unnötige CVT-Anweisungen in PTX-Binärdatei generiert von OpenCL

Question

Ich habe einen wirklich einfachen OpenCL-Code geschrieben und versucht, es auf einer Tesla K40m GPU auszuführen und die GFlops zu messen. Hier ist der Code, den ich geschrieben habe:

__kernel void test(__global float *GIn, __global float *GOut, int M, int N, int P) { 
    int X = get_global_id(0); 
    // Just a private variable 
    float temp = 1.0; 


    // Start of a new level of for loop 
    int baseIndex1 = (X) * 512; 
    temp += GIn[baseIndex1 + 0] * var; 
    temp += GIn[baseIndex1 + 1] * var; 
    temp += GIn[baseIndex1 + 2] * var; 
    temp += GIn[baseIndex1 + 3] * var; 
    temp += GIn[baseIndex1 + 4] * var; 
    temp += GIn[baseIndex1 + 5] * var; 
    temp += GIn[baseIndex1 + 6] * var; 
    temp += GIn[baseIndex1 + 7] * var; 
    temp += GIn[baseIndex1 + 8] * var; 
    temp += GIn[baseIndex1 + 9] * var; 
    temp += GIn[baseIndex1 + 10] * var; 
    ... 
    temp += GIn[baseIndex1 + 510] * var; 
    temp += GIn[baseIndex1 + 511] * var; 
    GOut[baseIndex1] = temp; 
} 

ich bereitgestellt habe diesen Kernel auf meiner GPU mit global_work_size von [1048576] und local_work_size von [128]. Die Gesamtzahl der Gleitkommaoperationen, die es pro Sekunde ausführen kann, liegt bei 1,6 GFlops, was extrem niedrig ist. Ich nehme an, ich mache nur einzelne Operationen und auch Speicher wird sequentiell gelesen. Ich habe beschlossen, bei der generierten PTX Code einen Blick zu nehmen:

.version 5.0 
.target sm_35, texmode_independent 
.address_size 64 

    // .globl test 
.func (.param .b64 func_retval0) get_global_id 
(
    .param .b32 get_global_id_param_0 
) 
; 

.entry test(
    .param .u64 .ptr .global .align 4 test_param_0, 
    .param .u64 .ptr .global .align 4 test_param_1, 
    .param .u32 test_param_2, 
    .param .u32 test_param_3, 
    .param .u32 test_param_4 
) 
{ 
    .reg .f32 %f<1537>; 
    .reg .b32 %r<515>; 
    .reg .b64 %rd<1543>; 


    ld.param.u64 %rd1, [test_param_0]; 
    ld.param.u64 %rd2, [test_param_1]; 
    mov.u32  %r1, 0; 
    // Callseq Start 0 
    { 
    .reg .b32 temp_param_reg; 
    // <end>} 
    .param .b32 param0; 
    st.param.b32 [param0+0], %r1; 
    .param .b64 retval0; 
    call.uni (retval0), 
    get_global_id, 
    (
    param0 
    ); 
    ld.param.b64 %rd3, [retval0+0]; 

    //{ 
    }// Callseq End 0 
    cvt.u32.u64 %r2, %rd3; 
    mul.lo.s32 %r3, %r2, 512; 
    cvt.s64.s32 %rd4, %r3; 
    shl.b64  %rd5, %rd4, 2; 
    add.s64  %rd6, %rd1, %rd5; 
    ld.global.f32 %f1, [%rd6]; 
    mul.f32  %f2, %f1, 0f3FC00000; 
    add.f32  %f3, %f2, 0f3F800000; 
    add.s32  %r4, %r3, 1; 
    cvt.s64.s32 %rd7, %r4; 
    shl.b64  %rd8, %rd7, 2; 
    add.s64  %rd9, %rd1, %rd8; 
    ld.global.f32 %f4, [%rd9]; 
    mul.f32  %f5, %f4, 0f3FC00000; 
    add.f32  %f6, %f3, %f5; 
    add.s32  %r5, %r3, 2; 
    cvt.s64.s32 %rd10, %r5; 
    shl.b64  %rd11, %rd10, 2; 
    add.s64  %rd12, %rd1, %rd11; 
    ld.global.f32 %f7, [%rd12]; 
    mul.f32  %f8, %f7, 0f3FC00000; 
    add.f32  %f9, %f6, %f8; 
    add.s32  %r6, %r3, 3; 
    cvt.s64.s32 %rd13, %r6; 
    shl.b64  %rd14, %rd13, 2; 
    add.s64  %rd15, %rd1, %rd14; 
    ld.global.f32 %f10, [%rd15]; 
    mul.f32  %f11, %f10, 0f3FC00000; 
    add.f32  %f12, %f9, %f11; 
    add.s32  %r7, %r3, 4; 
    cvt.s64.s32 %rd16, %r7; 
    shl.b64  %rd17, %rd16, 2; 
    add.s64  %rd18, %rd1, %rd17; 
    ld.global.f32 %f13, [%rd18]; 
    mul.f32  %f14, %f13, 0f3FC00000; 
    add.f32  %f15, %f12, %f14; 
    add.s32  %r8, %r3, 5; 
    cvt.s64.s32 %rd19, %r8; 
    shl.b64  %rd20, %rd19, 2; 
    add.s64  %rd21, %rd1, %rd20; 
    ld.global.f32 %f16, [%rd21]; 
    mul.f32  %f17, %f16, 0f3FC00000; 
    add.f32  %f18, %f15, %f17; 

Wie es in dem Code klar ist, ich habe unnötigen cvt und shl Anweisungen, die ich sind eine mögliche Ursache von Kopfnehmen.

Jetzt habe ich zwei Fragen hier: (1) Wie soll ich meinen Kernel umschreiben, um zwei erwähnte Anweisungen loszuwerden und den Kernel schneller zu machen? (2) Gibt es in meinem Code eine andere Quelle für Overhead, von der ich nichts weiß?

Answer 1

Bei var handelt es sich um einen doppelten Typ, bei dem es sich um eine Konvertierungsbefehlsquelle handeln könnte, da Float diese nicht direkt hinzufügen kann.

Verwenden Sie die gleiche Temperatur für das Hinzufügen von allem ist ein Pipeline-Stopper.

Zugriff auf Array mit einer Schrittlänge von 512 Floats könnte nur 1 Speicherkanal und sogar nur 1 Speicherbank auf einmal verwenden. Dies kann die Speicheroperationen auf bereits serialisierte Anweisungen pro Thread serialisieren.

Reduzieren Sie zwischen entfernten Elementen, nicht Nachbarn und nur Paaren oder vielleicht 4 Elementen pro Thread, um Speicherprobleme zu lösen.

Verwenden Sie mehrere Temps für Pipeline-Problem.

Setzen Sie f postfix für Floats, wenn sie nicht für doubles gedacht sind. Versuchen Sie zu vermeiden, doppelte und float wiederholt hinzuzufügen.

Verwenden eines anderen Speicherkanals pro Thread ist gut.

Den Compiler/Hardware einige Register umbenennen lassen ist gut.

Das Hinzufügen einer geringeren Anzahl von Werten in demselben Register bedeutet eine geringere Wahrscheinlichkeit für einen Rundungsfehler. Größer als der Mehrwert, was gut ist.

Verschiebung scheint Adressberechnung für float als Länge von 4 zu sein. Verschiebung um 2 nach links, um adr. Vielleicht ist der Puffer nicht ausgerichtet? Berechne den Basisindex plus Zeiger und füge dann andere Werte hinzu, anstatt die Basis und ihre Additionen an jeder Zeile neu zu berechnen, was langsam wird. Möglicherweise müssen Gin-Parameter das Schlüsselwort const beschränken, bevor Sie automatische Optimierungsüberlegungen eingeben.