gfx1201 Assembly For Dummies Part 1
Odd Rambles Masquerading As An Intro.
To live in the modern world is to experience capitalism viscerally. In your life, at least once you can expect the demi-god-bear that is capitalism to lift you to great heights of stock market ecstasy,and at least once you can expect to be dashed the rocks of some regulatory-capture-private-equity-backed bullshit. But, every once and a while capitalism holds each and every one of us in warm bear hug of soul-mate product.
If there is a consumer product that scratches one personal itch that you and a small number of relatively wealthy individuals have, the market will deliver. Converse with basset hounds on them? Cotton candy flavored toothpaste? For every jar there is a lid.
The "For Dummies" books are my soul-mate product. I like learning things, and I'm dumb. I feel seen, heard, maybe even cared about?
But somehow between "Sex for dummies" and "Existentialism for dummies" they never managed to write the guide on gfx1201 assembly?
Since capitalism couldn't provide, I must step in.
Here it is, finally...
gfx1201 Assembly For Dummies Part 1
NOTE: this blog is human authored, every word, 0% AI. Not because I'm a better writer than AI. Because our species has spent millions of years flapping our meat faces at each other and scratching out what we're thinking on rocks and paper and HTML documents. We, humans, have always built a shared understanding of world through our words. And I'm not here to buck that trend. When a human pretends that a robot's words are their own some great great ancestor looks out of Chauvet Cave and shakes their head in shame.
The code is pretty much all AI generated though, it's just faster than me, sorry. I also asked AI a bunch of questions while writing this, it has read every damn book on the planet, give me a break come on.
This is the first in a series of blog posts stepping through a basic add kernel written in gfx1201 assembly, you can find the complete code here. It's a very simple add kernel originally written in hip C++ then compiled to assembly and lightly modified.
The first post will just focus on understanding how the code gets packaged into something you can actually run, the code itself will be discussed in a future post. We have a .s file and we can only run a .hsaco file... Wait...
What The Heck Is HSACO?
I'm going to assume that you, like me, just want to run some code, well you have to pack it into a HSACO first.
HSACO stands for Heterogenous Systems Architecture Code Object.
HSA (the first three letters of HSACO) appears to have been a new open standard for how all of this was supposed to run at one time or another.
As far as I can tell, to understand gfx1201 assembly you don't need to know anything about HSA.
I'm not sure you ever need to know something about HSA?
There's an entire world of hardware and software that connects the code you want to run on the GPU to the "it's alive" moment when you program is actually doing things on the hardware. "The Runtime" sets things up on the host, it talks "The Driver" running in the linux kernel, which itself talks to firmware and the command processor on the actual GPU. HSA may be very important when you go spelunking into that stack, but I'm too scared to wander into that particular dungeon right now.
"The Runtime" and "The Driver" are all open source, check them out, be braver than me! ROCm/rocm-systems/tree/develop/projects/rocr-runtime, torvalds/linux/tree/master/drivers/gpu/drm/amd/amdgpu
But what is a HSACO really?
It's a regular ELF shared object file with some special nonsense packed in, that's it really.
Let's take a look at the first few lines of our vector_add.s kernel.
.amdgcn_target "amdgcn-amd-amdhsa--gfx1201"
.amdhsa_code_object_version 6
.text
.protected vector_add
.globl vector_add
.p2align 8
.type vector_add,@function
These are all assembler directives. To smash this stuff into a HSACO you need an assembler that understands the AMD specific directives. Here's what we're going to use for the smashing:
$ amdclang -x assembler -target amdgcn-amd-amdhsa -mcpu=gfx1201 kernels/vector_add.s -o thing.hsaco
Cool, smashed, what happened?
Regular Ass .text
The assembly is all loaded into a .text section, in both the .s and the ELF file that gets spit out, that's exactly how it would work with x86 or ARM assembly too.
.text
.protected vector_add
.globl vector_add
.p2align 8
.type vector_add,@function
vector_add:
s_load_b32 s2, s[0:1], 0x2c
s_load_b32 s3, s[0:1], 0x18
s_wait_kmcnt 0x0
The lines without a dot are real instructions.
The ones with dots are assembler directives.
.p2align is an assembler directive telling amdclang how to pack the HSACO, it won't end up in the HSACO.
vector_add is a label, it's not a directive and not an instruction, it's a marker, instructions can jump to it, and it's also used as the functions entry point (a call is basically a jump change my mind).
The .globl vector_add directive turns vector_add into a global symbol the linker can advertize.
The s_load_b32 is a real instruction and will end up as physical bytes in the artifact.
The .text section makes the most sense, it's pretty familiar if you've worked with any other flavor of assembly.
These are your instructions and they will end up in the .text section of your ELF file.
A few other sections are a little more surprising.
KD
.section .rodata,"a",@progbits
.p2align 6, 0x0
.amdhsa_kernel vector_add
.amdhsa_group_segment_fixed_size 0
.amdhsa_private_segment_fixed_size 0
.amdhsa_kernarg_size 288
.amdhsa_user_sgpr_count 2
.amdhsa_user_sgpr_dispatch_ptr 0
.amdhsa_user_sgpr_queue_ptr 0
.amdhsa_user_sgpr_kernarg_segment_ptr 1
.amdhsa_user_sgpr_dispatch_id 0
.amdhsa_user_sgpr_private_segment_size 0
.amdhsa_wavefront_size32 1
.amdhsa_uses_dynamic_stack 0
.amdhsa_enable_private_segment 0
.amdhsa_system_sgpr_workgroup_id_x 1
.amdhsa_system_sgpr_workgroup_id_y 0
.amdhsa_system_sgpr_workgroup_id_z 0
.amdhsa_system_sgpr_workgroup_info 0
.amdhsa_system_vgpr_workitem_id 0
.amdhsa_next_free_vgpr 6
.amdhsa_next_free_sgpr 8
.amdhsa_reserve_vcc 1
.amdhsa_float_round_mode_32 0
.amdhsa_float_round_mode_16_64 0
.amdhsa_float_denorm_mode_32 3
.amdhsa_float_denorm_mode_16_64 3
.amdhsa_fp16_overflow 0
.amdhsa_workgroup_processor_mode 1
.amdhsa_memory_ordered 1
.amdhsa_forward_progress 1
.amdhsa_inst_pref_size 2
.amdhsa_round_robin_scheduling 0
.amdhsa_exception_fp_ieee_invalid_op 0
.amdhsa_exception_fp_denorm_src 0
.amdhsa_exception_fp_ieee_div_zero 0
.amdhsa_exception_fp_ieee_overflow 0
.amdhsa_exception_fp_ieee_underflow 0
.amdhsa_exception_fp_ieee_inexact 0
.amdhsa_exception_int_div_zero 0
.end_amdhsa_kernel
This builds a Kernel Descriptor, you can see a whole bunch of .amdhsa_* directives that set up GPU specific registers.
Most of this is much lower level than we're going to get into, but roughly a kernel descriptor as the blob of stuff the GPU needs to actually launch a kernel.
Something has to tell the GPU things like where the actual code is, and how many registers that program needs; there's only so many registers on each Compute Unit, so the GPU decides how your program maps to CUs based partially on register requirements.
A KD is loaded into GPU memory and the host points at it when it wants to launch the kernel the KD is built for.
At a high level you can think of this block as just .amdhsa_kernel vector_add + compiler derived information about your kernel.
The .amdhsa_kernel directive ensures that the KD will contain a pointer to the physical bytes of the vector_add kernel.
The assembler writes the KD into the ELF's .rodata section (read only data), and creates symbol vector_add.kd so that linkers and loaders can find it.
"The Runtime" shuffles the KD over to the GPU (by black magic or pony express as far ask I know), and the host points at the KD on the GPU when it wants it to, you know, vector_add some stuff.
That's basically a KD: where the code is + blob of other stuff the GPU needs; it's everything the GPU needs to do the thing, and it needs to be loaded onto the GPU at runtime.
Weirdness
.text
.Lfunc_end0:
.size vector_add, .Lfunc_end0-vector_add
.p2alignl 7, 3214868480
.fill 96, 4, 3214868480
The next bit little bit is odd.
The .Lfunc_end0 is actually a label for the end of vector_add (we're back writing into the .text section because of the .text directive).
And the .size directive is there to do a little math to figure out how big the vector_add is as bytes of physical code; it also tells the assembler to update the vector_add symbol with this information.
The other directives tell the assembler to pad out our program with a bunch of something so it fits on alignment boundaries — that's probably important because the GPU can only fetch blocks of whatever code size IDK.
There's always a bunch of fiddly bits when you're dealing with hardware, this is one of them. The compiler seems to have fiddled it properly, cool beans.
Host Stuff
.amdgpu_metadata
---
amdhsa.kernels:
- .args:
- .address_space: global
.offset: 0
.size: 8
.value_kind: global_buffer
- .address_space: global
.offset: 8
.size: 8
.value_kind: global_buffer
- .address_space: global
.offset: 16
.size: 8
.value_kind: global_buffer
- .offset: 24
.size: 4
.value_kind: by_value
- .offset: 32
.size: 4
.value_kind: hidden_block_count_x
- .offset: 36
.size: 4
.value_kind: hidden_block_count_y
- .offset: 40
.size: 4
.value_kind: hidden_block_count_z
- .offset: 44
.size: 2
.value_kind: hidden_group_size_x
- .offset: 46
.size: 2
.value_kind: hidden_group_size_y
- .offset: 48
.size: 2
.value_kind: hidden_group_size_z
- .offset: 50
.size: 2
.value_kind: hidden_remainder_x
- .offset: 52
.size: 2
.value_kind: hidden_remainder_y
- .offset: 54
.size: 2
.value_kind: hidden_remainder_z
- .offset: 72
.size: 8
.value_kind: hidden_global_offset_x
- .offset: 80
.size: 8
.value_kind: hidden_global_offset_y
- .offset: 88
.size: 8
.value_kind: hidden_global_offset_z
- .offset: 96
.size: 2
.value_kind: hidden_grid_dims
.gfx1250_revision: B0
.group_segment_fixed_size: 0
.kernarg_segment_align: 8
.kernarg_segment_size: 288
.language: OpenCL C
.language_version:
- 2
- 0
.max_flat_workgroup_size: 1024
.name: vector_add
.private_segment_fixed_size: 0
.sgpr_count: 10
.sgpr_spill_count: 0
.symbol: vector_add.kd
.uniform_work_group_size: 1
.uses_dynamic_stack: false
.vgpr_count: 6
.vgpr_spill_count: 0
.wavefront_size: 32
.workgroup_processor_mode: 1
amdhsa.target: amdgcn-amd-amdhsa--gfx1201
amdhsa.version:
- 1
- 2
...
.end_amdgpu_metadata
This is actually a MessagePack that gets displayed in the ASM file like this.
"The Runtime" needs to know how kernel expects memory for its arguments to be packed.
The vector_add kennel will be passed a pointer to some blob of memory, it reads offsets into that memory for things like ARG1 and ARG2.
"The Runtime" packs the arguments you send it on the host into said blob then does some sort of incantation involving the KD from above to get stuff running.
This section exists so "The Runtime" knows how to pack that blob.
This whole section gets packed into a .note section in the ELF file that "The Runtime" can go grab when it needs to get of its lazy ass and actually run some code.
Yep, Well, That's About It
To learn anything about assembly you first need to be able to look at an actual assembly file that you can actually run.
I hope this helped.
A gfx1201 assembly file is pretty odd, but there's always three blobs you can expect: a host side blob with everything the runtime needs to know, a GPU side blob with everything the GPU needs to know, and your actual assembly (which is the interesting bit).
Stay tuned for next time where the spicy drama that is gfx1201 assembly kicks into high gear.