Yeah, all three of them are used. DYND_CUDA_HOST_DEVICE_CALLABLE is unavoidable. I can remove the others and sprinkle __device__ and __global__ throughout the code, but it keeping all three of them seems consistent.

Comments are fine, I'll add some. I don't care about the names, I'd be fine if you want to choose them. Maybe we just drop the 'CALLABLE' part?