копировать глобальную память потоками CUDA

Question

Мне нужно скопировать один массив в глобальной памяти в другой массив в глобальной памяти потоками CUDA (не с хоста).

Мой код выглядит следующим образом:

__global__ void copy_kernel(int *g_data1, int *g_data2, int n)
{
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  int start, end;
  start = some_func(idx);
  end = another_func(idx);
  unsigned int i;
  for (i = start; i < end; i++) {
      g_data2[i] = g_data1[idx];
  }
}

Это очень неэффективно, потому что для некоторого idx область [start, end] очень большая, что заставляет этот поток выдавать слишком много команд копирования. Есть ли способ реализовать это эффективно?

Спасибо,

Чжэн

Answer 1

То, как вы это написали, я предполагаю, что каждый поток пытается записать весь блок «начало» - «конец».Что действительно действительно неэффективно.

вам нужно сделать что-то вроде этого.

___shared___ unsigned sm_start[BLOCK_SIZE];
___shared___ unsigned sm_end[BLOCK_SIZE];
sm_start[threadIdx.x] = start;
sm_end[threadIdx.y] = end;
__syncthreads();
for (int n = 0; n < blockdDim.x; n++) {
g_data2 += sm_start[n];
unsigned lim = sm_end[n] - sm_start[n];
  for (int i = threadIdx.x; i < lim; i += blockDim.x) {
      g_data2[i] = g_data1[idx];
  }
}

Answer 2

попробуйте использовать это:

CUresult cuMemcpyDtoD(
    CUdeviceptr dst,
    CUdeviceptr src,
    unsigned int bytes   
)   

ОБНОВЛЕНИЕ:

Вы правы: http://forums.nvidia.com/index.php?showtopic=88745

Нет эффективного способа сделать это правильно, потому что дизайнCUDA хочет, чтобы вы использовали только небольшое количество данных в ядре.