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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to