1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* victory_matmul.cu - 256x256 tiled matrix multiply with shared memory.
*
* C = A * B for three 256x256 float matrices of small integer values. The
* inputs are integers in [0, 10), so every product and every partial sum is
* exactly representable in float (< 2^24); whatever order the GPU accumulates
* in, the result is exact. The host replays the multiply with 64-bit integers
* and verifies every element exactly, plus a checksum over C.
*
* Output lines are stable and greppable:
*
* DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
* MATMUL n=256 bad=<mismatches> csum=<sum>
* STATUS OK|FAIL
*
* Build with a static CUDA runtime (--cudart=static) so the guest only needs
* the driver's libcuda.so, not a CUDA toolkit installation:
*
* nvcc -O2 --cudart=static -o victory_matmul victory_matmul.cu
*
* Then, from the guest: /tmp/victory_matmul
*/
#include <cstdio>
#include <cstdlib>
#include <cstdint>
#include <cuda_runtime.h>
#define CHECK(x) \
do { \
cudaError_t e_ = (x); \
if (e_ != cudaSuccess) { \
printf("CUDA_ERROR %s at line %d\n", \
cudaGetErrorString(e_), __LINE__); \
printf("STATUS FAIL\n"); \
return 1; \
} \
} while (0)
#define TILE 16
#define N 256
__global__ void matmul_tiled(const float *A, const float *B, float *C, int n)
{
__shared__ float As[TILE][TILE];
__shared__ float Bs[TILE][TILE];
int row = blockIdx.y * TILE + threadIdx.y;
int col = blockIdx.x * TILE + threadIdx.x;
float acc = 0.0f;
for (int t = 0; t < n / TILE; t++) {
As[threadIdx.y][threadIdx.x] = A[row * n + t * TILE + threadIdx.x];
Bs[threadIdx.y][threadIdx.x] =
B[(t * TILE + threadIdx.y) * n + col];
__syncthreads();
for (int k = 0; k < TILE; k++)
acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];
__syncthreads();
}
C[row * n + col] = acc;
}
int main(void)
{
cudaDeviceProp prop;
const int n = N;
const size_t bytes = (size_t)n * n * sizeof(float);
float *h_a, *h_b, *h_c;
float *d_a, *d_b, *d_c;
double csum = 0.0;
int bad = 0, i, j, k;
dim3 grid(n / TILE, n / TILE);
dim3 block(TILE, TILE);
setvbuf(stdout, NULL, _IONBF, 0);
CHECK(cudaGetDeviceProperties(&prop, 0));
printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name,
prop.major, prop.minor, prop.multiProcessorCount,
(unsigned long long)prop.totalGlobalMem);
h_a = (float *)malloc(bytes);
h_b = (float *)malloc(bytes);
h_c = (float *)malloc(bytes);
if (!h_a || !h_b || !h_c)
return 1;
for (i = 0; i < n * n; i++) {
h_a[i] = (float)(i % 10);
h_b[i] = (float)((i * 3 + 1) % 10);
}
CHECK(cudaMalloc(&d_a, bytes));
CHECK(cudaMalloc(&d_b, bytes));
CHECK(cudaMalloc(&d_c, bytes));
CHECK(cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice));
matmul_tiled<<<grid, block>>>(d_a, d_b, d_c, n);
CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());
CHECK(cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost));
for (i = 0; i < n; i++) {
for (j = 0; j < n; j++) {
long long ref = 0;
for (k = 0; k < n; k++)
ref += (long long)h_a[i * n + k] *
(long long)h_b[k * n + j];
if (h_c[i * n + j] != (float)ref)
bad++;
csum += (double)h_c[i * n + j];
}
}
printf("MATMUL n=%d bad=%d csum=%.6f\n", n, bad, csum);
if (bad != 0) {
printf("STATUS FAIL\n");
return 1;
}
printf("STATUS OK\n");
CHECK(cudaFree(d_a));
CHECK(cudaFree(d_b));
CHECK(cudaFree(d_c));
free(h_a);
free(h_b);
free(h_c);
return 0;
}
|
