diff --git a/tests/compiler/CMakeLists.txt b/tests/compiler/CMakeLists.txt index 86b93f60f..c52625c76 100644 --- a/tests/compiler/CMakeLists.txt +++ b/tests/compiler/CMakeLists.txt @@ -119,6 +119,7 @@ add_hipcc_test(TestLdg.hip HIPCC_OPTIONS -fsyntax-only) add_hipcc_test(TestSwitchCase.hip HIPCC_OPTIONS -O1 -c) add_hipcc_test(TestHostSideHIPVectors.hip HIPCC_OPTIONS -fsyntax-only) add_hipcc_test(TestAlignAttr.hip HIPCC_OPTIONS -fsyntax-only) +add_hipcc_test(TestHipccInvalidBitWidth.cpp HIPCC_OPTIONS -O2) # Check __FAST_MATH__ is set for -ffast-math and preprocessor guards # using it are not hiding errors. diff --git a/tests/compiler/TestHipccInvalidBitWidth.cpp b/tests/compiler/TestHipccInvalidBitWidth.cpp new file mode 100644 index 000000000..194f40cb6 --- /dev/null +++ b/tests/compiler/TestHipccInvalidBitWidth.cpp @@ -0,0 +1,59 @@ +#include +#include + +__global__ void testWarpCalc(int* debug) { + int tid = threadIdx.x; + int bid = blockIdx.x; + int globalIdx = bid * blockDim.x + tid; + + // Do some computation to prevent optimization + int result = 0; + for(int i = 0; i < tid + 1; i++) { + result += i * globalIdx; + } + + // Store using atomic operation + atomicExch(&debug[globalIdx], result); +} + +int main() { + const int gridSize = 4; + const int blockSize = 64; + const int numThreads = gridSize * blockSize; + + // Allocate pinned memory + int* h_debug; + hipHostMalloc(&h_debug, numThreads * sizeof(int)); + memset(h_debug, 0, numThreads * sizeof(int)); + + // Allocate device memory + int* d_debug; + hipMalloc(&d_debug, numThreads * sizeof(int)); + hipMemset(d_debug, 0, numThreads * sizeof(int)); + + dim3 grid(gridSize); + dim3 block(blockSize); + + printf("Launching kernel with grid=%d, block=%d\n\n", gridSize, blockSize); + + // Use triple angle bracket syntax + testWarpCalc<<>>(d_debug); + hipDeviceSynchronize(); + + // Copy results back + hipMemcpy(h_debug, d_debug, numThreads * sizeof(int), hipMemcpyDeviceToHost); + + printf("Results for first few threads:\n"); + printf("GlobalIdx\tValue\n"); + + // Print first few entries + for (int i = 0; i < 8; i++) { + printf("%d\t\t%d\n", i, h_debug[i]); + } + + // Cleanup + hipHostFree(h_debug); + hipFree(d_debug); + + return 0; +} \ No newline at end of file