有效地從私有內存複製到本地內存

Question

我想將私有數據塊有效地複製到本地內存。每個工作項目需要複製3個字節的本地存儲器，所以目前我做的：

__kernel void apply(__local uchar* lmem) { 
    int offset = ...; 
    uchar data[3] = {1,2,3}; // just an example, is passed as an argument in fact 
    for (int j = 0; j < 3; j++) { 
     lmem[offset+j] = data[j]; 
    }  
} 

但是，我有一種感覺，這可以更有效的進行（因爲我需要寫連續3個字節）。因此，我嘗試使用memcpy：

__kernel void apply(__local uchar* lmem) { 
    int offset = ...; 
    uchar data[3] = {1,2,3}; 
    memcpy(&lmem[offset], data, 3); 
} 

不過，我得到以下錯誤消息：

__kernel void apply(__local uchar* lmem) { 
    int offset = ...; 
    uchar data[3] = {1,2,3} 
    event_t evt = async_work_group_copy((local uchar*) &lmem[offset], (uchar*) data, 3, 0); 
    wait_group_events(3, &evt); 
} 

這將導致錯誤： error: passing '__local uchar *' (aka '__local unsigned char *') to parameter of type 'void *' changes address space of pointer

我也使用async_work_group_copy試圖 note: candidate function not viable: 2nd argument ('uchar *' (aka 'unsigned char *')) is in address space 0, but parameter must be in address space 16776960

有沒有辦法讓e ach工作項能夠有效地將這3個字節從私有存儲器複製到本地存儲器，而不是一個接一個地複製這3個字節？

---

編輯：這是我如何分配的本地內存並將其傳遞給內核：

import pyopencl as cl 
... 
program = ... 
lmem = cl.LocalMemory(needed_size) 

applyKernel = program.apply 
applyKernel.set_scalar_arg_dtypes([None]) # None because it is a buffer 
applyKernel(queue, global_size, local_size, lmem) 

Answer 1

您可以使用vload3從私人數組加載（或任何記憶符它），然後使用vstore3存儲到本地陣列或者甚至全球陣列

__kernel void test(__global unsigned char * data) 
{ 
    int i=get_global_id(0); 
    int l=get_local_id(0); 
    unsigned char values[30]; 
    values[0]=1; 
    values[1]=2; 
    values[2]=3; 
    __local unsigned char testLocalArray[3*256]; 
    vstore3(vload3(0,values),l,testLocalArray); 
    barrier(CLK_LOCAL_MEM_FENCE); 
    data[i*3]=testLocalArray[l*3]; 
    data[i*3+1]=testLocalArray[l*3+1]; 
    data[i*3+2]=testLocalArray[l*3+2]; 
} 

這將產生1,2,3,1,2,3,1,2,3,1,2,3的數據陣列。但是一些硬件可能與vector3不完全兼容，所以它實際上可能會加載和存儲vector4，並會給出錯誤或錯誤的結果，除非您將每個塊填充1個字節。

由於vload或vstore是單指令，因此如果編譯器尚未自動執行此操作，它應該使用支持的單個指令所支持的任何硬件功能。

其ISA代碼R7-240 GPU：

s_mov_b32  m0, 0x00008000        // 00000000: BEFC03FF 00008000 
    s_buffer_load_dword s0, s[8:11], 0x04     // 00000008: C2000904 
    s_buffer_load_dword s1, s[8:11], 0x18     // 0000000C: C2008918 
    s_waitcnt  lgkmcnt(0)         // 00000010: BF8C007F 
    s_min_u32  s0, s0, 0x0000ffff       // 00000014: 8380FF00 0000FFFF 
    s_mul_i32  s0, s16, s0         // 0000001C: 93000010 
    v_mul_u32_u24 v1, v0, 3         // 00000020: D2160001 00010700 
    s_add_u32  s0, s0, s1         // 00000028: 80000100 
    v_mov_b32  v2, 1          // 0000002C: 7E040281 
    s_buffer_load_dword s1, s[12:15], 0x00     // 00000030: C2008D00 
    v_add_i32  v0, vcc, s0, v0        // 00000034: 4A000000 
    v_mov_b32  v3, 2          // 00000038: 7E060282 
    v_mov_b32  v4, 3          // 0000003C: 7E080283 
    v_mul_lo_i32 v0, v0, 3         // 00000040: D2D60000 00010700 
    ds_write_b8 v1, v2          // 00000048: D8780000 00000201 
    ds_write_b8 v1, v3 offset:1        // 00000050: D8780001 00000301 
    ds_write_b8 v1, v4 offset:2        // 00000058: D8780002 00000401 
    s_waitcnt  lgkmcnt(0)         // 00000060: BF8C007F 
    v_add_i32  v0, vcc, s1, v0        // 00000064: 4A000001 
    s_barrier             // 00000068: BF8A0000 
    ds_read_u8 v2, v1          // 0000006C: D8E80000 02000001 
    ds_read_u8 v3, v1 offset:1        // 00000074: D8E80001 03000001 
    ds_read_u8 v1, v1 offset:2        // 0000007C: D8E80002 01000001 
    s_waitcnt  lgkmcnt(2)         // 00000084: BF8C027F 
    v_bfe_u32  v2, v2, 0, 8        // 00000088: D2900002 02210102 
    s_waitcnt  lgkmcnt(1)         // 00000090: BF8C017F 
    v_bfe_u32  v3, v3, 0, 8        // 00000094: D2900003 02210103 
    s_waitcnt  lgkmcnt(0)         // 0000009C: BF8C007F 
    v_bfe_u32  v1, v1, 0, 8        // 000000A0: D2900001 02210101 
    buffer_store_byte v2, v0, s[4:7], 0 offen glc   // 000000A8: E0605000 80010200 
    buffer_store_byte v3, v0, s[4:7], 0 offen offset:1 glc // 000000B0: E0605001 80010300 
    buffer_store_byte v1, v0, s[4:7], 0 offen offset:2 glc // 000000B8: E0605002 80010100 

貌似現場仍然落後3指令。

對於RX550 GPU：

// 
    // &__OpenCL_test_kernel: 
    // 
    s_load_dword s0, s[4:5], 0x04      // 000000000100: C0020002 00000004 
    s_mov_b32  m0, 0x00010000       // 000000000108: BEFC00FF 00010000 
    s_waitcnt  lgkmcnt(0)        // 000000000110: BF8C007F 
    s_and_b32  s0, s0, 0x0000ffff      // 000000000114: 8600FF00 0000FFFF 
    s_mul_i32  s0, s0, s8        // 00000000011C: 92000800 
    s_load_dwordx2 s[2:3], s[6:7], 0x00     // 000000000120: C0060083 00000000 
    s_load_dwordx2 s[4:5], s[6:7], 0x30     // 000000000128: C0060103 00000030 
    v_mul_i32_i24 v1, v0, 3        // 000000000130: D1060001 00010700 
    v_mov_b32  v2, 1         // 000000000138: 7E040281 
    ds_write_b8 v1, v2         // 00000000013C: D83C0000 00000201 
    v_mov_b32  v2, 2         // 000000000144: 7E040282 
    ds_write_b8 v1, v2 offset:1       // 000000000148: D83C0001 00000201 
    v_mov_b32  v2, 3         // 000000000150: 7E040283 
    ds_write_b8 v1, v2 offset:2       // 000000000154: D83C0002 00000201 
    s_waitcnt  lgkmcnt(0)        // 00000000015C: BF8C007F 
    s_add_u32  s0, s0, s2        // 000000000160: 80000200 
    v_add_u32  v0, vcc, s0, v0       // 000000000164: 32000000 
    v_mul_lo_u32 v0, v0, 3        // 000000000168: D2850000 00010700 
    v_ashrrev_i32 v2, 31, v0        // 000000000170: 2204009F 
    v_add_u32  v9, vcc, s4, v0       // 000000000174: 32120004 
    v_mov_b32  v3, s5         // 000000000178: 7E060205 
    v_addc_u32 v10, vcc, v3, v2, vcc     // 00000000017C: 38140503 
    s_barrier            // 000000000180: BF8A0000 
    ds_read_u8 v5, v1         // 000000000184: D8740000 05000001 
    ds_read_u8 v6, v1 offset:1       // 00000000018C: D8740001 06000001 
    ds_read_u8 v1, v1 offset:2       // 000000000194: D8740002 01000001 
    v_add_u32  v3, vcc, v9, 1       // 00000000019C: D1196A03 00010309 
    v_addc_u32 v4, vcc, v10, 0, vcc     // 0000000001A4: D11C6A04 01A9010A 
    v_add_u32  v7, vcc, v9, 2       // 0000000001AC: D1196A07 00010509 
    v_addc_u32 v8, vcc, v10, 0, vcc     // 0000000001B4: D11C6A08 01A9010A 
    s_waitcnt  lgkmcnt(2)        // 0000000001BC: BF8C027F 
    flat_store_byte v[9:10], v5       // 0000000001C0: DC600000 00000509 
    s_waitcnt  lgkmcnt(2)        // 0000000001C8: BF8C027F 
    flat_store_byte v[3:4], v6       // 0000000001CC: DC600000 00000603 
    s_waitcnt  lgkmcnt(2)        // 0000000001D4: BF8C027F 
    flat_store_byte v[7:8], v1       // 0000000001D8: DC600000 00000107 
    s_endpgm 

這是比其他GPU結果有所不同，但仍然每VLOAD或VSTORE 3點的指令。也許它的vload和vstore更快。

唯一的優勢可能是缺少循環計數器。這可能會給硬件上的整數標量單元更多的空間來計算其他東西，這些東西肯定比循環版本更好。

這是同樣的GPU的循環版本：

s_load_dword s0, s[4:5], 0x04      // 000000000100: C0020002 00000004 
    s_mov_b32  m0, 0x00010000       // 000000000108: BEFC00FF 00010000 
    s_waitcnt  lgkmcnt(0)        // 000000000110: BF8C007F 
    s_and_b32  s0, s0, 0x0000ffff      // 000000000114: 8600FF00 0000FFFF 
    s_mul_i32  s0, s0, s8        // 00000000011C: 92000800 
    s_load_dwordx2 s[2:3], s[6:7], 0x00     // 000000000120: C0060083 00000000 
    s_waitcnt  lgkmcnt(0)        // 000000000128: BF8C007F 
    s_add_u32  s0, s0, s2        // 00000000012C: 80000200 
    s_load_dwordx2 s[2:3], s[6:7], 0x30     // 000000000130: C0060083 00000030 
    v_mul_i32_i24 v1, v0, 3        // 000000000138: D1060001 00010700 
    v_mov_b32  v2, 1         // 000000000140: 7E040281 
    v_add_u32  v0, vcc, s0, v0       // 000000000144: 32000000 
    v_mov_b32  v3, 2         // 000000000148: 7E060282 
    v_mul_lo_u32 v0, v0, 3        // 00000000014C: D2850000 00010700 
    v_mov_b32  v4, 3         // 000000000154: 7E080283 
    ds_write_b8 v1, v2         // 000000000158: D83C0000 00000201 
    ds_write_b8 v1, v3 offset:1       // 000000000160: D83C0001 00000301 
    ds_write_b8 v1, v4 offset:2       // 000000000168: D83C0002 00000401 
    v_ashrrev_i32 v2, 31, v0        // 000000000170: 2204009F 
    s_waitcnt  lgkmcnt(0)        // 000000000174: BF8C007F 
    v_add_u32  v9, vcc, s2, v0       // 000000000178: 32120002 
    v_mov_b32  v5, s3         // 00000000017C: 7E0A0203 
    v_addc_u32 v10, vcc, v5, v2, vcc     // 000000000180: 38140505 
    ds_write_b8 v1, v3 offset:1       // 000000000184: D83C0001 00000301 
    ds_write_b8 v1, v4 offset:2       // 00000000018C: D83C0002 00000401 
    s_waitcnt  lgkmcnt(0)        // 000000000194: BF8C007F 
    s_barrier            // 000000000198: BF8A0000 
    ds_read_u8 v5, v1         // 00000000019C: D8740000 05000001 
    ds_read_u8 v6, v1 offset:1       // 0000000001A4: D8740001 06000001 
    ds_read_u8 v1, v1 offset:2       // 0000000001AC: D8740002 01000001 
    v_add_u32  v3, vcc, v9, 1       // 0000000001B4: D1196A03 00010309 
    v_addc_u32 v4, vcc, v10, 0, vcc     // 0000000001BC: D11C6A04 01A9010A 
    v_add_u32  v7, vcc, v9, 2       // 0000000001C4: D1196A07 00010509 
    v_addc_u32 v8, vcc, v10, 0, vcc     // 0000000001CC: D11C6A08 01A9010A 
    s_waitcnt  lgkmcnt(2)        // 0000000001D4: BF8C027F 
    flat_store_byte v[9:10], v5       // 0000000001D8: DC600000 00000509 
    s_waitcnt  lgkmcnt(2)        // 0000000001E0: BF8C027F 
    flat_store_byte v[3:4], v6       // 0000000001E4: DC600000 00000603 
    s_waitcnt  lgkmcnt(2)        // 0000000001EC: BF8C027F 
    flat_store_byte v[7:8], v1       // 0000000001F0: DC600000 00000107 
    s_endpgm 

我不能在這裏找到循環計數器相關的指令，編譯器可能已經認識到這兩個vloadn和循環版本的模式，併產生相同的機器代碼。但這只是ISA，我不能說核心中真正發生了什麼。也許VLIW獲得更多，CPU收益更多，但最新的GPU可能並不多。

Answer 2

內核中的模式是編譯器優化代碼的明顯候選者，特別是因爲循環的上限定義爲常量值。 功能async_work_group_copy只是複製全球內存當地和之間，而不是私人。其主要目的是隱藏延遲（即在執行內存操作時執行其他操作）。在內核中調用「異步」函數，然後立即等待，可能會像沒有函數一樣慢。