1

私はこれを NVIDIA フォーラムに投稿しました。もう少し助けてもらえると思いました。

複数のケースで実行するためにコードを拡張しようとすると問題が発生します。私は最も一般的なケースを念頭に置いて開発してきましたが、今度はテストの時間であり、さまざまなケースですべてが機能することを確認する必要があります。現在、私のカーネルはループ内で実行されています (すべてを実行するために 1 つのカーネル呼び出しを行っていないのには理由があります)。行列の行全体の値を計算します。最も一般的なケースは、512 列 x 512 行です。512 x 512、1024 x 512、512 x 1024、およびその他の組み合わせのサイズのマトリックスを考慮する必要がありますが、最大のものは 1024 x 1024 マトリックスになります。私はかなり単純なカーネル呼び出しを使用しています:

launchKernel<<<1,512>>>(................)

このカーネルは、一般的な 512x512 および 512 x 1024 (それぞれ列、行) の場合には正常に機能しますが、1024 x 512 の場合には機能しません。この場合、実行には 1024 スレッドが必要です。私の素朴さで、1024 スレッドを起動する単純なカーネル呼び出しのさまざまなバージョンを試してきました。

launchKernel<<<2,512>>>(................)  // 2 blocks with 512 threads each ???
launchKernel<<<1,1024>>>(................) // 1 block with 1024 threads ???

私の問題は、スレッドとブロックの理解不足に関係していると思います

これは deviceQuery の出力です。ご覧のとおり、最大 1024 のスレッドを持つことができます。

C:\ProgramData\NVIDIA Corporation\NVIDIA GPU Computing SDK 4.1\C\bin\win64\Release\deviceQuery.exe Starting...

 CUDA Device Query (Runtime API) version (CUDART static linking)

Found 2 CUDA Capable device(s)

Device 0: "Tesla C2050"
  CUDA Driver Version / Runtime Version          4.2 / 4.1
  CUDA Capability Major/Minor version number:    2.0
  Total amount of global memory:                 2688 MBytes (2818572288 bytes)
  (14) Multiprocessors x (32) CUDA Cores/MP:     448 CUDA Cores
  GPU Clock Speed:                               1.15 GHz
  Memory Clock rate:                             1500.00 Mhz
  Memory Bus Width:                              384-bit
  L2 Cache Size:                                 786432 bytes
  Max Texture Dimension Size (x,y,z)             1D=(65536), 2D=(65536,65535), 3D=(2048,2048,2048)
  Max Layered Texture Size (dim) x layers        1D=(16384) x 2048, 2D=(16384,16384) x 2048
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total number of registers available per block: 32768
  Warp size:                                     32
  Maximum number of threads per block:           1024
  Maximum sizes of each dimension of a block:    1024 x 1024 x 64
  Maximum sizes of each dimension of a grid:     65535 x 65535 x 65535
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and execution:                 Yes with 2 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Concurrent kernel execution:                   Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support enabled:                Yes
  Device is using TCC driver mode:               No
  Device supports Unified Addressing (UVA):      No
  Device PCI Bus ID / PCI location ID:           40 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

Device 1: "Quadro 600"
  CUDA Driver Version / Runtime Version          4.2 / 4.1
  CUDA Capability Major/Minor version number:    2.1
  Total amount of global memory:                 1024 MBytes (1073741824 bytes)
  ( 2) Multiprocessors x (48) CUDA Cores/MP:     96 CUDA Cores
  GPU Clock Speed:                               1.28 GHz
  Memory Clock rate:                             800.00 Mhz
  Memory Bus Width:                              128-bit
  L2 Cache Size:                                 131072 bytes
  Max Texture Dimension Size (x,y,z)             1D=(65536), 2D=(65536,65535), 3D=(2048,2048,2048)
  Max Layered Texture Size (dim) x layers        1D=(16384) x 2048, 2D=(16384,16384) x 2048
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total number of registers available per block: 32768
  Warp size:                                     32
  Maximum number of threads per block:           1024
  Maximum sizes of each dimension of a block:    1024 x 1024 x 64
  Maximum sizes of each dimension of a grid:     65535 x 65535 x 65535
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and execution:                 Yes with 1 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Concurrent kernel execution:                   Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support enabled:                No
  Device is using TCC driver mode:               No
  Device supports Unified Addressing (UVA):      No
  Device PCI Bus ID / PCI location ID:           15 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 4.2, CUDA Runtime Version = 4.1, NumDevs = 2, Device = Tesla C2050, Device = Quadro 600

私は Tesla C2050 デバイスのみを使用しています。これは私のカーネルの削除されたバージョンです。

#define twoPi               6.283185307179586
#define speed_of_light      3.0E8
#define MaxSize             999

__global__ void calcRx4CPP4
(  
        const float *array1,  
        const double *array2,  
        const float scalar1,  
        const float scalar2,  
        const float scalar3,  
        const float scalar4,  
        const float scalar5,  
        const float scalar6,  
        const int scalar7,  
        const int scalar8,    
        float *outputArray1,
        float *outputArray2)  
{  

    float scalar9;  
    int idx;  
    double scalar10;
    double scalar11;  
    float sumReal, sumImag;  
    float real, imag;  

    float coeff1, coeff2, coeff3, coeff4;  

    sumReal = 0.0;  
    sumImag = 0.0;  

    // kk loop 1 .. 512 (scalar7)  
    idx = (blockIdx.x * blockDim.x) + threadIdx.x;  

    /* Declare the shared memory parameters */
    __shared__ float SharedArray1[MaxSize];
    __shared__ double SharedArray2[MaxSize];

    /* populate the arrays on shared memory */
    SharedArray1[idx] = array1[idx];  // first 512 elements
    SharedArray2[idx] = array2[idx];
    if (idx+blockDim.x < MaxSize){
        SharedArray1[idx+blockDim.x] = array1[idx+blockDim.x];
        SharedArray2[idx+blockDim.x] = array2[idx+blockDim.x];
    }            
    __syncthreads();

    // input scalars used here.
    scalar10 = ...;
    scalar11 = ...;

    for (int kk = 0; kk < scalar8; kk++)
    {  
        /* some calculations */
        // SharedArray1, SharedArray2 and scalar9 used here
        sumReal = ...;
        sumImag = ...;
    }  


    /* calculation of the exponential of a complex number */
    real = ...;
    imag = ...;
    coeff1 = (sumReal * real);  
    coeff2 = (sumReal * imag);  
    coeff3 = (sumImag * real);  
    coeff4 = (sumImag * imag);  

    outputArray1[idx] = (coeff1 - coeff4);  
    outputArray2[idx] = (coeff2 + coeff3);  


}  

ブロックあたりの最大スレッド数が 1024 であるため、単純なカーネル起動を引き続き使用できると思っていましたが、間違っていますか?

各カーネルを 1024 スレッドで正常に起動するにはどうすればよいですか?

4

2 に答える 2

5

ブロックごとのスレッド数を変えたくありません。CUDA Occupancy Calculator を使用して、カーネルのブロックあたりの最適なスレッド数を取得する必要があります。その数を取得したら、必要なスレッドの総数を取得するために必要な数のブロックを起動するだけです。特定のケースで必要なスレッドの数が、ブロックごとのスレッドの倍数であるとは限らない場合は、カーネルの先頭にコードを追加して、不要なスレッドを中止します。( if () return;)。次に、カーネルに必要な情報に応じて、追加のパラメーターを使用して、または x および y グリッド次元を使用して、マトリックスの次元を渡します (私はそれを調べていません)。

私の推測では、1024 スレッドで問題が発生している理由は、GPU がブロック内の多くのスレッドをサポートしているにもかかわらず、リソースの使用状況に基づいて各ブロックで使用できるスレッドの数に別の制限要因があるためです。あなたのカーネル。制限要因は、共有メモリまたはレジスタの使用です。Occupancy Calculator はどれかを教えてくれますが、その情報はカーネルを最適化したい場合にのみ重要です。

于 2012-05-04T03:20:46.927 に答える
3

1024スレッドで1つのブロックを使用すると、MaxSizeが999しかないため、データが間違ってしまうため、問題が発生します。

最後のスレッド#1023についてシミュレートしましょう

__shared__ float SharedArray1[999];     
__shared__ double SharedArray2[999];

/* populate the arrays on shared memory */     
SharedArray1[1023] = array1[1023]; 
SharedArray2[1023] = array2[1023];     

if (2047 < MaxSize)
{         
    SharedArray1[2047] = array1[2047];         
    SharedArray2[2047] = array2[2047];     
}                 
__syncthreads(); 

計算でこれらすべての要素を使用する場合、これは機能しないはずです。(計算コードは表示されていないので、その前提です)

于 2012-05-04T11:54:13.753 に答える