gpt4 book ai didi

assembly - 麻烦阅读和美国代码

转载 作者:行者123 更新时间:2023-12-02 22:03:20 26 4
gpt4 key购买 nike

我无法使用找到的文档阅读 AMD Southern Island GPU 的汇编语言 here .

这是一个示例 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 编译器生成的(dis)程序集:

 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 手册(参见第 12.6 节向量内存缓冲区指令)就写出了例如 *buffer_load_ubyte* 指令?
  • 我应该如何读取微码信息? (例如:*v_add_i32* 指令的第 161 页)
  • 是否有描述处理器架构的标准方法?

谢谢!

Ps:部分加分题在汇编代码注释中

最佳答案

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

了解 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 手册(请参阅第 12.6 节“向量内存缓冲区指令”)最终编写出一个 buffer_load_ubyte 指令?很难理解。您几乎需要通过示例来学习。

我应该如何读取微码信息?(例如:v_add_i32 指令的第 161 页)很难做到这一点。你几乎需要做一些试错。该手册实际上有几个地方不正确,我在论坛上向 AMD 报告了它。我构建了一个编译器 (asm4gcn),我不得不为此苦苦挣扎。我也引用了别人的项目。

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

关于assembly - 麻烦阅读和美国代码,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/16519269/

26 4 0
Copyright 2021 - 2024 cfsdn All Rights Reserved 蜀ICP备2022000587号
广告合作:1813099741@qq.com 6ren.com