https://github.com/wenju-he created
https://github.com/llvm/llvm-project/pull/197151
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
>From 2d10da58e2ecade7d99ec2900216161b7ed83636 Mon Sep 17 00:00:00 2001
From: Wenju He <[email protected]>
Date: Tue, 12 May 2026 08:44:55 +0200
Subject: [PATCH] [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);
+}
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
