concurrency - 同時カーネル実行を理解するcuda

Question

カーネルの同時実行がどのように機能するかを理解しようとしています。私はそれを理解しようとする簡単なプログラムを書きました。カーネルは、2 つのストリームを使用して 2D 配列を生成します。ストリームが 1 つあり、同時実行がない場合、正しい結果が得られます。2 つのストリームで試してみると、同時実行を試みると、間違った結果が得られます。これが正しいかどうか、またはカーネルのセットアップ方法がよくわからないため、メモリ転送に関係していると思います。プログラミングガイドでは、十分に説明されていません。私の目的のために、Matlab がカーネルを呼び出す必要があります。

私が理解しているように、メインプログラムは次のようになります。

ホスト上の固定メモリを割り当てます
1 つのストリームに必要な GPU のメモリを割り当てます (2 つのストリーム = ホストの合計メモリの半分)。
ストリームを作成する
ストリームをループする
cudaMemcpyAsync() を使用して、単一ストリームのメモリをホストからデバイスにコピーします
ストリームのカーネルを実行します
ストリームのメモリをホストにコピーします cudaMemcpyAsync()
- 各ストリームのデータのサイズとストリーム番号に基づくオフセットを使用して、各ストリームに必要な場所からメモリを参照することで、正しいことをしていると思います。
ストリームを破壊する
メモリを解放する

これが私が使用しようとしているコードです。

並行カーネル.cpp

__global__ void concurrentKernel(int const width, 
                                  int const streamIdx,
                                  double *array)
 {
     int thread = (blockIdx.x * blockDim.x) + threadIdx.x;;

     for (int i = 0; i < width; i ++)
     {
        array[thread*width+i] = thread+i*width+1;
//         array[thread*width+i+streamIdx] = thread+i*width+streamIdx*width/2;
     }

 }

並行MexFunction.cu

#include <stdio.h>
#include <math.h>
#include "mex.h"

/* Kernel function */
#include "concurrentKernel.cpp"


void mexFunction(int        nlhs,
                 mxArray    *plhs[],
                 int        nrhs,
                 mxArray    *prhs[])
{

    int const numberOfStreams = 2; // set number of streams to use here.
    cudaError_t cudaError;
    int offset;

    int width, height, fullSize, streamSize;
    width = 512;
    height = 512;
    fullSize = height*width;
    streamSize = (int)(fullSize/numberOfStreams);
    mexPrintf("fullSize: %d, streamSize: %d\n",fullSize, streamSize);

    /* Return the populated array */
    double *returnedArray;
    plhs[0] = mxCreateDoubleMatrix(height, width, mxREAL);
    returnedArray = mxGetPr(plhs[0]);

    cudaStream_t stream[numberOfStreams];
    for (int i = 0; i < numberOfStreams; i++)
    {
        cudaStreamCreate(&stream[i]);    
    }

    /* host memory */
    double *hostArray;
    cudaError = cudaMallocHost(&hostArray,sizeof(double)*fullSize);    // full size of array.
    if (cudaError != cudaSuccess) {mexPrintf("hostArray memory allocation failed\n********** Error: %s **********\n",cudaGetErrorString(cudaError)); return; }

    for (int i = 0; i < height; i++)
    {
        for (int j = 0; j < width; j++)
        {
            hostArray[i*width+j] = -1.0;
        }
    }

    /* device memory */
    double *deviceArray;
    cudaError = cudaMalloc( (void **)&deviceArray,sizeof(double)*streamSize);    // size of array for each stream.
    if (cudaError != cudaSuccess) {mexPrintf("deviceArray memory allocation failed\n********** Error: %s **********\n",cudaGetErrorString(cudaError)); return; }


    for (int i = 0; i < numberOfStreams; i++)
    {
        offset = i;//*streamSize;
        mexPrintf("offset: %d, element: %d\n",offset*sizeof(double),offset);

        cudaMemcpyAsync(deviceArray, hostArray+offset, sizeof(double)*streamSize, cudaMemcpyHostToDevice, stream[i]);
        if (cudaError != cudaSuccess) {mexPrintf("deviceArray memory allocation failed\n********** Error: %s **********\n",cudaGetErrorString(cudaError)); return; }

        concurrentKernel<<<1, 512, 0, stream[i]>>>(width, i, deviceArray);

        cudaMemcpyAsync(returnedArray+offset, deviceArray, sizeof(double)*streamSize, cudaMemcpyDeviceToHost, stream[i]);
        if (cudaError != cudaSuccess) {mexPrintf("returnedArray memory allocation failed\n********** Error: %s **********\n",cudaGetErrorString(cudaError)); return; }

        mexPrintf("returnedArray[offset]: %g, [end]: %g\n",returnedArray[offset/sizeof(double)],returnedArray[(i+1)*streamSize-1]);
    }


    for (int i = 0; i < numberOfStreams; i++)
    {
        cudaStreamDestroy(stream[i]);    
    }

    cudaFree(hostArray);
    cudaFree(deviceArray);

}

ストリームが 2 つある場合、結果はゼロの配列になり、メモリに問題があると思います。誰が私が間違っているのか説明できますか? Matlab からこれらをコンパイルして実行するのに助けが必要な場合は、そのためのコマンドを提供できます。

アップデート：

for (int i = 0; i < numberOfStreams; i++)
{
    offset = i*streamSize;
    mexPrintf("offset: %d, element: %d\n",offset*sizeof(double),offset);

    cudaMemcpyAsync(deviceArray, hostArray+offset, sizeof(double)*streamSize, cudaMemcpyHostToDevice, stream[i]);
    if (cudaError != cudaSuccess) {mexPrintf("deviceArray memory allocation failed\n********** Error: %s **********\n",cudaGetErrorString(cudaError)); return; }

    concurrentKernel<<<1, 512, 0, stream[i]>>>(width, i, deviceArray);


}
cudaDeviceSynchronize();


for (int i = 0; i < numberOfStreams; i++)
{
    offset = i*streamSize;
    mexPrintf("offset: %d, element: %d\n",offset*sizeof(double),offset);

    cudaMemcpyAsync(returnedArray+offset, deviceArray, sizeof(double)*streamSize, cudaMemcpyDeviceToHost, stream[i]);
    if (cudaError != cudaSuccess) {mexPrintf("returnedArray memory allocation failed\n********** Error: %s **********\n",cudaGetErrorString(cudaError)); return; }

    mexPrintf("returnedArray[offset]: %g, [end]: %g\n",returnedArray[offset/sizeof(double)],returnedArray[(i+1)*streamSize-1]);

    cudaStreamDestroy(stream[i]);    
}

score 6 · Accepted Answer

ストリームで使用している API は完全に非同期であるため、呼び出し元のホストスレッドに制御がすぐに返されることに注意する必要があります。非同期操作を実行している GPU とホストの間にある種の同期ポイントを挿入しない場合、ストリームでエンキューした操作が実際に終了するという保証はありません。あなたの例では、次のようなものが必要であることを意味します：

for (int i = 0; i < numberOfStreams; i++) 
{ 
    offset = i;//*streamSize; 
    mexPrintf("offset: %d, element: %d\n",offset*sizeof(double),offset); 

    cudaMemcpyAsync(deviceArray, hostArray+offset, sizeof(double)*streamSize, 
                    cudaMemcpyHostToDevice, stream[i]); 

    concurrentKernel<<<1, 512, 0, stream[i]>>>(width, i, deviceArray); 

    cudaMemcpyAsync(returnedArray+offset, deviceArray, sizeof(double)*streamSize,
                    cudaMemcpyDeviceToHost, stream[i]); 
} 

// Host thread waits here until both kernels and copies are finished
cudaDeviceSynchronize();

for (int i = 0; i < numberOfStreams; i++) 
{ 
    mexPrintf("returnedArray[offset]: %g, [end]: %g\n",returnedArray[offset/sizeof(double)],returnedArray[(i+1)*streamSize-1]); 
    cudaStreamDestroy(stream[i]);     
}

ここで重要なのは、ホストメモリで結果を検査する前に、両方のメモリ転送が完了していることを確認する必要があるということです。元のコードも更新もこれを行いません。

score 1 · Accepted Answer

また、異なる同時ストリームに対して deviceArray ポインターを再利用しているようです。現在のコードがそのまま機能する場合、@Tom が言及している誤った依存関係が原因で、ハードウェアがストリームを順番に実行する可能性があります。ストリームごとに個別の deviceArray が必要です。

/* device memory */
double *deviceArray[numberOfStreams];
for (int i = 0; i < numberOfStreams; i++)
{
    cudaError = cudaMalloc( (void **)&deviceArray[i],sizeof(double)*streamSize);    // size of array for each stream.
    if (cudaError != cudaSuccess) {mexPrintf("deviceArray memory allocation failed\n********** Error: %s **********\n",cudaGetErrorString(cudaError)); return; }
}

for (int i = 0; i < numberOfStreams; i++)
{
    offset = i;//*streamSize;
    mexPrintf("offset: %d, element: %d\n",offset*sizeof(double),offset);

    cudaMemcpyAsync(deviceArray[i], hostArray+offset, sizeof(double)*streamSize, cudaMemcpyHostToDevice, stream[i]);
    if (cudaError != cudaSuccess) {mexPrintf("deviceArray memory allocation failed\n********** Error: %s **********\n",cudaGetErrorString(cudaError)); return; }

    concurrentKernel<<<1, 512, 0, stream[i]>>>(width, i, deviceArray[i]); 

    cudaMemcpyAsync(returnedArray+offset, deviceArray[i], sizeof(double)*streamSize,
                    cudaMemcpyDeviceToHost, stream[i]);     
}

concurrency - 同時カーネル実行を理解するcuda

2 に答える 2

Related

Reference