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
136
137
138
|
// SPDX-License-Identifier: BSD-3-Clause
/*
* victory_strided.cu - strided/offset memory access over 4 MB of floats.
*
* Two kernels permute an n-element array (n = 1M floats, 4 MB) with a
* stride co-prime to n (idx = i*7 mod n), so every element is written exactly
* once and the host can verify every single position:
*
* out1[(i*7) % n] = in[i] + 1.0f (plain permutation)
* out2[(i*7) % n] = in[off + i] * 2.0f (offset base pointer)
*
* All values are multiples of 0.25, so both operations are exact in float
* and the host comparison is exact.
*
* Output lines are stable and greppable:
*
* DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
* STRIDED n=<count> stride=7 off=<offset> 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_strided victory_strided.cu
*
* Then, from the guest: /tmp/victory_strided
*/
#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 STRIDE 7
#define OFFSET 4097
__global__ void permute_add(const float *in, float *out, int n, int stride)
{
long long i = (long long)blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
long long idx = (i * (long long)stride) % (long long)n;
out[idx] = in[i] + 1.0f;
}
}
__global__ void permute_offset(const float *in, float *out, int n, int off,
int stride)
{
long long i = (long long)blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
long long idx = (i * (long long)stride) % (long long)n;
out[idx] = in[(long long)off + i] * 2.0f;
}
}
int main(void)
{
cudaDeviceProp prop;
const int n = 1 << 20; /* 1M floats = 4 MB */
const long long total = (long long)n + OFFSET;
size_t bytes = (size_t)n * sizeof(float);
size_t total_bytes = (size_t)total * sizeof(float);
float *h_in, *h_out1, *h_out2;
float *d_in, *d_out1, *d_out2;
double csum = 0.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_in = (float *)malloc(total_bytes);
h_out1 = (float *)malloc(bytes);
h_out2 = (float *)malloc(bytes);
if (!h_in || !h_out1 || !h_out2)
return 1;
for (i = 0; i < (int)total; i++)
h_in[i] = (float)(i % 1000) * 0.25f;
CHECK(cudaMalloc(&d_in, total_bytes));
CHECK(cudaMalloc(&d_out1, bytes));
CHECK(cudaMalloc(&d_out2, bytes));
CHECK(cudaMemcpy(d_in, h_in, total_bytes, cudaMemcpyHostToDevice));
permute_add<<<(n + 255) / 256, 256>>>(d_in, d_out1, n, STRIDE);
permute_offset<<<(n + 255) / 256, 256>>>(d_in, d_out2, n, OFFSET,
STRIDE);
CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());
CHECK(cudaMemcpy(h_out1, d_out1, bytes, cudaMemcpyDeviceToHost));
CHECK(cudaMemcpy(h_out2, d_out2, bytes, cudaMemcpyDeviceToHost));
for (i = 0; i < n; i++) {
long long idx = ((long long)i * STRIDE) % (long long)n;
if (h_out1[idx] != h_in[i] + 1.0f)
bad++;
if (h_out2[idx] != h_in[OFFSET + i] * 2.0f)
bad++;
csum += (double)h_out1[idx];
}
printf("STRIDED n=%d stride=%d off=%d bad=%d csum=%.6f\n", n, STRIDE,
OFFSET, bad, csum);
if (bad != 0) {
printf("STATUS FAIL\n");
return 1;
}
printf("STATUS OK\n");
CHECK(cudaFree(d_in));
CHECK(cudaFree(d_out1));
CHECK(cudaFree(d_out2));
free(h_in);
free(h_out1);
free(h_out2);
return 0;
}
|
