CUDA跨文件传递函数指针触发地址未对齐错误排查

我实现了一个获取CUDA设备函数指针的功能，在同一.cu文件的内核中使用该指针调用函数完全正常，但将指针传递给其他.cu文件的内核时，执行会触发misaligned address error，无法理解错误原因。

复现代码 #

file1.cu（运行正常） #

typedef float (*funcPtr)(float);

__device__ float device_relu(float x) {
    return fmaxf(0, x);
}

__device__ funcPtr ptr_relu = device_relu;

funcPtr get_activation_function_cuda(int activation) {
    funcPtr host_function;
    
    switch (activation) {
        case RELU:
            gpuErrchk( cudaMemcpyFromSymbol(&host_function, ptr_relu, sizeof(funcPtr)));
            break;
        // 其他激活函数分支
    }
    return host_function;
}

__global__ void kernel1(funcPtr f, ...) {
    // 内核逻辑
    x[...] = (*f)(x[...]); // 此处调用无问题
}

void func1(int activation) {
    // 初始化参数等
    funcPtr activation_function = get_activation_function_cuda(activation);
    kernel1<<<gridSize, blockSize>>>(activation_function, ...);
}

file2.cu（触发错误） #

#include ".../file1.h"

__global__ void kernel2(..., funcPtr d_f) {
    // 计算逻辑
    float tmp=0;
    for (int j=0; j < size_output; j++) {
        tmp += output[j]*weights[idx][j];
    }
    input[idx] = tmp*( (*d_f)(input_z[idx]) ); // 此处触发misaligned address error
}

void backward_dense_device(...) {
    dim3 gridSize1(i_div_up(size_input, BLOCKSIZE_x));
    dim3 blockSize1(BLOCKSIZE_x);

    funcPtr d_function = get_activation_function_cuda(RELU);
    kernel2<<<gridSize1, blockSize1>>>(..., d_function); 
}

编译流程 #

nvcc  -g  -c src/file1.cu -o build/file1.o
nvcc  -g  -c src/file2.cu -o build/file2.o
# 其他.c和.cu文件编译步骤
gcc  -c src/main.c -o build/main.cuda.o  -Wall -Wextra -std=gnu99 -g -O3 -DUSE_CUDA -lcuda -I/opt/cuda/include
nvcc  -ljpeg -Xcompiler -fopenmp -g  build/main.cuda.o build/file1.o build/file2.o -o build/main

问题原因 #

核心问题是CUDA的__device__变量和函数默认具有内部链接属性（仅在当前编译单元可见）：

1. file1.cu中定义的__device__ funcPtr ptr_relu和__device__ float device_relu(float)默认是static级别的符号，无法被其他编译单元（如file2.cu）正确解析。
2. 虽然get_activation_function_cuda能在file1.cu内部正确获取设备函数指针，但当该指针传递到file2.cu的内核时，file2.cu的编译/链接流程无法验证该指针指向的设备函数符号的有效性，导致运行时访问无效地址，触发对齐错误。

解决方案 #

有两种可行的修复方式：

方案1：显式声明外部链接的设备符号 #

在file1.h中添加对设备函数和变量的外部声明，确保跨编译单元可见：

// file1.h
typedef float (*funcPtr)(float);

// 声明外部可见的设备函数和变量
extern __device__ float device_relu(float x);
extern __device__ funcPtr ptr_relu;

// 声明get_activation_function_cuda函数
funcPtr get_activation_function_cuda(int activation);

file1.cu中保持原有定义不变，无需修改。

方案2：直接从设备函数符号获取指针，跳过__device__变量中转 #

修改get_activation_function_cuda，直接从device_relu符号拷贝指针，避免依赖跨编译单元的__device__变量：

funcPtr get_activation_function_cuda(int activation) {
    funcPtr host_function;
    
    switch (activation) {
        case RELU:
            // 直接从device_relu符号拷贝，而非ptr_relu
            gpuErrchk( cudaMemcpyFromSymbol(&host_function, device_relu, sizeof(funcPtr)));
            break;
        // 其他分支同理
    }
    return host_function;
}

同时需在file1.h中声明extern __device__ float device_relu(float x);，确保符号能被链接器正确解析。

验证说明 #

• 两种方案都能确保设备函数指针在跨编译单元传递时，指向有效的设备地址，避免运行时地址错误。
• 若选择方案2，可直接移除__device__ funcPtr ptr_relu变量，减少不必要的设备端变量开销。

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