Skip to content

Commit

Permalink
Add a test for InvalidBitWidth
Browse files Browse the repository at this point in the history
  • Loading branch information
pvelesko committed Jan 20, 2025
1 parent 656d183 commit 620cca2
Show file tree
Hide file tree
Showing 2 changed files with 60 additions and 0 deletions.
1 change: 1 addition & 0 deletions tests/compiler/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -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.
Expand Down
59 changes: 59 additions & 0 deletions tests/compiler/TestHipccInvalidBitWidth.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,59 @@
#include <hip/hip_runtime.h>
#include <stdio.h>

__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<<<grid, block>>>(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;
}

0 comments on commit 620cca2

Please sign in to comment.