[clang] [libclc] [llvm] [libclc] Add initial LIT tests (PR #87989)
via llvm-commits
llvm-commits at lists.llvm.org
Mon Jul 7 05:02:34 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Fraser Cormack (frasercrmck)
<details>
<summary>Changes</summary>
These tests aren't very meaningful and aren't immune to false positives,
but they do get the project building when running 'check-all' and so
enable libclc testing in CI.
---
Patch is 69.21 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/87989.diff
31 Files Affected:
- (modified) .ci/compute_projects.py (+4-2)
- (modified) clang/include/clang/Basic/DiagnosticDriverKinds.td (+2)
- (modified) clang/include/clang/Driver/CommonArgs.h (+3)
- (modified) clang/include/clang/Driver/Options.td (+2)
- (modified) clang/lib/Driver/ToolChains/AMDGPU.cpp (+2)
- (modified) clang/lib/Driver/ToolChains/CommonArgs.cpp (+37)
- (added) clang/test/Driver/Inputs/libclc/libclc.bc ()
- (added) clang/test/Driver/Inputs/libclc/subdir/libclc.bc ()
- (added) clang/test/Driver/opencl-libclc.cl (+9)
- (modified) libclc/CMakeLists.txt (+14-3)
- (added) libclc/test/CMakeLists.txt (+35)
- (removed) libclc/test/add_sat.cl (-11)
- (removed) libclc/test/as_type.cl (-11)
- (removed) libclc/test/convert.cl (-11)
- (removed) libclc/test/cos.cl (-11)
- (removed) libclc/test/cross.cl (-11)
- (removed) libclc/test/fabs.cl (-11)
- (added) libclc/test/geometric/cross.cl (+51)
- (removed) libclc/test/get_group_id.cl (-11)
- (added) libclc/test/integer/add_sat.cl (+32)
- (added) libclc/test/integer/sub_sat.cl (+72)
- (added) libclc/test/lit.cfg.py (+44)
- (added) libclc/test/lit.site.cfg.py.in (+23)
- (added) libclc/test/math/cos.cl (+296)
- (added) libclc/test/math/fabs.cl (+32)
- (added) libclc/test/math/rsqrt.cl (+48)
- (added) libclc/test/misc/as_type.cl (+31)
- (added) libclc/test/misc/convert.cl (+32)
- (removed) libclc/test/rsqrt.cl (-14)
- (removed) libclc/test/subsat.cl (-27)
- (added) libclc/test/work-item/get_group_id.cl (+33)
``````````diff
diff --git a/.ci/compute_projects.py b/.ci/compute_projects.py
index c3cf714ce6c10..4c268b7221663 100644
--- a/.ci/compute_projects.py
+++ b/.ci/compute_projects.py
@@ -46,6 +46,7 @@
"mlir",
"polly",
"flang",
+ "libclc",
},
"lld": {"bolt", "cross-project-tests"},
# TODO(issues/132795): LLDB should be enabled on clang changes.
@@ -75,7 +76,7 @@
# This mapping describes runtimes that should be tested when the key project is
# touched.
DEPENDENT_RUNTIMES_TO_TEST = {
- "clang": {"compiler-rt"},
+ "clang": {"compiler-rt", "libclc"},
"clang-tools-extra": {"libc"},
"libc": {"libc"},
".ci": {"compiler-rt", "libc"},
@@ -132,6 +133,7 @@
"lld": "check-lld",
"flang": "check-flang",
"libc": "check-libc",
+ "libclc": "check-libclc",
"lld": "check-lld",
"lldb": "check-lldb",
"mlir": "check-mlir",
@@ -139,7 +141,7 @@
"polly": "check-polly",
}
-RUNTIMES = {"libcxx", "libcxxabi", "libunwind", "compiler-rt", "libc"}
+RUNTIMES = {"libcxx", "libcxxabi", "libunwind", "compiler-rt", "libc", "libclc"}
def _add_dependencies(projects: Set[str], runtimes: Set[str]) -> Set[str]:
diff --git a/clang/include/clang/Basic/DiagnosticDriverKinds.td b/clang/include/clang/Basic/DiagnosticDriverKinds.td
index 34b6c0d7a8acd..8d07ade73ec89 100644
--- a/clang/include/clang/Basic/DiagnosticDriverKinds.td
+++ b/clang/include/clang/Basic/DiagnosticDriverKinds.td
@@ -392,6 +392,8 @@ def warn_drv_fraw_string_literals_in_cxx11 : Warning<
"ignoring '-f%select{no-|}0raw-string-literals', which is only valid for C and C++ standards before C++11">,
InGroup<UnusedCommandLineArgument>;
+def err_drv_libclc_not_found : Error<"no libclc library '%0' found in the clang resource directory">;
+
def err_drv_invalid_malign_branch_EQ : Error<
"invalid argument '%0' to -malign-branch=; each element must be one of: %1">;
diff --git a/clang/include/clang/Driver/CommonArgs.h b/clang/include/clang/Driver/CommonArgs.h
index 26aa3ccf84786..7e8ab82eb7863 100644
--- a/clang/include/clang/Driver/CommonArgs.h
+++ b/clang/include/clang/Driver/CommonArgs.h
@@ -215,6 +215,9 @@ void addOpenMPDeviceRTL(const Driver &D, const llvm::opt::ArgList &DriverArgs,
StringRef BitcodeSuffix, const llvm::Triple &Triple,
const ToolChain &HostTC);
+void addOpenCLBuiltinsLib(const Driver &D, const llvm::opt::ArgList &DriverArgs,
+ llvm::opt::ArgStringList &CC1Args);
+
void addOutlineAtomicsArgs(const Driver &D, const ToolChain &TC,
const llvm::opt::ArgList &Args,
llvm::opt::ArgStringList &CmdArgs,
diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td
index 54c71b066f9d4..cf24bcac9c07c 100644
--- a/clang/include/clang/Driver/Options.td
+++ b/clang/include/clang/Driver/Options.td
@@ -1425,6 +1425,8 @@ def openacc_macro_override_EQ
// End Clang specific/exclusive options for OpenACC.
+def libclc_lib_EQ : Joined<["--"], "libclc-lib=">, Group<opencl_Group>,
+ HelpText<"Namespec of libclc OpenCL bitcode library to link">;
def libomptarget_amdgpu_bc_path_EQ : Joined<["--"], "libomptarget-amdgpu-bc-path=">, Group<i_Group>,
HelpText<"Path to libomptarget-amdgcn bitcode library">;
def libomptarget_amdgcn_bc_path_EQ : Joined<["--"], "libomptarget-amdgcn-bc-path=">, Group<i_Group>,
diff --git a/clang/lib/Driver/ToolChains/AMDGPU.cpp b/clang/lib/Driver/ToolChains/AMDGPU.cpp
index b7564a0495da8..e6d1baa2a1caa 100644
--- a/clang/lib/Driver/ToolChains/AMDGPU.cpp
+++ b/clang/lib/Driver/ToolChains/AMDGPU.cpp
@@ -795,6 +795,8 @@ void AMDGPUToolChain::addClangTargetOptions(
CC1Args.push_back("-fvisibility=hidden");
CC1Args.push_back("-fapply-global-visibility-to-externs");
}
+
+ addOpenCLBuiltinsLib(getDriver(), DriverArgs, CC1Args);
}
void AMDGPUToolChain::addClangWarningOptions(ArgStringList &CC1Args) const {
diff --git a/clang/lib/Driver/ToolChains/CommonArgs.cpp b/clang/lib/Driver/ToolChains/CommonArgs.cpp
index bdd77ac84913c..36f335154e6bc 100644
--- a/clang/lib/Driver/ToolChains/CommonArgs.cpp
+++ b/clang/lib/Driver/ToolChains/CommonArgs.cpp
@@ -2942,6 +2942,43 @@ void tools::addHIPRuntimeLibArgs(const ToolChain &TC, Compilation &C,
}
}
+void tools::addOpenCLBuiltinsLib(const Driver &D,
+ const llvm::opt::ArgList &DriverArgs,
+ llvm::opt::ArgStringList &CC1Args) {
+ // Check whether user specifies a libclc bytecode library
+ const Arg *A = DriverArgs.getLastArg(options::OPT_libclc_lib_EQ);
+ if (!A)
+ return;
+
+ // Find device libraries in <LLVM_DIR>/lib/clang/<ver>/lib/libclc/
+ SmallString<128> LibclcPath(D.ResourceDir);
+ llvm::sys::path::append(LibclcPath, "lib", "libclc");
+
+ // If the namespec is of the form :filename, search for that file.
+ StringRef LibclcNamespec(A->getValue());
+ bool FilenameSearch = LibclcNamespec.consume_front(":");
+ SmallString<128> LibclcTargetFile(LibclcNamespec);
+
+ if (FilenameSearch && llvm::sys::fs::exists(LibclcTargetFile)) {
+ CC1Args.push_back("-mlink-builtin-bitcode");
+ CC1Args.push_back(DriverArgs.MakeArgString(LibclcTargetFile));
+ } else {
+ // Search the library paths for the file
+ if (!FilenameSearch)
+ LibclcTargetFile += ".bc";
+
+ llvm::sys::path::append(LibclcPath, LibclcTargetFile);
+ if (llvm::sys::fs::exists(LibclcPath)) {
+ CC1Args.push_back("-mlink-builtin-bitcode");
+ CC1Args.push_back(DriverArgs.MakeArgString(LibclcPath));
+ } else {
+ // Since the user requested a library, if we haven't one then report an
+ // error.
+ D.Diag(diag::err_drv_libclc_not_found) << LibclcTargetFile;
+ }
+ }
+}
+
void tools::addOutlineAtomicsArgs(const Driver &D, const ToolChain &TC,
const llvm::opt::ArgList &Args,
llvm::opt::ArgStringList &CmdArgs,
diff --git a/clang/test/Driver/Inputs/libclc/libclc.bc b/clang/test/Driver/Inputs/libclc/libclc.bc
new file mode 100644
index 0000000000000..e69de29bb2d1d
diff --git a/clang/test/Driver/Inputs/libclc/subdir/libclc.bc b/clang/test/Driver/Inputs/libclc/subdir/libclc.bc
new file mode 100644
index 0000000000000..e69de29bb2d1d
diff --git a/clang/test/Driver/opencl-libclc.cl b/clang/test/Driver/opencl-libclc.cl
new file mode 100644
index 0000000000000..185690768c75b
--- /dev/null
+++ b/clang/test/Driver/opencl-libclc.cl
@@ -0,0 +1,9 @@
+// RUN: %clang -### -target amdgcn-amd-amdhsa --no-offloadlib --libclc-lib=:%S/Inputs/libclc/libclc.bc %s 2>&1 | FileCheck %s
+// RUN: %clang -### -target amdgcn-amd-amdhsa --no-offloadlib --libclc-lib=:%S/Inputs/libclc/subdir/libclc.bc %s 2>&1 | FileCheck %s --check-prefix CHECK-SUBDIR
+
+// RUN: not %clang -### -target amdgcn-amd-amdhsa --no-offloadlib --libclc-lib=:%S/Inputs/libclc/subdir/not-here.bc %s 2>&1 | FileCheck %s --check-prefix CHECK-ERROR
+
+// CHECK: -mlink-builtin-bitcode{{.*}}Inputs{{/|\\\\}}libclc{{/|\\\\}}libclc.bc
+// CHECK-SUBDIR: -mlink-builtin-bitcode{{.*}}Inputs{{/|\\\\}}libclc{{/|\\\\}}subdir{{/|\\\\}}libclc.bc
+
+// CHECK-ERROR: no libclc library{{.*}}not-here.bc' found in the clang resource directory
diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt
index e2871d1b01a16..673309c88e2cd 100644
--- a/libclc/CMakeLists.txt
+++ b/libclc/CMakeLists.txt
@@ -63,6 +63,9 @@ if( LIBCLC_STANDALONE_BUILD OR CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DI
set( ${tool}_target )
endforeach()
endif()
+
+ # Setup the paths where libclc runtimes should be stored.
+ set( LIBCLC_OUTPUT_LIBRARY_DIR ${CMAKE_CURRENT_BINARY_DIR} )
else()
# In-tree configuration
set( LIBCLC_STANDALONE_BUILD FALSE )
@@ -82,10 +85,14 @@ else()
get_host_tool_path( llvm-link LLVM_LINK llvm-link_exe llvm-link_target )
get_host_tool_path( opt OPT opt_exe opt_target )
endif()
-endif()
-# Setup the paths where libclc runtimes should be stored.
-set( LIBCLC_OUTPUT_LIBRARY_DIR ${CMAKE_CURRENT_BINARY_DIR} )
+ # Setup the paths where libclc runtimes should be stored. By default, in an
+ # in-tree build we place the libraries in clang's resource driectory.
+ get_clang_resource_dir( LIBCLC_OUTPUT_DIR PREFIX ${LLVM_LIBRARY_OUTPUT_INTDIR}/.. )
+
+ # Note we do not adhere to LLVM_ENABLE_PER_TARGET_RUNTIME_DIR.
+ set( LIBCLC_OUTPUT_LIBRARY_DIR ${LIBCLC_OUTPUT_DIR}/lib/libclc )
+endif()
if( EXISTS ${LIBCLC_CUSTOM_LLVM_TOOLS_BINARY_DIR} )
message( WARNING "Using custom LLVM tools to build libclc: "
@@ -487,3 +494,7 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
)
endforeach( d )
endforeach( t )
+
+if( NOT LIBCLC_STANDALONE_BUILD )
+ add_subdirectory( test )
+endif()
diff --git a/libclc/test/CMakeLists.txt b/libclc/test/CMakeLists.txt
new file mode 100644
index 0000000000000..653b39c2821a7
--- /dev/null
+++ b/libclc/test/CMakeLists.txt
@@ -0,0 +1,35 @@
+set( LIBCLC_TEST_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR} )
+
+set( LIBCLC_TEST_TARGETS_ALL
+ amdgcn-mesa-mesa3d
+)
+
+foreach( target IN LISTS LIBCLC_TEST_TARGETS_ALL )
+ # If we haven't built this libclc target, don't build the tests
+ if( NOT TARGET prepare-${target} )
+ message( WARNING "libclc tests require target ${target}. Tests will not be built" )
+ # Add a dummy target
+ add_custom_target( check-libclc )
+ return()
+ endif()
+
+ list( APPEND LIBCLC_TEST_DEPS prepare-${target} )
+endforeach()
+
+list( APPEND LIBCLC_TEST_DEPS
+ ${clang_target}
+ FileCheck
+)
+
+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
+)
+
+add_lit_testsuite( check-libclc
+ "Running libclc regression tests"
+ ${CMAKE_CURRENT_BINARY_DIR}
+ DEPENDS ${LIBCLC_TEST_DEPS}
+)
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/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..4cb8c53bea5ee
--- /dev/null
+++ b/libclc/test/geometric/cross.cl
@@ -0,0 +1,51 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+// RUN: %clang -target amdgcn-mesa-mesa3d -O1 -emit-llvm -S -o - --libclc-lib=tahiti-amdgcn-mesa-mesa3d -fno-builtin --no-offloadlib %s | FileCheck %s
+
+// CHECK-LABEL: define protected amdgpu_kernel void @foo(
+// CHECK-SAME: ptr addrspace(1) noundef align 16 captures(none) [[F:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !kernel_arg_addr_space [[META6:![0-9]+]] !kernel_arg_access_qual [[META7:![0-9]+]] !kernel_arg_type [[META8:![0-9]+]] !kernel_arg_base_type [[META9:![0-9]+]] !kernel_arg_type_qual [[META10:![0-9]+]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[TMP0:%.*]] = load <4 x float>, ptr addrspace(1) [[F]], align 16, !tbaa [[TBAA11:![0-9]+]]
+// CHECK-NEXT: [[ARRAYIDX1_I:%.*]] = getelementptr inbounds nuw i8, ptr addrspace(1) [[F]], i64 16
+// CHECK-NEXT: [[TMP1:%.*]] = load <4 x float>, ptr addrspace(1) [[ARRAYIDX1_I]], align 16, !tbaa [[TBAA11]]
+// CHECK-NEXT: [[TMP2:%.*]] = extractelement <4 x float> [[TMP0]], i64 1
+// CHECK-NEXT: [[TMP3:%.*]] = extractelement <4 x float> [[TMP1]], i64 2
+// CHECK-NEXT: [[TMP4:%.*]] = extractelement <4 x float> [[TMP0]], i64 2
+// CHECK-NEXT: [[TMP5:%.*]] = extractelement <4 x float> [[TMP1]], i64 1
+// CHECK-NEXT: [[TMP6:%.*]] = fneg float [[TMP5]]
+// CHECK-NEXT: [[NEG_I_I:%.*]] = fmul float [[TMP4]], [[TMP6]]
+// CHECK-NEXT: [[TMP7:%.*]] = tail call float @llvm.fmuladd.f32(float [[TMP2]], float [[TMP3]], float [[NEG_I_I]])
+// CHECK-NEXT: [[TMP8:%.*]] = extractelement <4 x float> [[TMP1]], i64 0
+// CHECK-NEXT: [[TMP9:%.*]] = extractelement <4 x float> [[TMP0]], i64 0
+// CHECK-NEXT: [[TMP10:%.*]] = fneg float [[TMP3]]
+// CHECK-NEXT: [[NEG3_I_I:%.*]] = fmul float [[TMP9]], [[TMP10]]
+// CHECK-NEXT: [[TMP11:%.*]] = tail call float @llvm.fmuladd.f32(float [[TMP4]], float [[TMP8]], float [[NEG3_I_I]])
+// CHECK-NEXT: [[TMP12:%.*]] = fneg float [[TMP8]]
+// CHECK-NEXT: [[NEG6_I_I:%.*]] = fmul float [[TMP2]], [[TMP12]]
+// CHECK-NEXT: [[TMP13:%.*]] = tail call float @llvm.fmuladd.f32(float [[TMP9]], float [[TMP5]], float [[NEG6_I_I]])
+// CHECK-NEXT: [[TMP14:%.*]] = insertelement <4 x float> <float poison, float poison, float poison, float 0.000000e+00>, float [[TMP7]], i64 0
+// CHECK-NEXT: [[TMP15:%.*]] = insertelement <4 x float> [[TMP14]], float [[TMP11]], i64 1
+// CHECK-NEXT: [[VECINIT8_I_I:%.*]] = insertelement <4 x float> [[TMP15]], float [[TMP13]], i64 2
+// CHECK-NEXT: store <4 x float> [[VECINIT8_I_I]], ptr addrspace(1) [[F]], align 16, !tbaa [[TBAA11]]
+// CHECK-NEXT: ret void
+//
+__kernel void foo(__global float4 *f) {
+ *f = cross(f[0], f[1]);
+}
+//.
+// CHECK: [[META6]] = !{i32 1}
+// CHECK: [[META7]] = !{!"none"}
+// CHECK: [[META8]] = !{!"float4*"}
+// CHECK: [[META9]] = !{!"float __attribute__((ext_vector_type(4)))*"}
+// CHECK: [[META10]] = !{!""}
+// CHECK: [[TBAA11]] = !{[[META12:![0-9]+]], [[META12]], i64 0}
+// CHECK: [[META12]] = !{!"omnipotent char", [[META13:![0-9]+]], i64 0}
+// CHECK: [[META13]] = !{!"Simple C/C++ TBAA"}
+//.
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..ef5bf77b67d21
--- /dev/null
+++ b/libclc/test/integer/add_sat.cl
@@ -0,0 +1,32 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 5
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+// RUN: %clang -target amdgcn-mesa-mesa3d -O1 -emit-llvm -S -o - --libclc-lib=tahiti-amdgcn-mesa-mesa3d -fno-builtin --no-offloadlib %s | FileCheck %s
+
+// CHECK-LABEL: define protected amdgpu_kernel void @foo(
+// CHECK-SAME: ptr addrspace(1) noundef writeonly align 1 captures(none) initializes((0, 1)) [[A:%.*]], ptr addrspace(1) noundef readonly align 1 captures(none) [[B:%.*]], ptr addrspace(1) noundef readonly align 1 captures(none) [[C:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] !kernel_arg_addr_space [[META6:![0-9]+]] !kernel_arg_access_qual [[META7:![0-9]+]] !kernel_arg_type [[META8:![0-9]+]] !kernel_arg_base_type [[META8]] !kernel_arg_type_qual [[META9:![0-9]+]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[TMP0:%.*]] = load i8, ptr addrspace(1) [[B]], align 1, !tbaa [[TBAA10:![0-9]+]]
+// CHECK-NEXT: [[TMP1:%.*]] = load i8, ptr addrspace(1) [[C]], align 1, !tbaa [[TBAA10]]
+// CHECK-NEXT: [[ELT_SAT_I_I:%.*]] = tail call noundef i8 @llvm.sadd.sat.i8(i8 [[TMP0]], i8 [[TMP1]])
+// CHECK-NEXT: store i8 [[ELT_SAT_I_I]], ptr addrspace(1) [[A]], align 1, !tbaa [[TBAA10]]
+// CHECK-NEXT: ret void
+//
+__kernel void foo(__global char *a, __global char *b, __global char *c) {
+ *a = add_sat(*b, *c);
+}
+//.
+// CHECK: [[META6]] = !{i32 1, i32 1, i32 1}
+// CHECK: [[META7]] = !{!"none", !"none", !"none"}
+// CHECK: [[META8]] = !{!"char*", !"char*", !"char*"}
+// CHECK: [[META9]] = !{!"", !"", !""}
+// CHECK: [[TBAA10]] = !{[[META11:![0-9]+]], [[META11]], i64 0}
+// CHECK: [[META11]] = !{!"omnipotent char", [[META12:![0-9]+]], i64 0}
+// CHECK: [[META12]] = !{!"Simple C/C++ TBAA"}
+//.
diff --git a/libclc/test/integer/sub_sat.cl b/libclc/test/integer/sub_sat.cl
new file mode 100644
index 0000000000000..7c3f0a3aa306f
--- /dev/null
+++ b/libclc/test/integer/sub_sat.cl
@...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/87989
More information about the llvm-commits
mailing list