|
| 1 | +--- |
| 2 | +layout: default |
| 3 | +title: Architecture Guide |
| 4 | +nav_order: 2 |
| 5 | +--- |
| 6 | + |
| 7 | +# Architecture Guide |
| 8 | + |
| 9 | +This document describes the system architecture and design decisions for the SGEMM Optimization project. |
| 10 | + |
| 11 | +## System Overview |
| 12 | + |
| 13 | +The project implements five progressively optimized CUDA SGEMM (Single-precision General Matrix Multiply) kernels, demonstrating core GPU optimization techniques from a naive triple-loop to Tensor Core WMMA API usage. |
| 14 | + |
| 15 | +### Design Principles |
| 16 | + |
| 17 | +1. **Single Source of Truth**: All implementations follow specifications in `/specs/` |
| 18 | +2. **Spec-Driven Development (SDD)**: Specs before code, documentation synchronized with implementation |
| 19 | +3. **RAII Resource Management**: No raw `cudaFree()` calls, all resources use wrapper classes |
| 20 | +4. **Exception-Based Error Handling**: No `exit()` in library code |
| 21 | +5. **Header-Only Kernels**: All kernels are `.cuh` files for easy integration |
| 22 | + |
| 23 | +## Architecture Layers |
| 24 | + |
| 25 | +``` |
| 26 | +┌─────────────────────────────────────────────┐ |
| 27 | +│ Application Layer │ |
| 28 | +│ main.cu - Entry point & benchmark runner │ |
| 29 | +├─────────────────────────────────────────────┤ |
| 30 | +│ Framework Layer │ |
| 31 | +│ utils/benchmark.cuh - Benchmark framework │ |
| 32 | +│ utils/verify.cuh - Correctness checks │ |
| 33 | +│ utils/cuda_utils.cuh - RAII wrappers │ |
| 34 | +├─────────────────────────────────────────────┤ |
| 35 | +│ Kernel Layer │ |
| 36 | +│ kernels/naive_sgemm.cuh │ |
| 37 | +│ kernels/tiled_sgemm.cuh │ |
| 38 | +│ kernels/bank_conflict_free_sgemm.cuh │ |
| 39 | +│ kernels/double_buffer_sgemm.cuh │ |
| 40 | +│ kernels/tensor_core_sgemm.cuh │ |
| 41 | +├─────────────────────────────────────────────┤ |
| 42 | +│ CUDA Runtime Layer │ |
| 43 | +│ cuBLAS, CUDA Runtime API, WMMA API │ |
| 44 | +└─────────────────────────────────────────────┘ |
| 45 | +``` |
| 46 | + |
| 47 | +## Unified Kernel Interface |
| 48 | + |
| 49 | +All kernels implement the same interface template: |
| 50 | + |
| 51 | +```cpp |
| 52 | +template<int TILE_SIZE = 32> |
| 53 | +void sgemm_naive(const float* A, const float* B, float* C, |
| 54 | + int M, int N, int K, cudaStream_t stream = 0); |
| 55 | + |
| 56 | +template<int TILE_SIZE = 32> |
| 57 | +void sgemm_tiled(const float* A, const float* B, float* C, |
| 58 | + int M, int N, int K, cudaStream_t stream = 0); |
| 59 | + |
| 60 | +// ... and so on for each variant |
| 61 | +``` |
| 62 | +
|
| 63 | +### Interface Contract |
| 64 | +
|
| 65 | +| Parameter | Description | Constraints | |
| 66 | +|-----------|-------------|-------------| |
| 67 | +| `A` | Input matrix A (M×K), row-major | Device pointer, valid | |
| 68 | +| `B` | Input matrix B (K×N), row-major | Device pointer, valid | |
| 69 | +| `C` | Output matrix C (M×N), row-major | Device pointer, allocated | |
| 70 | +| `M` | Rows of A and C | M > 0 | |
| 71 | +| `N` | Columns of B and C | N > 0 | |
| 72 | +| `K` | Columns of A / Rows of B | K > 0 | |
| 73 | +| `stream` | CUDA stream (optional) | Default: 0 | |
| 74 | +
|
| 75 | +### Grid-Block Configuration |
| 76 | +
|
| 77 | +```cpp |
| 78 | +dim3 block(TILE_SIZE, TILE_SIZE); |
| 79 | +dim3 grid((N + TILE_SIZE - 1) / TILE_SIZE, |
| 80 | + (M + TILE_SIZE - 1) / TILE_SIZE); |
| 81 | +``` |
| 82 | + |
| 83 | +## Tensor Core Architecture |
| 84 | + |
| 85 | +The Tensor Core kernel uses NVIDIA's WMMA (Warp Matrix Multiply Accumulate) API: |
| 86 | + |
| 87 | +### Architecture-Specific Guards |
| 88 | + |
| 89 | +```cpp |
| 90 | +#if __CUDA_ARCH__ >= 700 |
| 91 | + // WMMA code for Volta and newer |
| 92 | + #include <mma.h> |
| 93 | +#else |
| 94 | + // Fallback to FP32 tiled kernel |
| 95 | +#endif |
| 96 | +``` |
| 97 | + |
| 98 | +### Alignment Requirements |
| 99 | + |
| 100 | +| Requirement | Detail | |
| 101 | +|-------------|--------| |
| 102 | +| M, K, N | Must be multiples of 16 for WMMA fast path | |
| 103 | +| Fallback | Automatic FP32 kernel for non-aligned sizes | |
| 104 | +| Precision | FP16 input, FP32 accumulation | |
| 105 | +| Verification | Relaxed tolerances: `rtol=5e-2, atol=1e-2` | |
| 106 | + |
| 107 | +## Error Handling Strategy |
| 108 | + |
| 109 | +All errors are handled via exceptions, not `exit()`: |
| 110 | + |
| 111 | +```cpp |
| 112 | +#define CUDA_CHECK(call) \ |
| 113 | + do { \ |
| 114 | + cudaError_t err = call; \ |
| 115 | + if (err != cudaSuccess) { \ |
| 116 | + throw std::runtime_error( \ |
| 117 | + std::string("CUDA error: ") + \ |
| 118 | + cudaGetErrorString(err)); \ |
| 119 | + } \ |
| 120 | + } while(0) |
| 121 | +``` |
| 122 | +
|
| 123 | +## Memory Management |
| 124 | +
|
| 125 | +RAII wrappers ensure no memory leaks: |
| 126 | +
|
| 127 | +```cpp |
| 128 | +struct DeviceBuffer { |
| 129 | + float* ptr; |
| 130 | + DeviceBuffer(size_t bytes) { |
| 131 | + CUDA_CHECK(cudaMalloc(&ptr, bytes)); |
| 132 | + } |
| 133 | + ~DeviceBuffer() { |
| 134 | + CUDA_CHECK(cudaFree(ptr)); |
| 135 | + } |
| 136 | + // Deleted copy/move constructors prevent misuse |
| 137 | +}; |
| 138 | +``` |
| 139 | + |
| 140 | +## Build System |
| 141 | + |
| 142 | +### CMake (Recommended) |
| 143 | + |
| 144 | +```cmake |
| 145 | +add_executable(sgemm_benchmark src/main.cu) |
| 146 | +target_link_libraries(sgemm_benchmark PRIVATE cublas) |
| 147 | +set_target_properties(sgemm_benchmark PROPERTIES |
| 148 | + CUDA_ARCHITECTURES ${GPU_ARCH}) |
| 149 | +``` |
| 150 | +
|
| 151 | +### GPU Architecture Support |
| 152 | +
|
| 153 | +| GPU | Architecture | Flag | |
| 154 | +|-----|-------------|------| |
| 155 | +| V100 | Volta | `sm_70` | |
| 156 | +| RTX 3090/A100 | Ampere | `sm_86` | |
| 157 | +| RTX 4090 | Ada | `sm_89` | |
| 158 | +| H100 | Hopper | `sm_90` | |
| 159 | +
|
| 160 | +## References |
| 161 | +
|
| 162 | +- [RFC 0001: Core Architecture](../specs/rfc/0001-core-architecture.md) |
| 163 | +- [RFC 0002: Implementation Roadmap](../specs/rfc/0002-implementation-roadmap.md) |
| 164 | +- [CUDA Programming Guide](https://docs.nvidia.com/cuda/) |
0 commit comments