Passing a pointer to __constant__ memory as a parameter for a __device__ function

[jsierra.siete]

Hello ! Im currently doing as the title says in a kernel. But I realised there is not much documentation whether this is something you should or shouldnt be allowed to do.

Im wondering if what Im doing is correct or it’s just randomly working.

__device__ __constant__ float constantData[DATA_SIZE];

__device__ void myFunction(const float* __restrict__ weights, ...) {

    some operations...
    x = x * weights[10];
    some other operations...
}

__global__ void someKernel(...) {
    some operations...
    myFunction(&constantData[100]);
    ...
}

If this is correct, then I assume there is an address space reserved exclusively to constant memory, let’s say from 00000000 to 0000FFFF for example. Should this also work if I split these functions into separated files ? Thanks in advance !

[Robert_Crovella]

There shouldn’t be a need to do this. __constant__ memory is defined at global scope, and it is accessible directly in device code, without the need to use a function parameter/argument for it.

Having said that, what you are showing should work. Don’t try to do the same thing passing it as a kernel argument from host code.

If you decide that you want to do this in a multi-module/multi-file way, then it’s just a matter of using the necessary extern declarations as you would with any global scope variable defined in another module, and also be sure to use relocatable device code linking (e.g. -rdc=true, or similar) when compiling/linking the modules together.

Yes, there is an address space reserved exclusively to constant memory.

$ cat t1998.cu
#include <cstdio>

__constant__ int data = 1;

__device__ void f(const int *d){

  printf("%d\n", *d);

}

__global__ void k(){

  f(&data);
}

int main(){

  k<<<1,1>>>();
  cudaDeviceSynchronize();
}
$ nvcc -o t1998 t1998.cu
$ compute-sanitizer ./t1998
========= COMPUTE-SANITIZER
1
========= ERROR SUMMARY: 0 errors
$