'inline' for global 函数以避免多重定义错误

Question

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

如果多个.cu文件包含kernel.cuh并调用vector_add[_kernel]，会导致link-time处的多个定义错误。在 C++ 中，可以使用 inline 限定符来避免此类错误。

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

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

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

下面是一个（not-so-minimal，抱歉）工作示例（在 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 类型的自然模板...

'inline' for global 函数以避免多重定义错误

'inline' for global functions to avoid multiple definition error

c++

linker

cuda

'inline' for __global__ 函数以避免多重定义错误

'inline' for __global__ functions to avoid multiple definition error

c++

linker

cuda

'inline' for global 函数以避免多重定义错误

'inline' for global functions to avoid multiple definition error