summaryrefslogtreecommitdiffstats
path: root/tools/cuda-ab/p2/src
diff options
context:
space:
mode:
Diffstat (limited to 'tools/cuda-ab/p2/src')
-rw-r--r--tools/cuda-ab/p2/src/common.h66
-rw-r--r--tools/cuda-ab/p2/src/p2_bench.cu95
-rw-r--r--tools/cuda-ab/p2/src/p2_churn.cu100
-rw-r--r--tools/cuda-ab/p2/src/p2_managed_probe.cu81
-rw-r--r--tools/cuda-ab/p2/src/p2_mixed.cu273
-rw-r--r--tools/cuda-ab/p2/src/p2_params.cu185
-rw-r--r--tools/cuda-ab/p2/src/p2_printf_err.cu100
-rw-r--r--tools/cuda-ab/p2/src/p2_seq5.cu127
-rw-r--r--tools/cuda-ab/p2/src/p2_shapes.cu111
-rw-r--r--tools/cuda-ab/p2/src/p2_smem.cu144
-rw-r--r--tools/cuda-ab/p2/src/p2_streams.cu109
11 files changed, 1391 insertions, 0 deletions
diff --git a/tools/cuda-ab/p2/src/common.h b/tools/cuda-ab/p2/src/common.h
new file mode 100644
index 0000000..e1e3137
--- /dev/null
+++ b/tools/cuda-ab/p2/src/common.h
@@ -0,0 +1,66 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * common.h - shared conventions for the P2 "arbitrary CUDA kernels" suite.
+ *
+ * Line conventions (p2/run_guest.sh filters on these):
+ *
+ * compared (stable) RESULT / CHECK / ERR / STATUS / P2PRINT / PROBE / MAIL
+ * informational DEVICE / INFO / TIMING / MEMINFO (ignored by compare)
+ *
+ * Every program accepts --emit-expected: it prints the deterministic
+ * machine-checkable lines a correct GPU run must produce, computed on the
+ * CPU only, without touching CUDA. p2/oracle.sh uses this fallback when the
+ * host has no usable GPU (e.g. the device is owned by a running guest), so
+ * the expected/ references are a documented deterministic computation
+ * rather than a captured bare-metal run.
+ *
+ * All verification in this suite is exact (integers, or bit-identical
+ * float/integer arithmetic and correctly-rounded fmaf), so `bad=0` is
+ * meaningful and no tolerances are used.
+ */
+#ifndef P2_COMMON_H
+#define P2_COMMON_H
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <cstdint>
+#include <cuda_runtime.h>
+
+#define P2_CHECK(call) \
+ do { \
+ cudaError_t e_ = (call); \
+ if (e_ != cudaSuccess) { \
+ printf("ERR cuda=%s code=%d line=%d\n", \
+ cudaGetErrorString(e_), (int)e_, __LINE__); \
+ printf("STATUS FAIL\n"); \
+ fflush(stdout); \
+ return 1; \
+ } \
+ } while (0)
+
+static inline int p2_emit_mode(int argc, char **argv)
+{
+ return (argc > 1 && strcmp(argv[1], "--emit-expected") == 0) ? 1 : 0;
+}
+
+static inline void p2_device_info(void)
+{
+ cudaDeviceProp prop;
+
+ if (cudaGetDeviceProperties(&prop, 0) == cudaSuccess)
+ printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name,
+ prop.major, prop.minor, prop.multiProcessorCount,
+ (unsigned long long)prop.totalGlobalMem);
+ else
+ printf("DEVICE unavailable\n");
+}
+
+static inline int p2_status(int bad)
+{
+ printf("STATUS %s\n", bad == 0 ? "OK" : "FAIL");
+ fflush(stdout);
+ return bad == 0 ? 0 : 1;
+}
+
+#endif /* P2_COMMON_H */
diff --git a/tools/cuda-ab/p2/src/p2_bench.cu b/tools/cuda-ab/p2/src/p2_bench.cu
new file mode 100644
index 0000000..2d1db09
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_bench.cu
@@ -0,0 +1,95 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_bench - compute-heavy FMA benchmark with a fixed, verifiable result.
+ *
+ * 8192 threads each run a 250,000-iteration serial chain of correctly
+ * rounded FMAs (__fmaf_rn on device, fmaf on host), producing 2.05e9 FMAs.
+ * Both implementations are IEEE-correctly-rounded, so the per-thread final
+ * value must match the host exactly (bad=0); the same value is printed for
+ * the A/B comparison. A TIMING line (ignored by the harness comparison)
+ * reports wall-clock ms and GFLOP/s for the performance demonstration.
+ * Device memory: 32 KB output (8 pages).
+ */
+#include "common.h"
+#include <math.h>
+
+#define BENCH_THREADS 8192
+#define BENCH_ITERS 250000
+#define BENCH_A 0.9999f
+#define BENCH_B 1.0e-6f
+
+__global__ void fma_chain(float *out, int iters, int n)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+ float x = 0.0f;
+
+ if (i >= n)
+ return;
+ for (int t = 0; t < iters; t++)
+ x = __fmaf_rn(x, BENCH_A, BENCH_B);
+ out[i] = x;
+}
+
+static float host_chain(int iters)
+{
+ float x = 0.0f;
+
+ for (int t = 0; t < iters; t++)
+ x = fmaf(x, BENCH_A, BENCH_B);
+ return x;
+}
+
+int main(int argc, char **argv)
+{
+ float ref = host_chain(BENCH_ITERS);
+
+ if (p2_emit_mode(argc, argv)) {
+ printf("RESULT bench threads=%d iters=%d bad=0 value=%.9g\n",
+ BENCH_THREADS, BENCH_ITERS, (double)ref);
+ printf("TIMING bench ms=0.000 gflops=0.00 (emit mode: no GPU)\n");
+ return p2_status(0);
+ }
+
+ p2_device_info();
+
+ float *d_out = NULL, *got = (float *)malloc(sizeof(float) * BENCH_THREADS);
+
+ if (!got)
+ return 1;
+ P2_CHECK(cudaMalloc(&d_out, sizeof(float) * BENCH_THREADS));
+
+ cudaEvent_t t0, t1;
+ float ms = 0.0f;
+
+ P2_CHECK(cudaEventCreate(&t0));
+ P2_CHECK(cudaEventCreate(&t1));
+ P2_CHECK(cudaEventRecord(t0));
+ fma_chain<<<BENCH_THREADS / 256, 256>>>(d_out, BENCH_ITERS,
+ BENCH_THREADS);
+ P2_CHECK(cudaGetLastError());
+ P2_CHECK(cudaEventRecord(t1));
+ P2_CHECK(cudaDeviceSynchronize());
+ P2_CHECK(cudaEventElapsedTime(&ms, t0, t1));
+ P2_CHECK(cudaMemcpy(got, d_out, sizeof(float) * BENCH_THREADS,
+ cudaMemcpyDeviceToHost));
+
+ int bad = 0;
+
+ for (int i = 0; i < BENCH_THREADS; i++)
+ if (got[i] != ref)
+ bad++;
+ printf("RESULT bench threads=%d iters=%d bad=%d value=%.9g\n",
+ BENCH_THREADS, BENCH_ITERS, bad, (double)got[0]);
+ if (ms > 0.0f) {
+ double flops = 2.0 * (double)BENCH_THREADS * (double)BENCH_ITERS;
+
+ printf("TIMING bench ms=%.3f gflops=%.2f\n", (double)ms,
+ flops / ((double)ms * 1.0e6));
+ }
+
+ cudaEventDestroy(t0);
+ cudaEventDestroy(t1);
+ cudaFree(d_out);
+ free(got);
+ return p2_status(bad);
+}
diff --git a/tools/cuda-ab/p2/src/p2_churn.cu b/tools/cuda-ab/p2/src/p2_churn.cu
new file mode 100644
index 0000000..7d5f9ef
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_churn.cu
@@ -0,0 +1,100 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_churn - cudaMalloc/cudaFree churn between launches.
+ *
+ * Six iterations allocate 1,2,3,4,1,2 buffers of 2048 ints (8 KB each),
+ * fill each with a seed-derived kernel, verify, then free all of them.
+ * A final fresh allocation + launch proves the allocator remains healthy.
+ * Peak live device memory: 32 KB (4 pages); 14 allocations overall.
+ */
+#include "common.h"
+
+#define CH_N 2048
+#define ITERS 6
+#define FINAL 7
+
+__global__ void fill(int *p, int seed, int n)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < n)
+ p[i] = seed * 1000 + i;
+}
+
+int main(int argc, char **argv)
+{
+ int *got = (int *)malloc(sizeof(int) * CH_N);
+ int bad = 0;
+
+ if (!got)
+ return 1;
+
+ if (p2_emit_mode(argc, argv)) {
+ for (int it = 0; it < ITERS; it++) {
+ int nbuf = 1 + (it % 4);
+
+ printf("RESULT churn iter=%d nbuf=%d n=%d bad=0\n", it,
+ nbuf, CH_N);
+ }
+ printf("RESULT churn final_buf=1 n=%d bad=0\n", CH_N);
+ printf("RESULT churn iters=%d allocs=%d bad=0\n", ITERS,
+ 1 + 2 + 3 + 4 + 1 + 2 + 1);
+ free(got);
+ return p2_status(0);
+ }
+
+ p2_device_info();
+
+ for (int it = 0; it < ITERS; it++) {
+ int nbuf = 1 + (it % 4);
+ int *d[4] = { NULL, NULL, NULL, NULL };
+ int sbad = 0, seed = it + 1;
+
+ for (int k = 0; k < nbuf; k++) {
+ P2_CHECK(cudaMalloc(&d[k], sizeof(int) * CH_N));
+ fill<<<(CH_N + 255) / 256, 256>>>(d[k], seed, CH_N);
+ P2_CHECK(cudaGetLastError());
+ P2_CHECK(cudaDeviceSynchronize());
+ P2_CHECK(cudaMemcpy(got, d[k], sizeof(int) * CH_N,
+ cudaMemcpyDeviceToHost));
+ for (int i = 0; i < CH_N; i++)
+ if (got[i] != seed * 1000 + i)
+ sbad++;
+ }
+ for (int k = 0; k < nbuf; k++)
+ P2_CHECK(cudaFree(d[k]));
+ printf("RESULT churn iter=%d nbuf=%d n=%d bad=%d\n", it, nbuf,
+ CH_N, sbad);
+ bad += sbad;
+ }
+
+ {
+ int *d = NULL;
+ int fbad = 0;
+
+ P2_CHECK(cudaMalloc(&d, sizeof(int) * CH_N));
+ fill<<<(CH_N + 255) / 256, 256>>>(d, FINAL, CH_N);
+ P2_CHECK(cudaGetLastError());
+ P2_CHECK(cudaDeviceSynchronize());
+ P2_CHECK(cudaMemcpy(got, d, sizeof(int) * CH_N,
+ cudaMemcpyDeviceToHost));
+ for (int i = 0; i < CH_N; i++)
+ if (got[i] != FINAL * 1000 + i)
+ fbad++;
+ P2_CHECK(cudaFree(d));
+ printf("RESULT churn final_buf=1 n=%d bad=%d\n", CH_N, fbad);
+ bad += fbad;
+ }
+
+ cudaError_t fin = cudaGetLastError();
+
+ if (fin != cudaSuccess) {
+ printf("ERR churn final_rc=%d name=%s\n", (int)fin,
+ cudaGetErrorString(fin));
+ bad++;
+ }
+ printf("RESULT churn iters=%d allocs=%d bad=%d\n", ITERS,
+ 1 + 2 + 3 + 4 + 1 + 2 + 1, bad);
+ free(got);
+ return p2_status(bad);
+}
diff --git a/tools/cuda-ab/p2/src/p2_managed_probe.cu b/tools/cuda-ab/p2/src/p2_managed_probe.cu
new file mode 100644
index 0000000..0d79352
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_managed_probe.cu
@@ -0,0 +1,81 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_managed_probe - EXPECTED-FAILURE PROBE (documentation, not acceptance).
+ *
+ * Managed (UVM) memory is a known unsupported path for the emulated vGPU.
+ * This probe attempts cudaMallocManaged, records the outcome as a PROBE
+ * line, clears any error, and then proves that ordinary cudaMalloc + launch
+ * still works afterwards. The emulated guest is expected to report a
+ * failure here; the host oracle may report success. p2/run_guest.sh treats
+ * this program as informational (content not diffed), but requires the
+ * program to complete with STATUS OK - i.e. cleanly report the managed
+ * behaviour and keep working - so a hang, crash or poisoned context still
+ * fails the suite.
+ *
+ * Device memory: 4 KB managed (if it works) + 4 KB normal.
+ */
+#include "common.h"
+
+__global__ void fillk(int *p, int v, int n)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < n)
+ p[i] = v + i;
+}
+
+int main(int argc, char **argv)
+{
+ if (p2_emit_mode(argc, argv)) {
+ printf("PROBE managed state=no-oracle-gpu note=emit-mode\n");
+ printf("STATUS OK\n");
+ return 0;
+ }
+
+ p2_device_info();
+
+ int *p = NULL;
+ cudaError_t e = cudaMallocManaged(&p, sizeof(int) * 4);
+
+ if (e == cudaSuccess) {
+ cudaError_t e2, e3;
+ int ok = 0;
+
+ fillk<<<1, 1>>>(p, 5, 4);
+ e2 = cudaGetLastError();
+ e3 = cudaDeviceSynchronize();
+ if (e2 == cudaSuccess && e3 == cudaSuccess)
+ ok = (p[0] == 5 && p[1] == 6 && p[2] == 7 &&
+ p[3] == 8) ? 1 : 0;
+ printf("PROBE managed alloc=ok write_read=%s\n",
+ ok ? "ok" : "bad");
+ cudaFree(p);
+ } else {
+ printf("PROBE managed alloc=fail rc=%d name=%s\n", (int)e,
+ cudaGetErrorString(e));
+ (void)cudaGetLastError();
+ }
+
+ int *d = NULL, bad = 0;
+ cudaError_t e4 = cudaMalloc(&d, sizeof(int) * 4);
+
+ if (e4 == cudaSuccess) {
+ int h[4];
+
+ fillk<<<1, 1>>>(d, 100, 4);
+ (void)cudaGetLastError();
+ cudaDeviceSynchronize();
+ cudaMemcpy(h, d, sizeof(h), cudaMemcpyDeviceToHost);
+ for (int i = 0; i < 4; i++)
+ if (h[i] != 100 + i)
+ bad++;
+ cudaFree(d);
+ printf("PROBE fallback rc=0 bad=%d\n", bad);
+ } else {
+ printf("PROBE fallback rc=%d name=%s\n", (int)e4,
+ cudaGetErrorString(e4));
+ bad = 1;
+ }
+
+ return p2_status(bad);
+}
diff --git a/tools/cuda-ab/p2/src/p2_mixed.cu b/tools/cuda-ab/p2/src/p2_mixed.cu
new file mode 100644
index 0000000..56f993d
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_mixed.cu
@@ -0,0 +1,273 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_mixed - mail + vadd + int reduction + 2D grid + 2 streams in one
+ * process, freeing each phase before the next so the peak live device
+ * footprint stays at the vadd phase (3 x 32 KB = 24 pages).
+ *
+ * phase 1: device-authored mail via a parameter-buffer kernel
+ * phase 2: vectorAdd n=8192 (element-exact fp32)
+ * phase 3: 4096 -> 1 int reduction (shared memory + atomicAdd)
+ * phase 4: 2D grid hash kernel (40x20, block 16x16, partial tiles)
+ * phase 5: 2 streams, addv + scale per stream (exact fp32)
+ */
+#include "common.h"
+
+#define MIX_MAIL "the little gpu man says p2 mixed"
+#define VN 8192
+#define RD_N 4096
+#define G2_NX 40
+#define G2_NY 20
+#define ST_N 1024
+#define ST_NS 2
+
+__global__ void mail_kernel(char *out, const char *msg, int len)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < len)
+ out[i] = msg[i];
+}
+
+__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 reduce_sum(const int *in, int *out, int n)
+{
+ __shared__ int sdata[256];
+ int tid = threadIdx.x;
+ int i = blockIdx.x * blockDim.x + tid;
+
+ sdata[tid] = (i < n) ? in[i] : 0;
+ __syncthreads();
+ for (int s = blockDim.x >> 1; s > 0; s >>= 1) {
+ if (tid < s)
+ sdata[tid] += sdata[tid + s];
+ __syncthreads();
+ }
+ if (tid == 0)
+ atomicAdd(out, sdata[0]);
+}
+
+__global__ void grid2d(int *out, int nx, int ny)
+{
+ int ix = blockIdx.x * blockDim.x + threadIdx.x;
+ int iy = blockIdx.y * blockDim.y + threadIdx.y;
+
+ if (ix < nx && iy < ny) {
+ int idx = iy * nx + ix;
+
+ out[idx] = 1000000 + idx;
+ }
+}
+
+__global__ void addv(float *x, float v, int n)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < n)
+ x[i] = x[i] + v;
+}
+
+__global__ void scale(float *x, float s, int n)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < n)
+ x[i] = x[i] * s;
+}
+
+static int phase_mail(void)
+{
+ char h[128];
+ char *d_out = NULL, *d_msg = NULL;
+ int len = (int)strlen(MIX_MAIL);
+ int bad = 0;
+
+ if (cudaMalloc(&d_out, len) != cudaSuccess)
+ return 1;
+ if (cudaMalloc(&d_msg, len) != cudaSuccess) {
+ cudaFree(d_out);
+ return 1;
+ }
+ cudaMemcpy(d_msg, MIX_MAIL, len, cudaMemcpyHostToDevice);
+ mail_kernel<<<1, 64>>>(d_out, d_msg, len);
+ cudaGetLastError();
+ cudaDeviceSynchronize();
+ cudaMemcpy(h, d_out, len, cudaMemcpyDeviceToHost);
+ bad = (memcmp(h, MIX_MAIL, len) != 0) ? 1 : 0;
+ if (bad == 0)
+ printf("MAIL %s\n", h);
+ cudaFree(d_out);
+ cudaFree(d_msg);
+ printf("RESULT mixed mail bad=%d\n", bad);
+ return bad;
+}
+
+static int phase_vadd(void)
+{
+ float *h_a = (float *)malloc(sizeof(float) * VN);
+ float *h_b = (float *)malloc(sizeof(float) * VN);
+ float *h_c = (float *)malloc(sizeof(float) * VN);
+ float *d_a = NULL, *d_b = NULL, *d_c = NULL;
+ int bad = 0;
+
+ if (!h_a || !h_b || !h_c)
+ return 1;
+ for (int i = 0; i < VN; i++) {
+ h_a[i] = (float)(i % 1000) * 0.5f;
+ h_b[i] = (float)(i % 997) * 0.25f;
+ }
+ if (cudaMalloc(&d_a, sizeof(float) * VN) != cudaSuccess)
+ return 1;
+ if (cudaMalloc(&d_b, sizeof(float) * VN) != cudaSuccess)
+ return 1;
+ if (cudaMalloc(&d_c, sizeof(float) * VN) != cudaSuccess)
+ return 1;
+ cudaMemcpy(d_a, h_a, sizeof(float) * VN, cudaMemcpyHostToDevice);
+ cudaMemcpy(d_b, h_b, sizeof(float) * VN, cudaMemcpyHostToDevice);
+ vadd<<<(VN + 255) / 256, 256>>>(d_a, d_b, d_c, VN);
+ cudaGetLastError();
+ cudaDeviceSynchronize();
+ cudaMemcpy(h_c, d_c, sizeof(float) * VN, cudaMemcpyDeviceToHost);
+ for (int i = 0; i < VN; i++)
+ if (h_c[i] != h_a[i] + h_b[i])
+ bad++;
+ cudaFree(d_a);
+ cudaFree(d_b);
+ cudaFree(d_c);
+ free(h_a);
+ free(h_b);
+ free(h_c);
+ printf("RESULT mixed vadd n=%d bad=%d\n", VN, bad);
+ return bad;
+}
+
+static int phase_reduce(void)
+{
+ int *h_in = (int *)malloc(sizeof(int) * RD_N);
+ int *d_in = NULL, *d_out = NULL;
+ int got = 0, exp = 0, bad = 0;
+
+ if (!h_in)
+ return 1;
+ for (int i = 0; i < RD_N; i++) {
+ h_in[i] = i;
+ exp += i;
+ }
+ if (cudaMalloc(&d_in, sizeof(int) * RD_N) != cudaSuccess)
+ return 1;
+ if (cudaMalloc(&d_out, sizeof(int)) != cudaSuccess)
+ return 1;
+ cudaMemcpy(d_in, h_in, sizeof(int) * RD_N, cudaMemcpyHostToDevice);
+ cudaMemset(d_out, 0, sizeof(int));
+ reduce_sum<<<(RD_N + 255) / 256, 256>>>(d_in, d_out, RD_N);
+ cudaGetLastError();
+ cudaDeviceSynchronize();
+ cudaMemcpy(&got, d_out, sizeof(int), cudaMemcpyDeviceToHost);
+ bad = (got != exp) ? 1 : 0;
+ cudaFree(d_in);
+ cudaFree(d_out);
+ free(h_in);
+ printf("RESULT mixed red n=%d sum=%d bad=%d\n", RD_N, got, bad);
+ return bad;
+}
+
+static int phase_grid2d(void)
+{
+ int n = G2_NX * G2_NY;
+ int *d_out = NULL, *got = (int *)malloc(sizeof(int) * n);
+ int bad = 0;
+
+ if (!got)
+ return 1;
+ if (cudaMalloc(&d_out, sizeof(int) * n) != cudaSuccess) {
+ free(got);
+ return 1;
+ }
+ grid2d<<<dim3(5, 2), dim3(16, 16)>>>(d_out, G2_NX, G2_NY);
+ cudaGetLastError();
+ cudaDeviceSynchronize();
+ cudaMemcpy(got, d_out, sizeof(int) * n, cudaMemcpyDeviceToHost);
+ for (int i = 0; i < n; i++)
+ if (got[i] != 1000000 + i)
+ bad++;
+ cudaFree(d_out);
+ free(got);
+ printf("RESULT mixed grid2d nx=%d ny=%d bad=%d\n", G2_NX, G2_NY,
+ bad);
+ return bad;
+}
+
+static int phase_streams(void)
+{
+ float *h = (float *)malloc(sizeof(float) * ST_N);
+ float *got = (float *)malloc(sizeof(float) * ST_N);
+ cudaStream_t st[ST_NS];
+ float *d_x[ST_NS];
+ int bad = 0;
+
+ if (!h || !got)
+ return 1;
+ for (int i = 0; i < ST_N; i++)
+ h[i] = (float)(i % 100);
+ for (int k = 0; k < ST_NS; k++) {
+ if (cudaStreamCreate(&st[k]) != cudaSuccess)
+ return 1;
+ if (cudaMalloc(&d_x[k], sizeof(float) * ST_N) != cudaSuccess)
+ return 1;
+ cudaMemcpyAsync(d_x[k], h, sizeof(float) * ST_N,
+ cudaMemcpyHostToDevice, st[k]);
+ addv<<<(ST_N + 255) / 256, 256, 0, st[k]>>>(d_x[k], 7.0f, ST_N);
+ scale<<<(ST_N + 255) / 256, 256, 0, st[k]>>>(d_x[k], 3.0f, ST_N);
+ cudaGetLastError();
+ }
+ for (int k = 0; k < ST_NS; k++) {
+ cudaStreamSynchronize(st[k]);
+ cudaMemcpy(got, d_x[k], sizeof(float) * ST_N,
+ cudaMemcpyDeviceToHost);
+ for (int i = 0; i < ST_N; i++)
+ if (got[i] != ((float)(i % 100) + 7.0f) * 3.0f)
+ bad++;
+ cudaStreamDestroy(st[k]);
+ cudaFree(d_x[k]);
+ }
+ free(h);
+ free(got);
+ printf("RESULT mixed stream nstreams=%d n=%d bad=%d\n", ST_NS, ST_N,
+ bad);
+ return bad;
+}
+
+int main(int argc, char **argv)
+{
+ if (p2_emit_mode(argc, argv)) {
+ printf("MAIL %s\n", MIX_MAIL);
+ printf("RESULT mixed mail bad=0\n");
+ printf("RESULT mixed vadd n=%d bad=0\n", VN);
+ printf("RESULT mixed red n=%d sum=%d bad=0\n", RD_N,
+ RD_N * (RD_N - 1) / 2);
+ printf("RESULT mixed grid2d nx=%d ny=%d bad=0\n", G2_NX, G2_NY);
+ printf("RESULT mixed stream nstreams=%d n=%d bad=0\n", ST_NS,
+ ST_N);
+ printf("RESULT mixed total bad=0\n");
+ return p2_status(0);
+ }
+
+ p2_device_info();
+
+ int bad = 0;
+
+ bad += phase_mail();
+ bad += phase_vadd();
+ bad += phase_reduce();
+ bad += phase_grid2d();
+ bad += phase_streams();
+
+ printf("RESULT mixed total bad=%d\n", bad);
+ return p2_status(bad);
+}
diff --git a/tools/cuda-ab/p2/src/p2_params.cu b/tools/cuda-ab/p2/src/p2_params.cu
new file mode 100644
index 0000000..5dd942c
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_params.cu
@@ -0,0 +1,185 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_params - kernel parameter-count and parameter-type coverage.
+ *
+ * k0..k12 have 0..12 value parameters respectively (plus the output
+ * pointer), mixing int, float, double and device-pointer types. Every
+ * kernel writes one int slot of a 13-int device buffer; the CPU reference
+ * mirrors each expression exactly and the buffer must match slot for slot.
+ *
+ * Device memory: out 52 B, p 8 B, pf 4 B (1 page).
+ */
+#include "common.h"
+
+#define NK 13
+
+__global__ void k0(int *out)
+{
+ out[0] = 7;
+}
+
+__global__ void k1(int *out, int a)
+{
+ out[1] = a + 1;
+}
+
+__global__ void k2(int *out, int a, int b)
+{
+ out[2] = a + b;
+}
+
+__global__ void k3(int *out, int a, int b, float f1)
+{
+ out[3] = a + b + (int)f1;
+}
+
+__global__ void k4(int *out, int a, int b, float f1, double c1)
+{
+ out[4] = a + b + (int)f1 + (int)c1;
+}
+
+__global__ void k5(int *out, int a, int b, float f1, double c1, int d)
+{
+ out[5] = a + b + (int)f1 + (int)c1 + d;
+}
+
+__global__ void k6(int *out, int a, int b, float f1, double c1, int d,
+ const int *p)
+{
+ out[6] = a + b + (int)f1 + (int)c1 + d + p[0];
+}
+
+__global__ void k7(int *out, int a, int b, float f1, double c1, int d,
+ const int *p, const int *q)
+{
+ out[7] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0];
+}
+
+__global__ void k8(int *out, int a, int b, float f1, double c1, int d,
+ const int *p, const int *q, int n)
+{
+ out[8] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n;
+}
+
+__global__ void k9(int *out, int a, int b, float f1, double c1, int d,
+ const int *p, const int *q, int n, float f2)
+{
+ out[9] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + (int)f2;
+}
+
+__global__ void k10(int *out, int a, int b, float f1, double c1, int d,
+ const int *p, const int *q, int n, float f2,
+ double c2)
+{
+ out[10] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
+ (int)f2 + (int)c2;
+}
+
+__global__ void k11(int *out, int a, int b, float f1, double c1, int d,
+ const int *p, const int *q, int n, float f2,
+ double c2, float s)
+{
+ out[11] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
+ (int)f2 + (int)c2 + (int)s;
+}
+
+__global__ void k12(int *out, int a, int b, float f1, double c1, int d,
+ const int *p, const int *q, int n, float f2,
+ double c2, float s, int e)
+{
+ out[12] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
+ (int)f2 + (int)c2 + (int)s + e;
+}
+
+/* CPU mirror of the 13 kernel expressions; keeps expected values exact. */
+static void expected(int *exp, int a, int b, float f1, double c1, int d,
+ const int *p, const int *q, int n, float f2, double c2,
+ float s, int e)
+{
+ exp[0] = 7;
+ exp[1] = a + 1;
+ exp[2] = a + b;
+ exp[3] = a + b + (int)f1;
+ exp[4] = a + b + (int)f1 + (int)c1;
+ exp[5] = a + b + (int)f1 + (int)c1 + d;
+ exp[6] = a + b + (int)f1 + (int)c1 + d + p[0];
+ exp[7] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0];
+ exp[8] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n;
+ exp[9] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n + (int)f2;
+ exp[10] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
+ (int)f2 + (int)c2;
+ exp[11] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
+ (int)f2 + (int)c2 + (int)s;
+ exp[12] = a + b + (int)f1 + (int)c1 + d + p[0] + q[0] + n +
+ (int)f2 + (int)c2 + (int)s + e;
+}
+
+static int checksum(const int *v)
+{
+ int sum = 0;
+
+ for (int i = 0; i < NK; i++)
+ sum += v[i];
+ return sum;
+}
+
+int main(int argc, char **argv)
+{
+ int a = 3, b = 5, d = 11, n = 2, e = 4;
+ float f1 = 6.5f, f2 = 8.5f, s = 2.5f;
+ double c1 = 2.5, c2 = 3.5;
+ int hp[2] = { 5, 6 }, hq[2] = { 6, 7 };
+ int exp[NK];
+
+ expected(exp, a, b, f1, c1, d, hp, hq, n, f2, c2, s, e);
+
+ if (p2_emit_mode(argc, argv)) {
+ printf("RESULT params kernels=%d nparam_max=12 bad=0 checksum=%d\n",
+ NK, checksum(exp));
+ return p2_status(0);
+ }
+
+ p2_device_info();
+
+ int *d_out = NULL, *got = (int *)malloc(sizeof(int) * NK);
+ int *dp = NULL, *dq = NULL;
+
+ if (!got)
+ return 1;
+ P2_CHECK(cudaMalloc(&d_out, sizeof(int) * NK));
+ P2_CHECK(cudaMalloc(&dp, sizeof(int) * 2));
+ P2_CHECK(cudaMalloc(&dq, sizeof(int) * 2));
+ P2_CHECK(cudaMemcpy(dp, hp, sizeof(int) * 2, cudaMemcpyHostToDevice));
+ P2_CHECK(cudaMemcpy(dq, hq, sizeof(int) * 2, cudaMemcpyHostToDevice));
+ P2_CHECK(cudaMemset(d_out, 0, sizeof(int) * NK));
+
+ k0<<<1, 1>>>(d_out);
+ k1<<<1, 1>>>(d_out, a);
+ k2<<<1, 1>>>(d_out, a, b);
+ k3<<<1, 1>>>(d_out, a, b, f1);
+ k4<<<1, 1>>>(d_out, a, b, f1, c1);
+ k5<<<1, 1>>>(d_out, a, b, f1, c1, d);
+ k6<<<1, 1>>>(d_out, a, b, f1, c1, d, dp);
+ k7<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq);
+ k8<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n);
+ k9<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2);
+ k10<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2, c2);
+ k11<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2, c2, s);
+ k12<<<1, 1>>>(d_out, a, b, f1, c1, d, dp, dq, n, f2, c2, s, e);
+ P2_CHECK(cudaGetLastError());
+ P2_CHECK(cudaDeviceSynchronize());
+ P2_CHECK(cudaMemcpy(got, d_out, sizeof(int) * NK,
+ cudaMemcpyDeviceToHost));
+
+ int bad = 0;
+
+ for (int i = 0; i < NK; i++)
+ if (got[i] != exp[i])
+ bad++;
+ printf("RESULT params kernels=%d nparam_max=12 bad=%d checksum=%d\n",
+ NK, bad, checksum(got));
+
+ cudaFree(d_out); cudaFree(dp); cudaFree(dq);
+ free(got);
+ return p2_status(bad);
+}
diff --git a/tools/cuda-ab/p2/src/p2_printf_err.cu b/tools/cuda-ab/p2/src/p2_printf_err.cu
new file mode 100644
index 0000000..d7f449e
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_printf_err.cu
@@ -0,0 +1,100 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_printf_err - device printf plus an intentional launch error path.
+ *
+ * 1. A 4-thread kernel prints deterministic P2PRINT lines from the device.
+ * 2. A kernel launched with gridDim.x == 0 must be rejected:
+ * cudaErrorInvalidConfiguration (9), and the error must be cleared by
+ * cudaGetLastError. A blockDim.x == 0 launch is exercised too.
+ * 3. A normal launch after the errors proves error recovery.
+ *
+ * Expected error strings come from the statically linked CUDA runtime, so
+ * they must be identical on host and guest.
+ * Device memory: out 4 B.
+ */
+#include "common.h"
+
+__global__ void hello(int nthreads)
+{
+ int i = threadIdx.x;
+
+ if (i < nthreads)
+ printf("P2PRINT hello i=%d\n", i);
+}
+
+__global__ void second(void)
+{
+ if (threadIdx.x == 0)
+ printf("P2PRINT second c=%c v=%d\n", 'Z', 7);
+}
+
+__global__ void simple(int *out)
+{
+ out[0] = 42;
+}
+
+int main(int argc, char **argv)
+{
+ if (p2_emit_mode(argc, argv)) {
+ printf("P2PRINT hello i=0\n");
+ printf("P2PRINT hello i=1\n");
+ printf("P2PRINT hello i=2\n");
+ printf("P2PRINT hello i=3\n");
+ printf("P2PRINT second c=Z v=7\n");
+ printf("ERR launch grid0 rc=%d name=%s\n",
+ (int)cudaErrorInvalidConfiguration,
+ cudaGetErrorString(cudaErrorInvalidConfiguration));
+ printf("ERR launch block0 rejected=1\n");
+ printf("ERR cleared rc=0\n");
+ printf("RESULT errmode grid0_bad=0 block0_bad=0 clear_bad=0 recovered=1\n");
+ return p2_status(0);
+ }
+
+ p2_device_info();
+
+ int *d_out = NULL;
+ P2_CHECK(cudaMalloc(&d_out, sizeof(int)));
+
+ hello<<<1, 4>>>(4);
+ P2_CHECK(cudaGetLastError());
+ P2_CHECK(cudaDeviceSynchronize());
+ second<<<1, 4>>>();
+ P2_CHECK(cudaGetLastError());
+ P2_CHECK(cudaDeviceSynchronize());
+
+ (void)cudaGetLastError();
+
+ int badgrid = 0, badblock = 0;
+ cudaError_t e_grid, e_block, cleared;
+ int got = 0;
+
+ simple<<<badgrid, 1>>>(d_out);
+ e_grid = cudaGetLastError();
+ printf("ERR launch grid0 rc=%d name=%s\n", (int)e_grid,
+ cudaGetErrorString(e_grid));
+
+ simple<<<1, badblock>>>(d_out);
+ e_block = cudaGetLastError();
+ printf("ERR launch block0 rejected=%d\n",
+ (e_block != cudaSuccess) ? 1 : 0);
+
+ cleared = cudaGetLastError();
+ printf("ERR cleared rc=%d\n", (int)cleared);
+
+ simple<<<1, 1>>>(d_out);
+ P2_CHECK(cudaGetLastError());
+ P2_CHECK(cudaDeviceSynchronize());
+ P2_CHECK(cudaMemcpy(&got, d_out, sizeof(int), cudaMemcpyDeviceToHost));
+
+ int grid0_bad = (e_grid != cudaErrorInvalidConfiguration) ? 1 : 0;
+ int block0_bad = (e_block == cudaSuccess) ? 1 : 0;
+ int clear_bad = (cleared != cudaSuccess) ? 1 : 0;
+ int recovered = (got == 42) ? 1 : 0;
+ int bad = grid0_bad + block0_bad + clear_bad + (recovered ? 0 : 1);
+
+ printf("RESULT errmode grid0_bad=%d block0_bad=%d clear_bad=%d recovered=%d\n",
+ grid0_bad, block0_bad, clear_bad, recovered);
+
+ cudaFree(d_out);
+ return p2_status(bad);
+}
diff --git a/tools/cuda-ab/p2/src/p2_seq5.cu b/tools/cuda-ab/p2/src/p2_seq5.cu
new file mode 100644
index 0000000..8a778f4
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_seq5.cu
@@ -0,0 +1,127 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_seq5 - five sequential kernel launches in one process on the same
+ * buffers, verifying the device state after every launch.
+ *
+ * step 1: c = a + b
+ * step 2: c = c + a
+ * step 3: c = c + b
+ * step 4: c = c * 2
+ * step 5: c = c + 1
+ *
+ * Inputs are small integers so every operation is exact in fp32; the CPU
+ * reference reproduces the same chain and must match element for element.
+ * Device memory: 3 x 4 KB pointers (3 pages total).
+ */
+#include "common.h"
+
+#define N 1024
+
+__global__ void kadd(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 kmul2(float *c, int n)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < n)
+ c[i] = c[i] * 2.0f;
+}
+
+__global__ void kaddv(float *c, float v, int n)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < n)
+ c[i] = c[i] + v;
+}
+
+static void make_inputs(float *a, float *b)
+{
+ for (int i = 0; i < N; i++) {
+ a[i] = (float)(i % 7);
+ b[i] = (float)(i % 5);
+ }
+}
+
+static void step_cpu(float *c, const float *a, const float *b, int step)
+{
+ for (int i = 0; i < N; i++) {
+ switch (step) {
+ case 1: c[i] = a[i] + b[i]; break;
+ case 2: c[i] = c[i] + a[i]; break;
+ case 3: c[i] = c[i] + b[i]; break;
+ case 4: c[i] = c[i] * 2.0f; break;
+ case 5: c[i] = c[i] + 1.0f; break;
+ }
+ }
+}
+
+int main(int argc, char **argv)
+{
+ float *h_a, *h_b, *h_c, *h_exp;
+ int bad = 0;
+
+ if (p2_emit_mode(argc, argv)) {
+ h_a = (float *)malloc(sizeof(float) * N);
+ h_b = (float *)malloc(sizeof(float) * N);
+ h_exp = (float *)malloc(sizeof(float) * N);
+ make_inputs(h_a, h_b);
+ for (int step = 1; step <= 5; step++) {
+ step_cpu(h_exp, h_a, h_b, step);
+ printf("RESULT seq step=%d n=%d bad=0\n", step, N);
+ }
+ printf("RESULT seq launches=5 n=%d bad=0\n", N);
+ free(h_a); free(h_b); free(h_exp);
+ return p2_status(0);
+ }
+
+ p2_device_info();
+ h_a = (float *)malloc(sizeof(float) * N);
+ h_b = (float *)malloc(sizeof(float) * N);
+ h_c = (float *)malloc(sizeof(float) * N);
+ h_exp = (float *)malloc(sizeof(float) * N);
+ if (!h_a || !h_b || !h_c || !h_exp)
+ return 1;
+ make_inputs(h_a, h_b);
+
+ float *d_a = NULL, *d_b = NULL, *d_c = NULL;
+
+ P2_CHECK(cudaMalloc(&d_a, sizeof(float) * N));
+ P2_CHECK(cudaMalloc(&d_b, sizeof(float) * N));
+ P2_CHECK(cudaMalloc(&d_c, sizeof(float) * N));
+ P2_CHECK(cudaMemcpy(d_a, h_a, sizeof(float) * N, cudaMemcpyHostToDevice));
+ P2_CHECK(cudaMemcpy(d_b, h_b, sizeof(float) * N, cudaMemcpyHostToDevice));
+
+ for (int step = 1; step <= 5; step++) {
+ int sbad = 0;
+
+ switch (step) {
+ case 1: kadd<<<(N + 255) / 256, 256>>>(d_a, d_b, d_c, N); break;
+ case 2: kadd<<<(N + 255) / 256, 256>>>(d_c, d_a, d_c, N); break;
+ case 3: kadd<<<(N + 255) / 256, 256>>>(d_c, d_b, d_c, N); break;
+ case 4: kmul2<<<(N + 255) / 256, 256>>>(d_c, N); break;
+ case 5: kaddv<<<(N + 255) / 256, 256>>>(d_c, 1.0f, N); break;
+ }
+ P2_CHECK(cudaGetLastError());
+ P2_CHECK(cudaDeviceSynchronize());
+ P2_CHECK(cudaMemcpy(h_c, d_c, sizeof(float) * N,
+ cudaMemcpyDeviceToHost));
+ step_cpu(h_exp, h_a, h_b, step);
+ for (int i = 0; i < N; i++)
+ if (h_c[i] != h_exp[i])
+ sbad++;
+ printf("RESULT seq step=%d n=%d bad=%d\n", step, N, sbad);
+ bad += sbad;
+ }
+ printf("RESULT seq launches=5 n=%d bad=%d\n", N, bad);
+
+ cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
+ free(h_a); free(h_b); free(h_c); free(h_exp);
+ return p2_status(bad);
+}
diff --git a/tools/cuda-ab/p2/src/p2_shapes.cu b/tools/cuda-ab/p2/src/p2_shapes.cu
new file mode 100644
index 0000000..93ae0e9
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_shapes.cu
@@ -0,0 +1,111 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_shapes - grid/block shape variation with a bound-checked index-hash
+ * kernel:
+ * case 1: 1x1x1 grid, 1x1x1 block (minimum launch)
+ * case 2: odd sizes (n=1013, block 33, 31 blocks)
+ * case 3: non-multiple n (n=1000, block 64, partial last block)
+ * case 4: 2D grid (40x20, block 16x16, partial tiles)
+ * case 5: 3D grid (24x8x8, block 8x4x2)
+ *
+ * Device memory: one 8 KB int buffer reused across cases (2 pages).
+ */
+#include "common.h"
+
+#define MAXN 2048
+
+struct shape {
+ int nx, ny, nz; /* logical extent */
+ int bx, by, bz; /* block dims */
+ int gx, gy, gz; /* grid dims */
+ int tag;
+};
+
+static const struct shape cases[] = {
+ { 1, 1, 1, 1, 1, 1, 1, 1, 1, 1 },
+ {1013, 1, 1, 33, 1, 1, 31, 1, 1, 2 },
+ {1000, 1, 1, 64, 1, 1, 16, 1, 1, 3 },
+ { 40, 20, 1, 16, 16, 1, 5, 2, 1, 4 },
+ { 24, 8, 8, 8, 4, 2, 3, 2, 4, 5 },
+};
+
+#define NCASES ((int)(sizeof(cases) / sizeof(cases[0])))
+
+__global__ void shape_kernel(int *out, int nx, int ny, int nz, int tag)
+{
+ int ix = blockIdx.x * blockDim.x + threadIdx.x;
+ int iy = blockIdx.y * blockDim.y + threadIdx.y;
+ int iz = blockIdx.z * blockDim.z + threadIdx.z;
+
+ if (ix < nx && iy < ny && iz < nz) {
+ int idx = (iz * ny + iy) * nx + ix;
+
+ out[idx] = tag * 1000000 + idx;
+ }
+}
+
+static void shape_expected(int *exp, const struct shape *s)
+{
+ for (int iz = 0; iz < s->nz; iz++)
+ for (int iy = 0; iy < s->ny; iy++)
+ for (int ix = 0; ix < s->nx; ix++) {
+ int idx = (iz * s->ny + iy) * s->nx + ix;
+
+ exp[idx] = s->tag * 1000000 + idx;
+ }
+}
+
+int main(int argc, char **argv)
+{
+ int bad = 0;
+
+ if (p2_emit_mode(argc, argv)) {
+ int *exp = (int *)malloc(sizeof(int) * MAXN);
+
+ for (int c = 0; c < NCASES; c++) {
+ int n = cases[c].nx * cases[c].ny * cases[c].nz;
+
+ shape_expected(exp, &cases[c]);
+ printf("RESULT shapes case=%d n=%d bad=0\n",
+ cases[c].tag, n);
+ }
+ printf("RESULT shapes cases=%d bad=0\n", NCASES);
+ free(exp);
+ return p2_status(0);
+ }
+
+ p2_device_info();
+
+ int *d_out = NULL, *h_got = (int *)malloc(sizeof(int) * MAXN);
+ int *h_exp = (int *)malloc(sizeof(int) * MAXN);
+
+ if (!h_got || !h_exp)
+ return 1;
+ P2_CHECK(cudaMalloc(&d_out, sizeof(int) * MAXN));
+
+ for (int c = 0; c < NCASES; c++) {
+ const struct shape *s = &cases[c];
+ int n = s->nx * s->ny * s->nz;
+ dim3 grid(s->gx, s->gy, s->gz), block(s->bx, s->by, s->bz);
+ int sbad = 0;
+
+ shape_kernel<<<grid, block>>>(d_out, s->nx, s->ny, s->nz,
+ s->tag);
+ P2_CHECK(cudaGetLastError());
+ P2_CHECK(cudaDeviceSynchronize());
+ P2_CHECK(cudaMemcpy(h_got, d_out, sizeof(int) * n,
+ cudaMemcpyDeviceToHost));
+ memset(h_exp, 0, sizeof(int) * n);
+ shape_expected(h_exp, s);
+ for (int i = 0; i < n; i++)
+ if (h_got[i] != h_exp[i])
+ sbad++;
+ printf("RESULT shapes case=%d n=%d bad=%d\n", s->tag, n, sbad);
+ bad += sbad;
+ }
+ printf("RESULT shapes cases=%d bad=%d\n", NCASES, bad);
+
+ cudaFree(d_out);
+ free(h_got); free(h_exp);
+ return p2_status(bad);
+}
diff --git a/tools/cuda-ab/p2/src/p2_smem.cu b/tools/cuda-ab/p2/src/p2_smem.cu
new file mode 100644
index 0000000..5db0d52
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_smem.cu
@@ -0,0 +1,144 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_smem - shared-memory kernels:
+ * reduce_sum: 4096 int32 -> single sum, 16 blocks x 256 threads, shared
+ * tree reduction plus a global atomicAdd.
+ * gemm_tiled: 32x32 int32 GEMM with 16x16 shared-memory tiles,
+ * grid(2,2) x block(16,16).
+ *
+ * All arithmetic is exact in int32 (inputs are small non-negative ints), so
+ * device results must equal the CPU reference bit for bit.
+ * Device memory: in 16 KB, out 4 B, then A/B/C 4 KB each (peak 5 pages).
+ */
+#include "common.h"
+
+#define R_N 4096
+#define TS 16
+#define M 32
+#define K 32
+#define NN 32
+
+__global__ void reduce_sum(const int *in, int *out, int n)
+{
+ __shared__ int sdata[256];
+ int tid = threadIdx.x;
+ int i = blockIdx.x * blockDim.x + tid;
+
+ sdata[tid] = (i < n) ? in[i] : 0;
+ __syncthreads();
+ for (int s = blockDim.x >> 1; s > 0; s >>= 1) {
+ if (tid < s)
+ sdata[tid] += sdata[tid + s];
+ __syncthreads();
+ }
+ if (tid == 0)
+ atomicAdd(out, sdata[0]);
+}
+
+__global__ void gemm_tiled(const int *A, const int *B, int *C)
+{
+ __shared__ int As[TS][TS], Bs[TS][TS];
+ int row = blockIdx.y * TS + threadIdx.y;
+ int col = blockIdx.x * TS + threadIdx.x;
+ int sum = 0;
+
+ for (int t = 0; t < K / TS; t++) {
+ As[threadIdx.y][threadIdx.x] = A[row * K + t * TS + threadIdx.x];
+ Bs[threadIdx.y][threadIdx.x] = B[(t * TS + threadIdx.y) * NN + col];
+ __syncthreads();
+ for (int j = 0; j < TS; j++)
+ sum += As[threadIdx.y][j] * Bs[j][threadIdx.x];
+ __syncthreads();
+ }
+ C[row * NN + col] = sum;
+}
+
+int main(int argc, char **argv)
+{
+ int *h_in = (int *)malloc(sizeof(int) * R_N);
+ int *h_A = (int *)malloc(sizeof(int) * M * K);
+ int *h_B = (int *)malloc(sizeof(int) * K * NN);
+ int *h_C = (int *)malloc(sizeof(int) * M * NN);
+ int red_sum = 0, gemm_sum = 0, bad = 0;
+
+ if (!h_in || !h_A || !h_B || !h_C)
+ return 1;
+ for (int i = 0; i < R_N; i++)
+ h_in[i] = i;
+ for (int i = 0; i < M * K; i++)
+ h_A[i] = i % 5;
+ for (int i = 0; i < K * NN; i++)
+ h_B[i] = i % 3;
+ for (int i = 0; i < R_N; i++)
+ red_sum += h_in[i];
+ for (int r = 0; r < M; r++)
+ for (int c = 0; c < NN; c++) {
+ int s = 0;
+
+ for (int t = 0; t < K; t++)
+ s += h_A[r * K + t] * h_B[t * NN + c];
+ h_C[r * NN + c] = s;
+ gemm_sum += s;
+ }
+
+ if (p2_emit_mode(argc, argv)) {
+ printf("RESULT smem reduce n=%d sum=%d bad=0\n", R_N, red_sum);
+ printf("RESULT smem gemm m=%d n=%d k=%d checksum=%d bad=0\n",
+ M, NN, K, gemm_sum);
+ free(h_in); free(h_A); free(h_B); free(h_C);
+ return p2_status(0);
+ }
+
+ p2_device_info();
+
+ int *d_in = NULL, *d_out = NULL, *d_A = NULL, *d_B = NULL, *d_C = NULL;
+ int got_sum = 0, got_gemm = 0;
+
+ P2_CHECK(cudaMalloc(&d_in, sizeof(int) * R_N));
+ P2_CHECK(cudaMalloc(&d_out, sizeof(int)));
+ P2_CHECK(cudaMemcpy(d_in, h_in, sizeof(int) * R_N,
+ cudaMemcpyHostToDevice));
+ P2_CHECK(cudaMemset(d_out, 0, sizeof(int)));
+ reduce_sum<<<(R_N + 255) / 256, 256>>>(d_in, d_out, R_N);
+ P2_CHECK(cudaGetLastError());
+ P2_CHECK(cudaDeviceSynchronize());
+ P2_CHECK(cudaMemcpy(&got_sum, d_out, sizeof(int),
+ cudaMemcpyDeviceToHost));
+ printf("RESULT smem reduce n=%d sum=%d bad=%d\n", R_N, got_sum,
+ got_sum == red_sum ? 0 : 1);
+ bad += (got_sum == red_sum) ? 0 : 1;
+ cudaFree(d_in); cudaFree(d_out);
+
+ P2_CHECK(cudaMalloc(&d_A, sizeof(int) * M * K));
+ P2_CHECK(cudaMalloc(&d_B, sizeof(int) * K * NN));
+ P2_CHECK(cudaMalloc(&d_C, sizeof(int) * M * NN));
+ P2_CHECK(cudaMemcpy(d_A, h_A, sizeof(int) * M * K,
+ cudaMemcpyHostToDevice));
+ P2_CHECK(cudaMemcpy(d_B, h_B, sizeof(int) * K * NN,
+ cudaMemcpyHostToDevice));
+ gemm_tiled<<<dim3(2, 2), dim3(TS, TS)>>>(d_A, d_B, d_C);
+ P2_CHECK(cudaGetLastError());
+ P2_CHECK(cudaDeviceSynchronize());
+
+ int *got_C = (int *)malloc(sizeof(int) * M * NN);
+
+ if (!got_C)
+ return 1;
+ P2_CHECK(cudaMemcpy(got_C, d_C, sizeof(int) * M * NN,
+ cudaMemcpyDeviceToHost));
+ int cbad = 0;
+
+ for (int i = 0; i < M * NN; i++) {
+ got_gemm += got_C[i];
+ if (got_C[i] != h_C[i])
+ cbad++;
+ }
+ printf("RESULT smem gemm m=%d n=%d k=%d checksum=%d bad=%d\n",
+ M, NN, K, got_gemm, cbad);
+ bad += cbad;
+
+ cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
+ free(got_C);
+ free(h_in); free(h_A); free(h_B); free(h_C);
+ return p2_status(bad);
+}
diff --git a/tools/cuda-ab/p2/src/p2_streams.cu b/tools/cuda-ab/p2/src/p2_streams.cu
new file mode 100644
index 0000000..09d6ed3
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_streams.cu
@@ -0,0 +1,109 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_streams - three created streams, two kernels per stream (scale then
+ * add), independent 1024-float buffers, per-stream events.
+ *
+ * Expected value per element: ((i % 100) * 2.0f) + 1.0f (exact in fp32).
+ * After all streams complete, each event must query as complete.
+ * Device memory: 3 x 4 KB (3 pages).
+ */
+#include "common.h"
+
+#define SN 1024
+#define NS 3
+
+__global__ void scale(float *x, float s, int n)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < n)
+ x[i] = x[i] * s;
+}
+
+__global__ void addv(float *x, float v, int n)
+{
+ int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+ if (i < n)
+ x[i] = x[i] + v;
+}
+
+static void fill(float *h, int n)
+{
+ for (int i = 0; i < n; i++)
+ h[i] = (float)(i % 100);
+}
+
+static float expected1(int i)
+{
+ return ((float)(i % 100) * 2.0f) + 1.0f;
+}
+
+int main(int argc, char **argv)
+{
+ float *h = (float *)malloc(sizeof(float) * SN);
+ float *got = (float *)malloc(sizeof(float) * SN);
+ int bad = 0;
+
+ if (!h || !got)
+ return 1;
+ fill(h, SN);
+
+ if (p2_emit_mode(argc, argv)) {
+ printf("RESULT streams nstreams=%d launches=%d n=%d bad=0\n",
+ NS, NS * 2, SN);
+ printf("RESULT streams events=%d bad=0\n", NS);
+ free(h); free(got);
+ return p2_status(0);
+ }
+
+ p2_device_info();
+
+ cudaStream_t st[NS];
+ cudaEvent_t ev[NS];
+ float *d_x[NS];
+
+ for (int k = 0; k < NS; k++) {
+ P2_CHECK(cudaStreamCreate(&st[k]));
+ P2_CHECK(cudaEventCreate(&ev[k]));
+ P2_CHECK(cudaMalloc(&d_x[k], sizeof(float) * SN));
+ P2_CHECK(cudaMemcpyAsync(d_x[k], h, sizeof(float) * SN,
+ cudaMemcpyHostToDevice, st[k]));
+ scale<<<(SN + 255) / 256, 256, 0, st[k]>>>(d_x[k], 2.0f, SN);
+ addv<<<(SN + 255) / 256, 256, 0, st[k]>>>(d_x[k], 1.0f, SN);
+ P2_CHECK(cudaGetLastError());
+ P2_CHECK(cudaEventRecord(ev[k], st[k]));
+ }
+
+ int ibad = 0;
+
+ for (int k = 0; k < NS; k++) {
+ cudaError_t q;
+ int sbad = 0;
+
+ P2_CHECK(cudaStreamSynchronize(st[k]));
+ P2_CHECK(cudaMemcpy(got, d_x[k], sizeof(float) * SN,
+ cudaMemcpyDeviceToHost));
+ for (int i = 0; i < SN; i++)
+ if (got[i] != expected1(i))
+ sbad++;
+ q = cudaEventQuery(ev[k]);
+ if (q != cudaSuccess)
+ sbad++;
+ printf("RESULT streams stream=%d query_rc=%d bad=%d\n", k, (int)q,
+ sbad);
+ ibad += sbad;
+ }
+ printf("RESULT streams nstreams=%d launches=%d n=%d bad=%d\n",
+ NS, NS * 2, SN, ibad);
+ printf("RESULT streams events=%d bad=%d\n", NS, ibad == 0 ? 0 : 1);
+ bad = ibad;
+
+ for (int k = 0; k < NS; k++) {
+ cudaEventDestroy(ev[k]);
+ cudaStreamDestroy(st[k]);
+ cudaFree(d_x[k]);
+ }
+ free(h); free(got);
+ return p2_status(bad);
+}
generated by cgit v1.2.3 (git 2.39.1) at 2026-10-12 11:35:22 +0000