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
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* vector_add.cu - CUDA A/B workload for the Open-IOV harness.
*
* One deterministic, machine-readable workload that can run on the host (bare
* metal) and in a guest (emulated vGPU). The output lines are stable so the
* two runs can be diffed:
*
* DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
* RESULT n=<count> bad=<mismatches> checksum=<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.
*/
#include <cstdio>
#include <cstdlib>
#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)
__global__ void vadd(const float *a, const float *b, float *c, int n)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
c[i] = a[i] + b[i];
}
int main(void)
{
cudaDeviceProp prop;
const int n = 1 << 20;
size_t bytes = (size_t)n * sizeof(float);
float *h_a, *h_b, *h_c;
float *d_a, *d_b, *d_c;
double sum = 0.0;
int bad = 0, i;
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; i++) {
h_a[i] = (float)(i % 1000) * 0.5f;
h_b[i] = (float)(i % 997) * 0.25f;
}
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));
vadd<<<(n + 255) / 256, 256>>>(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++) {
if (h_c[i] != h_a[i] + h_b[i])
bad++;
sum += h_c[i];
}
printf("RESULT n=%d bad=%d checksum=%.6f\n", n, bad, sum);
printf("STATUS %s\n", bad == 0 ? "OK" : "FAIL");
CHECK(cudaFree(d_a));
CHECK(cudaFree(d_b));
CHECK(cudaFree(d_c));
free(h_a);
free(h_b);
free(h_c);
return bad == 0 ? 0 : 1;
}
|
