1 files changed, 100 insertions, 0 deletions
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);
+}