Using __half atomicAdd in the anyhit module causes an OPTIX_ERROR_PIPELINE_LINK_ERROR

Optix version: 8.0.0
GPU: RTX A6000
System: Ubuntu 22.04
CUDA and NVCC Version: 12.1
Driver Version: 550.76
Geometry Primitive: Built-in sphere
Hi folks,

I found that Using __half atomicAdd for float16 in the anyhit module causes an OPTIX_ERROR_PIPELINE_LINK_ERROR.
Specifically, I have params.accum_buffer in global memory.
I want to do atomicAdd( &params.accum_buffer[idx], t); in my anyhit module. My scene is just a single GAS of built-in sphere.

when params.accum_buffer is float*, the program works quite well.
Then, I want to use float16 to reduce the memory movement cost. So I changed params.accum_buffer to __half, and change my atomicAdd to atomicAdd( &params.accum_bufferi[dx], __float2half(t));


However, the thread-unsafe params.accum_buffer[idx] += __float2half(t); does not have this error. Since my GPU A6000’s compute capability is 8.6, it should support __half atomicAdd. Moreover, I wrote a single test cu file without using OptiX, and it provided the correct answer.

Therefore, I was wondering whether __half atomicAdd is not supported in OptiX’s any-hit module?

Thank you!

Could you rerun that with enabled OptiX validation mode and a logging callback to see if there is any additional information provided by OptiX why the pipeline failed to link?
If the atomicAdd with half is not a supported instruction there should be some error about that in the logging output.
Example code here:

For clarification, you only used that inside an anyhit program so far, where it doesn’t work.
You’re not saying this only fails inside anyhit programs but works in other program domains?

Would you be able to provide a reproducer project for that, like by changing any of the OptiX SDK examples to fail the same way?

Thank you for the reply! I will try to create a reproducer project based on the OptiX sphere example and enable validation mode. I’ll keep you posted.