1

Visual Studio 2010 Professional 環境の Windows 7 64 ビットで CUDA 4.2 を実行しています。

まず、次のコードを実行しています。

// include the header files
#include <iostream>
#include <stdio.h>
#include <time.h>

#include "cuda.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

using namespace std; 

//kernel function
__global__ 
void dosomething(int *d_bPtr, int count, int* d_bStopPtr)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid==0)
       d_bStopPtr[tid]=0;
    else if(tid<count)
    {
       d_bPtr[tid]=tid;
// only if the arrary cell before it is 0, then change it to 0 too
        if (d_bStopPtr[tid-1]==0 )
           d_bStopPtr[tid]=0;

    }
}

int main()
{
    int count=100000;
// define the vectors
    thrust::host_vector <int> h_a(count);
    thrust::device_vector <int> d_b(count,0);
    int* d_bPtr=thrust::raw_pointer_cast(&d_b[0]);
    thrust::device_vector <int> d_bStop(count,1);
    int* d_bStopPtr=thrust::raw_pointer_cast(&d_bStop[0]);
// get the device property
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);

    int threadsPerBlock = prop.maxThreadsDim[0];
    int blocksPerGrid = min(prop.maxGridSize[0], (count + threadsPerBlock - 1) / threadsPerBlock);
//copy device to host
    thrust::copy(d_b.begin(),d_b.end(),h_a.begin());
    cout<<h_a[100]<<"\t"<<h_a[200]<<"\t"<<h_a[300]<<"\t"<<endl;
//run the kernel
    while(d_bStop[count-1])
    {
    dosomething<<<blocksPerGrid, threadsPerBlock>>>(d_bPtr,count,d_bStopPtr);
    }
//copy device back to host again
    thrust::copy(d_b.begin(),d_b.end(),h_a.begin());
    cout<<h_a[100]<<"\t"<<h_a[200]<<"\t"<<h_a[300]<<"\t"<<endl;
//wait to see the console output
    int x;
    cin>>x;
    return 0;
}

ただし、毎回while条件を確認する必要がありますが、遅いです。そこで、カーネル内でこのデバイス ベクターの状態を確認し、次のようにコードを変更することを考えています。

// include the header files
#include <iostream>
#include <stdio.h>
#include <time.h>

#include "cuda.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

using namespace std; 

//kernel function
__global__ 
void dosomething(int *d_bPtr, int count, int* d_bStopPtr)
{
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid==0)
    d_bStopPtr[tid]=0;
else if(tid<count)
    {
// if the last cell of the arrary is still not 0 yet, repeat
        while(d_bStopPtr[count-1])
        {
            d_bPtr[tid]=tid;
// only if the arrary cell before it is 0, then change it to 0 too
            if (d_bStopPtr[tid-1]==0 )
                d_bStopPtr[tid]=0;
        }
    }
}

int main()
{
    int count=100000;
// define the vectors
    thrust::host_vector <int> h_a(count);
    thrust::device_vector <int> d_b(count,0);
    int* d_bPtr=thrust::raw_pointer_cast(&d_b[0]);
    thrust::device_vector <int> d_bStop(count,1);
    int* d_bStopPtr=thrust::raw_pointer_cast(&d_bStop[0]);
// get the device property
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);

    int threadsPerBlock = prop.maxThreadsDim[0];
    int blocksPerGrid = min(prop.maxGridSize[0], (count + threadsPerBlock - 1) / threadsPerBlock);
//copy device to host
    thrust::copy(d_b.begin(),d_b.end(),h_a.begin());
    cout<<h_a[100]<<"\t"<<h_a[200]<<"\t"<<h_a[300]<<"\t"<<endl;
//run the kernel
    dosomething<<<blocksPerGrid, threadsPerBlock>>>(d_bPtr,count,d_bStopPtr);
//copy device back to host again
    thrust::copy(d_b.begin(),d_b.end(),h_a.begin());
    cout<<h_a[100]<<"\t"<<h_a[200]<<"\t"<<h_a[300]<<"\t"<<endl;
//wait to see the console output
    int x;
    cin>>x;
    return 0;
}

ただし、2 番目のバージョンでは常に、グラフィック カードとコンピューターがハングします。最初のバージョンのスピードアップを手伝ってくれませんか? カーネル内の状態を確認してから飛び出してカーネルを停止するにはどうすればよいですか?

4

1 に答える 1

2

基本的に、グローバルスレッドの同期動作を探しています。これは、GPU プログラミングでは禁物です。理想的には、各スレッドブロックは独立しており、独自のデータと処理に基づいて作業を完了できます。他のスレッドブロックの結果に依存して作業を完了するスレッドブロックを作成すると、デッドロック状態が発生する可能性があります。14 個の SM (スレッドブロック実行ユニット) を備えた GPU があり、100 個のスレッドブロックを作成するとします。ここで、スレッドブロック 0 ～ 13 が、スレッドブロック 99 がロックを解放するのを待っているとします (たとえば、特定の場所にゼロ値を書き込みます)。ここで、最初の 14 個のスレッドブロックが 14 個の SM で実行を開始し、おそらくループし、ロック値で回転するとします。スレッドブロック 99 が最初に実行されるか、まったく実行されないことを保証するメカニズムが GPU にありません。

「スレッドブロック 0 ～ 13 を強制的に削除する GMEM ストールについて」という質問には立ち入らないでください。スレッドブロック 99 がいつでも優先的に実行されることを保証するものはないからです。スレッドブロック 99 が実行されることを保証する唯一のことは、他のスレッドブロックの排出 (つまり、完了) です。しかし、他のスレッドブロックが回転していて、スレッドブロック 99 からの結果を待っている場合、それは決して起こらない可能性があります。

優れた上位互換性とスケーラブルな GPU コードは、独立した並列処理に依存します。そのため、少なくともスレッドブロック間レベルで、達成しようとしている作業を独立させるために、アルゴリズムを作り直すことをお勧めします。

グローバル スレッドの同期を行う必要がある場合、カーネルの起動は、これが確実に保証される唯一のポイントであるため、最初のアプローチが実用的なアプローチです。

これを支援するために、リダクション アルゴリズムが GPU にどのように実装されるかを調べることが役立つ場合があります。さまざまなタイプのリダクションにはすべてのスレッドに依存関係がありますが、中間結果を作成することで、作業を独立したピースに分割できます。次に、独立した部分をマルチカーネル アプローチ (またはその他のより高度なアプローチ) を使用して集約し、シリアル アルゴリズムに相当するものを高速化できます。

カーネルは実際にはあまり機能しません。1 つの配列をそのインデックスに等しく設定します。つまり、a[i] = i; そして、もう一方の配列をすべてゼロに設定します (順次ではありますが) b[i]=0;。最初のコードを「高速化」した例を示すには、次のようにします。

    // include the header files
#include <iostream>
#include <stdio.h>
#include <time.h>

#include "cuda.h"
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

using namespace std;

//kernel function
__global__
void dosomething(int *d_bPtr, int count, int* d_bStopPtr)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    while(tid<count)
    {
      d_bPtr[tid]=tid;
      while(d_bStopPtr[tid]!=0)
// only if the arrary cell before it is 0, then change it to 0 too
        if (tid==0) d_bStopPtr[tid] =0;
        else if (d_bStopPtr[tid-1]==0 )
               d_bStopPtr[tid]=0;
      tid += blockDim.x;
    }
}

int main()
{
    int count=100000;
// define the vectors
    thrust::host_vector <int> h_a(count);
    thrust::device_vector <int> d_b(count,0);
    int* d_bPtr=thrust::raw_pointer_cast(&d_b[0]);
    thrust::device_vector <int> d_bStop(count,1);
    int* d_bStopPtr=thrust::raw_pointer_cast(&d_bStop[0]);
// get the device property
    cudaDeviceProp prop;
    cudaGetDeviceProperties(&prop, 0);

//    int threadsPerBlock = prop.maxThreadsDim[0];
    int threadsPerBlock = 32;
//    int blocksPerGrid = min(prop.maxGridSize[0], (count + threadsPerBlock - 1) / threadsPerBlock);
    int blocksPerGrid = 1;
//copy device to host
    thrust::copy(d_b.begin(),d_b.end(),h_a.begin());
    cout<<h_a[100]<<"\t"<<h_a[200]<<"\t"<<h_a[300]<<"\t"<<endl;
//run the kernel
//    while(d_bStop[count-1])
//    {
    dosomething<<<blocksPerGrid, threadsPerBlock>>>(d_bPtr,count,d_bStopPtr);
//    }
//copy device back to host again
    cudaDeviceSynchronize();
    thrust::copy(d_b.begin(),d_b.end(),h_a.begin());
    cout<<h_a[100]<<"\t"<<h_a[200]<<"\t"<<h_a[300]<<"\t"<<endl;
//wait to see the console output
    int x;
    cin>>x;
    return 0;
}

私のマシンでは、これにより実行時間が 10 秒からほぼ瞬時 (1 秒未満) に短縮されます。32 スレッドのブロックを 1 つしか起動していないため、これは CUDA プログラミングの良い例ではないことに注意してください。機械を有効に活用するには、それだけでは十分ではありません。しかし、カーネルによって行われる作業は非常に些細なことなので、どのような例がよいのかわかりません。1 つの配列をそのインデックス a[i]=i; に設定するカーネルを作成できます。b[i]=0; および他の配列をゼロにします。すべて並行して。そうすればさらに高速になり、マシン全体をそのように使用できます。

于 2012-10-16T05:19:32.813 に答える