[libclc] [libclc] Move external-funcs.test to static file and use IR checks for .cl tests (PR #197151)
Wenju He via cfe-commits
cfe-commits at lists.llvm.org
Wed May 13 03:00:35 PDT 2026
https://github.com/wenju-he updated https://github.com/llvm/llvm-project/pull/197151
>From 2d10da58e2ecade7d99ec2900216161b7ed83636 Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Tue, 12 May 2026 08:44:55 +0200
Subject: [PATCH 1/7] [libclc] Move external-funcs.test to static file and use
IR checks for .cl tests
This PR supercedes #87989.
Moving external-funcs.test to static file simplifies test/CMakeLists.txt.
Static files follows llvm standard lit pattern and enables fine-grained
testing of missing symbols in specific libraries.
.cl test files uses %libclc_target and %check_prefix, which are replaced
with specific values during `ninja check-libclc` or
`llvm-lit build/runtimes/runtimes-${triple}-llvm-bins/libclc/test`.
This allows testing multiple triples in the same test file.
Add script libclc/test/update_libclc_tests.py, which wraps
utils/update_cc_test_checks.py to update CHECK lines in libclc .cl tests
for a given arch. Example usage:
`libclc/test/update_libclc_tests.py amdgpu`
Assisted-by: Claude Sonnet 4.6
---
libclc/test/AMDGPU/external-funcs.test | 1 +
libclc/test/AMDGPU/lit.local.cfg | 2 +
libclc/test/CMakeLists.txt | 42 ++-----
libclc/test/add_sat.cl | 11 --
libclc/test/as_type.cl | 11 --
libclc/test/conversion/convert.cl | 14 +++
libclc/test/convert.cl | 11 --
libclc/test/cos.cl | 11 --
libclc/test/cross.cl | 11 --
libclc/test/fabs.cl | 11 --
libclc/test/geometric/cross.cl | 31 +++++
libclc/test/get_group_id.cl | 11 --
libclc/test/integer/add_sat.cl | 14 +++
libclc/test/integer/sub_sat.cl | 44 +++++++
libclc/test/lit.cfg.py | 25 +++-
libclc/test/lit.site.cfg.py.in | 8 +-
libclc/test/math/cos.cl | 156 +++++++++++++++++++++++++
libclc/test/math/fabs.cl | 14 +++
libclc/test/math/rsqrt.cl | 45 +++++++
libclc/test/misc/as_type.cl | 14 +++
libclc/test/rsqrt.cl | 14 ---
libclc/test/subsat.cl | 27 -----
libclc/test/update_libclc_tests.py | 146 +++++++++++++++++++++++
libclc/test/work-item/get_group_id.cl | 15 +++
24 files changed, 534 insertions(+), 155 deletions(-)
create mode 100644 libclc/test/AMDGPU/external-funcs.test
create mode 100644 libclc/test/AMDGPU/lit.local.cfg
delete mode 100644 libclc/test/add_sat.cl
delete mode 100644 libclc/test/as_type.cl
create mode 100644 libclc/test/conversion/convert.cl
delete mode 100644 libclc/test/convert.cl
delete mode 100644 libclc/test/cos.cl
delete mode 100644 libclc/test/cross.cl
delete mode 100644 libclc/test/fabs.cl
create mode 100644 libclc/test/geometric/cross.cl
delete mode 100644 libclc/test/get_group_id.cl
create mode 100644 libclc/test/integer/add_sat.cl
create mode 100644 libclc/test/integer/sub_sat.cl
create mode 100644 libclc/test/math/cos.cl
create mode 100644 libclc/test/math/fabs.cl
create mode 100644 libclc/test/math/rsqrt.cl
create mode 100644 libclc/test/misc/as_type.cl
delete mode 100644 libclc/test/rsqrt.cl
delete mode 100644 libclc/test/subsat.cl
create mode 100755 libclc/test/update_libclc_tests.py
create mode 100644 libclc/test/work-item/get_group_id.cl
diff --git a/libclc/test/AMDGPU/external-funcs.test b/libclc/test/AMDGPU/external-funcs.test
new file mode 100644
index 0000000000000..73bddc02d11e1
--- /dev/null
+++ b/libclc/test/AMDGPU/external-funcs.test
@@ -0,0 +1 @@
+; RUN: llvm-nm -u "%libclc_library_dir/%libclc_target/libclc.bc" | FileCheck %s --allow-empty --implicit-check-not=" U "
diff --git a/libclc/test/AMDGPU/lit.local.cfg b/libclc/test/AMDGPU/lit.local.cfg
new file mode 100644
index 0000000000000..1ee232d629c16
--- /dev/null
+++ b/libclc/test/AMDGPU/lit.local.cfg
@@ -0,0 +1,2 @@
+if "AMDGCN" not in config.available_features:
+ config.unsupported = True
diff --git a/libclc/test/CMakeLists.txt b/libclc/test/CMakeLists.txt
index 4c54d93b26813..955235f32277d 100644
--- a/libclc/test/CMakeLists.txt
+++ b/libclc/test/CMakeLists.txt
@@ -1,36 +1,14 @@
set(LLVM_TOOLS_DIR ${LLVM_TOOLS_BINARY_DIR})
-set(LIBCLC_TEST_DEPS
- libclc-opencl-builtins
+configure_lit_site_cfg(
+ ${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.py.in
+ ${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg.py
+ MAIN_CONFIG
+ ${CMAKE_CURRENT_SOURCE_DIR}/lit.cfg.py
)
-umbrella_lit_testsuite_begin(check-libclc)
-
-# Testing unresolved symbols.
-# Skip nvptx and spirv targets.
-if(ARCH MATCHES amdgcn)
- foreach(tgt IN LISTS LIBCLC_UNRESOLVED_SYMBOL_TEST_TARGETS)
- set(target_file "$<TARGET_PROPERTY:${tgt},TARGET_FILE>")
-
- set(LIBCLC_TARGET_TEST_DIR ${CMAKE_CURRENT_BINARY_DIR}/${tgt})
- file(MAKE_DIRECTORY ${LIBCLC_TARGET_TEST_DIR})
- file(GENERATE OUTPUT ${LIBCLC_TARGET_TEST_DIR}/check-external-funcs.test
- CONTENT "; RUN: llvm-nm -u \"${target_file}\" | FileCheck %s --allow-empty\n\n; CHECK-NOT: {{.+}}\n"
- )
-
- configure_lit_site_cfg(
- ${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.py.in
- ${LIBCLC_TARGET_TEST_DIR}/lit.site.cfg.py
- MAIN_CONFIG
- ${CMAKE_CURRENT_SOURCE_DIR}/lit.cfg.py
- )
-
- add_lit_testsuite(check-libclc-external-funcs-${tgt} "Running ${tgt} unresolved symbols tests"
- ${LIBCLC_TARGET_TEST_DIR}
- DEPENDS ${LIBCLC_TEST_DEPS}
- )
- set_target_properties(check-libclc-external-funcs-${tgt} PROPERTIES FOLDER "libclc tests")
- endforeach()
-endif()
-
-umbrella_lit_testsuite_end(check-libclc)
+add_lit_testsuite(check-libclc "Running libclc tests"
+ ${CMAKE_CURRENT_BINARY_DIR}
+ DEPENDS libclc
+)
+set_target_properties(check-libclc PROPERTIES FOLDER "libclc tests")
diff --git a/libclc/test/add_sat.cl b/libclc/test/add_sat.cl
deleted file mode 100644
index 87c3d39df3542..0000000000000
--- a/libclc/test/add_sat.cl
+++ /dev/null
@@ -1,11 +0,0 @@
-//===----------------------------------------------------------------------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-__kernel void foo(__global char *a, __global char *b, __global char *c) {
- *a = add_sat(*b, *c);
-}
diff --git a/libclc/test/as_type.cl b/libclc/test/as_type.cl
deleted file mode 100644
index a926f48c4ea0c..0000000000000
--- a/libclc/test/as_type.cl
+++ /dev/null
@@ -1,11 +0,0 @@
-//===----------------------------------------------------------------------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-__kernel void foo(int4 *x, float4 *y) {
- *x = as_int4(*y);
-}
diff --git a/libclc/test/conversion/convert.cl b/libclc/test/conversion/convert.cl
new file mode 100644
index 0000000000000..00b524bb8f11d
--- /dev/null
+++ b/libclc/test/conversion/convert.cl
@@ -0,0 +1,14 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
+// REQUIRES: AMDGCN
+
+// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+
+// AMDGCN-LABEL: define hidden <4 x i32> @float4_to_int4(
+// AMDGCN-SAME: <4 x float> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
+// AMDGCN-NEXT: [[ENTRY:.*:]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = fptosi <4 x float> [[X]] to <4 x i32>
+// AMDGCN-NEXT: ret <4 x i32> [[TMP0]]
+//
+int4 float4_to_int4(float4 x) {
+ return convert_int4(x);
+}
diff --git a/libclc/test/convert.cl b/libclc/test/convert.cl
deleted file mode 100644
index 8eba608dc5f8c..0000000000000
--- a/libclc/test/convert.cl
+++ /dev/null
@@ -1,11 +0,0 @@
-//===----------------------------------------------------------------------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-__kernel void foo(int4 *x, float4 *y) {
- *x = convert_int4(*y);
-}
diff --git a/libclc/test/cos.cl b/libclc/test/cos.cl
deleted file mode 100644
index 92a998b3ba5f7..0000000000000
--- a/libclc/test/cos.cl
+++ /dev/null
@@ -1,11 +0,0 @@
-//===----------------------------------------------------------------------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-__kernel void foo(float4 *f) {
- *f = cos(*f);
-}
diff --git a/libclc/test/cross.cl b/libclc/test/cross.cl
deleted file mode 100644
index 90762d0d073a6..0000000000000
--- a/libclc/test/cross.cl
+++ /dev/null
@@ -1,11 +0,0 @@
-//===----------------------------------------------------------------------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-__kernel void foo(float4 *f) {
- *f = cross(f[0], f[1]);
-}
diff --git a/libclc/test/fabs.cl b/libclc/test/fabs.cl
deleted file mode 100644
index 3f5a964e0418a..0000000000000
--- a/libclc/test/fabs.cl
+++ /dev/null
@@ -1,11 +0,0 @@
-//===----------------------------------------------------------------------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-__kernel void foo(float *f) {
- *f = fabs(*f);
-}
diff --git a/libclc/test/geometric/cross.cl b/libclc/test/geometric/cross.cl
new file mode 100644
index 0000000000000..5aa7541b4d49c
--- /dev/null
+++ b/libclc/test/geometric/cross.cl
@@ -0,0 +1,31 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
+// REQUIRES: AMDGCN
+
+// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+
+// AMDGCN-LABEL: define hidden <4 x float> @test_float4(
+// AMDGCN-SAME: <4 x float> noundef [[X:%.*]], <4 x float> noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
+// AMDGCN-NEXT: [[ENTRY:.*:]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = extractelement <4 x float> [[X]], i64 1
+// AMDGCN-NEXT: [[TMP1:%.*]] = extractelement <4 x float> [[Y]], i64 2
+// AMDGCN-NEXT: [[TMP2:%.*]] = fmul contract float [[TMP0]], [[TMP1]]
+// AMDGCN-NEXT: [[TMP3:%.*]] = extractelement <4 x float> [[X]], i64 2
+// AMDGCN-NEXT: [[TMP4:%.*]] = extractelement <4 x float> [[Y]], i64 1
+// AMDGCN-NEXT: [[TMP5:%.*]] = fmul contract float [[TMP3]], [[TMP4]]
+// AMDGCN-NEXT: [[TMP6:%.*]] = fsub contract float [[TMP2]], [[TMP5]]
+// AMDGCN-NEXT: [[TMP7:%.*]] = extractelement <4 x float> [[Y]], i64 0
+// AMDGCN-NEXT: [[TMP8:%.*]] = fmul contract float [[TMP3]], [[TMP7]]
+// AMDGCN-NEXT: [[TMP9:%.*]] = extractelement <4 x float> [[X]], i64 0
+// AMDGCN-NEXT: [[TMP10:%.*]] = fmul contract float [[TMP9]], [[TMP1]]
+// AMDGCN-NEXT: [[TMP11:%.*]] = fsub contract float [[TMP8]], [[TMP10]]
+// AMDGCN-NEXT: [[TMP12:%.*]] = fmul contract float [[TMP9]], [[TMP4]]
+// AMDGCN-NEXT: [[TMP13:%.*]] = fmul contract float [[TMP0]], [[TMP7]]
+// AMDGCN-NEXT: [[TMP14:%.*]] = fsub contract float [[TMP12]], [[TMP13]]
+// AMDGCN-NEXT: [[TMP15:%.*]] = insertelement <4 x float> <float poison, float poison, float poison, float 0.000000e+00>, float [[TMP6]], i64 0
+// AMDGCN-NEXT: [[TMP16:%.*]] = insertelement <4 x float> [[TMP15]], float [[TMP11]], i64 1
+// AMDGCN-NEXT: [[TMP17:%.*]] = insertelement <4 x float> [[TMP16]], float [[TMP14]], i64 2
+// AMDGCN-NEXT: ret <4 x float> [[TMP17]]
+//
+float4 test_float4(float4 x, float4 y) {
+ return cross(x, y);
+}
diff --git a/libclc/test/get_group_id.cl b/libclc/test/get_group_id.cl
deleted file mode 100644
index c2349a0076889..0000000000000
--- a/libclc/test/get_group_id.cl
+++ /dev/null
@@ -1,11 +0,0 @@
-//===----------------------------------------------------------------------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-__kernel void foo(int *i) {
- i[get_group_id(0)] = 1;
-}
diff --git a/libclc/test/integer/add_sat.cl b/libclc/test/integer/add_sat.cl
new file mode 100644
index 0000000000000..4c30090953e99
--- /dev/null
+++ b/libclc/test/integer/add_sat.cl
@@ -0,0 +1,14 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
+// REQUIRES: AMDGCN
+
+// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+
+// AMDGCN-LABEL: define hidden noundef signext i8 @test_char(
+// AMDGCN-SAME: i8 noundef signext [[X:%.*]], i8 noundef signext [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
+// AMDGCN-NEXT: [[ENTRY:.*:]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call noundef i8 @llvm.sadd.sat.i8(i8 [[X]], i8 [[Y]])
+// AMDGCN-NEXT: ret i8 [[TMP0]]
+//
+char test_char(char x, char y) {
+ return add_sat(x, y);
+}
diff --git a/libclc/test/integer/sub_sat.cl b/libclc/test/integer/sub_sat.cl
new file mode 100644
index 0000000000000..6573afae45102
--- /dev/null
+++ b/libclc/test/integer/sub_sat.cl
@@ -0,0 +1,44 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
+// REQUIRES: AMDGCN
+
+// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+
+// AMDGCN-LABEL: define hidden noundef signext i8 @test_char(
+// AMDGCN-SAME: i8 noundef signext [[X:%.*]], i8 noundef signext [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
+// AMDGCN-NEXT: [[ENTRY:.*:]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call noundef i8 @llvm.ssub.sat.i8(i8 [[X]], i8 [[Y]])
+// AMDGCN-NEXT: ret i8 [[TMP0]]
+//
+char test_char(char x, char y) {
+ return sub_sat(x, y);
+}
+
+// AMDGCN-LABEL: define hidden noundef zeroext i8 @test_uchar(
+// AMDGCN-SAME: i8 noundef zeroext [[X:%.*]], i8 noundef zeroext [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// AMDGCN-NEXT: [[ENTRY:.*:]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call noundef i8 @llvm.usub.sat.i8(i8 [[X]], i8 [[Y]])
+// AMDGCN-NEXT: ret i8 [[TMP0]]
+//
+uchar test_uchar(uchar x, uchar y) {
+ return sub_sat(x, y);
+}
+
+// AMDGCN-LABEL: define hidden noundef i64 @test_long(
+// AMDGCN-SAME: i64 noundef [[X:%.*]], i64 noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// AMDGCN-NEXT: [[ENTRY:.*:]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call noundef i64 @llvm.ssub.sat.i64(i64 [[X]], i64 [[Y]])
+// AMDGCN-NEXT: ret i64 [[TMP0]]
+//
+long test_long(long x, long y) {
+ return sub_sat(x, y);
+}
+
+// AMDGCN-LABEL: define hidden noundef i64 @test_ulong(
+// AMDGCN-SAME: i64 noundef [[X:%.*]], i64 noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// AMDGCN-NEXT: [[ENTRY:.*:]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call noundef i64 @llvm.usub.sat.i64(i64 [[X]], i64 [[Y]])
+// AMDGCN-NEXT: ret i64 [[TMP0]]
+//
+ulong test_ulong(ulong x, ulong y) {
+ return sub_sat(x, y);
+}
diff --git a/libclc/test/lit.cfg.py b/libclc/test/lit.cfg.py
index de1865c2ab3b6..b447d20a30c77 100644
--- a/libclc/test/lit.cfg.py
+++ b/libclc/test/lit.cfg.py
@@ -6,27 +6,46 @@
import lit.formats
+from lit.llvm import llvm_config
+
# Configuration file for the 'lit' test runner.
# name: The name of this test suite.
config.name = "libclc"
# testFormat: The test format to use to interpret tests.
-config.test_format = lit.formats.ShTest(True)
+config.test_format = lit.formats.ShTest()
# suffixes: A list of file extensions to treat as test files.
-config.suffixes = [".test"]
+config.suffixes = [".cl", ".test"]
# Exclude certain directories from test discovery
config.excludes = ["CMakeLists.txt"]
# test_source_root: The root path where tests are located.
# For per-target tests, this is the target's test directory.
-config.test_source_root = config.libclc_obj_root
+config.test_source_root = os.path.dirname(__file__)
# test_exec_root: The root path where tests should be run.
config.test_exec_root = config.libclc_obj_root
+config.target_triple = config.libclc_target
+
+arch = config.libclc_arch.upper()
+config.available_features.add(arch)
+
+llvm_config.use_default_substitutions()
+
+llvm_config.use_clang()
+
+llvm_config.add_tool_substitutions(["llvm-nm"], config.llvm_tools_dir)
+
+config.substitutions.extend([
+ ("%libclc_library_dir", config.libclc_library_dir),
+ ("%libclc_target", config.libclc_target),
+ ("%check_prefix", arch)
+])
+
# Propagate PATH from environment
if "PATH" in os.environ:
config.environment["PATH"] = os.path.pathsep.join(
diff --git a/libclc/test/lit.site.cfg.py.in b/libclc/test/lit.site.cfg.py.in
index ca6ed0144ea9c..d24227d850e6f 100644
--- a/libclc/test/lit.site.cfg.py.in
+++ b/libclc/test/lit.site.cfg.py.in
@@ -2,8 +2,12 @@
import sys
-config.llvm_tools_dir = "@LLVM_TOOLS_DIR@"
-config.libclc_obj_root = "@LIBCLC_TARGET_TEST_DIR@"
+config.host_triple = "@LLVM_HOST_TRIPLE@"
+config.llvm_tools_dir = lit_config.substitute(path(r"@LLVM_TOOLS_DIR@"))
+config.libclc_arch = "@ARCH@"
+config.libclc_library_dir = lit_config.substitute(path(r"@LIBCLC_OUTPUT_LIBRARY_DIR@"))
+config.libclc_obj_root = lit_config.substitute(path(r"@CMAKE_CURRENT_BINARY_DIR@"))
+config.libclc_target = "@LIBCLC_TARGET@"
import lit.llvm
lit.llvm.initialize(lit_config, config)
diff --git a/libclc/test/math/cos.cl b/libclc/test/math/cos.cl
new file mode 100644
index 0000000000000..311024d7155e0
--- /dev/null
+++ b/libclc/test/math/cos.cl
@@ -0,0 +1,156 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
+// REQUIRES: AMDGCN
+
+// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+
+// AMDGCN-LABEL: define hidden float @test_float(
+// AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
+// AMDGCN-NEXT: [[ENTRY:.*:]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call float @llvm.fabs.f32(float [[X]])
+// AMDGCN-NEXT: [[TMP1:%.*]] = fcmp une float [[TMP0]], +inf
+// AMDGCN-NEXT: [[TMP2:%.*]] = select contract i1 [[TMP1]], float [[TMP0]], float +qnan
+// AMDGCN-NEXT: [[TMP3:%.*]] = fcmp ult float [[TMP2]], 1.310720e+05
+// AMDGCN-NEXT: br i1 [[TMP3]], label %[[BB108:.*]], label %[[BB4:.*]]
+// AMDGCN: [[BB4]]:
+// AMDGCN-NEXT: [[TMP5:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float [[TMP2]])
+// AMDGCN-NEXT: [[TMP6:%.*]] = extractvalue { float, i32 } [[TMP5]], 1
+// AMDGCN-NEXT: [[TMP7:%.*]] = extractvalue { float, i32 } [[TMP5]], 0
+// AMDGCN-NEXT: [[TMP8:%.*]] = fmul contract float [[TMP7]], f0x4B800000
+// AMDGCN-NEXT: [[TMP9:%.*]] = fptoui float [[TMP8]] to i32
+// AMDGCN-NEXT: [[TMP10:%.*]] = zext i32 [[TMP9]] to i64
+// AMDGCN-NEXT: [[TMP11:%.*]] = mul nuw i64 [[TMP10]], 4266746795
+// AMDGCN-NEXT: [[TMP12:%.*]] = trunc i64 [[TMP11]] to i32
+// AMDGCN-NEXT: [[TMP13:%.*]] = lshr i64 [[TMP11]], 32
+// AMDGCN-NEXT: [[TMP14:%.*]] = mul nuw nsw i64 [[TMP10]], 1011060801
+// AMDGCN-NEXT: [[TMP15:%.*]] = add nuw nsw i64 [[TMP13]], [[TMP14]]
+// AMDGCN-NEXT: [[TMP16:%.*]] = trunc i64 [[TMP15]] to i32
+// AMDGCN-NEXT: [[TMP17:%.*]] = lshr i64 [[TMP15]], 32
+// AMDGCN-NEXT: [[TMP18:%.*]] = mul nuw i64 [[TMP10]], 3680671129
+// AMDGCN-NEXT: [[TMP19:%.*]] = add nuw i64 [[TMP17]], [[TMP18]]
+// AMDGCN-NEXT: [[TMP20:%.*]] = trunc i64 [[TMP19]] to i32
+// AMDGCN-NEXT: [[TMP21:%.*]] = lshr i64 [[TMP19]], 32
+// AMDGCN-NEXT: [[TMP22:%.*]] = mul nuw i64 [[TMP10]], 4113882560
+// AMDGCN-NEXT: [[TMP23:%.*]] = add nuw i64 [[TMP21]], [[TMP22]]
+// AMDGCN-NEXT: [[TMP24:%.*]] = trunc i64 [[TMP23]] to i32
+// AMDGCN-NEXT: [[TMP25:%.*]] = lshr i64 [[TMP23]], 32
+// AMDGCN-NEXT: [[TMP26:%.*]] = mul nuw i64 [[TMP10]], 4230436817
+// AMDGCN-NEXT: [[TMP27:%.*]] = add nuw i64 [[TMP25]], [[TMP26]]
+// AMDGCN-NEXT: [[TMP28:%.*]] = trunc i64 [[TMP27]] to i32
+// AMDGCN-NEXT: [[TMP29:%.*]] = lshr i64 [[TMP27]], 32
+// AMDGCN-NEXT: [[TMP30:%.*]] = mul nuw nsw i64 [[TMP10]], 1313084713
+// AMDGCN-NEXT: [[TMP31:%.*]] = add nuw nsw i64 [[TMP29]], [[TMP30]]
+// AMDGCN-NEXT: [[TMP32:%.*]] = trunc i64 [[TMP31]] to i32
+// AMDGCN-NEXT: [[TMP33:%.*]] = lshr i64 [[TMP31]], 32
+// AMDGCN-NEXT: [[TMP34:%.*]] = mul nuw i64 [[TMP10]], 2734261102
+// AMDGCN-NEXT: [[TMP35:%.*]] = add nuw i64 [[TMP33]], [[TMP34]]
+// AMDGCN-NEXT: [[TMP36:%.*]] = trunc i64 [[TMP35]] to i32
+// AMDGCN-NEXT: [[TMP37:%.*]] = lshr i64 [[TMP35]], 32
+// AMDGCN-NEXT: [[TMP38:%.*]] = trunc nuw i64 [[TMP37]] to i32
+// AMDGCN-NEXT: [[TMP39:%.*]] = add nsw i32 [[TMP6]], 6
+// AMDGCN-NEXT: [[TMP40:%.*]] = icmp ugt i32 [[TMP39]], 63
+// AMDGCN-NEXT: [[TMP41:%.*]] = select i1 [[TMP40]], i32 [[TMP32]], i32 [[TMP38]]
+// AMDGCN-NEXT: [[TMP42:%.*]] = select i1 [[TMP40]], i32 [[TMP28]], i32 [[TMP36]]
+// AMDGCN-NEXT: [[TMP43:%.*]] = select i1 [[TMP40]], i32 [[TMP24]], i32 [[TMP32]]
+// AMDGCN-NEXT: [[TMP44:%.*]] = select i1 [[TMP40]], i32 [[TMP20]], i32 [[TMP28]]
+// AMDGCN-NEXT: [[TMP45:%.*]] = select i1 [[TMP40]], i32 [[TMP16]], i32 [[TMP24]]
+// AMDGCN-NEXT: [[TMP46:%.*]] = select i1 [[TMP40]], i32 [[TMP12]], i32 [[TMP20]]
+// AMDGCN-NEXT: [[TMP47:%.*]] = select i1 [[TMP40]], i32 -64, i32 0
+// AMDGCN-NEXT: [[TMP48:%.*]] = add i32 [[TMP47]], [[TMP39]]
+// AMDGCN-NEXT: [[TMP49:%.*]] = icmp ugt i32 [[TMP48]], 31
+// AMDGCN-NEXT: [[TMP50:%.*]] = select i1 [[TMP49]], i32 [[TMP42]], i32 [[TMP41]]
+// AMDGCN-NEXT: [[TMP51:%.*]] = select i1 [[TMP49]], i32 [[TMP43]], i32 [[TMP42]]
+// AMDGCN-NEXT: [[TMP52:%.*]] = select i1 [[TMP49]], i32 [[TMP44]], i32 [[TMP43]]
+// AMDGCN-NEXT: [[TMP53:%.*]] = select i1 [[TMP49]], i32 [[TMP45]], i32 [[TMP44]]
+// AMDGCN-NEXT: [[TMP54:%.*]] = select i1 [[TMP49]], i32 [[TMP46]], i32 [[TMP45]]
+// AMDGCN-NEXT: [[TMP55:%.*]] = select i1 [[TMP49]], i32 -32, i32 0
+// AMDGCN-NEXT: [[TMP56:%.*]] = add i32 [[TMP55]], [[TMP48]]
+// AMDGCN-NEXT: [[TMP57:%.*]] = icmp ugt i32 [[TMP56]], 31
+// AMDGCN-NEXT: [[TMP58:%.*]] = select i1 [[TMP57]], i32 [[TMP51]], i32 [[TMP50]]
+// AMDGCN-NEXT: [[TMP59:%.*]] = select i1 [[TMP57]], i32 [[TMP52]], i32 [[TMP51]]
+// AMDGCN-NEXT: [[TMP60:%.*]] = select i1 [[TMP57]], i32 [[TMP53]], i32 [[TMP52]]
+// AMDGCN-NEXT: [[TMP61:%.*]] = select i1 [[TMP57]], i32 [[TMP54]], i32 [[TMP53]]
+// AMDGCN-NEXT: [[TMP62:%.*]] = select i1 [[TMP57]], i32 -32, i32 0
+// AMDGCN-NEXT: [[TMP63:%.*]] = add i32 [[TMP62]], [[TMP56]]
+// AMDGCN-NEXT: [[TMP64:%.*]] = icmp eq i32 [[TMP63]], 0
+// AMDGCN-NEXT: [[TMP65:%.*]] = sub i32 32, [[TMP63]]
+// AMDGCN-NEXT: [[TMP66:%.*]] = tail call i32 @llvm.fshr.i32(i32 [[TMP58]], i32 [[TMP59]], i32 [[TMP65]])
+// AMDGCN-NEXT: [[TMP67:%.*]] = tail call i32 @llvm.fshr.i32(i32 [[TMP59]], i32 [[TMP60]], i32 [[TMP65]])
+// AMDGCN-NEXT: [[TMP68:%.*]] = tail call i32 @llvm.fshr.i32(i32 [[TMP60]], i32 [[TMP61]], i32 [[TMP65]])
+// AMDGCN-NEXT: [[TMP69:%.*]] = select i1 [[TMP64]], i32 [[TMP58]], i32 [[TMP66]]
+// AMDGCN-NEXT: [[TMP70:%.*]] = select i1 [[TMP64]], i32 [[TMP59]], i32 [[TMP67]]
+// AMDGCN-NEXT: [[TMP71:%.*]] = select i1 [[TMP64]], i32 [[TMP60]], i32 [[TMP68]]
+// AMDGCN-NEXT: [[TMP72:%.*]] = lshr i32 [[TMP69]], 29
+// AMDGCN-NEXT: [[TMP73:%.*]] = tail call i32 @llvm.fshl.i32(i32 [[TMP69]], i32 [[TMP70]], i32 2)
+// AMDGCN-NEXT: [[TMP74:%.*]] = tail call i32 @llvm.fshl.i32(i32 [[TMP70]], i32 [[TMP71]], i32 2)
+// AMDGCN-NEXT: [[TMP75:%.*]] = tail call i32 @llvm.fshl.i32(i32 [[TMP71]], i32 [[TMP61]], i32 2)
+// AMDGCN-NEXT: [[TMP76:%.*]] = and i32 [[TMP72]], 1
+// AMDGCN-NEXT: [[TMP77:%.*]] = sub nsw i32 0, [[TMP76]]
+// AMDGCN-NEXT: [[TMP78:%.*]] = shl i32 [[TMP72]], 31
+// AMDGCN-NEXT: [[TMP79:%.*]] = xor i32 [[TMP73]], [[TMP77]]
+// AMDGCN-NEXT: [[TMP80:%.*]] = xor i32 [[TMP74]], [[TMP77]]
+// AMDGCN-NEXT: [[TMP81:%.*]] = xor i32 [[TMP75]], [[TMP77]]
+// AMDGCN-NEXT: [[TMP82:%.*]] = tail call range(i32 0, 33) i32 @llvm.ctlz.i32(i32 [[TMP79]], i1 false)
+// AMDGCN-NEXT: [[DOTNEG_I_I_I_I:%.*]] = xor i32 [[TMP82]], -1
+// AMDGCN-NEXT: [[TMP83:%.*]] = sub nsw i32 31, [[TMP82]]
+// AMDGCN-NEXT: [[TMP84:%.*]] = tail call i32 @llvm.fshr.i32(i32 [[TMP79]], i32 [[TMP80]], i32 [[TMP83]])
+// AMDGCN-NEXT: [[TMP85:%.*]] = tail call i32 @llvm.fshr.i32(i32 [[TMP80]], i32 [[TMP81]], i32 [[TMP83]])
+// AMDGCN-NEXT: [[TMP86:%.*]] = lshr i32 [[TMP84]], 9
+// AMDGCN-NEXT: [[TMP87:%.*]] = shl nuw nsw i32 [[TMP82]], 23
+// AMDGCN-NEXT: [[REASS_SUB:%.*]] = sub nsw i32 [[TMP86]], [[TMP87]]
+// AMDGCN-NEXT: [[TMP88:%.*]] = add nsw i32 [[REASS_SUB]], 1056964608
+// AMDGCN-NEXT: [[TMP89:%.*]] = or i32 [[TMP88]], [[TMP78]]
+// AMDGCN-NEXT: [[TMP90:%.*]] = bitcast i32 [[TMP89]] to float
+// AMDGCN-NEXT: [[TMP91:%.*]] = tail call i32 @llvm.fshl.i32(i32 [[TMP84]], i32 [[TMP85]], i32 23)
+// AMDGCN-NEXT: [[TMP92:%.*]] = tail call range(i32 0, 33) i32 @llvm.ctlz.i32(i32 [[TMP91]], i1 false)
+// AMDGCN-NEXT: [[TMP93:%.*]] = xor i32 [[TMP92]], -1
+// AMDGCN-NEXT: [[TMP94:%.*]] = tail call i32 @llvm.fshr.i32(i32 [[TMP91]], i32 [[TMP85]], i32 [[TMP93]])
+// AMDGCN-NEXT: [[DOTNEG3_I_I_I_I:%.*]] = sub nuw nsw i32 [[DOTNEG_I_I_I_I]], [[TMP92]]
+// AMDGCN-NEXT: [[TMP95:%.*]] = lshr i32 [[TMP94]], 9
+// AMDGCN-NEXT: [[DOTNEG4_I_I_I_I:%.*]] = shl nsw i32 [[DOTNEG3_I_I_I_I]], 23
+// AMDGCN-NEXT: [[TMP96:%.*]] = add nsw i32 [[DOTNEG4_I_I_I_I]], 864026624
+// AMDGCN-NEXT: [[TMP97:%.*]] = or disjoint i32 [[TMP96]], [[TMP95]]
+// AMDGCN-NEXT: [[TMP98:%.*]] = or i32 [[TMP97]], [[TMP78]]
+// AMDGCN-NEXT: [[TMP99:%.*]] = bitcast i32 [[TMP98]] to float
+// AMDGCN-NEXT: [[TMP100:%.*]] = fmul float [[TMP90]], f0x3FC90FDA
+// AMDGCN-NEXT: [[TMP101:%.*]] = fneg contract float [[TMP100]]
+// AMDGCN-NEXT: [[TMP102:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[TMP90]], float f0x3FC90FDA, float [[TMP101]])
+// AMDGCN-NEXT: [[TMP103:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[TMP90]], float f0x33A22168, float [[TMP102]])
+// AMDGCN-NEXT: [[TMP104:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[TMP99]], float f0x3FC90FDA, float [[TMP103]])
+// AMDGCN-NEXT: [[TMP105:%.*]] = fadd float [[TMP100]], [[TMP104]]
+// AMDGCN-NEXT: [[TMP106:%.*]] = lshr i32 [[TMP69]], 30
+// AMDGCN-NEXT: [[TMP107:%.*]] = add nuw nsw i32 [[TMP76]], [[TMP106]]
+// AMDGCN-NEXT: br label %[[_Z3COSF_EXIT:.*]]
+// AMDGCN: [[BB108]]:
+// AMDGCN-NEXT: [[TMP109:%.*]] = fmul float [[TMP2]], f0x3F22F983
+// AMDGCN-NEXT: [[TMP110:%.*]] = tail call contract noundef float @llvm.rint.f32(float [[TMP109]])
+// AMDGCN-NEXT: [[TMP111:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[TMP110]], float f0xBFC90FDA, float [[TMP2]])
+// AMDGCN-NEXT: [[TMP112:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[TMP110]], float f0xB3A22168, float [[TMP111]])
+// AMDGCN-NEXT: [[TMP113:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[TMP110]], float f0xA7C234C4, float [[TMP112]])
+// AMDGCN-NEXT: [[TMP114:%.*]] = fptosi float [[TMP110]] to i32
+// AMDGCN-NEXT: br label %[[_Z3COSF_EXIT]]
+// AMDGCN: [[_Z3COSF_EXIT]]:
+// AMDGCN-NEXT: [[DOTSINK_I_I_I_I:%.*]] = phi float [ [[TMP113]], %[[BB108]] ], [ [[TMP105]], %[[BB4]] ]
+// AMDGCN-NEXT: [[TMP115:%.*]] = phi i32 [ [[TMP114]], %[[BB108]] ], [ [[TMP107]], %[[BB4]] ]
+// AMDGCN-NEXT: [[TMP116:%.*]] = fmul float [[DOTSINK_I_I_I_I]], [[DOTSINK_I_I_I_I]]
+// AMDGCN-NEXT: [[TMP117:%.*]] = tail call noundef float @llvm.fmuladd.f32(float [[TMP116]], float f0xB94C1982, float f0x3C0881C4)
+// AMDGCN-NEXT: [[TMP118:%.*]] = tail call noundef float @llvm.fmuladd.f32(float [[TMP116]], float [[TMP117]], float f0xBE2AAA9D)
+// AMDGCN-NEXT: [[TMP119:%.*]] = fmul float [[TMP116]], [[TMP118]]
+// AMDGCN-NEXT: [[TMP120:%.*]] = tail call noundef float @llvm.fmuladd.f32(float [[DOTSINK_I_I_I_I]], float [[TMP119]], float [[DOTSINK_I_I_I_I]])
+// AMDGCN-NEXT: [[TMP121:%.*]] = tail call noundef float @llvm.fmuladd.f32(float [[TMP116]], float f0x37D75334, float f0xBAB64F3B)
+// AMDGCN-NEXT: [[TMP122:%.*]] = tail call noundef float @llvm.fmuladd.f32(float [[TMP116]], float [[TMP121]], float f0x3D2AABF7)
+// AMDGCN-NEXT: [[TMP123:%.*]] = tail call noundef float @llvm.fmuladd.f32(float [[TMP116]], float [[TMP122]], float f0xBF000004)
+// AMDGCN-NEXT: [[TMP124:%.*]] = tail call noundef float @llvm.fmuladd.f32(float [[TMP116]], float [[TMP123]], float 1.000000e+00)
+// AMDGCN-NEXT: [[TMP125:%.*]] = and i32 [[TMP115]], 1
+// AMDGCN-NEXT: [[TMP126:%.*]] = icmp eq i32 [[TMP125]], 0
+// AMDGCN-NEXT: [[TMP127:%.*]] = shl i32 [[TMP115]], 30
+// AMDGCN-NEXT: [[TMP128:%.*]] = and i32 [[TMP127]], -2147483648
+// AMDGCN-NEXT: [[TMP129:%.*]] = fneg contract float [[TMP120]]
+// AMDGCN-NEXT: [[TMP130:%.*]] = select contract i1 [[TMP126]], float [[TMP124]], float [[TMP129]]
+// AMDGCN-NEXT: [[TMP131:%.*]] = bitcast float [[TMP130]] to i32
+// AMDGCN-NEXT: [[TMP132:%.*]] = xor i32 [[TMP128]], [[TMP131]]
+// AMDGCN-NEXT: [[TMP133:%.*]] = bitcast i32 [[TMP132]] to float
+// AMDGCN-NEXT: ret float [[TMP133]]
+//
+float test_float(float x) {
+ return cos(x);
+}
diff --git a/libclc/test/math/fabs.cl b/libclc/test/math/fabs.cl
new file mode 100644
index 0000000000000..8b2e34ad51594
--- /dev/null
+++ b/libclc/test/math/fabs.cl
@@ -0,0 +1,14 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
+// REQUIRES: AMDGCN
+
+// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+
+// AMDGCN-LABEL: define hidden noundef float @test_float(
+// AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
+// AMDGCN-NEXT: [[ENTRY:.*:]]
+// AMDGCN-NEXT: [[CALL:%.*]] = tail call fastcc float @llvm.fabs.f32(float noundef [[X]]) #[[ATTR2:[0-9]+]]
+// AMDGCN-NEXT: ret float [[CALL]]
+//
+float test_float(float x) {
+ return fabs(x);
+}
diff --git a/libclc/test/math/rsqrt.cl b/libclc/test/math/rsqrt.cl
new file mode 100644
index 0000000000000..f36e6a01ed514
--- /dev/null
+++ b/libclc/test/math/rsqrt.cl
@@ -0,0 +1,45 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
+// REQUIRES: AMDGCN
+
+// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+// AMDGCN-LABEL: define hidden noundef half @test_half(
+// AMDGCN-SAME: half noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
+// AMDGCN-NEXT: [[ENTRY:.*:]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract half @llvm.sqrt.f16(half [[X]]), !fpmath [[META13:![0-9]+]]
+// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract half 1.000000e+00, [[TMP0]], !fpmath [[META14:![0-9]+]]
+// AMDGCN-NEXT: ret half [[TMP1]]
+//
+half test_half(half x) {
+ return rsqrt(x);
+}
+
+// AMDGCN-LABEL: define hidden noundef float @test_float(
+// AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// AMDGCN-NEXT: [[ENTRY:.*:]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.sqrt.f32(float [[X]]), !fpmath [[META15:![0-9]+]]
+// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract float 1.000000e+00, [[TMP0]], !fpmath [[META16:![0-9]+]]
+// AMDGCN-NEXT: ret float [[TMP1]]
+//
+float test_float(float x) {
+ return rsqrt(x);
+}
+
+// AMDGCN-LABEL: define hidden noundef double @test_double(
+// AMDGCN-SAME: double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0]] {
+// AMDGCN-NEXT: [[ENTRY:.*:]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.sqrt.f64(double [[X]])
+// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract double 1.000000e+00, [[TMP0]]
+// AMDGCN-NEXT: ret double [[TMP1]]
+//
+double test_double(double x) {
+ return rsqrt(x);
+}
+//.
+// AMDGCN: [[META13]] = !{float 1.500000e+00}
+// AMDGCN: [[META14]] = !{float 1.000000e+00}
+// AMDGCN: [[META15]] = !{float 3.000000e+00}
+// AMDGCN: [[META16]] = !{float 2.500000e+00}
+//.
diff --git a/libclc/test/misc/as_type.cl b/libclc/test/misc/as_type.cl
new file mode 100644
index 0000000000000..d5c41b15162a6
--- /dev/null
+++ b/libclc/test/misc/as_type.cl
@@ -0,0 +1,14 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
+// REQUIRES: AMDGCN
+
+// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+
+// AMDGCN-LABEL: define hidden noundef <4 x i32> @test_float4_as_int4(
+// AMDGCN-SAME: <4 x float> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
+// AMDGCN-NEXT: [[ENTRY:.*:]]
+// AMDGCN-NEXT: [[ASTYPE:%.*]] = bitcast <4 x float> [[X]] to <4 x i32>
+// AMDGCN-NEXT: ret <4 x i32> [[ASTYPE]]
+//
+int4 test_float4_as_int4(float4 x) {
+ return as_int4(x);
+}
diff --git a/libclc/test/rsqrt.cl b/libclc/test/rsqrt.cl
deleted file mode 100644
index 4eebfe8ecf7f9..0000000000000
--- a/libclc/test/rsqrt.cl
+++ /dev/null
@@ -1,14 +0,0 @@
-//===----------------------------------------------------------------------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-#pragma OPENCL EXTENSION cl_khr_fp64 : enable
-
-__kernel void foo(float4 *x, double4 *y) {
- x[1] = rsqrt(x[0]);
- y[1] = rsqrt(y[0]);
-}
diff --git a/libclc/test/subsat.cl b/libclc/test/subsat.cl
deleted file mode 100644
index 5e6fbdcfbef9e..0000000000000
--- a/libclc/test/subsat.cl
+++ /dev/null
@@ -1,27 +0,0 @@
-//===----------------------------------------------------------------------===//
-//
-// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
-// See https://llvm.org/LICENSE.txt for license information.
-// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
-//
-//===----------------------------------------------------------------------===//
-
-__kernel void test_subsat_char(char *a, char x, char y) {
- *a = sub_sat(x, y);
- return;
-}
-
-__kernel void test_subsat_uchar(uchar *a, uchar x, uchar y) {
- *a = sub_sat(x, y);
- return;
-}
-
-__kernel void test_subsat_long(long *a, long x, long y) {
- *a = sub_sat(x, y);
- return;
-}
-
-__kernel void test_subsat_ulong(ulong *a, ulong x, ulong y) {
- *a = sub_sat(x, y);
- return;
-}
\ No newline at end of file
diff --git a/libclc/test/update_libclc_tests.py b/libclc/test/update_libclc_tests.py
new file mode 100755
index 0000000000000..40c48b82b8332
--- /dev/null
+++ b/libclc/test/update_libclc_tests.py
@@ -0,0 +1,146 @@
+#!/usr/bin/env python3
+"""A utility to wrap utils/update_cc_test_checks.py for updating CHECK lines in
+libclc .cl tests for a given architecture.
+
+The script accepts an architecture argument to determine the triple and check
+prefix. Supported arch values: amdgpu, amdgcn, nvptx64, spirv, spirv64.
+
+The script does 3 things:
+1. Replaces %libclc_target and %check_prefix in the .cl file for the arch.
+2. Runs update_cc_test_checks.py to update CHECK lines.
+3. Reverts the .cl file back to using %libclc_target and %check_prefix.
+
+Usage:
+
+% libclc/test/update_libclc_tests.py amdgpu
+
+"""
+
+import argparse
+import re
+import subprocess
+import sys
+from concurrent.futures import ThreadPoolExecutor, as_completed
+from pathlib import Path
+
+ARCH_TO_TRIPLE = {
+ "amdgpu": "amdgcn-amd-amdhsa-llvm",
+ "amdgcn": "amdgcn-amd-amdhsa-llvm",
+ "nvptx64": "nvptx64-nvidia-cuda",
+ "spirv": "spirv-unknown-mesa3d",
+ "spirv64": "spirv64-unknown-mesa3d",
+}
+
+SCRIPT_DIR = Path(__file__).parent
+REPO_ROOT = SCRIPT_DIR.parent.parent
+UPDATE_SCRIPT = REPO_ROOT / "llvm" / "utils" / "update_cc_test_checks.py"
+CLANG = REPO_ROOT / "build" / "bin" / "clang"
+
+
+def find_cl_files(test_dir: Path):
+ return list(test_dir.rglob("*.cl"))
+
+
+def replace_in_file(path: Path, triple: str, check_prefix: str):
+ content = path.read_bytes()
+ content = content.replace(b"%libclc_target", triple.encode())
+ content = content.replace(b"%check_prefix", check_prefix.encode())
+ path.write_bytes(content)
+
+
+def revert_in_file(path: Path, triple: str, check_prefix: str):
+ # Only revert in the RUN line context, not in generated CHECK lines.
+ content = path.read_bytes()
+ content = content.replace(
+ f"-target {triple}".encode(), b"-target %libclc_target"
+ )
+ content = content.replace(
+ f"--check-prefix={check_prefix}".encode(), b"--check-prefix=%check_prefix"
+ )
+ path.write_bytes(content)
+
+
+def file_requires_feature(path: Path, feature: str) -> bool:
+ try:
+ text = path.read_text(encoding="utf-8", errors="replace")
+ except Exception:
+ return False
+ for line in text.splitlines():
+ stripped = line.strip().lstrip("//").strip()
+ if stripped.startswith("REQUIRES:"):
+ rest = stripped[len("REQUIRES:"):]
+ features = [f.strip() for f in re.split(r",|\|\|", rest)]
+ if feature in features:
+ return True
+ return False
+
+
+def process_file(cl_file: Path, triple: str, check_prefix: str) -> bool:
+ replace_in_file(cl_file, triple, check_prefix)
+ cmd = [
+ sys.executable,
+ str(UPDATE_SCRIPT),
+ "--clang", str(CLANG),
+ str(cl_file),
+ ]
+ print(f" update: {cl_file.relative_to(REPO_ROOT)}")
+ result = subprocess.run(cmd, capture_output=True, text=True)
+ ok = result.returncode == 0
+ if not ok:
+ print(f" FAILED: {result.stderr.strip()}", file=sys.stderr)
+ revert_in_file(cl_file, triple, check_prefix)
+ return ok
+
+
+def main():
+ parser = argparse.ArgumentParser(
+ description="Update libclc FileCheck assertions for a given arch."
+ )
+ parser.add_argument(
+ "arch",
+ choices=list(ARCH_TO_TRIPLE.keys()),
+ help="Target arch: amdgpu, amdgcn, nvptx64, spirv, spirv64",
+ )
+ args = parser.parse_args()
+
+ arch = args.arch.lower()
+ triple = ARCH_TO_TRIPLE[arch]
+ # check_prefix matches REQUIRES feature: uppercase of canonical arch name
+ # amdgpu -> AMDGCN (same triple as amdgcn), others uppercased
+ if arch == "amdgpu":
+ check_prefix = "AMDGCN"
+ else:
+ check_prefix = arch.upper()
+
+ cl_files = find_cl_files(SCRIPT_DIR)
+ # Only process files that REQUIRES this feature
+ target_files = [f for f in cl_files if file_requires_feature(f, check_prefix)]
+
+ if not target_files:
+ print(f"No .cl files found with REQUIRES: {check_prefix}")
+ return
+
+ print(f"arch={arch} triple={triple} check_prefix={check_prefix}")
+ print(f"Processing {len(target_files)} file(s)...")
+
+ failed = []
+ with ThreadPoolExecutor() as executor:
+ futures = {
+ executor.submit(process_file, f, triple, check_prefix): f
+ for f in target_files
+ }
+ for future in as_completed(futures):
+ if not future.result():
+ failed.append(futures[future])
+
+ if failed:
+ print(f"\n{len(failed)} file(s) failed:", file=sys.stderr)
+ for f in failed:
+ print(f" {f}", file=sys.stderr)
+ sys.exit(1)
+ else:
+ print(f"Done. Updated {len(target_files)} file(s).")
+
+
+if __name__ == "__main__":
+ main()
diff --git a/libclc/test/work-item/get_group_id.cl b/libclc/test/work-item/get_group_id.cl
new file mode 100644
index 0000000000000..57e3d68787536
--- /dev/null
+++ b/libclc/test/work-item/get_group_id.cl
@@ -0,0 +1,15 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
+// REQUIRES: AMDGCN
+
+// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+
+// AMDGCN-LABEL: define hidden range(i64 0, 4294967296) i64 @test(
+// AMDGCN-SAME: ) local_unnamed_addr #[[ATTR0:[0-9]+]] {
+// AMDGCN-NEXT: [[ENTRY:.*:]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.workgroup.id.x()
+// AMDGCN-NEXT: [[TMP1:%.*]] = zext i32 [[TMP0]] to i64
+// AMDGCN-NEXT: ret i64 [[TMP1]]
+//
+size_t test() {
+ return get_group_id(0);
+}
>From 6518f001607e5792d01ec2cb4c2f899a6cb1f6a8 Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Tue, 12 May 2026 13:17:04 +0200
Subject: [PATCH 2/7] disable clang-format in test folder
---
libclc/test/.clang-format | 1 +
1 file changed, 1 insertion(+)
create mode 100644 libclc/test/.clang-format
diff --git a/libclc/test/.clang-format b/libclc/test/.clang-format
new file mode 100644
index 0000000000000..e3845288a2aec
--- /dev/null
+++ b/libclc/test/.clang-format
@@ -0,0 +1 @@
+DisableFormat: true
>From f3d10dbefaa60023085993ec94cf243a9b307ee2 Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Tue, 12 May 2026 13:56:38 +0200
Subject: [PATCH 3/7] add -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins
-Xclang -finclude-default-header
---
libclc/test/conversion/convert.cl | 2 +-
libclc/test/geometric/cross.cl | 2 +-
libclc/test/integer/add_sat.cl | 2 +-
libclc/test/integer/sub_sat.cl | 2 +-
libclc/test/math/cos.cl | 2 +-
libclc/test/math/fabs.cl | 2 +-
libclc/test/math/rsqrt.cl | 18 +++++++++---------
libclc/test/misc/as_type.cl | 2 +-
libclc/test/work-item/get_group_id.cl | 2 +-
9 files changed, 17 insertions(+), 17 deletions(-)
diff --git a/libclc/test/conversion/convert.cl b/libclc/test/conversion/convert.cl
index 00b524bb8f11d..eb673a493b403 100644
--- a/libclc/test/conversion/convert.cl
+++ b/libclc/test/conversion/convert.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// REQUIRES: AMDGCN
-// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden <4 x i32> @float4_to_int4(
// AMDGCN-SAME: <4 x float> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
diff --git a/libclc/test/geometric/cross.cl b/libclc/test/geometric/cross.cl
index 5aa7541b4d49c..79724183760e7 100644
--- a/libclc/test/geometric/cross.cl
+++ b/libclc/test/geometric/cross.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// REQUIRES: AMDGCN
-// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden <4 x float> @test_float4(
// AMDGCN-SAME: <4 x float> noundef [[X:%.*]], <4 x float> noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
diff --git a/libclc/test/integer/add_sat.cl b/libclc/test/integer/add_sat.cl
index 4c30090953e99..e223b28444835 100644
--- a/libclc/test/integer/add_sat.cl
+++ b/libclc/test/integer/add_sat.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// REQUIRES: AMDGCN
-// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden noundef signext i8 @test_char(
// AMDGCN-SAME: i8 noundef signext [[X:%.*]], i8 noundef signext [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
diff --git a/libclc/test/integer/sub_sat.cl b/libclc/test/integer/sub_sat.cl
index 6573afae45102..6ba40f75870db 100644
--- a/libclc/test/integer/sub_sat.cl
+++ b/libclc/test/integer/sub_sat.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// REQUIRES: AMDGCN
-// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden noundef signext i8 @test_char(
// AMDGCN-SAME: i8 noundef signext [[X:%.*]], i8 noundef signext [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
diff --git a/libclc/test/math/cos.cl b/libclc/test/math/cos.cl
index 311024d7155e0..40d6825a034a0 100644
--- a/libclc/test/math/cos.cl
+++ b/libclc/test/math/cos.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// REQUIRES: AMDGCN
-// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden float @test_float(
// AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
diff --git a/libclc/test/math/fabs.cl b/libclc/test/math/fabs.cl
index 8b2e34ad51594..3b6ec0f8b729e 100644
--- a/libclc/test/math/fabs.cl
+++ b/libclc/test/math/fabs.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// REQUIRES: AMDGCN
-// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden noundef float @test_float(
// AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
diff --git a/libclc/test/math/rsqrt.cl b/libclc/test/math/rsqrt.cl
index f36e6a01ed514..f0059742d18e6 100644
--- a/libclc/test/math/rsqrt.cl
+++ b/libclc/test/math/rsqrt.cl
@@ -1,15 +1,15 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// REQUIRES: AMDGCN
-// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
// AMDGCN-LABEL: define hidden noundef half @test_half(
// AMDGCN-SAME: half noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
// AMDGCN-NEXT: [[ENTRY:.*:]]
-// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract half @llvm.sqrt.f16(half [[X]]), !fpmath [[META13:![0-9]+]]
-// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract half 1.000000e+00, [[TMP0]], !fpmath [[META14:![0-9]+]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract half @llvm.sqrt.f16(half [[X]]), !fpmath [[META11:![0-9]+]]
+// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract half 1.000000e+00, [[TMP0]], !fpmath [[META12:![0-9]+]]
// AMDGCN-NEXT: ret half [[TMP1]]
//
half test_half(half x) {
@@ -19,8 +19,8 @@ half test_half(half x) {
// AMDGCN-LABEL: define hidden noundef float @test_float(
// AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0]] {
// AMDGCN-NEXT: [[ENTRY:.*:]]
-// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.sqrt.f32(float [[X]]), !fpmath [[META15:![0-9]+]]
-// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract float 1.000000e+00, [[TMP0]], !fpmath [[META16:![0-9]+]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.sqrt.f32(float [[X]]), !fpmath [[META13:![0-9]+]]
+// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract float 1.000000e+00, [[TMP0]], !fpmath [[META14:![0-9]+]]
// AMDGCN-NEXT: ret float [[TMP1]]
//
float test_float(float x) {
@@ -38,8 +38,8 @@ double test_double(double x) {
return rsqrt(x);
}
//.
-// AMDGCN: [[META13]] = !{float 1.500000e+00}
-// AMDGCN: [[META14]] = !{float 1.000000e+00}
-// AMDGCN: [[META15]] = !{float 3.000000e+00}
-// AMDGCN: [[META16]] = !{float 2.500000e+00}
+// AMDGCN: [[META11]] = !{float 1.500000e+00}
+// AMDGCN: [[META12]] = !{float 1.000000e+00}
+// AMDGCN: [[META13]] = !{float 3.000000e+00}
+// AMDGCN: [[META14]] = !{float 2.500000e+00}
//.
diff --git a/libclc/test/misc/as_type.cl b/libclc/test/misc/as_type.cl
index d5c41b15162a6..0ba62fefa4d28 100644
--- a/libclc/test/misc/as_type.cl
+++ b/libclc/test/misc/as_type.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// REQUIRES: AMDGCN
-// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden noundef <4 x i32> @test_float4_as_int4(
// AMDGCN-SAME: <4 x float> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
diff --git a/libclc/test/work-item/get_group_id.cl b/libclc/test/work-item/get_group_id.cl
index 57e3d68787536..cc26101c1b99f 100644
--- a/libclc/test/work-item/get_group_id.cl
+++ b/libclc/test/work-item/get_group_id.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// REQUIRES: AMDGCN
-// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden range(i64 0, 4294967296) i64 @test(
// AMDGCN-SAME: ) local_unnamed_addr #[[ATTR0:[0-9]+]] {
>From 7839a366682091baee84b184ff856bfa4333e425 Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Wed, 13 May 2026 05:44:26 +0200
Subject: [PATCH 4/7] mcpu, amdgpu-registered-target, ProcessPoolExecutor
---
libclc/CMakeLists.txt | 2 -
libclc/test/AMDGPU/external-funcs.test | 2 +-
libclc/test/AMDGPU/lit.local.cfg | 2 +-
libclc/test/CMakeLists.txt | 6 +++
libclc/test/conversion/convert.cl | 4 +-
libclc/test/geometric/cross.cl | 4 +-
libclc/test/integer/add_sat.cl | 4 +-
libclc/test/integer/sub_sat.cl | 4 +-
libclc/test/lit.cfg.py | 24 ++++++++---
libclc/test/lit.site.cfg.py.in | 1 +
libclc/test/math/cos.cl | 4 +-
libclc/test/math/fabs.cl | 4 +-
libclc/test/math/rsqrt.cl | 20 ++++-----
libclc/test/misc/as_type.cl | 4 +-
libclc/test/update_libclc_tests.py | 57 ++++++++++++++++++--------
libclc/test/work-item/get_group_id.cl | 4 +-
16 files changed, 95 insertions(+), 51 deletions(-)
diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt
index 3f84458336950..a169af7f73cf9 100644
--- a/libclc/CMakeLists.txt
+++ b/libclc/CMakeLists.txt
@@ -261,8 +261,6 @@ add_libclc_library(libclc-${LIBCLC_TARGET}
PARENT_TARGET libclc-opencl-builtins
)
-set(LIBCLC_UNRESOLVED_SYMBOL_TEST_TARGETS libclc-${LIBCLC_TARGET})
-
if(LLVM_INCLUDE_TESTS)
add_subdirectory(test)
endif()
diff --git a/libclc/test/AMDGPU/external-funcs.test b/libclc/test/AMDGPU/external-funcs.test
index 73bddc02d11e1..ab877fb60120d 100644
--- a/libclc/test/AMDGPU/external-funcs.test
+++ b/libclc/test/AMDGPU/external-funcs.test
@@ -1 +1 @@
-; RUN: llvm-nm -u "%libclc_library_dir/%libclc_target/libclc.bc" | FileCheck %s --allow-empty --implicit-check-not=" U "
+; RUN: llvm-nm -u "%library_dir/%target/libclc.bc" | FileCheck %s --allow-empty --implicit-check-not=" U "
diff --git a/libclc/test/AMDGPU/lit.local.cfg b/libclc/test/AMDGPU/lit.local.cfg
index 1ee232d629c16..68d2d4d5d0023 100644
--- a/libclc/test/AMDGPU/lit.local.cfg
+++ b/libclc/test/AMDGPU/lit.local.cfg
@@ -1,2 +1,2 @@
-if "AMDGCN" not in config.available_features:
+if "amdgpu-registered-target" not in config.available_features:
config.unsupported = True
diff --git a/libclc/test/CMakeLists.txt b/libclc/test/CMakeLists.txt
index 955235f32277d..f1b119df09f1b 100644
--- a/libclc/test/CMakeLists.txt
+++ b/libclc/test/CMakeLists.txt
@@ -1,5 +1,11 @@
set(LLVM_TOOLS_DIR ${LLVM_TOOLS_BINARY_DIR})
+if(ARCH STREQUAL amdgcn)
+ set(LIBCLC_CPU "gfx900")
+else()
+ set(LIBCLC_CPU "")
+endif()
+
configure_lit_site_cfg(
${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.py.in
${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg.py
diff --git a/libclc/test/conversion/convert.cl b/libclc/test/conversion/convert.cl
index eb673a493b403..d172caee1f282 100644
--- a/libclc/test/conversion/convert.cl
+++ b/libclc/test/conversion/convert.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
-// REQUIRES: AMDGCN
+// REQUIRES: amdgpu-registered-target
-// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden <4 x i32> @float4_to_int4(
// AMDGCN-SAME: <4 x float> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
diff --git a/libclc/test/geometric/cross.cl b/libclc/test/geometric/cross.cl
index 79724183760e7..c72930638eb74 100644
--- a/libclc/test/geometric/cross.cl
+++ b/libclc/test/geometric/cross.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
-// REQUIRES: AMDGCN
+// REQUIRES: amdgpu-registered-target
-// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden <4 x float> @test_float4(
// AMDGCN-SAME: <4 x float> noundef [[X:%.*]], <4 x float> noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
diff --git a/libclc/test/integer/add_sat.cl b/libclc/test/integer/add_sat.cl
index e223b28444835..2958eed22a47d 100644
--- a/libclc/test/integer/add_sat.cl
+++ b/libclc/test/integer/add_sat.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
-// REQUIRES: AMDGCN
+// REQUIRES: amdgpu-registered-target
-// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden noundef signext i8 @test_char(
// AMDGCN-SAME: i8 noundef signext [[X:%.*]], i8 noundef signext [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
diff --git a/libclc/test/integer/sub_sat.cl b/libclc/test/integer/sub_sat.cl
index 6ba40f75870db..713b6fce98e1f 100644
--- a/libclc/test/integer/sub_sat.cl
+++ b/libclc/test/integer/sub_sat.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
-// REQUIRES: AMDGCN
+// REQUIRES: amdgpu-registered-target
-// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden noundef signext i8 @test_char(
// AMDGCN-SAME: i8 noundef signext [[X:%.*]], i8 noundef signext [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
diff --git a/libclc/test/lit.cfg.py b/libclc/test/lit.cfg.py
index b447d20a30c77..2a3abcc385bd0 100644
--- a/libclc/test/lit.cfg.py
+++ b/libclc/test/lit.cfg.py
@@ -31,8 +31,21 @@
config.target_triple = config.libclc_target
-arch = config.libclc_arch.upper()
-config.available_features.add(arch)
+supported_test_architectures = ["amdgcn", "amdgpu"]
+
+def calculate_arch_features(arch_string):
+ features = []
+ for arch in arch_string.split():
+ if (
+ arch.lower() in supported_test_architectures
+ and config.libclc_arch.lower() in supported_test_architectures
+ ):
+ features.append(arch.lower() + "-registered-target")
+ return features
+
+llvm_config.feature_config(
+ [("--targets-built", calculate_arch_features)]
+)
llvm_config.use_default_substitutions()
@@ -41,9 +54,10 @@
llvm_config.add_tool_substitutions(["llvm-nm"], config.llvm_tools_dir)
config.substitutions.extend([
- ("%libclc_library_dir", config.libclc_library_dir),
- ("%libclc_target", config.libclc_target),
- ("%check_prefix", arch)
+ ("%library_dir", config.libclc_library_dir),
+ ("%target", config.libclc_target),
+ ("%cpu", config.libclc_cpu),
+ ("%check_prefix", config.libclc_arch.upper())
])
# Propagate PATH from environment
diff --git a/libclc/test/lit.site.cfg.py.in b/libclc/test/lit.site.cfg.py.in
index d24227d850e6f..e090e2e16f361 100644
--- a/libclc/test/lit.site.cfg.py.in
+++ b/libclc/test/lit.site.cfg.py.in
@@ -5,6 +5,7 @@ import sys
config.host_triple = "@LLVM_HOST_TRIPLE@"
config.llvm_tools_dir = lit_config.substitute(path(r"@LLVM_TOOLS_DIR@"))
config.libclc_arch = "@ARCH@"
+config.libclc_cpu = "@LIBCLC_CPU@"
config.libclc_library_dir = lit_config.substitute(path(r"@LIBCLC_OUTPUT_LIBRARY_DIR@"))
config.libclc_obj_root = lit_config.substitute(path(r"@CMAKE_CURRENT_BINARY_DIR@"))
config.libclc_target = "@LIBCLC_TARGET@"
diff --git a/libclc/test/math/cos.cl b/libclc/test/math/cos.cl
index 40d6825a034a0..69a7b3ccb8e4e 100644
--- a/libclc/test/math/cos.cl
+++ b/libclc/test/math/cos.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
-// REQUIRES: AMDGCN
+// REQUIRES: amdgpu-registered-target
-// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden float @test_float(
// AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
diff --git a/libclc/test/math/fabs.cl b/libclc/test/math/fabs.cl
index 3b6ec0f8b729e..773381b508e04 100644
--- a/libclc/test/math/fabs.cl
+++ b/libclc/test/math/fabs.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
-// REQUIRES: AMDGCN
+// REQUIRES: amdgpu-registered-target
-// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden noundef float @test_float(
// AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
diff --git a/libclc/test/math/rsqrt.cl b/libclc/test/math/rsqrt.cl
index f0059742d18e6..e47bbbdfe9283 100644
--- a/libclc/test/math/rsqrt.cl
+++ b/libclc/test/math/rsqrt.cl
@@ -1,15 +1,15 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
-// REQUIRES: AMDGCN
+// REQUIRES: amdgpu-registered-target
-// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %target -mcpu=%cpu -x cl -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
// AMDGCN-LABEL: define hidden noundef half @test_half(
// AMDGCN-SAME: half noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
// AMDGCN-NEXT: [[ENTRY:.*:]]
-// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract half @llvm.sqrt.f16(half [[X]]), !fpmath [[META11:![0-9]+]]
-// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract half 1.000000e+00, [[TMP0]], !fpmath [[META12:![0-9]+]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract half @llvm.sqrt.f16(half [[X]]), !fpmath [[META12:![0-9]+]]
+// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract half 1.000000e+00, [[TMP0]], !fpmath [[META13:![0-9]+]]
// AMDGCN-NEXT: ret half [[TMP1]]
//
half test_half(half x) {
@@ -19,8 +19,8 @@ half test_half(half x) {
// AMDGCN-LABEL: define hidden noundef float @test_float(
// AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0]] {
// AMDGCN-NEXT: [[ENTRY:.*:]]
-// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.sqrt.f32(float [[X]]), !fpmath [[META13:![0-9]+]]
-// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract float 1.000000e+00, [[TMP0]], !fpmath [[META14:![0-9]+]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.sqrt.f32(float [[X]]), !fpmath [[META14:![0-9]+]]
+// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract float 1.000000e+00, [[TMP0]], !fpmath [[META15:![0-9]+]]
// AMDGCN-NEXT: ret float [[TMP1]]
//
float test_float(float x) {
@@ -38,8 +38,8 @@ double test_double(double x) {
return rsqrt(x);
}
//.
-// AMDGCN: [[META11]] = !{float 1.500000e+00}
-// AMDGCN: [[META12]] = !{float 1.000000e+00}
-// AMDGCN: [[META13]] = !{float 3.000000e+00}
-// AMDGCN: [[META14]] = !{float 2.500000e+00}
+// AMDGCN: [[META12]] = !{float 1.500000e+00}
+// AMDGCN: [[META13]] = !{float 1.000000e+00}
+// AMDGCN: [[META14]] = !{float 3.000000e+00}
+// AMDGCN: [[META15]] = !{float 2.500000e+00}
//.
diff --git a/libclc/test/misc/as_type.cl b/libclc/test/misc/as_type.cl
index 0ba62fefa4d28..22fdac86215e6 100644
--- a/libclc/test/misc/as_type.cl
+++ b/libclc/test/misc/as_type.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
-// REQUIRES: AMDGCN
+// REQUIRES: amdgpu-registered-target
-// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden noundef <4 x i32> @test_float4_as_int4(
// AMDGCN-SAME: <4 x float> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
diff --git a/libclc/test/update_libclc_tests.py b/libclc/test/update_libclc_tests.py
index 40c48b82b8332..d16274542203f 100755
--- a/libclc/test/update_libclc_tests.py
+++ b/libclc/test/update_libclc_tests.py
@@ -6,9 +6,9 @@
prefix. Supported arch values: amdgpu, amdgcn, nvptx64, spirv, spirv64.
The script does 3 things:
-1. Replaces %libclc_target and %check_prefix in the .cl file for the arch.
+1. Replaces %target, %cpu, and %check_prefix in the .cl file for the arch.
2. Runs update_cc_test_checks.py to update CHECK lines.
-3. Reverts the .cl file back to using %libclc_target and %check_prefix.
+3. Reverts the .cl file back to using %target, %cpu, and %check_prefix.
Usage:
@@ -17,10 +17,11 @@
"""
import argparse
+import os
import re
import subprocess
import sys
-from concurrent.futures import ThreadPoolExecutor, as_completed
+from concurrent.futures import ProcessPoolExecutor, as_completed
from pathlib import Path
ARCH_TO_TRIPLE = {
@@ -31,6 +32,22 @@
"spirv64": "spirv64-unknown-mesa3d",
}
+ARCH_TO_CPU = {
+ "amdgpu": "gfx900",
+ "amdgcn": "gfx900",
+ "nvptx64": "",
+ "spirv": "",
+ "spirv64": "",
+}
+
+ARCH_TO_REQUIRES = {
+ "amdgpu": "amdgpu-registered-target",
+ "amdgcn": "amdgpu-registered-target",
+ "nvptx64": "nvptx-registered-target",
+ "spirv": "spirv-registered-target",
+ "spirv64": "spirv-registered-target",
+}
+
SCRIPT_DIR = Path(__file__).parent
REPO_ROOT = SCRIPT_DIR.parent.parent
UPDATE_SCRIPT = REPO_ROOT / "llvm" / "utils" / "update_cc_test_checks.py"
@@ -41,19 +58,25 @@ def find_cl_files(test_dir: Path):
return list(test_dir.rglob("*.cl"))
-def replace_in_file(path: Path, triple: str, check_prefix: str):
+def replace_in_file(path: Path, triple: str, cpu: str, check_prefix: str):
content = path.read_bytes()
- content = content.replace(b"%libclc_target", triple.encode())
+ content = content.replace(b"%target", triple.encode())
+ if cpu:
+ content = content.replace(b"%cpu", cpu.encode())
content = content.replace(b"%check_prefix", check_prefix.encode())
path.write_bytes(content)
-def revert_in_file(path: Path, triple: str, check_prefix: str):
+def revert_in_file(path: Path, triple: str, cpu: str, check_prefix: str):
# Only revert in the RUN line context, not in generated CHECK lines.
content = path.read_bytes()
content = content.replace(
- f"-target {triple}".encode(), b"-target %libclc_target"
+ f"-target {triple}".encode(), b"-target %target"
)
+ if cpu:
+ content = content.replace(
+ f"-mcpu={cpu}".encode(), b"-mcpu=%cpu"
+ )
content = content.replace(
f"--check-prefix={check_prefix}".encode(), b"--check-prefix=%check_prefix"
)
@@ -75,8 +98,8 @@ def file_requires_feature(path: Path, feature: str) -> bool:
return False
-def process_file(cl_file: Path, triple: str, check_prefix: str) -> bool:
- replace_in_file(cl_file, triple, check_prefix)
+def process_file(cl_file: Path, triple: str, cpu: str, check_prefix: str) -> bool:
+ replace_in_file(cl_file, triple, cpu, check_prefix)
cmd = [
sys.executable,
str(UPDATE_SCRIPT),
@@ -88,7 +111,7 @@ def process_file(cl_file: Path, triple: str, check_prefix: str) -> bool:
ok = result.returncode == 0
if not ok:
print(f" FAILED: {result.stderr.strip()}", file=sys.stderr)
- revert_in_file(cl_file, triple, check_prefix)
+ revert_in_file(cl_file, triple, cpu, check_prefix)
return ok
@@ -105,6 +128,7 @@ def main():
arch = args.arch.lower()
triple = ARCH_TO_TRIPLE[arch]
+ cpu = ARCH_TO_CPU[arch]
# check_prefix matches REQUIRES feature: uppercase of canonical arch name
# amdgpu -> AMDGCN (same triple as amdgcn), others uppercased
if arch == "amdgpu":
@@ -112,21 +136,22 @@ def main():
else:
check_prefix = arch.upper()
+ requires_feature = ARCH_TO_REQUIRES[arch]
cl_files = find_cl_files(SCRIPT_DIR)
- # Only process files that REQUIRES this feature
- target_files = [f for f in cl_files if file_requires_feature(f, check_prefix)]
+ target_files = [f for f in cl_files if file_requires_feature(f, requires_feature)]
if not target_files:
- print(f"No .cl files found with REQUIRES: {check_prefix}")
+ print(f"No .cl files found with REQUIRES: {requires_feature}")
return
- print(f"arch={arch} triple={triple} check_prefix={check_prefix}")
+ print(f"arch={arch} triple={triple} cpu={cpu} check_prefix={check_prefix} requires={requires_feature}")
print(f"Processing {len(target_files)} file(s)...")
failed = []
- with ThreadPoolExecutor() as executor:
+ num_workers = max(1, os.cpu_count() // 2)
+ with ProcessPoolExecutor(max_workers=num_workers) as executor:
futures = {
- executor.submit(process_file, f, triple, check_prefix): f
+ executor.submit(process_file, f, triple, cpu, check_prefix): f
for f in target_files
}
for future in as_completed(futures):
diff --git a/libclc/test/work-item/get_group_id.cl b/libclc/test/work-item/get_group_id.cl
index cc26101c1b99f..fc3b97be37bb6 100644
--- a/libclc/test/work-item/get_group_id.cl
+++ b/libclc/test/work-item/get_group_id.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
-// REQUIRES: AMDGCN
+// REQUIRES: amdgpu-registered-target
-// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden range(i64 0, 4294967296) i64 @test(
// AMDGCN-SAME: ) local_unnamed_addr #[[ATTR0:[0-9]+]] {
>From 716a595408f6dd27a983730bf24eb3c005f8af1f Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Wed, 13 May 2026 06:37:37 +0200
Subject: [PATCH 5/7] clang-format
---
libclc/test/lit.cfg.py | 20 +++++++------
libclc/test/update_libclc_tests.py | 47 +++++++++++++++---------------
2 files changed, 34 insertions(+), 33 deletions(-)
diff --git a/libclc/test/lit.cfg.py b/libclc/test/lit.cfg.py
index 2a3abcc385bd0..f55bcc1c42cc3 100644
--- a/libclc/test/lit.cfg.py
+++ b/libclc/test/lit.cfg.py
@@ -33,6 +33,7 @@
supported_test_architectures = ["amdgcn", "amdgpu"]
+
def calculate_arch_features(arch_string):
features = []
for arch in arch_string.split():
@@ -43,9 +44,8 @@ def calculate_arch_features(arch_string):
features.append(arch.lower() + "-registered-target")
return features
-llvm_config.feature_config(
- [("--targets-built", calculate_arch_features)]
-)
+
+llvm_config.feature_config([("--targets-built", calculate_arch_features)])
llvm_config.use_default_substitutions()
@@ -53,12 +53,14 @@ def calculate_arch_features(arch_string):
llvm_config.add_tool_substitutions(["llvm-nm"], config.llvm_tools_dir)
-config.substitutions.extend([
- ("%library_dir", config.libclc_library_dir),
- ("%target", config.libclc_target),
- ("%cpu", config.libclc_cpu),
- ("%check_prefix", config.libclc_arch.upper())
-])
+config.substitutions.extend(
+ [
+ ("%library_dir", config.libclc_library_dir),
+ ("%target", config.libclc_target),
+ ("%cpu", config.libclc_cpu),
+ ("%check_prefix", config.libclc_arch.upper()),
+ ]
+)
# Propagate PATH from environment
if "PATH" in os.environ:
diff --git a/libclc/test/update_libclc_tests.py b/libclc/test/update_libclc_tests.py
index d16274542203f..e3de7f33e526c 100755
--- a/libclc/test/update_libclc_tests.py
+++ b/libclc/test/update_libclc_tests.py
@@ -25,27 +25,27 @@
from pathlib import Path
ARCH_TO_TRIPLE = {
- "amdgpu": "amdgcn-amd-amdhsa-llvm",
- "amdgcn": "amdgcn-amd-amdhsa-llvm",
- "nvptx64": "nvptx64-nvidia-cuda",
- "spirv": "spirv-unknown-mesa3d",
- "spirv64": "spirv64-unknown-mesa3d",
+ "amdgpu": "amdgcn-amd-amdhsa-llvm",
+ "amdgcn": "amdgcn-amd-amdhsa-llvm",
+ "nvptx64": "nvptx64-nvidia-cuda",
+ "spirv": "spirv-unknown-mesa3d",
+ "spirv64": "spirv64-unknown-mesa3d",
}
ARCH_TO_CPU = {
- "amdgpu": "gfx900",
- "amdgcn": "gfx900",
- "nvptx64": "",
- "spirv": "",
- "spirv64": "",
+ "amdgpu": "gfx900",
+ "amdgcn": "gfx900",
+ "nvptx64": "",
+ "spirv": "",
+ "spirv64": "",
}
ARCH_TO_REQUIRES = {
- "amdgpu": "amdgpu-registered-target",
- "amdgcn": "amdgpu-registered-target",
- "nvptx64": "nvptx-registered-target",
- "spirv": "spirv-registered-target",
- "spirv64": "spirv-registered-target",
+ "amdgpu": "amdgpu-registered-target",
+ "amdgcn": "amdgpu-registered-target",
+ "nvptx64": "nvptx-registered-target",
+ "spirv": "spirv-registered-target",
+ "spirv64": "spirv-registered-target",
}
SCRIPT_DIR = Path(__file__).parent
@@ -70,13 +70,9 @@ def replace_in_file(path: Path, triple: str, cpu: str, check_prefix: str):
def revert_in_file(path: Path, triple: str, cpu: str, check_prefix: str):
# Only revert in the RUN line context, not in generated CHECK lines.
content = path.read_bytes()
- content = content.replace(
- f"-target {triple}".encode(), b"-target %target"
- )
+ content = content.replace(f"-target {triple}".encode(), b"-target %target")
if cpu:
- content = content.replace(
- f"-mcpu={cpu}".encode(), b"-mcpu=%cpu"
- )
+ content = content.replace(f"-mcpu={cpu}".encode(), b"-mcpu=%cpu")
content = content.replace(
f"--check-prefix={check_prefix}".encode(), b"--check-prefix=%check_prefix"
)
@@ -91,7 +87,7 @@ def file_requires_feature(path: Path, feature: str) -> bool:
for line in text.splitlines():
stripped = line.strip().lstrip("//").strip()
if stripped.startswith("REQUIRES:"):
- rest = stripped[len("REQUIRES:"):]
+ rest = stripped[len("REQUIRES:") :]
features = [f.strip() for f in re.split(r",|\|\|", rest)]
if feature in features:
return True
@@ -103,7 +99,8 @@ def process_file(cl_file: Path, triple: str, cpu: str, check_prefix: str) -> boo
cmd = [
sys.executable,
str(UPDATE_SCRIPT),
- "--clang", str(CLANG),
+ "--clang",
+ str(CLANG),
str(cl_file),
]
print(f" update: {cl_file.relative_to(REPO_ROOT)}")
@@ -144,7 +141,9 @@ def main():
print(f"No .cl files found with REQUIRES: {requires_feature}")
return
- print(f"arch={arch} triple={triple} cpu={cpu} check_prefix={check_prefix} requires={requires_feature}")
+ print(
+ f"arch={arch} triple={triple} cpu={cpu} check_prefix={check_prefix} requires={requires_feature}"
+ )
print(f"Processing {len(target_files)} file(s)...")
failed = []
>From 55fed6935cc1c778642f4aacd5e93f2c59be5d8b Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Wed, 13 May 2026 11:51:31 +0200
Subject: [PATCH 6/7] check AMDGPU in config.root.targets
---
libclc/test/AMDGPU/lit.local.cfg | 2 +-
libclc/test/CMakeLists.txt | 6 +++---
libclc/test/lit.cfg.py | 16 +++++++++++-----
libclc/test/lit.site.cfg.py.in | 4 ++--
libclc/test/math/fabs.cl | 2 +-
libclc/test/math/rsqrt.cl | 16 ++++++++--------
6 files changed, 26 insertions(+), 20 deletions(-)
diff --git a/libclc/test/AMDGPU/lit.local.cfg b/libclc/test/AMDGPU/lit.local.cfg
index 68d2d4d5d0023..7c492428aec76 100644
--- a/libclc/test/AMDGPU/lit.local.cfg
+++ b/libclc/test/AMDGPU/lit.local.cfg
@@ -1,2 +1,2 @@
-if "amdgpu-registered-target" not in config.available_features:
+if not "AMDGPU" in config.root.targets:
config.unsupported = True
diff --git a/libclc/test/CMakeLists.txt b/libclc/test/CMakeLists.txt
index f1b119df09f1b..0cbbc2838d624 100644
--- a/libclc/test/CMakeLists.txt
+++ b/libclc/test/CMakeLists.txt
@@ -1,9 +1,9 @@
set(LLVM_TOOLS_DIR ${LLVM_TOOLS_BINARY_DIR})
-if(ARCH STREQUAL amdgcn)
- set(LIBCLC_CPU "gfx900")
+if(LIBCLC_TARGET_ARCH IN_LIST LIBCLC_ARCHS_AMDGPU)
+ set(LIBCLC_TARGET_CPU "gfx900")
else()
- set(LIBCLC_CPU "")
+ set(LIBCLC_TARGET_CPU "")
endif()
configure_lit_site_cfg(
diff --git a/libclc/test/lit.cfg.py b/libclc/test/lit.cfg.py
index f55bcc1c42cc3..22715411ea9d6 100644
--- a/libclc/test/lit.cfg.py
+++ b/libclc/test/lit.cfg.py
@@ -19,8 +19,11 @@
# suffixes: A list of file extensions to treat as test files.
config.suffixes = [".cl", ".test"]
-# Exclude certain directories from test discovery
-config.excludes = ["CMakeLists.txt"]
+# Exclude certain directories and files from test discovery
+config.excludes = [
+ "CMakeLists.txt",
+ "update_libclc_tests.py",
+]
# test_source_root: The root path where tests are located.
# For per-target tests, this is the target's test directory.
@@ -33,15 +36,18 @@
supported_test_architectures = ["amdgcn", "amdgpu"]
+config.targets = set()
+
def calculate_arch_features(arch_string):
features = []
for arch in arch_string.split():
if (
arch.lower() in supported_test_architectures
- and config.libclc_arch.lower() in supported_test_architectures
+ and config.libclc_target_arch.lower() in supported_test_architectures
):
features.append(arch.lower() + "-registered-target")
+ config.targets.add(arch.upper())
return features
@@ -57,8 +63,8 @@ def calculate_arch_features(arch_string):
[
("%library_dir", config.libclc_library_dir),
("%target", config.libclc_target),
- ("%cpu", config.libclc_cpu),
- ("%check_prefix", config.libclc_arch.upper()),
+ ("%cpu", config.libclc_target_cpu),
+ ("%check_prefix", config.libclc_target_arch.upper()),
]
)
diff --git a/libclc/test/lit.site.cfg.py.in b/libclc/test/lit.site.cfg.py.in
index e090e2e16f361..3c9acfc0551f0 100644
--- a/libclc/test/lit.site.cfg.py.in
+++ b/libclc/test/lit.site.cfg.py.in
@@ -4,11 +4,11 @@ import sys
config.host_triple = "@LLVM_HOST_TRIPLE@"
config.llvm_tools_dir = lit_config.substitute(path(r"@LLVM_TOOLS_DIR@"))
-config.libclc_arch = "@ARCH@"
-config.libclc_cpu = "@LIBCLC_CPU@"
config.libclc_library_dir = lit_config.substitute(path(r"@LIBCLC_OUTPUT_LIBRARY_DIR@"))
config.libclc_obj_root = lit_config.substitute(path(r"@CMAKE_CURRENT_BINARY_DIR@"))
config.libclc_target = "@LIBCLC_TARGET@"
+config.libclc_target_arch = "@LIBCLC_TARGET_ARCH@"
+config.libclc_target_cpu = "@LIBCLC_TARGET_CPU@"
import lit.llvm
lit.llvm.initialize(lit_config, config)
diff --git a/libclc/test/math/fabs.cl b/libclc/test/math/fabs.cl
index 773381b508e04..345f2b25b040b 100644
--- a/libclc/test/math/fabs.cl
+++ b/libclc/test/math/fabs.cl
@@ -6,7 +6,7 @@
// AMDGCN-LABEL: define hidden noundef float @test_float(
// AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
// AMDGCN-NEXT: [[ENTRY:.*:]]
-// AMDGCN-NEXT: [[CALL:%.*]] = tail call fastcc float @llvm.fabs.f32(float noundef [[X]]) #[[ATTR2:[0-9]+]]
+// AMDGCN-NEXT: [[CALL:%.*]] = tail call float @llvm.fabs.f32(float noundef [[X]]) #[[ATTR2:[0-9]+]]
// AMDGCN-NEXT: ret float [[CALL]]
//
float test_float(float x) {
diff --git a/libclc/test/math/rsqrt.cl b/libclc/test/math/rsqrt.cl
index e47bbbdfe9283..a40f7defdb607 100644
--- a/libclc/test/math/rsqrt.cl
+++ b/libclc/test/math/rsqrt.cl
@@ -8,8 +8,8 @@
// AMDGCN-LABEL: define hidden noundef half @test_half(
// AMDGCN-SAME: half noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
// AMDGCN-NEXT: [[ENTRY:.*:]]
-// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract half @llvm.sqrt.f16(half [[X]]), !fpmath [[META12:![0-9]+]]
-// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract half 1.000000e+00, [[TMP0]], !fpmath [[META13:![0-9]+]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract half @llvm.sqrt.f16(half [[X]]), !fpmath [[META11:![0-9]+]]
+// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract half 1.000000e+00, [[TMP0]], !fpmath [[META12:![0-9]+]]
// AMDGCN-NEXT: ret half [[TMP1]]
//
half test_half(half x) {
@@ -19,8 +19,8 @@ half test_half(half x) {
// AMDGCN-LABEL: define hidden noundef float @test_float(
// AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0]] {
// AMDGCN-NEXT: [[ENTRY:.*:]]
-// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.sqrt.f32(float [[X]]), !fpmath [[META14:![0-9]+]]
-// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract float 1.000000e+00, [[TMP0]], !fpmath [[META15:![0-9]+]]
+// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.sqrt.f32(float [[X]]), !fpmath [[META13:![0-9]+]]
+// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract float 1.000000e+00, [[TMP0]], !fpmath [[META14:![0-9]+]]
// AMDGCN-NEXT: ret float [[TMP1]]
//
float test_float(float x) {
@@ -38,8 +38,8 @@ double test_double(double x) {
return rsqrt(x);
}
//.
-// AMDGCN: [[META12]] = !{float 1.500000e+00}
-// AMDGCN: [[META13]] = !{float 1.000000e+00}
-// AMDGCN: [[META14]] = !{float 3.000000e+00}
-// AMDGCN: [[META15]] = !{float 2.500000e+00}
+// AMDGCN: [[META11]] = !{float 1.500000e+00}
+// AMDGCN: [[META12]] = !{float 1.000000e+00}
+// AMDGCN: [[META13]] = !{float 3.000000e+00}
+// AMDGCN: [[META14]] = !{float 2.500000e+00}
//.
>From c76ec4e065a911694b86a0a4625f2355018c864c Mon Sep 17 00:00:00 2001
From: Wenju He <wenju.he at intel.com>
Date: Wed, 13 May 2026 18:00:22 +0800
Subject: [PATCH 7/7] Update libclc/test/conversion/convert.cl
Co-authored-by: Matt Arsenault <arsenm2 at gmail.com>
---
libclc/test/conversion/convert.cl | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/libclc/test/conversion/convert.cl b/libclc/test/conversion/convert.cl
index d172caee1f282..2e2a79e4ebb1e 100644
--- a/libclc/test/conversion/convert.cl
+++ b/libclc/test/conversion/convert.cl
@@ -1,7 +1,7 @@
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
// REQUIRES: amdgpu-registered-target
-// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
+// RUN: %clang --target=%target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix
// AMDGCN-LABEL: define hidden <4 x i32> @float4_to_int4(
// AMDGCN-SAME: <4 x float> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {
More information about the cfe-commits
mailing list