[clang] [cuda][[HIP] `__constant__` should imply constant (PR #110182)

Alex Voicu via cfe-commits cfe-commits at lists.llvm.org
Thu Sep 26 16:38:27 PDT 2024


AlexVlx wrote:

> > `__constant__` may not necessarily be `const` for IR purposes. I.e. IR may not rely on the 'known' values, as seen in IR, as the data may actually be populated by the host via CUDA API calls `cudaMemcpyToSymbol` before the GPU kernel launch.
> 
> But since this is marked `externally_initialised` on the device side, one would assume that those semantics still hold i.e. there's no known value to assume?

Oh, I guess the concern might be around doing a `cuda/hipMemcpyToSymbol` during a kernel's execution (i.e. `__global__ foo()` is executing, host sets a __constant__, notifies foo, and expects foo to observe the updated constant)? I admit that I'm not certain if that is allowed, if it is then yes, this is problematic / not viable.

https://github.com/llvm/llvm-project/pull/110182


More information about the cfe-commits mailing list