從OpenCL生成的PTX二進制文件中不必要的CVT指令

我寫了一個非常簡單的OpenCL代碼，並試圖在Tesla K40m GPU上執行它並測量GFlops。這是我寫的代碼：我已經部署在我的GPU與[1048576] global_work_size和[128]的local_work_size這個內核從OpenCL生成的PTX二進制文件中不必要的CVT指令

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

。它每秒可以執行的浮點操作總數約爲1.6 GFlops，這是非常低的。我假設我只是在進行單個操作，並且內存正在被順序讀取。我決定看看生成的PTX代碼：

.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;

由於它的代碼裏面清楚了，我有必要CVT和SHL說明，我以爲是架空的一個潛在原因。

現在我在這裏有兩個問題：（1）我應該如何重寫我的內核以擺脫兩條提到的指令並使內核執行得更快？（2）我的代碼中是否還有其他開銷源，我不知道？

來源

2017-08-14 saman

你是內存限制的，所以FLOPS會受到影響。您需要合併全局內存訪問才能獲得全部內存帶寬。 – Dithermaster

isf var是double類型，可以將指令源轉換爲float不能直接添加。

使用相同的溫度來添加所有東西是一個管道塞。

訪問陣列的步幅爲512浮點數，一次只能使用1個內存通道，甚至只能使用1個內存條。這個5月將每個線程的序列化指令序列化爲內存操作。

減少遠項目之間的縮減，而不是相鄰項目，並且每個線程只有4個項目來解決內存問題。

使用多個臨時管道問題。

將f後綴設爲浮點數，如果它們不適用於雙精度浮點數。儘量避免重複添加double和float。

每個線程使用不同的內存通道是很好的。

讓編譯器/硬件重命名一些寄存器是好的。

在相同寄存器上添加較少數值意味着較小的舍入誤差概率大於增加的值，這是很好的。

轉移似乎是float的地址計算爲4的長度。將左移2以獲得adr。也許緩衝區不對齊？計算基礎索引加指針，然後添加其他值，而不是重新計算基數，並在每行減少基數。在輸入自動優化考慮事項之前，可能是gin參數需要限制或const關鍵字

來源

2017-08-14 21:58:22

感謝@huseyin提示。我會考慮他們。 – saman

從OpenCL生成的PTX二進制文件中不必要的CVT指令

回答

相關問題