CUDA Pro Tip: Increase Performance with Vectorized Memory Access

Keywords: sass

Reference CUDA Pro Tip: Increase Performance with Vectorized Memory Access
I tried it, as follows. (Forget it, don't write it, test a bunch of data, write it up in trouble.)
This article mainly talks about vectorized load - vector loading (I think SIMD on CPU is a principle that loads more than one data at a time). Vector loading can be achieved through built-in variables provided by CUDA, such as int2, int4, float2, etc.

The code before optimization is as follows

__global__ void device_copy_scalar_kernel(int* d_in, int* d_out, int N) { 
  int idx = blockIdx.x * blockDim.x + threadIdx.x; 
  for (int i = idx; i < N; i += blockDim.x * gridDim.x) { 
    d_out[i] = d_in[i]; 
  } 
} 

void device_copy_scalar(int* d_in, int* d_out, int N) 
{ 
  int threads = 128; 
  int blocks = min((N + threads-1) / threads, MAX_BLOCKS);  
  device_copy_scalar_kernel<<<blocks, threads>>>(d_in, d_out, N); 
}

Optimized Code Version 1

__global__ void device_copy_vector2_kernel(int* d_in, int* d_out, int N) {
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  for (int i = idx; i < N/2; i += blockDim.x * gridDim.x) {
    reinterpret_cast<int2*>(d_out)[i] = reinterpret_cast<int2*>(d_in)[i];
  }

  //process remaining elements
  for(int i = idx + N/2 * 2; i<N; i += blockIdx.x*blockDim.x+threadIdx.x)
   d_out[i] = d_in[i];
}

void device_copy_vector2(int* d_in, int* d_out, int n) {
  threads = 128; 
  blocks = min((N/2 + threads-1) / threads, MAX_BLOCKS); 

  device_copy_vector2_kernel<<<blocks, threads>>>(d_in, d_out, N);
}

Optimized Code Version 2

__global__ void device_copy_vector4_kernel(int* d_in, int* d_out, int N) {
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  for(int i = idx; i < N/4; i += blockDim.x * gridDim.x) {
    reinterpret_cast<int4*>(d_out)[i] = reinterpret_cast<int4*>(d_in)[i];
  }

  //process remaining elements
  for(int i = idx + N/4 * 4; i<N; i += blockIdx.x*blockDim.x+threadIdx.x)
   d_out[i] = d_in[i];
}

void device_copy_vector4(int* d_in, int* d_out, int N) {
  int threads = 128;
  int blocks = min((N/4 + threads-1) / threads, MAX_BLOCKS);

  device_copy_vector4_kernel<<<blocks, threads>>>(d_in, d_out, N);
}

sass code can be viewed through cuobjdump to see the changes of loading instructions in different versions.

Posted by davidlenehan on Tue, 12 Feb 2019 13:00:18 -0800