Inspecting the difference between kernel arguments and __constant__ variables.

symarg.cu

//
//
//

#define KERNEL_QUALIFIERS   extern "C" __global__

//
//
//

KERNEL_QUALIFIERS
void
argKernel(const int* const argIn, int* const argOut)
{
  argOut[threadIdx.x] = argIn[threadIdx.x];
}

//
//
//

__constant__ const int* symIn;
__constant__       int* symOut;

KERNEL_QUALIFIERS
void
symKernel()
{
  symOut[threadIdx.x] = symIn[threadIdx.x];
}

//
//
//

typedef struct
{
  const int* in;
  int*       out;
} KernelEnv;

KERNEL_QUALIFIERS
void
envKernel(const KernelEnv kenv)
{
  kenv.out[threadIdx.x] = kenv.in[threadIdx.x];
}

Architectures before sm_20 pass kernel parameters via shared: nvcc -m 32 -arch sm_12 -Xptxas=-v,-abi=no -cubin symarg.cu

    code for sm_12
        Function : symKernel
    /*0000*/     /*0x1004000500000003*/     MVI.U16 R0H, 0x4;
    /*0008*/     /*0x6101000500000780*/     IMAD.U16 R1, R0L, R0H, c [0x0] [0x0];
    /*0010*/     /*0x6101000900004780*/     IMAD.U16 R2, R0L, R0H, c [0x0] [0x1];
    /*0018*/     /*0xd00e020180c00780*/     GLD.U32 R0, global14 [R1];
    /*0020*/     /*0xd00e0401a0c00781*/     GST.U32 global14 [R2], R0;
        ..........................


        Function : argKernel
    /*0000*/     /*0x1100e804        */     MOV32 R1, g [0x4];
    /*0004*/     /*0x1100ea08        */     MOV32 R2, g [0x5];
    /*0008*/     /*0x6004000500000003*/     IMAD32I.U16 R1, R0L, 0x4, R1;
    /*0010*/     /*0x6004000900000003*/     IMAD32I.U16 R2, R0L, 0x4, R2;
    /*0018*/     /*0xd00e020180c00780*/     GLD.U32 R0, global14 [R1];
    /*0020*/     /*0xd00e0401a0c00781*/     GST.U32 global14 [R2], R0;
        ..........................


        Function : envKernel
    /*0000*/     /*0x1100e804        */     MOV32 R1, g [0x4];
    /*0004*/     /*0x1100ea08        */     MOV32 R2, g [0x5];
    /*0008*/     /*0x6004000500000003*/     IMAD32I.U16 R1, R0L, 0x4, R1;
    /*0010*/     /*0x6004000900000003*/     IMAD32I.U16 R2, R0L, 0x4, R2;
    /*0018*/     /*0xd00e020180c00780*/     GLD.U32 R0, global14 [R1];
    /*0020*/     /*0xd00e0401a0c00781*/     GST.U32 global14 [R2], R0;
        .............................