cuda-ab: victory-series coverage battery (8 new clients) + soak runner list

victory_saxpy, victory_reduce (shared+atomics), victory_matmul (tiled
shared), victory_strided, victory_memset2d (pitched), victory_atomics,
victory_streams_events (3 streams + events), victory_math (mixed
int/float/double). Each verifies via readback and prints STATUS OK|FAIL.
soak.sh now runs the full 14-client set (seq or sustained).

1 files changed, 129 insertions, 0 deletions

diff --git a/tools/cuda-ab/victory_atomics.cu b/tools/cuda-ab/victory_atomics.cu
new file mode 100644
index 0000000..7e5eedf
--- /dev/null
+++ b/tools/cuda-ab/victory_atomics.cu
@@ -0,0 +1,129 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * victory_atomics.cu - global atomicAdd/atomicMax histogram over 4 MB.
+ *
+ * A deterministic 4M-byte buffer (every value 0..255) is histogrammed on the
+ * device into a 256-bin counter array with atomicAdd, while a global maximum
+ * is tracked with atomicMax. The host recomputes the exact histogram and
+ * maximum and verifies every bin, the max, and the total count.
+ *
+ * Output lines are stable and greppable:
+ *
+ *   DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
+ *   ATOMICS n=<count> bins=256 max=<value> bad=<mismatches> total=<sum>
+ *   STATUS OK|FAIL
+ *
+ * Build with a static CUDA runtime (--cudart=static) so the guest only needs
+ * the driver's libcuda.so, not a CUDA toolkit installation:
+ *
+ *   nvcc -O2 --cudart=static -o victory_atomics victory_atomics.cu
+ *
+ * Then, from the guest:  /tmp/victory_atomics
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstdint>
+#include <cstring>
+#include <cuda_runtime.h>
+
+#define CHECK(x)                                                         \
+	do {                                                             \
+		cudaError_t e_ = (x);                                    \
+		if (e_ != cudaSuccess) {                                 \
+			printf("CUDA_ERROR %s at line %d\n",             \
+			       cudaGetErrorString(e_), __LINE__);        \
+			printf("STATUS FAIL\n");                         \
+			return 1;                                        \
+		}                                                        \
+	} while (0)
+
+#define BINS 256
+
+__global__ void histogram(const unsigned char *data, unsigned int *hist,
+			  unsigned int *gmax, int n)
+{
+	int i = blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n) {
+		unsigned int v = data[i];
+
+		atomicAdd(&hist[v], 1u);
+		atomicMax(gmax, v);
+	}
+}
+
+int main(void)
+{
+	cudaDeviceProp prop;
+	const int n = 1 << 22; /* 4M bytes */
+	unsigned char *h_data;
+	unsigned int *h_hist;
+	unsigned int h_max = 0, h_total = 0;
+	unsigned char *d_data;
+	unsigned int *d_hist, *d_max;
+	unsigned int ref[BINS];
+	unsigned int ref_max = 0;
+	int bad = 0, i;
+
+	setvbuf(stdout, NULL, _IONBF, 0);
+
+	CHECK(cudaGetDeviceProperties(&prop, 0));
+	printf("DEVICE %s cc=%d.%d sm=%d mem=%llu\n", prop.name,
+	       prop.major, prop.minor, prop.multiProcessorCount,
+	       (unsigned long long)prop.totalGlobalMem);
+
+	h_data = (unsigned char *)malloc(n);
+	h_hist = (unsigned int *)malloc(BINS * sizeof(unsigned int));
+	if (!h_data || !h_hist)
+		return 1;
+
+	for (i = 0; i < n; i++)
+		h_data[i] = (unsigned char)(((long long)i * 37 + i / 256) & 0xFF);
+
+	memset(ref, 0, sizeof(ref));
+	for (i = 0; i < n; i++) {
+		ref[h_data[i]]++;
+		if (h_data[i] > ref_max)
+			ref_max = h_data[i];
+	}
+
+	CHECK(cudaMalloc(&d_data, n));
+	CHECK(cudaMalloc(&d_hist, BINS * sizeof(unsigned int)));
+	CHECK(cudaMalloc(&d_max, sizeof(unsigned int)));
+	CHECK(cudaMemcpy(d_data, h_data, n, cudaMemcpyHostToDevice));
+	CHECK(cudaMemset(d_hist, 0, BINS * sizeof(unsigned int)));
+	CHECK(cudaMemset(d_max, 0, sizeof(unsigned int)));
+
+	histogram<<<(n + 255) / 256, 256>>>(d_data, d_hist, d_max, n);
+	CHECK(cudaGetLastError());
+	CHECK(cudaDeviceSynchronize());
+	CHECK(cudaMemcpy(h_hist, d_hist, BINS * sizeof(unsigned int),
+			 cudaMemcpyDeviceToHost));
+	CHECK(cudaMemcpy(&h_max, d_max, sizeof(unsigned int),
+			 cudaMemcpyDeviceToHost));
+
+	for (i = 0; i < BINS; i++) {
+		if (h_hist[i] != ref[i])
+			bad++;
+		h_total += h_hist[i];
+	}
+	if (h_max != ref_max)
+		bad++;
+
+	printf("ATOMICS n=%d bins=256 max=%u bad=%d total=%u\n", n, h_max, bad,
+	       h_total);
+
+	if (bad != 0 || h_total != (unsigned int)n) {
+		printf("STATUS FAIL\n");
+		return 1;
+	}
+
+	printf("STATUS OK\n");
+	CHECK(cudaFree(d_data));
+	CHECK(cudaFree(d_hist));
+	CHECK(cudaFree(d_max));
+	free(h_data);
+	free(h_hist);
+	return 0;
+}