0

私のアプリでは、各スレッドに独自のデータマトリックスが必要です。たとえば、私にはTスレッドがあり、各スレッドは異なるマトリックスで動作するとしますD[M][N]。

私の質問：データ構造をどのように整理するのですか？

私の解決策：要素の配列を定義しますA。T*M*Nバンクの競合を避けるために、最初D[0][0] Tに各スレッドの要素時間を格納し、次にD[0][1]...D[0][M-1]などD[1][0]を格納します（行列のようにこの配列を見ると、T * (M*N)スレッドごとに1つの列があります）。このようにして、異なるメモリバンクの異なるスレッドに同じ要素があります。これに対応して、次の方法D[i][j]でスレッドの要素にアクセスします。xD[i][j](x) == A[T * (M * i + j) + x]

私の問題：複雑なインデックスを計算するのは計算コストがかかります。

PS私はNvidiaTeslaC2075（CUDA 2.0）を持っています。

4

1 に答える 1

0

MとNは数百になる可能性があるとあなたは言います。そのため、共有メモリをあまり使用できなくなります（あるとしても）。グローバルメモリ消費も注意深く監視できます（ただし、Teslaには大量のメモリがあります）。200x200 x 3584スレッド（私の意見ではC2075の最小値）x sizeof（int）-547MBのデータになります。

グローバルメモリアクセスパターンの動作は異なります。グローバルメモリは、32、64、および128Bのセグメントに分割されます。読み取りのコストは、ワープごとの異なるセグメントアクセスの数とほぼ同じです。要するに、それは通常、要約すると、アクセスが分散しているほど、悪化します。

したがって、すべてのスレッドが同じインデックスで独自のマトリックスにアクセスしない限り（少なくともほとんどの場合）、メモリ編成は効率的ではありません。ただし、上記が当てはまる場合は、説明しているレイアウトが機能する可能性があります。

また、アクセスパターンが分散している場合は、L1キャッシュを無効にすると役立つ場合があります。これは、L1キャッシュラインが128Bであるのに対し、L2は32Bしかないためです。したがって、オーバーフェッチを減らすことができます。少なくとも-試してみてください:)

配列にアクセスする際の苦痛を和らげるために、私は次のようなことをします。

//let the kernel dimentions be known at compile time - you can safe some computation and registers
//assuming one-dimentional kernels

static const int blockSize = ...; //equivalent to blockDim
static const int gridSize = ...; //equivalent to gridDim
static const int rowSize = blockSize * gridSize;

template <typename T, int M, int N>
class MyMatrix {
private:
  T* data; //flattened array in global memory
  int tid;
public:
  __device__ inline MyMatrix(T* dataIsHere) : data(dataIsHere) {
    tid = threadIdx.x+blockDim.x*blockIdx.x;
  }
  __device__ inline T& operator()(int x, int y) {
    return data[(y*M+x)*rowSize+tid];
  }
}

//assuming the matrix size is 200x200 and consists of ints:

__global__ void myKernel(int* flattenedMatrices) {
  MyMatrix<int,200,200> matrix(flattenedMatrices);

  ...

  matrix(2,4) = .... // happily access the matrix for both loads and stores
}
于 2012-12-05T22:33:19.283 に答える