Incorrect execution result of running CUDA to OpenMP compilation?
Nobody has claimed this yet.
- Dominant language
- C++
- Stars
- 624
- Forks
- 170
- PR merge metrics
- No merged PRs in 30d
Description
I have a simple CUDA program here that performs a simple reduction:
```
#include
#include
// CUDA kernel for performing reduction (sum) of an array
__global__ void reduceSum(int *g_input, int *g_output, int n) {
extern __shared__ int s_data[];
// Each thread loads one element from global to shared memory
unsigned int tid = threadIdx.x;
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) s_data[tid] = g_input[i];
else s_data[tid] = 0;
__syncthreads();
// Do reduction in shared memory
for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) {
if (tid < s) {
s_data[tid] += s_data[tid + s];
}
__syncthreads();
}
// Write result for this block to global memory
if (tid == 0) g_output[blockIdx.x] = s_data[0];
}
int main() {
int n = 1024;
int size = n * sizeof(int);
int *h_input, *h_output;
int *d_input, *d_output;
// Allocate host memory
h_input = (int*)malloc(size);
h_output = (int*)malloc(sizeof(int));
// Initialize input array
for(int i = 0; i < n; i++) {
h_input[i] = 1; // Example: fill with 1 for simplicity
}
// Allocate device memory
cudaMalloc((void **)&d_input, size);
cudaMalloc((void **)&d_output, sizeof(int));
// Copy from host to device
cudaMemcpy(d_input, h_input, size, cudaMemcpyHostToDevice);
// Launch the kernel
int threadsPerBlock = 256;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
reduceSum<<>>(d_input, d_output, n);
// Copy result back to host
cudaMemcpy(h_output, d_output, sizeof(int), cudaMemcpyDeviceToHost);
printf("Sum is %d\n", *h_output);
// Cleanup
free(h_input);
free(h_output);
cudaFree(d_input);
cudaFree(d_output);
return 0;
}
```
Then I compile it using this command (which I believe is the correct one to generate OpenMP with all the optimization):
```
cgeist --cuda-gpu-arch=sm_75 --cuda-lower --cpuify="distribute.mincut" -scal-rep=0 -raise-scf-to-affine --inner-serialize=1 --function=* -O2 -I/home/ericxu233/CUDAtoX/Polygeist/build/projects/openmp/src/runtime -L/home/ericxu233/CUDAtoX/Polygeist/build/projects/openmp/libomptarget/ -resource-dir=/home/ericxu233/CUDAtoX/Polygeist/build/lib/clang/18 simple.cu -o simple
```
When I try to run the executable ./simple it shows that:
```
Sum is 0
```
but in fact according to the source code the sum should be 256 and my CUDA GPU run confirms this. Moreover, when I remove the optimization options "-scal-rep=0 -raise-scf-to-affine --inner-serialize= -O2", the executable run results in a segmentation fault.
I was wondering what am I doing wrong here? Is the CUDA to OpenMP flow not properly supported anymore?
Contributor guide
No contributing guide indexed for this repository
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Research direction
Reproduce simple.cu with the exact cgeist command and runtime paths shown, then compare the optimized run with the command that omits the optimization options. Inspect the generated execution path and its handling of the reduction; done means the executable returns 256 without optimization-dependent crashes.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- compilers
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Stale
- Clarity
- Needs clarification
- Newbie friendliness
- 25/100