从内核调用内核

Question

后续Q来自：CUDA: Calling a __device__ function from a kernel

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

// 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体系结构或更高版本的。

终端方式您可以从termial运行上面的父内核(最终将运行子内核)。已在Linux计算机上验证。

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

希望能有所帮助。谢谢！1：http://developer.download.nvidia.com/assets/cuda/docs/TechBrief_Dynamic_Parallelism_in_CUDA_v2.pdf