To what memory region do kernel parameters go?

[DoZo1971]

Maybe a silly question but…

Do kernel parameters always reside in shared memory?

Say for a kernel like this:

[codebox]global void CudaCSGSphereKernel

(const float3 *p_CudaPointArray,

float *p_CudaResultArray,

const int p_Size,

const int p_NumberOfNodes,

const int p_Label,

const float p_Radius,

const float3 p_Center)

{

...

}[/codebox]

… the .ptx looks like this…

The [codebox].reg .u16 %rh<4>;

.reg .u32 %r<16>;

.reg .f32 %f<17>;

.reg .pred %p<3>;

.param .u32 __cudaparm__Z19CudaCSGSphereKernelPK6float3PfiiifS__p_CudaPo

intArray;

.param .u32 __cudaparm__Z19CudaCSGSphereKernelPK6float3PfiiifS__p_CudaRe

sultArray;

.param .s32 __cudaparm__Z19CudaCSGSphereKernelPK6float3PfiiifS__p_Size;

.param .s32 __cudaparm__Z19CudaCSGSphereKernelPK6float3PfiiifS__p_Number

OfNodes;

.param .s32 __cudaparm__Z19CudaCSGSphereKernelPK6float3PfiiifS__p_Label;

.param .f32 __cudaparm__Z19CudaCSGSphereKernelPK6float3PfiiifS__p_Radius

;

.param .align 4 .b8 __cudaparm__Z19CudaCSGSphereKernelPK6float3PfiiifS__p_Center

[12];

.loc 14 116 0[/codebox]

I would think these parameters are copied to shared memory and fast-fetched for every individual thread (since they are constant). But is this true?

Kind regards,

Daniel Dekkers

[jph4599]

I believe they are passed through shared memory…from my interpretation, they should be at least…

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.

I recall reading something about the possibility to use constant memory instead but I’m not sure about that.