问题背景 #

已编译导出设备函数的DLL，尝试在CUDA核函数中调用该设备函数时，内核启动行kernel <<<10, 1024 >>> (a_d, c_d);报expected an expression错误，注释掉设备函数调用后编译正常。

DLL源码（myfunc.cu）：

#include "cuda_runtime.h"

__declspec(dllexport) __device__ double myfunc(double x) {
    return x * x;  // Example function: square the input
}

编译命令：

nvcc -shared -o myfunc.dll myfunc.cu

调用程序（kernel.cu）：

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

__device__ double myfunc(double x);

__global__ void kernel(double* a, double* c) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    c[idx] = myfunc(a[idx]);
}

int main() {
    int num_elements = 10000;
    double* a_h, * c_h;
    double* a_d, * c_d;

    // 分配主机内存
    cudaMallocHost(&a_h, num_elements * sizeof(double));
    cudaMallocHost(&c_h, num_elements * sizeof(double));

    // 分配设备内存
    cudaMalloc(&a_d, num_elements * sizeof(double));
    cudaMalloc(&c_d, num_elements * sizeof(double));

    // 主机到设备数据拷贝
    cudaMemcpy(a_d, a_h, num_elements * sizeof(double), cudaMemcpyHostToDevice);

    // 启动内核
    kernel <<<10, 1024 >>> (a_d, c_d);

    // 设备到主机数据拷贝
    cudaMemcpy(c_h, c_d, num_elements * sizeof(double), cudaMemcpyDeviceToHost);

    // 释放内存
    cudaFree(a_d);
    cudaFree(c_d);
    cudaFreeHost(a_h);
    cudaFreeHost(c_h);

    return 0;
}

错误原因 #

Windows的__declspec(dllexport)是为CPU函数设计的导出机制，CUDA设备函数无法通过普通DLL导出：

1. 设备函数运行在GPU上，编译后生成的是PTX或GPU二进制代码，不属于DLL的CPU代码段，链接器无法正确解析其引用。
2. 当核函数依赖未正确解析的设备函数时，编译器会进入异常状态，误报内核启动语法<<<>>>的错误（实际是链接层面的设备函数未定义问题）。

正确实现方式 #

方法1：使用CUDA静态库（推荐） #

将设备函数编译为静态库，让nvcc在编译主程序时直接合并设备函数代码，避免链接问题：

1. 编译静态库：
   执行命令生成静态库：

   nvcc -lib -o myfunc.lib myfunc.cu
   

   或者分两步：

   nvcc -c myfunc.cu -o myfunc.obj
   lib /OUT:myfunc.lib myfunc.obj
   

2. 编译主程序：
   • 保留kernel.cu中__device__ double myfunc(double x);的声明
   • 在Visual Studio中，将myfunc.lib添加到链接器输入
   • 直接编译运行即可，此时设备函数的代码会被正确嵌入到核函数中，内核启动语法也能正常解析。

方法2：使用CUDA动态库（复杂场景） #

如果需要动态加载设备函数，需使用CUDA Driver API加载包含设备函数的模块，步骤如下：

1. 编译PTX或二进制模块：
   先编译生成PTX文件：

   nvcc -ptx myfunc.cu -o myfunc.ptx
   

2. 在主程序中加载模块：
   使用CUDA Driver API的cuModuleLoad加载PTX，通过cuModuleGetFunction获取设备函数的句柄，再在核函数中调用（需结合动态并行或显式函数调用）。这种方式复杂度较高，仅适合动态扩展GPU功能的场景。

内容的提问来源于stack exchange，提问作者Rich Tanenbaum