2017-09-29 70 views
0

我正在學習在python中使用opencl,我想優化一個函數。我瞭解到,這可以通過將全局內存存儲在本地內存中來完成。然而,它不應該像它應該那樣工作,持續時間是其兩倍。這做得好嗎?我可以更優化此代碼嗎?OpenCl簡單不成功優化

__kernel void sumOP( __global float *input, 
        __global float *weights, 
        int layer_size, 
        __global float *partialSums,__local float* cache) 

{ 

    private const int i = get_global_id(0); 
    private const int in_layer_s = layer_size; 
    private const int item_id = get_local_id(0);  
    private const int group_id = get_group_id(0); 
    private const int group_count = get_num_groups(0); 

    const int localsize = get_local_size(0); 

      for (int x = 0; x < in_layer_s; x++) 
      { 
       cache[x] = weights[i*in_layer_s + x]; 
      } 


      float total1 = 0; 

      for (int x = 0; x < in_layer_s; x++) 
      { 
       total1 += cache[x] *input[x]; 
      } 
      partialSums[i] = sigmoid(total1); 

} 

Python的通話

l = opencl.LocalMemory(len(inputs)) 
event = program.sumOP(queue, output.shape, np.random.randn(6,).shape, inputs.data, weights.data,np.int32(len(inputs)),output.data,l) 

感謝一些建議

+1

使用本地內存進行優化的一般想法適用於工作組中的工作項都使用全局內存中的相似值的情況。您不必多次讀取這些內容,而是將它們緩存在更快(但更小)的本地內存中,以便在工作組中重新使用。你的內核是否需要這個?如果是這樣,你的內核的第一部分應該(並行)分擔加載它們的負擔,有障礙,然後做計算。最小限度地,你的代碼缺少障礙。 – Dithermaster

+1

同樣,寫入同一個「cache [x]」地址的組的所有工作項在競態條件方面都不好。應該像'cache [i * k + x]'或者只是'cache [i]'。 –

回答

0

另外一組的所有工作項做寫作,以相同的共享內存地址cache[x]數據寫競爭條件(如Dithermaster說)和缺少barrier()函數,可以在修復之後添加一些優化:

內核中的第一個循環

 for (int x = 0; x < in_layer_s; x++) 
     { 
      cache[x] = weights[i*in_layer_s + x]; 
     } 

爲每個工作項掃描一個不同的內存區域,一次掃描一個元素。這在全局內存性能方面可能是錯誤的,因爲每個工作項在自己的循環中都可能使用相同的內存通道或甚至是相同的內存組,因此所有工作項都可以串行訪問該通道或庫。如果in_layer_s獲得更大的值,尤其是如果它的權力爲2,則情況會更糟糕。要解決此問題,所有工作項都應該與其鄰居訪問連續的地址。當工作項目統一訪問全局內存時,GPU可以更好地工作。在本地存儲器上,隨機訪問或工作項目之間的差距不太成問題。這就是爲什麼它建議使用統一的保存/加載全局,同時做本地隨機/分散/收集。在內核

 for (int x = 0; x < in_layer_s; x++) 
     { 
      total1 += cache[x] *input[x]; 
     } 

第二環路僅使用單個累加器。這是一個依賴鏈,需要在完成下一個循環之前完成每個循環。至少使用2個臨時「總」變量並展開循環。在這裏,如果in_layer_s足夠小,可以將input數組移動到本地或常量內存中以更快地訪問它(由於所有工作項都訪問相同的輸入數組,所有工作項都會重複)(可能是一半輸入到常量內存,另一半是到本地內存增加總帶寬)


weights[i*in_layer_s + x];是一個結構數組?如果是的話,你可以通過使其成爲數組的結構並完全擺脫第一個循環的優化來實現加速,主機端的代碼膨脹會增加,但如果優先級是速度,那麼數組的結構在gpu端更快且可讀。這也使得從主機端僅將必要的權重數據(SOA陣列)上傳到gpu成爲可能,進一步減少了總延遲(上傳+計算+下載)。


您也可以嘗試異步本地< - >全局轉移功能,使負載和計算重疊的每一個工作項組,隱藏更延遲作爲最後的手段。 https://www.khronos.org/registry/OpenCL/sdk/1.0/docs/man/xhtml/async_work_group_copy.html

+0

謝謝你的建議,我認爲它會提高速度。我只是不知道如何實現這一點,但我瞭解了更多關於opencl.Thanks – Ertryw