cuda - 从内核调用内核

Question

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

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

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

score 8 · Accepted Answer

这个问题被问了很久。当我用谷歌搜索同样的问题时，我到了这个页面。好像我得到了解决方案。

解决方案：

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

__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

希望能帮助到你。谢谢！[1]：http: //developer.download.nvidia.com/assets/cuda/docs/TechBrief_Dynamic_Parallelism_in_CUDA_v2.pdf

score 5 · Accepted Answer

编辑（2016）：

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

原答案：

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

cuda - 从内核调用内核

2 回答 2

Related

Reference