[libclc] Add initial LIT tests #87989
Conversation
|
✅ With the latest revision this PR passed the Python code formatter. |
frasercrmck
force-pushed
the
libclc-tests
branch
2 times, most recently
from
c53dc35 to
9048dbc
Compare
libclc/test/get_group_id.cl
Outdated
| @@ -1,3 +1,6 @@ | |||
| __kernel void foo(int *i) { | |||
| // RUN: %clang -emit-llvm -S -o - %s | FileCheck %s | |||
There was a problem hiding this comment.
Can we use update_cc_test_checks? How will checks for different targets co-exist? These all need some kind of explicit target to avoid host dependence
There was a problem hiding this comment.
Yes, I was hoping to explore in that direction. I've kind of just copied the skeleton of the LIT tests we have downstream, but as you say, this isn't really workable as it assumes all tests will be able to use the same CHECK.
I think, for this to work, we'd have to configure LIT with all of the libclc targets (or, a subset we're interested in), then RUN each test multiple times with a unique FileCheck check prefix for each target. Then when running update_cc_test_checks it'd run the test on all targets.
I think we'd need a macro or substitution to expand a single RUN line to multiple architectures, to keep things maintainable.
We could also allow running the tests on just the one or a subset of architectures locally, using a --param to control the architectures to test.
Does that sound like a good direction? to take this in?
There was a problem hiding this comment.
Ach I misunderstood how update_cc_test_checks works. It doesn't actually go through the regular LIT infrastructure, so we can't do anything involving custom substitutions to configure which target is being run.
If we wanted to be able to use update_cc_test_checks I think we'd need to have every target explicit in every test file, which brings its own inflexibilities and maintenance problems.
There was a problem hiding this comment.
The target absolutely should be explicit in any testing files. Pretending these tests can be generic is going to be an intractable problem. This is no different than any clang codegen test
There was a problem hiding this comment.
I wasn't so much thinking that it would pretend to be generic, but rather that it'd use a custom test format (something like analyzer_test.py) to run the test once for each target we build/test, substituting in the target triple and target-specific check prefix each time.
That way we could test all or a subset of all targets without worrying about the test failing on a target we haven't built.
In that sense these are a little different to clang codegen tests - aren't they? clang can always compile for any of its supported targets. With libclc, however, the user is allowed to build only a handful of the targets, and each of those targets can have drastically different builtin implementations through specialization.
We don't want the tests to fail if the user hasn't built certain targets, so I think we need some kind of configuration that allows tests to be skipped. But update_cc_test_checks doesn't handle features or %if or anything dynamic like that. It's really only suitable for the simplest of clang tests. So if we wanted to have each test file run on multiple targets, I think update_cc_test_checks could only be used if the user built all libclc targets, and each test had one RUN line for each target. I think we have ten targets so it might get a bit unwieldy, is all I'm saying. I do want us to be able to use that script as it's by far the best way of maintaining tests that we've got.
There was a problem hiding this comment.
Fair enough, I hadn't really appreciated that update_cc_test_checks was orthogonal to all of these other LIT concepts like REQUIRES and features more broadly, but that's due to my misunderstanding of how that script worked.
Those clang codegen tests that do have REQUIRES are relegated to their own target-specific source files. What we want for libclc I think is ideally to be testing every libclc builtin with every target. I was hoping we could use a single set of test files to avoid duplication, but that doesn't seem possible.
So each libclc target would have their own copy of each test source file: test/amdgcn/smoothstep.cl, test/nvidia/smoothstep.cl, etc? That's not so bad as long as it's easy enough to bring up a new libclc target (maybe with a generator script, though maybe just copy/pasting a directory of tests, removing the old checks, and running update_cc_test_checks is sufficient)? I suppose the amdgpu tests would be amalgamated and would RUN over each sub-target (r600, amdgcn-mesa3d, etc., as appropriate). Similar for NVIDIA and other groups.
I understand the idea of having libclc targets linked to the backend targets. Some, like clspv/clspv64 and spirv/spirv64, aren't intrinsically tied to any specific backend. clspv targets seem to use the "generic" spir/spir64 triples so could probably be unconditionally enabled, and spirv targets are compiled to those same spir/spir64 triples before being compiled to SPIR-V, which we can't test here. My concern is that this might be a breaking change to some downstream users (we don't really know what they're doing with libclc). Either way I'm not sure this needs to concern this initial round of tests - we can probably stomach the complexity.
There was a problem hiding this comment.
You can have different target coexist in the same test files when appropriate and just multi-list REQUIRES. Clang does this regularly (OpenMP in particular has a lot of multi targeted codegen tests)
There was a problem hiding this comment.
That's true, but I'm still sceptical that we can make the tests simple enough for update_cc_test_checks to be viable. We'd have to embed libclc further into the clang driver for it to work, given that it currently doesn't know where to find libclc libraries for each target, and we can't express the path to libclc libraries in LIT in such a way that update_cc_test_checks would handle it. I'm not convinced clang should necessarily know about libclc. I don't think we should be tying libclc to any specific compilation model like OpenCL C.
There was a problem hiding this comment.
It is essential that clang should know about libclc. libclc exists purely as an extension of the compiler. From the user perspective the opencl builtin library is invisible and always available, which requires clang to handle it without any additional user configuration
Ideally it should universally be available in the clang resource directory
There was a problem hiding this comment.
With #146503, clang does at last know about libclc. I've updated this PR to include that one, just to explore how this might work in CI.
I've updated these tests to use update_cc_test_checks.py, just on the amdgcn-mesa-mesa3d triple for now. There's good and bad to this. I think the good is fairly obvious - thanks for the suggestion - but the downside is that now libclc would ideally be tested every time clang and llvm are modified. I'm not sure whether or not there'll be the appetite for the impact on CI times, let alone on having people have to update the libclc test checks every so often.
We could of course simply not test libclc every time clang/llvm are changed, acknowledging that the test checks will go stale and will need to be periodically updated. I'm not sure of the best approach.
llvmbot
added
clang
|
@llvm/pr-subscribers-clang-driver @llvm/pr-subscribers-backend-amdgpu Author: Fraser Cormack (frasercrmck) ChangesThese tests aren't very meaningful and aren't immune to false positives, 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:
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]
|
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
libclc/test/math/fabs.cl
Outdated
There was a problem hiding this comment.
This should specify the -mcpu, and the driver should figure out the appropriate library. It shouldn't be given an explicit subtarget name
There was a problem hiding this comment.
Agreed. I'd like to do that as a follow-up to the aforementioned clang PR.
To be clear, do you still imagine a user specifying --libclc-lib=amdgcn-mesa-mesa3d and the driver infers the tahiti part from -mcpu? Or does the user just specify --libclc-lib with no value to signify they want libclc, and the driver uses both -target and -mcpu to come to the right library name?
There was a problem hiding this comment.
Doesn't clang manage this now?
There was a problem hiding this comment.
No not yet. I haven't had much time to work on libclc recently.
I started a patch that uses -mcpu to find the libclc lib name (for AMDGPU) but ran into difficulties where libclc has, e.g., tonga.bc but not gfx802.bc. I wasn't sure how best to standardize the naming conventions - whether libclc should provide a consistent alias scheme or whether clang should special-case some names.
Any thoughts?
There was a problem hiding this comment.
probably should switch the libclc build to use the gfxXYZ naming scheme instead of the legacy names
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.
frasercrmck
force-pushed
the
libclc-tests
branch
from
c25375f to
35626cb
Compare
|
|
||
| #pragma OPENCL EXTENSION cl_khr_fp64 : enable | ||
|
|
||
| #if defined(cl_khr_fp64) |
There was a problem hiding this comment.
This is set by the target, which is already hardcode here so this is dead? Also the extension was already used/enabled above?
libclc/test/math/fabs.cl
Outdated
There was a problem hiding this comment.
Doesn't clang manage this now?
| __kernel void foo(__global float4 *x, __global double4 *y) { | ||
| x[1] = rsqrt(x[0]); | ||
| y[1] = rsqrt(y[0]); | ||
| } |
There was a problem hiding this comment.
Better to test one case per function, with a meaningful function name.
Plus I think it would be more useful to test the scalar cases for half, float and double. Also plus to not put it in a kernel
|
|
||
| #endif | ||
| //. | ||
| // CHECK: [[META6]] = !{i32 1, i32 1} |
There was a problem hiding this comment.
Is there a way to only check the relevant metadata? i.e. the !fpmath cases are useful but the other stuff isn't really
| // | ||
| //===----------------------------------------------------------------------===// | ||
|
|
||
| // RUN: %clang -target amdgcn-mesa-mesa3d -O1 -emit-llvm -S -o - --libclc-lib=tahiti-amdgcn-mesa-mesa3d -fno-builtin --no-offloadlib %s | FileCheck %s |
There was a problem hiding this comment.
Pretty sure you shouldn't need -fno-builtin or --no-offloadlib. Particularly -fno-builtin. OpenCL isn't really an offload language, but did that end up as the alias for nogpulib?
There was a problem hiding this comment.
Yes, no-offloadlib is the alias for nogpulib. I'll try without them but yeah if we're generating LLVM IR then it's not going to try and link in the libdevice, is it?
There was a problem hiding this comment.
It will be linked in, that's the point of the test? And yes, it will always be linked especially when targeting IR
3 participants
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.