where does the parameters of kernels store?

[cmmfir]

For example, global void mykernel(float a, int *p) {…}; where does a and p store in? registers or shared memory? Thanks!

[nitin.life]

in registers ( as far as my kernels are concerned)…

[jph4599]

I believe it is shared memory:

http://forums.nvidia.com/index.php?showtopic=68779

CUDA stores kernel parameters in shared memory. The limit is 256 bytes total.

http://forums.nvidia.com/index.php?showtopic=94009

From the 2.1 Programming guide section 4.2.3 Execution Configuration

As an example, a function declared as

global void Func(float* parameter);

must be called like this:

Func<<< Dg, Db, Ns >>>(parameter);

The arguments to the execution configuration are evaluated before the actual

function arguments and like the function arguments, are currently passed via shared

memory to the device.

4.2.1.4 Restrictions

global function parameters are currently passed via shared memory to the

device and limited to 256 bytes.

http://forums.nvidia.com/index.php?showtopic=40497

The actual kernel parameters are passed into shared memory, which is 16k (i.e. less than 32x32x32xsizeof(z)). Things may work if you instead pass a pointer (i.e. 4bytes) to your kernel.

[nitin.life]

Your right JPH4599, Mybad.

its strange though, I always counted my registers including these arguments and the count turned out to be correct. I guess I have to recheck my kernel… :">

[tmurray]

arguments are never stored in registers at kernel launch time.

[Gregory_Diamos]

If you look at the code generated for some applications, there are several ld.param PTX instructions which define ‘param’ as an address space. This could be mapped to anything at runtime (it looks like the current implementation maps it to shared memory), but this could be easily changed in the future.

[cmmfir]

I believe it is shared memory:

http://forums.nvidia.com/index.php?showtopic=68779

http://forums.nvidia.com/index.php?showtopic=94009

From the 2.1 Programming guide section 4.2.3 Execution Configuration

4.2.1.4 Restrictions

http://forums.nvidia.com/index.php?showtopic=40497

Thanks all. If they are in shared memory, I’m just curious if this kind of kernel works.

global void mykernel(float *A, int lda, float *B, int ldb)

{

A += lda;

B += ldb;

//read A[0] B[0]

...    

}

[wakensky]

in shared memory. Programer guide had mentioned that, and you can see the usage of shared memory with compile option --ptx or output details to cubin file