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

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

以下面的代码为例：

#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(),print_func());
//#endif
}

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

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

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

解决方法

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

#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(),print_func());
    }
}

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