75 lines
1.9 KiB
ArmAsm
75 lines
1.9 KiB
ArmAsm
.amdgcn_target "amdgcn-amd-amdhsa--@TARGET@"
|
|
.text
|
|
.globl probe
|
|
.p2align 8
|
|
.type probe,@function
|
|
probe:
|
|
s_load_dwordx4 s[4:7], s[0:1], 0x0
|
|
v_lshlrev_b32 v1, 4, v0
|
|
v_lshlrev_b32 v2, 6, v0
|
|
s_waitcnt lgkmcnt(0)
|
|
global_load_dwordx4 v[4:7], v1, s[4:5]
|
|
s_waitcnt vmcnt(0)
|
|
v_mov_b32 v10, 0
|
|
v_mov_b32 v11, 0
|
|
v_mov_b32 v12, 0
|
|
v_mov_b32 v13, 0
|
|
v_mov_b32 v14, 0
|
|
v_mov_b32 v15, 0
|
|
v_mov_b32 v16, 0
|
|
v_mov_b32 v17, 0
|
|
v_mov_b32 v18, 0
|
|
v_mov_b32 v19, 0
|
|
v_mov_b32 v20, 0
|
|
v_mov_b32 v21, 0
|
|
v_mov_b32 v22, 0
|
|
v_mov_b32 v23, 0
|
|
v_mov_b32 v24, 0
|
|
v_mov_b32 v25, 0
|
|
@BODY@
|
|
global_store_dwordx4 v2, v[10:13], s[6:7]
|
|
global_store_dwordx4 v2, v[14:17], s[6:7] offset:16
|
|
global_store_dwordx4 v2, v[18:21], s[6:7] offset:32
|
|
global_store_dwordx4 v2, v[22:25], s[6:7] offset:48
|
|
s_endpgm
|
|
.Lend:
|
|
.size probe, .Lend-probe
|
|
|
|
.rodata
|
|
.p2align 6
|
|
.amdhsa_kernel probe
|
|
.amdhsa_user_sgpr_kernarg_segment_ptr 1
|
|
.amdhsa_kernarg_size 16
|
|
.amdhsa_group_segment_fixed_size @LDS@
|
|
.amdhsa_next_free_vgpr 128
|
|
.amdhsa_next_free_sgpr 48
|
|
.amdhsa_wavefront_size32 @WAVE32@
|
|
.amdhsa_float_denorm_mode_32 @DENORM@
|
|
.amdhsa_float_denorm_mode_16_64 @DENORM16@
|
|
.amdhsa_ieee_mode @IEEE@
|
|
.amdhsa_dx10_clamp @DX10_CLAMP@
|
|
.amdhsa_float_round_mode_32 @ROUND32@
|
|
.amdhsa_float_round_mode_16_64 @ROUND16@
|
|
.amdhsa_fp16_overflow @FP16_OVERFLOW@
|
|
.end_amdhsa_kernel
|
|
|
|
.amdgpu_metadata
|
|
---
|
|
amdhsa.version: [ 1, 2 ]
|
|
amdhsa.kernels:
|
|
- .name: probe
|
|
.symbol: probe.kd
|
|
.kernarg_segment_size: 16
|
|
.group_segment_fixed_size: @LDS@
|
|
.private_segment_fixed_size: 0
|
|
.kernarg_segment_align: 8
|
|
.wavefront_size: @WAVESIZE@
|
|
.sgpr_count: 48
|
|
.vgpr_count: 64
|
|
.max_flat_workgroup_size: 1024
|
|
.args:
|
|
- { .size: 8, .offset: 0, .value_kind: global_buffer, .address_space: global }
|
|
- { .size: 8, .offset: 8, .value_kind: global_buffer, .address_space: global }
|
|
...
|
|
.end_amdgpu_metadata
|