2

我无法使用此处找到的文档阅读 AMD 南岛 GPU 的汇编语言。

这是一个示例 OpenCL 代码：

 1    __attribute__((reqd_work_group_size(256, 1, 1)))
 2    void kernel foo(global uchar* data) {
 3        const uint block_size = get_local_size(0);
 4        const uint lid = get_local_id(0);
 5        
 6        data[lid] = max(data[lid], data[lid + block_size]);
 7    }

以及 AMD 编译器生成的（反）汇编：

 1    shader main
 2      asic(SI_ASIC)
 3      type(CS)
 4    
 5      s_buffer_load_dword  s0, s[8:11], 0x00                 // what is the purpose of s[8:11] ?
 6      s_waitcnt     lgkmcnt(0)
 7      v_add_i32     v1, vcc, s0, v0                          // I guess v0 initially contains the local IDs ?
 8      v_add_i32     v0, vcc, s0, v0                          // wouldn't a v_mov v0, v1 performs better ?
 9      buffer_load_ubyte  v2, v1, s[4:7], 0 offen             // s[4:7] ?
10      buffer_load_ubyte  v0, v0, s[4:7], 0 offen offset:256
11      s_waitcnt     vmcnt(0)
12      v_max_u32     v0, v2, v0
13      buffer_store_byte  v0, v1, s[4:7], 0 offen glc
14      s_endpgm
15    end
  • 我不明白的是，我怎么可能最终编写了一条只有 ISA 手册的 *buffer_load_ubyte* 指令（请参阅第 12.6 节矢量内存缓冲区指令）？
  • 我应该如何阅读微码信息？（例如：*v_add_i32* 指令的第 161 页）
  • 是否有任何标准方法来描述处理器架构？

谢谢 ！

Ps：一些额外的问题在汇编代码注释中

4

2 回答 2

3

我迟到了大约一年，但也许这对其他人有帮助。

了解 AMD GCN 都归结为：

Address = BASE + offset + lane

在内核启动时，会预加载以下寄存器

  • s[4:7] 是基地址
  • s[8:11] 是指向参数的指针
  • 偏移量 v0 预加载车道号 (0-63)

简而言之，我们将 (1) 获取地址 *data (2) 获取 *data 处的值 (3) 获取 offset:256 处的值 (4) 将两个数字相加 (5) 将结果存储到*数据。

shader main
  asic(SI_ASIC)
  type(CS)

  s_buffer_load_dword  s0, s[8:11], 0x00  // s[8:11] is the pointer to the params
  s_waitcnt     lgkmcnt(0)           // wait for s0 to be filled
  v_add_i32     v1, vcc, s0, v0      // s0=offset v0=lane We just need the base now.
  v_add_i32     v0, vcc, s0, v0      // wouldn't a v_mov v0, v1 performs better
  buffer_load_ubyte v2, v1, s[4:7], 0 offen //Get value at Base(s[4:7]) + v1(offset & lane) 
  buffer_load_ubyte v0, v0, s[4:7], 0 offen offset:256 // like above but address+256
  s_waitcnt     vmcnt(0)                        //wait for the memory transfer to complete
  v_max_u32     v0, v2, v0                      // do the MAX operation
  buffer_store_byte v0, v1, s[4:7], 0 offen glc //save v0 using the base+v1(offset+lane)
  s_endpgm                                      //stop kernel
end

在第 10 行使用 "v_mov v0, v1" 而不是 "v_add_i32 v0, vcc, s0, v0" 不会让它更快，因为 v_mov 和 v_add 都需要相同的时间。但是，如果这是在 CPU 上，则 v_mov 会更慢，因为它取决于先前的指令并且它不能一次执行多个指令。gpu 不能在一个核心上同时执行多条指令，因此两种方式的速度相同。

我不明白的是，我怎么可能最终编写了一个只有 ISA 手册的 buffer_load_ubyte 指令（参见第 12.6 节向量内存缓冲区指令）？这有点难以理解。你几乎需要通过例子来学习。

我应该如何阅读微码信息？（例如：第 161 页的v_add_i32指令）很难做到这一点。你几乎需要做一些跟踪和错误。该手册实际上在一些地方是不正确的，我在论坛上向 AMD 报告了它。我构建了一个编译器（asm4gcn），我不得不为此苦苦挣扎。我也参考了其他人的项目。

是否有任何标准的方式来描述处理器架构？我想它们在某些方面都是不同的。但大多数都有缓存部分、浮点单元、控制流单元和寄存器。

于 2015-07-16T06:22:46.790 回答
2

在这里为全世界发布。这是我正在上课的一些 SI 组件。帮助了解 ABI 发生了什么，假定此代码在 Multi2sim 之上运行。

.global vector_add
.args
    i32* src1 0 uav10 RO
    i32* src2 16 uav11 RO
    i32* dst 32 uav12 RW
.metadata
    uavprivate = 0
    hwregion = 0
    hwlocal = 0
    FloatMode = 192
    IeeeMode = 0
    # Loads UAV table in s2...s3
    userElements[0] = PTR_UAV_TABLE, 0, s[2:3]
    # Loads constant buffer 0 descriptor in s4...s7
    userElements[1] = IMM_CONST_BUFFER, 0, s[4:7]
    # Loads constant buffer 1 descriptor in s8...s11
    userElements[2] = IMM_CONST_BUFFER, 1, s[8:11]
    # Forces wg_id[0] (work-group ID in dimension 0) to be available in s12
    COMPUTE_PGM_RSRC2:USER_SGPR = 12
    COMPUTE_PGM_RSRC2:TGID_X_EN = 1
.text
    # Load lsize[0] into s0
    s_buffer_load_dword s0, s[4:7], 0x04
    # Load src1, src2, and dst base addresses (arguments) from CB1
    s_buffer_load_dword s4, s[8:11], 0x00
    s_buffer_load_dword s5, s[8:11], 0x04
    s_buffer_load_dword s6, s[8:11], 0x08
    # Load UAVs from UAV table
    s_load_dwordx4 s[20:23], s[2:3], 0x50
    s_load_dwordx4 s[24:27], s[2:3], 0x58
    s_load_dwordx4 s[28:31], s[2:3], 0x60
    # Waits for memory operations to complete
    s_waitcnt lgkmcnt(0)
    # v1 <= lsize[0]
    v_mov_b32 v1, s0
    # v1 <= lsize[0] * wg_id[0]
    v_mul_i32_i24 v1, s12, v1
    # v2 <= lsize[0] * wg_id[0] + lid[0] = gid[0]
    v_add_i32 v2, vcc, v0, v1
    # v3 <= gid[0] * 4
    v_lshlrev_b32 v3, 2, v2
    # Calcaulte effective addresses
    v_add_i32 v10, vcc, s4, v3
    v_add_i32 v11, vcc, s5, v3
    v_add_i32 v12, vcc, s6, v3
    # Load src1[id] and src2[id]
    tbuffer_load_format_x v20, v10, s[20:23], 0 offen format:[BUF_DATA_FORMAT_32, BUF_NUM_FORMAT_FLOAT]
    tbuffer_load_format_x v21, v11, s[24:27], 0 offen format:[BUF_DATA_FORMAT_32,BUF_NUM_FORMAT_FLOAT]
    # Waits for memory operations to complete
    s_waitcnt vmcnt(0)
    # Add source elements
    v_add_i32 v22, vcc, v20, v21
    # Store result in dst[id]
    tbuffer_store_format_x v22, v12, s[28:31], 0 offen format:[BUF_DATA_FORMAT_32,BUF_NUM_FORMAT_FLOAT]
    # End program
    s_endpgm

这是一个简单的向量加法内核，需要 3 个参数。src1,src2和dst.

于 2014-06-08T21:59:23.610 回答