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

1 files changed, 81 insertions, 5 deletions

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);
 }