diff options
| -rw-r--r-- | tools/cuda-ab/memset_probe.cu | 87 | ||||
| -rwxr-xr-x | tools/cuda-ab/run/inplace_scale | bin | 0 -> 820936 bytes | |||
| -rwxr-xr-x | tools/cuda-ab/run/memset_probe | bin | 0 -> 820840 bytes | |||
| -rwxr-xr-x | tools/cuda-ab/run/wedge_recover | bin | 0 -> 821024 bytes | |||
| -rw-r--r-- | tools/cuda-ab/wedge_recover.cu | 116 |
5 files changed, 203 insertions, 0 deletions
diff --git a/tools/cuda-ab/memset_probe.cu b/tools/cuda-ab/memset_probe.cu new file mode 100644 index 0000000..a6a12e0 --- /dev/null +++ b/tools/cuda-ab/memset_probe.cu @@ -0,0 +1,87 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * memset_probe.cu - isolate the memcpy_roundtrip memset-phase failure. + * + * cudaMemset(8 MB, 0x5A) -> D2H -> byte verify, with a run-length report: + * the first bad offset, and a coverage map at 4 KB granularity (correct / + * wrong / untouched-malloc-prefill) so "fill under-delivered" is + * distinguishable from "readback delivered wrong pages". + * + * Build: nvcc -O2 --cudart=static -o memset_probe memset_probe.cu + */ + +#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) + +#define PAGE 4096 +#define MemsetByte 0x5A + +int main(void) +{ + const size_t bytes = 8u << 20; + unsigned char *h; + unsigned char *d; + size_t i; + int bad = 0; + size_t first_bad = (size_t)-1; + size_t n_pages, p; + + setvbuf(stdout, NULL, _IONBF, 0); + h = (unsigned char *)malloc(bytes); + if (!h) + return 1; + memset(h, 0xEE, bytes); + + CHECK(cudaMalloc(&d, bytes)); + printf("DPTR d=%p bytes=%zu\n", (void *)d, bytes); + + CHECK(cudaMemset(d, MemsetByte, bytes)); + CHECK(cudaDeviceSynchronize()); + CHECK(cudaMemcpy(h, d, bytes, cudaMemcpyDeviceToHost)); + + for (i = 0; i < bytes; i++) { + if (h[i] != (unsigned char)MemsetByte) { + if (first_bad == (size_t)-1) + first_bad = i; + bad++; + } + } + printf("FIRST_BAD off=%zu (0x%zx) bad=%d total=%zu\n", + first_bad, first_bad, bad, bytes); + n_pages = bytes / PAGE; + for (p = 0; p < n_pages; p++) { + int good = 1, prefill = 1; + + for (i = 0; i < PAGE; i++) { + if (h[p * PAGE + i] != (unsigned char)MemsetByte) + good = 0; + if (h[p * PAGE + i] != 0xEE) + prefill = 0; + if (!good && !prefill) + break; + } + if (!good) { + printf("PAGE %zu %s b0=%02x b1=%02x b2=%02x b3=%02x b4095=%02x\n", + p, prefill ? "PREFILL" : "OTHER", + h[p * PAGE], h[p * PAGE + 1], + h[p * PAGE + 2], h[p * PAGE + 3], + h[p * PAGE + 4095]); + if (p > 24) + break; + } + } + printf("STATUS %s\n", bad == 0 ? "OK" : "FAIL"); + return 0; +} diff --git a/tools/cuda-ab/run/inplace_scale b/tools/cuda-ab/run/inplace_scale Binary files differnew file mode 100755 index 0000000..bb47f42 --- /dev/null +++ b/tools/cuda-ab/run/inplace_scale diff --git a/tools/cuda-ab/run/memset_probe b/tools/cuda-ab/run/memset_probe Binary files differnew file mode 100755 index 0000000..bd9df5e --- /dev/null +++ b/tools/cuda-ab/run/memset_probe diff --git a/tools/cuda-ab/run/wedge_recover b/tools/cuda-ab/run/wedge_recover Binary files differnew file mode 100755 index 0000000..2238fc1 --- /dev/null +++ b/tools/cuda-ab/run/wedge_recover diff --git a/tools/cuda-ab/wedge_recover.cu b/tools/cuda-ab/wedge_recover.cu new file mode 100644 index 0000000..deba670 --- /dev/null +++ b/tools/cuda-ab/wedge_recover.cu @@ -0,0 +1,116 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * wedge_recover.cu - reproduce the per-launch no-execute (wedge) and test + * whether an idle gap heals a wedged context. + * + * Each iteration launches a small vecadd, syncs, reads back with two separate + * D2H copies (two independent readback paths so a single copy flake is not + * counted), verifies, and prints one line. On a verification failure the + * launch is retried after a configurable sleep (default 20 s): if the wedge + * heals by idling, retried iterations pass; if it needs a fresh context they + * keep failing. Latency per iteration is printed so the fix knob is visible. + * + * Build: nvcc -O2 --cudart=static -o wedge_recover wedge_recover.cu + */ + +#include <cstdio> +#include <cstdlib> +#include <cuda_runtime.h> +#include <unistd.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(int argc, char **argv) +{ + const int n = 1 << 15; + size_t bytes = (size_t)n * sizeof(float); + float *ha, *hb, *hc, *da, *db, *dc; + int iters = argc > 1 ? atoi(argv[1]) : 40; + int sleep_s = argc > 2 ? atoi(argv[2]) : 20; + int fails = 0, heals = 0, i; + + setvbuf(stdout, NULL, _IONBF, 0); + ha = (float *)malloc(bytes); + hb = (float *)malloc(bytes); + hc = (float *)malloc(bytes); + if (!ha || !hb || !hc) + return 1; + for (i = 0; i < n; i++) { + ha[i] = (float)(i % 977) * 0.25f; + hb[i] = (float)(i % 251) * 0.5f; + } + CHECK(cudaMalloc(&da, bytes)); + CHECK(cudaMalloc(&db, bytes)); + CHECK(cudaMalloc(&dc, bytes)); + CHECK(cudaMemcpy(da, ha, bytes, cudaMemcpyHostToDevice)); + CHECK(cudaMemcpy(db, hb, bytes, cudaMemcpyHostToDevice)); + + for (i = 0; i < iters; i++) { + int attempt = 0; + long long ms[2] = { 0, 0 }; + + for (;;) { + int bad = 0, j; + cudaEvent_t e0, e1; + float ems = 0; + + cudaEventCreate(&e0); + cudaEventCreate(&e1); + cudaEventRecord(e0); + vadd<<<(n + 255) / 256, 256>>>(da, db, dc, n); + CHECK(cudaGetLastError()); + cudaEventRecord(e1); + CHECK(cudaDeviceSynchronize()); + cudaEventElapsedTime(&ems, e0, e1); + cudaEventDestroy(e0); + cudaEventDestroy(e1); + CHECK(cudaMemcpy(hc, dc, bytes, cudaMemcpyDeviceToHost)); + for (j = 0; j < n; j++) { + float expect = ha[j] + hb[j]; + + if (hc[j] != expect) { + bad++; + } + } + ms[attempt & 1] = (long long)(ems * 1000.0f); + if (bad == 0) { + if (attempt > 0) + heals++; + printf("ITER %d OK attempt=%d launch_us=%lld\n", + i, attempt, ms[attempt & 1]); + break; + } + fails++; + printf("ITER %d BAD bad=%d attempt=%d launch_us=%lld\n", + i, bad, attempt, ms[attempt & 1]); + if (attempt >= 3) { + printf("ITER %d GIVEUP\n", i); + break; + } + attempt++; + printf("ITER %d SLEEP %ds before retry\n", i, sleep_s); + sleep(sleep_s); + } + } + printf("SUMMARY iters=%d fails=%d heals=%d sleep_s=%d\n", + iters, fails, heals, sleep_s); + printf("STATUS %s\n", fails == 0 ? "OK" : "FAIL"); + return fails == 0 ? 0 : 1; +} |