__CUDA_ARCH__ 带有 Thrust 执行策略的标志

Question

我有一个 __host__ __device__ 函数，它是调用 thrust 库的 "sort" 函数的包装器。在这个包装器中，我使用 __CUDA_ARCH__ 标志将执行策略设置为 "thrust::device"（当从主机调用时）和 "thrust::seq"（当从设备调用时）。以下代码生成运行时错误 -

#ifndef __CUDA_ARCH__
    thrust::stable_sort(thrust::device, data, data + num, customGreater<T>());
#else
    thrust::stable_sort(thrust::seq, data, data + num, customGreater<T>());
#endif

错误是-

意外的标准异常： What() is:merge_sort：第二步失败：无效的设备功能

据我理解，CUDA_ARCH可以用于条件编译。我请求帮助理解为什么会抛出此错误。

Answer 1

您似乎正在使用某些算法（包括排序）的 this issue. In a nutshell, thrust uses CUB 功能。您在代码中使用了 __CUDA_ARCH__ 宏，它环绕着使用 CUB 的推力算法调用，这会干扰期望能够将此宏用于所有路径的 CUB 代码。

一个可能的解决方法是 "your own dispatch":

$ cat t142.cu
#include <iostream>
#include <thrust/sort.h>
#include <thrust/execution_policy.h>


template <typename T>
struct customGreater {
__host__ __device__ bool operator()(T &t1, T &t2){
   return (t1 > t2);}
};

template <typename T>
__host__ __device__
void my_sort_wrapper(T *data, size_t num){
    int hostdev = 0;  // 0=device code
#ifndef __CUDA_ARCH__
    hostdev = 1;  // 1=host code
#endif
    if (hostdev == 0) thrust::stable_sort(thrust::seq, data, data + num, customGreater<T>());
    else thrust::stable_sort(thrust::device, data, data + num, customGreater<T>());

}

template <typename T>
__global__ void my_dev_sort(T *data, size_t num){
  my_sort_wrapper(data, num);
}
typedef int mytype;
const size_t sz = 10;
int main(){
  mytype *d_data;
  cudaMalloc(&d_data, sz*sizeof(mytype));
  cudaMemset(d_data, 0, sz*sizeof(mytype));
  my_sort_wrapper(d_data, sz);
  my_dev_sort<<<1,1>>>(d_data, sz);
  cudaDeviceSynchronize();
}
$ nvcc t142.cu -o t142
$ cuda-memcheck ./t142
========= CUDA-MEMCHECK
========= ERROR SUMMARY: 0 errors
$

通过这种实现，__CUDA_ARCH__ 宏的使用不会干扰推力算法的编译。

另一种可能的解决方法是对这两种情况简单地使用 thrust::device 策略（不分派 - 仅调用推力算法）。除了 CUDA 动态并行的情况外，thrust::device 在设备代码中使用时将 "decay" 到 thrust::seq。

我预计只有当推力算法在底层实现中使用 CUB 功能时，这些建议才会 necessary/relevant。

如果您不喜欢这种行为，可以提交 thrust issue。

Answer 2

很遗憾，我们无法在 Thrust 中修复此问题。这里的麻烦是 NVCC 编译器需要在主机编译期间查看所有 __global__ 函数模板实例化（例如，当 __CUDA_ARCH__ 未定义时），否则内核将被视为未使用并被丢弃。有关详细信息，请参阅 this CUB GitHub issue。

正如罗伯特所建议的那样，像这样的解决方法应该没问题：

#include <iostream>
#include <thrust/sort.h>
#include <thrust/execution_policy.h>

template <typename T>
struct customGreater {
__host__ __device__ bool operator()(T &t1, T &t2){
   return (t1 > t2);}
};

#if defined(__CUDA_ARCH__)
  #define DEVICE_COMPILATION 1
#else
  #define DEVICE_COMPILATION 0
#endif

template <typename T>
__host__ __device__
void my_sort(T *data, size_t num){
  if (DEVICE_COMPILATION)
    thrust::stable_sort(thrust::device, data, data + num, customGreater<T>());
  else
    thrust::stable_sort(thrust::seq, data, data + num, customGreater<T>());
}

template <typename T>
__global__ void my_dev_sort(T *data, size_t num){
  my_sort(data, num);
}
typedef int mytype;
const size_t sz = 10;
int main(){
  mytype *d_data;
  cudaMallocManaged(&d_data, sz*sizeof(mytype));
  cudaMemset(d_data, 0, sz*sizeof(mytype));
  my_sort(d_data, sz);
  my_dev_sort<<<1,1>>>(d_data, sz);
  cudaFree(d_data);
  cudaDeviceSynchronize();
}

__CUDA_ARCH__ 带有 Thrust 执行策略的标志

__CUDA_ARCH__ flag with Thrust execution policy

cuda

gpu

thrust