[CudaTarget] Why do we need to cast the SPMD grid keywords?

[kaushikcfd]

For example an arange(0, 10) kernel with the CudaTarget looks like --

// Grid size: (10, 1)
extern "C" __global__ void __launch_bounds__(1) arange(__global__ int *__restrict__ y)
{
  y[((int32_t) blockIdx.x)] = ((int32_t) blockIdx.x);
}

Notice the (int32_t) casts. I find those to be bringing down the SNR pretty significantly. Any reason we need those?

[inducer]

Loopy codegen expects integer math to be signed. (It'll generate a - b as a + (-1)*b for instance.) dim3 (the type of these {block,grid}{Dim,Idx} variables) is based on unsigned integers:

https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#dim3

[inducer]

We could package the casts up in macros, like we do in CL though, if you like the SNR of that better.

[kaushikcfd]

Thanks!