GCNを読む

はじめに

現在AMDの主流GPUはRDNAアーキテクチャのものですが、今回取り上げるのは旧世代のGCNアーキテクチャです。命令セット自体は(その一つ昔のTeraScaleに比べれば)似通っているので見てみるのも悪くないと思います。
次の順で書いていきます。

• OpenCL Cでカーネルを記述、コンパイルし、バイナリ列を保存する
• バイナリ列をディスアセンブル、およびその結果をアセンブルする
• アセンブル結果を期待通り実行できることを確認する
• ディスアセンブル結果を読む

OpenCL Cでカーネルを記述、コンパイルし、バイナリ列を保存する

__attribute__((reqd_work_group_size(2, 2, 1)))
kernel void Example(global float *dst, uint a, global float *b)
{
	size_t gid = get_global_id(0);
	uint lid_x = (uint)get_local_id(0);
	uint lid_y = (uint)get_local_id(1);

	dst[gid] = (lid_x * a) + (lid_y * (*b));
}

global_work_sizeを(2, 2)として2*2要素の処理を行い、
dst[0] := 0;
dst[1] := a;
dst[2] := b;
dst[3] := a + b;
という結果となることを期待します。

次のようなプログラムを用いてバイナリ列を保存します。

この方法で得られるデータのフォーマットはCode object V2と呼ばれる古いもので、現在LLVMでは生成しないようです。
Radeon GPU Analyzer version 2.7.1.9のclang version 16.0.0コンパイル出力ではAMDHSA Code Object V4となるようです。このV4をOpenCL 2.0 AMD-APP (3584.0)でclCreateProgramWithBinaryしても、その後のclBuildProgramでCL_BUILD_PROGRAM_FAILURE, Error while BRIG Codegen phase: the binary is incompleteを返すので動かせません。

プログラムのソースコード

#include <exception>
#include <format>
#include <fstream>
#include <iostream>
#include <vector>

#ifdef _WIN32
#include <fcntl.h>
#include <io.h>
#endif

#define CL_HPP_TARGET_OPENCL_VERSION 200
#define CL_HPP_MINIMUM_OPENCL_VERSION 200
#define CL_HPP_ENABLE_EXCEPTIONS
#include <CL/opencl.hpp>

void PrintDevices();
void Compile(std::vector<unsigned char>& binary, size_t platform_idx, size_t device_idx, const char *file_path);

int main(int argc, char *argv[])
{
#ifdef _WIN32
	if (-1 == _setmode(_fileno(stdout), _O_BINARY))
	{
		std::cerr << "Cannot set stdout to binary mode." << std::endl;
		return -1;
	}
#endif

	if (argc < 4)
	{
		PrintDevices();
		std::cout << argv[0] << " Platform-Index Device-Index Source-File-Path" << std::endl;
		return 0;
	}

	std::vector<unsigned char> binary;

	try
	{
		Compile(binary, atoi(argv[1]), atoi(argv[2]), argv[3]);
	}
	catch (cl::BuildError& e)
	{
		std::cerr << e.getBuildLog().at(0).second.c_str() << std::endl;
		return -1;
	}
	catch (std::exception& e)
	{
		std::cerr << e.what() << std::endl;
		return -1;
	}
	catch (...)
	{
		return -1;
	}

	std::cout.write(reinterpret_cast<char *>(binary.data()), binary.size());
	std::cout.flush();
	return 0;
}

void PrintDevices()
{
	std::vector<cl::Platform> platforms;
	cl::Platform::get(&platforms);

	std::cout << "Platform-Index, Device-Index: Device-Name" << std::endl;

	for (int plat_index=0; plat_index<platforms.size(); ++plat_index)
	{
		std::vector<cl::Device> devices;

		platforms[plat_index].getDevices(CL_DEVICE_TYPE_ALL, &devices);

		for (int dev_index=0; dev_index<devices.size(); ++dev_index)
		{
			std::string device_name = devices[dev_index].getInfo<CL_DEVICE_NAME>();
			std::cout << std::format("{}, {}: {}", plat_index, dev_index, device_name) << std::endl;
		}
	}
}

void Compile(std::vector<unsigned char>& binary, size_t platform_idx, size_t device_idx, const char *file_path)
{
	std::string program_string;
	{
		std::ifstream ifs(file_path);
		program_string = std::string((std::istreambuf_iterator<char>(ifs)), std::istreambuf_iterator<char>());
	}

	std::vector<cl::Platform> platforms;
	cl::Platform::get(&platforms);

	std::vector<cl::Device> devices;
	platforms.at(platform_idx).getDevices(CL_DEVICE_TYPE_ALL, &devices);

	cl::Context ctx(devices.at(device_idx));
	cl::Program program(ctx, program_string, false);
	program.build(devices.at(device_idx), "-g -cl-std=CL2.0");
	binary = program.getInfo<CL_PROGRAM_BINARIES>().at(0);
}

バイナリ列をディスアセンブル、およびその結果をアセンブルする

CLRadeonExtender を使います。
clrxdisasm -aC example.bin > example.asm
clrxasm -o reasm_example.bin example.asm
とでもすればよいでしょう。
example.binに比べてreasm_example.binはファイスサイズが6kBほど小さくなっていますが、気にしないでおきます。

アセンブル結果を期待通り実行できることを確認する

次のようなプログラムを用いて確認します。

プログラムのソースコード

#include <cstdio>
#include <cstdlib>
#include <exception>
#include <filesystem>
#include <format>
#include <fstream>
#include <iostream>
#include <vector>

#define CL_HPP_TARGET_OPENCL_VERSION 200
#define CL_HPP_MINIMUM_OPENCL_VERSION 200
#define CL_HPP_ENABLE_EXCEPTIONS
#include <CL/opencl.hpp>

void PrintDevices();
void RunExampleKernel(size_t platform_idx, size_t device_idx, const char *file_path);

int main(int argc, char *argv[])
{
	if (argc < 4)
	{
		PrintDevices();
		std::cout << argv[0] << " Platform-Index Device-Index Binary-File-Path" << std::endl;
		return 0;
	}

	try
	{
		RunExampleKernel(std::atoi(argv[1]), std::atoi(argv[2]), argv[3]);
	}
	catch (std::exception& e)
	{
		std::cerr << e.what() << std::endl;
		return -1;
	}
	catch (...)
	{
		std::cerr << "unknown error." << std::endl;
		return -1;
	}
}

void PrintDevices()
{
	std::vector<cl::Platform> platforms;
	cl::Platform::get(&platforms);

	std::cout << "Platform-Index, Device-Index: Device-Name" << std::endl;

	for (int plat_index=0; plat_index<platforms.size(); ++plat_index)
	{
		std::vector<cl::Device> devices;

		platforms[plat_index].getDevices(CL_DEVICE_TYPE_ALL, &devices);

		for (int dev_index=0; dev_index<devices.size(); ++dev_index)
		{
			std::string device_name = devices[dev_index].getInfo<CL_DEVICE_NAME>();
			std::cout << std::format("{}, {}: {}", plat_index, dev_index, device_name) << std::endl;
		}
	}
}

void RunExampleKernel(size_t platform_idx, size_t device_idx, const char *file_path)
{
	cl::Program::Binaries binaries;
	{
		std::vector<unsigned char> d;

		int filesz = std::filesystem::file_size(file_path);
		FILE *f = std::fopen(file_path, "rb");
		d.resize(filesz);
		std::fread(d.data(), filesz, 1, f);
		std::fclose(f);
		binaries.push_back(d);
	}

	std::vector<cl::Platform> platforms;
	cl::Platform::get(&platforms);

	std::vector<cl::Device> devices;
	platforms.at(platform_idx).getDevices(CL_DEVICE_TYPE_ALL, &devices);

	cl::Context ctx(devices.at(device_idx));
	cl::CommandQueue command_queue(ctx);
	cl::Program program(ctx, devices, binaries);
	program.build();
	cl::Kernel kernel(program, "Example");

	static const size_t kResultElements = 4;
	static const size_t kResultBufferSize = sizeof (cl_float) * kResultElements;
	static const size_t kInputBufferSize = sizeof (cl_float);
	cl::Buffer result_buffer(ctx, CL_MEM_READ_WRITE, kResultBufferSize);
	cl::Buffer input_buffer(ctx, CL_MEM_READ_ONLY, kInputBufferSize);

	command_queue.enqueueFillBuffer(result_buffer, cl_float(0), 0, kResultBufferSize);
	command_queue.finish();

	{
		void *p = command_queue.enqueueMapBuffer(input_buffer, CL_TRUE, CL_MAP_WRITE_INVALIDATE_REGION, 0, kInputBufferSize);
		*static_cast<cl_float *>(p) = cl_float(0.5f);
		command_queue.enqueueUnmapMemObject(input_buffer, p);
	}

	kernel.setArg(0, result_buffer);
	kernel.setArg(1, cl_uint(10));
	kernel.setArg(2, input_buffer);
	command_queue.enqueueNDRangeKernel(kernel, cl::NDRange(0, 0), cl::NDRange(2, 2));
	command_queue.finish();

	{
		void *p = command_queue.enqueueMapBuffer(result_buffer, CL_TRUE, CL_MAP_READ, 0, kResultBufferSize);

		for (int i=0; i<4; ++i)
			std::cout << static_cast<cl_float *>(p)[i] << std::endl;

		command_queue.enqueueUnmapMemObject(result_buffer, p);
	}
}

期待通りの出力を確認できたと思います。

ディスアセンブル結果を読む

ディスアセンブル結果は次の通りです。

/* Disassembling 'example.bin' */
.amdcl2
.gpu GFX900
.64bit
.arch_minor 0
.arch_stepping 0
.driver_version 223600
.compile_options " -cl-std=CL2.0 -D__AMD__=1 -D__gfx900__=1 -D__gfx900=1 -DFP_FAST_FMAF=1 -DFP_FAST_FMA=1 -D__IMAGE_SUPPORT__=1 -DCL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE=6558465280 -cl-denorms-are-zero -m64 -Dcl_khr_fp64=1 -Dcl_amd_fp64=1 -Dcl_khr_global_int32_base_atomics=1 -Dcl_khr_global_int32_extended_atomics=1 -Dcl_khr_local_int32_base_atomics=1 -Dcl_khr_local_int32_extended_atomics=1 -Dcl_khr_int64_base_atomics=1 -Dcl_khr_int64_extended_atomics=1 -Dcl_khr_3d_image_writes=1 -Dcl_khr_byte_addressable_store=1 -Dcl_khr_fp16=1 -Dcl_khr_gl_sharing=1 -Dcl_amd_device_attribute_query=1 -Dcl_amd_vec3=1 -Dcl_amd_printf=1 -Dcl_amd_media_ops=1 -Dcl_amd_media_ops2=1 -Dcl_amd_popcnt=1 -Dcl_khr_d3d10_sharing=1 -Dcl_khr_d3d11_sharing=1 -Dcl_khr_dx9_media_sharing=1 -Dcl_khr_image2d_from_buffer=1 -Dcl_khr_subgroups=1 -Dcl_khr_gl_event=1 -Dcl_khr_depth_images=1 -Dcl_khr_mipmap_image=1 -Dcl_khr_mipmap_image_writes=1 -Dcl_amd_copy_buffer_p2p=1 -Dcl_amd_planar_yuv=1 "
.acl_version "AMD-COMP-LIB-v0.8 (0.0.SC_BUILD_NUMBER)"
.kernel Example
    .config
        .dims xyz
        .cws 2, 2, 1
        .sgprsnum 17
        .vgprsnum 5
        .floatmode 0xc0
        .pgmrsrc1 0x00ac0081
        .pgmrsrc2 0x00001394
        .dx10clamp
        .ieeemode
        .useargs
        .usesetup
        .priority 0
        .arg _.global_offset_0, "size_t", long
        .arg _.global_offset_1, "size_t", long
        .arg _.global_offset_2, "size_t", long
        .arg _.printf_buffer, "size_t", void*, global, , rdonly
        .arg _.vqueue_pointer, "size_t", long
        .arg _.aqlwrap_pointer, "size_t", long
        .arg dst, "float*", float*, global, 
        .arg a, "uint", uint
        .arg b, "float*", float*, global, , rdonly
    .text
/*000000000000: c0060002 0000000c*/ s_load_dwordx2  s[0:1], s[4:5], 0xc
/*000000000008: 6804040c         */ v_add_u32       v2, s12, v2
/*00000000000c: 8e02810b         */ s_lshl_b32      s2, s11, 1
/*000000000010: 8e03810a         */ s_lshl_b32      s3, s10, 1
/*000000000014: bf8cc07f         */ s_waitcnt       lgkmcnt(0)
/*000000000018: d2850002 00020401*/ v_mul_lo_u32    v2, s1, v2
/*000000000020: c0060103 00000040*/ s_load_dwordx2  s[4:5], s[6:7], 0x40
/*000000000028: c0020043 00000038*/ s_load_dword    s1, s[6:7], 0x38
/*000000000030: d1ff0002 040a0202*/ v_add3_u32      v2, s2, v1, v2
/*000000000038: d2850002 00020400*/ v_mul_lo_u32    v2, s0, v2
/*000000000040: bf8cc07f         */ s_waitcnt       lgkmcnt(0)
/*000000000044: c0020002 00000000*/ s_load_dword    s0, s[4:5], 0x0
/*00000000004c: c0060103 00000030*/ s_load_dwordx2  s[4:5], s[6:7], 0x30
/*000000000054: d2850004 00000300*/ v_mul_lo_u32    v4, v0, s1
/*00000000005c: d1ff0002 040a0003*/ v_add3_u32      v2, s3, v0, v2
/*000000000064: 7e060280         */ v_mov_b32       v3, 0
/*000000000068: d28f0002 00020482*/ v_lshlrev_b64   v[2:3], 2, v[2:3]
/*000000000070: 7e000d04         */ v_cvt_f32_u32   v0, v4
/*000000000074: 7e020d01         */ v_cvt_f32_u32   v1, v1
/*000000000078: bf8cc07f         */ s_waitcnt       lgkmcnt(0)
/*00000000007c: 32040404         */ v_add_co_u32    v2, vcc, s4, v2
/*000000000080: 7e080205         */ v_mov_b32       v4, s5
/*000000000084: 38060704         */ v_addc_co_u32   v3, vcc, v4, v3, vcc
/*000000000088: d1c10004 04000101*/ v_mad_f32       v4, v1, s0, v0
/*000000000090: dc708000 007f0402*/ global_store_dword v[2:3], v4, off
/*000000000098: bf810000         */ s_endpgm

まずはCLRadeonExtender Assembler Pseudo-OperationsとCLRadeonExtender Assembler AMD Catalyst OpenCL 2.0 handling、AMDGPU Compute Application Binary Interfaceを参照し、重要なところから見ていきます。

• プラットフォームから受け取るデータの設定
• 使用する、割当てるレジスタ数の設定
• 引数の設定
• その他の設定

その後、命令列を読みます。

プラットフォームから受け取るデータの設定

カーネルがプラットフォームから受け取るデータはSGPR(スカラ汎用レジスタ)とVGPR(ベクタ汎用レジスタにまとめてあり、どのデータをまとめておくのかを設定できます。
設定は

• .dims xyz
• .useargs
• .usesetup
• .pgmrsrc2 0x00001394
  で行われています。
  .dims xyzはamd_compute_pgm_rsrc2_t::enable_sgpr_workgroup_id_x, _y, _zを立て、amd_compute_pgm_rsrc2_t::enable_vgpr_workitem_id:=2とすることを意味します。ここからenable_sgpr_workgroup_id_zをクリアする場合は.dims xy, xyzとなります。enable_vgpr_workitem_id:=0としたい場合は.dims xyz, xです。
  .useargsはenable_sgpr_private_segment_bufferとenable_sgpr_kernarg_segment_ptrを立てることを意味します。
  .usesetupは.useargsに加えてenable_sgpr_dispatch_ptrを立てることを意味します。
  設定を一切行わなくてもVega 10ではenable_sgpr_flat_scratch_initが立ちます。
  .pgmrsrc2 0x00001394、.dimsでの設定を含んでいます。それに加えてamd_compute_pgm_rsrc2_t::user_sgpr_count:=10となっています。
  これは

項目	レジスタ数
sgpr_private_segment_buffer	4
sgpr_dispatch_ptr	2
sgpr_queue_ptr	2
sgpr_kernarg_segment_ptr	2
sgpr_dispatch_id	2
sgpr_flat_scratch_init	2
sgpr_private_segment_size	1
sgpr_grid_workgroup_count_X	1
sgpr_grid_workgroup_count_Y	1
sgpr_grid_workgroup_count_Z	1

のうち有効にしたい項目のレジスタ数の総和と等しくします。
今回は

• sgpr_private_segment_buffer
• sgpr_dispatch_ptr
• sgpr_kernarg_segment_ptr
• sgpr_flat_scratch_init
  の四つだけが有効なので、SGPRにまとめる際に用いるレジスタ数は4+2+2+2の10で、これら項目はそれぞれsgpr0..sgpr3, sgpr4..sgpr5, sgpr6..sgpr7, sgpr8..sgpr9に書き込まれます。値はリトルエンディアンです。
  これらに加えamd_compute_pgm_rsrc2_t::enable_sgpr_workgroup_id_x, _y, _zが立っているので、これらはそれぞれsgpr10, sgpr11, sgpr12に書き込まれます。
  amd_compute_pgm_rsrc2_t::enable_vgpr_workitem_id:=2なので、local idのx, y, zがそれぞれv0, v1, v2に書き込まれます。

使用する、割当てるレジスタ数の設定

reserved_vgpr_firstとreserved_vgpr_count、reserved_sgpr_firstとreserved_sgpr_countの意味合いが逆転しているように思えるのですが、気にしないでおきましょう。
まずは使用するレジスタ数ついてです。
.sgprsnum 17 reserved_sgpr_countの値です。今回wavefront_sgpr_countは19で、これはreserved_sgpr_countの値にVCCレジスタ(64bitなのでレジスタ2個)の2を加えたものだと思われます。つまりFLAT_SCRATCHレジスタとXNACK_MASKレジスタを割当てないらしいということです。
.vgprsnum 5 workitem_vgpr_count及びreserved_vgpr_firstの値です。
次は割当てるレジスタ数ついてです。
.pgmrsrc1 0x00ac0081 granulated_workitem_vgpr_count:=1, granulated_wavefront_sgpr_count:=2です。
これはレジスタの割当単位から1を減じた値で、今回のgfx900の割当単位はVGPRで4個、SGPRは16個なので、VGPRは4*(1+1)の8個、SGPRは16*(2+1)の48個割当てるよう求めます。この設定値にはトラップハンドラが有効な場合に使用するSGPRを含みません。

引数の設定

sgpr_kernarg_segment_ptrが指す場所に書き込まれる引数を設定します。
.arg _.global_offset_0, "size_t", long
.arg _.global_offset_1, "size_t", long
.arg _.global_offset_2, "size_t", long
.arg _.printf_buffer, "size_t", void*, global, , rdonly
.arg _.vqueue_pointer, "size_t", long
.arg _.aqlwrap_pointer, "size_t", long
.arg dst, "float*", float*, global,
.arg a, "uint", uint
.arg b, "float*", float*, global, , rdonly
dst, a, bのオフセットは48, 56, 64です。

その他の設定

.amdcl2 おそらくファイルがCode Object V2であることを表しています。
.gpu GFX900 ターゲットGPUがGFX900(Vega10のHBCC無効)であることを表しています。HBCC有効な状態でコンパイルしたデータを-aCオプションでアセンブルした場合GFX901になります。その場合xnackのフラグが立たない不具合があります。
.64bit 64bitアドレスを使うことを表しています。
.arch_minor 0 .arch_stepping 0 アーキテクチャマイナ、ステッピング番号です。GFX901だとそれぞれ0と1になります。
.driver_version 223600 ドライバに要求するバージョンが2236.0以上であることを表しているようです。

cl2GPUDeviceTypeMinDriverVersion

static const cxuint cl2GPUDeviceTypeMinDriverVersion[] =
{
    UINT_MAX, // CAPE_VERDE = 0, ///< Radeon HD7700
    UINT_MAX, // PITCAIRN, ///< Radeon HD7800
    UINT_MAX, // TAHITI, ///< Radeon HD7900
    UINT_MAX, // OLAND, ///< Radeon R7 250
    180005U, // BONAIRE, ///< Radeon R7 260
    180005U, // SPECTRE, ///< Kaveri
    180005U, // SPOOKY, ///< Kaveri
    180005U, // KALINDI, ///< ???  GCN1.1
    UINT_MAX, // HAINAN, ///< ????  GCN1.0
    180005U, // HAWAII, ///< Radeon R9 290
    180005U, // ICELAND, ///< ???
    180005U, // TONGA, ///< Radeon R9 285
    180005U, // MULLINS, ///< ???
    180005U, // FIJI,  ///< Radeon Fury
    180005U, // CARRIZO, ///< APU
    191205U, // DUMMY,
    200406U, // GOOSE,
    200406U, // HORSE,
    200406U, // STONEY,
    200406U, // ELLESMERE,
    200406U, // BAFFIN,
    223600U, // GFX804,
    223600U, // GFX900,
    226400U, // GFX901,
    252700U, // GFX902,
    252700U, // GFX903,
    252700U, // GFX904,
    252700U, // GFX905,
    258000U, // GFX906,
    258000U  // GFX907,
};

.compile_options コンパイルオプションです。記述しなくても問題なく動くようです。
.acl_version ACLバージョンです。記述しなくてもアセンブラが適当に付加します。
.cws 2, 2, 1 reqd_work_group_sizeです。
.floatmode 0xc0 .pgmrsrc1 0x00ac0081にあるものと同じです。デノーマル数については32bitの場合0へ丸め、16bitと64bitの場合は丸め無し、正規数は最も近い偶数への丸めです。

命令列を読む

次のように書いたものをコンパイルしました。

__attribute__((reqd_work_group_size(2, 2, 1)))
kernel void Example(global float *dst, uint a, global float *b)
{
	size_t gid = get_global_id(0);
	uint lid_x = (uint)get_local_id(0);
	uint lid_y = (uint)get_local_id(1);

	dst[gid] = (lid_x * a) + (lid_y * (*b));
}

s_load_dwordx2 s[0:1], s[4:5], 0xc
s[4:5]はArchitected Queuing Language (AQL) kernel dispatch packetを指しています。これはHeterogeneous System Architecture Foundation: StandardsにあるHSA Platform System Architecture Specificationに詳細があります。そのオフセット12バイトから64ビット読むので、
s0 = grid_size_x
s1 = grid_size_y
となります。

v_add_u32 v2, s12, v2
v2にはlocal_id_z、s12にはgroup_id_zが入っています。reqd_work_group_size(2, 2, 1)によりlocal_id_zは常に0なので、
v2 = group_id_z
となります。v_mov_b32 v2, s12と書くのと同じです。

s_lshl_b32 s2, s11, 1
s_lshl_b32 s3, s10, 1
s11にはgroup_id_y、s10にはgroup_id_xが入っています。1bit左シフトで2倍になるので、
s2 = group_id_y * 2
s3 = group_id_x * 2
となります。

s_waitcnt lgkmcnt(0)
異なるプリフィクス命令のメモリアクセスは発行順に完了する保証がありません。
lgkm...LDS, GDS, Constant, Messageに関わる命令の完了待ち数の総和が0以下になるまで待ちます。スカラ命令の場合はConstant扱いのようです。
ここでは具体的にはs_load_dwordx2 s[0:1], s[4:5], 0xcが完了するまで待ちます。待たない場合、s1への読込が完了する前に次のv_mul_lo_u32 v2, s1, v2を実行するかもしれません。

v_mul_lo_u32 v2, s1, v2
v2 = grid_size_y * group_id_z

s_load_dwordx2 s[4:5], s[6:7], 0x40
s[6:7]は引数のある場所を指しています。オフセット64には引数bがあるので、
s[4:5] = b
となります。

s_load_dword s1, s[6:7], 0x38
s1 = a

v_add3_u32 v2, s2, v1, v2
v2 = (group_id_y * 2) + (local_id_y) + (grid_size_y * group_id_z)

v_mul_lo_u32 v2, s0, v2
v2 = (grid_size_x) * ((group_id_y * 2) + (local_id_y) + (grid_size_y * group_id_z))

s_waitcnt lgkmcnt(0)
後続のv_mul_lo_u32 v4, v0, s1のためにs_load_dword s1, s[6:7], 0x38の完了を待ちます。

s_load_dword s0, s[4:5], 0x0
s0 = *b

s_load_dwordx2 s[4:5], s[6:7], 0x30
s[4:5] = dst

v_mul_lo_u32 v4, v0, s1
v4 = local_id_x * a

v_add3_u32 v2, s3, v0, v2
v2 = (group_id_x * 2) + (local_id_x) + (grid_size_x * ((group_id_y * 2) + local_id_y + (grid_size_y * group_id_z)))
書換えると
(grid_size_x * grid_size_y) * group_id_z
+ grid_size_x * (group_size_y * group_id_y)
+ grid_size_x * local_id_y
+ group_size_x * group_id_x
+ local_id_x
つまり、v2はglobal_idです。size_tなのに32bitで計算していますがデバイスメモリは8GiBなのでclGetDeviceInfo(,CL_DEVICE_MAX_MEM_ALLOC_SIZE で得られる値は16GiB未満なので問題ありません。

v_mov_b32 v3, 0
v3 = 0

v_lshlrev_b64 v[2:3], 2, v[2:3]
v[2:3] = v[2:3] * 4
v2はuint型配列の添え字を意味する値だったので、左2ビットシフトで4倍してオフセットアドレスにします。

v_cvt_f32_u32 v0, v4
v0 = float(local_id_x * a)

v_cvt_f32_u32 v1, v1
v1 = float(local_id_y)

s_waitcnt lgkmcnt(0)
s_load_dword s0, s[4:5], 0x0とs_load_dwordx2 s[4:5], s[6:7], 0x30の完了を待ちます。

v_add_co_u32 v2, vcc, s4, v2
vcc:v2 = dst.lo + オフセット.lo

v_mov_b32 v4, s5
v4 = dst.hi

v_addc_co_u32 v3, vcc, v4, v3, vcc
vcc:v3 = dst.hi + オフセット.hi + vcc

v_mad_f32 v4, v1, s0, v0
v4 = (float(local_id_y)) * (*b) + (float(local_id_x * a))

global_store_dword v[2:3], v4, off
dst[gid] = (lid_x * a) + (lid_y * (*b));
offはsgprによるオフセット指定をしないという意味です。

s_endpgm
プログラムの終了。事前のglobal_store_dwordの完了待ちは不要です。