1 files changed, 32 insertions, 4 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);