[flang-commits] [flang] [flang][cuda] Pass assumed-shape arrays of bind(c) attributes(global) kernels by base address (PR #199313)

[Valentin Clement バレンタイン クレメン via flang-commits]

clementval wrote:

> > Is that documented anywhere is the spec? Can you add an example of what is the actual C kernel?
> 
> The CUDA Fortran Programming Guide does not document this case. The reference compiler passes the data pointer when lowering `dimension(:), device` under `bind(c) attributes(global)`, and the fix matches that.
> 
> Example:
> 
> ```
> extern "C"
> __global__ void mykernel(float *d, int n) {
>     int tid = blockDim.x * blockIdx.x + threadIdx.x;
>     if (tid < n) d[tid] = (float)tid;
> }
> ```

That's probably something worth documenting since it is not in the spec and looks like it diverge from standard bind c lowering. Maybe we can start a CUDA.md in flang/docs to keep track of such special case. 

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