__global__ 函数的“内联”以避免多重定义错误

Question

我有一个 CUDA 模板库，其中一个函数实际上 不是 模板，但 是 在 .cuh 标头中定义的。 （vector_add_kernel 在下面的kernel.cuh。）

如果多个.cu文件包含kernel.cuh并调用vector_add[_kernel]，将导致链接时出现多个定义错误。在 C++ 中，可以使用 inline 限定符来避免此类错误。

但是，inline __global__ ... - 在我的系统上防止多重定义错误 - 导致警告inline 限定符已被忽略。

问：有没有更好的方法来避免多重定义错误，或者只针对这个函数抑制这个警告？ inline __global__ 是否安全，或者其他主机编译器真的会忽略它吗？

我可以简单地将vector_add_kernel 移动到一个单独的.cu 文件中，但这将是唯一 非头文件。我也可以模板vector_add_kernel，但在我的库中这没什么意义。

一个（不太简单，抱歉）工作示例（在 Debian 上使用 CUDA 7.0、gcc 4.7.2 测试）如下。

为了澄清，main.cu 是一些用户的代码； lib.cu 是一些外部库不属于我； kernel.cuh 是我的模板库的一部分。所以，外部的lib 和用户的main 都在使用我的模板库kernel.cuh - 但是是分开的。

main.cu:

#include "lib.hpp"
#include "kernel.cuh"

#include <thrust/device_vector.h>
#include <thrust/host_vector.h>

#include <cstddef>
#include <cstdlib>
#include <iostream>

int main(void)
{
    const size_t N = 1u << 7;

    float* a = (float*) malloc(N * sizeof(float));
    float* b = (float*) malloc(N * sizeof(float));
    float* c = (float*) malloc(N * sizeof(float));

    for (int i = 0; i < N; ++i) {
        a[i] = b[i] = 2.0f * i;
    }

    lib_vector_add(a, b, c, N);
    for (int i = 0; i < N; ++i) {
        if (c[i] != 2.0f * i + 2.0f * i)
            std::cout << "Error, lib, element " << i << std::endl;
    }

    thrust::device_vector<float> d_a(a, a + N);
    thrust::device_vector<float> d_b(b, b + N);
    thrust::device_vector<float> d_c(N);

    vector_add(d_a, d_b, d_c);
    thrust::host_vector<float> h_c = d_c;
    for (int i = 0; i < N; ++i) {
        if (h_c[i] != 2.0f * i + 2.0f * i)
            std::cout << "Error, element " << i << std::endl;
    }
}

lib.cu,

#include <kernel.cuh>

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

void lib_vector_add(float* a, float* b, float* c, size_t N)
{
    thrust::host_vector<float> h_a(a, a + N);
    thrust::host_vector<float> h_b(b, b + N);

    thrust::device_vector<float> d_a = h_a;
    thrust::device_vector<float> d_b = h_b;
    thrust::device_vector<float> d_c(N);

    vector_add(d_a, d_b, d_c);

    thrust::host_vector<float> h_c = d_c;
    for (int i = 0; i < N; ++i)
    {
        c[i] = h_c[i];
    }
}

lib.hpp,

#pragma once

#include <cstddef>

void lib_vector_add(float*, float*, float*, size_t);

kernel.cuh - 这种形式会导致链接器错误。取消注释第一个 inline 以获得有效代码。

#pragma once

#include <thrust/device_vector.h>
#include <cstddef>

// inline keyword avoids multiple definition errors, but produces warnings.
// UNCOMMENT TO GET A WORKING EXECUTABLE.
// inline
__global__ void vector_add_kernel(
    const float *const a,
    const float *const b,
    float *const c,
    const size_t N)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    while (tid < N)
    {
        c[tid] = a[tid] + b[tid];
        tid += blockDim.x * gridDim.x;
    }
}

// inline produces no warnings.
inline
void vector_add(
    const thrust::device_vector<float>& d_a,
    const thrust::device_vector<float>& d_b,
    thrust::device_vector<float>& d_c)
{
    const float *const a_ptr = thrust::raw_pointer_cast(d_a.data());
    const float *const b_ptr = thrust::raw_pointer_cast(d_b.data());
    float *const c_ptr = thrust::raw_pointer_cast(d_c.data());

    const size_t N = d_a.size();

    dim3 block(128);
    dim3 grid((N + 127) / 128);

    vector_add_kernel<<<grid, block>>>(a_ptr, b_ptr, c_ptr, N);
}

Makefile

OBJS = main.o lib.o
DEPS = kernel.cuh
CU_ARCH = -gencode arch=compute_20,code=sm_20

all: app

app: $(OBJS)
    nvcc $(CU_ARCH) $(OBJS) -o app

%.o: %.cu $(DEPS)
    nvcc $(CU_ARCH) -dc -I./ $< -o $@

clean:
    -rm *.o

Answer 1

如果你想保持你当前的代码组织，你有一个非常简单的解决方案，就是声明你的内核static（代替你的inline关键字）。这将防止链接器抱怨，但会生成与包含 kernel.cuh 的编译单元（目标文件）一样多的不同版本的内核。

另一个解决方案是模板化你的内核。我知道您已经排除了这种可能性，但您应该重新考虑一下，因为您的内核是 float 输入参数类型的自然模板...