Skip to content

[SYCL][ABI Break] Remove "cl" namespace #6518

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
steffenlarsen merged 12 commits into from
Aug 11, 2022
Merged

[SYCL][ABI Break] Remove "cl" namespace #6518

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
12 commits
Select commit Hold shift + click to select a range
ce49e08
[SYCL][ABI Break] Remove "cl" namespace
aelovikov-intel Jul 27, 2022
64d972a
Merge remote-tracking branch 'origin/sycl' into ns-cl
aelovikov-intel Aug 8, 2022
4e54eb1
Update a comment per code review
aelovikov-intel Aug 8, 2022
678f522
Remove TODOs regarding __host_std versioning
aelovikov-intel Aug 8, 2022
684216f
Align llvm-spirv/lib/SPIRV/SPIRVUtil.cpp change with upstream PR
aelovikov-intel Aug 8, 2022
412ce44
Merge remote-tracking branch 'origin/sycl' into ns-cl
aelovikov-intel Aug 8, 2022
f420c43
Delete clang/test/SemaSYCL/intel-fpga-pipeline-ast.cpp
aelovikov-intel Aug 8, 2022
9bcc2d4
Merge remote-tracking branch 'origin/sycl' into ns-cl
aelovikov-intel Aug 10, 2022
48b7958
Merge remote-tracking branch 'origin/sycl' into ns-cl
aelovikov-intel Aug 10, 2022
eaf3b39
Add a basic test
aelovikov-intel Aug 10, 2022
930fecd
Merge branch 'sycl' into ns-cl
aelovikov-intel Aug 10, 2022
1265739
Merge remote-tracking branch 'origin/sycl' into ns-cl
aelovikov-intel Aug 11, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
The table of contents is too big for display.
Diff view
Diff view
  •  
  •  
  •  
2 changes: 1 addition & 1 deletion clang/lib/CodeGen/CGSYCLRuntime.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,8 +44,8 @@ static bool isPFWI(const FunctionDecl &FD) {
if (!MD)
return false;
static std::array<Util::DeclContextDesc, 3> Scopes = {
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"},
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"},
Util::DeclContextDesc{clang::Decl::Kind::Namespace, "_V1"},
Comment on lines 46 to +48
Copy link
Contributor

@bader bader Aug 7, 2022

Choose a reason for hiding this comment

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

Do you we have SYCL 1.2.1 compatibility tests?
https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:headers-and-namespaces
We should have them in SYCL-CTS and make sure that DPC++ complies.

Copy link
Contributor Author

@aelovikov-intel aelovikov-intel Aug 8, 2022

Choose a reason for hiding this comment

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

Are we talking source-compatible or binary-compatible? The former is being done with namespace alias as part of this PR (I might need to add an explicit test for that though), while the latter shouldn't (can't?) be done.

Copy link
Contributor

@bader bader Aug 9, 2022

Choose a reason for hiding this comment

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

I think we are talking about source-compatibility. Yes, please, add SYCL-1.2.1 compatibility test to make sure that when <CL/sycl.hpp> is included, we can still use data types defined in cl::sycl namespace in device code.

Copy link
Contributor

@elizabethandrews elizabethandrews Aug 9, 2022

Choose a reason for hiding this comment

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

I think we would need to account for both scopes (here and in Sema) for compatibility right?

Copy link
Contributor

@elizabethandrews elizabethandrews Aug 9, 2022 •
edited
Loading

Choose a reason for hiding this comment

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

We need to remove these name based checks. Now that we're modifying namespaces, this isn't future proof at all. We should instead modify sycl_special attribute to accept an argument and use that to do these checks where required.

Copy link
Contributor Author

@aelovikov-intel aelovikov-intel Aug 9, 2022

Choose a reason for hiding this comment

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

I think we would need to account for both scopes (here and in Sema) for compatibility right?

I don't think we do, but I'd like to hear your reasoning why you think we might need that.

We need to remove these name based checks. ... We should instead modify sycl_special attribute...

That is way above my comfort level with the FE - I guess I'd need someone from FE to make that part of the change.
Also, IMO, it should be done in a separate patch (either before or after this PR).

Copy link
Contributor

@elizabethandrews elizabethandrews Aug 9, 2022

Choose a reason for hiding this comment

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

I don't think we do, but I'd like to hear your reasoning why you think we might need that.

If the cl/sycl header is used as @bader mentioned above and users declare types using cl::sycl::aspect, etc, wouldn't compiler checks like isDeviceAspectType fail since it checks hard-coded scope which is now sycl::_V1::aspect?

That is way above my comfort level with the FE - I guess I'd need someone from FE to make that part of the change.
Also, IMO, it should be done in a separate patch (either before or after this PR).

I can help with that. What is the timeline for this change like?

Copy link
Contributor Author

@aelovikov-intel aelovikov-intel Aug 9, 2022

Choose a reason for hiding this comment

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

Using cl/sycl.h doesn't change ABI, only introduce namespace alias (https://godbolt.org/z/x17Toq5vY).

I can help with that. What is the timeline for this change like?

I don't think there are hard requirement other than the 2023 release, but this PR gets merge conflicts pretty easily and I really wouldn't want it to sit uncommitted for long.

Copy link
Contributor

@elizabethandrews elizabethandrews Aug 9, 2022

Choose a reason for hiding this comment

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

Ok Thanks for explaining. I guess this should work then, as long as we don't use older versions of cl/sycl.h

Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, "group"}};
if (!Util::matchQualifiedTypeName(MD->getParent(), Scopes))
return false;
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/Sema/SemaDeclAttr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10426,8 +10426,8 @@ bool isDeviceAspectType(const QualType Ty) {
return false;

std::array<std::pair<Decl::Kind, StringRef>, 3> Scopes = {
MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
MakeDeclContextDesc(Decl::Kind::Enum, "aspect")};

const auto *Ctx = cast<DeclContext>(ET->getDecl());
Expand Down
28 changes: 14 additions & 14 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -925,13 +925,13 @@ class MarkWIScopeFnVisitor : public RecursiveASTVisitor<MarkWIScopeFnVisitor> {
return true;
QualType Ty = Ctx.getRecordType(Call->getRecordDecl());
if (!Util::isSyclType(Ty, "group", true /*Tmpl*/))
// not a member of cl::sycl::group - continue search
// not a member of sycl::group - continue search
return true;
auto Name = Callee->getName();
if (((Name != "parallel_for_work_item") && (Name != "wait_for")) ||
Callee->hasAttr<SYCLScopeAttr>())
return true;
// it is a call to cl::sycl::group::parallel_for_work_item/wait_for -
// it is a call to sycl::group::parallel_for_work_item/wait_for -
// mark the callee
Callee->addAttr(
SYCLScopeAttr::CreateImplicit(Ctx, SYCLScopeAttr::Level::WorkItem));
Expand Down Expand Up @@ -4215,7 +4215,7 @@ static const char *paramKind2Str(KernelParamKind K) {
// VB,
// std::array<T1, N>& VC, int param, T2 ... varargs) {
// ...
// deviceQueue.submit([&](cl::sycl::handler& cgh) {
// deviceQueue.submit([&](sycl::handler& cgh) {
// ...
// cgh.parallel_for<class SimpleVadd<T1, N, T2...>>(...)
// ...
Expand Down Expand Up @@ -4650,8 +4650,8 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
FwdDeclEmitter.Visit(K.NameType);
O << "\n";

O << "__SYCL_INLINE_NAMESPACE(cl) {\n";
O << "namespace sycl {\n";
O << "__SYCL_INLINE_VER_NAMESPACE(_V1) {\n";
O << "namespace detail {\n";

// Generate declaration of variable of type __sycl_device_global_registration
Expand Down Expand Up @@ -4795,8 +4795,8 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
}
O << "\n";
O << "} // namespace detail\n";
O << "} // __SYCL_INLINE_VER_NAMESPACE(_V1)\n";
O << "} // namespace sycl\n";
O << "} // __SYCL_INLINE_NAMESPACE(cl)\n";
O << "\n";
}

Expand Down Expand Up @@ -5076,8 +5076,8 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
DeviceGlobOS << "\");\n";
} else {
EmittedFirstSpecConstant = true;
OS << "__SYCL_INLINE_NAMESPACE(cl) {\n";
OS << "namespace sycl {\n";
OS << "__SYCL_INLINE_VER_NAMESPACE(_V1) {\n";
OS << "namespace detail {\n";
OS << "template<>\n";
OS << "inline const char *get_spec_constant_symbolic_ID_impl<";
Expand All @@ -5095,8 +5095,8 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
OS << "\";\n";
OS << "}\n";
OS << "} // namespace detail\n";
OS << "} // __SYCL_INLINE_VER_NAMESPACE(_V1)\n";
OS << "} // namespace sycl\n";
OS << "} // __SYCL_INLINE_NAMESPACE(cl)\n";
}
}

Expand Down Expand Up @@ -5132,8 +5132,8 @@ bool Util::isSyclSpecialType(const QualType Ty) {

bool Util::isSyclSpecConstantType(QualType Ty) {
std::array<DeclContextDesc, 6> Scopes = {
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "oneapi"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "experimental"),
Expand All @@ -5144,25 +5144,25 @@ bool Util::isSyclSpecConstantType(QualType Ty) {

bool Util::isSyclSpecIdType(QualType Ty) {
std::array<DeclContextDesc, 3> Scopes = {
Util::MakeDeclContextDesc(clang::Decl::Kind::Namespace, "cl"),
Util::MakeDeclContextDesc(clang::Decl::Kind::Namespace, "sycl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization,
"specialization_id")};
return matchQualifiedTypeName(Ty, Scopes);
}

bool Util::isSyclKernelHandlerType(QualType Ty) {
std::array<DeclContextDesc, 3> Scopes = {
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
Util::MakeDeclContextDesc(Decl::Kind::CXXRecord, "kernel_handler")};
return matchQualifiedTypeName(Ty, Scopes);
}

bool Util::isSyclAccessorNoAliasPropertyType(QualType Ty) {
std::array<DeclContextDesc, 7> Scopes = {
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "oneapi"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "property"),
Expand All @@ -5174,8 +5174,8 @@ bool Util::isSyclAccessorNoAliasPropertyType(QualType Ty) {

bool Util::isSyclBufferLocationType(QualType Ty) {
std::array<DeclContextDesc, 7> Scopes = {
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "intel"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "property"),
Expand All @@ -5189,16 +5189,16 @@ bool Util::isSyclType(QualType Ty, StringRef Name, bool Tmpl) {
Decl::Kind ClassDeclKind =
Tmpl ? Decl::Kind::ClassTemplateSpecialization : Decl::Kind::CXXRecord;
std::array<DeclContextDesc, 3> Scopes = {
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
Util::MakeDeclContextDesc(ClassDeclKind, Name)};
return matchQualifiedTypeName(Ty, Scopes);
}

bool Util::isAccessorPropertyListType(QualType Ty) {
std::array<DeclContextDesc, 5> Scopes = {
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "cl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "sycl"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "_V1"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "ext"),
Util::MakeDeclContextDesc(Decl::Kind::Namespace, "oneapi"),
Util::MakeDeclContextDesc(Decl::Kind::ClassTemplateSpecialization,
Expand Down
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/Inputs/sycl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,8 +10,8 @@ __spirv_ControlBarrier(int, int, int) noexcept;
#endif

// Dummy runtime classes to model SYCL API.
inline namespace cl {
namespace sycl {
inline namespace _V1 {
struct sampler_impl {
#ifdef __SYCL_DEVICE_ONLY__
__ocl_sampler_t m_Sampler;
Expand Down Expand Up @@ -491,7 +491,7 @@ class __attribute__((sycl_special_class)) stream {
void __finalize() {}

private:
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read_write> Acc;
sycl::accessor<char, 1, sycl::access::mode::read_write> Acc;
int FlushBufferSize;
};

Expand Down Expand Up @@ -621,5 +621,5 @@ class image {
}
};

} // inline namespace _V1
} // namespace sycl
} // namespace cl
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/accessor-readonly-invalid-lib.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,8 +3,8 @@
// Test which verifies that readonly attribute is generated for unexpected access mode value.

// Dummy library with unexpected access::mode enum value.
namespace cl {
namespace sycl {
inline namespace _V1 {

namespace access {

Expand Down Expand Up @@ -71,16 +71,16 @@ class __attribute__((sycl_special_class)) accessor {
range<dimensions> MemRange, id<dimensions> Offset) {}
};

} // inline namespace _V1
} // namespace sycl
} // namespace cl

template <typename name, typename Func>
__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
kernelFunc();
}

int main() {
cl::sycl::accessor<int, 1, cl::sycl::access::mode::read> Acc;
sycl::accessor<int, 1, sycl::access::mode::read> Acc;
// CHECK: spir_kernel{{.*}}fake_kernel
// CHECK-SAME: readonly
kernel_single_task<class fake_kernel>([=]() {
Expand Down
22 changes: 11 additions & 11 deletions clang/test/CodeGenSYCL/accessor-readonly.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,9 +5,9 @@
#include "Inputs/sycl.hpp"

// CHECK-NOT: spir_kernel{{.*}}f0_kernel{{.*}}readonly
void f0(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::buffer<int, 1> &out_buf) {
myQueue.submit([&](cl::sycl::handler &cgh) {
auto write_acc = out_buf.get_access<cl::sycl::access::mode::write>(cgh);
void f0(sycl::queue &myQueue, sycl::buffer<int, 1> &in_buf, sycl::buffer<int, 1> &out_buf) {
myQueue.submit([&](sycl::handler &cgh) {
auto write_acc = out_buf.get_access<sycl::access::mode::write>(cgh);
cgh.single_task<class f0_kernel>([write_acc] {});
});
}
Expand All @@ -16,10 +16,10 @@ void f0(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::bu
// CHECK-NOT: readonly
// CHECK-SAME: %_arg_write_acc{{.*}}%_arg_write_acc1{{.*}}%_arg_write_acc2{{.*}}%_arg_write_acc3
// CHECK-SAME: readonly align 4 %_arg_read_acc
void f1(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::buffer<int, 1> &out_buf) {
myQueue.submit([&](cl::sycl::handler &cgh) {
auto write_acc = out_buf.get_access<cl::sycl::access::mode::write>(cgh);
auto read_acc = in_buf.get_access<cl::sycl::access::mode::read>(cgh);
void f1(sycl::queue &myQueue, sycl::buffer<int, 1> &in_buf, sycl::buffer<int, 1> &out_buf) {
myQueue.submit([&](sycl::handler &cgh) {
auto write_acc = out_buf.get_access<sycl::access::mode::write>(cgh);
auto read_acc = in_buf.get_access<sycl::access::mode::read>(cgh);
cgh.single_task<class f1_kernel>([write_acc, read_acc] {});
});
}
Expand All @@ -28,10 +28,10 @@ void f1(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::bu
// CHECK-SAME: readonly align 4 %_arg_read_acc
// CHECK-NOT: readonly
// CHECK-SAME: %_arg_write_acc
void f2(cl::sycl::queue &myQueue, cl::sycl::buffer<int, 1> &in_buf, cl::sycl::buffer<int, 1> &out_buf) {
myQueue.submit([&](cl::sycl::handler &cgh) {
auto read_acc = in_buf.get_access<cl::sycl::access::mode::read>(cgh);
auto write_acc = out_buf.get_access<cl::sycl::access::mode::write>(cgh);
void f2(sycl::queue &myQueue, sycl::buffer<int, 1> &in_buf, sycl::buffer<int, 1> &out_buf) {
myQueue.submit([&](sycl::handler &cgh) {
auto read_acc = in_buf.get_access<sycl::access::mode::read>(cgh);
auto write_acc = out_buf.get_access<sycl::access::mode::write>(cgh);
cgh.single_task<class f2_kernel>([read_acc, write_acc] {});
});
}
14 changes: 7 additions & 7 deletions clang/test/CodeGenSYCL/accessor_inheritance.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,26 +3,26 @@

struct Base {
int A, B;
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> AccField;
sycl::accessor<char, 1, sycl::access::mode::read> AccField;
};

struct Captured : Base,
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> {
sycl::accessor<char, 1, sycl::access::mode::read> {
int C;
};

int main() {
Captured Obj;
cl::sycl::kernel_single_task<class kernel>(
sycl::kernel_single_task<class kernel>(
[=]() {
Obj.use();
});
return 0;
}

// Check kernel parameters
// CHECK: %[[RANGE_TYPE:"struct.*cl::sycl::range"]]
// CHECK: %[[ID_TYPE:"struct.*cl::sycl::id"]]
// CHECK: %[[RANGE_TYPE:"struct.*sycl::_V1::range"]]
// CHECK: %[[ID_TYPE:"struct.*sycl::_V1::id"]]
// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE6kernel
// CHECK-SAME: i32 noundef [[ARG_A:%[a-zA-Z0-9_]+]],
// CHECK-SAME: i32 noundef [[ARG_B:%[a-zA-Z0-9_]+]],
Expand Down Expand Up @@ -71,10 +71,10 @@ int main() {
// Check accessors initialization
// CHECK: [[ACC_FIELD:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Base, ptr addrspace(4) [[GEP]], i32 0, i32 2
// Default constructor call
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(ptr addrspace(4) {{[^,]*}} [[ACC_FIELD]])
// CHECK: call spir_func void @_ZN4sycl3_V18accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC1Ev(ptr addrspace(4) {{[^,]*}} [[ACC_FIELD]])
// CHECK: [[GEP1:%[a-zA-Z0-9_]+]] = getelementptr inbounds i8, ptr addrspace(4) [[GEP]], i64 20
// Default constructor call
// CHECK: call spir_func void @_ZN2cl4sycl8accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC2Ev(ptr addrspace(4) {{[^,]*}} [[GEP1]])
// CHECK: call spir_func void @_ZN4sycl3_V18accessorIcLi1ELNS0_6access4modeE1024ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEC2Ev(ptr addrspace(4) {{[^,]*}} [[GEP1]])

// CHECK C field initialization
// CHECK: [[FIELD_C:%[a-zA-Z0-9_]+]] = getelementptr inbounds %struct{{.*}}Captured, ptr addrspace(4) [[GEP]], i32 0, i32 2
Expand Down
24 changes: 12 additions & 12 deletions clang/test/CodeGenSYCL/accessor_no_alias_property.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,23 +5,23 @@
#include "Inputs/sycl.hpp"

int main() {
cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
cl::sycl::access::target::global_buffer,
cl::sycl::access::placeholder::false_t,
cl::sycl::ext::oneapi::accessor_property_list<
cl::sycl::ext::oneapi::property::no_alias::instance<true>>>
sycl::accessor<int, 1, sycl::access::mode::read_write,
sycl::access::target::global_buffer,
sycl::access::placeholder::false_t,
sycl::ext::oneapi::accessor_property_list<
sycl::ext::oneapi::property::no_alias::instance<true>>>
accessorA;

cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
cl::sycl::access::target::global_buffer,
cl::sycl::access::placeholder::false_t,
cl::sycl::ext::oneapi::accessor_property_list<
cl::sycl::ext::intel::property::buffer_location::instance<1>>>
sycl::accessor<int, 1, sycl::access::mode::read_write,
sycl::access::target::global_buffer,
sycl::access::placeholder::false_t,
sycl::ext::oneapi::accessor_property_list<
sycl::ext::intel::property::buffer_location::instance<1>>>
accessorB;

// Check that noalias parameter attribute is emitted when no_alias accessor property is used
// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE16kernel_function1({{.*}} noalias {{.*}} %_arg_accessorA, {{.*}})
cl::sycl::kernel_single_task<class kernel_function1>(
sycl::kernel_single_task<class kernel_function1>(
[=]() {
accessorA.use();
});
Expand All @@ -30,7 +30,7 @@ int main() {
// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE16kernel_function2
// CHECK-NOT: noalias
// CHECK-SAME: {
cl::sycl::kernel_single_task<class kernel_function2>(
sycl::kernel_single_task<class kernel_function2>(
[=]() {
accessorB.use();
});
Expand Down
Loading