CUDAフレームワークを使用してGPUで動作できるように、メモリの連続したチャンクに格納されているデータに対して計算やその他の操作を実行するライブラリを書き直しています。データは、4次元グリッド上に存在する情報を表します。グリッドの合計サイズは、数千から数百万のグリッドポイントの範囲になります。各方向に沿って、グリッドには8個から数百個のポイントが含まれる場合があります。私の質問は、グリッドのサブセットに操作を実装するための最良の方法は何かについてです。たとえば、グリッドが[0、nx)x [0、ny)x [0、nz)x [0、nq)であり、インデックスが[1に属するすべてのポイントを乗算する変換を実装するとします。 、nx-1)x [1、ny-1)x [1、nz-1)x [0、nq-1)マイナス1。
今、私がしているのはネストされたループです。これはコードの骨組みです
{
int nx,ny,nz,nq;
nx=10,ny=10,nz=10,nq=10;
typedef thrust::device_vector<double> Array;
Array A(nx*ny*nz*nq);
thrust::fill(A.begin(), A.end(), (double) 1);
for (auto q=1; q<nq-1; ++q){
for (auto k=1; k<nz-1; ++k){
for (auto j=1; j<ny-1; ++j){
int offset1=+1+j*nx+k*nx*ny+q*nx*ny*nz;
int offset2=offset1+nx-2;
thrust::transform(A.begin()+offset1,
A.begin()+offset2,
thrust::negate<double>());
}
}
}
}
ただし、これが最も効率的な方法であるかどうか疑問に思います。この場合、最大でnx-2スレッドしか同時に実行できないように思われるからです。したがって、おそらくより良い方法は、シーケンスイテレータを生成し(配列に沿って線形位置を返す)、zipイテレータを使用して配列に圧縮し、タプルの2番目の要素を調べるファンクタを定義することだと考えていました(位置の値)、その値が許容範囲内にある場合は、タプルの最初の要素を変更します。ただし、それを行うためのより良い方法があるかもしれません。私はCUDAを初めて使用しますが、さらに悪いことに、Fortranで歯を切ったので、forループボックスの外側を考えるのは難しいです...
最も効率的な方法がわかりません。私はあなたのスケルトンコードよりも効率的だと思うものを提案することができます。
本文中のあなたの提案は正しい方向に向かっています。潜在的にかなりの回数繰り返されるネストされたforループのセットを使用するのではなく、1回のスラスト呼び出しですべてを実行するように努める必要があります。ただし、1回の推力呼び出しで、操作対象の「立方体」ボリューム内のインデックスの配列値のみを変更する必要があります。
ただし、ご提案のとおり、生成されたインデックスを有効なインデックスボリュームに対してテストする方法は使用したくありません。これには、少量のグリッドのみを変更したい場合でも、アレイと同じ大きさのグリッドを起動する必要があります。
代わりに、変更する必要のある要素数をカバーするのに十分な大きさの操作を起動し、線形インデックス-> 4Dインデックス->調整された線形インデックス変換を行うファンクターを作成します。次に、そのファンクターは変換イテレーター内で動作して、0、1、2などで始まる通常の線形シーケンスを、変更するボリューム内で開始および留まるシーケンスに変換します。次に、この変更されたシーケンスで順列イテレータを使用して、変更する配列の値を選択します。
これは、64x64x64x64の配列と62x62x62x62の変更されたボリュームについて、ネストされたループメソッド(1)と私のもの(2)のタイミングの違いを示す例です。
$ cat t39.cu
#include <thrust/device_vector.h>
#include <thrust/transform.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/functional.h>
#include <thrust/equal.h>
#include <cassert>
#include <iostream>
struct my_idx
{
int nx, ny, nz, nq, lx, ly, lz, lq, dx, dy, dz, dq;
my_idx(int _nx, int _ny, int _nz, int _nq, int _lx, int _ly, int _lz, int _lq, int _hx, int _hy, int _hz, int _hq) {
nx = _nx;
ny = _ny;
nz = _nz;
nq = _nq;
lx = _lx;
ly = _ly;
lz = _lz;
lq = _lq;
dx = _hx - lx;
dy = _hy - ly;
dz = _hz - lz;
dq = _hq - lq;
// could do a lot of assert checking here
}
__host__ __device__
int operator()(int idx){
int rx = idx / dx;
int ix = idx - (rx * dx);
int ry = rx / dy;
int iy = rx - (ry * dy);
int rz = ry / dz;
int iz = ry - (rz * dz);
int rq = rz / dq;
int iq = rz - (rq * dq);
return (((iq+lq)*nz+iz+lz)*ny+iy+ly)*nx+ix+lx;
}
};
#include <time.h>
#include <sys/time.h>
#define USECPSEC 1000000ULL
unsigned long long dtime_usec(unsigned long long start){
timeval tv;
gettimeofday(&tv, 0);
return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start;
}
int main()
{
int nx,ny,nz,nq,lx,ly,lz,lq,hx,hy,hz,hq;
nx=64,ny=64,nz=64,nq=64;
lx=1,ly=1,lz=1,lq=1;
hx=nx-1,hy=ny-1,hz=nz-1,hq=nq-1;
thrust::device_vector<double> A(nx*ny*nz*nq);
thrust::device_vector<double> B(nx*ny*nz*nq);
thrust::fill(A.begin(), A.end(), (double) 1);
thrust::fill(B.begin(), B.end(), (double) 1);
// method 1
unsigned long long m1_time = dtime_usec(0);
for (auto q=lq; q<hq; ++q){
for (auto k=lz; k<hz; ++k){
for (auto j=ly; j<hy; ++j){
int offset1=lx+j*nx+k*nx*ny+q*nx*ny*nz;
int offset2=offset1+(hx-lx);
thrust::transform(A.begin()+offset1,
A.begin()+offset2, A.begin()+offset1,
thrust::negate<double>());
}
}
}
cudaDeviceSynchronize();
m1_time = dtime_usec(m1_time);
// method 2
unsigned long long m2_time = dtime_usec(0);
auto p = thrust::make_permutation_iterator(B.begin(), thrust::make_transform_iterator(thrust::counting_iterator<int>(0), my_idx(nx, ny, nz, nq, lx, ly, lz, lq, hx, hy, hz, hq)));
thrust::transform(p, p+(hx-lx)*(hy-ly)*(hz-lz)*(hq-lq), p, thrust::negate<double>());
cudaDeviceSynchronize();
m2_time = dtime_usec(m2_time);
if (thrust::equal(A.begin(), A.end(), B.begin()))
std::cout << "method 1 time: " << m1_time/(float)USECPSEC << "s method 2 time: " << m2_time/(float)USECPSEC << "s" << std::endl;
else
std::cout << "mismatch error" << std::endl;
}
$ nvcc -std=c++11 t39.cu -o t39
$ ./t39
method 1 time: 1.6005s method 2 time: 0.013182s
$
この記事はインターネットから収集されたものであり、転載の際にはソースを示してください。
侵害の場合は、連絡してください[email protected]
コメントを追加