cuda核函数中动态分配显存

cuda核函数中动态分配显存

​ cuda在核函数kernel中可调用cudaMalloc动态分配显存，不过该实现该功能需要设备具有动态并行的能力，即GPU的架构需要>sm_35。默认的nvcc编译方式选择的架构是sm_30，其不支持动态并行特性，如果在编译时不选择设备计算能力会报如下错误，即在核函数中无法调用主机调用函数。

calling a __host__ function("cudaMalloc") from a __global__ function("xxx")is not allowed

​ 解决方案是在编译时加入编译选项，指定设备计算能力，并选择独立编译模式。.cu代码如下, 该代码在核函数中通过索引为(10,10)的线程分配全局内存空间。线程同步后，通过索引为(11,11)的线程来访问该全局内存，来确定我们分配的全局内存是有效的。

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


__global__ void malloc_test(double **a)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    int j = threadIdx.y + blockDim.y * blockIdx.y;
    if(i==10&&j==10)
    {
        int N = 10000;
        cudaMalloc((void**)a, sizeof(double)*N); //分配的地址记录在(*a)
        for(int i=0;i<N;i++)
        {
            (*a)[i] = i;
        } 
    }
    __syncthreads();
    if(i==11&&j==11)
    {
        printf("%f\n",(*a)[500]);
    }

}

int main()
{
    double **a;
    cudaMalloc((void**)&a, sizeof(double *)); //为a分配显存空间
    dim3 blockdim(3,3);
    dim3 griddim(6,6);
    malloc_test<<<griddim, blockdim>>>(a);
    cudaDeviceSynchronize();
    return 0;
}

编译方式makefile：

target = dynamic.out
objs = dynamic.o
NVCC = nvcc
CC = g++
CUGENCODE_FLAGS =  -arch=sm_50 -rdc=true  -std=c++11
LDFLAGS  = -L/usr/local/cuda/lib64/  -lcudart   -lcudadevrt

$(target):$(objs)
	$(NVCC) -arch=sm_50 -dlink -o dynamic_link.o dynamic.o 
	$(CC) dynamic_link.o dynamic.o $(LDFLAGS) -o $(target) 

%.o:%.cu
	$(NVCC) -c  $(CUGENCODE_FLAGS) $^

.PHONY:clean  
clean:
	rm -f $(target) $(objs) dynamic_link.o

命令行编译方式：

nvcc -c  -arch=sm_50 -rdc=true  -std=c++11 dynamic.cu
nvcc -arch=sm_50 -dlink -o dynamic_link.o dynamic.o
g++ dynamic_link.o dynamic.o -L/usr/local/cuda/lib64/ -lcudart -lcudadevrt -o dynamic.out

结果输出：