5

ブロックとグリッドのサイズを知るために助けが必要です。ユークリッド距離、マンハッタン、ピアソン、コサインなど、scipyに基づいてメトリック計算を実行するPythonアプリを構築しています。

プロジェクトはPycudaDistancesです。

小さなアレイで非常にうまく機能するようです。もっと徹底的なテストを行ったところ、残念ながらうまくいきませんでした。movielens set（http://www.grouplens.org/node/73）をダウンロードしました。

100kを使用Movielensして、形状（943、1682）の配列を宣言しました。つまり、ユーザーは943本と1682本の映画が評価されます。分類子ユーザーによるものではないフィルムは、値を0に構成しました。

はるかに大きな配列アルゴリズムでは、機能しなくなります。次のエラーが発生します。

pycuda._driver.LogicError：cuFuncSetBlockShapeが失敗しました：値が無効です。

このエラーを調べて、512スレッドをサポートするAndrewに、ブロックとグリッドを操作する必要がある、より大きなブロックを結合して操作するように指示する説明を見つけました。

アルゴリズムのユークリッド距離配列を小さな配列から巨大な配列まで機能するように適合させるための助けが必要でした。

def euclidean_distances(X, Y=None, inverse=True):
    X, Y = check_pairwise_arrays(X,Y)
    rows = X.shape[0]
    cols = Y.shape[0]
    solution = numpy.zeros((rows, cols))
    solution = solution.astype(numpy.float32)

    kernel_code_template = """
    #include <math.h>
    
    __global__ void euclidean(float *x, float *y, float *solution) {

        int idx = threadIdx.x + blockDim.x * blockIdx.x;
        int idy = threadIdx.y + blockDim.y * blockIdx.y;
        
        float result = 0.0;
        
        for(int iter = 0; iter < %(NDIM)s; iter++) {
            
            float x_e = x[%(NDIM)s * idy + iter];
            float y_e = y[%(NDIM)s * idx + iter];
            result += pow((x_e - y_e), 2);
        }
        int pos = idx + %(NCOLS)s * idy;
        solution[pos] = sqrt(result);
    }
    """
    kernel_code = kernel_code_template % {
        'NCOLS': cols,
        'NDIM': X.shape[1]
    }

    mod = SourceModule(kernel_code)

    func = mod.get_function("euclidean")
    func(drv.In(X), drv.In(Y), drv.Out(solution), block=(cols, rows, 1))

    return numpy.divide(1.0, (1.0 + solution)) if inverse else solution

詳細については、https ：//github.com/vinigracindo/pycudaDistances/blob/master/distances.pyを参照してください。

4

2 に答える 2

16

カーネルの実行パラメータのサイズを設定するには、次の2つのことを（この順序で）行う必要があります。

1.ブロックサイズを決定します

ブロックサイズは、主にハードウェアの制限とパフォーマンスによって決まります。詳細については、この回答を読むことをお勧めしますが、非常に短い要約では、GPUには、実行できるブロックあたりのスレッドの総数に制限があり、有限のレジスタファイル、共有およびローカルメモリサイズがあります。選択するブロックの寸法は、これらの制限内に収まっている必要があります。そうでない場合、カーネルは実行されません。ブロックサイズもカーネルのパフォーマンスに影響を与える可能性があり、最適なパフォーマンスを提供するブロックサイズが見つかります。ブロックサイズは常にワープサイズのラウンド倍である必要があります。これは、これまでにリリースされたすべてのCUDA互換ハードウェアで32です。

2.グリッドサイズを決定します

示した種類のカーネルの場合、必要なブロックの数は、入力データの量と各ブロックの次元に直接関係しています。

たとえば、入力配列のサイズが943x1682で、ブロックサイズが16x16の場合、59 x 106グリッドが必要になり、カーネルの起動時に944x1696スレッドが生成されます。この場合、入力データサイズはブロックサイズの丸め倍数ではないため、カーネルを変更して、範囲外に読み取られないようにする必要があります。1つのアプローチは次のようになります。

__global__ void euclidean(float *x, float *y, float *solution) {
    int idx = threadIdx.x + blockDim.x * blockIdx.x;
    int idy = threadIdx.y + blockDim.y * blockIdx.y;

     if ( ( idx < %(NCOLS)s ) && ( idy < %(NDIM)s ) ) {

        .....
     }
}

カーネルを起動するPythonコードは、次のようになります。

bdim = (16, 16, 1)
dx, mx = divmod(cols, bdim[0])
dy, my = divmod(rows, bdim[1])

gdim = ( (dx + (mx>0)) * bdim[0], (dy + (my>0)) * bdim[1]) )
func(drv.In(X), drv.In(Y), drv.Out(solution), block=bdim, grid=gdim)

この質問と回答は、このプロセスがどのように機能するかを理解するのにも役立ちます。

上記のコードはすべてブラウザで記述されており、テストされていないことに注意してください。自己責任で使用してください。

また、これはコードの非常に短い読み取りに基づいており、質問でコードがどのように呼び出されるかについて実際には何も説明していないため、正しくない可能性があることにも注意してください。

于 2013-01-25T11:48:40.000 に答える
2

受け入れられた答えは原則として正しいですが、talonmiesがリストしたコードは完全には正しくありません。行： あるgdim = ( (dx + (mx>0)) * bdim[0], (dy + (my>0)) * bdim[1]) ) べき： gdim = ( (dx + (mx>0)), (dy + (my>0)) ) 明らかな余分な括弧に加えて、gdimはあなたが望むものよりもはるかに多くのスレッドを生成します。talonmiesは、スレッドはブロックサイズ*グリッドサイズであると彼のテキストで正しく説明していました。ただし、彼がリストしたgdimは、必要な正しいグリッドサイズではなく、合計スレッドを提供します。

于 2017-12-28T19:23:21.690 に答える