diff options
Diffstat (limited to 'tools/cuda-ab/memset2d_probe.cu')
| -rw-r--r-- | tools/cuda-ab/memset2d_probe.cu | 157 |
1 files changed, 157 insertions, 0 deletions
diff --git a/tools/cuda-ab/memset2d_probe.cu b/tools/cuda-ab/memset2d_probe.cu new file mode 100644 index 0000000..acf43d4 --- /dev/null +++ b/tools/cuda-ab/memset2d_probe.cu @@ -0,0 +1,157 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * memset2d_probe.cu - per-phase diagnostics for victory_memset2d. + * + * victory_memset2d reports one aggregate bad count (~12.5M) that mixes all + * three phases. This probe runs the SAME three phases and reports each + * phase's bad count, the first mismatching byte (offset, got, want), and a + * per-phase 1MB-chunk histogram so the failing phase and the failing region + * are attributable in one run. + * + * Phases: + * P1 memset : cudaMemset(d_img, 0xA5, pitch*h); D2H 2D readback + * P2 roundtrip: H2D 2D pattern; D2H 2D readback + * P3 memset2d : cudaMemset2D quadrant 0x3C; D2H 2D readback + * + * Output lines: + * DEVICE ... + * PHASE p=1 bad=<n> first_off=<i> got=<g> want=<w> + * PHIST p=1 mb=<m> ok=<n> bad=<n> + * MEMSET2D2 bad_total=<n> + * STATUS OK|FAIL + */ + +#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 IMG_W 4096 +#define IMG_H 1024 +#define FILL_BYTE 0xA5 +#define QUAD_BYTE 0x3C + +static int phase_check(const char *tag, const unsigned char *back, + const unsigned char *want, size_t n, int p) +{ + size_t i, first = (size_t)-1; + int bad = 0, got = 0, wantv = 0; + size_t mb; + + for (i = 0; i < n; i++) { + if (back[i] != want[i]) { + if (first == (size_t)-1) { + first = i; + got = back[i]; + wantv = want[i]; + } + bad++; + } + } + printf("PHASE p=%d bad=%d first_off=%zu got=0x%02x want=0x%02x\n", + p, bad, first, got, wantv); + if (bad) { + size_t nb = (n + (1u << 20) - 1) >> 20; + + for (mb = 0; mb < nb && mb < 8; mb++) { + size_t lo = mb << 20, hi = lo + (1u << 20); + size_t j; + int bok = 0, bbad = 0; + + if (hi > n) + hi = n; + for (j = lo; j < hi; j++) { + if (back[j] == want[j]) + bok++; + else + bbad++; + } + printf("PHIST p=%d mb=%zu ok=%d bad=%d\n", + p, mb, bok, bbad); + } + } + (void)tag; + return bad; +} + +int main(void) +{ + cudaDeviceProp prop; + const size_t w = IMG_W, h = IMG_H; + unsigned char *h_img, *h_back, *h_want; + unsigned char *d_img; + size_t pitch = 0; + int bad = 0, row, col; + + 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_img = (unsigned char *)malloc(w * h); + h_back = (unsigned char *)malloc(w * h); + h_want = (unsigned char *)malloc(w * h); + if (!h_img || !h_back || !h_want) + return 1; + + CHECK(cudaMallocPitch(&d_img, &pitch, w, h)); + printf("PITCH pitch=%zu\n", pitch); + printf("DPTRS h_img=%p h_back=%p h_want=%p d_img=%p\n", (void *)h_img, + (void *)h_back, (void *)h_want, (void *)d_img); + + /* P1: full cudaMemset over the pitched allocation, 2D readback. */ + CHECK(cudaMemset(d_img, FILL_BYTE, pitch * h)); + memset(h_back, 0, w * h); + CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h, + cudaMemcpyDeviceToHost)); + memset(h_want, FILL_BYTE, w * h); + bad += phase_check("P1", h_back, h_want, w * h, 1); + + /* P2: pitched H2D pattern -> D2H round trip. */ + for (row = 0; row < (int)h; row++) + for (col = 0; col < (int)w; col++) + h_img[row * w + col] = (unsigned char)((row + col) & 0xFF); + CHECK(cudaMemcpy2D(d_img, pitch, h_img, w, w, h, + cudaMemcpyHostToDevice)); + memset(h_back, 0, w * h); + CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h, + cudaMemcpyDeviceToHost)); + bad += phase_check("P2", h_back, h_img, w * h, 2); + + /* P3: cudaMemset2D of the upper-left quadrant only. */ + CHECK(cudaMemset2D(d_img, pitch, QUAD_BYTE, w / 2, h / 2)); + for (row = 0; row < (int)h; row++) + for (col = 0; col < (int)w; col++) { + unsigned char expect = h_img[row * w + col]; + + if (row < (int)(h / 2) && col < (int)(w / 2)) + expect = QUAD_BYTE; + h_want[row * w + col] = expect; + } + memset(h_back, 0, w * h); + CHECK(cudaMemcpy2D(h_back, w, d_img, pitch, w, h, + cudaMemcpyDeviceToHost)); + bad += phase_check("P3", h_back, h_want, w * h, 3); + + printf("MEMSET2D2 bad_total=%d\n", bad); + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + printf("STATUS OK\n"); + return 0; +} |