从内核调用内核

A跟进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); 
    } 
}

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

来源

2012-07-31 Doug

编辑（2016）：

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

原来的答案：

你将不得不等待，直到今年结束的时候，新一代的硬件是可用的。目前没有任何CUDA设备可以从其他内核启动内核 - 目前它不受支持。

来源

2012-07-31 19:15:38 talonmies

是否可以在最新版本的CUDA（v6.5）和具有计算能力3.0的NVIDIA Grid K520上执行此操作？ – 2015-03-25 19:09:49

@talonmies，我已经等了4年你的答案：D希望有这样的解决方案:-)你能指出我吗？ – Nabin 2016-03-11 05:41:05

没关系。我找到了什么 – Nabin 2016-03-11 05:56:59

我知道问这个问题已经很久了。当我搜索同样的问题时，我到了这个页面。似乎我得到了解决方案。

解决方案：

我伸手here莫名其妙，看到清凉的做法从另一个内核中启动内核。

__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

希望它有帮助。谢谢！

来源

2016-03-11 06:02:51 Nabin

从内核调用内核

回答

相关问题