scrittura alla memoria globale o locale aumenta il tempo di esecuzione del kernel per 10000%

Question

Ho il seguente kernel OpenCL:

kernel void ndft(
    global float *re, global float *im, int num_values, 
    global float *spectrum_re, global float *spectrum_im, 
    global float *spectrum_abs, 
    global float *sin_array, global float *cos_array, 
    float sqrt_num_values_reciprocal) 
{ 
    // MATH MAGIC - DISREGARD FROM HERE ----------- 

    float x; 
    float y; 
    float sum_re = 0; 
    float sum_im = 0; 

    size_t thread_id = get_global_id(0); 
    //size_t local_id = get_local_id(0); 

    // num_values = 24 (live environment), 48 (test) 
    for (int i = 0; i < num_values; i++) 
    { 
     x = cos_array[thread_id * num_values + i] * sqrt_num_values_reciprocal; 
     y = sin_array[thread_id * num_values + i] * sqrt_num_values_reciprocal; 
     sum_re = sum_re + re[i] * x + im[i] * y; 
     sum_im = sum_im - re[i] * y + x * im[i]; 
    } 

    // MATH MAGIC DONE ---------------------------- 

    //spectrum_re[thread_id] = sum_re; 
    //spectrum_im[thread_id] = sum_im; 
    //spectrum_abs[thread_id] = hypot(sum_re, sum_im); 
    float asdf = hypot(sum_re, sum_im); // this is just a dummy calculation 
}

In questo modo, il tempo di esecuzione è su di noi 15 (dimensioni gruppo di lavoro = 567, 14 gruppi di lavoro , per un totale di 7938 thread).

Tuttavia, naturalmente, in qualche modo bisogno di recuperare i risultati dell'operazione, che è quello che le ultime righe sono per, (commentata). Appena ho eseguire un singolo di quelle operazioni di memoria (e non importa se è spectrum_Xglobal, come nell'esempio, o local), il tempo exeuction del kernel aumenta a ~ 1,4 a 1,5 ms.

ho pensato l'aumento del tempo di esecuzione è una sorta di overhead fisso, quindi sarebbe solo accumulare più dati, in modo che la quantità relativa di tempo perso a causa di questo effetto minimizza. Ma quando raddoppio il mio numero di thread (cioè il doppio della quantità di dati), anche il tempo di esecuzione raddoppia (a 2,8 ~ 3,0 ms).

Ho scoperto che anche se ho solo decommentato uno di quelle righe, ho lo stesso tempo di esecuzione come se disattivalo tutte e tre. Anche se aggiungo uno if (thread_id == 0) ed eseguo, ho lo stesso tempo di esecuzione. Tuttavia, è troppo lento in questo modo (il limite massimo per la mia applicazione è di circa 30 us). Perfino si perfeziona circa 5 volte più velocemente quando lo eseguo nel normale codice C sulla mia CPU.

Ora sono ovviamente facendo qualcosa di sbagliato, ma io non sono sicuro da dove cominciare a cercare una soluzione.

---

Come ho commentato in risposta talonmies', ho anche fatto il seguente:

Dal codice di cui sopra, ho fatto gli ultimi 4 linee assomigliano

//spectrum_re[thread_id] = sum_re; 
//spectrum_im[thread_id] = sum_im; 
spectrum_abs[thread_id] = hypot(sum_re, sum_im); 
//float asdf = hypot(sum_re, sum_im);

Come previsto, il tempo di esecuzione ~ 1,8 ms. ci

// 
// Generated by NVIDIA NVVM Compiler 
// Compiler built on Tue Apr 03 12:42:39 2012 (1333449759) 
// Driver 
// 

.version 3.0 
.target sm_21, texmode_independent 
.address_size 32 


.entry ndft(
    .param .u32 .ptr .global .align 4 ndft_param_0, 
    .param .u32 .ptr .global .align 4 ndft_param_1, 
    .param .u32 ndft_param_2, 
    .param .u32 .ptr .global .align 4 ndft_param_3, 
    .param .u32 .ptr .global .align 4 ndft_param_4, 
    .param .u32 .ptr .global .align 4 ndft_param_5, 
    .param .u32 .ptr .global .align 4 ndft_param_6, 
    .param .u32 .ptr .global .align 4 ndft_param_7, 
    .param .f32 ndft_param_8 
) 
{ 
    .reg .f32 %f; 
    .reg .pred %p; 
    .reg .s32 %r; 


    ld.param.u32 %r3, [ndft_param_2]; 
    // inline asm 
    mov.u32  %r18, %envreg3; 
    // inline asm 
    // inline asm 
    mov.u32  %r19, %ntid.x; 
    // inline asm 
    // inline asm 
    mov.u32  %r20, %ctaid.x; 
    // inline asm 
    // inline asm 
    mov.u32  %r21, %tid.x; 
    // inline asm 
    add.s32  %r22, %r21, %r18; 
    mad.lo.s32 %r11, %r20, %r19, %r22; 
    setp.gt.s32  %p1, %r3, 0; 
    @%p1 bra BB0_2; 

    mov.f32  %f46, 0f00000000; 
    mov.f32  %f45, %f46; 
    bra.uni  BB0_4; 

BB0_2: 
    ld.param.u32 %r38, [ndft_param_2]; 
    mul.lo.s32 %r27, %r38, %r11; 
    shl.b32  %r28, %r27, 2; 
    ld.param.u32 %r40, [ndft_param_6]; 
    add.s32  %r12, %r40, %r28; 
    ld.param.u32 %r41, [ndft_param_7]; 
    add.s32  %r13, %r41, %r28; 
    mov.f32  %f46, 0f00000000; 
    mov.f32  %f45, %f46; 
    mov.u32  %r43, 0; 
    mov.u32  %r42, %r43; 

BB0_3: 
    add.s32  %r29, %r13, %r42; 
    ld.global.f32 %f18, [%r29]; 
    ld.param.f32 %f44, [ndft_param_8]; 
    mul.f32  %f19, %f18, %f44; 
    add.s32  %r30, %r12, %r42; 
    ld.global.f32 %f20, [%r30]; 
    mul.f32  %f21, %f20, %f44; 
    ld.param.u32 %r35, [ndft_param_0]; 
    add.s32  %r31, %r35, %r42; 
    ld.global.f32 %f22, [%r31]; 
    fma.rn.f32 %f23, %f22, %f19, %f46; 
    ld.param.u32 %r36, [ndft_param_1]; 
    add.s32  %r32, %r36, %r42; 
    ld.global.f32 %f24, [%r32]; 
    fma.rn.f32 %f46, %f24, %f21, %f23; 
    neg.f32  %f25, %f22; 
    fma.rn.f32 %f26, %f25, %f21, %f45; 
    fma.rn.f32 %f45, %f24, %f19, %f26; 
    add.s32  %r42, %r42, 4; 
    add.s32  %r43, %r43, 1; 
    ld.param.u32 %r37, [ndft_param_2]; 
    setp.lt.s32  %p2, %r43, %r37; 
    @%p2 bra BB0_3; 

BB0_4: 
    // inline asm 
    abs.f32  %f27, %f46; 
    // inline asm 
    // inline asm 
    abs.f32  %f29, %f45; 
    // inline asm 
    setp.gt.f32  %p3, %f27, %f29; 
    selp.f32 %f8, %f29, %f27, %p3; 
    selp.f32 %f32, %f27, %f29, %p3; 
    // inline asm 
    abs.f32  %f31, %f32; 
    // inline asm 
    setp.gt.f32  %p4, %f31, 0f7E800000; 
    mov.f32  %f47, %f32; 
    @%p4 bra BB0_6; 

    mov.f32  %f48, %f8; 
    bra.uni  BB0_7; 

BB0_6: 
    mov.f32  %f33, 0f3E800000; 
    mul.rn.f32 %f10, %f8, %f33; 
    mul.rn.f32 %f47, %f32, %f33; 
    mov.f32  %f48, %f10; 

BB0_7: 
    mov.f32  %f13, %f48; 
    // inline asm 
    div.approx.f32 %f34, %f13, %f47; 
    // inline asm 
    mul.rn.f32 %f39, %f34, %f34; 
    add.f32  %f38, %f39, 0f3F800000; 
    // inline asm 
    sqrt.approx.f32  %f37, %f38;  // <-- this is part of hypot() 
    // inline asm 
    mul.rn.f32 %f40, %f32, %f37; 
    add.f32  %f41, %f32, %f8; 
    setp.eq.f32  %p5, %f32, 0f00000000; 
    selp.f32 %f42, %f41, %f40, %p5; 
    setp.eq.f32  %p6, %f32, 0f7F800000; 
    setp.eq.f32  %p7, %f8, 0f7F800000; 
    or.pred  %p8, %p6, %p7; 
    selp.f32 %f43, 0f7F800000, %f42, %p8; 
    shl.b32  %r33, %r11, 2; 
    ld.param.u32 %r39, [ndft_param_5]; 
    add.s32  %r34, %r39, %r33; 
    st.global.f32 [%r34], %f43; // <-- stores the hypot's result in spectrum_abs 
    ret; 
} 

Infatti tutte le mie operazioni di calcolo sono - un sacco di aggiunge/MULTS, nonché una sqrt per la funzione hypot: Il codice assembler generato per il mio sistema è. Dal codice asm sopra, ho rimosso la penultima riga:

st.global.f32 [%r34], %f43;

che è la linea che memorizza effettivamente i dati nella matrice globale spectrum_abs. Quindi ho usato clCreateProgramWithBinary e ho usato il file di codice asm modificato come input. Il tempo di esecuzione è sceso a 20 us.

Answer 1

direi che si sta vedendo gli effetti di ottimizzazione del compilatore.

Il compilatore NVIDIA è molto aggressivo ad eliminare "codice morto" che non partecipa direttamente in una scrittura a memoria globale. Quindi nel tuo kernel, se non scrivi sum_re o sum_im, il compilatore ottimizzerà l'intero ciclo di calcolo (e probabilmente tutto il resto) e lascerà il tuo con un kernel vuoto che contiene nient'altro che un no-op. Il tempo di esecuzione di 15 microsecondi che si sta vedendo è principalmente l'overhead di lancio del kernel e non molto altro. Quando si decommenta una memoria globale, il compilatore lascia tutto il codice di calcolo in posizione e si vede il vero tempo di esecuzione del codice.

Quindi la vera domanda che dovresti probabilmente porsi è come ottimizzare quel kernel per ridurne il tempo di esecuzione dagli 1,5 millisecondi attualmente necessari al tuo obiettivo (molto ambizioso) di 30 microsecondi.

---

Nonostante lo scetticismo espresso alla risposta originale, ecco un caso Repro completa che supporta l'affermazione che questo è un effetto correlato compilatore:

#include <iostream> 
#include <OpenCL/opencl.h> 

size_t source_size; 
const char * source_str = 
"kernel void ndft(                 \n" \ 
" global float *re, global float *im, int num_values,        \n" \ 
" global float *spectrum_re, global float *spectrum_im,        \n" \ 
" global float *spectrum_abs,              \n" \ 
" global float *sin_array, global float *cos_array,         \n" \ 
" float sqrt_num_values_reciprocal)             \n" \ 
"{                      \n" \ 
" // MATH MAGIC - DISREGARD FROM HERE -----------         \n" \ 
"                      \n" \ 
" float x;                   \n" \ 
" float y;                   \n" \ 
" float sum_re = 0;                 \n" \ 
" float sum_im = 0;                 \n" \ 
"                      \n" \ 
" size_t thread_id = get_global_id(0);            \n" \ 
"                      \n" \ 
" for (int i = 0; i < num_values; i++)            \n" \ 
" {                     \n" \ 
"  x = cos_array[thread_id * num_values + i] * sqrt_num_values_reciprocal;  \n" \ 
"  y = sin_array[thread_id * num_values + i] * sqrt_num_values_reciprocal;  \n" \ 
"  sum_re += re[i] * x + im[i] * y;            \n" \ 
"  sum_im -= re[i] * y + x * im[i];            \n" \ 
" }                     \n" \ 
"                      \n" \ 
" // MATH MAGIC DONE ----------------------------         \n" \ 
"                      \n" \ 
" //spectrum_re[thread_id] = sum_re;            \n" \ 
" //spectrum_im[thread_id] = sum_im;            \n" \ 
" //spectrum_abs[thread_id] = hypot(sum_re, sum_im);        \n" \ 
"}                      \n"; 

int main(void) 
{ 
    int err; 

    cl_device_id device_id; 
    clGetDeviceIDs(NULL, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL); 
    cl_context context = clCreateContext(0, 1, &device_id, NULL, NULL, &err); 
    cl_program program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &err); 

    err = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL); 

    cl_uint program_num_devices; 
    clGetProgramInfo(program, CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint), &program_num_devices, NULL); 

    size_t * binaries_sizes = new size_t[program_num_devices]; 
    clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES, program_num_devices*sizeof(size_t), binaries_sizes, NULL); 

    char **binaries = new char*[program_num_devices]; 
    for (size_t i = 0; i < program_num_devices; i++) 
     binaries[i] = new char[binaries_sizes[i]+1]; 

    clGetProgramInfo(program, CL_PROGRAM_BINARIES, program_num_devices*sizeof(size_t), binaries, NULL); 
    for (size_t i = 0; i < program_num_devices; i++) 
    { 
     binaries[i][binaries_sizes[i]] = '\0'; 
     std::cout << "Program " << i << ":" << std::endl; 
     std::cout << binaries[i]; 
    } 
    return 0; 
} 

Quando compilato ed eseguito, emette il follow Codice PTX dal runtime OpenCL:

Program 0: 
bplist00?^clBinaryDriver\clBinaryData_clBinaryVersionWCLH 1.0O!.version 1.5 
.target sm_12 
.target texmode_independent 

.reg .b32 r<126>; /* define r0..125 */ 
.reg .b64 x<126>; /* define r0..125 */ 
.reg .b32 f<128>; /* define f0..127 */ 
.reg .pred p<32>; /* define p0..31 */ 
.reg .u32 sp; 

.reg .b8 wb0,wb1,wb2,wb3; /* 8-bit write buffer */ 
.reg .b16 ws0,ws1,ws2,ws3; /* 16-bit write buffer */ 
.reg .b32 tb0,tb1,tb2,tb3; /* read tex buffer */ 
.reg .b64 vl0,vl1; /* 64-bit vector buffer */ 
.reg .b16 cvt16_0,cvt16_1; /* tmps for conversions */ 


.const .align 1 .b8 ndft_gid_base[52]; 
.local .align 16 .b8 ndft_stack[8]; 
.entry ndft(
    .param.b32 ndft_0 /* re */, 
    .param.b32 ndft_1 /* im */, 
    .param.b32 ndft_2 /* num_values */, 
    .param.b32 ndft_3 /* spectrum_re */, 
    .param.b32 ndft_4 /* spectrum_im */, 
    .param.b32 ndft_5 /* spectrum_abs */, 
    .param.b32 ndft_6 /* sin_array */, 
    .param.b32 ndft_7 /* cos_array */, 
    .param.f32 ndft_8 /* sqrt_num_values_reciprocal */ 
) { 
    mov.u32 sp, ndft_stack; 
    mov.u32 r0, 4294967295; 
    ld.param.u32 r1, [ndft_2 + 0]; 
LBB1_1: 
    add.u32 r0, r0, 1; 
    setp.lt.s32 p0, r0, r1; 
    @p0 bra LBB1_1; 
LBB1_2: 
    ret; 
} 

ie. uno stub del kernel che non contiene alcun ciclo di calcolo. Quando i tre memoria globale, scrive nelle ultime tre righe del kernel non sono commentate, emette questo:

Program 0: 
S.version 1.5inaryDriver\clBinaryData_clBinaryVersionWCLH 1.0O 
.target sm_12 
.target texmode_independent 

.reg .b32 r<126>; /* define r0..125 */ 
.reg .b64 x<126>; /* define r0..125 */ 
.reg .b32 f<128>; /* define f0..127 */ 
.reg .pred p<32>; /* define p0..31 */ 
.reg .u32 sp; 

.reg .b8 wb0,wb1,wb2,wb3; /* 8-bit write buffer */ 
.reg .b16 ws0,ws1,ws2,ws3; /* 16-bit write buffer */ 
.reg .b32 tb0,tb1,tb2,tb3; /* read tex buffer */ 
.reg .b64 vl0,vl1; /* 64-bit vector buffer */ 
.reg .b16 cvt16_0,cvt16_1; /* tmps for conversions */ 


.const .align 1 .b8 ndft_gid_base[52]; 
.local .align 16 .b8 ndft_stack[8]; 
.entry ndft(
    .param.b32 ndft_0 /* re */, 
    .param.b32 ndft_1 /* im */, 
    .param.b32 ndft_2 /* num_values */, 
    .param.b32 ndft_3 /* spectrum_re */, 
    .param.b32 ndft_4 /* spectrum_im */, 
    .param.b32 ndft_5 /* spectrum_abs */, 
    .param.b32 ndft_6 /* sin_array */, 
    .param.b32 ndft_7 /* cos_array */, 
    .param.f32 ndft_8 /* sqrt_num_values_reciprocal */ 
) { 
    mov.u32 sp, ndft_stack; 
    cvt.u32.u16 r0, %tid.x; 
    cvt.u32.u16 r1, %ntid.x; 
    cvt.u32.u16 r2, %ctaid.x; 
    mad24.lo.u32 r0, r2, r1, r0; 
    mov.u32 r1, 0; 
    shl.b32 r2, r1, 2; 
    mov.u32 r3, ndft_gid_base; 
    add.u32 r2, r2, r3; 
    ld.const.u32 r2, [r2 + 40]; 
    add.u32 r0, r0, r2; 
    ld.param.u32 r2, [ndft_2 + 0]; 
    mul.lo.u32 r3, r0, r2; 
    shl.b32 r3, r3, 2; 
    mov.f32 f0, 0f00000000 /* 0.000000e+00 */; 
    ld.param.f32 f1, [ndft_8 + 0]; 
    ld.param.u32 r4, [ndft_7 + 0]; 
    ld.param.u32 r5, [ndft_6 + 0]; 
    ld.param.u32 r6, [ndft_5 + 0]; 
    ld.param.u32 r7, [ndft_4 + 0]; 
    ld.param.u32 r8, [ndft_3 + 0]; 
    ld.param.u32 r9, [ndft_1 + 0]; 
    ld.param.u32 r10, [ndft_0 + 0]; 
    mov.u32 r11, r1; 
    mov.f32 f2, f0; 
LBB1_1: 
    setp.ge.s32 p0, r11, r2; 
    @!p0 bra LBB1_7; 
LBB1_2: 
    shl.b32 r1, r0, 2; 
    add.u32 r2, r8, r1; 
    st.global.f32 [r2+0], f0; 
    add.u32 r1, r7, r1; 
    st.global.f32 [r1+0], f2; 
    abs.f32 f1, f2; 
    abs.f32 f0, f0; 
    setp.gt.f32 p0, f0, f1; 
    selp.f32 f2, f0, f1, p0; 
    abs.f32 f3, f2; 
    mov.f32 f4, 0f7E800000 /* 8.507059e+37 */; 
    setp.gt.f32 p1, f3, f4; 
    selp.f32 f0, f1, f0, p0; 
    shl.b32 r0, r0, 2; 
    add.u32 r0, r6, r0; 
    @!p1 bra LBB1_8; 
LBB1_3: 
    mul.rn.f32 f3, f2, 0f3E800000 /* 2.500000e-01 */; 
    mul.rn.f32 f1, f0, 0f3E800000 /* 2.500000e-01 */; 
LBB1_4: 
    mov.f32 f4, 0f00000000 /* 0.000000e+00 */; 
    setp.eq.f32 p0, f2, f4; 
    @!p0 bra LBB1_9; 
LBB1_5: 
    add.f32 f1, f2, f0; 
LBB1_6: 
    mov.f32 f3, 0f7F800000 /* inf */; 
    setp.eq.f32 p0, f0, f3; 
    setp.eq.f32 p1, f2, f3; 
    or.pred p0, p1, p0; 
    selp.f32 f0, f3, f1, p0; 
    st.global.f32 [r0+0], f0; 
    ret; 
LBB1_7: 
    add.u32 r12, r3, r1; 
    add.u32 r13, r4, r12; 
    ld.global.f32 f3, [r13+0]; 
    mul.rn.f32 f3, f3, f1; 
    add.u32 r13, r9, r1; 
    ld.global.f32 f4, [r13+0]; 
    mul.rn.f32 f5, f3, f4; 
    add.u32 r12, r5, r12; 
    ld.global.f32 f6, [r12+0]; 
    mul.rn.f32 f6, f6, f1; 
    add.u32 r12, r10, r1; 
    ld.global.f32 f7, [r12+0]; 
    mul.rn.f32 f8, f7, f6; 
    add.f32 f5, f8, f5; 
    sub.f32 f2, f2, f5; 
    mul.rn.f32 f4, f4, f6; 
    mul.rn.f32 f3, f7, f3; 
    add.f32 f3, f3, f4; 
    add.f32 f0, f0, f3; 
    add.u32 r11, r11, 1; 
    add.u32 r1, r1, 4; 
    bra LBB1_1; 
LBB1_8: 
    mov.f32 f1, f0; 
    mov.f32 f3, f2; 
    bra LBB1_4; 
LBB1_9: 
    div.approx.f32 f1, f1, f3; 
    mul.rn.f32 f1, f1, f1; 
    add.f32 f1, f1, 0f3F800000 /* 1.000000e+00 */; 
    sqrt.approx.ftz.f32 f1, f1; 
    mul.rn.f32 f1, f2, f1; 
    bra LBB1_6; 
} 

Penso che questa è la prova abbastanza inconfutabile che è l'ottimizzazione del compilatore, che è causa la differenza di tempo di esecuzione, e dipende solo se le scritture di memoria sono incluse nel codice del kernel o meno.

---

Credo che la domanda finale diventa allora perché questo è così lento (a prescindere dal dibattito sul fatto che questo è causato da ottimizzazione del compilatore o meno). Il tempo di esecuzione di 1,5 millisecondi che stai vedendo è un vero riflesso delle prestazioni del codice e la vera domanda è perché. Dalla mia lettura del codice del kernel, la risposta sembra risiedere in schemi di accesso alla memoria che sono piuttosto orribili per la GPU. All'interno del ciclo di calcolo si dispone di una memoria a due mondiale legge con molto grandi passi avanti, come questo:

x = cos_array[thread_id * num_values + i] * sqrt_num_values_reciprocal; 

accordo con il commento nel codice num_values è o 24 o 48. Ciò significa che la memoria si legge non può assolutamente coalesce e la cache L1 su una GPU Fermi non sarà di grande aiuto. Ciò avrà un enorme impatto negativo sull'utilizzo della larghezza di banda della memoria e renderà il codice molto lento. Se si è bloccati con l'ordinamento dei dati di input, una soluzione più rapida sarebbe quella di utilizzare una distorsione per eseguire il calcolo di un output (quindi una riduzione del warp wide alla somma finale). Ciò ridurrà il passo di lettura da 24 o 48 a 1 e fonderà le letture di memoria globale da questi due grandi array di input.

All'interno del ciclo v'è anche ripetuto recupera alla memoria globale per 24 o 48 elementi di re e im:

sum_re += re[i] * x + im[i] * y; 
    sum_im -= re[i] * y + x * im[i]; 

Questo non è necessaria, e spreca un sacco di banda di memoria globale o l'efficienza della cache (la La GPU non ha abbastanza registri per permettere al compilatore di tenere l'intero array nel registro). Sarebbe molto meglio che ogni gruppo di lavoro legga questi due array negli array di memoria __local una volta e utilizzi la copia di memoria locale all'interno del ciclo di elaborazione. Se ogni gruppo di lavoro esegue calcoli più volte, anziché solo una volta, è possibile potenzialmente risparmiare un sacco di larghezza di banda della memoria globale e ammortizzare la lettura iniziale fino a renderlo quasi libero.