c++ - 可在 CPU 或 GPU 中灵活运行的 CUDA 推力仿函数

Question

这可能是一个愚蠢的问题，但我似乎无法找到任何与之相关的资源，因此非常感谢您的意见。

假设我有一些仿函数

struct AddOne {
    thrust::device_ptr<int> numbers;

    __device__
    void operator()(int i) {
        numbers[i] = numbers[i] + 1;
    }
}

我可以调用

AddOne addOneFunctor;
thrust::device_vector<int> idx(100), numbers(100);
addOneFunctor.numbers = numbers.data();
thrust::sequence(idx.begin(), idx.end(), 0);
thrust::for_each(thrust::device, idx.begin(), idx.end(), addOneFunctor);

是否可以编写上述代码，以便可以在编译时或理想情况下在运行时更改执行策略？

例如将结构更改为

struct AddOne {
    thrust::pointer<int> numbers;

     __host__ __device__
     void operator()(int i) {
         numbers[i] = numbers[i] + 1;
     }
}

所以它可以像这样运行

AddOne addOneFunctor;
std::vector<int> idx(100), numbers(100);
addOneFunctor.numbers = numbers.data();
thrust::sequence(idx.begin(), idx.end(), 0);
thrust::for_each(thrust::cpp::par, idx.begin(), idx.end(), addOneFunctor);

底线是:我想要一个单一的代码库，我可以在其中决定是使用 thrust::device_vector 还是某种宿主 vector (例如 std::vectors) 并在 GPU(使用 thrust::device 执行策略)或 CPU(使用 thrust::cpp::par 或类似政策)。

PS:我现在想避免使用 PGI。

Answer 1

是的，有可能，几乎与您描述的完全一样。

这是一个完整的示例:

$ cat t1205.cu
#include <thrust/execution_policy.h>
#include <thrust/for_each.h>
#include <thrust/device_vector.h>
#include <thrust/sequence.h>
#include <iostream>
#include <vector>

struct AddOne {
    int *numbers;
     template <typename T>
     __host__ __device__
     void operator()(T &i) {
         numbers[i] = numbers[i] + 1;
     }
};

int main(){

  AddOne addOneFunctor;
  std::vector<int> idx(100), numbers(100);
  addOneFunctor.numbers = thrust::raw_pointer_cast(numbers.data());
  thrust::sequence(idx.begin(), idx.end(), 0);
  thrust::for_each(thrust::cpp::par, idx.begin(), idx.end(), addOneFunctor);
  for (int i = 0; i < 5; i++)
    std::cout << numbers[i] << ",";
  std::cout << std::endl;

  thrust::device_vector<int> didx(100), dnumbers(100);
  addOneFunctor.numbers = thrust::raw_pointer_cast(dnumbers.data());
  thrust::sequence(didx.begin(), didx.end(), 0);
  thrust::for_each(thrust::device, didx.begin(), didx.end(), addOneFunctor);
  for (int i = 0; i < 5; i++)
    std::cout << dnumbers[i] << ",";
  std::cout << std::endl;
}
$ nvcc -o t1205 t1205.cu
$ ./t1205
1,1,1,1,1,
1,1,1,1,1,
$

请注意，该算法是 thrust::sequence 而不是 thrust::seq。

使用 CUDA 8RC

作为@m.s.指出，没有必要在上述代码的算法上明确使用执行策略——您可以删除它们，它的工作方式相同。然而，执行策略的正式使用允许将上面的示例扩展到您不使用容器，而是使用普通主机和设备数据的情况，因此它可能仍然具有一些值(value)，具体取决于您的总体目标。