3

あらゆる種類のサイズの SDK マトリックス転置サンプルのバリエーションを取得しようとしていました。簡単に言うと、入力配列 (double *a) を取得し、それをより大きな行列 (double *tab) の 2 つの異なる部分 (オフセットが異なることに気付くでしょう) に書き込む必要があります。データを行優先形式で保存しているので、インデックス作成にこのマクロを使用しています。

#define IDX2L(i,j,ld) (((i)*ld))+(j)) // 0 based index +row-major format

これは私が使用する単純なコードです。

__global__ void cuda_a_Coalesced(double *tab, int tab_rows, int a_rows, double *a)
{
    __shared__  double tile[16*(16+1)]; 
    int col = threadIdx.x + blockIdx.x * blockDim.x;
    int row = threadIdx.y + blockIdx.y * blockDim.y;

    int col_2, row_2;
    int a_cols=tab_rows-a_rows; // tab_rows-a_rows is the number of columns of a
    int tab_cols=2*tab_rows+2;  // 2*tab_rows+2 is the number of columns of tab

    if( (col<a_cols) && (row<a_rows) ) 
    {
        // Load the data into shared mem
        tile[threadIdx.x+threadIdx.y*(16+1)]=a[IDX2L(row,col,a_cols)];

        // Normal copy (+ offsets)
        tab[IDX2L(row,col+tab_rows+a_rows,tab_cols)]= tile[threadIdx.x+threadIdx.y*(16+1)];

        // New idx
        col_2 = blockIdx.y * blockDim.y + threadIdx.x;
        row_2 = blockIdx.x * blockDim.x + threadIdx.y;
    }
    __syncthreads();

    if( (row_2<a_cols) && (col_2<a_rows) )
        // Transpose (+ other offsets)
        tab[IDX2L(row_2+a_rows,col_2+tab_rows,tab_cols)]= -tile[threadIdx.y+threadIdx.x*(16+1)];

}

起動パラメータは次のとおりです。

b1=(int)ceil((float)a_cols/16);
b2=(int)ceil((float)a_rows/16);
dim bck(b1,b2):dim th(16,16);

cuda_a_Coalesced<<<bck,th>>>(tab,tab_rows,a_rows,a);

通常のコピーは、サイズに関係なく常に良好に実行されます。転置コピーは、ブロック サイズの整数倍のサイズに対してのみ機能します (SDK サンプルのように)。転置コピーが失敗した場合、正確に予測または追跡できない方法で、操作の一部が正しく、他の部分は正しくありません。アイデアは共有メモリのインデックスを変更して、転置を出力行列の合体形式で書き込むことができることに注意してください (行の主な形式のため)。

コードがその種のサイズでのみ機能する理由を誰かが教えてくれますか?

この状況を解決するためのトリックはありますか？

4

1 に答える 1

2

この問題は、col_2 と row_2 の値が if() ステートメント内で割り当てられ、すべてのスレッドがアクセスしていなかったため、いくつかの未定義のスレッドが原因でした。

この状況を解決するために、これらの変数を宣言するときに col_2 と row_2 の値を指定し、前述の if() 内に配置されていた同義計算を削除できます。

__shared__  double tile[16*(16+1)];

int col = threadIdx.x + blockIdx.x * blockDim.x;
int row = threadIdx.y + blockIdx.y * blockDim.y;

int col_2 = blockIdx.y * blockDim.y + threadIdx.x;
int row_2 = blockIdx.x * blockDim.x + threadIdx.y;

int a_cols=tab_rows-a_rows; 
int tab_cols=2*tab_rows+2;

したがって、残りのコードは次のようになります。

if( (col<a_cols) && (row<a_rows) ) 
{
    // Load the data into shared mem
    tile[threadIdx.x+threadIdx.y*(16+1)]=a[IDX2L(row,col,a_cols)];
    // Normal copy (+ offsets)
    tab[IDX2L(row,col+tab_rows+a_rows,tab_cols)]= tile[threadIdx.x+threadIdx.y*(16+1)];
}
__syncthreads();

if( (row_2<a_cols) && (col_2<a_rows) )
    // Transpose (+ other offsets)
    tab[IDX2L(row_2+a_rows,col_2+tab_rows,tab_cols)]= -tile[threadIdx.y+threadIdx.x*(16+1)];
于 2012-11-15T13:27:10.670 に答える