Difference between 4 sequential 32-bit loads and 1 128-bit (4x32-bit) vector load.

mem4.cu

#define KERNEL_QUALIFIERS  extern "C" __global__

//
//
//

#define REPEAT1()                               \
  REPS(0) 

#define REPEAT4()                               \
  REPS(0)                                       \
  REPS(1)                                       \
  REPS(2)                                       \
  REPS(3) 

//
//
//

KERNEL_QUALIFIERS
void
mem4(const float* const fin, float* const fout)
{
  const unsigned int tidx = threadIdx.x * 4;

#undef  REPS
#define REPS(n)  const float f##n = fin[tidx+n];

  REPEAT4();

#undef  REPS
#define REPS(n)  fout[tidx+n] = f##n;

  REPEAT4();
}

//
//
//

KERNEL_QUALIFIERS
void
mem1(const float4* const fin, float4* const fout)
{
  const unsigned int tidx = threadIdx.x;

#undef  REPS
#define REPS(n)  const float4 f##n = fin[tidx+n];

  REPEAT1();

#undef  REPS
#define REPS(n)  fout[tidx+n] = f##n;

  REPEAT1();
}

//
//
//

Compiled with: nvcc -m 32 -arch sm_35 -Xptxas=-v,-abi=no -cubin mem4.cu

Dumped with: nvdisasm -c mem4.cubin

//--------------------- .text.mem4                --------------------------
mem4:
.text.mem4:
        /*0008*/                S2R R4, SR_TID.X;
        /*0010*/                ISCADD R5, R4, c[0x0][0x140], 0x4;
        /*0018*/                LD R3, [R5];
        /*0020*/                ISCADD R4, R4, c[0x0][0x144], 0x4;
        /*0028*/                LD R2, [R5+0x4];
        /*0030*/                LD R1, [R5+0x8];
        /*0038*/                LD R0, [R5+0xc];
        /*0048*/                ST [R4], R3;
        /*0050*/                ST [R4+0x4], R2;
        /*0058*/                ST [R4+0x8], R1;
        /*0060*/                ST [R4+0xc], R0;
        /*0068*/                EXIT ;
.L_1:
        /*0070*/                BRA `(.L_1);
.L_19:

//--------------------- .text.mem1                --------------------------
mem1:
.text.mem1:
        /*0008*/                S2R R0, SR_TID.X;
        /*0010*/                ISCADD R1, R0, c[0x0][0x140], 0x4;
        /*0018*/                ISCADD R0, R0, c[0x0][0x144], 0x4;
        /*0020*/                LDG.128 R4, [R1];
        /*0028*/                TEXDEPBAR 0x0;
        /*0030*/                ST.128 [R0], R4;
        /*0038*/                EXIT ;
.L_2:
        /*0040*/                BRA `(.L_2);
.L_20: