用CUDA计算递归数组
用CUDA计算递归数组
提问于 2015-04-07 22:50:55
我有一个递归数组定义。甭管他们
A(x, y + 1) = f(A(x - 1, y), A(x, y), A(x + 1, y))第一层初始化
A(x, 0) = g(x)我想用CUDA逐层计算这样的数组。问题是，做这些事情的首选方式是什么。单个内核是否应该在每一步的同步中为A(tid, y)的y计算数组[1, height)？或者它应该只计算单点，但应该被多次调用？还是把一个问题分解成更大的独立部分更好？例如，如果完成了菱形的前一层操作，则该数组可以被菱形拆分，这样就可以独立计算整个菱形(在菱形内没有同步)。
如果层是2D而不是一维，事情会变得不同吗？
我计划每秒钟计算一个宽度~ 10000 (可能更少)和高度44100的数组。问题实际上是3D (200x50x44100) (如果它重要的话)。我只是为了简单起见把它设计成2D的。
回答 1
Stack Overflow用户
回答已采纳
发布于 2015-04-09 02:29:57
一种简单的方法可能只是从这里概述的内容开始：
“单个内核是否应该在[1，高度]中为y计算一个数组A(tid，y)，在每一步上同步？”
这应该很容易实现。
大约有10,000的x“宽度”，以使GPU能够合理地处理这么多线程。
对于一个复杂的f()函数，能够每秒进行44100次迭代(平均迭代时间为22 us)可能是一项挑战。然而，对于一个相当简单的f()函数，根据下面的快速测试，它似乎是可能的。我们受益于这样的事实:通过像这样迭代地启动内核，内核的启动开销大部分是隐藏的。
下面是用proof编写的示例代码，以演示概念的证明：
$ cat t708.cu
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>
#include <thrust/for_each.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/copy.h>
#include <stdlib.h>
#include <iostream>
#define DSIZE 10000
#define YSIZE 2
#define NUM_ITER 44100
#define AVG_SIZE 3
#define DISP_WIDTH 5
struct f
{
template <typename T>
__host__ __device__
void operator()(T t) {
thrust::get<AVG_SIZE>(t) = thrust::get<0>(t);
thrust::get<AVG_SIZE>(t) += thrust::get<1>(t);
thrust::get<AVG_SIZE>(t) += thrust::get<2>(t);
thrust::get<AVG_SIZE>(t) /= AVG_SIZE;}
};
int main(){
thrust::host_vector<float> h_A(DSIZE);
for (int i =0; i < DSIZE; i++) h_A[i] = rand()/(float)RAND_MAX; // A(x, 0) = g(x)
thrust::device_vector<float> d_A[YSIZE];
d_A[0].resize(h_A.size());
d_A[1].resize(h_A.size());
thrust::copy(h_A.begin(), h_A.end(), d_A[0].begin());
thrust::copy(h_A.begin(), h_A.end(), d_A[1].begin());
std::cout << "input left end: " << std::endl;
thrust::copy(d_A[0].begin(), d_A[0].begin()+DISP_WIDTH, std::ostream_iterator<float>(std::cout, ","));
std::cout << std::endl << "input right end: " << std::endl;
thrust::copy(d_A[0].end() - DISP_WIDTH, d_A[0].end(), std::ostream_iterator<float>(std::cout, ","));
std::cout << std::endl;
cudaEvent_t start, stop;
cudaEventCreate(&start); cudaEventCreate(&stop);
int cur = 0;
int nxt = 1;
cudaEventRecord(start, 0);
for (int i = 0; i < NUM_ITER; i++){
thrust::for_each(thrust::make_zip_iterator(thrust::make_tuple(d_A[cur].begin(), d_A[cur].begin()+1, d_A[cur].begin()+2, d_A[nxt].begin()+1)), thrust::make_zip_iterator(thrust::make_tuple(d_A[cur].end()-2, d_A[cur].end()-1, d_A[cur].end(), d_A[nxt].end()-1)), f());
cur = (cur==0) ? 1:0; // modify for a full storage in y
nxt = (nxt==0) ? 1:0;}
cudaDeviceSynchronize();
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
float et;
cudaEventElapsedTime(&et, start, stop);
std::cout << "elapsed time: " << et << "ms" << std::endl << "output left end: " << std::endl;
thrust::copy(d_A[cur].begin(), d_A[cur].begin()+DISP_WIDTH, std::ostream_iterator<float>(std::cout, ","));
std::cout << std::endl << "output right end: " << std::endl;
thrust::copy(d_A[cur].end() - DISP_WIDTH, d_A[cur].end(), std::ostream_iterator<float>(std::cout, ","));
std::cout << std::endl;
return 0;
}
$ nvcc -O3 -o t708 t708.cu
$ ./t708
input left end:
0.840188,0.394383,0.783099,0.79844,0.911647,
input right end:
0.865333,0.828169,0.311025,0.373209,0.888766,
elapsed time: 368.337ms
output left end:
0.840188,0.838681,0.837174,0.835667,0.83416,
output right end:
0.881355,0.883207,0.88506,0.886913,0.888766,
$备注：
- ~370 8us的44100次迭代的总体执行时间表明，平均循环时间约为8us。
- 这是运行在一个Fedora20系统，与CUDA 7和一个Quadro5000图形处理器。
- 我包括了输入和输出数据集的左侧和右侧的显示，以验证“平均”
f()函数(在本例中为函子)。这实际上是一种“放松”形式，因此我们期望数据集的左侧收敛到左侧的值(这不会改变)，而数据集的右侧将收敛到右侧的值，两者之间有一条近似直线。 - 我只保留y(n)和y(n-1)数据集，用于计算和计时。如果您需要所有的y(0..44100)，这是一个相当容易的代码修改。
- 如果您不熟悉推力，快速启动指南可能会帮助您达到速度，并且您会发现这个问题覆盖了一些类似的算法，它应该让您了解如何将该推力方法转换为等效的库达内核方法。
- 在不太明显的情况下，“每个步骤上的同步”都是由cuda内核调用完成的，这是对
thrust::for_each调用中隐含的一种设备范围的同步。
页面原文内容由Stack Overflow提供。腾讯云小微IT领域专用引擎提供翻译支持
原文链接：
https://stackoverflow.com/questions/29502757
复制相关文章
相似问题
腾讯云开发者
