From d2ed973482fdd800013658e83a61709b29e0a80f Mon Sep 17 00:00:00 2001 From: V3n3RiX Date: Thu, 27 Jun 2024 07:59:40 +0100 Subject: gentoo auto-resync : 27:06:2024 - 07:59:39 --- .../rocm-device-libs-5.5.0-test-bitcode-dir.patch | 37 +++++++++++++ ...ocm-device-libs-5.5.1-remove-gfx700-tests.patch | 34 ++++++++++++ .../rocm-device-libs-6.0.0-add-gws-attribute.patch | 22 ++++++++ .../rocm-device-libs-6.1.0-fix-llvm-link.patch | 28 ++++++++++ .../rocm-device-libs-6.1.0-fix-test-failures.patch | 64 ++++++++++++++++++++++ ...rocm-device-libs-6.1.0-fix-test-failures2.patch | 44 +++++++++++++++ 6 files changed, 229 insertions(+) create mode 100644 dev-libs/rocm-device-libs/files/rocm-device-libs-5.5.0-test-bitcode-dir.patch create mode 100644 dev-libs/rocm-device-libs/files/rocm-device-libs-5.5.1-remove-gfx700-tests.patch create mode 100644 dev-libs/rocm-device-libs/files/rocm-device-libs-6.0.0-add-gws-attribute.patch create mode 100644 dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-llvm-link.patch create mode 100644 dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures.patch create mode 100644 dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures2.patch (limited to 'dev-libs/rocm-device-libs/files') diff --git a/dev-libs/rocm-device-libs/files/rocm-device-libs-5.5.0-test-bitcode-dir.patch b/dev-libs/rocm-device-libs/files/rocm-device-libs-5.5.0-test-bitcode-dir.patch new file mode 100644 index 000000000000..2a307786edd5 --- /dev/null +++ b/dev-libs/rocm-device-libs/files/rocm-device-libs-5.5.0-test-bitcode-dir.patch @@ -0,0 +1,37 @@ +`--rocm-device-lib-path` is needed to execute the tests because they are not +installed to system yet +=================================================================== +--- ROCm-Device-Libs-rocm-6.0.0.orig/test/compile/CMakeLists.txt ++++ ROCm-Device-Libs-rocm-6.0.0/test/compile/CMakeLists.txt +@@ -30,6 +30,7 @@ function(add_compile_test test_name func + -DFILECHECK_BIN=${FILECHECK_BIN} + -DOUTPUT_FILE=output.${name}.${test_cpu}.s + -DINPUT_FILE=${CMAKE_CURRENT_SOURCE_DIR}/${func_name}.cl ++ -DAMDGCN_BITCODES=${PROJECT_BINARY_DIR}/lib/amdgcn/bitcode + -DTEST_CPU=${test_cpu} + -DEXTRA_CHECK_PREFIX=${extra_check_prefixes} + -P ${script}) +Index: ROCm-Device-Libs-rocm-6.0.0/test/compile/RunConstantFoldTest.cmake +=================================================================== +--- ROCm-Device-Libs-rocm-6.0.0.orig/test/compile/RunConstantFoldTest.cmake ++++ ROCm-Device-Libs-rocm-6.0.0/test/compile/RunConstantFoldTest.cmake +@@ -16,6 +16,7 @@ execute_process(COMMAND + -target amdgcn-amd-amdhsa -mcpu=${TEST_CPU} + -Xclang -finclude-default-header + --rocm-path=${BINARY_DIR} ++ --rocm-device-lib-path=${AMDGCN_BITCODES} + -mllvm -amdgpu-simplify-libcall=0 + -o ${OUTPUT_FILE} ${INPUT_FILE} + RESULT_VARIABLE CLANG_RESULT +Index: ROCm-Device-Libs-rocm-6.0.0/test/compile/RunCompileTest.cmake +=================================================================== +--- ROCm-Device-Libs-rocm-6.0.0.orig/test/compile/RunCompileTest.cmake ++++ ROCm-Device-Libs-rocm-6.0.0/test/compile/RunCompileTest.cmake +@@ -16,6 +16,7 @@ execute_process(COMMAND + -target amdgcn-amd-amdhsa -mcpu=${TEST_CPU} + -Xclang -finclude-default-header + --rocm-path=${BINARY_DIR} ++ --rocm-device-lib-path=${AMDGCN_BITCODES} + -mllvm -amdgpu-simplify-libcall=0 + -o ${OUTPUT_FILE} ${INPUT_FILE} + RESULT_VARIABLE CLANG_RESULT diff --git a/dev-libs/rocm-device-libs/files/rocm-device-libs-5.5.1-remove-gfx700-tests.patch b/dev-libs/rocm-device-libs/files/rocm-device-libs-5.5.1-remove-gfx700-tests.patch new file mode 100644 index 000000000000..0fb5e44e218c --- /dev/null +++ b/dev-libs/rocm-device-libs/files/rocm-device-libs-5.5.1-remove-gfx700-tests.patch @@ -0,0 +1,34 @@ +https://github.com/ROCm/ROCm-Device-Libs/issues/86 +https://github.com/ROCm/ROCm-Device-Libs/pull/89 +From 936a78464491c95c7cfffde08491bfe1a48c7177 Mon Sep 17 00:00:00 2001 +From: Brian Sumner +Date: Tue, 9 May 2023 07:55:03 -0700 +Subject: [PATCH] Remove gfx700 from tests + +Change-Id: I23d6639cb7c04246cc5be86d6e172e32546b3b90 +--- + test/compile/CMakeLists.txt | 4 ++-- + 1 file changed, 2 insertions(+), 2 deletions(-) + +diff --git a/test/compile/CMakeLists.txt b/test/compile/CMakeLists.txt +index 9af0b1a..a789222 100644 +--- a/test/compile/CMakeLists.txt ++++ b/test/compile/CMakeLists.txt +@@ -54,12 +54,12 @@ foreach(gpu gfx900 gfx1030) + add_constant_fold_test(lgamma_r ${gpu}) + endforeach() + +-foreach(gpu gfx700 gfx803) ++foreach(gpu gfx803) + add_isa_test(asin ${gpu}) + add_isa_test(atan2 ${gpu}) + add_isa_test(atan2pi ${gpu}) + endforeach() + +-foreach(gpu gfx600 gfx700) ++foreach(gpu gfx600) + add_isa_test(frexp ${gpu}) + endforeach() +-- +2.41.0 + diff --git a/dev-libs/rocm-device-libs/files/rocm-device-libs-6.0.0-add-gws-attribute.patch b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.0.0-add-gws-attribute.patch new file mode 100644 index 000000000000..1aaecbb0c1e6 --- /dev/null +++ b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.0.0-add-gws-attribute.patch @@ -0,0 +1,22 @@ +Fix compatilibity with Clang-18 +https://github.com/ROCm/ROCm-Device-Libs/issues/96 + +Backports https://github.com/ROCm/llvm-project/commit/688c78d85caf499957db175811f8b00c7c818f83 +--- a/ockl/src/cg.cl ++++ b/ockl/src/cg.cl +@@ -84,13 +84,13 @@ multi_grid_sync(__global struct mg_sync *s, uint members) + } + } + +-void ++__attribute__((target("gws"))) void + __ockl_gws_init(uint nwm1, uint rid) + { + __builtin_amdgcn_ds_gws_init(nwm1, rid); + } + +-void ++__attribute__((target("gws"))) void + __ockl_gws_barrier(uint nwm1, uint rid) + { + __builtin_amdgcn_ds_gws_barrier(nwm1, rid); \ No newline at end of file diff --git a/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-llvm-link.patch b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-llvm-link.patch new file mode 100644 index 000000000000..7a08dc4a1d0d --- /dev/null +++ b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-llvm-link.patch @@ -0,0 +1,28 @@ +https://github.com/ROCm/llvm-project/pull/68 +From 1c7e7f872980a5b15fb3d85f8780e78ce3b715b1 Mon Sep 17 00:00:00 2001 +From: Selene +Date: Wed, 1 May 2024 17:38:33 +0800 +Subject: [PATCH] Allow link to llvm shared library for current distros + +--- + amd/device-libs/utils/prepare-builtins/CMakeLists.txt | 6 +++++- + 1 file changed, 5 insertions(+), 1 deletion(-) + +diff --git a/amd/device-libs/utils/prepare-builtins/CMakeLists.txt b/amd/device-libs/utils/prepare-builtins/CMakeLists.txt +index 63661962860a..079dc08e3419 100644 +--- a/utils/prepare-builtins/CMakeLists.txt ++++ b/utils/prepare-builtins/CMakeLists.txt +@@ -26,5 +26,9 @@ set_target_properties(prepare-builtins PROPERTIES + CXX_STANDARD_REQUIRED Yes + CXX_EXTENSIONS No) + llvm_update_compile_flags(prepare-builtins) +-llvm_map_components_to_libnames(llvm_libs support core bitreader bitwriter) ++if (LLVM_LINK_LLVM_DYLIB) ++ set(llvm_libs LLVM) ++else() ++ llvm_map_components_to_libnames(llvm_libs support core bitreader bitwriter) ++endif() + target_link_libraries(prepare-builtins ${llvm_libs}) +-- +2.44.0 + diff --git a/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures.patch b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures.patch new file mode 100644 index 000000000000..856b5cf7e392 --- /dev/null +++ b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures.patch @@ -0,0 +1,64 @@ +Modified from https://github.com/ROCm/llvm-project/commit/7c2188cbc193f2b4dd5394f17404b44340001f30.patch +From 7c2188cbc193f2b4dd5394f17404b44340001f30 Mon Sep 17 00:00:00 2001 +From: Matt Arsenault +Date: Thu, 4 Jan 2024 19:34:47 +0700 +Subject: [PATCH] device-libs: Fix input file path test failures + +The test file input paths broke with the move to the mono-repo. Some of +the constant folding values are host dependent, so update the values to +what works for me now. Not really sure what else I can do about these. + +Change-Id: Ic764f637bb9532fcede9bfb9ce3886a2b7d467d8 +--- + amd/device-libs/test/compile/frexp.cl | 10 ++++++---- + amd/device-libs/test/compile/lgamma_r.cl | 2 +- + 3 files changed, 8 insertions(+), 9 deletions(-) + +diff --git a/test/compile/frexp.cl b/test/compile/frexp.cl +index 780c541500021b..b3181ce87db711 100644 +--- a/test/compile/frexp.cl ++++ b/test/compile/frexp.cl +@@ -5,11 +5,11 @@ + // later. + + // GCN-LABEL: {{^}}test_frexp_f32: +-// GFX600-DAG: s_movk_i32 [[CLASS_MASK:s[0-9]+]], 0x1f8 ++// GFX600-DAG: s_mov_b32 [[INF:s[0-9]+]], 0x7f80000 + // GFX600-DAG: v_frexp_mant_f32{{(_e32)?}} [[MANT:v[0-9]+]], [[SRC:v[0-9]+]] + // GFX600-DAG: v_frexp_exp_i32_f32{{(_e32)?}} [[EXP:v[0-9]+]], [[SRC:v[0-9]+]] + +-// GFX600-DAG: v_cmp_class_f32{{(_e64)?}} [[CMP:(vcc|s{{\[[0-9]+:[0-9]+\]}})]], [[SRC]], [[CLASS_MASK]] ++// GFX600-DAG: v_cmp_lt_f32{{(_e64)?}} [[CMP:(vcc|s{{\[[0-9]+:[0-9]+\]}})]], |[[SRC]]|, [[INF]] + + // GFX600-DAG: v_cndmask_b32{{(_e32)?|(e64)?}} v{{[0-9]+}}, [[SRC]], [[MANT]], [[CMP]] + // GFX600-DAG: v_cndmask_b32{{(_e32)?|(e64)?}} v{{[0-9]+}}, 0, [[EXP]], [[CMP]] +@@ -30,12 +30,14 @@ kernel void test_frexp_f32(global float* restrict out0, + } + + // GCN-LABEL: {{^}}test_frexp_f64: ++// GFX600: s_mov_b32 s{{[0-9]+}}, 0{{$}} + +-// GFX600-DAG: s_movk_i32 [[CLASS_MASK:s[0-9]+]], 0x1f8 ++// GFX600-DAG: s_mov_b32 s[[INF_LO:[0-9]+]], 0{{$}} ++// GFX600-DAG: s_mov_b32 s[[INF_HI:[0-9]+]], 0x7ff00000{{$}} + // GFX600-DAG: v_frexp_mant_f64{{(_e32)?}} v{{\[}}[[MANT_LO:[0-9]+]]:[[MANT_HI:[0-9]+]]{{\]}}, [[SRC:v\[[0-9]+:[0-9]+\]]] + // GFX600-DAG: v_frexp_exp_i32_f64{{(_e32)?}} [[EXP:v[0-9]+]], [[SRC:v\[[0-9]+:[0-9]+\]]] + +-// GFX600-DAG: v_cmp_class_f64{{(_e64)?}} [[CMP:(vcc|s{{\[[0-9]+:[0-9]+\]}})]], [[SRC]], [[CLASS_MASK]] ++// GFX600-DAG: v_cmp_lt_f64{{(_e64)?}} [[CMP:(vcc|s{{\[[0-9]+:[0-9]+\]}})]], |[[SRC]]|, s{{\[}}[[INF_LO]]:[[INF_HI]]{{\]}} + + // GFX600-DAG: v_cndmask_b32{{(_e32)?|(e64)?}} v{{[0-9]+}}, v{{[0-9]+}}, v[[MANT_LO]], [[CMP]] + // GFX600-DAG: v_cndmask_b32{{(_e32)?|(e64)?}} v{{[0-9]+}}, v{{[0-9]+}}, v[[MANT_HI]], [[CMP]] +diff --git a/test/compile/lgamma_r.cl b/test/compile/lgamma_r.cl +index 1e1984226cd55c..56d1ba15f761f2 100644 +--- a/test/compile/lgamma_r.cl ++++ b/test/compile/lgamma_r.cl +@@ -66,7 +66,7 @@ kernel void constant_fold_lgamma_r_f32(volatile global float* out, + out[0] = test_lgamma_r(0x1.0p-127f, sign_out); + + // CONSTANTFOLD-NEXT: store volatile i32 1, +- // CONSTANTFOLD-NEXT: store volatile float 0x419DE28040000000, ++ // CONSTANTFOLD-NEXT: store volatile float 0x419DE28060000000, + out[0] = test_lgamma_r(nextafter(0x1.0p+23f, __builtin_inff()), sign_out); + + // CONSTANTFOLD-NEXT: store volatile i32 1, diff --git a/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures2.patch b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures2.patch new file mode 100644 index 000000000000..c0559e9468f6 --- /dev/null +++ b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures2.patch @@ -0,0 +1,44 @@ +Modified from https://github.com/ROCm/llvm-project/commit/794ebeffcafbf6f4d86cb1bfd7a5a0d1d30f1fc7 +From 794ebeffcafbf6f4d86cb1bfd7a5a0d1d30f1fc7 Mon Sep 17 00:00:00 2001 +From: Matt Arsenault +Date: Thu, 4 Jan 2024 20:40:30 +0700 +Subject: [PATCH] device-libs: Fix native_rsqrt test failures for f16-as-f32 + case + +At some point the code quality regressed for these. + +Change-Id: Ib5c4c1d093a6056dd6213ed6eaf157ad806101fb +--- + amd/device-libs/test/compile/native_rsqrt.cl | 13 +++++++++---- + 1 file changed, 9 insertions(+), 4 deletions(-) + +diff --git a/test/compile/native_rsqrt.cl b/test/compile/native_rsqrt.cl +index 94ef375e4d522f..291cc6f32c3351 100644 +--- a/test/compile/native_rsqrt.cl ++++ b/test/compile/native_rsqrt.cl +@@ -6,16 +6,21 @@ + + half __ocml_native_rsqrt_f16(half); + +-// GCN: {{^}}test_native_rsqrt_f16: +-// GFX600: v_rsq_f32 +-// GFX700: v_rsq_f32 ++// FIXME: Promoted case using full expansion ++// GCN-LABEL: {{^}}test_native_rsqrt_f16: ++// GFX600: v_sqrt_f32 ++// GFX600: v_rcp_f32 ++ ++// GFX700: v_sqrt_f32 ++// GFX700: v_rcp_f32 ++ + // GFX803: v_rsq_f16 + kernel void test_native_rsqrt_f16(global half* restrict out, global half* restrict in) { + int id = get_local_id(0); + out[id] = __ocml_native_rsqrt_f16(in[id]); + } + +-// GCN: {{^}}test_native_rsqrt_f32: ++// GCN-LABEL: {{^}}test_native_rsqrt_f32: + // GCN: v_rsq_f32 + kernel void test_native_rsqrt_f32(global float* restrict out, global float* restrict in) { + int id = get_local_id(0); -- cgit v1.2.3