内核调用期间仅用于设备计算的 CUDA 内存（类型）（计算 1.1 或 1.2）答案

【问题标题】：CUDA memory (type) for device only calculation during kernel calls (compute 1.1 or 1.2)内核调用期间仅用于设备计算的 CUDA 内存（类型）（计算 1.1 或 1.2）
【发布时间】：2013-04-08 09:52:11
【问题描述】：

我目前正在学习 CUDA，我的算法必须根据一些输入数据进行一些繁重的计算。这些计算是在一个最多旋转 1024 轮的循环中进行的。只要每个内核有少量线程（

我的解决方案是将繁重的计算拆分为多个内核调用：

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

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

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

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)
}

【问题讨论】：

如果你有足够的块，只需减小网格大小并多次启动内核，在内核内的块号上添加适当的偏移量，就会容易得多。
当您所能提供的只是包含大量“做某事”的模板实例化的伪代码时，很难就性能提出建议。你能把你真正想知道的更具体一点吗？
代码根本不重要，在这里发布太错误了。我在问我必须做什么/我可以使用哪种设备内存在内核调用之间保存数据（读取数据/写入数据的行）。性能与那段记忆有关。
如果内核超时确实是个问题，您可能会认真考虑在未连接显示器的 GPU 上运行它以避免超时。否则我会认真考虑 tera 的建议，除非线程之间存在阻止它的依赖关系。否则，您可能会遇到不必要的带宽瓶颈。
好的，我会尽量减小网格大小。但是与网格大小无关，是否有可能使用设备内存来执行此操作。什么时候是，怎么做？由于我正在学习过程中，很高兴知道。

标签： c++ c memory cuda

【解决方案1】：

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

#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 读取和写入数据期间进行合并访问。

【讨论】：