CUDA在核函数内调用核函数(动态并行 Dynamic Parallelism)

时间：2024-03-31 13:06:47

CUDA 5.0之后支持global函数内调用global函数，也就是核函数内调用核函数，即核函数的嵌套调用，也可以实现递归调用（暂未测试）。需要保证GPU计算能力3.5及以上。

然后进行如下设置：1.在项目属性中, 设定 CUDA C/C++==>common 的Generate Relocatable Device Code为True(-rdc=true). 

 2. 附加依赖项添加  cudadevrt.lib

以下是自己写的测试代码：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>
#include <Windows.h>
#include<stdio.h>

using namespace std;


__global__ void AplusThree(int *ret)
{
ret[threadIdx.x] += 3;
//printf("ֵ:%d", ret[threadIdx.x]);
printf("from AplusThree function\n");
}

__global__ void AplusB(int *ret, int a, int b)
{
ret[threadIdx.x] += a + b + threadIdx.x;
AplusThree << < 1, 5 >> >(ret);
//printf("ֵ:%d", ret[threadIdx.x]);
}


void test()
{
int ret[5] = { 1,2,3,4,5 };
int a[5] = { 0 };
int *dev_ret;
cudaError_t cudaStatus;
cudaStatus = cudaMalloc((void**)&dev_ret, sizeof(int) * 5);
cudaStatus = cudaMemcpy(dev_ret, ret, sizeof(int) * 5, cudaMemcpyHostToDevice);
AplusB << < 1, 5 >> >(dev_ret, 10, 100);
cudaStatus = cudaMemcpy(a, dev_ret, sizeof(int) * 5, cudaMemcpyDeviceToHost);
cudaDeviceSynchronize();


for (int i = 0; i<5; i++)
{
cout << "A+B = " << a[i] << endl;
}


cudaFree(dev_ret);
}


int main()
{
cudaEvent_t start, stop;
float elapseTime = 0;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);

test();

cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapseTime, start, stop);
cout << elapseTime<<" ms" << endl;


system("pause");
return 0;

}

测试结果如下:

CUDA在核函数内调用核函数(动态并行 Dynamic Parallelism)


以下是转载的其他博客：

一、

error : calling a global function(“childKernel”) from a global function(“kernel”) is only allowed on the compute_35 architecture or above

  • 原因及解决方法：

这是因为默认计算能力被设定成了sm_20,compute_20,从而阻止你使用动态并行.

解决方案: 
在您的项目属性中, 设定CUDA C/C++中的代码生成为:compute_50,sm_50 
这样即可让您的5.0的卡, 支持动态并行, 也就是您说的核函数调用核函数.

二、

改成compute_50,sm_50了，然后从网上找了一个例子，还是报错：错误 17 error : kernel launch from device or global functions requires separate compilation mode 
请问这个“独立编译模式”要怎么配置？

  • 原因及解决方法： 
    - 
    其实只需要将RDC(可重定位设备代码)打开即可.

方式1: 
在项目属性中, 设定 CUDA C/C++==>common 的Generate Relocatable Device Code为True(-rdc=true). 
同时所有所有cu文件中的该属性为”继承自项目”.

方式2: 
手工将你所有的cu文件该属性设定为”真”. 
然后可以不管项目属性.

任选一种.

三、

error LNK2001: unresolved external symbol ___fatbinwrap_66_tmpxft_00000b3c_00000000_17_cuda_device_runtime_compute_52_cpp1_ii_8b1a5d37.

  • 原因及解决方法：

    cuda从5.0版本之后开始支持dynamic parallelism，即可以在global函数中调用其他global函数，因此可以实现核函数中再调用核函数。 
    dynamic parallelism（动态并行）的软硬件条件有：

    1. cuda toolKit 版本5.0或以上；
    2. GPU compute capability（计算能力）3.5及以上。

此时就可以在核函数中调用另一个核函数了，也可以实现递归调用。但这时会出现如下link error： 
error LNK2001: unresolved external symbol ___fatbinwrap_66_tmpxft_00000b3c_00000000_17_cuda_device_runtime_compute_52_cpp1_ii_8b1a5d37.

这是因为动态并行还需要附加另一个库：cudadevrt.lib。加入即可。


相关文章