глобальная и разделяемая память в CUDA

Question

У меня есть два ядра CUDA, которые вычисляют подобные вещи. Один использует глобальную память (myfun - это функция устройства, которая много читает из глобальной памяти и выполняет вычисления). Второе ядро ​​передает этот кусок данных из глобальной памяти в общую память, чтобы данные могли быть разделены между различными потоками блока. Мое ядро, которое использует глобальную память, намного быстрее, чем ядро ​​с общей памятью. Каковы возможные причины?

loadArray просто копирует небольшую часть d_x в m.

__global__ void mykernel(float *d_x, float *d_y, int *d_z, float *d_u, int N, int K, int D)
{

  int tid = blockIdx.x*blockDim.x + threadIdx.x;
  int index = 0;
  float max_s = 1e+37F;


  if (tid < N)
    {

      for (int i = 0; i < K; i++)
        {

          float s = myfun(&d_x[i*D], d_y, tid);

          if (s > max_s)
            {
              max_s = s;
              index = i;
            }
        }

      d_z[tid] = index;
      d_u[tid] = max_s;
    }
}

Использование общей памяти:

__global__ void mykernel(float *d_x, float *d_y, int *d_z, float *d_u, int N, int K)
{
  int tid = blockIdx.x*blockDim.x + threadIdx.x;
  int index = 0;
  float max_s = 1e+37F;

  extern __shared__ float m[];
  if( threadIdx.x == 0 )
    loadArray( m, d_x );
  __syncthreads();

  if (tid < N)
    {

      for (int i = 0; i < K; i++)
        {

          float s = myfun(m, d_y, tid);

          if (s > max_s)
            {
              max_s = s;
              index = i;
            }
        }

      d_z[tid] = index;
      d_u[tid] = max_s;
    }
}

Answer 1

Проблема в том, что только первый поток в каждом блоке читает из глобальной памяти в общую память, это намного медленнее, чем одновременное чтение всех потоков из глобальной памяти.

Использование разделяемой памяти является преимуществом, когда одному потоку необходим доступ к соседним элементам из глобальной памяти, но здесь это не так.

Answer 2

IMO, если у вас есть параллельный nsight , установленный, например, на компьютере с Windows и отслеживающий выполнение, у вас может быть больше идей.Либо запустите cudaprof в своем приложении, чтобы попытаться выяснить, где возможны задержки.