Numba python CUDAとcuBLASの単純な操作での速度の差

Question

私はいくつかのコードをプロファイリングしていますが、パフォーマンスの不一致を理解することはできません。私は2つの配列（インプレース）の間で単純な要素別の追加をしようとしています。これはnumbaを使ってCUDAカーネルです：

from numba import cuda 

@cuda.jit('void(float32[:], float32[:])') 
def cuda_add(x, y): 

    ix = cuda.threadIdx.x + cuda.blockIdx.x * cuda.blockDim.x 
    stepSize = cuda.gridDim.x * cuda.blockDim.x 
    while ix < v0.shape[0]: 
     y[ix] += x[ix] 
     ix += stepSize 

私はパフォーマンスは大丈夫だと思ったが、その後、私はCUBLASメソッドにそれを比較した：

from accelerate.cuda.blas import Blas 

blas = Blas() 
blas.axpy(1.0, X, Y) 

BLAS方法の性能が約25％高速であります大規模な配列（20M要素）の場合コンパイルされたPTXコードがすでにキャッシュされているので、これを呼び出すことでcuda.jitカーネルを「ウォーミングアップ」した後です（この問題があるかどうかはわかりませんが、問題ではないことを確認するだけです）。

レベル3のマトリックスマトリックス操作でこのパフォーマンスの違いを理解できましたが、これは簡単な追加です。私はcuda.jitコードからより多くのパフォーマンスを絞るために何かできることはありますか？最適化したい実際のコードはblas.axpyに渡すことができない2次元配列なので、私は尋ねています。

EDIT実行コードやその他の必要なパッケージ：

import numpy as np 

def main(): 
    n = 20 * 128 * 128 * 64 
    x = np.random.rand(n).astype(np.float32) 
    y = np.random.rand(n).astype(np.float32) 

    ## Create necessary GPU arrays 
    d_x = cuda.to_device(x) 
    d_y = cuda.to_device(y) 

    ## My function 
    cuda_add[1024, 64](d_x , d_y) 

    ## cuBLAS function 
    blas = Blas() 
    blas.axpy(1.0, d_x , d_y) 

Answer 1

非常に短い答えはノーです。 CUBLASは、numba CUDA方言が現在サポートしていないメモリ結合コードのパフォーマンスを向上させるために、多くのこと（テクスチャ、ベクトルタイプ）を活用しています。

私はCUDAでこれを書き上げ：

__device__ float4 add(float4 x, float4 y) 
{ 
    x.x += y.x; x.y += y.y; x.z += y.z; x.w += y.w; 
    return x; 
} 

__global__ void mykern(float* x, float* y, int N) 
{ 
    float4* x4 = reinterpret_cast<float4*>(x); 
    float4* y4 = reinterpret_cast<float4*>(y); 

    int strid = gridDim.x * blockDim.x; 
    int tid = threadIdx.x + blockDim.x * blockIdx.x; 

    for(; tid < N/4; tid += strid) { 
     float4 valx = x4[tid]; 
     float4 valy = y4[tid]; 
     y4[tid] = add(valx, valy); 
    }  
} 

と私のベンチマークは、それがCUBLASの約5％以内であることを示しているが、私はあなたが現時点でnumbaでそれを行うことができると信じていません。

saxpyを2D配列で実行できないとのご意見はお控えください。配列がメモリ内で連続していて（同じでなければならないと思われる）同じレイアウト（つまり、転置を追加しようとしていない）の場合、saxpyを2次元配列に使用します。