有效地从私有内存复制到本地内存

我想将私有数据块有效地复制到本地内存。每个工作项目需要复制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)

来源

2017-05-24 HyperZ

您可以使用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可能并不多。

来源

2017-05-25 08:32:33

将此应用于我的内核不起作用。但是，它编译时在运行时会出现分段错误。我认为问题是你的解决方案写入本地数组，而我需要写入本地缓冲区？ – HyperZ

如果vector4版本不起作用，那么它必须是“缓冲区”与“数组”问题。 –

谢谢。由于我的本地缓冲区的大小是3的倍数，所以我分配了一个更大的大小，这样我就可以从它上面执行vload4，而无需访问它。但是，这仍然会崩溃，因此它必须与缓冲区问题有关。 – HyperZ

内核中的模式是编译器优化代码的明显候选者，特别是因为循环的上限定义为常量值。功能async_work_group_copy只是复制全球内存当地和之间，而不是私人。其主要目的是隐藏延迟（即在执行内存操作时执行其他操作）。在内核中调用“异步”函数，然后立即等待，可能会像没有函数一样慢。

来源

2017-05-30 08:05:16 Ruyk

有效地从私有内存复制到本地内存

回答

相关问题