推力：为什么总是执行主机代码，尽管有 __CUDA_ARCH__答案

【问题标题】：Thrust: why always host code is executed in spite of __CUDA_ARCH__推力：为什么总是执行主机代码，尽管有 __CUDA_ARCH__
【发布时间】：2016-10-07 08:51:58
【问题描述】：

我尝试在代码中定义两个分支：一个用于 CUDA 执行，另一个 - 没有它（考虑到未来的 OMP）。但是当我使用宏__CUDA_ARCH__ 时，它看起来好像总是在执行主机代码。但我认为 Thrust 默认使用 CUDA（以及设备代码的分支）。我的代码有什么问题？这里是：

#include <thrust/transform.h>                                 
#include <thrust/functional.h>                                
#include <thrust/iterator/counting_iterator.h>                
#include <stdio.h>                                            

struct my_op                                                  
{                                                             
    my_op(int init_const) : constanta(init_const) {}      
    __host__ __device__ int operator()(const int &x) const
    {                                                     
        #if defined(__CUDA_ARCH__)                    
            return 2 * x * constanta;    // never executed - why?
        #else                                     
            return x * constanta;        // always executed                 
        #endif                       
    }                                                     

private:                                                      
    int constanta;                                        
};                                                            

int main()                                                    
{                                                             
 int data[7] = { 0, 0, 0, 0, 0, 0, 0 };                        
 thrust::counting_iterator<int> first(10);                     
 thrust::counting_iterator<int> last = first + 7;              

 int init_value = 1;                                           
 my_op op(init_value);                                         

 thrust::transform(first, last, data, op);                     
 for each (int el in data)                                     
    std::cout << el << " ";                               

 std::cout << std::endl;                                       
}

我希望“变换”将向量定义为乘以 2*constanta，但我看到使用了主机代码 - 输出是“10 11 12 13 14 15 16”，而不是“20 22 24 26 28 30 32” （如预期的那样）。

为什么？

【问题讨论】：

标签： cuda thrust

【解决方案1】：

推力正在选择主机路径，因为提供给推力变换操作的数据项之一在主机内存中：

 thrust::transform(first, last, data, op); 
                                ^^^^

如果您想让推力算法在设备上运行，一般来说，您传入/传出的所有容器数据也必须驻留在设备内存中。

以下是对代码的修改，表明如果我们将data 替换为设备驻留容器，推力将遵循设备路径：

$ cat t13.cu
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/device_vector.h>
#include <stdio.h>

struct my_op
{
    my_op(int init_const) : constanta(init_const) {}
    __host__ __device__ int operator()(const int &x) const
    {
        #if defined(__CUDA_ARCH__)
            return 2 * x * constanta;    // never executed - why?
        #else
            return x * constanta;        // always executed
        #endif
    }

private:
    int constanta;
};

int main()
{
// int data[7] = { 0, 0, 0, 0, 0, 0, 0 };
 thrust::counting_iterator<int> first(10);
 thrust::counting_iterator<int> last = first + 7;
 thrust::device_vector<int> d_data(7);

 int init_value = 1;
 my_op op(init_value);

 thrust::transform(first, last, d_data.begin(), op);
 for (int el = 0; el < 7; el++) {
    int dat = d_data[el];
    std::cout << dat  << " ";    }

 std::cout << std::endl;
}
$ nvcc -arch=sm_61 -o t13 t13.cu
$ ./t13
20 22 24 26 28 30 32
$

您可能想阅读thrust quick start guide 以了解推力算法调度。

【讨论】：