To what memory region do kernel parameters go?

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)




… 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


.param .u32 __cudaparm__Z19CudaCSGSphereKernelPK6float3PfiiifS__p_CudaRe


.param .s32 __cudaparm__Z19CudaCSGSphereKernelPK6float3PfiiifS__p_Size;

.param .s32 __cudaparm__Z19CudaCSGSphereKernelPK6float3PfiiifS__p_Number


.param .s32 __cudaparm__Z19CudaCSGSphereKernelPK6float3PfiiifS__p_Label;

.param .f32 __cudaparm__Z19CudaCSGSphereKernelPK6float3PfiiifS__p_Radius


.param .align 4 .b8 __cudaparm__Z19CudaCSGSphereKernelPK6float3PfiiifS__p_Center


.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

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 Restrictions

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