cuda - cudaトリプルリーマン和

Question

cuda経由でトリプルリーマン和を実行しようとしています。ループの入れ子を避けるために、合計反復子に多次元グリッド反復子を使用しようとしています。2.0 telsa カードを使用しているため、ネストされたカーネルを使用できません。

必要な x、y、z 変数のそれぞれについて、完全な 0 -> N 反復を取得しているようには見えません。

__global__ void test(){
uint xIteration = blockDim.x * blockIdx.x + threadIdx.x;
uint yIteration = blockDim.y * blockIdx.y + threadIdx.y;
uint zIteration = blockDim.z * blockIdx.z + threadIdx.z;
printf("x: %d * %d + %d = %d\n y: %d * %d + %d = %d\n z: %d * %d + %d = %d\n", blockDim.x, blockIdx.x, threadIdx.x, xIteration, blockDim.y, blockIdx.y, threadIdx.y, yIteration,  blockDim.z, blockIdx.z, threadIdx.z, zIteration);
}

---- によって呼び出された -----

int totalIterations = 128; // N value for single sum (i = 0; i < N)
dim3 threadsPerBlock(8,8,8);
dim3 blocksPerGrid((totalIterations + threadsPerBlock.x - 1) / threadsPerBlock.x, 
                   (totalIterations + threadsPerBlock.y - 1) / threadsPerBlock.y, 
                   (totalIterations + threadsPerBlock.z - 1) / threadsPerBlock.z);
test<<<blocksPerGrid, threadsPerBlock>>>();

----出力-----

出力が切り捨てられました。現在、0 < x、y、z < 7 のすべての順列を取得していますが、totalIterations が 128 の場合、0 < x、y、z < 127 が必要です。たとえば、この実行では、40 < z < 49 です。、ここで、0 <= z <= 127 である必要があります。多次元グリッドに関する私の理解は間違っている可能性がありますが、リーマンの場合、各反復子、x、y、および z は 0 から 127 の値を持つ必要があります。

また、totalIterations > 128、ex 1024 にすると、プログラムは cudaError コード 6 で終了します。これは、起動タイマーの期限切れであると理解しています。カーネルは印刷だけを行っているため、タイムアウトする理由がわかりません。これをセカンダリデバイスで実行すると、問題が解決するようです。テスラの 1 つを使用して X を実行していますが、計算用に両方のテスラを解放するための新しいディスプレイデバイスになる geforce が郵送されています。

printf(...) は、合計される関数の実行に置き換えられます。

アイデアは、のシリアルコードバージョンを置き換えることです

for (int i = 0...)
    for (int j = 0 ..)
       for (int k = 0...)

また、潜在的に巨大な (数百万 x 数百万 x 数百万) 3D 配列を作成してから縮小するのはメモリ効率が悪いように見えるため、関数値を格納する方法もわかりませんが、何らかの方法で関数値をある種の共有配列に連結します。変数。

---- デバイス情報 (これらのカードが 2 枚あり、両方の出力は同じです ----

Device 1: "Tesla C2050"
  CUDA Driver Version / Runtime Version          5.0 / 5.0
  CUDA Capability Major/Minor version number:    2.0
  Total amount of global memory:                 2687 MBytes (2817982464 bytes)
  (14) Multiprocessors x ( 32) CUDA Cores/MP:    448 CUDA Cores
  GPU Clock rate:                                1147 MHz (1.15 GHz)
  Memory Clock rate:                             1500 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 multiprocessor:  1536
  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:                     No
  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):      Yes
  Device PCI Bus ID / PCI location ID:           132 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

score 1 · Accepted Answer

デバイスコードで printf を使用して、(x,y,z) 配列のすべての要素がスレッドによって処理されたことを検証することは、x,y,z の値が大きい場合には賢明ではないと既に述べたように思います。

すべての要素 x、y、z がスレッドに触れていることを証明するために、コードに基づいて次を作成しました。

#include <stdio.h>
#define DATAVAL 1
#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)

__global__ void test(int *data, int dim){
  uint xIteration = blockDim.x * blockIdx.x + threadIdx.x;
  uint yIteration = blockDim.y * blockIdx.y + threadIdx.y;
  uint zIteration = blockDim.z * blockIdx.z + threadIdx.z;

  data[((((zIteration*dim)+yIteration)*dim)+xIteration)]=DATAVAL;
}

int main(){
  int *testdata;
  int *result;
  int totalIterations = 128; // N value for single sum (i = 0; i < N)
  int testsize = totalIterations*totalIterations*totalIterations;
  dim3 threadsPerBlock(8,8,8);
  dim3 blocksPerGrid((totalIterations + threadsPerBlock.x - 1) / threadsPerBlock.x,  (totalIterations + threadsPerBlock.y - 1) / threadsPerBlock.y,  (totalIterations + threadsPerBlock.z - 1) / threadsPerBlock.z);
  cudaMalloc(&testdata, testsize*sizeof(int));
  cudaCheckErrors("cudaMalloc fail");
  cudaMemset(testdata, 0, testsize*sizeof(int));
  cudaCheckErrors("cudaMemset fail");
  result=(int *)malloc(testsize*sizeof(int));
  if (result == 0) {printf("malloc fail \n"); return 1;}
  memset(result, 0, testsize*sizeof(int));
  test<<<blocksPerGrid, threadsPerBlock>>>(testdata, totalIterations);
  cudaDeviceSynchronize();
  cudaCheckErrors("Kernel launch failure");
  cudaMemcpy(result, testdata, testsize*sizeof(int), cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy failure");

  for (unsigned i=0; i<testsize; i++)
    if (result[i] != DATAVAL) {printf("fail! \n"); return 1;}

  printf("Success \n");
  return 0;

}

cuda - cudaトリプルリーマン和

1 に答える 1

Related

Reference