From 185db16e37211eb65b6be785a0adaf18d2ce91a4 Mon Sep 17 00:00:00 2001 From: "Podchishchaeva, Mariya" Date: Fri, 22 Nov 2024 06:22:29 -0800 Subject: [PATCH 1/2] [clang][SYCL] Emit canonical types in integration footer A recent change to sycl headers revealed a bug in integration footer - missing nns in case of an explicit template specialization in some of device_globa's template arguments. --- clang/lib/Sema/SemaSYCL.cpp | 1 + ...nt_footer_with_explicit_specialization.cpp | 29 +++++++++++++++++++ 2 files changed, 30 insertions(+) create mode 100644 clang/test/CodeGenSYCL/int_footer_with_explicit_specialization.cpp diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 242e6c8a9d7d4..ad7da52d635d0 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -6969,6 +6969,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) { Policy.adjustForCPlusPlusFwdDecl(); Policy.SuppressTypedefs = true; Policy.SuppressUnwrittenScope = true; + Policy.PrintCanonicalTypes = true; llvm::SmallSet Visited; bool EmittedFirstSpecConstant = false; diff --git a/clang/test/CodeGenSYCL/int_footer_with_explicit_specialization.cpp b/clang/test/CodeGenSYCL/int_footer_with_explicit_specialization.cpp new file mode 100644 index 0000000000000..7b98a70e165c9 --- /dev/null +++ b/clang/test/CodeGenSYCL/int_footer_with_explicit_specialization.cpp @@ -0,0 +1,29 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -emit-llvm %s -o - +// RUN: FileCheck -input-file=%t.footer.h %s --check-prefix=CHECK-FOOTER + +#include "sycl.hpp" + +namespace sycl { +template struct X {}; +template <> struct X {}; +namespace detail { +struct Y {}; +} // namespace detail +template <> struct X {}; +} // namespace sycl + +using namespace sycl; +template > struct Arg1 { T val; }; + +using namespace sycl::ext::oneapi; +template +device_global dev_global; + +SYCL_EXTERNAL auto foo() { + (void)dev_global>; +} + +// CHECK-FOOTER: __sycl_device_global_registration::__sycl_device_global_registration() noexcept { +// CHECK-FOOTER-NEXT: device_global_map::add((void *)&::dev_global>>, "_Z10dev_globalI4Arg1IiN4sycl1XINS1_6detail1YEEEEE"); +// CHECK-FOOTER-NEXT: } +// CHECK-FOOTER-NEXT: } // namespace (unnamed) From bb9843e268d7091260b92d073388bd97108e28fe Mon Sep 17 00:00:00 2001 From: "Podchishchaeva, Mariya" Date: Mon, 25 Nov 2024 01:22:51 -0800 Subject: [PATCH 2/2] Add test comment --- .../CodeGenSYCL/int_footer_with_explicit_specialization.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/clang/test/CodeGenSYCL/int_footer_with_explicit_specialization.cpp b/clang/test/CodeGenSYCL/int_footer_with_explicit_specialization.cpp index 7b98a70e165c9..ec066b4bdaec1 100644 --- a/clang/test/CodeGenSYCL/int_footer_with_explicit_specialization.cpp +++ b/clang/test/CodeGenSYCL/int_footer_with_explicit_specialization.cpp @@ -1,6 +1,9 @@ // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -fsycl-int-footer=%t.footer.h -emit-llvm %s -o - // RUN: FileCheck -input-file=%t.footer.h %s --check-prefix=CHECK-FOOTER +// This test checks that integration footer is emitted correctly when a +// device_global has an explicit template specialization in template arguments. + #include "sycl.hpp" namespace sycl {