tools/cuda-ab: P2 arbitrary-kernel suite expansion (regs/streams/printf), smiley push wiring, near-native perf harness (offline readiness; guest validation pending)

3 files changed, 292 insertions, 9 deletions

diff --git a/tools/cuda-ab/p2/src/p2_printf_err.cu b/tools/cuda-ab/p2/src/p2_printf_err.cu
index d7f449e..dea2fd5 100644
--- a/tools/cuda-ab/p2/src/p2_printf_err.cu
+++ b/tools/cuda-ab/p2/src/p2_printf_err.cu
@@ -1,18 +1,23 @@
 // SPDX-License-Identifier: BSD-3-Clause
 /*
- * p2_printf_err - device printf plus an intentional launch error path.
+ * p2_printf_err - device printf, a passing device assert, and 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:
+ *   2. A kernel with assert(n == 4) exercises the device-side assert
+ *      machinery on its passing path (a fired assert is deliberately not
+ *      tested: it traps the kernel and poisons the context).
+ *   3. 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.
+ *   4. 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"
+#include <cassert>
 
 __global__ void hello(int nthreads)
 {
@@ -33,6 +38,13 @@ __global__ void simple(int *out)
 	out[0] = 42;
 }
 
+__global__ void assertk(int *out, int n)
+{
+	assert(n == 4);
+	if (threadIdx.x == 0)
+		out[0] = n;
+}
+
 int main(int argc, char **argv)
 {
 	if (p2_emit_mode(argc, argv)) {
@@ -47,6 +59,7 @@ int main(int argc, char **argv)
 		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");
+		printf("RESULT assert n=4 bad=0\n");
 		return p2_status(0);
 	}
 
@@ -90,10 +103,25 @@ int main(int argc, char **argv)
 	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);
+	int assert_bad = 0;
+
+	assertk<<<1, 1>>>(d_out, 4);
+	P2_CHECK(cudaGetLastError());
+	P2_CHECK(cudaDeviceSynchronize());
+	{
+		int agot = 0;
+
+		P2_CHECK(cudaMemcpy(&agot, d_out, sizeof(int),
+				    cudaMemcpyDeviceToHost));
+		assert_bad = (agot == 4) ? 0 : 1;
+	}
+
+	int bad = grid0_bad + block0_bad + clear_bad + (recovered ? 0 : 1) +
+		  assert_bad;
 
 	printf("RESULT errmode grid0_bad=%d block0_bad=%d clear_bad=%d recovered=%d\n",
 	       grid0_bad, block0_bad, clear_bad, recovered);
+	printf("RESULT assert n=4 bad=%d\n", assert_bad);
 
 	cudaFree(d_out);
 	return p2_status(bad);
diff --git a/tools/cuda-ab/p2/src/p2_regs.cu b/tools/cuda-ab/p2/src/p2_regs.cu
new file mode 100644
index 0000000..d463295
--- /dev/null
+++ b/tools/cuda-ab/p2/src/p2_regs.cu
@@ -0,0 +1,179 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * p2_regs - register-usage / launch-bounds variants.
+ *
+ * Three kernels cover different register-allocation regimes while staying
+ * exact (unsigned wraparound is well defined; no floating point):
+ *
+ *   regs_lo     __launch_bounds__(256, 4) - strongly capped register use
+ *               (<= 64 registers/thread on the 65536-register file)
+ *   regs_hi     64 independent accumulators x 128 update steps - high
+ *               register pressure (the compiler may spill; correctness is
+ *               unaffected by where the values live)
+ *   regs_bounds __launch_bounds__(128, 4) - a second occupancy target with a
+ *               different block size
+ *
+ * Every kernel writes one unsigned per element of a 1024-element buffer; the
+ * CPU mirrors the exact integer expressions, so the buffer must match
+ * element for element. INFO lines report the compiler's register counts and
+ * the queried occupancy (informational; ignored by the A/B comparison).
+ *
+ * Device memory: one 4 KB unsigned buffer (1 page).
+ */
+#include "common.h"
+
+#define RG_N 1024
+
+__global__ void __launch_bounds__(256, 4) regs_lo(unsigned *out, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		out[i] = ((unsigned)i ^ 0x5a5au) + 3u;
+}
+
+__global__ void regs_hi(unsigned *out, int n, unsigned seed)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+	unsigned acc[64];
+	unsigned s = 0;
+
+	if (i >= n)
+		return;
+#pragma unroll
+	for (int k = 0; k < 64; k++)
+		acc[k] = seed + (unsigned)k * ((unsigned)i + 1u);
+	for (int t = 0; t < 128; t++) {
+#pragma unroll
+		for (int k = 0; k < 64; k++)
+			acc[k] = acc[k] * 2654435761u + ((unsigned)t ^ (unsigned)k);
+	}
+#pragma unroll
+	for (int k = 0; k < 64; k++)
+		s += acc[k];
+	out[i] = s;
+}
+
+__global__ void __launch_bounds__(128, 4) regs_bounds(unsigned *out, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		out[i] = (unsigned)(i * 5) + 7u;
+}
+
+static unsigned cpu_lo(int i)
+{
+	return ((unsigned)i ^ 0x5a5au) + 3u;
+}
+
+static unsigned cpu_hi(int i, unsigned seed)
+{
+	unsigned acc[64];
+	unsigned s = 0;
+
+	for (int k = 0; k < 64; k++)
+		acc[k] = seed + (unsigned)k * ((unsigned)i + 1u);
+	for (int t = 0; t < 128; t++)
+		for (int k = 0; k < 64; k++)
+			acc[k] = acc[k] * 2654435761u + ((unsigned)t ^ (unsigned)k);
+	for (int k = 0; k < 64; k++)
+		s += acc[k];
+	return s;
+}
+
+static unsigned cpu_bounds(int i)
+{
+	return (unsigned)(i * 5) + 7u;
+}
+
+template <typename T>
+static void info_attrs(const char *name, T fn, int block)
+{
+	cudaFuncAttributes a;
+	int occ = -1;
+
+	if (cudaFuncGetAttributes(&a, fn) != cudaSuccess)
+		return;
+	(void)cudaOccupancyMaxActiveBlocksPerMultiprocessor(&occ, fn, block, 0);
+	printf("INFO regs %s numregs=%d maxthreads=%d occ_blocks_per_sm=%d\n",
+	       name, a.numRegs, a.maxThreadsPerBlock, occ);
+}
+
+int main(int argc, char **argv)
+{
+	static const unsigned seed = 0x12345678u;
+	int bad = 0;
+
+	if (p2_emit_mode(argc, argv)) {
+		printf("RESULT regs kernel=lo n=%d bad=0\n", RG_N);
+		printf("RESULT regs kernel=hi n=%d bad=0\n", RG_N);
+		printf("RESULT regs kernel=bounds n=%d bad=0\n", RG_N);
+		printf("RESULT regs kernels=3 bad=0\n");
+		return p2_status(0);
+	}
+
+	p2_device_info();
+
+	unsigned *d_out = NULL;
+	unsigned *got = (unsigned *)malloc(sizeof(unsigned) * RG_N);
+
+	if (!got)
+		return 1;
+	P2_CHECK(cudaMalloc(&d_out, sizeof(unsigned) * RG_N));
+
+	regs_lo<<<(RG_N + 255) / 256, 256>>>(d_out, RG_N);
+	P2_CHECK(cudaGetLastError());
+	P2_CHECK(cudaDeviceSynchronize());
+	P2_CHECK(cudaMemcpy(got, d_out, sizeof(unsigned) * RG_N,
+			    cudaMemcpyDeviceToHost));
+	{
+		int sbad = 0;
+
+		for (int i = 0; i < RG_N; i++)
+			if (got[i] != cpu_lo(i))
+				sbad++;
+		printf("RESULT regs kernel=lo n=%d bad=%d\n", RG_N, sbad);
+		bad += sbad;
+	}
+
+	regs_hi<<<(RG_N + 255) / 256, 256>>>(d_out, RG_N, seed);
+	P2_CHECK(cudaGetLastError());
+	P2_CHECK(cudaDeviceSynchronize());
+	P2_CHECK(cudaMemcpy(got, d_out, sizeof(unsigned) * RG_N,
+			    cudaMemcpyDeviceToHost));
+	{
+		int sbad = 0;
+
+		for (int i = 0; i < RG_N; i++)
+			if (got[i] != cpu_hi(i, seed))
+				sbad++;
+		printf("RESULT regs kernel=hi n=%d bad=%d\n", RG_N, sbad);
+		bad += sbad;
+	}
+
+	regs_bounds<<<(RG_N + 127) / 128, 128>>>(d_out, RG_N);
+	P2_CHECK(cudaGetLastError());
+	P2_CHECK(cudaDeviceSynchronize());
+	P2_CHECK(cudaMemcpy(got, d_out, sizeof(unsigned) * RG_N,
+			    cudaMemcpyDeviceToHost));
+	{
+		int sbad = 0;
+
+		for (int i = 0; i < RG_N; i++)
+			if (got[i] != cpu_bounds(i))
+				sbad++;
+		printf("RESULT regs kernel=bounds n=%d bad=%d\n", RG_N, sbad);
+		bad += sbad;
+	}
+
+	info_attrs("lo", regs_lo, 256);
+	info_attrs("hi", regs_hi, 256);
+	info_attrs("bounds", regs_bounds, 128);
+
+	printf("RESULT regs kernels=3 bad=%d\n", bad);
+
+	cudaFree(d_out);
+	free(got);
+	return p2_status(bad);
+}
diff --git a/tools/cuda-ab/p2/src/p2_streams.cu b/tools/cuda-ab/p2/src/p2_streams.cu
index 09d6ed3..12ee56c 100644
--- a/tools/cuda-ab/p2/src/p2_streams.cu
+++ b/tools/cuda-ab/p2/src/p2_streams.cu
@@ -1,16 +1,23 @@
 // 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.
+ * p2_streams - created streams, two kernels per stream, independent 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).
+ *   phase 1  three streams, each: scale (x2) then addv (+1), dependent pair
+ *   phase 2  four streams, each: two *independent* fill kernels on separate
+ *            buffers - up to eight kernels in flight simultaneously; all
+ *            results verified and every stream event required complete
+ *
+ * Expected value phase 1 per element: ((i % 100) * 2.0f) + 1.0f (exact in
+ * fp32). Phase 2 values are small integers (exact in int32).
+ * Device memory: phase 1 3 x 4 KB, phase 2 8 x 4 KB (peak 32 KB, 8 pages).
  */
 #include "common.h"
 
 #define SN 1024
 #define NS 3
+#define IN_NS 4
+#define IN_N 1024
 
 __global__ void scale(float *x, float s, int n)
 {
@@ -28,6 +35,14 @@ __global__ void addv(float *x, float v, int n)
 		x[i] = x[i] + v;
 }
 
+__global__ void fillk(int *x, int base, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n)
+		x[i] = base * 1000 + i;
+}
+
 static void fill(float *h, int n)
 {
 	for (int i = 0; i < n; i++)
@@ -39,6 +54,61 @@ static float expected1(int i)
 	return ((float)(i % 100) * 2.0f) + 1.0f;
 }
 
+/* Phase 2: independent kernels submitted across four streams. No kernel
+ * depends on any other, so the driver is free to overlap all eight; every
+ * buffer is verified element-exactly and every stream event must be complete
+ * after the stream is synchronized. */
+static int indep_phase(void)
+{
+	int *got = (int *)malloc(sizeof(int) * IN_N);
+	cudaStream_t st[IN_NS];
+	cudaEvent_t ev[IN_NS];
+	int *d_a[IN_NS], *d_b[IN_NS];
+	int vbad = 0, ebad = 0;
+
+	if (!got)
+		return 1;
+
+	for (int k = 0; k < IN_NS; k++) {
+		P2_CHECK(cudaStreamCreate(&st[k]));
+		P2_CHECK(cudaEventCreate(&ev[k]));
+		P2_CHECK(cudaMalloc(&d_a[k], sizeof(int) * IN_N));
+		P2_CHECK(cudaMalloc(&d_b[k], sizeof(int) * IN_N));
+		fillk<<<(IN_N + 255) / 256, 256, 0, st[k]>>>(d_a[k], k + 2,
+							     IN_N);
+		fillk<<<(IN_N + 255) / 256, 256, 0, st[k]>>>(d_b[k], k + 9,
+							     IN_N);
+		P2_CHECK(cudaGetLastError());
+		P2_CHECK(cudaEventRecord(ev[k], st[k]));
+	}
+
+	for (int k = 0; k < IN_NS; k++) {
+		P2_CHECK(cudaStreamSynchronize(st[k]));
+		P2_CHECK(cudaMemcpy(got, d_a[k], sizeof(int) * IN_N,
+				    cudaMemcpyDeviceToHost));
+		for (int i = 0; i < IN_N; i++)
+			if (got[i] != (k + 2) * 1000 + i)
+				vbad++;
+		P2_CHECK(cudaMemcpy(got, d_b[k], sizeof(int) * IN_N,
+				    cudaMemcpyDeviceToHost));
+		for (int i = 0; i < IN_N; i++)
+			if (got[i] != (k + 9) * 1000 + i)
+				vbad++;
+		if (cudaEventQuery(ev[k]) != cudaSuccess)
+			ebad++;
+		cudaEventDestroy(ev[k]);
+		cudaStreamDestroy(st[k]);
+		cudaFree(d_a[k]);
+		cudaFree(d_b[k]);
+	}
+
+	printf("RESULT streams indep nstreams=%d launches=%d n=%d bad=%d\n",
+	       IN_NS, IN_NS * 2, IN_N, vbad);
+	printf("RESULT streams indep_events=%d bad=%d\n", IN_NS, ebad);
+	free(got);
+	return vbad + ebad;
+}
+
 int main(int argc, char **argv)
 {
 	float *h = (float *)malloc(sizeof(float) * SN);
@@ -53,6 +123,9 @@ int main(int argc, char **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);
+		printf("RESULT streams indep nstreams=%d launches=%d n=%d bad=0\n",
+		       IN_NS, IN_NS * 2, IN_N);
+		printf("RESULT streams indep_events=%d bad=0\n", IN_NS);
 		free(h); free(got);
 		return p2_status(0);
 	}
@@ -104,6 +177,9 @@ int main(int argc, char **argv)
 		cudaStreamDestroy(st[k]);
 		cudaFree(d_x[k]);
 	}
+
+	bad += indep_phase();
+
 	free(h); free(got);
 	return p2_status(bad);
 }