cuda - CUDA Thrust を使用して複数の 1D 移動平均を並行して実行する

Question

私は何の能力もないプログラマーではありません。CUDA に興味があるので、少し読んでいます。Thrust を使用して移動平均を行う例に出くわしました。

このような例は、実行され、ほとんど正しく動作します。ただし、移動平均操作を 1 回しか実行しないという意味では簡単です。

352すべてが同じデータストリームで動作しているこれらの移動平均演算を並列に実行することについて、どのように言えばよいでしょうか? 私の考えでは、プログラムの流れは次のようになります。

データを生成し、1 つの CUDA コアに送信します。(既存のコードと同じですが、長さを1000or10000の代わりに考える30)
351それが入っているCUDAコアからGTX 465の他のすべてのCUDAコアにコピーします
各 CUDA コアに、平均するデータ項目の数を伝えます。( 4, 5, 6,..., 352, 353, 354)
各コアで平均を並行して実行するようにデバイスに指示します
各コアから結果を読み戻す

私はこのコードを取得します

// compute SMA using standard summation
simple_moving_average(data, w, averages);

すべてが実現しますが、Thrust にこれらの多くを並行して実行させるにはどうすればよいでしょうか?

ここでの私の関心は、株式データのようなものです。GOOG の価格を見ている場合は、すべてのコアを使用して GPU に入れ、そのままにしておきます。その後、データをロードせずに、各コアから結果を読み取るだけで、自由に多くの処理を行うことができます。注: すべてのコアで GOOG を使用したくない場合があります。いくつかのコアはGOOGかもしれませんし、他のシンボルは他のシンボルかもしれませんが、それについては後で説明します. 各コアに十分なスペースがある場合、グローバルメモリにストックデータが必要ないと考えているだけです。

CUDA と Thrust の場合、これはかなり簡単だと思いますか?

score 3 · Accepted Answer

arrayfire でこれを行う方法は次のとおりです。私はこのライブラリとはまったく関係がないことに注意してください。
これは推力でも実行できると確信していますが、これは arrayfire を使用するとはるかに簡単であることがわかりました。また、ライブラリが無料である場合、なぜそれを推力の代わりに使用できないのでしょうか?

arrayfire では、matrix を使用して複数の SMA 操作を並行して実行できます。

unsigned n_SMAs = 1000;   // # of SMA indicators to evaluate 
unsigned len = 2000;      // # of stock prices per indicator
unsigned w = 6; // window size

// generate stock prices: [0..10] 
af::array data = af::randu(n_SMAs, len) * 10;

// compute inclusive prefix sums along colums of the matrix
af::array s = af::accum(data, 1);

// compute the average
af::array avg = (s.cols(w, af::end) - s.cols(0, af::end - w)) / w;
af::eval(avg);

std::cout << avg.dims() << "\n" << avg << "\n";

それがあなたが探しているものかどうか教えてください。これは私があなたの質問をどのように理解したかです: 複数の SMA 指標を並行して計算する

score 1 · Accepted Answer

あなたは次の 2 つの状況に関心があると理解しています。

長い一連のアイテムがあり、さまざまな数のアイテムを平均することによって、つまり、移動平均ウィンドウにさまざまな長さを使用して、特定の数の平均を計算したいと考えています。これは、元の質問から私が理解していることです。
メモリに連続して保存された一連のシーケンスがあり、 size の固定平均化ウィンドウでそれらを並列に平均化したいとします2 * RADIUS + 1。これは @asm によって提案された ArrayFire コードが行うことです - あなたはそれを受け入れました。

CUDA Thrust を使用する代わりに、独自の CUDA カーネルを作成して上記の操作を行う方が簡単だと思います。以下は、@asm によって提案された ArrayFire コードと同じように動作する完全に動作する例で、ケース #2 をカバーしています。ケース＃1をカバーするように変更するのは簡単です。

#include <thrust/device_vector.h>

#define RADIUS        3
#define BLOCK_SIZE_X  8
#define BLOCK_SIZE_Y  8

/*******************/
/* iDivUp FUNCTION */
/*******************/
int iDivUp(int a, int b){ return ((a % b) != 0) ? (a / b + 1) : (a / b); }

/********************/
/* CUDA ERROR CHECK */
/********************/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

/**********/
/* KERNEL */
/**********/
__global__ void moving_average(unsigned int *in, unsigned int *out, unsigned int M, unsigned int N) {

    __shared__ unsigned int temp[BLOCK_SIZE_Y][BLOCK_SIZE_X + 2 * RADIUS];

    unsigned int gindexx = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int gindexy = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int gindex  = gindexy * N + gindexx;

    unsigned int lindexx = threadIdx.x + RADIUS;
    unsigned int lindexy = threadIdx.y;

    // --- Read input elements into shared memory
    temp[lindexy][lindexx] = ((gindexx < N)&&(gindexy < M))? in[gindex] : 0;
    if (threadIdx.x < RADIUS) {
        temp[lindexy][threadIdx.x] = ((gindexx >= RADIUS)&&(gindexx < (N + RADIUS))&&(gindexy < M)) ? in[gindex - RADIUS] : 0;
        temp[lindexy][threadIdx.x + (RADIUS + min(BLOCK_SIZE_X, N - blockIdx.x * BLOCK_SIZE_X))] = (((gindexx + min(BLOCK_SIZE_X, N - blockIdx.x * BLOCK_SIZE_X)) < N)&&(gindexy < M))? in[gindexy * N + gindexx + min(BLOCK_SIZE_X, N - blockIdx.x * BLOCK_SIZE_X)] : 0;
        if ((threadIdx.y == 0)&&(gindexy < M)&&((gindexx + BLOCK_SIZE_X) < N)&&(gindexy < M)) printf("Inside 2 - tidx = %i; bidx = %i; tidy = %i; bidy = %i; lindexx = %i; temp = %i\n", threadIdx.x, blockIdx.x, threadIdx.y, blockIdx.y, threadIdx.x + (RADIUS + BLOCK_SIZE_X), temp[lindexy][threadIdx.x + (RADIUS + BLOCK_SIZE_X)]);
    }

    __syncthreads();

    // --- Apply the stencil
    unsigned int result = 0;
    for (int offset = -RADIUS ; offset <= RADIUS ; offset++) {
        result += temp[lindexy][lindexx + offset];
    }

    // --- Store the result
    out[gindexy * N + gindexx] = result;
}

/********/
/* MAIN */
/********/
int main() {

    const unsigned int M        = 2;
    const unsigned int N        = 4 + 2 * RADIUS;

    const unsigned int constant = 3;

    thrust::device_vector<unsigned int> d_in(M * N, constant);
    thrust::device_vector<unsigned int> d_out(M * N);

    dim3 GridSize(iDivUp(N, BLOCK_SIZE_X), iDivUp(M, BLOCK_SIZE_Y));
    dim3 BlockSize(BLOCK_SIZE_X, BLOCK_SIZE_Y);
    moving_average<<<GridSize, BlockSize>>>(thrust::raw_pointer_cast(d_in.data()), thrust::raw_pointer_cast(d_out.data()), M, N);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    thrust::host_vector<unsigned int> h_out = d_out;

    for (int j=0; j<M; j++) {
        for (int i=0; i<N; i++)
            printf("Element j = %i; i = %i; h_out = %i\n", j, i, h_out[N*j+i]);
    }

    return 0;

}

cuda - CUDA Thrust を使用して複数の 1D 移動平均を並行して実行する

2 に答える 2

Related

Reference