https://github.com/wenju-he updated https://github.com/llvm/llvm-project/pull/197151
>From 2d10da58e2ecade7d99ec2900216161b7ed83636 Mon Sep 17 00:00:00 2001 From: Wenju He <[email protected]> Date: Tue, 12 May 2026 08:44:55 +0200 Subject: [PATCH 1/5] [libclc] Move external-funcs.test to static file and use IR checks for .cl tests This PR supercedes #87989. Moving external-funcs.test to static file simplifies test/CMakeLists.txt. Static files follows llvm standard lit pattern and enables fine-grained testing of missing symbols in specific libraries. .cl test files uses %libclc_target and %check_prefix, which are replaced with specific values during `ninja check-libclc` or `llvm-lit build/runtimes/runtimes-${triple}-llvm-bins/libclc/test`. This allows testing multiple triples in the same test file. Add script libclc/test/update_libclc_tests.py, which wraps utils/update_cc_test_checks.py to update CHECK lines in libclc .cl tests for a given arch. Example usage: `libclc/test/update_libclc_tests.py amdgpu` Assisted-by: Claude Sonnet 4.6 --- libclc/test/AMDGPU/external-funcs.test | 1 + libclc/test/AMDGPU/lit.local.cfg | 2 + libclc/test/CMakeLists.txt | 42 ++----- libclc/test/add_sat.cl | 11 -- libclc/test/as_type.cl | 11 -- libclc/test/conversion/convert.cl | 14 +++ libclc/test/convert.cl | 11 -- libclc/test/cos.cl | 11 -- libclc/test/cross.cl | 11 -- libclc/test/fabs.cl | 11 -- libclc/test/geometric/cross.cl | 31 +++++ libclc/test/get_group_id.cl | 11 -- libclc/test/integer/add_sat.cl | 14 +++ libclc/test/integer/sub_sat.cl | 44 +++++++ libclc/test/lit.cfg.py | 25 +++- libclc/test/lit.site.cfg.py.in | 8 +- libclc/test/math/cos.cl | 156 +++++++++++++++++++++++++ libclc/test/math/fabs.cl | 14 +++ libclc/test/math/rsqrt.cl | 45 +++++++ libclc/test/misc/as_type.cl | 14 +++ libclc/test/rsqrt.cl | 14 --- libclc/test/subsat.cl | 27 ----- libclc/test/update_libclc_tests.py | 146 +++++++++++++++++++++++ libclc/test/work-item/get_group_id.cl | 15 +++ 24 files changed, 534 insertions(+), 155 deletions(-) create mode 100644 libclc/test/AMDGPU/external-funcs.test create mode 100644 libclc/test/AMDGPU/lit.local.cfg delete mode 100644 libclc/test/add_sat.cl delete mode 100644 libclc/test/as_type.cl create mode 100644 libclc/test/conversion/convert.cl delete mode 100644 libclc/test/convert.cl delete mode 100644 libclc/test/cos.cl delete mode 100644 libclc/test/cross.cl delete mode 100644 libclc/test/fabs.cl create mode 100644 libclc/test/geometric/cross.cl delete mode 100644 libclc/test/get_group_id.cl create mode 100644 libclc/test/integer/add_sat.cl create mode 100644 libclc/test/integer/sub_sat.cl create mode 100644 libclc/test/math/cos.cl create mode 100644 libclc/test/math/fabs.cl create mode 100644 libclc/test/math/rsqrt.cl create mode 100644 libclc/test/misc/as_type.cl delete mode 100644 libclc/test/rsqrt.cl delete mode 100644 libclc/test/subsat.cl create mode 100755 libclc/test/update_libclc_tests.py create mode 100644 libclc/test/work-item/get_group_id.cl diff --git a/libclc/test/AMDGPU/external-funcs.test b/libclc/test/AMDGPU/external-funcs.test new file mode 100644 index 0000000000000..73bddc02d11e1 --- /dev/null +++ b/libclc/test/AMDGPU/external-funcs.test @@ -0,0 +1 @@ +; RUN: llvm-nm -u "%libclc_library_dir/%libclc_target/libclc.bc" | FileCheck %s --allow-empty --implicit-check-not=" U " diff --git a/libclc/test/AMDGPU/lit.local.cfg b/libclc/test/AMDGPU/lit.local.cfg new file mode 100644 index 0000000000000..1ee232d629c16 --- /dev/null +++ b/libclc/test/AMDGPU/lit.local.cfg @@ -0,0 +1,2 @@ +if "AMDGCN" not in config.available_features: + config.unsupported = True diff --git a/libclc/test/CMakeLists.txt b/libclc/test/CMakeLists.txt index 4c54d93b26813..955235f32277d 100644 --- a/libclc/test/CMakeLists.txt +++ b/libclc/test/CMakeLists.txt @@ -1,36 +1,14 @@ set(LLVM_TOOLS_DIR ${LLVM_TOOLS_BINARY_DIR}) -set(LIBCLC_TEST_DEPS - libclc-opencl-builtins +configure_lit_site_cfg( + ${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.py.in + ${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg.py + MAIN_CONFIG + ${CMAKE_CURRENT_SOURCE_DIR}/lit.cfg.py ) -umbrella_lit_testsuite_begin(check-libclc) - -# Testing unresolved symbols. -# Skip nvptx and spirv targets. -if(ARCH MATCHES amdgcn) - foreach(tgt IN LISTS LIBCLC_UNRESOLVED_SYMBOL_TEST_TARGETS) - set(target_file "$<TARGET_PROPERTY:${tgt},TARGET_FILE>") - - set(LIBCLC_TARGET_TEST_DIR ${CMAKE_CURRENT_BINARY_DIR}/${tgt}) - file(MAKE_DIRECTORY ${LIBCLC_TARGET_TEST_DIR}) - file(GENERATE OUTPUT ${LIBCLC_TARGET_TEST_DIR}/check-external-funcs.test - CONTENT "; RUN: llvm-nm -u \"${target_file}\" | FileCheck %s --allow-empty\n\n; CHECK-NOT: {{.+}}\n" - ) - - configure_lit_site_cfg( - ${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.py.in - ${LIBCLC_TARGET_TEST_DIR}/lit.site.cfg.py - MAIN_CONFIG - ${CMAKE_CURRENT_SOURCE_DIR}/lit.cfg.py - ) - - add_lit_testsuite(check-libclc-external-funcs-${tgt} "Running ${tgt} unresolved symbols tests" - ${LIBCLC_TARGET_TEST_DIR} - DEPENDS ${LIBCLC_TEST_DEPS} - ) - set_target_properties(check-libclc-external-funcs-${tgt} PROPERTIES FOLDER "libclc tests") - endforeach() -endif() - -umbrella_lit_testsuite_end(check-libclc) +add_lit_testsuite(check-libclc "Running libclc tests" + ${CMAKE_CURRENT_BINARY_DIR} + DEPENDS libclc +) +set_target_properties(check-libclc PROPERTIES FOLDER "libclc tests") diff --git a/libclc/test/add_sat.cl b/libclc/test/add_sat.cl deleted file mode 100644 index 87c3d39df3542..0000000000000 --- a/libclc/test/add_sat.cl +++ /dev/null @@ -1,11 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -__kernel void foo(__global char *a, __global char *b, __global char *c) { - *a = add_sat(*b, *c); -} diff --git a/libclc/test/as_type.cl b/libclc/test/as_type.cl deleted file mode 100644 index a926f48c4ea0c..0000000000000 --- a/libclc/test/as_type.cl +++ /dev/null @@ -1,11 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -__kernel void foo(int4 *x, float4 *y) { - *x = as_int4(*y); -} diff --git a/libclc/test/conversion/convert.cl b/libclc/test/conversion/convert.cl new file mode 100644 index 0000000000000..00b524bb8f11d --- /dev/null +++ b/libclc/test/conversion/convert.cl @@ -0,0 +1,14 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 +// REQUIRES: AMDGCN + +// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix + +// AMDGCN-LABEL: define hidden <4 x i32> @float4_to_int4( +// AMDGCN-SAME: <4 x float> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { +// AMDGCN-NEXT: [[ENTRY:.*:]] +// AMDGCN-NEXT: [[TMP0:%.*]] = fptosi <4 x float> [[X]] to <4 x i32> +// AMDGCN-NEXT: ret <4 x i32> [[TMP0]] +// +int4 float4_to_int4(float4 x) { + return convert_int4(x); +} diff --git a/libclc/test/convert.cl b/libclc/test/convert.cl deleted file mode 100644 index 8eba608dc5f8c..0000000000000 --- a/libclc/test/convert.cl +++ /dev/null @@ -1,11 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -__kernel void foo(int4 *x, float4 *y) { - *x = convert_int4(*y); -} diff --git a/libclc/test/cos.cl b/libclc/test/cos.cl deleted file mode 100644 index 92a998b3ba5f7..0000000000000 --- a/libclc/test/cos.cl +++ /dev/null @@ -1,11 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -__kernel void foo(float4 *f) { - *f = cos(*f); -} diff --git a/libclc/test/cross.cl b/libclc/test/cross.cl deleted file mode 100644 index 90762d0d073a6..0000000000000 --- a/libclc/test/cross.cl +++ /dev/null @@ -1,11 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -__kernel void foo(float4 *f) { - *f = cross(f[0], f[1]); -} diff --git a/libclc/test/fabs.cl b/libclc/test/fabs.cl deleted file mode 100644 index 3f5a964e0418a..0000000000000 --- a/libclc/test/fabs.cl +++ /dev/null @@ -1,11 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -__kernel void foo(float *f) { - *f = fabs(*f); -} diff --git a/libclc/test/geometric/cross.cl b/libclc/test/geometric/cross.cl new file mode 100644 index 0000000000000..5aa7541b4d49c --- /dev/null +++ b/libclc/test/geometric/cross.cl @@ -0,0 +1,31 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 +// REQUIRES: AMDGCN + +// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix + +// AMDGCN-LABEL: define hidden <4 x float> @test_float4( +// AMDGCN-SAME: <4 x float> noundef [[X:%.*]], <4 x float> noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { +// AMDGCN-NEXT: [[ENTRY:.*:]] +// AMDGCN-NEXT: [[TMP0:%.*]] = extractelement <4 x float> [[X]], i64 1 +// AMDGCN-NEXT: [[TMP1:%.*]] = extractelement <4 x float> [[Y]], i64 2 +// AMDGCN-NEXT: [[TMP2:%.*]] = fmul contract float [[TMP0]], [[TMP1]] +// AMDGCN-NEXT: [[TMP3:%.*]] = extractelement <4 x float> [[X]], i64 2 +// AMDGCN-NEXT: [[TMP4:%.*]] = extractelement <4 x float> [[Y]], i64 1 +// AMDGCN-NEXT: [[TMP5:%.*]] = fmul contract float [[TMP3]], [[TMP4]] +// AMDGCN-NEXT: [[TMP6:%.*]] = fsub contract float [[TMP2]], [[TMP5]] +// AMDGCN-NEXT: [[TMP7:%.*]] = extractelement <4 x float> [[Y]], i64 0 +// AMDGCN-NEXT: [[TMP8:%.*]] = fmul contract float [[TMP3]], [[TMP7]] +// AMDGCN-NEXT: [[TMP9:%.*]] = extractelement <4 x float> [[X]], i64 0 +// AMDGCN-NEXT: [[TMP10:%.*]] = fmul contract float [[TMP9]], [[TMP1]] +// AMDGCN-NEXT: [[TMP11:%.*]] = fsub contract float [[TMP8]], [[TMP10]] +// AMDGCN-NEXT: [[TMP12:%.*]] = fmul contract float [[TMP9]], [[TMP4]] +// AMDGCN-NEXT: [[TMP13:%.*]] = fmul contract float [[TMP0]], [[TMP7]] +// AMDGCN-NEXT: [[TMP14:%.*]] = fsub contract float [[TMP12]], [[TMP13]] +// AMDGCN-NEXT: [[TMP15:%.*]] = insertelement <4 x float> <float poison, float poison, float poison, float 0.000000e+00>, float [[TMP6]], i64 0 +// AMDGCN-NEXT: [[TMP16:%.*]] = insertelement <4 x float> [[TMP15]], float [[TMP11]], i64 1 +// AMDGCN-NEXT: [[TMP17:%.*]] = insertelement <4 x float> [[TMP16]], float [[TMP14]], i64 2 +// AMDGCN-NEXT: ret <4 x float> [[TMP17]] +// +float4 test_float4(float4 x, float4 y) { + return cross(x, y); +} diff --git a/libclc/test/get_group_id.cl b/libclc/test/get_group_id.cl deleted file mode 100644 index c2349a0076889..0000000000000 --- a/libclc/test/get_group_id.cl +++ /dev/null @@ -1,11 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -__kernel void foo(int *i) { - i[get_group_id(0)] = 1; -} diff --git a/libclc/test/integer/add_sat.cl b/libclc/test/integer/add_sat.cl new file mode 100644 index 0000000000000..4c30090953e99 --- /dev/null +++ b/libclc/test/integer/add_sat.cl @@ -0,0 +1,14 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 +// REQUIRES: AMDGCN + +// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix + +// AMDGCN-LABEL: define hidden noundef signext i8 @test_char( +// AMDGCN-SAME: i8 noundef signext [[X:%.*]], i8 noundef signext [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { +// AMDGCN-NEXT: [[ENTRY:.*:]] +// AMDGCN-NEXT: [[TMP0:%.*]] = tail call noundef i8 @llvm.sadd.sat.i8(i8 [[X]], i8 [[Y]]) +// AMDGCN-NEXT: ret i8 [[TMP0]] +// +char test_char(char x, char y) { + return add_sat(x, y); +} diff --git a/libclc/test/integer/sub_sat.cl b/libclc/test/integer/sub_sat.cl new file mode 100644 index 0000000000000..6573afae45102 --- /dev/null +++ b/libclc/test/integer/sub_sat.cl @@ -0,0 +1,44 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 +// REQUIRES: AMDGCN + +// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix + +// AMDGCN-LABEL: define hidden noundef signext i8 @test_char( +// AMDGCN-SAME: i8 noundef signext [[X:%.*]], i8 noundef signext [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { +// AMDGCN-NEXT: [[ENTRY:.*:]] +// AMDGCN-NEXT: [[TMP0:%.*]] = tail call noundef i8 @llvm.ssub.sat.i8(i8 [[X]], i8 [[Y]]) +// AMDGCN-NEXT: ret i8 [[TMP0]] +// +char test_char(char x, char y) { + return sub_sat(x, y); +} + +// AMDGCN-LABEL: define hidden noundef zeroext i8 @test_uchar( +// AMDGCN-SAME: i8 noundef zeroext [[X:%.*]], i8 noundef zeroext [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] { +// AMDGCN-NEXT: [[ENTRY:.*:]] +// AMDGCN-NEXT: [[TMP0:%.*]] = tail call noundef i8 @llvm.usub.sat.i8(i8 [[X]], i8 [[Y]]) +// AMDGCN-NEXT: ret i8 [[TMP0]] +// +uchar test_uchar(uchar x, uchar y) { + return sub_sat(x, y); +} + +// AMDGCN-LABEL: define hidden noundef i64 @test_long( +// AMDGCN-SAME: i64 noundef [[X:%.*]], i64 noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] { +// AMDGCN-NEXT: [[ENTRY:.*:]] +// AMDGCN-NEXT: [[TMP0:%.*]] = tail call noundef i64 @llvm.ssub.sat.i64(i64 [[X]], i64 [[Y]]) +// AMDGCN-NEXT: ret i64 [[TMP0]] +// +long test_long(long x, long y) { + return sub_sat(x, y); +} + +// AMDGCN-LABEL: define hidden noundef i64 @test_ulong( +// AMDGCN-SAME: i64 noundef [[X:%.*]], i64 noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR0]] { +// AMDGCN-NEXT: [[ENTRY:.*:]] +// AMDGCN-NEXT: [[TMP0:%.*]] = tail call noundef i64 @llvm.usub.sat.i64(i64 [[X]], i64 [[Y]]) +// AMDGCN-NEXT: ret i64 [[TMP0]] +// +ulong test_ulong(ulong x, ulong y) { + return sub_sat(x, y); +} diff --git a/libclc/test/lit.cfg.py b/libclc/test/lit.cfg.py index de1865c2ab3b6..b447d20a30c77 100644 --- a/libclc/test/lit.cfg.py +++ b/libclc/test/lit.cfg.py @@ -6,27 +6,46 @@ import lit.formats +from lit.llvm import llvm_config + # Configuration file for the 'lit' test runner. # name: The name of this test suite. config.name = "libclc" # testFormat: The test format to use to interpret tests. -config.test_format = lit.formats.ShTest(True) +config.test_format = lit.formats.ShTest() # suffixes: A list of file extensions to treat as test files. -config.suffixes = [".test"] +config.suffixes = [".cl", ".test"] # Exclude certain directories from test discovery config.excludes = ["CMakeLists.txt"] # test_source_root: The root path where tests are located. # For per-target tests, this is the target's test directory. -config.test_source_root = config.libclc_obj_root +config.test_source_root = os.path.dirname(__file__) # test_exec_root: The root path where tests should be run. config.test_exec_root = config.libclc_obj_root +config.target_triple = config.libclc_target + +arch = config.libclc_arch.upper() +config.available_features.add(arch) + +llvm_config.use_default_substitutions() + +llvm_config.use_clang() + +llvm_config.add_tool_substitutions(["llvm-nm"], config.llvm_tools_dir) + +config.substitutions.extend([ + ("%libclc_library_dir", config.libclc_library_dir), + ("%libclc_target", config.libclc_target), + ("%check_prefix", arch) +]) + # Propagate PATH from environment if "PATH" in os.environ: config.environment["PATH"] = os.path.pathsep.join( diff --git a/libclc/test/lit.site.cfg.py.in b/libclc/test/lit.site.cfg.py.in index ca6ed0144ea9c..d24227d850e6f 100644 --- a/libclc/test/lit.site.cfg.py.in +++ b/libclc/test/lit.site.cfg.py.in @@ -2,8 +2,12 @@ import sys -config.llvm_tools_dir = "@LLVM_TOOLS_DIR@" -config.libclc_obj_root = "@LIBCLC_TARGET_TEST_DIR@" +config.host_triple = "@LLVM_HOST_TRIPLE@" +config.llvm_tools_dir = lit_config.substitute(path(r"@LLVM_TOOLS_DIR@")) +config.libclc_arch = "@ARCH@" +config.libclc_library_dir = lit_config.substitute(path(r"@LIBCLC_OUTPUT_LIBRARY_DIR@")) +config.libclc_obj_root = lit_config.substitute(path(r"@CMAKE_CURRENT_BINARY_DIR@")) +config.libclc_target = "@LIBCLC_TARGET@" import lit.llvm lit.llvm.initialize(lit_config, config) diff --git a/libclc/test/math/cos.cl b/libclc/test/math/cos.cl new file mode 100644 index 0000000000000..311024d7155e0 --- /dev/null +++ b/libclc/test/math/cos.cl @@ -0,0 +1,156 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 +// REQUIRES: AMDGCN + +// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix + +// AMDGCN-LABEL: define hidden float @test_float( +// AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { +// AMDGCN-NEXT: [[ENTRY:.*:]] +// AMDGCN-NEXT: [[TMP0:%.*]] = tail call float @llvm.fabs.f32(float [[X]]) +// AMDGCN-NEXT: [[TMP1:%.*]] = fcmp une float [[TMP0]], +inf +// AMDGCN-NEXT: [[TMP2:%.*]] = select contract i1 [[TMP1]], float [[TMP0]], float +qnan +// AMDGCN-NEXT: [[TMP3:%.*]] = fcmp ult float [[TMP2]], 1.310720e+05 +// AMDGCN-NEXT: br i1 [[TMP3]], label %[[BB108:.*]], label %[[BB4:.*]] +// AMDGCN: [[BB4]]: +// AMDGCN-NEXT: [[TMP5:%.*]] = tail call { float, i32 } @llvm.frexp.f32.i32(float [[TMP2]]) +// AMDGCN-NEXT: [[TMP6:%.*]] = extractvalue { float, i32 } [[TMP5]], 1 +// AMDGCN-NEXT: [[TMP7:%.*]] = extractvalue { float, i32 } [[TMP5]], 0 +// AMDGCN-NEXT: [[TMP8:%.*]] = fmul contract float [[TMP7]], f0x4B800000 +// AMDGCN-NEXT: [[TMP9:%.*]] = fptoui float [[TMP8]] to i32 +// AMDGCN-NEXT: [[TMP10:%.*]] = zext i32 [[TMP9]] to i64 +// AMDGCN-NEXT: [[TMP11:%.*]] = mul nuw i64 [[TMP10]], 4266746795 +// AMDGCN-NEXT: [[TMP12:%.*]] = trunc i64 [[TMP11]] to i32 +// AMDGCN-NEXT: [[TMP13:%.*]] = lshr i64 [[TMP11]], 32 +// AMDGCN-NEXT: [[TMP14:%.*]] = mul nuw nsw i64 [[TMP10]], 1011060801 +// AMDGCN-NEXT: [[TMP15:%.*]] = add nuw nsw i64 [[TMP13]], [[TMP14]] +// AMDGCN-NEXT: [[TMP16:%.*]] = trunc i64 [[TMP15]] to i32 +// AMDGCN-NEXT: [[TMP17:%.*]] = lshr i64 [[TMP15]], 32 +// AMDGCN-NEXT: [[TMP18:%.*]] = mul nuw i64 [[TMP10]], 3680671129 +// AMDGCN-NEXT: [[TMP19:%.*]] = add nuw i64 [[TMP17]], [[TMP18]] +// AMDGCN-NEXT: [[TMP20:%.*]] = trunc i64 [[TMP19]] to i32 +// AMDGCN-NEXT: [[TMP21:%.*]] = lshr i64 [[TMP19]], 32 +// AMDGCN-NEXT: [[TMP22:%.*]] = mul nuw i64 [[TMP10]], 4113882560 +// AMDGCN-NEXT: [[TMP23:%.*]] = add nuw i64 [[TMP21]], [[TMP22]] +// AMDGCN-NEXT: [[TMP24:%.*]] = trunc i64 [[TMP23]] to i32 +// AMDGCN-NEXT: [[TMP25:%.*]] = lshr i64 [[TMP23]], 32 +// AMDGCN-NEXT: [[TMP26:%.*]] = mul nuw i64 [[TMP10]], 4230436817 +// AMDGCN-NEXT: [[TMP27:%.*]] = add nuw i64 [[TMP25]], [[TMP26]] +// AMDGCN-NEXT: [[TMP28:%.*]] = trunc i64 [[TMP27]] to i32 +// AMDGCN-NEXT: [[TMP29:%.*]] = lshr i64 [[TMP27]], 32 +// AMDGCN-NEXT: [[TMP30:%.*]] = mul nuw nsw i64 [[TMP10]], 1313084713 +// AMDGCN-NEXT: [[TMP31:%.*]] = add nuw nsw i64 [[TMP29]], [[TMP30]] +// AMDGCN-NEXT: [[TMP32:%.*]] = trunc i64 [[TMP31]] to i32 +// AMDGCN-NEXT: [[TMP33:%.*]] = lshr i64 [[TMP31]], 32 +// AMDGCN-NEXT: [[TMP34:%.*]] = mul nuw i64 [[TMP10]], 2734261102 +// AMDGCN-NEXT: [[TMP35:%.*]] = add nuw i64 [[TMP33]], [[TMP34]] +// AMDGCN-NEXT: [[TMP36:%.*]] = trunc i64 [[TMP35]] to i32 +// AMDGCN-NEXT: [[TMP37:%.*]] = lshr i64 [[TMP35]], 32 +// AMDGCN-NEXT: [[TMP38:%.*]] = trunc nuw i64 [[TMP37]] to i32 +// AMDGCN-NEXT: [[TMP39:%.*]] = add nsw i32 [[TMP6]], 6 +// AMDGCN-NEXT: [[TMP40:%.*]] = icmp ugt i32 [[TMP39]], 63 +// AMDGCN-NEXT: [[TMP41:%.*]] = select i1 [[TMP40]], i32 [[TMP32]], i32 [[TMP38]] +// AMDGCN-NEXT: [[TMP42:%.*]] = select i1 [[TMP40]], i32 [[TMP28]], i32 [[TMP36]] +// AMDGCN-NEXT: [[TMP43:%.*]] = select i1 [[TMP40]], i32 [[TMP24]], i32 [[TMP32]] +// AMDGCN-NEXT: [[TMP44:%.*]] = select i1 [[TMP40]], i32 [[TMP20]], i32 [[TMP28]] +// AMDGCN-NEXT: [[TMP45:%.*]] = select i1 [[TMP40]], i32 [[TMP16]], i32 [[TMP24]] +// AMDGCN-NEXT: [[TMP46:%.*]] = select i1 [[TMP40]], i32 [[TMP12]], i32 [[TMP20]] +// AMDGCN-NEXT: [[TMP47:%.*]] = select i1 [[TMP40]], i32 -64, i32 0 +// AMDGCN-NEXT: [[TMP48:%.*]] = add i32 [[TMP47]], [[TMP39]] +// AMDGCN-NEXT: [[TMP49:%.*]] = icmp ugt i32 [[TMP48]], 31 +// AMDGCN-NEXT: [[TMP50:%.*]] = select i1 [[TMP49]], i32 [[TMP42]], i32 [[TMP41]] +// AMDGCN-NEXT: [[TMP51:%.*]] = select i1 [[TMP49]], i32 [[TMP43]], i32 [[TMP42]] +// AMDGCN-NEXT: [[TMP52:%.*]] = select i1 [[TMP49]], i32 [[TMP44]], i32 [[TMP43]] +// AMDGCN-NEXT: [[TMP53:%.*]] = select i1 [[TMP49]], i32 [[TMP45]], i32 [[TMP44]] +// AMDGCN-NEXT: [[TMP54:%.*]] = select i1 [[TMP49]], i32 [[TMP46]], i32 [[TMP45]] +// AMDGCN-NEXT: [[TMP55:%.*]] = select i1 [[TMP49]], i32 -32, i32 0 +// AMDGCN-NEXT: [[TMP56:%.*]] = add i32 [[TMP55]], [[TMP48]] +// AMDGCN-NEXT: [[TMP57:%.*]] = icmp ugt i32 [[TMP56]], 31 +// AMDGCN-NEXT: [[TMP58:%.*]] = select i1 [[TMP57]], i32 [[TMP51]], i32 [[TMP50]] +// AMDGCN-NEXT: [[TMP59:%.*]] = select i1 [[TMP57]], i32 [[TMP52]], i32 [[TMP51]] +// AMDGCN-NEXT: [[TMP60:%.*]] = select i1 [[TMP57]], i32 [[TMP53]], i32 [[TMP52]] +// AMDGCN-NEXT: [[TMP61:%.*]] = select i1 [[TMP57]], i32 [[TMP54]], i32 [[TMP53]] +// AMDGCN-NEXT: [[TMP62:%.*]] = select i1 [[TMP57]], i32 -32, i32 0 +// AMDGCN-NEXT: [[TMP63:%.*]] = add i32 [[TMP62]], [[TMP56]] +// AMDGCN-NEXT: [[TMP64:%.*]] = icmp eq i32 [[TMP63]], 0 +// AMDGCN-NEXT: [[TMP65:%.*]] = sub i32 32, [[TMP63]] +// AMDGCN-NEXT: [[TMP66:%.*]] = tail call i32 @llvm.fshr.i32(i32 [[TMP58]], i32 [[TMP59]], i32 [[TMP65]]) +// AMDGCN-NEXT: [[TMP67:%.*]] = tail call i32 @llvm.fshr.i32(i32 [[TMP59]], i32 [[TMP60]], i32 [[TMP65]]) +// AMDGCN-NEXT: [[TMP68:%.*]] = tail call i32 @llvm.fshr.i32(i32 [[TMP60]], i32 [[TMP61]], i32 [[TMP65]]) +// AMDGCN-NEXT: [[TMP69:%.*]] = select i1 [[TMP64]], i32 [[TMP58]], i32 [[TMP66]] +// AMDGCN-NEXT: [[TMP70:%.*]] = select i1 [[TMP64]], i32 [[TMP59]], i32 [[TMP67]] +// AMDGCN-NEXT: [[TMP71:%.*]] = select i1 [[TMP64]], i32 [[TMP60]], i32 [[TMP68]] +// AMDGCN-NEXT: [[TMP72:%.*]] = lshr i32 [[TMP69]], 29 +// AMDGCN-NEXT: [[TMP73:%.*]] = tail call i32 @llvm.fshl.i32(i32 [[TMP69]], i32 [[TMP70]], i32 2) +// AMDGCN-NEXT: [[TMP74:%.*]] = tail call i32 @llvm.fshl.i32(i32 [[TMP70]], i32 [[TMP71]], i32 2) +// AMDGCN-NEXT: [[TMP75:%.*]] = tail call i32 @llvm.fshl.i32(i32 [[TMP71]], i32 [[TMP61]], i32 2) +// AMDGCN-NEXT: [[TMP76:%.*]] = and i32 [[TMP72]], 1 +// AMDGCN-NEXT: [[TMP77:%.*]] = sub nsw i32 0, [[TMP76]] +// AMDGCN-NEXT: [[TMP78:%.*]] = shl i32 [[TMP72]], 31 +// AMDGCN-NEXT: [[TMP79:%.*]] = xor i32 [[TMP73]], [[TMP77]] +// AMDGCN-NEXT: [[TMP80:%.*]] = xor i32 [[TMP74]], [[TMP77]] +// AMDGCN-NEXT: [[TMP81:%.*]] = xor i32 [[TMP75]], [[TMP77]] +// AMDGCN-NEXT: [[TMP82:%.*]] = tail call range(i32 0, 33) i32 @llvm.ctlz.i32(i32 [[TMP79]], i1 false) +// AMDGCN-NEXT: [[DOTNEG_I_I_I_I:%.*]] = xor i32 [[TMP82]], -1 +// AMDGCN-NEXT: [[TMP83:%.*]] = sub nsw i32 31, [[TMP82]] +// AMDGCN-NEXT: [[TMP84:%.*]] = tail call i32 @llvm.fshr.i32(i32 [[TMP79]], i32 [[TMP80]], i32 [[TMP83]]) +// AMDGCN-NEXT: [[TMP85:%.*]] = tail call i32 @llvm.fshr.i32(i32 [[TMP80]], i32 [[TMP81]], i32 [[TMP83]]) +// AMDGCN-NEXT: [[TMP86:%.*]] = lshr i32 [[TMP84]], 9 +// AMDGCN-NEXT: [[TMP87:%.*]] = shl nuw nsw i32 [[TMP82]], 23 +// AMDGCN-NEXT: [[REASS_SUB:%.*]] = sub nsw i32 [[TMP86]], [[TMP87]] +// AMDGCN-NEXT: [[TMP88:%.*]] = add nsw i32 [[REASS_SUB]], 1056964608 +// AMDGCN-NEXT: [[TMP89:%.*]] = or i32 [[TMP88]], [[TMP78]] +// AMDGCN-NEXT: [[TMP90:%.*]] = bitcast i32 [[TMP89]] to float +// AMDGCN-NEXT: [[TMP91:%.*]] = tail call i32 @llvm.fshl.i32(i32 [[TMP84]], i32 [[TMP85]], i32 23) +// AMDGCN-NEXT: [[TMP92:%.*]] = tail call range(i32 0, 33) i32 @llvm.ctlz.i32(i32 [[TMP91]], i1 false) +// AMDGCN-NEXT: [[TMP93:%.*]] = xor i32 [[TMP92]], -1 +// AMDGCN-NEXT: [[TMP94:%.*]] = tail call i32 @llvm.fshr.i32(i32 [[TMP91]], i32 [[TMP85]], i32 [[TMP93]]) +// AMDGCN-NEXT: [[DOTNEG3_I_I_I_I:%.*]] = sub nuw nsw i32 [[DOTNEG_I_I_I_I]], [[TMP92]] +// AMDGCN-NEXT: [[TMP95:%.*]] = lshr i32 [[TMP94]], 9 +// AMDGCN-NEXT: [[DOTNEG4_I_I_I_I:%.*]] = shl nsw i32 [[DOTNEG3_I_I_I_I]], 23 +// AMDGCN-NEXT: [[TMP96:%.*]] = add nsw i32 [[DOTNEG4_I_I_I_I]], 864026624 +// AMDGCN-NEXT: [[TMP97:%.*]] = or disjoint i32 [[TMP96]], [[TMP95]] +// AMDGCN-NEXT: [[TMP98:%.*]] = or i32 [[TMP97]], [[TMP78]] +// AMDGCN-NEXT: [[TMP99:%.*]] = bitcast i32 [[TMP98]] to float +// AMDGCN-NEXT: [[TMP100:%.*]] = fmul float [[TMP90]], f0x3FC90FDA +// AMDGCN-NEXT: [[TMP101:%.*]] = fneg contract float [[TMP100]] +// AMDGCN-NEXT: [[TMP102:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[TMP90]], float f0x3FC90FDA, float [[TMP101]]) +// AMDGCN-NEXT: [[TMP103:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[TMP90]], float f0x33A22168, float [[TMP102]]) +// AMDGCN-NEXT: [[TMP104:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[TMP99]], float f0x3FC90FDA, float [[TMP103]]) +// AMDGCN-NEXT: [[TMP105:%.*]] = fadd float [[TMP100]], [[TMP104]] +// AMDGCN-NEXT: [[TMP106:%.*]] = lshr i32 [[TMP69]], 30 +// AMDGCN-NEXT: [[TMP107:%.*]] = add nuw nsw i32 [[TMP76]], [[TMP106]] +// AMDGCN-NEXT: br label %[[_Z3COSF_EXIT:.*]] +// AMDGCN: [[BB108]]: +// AMDGCN-NEXT: [[TMP109:%.*]] = fmul float [[TMP2]], f0x3F22F983 +// AMDGCN-NEXT: [[TMP110:%.*]] = tail call contract noundef float @llvm.rint.f32(float [[TMP109]]) +// AMDGCN-NEXT: [[TMP111:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[TMP110]], float f0xBFC90FDA, float [[TMP2]]) +// AMDGCN-NEXT: [[TMP112:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[TMP110]], float f0xB3A22168, float [[TMP111]]) +// AMDGCN-NEXT: [[TMP113:%.*]] = tail call contract noundef float @llvm.fma.f32(float [[TMP110]], float f0xA7C234C4, float [[TMP112]]) +// AMDGCN-NEXT: [[TMP114:%.*]] = fptosi float [[TMP110]] to i32 +// AMDGCN-NEXT: br label %[[_Z3COSF_EXIT]] +// AMDGCN: [[_Z3COSF_EXIT]]: +// AMDGCN-NEXT: [[DOTSINK_I_I_I_I:%.*]] = phi float [ [[TMP113]], %[[BB108]] ], [ [[TMP105]], %[[BB4]] ] +// AMDGCN-NEXT: [[TMP115:%.*]] = phi i32 [ [[TMP114]], %[[BB108]] ], [ [[TMP107]], %[[BB4]] ] +// AMDGCN-NEXT: [[TMP116:%.*]] = fmul float [[DOTSINK_I_I_I_I]], [[DOTSINK_I_I_I_I]] +// AMDGCN-NEXT: [[TMP117:%.*]] = tail call noundef float @llvm.fmuladd.f32(float [[TMP116]], float f0xB94C1982, float f0x3C0881C4) +// AMDGCN-NEXT: [[TMP118:%.*]] = tail call noundef float @llvm.fmuladd.f32(float [[TMP116]], float [[TMP117]], float f0xBE2AAA9D) +// AMDGCN-NEXT: [[TMP119:%.*]] = fmul float [[TMP116]], [[TMP118]] +// AMDGCN-NEXT: [[TMP120:%.*]] = tail call noundef float @llvm.fmuladd.f32(float [[DOTSINK_I_I_I_I]], float [[TMP119]], float [[DOTSINK_I_I_I_I]]) +// AMDGCN-NEXT: [[TMP121:%.*]] = tail call noundef float @llvm.fmuladd.f32(float [[TMP116]], float f0x37D75334, float f0xBAB64F3B) +// AMDGCN-NEXT: [[TMP122:%.*]] = tail call noundef float @llvm.fmuladd.f32(float [[TMP116]], float [[TMP121]], float f0x3D2AABF7) +// AMDGCN-NEXT: [[TMP123:%.*]] = tail call noundef float @llvm.fmuladd.f32(float [[TMP116]], float [[TMP122]], float f0xBF000004) +// AMDGCN-NEXT: [[TMP124:%.*]] = tail call noundef float @llvm.fmuladd.f32(float [[TMP116]], float [[TMP123]], float 1.000000e+00) +// AMDGCN-NEXT: [[TMP125:%.*]] = and i32 [[TMP115]], 1 +// AMDGCN-NEXT: [[TMP126:%.*]] = icmp eq i32 [[TMP125]], 0 +// AMDGCN-NEXT: [[TMP127:%.*]] = shl i32 [[TMP115]], 30 +// AMDGCN-NEXT: [[TMP128:%.*]] = and i32 [[TMP127]], -2147483648 +// AMDGCN-NEXT: [[TMP129:%.*]] = fneg contract float [[TMP120]] +// AMDGCN-NEXT: [[TMP130:%.*]] = select contract i1 [[TMP126]], float [[TMP124]], float [[TMP129]] +// AMDGCN-NEXT: [[TMP131:%.*]] = bitcast float [[TMP130]] to i32 +// AMDGCN-NEXT: [[TMP132:%.*]] = xor i32 [[TMP128]], [[TMP131]] +// AMDGCN-NEXT: [[TMP133:%.*]] = bitcast i32 [[TMP132]] to float +// AMDGCN-NEXT: ret float [[TMP133]] +// +float test_float(float x) { + return cos(x); +} diff --git a/libclc/test/math/fabs.cl b/libclc/test/math/fabs.cl new file mode 100644 index 0000000000000..8b2e34ad51594 --- /dev/null +++ b/libclc/test/math/fabs.cl @@ -0,0 +1,14 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 +// REQUIRES: AMDGCN + +// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix + +// AMDGCN-LABEL: define hidden noundef float @test_float( +// AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { +// AMDGCN-NEXT: [[ENTRY:.*:]] +// AMDGCN-NEXT: [[CALL:%.*]] = tail call fastcc float @llvm.fabs.f32(float noundef [[X]]) #[[ATTR2:[0-9]+]] +// AMDGCN-NEXT: ret float [[CALL]] +// +float test_float(float x) { + return fabs(x); +} diff --git a/libclc/test/math/rsqrt.cl b/libclc/test/math/rsqrt.cl new file mode 100644 index 0000000000000..f36e6a01ed514 --- /dev/null +++ b/libclc/test/math/rsqrt.cl @@ -0,0 +1,45 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 +// REQUIRES: AMDGCN + +// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix + +#pragma OPENCL EXTENSION cl_khr_fp16 : enable + +// AMDGCN-LABEL: define hidden noundef half @test_half( +// AMDGCN-SAME: half noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { +// AMDGCN-NEXT: [[ENTRY:.*:]] +// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract half @llvm.sqrt.f16(half [[X]]), !fpmath [[META13:![0-9]+]] +// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract half 1.000000e+00, [[TMP0]], !fpmath [[META14:![0-9]+]] +// AMDGCN-NEXT: ret half [[TMP1]] +// +half test_half(half x) { + return rsqrt(x); +} + +// AMDGCN-LABEL: define hidden noundef float @test_float( +// AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0]] { +// AMDGCN-NEXT: [[ENTRY:.*:]] +// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.sqrt.f32(float [[X]]), !fpmath [[META15:![0-9]+]] +// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract float 1.000000e+00, [[TMP0]], !fpmath [[META16:![0-9]+]] +// AMDGCN-NEXT: ret float [[TMP1]] +// +float test_float(float x) { + return rsqrt(x); +} + +// AMDGCN-LABEL: define hidden noundef double @test_double( +// AMDGCN-SAME: double noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0]] { +// AMDGCN-NEXT: [[ENTRY:.*:]] +// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract double @llvm.sqrt.f64(double [[X]]) +// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract double 1.000000e+00, [[TMP0]] +// AMDGCN-NEXT: ret double [[TMP1]] +// +double test_double(double x) { + return rsqrt(x); +} +//. +// AMDGCN: [[META13]] = !{float 1.500000e+00} +// AMDGCN: [[META14]] = !{float 1.000000e+00} +// AMDGCN: [[META15]] = !{float 3.000000e+00} +// AMDGCN: [[META16]] = !{float 2.500000e+00} +//. diff --git a/libclc/test/misc/as_type.cl b/libclc/test/misc/as_type.cl new file mode 100644 index 0000000000000..d5c41b15162a6 --- /dev/null +++ b/libclc/test/misc/as_type.cl @@ -0,0 +1,14 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 +// REQUIRES: AMDGCN + +// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix + +// AMDGCN-LABEL: define hidden noundef <4 x i32> @test_float4_as_int4( +// AMDGCN-SAME: <4 x float> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { +// AMDGCN-NEXT: [[ENTRY:.*:]] +// AMDGCN-NEXT: [[ASTYPE:%.*]] = bitcast <4 x float> [[X]] to <4 x i32> +// AMDGCN-NEXT: ret <4 x i32> [[ASTYPE]] +// +int4 test_float4_as_int4(float4 x) { + return as_int4(x); +} diff --git a/libclc/test/rsqrt.cl b/libclc/test/rsqrt.cl deleted file mode 100644 index 4eebfe8ecf7f9..0000000000000 --- a/libclc/test/rsqrt.cl +++ /dev/null @@ -1,14 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -#pragma OPENCL EXTENSION cl_khr_fp64 : enable - -__kernel void foo(float4 *x, double4 *y) { - x[1] = rsqrt(x[0]); - y[1] = rsqrt(y[0]); -} diff --git a/libclc/test/subsat.cl b/libclc/test/subsat.cl deleted file mode 100644 index 5e6fbdcfbef9e..0000000000000 --- a/libclc/test/subsat.cl +++ /dev/null @@ -1,27 +0,0 @@ -//===----------------------------------------------------------------------===// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -__kernel void test_subsat_char(char *a, char x, char y) { - *a = sub_sat(x, y); - return; -} - -__kernel void test_subsat_uchar(uchar *a, uchar x, uchar y) { - *a = sub_sat(x, y); - return; -} - -__kernel void test_subsat_long(long *a, long x, long y) { - *a = sub_sat(x, y); - return; -} - -__kernel void test_subsat_ulong(ulong *a, ulong x, ulong y) { - *a = sub_sat(x, y); - return; -} \ No newline at end of file diff --git a/libclc/test/update_libclc_tests.py b/libclc/test/update_libclc_tests.py new file mode 100755 index 0000000000000..40c48b82b8332 --- /dev/null +++ b/libclc/test/update_libclc_tests.py @@ -0,0 +1,146 @@ +#!/usr/bin/env python3 +"""A utility to wrap utils/update_cc_test_checks.py for updating CHECK lines in +libclc .cl tests for a given architecture. + +The script accepts an architecture argument to determine the triple and check +prefix. Supported arch values: amdgpu, amdgcn, nvptx64, spirv, spirv64. + +The script does 3 things: +1. Replaces %libclc_target and %check_prefix in the .cl file for the arch. +2. Runs update_cc_test_checks.py to update CHECK lines. +3. Reverts the .cl file back to using %libclc_target and %check_prefix. + +Usage: + +% libclc/test/update_libclc_tests.py amdgpu + +""" + +import argparse +import re +import subprocess +import sys +from concurrent.futures import ThreadPoolExecutor, as_completed +from pathlib import Path + +ARCH_TO_TRIPLE = { + "amdgpu": "amdgcn-amd-amdhsa-llvm", + "amdgcn": "amdgcn-amd-amdhsa-llvm", + "nvptx64": "nvptx64-nvidia-cuda", + "spirv": "spirv-unknown-mesa3d", + "spirv64": "spirv64-unknown-mesa3d", +} + +SCRIPT_DIR = Path(__file__).parent +REPO_ROOT = SCRIPT_DIR.parent.parent +UPDATE_SCRIPT = REPO_ROOT / "llvm" / "utils" / "update_cc_test_checks.py" +CLANG = REPO_ROOT / "build" / "bin" / "clang" + + +def find_cl_files(test_dir: Path): + return list(test_dir.rglob("*.cl")) + + +def replace_in_file(path: Path, triple: str, check_prefix: str): + content = path.read_bytes() + content = content.replace(b"%libclc_target", triple.encode()) + content = content.replace(b"%check_prefix", check_prefix.encode()) + path.write_bytes(content) + + +def revert_in_file(path: Path, triple: str, check_prefix: str): + # Only revert in the RUN line context, not in generated CHECK lines. + content = path.read_bytes() + content = content.replace( + f"-target {triple}".encode(), b"-target %libclc_target" + ) + content = content.replace( + f"--check-prefix={check_prefix}".encode(), b"--check-prefix=%check_prefix" + ) + path.write_bytes(content) + + +def file_requires_feature(path: Path, feature: str) -> bool: + try: + text = path.read_text(encoding="utf-8", errors="replace") + except Exception: + return False + for line in text.splitlines(): + stripped = line.strip().lstrip("//").strip() + if stripped.startswith("REQUIRES:"): + rest = stripped[len("REQUIRES:"):] + features = [f.strip() for f in re.split(r",|\|\|", rest)] + if feature in features: + return True + return False + + +def process_file(cl_file: Path, triple: str, check_prefix: str) -> bool: + replace_in_file(cl_file, triple, check_prefix) + cmd = [ + sys.executable, + str(UPDATE_SCRIPT), + "--clang", str(CLANG), + str(cl_file), + ] + print(f" update: {cl_file.relative_to(REPO_ROOT)}") + result = subprocess.run(cmd, capture_output=True, text=True) + ok = result.returncode == 0 + if not ok: + print(f" FAILED: {result.stderr.strip()}", file=sys.stderr) + revert_in_file(cl_file, triple, check_prefix) + return ok + + +def main(): + parser = argparse.ArgumentParser( + description="Update libclc FileCheck assertions for a given arch." + ) + parser.add_argument( + "arch", + choices=list(ARCH_TO_TRIPLE.keys()), + help="Target arch: amdgpu, amdgcn, nvptx64, spirv, spirv64", + ) + args = parser.parse_args() + + arch = args.arch.lower() + triple = ARCH_TO_TRIPLE[arch] + # check_prefix matches REQUIRES feature: uppercase of canonical arch name + # amdgpu -> AMDGCN (same triple as amdgcn), others uppercased + if arch == "amdgpu": + check_prefix = "AMDGCN" + else: + check_prefix = arch.upper() + + cl_files = find_cl_files(SCRIPT_DIR) + # Only process files that REQUIRES this feature + target_files = [f for f in cl_files if file_requires_feature(f, check_prefix)] + + if not target_files: + print(f"No .cl files found with REQUIRES: {check_prefix}") + return + + print(f"arch={arch} triple={triple} check_prefix={check_prefix}") + print(f"Processing {len(target_files)} file(s)...") + + failed = [] + with ThreadPoolExecutor() as executor: + futures = { + executor.submit(process_file, f, triple, check_prefix): f + for f in target_files + } + for future in as_completed(futures): + if not future.result(): + failed.append(futures[future]) + + if failed: + print(f"\n{len(failed)} file(s) failed:", file=sys.stderr) + for f in failed: + print(f" {f}", file=sys.stderr) + sys.exit(1) + else: + print(f"Done. Updated {len(target_files)} file(s).") + + +if __name__ == "__main__": + main() diff --git a/libclc/test/work-item/get_group_id.cl b/libclc/test/work-item/get_group_id.cl new file mode 100644 index 0000000000000..57e3d68787536 --- /dev/null +++ b/libclc/test/work-item/get_group_id.cl @@ -0,0 +1,15 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 +// REQUIRES: AMDGCN + +// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix + +// AMDGCN-LABEL: define hidden range(i64 0, 4294967296) i64 @test( +// AMDGCN-SAME: ) local_unnamed_addr #[[ATTR0:[0-9]+]] { +// AMDGCN-NEXT: [[ENTRY:.*:]] +// AMDGCN-NEXT: [[TMP0:%.*]] = tail call i32 @llvm.amdgcn.workgroup.id.x() +// AMDGCN-NEXT: [[TMP1:%.*]] = zext i32 [[TMP0]] to i64 +// AMDGCN-NEXT: ret i64 [[TMP1]] +// +size_t test() { + return get_group_id(0); +} >From 6518f001607e5792d01ec2cb4c2f899a6cb1f6a8 Mon Sep 17 00:00:00 2001 From: Wenju He <[email protected]> Date: Tue, 12 May 2026 13:17:04 +0200 Subject: [PATCH 2/5] disable clang-format in test folder --- libclc/test/.clang-format | 1 + 1 file changed, 1 insertion(+) create mode 100644 libclc/test/.clang-format diff --git a/libclc/test/.clang-format b/libclc/test/.clang-format new file mode 100644 index 0000000000000..e3845288a2aec --- /dev/null +++ b/libclc/test/.clang-format @@ -0,0 +1 @@ +DisableFormat: true >From f3d10dbefaa60023085993ec94cf243a9b307ee2 Mon Sep 17 00:00:00 2001 From: Wenju He <[email protected]> Date: Tue, 12 May 2026 13:56:38 +0200 Subject: [PATCH 3/5] add -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header --- libclc/test/conversion/convert.cl | 2 +- libclc/test/geometric/cross.cl | 2 +- libclc/test/integer/add_sat.cl | 2 +- libclc/test/integer/sub_sat.cl | 2 +- libclc/test/math/cos.cl | 2 +- libclc/test/math/fabs.cl | 2 +- libclc/test/math/rsqrt.cl | 18 +++++++++--------- libclc/test/misc/as_type.cl | 2 +- libclc/test/work-item/get_group_id.cl | 2 +- 9 files changed, 17 insertions(+), 17 deletions(-) diff --git a/libclc/test/conversion/convert.cl b/libclc/test/conversion/convert.cl index 00b524bb8f11d..eb673a493b403 100644 --- a/libclc/test/conversion/convert.cl +++ b/libclc/test/conversion/convert.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 // REQUIRES: AMDGCN -// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden <4 x i32> @float4_to_int4( // AMDGCN-SAME: <4 x float> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { diff --git a/libclc/test/geometric/cross.cl b/libclc/test/geometric/cross.cl index 5aa7541b4d49c..79724183760e7 100644 --- a/libclc/test/geometric/cross.cl +++ b/libclc/test/geometric/cross.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 // REQUIRES: AMDGCN -// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden <4 x float> @test_float4( // AMDGCN-SAME: <4 x float> noundef [[X:%.*]], <4 x float> noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { diff --git a/libclc/test/integer/add_sat.cl b/libclc/test/integer/add_sat.cl index 4c30090953e99..e223b28444835 100644 --- a/libclc/test/integer/add_sat.cl +++ b/libclc/test/integer/add_sat.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 // REQUIRES: AMDGCN -// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden noundef signext i8 @test_char( // AMDGCN-SAME: i8 noundef signext [[X:%.*]], i8 noundef signext [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { diff --git a/libclc/test/integer/sub_sat.cl b/libclc/test/integer/sub_sat.cl index 6573afae45102..6ba40f75870db 100644 --- a/libclc/test/integer/sub_sat.cl +++ b/libclc/test/integer/sub_sat.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 // REQUIRES: AMDGCN -// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden noundef signext i8 @test_char( // AMDGCN-SAME: i8 noundef signext [[X:%.*]], i8 noundef signext [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { diff --git a/libclc/test/math/cos.cl b/libclc/test/math/cos.cl index 311024d7155e0..40d6825a034a0 100644 --- a/libclc/test/math/cos.cl +++ b/libclc/test/math/cos.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 // REQUIRES: AMDGCN -// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden float @test_float( // AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { diff --git a/libclc/test/math/fabs.cl b/libclc/test/math/fabs.cl index 8b2e34ad51594..3b6ec0f8b729e 100644 --- a/libclc/test/math/fabs.cl +++ b/libclc/test/math/fabs.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 // REQUIRES: AMDGCN -// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden noundef float @test_float( // AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { diff --git a/libclc/test/math/rsqrt.cl b/libclc/test/math/rsqrt.cl index f36e6a01ed514..f0059742d18e6 100644 --- a/libclc/test/math/rsqrt.cl +++ b/libclc/test/math/rsqrt.cl @@ -1,15 +1,15 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 // REQUIRES: AMDGCN -// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix #pragma OPENCL EXTENSION cl_khr_fp16 : enable // AMDGCN-LABEL: define hidden noundef half @test_half( // AMDGCN-SAME: half noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { // AMDGCN-NEXT: [[ENTRY:.*:]] -// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract half @llvm.sqrt.f16(half [[X]]), !fpmath [[META13:![0-9]+]] -// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract half 1.000000e+00, [[TMP0]], !fpmath [[META14:![0-9]+]] +// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract half @llvm.sqrt.f16(half [[X]]), !fpmath [[META11:![0-9]+]] +// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract half 1.000000e+00, [[TMP0]], !fpmath [[META12:![0-9]+]] // AMDGCN-NEXT: ret half [[TMP1]] // half test_half(half x) { @@ -19,8 +19,8 @@ half test_half(half x) { // AMDGCN-LABEL: define hidden noundef float @test_float( // AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0]] { // AMDGCN-NEXT: [[ENTRY:.*:]] -// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.sqrt.f32(float [[X]]), !fpmath [[META15:![0-9]+]] -// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract float 1.000000e+00, [[TMP0]], !fpmath [[META16:![0-9]+]] +// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.sqrt.f32(float [[X]]), !fpmath [[META13:![0-9]+]] +// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract float 1.000000e+00, [[TMP0]], !fpmath [[META14:![0-9]+]] // AMDGCN-NEXT: ret float [[TMP1]] // float test_float(float x) { @@ -38,8 +38,8 @@ double test_double(double x) { return rsqrt(x); } //. -// AMDGCN: [[META13]] = !{float 1.500000e+00} -// AMDGCN: [[META14]] = !{float 1.000000e+00} -// AMDGCN: [[META15]] = !{float 3.000000e+00} -// AMDGCN: [[META16]] = !{float 2.500000e+00} +// AMDGCN: [[META11]] = !{float 1.500000e+00} +// AMDGCN: [[META12]] = !{float 1.000000e+00} +// AMDGCN: [[META13]] = !{float 3.000000e+00} +// AMDGCN: [[META14]] = !{float 2.500000e+00} //. diff --git a/libclc/test/misc/as_type.cl b/libclc/test/misc/as_type.cl index d5c41b15162a6..0ba62fefa4d28 100644 --- a/libclc/test/misc/as_type.cl +++ b/libclc/test/misc/as_type.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 // REQUIRES: AMDGCN -// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden noundef <4 x i32> @test_float4_as_int4( // AMDGCN-SAME: <4 x float> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { diff --git a/libclc/test/work-item/get_group_id.cl b/libclc/test/work-item/get_group_id.cl index 57e3d68787536..cc26101c1b99f 100644 --- a/libclc/test/work-item/get_group_id.cl +++ b/libclc/test/work-item/get_group_id.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 // REQUIRES: AMDGCN -// RUN: %clang -target %libclc_target -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden range(i64 0, 4294967296) i64 @test( // AMDGCN-SAME: ) local_unnamed_addr #[[ATTR0:[0-9]+]] { >From 7839a366682091baee84b184ff856bfa4333e425 Mon Sep 17 00:00:00 2001 From: Wenju He <[email protected]> Date: Wed, 13 May 2026 05:44:26 +0200 Subject: [PATCH 4/5] mcpu, amdgpu-registered-target, ProcessPoolExecutor --- libclc/CMakeLists.txt | 2 - libclc/test/AMDGPU/external-funcs.test | 2 +- libclc/test/AMDGPU/lit.local.cfg | 2 +- libclc/test/CMakeLists.txt | 6 +++ libclc/test/conversion/convert.cl | 4 +- libclc/test/geometric/cross.cl | 4 +- libclc/test/integer/add_sat.cl | 4 +- libclc/test/integer/sub_sat.cl | 4 +- libclc/test/lit.cfg.py | 24 ++++++++--- libclc/test/lit.site.cfg.py.in | 1 + libclc/test/math/cos.cl | 4 +- libclc/test/math/fabs.cl | 4 +- libclc/test/math/rsqrt.cl | 20 ++++----- libclc/test/misc/as_type.cl | 4 +- libclc/test/update_libclc_tests.py | 57 ++++++++++++++++++-------- libclc/test/work-item/get_group_id.cl | 4 +- 16 files changed, 95 insertions(+), 51 deletions(-) diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt index 3f84458336950..a169af7f73cf9 100644 --- a/libclc/CMakeLists.txt +++ b/libclc/CMakeLists.txt @@ -261,8 +261,6 @@ add_libclc_library(libclc-${LIBCLC_TARGET} PARENT_TARGET libclc-opencl-builtins ) -set(LIBCLC_UNRESOLVED_SYMBOL_TEST_TARGETS libclc-${LIBCLC_TARGET}) - if(LLVM_INCLUDE_TESTS) add_subdirectory(test) endif() diff --git a/libclc/test/AMDGPU/external-funcs.test b/libclc/test/AMDGPU/external-funcs.test index 73bddc02d11e1..ab877fb60120d 100644 --- a/libclc/test/AMDGPU/external-funcs.test +++ b/libclc/test/AMDGPU/external-funcs.test @@ -1 +1 @@ -; RUN: llvm-nm -u "%libclc_library_dir/%libclc_target/libclc.bc" | FileCheck %s --allow-empty --implicit-check-not=" U " +; RUN: llvm-nm -u "%library_dir/%target/libclc.bc" | FileCheck %s --allow-empty --implicit-check-not=" U " diff --git a/libclc/test/AMDGPU/lit.local.cfg b/libclc/test/AMDGPU/lit.local.cfg index 1ee232d629c16..68d2d4d5d0023 100644 --- a/libclc/test/AMDGPU/lit.local.cfg +++ b/libclc/test/AMDGPU/lit.local.cfg @@ -1,2 +1,2 @@ -if "AMDGCN" not in config.available_features: +if "amdgpu-registered-target" not in config.available_features: config.unsupported = True diff --git a/libclc/test/CMakeLists.txt b/libclc/test/CMakeLists.txt index 955235f32277d..f1b119df09f1b 100644 --- a/libclc/test/CMakeLists.txt +++ b/libclc/test/CMakeLists.txt @@ -1,5 +1,11 @@ set(LLVM_TOOLS_DIR ${LLVM_TOOLS_BINARY_DIR}) +if(ARCH STREQUAL amdgcn) + set(LIBCLC_CPU "gfx900") +else() + set(LIBCLC_CPU "") +endif() + configure_lit_site_cfg( ${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.py.in ${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg.py diff --git a/libclc/test/conversion/convert.cl b/libclc/test/conversion/convert.cl index eb673a493b403..d172caee1f282 100644 --- a/libclc/test/conversion/convert.cl +++ b/libclc/test/conversion/convert.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 -// REQUIRES: AMDGCN +// REQUIRES: amdgpu-registered-target -// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden <4 x i32> @float4_to_int4( // AMDGCN-SAME: <4 x float> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { diff --git a/libclc/test/geometric/cross.cl b/libclc/test/geometric/cross.cl index 79724183760e7..c72930638eb74 100644 --- a/libclc/test/geometric/cross.cl +++ b/libclc/test/geometric/cross.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 -// REQUIRES: AMDGCN +// REQUIRES: amdgpu-registered-target -// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden <4 x float> @test_float4( // AMDGCN-SAME: <4 x float> noundef [[X:%.*]], <4 x float> noundef [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { diff --git a/libclc/test/integer/add_sat.cl b/libclc/test/integer/add_sat.cl index e223b28444835..2958eed22a47d 100644 --- a/libclc/test/integer/add_sat.cl +++ b/libclc/test/integer/add_sat.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 -// REQUIRES: AMDGCN +// REQUIRES: amdgpu-registered-target -// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden noundef signext i8 @test_char( // AMDGCN-SAME: i8 noundef signext [[X:%.*]], i8 noundef signext [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { diff --git a/libclc/test/integer/sub_sat.cl b/libclc/test/integer/sub_sat.cl index 6ba40f75870db..713b6fce98e1f 100644 --- a/libclc/test/integer/sub_sat.cl +++ b/libclc/test/integer/sub_sat.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 -// REQUIRES: AMDGCN +// REQUIRES: amdgpu-registered-target -// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden noundef signext i8 @test_char( // AMDGCN-SAME: i8 noundef signext [[X:%.*]], i8 noundef signext [[Y:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { diff --git a/libclc/test/lit.cfg.py b/libclc/test/lit.cfg.py index b447d20a30c77..2a3abcc385bd0 100644 --- a/libclc/test/lit.cfg.py +++ b/libclc/test/lit.cfg.py @@ -31,8 +31,21 @@ config.target_triple = config.libclc_target -arch = config.libclc_arch.upper() -config.available_features.add(arch) +supported_test_architectures = ["amdgcn", "amdgpu"] + +def calculate_arch_features(arch_string): + features = [] + for arch in arch_string.split(): + if ( + arch.lower() in supported_test_architectures + and config.libclc_arch.lower() in supported_test_architectures + ): + features.append(arch.lower() + "-registered-target") + return features + +llvm_config.feature_config( + [("--targets-built", calculate_arch_features)] +) llvm_config.use_default_substitutions() @@ -41,9 +54,10 @@ llvm_config.add_tool_substitutions(["llvm-nm"], config.llvm_tools_dir) config.substitutions.extend([ - ("%libclc_library_dir", config.libclc_library_dir), - ("%libclc_target", config.libclc_target), - ("%check_prefix", arch) + ("%library_dir", config.libclc_library_dir), + ("%target", config.libclc_target), + ("%cpu", config.libclc_cpu), + ("%check_prefix", config.libclc_arch.upper()) ]) # Propagate PATH from environment diff --git a/libclc/test/lit.site.cfg.py.in b/libclc/test/lit.site.cfg.py.in index d24227d850e6f..e090e2e16f361 100644 --- a/libclc/test/lit.site.cfg.py.in +++ b/libclc/test/lit.site.cfg.py.in @@ -5,6 +5,7 @@ import sys config.host_triple = "@LLVM_HOST_TRIPLE@" config.llvm_tools_dir = lit_config.substitute(path(r"@LLVM_TOOLS_DIR@")) config.libclc_arch = "@ARCH@" +config.libclc_cpu = "@LIBCLC_CPU@" config.libclc_library_dir = lit_config.substitute(path(r"@LIBCLC_OUTPUT_LIBRARY_DIR@")) config.libclc_obj_root = lit_config.substitute(path(r"@CMAKE_CURRENT_BINARY_DIR@")) config.libclc_target = "@LIBCLC_TARGET@" diff --git a/libclc/test/math/cos.cl b/libclc/test/math/cos.cl index 40d6825a034a0..69a7b3ccb8e4e 100644 --- a/libclc/test/math/cos.cl +++ b/libclc/test/math/cos.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 -// REQUIRES: AMDGCN +// REQUIRES: amdgpu-registered-target -// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden float @test_float( // AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { diff --git a/libclc/test/math/fabs.cl b/libclc/test/math/fabs.cl index 3b6ec0f8b729e..773381b508e04 100644 --- a/libclc/test/math/fabs.cl +++ b/libclc/test/math/fabs.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 -// REQUIRES: AMDGCN +// REQUIRES: amdgpu-registered-target -// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden noundef float @test_float( // AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { diff --git a/libclc/test/math/rsqrt.cl b/libclc/test/math/rsqrt.cl index f0059742d18e6..e47bbbdfe9283 100644 --- a/libclc/test/math/rsqrt.cl +++ b/libclc/test/math/rsqrt.cl @@ -1,15 +1,15 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 -// REQUIRES: AMDGCN +// REQUIRES: amdgpu-registered-target -// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %target -mcpu=%cpu -x cl -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix #pragma OPENCL EXTENSION cl_khr_fp16 : enable // AMDGCN-LABEL: define hidden noundef half @test_half( // AMDGCN-SAME: half noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { // AMDGCN-NEXT: [[ENTRY:.*:]] -// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract half @llvm.sqrt.f16(half [[X]]), !fpmath [[META11:![0-9]+]] -// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract half 1.000000e+00, [[TMP0]], !fpmath [[META12:![0-9]+]] +// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract half @llvm.sqrt.f16(half [[X]]), !fpmath [[META12:![0-9]+]] +// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract half 1.000000e+00, [[TMP0]], !fpmath [[META13:![0-9]+]] // AMDGCN-NEXT: ret half [[TMP1]] // half test_half(half x) { @@ -19,8 +19,8 @@ half test_half(half x) { // AMDGCN-LABEL: define hidden noundef float @test_float( // AMDGCN-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0]] { // AMDGCN-NEXT: [[ENTRY:.*:]] -// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.sqrt.f32(float [[X]]), !fpmath [[META13:![0-9]+]] -// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract float 1.000000e+00, [[TMP0]], !fpmath [[META14:![0-9]+]] +// AMDGCN-NEXT: [[TMP0:%.*]] = tail call contract float @llvm.sqrt.f32(float [[X]]), !fpmath [[META14:![0-9]+]] +// AMDGCN-NEXT: [[TMP1:%.*]] = fdiv contract float 1.000000e+00, [[TMP0]], !fpmath [[META15:![0-9]+]] // AMDGCN-NEXT: ret float [[TMP1]] // float test_float(float x) { @@ -38,8 +38,8 @@ double test_double(double x) { return rsqrt(x); } //. -// AMDGCN: [[META11]] = !{float 1.500000e+00} -// AMDGCN: [[META12]] = !{float 1.000000e+00} -// AMDGCN: [[META13]] = !{float 3.000000e+00} -// AMDGCN: [[META14]] = !{float 2.500000e+00} +// AMDGCN: [[META12]] = !{float 1.500000e+00} +// AMDGCN: [[META13]] = !{float 1.000000e+00} +// AMDGCN: [[META14]] = !{float 3.000000e+00} +// AMDGCN: [[META15]] = !{float 2.500000e+00} //. diff --git a/libclc/test/misc/as_type.cl b/libclc/test/misc/as_type.cl index 0ba62fefa4d28..22fdac86215e6 100644 --- a/libclc/test/misc/as_type.cl +++ b/libclc/test/misc/as_type.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 -// REQUIRES: AMDGCN +// REQUIRES: amdgpu-registered-target -// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden noundef <4 x i32> @test_float4_as_int4( // AMDGCN-SAME: <4 x float> noundef [[X:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] { diff --git a/libclc/test/update_libclc_tests.py b/libclc/test/update_libclc_tests.py index 40c48b82b8332..d16274542203f 100755 --- a/libclc/test/update_libclc_tests.py +++ b/libclc/test/update_libclc_tests.py @@ -6,9 +6,9 @@ prefix. Supported arch values: amdgpu, amdgcn, nvptx64, spirv, spirv64. The script does 3 things: -1. Replaces %libclc_target and %check_prefix in the .cl file for the arch. +1. Replaces %target, %cpu, and %check_prefix in the .cl file for the arch. 2. Runs update_cc_test_checks.py to update CHECK lines. -3. Reverts the .cl file back to using %libclc_target and %check_prefix. +3. Reverts the .cl file back to using %target, %cpu, and %check_prefix. Usage: @@ -17,10 +17,11 @@ """ import argparse +import os import re import subprocess import sys -from concurrent.futures import ThreadPoolExecutor, as_completed +from concurrent.futures import ProcessPoolExecutor, as_completed from pathlib import Path ARCH_TO_TRIPLE = { @@ -31,6 +32,22 @@ "spirv64": "spirv64-unknown-mesa3d", } +ARCH_TO_CPU = { + "amdgpu": "gfx900", + "amdgcn": "gfx900", + "nvptx64": "", + "spirv": "", + "spirv64": "", +} + +ARCH_TO_REQUIRES = { + "amdgpu": "amdgpu-registered-target", + "amdgcn": "amdgpu-registered-target", + "nvptx64": "nvptx-registered-target", + "spirv": "spirv-registered-target", + "spirv64": "spirv-registered-target", +} + SCRIPT_DIR = Path(__file__).parent REPO_ROOT = SCRIPT_DIR.parent.parent UPDATE_SCRIPT = REPO_ROOT / "llvm" / "utils" / "update_cc_test_checks.py" @@ -41,19 +58,25 @@ def find_cl_files(test_dir: Path): return list(test_dir.rglob("*.cl")) -def replace_in_file(path: Path, triple: str, check_prefix: str): +def replace_in_file(path: Path, triple: str, cpu: str, check_prefix: str): content = path.read_bytes() - content = content.replace(b"%libclc_target", triple.encode()) + content = content.replace(b"%target", triple.encode()) + if cpu: + content = content.replace(b"%cpu", cpu.encode()) content = content.replace(b"%check_prefix", check_prefix.encode()) path.write_bytes(content) -def revert_in_file(path: Path, triple: str, check_prefix: str): +def revert_in_file(path: Path, triple: str, cpu: str, check_prefix: str): # Only revert in the RUN line context, not in generated CHECK lines. content = path.read_bytes() content = content.replace( - f"-target {triple}".encode(), b"-target %libclc_target" + f"-target {triple}".encode(), b"-target %target" ) + if cpu: + content = content.replace( + f"-mcpu={cpu}".encode(), b"-mcpu=%cpu" + ) content = content.replace( f"--check-prefix={check_prefix}".encode(), b"--check-prefix=%check_prefix" ) @@ -75,8 +98,8 @@ def file_requires_feature(path: Path, feature: str) -> bool: return False -def process_file(cl_file: Path, triple: str, check_prefix: str) -> bool: - replace_in_file(cl_file, triple, check_prefix) +def process_file(cl_file: Path, triple: str, cpu: str, check_prefix: str) -> bool: + replace_in_file(cl_file, triple, cpu, check_prefix) cmd = [ sys.executable, str(UPDATE_SCRIPT), @@ -88,7 +111,7 @@ def process_file(cl_file: Path, triple: str, check_prefix: str) -> bool: ok = result.returncode == 0 if not ok: print(f" FAILED: {result.stderr.strip()}", file=sys.stderr) - revert_in_file(cl_file, triple, check_prefix) + revert_in_file(cl_file, triple, cpu, check_prefix) return ok @@ -105,6 +128,7 @@ def main(): arch = args.arch.lower() triple = ARCH_TO_TRIPLE[arch] + cpu = ARCH_TO_CPU[arch] # check_prefix matches REQUIRES feature: uppercase of canonical arch name # amdgpu -> AMDGCN (same triple as amdgcn), others uppercased if arch == "amdgpu": @@ -112,21 +136,22 @@ def main(): else: check_prefix = arch.upper() + requires_feature = ARCH_TO_REQUIRES[arch] cl_files = find_cl_files(SCRIPT_DIR) - # Only process files that REQUIRES this feature - target_files = [f for f in cl_files if file_requires_feature(f, check_prefix)] + target_files = [f for f in cl_files if file_requires_feature(f, requires_feature)] if not target_files: - print(f"No .cl files found with REQUIRES: {check_prefix}") + print(f"No .cl files found with REQUIRES: {requires_feature}") return - print(f"arch={arch} triple={triple} check_prefix={check_prefix}") + print(f"arch={arch} triple={triple} cpu={cpu} check_prefix={check_prefix} requires={requires_feature}") print(f"Processing {len(target_files)} file(s)...") failed = [] - with ThreadPoolExecutor() as executor: + num_workers = max(1, os.cpu_count() // 2) + with ProcessPoolExecutor(max_workers=num_workers) as executor: futures = { - executor.submit(process_file, f, triple, check_prefix): f + executor.submit(process_file, f, triple, cpu, check_prefix): f for f in target_files } for future in as_completed(futures): diff --git a/libclc/test/work-item/get_group_id.cl b/libclc/test/work-item/get_group_id.cl index cc26101c1b99f..fc3b97be37bb6 100644 --- a/libclc/test/work-item/get_group_id.cl +++ b/libclc/test/work-item/get_group_id.cl @@ -1,7 +1,7 @@ // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6 -// REQUIRES: AMDGCN +// REQUIRES: amdgpu-registered-target -// RUN: %clang -target %libclc_target -x cl -cl-std=CL3.0 -Xclang -fdeclare-opencl-builtins -Xclang -finclude-default-header -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix +// RUN: %clang -target %target -mcpu=%cpu -cl-std=CL3.0 -O2 -emit-llvm -S -o - %s | FileCheck %s --check-prefix=%check_prefix // AMDGCN-LABEL: define hidden range(i64 0, 4294967296) i64 @test( // AMDGCN-SAME: ) local_unnamed_addr #[[ATTR0:[0-9]+]] { >From 716a595408f6dd27a983730bf24eb3c005f8af1f Mon Sep 17 00:00:00 2001 From: Wenju He <[email protected]> Date: Wed, 13 May 2026 06:37:37 +0200 Subject: [PATCH 5/5] clang-format --- libclc/test/lit.cfg.py | 20 +++++++------ libclc/test/update_libclc_tests.py | 47 +++++++++++++++--------------- 2 files changed, 34 insertions(+), 33 deletions(-) diff --git a/libclc/test/lit.cfg.py b/libclc/test/lit.cfg.py index 2a3abcc385bd0..f55bcc1c42cc3 100644 --- a/libclc/test/lit.cfg.py +++ b/libclc/test/lit.cfg.py @@ -33,6 +33,7 @@ supported_test_architectures = ["amdgcn", "amdgpu"] + def calculate_arch_features(arch_string): features = [] for arch in arch_string.split(): @@ -43,9 +44,8 @@ def calculate_arch_features(arch_string): features.append(arch.lower() + "-registered-target") return features -llvm_config.feature_config( - [("--targets-built", calculate_arch_features)] -) + +llvm_config.feature_config([("--targets-built", calculate_arch_features)]) llvm_config.use_default_substitutions() @@ -53,12 +53,14 @@ def calculate_arch_features(arch_string): llvm_config.add_tool_substitutions(["llvm-nm"], config.llvm_tools_dir) -config.substitutions.extend([ - ("%library_dir", config.libclc_library_dir), - ("%target", config.libclc_target), - ("%cpu", config.libclc_cpu), - ("%check_prefix", config.libclc_arch.upper()) -]) +config.substitutions.extend( + [ + ("%library_dir", config.libclc_library_dir), + ("%target", config.libclc_target), + ("%cpu", config.libclc_cpu), + ("%check_prefix", config.libclc_arch.upper()), + ] +) # Propagate PATH from environment if "PATH" in os.environ: diff --git a/libclc/test/update_libclc_tests.py b/libclc/test/update_libclc_tests.py index d16274542203f..e3de7f33e526c 100755 --- a/libclc/test/update_libclc_tests.py +++ b/libclc/test/update_libclc_tests.py @@ -25,27 +25,27 @@ from pathlib import Path ARCH_TO_TRIPLE = { - "amdgpu": "amdgcn-amd-amdhsa-llvm", - "amdgcn": "amdgcn-amd-amdhsa-llvm", - "nvptx64": "nvptx64-nvidia-cuda", - "spirv": "spirv-unknown-mesa3d", - "spirv64": "spirv64-unknown-mesa3d", + "amdgpu": "amdgcn-amd-amdhsa-llvm", + "amdgcn": "amdgcn-amd-amdhsa-llvm", + "nvptx64": "nvptx64-nvidia-cuda", + "spirv": "spirv-unknown-mesa3d", + "spirv64": "spirv64-unknown-mesa3d", } ARCH_TO_CPU = { - "amdgpu": "gfx900", - "amdgcn": "gfx900", - "nvptx64": "", - "spirv": "", - "spirv64": "", + "amdgpu": "gfx900", + "amdgcn": "gfx900", + "nvptx64": "", + "spirv": "", + "spirv64": "", } ARCH_TO_REQUIRES = { - "amdgpu": "amdgpu-registered-target", - "amdgcn": "amdgpu-registered-target", - "nvptx64": "nvptx-registered-target", - "spirv": "spirv-registered-target", - "spirv64": "spirv-registered-target", + "amdgpu": "amdgpu-registered-target", + "amdgcn": "amdgpu-registered-target", + "nvptx64": "nvptx-registered-target", + "spirv": "spirv-registered-target", + "spirv64": "spirv-registered-target", } SCRIPT_DIR = Path(__file__).parent @@ -70,13 +70,9 @@ def replace_in_file(path: Path, triple: str, cpu: str, check_prefix: str): def revert_in_file(path: Path, triple: str, cpu: str, check_prefix: str): # Only revert in the RUN line context, not in generated CHECK lines. content = path.read_bytes() - content = content.replace( - f"-target {triple}".encode(), b"-target %target" - ) + content = content.replace(f"-target {triple}".encode(), b"-target %target") if cpu: - content = content.replace( - f"-mcpu={cpu}".encode(), b"-mcpu=%cpu" - ) + content = content.replace(f"-mcpu={cpu}".encode(), b"-mcpu=%cpu") content = content.replace( f"--check-prefix={check_prefix}".encode(), b"--check-prefix=%check_prefix" ) @@ -91,7 +87,7 @@ def file_requires_feature(path: Path, feature: str) -> bool: for line in text.splitlines(): stripped = line.strip().lstrip("//").strip() if stripped.startswith("REQUIRES:"): - rest = stripped[len("REQUIRES:"):] + rest = stripped[len("REQUIRES:") :] features = [f.strip() for f in re.split(r",|\|\|", rest)] if feature in features: return True @@ -103,7 +99,8 @@ def process_file(cl_file: Path, triple: str, cpu: str, check_prefix: str) -> boo cmd = [ sys.executable, str(UPDATE_SCRIPT), - "--clang", str(CLANG), + "--clang", + str(CLANG), str(cl_file), ] print(f" update: {cl_file.relative_to(REPO_ROOT)}") @@ -144,7 +141,9 @@ def main(): print(f"No .cl files found with REQUIRES: {requires_feature}") return - print(f"arch={arch} triple={triple} cpu={cpu} check_prefix={check_prefix} requires={requires_feature}") + print( + f"arch={arch} triple={triple} cpu={cpu} check_prefix={check_prefix} requires={requires_feature}" + ) print(f"Processing {len(target_files)} file(s)...") failed = [] _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
