1 files changed, 138 insertions, 0 deletions
diff --git a/tools/cuda-ab/victory_strided.cu b/tools/cuda-ab/victory_strided.cu
new file mode 100644
index 0000000..2fb8739
--- /dev/null
+++ b/tools/cuda-ab/victory_strided.cu
@@ -0,0 +1,138 @@
+// SPDX-License-Identifier: BSD-3-Clause
+/*
+ * victory_strided.cu - strided/offset memory access over 4 MB of floats.
+ *
+ * Two kernels permute an n-element array (n = 1M floats, 4 MB) with a
+ * stride co-prime to n (idx = i*7 mod n), so every element is written exactly
+ * once and the host can verify every single position:
+ *
+ *   out1[(i*7) % n] = in[i] + 1.0f                 (plain permutation)
+ *   out2[(i*7) % n] = in[off + i] * 2.0f           (offset base pointer)
+ *
+ * All values are multiples of 0.25, so both operations are exact in float
+ * and the host comparison is exact.
+ *
+ * Output lines are stable and greppable:
+ *
+ *   DEVICE <name> cc=<major>.<minor> sm=<count> mem=<bytes>
+ *   STRIDED n=<count> stride=7 off=<offset> bad=<mismatches> csum=<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_strided victory_strided.cu
+ *
+ * Then, from the guest:  /tmp/victory_strided
+ */
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstdint>
+#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 STRIDE 7
+#define OFFSET 4097
+
+__global__ void permute_add(const float *in, float *out, int n, int stride)
+{
+	long long i = (long long)blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n) {
+		long long idx = (i * (long long)stride) % (long long)n;
+
+		out[idx] = in[i] + 1.0f;
+	}
+}
+
+__global__ void permute_offset(const float *in, float *out, int n, int off,
+			       int stride)
+{
+	long long i = (long long)blockIdx.x * blockDim.x + threadIdx.x;
+
+	if (i < n) {
+		long long idx = (i * (long long)stride) % (long long)n;
+
+		out[idx] = in[(long long)off + i] * 2.0f;
+	}
+}
+
+int main(void)
+{
+	cudaDeviceProp prop;
+	const int n = 1 << 20; /* 1M floats = 4 MB */
+	const long long total = (long long)n + OFFSET;
+	size_t bytes = (size_t)n * sizeof(float);
+	size_t total_bytes = (size_t)total * sizeof(float);
+	float *h_in, *h_out1, *h_out2;
+	float *d_in, *d_out1, *d_out2;
+	double csum = 0.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_in = (float *)malloc(total_bytes);
+	h_out1 = (float *)malloc(bytes);
+	h_out2 = (float *)malloc(bytes);
+	if (!h_in || !h_out1 || !h_out2)
+		return 1;
+
+	for (i = 0; i < (int)total; i++)
+		h_in[i] = (float)(i % 1000) * 0.25f;
+
+	CHECK(cudaMalloc(&d_in, total_bytes));
+	CHECK(cudaMalloc(&d_out1, bytes));
+	CHECK(cudaMalloc(&d_out2, bytes));
+	CHECK(cudaMemcpy(d_in, h_in, total_bytes, cudaMemcpyHostToDevice));
+
+	permute_add<<<(n + 255) / 256, 256>>>(d_in, d_out1, n, STRIDE);
+	permute_offset<<<(n + 255) / 256, 256>>>(d_in, d_out2, n, OFFSET,
+						 STRIDE);
+	CHECK(cudaGetLastError());
+	CHECK(cudaDeviceSynchronize());
+	CHECK(cudaMemcpy(h_out1, d_out1, bytes, cudaMemcpyDeviceToHost));
+	CHECK(cudaMemcpy(h_out2, d_out2, bytes, cudaMemcpyDeviceToHost));
+
+	for (i = 0; i < n; i++) {
+		long long idx = ((long long)i * STRIDE) % (long long)n;
+
+		if (h_out1[idx] != h_in[i] + 1.0f)
+			bad++;
+		if (h_out2[idx] != h_in[OFFSET + i] * 2.0f)
+			bad++;
+		csum += (double)h_out1[idx];
+	}
+
+	printf("STRIDED n=%d stride=%d off=%d bad=%d csum=%.6f\n", n, STRIDE,
+	       OFFSET, bad, csum);
+
+	if (bad != 0) {
+		printf("STATUS FAIL\n");
+		return 1;
+	}
+
+	printf("STATUS OK\n");
+	CHECK(cudaFree(d_in));
+	CHECK(cudaFree(d_out1));
+	CHECK(cudaFree(d_out2));
+	free(h_in);
+	free(h_out1);
+	free(h_out2);
+	return 0;
+}