differentiate reduction function with CUDA atomicAdd
- Dominant language
- LLVM
- Stars
- 1.7k
- Forks
- 188
- Avg merge
- 1d 22h
- Merged PRs (30d)
- 26
Description
Hi, I met the issue (shown below) when differentiating a reduction function with atomicAdd. It seems CUDA `atomicAdd` is not supported in enzyme AD. I am wondering if anyone can please help give any suggestions? Thanks!
```
in Mode: ReverseModeCombined
cannot handle unknown instruction
%22 = atomicrmw fadd float* %0, float %21 seq_cst, align 4, !dbg !45fatal error: error in backend: unknown value
```
The device function to differentiate:
```cuda
__device__ void compute_loss(float * loss,
float * out_feat,
const uint32_t N){
const uint32_t i = threadIdx.x + blockIdx.x * blockDim.x;
float local_loss = out_feat[i * 2] * out_feat[i * 2] + out_feat[i * 2 + 1] * out_feat[i * 2 + 1];
atomicAdd(loss, local_loss / N);
}
```
Contributor guide
Assessment
This issue has not been assessed yet.