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
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* victory_atomics.cu - global atomicAdd/atomicMax histogram over 4 MB.
*
* A deterministic 4M-byte buffer (every value 0..255) is histogrammed on the
* device into a 256-bin counter array with atomicAdd, while a global maximum
* is tracked with atomicMax. The host recomputes the exact histogram and
* maximum and verifies every bin, the max, and the total count.
*
* Output lines are stable and greppable:
*
* DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
* ATOMICS n=<count> bins=256 max=<value> bad=<mismatches> total=<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_atomics victory_atomics.cu
*
* Then, from the guest: /tmp/victory_atomics
*/
#include <cstdio>
#include <cstdlib>
#include <cstdint>
#include <cstring>
#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 BINS 256
__global__ void histogram(const unsigned char *data, unsigned int *hist,
unsigned int *gmax, int n)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
unsigned int v = data[i];
atomicAdd(&hist[v], 1u);
atomicMax(gmax, v);
}
}
int main(void)
{
cudaDeviceProp prop;
const int n = 1 << 22; /* 4M bytes */
unsigned char *h_data;
unsigned int *h_hist;
unsigned int h_max = 0, h_total = 0;
unsigned char *d_data;
unsigned int *d_hist, *d_max;
unsigned int ref[BINS];
unsigned int ref_max = 0;
int bad = 0, i;
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_data = (unsigned char *)malloc(n);
h_hist = (unsigned int *)malloc(BINS * sizeof(unsigned int));
if (!h_data || !h_hist)
return 1;
for (i = 0; i < n; i++)
h_data[i] = (unsigned char)(((long long)i * 37 + i / 256) & 0xFF);
memset(ref, 0, sizeof(ref));
for (i = 0; i < n; i++) {
ref[h_data[i]]++;
if (h_data[i] > ref_max)
ref_max = h_data[i];
}
CHECK(cudaMalloc(&d_data, n));
CHECK(cudaMalloc(&d_hist, BINS * sizeof(unsigned int)));
CHECK(cudaMalloc(&d_max, sizeof(unsigned int)));
CHECK(cudaMemcpy(d_data, h_data, n, cudaMemcpyHostToDevice));
CHECK(cudaMemset(d_hist, 0, BINS * sizeof(unsigned int)));
CHECK(cudaMemset(d_max, 0, sizeof(unsigned int)));
histogram<<<(n + 255) / 256, 256>>>(d_data, d_hist, d_max, n);
CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());
CHECK(cudaMemcpy(h_hist, d_hist, BINS * sizeof(unsigned int),
cudaMemcpyDeviceToHost));
CHECK(cudaMemcpy(&h_max, d_max, sizeof(unsigned int),
cudaMemcpyDeviceToHost));
for (i = 0; i < BINS; i++) {
if (h_hist[i] != ref[i])
bad++;
h_total += h_hist[i];
}
if (h_max != ref_max)
bad++;
printf("ATOMICS n=%d bins=256 max=%u bad=%d total=%u\n", n, h_max, bad,
h_total);
if (bad != 0 || h_total != (unsigned int)n) {
printf("STATUS FAIL\n");
return 1;
}
printf("STATUS OK\n");
CHECK(cudaFree(d_data));
CHECK(cudaFree(d_hist));
CHECK(cudaFree(d_max));
free(h_data);
free(h_hist);
return 0;
}
|
