サム3D行列の合計

Question

A [x] [y] = sum {z = 0からz = n}まで{B [x] [y] [z] + C [x]行列Aは次元[高さ] [幅]、行列BはCは次元[高さ] [幅] [n]を有する。

index = 0; 
for (z = 0; z<n; ++z) 
    for(y = 0; y<width; ++y) 
     for(x = 0; x<height; ++x) { 
      matrix[index] = value; 
      index++; 
     } 

Q1：

値は、のようなもので、メモリにマッピングされ、このクーダカーネルは大丈夫でしょうか？

idx = blockIdx.x*blockDim.x + threadIdx.x; 
idy = blockIdx.y*blockDim.y + threadIdx.y; 

for(z=0; z<n; z++){ 
    A[idx*width+idy] += B[idx*width+idy+z*width*height] + C[idx*width+idy+z*width*height]; 
} 

Q2：これは計算を高速化しますか？

idx = blockIdx.x*blockDim.x + threadIdx.x; 
idy = blockIdx.y*blockDim.y + threadIdx.y; 
idz = blockIdx.z*blockDim.z + threadIdx.z; 

int stride_x = blockDim.x * gridDim.x; 
int stride_y = blockDim.y * gridDim.y; 
int stride_z = blockDim.z * gridDim.z; 

while (idx < height && idy < width && idz < n) { 
    atomicAdd(&(A[idx*width+idy]), B[idx*width+idy+idz*width*height] + C[idx*width+idy+idz*width*height]); 
    idx += stride_x; 
    idy += stride_y; 
    idz += stride_z; 
} 

Answer 1

Q1：非常に大きな行列を使用するときに問題がある可能性があります：あなたは答え

備考を知っている行列でそれをテストします。適切な増分でwhileループを使用します。 Example by Cudaはいつものように参考書です。

ネストループを実装する例は、For nested loops with CUDAです。そこにはwhileループが実装されています。

marina.kは競合状態です。それは、原子操作がコードを遅くする傾向があるので、アプローチ1を好むでしょう。

Answer 2

最初のカーネルは問題ありません。しかし、マトリックスBとCへのアクセスを統合していません。

第2カーネル機能については、 1つのスレッドだけでなく、A[idx*width+idy]のアドレスに書き込むことができるデータ競合の原因があります。 AttomicAdd

一般的な質問： 私は実験がより良いことを示すと思います。それはあなたが持っている典型的なマトリックスサイズに依存します。 Fermi < 1024の最大スレッドブロックサイズと、行列のサイズが大きい場合は、多くのスレッドブロックが宝石であることに注意してください。通常は（スレッドブロックが多い）処理が遅くなります。 ArrayFireで

Answer 3

実簡単：スレッドブロック= 3、スレッドブロックのグリッドの最大次元の

array A = randu(nx,ny,nz); 
array B = sum(A,2); // sum along 3rd dimension 
print(B);