从内核调用内核

Question

后续问答来自：CUDA：从内核调用__device__函数

我正在尝试加快排序操作。简化的伪版本如下：

// some costly swap operation
__device__ swap(float* ptrA, float* ptrB){
  float saveData;         // swap some 
  saveData= *Adata;       //   big complex
  *Adata= *Bdata          //     data chunk
  *Bdata= saveData;
}

// a rather simple sort operation
__global__ sort(float data[]){
  for (i=0; i<limit: i++){
  find left swap point
  find right swap point
  swap<<<1,1>>>(left, right);
  }
}

（注意：这个简单的版本没有在块中显示还原技术。）这个想法是很容易（快速）识别交换点。交换操作成本高（缓慢）。因此，使用一个块来查找/标识交换点。使用其他块执行交换操作。即并行进行实际交换。这听起来像一个不错的计划。但是，如果编译器内联设备调用，则不会发生并行交换。有没有办法告诉编译器不要内联设备调用？

Answer 1

编辑（2016）：

动态并行是在第二代开普勒架构 GPU 中引入的。在计算能力 3.5 及更高版本的设备上支持在设备中启动内核。

原答案：

您将不得不等到今年年底下一代硬件可用时。当前没有 CUDA 设备可以从其他内核启动内核 - 目前不受支持。

Answer 2

我知道这个问题已经问了很长时间了。当我用谷歌搜索同样的问题时，我到了这个页面。好像我得到了解决方案。

解决方案：

我以某种方式到达这里并看到了从另一个内核中启动内核的酷方法。

__global__ void kernel_child(float *var1, int N){
    //do data operations here
}


__global__ void kernel_parent(float *var1, int N)
{
    kernel_child<<<1,2>>>(var1,N);
} 

cuda 5.0 及更高版本的动态并行性使这成为可能。此外，在运行时请确保您使用compute_35 架构或更高版本。

终端方式 您可以从终端运行上述父内核（最终将运行子内核）。在 Linux 机器上验证。

$ nvcc -arch=sm_35 -rdc=true yourFile.cu
$ ./a.out

希望能帮助到你。谢谢！