From 1772e394e2b835afcac45ac43c9ae0b0abd71d1b Mon Sep 17 00:00:00 2001 From: Harald van Dijk Date: Mon, 14 Apr 2025 13:24:36 +0100 Subject: [PATCH 1/6] [SYCL][NativeCPU] Copy over more host/aux target data. * Use copyAuxTarget instead of copying over each field. * Copy over TargetOpts as well. * Prevent recording the target flags for builtins. --- clang/lib/Basic/Targets/NativeCPU.cpp | 51 ++++----------------------- clang/lib/Basic/Targets/NativeCPU.h | 2 ++ clang/lib/Driver/Compilation.cpp | 4 +-- clang/test/Driver/sycl-native-cpu.cpp | 2 +- libclc/utils/prepare-builtins.cpp | 7 +++- 5 files changed, 17 insertions(+), 49 deletions(-) diff --git a/clang/lib/Basic/Targets/NativeCPU.cpp b/clang/lib/Basic/Targets/NativeCPU.cpp index 685cf9093645..226c65d03a84 100644 --- a/clang/lib/Basic/Targets/NativeCPU.cpp +++ b/clang/lib/Basic/Targets/NativeCPU.cpp @@ -60,50 +60,13 @@ NativeCPUTargetInfo::NativeCPUTargetInfo(const llvm::Triple &, }()); if (HostTriple.getArch() != llvm::Triple::UnknownArch) { HostTarget = AllocateTarget(HostTriple, Opts); + copyAuxTarget(&*HostTarget); + } +} - // Copy properties from host target. - BoolWidth = HostTarget->getBoolWidth(); - BoolAlign = HostTarget->getBoolAlign(); - IntWidth = HostTarget->getIntWidth(); - IntAlign = HostTarget->getIntAlign(); - HalfWidth = HostTarget->getHalfWidth(); - HalfAlign = HostTarget->getHalfAlign(); - FloatWidth = HostTarget->getFloatWidth(); - FloatAlign = HostTarget->getFloatAlign(); - DoubleWidth = HostTarget->getDoubleWidth(); - DoubleAlign = HostTarget->getDoubleAlign(); - LongWidth = HostTarget->getLongWidth(); - LongAlign = HostTarget->getLongAlign(); - LongLongWidth = HostTarget->getLongLongWidth(); - LongLongAlign = HostTarget->getLongLongAlign(); - PointerWidth = HostTarget->getPointerWidth(LangAS::Default); - PointerAlign = HostTarget->getPointerAlign(LangAS::Default); - MinGlobalAlign = HostTarget->getMinGlobalAlign(/*TypeSize=*/0, - /*HasNonWeakDef=*/true); - NewAlign = HostTarget->getNewAlign(); - DefaultAlignForAttributeAligned = - HostTarget->getDefaultAlignForAttributeAligned(); - SizeType = HostTarget->getSizeType(); - PtrDiffType = HostTarget->getPtrDiffType(LangAS::Default); - IntMaxType = HostTarget->getIntMaxType(); - WCharType = HostTarget->getWCharType(); - WIntType = HostTarget->getWIntType(); - Char16Type = HostTarget->getChar16Type(); - Char32Type = HostTarget->getChar32Type(); - Int64Type = HostTarget->getInt64Type(); - SigAtomicType = HostTarget->getSigAtomicType(); - ProcessIDType = HostTarget->getProcessIDType(); - - UseBitFieldTypeAlignment = HostTarget->useBitFieldTypeAlignment(); - UseZeroLengthBitfieldAlignment = - HostTarget->useZeroLengthBitfieldAlignment(); - UseExplicitBitFieldAlignment = HostTarget->useExplicitBitFieldAlignment(); - ZeroLengthBitfieldBoundary = HostTarget->getZeroLengthBitfieldBoundary(); - - // This is a bit of a lie, but it controls __GCC_ATOMIC_XXX_LOCK_FREE, and - // we need those macros to be identical on host and device, because (among - // other things) they affect which standard library classes are defined, - // and we need all classes to be defined on both the host and device. - MaxAtomicInlineWidth = HostTarget->getMaxAtomicInlineWidth(); +void NativeCPUTargetInfo::setAuxTarget(const TargetInfo *Aux) { + if (Aux) { + copyAuxTarget(Aux); + getTargetOpts() = Aux->getTargetOpts(); } } diff --git a/clang/lib/Basic/Targets/NativeCPU.h b/clang/lib/Basic/Targets/NativeCPU.h index 44106cd8d028..cb2c71ebe39b 100644 --- a/clang/lib/Basic/Targets/NativeCPU.h +++ b/clang/lib/Basic/Targets/NativeCPU.h @@ -57,6 +57,8 @@ class LLVM_LIBRARY_VISIBILITY NativeCPUTargetInfo final : public TargetInfo { } protected: + void setAuxTarget(const TargetInfo *Aux) override; + ArrayRef getGCCRegNames() const override { return {}; } ArrayRef getGCCRegAliases() const override { diff --git a/clang/lib/Driver/Compilation.cpp b/clang/lib/Driver/Compilation.cpp index 1a91cbaa998d..a07e81892372 100644 --- a/clang/lib/Driver/Compilation.cpp +++ b/clang/lib/Driver/Compilation.cpp @@ -7,7 +7,6 @@ //===----------------------------------------------------------------------===// #include "clang/Driver/Compilation.h" -#include "ToolChains/SYCL.h" #include "clang/Basic/LLVM.h" #include "clang/Driver/Action.h" #include "clang/Driver/Driver.h" @@ -128,8 +127,7 @@ Compilation::getArgsForToolChain(const ToolChain *TC, StringRef BoundArch, if (DeviceOffloadKind == Action::OFK_OpenMP || DeviceOffloadKind == Action::OFK_SYCL) { const ToolChain *HostTC = getSingleOffloadToolChain(); - bool SameTripleAsHost = (TC->getTriple() == HostTC->getTriple()) || - isSYCLNativeCPU(TC->getTriple()); + bool SameTripleAsHost = (TC->getTriple() == HostTC->getTriple()); OffloadArgs = TC->TranslateOffloadTargetArgs( *TranslatedArgs, SameTripleAsHost, AllocatedArgs, DeviceOffloadKind); } diff --git a/clang/test/Driver/sycl-native-cpu.cpp b/clang/test/Driver/sycl-native-cpu.cpp index dccb20cec166..6e2cb1f7946d 100644 --- a/clang/test/Driver/sycl-native-cpu.cpp +++ b/clang/test/Driver/sycl-native-cpu.cpp @@ -25,7 +25,7 @@ // CHECK-OPTS: clang{{.*}}"-triple" "[[TRIPLE]]"{{.*}}"-fsycl-is-device" // CHECK-OPTS-NOT: -sycl-opt // CHECK-OPTS-SAME: "-Wno-override-module" "-mllvm" "-sycl-native-cpu-backend" -// CHECK-OPTS-SAME: "-target-feature" "+v9.4a" +// CHECK-OPTS-SAME: "-aux-target-feature" "+v9.4a" // RUN: %clangxx -fsycl -fsycl-targets=spir64 %s -### 2>&1 | FileCheck -check-prefix=CHECK-NONATIVECPU %s // CHECK-NONATIVECPU-NOT: "-D" "__SYCL_NATIVE_CPU__" diff --git a/libclc/utils/prepare-builtins.cpp b/libclc/utils/prepare-builtins.cpp index cf20ffae91d4..28fabd300780 100644 --- a/libclc/utils/prepare-builtins.cpp +++ b/libclc/utils/prepare-builtins.cpp @@ -116,7 +116,12 @@ int main(int argc, char **argv) { // functions were inlined prior to incompatible functions pass. Now that the // inliner runs later in the pipeline we have to remove all of the target // features, so libclc functions will not be earmarked for deletion. - if (M->getTargetTriple().str().find("amdgcn") != std::string::npos) { + // + // NativeCPU uses the same builtins for multiple host targets and should + // likewise not have features that limit the builtins to any particular + // target. + if (M->getTargetTriple().str().find("amdgcn") != std::string::npos || + M->getTargetTriple().str() != "native_cpu") { AttributeMask AM; AM.addAttribute("target-features"); AM.addAttribute("target-cpu"); From 09d9360447723eee7d53a884c6cf2b6d71481426 Mon Sep 17 00:00:00 2001 From: Harald van Dijk Date: Fri, 18 Apr 2025 13:45:54 +0100 Subject: [PATCH 2/6] Assert that Aux is non-nullptr, add test. --- clang/lib/Basic/Targets/NativeCPU.cpp | 7 +++---- sycl/test/native_cpu/target-features.cpp | 12 ++++++++++++ 2 files changed, 15 insertions(+), 4 deletions(-) create mode 100644 sycl/test/native_cpu/target-features.cpp diff --git a/clang/lib/Basic/Targets/NativeCPU.cpp b/clang/lib/Basic/Targets/NativeCPU.cpp index 226c65d03a84..72ef1f2c40be 100644 --- a/clang/lib/Basic/Targets/NativeCPU.cpp +++ b/clang/lib/Basic/Targets/NativeCPU.cpp @@ -65,8 +65,7 @@ NativeCPUTargetInfo::NativeCPUTargetInfo(const llvm::Triple &, } void NativeCPUTargetInfo::setAuxTarget(const TargetInfo *Aux) { - if (Aux) { - copyAuxTarget(Aux); - getTargetOpts() = Aux->getTargetOpts(); - } + assert(Aux && "Cannot invoke setAuxTarget without a valid auxiliary target!"); + copyAuxTarget(Aux); + getTargetOpts() = Aux->getTargetOpts(); } diff --git a/sycl/test/native_cpu/target-features.cpp b/sycl/test/native_cpu/target-features.cpp new file mode 100644 index 000000000000..d1c7737147f8 --- /dev/null +++ b/sycl/test/native_cpu/target-features.cpp @@ -0,0 +1,12 @@ +// RUN: %clang_cc1 -triple native_cpu -aux-triple x86_64-unknown-linux-gnu -fsycl-is-device -fsycl-is-native-cpu -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,NOAVX +// RUN: %clang_cc1 -triple native_cpu -aux-triple x86_64-unknown-linux-gnu -aux-target-cpu skylake -fsycl-is-device -fsycl-is-native-cpu -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,AVX +// RUN: %clang_cc1 -triple native_cpu -aux-triple x86_64-unknown-linux-gnu -aux-target-feature +avx -fsycl-is-device -fsycl-is-native-cpu -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,AVX + +#if __SYCL_DEVICE_ONLY__ +SYCL_EXTERNAL void foo() {} +#endif + +// CHECK: define void @_Z3foov() [[FOO_ATTRS:#[0-9]+]] { +// CHECK: [[FOO_ATTRS]] = { +// NOAVX-NOT: "target-features"="{{[^"]*}}+avx{{[^"]*}}" +// AVX-SAME: "target-features"="{{[^"]*}}+avx{{[^"]*}}" From 997f12033d0c2ebf672efff43f50a82164b672c5 Mon Sep 17 00:00:00 2001 From: Harald van Dijk Date: Fri, 18 Apr 2025 14:28:05 +0100 Subject: [PATCH 3/6] Require native_cpu for new test. --- sycl/test/native_cpu/target-features.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/test/native_cpu/target-features.cpp b/sycl/test/native_cpu/target-features.cpp index d1c7737147f8..1dc0bbc3ba19 100644 --- a/sycl/test/native_cpu/target-features.cpp +++ b/sycl/test/native_cpu/target-features.cpp @@ -1,3 +1,4 @@ +// REQUIRES: native_cpu // RUN: %clang_cc1 -triple native_cpu -aux-triple x86_64-unknown-linux-gnu -fsycl-is-device -fsycl-is-native-cpu -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,NOAVX // RUN: %clang_cc1 -triple native_cpu -aux-triple x86_64-unknown-linux-gnu -aux-target-cpu skylake -fsycl-is-device -fsycl-is-native-cpu -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,AVX // RUN: %clang_cc1 -triple native_cpu -aux-triple x86_64-unknown-linux-gnu -aux-target-feature +avx -fsycl-is-device -fsycl-is-native-cpu -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,AVX From 8b6cfe3838754c9cc39b67461ade09665cd58312 Mon Sep 17 00:00:00 2001 From: Harald van Dijk Date: Fri, 18 Apr 2025 15:14:32 +0100 Subject: [PATCH 4/6] Move new test, check IR differently to hopefully work with --hip. --- .../native_cpu/target-features.cpp | 18 ++++++++++++++++++ sycl/test/native_cpu/target-features.cpp | 13 ------------- 2 files changed, 18 insertions(+), 13 deletions(-) create mode 100644 sycl/test/check_device_code/native_cpu/target-features.cpp delete mode 100644 sycl/test/native_cpu/target-features.cpp diff --git a/sycl/test/check_device_code/native_cpu/target-features.cpp b/sycl/test/check_device_code/native_cpu/target-features.cpp new file mode 100644 index 000000000000..81fab20aa79b --- /dev/null +++ b/sycl/test/check_device_code/native_cpu/target-features.cpp @@ -0,0 +1,18 @@ +// REQUIRES: native_cpu +// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=native_cpu -mllvm -sycl-native-dump-device-ir %s 2>&1 | FileCheck %s --check-prefixes=CHECK,NOAVX +// RUN: %clangxx --target=x86_64-unknown-linux-gnu -march=skylake -fsycl -fsycl-targets=native_cpu -mllvm -sycl-native-dump-device-ir %s 2>&1 | FileCheck %s --check-prefixes=CHECK,AVX +// RUN: %clangxx --target=x86_64-unknown-linux-gnu -mavx -fsycl -fsycl-targets=native_cpu -mllvm -sycl-native-dump-device-ir %s 2>&1 | FileCheck %s --check-prefixes=CHECK,AVX + +#include +using namespace sycl; + +class Test; +int main() { + sycl::queue deviceQueue; + deviceQueue.submit([&](handler &h) { h.single_task([=] {}); }); +} + +// CHECK: define void @_ZTS4Test.NativeCPUKernel({{.*}}) [[ATTRS:#[0-9]+]] +// CHECK: [[ATTRS]] = { +// NOAVX-NOT: "target-features"="{{[^"]*}}+avx{{[^"]*}}" +// AVX-SAME: "target-features"="{{[^"]*}}+avx{{[^"]*}}" diff --git a/sycl/test/native_cpu/target-features.cpp b/sycl/test/native_cpu/target-features.cpp deleted file mode 100644 index 1dc0bbc3ba19..000000000000 --- a/sycl/test/native_cpu/target-features.cpp +++ /dev/null @@ -1,13 +0,0 @@ -// REQUIRES: native_cpu -// RUN: %clang_cc1 -triple native_cpu -aux-triple x86_64-unknown-linux-gnu -fsycl-is-device -fsycl-is-native-cpu -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,NOAVX -// RUN: %clang_cc1 -triple native_cpu -aux-triple x86_64-unknown-linux-gnu -aux-target-cpu skylake -fsycl-is-device -fsycl-is-native-cpu -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,AVX -// RUN: %clang_cc1 -triple native_cpu -aux-triple x86_64-unknown-linux-gnu -aux-target-feature +avx -fsycl-is-device -fsycl-is-native-cpu -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,AVX - -#if __SYCL_DEVICE_ONLY__ -SYCL_EXTERNAL void foo() {} -#endif - -// CHECK: define void @_Z3foov() [[FOO_ATTRS:#[0-9]+]] { -// CHECK: [[FOO_ATTRS]] = { -// NOAVX-NOT: "target-features"="{{[^"]*}}+avx{{[^"]*}}" -// AVX-SAME: "target-features"="{{[^"]*}}+avx{{[^"]*}}" From 57becc9a5040ea3ef4b3ab6d126cb0c07ce76fe8 Mon Sep 17 00:00:00 2001 From: Harald van Dijk Date: Fri, 18 Apr 2025 16:20:17 +0100 Subject: [PATCH 5/6] ELoosen CHECK to support NATIVECPU_USE_OCK=OFF. --- sycl/test/check_device_code/native_cpu/target-features.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/check_device_code/native_cpu/target-features.cpp b/sycl/test/check_device_code/native_cpu/target-features.cpp index 81fab20aa79b..561c636808c2 100644 --- a/sycl/test/check_device_code/native_cpu/target-features.cpp +++ b/sycl/test/check_device_code/native_cpu/target-features.cpp @@ -12,7 +12,7 @@ int main() { deviceQueue.submit([&](handler &h) { h.single_task([=] {}); }); } -// CHECK: define void @_ZTS4Test.NativeCPUKernel({{.*}}) [[ATTRS:#[0-9]+]] +// CHECK: void @_ZTS4Test.NativeCPUKernel({{.*}} [[ATTRS:#[0-9]+]] // CHECK: [[ATTRS]] = { // NOAVX-NOT: "target-features"="{{[^"]*}}+avx{{[^"]*}}" // AVX-SAME: "target-features"="{{[^"]*}}+avx{{[^"]*}}" From 1d4cdef2276ed0f30bcf52bc9a6c7c7c0145bd59 Mon Sep 17 00:00:00 2001 From: Harald van Dijk Date: Tue, 22 Apr 2025 12:17:45 +0100 Subject: [PATCH 6/6] Move test to clang/test/CodeGenSYCL. --- .../CodeGenSYCL/native_cpu_target_features.cpp | 17 +++++++++++++++++ .../native_cpu/target-features.cpp | 18 ------------------ 2 files changed, 17 insertions(+), 18 deletions(-) create mode 100644 clang/test/CodeGenSYCL/native_cpu_target_features.cpp delete mode 100644 sycl/test/check_device_code/native_cpu/target-features.cpp diff --git a/clang/test/CodeGenSYCL/native_cpu_target_features.cpp b/clang/test/CodeGenSYCL/native_cpu_target_features.cpp new file mode 100644 index 000000000000..2005ab84e52c --- /dev/null +++ b/clang/test/CodeGenSYCL/native_cpu_target_features.cpp @@ -0,0 +1,17 @@ +// RUN: %clang_cc1 -triple native_cpu -aux-triple x86_64-unknown-linux-gnu -fsycl-is-device -emit-llvm -fsycl-is-native-cpu -o - %s | FileCheck %s --check-prefixes=CHECK,NOAVX +// RUN: %clang_cc1 -triple native_cpu -aux-triple x86_64-unknown-linux-gnu -aux-target-cpu skylake -fsycl-is-device -emit-llvm -fsycl-is-native-cpu -o - %s | FileCheck %s --check-prefixes=CHECK,AVX +// RUN: %clang_cc1 -triple native_cpu -aux-triple x86_64-unknown-linux-gnu -aux-target-feature +avx -fsycl-is-device -emit-llvm -fsycl-is-native-cpu -o - %s | FileCheck %s --check-prefixes=CHECK,AVX + +#include "Inputs/sycl.hpp" +using namespace sycl; + +class Test; +int main() { + sycl::queue deviceQueue; + deviceQueue.submit([&](handler &h) { h.single_task([=] {}); }); +} + +// CHECK: void @_ZTS4Test() [[ATTRS:#[0-9]+]] +// CHECK: [[ATTRS]] = { +// NOAVX-NOT: "target-features"="{{[^"]*}}+avx{{[^"]*}}" +// AVX-SAME: "target-features"="{{[^"]*}}+avx{{[^"]*}}" diff --git a/sycl/test/check_device_code/native_cpu/target-features.cpp b/sycl/test/check_device_code/native_cpu/target-features.cpp deleted file mode 100644 index 561c636808c2..000000000000 --- a/sycl/test/check_device_code/native_cpu/target-features.cpp +++ /dev/null @@ -1,18 +0,0 @@ -// REQUIRES: native_cpu -// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl -fsycl-targets=native_cpu -mllvm -sycl-native-dump-device-ir %s 2>&1 | FileCheck %s --check-prefixes=CHECK,NOAVX -// RUN: %clangxx --target=x86_64-unknown-linux-gnu -march=skylake -fsycl -fsycl-targets=native_cpu -mllvm -sycl-native-dump-device-ir %s 2>&1 | FileCheck %s --check-prefixes=CHECK,AVX -// RUN: %clangxx --target=x86_64-unknown-linux-gnu -mavx -fsycl -fsycl-targets=native_cpu -mllvm -sycl-native-dump-device-ir %s 2>&1 | FileCheck %s --check-prefixes=CHECK,AVX - -#include -using namespace sycl; - -class Test; -int main() { - sycl::queue deviceQueue; - deviceQueue.submit([&](handler &h) { h.single_task([=] {}); }); -} - -// CHECK: void @_ZTS4Test.NativeCPUKernel({{.*}} [[ATTRS:#[0-9]+]] -// CHECK: [[ATTRS]] = { -// NOAVX-NOT: "target-features"="{{[^"]*}}+avx{{[^"]*}}" -// AVX-SAME: "target-features"="{{[^"]*}}+avx{{[^"]*}}"