Kernel parameter passing in CUDA?

Question

I have a newbie doubt regarding how do CUDA kernels work.

If have the following code (which use the function cuPrintf taken from here):

#include "cuPrintf.cu"

__global__ void testKernel(int param){
    cuPrintf("Param value: %d\n", param);
}

int main(void){

    // initialize cuPrintf
    cudaPrintfInit();

    int a = 456;    

    testKernel<<<4,1>>>(a);

    // display the device's greeting
    cudaPrintfDisplay();

    // clean up after cuPrintf
    cudaPrintfEnd();
}

The output of the execution is:

Param value: 456
Param value: 456
Param value: 456
Param value: 456

I cannot get how the kernel can read the correct value of the parameter I pass, isn't it allocated in the host memory? Can the GPU read from the host memory?

Thanks,

Andrea

Answer 1

According to the section E.2.5.2. Function Parameters in CUDA C Programming Guide

__global__ function parameters are passed to the device:

• via shared memory and are limited to 256 bytes on devices of compute capability 1.x,
• via constant memory and are limited to 4 KB on devices of compute capability 2.x and higher.

Answer 2

The declaration void testKernel(int param) says that param is passed by value, not by reference. In other words, the stack contains a copy of a's value, not a pointer to a. CUDA copies the stack to the kernel running on the GPU.