llvm / llvm/Polygeist

Incorrect execution result of running CUDA to OpenMP compilation?

Open
#397 1 comment 0 reactions 0 assignees View on GitHub

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

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.