2

我目前正在学习 CUDA，我的算法必须根据一些输入数据进行一些繁重的计算。这些计算是在一个最多旋转 1024 轮的循环中进行的。只要每个内核有少量线程（< 100´000），一切都可以正常工作，但是如果我想使用更多线程，内核将被窗口中断，因为它需要很长时间才能完成。

我的解决方案是在几个内核调用中拆分繁重的计算：

  1. 主内核，准备输入数据并计算前 x 轮（循环展开）。每个输入只会调用一次。
  2. 工作内核，执行接下来的 x 轮（循环展开）。这将根据需要经常调用以计算所有所需的回合。

在每个内核调用（一个main，许多工作）之间，我必须保存 16 + length 个字节的数据，这些数据将在下一次调用中使用（长度是输入的长度，每个main调用都是固定的）。主内核将初始写入这些字节，工作内核将读取它们，运行下一个计算并用新结果写入原始数据。我只需要设备上的那些数据，不需要主机访问。我必须为此使用哪种内存？至少它必须是全局内存，因为它是内核调用期间唯一持久的可写内存，不是吗？但是，然后呢？您能否就如何继续使用正确的内存（和最佳性能）给我一个建议？

在“伪代码”中，它可能如下所示：

prepare memory to hold threads * (16 + length) bytes

for length = 1 to x step 1
  call mainKernel
  rounds = 1024 - rounds_done_in_main
  for rounds to 0 step rounds_done_in_work
    call workKernel
  end for
end for

cleanup memory

--------

template <unsigned char length> __global__ mainKernel() {
  unsigned char input[length];
  unsigned char output[16];
  const int tid = ...;

  devPrepareInput<length>(input);

  calc round 1: doSomething<length>(output, input)
  calc round 2: doSomething<length>(output, output + input) // '+' == append

  write data to memory based on tid // data == output + input
}

template <unsigned char length, remaining rounds> __global__ workKernel() {
  unsigned char *input;
  unsigned char *output;
  const int tid = ...;

  read data from memory based on tid
  ouput = data
  input = data+16

  if rounds >= 1
    calc round x  : doSomething<length>(output, output + input)
  if rounds >= 2
    calc round x+1: doSomething<length>(output, output + input) // '+' == append

  if rounds == x // x is the number of rounds in the last work call
    do final steps on output
  else
    write ouput + input to memory based on tid (for next call)
}
4

1 回答 1

1

是的，您可以使用设备内存来执行此操作。声明的变量__device__提供了一个缓冲区的静态声明，内核可以直接使用，不需要任何cudaMemcpy操作，也不需要将指针显式传递给内核。由于它具有应用程序的生命周期，因此其中的数据将从一个内核调用持续到下一个内核调用。

#define NUM_THREADS 1024
#define DATA_PER_THREAD 16
__device__ int temp_data[NUM_THREADS*DATA_PER_THREAD];

__global__ my_kernel1(...){
  int my_data[DATA_PER_THREAD] = {0};
  int idx = threadIdx.x + blockDim.x * blockIdx.x;
  // perform calculations

  // write out temp data
  for (int i = 0; i < DATA_PER_THREAD; i++) temp_data[i + (idx * DATA_PER_THREAD)] = my_data[i];
  }

__global__ my_kernel2(...){
  int my_data[DATA_PER_THREAD];
  // read in temp data
  for (int i = 0; i < DATA_PER_THREAD; i++) my_data[i] = temp_data[i + (idx * DATA_PER_THREAD)];
  // perform calculations

  }

有多种方法可以根据您在内核中的使用模式进行优化。数据传入/传出my_data并不是真正必要的。显然，您的内核代码可以通过适当的索引temp_data直接访问。my_data

如果您确实决定加载/存储它，您可以交错数据以允许在for-loop 读取和写入数据期间进行合并访问。

于 2014-02-10T05:57:25.443 回答