Publish matrix_mult kernel and performance card
Browse filesCUDA source, short description, performance table, and plot.
- README.md +40 -0
- kernel.cu +30 -0
- performance.svg +8 -0
README.md
ADDED
|
@@ -0,0 +1,40 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
---
|
| 2 |
+
license: apache-2.0
|
| 3 |
+
tags:
|
| 4 |
+
- cuda
|
| 5 |
+
- kernel
|
| 6 |
+
- gpu-optimization
|
| 7 |
+
- hpc
|
| 8 |
+
---
|
| 9 |
+
|
| 10 |
+
# matrix_mult
|
| 11 |
+
|
| 12 |
+
CUDA matrix multiplication kernel (compile-only baseline).
|
| 13 |
+
|
| 14 |
+
This repository contains the standalone CUDA source for the `matrix_mult` lane from
|
| 15 |
+
the PyC kernel lab. It is a source artifact for inspection and benchmarking;
|
| 16 |
+
it is not a precompiled binary and the result below is not a universal ranking.
|
| 17 |
+
|
| 18 |
+
## Performance
|
| 19 |
+
|
| 20 |
+
| Kernel | GPU / architecture | Shape | Best recorded result | Evidence |
|
| 21 |
+
|---|---|---|---|---|
|
| 22 |
+
| `matrix_mult` | not recorded | not recorded | Not measured in the published campaign | No published performance receipt was found for this lane. |
|
| 23 |
+
|
| 24 |
+
|
| 25 |
+
|
| 26 |
+
The result is reported with the original campaign's timing and correctness
|
| 27 |
+
context. Compare kernels only when GPU, CUDA version, matrix shape, warmup,
|
| 28 |
+
repeats, and reference/correctness mode match.
|
| 29 |
+
|
| 30 |
+
## Source
|
| 31 |
+
|
| 32 |
+
- `kernel.cu` — copied from `kernels/prototypes/baseline/matmul/kernel.cu`.
|
| 33 |
+
- Original lane tags: `cuda, matmul`.
|
| 34 |
+
|
| 35 |
+
## Build/run contract
|
| 36 |
+
|
| 37 |
+
```text
|
| 38 |
+
{nvcc} -O3 -c {source} -o {build_dir}/{name}.o
|
| 39 |
+
(compile-only)
|
| 40 |
+
```
|
kernel.cu
ADDED
|
@@ -0,0 +1,30 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
// /kernel/matrix_mult.cu
|
| 2 |
+
#include <cuda_runtime.h>
|
| 3 |
+
|
| 4 |
+
__global__ void matrix_mult_kernel(float* a, float* b, float* c, int m, int n, int k) {
|
| 5 |
+
int row = blockIdx.y * blockDim.y + threadIdx.y;
|
| 6 |
+
int col = blockIdx.x * blockDim.x + threadIdx.x;
|
| 7 |
+
if (row < m && col < n) {
|
| 8 |
+
float sum = 0.0f;
|
| 9 |
+
for (int i = 0; i < k; i++) {
|
| 10 |
+
sum += a[row * k + i] * b[i * n + col];
|
| 11 |
+
}
|
| 12 |
+
c[row * n + col] = sum;
|
| 13 |
+
}
|
| 14 |
+
}
|
| 15 |
+
|
| 16 |
+
void matrix_multiply(float* a, float* b, float* c, int m, int n, int k) {
|
| 17 |
+
float *d_a, *d_b, *d_c;
|
| 18 |
+
cudaMalloc(&d_a, m * k * sizeof(float));
|
| 19 |
+
cudaMalloc(&d_b, k * n * sizeof(float));
|
| 20 |
+
cudaMalloc(&d_c, m * n * sizeof(float));
|
| 21 |
+
cudaMemcpy(d_a, a, m * k * sizeof(float), cudaMemcpyHostToDevice);
|
| 22 |
+
cudaMemcpy(d_b, b, k * n * sizeof(float), cudaMemcpyHostToDevice);
|
| 23 |
+
|
| 24 |
+
dim3 threads(16, 16);
|
| 25 |
+
dim3 blocks((n + threads.x - 1) / threads.x, (m + threads.y - 1) / threads.y);
|
| 26 |
+
matrix_mult_kernel<<<blocks, threads>>>(d_a, d_b, d_c, m, n, k);
|
| 27 |
+
|
| 28 |
+
cudaMemcpy(c, d_c, m * n * sizeof(float), cudaMemcpyDeviceToHost);
|
| 29 |
+
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
|
| 30 |
+
}
|
performance.svg
ADDED
|