Why do we need a new bitcode file? Could we just not implement the unsafe atomic add in the existing one with some JIT reflection?

Unfortunately, clang will choke on these builtins if there is no appropriate-mcpu specified.

Compiling it without specifying gfx90a target arch results in the following error:

/home/soproni/Projects/AdaptiveCpp/src/libkernel/sscp/amdgpu/atomic_gfx90a.cpp:18:10: error: '__builtin_amdgcn_global_atomic_fadd_f64' needs target feature gfx90a-insts   18 |   return __builtin_amdgcn_global_atomic_fadd_f64(ptr, x);      |

Which is as far as I can tell is due to checking the sub-target compatibility of the builtin by the frontend here:
https://github.com/intel/llvm/blob/sycl/clang/lib/CodeGen/CodeGenFunction.cpp#L3190

There are some exceptions made for compiling when targeting--hipstdpar however passing this when compiling the bitcode library results in an error about amdgcn not being a valid target for host compilation. Even after adding-Xclang -fcuda-is-device to the the compilation arguments the same issue persists. I am not sure what--hipstdpar changes that results in this new behavior. Do you think it make sense to dig further into the LLVM source to figure out if we could trick the fronted to emitting IR for these builtins?

Additionally the--hipstdpar flag has been only merged into llvm 18.1