From cafc460bab4cab32968a6656242f34f2dc392486 Mon Sep 17 00:00:00 2001 From: Arthur Hanson Rasmusson Date: Tue, 20 Oct 2026 08:48:46 -0400 Subject: cuda-ab: launch_count - single-launch execution-count repro (each forwarded launch runs its engine work twice; the multi-pass class root cause) --- tools/cuda-ab/launch_count.cu | 25 +++++++++++++++++++++++++ 1 file changed, 25 insertions(+) create mode 100644 tools/cuda-ab/launch_count.cu (limited to 'tools') diff --git a/tools/cuda-ab/launch_count.cu b/tools/cuda-ab/launch_count.cu new file mode 100644 index 0000000..02db1d6 --- /dev/null +++ b/tools/cuda-ab/launch_count.cu @@ -0,0 +1,25 @@ +// SPDX-License-Identifier: BSD-3-Clause +// launch_count.cu - single-launch execution-count repro (r342). +// +// One launch, one atomic increment. The engine-written page witness showed the +// increment applied TWICE (1 -> 2; the 100-increment variant -> 200): every +// forwarded launch executes its engine work twice (one submission, one ring +// publish, completion OK). This is the root cause of the multi-pass content +// class (victory_reduce exactly 2x, scale/saxpy/smiley/multikernel wrong). +// Run in the guest; expected LAUNCH_COUNT=1. +// for a single launch, the launch path executes work more than once. +#include +#include +__global__ void inc(unsigned long long *c) { if (threadIdx.x == 0 && blockIdx.x == 0) atomicAdd(c, 1ULL); } +int main(void) { + unsigned long long *d; unsigned long long h = 0; + if (cudaMalloc(&d, 8) != cudaSuccess) { printf("CUDA_ERROR malloc\nSTATUS FAIL\n"); return 1; } + if (cudaMemset(d, 0, 8) != cudaSuccess) { printf("CUDA_ERROR memset\nSTATUS FAIL\n"); return 1; } + inc<<<1, 32>>>(d); + if (cudaGetLastError() != cudaSuccess) { printf("CUDA_ERROR launch\nSTATUS FAIL\n"); return 1; } + if (cudaDeviceSynchronize() != cudaSuccess) { printf("CUDA_ERROR sync\nSTATUS FAIL\n"); return 1; } + if (cudaMemcpy(&h, d, 8, cudaMemcpyDeviceToHost) != cudaSuccess) { printf("CUDA_ERROR d2h\nSTATUS FAIL\n"); return 1; } + printf("LAUNCH_COUNT=%llu (expect 1)\n", h); + printf("STATUS %s\n", h == 1 ? "OK" : "FAIL"); + return h == 1 ? 0 : 1; +} -- cgit v1.2.3