6

次のスクリプトは、ベンチマーク用に設定されています。ユークリッド L2 ノルムを使用して N ポイント間の距離を計算します。3 つの異なるルーチンが実装されています。

  1. scipy.spatial.distance.pdist関数を使用した高度なソリューション。
  2. かなり低レベルの OpenMP を利用したscipy.weave.inlineソリューション。
  3. pyCUDA を利用した GPGPU ソリューション。

GTX660 (2GB RAM) を使用した i5-3470 (16GB RAM) でのベンチマーク結果は次のとおりです。

    ------------
    Scipy Pdist
    Execution time: 3.01975 s
    Frist five elements: [ 0.74968684  0.71457213  0.833188    0.48084545  0.86407363]
    Last five elements: [ 0.65717077  0.76850474  0.29652017  0.856179    0.56074625]
    ------------
    Weave Inline
    Execution time: 2.48705 s
    Frist five elements: [ 0.74968684  0.71457213  0.83318806  0.48084542  0.86407363]
    Last five elements: [ 0.65717083  0.76850474  0.29652017  0.856179    0.56074625]
    ------------
    pyCUDA
    CUDA clock timing:  0.713028930664
    Execution time: 2.04364 s
    Frist five elements: [ 0.74968684  0.71457213  0.83318806  0.48084542  0.86407363]
    Last five elements: [ 0.65717083  0.76850468  0.29652017  0.856179    0.56074625]
    ------------

pyCUDA のパフォーマンスには少しがっかりしています。私は CUDA を初めて使用するので、ここで何かが欠けている可能性があります。では、問題の核心はどこにあるのでしょうか? グローバル メモリ帯域幅の限界に達していますか? ブロックサイズとグリッドサイズの選択が悪い?

import numpy,time,math
import pycuda.autoinit
import pycuda.driver as drv
from pycuda.compiler import SourceModule
from scipy.spatial.distance import pdist
from scipy import weave

def weave_solution(x):
    """
    OpenMP powered weave inline.
    """
    N,DIM     = numpy.shape(x)
    L         = ((N-1)**2+(N-1))/2
    solution  = numpy.zeros(L).astype(numpy.float32)
    ncpu      = 4
    weave_omp = {'headers'           : ['<omp.h>'],
                 'extra_compile_args': ['-fopenmp'],
                 'extra_link_args'   : ['-lgomp']}
    code = \
         r'''
         omp_set_num_threads(ncpu);
         #pragma omp parallel
         {            
            int j,d,pos;
            float r=0.0;

            #pragma omp for
               for (int i=0; i<(N-1); i++){
                  for (j=(i+1); j<N; j++){
                     r = 0.0;
                     for (d=0; d<DIM; d++){
                        r += (x[i*DIM+d]-x[j*DIM+d])*(x[i*DIM+d]-x[j*DIM+d]);
                     }
                     pos = (i*N+j)-(i*(i+1)/2)-i-1;
                     solution[pos] = sqrt(r);
                  }
               }

         }
         '''
    weave.inline(code,['x','N','DIM','solution','ncpu'],**weave_omp)
    return numpy.array(solution)

def scipy_solution(x):
    """
    SciPy High-level function
    """
    return pdist(x).astype(numpy.float32)

def cuda_solution(x):
    """
    pyCUDA
    """
    N,DIM     = numpy.shape(x)
    N         = numpy.int32(N)
    DIM       = numpy.int32(DIM)    
    L         = ((N-1)**2+(N-1))/2
    solution  = numpy.zeros(L).astype(numpy.float32)

    start = drv.Event()
    end   = drv.Event()       

    mod = SourceModule("""
    __global__ void distance(float *x,int N,int DIM,float *solution){

    const int i = blockDim.x * blockIdx.x + threadIdx.x;

    int j,d,pos;
    float r=0.0;

    if ( i < (N-1) ){

       for (j=(i+1); j<N; j++){

          r = 0.0;
          for (d=0; d<DIM; d++){
             r += (x[i*DIM+d]-x[j*DIM+d])*(x[i*DIM+d]-x[j*DIM+d]);
          }
          pos = (i*N+j)-(i*(i+1)/2)-i-1;
          solution[pos] = sqrt(r);
       }

    }
    }
    """)


    func = mod.get_function("distance")

    start.record()
    func(drv.In(x),N,DIM,drv.Out(solution),block=(192,1,1),grid=(192,1))
    end.record()
    end.synchronize()
    secs = start.time_till(end)*1e-3

    print "CUDA clock timing: ",secs
    return solution

if __name__ == '__main__':

    # Set up data points
    N   = 25000
    DIM = 3
    x   = numpy.random.rand(N,DIM).astype(numpy.float32)

    print "-"*12
    # Scipy solution
    print "Scipy Pdist"
    stime = time.time()
    spsolution = scipy_solution(x)
    stime = time.time()-stime
    print "Execution time: {0:.5f} s".format(stime)
    print "Frist five elements:", spsolution[:5]
    print "Last five elements:", spsolution[-5:]    
    print "-"*12

    # Weave solution
    print "Weave Inline"
    wtime = time.time()
    wsolution = weave_solution(x)
    wtime = time.time()-wtime
    print "Execution time: {0:.5f} s".format(wtime)
    print "Frist five elements:", wsolution[:5]
    print "Last five elements:", wsolution[-5:]
    print "-"*12

    # pyCUDA solution
    print "pyCUDA"
    ctime = time.time()
    csolution = cuda_solution(x)
    ctime = time.time()-ctime
    print "Execution time: {0:.5f} s".format(ctime)
    print "Frist five elements:", csolution[:5]
    print "Last five elements:", csolution[-5:]    
    print "-"*12

編集:

ハッシュバン行を追加しました

#!/usr/bin/env python

ファイルの先頭に追加し、実行可能にしました。weave.inlineとを使用して計算をコメントアウトした後scipy.spatial.distance.pdist、NVIDIA Visual Profiler は次の結果を表示します。

NVIDIA ビジュアル プロファイラー

4

1 に答える 1

4

現在、それぞれがN-1の位置を更新する192のスレッドがあり、より多くのブロック/スレッドを簡単に起動できます。

このループfor (j=(i+1); j<N; j++){の代わりに、内側のループだけを実行するN-1スレッドに置き換えます。

さらに詳しく知りたい場合は、N-1 * DIMスレッドごとに内部ループでステートメントを実行し、結果を共有メモリに格納して、最後にその削減を行うことができます。CUDAでの並列削減の最適化を参照してください

この行を見てください:

r += (x[i*DIM+d]-x[j*DIM+d])*(x[i*DIM+d]-x[j*DIM+d]);

メモリトランザクションとパターンは均一ではなく、合体しています。また、pycudaが-O3をnvccに渡すかどうかわからないため、nvccが式をここに示されている4つではなく2つのメモリトランザクションに最適化するかどうかもわかりません。(x[i*DIM+d]-x[j*DIM+d])レジスタ変数に入れて確認し、自分で二乗してください。

それ以外の場合は#pragma unroll、可能であれば、各forループの前に配置して展開することもできます。

于 2012-12-30T21:01:13.213 に答える