为什么在 __CUDA_ARCH__ 中包装主机代码会导致无效的设备函子错误？答案

【问题标题】：Why does wrapping host code in __CUDA_ARCH__ cause an invalid device functor error?为什么在 __CUDA_ARCH__ 中包装主机代码会导致无效的设备函子错误？
【发布时间】：2021-07-28 13:45:18
【问题描述】：

以如下代码为例：

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

struct print_func {
    __device__ __host__ void operator()(int i) {
        printf("%d, ", i);
    }
};

struct functor {
    __device__ __host__ bool operator()(int i) {
        return i % 2 == 0;
    }
};

int main() {
    thrust::device_vector<int> vec(10);
    thrust::sequence(vec.begin(), vec.end());

//#ifndef __CUDA_ARCH__
    auto newLast = thrust::remove_if(vec.begin(), vec.end(), functor());
    vec.resize(thrust::distance(vec.begin(), newLast));
    thrust::for_each(vec.begin(), vec.end(), print_func());
//#endif
}

如果你取消注释预处理器条件（理论上应该没有影响，因为__CUDA_ARCH__没有在主机端定义），突然抛出CUDA error 98: invalid device function运行时错误。

为什么会这样，我该如何正确解决这个问题？

对于一些额外的上下文，我在尝试从单个 __host__ __device__ 函数实现单独的主机和设备代码时遇到了这个问题。

【问题讨论】：

Can't reproduce this on Godbolt。也许问题出在您的个人系统上？
@einpoklum：Godbolt 如何产生 runtime 错误？
“对于一些额外的上下文，我在尝试从单个 __host__ __device__ 函数实现单独的主机和设备代码时遇到了这个问题”——你不能这样做。语言不允许
而您的预处理器定义正在破坏编译轨迹。尽管您认为会发生什么情况，但该推力代码会同时发出主机代码和设备代码，这两者都必须进行编译。您的预处理器节正在阻止发出一些必要的运行时样板并破坏一切
@talonmies 在 cuda 中绝对有可能拥有 __host__ __device__ 函数，并且从 3.0 版开始，我相信您可以使用 __CUDA_ARCH__ 在所述函数中将设备和主机代码分开。 This is the earliest example of this I can find.

标签： c++ cuda thrust

【解决方案1】：

根据 Robert Crovella 非常有启发性的评论，这个问题是因为在编译步骤中，thrust 没有机会包含其设备端代码。一个基本的修复将类似于以下内容：

#include <thrust/device_vector.h>

struct print_func {
    __device__ __host__ void operator()(int i) {
        printf("%d, ", i);
    }
};

struct functor {
    __device__ __host__ bool operator()(int i) {
        return i % 2 == 0;
    }
};

int main() {
    thrust::device_vector<int> vec(10);
    thrust::sequence(vec.begin(), vec.end());

    if (THRUST_IS_HOST_CODE) {
        auto newLast = thrust::remove_if(vec.begin(), vec.end(), functor());
        vec.resize(thrust::distance(vec.begin(), newLast));
        thrust::for_each(vec.begin(), vec.end(), print_func());
    }
}

但是，这种策略仍然存在局限性。例如，如果您调用任何仅限设备的函数，这将不会编译。因此，应该避免这种实现方式。

【讨论】：