From 2dccbf8c7fb8b56b79a4ea9a212d3c998bf3c5db Mon Sep 17 00:00:00 2001 From: Alyssa Rosenzweig Date: Tue, 27 May 2025 15:54:47 -0400 Subject: [PATCH] agx: add XML description of the AGX2 ISA for our disassembler and later more stuff. should cover everything mesa uses, not 100% complete to what applegpu has but mostly there. Signed-off-by: Alyssa Rosenzweig Part-of: --- src/asahi/isa/AGX2.xml | 1620 ++++++++++++++++++++++++++++++++++++++++ 1 file changed, 1620 insertions(+) create mode 100644 src/asahi/isa/AGX2.xml diff --git a/src/asahi/isa/AGX2.xml b/src/asahi/isa/AGX2.xml new file mode 100644 index 00000000000..55ab3fa831d --- /dev/null +++ b/src/asahi/isa/AGX2.xml @@ -0,0 +1,1620 @@ + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + + none + cache + discard + + + + center + sample + centroid + sample + + + + incoherent + coherent1 + coherent2 + coherent + + + + add + sub + xchg + cmpxchg + umin + imin + umax + imax + and + or + xor + + + + + + + + + + + + 3 + + + 01 + + + + + + + + + + + + + + + + + + + + 11 + + + + + + 01 + + + + + + 0 + + + + + + + + + + + + + + + + + 1111 + + + + + + 0011 + + + + + + 1111 + 4 + + + + 0111 + 3 + + + + 0011 + 2 + + + + 0001 + + + + + + + + 1111 + + + + + + 0011 + + + + + + 1111 + 4 + + + + 0111 + 3 + + + + 0011 + 2 + + + + 0001 + + + + + + + + + + 0 + + + + 1 + + + + + + + + + + + + + 0100 + + + 2 + + + + + 1100 + + + 5 + + + + 0010 + + + + + 0001 + + + + + 1001 + + 2 + + + + 1101 + + 2 + + + + 0110 + + + + + 0101 + + + + + + + + + + + + + + + + + + + + + 1 + + + + + + + + + + + 11 + + + + + + + + + + + + + + + + 0011 + 2 + + + + + 1 + + + + + + + + + + + + 0 + + + + + 1 + + + + + + + 0 + + + + + 1 + + + + + + + + + + + 11 + + + + + 01 + + + + + 0 + + + + + + + 4 + + + 0 + + + + + 1 + + + + + + + + + + 0 + + + + + + + 1 + + + + + + 00 + + + + + 10 + + + + + 1 + + + + + + + + + + + + + + + 1 + 0 + + + + + 1 + 1 + + + + + + + + + + + + 1 + + + + + + 100 + + + + + + + + + + + + + + 11 + + + + + + + + + + + 01 + + + + + + + + + + + + + + + + + + + + + + 01 + + + + + + + + + + + + + + + + + + + + + + 01 + + + + + + + + + + + + + + + + + + + + + + + 01 + + + + + + + + + + + + + + + + + + + + + + + 1 + + + + + + + + + + + + + + 0 + + + + + 1 + + + + + + + 1 + + + + + 0 + + + + + + + 0 + + + + + 1 + + 2 + + + + + + + + + + 01 + + + + + + 10 + + + + + 11 + + 3 + + + + + + + + + + 1 + + + + + + + + 00 + + + + + + 01 + + + + + + 10 + + + + + + 11 + + + + + + + + 0 + + + + + 1 + + + + + + threadgroup_in_grid_x + threadgroup_in_grid_y + threadgroup_in_grid_z + threads_per_threadgroup_x + threads_per_threadgroup_y + threads_per_threadgroup_z + dispatch_threads_per_threadgroup_x + dispatch_threads_per_threadgroup_y + dispatch_threads_per_threadgroup_z + threadgroup_width + threadgroup_height + samples_log2 + core_index + vm_slot + thread_in_threadgroup_x + thread_in_threadgroup_y + thread_in_threadgroup_z + thread_index_in_threadgroup + thread_index_in_simdgroup + simdgroup_index_in_threadgroup + active_thread_index_in_quadgroup + total_active_threads_in_quadgroup + active_thread_index_in_simdgroup + total_active_threads_in_simdgroup + internal_coverage_mask + backfacing + is_active_thread + thread_in_grid_x + thread_in_grid_y + thread_in_grid_z + input_sample_mask + helper_op + helper_arg_l + helper_arg_h + + + + rtz + rte + + + + a + b + + + + + lsl 1 + lsl 2 + lsl 3 + lsl 4 + + + + sx + zx + + + + ieq + ult + ugt + seq + slt + sgt + + ine + ugte + ulte + nseq + sgte + slte + + + + feq + flt + fgt + fltn + fgte + flte + fgtn + + fneq + fnlt + fngt + fnltn + fngte + fnlte + fngtn + + + + + none + x + y + xy + z + xz + yz + xyz + w + xw + yw + xyw + zw + xzw + yzw + xyzw + + + + i8 + i16 + i32 + f16 + u8norm + s8norm + u16norm + s16norm + rgb10a2 + srgba8 + rg11b10f + rgb9e5 + rg11b10f_rtz + rgb9e5_rtz + + + + 1d + 1d_array + 2d + 2d_array + 2d_ms + 3d + cube + cube_array + 2d_ms_array + + + + none + gather_r + gather_g + gather_b + gather_a + + + + auto_lod + auto_lod_bias + lod_min + lod_grad + auto_lod_bias + lod_min + auto_lod_bias_min + lod_grad_min + auto_lod_bias_min + + + + 01100010 + + + + + + 11100010 + + + + + + 01110010 + + + + + + 0001110 + negate + + + + + + + + + + + + 0011110 + + + + + + + + + 0011110 + negate + + + + + + + + + + + + + 10111110 + mode + + + + + + + + + + + + + + + + + + + 0101110 + opcode + + + + + + + + + + + 0101110 + opcode + + + + + + + + + + + + + + + + + 111010 + + + + + + + + + 110110 + + + + + + + + + opcode + + + + + + + + + + + opcode + + + + + + + + + + + 1111110 + + + + opcode + + + + + + + + + + + 1111110 + + opcode + + + + + + + + + + + + + + + + + + + + 0111110 + + opcode + + + + + + + + + + + 0111110 + + + + + + + + + + 001010 + opcode + + + + + + + + + + + + + + + + + + + + + + 0000000010001000 + + + + 0000000000001000 + + + + + + + + + opcode + + + + + + + + opcode + + + + + + + + + 1010010 + 11 + + + + + + + 1010010 + opcode + + + + + + + + + + + + + + 1000010 + + + + + + 1000010 + opcode + + + + + + + + + + + + + + 0010010 + + + + + + + + + + 0000010 + + + + + + + + + + + 0110010 + 111 + + + + + + quad + simd + + + + 0110010 + + + + + + + + + + 0100010 + + + + + + + + + + none + xor + up + down + + + + 1101111 + + + + + + + + + + + opcode + + + + + + + + + 00111000 + + + + + 100001010001 + + + + 0010001 + + + 1 + + + + + + + + 0000101 + 11 + + + + + + + + + + + + + 010101 + 111 + + + + + + + + + + + 011001 + 11 + + + + + + + + + + 00110101 + 1 + + 100 + + + + + + + + + + + 10110101 + 11 + 010 + + + + + + + 10110101 + 1 + + 100 + + + + + + + + + + + 1000101 + 11 + 11 + + 0001 + 1 + + + + + + + + + + 1000101 + 11 + 11 + + + + + + + + + + + + 111101101 + + + + + + + 110101101 + + + + + + + 1000101 + 11 + + + + + + + + + + + + + 1101001 + + + + + + + + + 0101001 + + + + + + + + + 01101000 + + + + 00101000 + + + + + 0000000000100000 + + + + + + 01110101 + 1 + op + 01 + + + + + 1 + 10000 + 00010000 + + + + + + + + + + + 11110101 + b + a + c + + + + + + + + + + + + + + + + + + and + fadd + or + fmul + xor + fmin + fmax + iadd + smin + smax + umin + umax + + + + 01101111 + 1 + + simd + + + + + + + + + + + + 1001001 + 1 + + + + + + + + + + 0001001 + + + + + + + + + + + 11000001 + + + + + + + 01000001 + + + + + + + 00100001 + + + + + + + + + + + 01100001 + + + + + + + + + + + + 10100001 + + + + + + + 0110001 + intcoord + + + 11 + + + + + + + + + + + + + + + + + + + + + + + + + 11110001 + + + + + + 111 + + + + + + + 10110001 + + + 111 + + + + + + + + + + + +