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 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits