请注意，我绝对是CUDA的初学者，下面的所有内容都是未经测试的伪代码。我来自JavaScript，而我的C++也非常使用rust ，因此对我的无知表示歉意:)
我正在尝试使用CUDA对许多不同的外汇策略进行回测。
使用Thrust，我从一个类(伪代码)实例化了1000个对象:
#include <stdio.h>
#include <thrust/device_ptr.h>
#include <thrust/device_new.h>
#define N 1000
typedef struct dataPoint {
...
} dataPoint;
class Strategy {
public:
__device__ __host__ Strategy(...) {
...
}
__device__ __host__ void backtest(dataPoint data) {
...
}
};
int main() {
dataPoint data[100000];
thrust::device_ptr<Strategy> strategies[1000];
int i;
// Instantiate 5000 strategies.
for (i=0; i<1000; i++) {
thrust::device_ptr<Strategy> strategies[i] = thrust::device_new<Strategy>(...);
}
// Iterate over all 100000 data points.
for (i=0; i<100000; i++) {
// Somehow run .backtest(data[j]) on each strategy here.
// i.e. Run backtest() in parallel for all 1000
// strategy objects here.
}
}
现在，我们想对
.backtest()中每个项目的每个对象运行data方法。程序上，我将执行以下操作:// Iterate over all 100000 data points.
for (j=0; j<100000; j++) {
// Iterate over all 1000 strategies.
for (i=0; i<1000; i++) {
strategies[i].backtest(data[j]);
}
}
我如何使用CUDA做到这一点，以便
.backtest()对所有策略的所有策略并行运行，每次迭代j通过数据？如果我必须完全重新整理所有内容，就可以了-我愿意接受任何需要的东西。如果类无法做到这一点，那就去吧。
最佳答案
典型的推力代码经常使用某些C++习惯用法(例如函子)，因此，如果您的C++使用rust ，则可能需要阅读有关C++函子的信息。您可能还需要查看thrust quick start guide，以讨论函子以及我们现在将使用的奇特迭代器。
通常，至少从表达的 Angular 来看，我认为推力非常适合您的问题描述。考虑到推力表达对于这些类型问题的灵活性，可能有很多方法可以给猫蒙皮。我将尝试介绍一些与您的伪代码“接近”的东西。但是毫无疑问，有很多方法可以实现这一点。
首先，在推力方面，我们通常尝试避免for循环。这将非常慢，因为它们通常在每次迭代时都涉及主机代码和设备代码的交互(例如，在每次迭代时在后台调用CUDA内核)。如果可能的话，我们更喜欢使用推力算法，因为这些算法通常会在引擎盖下“转换”为一个或少量CUDA内核。
推力中最基本的算法之一是transform。它具有多种风格，但基本上是接受输入数据并对每个元素进行任意操作。
使用基本推力变换操作，我们可以初始化您的数据以及您的策略，而无需求助于for循环。我们将为每种类型的对象(dataPoint，Strategy)构造一个适当长度的设备 vector ，然后使用thrust::transform初始化每个 vector 。
剩下的任务是针对每个dataPoint执行每个Strategy。理想情况下，我们也希望并行执行此操作。不仅针对您建议的每个for-loop迭代，还针对每个Strategy和每个dataPoint，全部“一次”(即在单个算法调用中)。
实际上，我们可以想到一个矩阵，一个轴由dataPoint(在您的示例中为100000)组成，另一轴由Strategy(在您的示例中为1000)组成。对于此矩阵中的每个点，我们都可以设想它针对该Strategy保留了dataPoint的应用结果。
在推力方面，我们通常更喜欢将2D概念实现为一维。因此，我们的结果空间等于dataPoint数量乘以Strategy数量的乘积。我们将创建此大小的result device_vector(在您的示例中为100000 * 1000)以保存结果。
为了演示起见，由于您对要执行的算术类型的指导很少，因此我们假定以下内容:
Strategy应用dataPoint的结果是float。 dataPoint是一个结构，由int(dtype-在此示例中忽略)和float(dval)组成。对于dval，dataPoint(i)将包含1.0f + i*10.0f。 Strategy由multiplier和adder组成，将按以下方式使用:Strategy(i) = multiplier(i) * dval + adder(i);
Strategy应用于dataPoint包括以下步骤:检索与dval关联的dataPoint，并将其代入上述项目3给出的方程式中。该公式在backtest类的Strategy方法中捕获。 backtest方法将dataPoint类型的对象作为其参数，从中将检索适当的dval。 我们还需要介绍一些其他概念。 2D结果矩阵的1D实现将需要我们提供适当的索引方法，以便在2D矩阵的每个点上，根据线性尺寸，我们可以确定在该位置使用哪个
Strategy和哪个dataPoint来计算result。点。在推力方面，我们可以结合使用复杂的迭代器来执行此操作。简而言之，从“由内而外”，我们将从转换迭代器开始，该迭代器采用索引映射函子和
thrust::counting_iterator提供的线性序列，以便为每个索引(每个矩阵维)创建一个映射。每个映射函子中的算法会将result的线性索引转换为矩阵行和列的适当重复索引。给定此转换迭代器以创建重复的行或列索引，我们将该索引传递给置换迭代器，该置换迭代器为指示的每个行/列选择适当的dataPoint或Strategy。然后将这两个项目(dataPoint，Strategy)压缩为zip_iterator。然后将zip_iterator传递给run_strat函子，该函子实际上计算应用于给定Strategy的给定dataPoint。以下是概述上述概念的示例代码:
#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/transform.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <math.h>
#define TOL 0.00001f
// number of strategies
#define N 1000
// number of data points
#define DSIZE 100000
// could use int instead of size_t here, for these problem dimensions
typedef size_t idx_t;
struct dataPoint {
int dtype;
float dval;
};
class Strategy {
float multiplier;
float adder;
idx_t id;
public:
__device__ __host__ Strategy(){
id = 0;
multiplier = 0.0f;
adder = 0.0f;
}
__device__ __host__ Strategy(idx_t _id) {
id = _id;
multiplier = 1.0f + ((float)id)/(float)N;
adder = (float)id;
}
__device__ __host__ float backtest(dataPoint data) {
return multiplier*data.dval+adder;
}
};
// functor to initialize dataPoint
struct data_init
{
__host__ __device__
dataPoint operator()(idx_t id){
dataPoint temp;
temp.dtype = id;
temp.dval = 1.0f + id * 10.0f;
return temp;
}
};
// functor to initialize Strategy
struct strat_init
{
__host__ __device__
Strategy operator()(idx_t id){
Strategy temp(id);
return temp;
}
};
// functor to "test" a Strategy against a dataPoint, using backtest method
struct run_strat
{
template <typename T>
__host__ __device__
float operator()(idx_t id, T t){
return (thrust::get<0>(t)).backtest(thrust::get<1>(t));
}
};
// mapping functor to generate "row" (Strategy) index from linear index
struct strat_mapper : public thrust::unary_function<idx_t, idx_t>
{
__host__ __device__
idx_t operator()(idx_t id){
return id/DSIZE;
}
};
// mapping functor to generate "column" (dataPoint) index from linear index
struct data_mapper : public thrust::unary_function<idx_t, idx_t>
{
__host__ __device__
idx_t operator()(idx_t id){
return id%DSIZE;
}
};
int main() {
// initialize data
thrust::device_vector<dataPoint> data(DSIZE);
thrust::transform(thrust::counting_iterator<idx_t>(0), thrust::counting_iterator<idx_t>(DSIZE), data.begin(), data_init());
// initialize strategies
thrust::device_vector<Strategy> strategies(N);
thrust::transform(thrust::counting_iterator<idx_t>(0), thrust::counting_iterator<idx_t>(N), strategies.begin(), strat_init());
// test each data point against each strategy
// Somehow run .backtest(data[j]) on each strategy here.
// i.e. Run backtest() in parallel for all 1000
// strategy objects here.
// allocate space for results for each datapoint against each strategy
thrust::device_vector<float> result(DSIZE*N);
thrust::transform(thrust::counting_iterator<idx_t>(0), thrust::counting_iterator<idx_t>(DSIZE*N), thrust::make_zip_iterator(thrust::make_tuple(thrust::make_permutation_iterator(strategies.begin(), thrust::make_transform_iterator(thrust::counting_iterator<idx_t>(0), strat_mapper())), thrust::make_permutation_iterator(data.begin(), thrust::make_transform_iterator(thrust::counting_iterator<idx_t>(0), data_mapper())))), result.begin(), run_strat());
// validation
// this would have to be changed if you change initialization of dataPoint
// or Strategy
thrust::host_vector<float> h_result = result;
for (int j = 0; j < N; j++){
float m = 1.0f + (float)j/(float)N;
float a = j;
for (int i = 0; i < DSIZE; i++){
float d = 1.0f + i*10.0f;
if (fabsf(h_result[j*DSIZE+i] - (m*d+a))/(m*d+a) > TOL) {std::cout << "mismatch at: " << i << "," << j << " was: " << h_result[j*DSIZE+i] << " should be: " << m*d+a << std::endl; return 1;}}}
return 0;
}
笔记:
如上所述的
transform操作使用counting_iterator作为第一个参数(和第二个参数)，但这实际上被忽略并且“虚拟”使用，只是为了适当地确定问题的大小。可以通过更简单的实现将其消除，但是在我看来，最简单的方法(无需进一步使代码困惑)将是使用C++ 11 auto定义zip_iterator，然后将其自身传递，以及一个偏移版本其中，对于thrust::transform，使用仅使用一个输入 vector 而不是2的版本。我认为这不会在性能上产生太大差异，我觉得这更容易解析，但也许没有。 关于c++ - 如何在CUDA中为多个Thrust对象成员函数调用内核函数？，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/35998223/