Skip to content

Commit f1d7d6f

Browse files
SHANAjamSHANAjam
committed
xbar5090 0.2.2 update: src/xbar5090/data/xbar_l2_check.cu
1 parent 4b0280e commit f1d7d6f

1 file changed

Lines changed: 184 additions & 6 deletions

File tree

‎src/xbar5090/data/xbar_l2_check.cu‎

Lines changed: 184 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1,14 +1,27 @@
1-
// XBAR L2 data-integrity checker (CUDA source).
1+
// XBAR L2 data-integrity checker (CUDA source, full host + device).
22
//
3-
// This is the original kernel published by Loong0x00 in LACT issue #1147:
4-
// https://github.com/ilya-zlobintsev/LACT/issues/1147
3+
// This is the original L2 checker published by Loong0x00 in LACT issue #1147
4+
// (https://github.com/ilya-zlobintsev/LACT/issues/1147), with a small
5+
// maintainable host wrapper:
6+
// - --mb controls the checker buffer size (default 32 MiB).
7+
// - Error messages include the numeric CUDA error code so failures such as
8+
// "cudaMalloc data: (null)" on some systems are actionable instead of
9+
// opaque.
10+
// - load mode runs a concurrent stress kernel to add L2/global traffic.
511
//
6-
// The bundled xbar_l2_check.exe is built from this source. Keeping the source
7-
// in the repository makes the trust chain auditable.
12+
// The PTX-specific L2 load is unchanged:
13+
// ld.global.cg.u32 bypasses L1 allocation and exercises the L2/global path.
814

915
#include <cstdint>
16+
#include <cstdio>
17+
#include <cstdlib>
18+
#include <cstring>
19+
#include <chrono>
20+
#include <string>
1021

11-
__device__ __forceinline__ uint32_t pattern(uint32_t i) {
22+
#include <cuda_runtime.h>
23+
24+
__host__ __device__ __forceinline__ uint32_t pattern(uint32_t i) {
1225
return i * 747796405u + 2891336453u;
1326
}
1427

@@ -24,3 +37,168 @@ __global__ void xbar_l2_check(const uint32_t *data, uint32_t mask,
2437
atomicAdd(errors, 1ULL);
2538
}
2639
}
40+
41+
__global__ void stress_kernel(uint32_t *dummy, uint32_t mask,
42+
uint32_t stress_iters) {
43+
uint32_t s = blockIdx.x * blockDim.x + threadIdx.x;
44+
for (uint32_t n = 0; n < stress_iters; ++n) {
45+
s = s * 1664525u + 1013904223u;
46+
uint32_t i = s & mask;
47+
dummy[i] = dummy[i] * 1664525u + 1013904223u;
48+
}
49+
}
50+
51+
static int gpu_blocks = 1360;
52+
static int gpu_threads = 256;
53+
static int gpu_rounds = 1;
54+
static std::string gpu_mode = "idle";
55+
static int stress_iters = 100000;
56+
static int mb = 32;
57+
58+
static bool parse_args(int argc, char **argv) {
59+
for (int i = 1; i < argc; ++i) {
60+
std::string arg = argv[i];
61+
auto next = [&](const char *name) -> const char * {
62+
if (i + 1 >= argc) {
63+
fprintf(stderr, "missing value for %s\n", name);
64+
return nullptr;
65+
}
66+
return argv[++i];
67+
};
68+
if (arg == "--blocks") {
69+
const char *v = next("--blocks");
70+
if (!v) return false;
71+
gpu_blocks = atoi(v);
72+
if (gpu_blocks <= 0) { fprintf(stderr, "bad --blocks\n"); return false; }
73+
} else if (arg == "--threads") {
74+
const char *v = next("--threads");
75+
if (!v) return false;
76+
gpu_threads = atoi(v);
77+
if (gpu_threads <= 0) { fprintf(stderr, "bad --threads\n"); return false; }
78+
} else if (arg == "--rounds") {
79+
const char *v = next("--rounds");
80+
if (!v) return false;
81+
gpu_rounds = atoi(v);
82+
if (gpu_rounds <= 0) { fprintf(stderr, "bad --rounds\n"); return false; }
83+
} else if (arg == "--mode") {
84+
const char *v = next("--mode");
85+
if (!v) return false;
86+
gpu_mode = v;
87+
if (gpu_mode != "idle" && gpu_mode != "load") {
88+
fprintf(stderr, "bad --mode (idle|load)\n");
89+
return false;
90+
}
91+
} else if (arg == "--stress-iters") {
92+
const char *v = next("--stress-iters");
93+
if (!v) return false;
94+
stress_iters = atoi(v);
95+
if (stress_iters <= 0) { fprintf(stderr, "bad --stress-iters\n"); return false; }
96+
} else if (arg == "--mb") {
97+
const char *v = next("--mb");
98+
if (!v) return false;
99+
mb = atoi(v);
100+
if (mb <= 0 || mb > 4096) { fprintf(stderr, "bad --mb\n"); return false; }
101+
} else {
102+
fprintf(stderr, "unknown arg: %s\n", arg.c_str());
103+
return false;
104+
}
105+
}
106+
return true;
107+
}
108+
109+
static const char *err_string(cudaError_t err) {
110+
// cudaGetErrorString can return NULL for a not-yet-initialized runtime on
111+
// some systems; fall back to a numeric representation.
112+
const char *s = cudaGetErrorString(err);
113+
return s ? s : "(unknown error)";
114+
}
115+
116+
#define CUDA_CHECK(label, expr) \
117+
do { \
118+
cudaError_t err = (expr); \
119+
if (err != cudaSuccess) { \
120+
fprintf(stderr, "CUDA error: %s: %s [code=%d]\n", \
121+
label, err_string(err), (int)err); \
122+
return 1; \
123+
} \
124+
} while (0)
125+
126+
int main(int argc, char **argv) {
127+
if (!parse_args(argc, argv))
128+
return 2;
129+
130+
size_t bytes = (size_t)mb * 1024 * 1024;
131+
size_t words = bytes / sizeof(uint32_t);
132+
if ((words & (words - 1)) != 0) {
133+
fprintf(stderr, "--mb must make a power-of-two word count\n");
134+
return 2;
135+
}
136+
uint32_t mask = (uint32_t)(words - 1);
137+
138+
uint32_t *host_data = (uint32_t *)malloc(bytes);
139+
if (!host_data) {
140+
fprintf(stderr, "host malloc failed\n");
141+
return 1;
142+
}
143+
for (size_t i = 0; i < words; ++i)
144+
host_data[i] = pattern((uint32_t)i);
145+
146+
uint32_t *data = nullptr;
147+
CUDA_CHECK("cudaMalloc data", cudaMalloc(&data, bytes));
148+
unsigned long long *errors = nullptr;
149+
CUDA_CHECK("cudaMalloc errors", cudaMalloc(&errors, sizeof(*errors)));
150+
uint32_t *dummy = nullptr;
151+
if (gpu_mode == "load") {
152+
size_t dummy_words = (size_t)gpu_blocks * gpu_threads;
153+
if (dummy_words == 0)
154+
dummy_words = 1;
155+
size_t dummy_bytes = dummy_words * sizeof(uint32_t);
156+
CUDA_CHECK("cudaMalloc dummy", cudaMalloc(&dummy, dummy_bytes));
157+
CUDA_CHECK("cudaMemset dummy", cudaMemset(dummy, 0, dummy_bytes));
158+
}
159+
160+
CUDA_CHECK("cudaMemcpy data", cudaMemcpy(data, host_data, bytes, cudaMemcpyHostToDevice));
161+
CUDA_CHECK("cudaMemset errors", cudaMemset(errors, 0, sizeof(*errors)));
162+
163+
cudaStream_t stream_a = nullptr, stream_b = nullptr;
164+
if (gpu_mode == "load") {
165+
CUDA_CHECK("cudaStreamCreate", cudaStreamCreateWithFlags(&stream_a, cudaStreamNonBlocking));
166+
CUDA_CHECK("cudaStreamCreate", cudaStreamCreateWithFlags(&stream_b, cudaStreamNonBlocking));
167+
}
168+
169+
auto start = std::chrono::high_resolution_clock::now();
170+
for (int round = 0; round < gpu_rounds; ++round) {
171+
CUDA_CHECK("cudaMemset errors", cudaMemset(errors, 0, sizeof(*errors)));
172+
if (gpu_mode == "load") {
173+
size_t dummy_words = (size_t)gpu_blocks * gpu_threads;
174+
uint32_t dummy_mask = (uint32_t)(dummy_words - 1);
175+
stress_kernel<<<gpu_blocks, gpu_threads, 0, stream_a>>>(
176+
dummy, dummy_mask, (uint32_t)stress_iters);
177+
CUDA_CHECK("stress launch", cudaGetLastError());
178+
xbar_l2_check<<<gpu_blocks, gpu_threads, 0, stream_b>>>(
179+
data, mask, errors);
180+
CUDA_CHECK("checker launch", cudaGetLastError());
181+
CUDA_CHECK("stress sync", cudaStreamSynchronize(stream_a));
182+
CUDA_CHECK("device sync", cudaStreamSynchronize(stream_b));
183+
} else {
184+
xbar_l2_check<<<gpu_blocks, gpu_threads>>>(data, mask, errors);
185+
CUDA_CHECK("checker launch", cudaGetLastError());
186+
CUDA_CHECK("device sync", cudaDeviceSynchronize());
187+
}
188+
unsigned long long host_errors = 0;
189+
CUDA_CHECK("cudaMemcpy errors",
190+
cudaMemcpy(&host_errors, errors, sizeof(host_errors), cudaMemcpyDeviceToHost));
191+
auto end = std::chrono::high_resolution_clock::now();
192+
double elapsed_ms = std::chrono::duration<double, std::milli>(end - start).count();
193+
printf("round=%u mode=%s errors=%llu elapsed_ms=%.3f\n",
194+
round, gpu_mode.c_str(), host_errors, elapsed_ms);
195+
}
196+
197+
if (stream_a) cudaStreamDestroy(stream_a);
198+
if (stream_b) cudaStreamDestroy(stream_b);
199+
if (dummy) cudaFree(dummy);
200+
cudaFree(errors);
201+
cudaFree(data);
202+
free(host_data);
203+
return 0;
204+
}

0 commit comments

Comments
 (0)