Numba Python CUDA vs. cuBLAS Geschwindigkeit Unterschied auf einfache Operationen

Question

Ich bin Profiling einige Code und kann nicht feststellen, eine Leistungsdiskrepanz. Ich versuche eine einfache elementweise Addition zwischen zwei Arrays (in-place) zu machen. Dies ist der CUDA Kernel numba:

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 

ich die Leistung dachte, war in Ordnung, aber dann vergleichen ich es in die cuBLAS Methode:

from accelerate.cuda.blas import Blas 

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

Die Leistung der BLAS Methode ist in etwa 25% schneller für große Arrays (20M Elemente). Dies geschieht nach dem "Aufwärmen" des cuda.jit-Kernels durch vorheriges Aufrufen, so dass der kompilierte PTX-Code bereits zwischengespeichert ist (nicht sicher, ob dies von Bedeutung ist, aber nur um sicherzustellen, dass das nicht das Problem war).

Ich könnte diesen Leistungsunterschied für Ebene 3 Matrix-Matrix-Operationen verstehen, aber das ist eine einfache Ergänzung. Gibt es etwas, was ich tun kann, um mehr Leistung aus dem cuda.jit-Code herauszuholen? Ich frage, weil der wirkliche Code, den ich optimieren möchte, ein 2D-Array ist, das nicht an blas.axpy übergeben werden kann.

EDIT Execution-Code und weitere Pakete:

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

Die sehr kurze Antwort ist nein. CUBLAS nutzt eine Reihe von Dingen (Texturen, Vektortypen), um die Performance von speichergebundenem Code, wie ihn der cua-Dialekt von numba derzeit nicht unterstützt, zu verbessern.

gestrichelte ich dies in CUDA aus:

__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); 
    }  
} 

und mein Benchmarking zeigt es innerhalb etwa 5% der CUBLAS sein, aber ich glaube nicht, dass Sie in numba im Moment tun können.

Nebenbei verstehe ich nicht Ihre Bemerkung über saxpy auf einem 2D-Array nicht ausführen zu können. Wenn Sie Arrays im Speicher zusammenhängend sind (wie ich vermute, dass sie sein müssen) und das gleiche Layout haben (d. H. Nicht versuchen, eine Transponierte hinzuzufügen), können Sie saxpy auf einem 2D-Array verwenden.