1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
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 <cstdio>
#include <cuda_runtime.h>
__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;
}
|
