Author: Wenju He Date: 2026-05-13T18:21:36+08:00 New Revision: f6007cf3ada595da3e003e0bc7f8dc1a86bff27b
URL: https://github.com/llvm/llvm-project/commit/f6007cf3ada595da3e003e0bc7f8dc1a86bff27b DIFF: https://github.com/llvm/llvm-project/commit/f6007cf3ada595da3e003e0bc7f8dc1a86bff27b.diff LOG: [libclc] Move external-funcs.test to static file and use IR checks for .cl tests (#197151) 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 check of missing symbols in specific libraries. .cl test files uses %target, %cpu 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 checking outputs of 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 --------- Co-authored-by: Matt Arsenault <[email protected]> Added: libclc/test/.clang-format libclc/test/AMDGPU/external-funcs.test libclc/test/AMDGPU/lit.local.cfg libclc/test/conversion/convert.cl libclc/test/geometric/cross.cl libclc/test/integer/add_sat.cl libclc/test/integer/sub_sat.cl libclc/test/math/cos.cl libclc/test/math/fabs.cl libclc/test/math/rsqrt.cl libclc/test/misc/as_type.cl libclc/test/update_libclc_tests.py libclc/test/work-item/get_group_id.cl Modified: libclc/CMakeLists.txt libclc/test/CMakeLists.txt libclc/test/lit.cfg.py libclc/test/lit.site.cfg.py.in Removed: libclc/test/add_sat.cl libclc/test/as_type.cl libclc/test/convert.cl libclc/test/cos.cl libclc/test/cross.cl libclc/test/fabs.cl libclc/test/get_group_id.cl libclc/test/rsqrt.cl libclc/test/subsat.cl ################################################################################ diff --git a/libclc/CMakeLists.txt b/libclc/CMakeLists.txt index 0d5a20206ee19..5630c8c59cca4 100644 --- a/libclc/CMakeLists.txt +++ b/libclc/CMakeLists.txt @@ -229,8 +229,6 @@ libclc_add_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/.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 diff --git a/libclc/test/AMDGPU/external-funcs.test b/libclc/test/AMDGPU/external-funcs.test new file mode 100644 index 0000000000000..ab877fb60120d --- /dev/null +++ b/libclc/test/AMDGPU/external-funcs.test @@ -0,0 +1 @@ +; 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 new file mode 100644 index 0000000000000..7c492428aec76 --- /dev/null +++ b/libclc/test/AMDGPU/lit.local.cfg @@ -0,0 +1,2 @@ +if not "AMDGPU" in config.root.targets: + config.unsupported = True diff --git a/libclc/test/CMakeLists.txt b/libclc/test/CMakeLists.txt index 4c54d93b26813..0cbbc2838d624 100644 --- a/libclc/test/CMakeLists.txt +++ b/libclc/test/CMakeLists.txt @@ -1,36 +1,20 @@ set(LLVM_TOOLS_DIR ${LLVM_TOOLS_BINARY_DIR}) -set(LIBCLC_TEST_DEPS - libclc-opencl-builtins -) - -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() +if(LIBCLC_TARGET_ARCH IN_LIST LIBCLC_ARCHS_AMDGPU) + set(LIBCLC_TARGET_CPU "gfx900") +else() + set(LIBCLC_TARGET_CPU "") endif() -umbrella_lit_testsuite_end(check-libclc) +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 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..2e2a79e4ebb1e --- /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: amdgpu-registered-target + +// 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]+]] { +// 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..41be4176f76cf --- /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: amdgpu-registered-target + +// 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]+]] { +// 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..21d17423e3f68 --- /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: amdgpu-registered-target + +// 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]+]] { +// 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..15096d57d66d7 --- /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: amdgpu-registered-target + +// 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]+]] { +// 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..22715411ea9d6 100644 --- a/libclc/test/lit.cfg.py +++ b/libclc/test/lit.cfg.py @@ -6,27 +6,68 @@ 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"] +# Exclude certain directories and files from test discovery +config.excludes = [ + "CMakeLists.txt", + "update_libclc_tests.py", +] # 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 + +supported_test_architectures = ["amdgcn", "amdgpu"] + +config.targets = set() + + +def calculate_arch_features(arch_string): + features = [] + for arch in arch_string.split(): + if ( + arch.lower() in supported_test_architectures + and config.libclc_target_arch.lower() in supported_test_architectures + ): + features.append(arch.lower() + "-registered-target") + config.targets.add(arch.upper()) + return features + + +llvm_config.feature_config([("--targets-built", calculate_arch_features)]) + +llvm_config.use_default_substitutions() + +llvm_config.use_clang() + +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_target_cpu), + ("%check_prefix", config.libclc_target_arch.upper()), + ] +) + # 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..3c9acfc0551f0 100644 --- a/libclc/test/lit.site.cfg.py.in +++ b/libclc/test/lit.site.cfg.py.in @@ -2,8 +2,13 @@ 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_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@" +config.libclc_target_arch = "@LIBCLC_TARGET_ARCH@" +config.libclc_target_cpu = "@LIBCLC_TARGET_CPU@" 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..99d3acc2d714d --- /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: amdgpu-registered-target + +// 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]+]] { +// 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..e512abd80468b --- /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: amdgpu-registered-target + +// 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]+]] { +// AMDGCN-NEXT: [[ENTRY:.*:]] +// AMDGCN-NEXT: [[CALL:%.*]] = tail call 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..339c93ea67cf1 --- /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: amdgpu-registered-target + +// 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 [[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) { + 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 [[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) { + 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: [[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 new file mode 100644 index 0000000000000..25f9df2e11137 --- /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: amdgpu-registered-target + +// 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]+]] { +// 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..c0519088dd782 --- /dev/null +++ b/libclc/test/update_libclc_tests.py @@ -0,0 +1,170 @@ +#!/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 %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 %target, %cpu, and %check_prefix. + +Usage: + +% libclc/test/update_libclc_tests.py amdgpu + +""" + +import argparse +import os +import re +import subprocess +import sys +from concurrent.futures import ProcessPoolExecutor, 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", +} + +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" +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, cpu: str, check_prefix: str): + content = path.read_bytes() + 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, 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") + 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" + ) + 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, cpu: str, check_prefix: str) -> bool: + replace_in_file(cl_file, triple, cpu, 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, cpu, 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] + 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": + check_prefix = "AMDGCN" + else: + check_prefix = arch.upper() + + requires_feature = ARCH_TO_REQUIRES[arch] + cl_files = find_cl_files(SCRIPT_DIR) + 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: {requires_feature}") + return + + print( + f"arch={arch} triple={triple} cpu={cpu} check_prefix={check_prefix} requires={requires_feature}" + ) + print(f"Processing {len(target_files)} file(s)...") + + failed = [] + num_workers = max(1, os.cpu_count() // 2) + with ProcessPoolExecutor(max_workers=num_workers) as executor: + futures = { + executor.submit(process_file, f, triple, cpu, 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..1818662c60be1 --- /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: amdgpu-registered-target + +// 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]+]] { +// 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); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
