Модификация базового примера VECADD для использования общей памяти

Question

Я написал следующее ядро ​​для использования разделяемой памяти в базовом примере CUDA vecadd (сумма двух векторов).Код работает, но истекшее время выполнения ядра совпадает с основным исходным кодом.Может кто-нибудь предложить мне способ ускорить такой код?

__global__ void vecAdd(float *in1, float *in2, float *out,long int len) 
{
 __shared__ float s_in1[THREADS_PER_BLOCK];
 __shared__ float s_in2[THREADS_PER_BLOCK];

 unsigned int xIndex = blockIdx.x * THREADS_PER_BLOCK + threadIdx.x;

 s_in1[threadIdx.x]=in1[xIndex];
 s_in2[threadIdx.x]=in2[xIndex];

 out[xIndex]=s_in1[threadIdx.x]+s_in2[threadIdx.x];
}

Answer 1

Может кто-нибудь предложить мне способ ускорить такой код

Практически нет полезных оптимизаций для такой операции, как сложение векторов.Из-за характера вычислений код может рассчитывать только на достижение 50% пиковой арифметической пропускной способности, а требование для трех транзакций памяти на FLOP делает эту операцию встроенной с ограниченной пропускной способностью памяти.

В результатеэто:

__global__ void vecAdd(float *in1, float *in2, float *out, unsigned int len) 
{
 unsigned int xIndex = blockIdx.x * blockDim.x + threadIdx.x;

 if (xIndex < len) {
  float x = in1[xIndex];
  float y = in2[xIndex];
  out[xIndex] = x + y;
 }
}

- наиболее эффективный вариант на новейшем оборудовании, если размер блока выбран для максимальной загрузки, а len достаточно велик, например:

  int minGrid, minBlockSize;
  cudaOccupancyMaxPotentialBlockSize(&minGrid, &minBlockSize, vecAdd);
  int nblocks = (len / minBlockSize) + ((len % minBlockSize > 0) ? 1 : 0);
  vecAdd<<<nblocks, minBlockSize>>>(x, y, z, len);