summaryrefslogtreecommitdiff
path: root/dev-libs/rocm-device-libs/files
diff options
context:
space:
mode:
authorV3n3RiX <venerix@koprulu.sector>2024-06-27 07:59:40 +0100
committerV3n3RiX <venerix@koprulu.sector>2024-06-27 07:59:40 +0100
commitd2ed973482fdd800013658e83a61709b29e0a80f (patch)
tree57ea7666a57b5a05a4c8866e4915e90b4a6e7c94 /dev-libs/rocm-device-libs/files
parent9f6a82a85d400d6ae7de04c43cee88dbc6bc4da0 (diff)
gentoo auto-resync : 27:06:2024 - 07:59:39
Diffstat (limited to 'dev-libs/rocm-device-libs/files')
-rw-r--r--dev-libs/rocm-device-libs/files/rocm-device-libs-5.5.0-test-bitcode-dir.patch37
-rw-r--r--dev-libs/rocm-device-libs/files/rocm-device-libs-5.5.1-remove-gfx700-tests.patch34
-rw-r--r--dev-libs/rocm-device-libs/files/rocm-device-libs-6.0.0-add-gws-attribute.patch22
-rw-r--r--dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-llvm-link.patch28
-rw-r--r--dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures.patch64
-rw-r--r--dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures2.patch44
6 files changed, 229 insertions, 0 deletions
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 <brian.sumner@amd.com>
+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 <lixueying@mail.bnu.edu.cn>
+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 <Matthew.Arsenault@amd.com>
+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 <Matthew.Arsenault@amd.com>
+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);