Go to:
Gentoo Home
Documentation
Forums
Lists
Bugs
Planet
Store
Wiki
Get Gentoo!
Gentoo's Bugzilla – Attachment 900599 Details for
Bug 938163
dev-libs/rocm-device-libs-6.1.2 fails test - 6 - compile_frexp_gfx600 (Failed)
Home
|
New
–
[Ex]
|
Browse
|
Search
|
Privacy Policy
|
[?]
|
Reports
|
Requests
|
Help
|
New Account
|
Log In
[x]
|
Forgot Password
Login:
[x]
LastTest.log
LastTest.log (text/plain), 121.95 KB, created by
Toralf Förster
on 2024-08-18 20:32:48 UTC
(
hide
)
Description:
LastTest.log
Filename:
MIME Type:
Creator:
Toralf Förster
Created:
2024-08-18 20:32:48 UTC
Size:
121.95 KB
patch
obsolete
>Start testing: Aug 18 19:10 UTC >---------------------------------------------------------- >1/21 Testing: constant_fold_lgamma_r__gfx900 >1/21 Test: constant_fold_lgamma_r__gfx900 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.lgamma_r.gfx900.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/lgamma_r.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx900" "-DEXTRA_CHECK_PREFIX=CHECK" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunConstantFoldTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"constant_fold_lgamma_r__gfx900" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.18 sec >---------------------------------------------------------- >Test Passed. >"constant_fold_lgamma_r__gfx900" end time: Aug 18 19:10 UTC >"constant_fold_lgamma_r__gfx900" time elapsed: 00:00:00 >---------------------------------------------------------- > >3/21 Testing: compile_asin__gfx803 >3/21 Test: compile_asin__gfx803 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.asin.gfx803.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/asin.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx803" "-DEXTRA_CHECK_PREFIX=GFX803,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_asin__gfx803" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.18 sec >---------------------------------------------------------- >Test Passed. >"compile_asin__gfx803" end time: Aug 18 19:10 UTC >"compile_asin__gfx803" time elapsed: 00:00:00 >---------------------------------------------------------- > >2/21 Testing: constant_fold_lgamma_r__gfx1030 >2/21 Test: constant_fold_lgamma_r__gfx1030 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.lgamma_r.gfx1030.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/lgamma_r.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx1030" "-DEXTRA_CHECK_PREFIX=CHECK" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunConstantFoldTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"constant_fold_lgamma_r__gfx1030" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.20 sec >---------------------------------------------------------- >Test Passed. >"constant_fold_lgamma_r__gfx1030" end time: Aug 18 19:10 UTC >"constant_fold_lgamma_r__gfx1030" time elapsed: 00:00:00 >---------------------------------------------------------- > >4/21 Testing: compile_atan2__gfx803 >4/21 Test: compile_atan2__gfx803 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.atan2.gfx803.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/atan2.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx803" "-DEXTRA_CHECK_PREFIX=GFX803,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_atan2__gfx803" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.20 sec >---------------------------------------------------------- >Test Passed. >"compile_atan2__gfx803" end time: Aug 18 19:10 UTC >"compile_atan2__gfx803" time elapsed: 00:00:00 >---------------------------------------------------------- > >5/21 Testing: compile_atan2pi__gfx803 >5/21 Test: compile_atan2pi__gfx803 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.atan2pi.gfx803.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/atan2pi.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx803" "-DEXTRA_CHECK_PREFIX=GFX803,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_atan2pi__gfx803" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.20 sec >---------------------------------------------------------- >Test Passed. >"compile_atan2pi__gfx803" end time: Aug 18 19:10 UTC >"compile_atan2pi__gfx803" time elapsed: 00:00:00 >---------------------------------------------------------- > >6/21 Testing: compile_frexp__gfx600 >6/21 Test: compile_frexp__gfx600 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.frexp.gfx600.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/frexp.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx600" "-DEXTRA_CHECK_PREFIX=GFX600,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_frexp__gfx600" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- >CMake Error at /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake:37 (message): > Error in test output: > /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/frexp.cl:8:16: > error: GFX600-DAG: expected string not found in input > > // GFX600-DAG: s_mov_b32 [[INF:s[0-9]+]], 0x7f80000 > > ^ > > output.frexp.gfx600.s:8:16: note: scanning from here > > test_frexp_f32: ; @test_frexp_f32 > > ^ > > output.frexp.gfx600.s:11:2: note: possible intended match here > > s_mov_b32 s3, 0x100f000 > ^ > > > /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/frexp.cl:35:16: > error: GFX600-DAG: expected string not found in input > > // GFX600-DAG: s_mov_b32 s[[INF_LO:[0-9]+]], 0{{$}} > > ^ > > output.frexp.gfx600.s:104:17: note: scanning from here > > s_mov_b32 s2, 0 > ^ > > output.frexp.gfx600.s:106:2: note: possible intended match here > > v_mov_b32_e32 v2, 0 > ^ > > > > Input file: output.frexp.gfx600.s > > Check file: > /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/frexp.cl > > > > > -dump-input=help explains the following input dump. > > > > Input was: > > <<<<<< > > 1: .text > 2: .amdgcn_target "amdgcn-amd-amdhsa--gfx600" > 3: .amdhsa_code_object_version 5 > 4: .protected test_frexp_f32 ; -- Begin function test_frexp_f32 > 5: .globl test_frexp_f32 > 6: .p2align 8 > 7: .type test_frexp_f32,@function > 8: test_frexp_f32: ; @test_frexp_f32 > > label:7'0 ^~~~~~~~~~~~~~~ > > label:7'1 ^~~~~~~~~~~~~~~ > > dag:8'0 X~~~~~~~~~~~~~~~~~~ error: no match found > > 9: ; %bb.0: > > dag:8'0 ~~~~~~~~~ > > 10: s_load_dwordx2 s[0:1], s[4:5], 0x4 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 11: s_mov_b32 s3, 0x100f000 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~ > > dag:8'1 ? possible intended match > > 12: s_mov_b32 s2, 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~ > > 13: v_lshlrev_b32_e32 v0, 2, v0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 14: v_mov_b32_e32 v1, 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~ > > 15: s_waitcnt lgkmcnt(0) > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 16: buffer_load_dword v2, v[0:1], s[0:3], 0 addr64 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 17: s_load_dwordx4 s[4:7], s[4:5], 0x0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 18: s_movk_i32 s8, 0x1f8 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 19: s_waitcnt lgkmcnt(0) > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 20: s_mov_b64 s[0:1], s[4:5] > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 21: s_waitcnt vmcnt(0) > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~ > > 22: v_frexp_mant_f32_e32 v4, v2 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 23: v_cmp_class_f32_e64 vcc, v2, s8 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 24: v_frexp_exp_i32_f32_e32 v3, v2 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 25: v_cndmask_b32_e32 v2, v2, v4, vcc > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 26: v_cndmask_b32_e32 v3, 0, v3, vcc > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 27: buffer_store_dword v2, v[0:1], s[0:3], 0 addr64 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 28: s_mov_b64 s[0:1], s[6:7] > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 29: buffer_store_dword v3, v[0:1], s[0:3], 0 addr64 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 30: s_endpgm > > dag:8'0 ~~~~~~~~~~ > > 31: .section .rodata,"a",@progbits > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 32: .p2align 6, 0x0 > > dag:8'0 ~~~~~~~~~~~~~~~~~ > > 33: .amdhsa_kernel test_frexp_f32 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 34: .amdhsa_group_segment_fixed_size 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 35: .amdhsa_private_segment_fixed_size 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 36: .amdhsa_kernarg_size 24 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~ > > 37: .amdhsa_user_sgpr_count 6 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 38: .amdhsa_user_sgpr_private_segment_buffer 1 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 39: .amdhsa_user_sgpr_dispatch_ptr 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 40: .amdhsa_user_sgpr_queue_ptr 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 41: .amdhsa_user_sgpr_kernarg_segment_ptr 1 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 42: .amdhsa_user_sgpr_dispatch_id 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 43: .amdhsa_user_sgpr_flat_scratch_init 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 44: .amdhsa_user_sgpr_private_segment_size 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 45: .amdhsa_uses_dynamic_stack 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 46: .amdhsa_system_sgpr_private_segment_wavefront_offset 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 47: .amdhsa_system_sgpr_workgroup_id_x 1 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 48: .amdhsa_system_sgpr_workgroup_id_y 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 49: .amdhsa_system_sgpr_workgroup_id_z 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 50: .amdhsa_system_sgpr_workgroup_info 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 51: .amdhsa_system_vgpr_workitem_id 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 52: .amdhsa_next_free_vgpr 5 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 53: .amdhsa_next_free_sgpr 9 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 54: .amdhsa_float_round_mode_32 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 55: .amdhsa_float_round_mode_16_64 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 56: .amdhsa_float_denorm_mode_32 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 57: .amdhsa_float_denorm_mode_16_64 3 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 58: .amdhsa_dx10_clamp 1 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 59: .amdhsa_ieee_mode 1 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~ > > 60: .amdhsa_exception_fp_ieee_invalid_op 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 61: .amdhsa_exception_fp_denorm_src 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 62: .amdhsa_exception_fp_ieee_div_zero 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 63: .amdhsa_exception_fp_ieee_overflow 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 64: .amdhsa_exception_fp_ieee_underflow 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 65: .amdhsa_exception_fp_ieee_inexact 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 66: .amdhsa_exception_int_div_zero 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 67: .end_amdhsa_kernel > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~ > > 68: .text > > dag:8'0 ~~~~~~~ > > 69: .Lfunc_end0: > > dag:8'0 ~~~~~~~~~~~~~ > > 70: .size test_frexp_f32, .Lfunc_end0-test_frexp_f32 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 71: ; -- End function > > dag:8'0 ~~~~~~~~~~~~~~~~~~~ > > 72: .section .AMDGPU.csdata,"",@progbits > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 73: ; Kernel info: > > dag:8'0 ~~~~~~~~~~~~~~~ > > 74: ; codeLenInByte = 104 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 75: ; NumSgprs: 11 > > dag:8'0 ~~~~~~~~~~~~~~~ > > 76: ; NumVgprs: 5 > > dag:8'0 ~~~~~~~~~~~~~~ > > 77: ; ScratchSize: 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~ > > 78: ; MemoryBound: 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~ > > 79: ; FloatMode: 192 > > dag:8'0 ~~~~~~~~~~~~~~~~~ > > 80: ; IeeeMode: 1 > > dag:8'0 ~~~~~~~~~~~~~~ > > 81: ; LDSByteSize: 0 bytes/workgroup (compile time only) > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 82: ; SGPRBlocks: 1 > > dag:8'0 ~~~~~~~~~~~~~~~~ > > 83: ; VGPRBlocks: 1 > > dag:8'0 ~~~~~~~~~~~~~~~~ > > 84: ; NumSGPRsForWavesPerEU: 11 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 85: ; NumVGPRsForWavesPerEU: 5 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 86: ; Occupancy: 10 > > dag:8'0 ~~~~~~~~~~~~~~~~ > > 87: ; WaveLimiterHint : 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 88: ; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 89: ; COMPUTE_PGM_RSRC2:USER_SGPR: 6 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 90: ; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 91: ; COMPUTE_PGM_RSRC2:TGID_X_EN: 1 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 92: ; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 93: ; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 94: ; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 95: .text > > dag:8'0 ~~~~~~~ > > 96: .protected test_frexp_f64 ; -- Begin function test_frexp_f64 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 97: .globl test_frexp_f64 > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~ > > 98: .p2align 8 > > dag:8'0 ~~~~~~~~~~~~ > > 99: .type test_frexp_f64,@function > > dag:8'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 100: test_frexp_f64: ; @test_frexp_f64 > > label:32 ^~~~~~~~~~~~~~~ > > dag:8'0 ~~~~~~~~~~~~~~~ > > 101: ; %bb.0: > 102: s_load_dwordx2 s[0:1], s[4:5], 0x4 > 103: s_mov_b32 s3, 0x100f000 > 104: s_mov_b32 s2, 0 > > check:33 ^~~~~~~~~~~~~~~ > > dag:35'0 X error: no match found > > 105: v_lshlrev_b32_e32 v1, 3, v0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 106: v_mov_b32_e32 v2, 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~ > > dag:35'1 ? possible intended match > > 107: s_waitcnt lgkmcnt(0) > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 108: buffer_load_dwordx2 v[3:4], v[1:2], s[0:3], 0 addr64 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 109: s_load_dwordx4 s[4:7], s[4:5], 0x0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 110: s_movk_i32 s8, 0x1f8 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 111: s_waitcnt lgkmcnt(0) > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 112: s_mov_b64 s[0:1], s[4:5] > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 113: s_waitcnt vmcnt(0) > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~ > > 114: v_frexp_mant_f64_e32 v[5:6], v[3:4] > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 115: v_cmp_class_f64_e64 vcc, v[3:4], s8 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 116: v_frexp_exp_i32_f64_e32 v7, v[3:4] > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 117: v_cndmask_b32_e32 v4, v4, v6, vcc > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 118: v_cndmask_b32_e32 v3, v3, v5, vcc > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 119: v_cndmask_b32_e32 v7, 0, v7, vcc > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 120: buffer_store_dwordx2 v[3:4], v[1:2], s[0:3], 0 addr64 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 121: s_mov_b64 s[0:1], s[6:7] > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 122: v_lshlrev_b32_e32 v1, 2, v0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 123: buffer_store_dword v7, v[1:2], s[0:3], 0 addr64 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 124: s_endpgm > > dag:35'0 ~~~~~~~~~~ > > 125: .section .rodata,"a",@progbits > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 126: .p2align 6, 0x0 > > dag:35'0 ~~~~~~~~~~~~~~~~~ > > 127: .amdhsa_kernel test_frexp_f64 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 128: .amdhsa_group_segment_fixed_size 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 129: .amdhsa_private_segment_fixed_size 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 130: .amdhsa_kernarg_size 24 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~ > > 131: .amdhsa_user_sgpr_count 6 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 132: .amdhsa_user_sgpr_private_segment_buffer 1 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 133: .amdhsa_user_sgpr_dispatch_ptr 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 134: .amdhsa_user_sgpr_queue_ptr 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 135: .amdhsa_user_sgpr_kernarg_segment_ptr 1 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 136: .amdhsa_user_sgpr_dispatch_id 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 137: .amdhsa_user_sgpr_flat_scratch_init 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 138: .amdhsa_user_sgpr_private_segment_size 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 139: .amdhsa_uses_dynamic_stack 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 140: .amdhsa_system_sgpr_private_segment_wavefront_offset 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 141: .amdhsa_system_sgpr_workgroup_id_x 1 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 142: .amdhsa_system_sgpr_workgroup_id_y 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 143: .amdhsa_system_sgpr_workgroup_id_z 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 144: .amdhsa_system_sgpr_workgroup_info 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 145: .amdhsa_system_vgpr_workitem_id 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 146: .amdhsa_next_free_vgpr 8 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 147: .amdhsa_next_free_sgpr 9 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 148: .amdhsa_float_round_mode_32 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 149: .amdhsa_float_round_mode_16_64 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 150: .amdhsa_float_denorm_mode_32 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 151: .amdhsa_float_denorm_mode_16_64 3 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 152: .amdhsa_dx10_clamp 1 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 153: .amdhsa_ieee_mode 1 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~ > > 154: .amdhsa_exception_fp_ieee_invalid_op 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 155: .amdhsa_exception_fp_denorm_src 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 156: .amdhsa_exception_fp_ieee_div_zero 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 157: .amdhsa_exception_fp_ieee_overflow 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 158: .amdhsa_exception_fp_ieee_underflow 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 159: .amdhsa_exception_fp_ieee_inexact 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 160: .amdhsa_exception_int_div_zero 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 161: .end_amdhsa_kernel > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~ > > 162: .text > > dag:35'0 ~~~~~~~ > > 163: .Lfunc_end1: > > dag:35'0 ~~~~~~~~~~~~~ > > 164: .size test_frexp_f64, .Lfunc_end1-test_frexp_f64 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 165: ; -- End function > > dag:35'0 ~~~~~~~~~~~~~~~~~~~ > > 166: .section .AMDGPU.csdata,"",@progbits > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 167: ; Kernel info: > > dag:35'0 ~~~~~~~~~~~~~~~ > > 168: ; codeLenInByte = 112 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 169: ; NumSgprs: 11 > > dag:35'0 ~~~~~~~~~~~~~~~ > > 170: ; NumVgprs: 8 > > dag:35'0 ~~~~~~~~~~~~~~ > > 171: ; ScratchSize: 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~ > > 172: ; MemoryBound: 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~ > > 173: ; FloatMode: 192 > > dag:35'0 ~~~~~~~~~~~~~~~~~ > > 174: ; IeeeMode: 1 > > dag:35'0 ~~~~~~~~~~~~~~ > > 175: ; LDSByteSize: 0 bytes/workgroup (compile time only) > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 176: ; SGPRBlocks: 1 > > dag:35'0 ~~~~~~~~~~~~~~~~ > > 177: ; VGPRBlocks: 1 > > dag:35'0 ~~~~~~~~~~~~~~~~ > > 178: ; NumSGPRsForWavesPerEU: 11 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 179: ; NumVGPRsForWavesPerEU: 8 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 180: ; Occupancy: 10 > > dag:35'0 ~~~~~~~~~~~~~~~~ > > 181: ; WaveLimiterHint : 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 182: ; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 183: ; COMPUTE_PGM_RSRC2:USER_SGPR: 6 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 184: ; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 185: ; COMPUTE_PGM_RSRC2:TGID_X_EN: 1 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 186: ; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 187: ; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 188: ; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 189: .hidden __oclc_ABI_version ; @__oclc_ABI_version > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 190: .type __oclc_ABI_version,@object > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 191: .section .rodata,"a",@progbits > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 192: .weak __oclc_ABI_version > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 193: .p2align 2, 0x0 > > dag:35'0 ~~~~~~~~~~~~~~~~~ > > 194: __oclc_ABI_version: > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~ > > 195: .long 500 ; 0x1f4 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~ > > 196: .size __oclc_ABI_version, 4 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 197: > > dag:35'0 ~ > > 198: .ident "clang version 18.1.8" > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 199: .section ".note.GNU-stack","",@progbits > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 200: .addrsig > > dag:35'0 ~~~~~~~~~~ > > 201: .amdgpu_metadata > > dag:35'0 ~~~~~~~~~~~~~~~~~~ > > 202: --- > > dag:35'0 ~~~~ > > 203: amdhsa.kernels: > > dag:35'0 ~~~~~~~~~~~~~~~~ > > 204: - .args: > > dag:35'0 ~~~~~~~~~~ > > 205: - .actual_access: write_only > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 206: .address_space: global > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~ > > 207: .is_restrict: true > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~ > > 208: .offset: 0 > > dag:35'0 ~~~~~~~~~~~~ > > 209: .size: 8 > > dag:35'0 ~~~~~~~~~~ > > 210: .type_name: 'float*' > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 211: .value_kind: global_buffer > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 212: - .actual_access: write_only > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 213: .address_space: global > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~ > > 214: .is_restrict: true > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~ > > 215: .offset: 8 > > dag:35'0 ~~~~~~~~~~~~ > > 216: .size: 8 > > dag:35'0 ~~~~~~~~~~ > > 217: .type_name: 'int*' > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~ > > 218: .value_kind: global_buffer > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 219: - .actual_access: read_only > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 220: .address_space: global > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~ > > 221: .is_restrict: true > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~ > > 222: .offset: 16 > > dag:35'0 ~~~~~~~~~~~~~ > > 223: .size: 8 > > dag:35'0 ~~~~~~~~~~ > > 224: .type_name: 'float*' > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 225: .value_kind: global_buffer > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 226: .group_segment_fixed_size: 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 227: .kernarg_segment_align: 8 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 228: .kernarg_segment_size: 24 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 229: .language: OpenCL C > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~ > > 230: .language_version: > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~ > > 231: - 2 > > dag:35'0 ~~~~~ > > 232: - 0 > > dag:35'0 ~~~~~ > > 233: .max_flat_workgroup_size: 256 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 234: .name: test_frexp_f32 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~ > > 235: .private_segment_fixed_size: 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 236: .sgpr_count: 11 > > dag:35'0 ~~~~~~~~~~~~~~~~~ > > 237: .sgpr_spill_count: 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 238: .symbol: test_frexp_f32.kd > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 239: .uses_dynamic_stack: false > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 240: .vgpr_count: 5 > > dag:35'0 ~~~~~~~~~~~~~~~~ > > 241: .vgpr_spill_count: 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 242: .wavefront_size: 64 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~ > > 243: - .args: > > dag:35'0 ~~~~~~~~~~ > > 244: - .actual_access: write_only > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 245: .address_space: global > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~ > > 246: .is_restrict: true > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~ > > 247: .offset: 0 > > dag:35'0 ~~~~~~~~~~~~ > > 248: .size: 8 > > dag:35'0 ~~~~~~~~~~ > > 249: .type_name: 'double*' > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~ > > 250: .value_kind: global_buffer > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 251: - .actual_access: write_only > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 252: .address_space: global > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~ > > 253: .is_restrict: true > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~ > > 254: .offset: 8 > > dag:35'0 ~~~~~~~~~~~~ > > 255: .size: 8 > > dag:35'0 ~~~~~~~~~~ > > 256: .type_name: 'int*' > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~ > > 257: .value_kind: global_buffer > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 258: - .actual_access: read_only > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 259: .address_space: global > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~ > > 260: .is_restrict: true > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~ > > 261: .offset: 16 > > dag:35'0 ~~~~~~~~~~~~~ > > 262: .size: 8 > > dag:35'0 ~~~~~~~~~~ > > 263: .type_name: 'double*' > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~ > > 264: .value_kind: global_buffer > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 265: .group_segment_fixed_size: 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 266: .kernarg_segment_align: 8 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 267: .kernarg_segment_size: 24 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 268: .language: OpenCL C > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~ > > 269: .language_version: > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~ > > 270: - 2 > > dag:35'0 ~~~~~ > > 271: - 0 > > dag:35'0 ~~~~~ > > 272: .max_flat_workgroup_size: 256 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 273: .name: test_frexp_f64 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~ > > 274: .private_segment_fixed_size: 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 275: .sgpr_count: 11 > > dag:35'0 ~~~~~~~~~~~~~~~~~ > > 276: .sgpr_spill_count: 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 277: .symbol: test_frexp_f64.kd > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 278: .uses_dynamic_stack: false > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 279: .vgpr_count: 8 > > dag:35'0 ~~~~~~~~~~~~~~~~ > > 280: .vgpr_spill_count: 0 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 281: .wavefront_size: 64 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~ > > 282: amdhsa.target: amdgcn-amd-amdhsa--gfx600 > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 283: amdhsa.version: > > dag:35'0 ~~~~~~~~~~~~~~~~ > > 284: - 1 > > dag:35'0 ~~~~~ > > 285: - 2 > > dag:35'0 ~~~~~ > > 286: ... > > dag:35'0 ~~~~ > > 287: > > dag:35'0 ~ > > 288: .end_amdgpu_metadata > > dag:35'0 ~~~~~~~~~~~~~~~~~~~~~~ > > >>>>>> > > > ><end of output> >Test time = 0.21 sec >---------------------------------------------------------- >Test Failed. >"compile_frexp__gfx600" end time: Aug 18 19:10 UTC >"compile_frexp__gfx600" time elapsed: 00:00:00 >---------------------------------------------------------- > >8/21 Testing: compile_native_rcp__gfx600 >8/21 Test: compile_native_rcp__gfx600 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.native_rcp.gfx600.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/native_rcp.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx600" "-DEXTRA_CHECK_PREFIX=GFX600,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_native_rcp__gfx600" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.19 sec >---------------------------------------------------------- >Test Passed. >"compile_native_rcp__gfx600" end time: Aug 18 19:10 UTC >"compile_native_rcp__gfx600" time elapsed: 00:00:00 >---------------------------------------------------------- > >7/21 Testing: compile_fract__gfx600 >7/21 Test: compile_fract__gfx600 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.fract.gfx600.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/fract.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx600" "-DEXTRA_CHECK_PREFIX=GFX600,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_fract__gfx600" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- >CMake Error at /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake:37 (message): > Error in test output: > /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/fract.cl:4:12: > error: GFX600: expected string not found in input > > // GFX600: v_cvt_f32_f16 > > ^ > > output.fract.gfx600.s:4:27: note: scanning from here > > .protected test_fract_f16 ; -- Begin function test_fract_f16 > ^ > > output.fract.gfx600.s:4:54: note: possible intended match here > > .protected test_fract_f16 ; -- Begin function test_fract_f16 > ^ > > > > Input file: output.fract.gfx600.s > > Check file: > /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/fract.cl > > > > > -dump-input=help explains the following input dump. > > > > Input was: > > <<<<<< > > 1: .text > 2: .amdgcn_target "amdgcn-amd-amdhsa--gfx600" > 3: .amdhsa_code_object_version 5 > 4: .protected test_fract_f16 ; -- Begin function test_fract_f16 > > label:3'0 ^~~~~~~~~~~~~~ > > label:3'1 ^~~~~~~~~~~~~~ > > check:4'0 X~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ error: no match found > > check:4'1 ? possible intended match > > 5: .globl test_fract_f16 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~ > > 6: .p2align 8 > > check:4'0 ~~~~~~~~~~~~ > > 7: .type test_fract_f16,@function > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 8: test_fract_f16: ; @test_fract_f16 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 9: ; %bb.0: > > check:4'0 ~~~~~~~~~ > > 10: s_add_u32 s0, s0, s13 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~ > > 11: s_addc_u32 s1, s1, 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 12: buffer_load_ushort v2, off, s[0:3], 0 offset:4 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 13: s_load_dwordx2 s[4:5], s[6:7], 0x2 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 14: v_lshlrev_b32_e32 v0, 1, v0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 15: v_mov_b32_e32 v1, 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~ > > 16: s_mov_b32 s7, 0x100f000 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~ > > 17: s_mov_b32 s6, 0 > > check:4'0 ~~~~~~~~~~~~~~~~~ > > 18: s_waitcnt vmcnt(0) lgkmcnt(0) > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 19: buffer_store_short v2, v[0:1], s[4:7], 0 addr64 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 20: s_endpgm > > check:4'0 ~~~~~~~~~~ > > 21: .section .rodata,"a",@progbits > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 22: .p2align 6, 0x0 > > check:4'0 ~~~~~~~~~~~~~~~~~ > > 23: .amdhsa_kernel test_fract_f16 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 24: .amdhsa_group_segment_fixed_size 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 25: .amdhsa_private_segment_fixed_size 6 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 26: .amdhsa_kernarg_size 280 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 27: .amdhsa_user_sgpr_count 10 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 28: .amdhsa_user_sgpr_private_segment_buffer 1 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 29: .amdhsa_user_sgpr_dispatch_ptr 1 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 30: .amdhsa_user_sgpr_queue_ptr 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 31: .amdhsa_user_sgpr_kernarg_segment_ptr 1 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 32: .amdhsa_user_sgpr_dispatch_id 1 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 33: .amdhsa_user_sgpr_flat_scratch_init 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 34: .amdhsa_user_sgpr_private_segment_size 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 35: .amdhsa_uses_dynamic_stack 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 36: .amdhsa_system_sgpr_private_segment_wavefront_offset 1 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 37: .amdhsa_system_sgpr_workgroup_id_x 1 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 38: .amdhsa_system_sgpr_workgroup_id_y 1 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 39: .amdhsa_system_sgpr_workgroup_id_z 1 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 40: .amdhsa_system_sgpr_workgroup_info 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 41: .amdhsa_system_vgpr_workitem_id 2 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 42: .amdhsa_next_free_vgpr 3 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 43: .amdhsa_next_free_sgpr 14 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 44: .amdhsa_reserve_vcc 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~ > > 45: .amdhsa_float_round_mode_32 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 46: .amdhsa_float_round_mode_16_64 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 47: .amdhsa_float_denorm_mode_32 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 48: .amdhsa_float_denorm_mode_16_64 3 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 49: .amdhsa_dx10_clamp 1 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 50: .amdhsa_ieee_mode 1 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~ > > 51: .amdhsa_exception_fp_ieee_invalid_op 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 52: .amdhsa_exception_fp_denorm_src 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 53: .amdhsa_exception_fp_ieee_div_zero 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 54: .amdhsa_exception_fp_ieee_overflow 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 55: .amdhsa_exception_fp_ieee_underflow 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 56: .amdhsa_exception_fp_ieee_inexact 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 57: .amdhsa_exception_int_div_zero 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 58: .end_amdhsa_kernel > > check:4'0 ~~~~~~~~~~~~~~~~~~~~ > > 59: .text > > check:4'0 ~~~~~~~ > > 60: .Lfunc_end0: > > check:4'0 ~~~~~~~~~~~~~ > > 61: .size test_fract_f16, .Lfunc_end0-test_fract_f16 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 62: ; -- End function > > check:4'0 ~~~~~~~~~~~~~~~~~~~ > > 63: .section .AMDGPU.csdata,"",@progbits > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 64: ; Kernel info: > > check:4'0 ~~~~~~~~~~~~~~~ > > 65: ; codeLenInByte = 56 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~ > > 66: ; NumSgprs: 14 > > check:4'0 ~~~~~~~~~~~~~~~ > > 67: ; NumVgprs: 3 > > check:4'0 ~~~~~~~~~~~~~~ > > 68: ; ScratchSize: 6 > > check:4'0 ~~~~~~~~~~~~~~~~~ > > 69: ; MemoryBound: 0 > > check:4'0 ~~~~~~~~~~~~~~~~~ > > 70: ; FloatMode: 192 > > check:4'0 ~~~~~~~~~~~~~~~~~ > > 71: ; IeeeMode: 1 > > check:4'0 ~~~~~~~~~~~~~~ > > 72: ; LDSByteSize: 0 bytes/workgroup (compile time only) > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 73: ; SGPRBlocks: 1 > > check:4'0 ~~~~~~~~~~~~~~~~ > > 74: ; VGPRBlocks: 0 > > check:4'0 ~~~~~~~~~~~~~~~~ > > 75: ; NumSGPRsForWavesPerEU: 14 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 76: ; NumVGPRsForWavesPerEU: 3 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 77: ; Occupancy: 10 > > check:4'0 ~~~~~~~~~~~~~~~~ > > 78: ; WaveLimiterHint : 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 79: ; COMPUTE_PGM_RSRC2:SCRATCH_EN: 1 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 80: ; COMPUTE_PGM_RSRC2:USER_SGPR: 10 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 81: ; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 82: ; COMPUTE_PGM_RSRC2:TGID_X_EN: 1 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 83: ; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 84: ; COMPUTE_PGM_RSRC2:TGID_Z_EN: 1 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 85: ; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2 > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 86: .text > > check:4'0 ~~~~~~~ > > 87: .protected test_fract_f32 ; -- Begin function test_fract_f32 > > label:43 ^~~~~~~~~~~~~~ > > check:4'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 88: .globl test_fract_f32 > 89: .p2align 8 > 90: .type test_fract_f32,@function > 91: test_fract_f32: ; @test_fract_f32 > 92: ; %bb.0: > 93: s_load_dwordx2 s[0:1], s[4:5], 0x4 > 94: s_mov_b32 s3, 0x100f000 > 95: s_mov_b32 s2, 0 > 96: v_lshlrev_b32_e32 v0, 2, v0 > 97: v_mov_b32_e32 v1, 0 > 98: s_waitcnt lgkmcnt(0) > 99: buffer_load_dword v2, v[0:1], s[0:3], 0 addr64 > 100: s_load_dwordx4 s[4:7], s[4:5], 0x0 > 101: s_mov_b32 s8, 0x7f800000 > 102: s_waitcnt lgkmcnt(0) > 103: s_mov_b64 s[0:1], s[6:7] > 104: s_mov_b64 s[6:7], s[2:3] > 105: s_waitcnt vmcnt(0) > 106: v_floor_f32_e32 v3, v2 > > dag:44 ^~~~~~~~~~~ > > 107: v_sub_f32_e32 v4, v2, v3 > > dag:45 ^~~~~~~~~ > > 108: buffer_store_dword v3, v[0:1], s[0:3], 0 addr64 > 109: s_waitcnt expcnt(0) > 110: v_min_f32_e32 v3, 0x3f7fffff, v4 > > dag:46 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 111: v_cmp_u_f32_e32 vcc, v2, v2 > > dag:47 ^~~~~~~~~~~ > > 112: v_cndmask_b32_e32 v3, v3, v2, vcc > > dag:48 ^~~~~~~~~~~~~ > > 113: v_cmp_neq_f32_e64 vcc, |v2|, s8 > > dag:49 ^~~~~~~~~~~~~ > > 114: v_cndmask_b32_e32 v2, 0, v3, vcc > > dag:50 ^~~~~~~~~~~~~ > > 115: buffer_store_dword v2, v[0:1], s[4:7], 0 addr64 > 116: s_endpgm > 117: .section .rodata,"a",@progbits > 118: .p2align 6, 0x0 > 119: .amdhsa_kernel test_fract_f32 > 120: .amdhsa_group_segment_fixed_size 0 > 121: .amdhsa_private_segment_fixed_size 0 > 122: .amdhsa_kernarg_size 24 > 123: .amdhsa_user_sgpr_count 6 > 124: .amdhsa_user_sgpr_private_segment_buffer 1 > 125: .amdhsa_user_sgpr_dispatch_ptr 0 > 126: .amdhsa_user_sgpr_queue_ptr 0 > 127: .amdhsa_user_sgpr_kernarg_segment_ptr 1 > 128: .amdhsa_user_sgpr_dispatch_id 0 > 129: .amdhsa_user_sgpr_flat_scratch_init 0 > 130: .amdhsa_user_sgpr_private_segment_size 0 > 131: .amdhsa_uses_dynamic_stack 0 > 132: .amdhsa_system_sgpr_private_segment_wavefront_offset 0 > 133: .amdhsa_system_sgpr_workgroup_id_x 1 > 134: .amdhsa_system_sgpr_workgroup_id_y 0 > 135: .amdhsa_system_sgpr_workgroup_id_z 0 > 136: .amdhsa_system_sgpr_workgroup_info 0 > 137: .amdhsa_system_vgpr_workitem_id 0 > 138: .amdhsa_next_free_vgpr 5 > 139: .amdhsa_next_free_sgpr 9 > 140: .amdhsa_float_round_mode_32 0 > 141: .amdhsa_float_round_mode_16_64 0 > 142: .amdhsa_float_denorm_mode_32 0 > 143: .amdhsa_float_denorm_mode_16_64 3 > 144: .amdhsa_dx10_clamp 1 > 145: .amdhsa_ieee_mode 1 > 146: .amdhsa_exception_fp_ieee_invalid_op 0 > 147: .amdhsa_exception_fp_denorm_src 0 > 148: .amdhsa_exception_fp_ieee_div_zero 0 > 149: .amdhsa_exception_fp_ieee_overflow 0 > 150: .amdhsa_exception_fp_ieee_underflow 0 > 151: .amdhsa_exception_fp_ieee_inexact 0 > 152: .amdhsa_exception_int_div_zero 0 > 153: .end_amdhsa_kernel > 154: .text > 155: .Lfunc_end1: > 156: .size test_fract_f32, .Lfunc_end1-test_fract_f32 > 157: ; -- End function > 158: .section .AMDGPU.csdata,"",@progbits > 159: ; Kernel info: > 160: ; codeLenInByte = 124 > 161: ; NumSgprs: 11 > 162: ; NumVgprs: 5 > 163: ; ScratchSize: 0 > 164: ; MemoryBound: 0 > 165: ; FloatMode: 192 > 166: ; IeeeMode: 1 > 167: ; LDSByteSize: 0 bytes/workgroup (compile time only) > 168: ; SGPRBlocks: 1 > 169: ; VGPRBlocks: 1 > 170: ; NumSGPRsForWavesPerEU: 11 > 171: ; NumVGPRsForWavesPerEU: 5 > 172: ; Occupancy: 10 > 173: ; WaveLimiterHint : 0 > 174: ; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0 > 175: ; COMPUTE_PGM_RSRC2:USER_SGPR: 6 > 176: ; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0 > 177: ; COMPUTE_PGM_RSRC2:TGID_X_EN: 1 > 178: ; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0 > 179: ; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0 > 180: ; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0 > 181: .text > 182: .protected test_fract_f64 ; -- Begin function test_fract_f64 > > label:68'0 ^~~~~~~~~~~~~~ > > label:68'1 ^~~~~~~~~~~~~~ > > 183: .globl test_fract_f64 > 184: .p2align 8 > 185: .type test_fract_f64,@function > 186: test_fract_f64: ; @test_fract_f64 > 187: ; %bb.0: > 188: s_load_dwordx2 s[0:1], s[4:5], 0x4 > 189: s_mov_b32 s3, 0x100f000 > 190: s_mov_b32 s2, 0 > 191: v_lshlrev_b32_e32 v0, 3, v0 > 192: v_mov_b32_e32 v1, 0 > 193: s_waitcnt lgkmcnt(0) > 194: buffer_load_dwordx2 v[2:3], v[0:1], s[0:3], 0 addr64 > 195: v_mov_b32_e32 v4, -1 > 196: v_mov_b32_e32 v5, 0x3fefffff > 197: s_load_dwordx4 s[4:7], s[4:5], 0x0 > 198: s_mov_b32 s8, -1 > 199: s_mov_b32 s9, 0x3fefffff > 200: s_mov_b32 s10, 0 > 201: s_mov_b32 s11, 0x7ff00000 > 202: s_waitcnt lgkmcnt(0) > 203: s_mov_b64 s[0:1], s[6:7] > 204: s_mov_b64 s[6:7], s[2:3] > 205: s_waitcnt vmcnt(0) > 206: v_fract_f64_e32 v[6:7], v[2:3] > > check:71 ^~~~~~~~~~~~~~~ > > 207: v_cmp_class_f64_e64 vcc, v[2:3], 3 > > check:72 ^~~~~~~~~~~~~~~~~~~ > > 208: v_min_f64 v[4:5], v[6:7], v[4:5] > > check:73 ^~~~~~~~~ > > 209: v_cndmask_b32_e32 v4, v4, v2, vcc > > check:74 ^~~~~~~~~~~~~ > > 210: v_cndmask_b32_e32 v5, v5, v3, vcc > > check:75 ^~~~~~~~~~~~~ > > 211: v_add_f64 v[4:5], v[2:3], -v[4:5] > > check:76 ^~~~~~~~~ > > 212: v_cmp_u_f64_e32 vcc, v[2:3], v[2:3] > > check:77 ^~~~~~~~~~~ > > 213: v_add_f64 v[6:7], v[2:3], -v[4:5] > > check:78 ^~~~~~~~~ > > 214: buffer_store_dwordx2 v[4:5], v[0:1], s[0:3], 0 addr64 > 215: s_waitcnt expcnt(0) > 216: v_min_f64 v[4:5], v[6:7], s[8:9] > > check:79 ^~~~~~~~~ > > 217: v_cndmask_b32_e32 v4, v4, v2, vcc > 218: v_cndmask_b32_e32 v5, v5, v3, vcc > 219: v_cmp_neq_f64_e64 vcc, |v[2:3]|, s[10:11] > > check:80 ^~~~~~~~~~~~~ > > 220: v_cndmask_b32_e32 v3, 0, v5, vcc > 221: v_cndmask_b32_e32 v2, 0, v4, vcc > 222: buffer_store_dwordx2 v[2:3], v[0:1], s[4:7], 0 addr64 > 223: s_endpgm > 224: .section .rodata,"a",@progbits > 225: .p2align 6, 0x0 > 226: .amdhsa_kernel test_fract_f64 > 227: .amdhsa_group_segment_fixed_size 0 > 228: .amdhsa_private_segment_fixed_size 0 > 229: .amdhsa_kernarg_size 24 > 230: .amdhsa_user_sgpr_count 6 > 231: .amdhsa_user_sgpr_private_segment_buffer 1 > 232: .amdhsa_user_sgpr_dispatch_ptr 0 > 233: .amdhsa_user_sgpr_queue_ptr 0 > 234: .amdhsa_user_sgpr_kernarg_segment_ptr 1 > 235: .amdhsa_user_sgpr_dispatch_id 0 > 236: .amdhsa_user_sgpr_flat_scratch_init 0 > 237: .amdhsa_user_sgpr_private_segment_size 0 > 238: .amdhsa_uses_dynamic_stack 0 > 239: .amdhsa_system_sgpr_private_segment_wavefront_offset 0 > 240: .amdhsa_system_sgpr_workgroup_id_x 1 > 241: .amdhsa_system_sgpr_workgroup_id_y 0 > 242: .amdhsa_system_sgpr_workgroup_id_z 0 > 243: .amdhsa_system_sgpr_workgroup_info 0 > 244: .amdhsa_system_vgpr_workitem_id 0 > 245: .amdhsa_next_free_vgpr 8 > 246: .amdhsa_next_free_sgpr 12 > 247: .amdhsa_float_round_mode_32 0 > 248: .amdhsa_float_round_mode_16_64 0 > 249: .amdhsa_float_denorm_mode_32 0 > 250: .amdhsa_float_denorm_mode_16_64 3 > 251: .amdhsa_dx10_clamp 1 > 252: .amdhsa_ieee_mode 1 > 253: .amdhsa_exception_fp_ieee_invalid_op 0 > 254: .amdhsa_exception_fp_denorm_src 0 > 255: .amdhsa_exception_fp_ieee_div_zero 0 > 256: .amdhsa_exception_fp_ieee_overflow 0 > 257: .amdhsa_exception_fp_ieee_underflow 0 > 258: .amdhsa_exception_fp_ieee_inexact 0 > 259: .amdhsa_exception_int_div_zero 0 > 260: .end_amdhsa_kernel > 261: .text > 262: .Lfunc_end2: > 263: .size test_fract_f64, .Lfunc_end2-test_fract_f64 > 264: ; -- End function > 265: .section .AMDGPU.csdata,"",@progbits > 266: ; Kernel info: > 267: ; codeLenInByte = 196 > 268: ; NumSgprs: 14 > 269: ; NumVgprs: 8 > 270: ; ScratchSize: 0 > 271: ; MemoryBound: 0 > 272: ; FloatMode: 192 > 273: ; IeeeMode: 1 > 274: ; LDSByteSize: 0 bytes/workgroup (compile time only) > 275: ; SGPRBlocks: 1 > 276: ; VGPRBlocks: 1 > 277: ; NumSGPRsForWavesPerEU: 14 > 278: ; NumVGPRsForWavesPerEU: 8 > 279: ; Occupancy: 10 > 280: ; WaveLimiterHint : 0 > 281: ; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0 > 282: ; COMPUTE_PGM_RSRC2:USER_SGPR: 6 > 283: ; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0 > 284: ; COMPUTE_PGM_RSRC2:TGID_X_EN: 1 > 285: ; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0 > 286: ; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0 > 287: ; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0 > 288: .hidden __oclc_ABI_version ; @__oclc_ABI_version > 289: .type __oclc_ABI_version,@object > 290: .section .rodata,"a",@progbits > 291: .weak __oclc_ABI_version > 292: .p2align 2, 0x0 > 293: __oclc_ABI_version: > 294: .long 500 ; 0x1f4 > 295: .size __oclc_ABI_version, 4 > 296: > 297: .ident "clang version 18.1.8" > 298: .section ".note.GNU-stack","",@progbits > 299: .addrsig > 300: .amdgpu_metadata > 301: --- > 302: amdhsa.kernels: > 303: - .args: > 304: - .actual_access: write_only > 305: .address_space: global > 306: .is_restrict: true > 307: .offset: 0 > 308: .size: 8 > 309: .type_name: 'half*' > 310: .value_kind: global_buffer > 311: - .actual_access: write_only > 312: .address_space: global > 313: .is_restrict: true > 314: .offset: 8 > 315: .size: 8 > 316: .type_name: 'half*' > 317: .value_kind: global_buffer > 318: - .actual_access: read_only > 319: .address_space: global > 320: .is_restrict: true > 321: .offset: 16 > 322: .size: 8 > 323: .type_name: 'half*' > 324: .value_kind: global_buffer > 325: - .offset: 24 > 326: .size: 4 > 327: .value_kind: hidden_block_count_x > 328: - .offset: 28 > 329: .size: 4 > 330: .value_kind: hidden_block_count_y > 331: - .offset: 32 > 332: .size: 4 > 333: .value_kind: hidden_block_count_z > 334: - .offset: 36 > 335: .size: 2 > 336: .value_kind: hidden_group_size_x > 337: - .offset: 38 > 338: .size: 2 > 339: .value_kind: hidden_group_size_y > 340: - .offset: 40 > 341: .size: 2 > 342: .value_kind: hidden_group_size_z > 343: - .offset: 42 > 344: .size: 2 > 345: .value_kind: hidden_remainder_x > 346: - .offset: 44 > 347: .size: 2 > 348: .value_kind: hidden_remainder_y > 349: - .offset: 46 > 350: .size: 2 > 351: .value_kind: hidden_remainder_z > 352: - .offset: 64 > 353: .size: 8 > 354: .value_kind: hidden_global_offset_x > 355: - .offset: 72 > 356: .size: 8 > 357: .value_kind: hidden_global_offset_y > 358: - .offset: 80 > 359: .size: 8 > 360: .value_kind: hidden_global_offset_z > 361: - .offset: 88 > 362: .size: 2 > 363: .value_kind: hidden_grid_dims > 364: - .offset: 104 > 365: .size: 8 > 366: .value_kind: hidden_hostcall_buffer > 367: - .offset: 112 > 368: .size: 8 > 369: .value_kind: hidden_multigrid_sync_arg > 370: - .offset: 120 > 371: .size: 8 > 372: .value_kind: hidden_heap_v1 > 373: - .offset: 128 > 374: .size: 8 > 375: .value_kind: hidden_default_queue > 376: - .offset: 136 > 377: .size: 8 > 378: .value_kind: hidden_completion_action > 379: - .offset: 216 > 380: .size: 4 > 381: .value_kind: hidden_private_base > 382: - .offset: 220 > 383: .size: 4 > 384: .value_kind: hidden_shared_base > 385: - .offset: 224 > 386: .size: 8 > 387: .value_kind: hidden_queue_ptr > 388: .group_segment_fixed_size: 0 > 389: .kernarg_segment_align: 8 > 390: .kernarg_segment_size: 280 > 391: .language: OpenCL C > 392: .language_version: > 393: - 2 > 394: - 0 > 395: .max_flat_workgroup_size: 256 > 396: .name: test_fract_f16 > 397: .private_segment_fixed_size: 6 > 398: .sgpr_count: 14 > 399: .sgpr_spill_count: 0 > 400: .symbol: test_fract_f16.kd > 401: .uses_dynamic_stack: false > 402: .vgpr_count: 3 > 403: .vgpr_spill_count: 0 > 404: .wavefront_size: 64 > 405: - .args: > 406: - .actual_access: write_only > 407: .address_space: global > 408: .is_restrict: true > 409: .offset: 0 > 410: .size: 8 > 411: .type_name: 'float*' > 412: .value_kind: global_buffer > 413: - .actual_access: write_only > 414: .address_space: global > 415: .is_restrict: true > 416: .offset: 8 > 417: .size: 8 > 418: .type_name: 'float*' > 419: .value_kind: global_buffer > 420: - .actual_access: read_only > 421: .address_space: global > 422: .is_restrict: true > 423: .offset: 16 > 424: .size: 8 > 425: .type_name: 'float*' > 426: .value_kind: global_buffer > 427: .group_segment_fixed_size: 0 > 428: .kernarg_segment_align: 8 > 429: .kernarg_segment_size: 24 > 430: .language: OpenCL C > 431: .language_version: > 432: - 2 > 433: - 0 > 434: .max_flat_workgroup_size: 256 > 435: .name: test_fract_f32 > 436: .private_segment_fixed_size: 0 > 437: .sgpr_count: 11 > 438: .sgpr_spill_count: 0 > 439: .symbol: test_fract_f32.kd > 440: .uses_dynamic_stack: false > 441: .vgpr_count: 5 > 442: .vgpr_spill_count: 0 > 443: .wavefront_size: 64 > 444: - .args: > 445: - .actual_access: write_only > 446: .address_space: global > 447: .is_restrict: true > 448: .offset: 0 > 449: .size: 8 > 450: .type_name: 'double*' > 451: .value_kind: global_buffer > 452: - .actual_access: write_only > 453: .address_space: global > 454: .is_restrict: true > 455: .offset: 8 > 456: .size: 8 > 457: .type_name: 'double*' > 458: .value_kind: global_buffer > 459: - .actual_access: read_only > 460: .address_space: global > 461: .is_restrict: true > 462: .offset: 16 > 463: .size: 8 > 464: .type_name: 'double*' > 465: .value_kind: global_buffer > 466: .group_segment_fixed_size: 0 > 467: .kernarg_segment_align: 8 > 468: .kernarg_segment_size: 24 > 469: .language: OpenCL C > 470: .language_version: > 471: - 2 > 472: - 0 > 473: .max_flat_workgroup_size: 256 > 474: .name: test_fract_f64 > 475: .private_segment_fixed_size: 0 > 476: .sgpr_count: 14 > 477: .sgpr_spill_count: 0 > 478: .symbol: test_fract_f64.kd > 479: .uses_dynamic_stack: false > 480: .vgpr_count: 8 > 481: .vgpr_spill_count: 0 > 482: .wavefront_size: 64 > 483: amdhsa.target: amdgcn-amd-amdhsa--gfx600 > 484: amdhsa.version: > 485: - 1 > 486: - 2 > 487: ... > 488: > 489: .end_amdgpu_metadata > > >>>>>> > > > ><end of output> >Test time = 0.22 sec >---------------------------------------------------------- >Test Failed. >"compile_fract__gfx600" end time: Aug 18 19:10 UTC >"compile_fract__gfx600" time elapsed: 00:00:00 >---------------------------------------------------------- > >9/21 Testing: compile_native_rsqrt__gfx600 >9/21 Test: compile_native_rsqrt__gfx600 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.native_rsqrt.gfx600.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/native_rsqrt.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx600" "-DEXTRA_CHECK_PREFIX=GFX600,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_native_rsqrt__gfx600" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.18 sec >---------------------------------------------------------- >Test Passed. >"compile_native_rsqrt__gfx600" end time: Aug 18 19:10 UTC >"compile_native_rsqrt__gfx600" time elapsed: 00:00:00 >---------------------------------------------------------- > >10/21 Testing: compile_native_log__gfx600 >10/21 Test: compile_native_log__gfx600 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.native_log.gfx600.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/native_log.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx600" "-DEXTRA_CHECK_PREFIX=GFX600,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_native_log__gfx600" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.16 sec >---------------------------------------------------------- >Test Passed. >"compile_native_log__gfx600" end time: Aug 18 19:10 UTC >"compile_native_log__gfx600" time elapsed: 00:00:00 >---------------------------------------------------------- > >11/21 Testing: compile_native_exp__gfx600 >11/21 Test: compile_native_exp__gfx600 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.native_exp.gfx600.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/native_exp.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx600" "-DEXTRA_CHECK_PREFIX=GFX600,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_native_exp__gfx600" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.17 sec >---------------------------------------------------------- >Test Passed. >"compile_native_exp__gfx600" end time: Aug 18 19:10 UTC >"compile_native_exp__gfx600" time elapsed: 00:00:00 >---------------------------------------------------------- > >12/21 Testing: compile_fract__gfx700 >12/21 Test: compile_fract__gfx700 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.fract.gfx700.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/fract.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx700" "-DEXTRA_CHECK_PREFIX=GFX700,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_fract__gfx700" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- >CMake Error at /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake:37 (message): > Error in test output: > /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/fract.cl:16:12: > error: GFX700: expected string not found in input > > // GFX700: flat_load_ushort [[VAL:v[0-9]+]] > > ^ > > output.fract.gfx700.s:4:27: note: scanning from here > > .protected test_fract_f16 ; -- Begin function test_fract_f16 > ^ > > output.fract.gfx700.s:23:2: note: possible intended match here > > flat_store_short v[0:1], v2 > ^ > > > > Input file: output.fract.gfx700.s > > Check file: > /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/fract.cl > > > > > -dump-input=help explains the following input dump. > > > > Input was: > > <<<<<< > > 1: .text > 2: .amdgcn_target "amdgcn-amd-amdhsa--gfx700" > 3: .amdhsa_code_object_version 5 > 4: .protected test_fract_f16 ; -- Begin function test_fract_f16 > > label:3'0 ^~~~~~~~~~~~~~ > > label:3'1 ^~~~~~~~~~~~~~ > > check:16'0 X~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ error: no match found > > 5: .globl test_fract_f16 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~ > > 6: .p2align 8 > > check:16'0 ~~~~~~~~~~~~ > > 7: .type test_fract_f16,@function > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 8: test_fract_f16: ; @test_fract_f16 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 9: ; %bb.0: > > check:16'0 ~~~~~~~~~ > > 10: s_add_i32 s10, s10, s15 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~ > > 11: s_lshr_b32 flat_scratch_hi, s10, 8 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 12: s_add_u32 s0, s0, s15 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~ > > 13: s_addc_u32 s1, s1, 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 14: buffer_load_ushort v2, off, s[0:3], 0 offset:4 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 15: s_load_dwordx2 s[4:5], s[6:7], 0x2 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 16: v_lshlrev_b32_e32 v0, 1, v0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 17: s_mov_b32 flat_scratch_lo, s11 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 18: s_waitcnt lgkmcnt(0) > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 19: v_mov_b32_e32 v1, s5 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 20: v_add_i32_e32 v0, vcc, s4, v0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 21: v_addc_u32_e32 v1, vcc, 0, v1, vcc > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 22: s_waitcnt vmcnt(0) > > check:16'0 ~~~~~~~~~~~~~~~~~~~~ > > 23: flat_store_short v[0:1], v2 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > check:16'1 ? possible intended match > > 24: s_endpgm > > check:16'0 ~~~~~~~~~~ > > 25: .section .rodata,"a",@progbits > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 26: .p2align 6, 0x0 > > check:16'0 ~~~~~~~~~~~~~~~~~ > > 27: .amdhsa_kernel test_fract_f16 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 28: .amdhsa_group_segment_fixed_size 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 29: .amdhsa_private_segment_fixed_size 6 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 30: .amdhsa_kernarg_size 280 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 31: .amdhsa_user_sgpr_count 12 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 32: .amdhsa_user_sgpr_private_segment_buffer 1 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 33: .amdhsa_user_sgpr_dispatch_ptr 1 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 34: .amdhsa_user_sgpr_queue_ptr 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 35: .amdhsa_user_sgpr_kernarg_segment_ptr 1 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 36: .amdhsa_user_sgpr_dispatch_id 1 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 37: .amdhsa_user_sgpr_flat_scratch_init 1 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 38: .amdhsa_user_sgpr_private_segment_size 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 39: .amdhsa_uses_dynamic_stack 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 40: .amdhsa_system_sgpr_private_segment_wavefront_offset 1 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 41: .amdhsa_system_sgpr_workgroup_id_x 1 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 42: .amdhsa_system_sgpr_workgroup_id_y 1 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 43: .amdhsa_system_sgpr_workgroup_id_z 1 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 44: .amdhsa_system_sgpr_workgroup_info 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 45: .amdhsa_system_vgpr_workitem_id 2 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 46: .amdhsa_next_free_vgpr 3 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 47: .amdhsa_next_free_sgpr 16 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 48: .amdhsa_float_round_mode_32 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 49: .amdhsa_float_round_mode_16_64 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 50: .amdhsa_float_denorm_mode_32 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 51: .amdhsa_float_denorm_mode_16_64 3 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 52: .amdhsa_dx10_clamp 1 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 53: .amdhsa_ieee_mode 1 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~ > > 54: .amdhsa_exception_fp_ieee_invalid_op 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 55: .amdhsa_exception_fp_denorm_src 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 56: .amdhsa_exception_fp_ieee_div_zero 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 57: .amdhsa_exception_fp_ieee_overflow 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 58: .amdhsa_exception_fp_ieee_underflow 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 59: .amdhsa_exception_fp_ieee_inexact 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 60: .amdhsa_exception_int_div_zero 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 61: .end_amdhsa_kernel > > check:16'0 ~~~~~~~~~~~~~~~~~~~~ > > 62: .text > > check:16'0 ~~~~~~~ > > 63: .Lfunc_end0: > > check:16'0 ~~~~~~~~~~~~~ > > 64: .size test_fract_f16, .Lfunc_end0-test_fract_f16 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 65: ; -- End function > > check:16'0 ~~~~~~~~~~~~~~~~~~~ > > 66: .section .AMDGPU.csdata,"",@progbits > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 67: ; Kernel info: > > check:16'0 ~~~~~~~~~~~~~~~ > > 68: ; codeLenInByte = 68 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~ > > 69: ; NumSgprs: 20 > > check:16'0 ~~~~~~~~~~~~~~~ > > 70: ; NumVgprs: 3 > > check:16'0 ~~~~~~~~~~~~~~ > > 71: ; ScratchSize: 6 > > check:16'0 ~~~~~~~~~~~~~~~~~ > > 72: ; MemoryBound: 0 > > check:16'0 ~~~~~~~~~~~~~~~~~ > > 73: ; FloatMode: 192 > > check:16'0 ~~~~~~~~~~~~~~~~~ > > 74: ; IeeeMode: 1 > > check:16'0 ~~~~~~~~~~~~~~ > > 75: ; LDSByteSize: 0 bytes/workgroup (compile time only) > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 76: ; SGPRBlocks: 2 > > check:16'0 ~~~~~~~~~~~~~~~~ > > 77: ; VGPRBlocks: 0 > > check:16'0 ~~~~~~~~~~~~~~~~ > > 78: ; NumSGPRsForWavesPerEU: 20 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 79: ; NumVGPRsForWavesPerEU: 3 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 80: ; Occupancy: 10 > > check:16'0 ~~~~~~~~~~~~~~~~ > > 81: ; WaveLimiterHint : 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 82: ; COMPUTE_PGM_RSRC2:SCRATCH_EN: 1 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 83: ; COMPUTE_PGM_RSRC2:USER_SGPR: 12 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 84: ; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 85: ; COMPUTE_PGM_RSRC2:TGID_X_EN: 1 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 86: ; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 87: ; COMPUTE_PGM_RSRC2:TGID_Z_EN: 1 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 88: ; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2 > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 89: .text > > check:16'0 ~~~~~~~ > > 90: .protected test_fract_f32 ; -- Begin function test_fract_f32 > > label:43 ^~~~~~~~~~~~~~ > > check:16'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 91: .globl test_fract_f32 > 92: .p2align 8 > 93: .type test_fract_f32,@function > 94: test_fract_f32: ; @test_fract_f32 > 95: ; %bb.0: > 96: s_load_dwordx2 s[0:1], s[4:5], 0x4 > 97: v_lshlrev_b32_e32 v2, 2, v0 > 98: s_waitcnt lgkmcnt(0) > 99: v_mov_b32_e32 v1, s1 > 100: v_add_i32_e32 v0, vcc, s0, v2 > 101: v_addc_u32_e32 v1, vcc, 0, v1, vcc > 102: flat_load_dword v3, v[0:1] > 103: s_load_dwordx4 s[0:3], s[4:5], 0x0 > 104: s_mov_b32 s4, 0x7f800000 > 105: s_waitcnt lgkmcnt(0) > 106: v_mov_b32_e32 v1, s3 > 107: v_add_i32_e32 v0, vcc, s2, v2 > 108: v_addc_u32_e32 v1, vcc, 0, v1, vcc > 109: v_mov_b32_e32 v4, s1 > 110: s_waitcnt vmcnt(0) > 111: v_floor_f32_e32 v5, v3 > 112: v_fract_f32_e32 v6, v3 > 113: v_cmp_neq_f32_e64 vcc, |v3|, s4 > 114: v_cndmask_b32_e32 v3, 0, v6, vcc > 115: flat_store_dword v[0:1], v5 > 116: v_add_i32_e32 v0, vcc, s0, v2 > 117: v_addc_u32_e32 v1, vcc, 0, v4, vcc > 118: flat_store_dword v[0:1], v3 > 119: s_endpgm > 120: .section .rodata,"a",@progbits > 121: .p2align 6, 0x0 > 122: .amdhsa_kernel test_fract_f32 > 123: .amdhsa_group_segment_fixed_size 0 > 124: .amdhsa_private_segment_fixed_size 0 > 125: .amdhsa_kernarg_size 24 > 126: .amdhsa_user_sgpr_count 6 > 127: .amdhsa_user_sgpr_private_segment_buffer 1 > 128: .amdhsa_user_sgpr_dispatch_ptr 0 > 129: .amdhsa_user_sgpr_queue_ptr 0 > 130: .amdhsa_user_sgpr_kernarg_segment_ptr 1 > 131: .amdhsa_user_sgpr_dispatch_id 0 > 132: .amdhsa_user_sgpr_flat_scratch_init 0 > 133: .amdhsa_user_sgpr_private_segment_size 0 > 134: .amdhsa_uses_dynamic_stack 0 > 135: .amdhsa_system_sgpr_private_segment_wavefront_offset 0 > 136: .amdhsa_system_sgpr_workgroup_id_x 1 > 137: .amdhsa_system_sgpr_workgroup_id_y 0 > 138: .amdhsa_system_sgpr_workgroup_id_z 0 > 139: .amdhsa_system_sgpr_workgroup_info 0 > 140: .amdhsa_system_vgpr_workitem_id 0 > 141: .amdhsa_next_free_vgpr 7 > 142: .amdhsa_next_free_sgpr 6 > 143: .amdhsa_reserve_flat_scratch 0 > 144: .amdhsa_float_round_mode_32 0 > 145: .amdhsa_float_round_mode_16_64 0 > 146: .amdhsa_float_denorm_mode_32 0 > 147: .amdhsa_float_denorm_mode_16_64 3 > 148: .amdhsa_dx10_clamp 1 > 149: .amdhsa_ieee_mode 1 > 150: .amdhsa_exception_fp_ieee_invalid_op 0 > 151: .amdhsa_exception_fp_denorm_src 0 > 152: .amdhsa_exception_fp_ieee_div_zero 0 > 153: .amdhsa_exception_fp_ieee_overflow 0 > 154: .amdhsa_exception_fp_ieee_underflow 0 > 155: .amdhsa_exception_fp_ieee_inexact 0 > 156: .amdhsa_exception_int_div_zero 0 > 157: .end_amdhsa_kernel > 158: .text > 159: .Lfunc_end1: > 160: .size test_fract_f32, .Lfunc_end1-test_fract_f32 > 161: ; -- End function > 162: .section .AMDGPU.csdata,"",@progbits > 163: ; Kernel info: > 164: ; codeLenInByte = 116 > 165: ; NumSgprs: 8 > 166: ; NumVgprs: 7 > 167: ; ScratchSize: 0 > 168: ; MemoryBound: 0 > 169: ; FloatMode: 192 > 170: ; IeeeMode: 1 > 171: ; LDSByteSize: 0 bytes/workgroup (compile time only) > 172: ; SGPRBlocks: 0 > 173: ; VGPRBlocks: 1 > 174: ; NumSGPRsForWavesPerEU: 8 > 175: ; NumVGPRsForWavesPerEU: 7 > 176: ; Occupancy: 10 > 177: ; WaveLimiterHint : 0 > 178: ; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0 > 179: ; COMPUTE_PGM_RSRC2:USER_SGPR: 6 > 180: ; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0 > 181: ; COMPUTE_PGM_RSRC2:TGID_X_EN: 1 > 182: ; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0 > 183: ; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0 > 184: ; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0 > 185: .text > 186: .protected test_fract_f64 ; -- Begin function test_fract_f64 > > label:68'0 ^~~~~~~~~~~~~~ > > label:68'1 ^~~~~~~~~~~~~~ > > 187: .globl test_fract_f64 > 188: .p2align 8 > 189: .type test_fract_f64,@function > 190: test_fract_f64: ; @test_fract_f64 > 191: ; %bb.0: > 192: s_load_dwordx2 s[0:1], s[4:5], 0x4 > 193: v_lshlrev_b32_e32 v6, 3, v0 > 194: s_load_dwordx4 s[4:7], s[4:5], 0x0 > 195: s_waitcnt lgkmcnt(0) > 196: v_mov_b32_e32 v1, s1 > 197: v_add_i32_e32 v0, vcc, s0, v6 > 198: v_addc_u32_e32 v1, vcc, 0, v1, vcc > 199: flat_load_dwordx2 v[0:1], v[0:1] > > check:83'0 ^~~~~~~~~~~~~~~~~~~~~~~~ > > check:83'1 ^~~~~~ captured var "VAL" > > 200: s_mov_b32 s0, 0 > > dag:89'0 ^~~~~~~~~~~~~~~ > > dag:89'1 ^ captured var "INF_LO" > > 201: s_mov_b32 s1, 0x7ff00000 > > dag:88'0 ^~~~~~~~~~~~~~~~~~~~~~~~ > > dag:88'1 ^ captured var "INF_HI" > > 202: v_mov_b32_e32 v5, s7 > 203: v_mov_b32_e32 v7, s5 > 204: s_waitcnt vmcnt(0) > 205: v_fract_f64_e32 v[2:3], v[0:1] > > dag:86'0 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > dag:86'1 with "VAL" equal to "v\\[0:1\\]" > > dag:86'2 ^ captured var "FRACT_LO" > > dag:86'3 ^ captured var "FRACT_HI" > > 206: v_cmp_neq_f64_e64 vcc, |v[0:1]|, s[0:1] > > dag:90'0 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > dag:90'1 with "VAL" equal to "v\\[0:1\\]" > > dag:90'2 with "INF_LO" equal to "0" > > dag:90'3 with "INF_HI" equal to "1" > > dag:90'4 ^~~ captured var "FINITE" > > 207: v_floor_f64_e32 v[0:1], v[0:1] > > dag:84'0 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > dag:84'1 with "VAL" equal to "v\\[0:1\\]" > > dag:84'2 ^~~~~~ captured var "FLOOR" > > 208: v_add_i32_e64 v4, s[0:1], s6, v6 > 209: v_addc_u32_e64 v5, s[0:1], 0, v5, s[0:1] > 210: v_cndmask_b32_e32 v3, 0, v3, vcc > > dag:93'0 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ > > dag:93'1 with "FRACT_HI" equal to "3" > > dag:93'2 ^ captured var "SELECT1" > > 211: v_cndmask_b32_e32 v2, 0, v2, vcc > > dag:92'0 ^~~~~~~~~~~~~~~~~~~~~~~~~~~ > > dag:92'1 with "FRACT_LO" equal to "2" > > dag:92'2 ^ captured var "SELECT0" > > 212: flat_store_dwordx2 v[4:5], v[0:1] > > check:94'0 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > check:94'1 with "FLOOR" equal to "v\\[0:1\\]" > > 213: v_add_i32_e32 v0, vcc, s4, v6 > 214: v_addc_u32_e32 v1, vcc, 0, v7, vcc > 215: flat_store_dwordx2 v[0:1], v[2:3] > > check:95'0 ^~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > check:95'1 with "SELECT0" equal to "2" > > check:95'2 with "SELECT1" equal to "3" > > 216: s_endpgm > 217: .section .rodata,"a",@progbits > 218: .p2align 6, 0x0 > 219: .amdhsa_kernel test_fract_f64 > 220: .amdhsa_group_segment_fixed_size 0 > 221: .amdhsa_private_segment_fixed_size 0 > 222: .amdhsa_kernarg_size 24 > 223: .amdhsa_user_sgpr_count 6 > 224: .amdhsa_user_sgpr_private_segment_buffer 1 > 225: .amdhsa_user_sgpr_dispatch_ptr 0 > 226: .amdhsa_user_sgpr_queue_ptr 0 > 227: .amdhsa_user_sgpr_kernarg_segment_ptr 1 > 228: .amdhsa_user_sgpr_dispatch_id 0 > 229: .amdhsa_user_sgpr_flat_scratch_init 0 > 230: .amdhsa_user_sgpr_private_segment_size 0 > 231: .amdhsa_uses_dynamic_stack 0 > 232: .amdhsa_system_sgpr_private_segment_wavefront_offset 0 > 233: .amdhsa_system_sgpr_workgroup_id_x 1 > 234: .amdhsa_system_sgpr_workgroup_id_y 0 > 235: .amdhsa_system_sgpr_workgroup_id_z 0 > 236: .amdhsa_system_sgpr_workgroup_info 0 > 237: .amdhsa_system_vgpr_workitem_id 0 > 238: .amdhsa_next_free_vgpr 8 > 239: .amdhsa_next_free_sgpr 8 > 240: .amdhsa_reserve_flat_scratch 0 > 241: .amdhsa_float_round_mode_32 0 > 242: .amdhsa_float_round_mode_16_64 0 > 243: .amdhsa_float_denorm_mode_32 0 > 244: .amdhsa_float_denorm_mode_16_64 3 > 245: .amdhsa_dx10_clamp 1 > 246: .amdhsa_ieee_mode 1 > 247: .amdhsa_exception_fp_ieee_invalid_op 0 > 248: .amdhsa_exception_fp_denorm_src 0 > 249: .amdhsa_exception_fp_ieee_div_zero 0 > 250: .amdhsa_exception_fp_ieee_overflow 0 > 251: .amdhsa_exception_fp_ieee_underflow 0 > 252: .amdhsa_exception_fp_ieee_inexact 0 > 253: .amdhsa_exception_int_div_zero 0 > 254: .end_amdhsa_kernel > 255: .text > 256: .Lfunc_end2: > 257: .size test_fract_f64, .Lfunc_end2-test_fract_f64 > 258: ; -- End function > 259: .section .AMDGPU.csdata,"",@progbits > 260: ; Kernel info: > 261: ; codeLenInByte = 128 > 262: ; NumSgprs: 10 > 263: ; NumVgprs: 8 > 264: ; ScratchSize: 0 > 265: ; MemoryBound: 0 > 266: ; FloatMode: 192 > 267: ; IeeeMode: 1 > 268: ; LDSByteSize: 0 bytes/workgroup (compile time only) > 269: ; SGPRBlocks: 1 > 270: ; VGPRBlocks: 1 > 271: ; NumSGPRsForWavesPerEU: 10 > 272: ; NumVGPRsForWavesPerEU: 8 > 273: ; Occupancy: 10 > 274: ; WaveLimiterHint : 0 > 275: ; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0 > 276: ; COMPUTE_PGM_RSRC2:USER_SGPR: 6 > 277: ; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0 > 278: ; COMPUTE_PGM_RSRC2:TGID_X_EN: 1 > 279: ; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0 > 280: ; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0 > 281: ; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0 > 282: .hidden __oclc_ABI_version ; @__oclc_ABI_version > 283: .type __oclc_ABI_version,@object > 284: .section .rodata,"a",@progbits > 285: .weak __oclc_ABI_version > 286: .p2align 2, 0x0 > 287: __oclc_ABI_version: > 288: .long 500 ; 0x1f4 > 289: .size __oclc_ABI_version, 4 > 290: > 291: .ident "clang version 18.1.8" > 292: .section ".note.GNU-stack","",@progbits > 293: .addrsig > 294: .amdgpu_metadata > 295: --- > 296: amdhsa.kernels: > 297: - .args: > 298: - .actual_access: write_only > 299: .address_space: global > 300: .is_restrict: true > 301: .offset: 0 > 302: .size: 8 > 303: .type_name: 'half*' > 304: .value_kind: global_buffer > 305: - .actual_access: write_only > 306: .address_space: global > 307: .is_restrict: true > 308: .offset: 8 > 309: .size: 8 > 310: .type_name: 'half*' > 311: .value_kind: global_buffer > 312: - .actual_access: read_only > 313: .address_space: global > 314: .is_restrict: true > 315: .offset: 16 > 316: .size: 8 > 317: .type_name: 'half*' > 318: .value_kind: global_buffer > 319: - .offset: 24 > 320: .size: 4 > 321: .value_kind: hidden_block_count_x > 322: - .offset: 28 > 323: .size: 4 > 324: .value_kind: hidden_block_count_y > 325: - .offset: 32 > 326: .size: 4 > 327: .value_kind: hidden_block_count_z > 328: - .offset: 36 > 329: .size: 2 > 330: .value_kind: hidden_group_size_x > 331: - .offset: 38 > 332: .size: 2 > 333: .value_kind: hidden_group_size_y > 334: - .offset: 40 > 335: .size: 2 > 336: .value_kind: hidden_group_size_z > 337: - .offset: 42 > 338: .size: 2 > 339: .value_kind: hidden_remainder_x > 340: - .offset: 44 > 341: .size: 2 > 342: .value_kind: hidden_remainder_y > 343: - .offset: 46 > 344: .size: 2 > 345: .value_kind: hidden_remainder_z > 346: - .offset: 64 > 347: .size: 8 > 348: .value_kind: hidden_global_offset_x > 349: - .offset: 72 > 350: .size: 8 > 351: .value_kind: hidden_global_offset_y > 352: - .offset: 80 > 353: .size: 8 > 354: .value_kind: hidden_global_offset_z > 355: - .offset: 88 > 356: .size: 2 > 357: .value_kind: hidden_grid_dims > 358: - .offset: 104 > 359: .size: 8 > 360: .value_kind: hidden_hostcall_buffer > 361: - .offset: 112 > 362: .size: 8 > 363: .value_kind: hidden_multigrid_sync_arg > 364: - .offset: 120 > 365: .size: 8 > 366: .value_kind: hidden_heap_v1 > 367: - .offset: 128 > 368: .size: 8 > 369: .value_kind: hidden_default_queue > 370: - .offset: 136 > 371: .size: 8 > 372: .value_kind: hidden_completion_action > 373: - .offset: 216 > 374: .size: 4 > 375: .value_kind: hidden_private_base > 376: - .offset: 220 > 377: .size: 4 > 378: .value_kind: hidden_shared_base > 379: - .offset: 224 > 380: .size: 8 > 381: .value_kind: hidden_queue_ptr > 382: .group_segment_fixed_size: 0 > 383: .kernarg_segment_align: 8 > 384: .kernarg_segment_size: 280 > 385: .language: OpenCL C > 386: .language_version: > 387: - 2 > 388: - 0 > 389: .max_flat_workgroup_size: 256 > 390: .name: test_fract_f16 > 391: .private_segment_fixed_size: 6 > 392: .sgpr_count: 20 > 393: .sgpr_spill_count: 0 > 394: .symbol: test_fract_f16.kd > 395: .uses_dynamic_stack: false > 396: .vgpr_count: 3 > 397: .vgpr_spill_count: 0 > 398: .wavefront_size: 64 > 399: - .args: > 400: - .actual_access: write_only > 401: .address_space: global > 402: .is_restrict: true > 403: .offset: 0 > 404: .size: 8 > 405: .type_name: 'float*' > 406: .value_kind: global_buffer > 407: - .actual_access: write_only > 408: .address_space: global > 409: .is_restrict: true > 410: .offset: 8 > 411: .size: 8 > 412: .type_name: 'float*' > 413: .value_kind: global_buffer > 414: - .actual_access: read_only > 415: .address_space: global > 416: .is_restrict: true > 417: .offset: 16 > 418: .size: 8 > 419: .type_name: 'float*' > 420: .value_kind: global_buffer > 421: .group_segment_fixed_size: 0 > 422: .kernarg_segment_align: 8 > 423: .kernarg_segment_size: 24 > 424: .language: OpenCL C > 425: .language_version: > 426: - 2 > 427: - 0 > 428: .max_flat_workgroup_size: 256 > 429: .name: test_fract_f32 > 430: .private_segment_fixed_size: 0 > 431: .sgpr_count: 8 > 432: .sgpr_spill_count: 0 > 433: .symbol: test_fract_f32.kd > 434: .uses_dynamic_stack: false > 435: .vgpr_count: 7 > 436: .vgpr_spill_count: 0 > 437: .wavefront_size: 64 > 438: - .args: > 439: - .actual_access: write_only > 440: .address_space: global > 441: .is_restrict: true > 442: .offset: 0 > 443: .size: 8 > 444: .type_name: 'double*' > 445: .value_kind: global_buffer > 446: - .actual_access: write_only > 447: .address_space: global > 448: .is_restrict: true > 449: .offset: 8 > 450: .size: 8 > 451: .type_name: 'double*' > 452: .value_kind: global_buffer > 453: - .actual_access: read_only > 454: .address_space: global > 455: .is_restrict: true > 456: .offset: 16 > 457: .size: 8 > 458: .type_name: 'double*' > 459: .value_kind: global_buffer > 460: .group_segment_fixed_size: 0 > 461: .kernarg_segment_align: 8 > 462: .kernarg_segment_size: 24 > 463: .language: OpenCL C > 464: .language_version: > 465: - 2 > 466: - 0 > 467: .max_flat_workgroup_size: 256 > 468: .name: test_fract_f64 > 469: .private_segment_fixed_size: 0 > 470: .sgpr_count: 10 > 471: .sgpr_spill_count: 0 > 472: .symbol: test_fract_f64.kd > 473: .uses_dynamic_stack: false > 474: .vgpr_count: 8 > 475: .vgpr_spill_count: 0 > 476: .wavefront_size: 64 > 477: amdhsa.target: amdgcn-amd-amdhsa--gfx700 > 478: amdhsa.version: > 479: - 1 > 480: - 2 > 481: ... > 482: > 483: .end_amdgpu_metadata > > >>>>>> > > > ><end of output> >Test time = 0.26 sec >---------------------------------------------------------- >Test Failed. >"compile_fract__gfx700" end time: Aug 18 19:10 UTC >"compile_fract__gfx700" time elapsed: 00:00:00 >---------------------------------------------------------- > >13/21 Testing: compile_native_rcp__gfx700 >13/21 Test: compile_native_rcp__gfx700 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.native_rcp.gfx700.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/native_rcp.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx700" "-DEXTRA_CHECK_PREFIX=GFX700,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_native_rcp__gfx700" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.18 sec >---------------------------------------------------------- >Test Passed. >"compile_native_rcp__gfx700" end time: Aug 18 19:10 UTC >"compile_native_rcp__gfx700" time elapsed: 00:00:00 >---------------------------------------------------------- > >14/21 Testing: compile_native_rsqrt__gfx700 >14/21 Test: compile_native_rsqrt__gfx700 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.native_rsqrt.gfx700.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/native_rsqrt.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx700" "-DEXTRA_CHECK_PREFIX=GFX700,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_native_rsqrt__gfx700" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.17 sec >---------------------------------------------------------- >Test Passed. >"compile_native_rsqrt__gfx700" end time: Aug 18 19:10 UTC >"compile_native_rsqrt__gfx700" time elapsed: 00:00:00 >---------------------------------------------------------- > >15/21 Testing: compile_native_log__gfx700 >15/21 Test: compile_native_log__gfx700 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.native_log.gfx700.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/native_log.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx700" "-DEXTRA_CHECK_PREFIX=GFX700,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_native_log__gfx700" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.18 sec >---------------------------------------------------------- >Test Passed. >"compile_native_log__gfx700" end time: Aug 18 19:10 UTC >"compile_native_log__gfx700" time elapsed: 00:00:00 >---------------------------------------------------------- > >16/21 Testing: compile_native_exp__gfx700 >16/21 Test: compile_native_exp__gfx700 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.native_exp.gfx700.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/native_exp.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx700" "-DEXTRA_CHECK_PREFIX=GFX700,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_native_exp__gfx700" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.19 sec >---------------------------------------------------------- >Test Passed. >"compile_native_exp__gfx700" end time: Aug 18 19:10 UTC >"compile_native_exp__gfx700" time elapsed: 00:00:00 >---------------------------------------------------------- > >17/21 Testing: compile_fract__gfx803 >17/21 Test: compile_fract__gfx803 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.fract.gfx803.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/fract.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx803" "-DEXTRA_CHECK_PREFIX=GFX803,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_fract__gfx803" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.19 sec >---------------------------------------------------------- >Test Passed. >"compile_fract__gfx803" end time: Aug 18 19:10 UTC >"compile_fract__gfx803" time elapsed: 00:00:00 >---------------------------------------------------------- > >18/21 Testing: compile_native_rcp__gfx803 >18/21 Test: compile_native_rcp__gfx803 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.native_rcp.gfx803.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/native_rcp.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx803" "-DEXTRA_CHECK_PREFIX=GFX803,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_native_rcp__gfx803" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.20 sec >---------------------------------------------------------- >Test Passed. >"compile_native_rcp__gfx803" end time: Aug 18 19:10 UTC >"compile_native_rcp__gfx803" time elapsed: 00:00:00 >---------------------------------------------------------- > >19/21 Testing: compile_native_rsqrt__gfx803 >19/21 Test: compile_native_rsqrt__gfx803 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.native_rsqrt.gfx803.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/native_rsqrt.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx803" "-DEXTRA_CHECK_PREFIX=GFX803,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_native_rsqrt__gfx803" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- >CMake Error at /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake:37 (message): > Error in test output: > /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/native_rsqrt.cl:17:12: > error: GFX803: expected string not found in input > > // GFX803: v_rsq_f16 > > ^ > > output.native_rsqrt.gfx803.s:8:23: note: scanning from here > > test_native_rsqrt_f16: ; @test_native_rsqrt_f16 > > ^ > > output.native_rsqrt.gfx803.s:20:2: note: possible intended match here > > v_rcp_f16_e32 v3, v0 > ^ > > > > Input file: output.native_rsqrt.gfx803.s > > Check file: > /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/native_rsqrt.cl > > > > > -dump-input=help explains the following input dump. > > > > Input was: > > <<<<<< > > 1: .text > 2: .amdgcn_target "amdgcn-amd-amdhsa--gfx803" > 3: .amdhsa_code_object_version 5 > 4: .protected test_native_rsqrt_f16 ; -- Begin function test_native_rsqrt_f16 > 5: .globl test_native_rsqrt_f16 > 6: .p2align 8 > 7: .type test_native_rsqrt_f16,@function > 8: test_native_rsqrt_f16: ; @test_native_rsqrt_f16 > > label:10'0 ^~~~~~~~~~~~~~~~~~~~~~ > > label:10'1 ^~~~~~~~~~~~~~~~~~~~~~ > > check:17'0 X~~~~~~~~~~~~~~~~~~~~~~~~~ error: no match found > > 9: ; %bb.0: > > check:17'0 ~~~~~~~~~ > > 10: s_load_dwordx4 s[0:3], s[4:5], 0x0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 11: v_lshlrev_b32_e32 v2, 1, v0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 12: s_waitcnt lgkmcnt(0) > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 13: v_mov_b32_e32 v1, s3 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 14: v_add_u32_e32 v0, vcc, s2, v2 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 15: v_addc_u32_e32 v1, vcc, 0, v1, vcc > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 16: flat_load_ushort v0, v[0:1] > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 17: v_mov_b32_e32 v1, s1 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 18: s_waitcnt vmcnt(0) > > check:17'0 ~~~~~~~~~~~~~~~~~~~~ > > 19: v_sqrt_f16_e32 v0, v0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~ > > 20: v_rcp_f16_e32 v3, v0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~ > > check:17'1 ? possible intended match > > 21: v_add_u32_e32 v0, vcc, s0, v2 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 22: v_addc_u32_e32 v1, vcc, 0, v1, vcc > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 23: flat_store_short v[0:1], v3 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 24: s_endpgm > > check:17'0 ~~~~~~~~~~ > > 25: .section .rodata,"a",@progbits > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 26: .p2align 6, 0x0 > > check:17'0 ~~~~~~~~~~~~~~~~~ > > 27: .amdhsa_kernel test_native_rsqrt_f16 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 28: .amdhsa_group_segment_fixed_size 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 29: .amdhsa_private_segment_fixed_size 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 30: .amdhsa_kernarg_size 16 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~ > > 31: .amdhsa_user_sgpr_count 6 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 32: .amdhsa_user_sgpr_private_segment_buffer 1 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 33: .amdhsa_user_sgpr_dispatch_ptr 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 34: .amdhsa_user_sgpr_queue_ptr 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 35: .amdhsa_user_sgpr_kernarg_segment_ptr 1 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 36: .amdhsa_user_sgpr_dispatch_id 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 37: .amdhsa_user_sgpr_flat_scratch_init 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 38: .amdhsa_user_sgpr_private_segment_size 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 39: .amdhsa_uses_dynamic_stack 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 40: .amdhsa_system_sgpr_private_segment_wavefront_offset 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 41: .amdhsa_system_sgpr_workgroup_id_x 1 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 42: .amdhsa_system_sgpr_workgroup_id_y 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 43: .amdhsa_system_sgpr_workgroup_id_z 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 44: .amdhsa_system_sgpr_workgroup_info 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 45: .amdhsa_system_vgpr_workitem_id 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 46: .amdhsa_next_free_vgpr 4 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 47: .amdhsa_next_free_sgpr 6 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 48: .amdhsa_reserve_flat_scratch 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 49: .amdhsa_float_round_mode_32 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 50: .amdhsa_float_round_mode_16_64 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 51: .amdhsa_float_denorm_mode_32 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 52: .amdhsa_float_denorm_mode_16_64 3 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 53: .amdhsa_dx10_clamp 1 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 54: .amdhsa_ieee_mode 1 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~ > > 55: .amdhsa_exception_fp_ieee_invalid_op 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 56: .amdhsa_exception_fp_denorm_src 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 57: .amdhsa_exception_fp_ieee_div_zero 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 58: .amdhsa_exception_fp_ieee_overflow 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 59: .amdhsa_exception_fp_ieee_underflow 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 60: .amdhsa_exception_fp_ieee_inexact 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 61: .amdhsa_exception_int_div_zero 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 62: .end_amdhsa_kernel > > check:17'0 ~~~~~~~~~~~~~~~~~~~~ > > 63: .text > > check:17'0 ~~~~~~~ > > 64: .Lfunc_end0: > > check:17'0 ~~~~~~~~~~~~~ > > 65: .size test_native_rsqrt_f16, .Lfunc_end0-test_native_rsqrt_f16 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 66: ; -- End function > > check:17'0 ~~~~~~~~~~~~~~~~~~~ > > 67: .section .AMDGPU.csdata,"",@progbits > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 68: ; Kernel info: > > check:17'0 ~~~~~~~~~~~~~~~ > > 69: ; codeLenInByte = 72 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~ > > 70: ; NumSgprs: 8 > > check:17'0 ~~~~~~~~~~~~~~ > > 71: ; NumVgprs: 4 > > check:17'0 ~~~~~~~~~~~~~~ > > 72: ; ScratchSize: 0 > > check:17'0 ~~~~~~~~~~~~~~~~~ > > 73: ; MemoryBound: 0 > > check:17'0 ~~~~~~~~~~~~~~~~~ > > 74: ; FloatMode: 192 > > check:17'0 ~~~~~~~~~~~~~~~~~ > > 75: ; IeeeMode: 1 > > check:17'0 ~~~~~~~~~~~~~~ > > 76: ; LDSByteSize: 0 bytes/workgroup (compile time only) > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 77: ; SGPRBlocks: 0 > > check:17'0 ~~~~~~~~~~~~~~~~ > > 78: ; VGPRBlocks: 0 > > check:17'0 ~~~~~~~~~~~~~~~~ > > 79: ; NumSGPRsForWavesPerEU: 8 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 80: ; NumVGPRsForWavesPerEU: 4 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 81: ; Occupancy: 10 > > check:17'0 ~~~~~~~~~~~~~~~~ > > 82: ; WaveLimiterHint : 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 83: ; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 84: ; COMPUTE_PGM_RSRC2:USER_SGPR: 6 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 85: ; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 86: ; COMPUTE_PGM_RSRC2:TGID_X_EN: 1 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 87: ; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 88: ; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 89: ; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 90: .text > > check:17'0 ~~~~~~~ > > 91: .protected test_native_rsqrt_f32 ; -- Begin function test_native_rsqrt_f32 > > check:17'0 > ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > > 92: .globl test_native_rsqrt_f32 > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 93: .p2align 8 > > check:17'0 ~~~~~~~~~~~~ > > 94: .type test_native_rsqrt_f32,@function > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ > > 95: test_native_rsqrt_f32: ; @test_native_rsqrt_f32 > > label:23 ^~~~~~~~~~~~~~~~~~~~~~ > > check:17'0 ~~~~~~~~~~~~~~~~~~~~~~ > > 96: ; %bb.0: > 97: s_load_dwordx4 s[0:3], s[4:5], 0x0 > 98: v_lshlrev_b32_e32 v2, 2, v0 > 99: s_waitcnt lgkmcnt(0) > 100: v_mov_b32_e32 v1, s3 > 101: v_add_u32_e32 v0, vcc, s2, v2 > 102: v_addc_u32_e32 v1, vcc, 0, v1, vcc > 103: flat_load_dword v0, v[0:1] > 104: v_mov_b32_e32 v1, s1 > 105: s_waitcnt vmcnt(0) > 106: v_rsq_f32_e32 v3, v0 > > check:24 ^~~~~~~~~ > > 107: v_add_u32_e32 v0, vcc, s0, v2 > 108: v_addc_u32_e32 v1, vcc, 0, v1, vcc > 109: flat_store_dword v[0:1], v3 > 110: s_endpgm > 111: .section .rodata,"a",@progbits > 112: .p2align 6, 0x0 > 113: .amdhsa_kernel test_native_rsqrt_f32 > 114: .amdhsa_group_segment_fixed_size 0 > 115: .amdhsa_private_segment_fixed_size 0 > 116: .amdhsa_kernarg_size 16 > 117: .amdhsa_user_sgpr_count 6 > 118: .amdhsa_user_sgpr_private_segment_buffer 1 > 119: .amdhsa_user_sgpr_dispatch_ptr 0 > 120: .amdhsa_user_sgpr_queue_ptr 0 > 121: .amdhsa_user_sgpr_kernarg_segment_ptr 1 > 122: .amdhsa_user_sgpr_dispatch_id 0 > 123: .amdhsa_user_sgpr_flat_scratch_init 0 > 124: .amdhsa_user_sgpr_private_segment_size 0 > 125: .amdhsa_uses_dynamic_stack 0 > 126: .amdhsa_system_sgpr_private_segment_wavefront_offset 0 > 127: .amdhsa_system_sgpr_workgroup_id_x 1 > 128: .amdhsa_system_sgpr_workgroup_id_y 0 > 129: .amdhsa_system_sgpr_workgroup_id_z 0 > 130: .amdhsa_system_sgpr_workgroup_info 0 > 131: .amdhsa_system_vgpr_workitem_id 0 > 132: .amdhsa_next_free_vgpr 4 > 133: .amdhsa_next_free_sgpr 6 > 134: .amdhsa_reserve_flat_scratch 0 > 135: .amdhsa_float_round_mode_32 0 > 136: .amdhsa_float_round_mode_16_64 0 > 137: .amdhsa_float_denorm_mode_32 0 > 138: .amdhsa_float_denorm_mode_16_64 3 > 139: .amdhsa_dx10_clamp 1 > 140: .amdhsa_ieee_mode 1 > 141: .amdhsa_exception_fp_ieee_invalid_op 0 > 142: .amdhsa_exception_fp_denorm_src 0 > 143: .amdhsa_exception_fp_ieee_div_zero 0 > 144: .amdhsa_exception_fp_ieee_overflow 0 > 145: .amdhsa_exception_fp_ieee_underflow 0 > 146: .amdhsa_exception_fp_ieee_inexact 0 > 147: .amdhsa_exception_int_div_zero 0 > 148: .end_amdhsa_kernel > 149: .text > 150: .Lfunc_end1: > 151: .size test_native_rsqrt_f32, .Lfunc_end1-test_native_rsqrt_f32 > 152: ; -- End function > 153: .section .AMDGPU.csdata,"",@progbits > 154: ; Kernel info: > 155: ; codeLenInByte = 68 > 156: ; NumSgprs: 8 > 157: ; NumVgprs: 4 > 158: ; ScratchSize: 0 > 159: ; MemoryBound: 0 > 160: ; FloatMode: 192 > 161: ; IeeeMode: 1 > 162: ; LDSByteSize: 0 bytes/workgroup (compile time only) > 163: ; SGPRBlocks: 0 > 164: ; VGPRBlocks: 0 > 165: ; NumSGPRsForWavesPerEU: 8 > 166: ; NumVGPRsForWavesPerEU: 4 > 167: ; Occupancy: 10 > 168: ; WaveLimiterHint : 0 > 169: ; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0 > 170: ; COMPUTE_PGM_RSRC2:USER_SGPR: 6 > 171: ; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0 > 172: ; COMPUTE_PGM_RSRC2:TGID_X_EN: 1 > 173: ; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0 > 174: ; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0 > 175: ; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0 > 176: .hidden __oclc_ABI_version ; @__oclc_ABI_version > 177: .type __oclc_ABI_version,@object > 178: .section .rodata,"a",@progbits > 179: .weak __oclc_ABI_version > 180: .p2align 2, 0x0 > 181: __oclc_ABI_version: > 182: .long 500 ; 0x1f4 > 183: .size __oclc_ABI_version, 4 > 184: > 185: .ident "clang version 18.1.8" > 186: .section ".note.GNU-stack","",@progbits > 187: .addrsig > 188: .amdgpu_metadata > 189: --- > 190: amdhsa.kernels: > 191: - .args: > 192: - .actual_access: write_only > 193: .address_space: global > 194: .is_restrict: true > 195: .offset: 0 > 196: .size: 8 > 197: .type_name: 'half*' > 198: .value_kind: global_buffer > 199: - .actual_access: read_only > 200: .address_space: global > 201: .is_restrict: true > 202: .offset: 8 > 203: .size: 8 > 204: .type_name: 'half*' > 205: .value_kind: global_buffer > 206: .group_segment_fixed_size: 0 > 207: .kernarg_segment_align: 8 > 208: .kernarg_segment_size: 16 > 209: .language: OpenCL C > 210: .language_version: > 211: - 2 > 212: - 0 > 213: .max_flat_workgroup_size: 256 > 214: .name: test_native_rsqrt_f16 > 215: .private_segment_fixed_size: 0 > 216: .sgpr_count: 8 > 217: .sgpr_spill_count: 0 > 218: .symbol: test_native_rsqrt_f16.kd > 219: .uses_dynamic_stack: false > 220: .vgpr_count: 4 > 221: .vgpr_spill_count: 0 > 222: .wavefront_size: 64 > 223: - .args: > 224: - .actual_access: write_only > 225: .address_space: global > 226: .is_restrict: true > 227: .offset: 0 > 228: .size: 8 > 229: .type_name: 'float*' > 230: .value_kind: global_buffer > 231: - .actual_access: read_only > 232: .address_space: global > 233: .is_restrict: true > 234: .offset: 8 > 235: .size: 8 > 236: .type_name: 'float*' > 237: .value_kind: global_buffer > 238: .group_segment_fixed_size: 0 > 239: .kernarg_segment_align: 8 > 240: .kernarg_segment_size: 16 > 241: .language: OpenCL C > 242: .language_version: > 243: - 2 > 244: - 0 > 245: .max_flat_workgroup_size: 256 > 246: .name: test_native_rsqrt_f32 > 247: .private_segment_fixed_size: 0 > 248: .sgpr_count: 8 > 249: .sgpr_spill_count: 0 > 250: .symbol: test_native_rsqrt_f32.kd > 251: .uses_dynamic_stack: false > 252: .vgpr_count: 4 > 253: .vgpr_spill_count: 0 > 254: .wavefront_size: 64 > 255: amdhsa.target: amdgcn-amd-amdhsa--gfx803 > 256: amdhsa.version: > 257: - 1 > 258: - 2 > 259: ... > 260: > 261: .end_amdgpu_metadata > > >>>>>> > > > ><end of output> >Test time = 0.22 sec >---------------------------------------------------------- >Test Failed. >"compile_native_rsqrt__gfx803" end time: Aug 18 19:10 UTC >"compile_native_rsqrt__gfx803" time elapsed: 00:00:00 >---------------------------------------------------------- > >20/21 Testing: compile_native_log__gfx803 >20/21 Test: compile_native_log__gfx803 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.native_log.gfx803.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/native_log.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx803" "-DEXTRA_CHECK_PREFIX=GFX803,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_native_log__gfx803" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.18 sec >---------------------------------------------------------- >Test Passed. >"compile_native_log__gfx803" end time: Aug 18 19:10 UTC >"compile_native_log__gfx803" time elapsed: 00:00:00 >---------------------------------------------------------- > >21/21 Testing: compile_native_exp__gfx803 >21/21 Test: compile_native_exp__gfx803 >Command: "/usr/bin/cmake" "-DCLANG_BIN=/usr/lib/llvm/18/bin/clang-18" "-DBINARY_DIR=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build" "-DFILECHECK_BIN=/usr/lib/llvm/18/bin/FileCheck" "-DOUTPUT_FILE=output.native_exp.gfx803.s" "-DINPUT_FILE=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/native_exp.cl" "-DAMDGCN_BITCODES=/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/lib/amdgcn/bitcode" "-DTEST_CPU=gfx803" "-DEXTRA_CHECK_PREFIX=GFX803,GCN" "-P" "/var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs/test/compile/RunCompileTest.cmake" >Directory: /var/tmp/portage/dev-libs/rocm-device-libs-6.1.2/work/llvm-project-rocm-6.1.2/amd/device-libs_build/test/compile >"compile_native_exp__gfx803" start time: Aug 18 19:10 UTC >Output: >---------------------------------------------------------- ><end of output> >Test time = 0.16 sec >---------------------------------------------------------- >Test Passed. >"compile_native_exp__gfx803" end time: Aug 18 19:10 UTC >"compile_native_exp__gfx803" time elapsed: 00:00:00 >---------------------------------------------------------- > >End testing: Aug 18 19:10 UTC
<b>You cannot view the attachment while viewing its details because your browser does not support IFRAMEs. <a href="attachment.cgi?id=900599">View the attachment on a separate page</a>.</b>
View Attachment As Raw
Actions:
View
Attachments on
bug 938163
:
900598
| 900599 |
900600
|
900601
|
900602
|
900603
|
900604
|
900605
|
900606
|
900607
