Skip to content

[libclc] Add initial LIT tests #87989

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 1 commit into
base: main
Choose a base branch
from
Open

[libclc] Add initial LIT tests #87989

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
16 changes: 13 additions & 3 deletions libclc/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -366,7 +366,10 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
COMMAND ${LLVM_SPIRV} ${spvflags} -o ${spv_suffix} ${builtins_link_lib}
DEPENDS ${builtins_link_lib}
)
add_custom_target( "prepare-${spv_suffix}" ALL DEPENDS "${spv_suffix}" )
add_custom_target( prepare-${arch_suffix} ALL DEPENDS ${spv_suffix} )
set_target_properties( prepare-${arch_suffix}
PROPERTIES TARGET_FILE ${spv_suffix}
)
install( FILES ${CMAKE_CURRENT_BINARY_DIR}/${spv_suffix}
DESTINATION "${CMAKE_INSTALL_DATADIR}/clc" )
else()
Expand All @@ -392,7 +395,10 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
add_custom_command( OUTPUT ${obj_suffix}
COMMAND prepare_builtins -o ${obj_suffix} ${builtins_opt_lib}
DEPENDS ${builtins_opt_lib} prepare_builtins )
add_custom_target( prepare-${obj_suffix} ALL DEPENDS ${obj_suffix} )
add_custom_target( prepare-${arch_suffix} ALL DEPENDS ${obj_suffix} )
set_target_properties( prepare-${arch_suffix}
PROPERTIES TARGET_FILE ${obj_suffix}
)

# nvptx-- targets don't include workitem builtins
if( NOT clang_triple MATCHES ".*ptx.*--$" )
Expand All @@ -406,9 +412,13 @@ foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
set( alias_suffix "${a}-${clang_triple}.bc" )
add_custom_target( ${alias_suffix} ALL
COMMAND ${CMAKE_COMMAND} -E create_symlink ${obj_suffix} ${alias_suffix}
DEPENDS prepare-${obj_suffix} )
DEPENDS prepare-${arch_suffix} )
install( FILES ${CMAKE_CURRENT_BINARY_DIR}/${alias_suffix} DESTINATION "${CMAKE_INSTALL_DATADIR}/clc" )
endforeach( a )
endif()
endforeach( d )
endforeach( t )

if( NOT LIBCLC_STANDALONE_BUILD )
add_subdirectory( test )
endif()
46 changes: 46 additions & 0 deletions libclc/test/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,46 @@
set(LIBCLC_TEST_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR})

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
)

list(APPEND LIBCLC_TEST_DEPS
libclc::clang
FileCheck count not
)

add_custom_target(check-libclc)

foreach( t ${LIBCLC_TARGETS_TO_BUILD} )
foreach( d ${${t}_devices} )

# The tests use LLVM IR, so we can't test SPIR-V libraries here.
string( REPLACE "-" ";" TRIPLE ${t} )
list( GET TRIPLE 0 ARCH )
if( ARCH STREQUAL spirv OR ARCH STREQUAL spirv64 )
continue()
endif()

get_libclc_device_info(
TRIPLE ${t}
DEVICE ${d}
CPU cpu
ARCH_SUFFIX arch_suffix
CLANG_TRIPLE clang_triple
)

add_lit_testsuite(check-libclc-${arch_suffix}
"Running libclc ${arch_suffix} regression tests"
${CMAKE_CURRENT_BINARY_DIR}
PARAMS "target=${clang_triple}" cpu=${cpu}
"builtins=$<TARGET_PROPERTY:prepare-${arch_suffix},TARGET_FILE>"
DEPENDS prepare-${arch_suffix} ${LIBCLC_TEST_DEPS}
)

add_dependencies( check-libclc check-libclc-${arch_suffix} )

endforeach()
endforeach()
3 changes: 3 additions & 0 deletions libclc/test/add_sat.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,3 +1,6 @@
// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s

// CHECK: foo
__kernel void foo(__global char *a, __global char *b, __global char *c) {
*a = add_sat(*b, *c);
}
5 changes: 4 additions & 1 deletion libclc/test/as_type.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,3 +1,6 @@
__kernel void foo(int4 *x, float4 *y) {
// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s

// CHECK: foo
__kernel void foo(__global int4 *x, __global float4 *y) {
*x = as_int4(*y);
}
5 changes: 4 additions & 1 deletion libclc/test/convert.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,3 +1,6 @@
__kernel void foo(int4 *x, float4 *y) {
// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s

// CHECK: foo
__kernel void foo(__global int4 *x, __global float4 *y) {
*x = convert_int4(*y);
}
5 changes: 4 additions & 1 deletion libclc/test/cos.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,3 +1,6 @@
__kernel void foo(float4 *f) {
// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s

// CHECK: foo
__kernel void foo(__global float4 *f) {
*f = cos(*f);
}
5 changes: 4 additions & 1 deletion libclc/test/cross.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,3 +1,6 @@
__kernel void foo(float4 *f) {
// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s

// CHECK: foo
__kernel void foo(__global float4 *f) {
*f = cross(f[0], f[1]);
}
5 changes: 4 additions & 1 deletion libclc/test/fabs.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,3 +1,6 @@
__kernel void foo(float *f) {
// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s

// CHECK: foo
__kernel void foo(__global float *f) {
*f = fabs(*f);
}
5 changes: 4 additions & 1 deletion libclc/test/get_group_id.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,3 +1,6 @@
__kernel void foo(int *i) {
// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Contributor

@arsenm arsenm Apr 17, 2024 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This just requires additional REQUIRES support in lit for the built targets when running the test, the same as other built-backend dependent codegen tests. If you're actually updating the tests, you kind of need to have built all the targets

We don't want the tests to fail if the user hasn't built certain targets,

Right, they should be skipped UNSUPPORTED.

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.

This is identical to clang codegen tests. Not every builtin emits identical IR, not every builtin is supported on every target, and some clang tests directly emit ISA and do have REQUIRES directives to only run if the backend built. It's something of an accident that clang can always emit target intrinsics without the corresponding backend being built (which in the past were stalled efforts to fix)

Personally I think the build complexity of libclc having a parallel enabled target mechanism is not worth it, and should just be tied to the built backend support for the top level targets. It doesn't take all that long to build for the handful of triples

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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)

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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


// CHECK: foo
__kernel void foo(__global int *i) {
i[get_group_id(0)] = 1;
}
57 changes: 57 additions & 0 deletions libclc/test/lit.cfg.py
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,57 @@
import os

import lit.formats
import lit.util

from lit.llvm import llvm_config
from lit.llvm.subst import ToolSubst
from lit.llvm.subst import FindTool

# 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(not llvm_config.use_lit_shell)

# suffixes: A list of file extensions to treat as test files.
config.suffixes = [
".cl",
]

# test_source_root: The root path where tests are located.
config.test_source_root = os.path.join(os.path.dirname(__file__))

# test_exec_root: The root path where tests should be run.
config.test_exec_root = os.path.join(config.test_run_dir, "test")

llvm_config.use_default_substitutions()

target = lit_config.params.get("target", "")
builtins = lit_config.params.get("builtins", "")

clang_flags = [
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ideally the clang driver would know to do all this with the provided path

"-fno-builtin",
"-target",
target,
"-Xclang",
"-mlink-builtin-bitcode",
"-Xclang",
os.path.join(config.libclc_lib_dir, builtins),
"-nogpulib",
]

cpu = lit_config.params.get("cpu", "")
if cpu:
clang_flags.append(f"-mcpu={cpu}")

llvm_config.use_clang(additional_flags=clang_flags)

tools = [
"llvm-dis",
"not",
]
tool_dirs = [config.llvm_tools_dir]

llvm_config.add_tool_substitutions(tools, tool_dirs)
24 changes: 24 additions & 0 deletions libclc/test/lit.site.cfg.py.in
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,24 @@
@LIT_SITE_CFG_IN_HEADER@

import sys

config.llvm_src_root = path(r"@LLVM_SOURCE_DIR@")
config.llvm_obj_root = path(r"@LLVM_BINARY_DIR@")
config.llvm_tools_dir = lit_config.substitute(path(r"@LLVM_TOOLS_DIR@"))
config.llvm_libs_dir = lit_config.substitute(path(r"@LLVM_LIBS_DIR@"))
config.llvm_shlib_dir = lit_config.substitute(path(r"@SHLIBDIR@"))
config.lit_tools_dir = path(r"@LLVM_LIT_TOOLS_DIR@")
config.host_triple = "@LLVM_HOST_TRIPLE@"
config.target_triple = "@LLVM_TARGET_TRIPLE@"
config.host_arch = "@HOST_ARCH@"
config.python_executable = "@Python3_EXECUTABLE@"
config.libclc_src_dir = path(r"@LIBCLC_SOURCE_DIR@")
config.libclc_lib_dir = path(r"@LIBCLC_BINARY_DIR@")
config.test_run_dir = path(r"@LIBCLC_BINARY_DIR@")

import lit.llvm
lit.llvm.initialize(lit_config, config)

# Let the main config do the real work.
lit_config.load_config(
config, os.path.join(config.libclc_src_dir, "test/lit.cfg.py"))
8 changes: 6 additions & 2 deletions libclc/test/rsqrt.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,10 @@
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
// RUN: %clang -emit-llvm -S %s

__kernel void foo(float4 *x, double4 *y) {
#if defined(cl_khr_fp64)

__kernel void foo(__global float4 *x, __global double4 *y) {
x[1] = rsqrt(x[0]);
y[1] = rsqrt(y[0]);
}

#endif
11 changes: 7 additions & 4 deletions libclc/test/subsat.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,19 +1,22 @@
__kernel void test_subsat_char(char *a, char x, char y) {
// RUN: %clang -emit-llvm -S -o - %s | FileCheck %s

// CHECK: test_subsat_char
__kernel void test_subsat_char(__global char *a, char x, char y) {
*a = sub_sat(x, y);
return;
}

__kernel void test_subsat_uchar(uchar *a, uchar x, uchar y) {
__kernel void test_subsat_uchar(__global uchar *a, uchar x, uchar y) {
*a = sub_sat(x, y);
return;
}

__kernel void test_subsat_long(long *a, long x, long y) {
__kernel void test_subsat_long(__global long *a, long x, long y) {
*a = sub_sat(x, y);
return;
}

__kernel void test_subsat_ulong(ulong *a, ulong x, ulong y) {
__kernel void test_subsat_ulong(__global ulong *a, ulong x, ulong y) {
*a = sub_sat(x, y);
return;
}