diff options
Diffstat (limited to 'tools/cuda-ab/streams_events_probe.cu')
| -rw-r--r-- | tools/cuda-ab/streams_events_probe.cu | 170 |
1 files changed, 170 insertions, 0 deletions
diff --git a/tools/cuda-ab/streams_events_probe.cu b/tools/cuda-ab/streams_events_probe.cu new file mode 100644 index 0000000..94aef69 --- /dev/null +++ b/tools/cuda-ab/streams_events_probe.cu @@ -0,0 +1,170 @@ +// SPDX-License-Identifier: BSD-3-Clause +/* + * streams_events_probe.cu - instrumented victory_streams_events. + * + * Same 3-stream event chain (c0 = a0+b0; c1 = (a1+b1)+c0; c2 = (a2+b2)+c1) + * but with per-buffer diagnostics: device pointers, per-buffer bad counts, + * and the first 8 expected/actual values of every output so a stale-input + * vs stale-readback failure is distinguishable in one run. + * + * Output lines: + * DEVICE ... + * DPTRS da0..dc2 (device), ha0..hc2 (host) + * BUF s=<n> bad=<n> + * FIRST s=<n> exp=<v0..v7> got=<v0..v7> + * STREAMS3 n=<count> bad=<mismatches> + * STATUS OK|FAIL + */ + +#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 NSTREAMS 3 + +__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]; +} + +__global__ void vadd_into(float *c, const float *s, int n) +{ + int i = blockIdx.x * blockDim.x + threadIdx.x; + + if (i < n) + c[i] = c[i] + s[i]; +} + +int main(void) +{ + cudaDeviceProp prop; + const int n = 1 << 16; + size_t bytes = (size_t)n * sizeof(float); + float *h_a[NSTREAMS], *h_b[NSTREAMS], *h_c[NSTREAMS]; + float *d_a[NSTREAMS], *d_b[NSTREAMS], *d_c[NSTREAMS]; + cudaStream_t st[NSTREAMS]; + cudaEvent_t ev[NSTREAMS], t0, t1; + float ms = 0.0f; + int bad = 0, i, s; + const int grid = (n + 255) / 256; + + 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); + + for (s = 0; s < NSTREAMS; s++) { + h_a[s] = (float *)malloc(bytes); + h_b[s] = (float *)malloc(bytes); + h_c[s] = (float *)malloc(bytes); + if (!h_a[s] || !h_b[s] || !h_c[s]) + return 1; + for (i = 0; i < n; i++) { + h_a[s][i] = (float)((i + s) % 1000) * 0.5f; + h_b[s][i] = (float)((i + 2 * s) % 997) * 0.25f; + } + CHECK(cudaMalloc(&d_a[s], bytes)); + CHECK(cudaMalloc(&d_b[s], bytes)); + CHECK(cudaMalloc(&d_c[s], bytes)); + CHECK(cudaStreamCreate(&st[s])); + CHECK(cudaEventCreate(&ev[s])); + } + CHECK(cudaEventCreate(&t0)); + CHECK(cudaEventCreate(&t1)); + + printf("DPTRS"); + for (s = 0; s < NSTREAMS; s++) + printf(" da%d=%p db%d=%p dc%d=%p", s, (void *)d_a[s], s, + (void *)d_b[s], s, (void *)d_c[s]); + printf("\n"); + printf("HPTRS"); + for (s = 0; s < NSTREAMS; s++) + printf(" ha%d=%p hb%d=%p hc%d=%p", s, (void *)h_a[s], s, + (void *)h_b[s], s, (void *)h_c[s]); + printf("\n"); + + CHECK(cudaEventRecord(t0)); + + for (s = 0; s < NSTREAMS; s++) { + CHECK(cudaMemcpyAsync(d_a[s], h_a[s], bytes, + cudaMemcpyHostToDevice, st[s])); + CHECK(cudaMemcpyAsync(d_b[s], h_b[s], bytes, + cudaMemcpyHostToDevice, st[s])); + vadd<<<grid, 256, 0, st[s]>>>(d_a[s], d_b[s], d_c[s], n); + CHECK(cudaGetLastError()); + CHECK(cudaEventRecord(ev[s], st[s])); + } + + CHECK(cudaStreamWaitEvent(st[1], ev[0], 0)); + vadd_into<<<grid, 256, 0, st[1]>>>(d_c[1], d_c[0], n); + CHECK(cudaGetLastError()); + CHECK(cudaEventRecord(ev[1], st[1])); + + CHECK(cudaStreamWaitEvent(st[2], ev[1], 0)); + vadd_into<<<grid, 256, 0, st[2]>>>(d_c[2], d_c[1], n); + CHECK(cudaGetLastError()); + + for (s = 0; s < NSTREAMS; s++) + CHECK(cudaMemcpyAsync(h_c[s], d_c[s], bytes, + cudaMemcpyDeviceToHost, st[s])); + CHECK(cudaEventRecord(t1, st[2])); + + for (s = 0; s < NSTREAMS; s++) + CHECK(cudaStreamSynchronize(st[s])); + CHECK(cudaEventElapsedTime(&ms, t0, t1)); + + for (s = 0; s < NSTREAMS; s++) { + int b = 0; + float exp0[8], got0[8]; + + for (i = 0; i < n; i++) { + float c0 = h_a[0][i] + h_b[0][i]; + float c1 = h_a[1][i] + h_b[1][i] + c0; + float c2 = h_a[2][i] + h_b[2][i] + c1; + float want = (s == 0) ? c0 : (s == 1) ? c1 : c2; + + if (i < 8) { + exp0[i] = want; + got0[i] = h_c[s][i]; + } + if (h_c[s][i] != want) + b++; + } + bad += b; + printf("BUF s=%d bad=%d\n", s, b); + printf("FIRST s=%d exp=", s); + for (i = 0; i < 8; i++) + printf("%.4f ", exp0[i]); + printf(" got="); + for (i = 0; i < 8; i++) + printf("%.4f ", got0[i]); + printf("\n"); + } + + printf("STREAMS3 n=%d bad=%d\n", n, bad); + printf("TIMING ms=%.3f\n", ms); + + if (bad != 0) { + printf("STATUS FAIL\n"); + return 1; + } + printf("STATUS OK\n"); + return 0; +} |