EnzymeAD / EnzymeAD/Enzyme

differentiate reduction function with CUDA atomicAdd

Open
#1,288 9 comments 0 reactions 0 assignees View on GitHub
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.