device constant memory variables declared "__device__ const" appear as .global

Accelerated Computing CUDA CUDA Programming and Performance
1

In the following code,

__device__ const int arr[] = { 1, 3, 4, 6 };

__global__ void hello_world(int *out, int in) {

	*out = arr[in];

}

the $arr$ variable is translated into a “.global” statement in ptx, which the ptx guide confirms to be global memory. On p. 34 (version 1.4) the guide says .global and .const memory is initializable. Indeed, I can change the statement to “constant const int arr…” and it works (translates into .const in ptx). However, are both of these going to be slow if I have a lot of divergent accesses (in which case I should copy it to shared)?

Thanks,

Nicholas

2

device const does not equal constant. And yes, constant memory is slow for any case where a half-warp does not access the same element.