Skip to content

[SYCL] Fix one more bug in int-header generation #20706

New issue

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

Sign up for GitHub

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

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Nov 21, 2025
+42 −2
Merged

[SYCL] Fix one more bug in int-header generation #20706

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
9a01c12
[SYCL] Fix one more bug in int-header generation
AlexeySachkov Nov 20, 2025
99174d6
Add one more test case
AlexeySachkov Nov 21, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
9 changes: 7 additions & 2 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6814,13 +6814,18 @@ class FreeFunctionPrinter {
QualType T = Param->getType();
QualType CT = T.getCanonicalType();

auto *TST = dyn_cast<TemplateSpecializationType>(T.getTypePtr());
auto *CTST = dyn_cast<TemplateSpecializationType>(CT.getTypePtr());
const auto *TST = dyn_cast<TemplateSpecializationType>(T.getTypePtr());
const auto *CTST = dyn_cast<TemplateSpecializationType>(CT.getTypePtr());
if (!TST || !CTST) {
ParmListOstream << T.getAsString(Policy);
continue;
}

const TemplateSpecializationType *TSTAsNonAlias =
TST->getAsNonAliasTemplateSpecializationType();
Copy link
Contributor

@Fznamznon Fznamznon Nov 20, 2025

Choose a reason for hiding this comment

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

Does that work for nested aliases? i.e. alias for alias?

Copy link

@jiafuzha jiafuzha Nov 21, 2025

Choose a reason for hiding this comment

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

Just a note here. Nested alias is needed for pytorch. Please check below code from https://github.com/pytorch/pytorch/blob/a7f3b10866098c452d89cd7a30bc4ce5713b8319/aten/src/ATen/core/TensorAccessor.h#L45

template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
using TensorAccessor = torch::headeronly::detail::TensorAccessor<c10::IntArrayRef, T, N, PtrTraits, index_t>;
...
template<typename T, size_t N, template <typename U> class PtrTraits = DefaultPtrTraits, typename index_t = int64_t>
using GenericPackedTensorAccessor = torch::headeronly::detail::GenericPackedTensorAccessor<TensorAccessor<T, N-1, PtrTraits, index_t>, detail::IndexBoundsCheck<N, index_t>, T, N, PtrTraits, index_t>;

Copy link
Contributor Author

@AlexeySachkov AlexeySachkov Nov 21, 2025

Choose a reason for hiding this comment

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

Does that work for nested aliases? i.e. alias for alias?

Yes, looking at the implementation it unwraps all aliases, I've also added a test case in 99174d6

if (TSTAsNonAlias)
TST = TSTAsNonAlias;

TemplateName CTN = CTST->getTemplateName();
CTN.getAsTemplateDecl()->printQualifiedName(ParmListOstream);
ParmListOstream << "<";
Expand Down
25 changes: 25 additions & 0 deletions clang/test/CodeGenSYCL/free-function-kernel-type-alias-arg.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,11 +23,18 @@ struct Bar {};
template<typename T1>
using BarUsing = Bar<T1, float>;

template<typename T1, typename T2>
using BarUsing2 = Bar<Foo<T2>, T1>;

class Baz {
public:
using type = BarUsing<double>;
};

template <typename T1, typename T2, typename T3 = BarUsing2<T1, T2>,
typename T4 = BarUsing<T2>>
struct AliasAsDefaultArg {};

} // namespace ns

[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
Expand Down Expand Up @@ -57,7 +64,25 @@ template void bar_using(ns::BarUsing<int>);

// CHECK: template <typename T> void bar_using(ns::Bar<T, float>);

template<typename T1, typename T2>
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
void bar_using2(ns::BarUsing2<T1, T2> Arg) {}
template void bar_using2(ns::BarUsing2<int, float>);

// CHECK: template <typename T1, typename T2> void bar_using2(ns::Bar<ns::Foo<T2>, T1>);

[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
void baz_type(ns::Baz::type Arg) {}

// CHECK: void baz_type(ns::Bar<double, float> Arg);

#if 0
// This test case fails, but it is added here in advance to add a record of a
Copy link
Contributor Author

@AlexeySachkov AlexeySachkov Nov 20, 2025

Choose a reason for hiding this comment

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

Note for reviewers: this is an additional test case that I came up with, it was not (at least yet) reported by any customer. I do have a work-in-progress patch for it, but it requires bigger changes to address.

I'm adding the known failing test case right now, proceeding with the simpler patch to unblock the reporter. The extra issue that I found will be addressed in a separate PR

// known bug.
template<typename T1, typename T2>
[[__sycl_detail__::add_ir_attributes_function("sycl-nd-range-kernel", 2)]]
void alias_as_default_template_arg(ns::AliasAsDefaultArg<T1, T2> Arg) {}
template void alias_as_default_template_arg(ns::AliasAsDefaultArg<int, float>);

// CHECK-DISABLED: template <typename T1, typename T2> void alias_as_default_template_arg(ns::AliasAsDefaultArg<T1, T2, ns::Bar<ns::Foo<T2>, T1>, ns::Bar<T2, float>>);
#endif
Loading