From d7e3e046b2e4cf6f2ffcc9bb110bb62e3d284d31 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Mon, 29 Mar 2021 15:47:12 -0700 Subject: [PATCH 01/12] 25700: Add clang implementation for accessor property no_alias --- clang/lib/CodeGen/CodeGenModule.cpp | 7 ++++ clang/lib/Sema/SemaSYCL.cpp | 34 +++++++++++++++++++ clang/test/CodeGenSYCL/Inputs/sycl.hpp | 9 +++++ .../accessor_no_alias_property.cpp | 20 +++++++++++ 4 files changed, 70 insertions(+) create mode 100644 clang/test/CodeGenSYCL/accessor_no_alias_property.cpp diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index be092a0fd97b2..0fd3889c81f53 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1580,6 +1580,9 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, // MDNode for the intel_buffer_location attribute. SmallVector argSYCLBufferLocationAttr; + // MDNode for accessor no_alias property + SmallVector argSYCLIntelAccessorNoAliasProperty; + // MDNode for listing ESIMD kernel pointer arguments originating from // accessors SmallVector argESIMDAccPtrs; @@ -1687,6 +1690,10 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, SYCLBufferLocationAttr->getLocationID())) : llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(-1))); + //Sindhu + //auto *SYCLIntelAccessorNoAliasPropertyAttr = + // parm->getAttr(); + if (FD->hasAttr()) argESIMDAccPtrs.push_back(llvm::ConstantAsMetadata::get( CGF->Builder.getInt1(parm->hasAttr()))); diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index a76d5c3bc7adb..2ac5a893165c2 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -90,6 +90,10 @@ class Util { /// accessor_property_list class. static bool isAccessorPropertyListType(const QualType &Ty); + /// Checks whether given clang type is a full specialization of the SYCL + /// no_alias class. + static bool isSyclAccessorNoAliasPropertyType(const QualType &Ty); + /// Checks whether given clang type is a full specialization of the SYCL /// buffer_location class. static bool isSyclBufferLocationType(const QualType &Ty); @@ -1742,11 +1746,27 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin(); Prop != TemplArg.pack_end(); ++Prop) { QualType PropTy = Prop->getAsType(); + ASTContext &Ctx = SemaRef.getASTContext(); + //const auto *PropDecl = + // cast(PropTy->getAsRecordDecl()); + //const auto NoAliasLoc = PropDecl->getTemplateArgs()[0]; + //int LocationID = static_cast(NoAliasLoc.getAsIntegral().getExtValue()); + + if (Util::isSyclAccessorNoAliasPropertyType(PropTy)) + Param->addAttr(SYCLIntelKernelArgsRestrictAttr::CreateImplicit(Ctx, Loc)); + + //handleNoAliasProperty(Param, PropTy, Loc); if (Util::isSyclBufferLocationType(PropTy)) handleBufferLocationProperty(Param, PropTy, Loc); } } + void handleNoAliasProperty(ParmVarDecl *Param, QualType PropTy, + SourceLocation Loc) { + //if (Param->hasAttr) + // Param->addAttr(SYCLIntelKernelArgsRestrictAttr);::CreateImplicit(Ctx + + } // Obtain an integer value stored in a template parameter of buffer_location // property to pass it to buffer_location kernel attribute void handleBufferLocationProperty(ParmVarDecl *Param, QualType PropTy, @@ -4313,6 +4333,20 @@ bool Util::isSyclKernelHandlerType(const QualType &Ty) { return matchQualifiedTypeName(Ty, Scopes); } +bool Util::isSyclAccessorNoAliasPropertyType(const QualType &Ty) { + const StringRef &PropertyName = "no_alias"; + const StringRef &InstanceName = "instance"; + std::array Scopes = { + Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"}, + Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "ONEAPI"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "property"}, + Util::DeclContextDesc{Decl::Kind::CXXRecord, PropertyName}, + Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, + InstanceName}}; + return matchQualifiedTypeName(Ty, Scopes); +} + bool Util::isSyclBufferLocationType(const QualType &Ty) { const StringRef &PropertyName = "buffer_location"; const StringRef &InstanceName = "instance"; diff --git a/clang/test/CodeGenSYCL/Inputs/sycl.hpp b/clang/test/CodeGenSYCL/Inputs/sycl.hpp index fc3ca2c146ad6..aa38dbd9f15ef 100644 --- a/clang/test/CodeGenSYCL/Inputs/sycl.hpp +++ b/clang/test/CodeGenSYCL/Inputs/sycl.hpp @@ -103,6 +103,15 @@ struct buffer_location { } // namespace property } // namespace INTEL +namespace ONEAPI { +namespace property { +// Compile time known accessor property +struct no_alias { + template class instance {}; +}; +} // namespace property +} // namespace ONEAPI + namespace ONEAPI { template class accessor_property_list {}; diff --git a/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp b/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp new file mode 100644 index 0000000000000..1b47d5f07865e --- /dev/null +++ b/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp @@ -0,0 +1,20 @@ +// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s + +// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function{{.*}} !kernel_arg_buffer_location ![[MDBL:[0-9]+]] +// CHECK: ![[MDBL]] = !{i32 3, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 2, i32 -1, i32 -1, i32 -1, i32 2, i32 -1, i32 -1, i32 -1, i32 -1} + +#include "Inputs/sycl.hpp" + +int main() { + cl::sycl::accessor>> + accessorA; + cl::sycl::kernel_single_task( + [=]() { + accessorA.use(); + }); + return 0; +} From 0b55f906ff29a6b67a3164154eb079a4457ac67f Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 1 Apr 2021 10:22:45 -0700 Subject: [PATCH 02/12] WIP edit 2 --- clang/lib/CodeGen/CodeGenModule.cpp | 18 ++++++++++------- clang/lib/Sema/SemaSYCL.cpp | 20 +++++++------------ .../accessor_no_alias_property.cpp | 3 +-- 3 files changed, 19 insertions(+), 22 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 0fd3889c81f53..9b2bd2ffbcd69 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1581,7 +1581,7 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, SmallVector argSYCLBufferLocationAttr; // MDNode for accessor no_alias property - SmallVector argSYCLIntelAccessorNoAliasProperty; + SmallVector argAccessorNoAliasPropertyAttr; // MDNode for listing ESIMD kernel pointer arguments originating from // accessors @@ -1690,10 +1690,10 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, SYCLBufferLocationAttr->getLocationID())) : llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(-1))); - //Sindhu - //auto *SYCLIntelAccessorNoAliasPropertyAttr = - // parm->getAttr(); - + if (parm->hasAttr()) + argAccessorNoAliasPropertyAttr.push_back(llvm::ConstantAsMetadata::get( + CGF->Builder.getInt1(parm->hasAttr()))); + if (FD->hasAttr()) argESIMDAccPtrs.push_back(llvm::ConstantAsMetadata::get( CGF->Builder.getInt1(parm->hasAttr()))); @@ -1701,10 +1701,14 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, bool IsEsimdFunction = FD && FD->hasAttr(); - if (LangOpts.SYCLIsDevice && !IsEsimdFunction) + if (LangOpts.SYCLIsDevice && !IsEsimdFunction) { + + Fn->setMetadata(llvm::Attribute::NoAlias, + llvm::MDNode::get(VMContext, argAccessorNoAliasPropertyAttr)); + Fn->setMetadata("kernel_arg_buffer_location", llvm::MDNode::get(VMContext, argSYCLBufferLocationAttr)); - else { + } else { Fn->setMetadata("kernel_arg_addr_space", llvm::MDNode::get(VMContext, addressQuals)); Fn->setMetadata("kernel_arg_access_qual", diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 2ac5a893165c2..bdcfdae68eff7 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1746,27 +1746,21 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin(); Prop != TemplArg.pack_end(); ++Prop) { QualType PropTy = Prop->getAsType(); - ASTContext &Ctx = SemaRef.getASTContext(); - //const auto *PropDecl = - // cast(PropTy->getAsRecordDecl()); - //const auto NoAliasLoc = PropDecl->getTemplateArgs()[0]; - //int LocationID = static_cast(NoAliasLoc.getAsIntegral().getExtValue()); - if (Util::isSyclAccessorNoAliasPropertyType(PropTy)) - Param->addAttr(SYCLIntelKernelArgsRestrictAttr::CreateImplicit(Ctx, Loc)); - - //handleNoAliasProperty(Param, PropTy, Loc); + handleNoAliasProperty(Param, PropTy, Loc); if (Util::isSyclBufferLocationType(PropTy)) handleBufferLocationProperty(Param, PropTy, Loc); } } - void handleNoAliasProperty(ParmVarDecl *Param, QualType PropTy, + void handleNoAliasProperty(ParmVarDecl *Param, QualType PropTy, SourceLocation Loc) { - //if (Param->hasAttr) - // Param->addAttr(SYCLIntelKernelArgsRestrictAttr);::CreateImplicit(Ctx - + if (PropTy.isRestrictQualified()) { + ASTContext &Ctx = SemaRef.getASTContext(); + Param->addAttr(RestrictAttr::CreateImplicit(Ctx, Loc)); + } } + // Obtain an integer value stored in a template parameter of buffer_location // property to pass it to buffer_location kernel attribute void handleBufferLocationProperty(ParmVarDecl *Param, QualType PropTy, diff --git a/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp b/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp index 1b47d5f07865e..54a708c3256db 100644 --- a/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp +++ b/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp @@ -1,7 +1,6 @@ // RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s -// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function{{.*}} !kernel_arg_buffer_location ![[MDBL:[0-9]+]] -// CHECK: ![[MDBL]] = !{i32 3, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 2, i32 -1, i32 -1, i32 -1, i32 2, i32 -1, i32 -1, i32 -1, i32 -1} +// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function{{.*}} !noalias #include "Inputs/sycl.hpp" From c7441a10eec555c78874e321ea4667ecec4a7eab Mon Sep 17 00:00:00 2001 From: schittir Date: Thu, 1 Apr 2021 12:01:54 -0700 Subject: [PATCH 03/12] Update clang/lib/CodeGen/CodeGenModule.cpp Co-authored-by: Aaron Ballman --- clang/lib/CodeGen/CodeGenModule.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 9b2bd2ffbcd69..60553acc6794a 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1580,7 +1580,7 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, // MDNode for the intel_buffer_location attribute. SmallVector argSYCLBufferLocationAttr; - // MDNode for accessor no_alias property + // MDNode for accessor no_alias property. SmallVector argAccessorNoAliasPropertyAttr; // MDNode for listing ESIMD kernel pointer arguments originating from From 5b3446d6982af82e58f5f88e15c3a9de3fefc6bf Mon Sep 17 00:00:00 2001 From: schittir Date: Thu, 1 Apr 2021 12:05:00 -0700 Subject: [PATCH 04/12] Update clang/lib/CodeGen/CodeGenModule.cpp Co-authored-by: Aaron Ballman --- clang/lib/CodeGen/CodeGenModule.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 60553acc6794a..103cefb9fd7c3 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1691,8 +1691,8 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, : llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(-1))); if (parm->hasAttr()) - argAccessorNoAliasPropertyAttr.push_back(llvm::ConstantAsMetadata::get( - CGF->Builder.getInt1(parm->hasAttr()))); + argAccessorNoAliasPropertyAttr.push_back(llvm::ConstantAsMetadata::get( + CGF->Builder.getInt1(true))); if (FD->hasAttr()) argESIMDAccPtrs.push_back(llvm::ConstantAsMetadata::get( From 57ca989796fa9082ed945a8fbb19bca2b09180a7 Mon Sep 17 00:00:00 2001 From: schittir Date: Thu, 1 Apr 2021 12:05:39 -0700 Subject: [PATCH 05/12] Update clang/lib/Sema/SemaSYCL.cpp Co-authored-by: Aaron Ballman --- clang/lib/Sema/SemaSYCL.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index bdcfdae68eff7..d6b5ef8399d07 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -4331,8 +4331,8 @@ bool Util::isSyclAccessorNoAliasPropertyType(const QualType &Ty) { const StringRef &PropertyName = "no_alias"; const StringRef &InstanceName = "instance"; std::array Scopes = { - Util::DeclContextDesc{clang::Decl::Kind::Namespace, "cl"}, - Util::DeclContextDesc{clang::Decl::Kind::Namespace, "sycl"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "cl"}, + Util::DeclContextDesc{Decl::Kind::Namespace, "sycl"}, Util::DeclContextDesc{Decl::Kind::Namespace, "ONEAPI"}, Util::DeclContextDesc{Decl::Kind::Namespace, "property"}, Util::DeclContextDesc{Decl::Kind::CXXRecord, PropertyName}, From 9bbe8f97e26c494bfb3e9a5c50a737d6f8f69af3 Mon Sep 17 00:00:00 2001 From: schittir Date: Thu, 1 Apr 2021 12:09:47 -0700 Subject: [PATCH 06/12] Update clang/lib/Sema/SemaSYCL.cpp Co-authored-by: Aaron Ballman --- clang/lib/Sema/SemaSYCL.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index d6b5ef8399d07..86e660ad64a76 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -4327,9 +4327,9 @@ bool Util::isSyclKernelHandlerType(const QualType &Ty) { return matchQualifiedTypeName(Ty, Scopes); } -bool Util::isSyclAccessorNoAliasPropertyType(const QualType &Ty) { - const StringRef &PropertyName = "no_alias"; - const StringRef &InstanceName = "instance"; +bool Util::isSyclAccessorNoAliasPropertyType(QualType Ty) { + StringRef PropertyName = "no_alias"; + StringRef InstanceName = "instance"; std::array Scopes = { Util::DeclContextDesc{Decl::Kind::Namespace, "cl"}, Util::DeclContextDesc{Decl::Kind::Namespace, "sycl"}, From c38d9f4d21309d6da362193ef0b8263bbfd7b4e0 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 1 Apr 2021 15:50:41 -0700 Subject: [PATCH 07/12] 1. Fix format 2. replace "llvm::Attribute::NoAlias with "noalias" 3. replace variable names with string literals --- clang/lib/CodeGen/CodeGenModule.cpp | 4 ++-- clang/lib/Sema/SemaSYCL.cpp | 10 ++++------ 2 files changed, 6 insertions(+), 8 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 103cefb9fd7c3..578ec53306491 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1703,8 +1703,8 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, if (LangOpts.SYCLIsDevice && !IsEsimdFunction) { - Fn->setMetadata(llvm::Attribute::NoAlias, - llvm::MDNode::get(VMContext, argAccessorNoAliasPropertyAttr)); + Fn->setMetadata("noalias", llvm::MDNode::get( + VMContext, argAccessorNoAliasPropertyAttr)); Fn->setMetadata("kernel_arg_buffer_location", llvm::MDNode::get(VMContext, argSYCLBufferLocationAttr)); diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 86e660ad64a76..69e23e9dce64f 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -92,7 +92,7 @@ class Util { /// Checks whether given clang type is a full specialization of the SYCL /// no_alias class. - static bool isSyclAccessorNoAliasPropertyType(const QualType &Ty); + static bool isSyclAccessorNoAliasPropertyType(QualType Ty); /// Checks whether given clang type is a full specialization of the SYCL /// buffer_location class. @@ -1754,7 +1754,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { } void handleNoAliasProperty(ParmVarDecl *Param, QualType PropTy, - SourceLocation Loc) { + SourceLocation Loc) { if (PropTy.isRestrictQualified()) { ASTContext &Ctx = SemaRef.getASTContext(); Param->addAttr(RestrictAttr::CreateImplicit(Ctx, Loc)); @@ -4328,16 +4328,14 @@ bool Util::isSyclKernelHandlerType(const QualType &Ty) { } bool Util::isSyclAccessorNoAliasPropertyType(QualType Ty) { - StringRef PropertyName = "no_alias"; - StringRef InstanceName = "instance"; std::array Scopes = { Util::DeclContextDesc{Decl::Kind::Namespace, "cl"}, Util::DeclContextDesc{Decl::Kind::Namespace, "sycl"}, Util::DeclContextDesc{Decl::Kind::Namespace, "ONEAPI"}, Util::DeclContextDesc{Decl::Kind::Namespace, "property"}, - Util::DeclContextDesc{Decl::Kind::CXXRecord, PropertyName}, + Util::DeclContextDesc{Decl::Kind::CXXRecord, "no_alias"}, Util::DeclContextDesc{Decl::Kind::ClassTemplateSpecialization, - InstanceName}}; + "instance"}}; return matchQualifiedTypeName(Ty, Scopes); } From 75b3fb11e6930a0a296a7e9127d63c4296a680b9 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Mon, 5 Apr 2021 13:42:38 -0700 Subject: [PATCH 08/12] Fix formatting again --- clang/lib/CodeGen/CodeGenModule.cpp | 4 ++-- clang/lib/Sema/SemaSYCL.cpp | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 578ec53306491..775f12e86b3e6 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1691,8 +1691,8 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, : llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(-1))); if (parm->hasAttr()) - argAccessorNoAliasPropertyAttr.push_back(llvm::ConstantAsMetadata::get( - CGF->Builder.getInt1(true))); + argAccessorNoAliasPropertyAttr.push_back( + llvm::ConstantAsMetadata::get(CGF->Builder.getInt1(true))); if (FD->hasAttr()) argESIMDAccPtrs.push_back(llvm::ConstantAsMetadata::get( diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 69e23e9dce64f..3712cef2a6a25 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1746,7 +1746,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin(); Prop != TemplArg.pack_end(); ++Prop) { QualType PropTy = Prop->getAsType(); - if (Util::isSyclAccessorNoAliasPropertyType(PropTy)) + if (Util::isSyclAccessorNoAliasPropertyType(PropTy)) handleNoAliasProperty(Param, PropTy, Loc); if (Util::isSyclBufferLocationType(PropTy)) handleBufferLocationProperty(Param, PropTy, Loc); From d2642be62b39c2b05102002398bd7a5457986151 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 6 Apr 2021 10:04:51 -0700 Subject: [PATCH 09/12] Fix failing lit tests --- clang/test/CodeGenSYCL/disable_loop_pipelining.cpp | 7 +++---- clang/test/CodeGenSYCL/initiation_interval.cpp | 9 ++++----- 2 files changed, 7 insertions(+), 9 deletions(-) diff --git a/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp b/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp index 33f83f6b1961f..85c51b090616c 100644 --- a/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp +++ b/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp @@ -29,8 +29,7 @@ int main() { return 0; } -// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel1"() #0 !kernel_arg_buffer_location ![[NUM4:[0-9]+]] !disable_loop_pipelining ![[NUM5:[0-9]+]] -// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel2"() #0 !kernel_arg_buffer_location ![[NUM4]] -// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel3"() #0 !kernel_arg_buffer_location ![[NUM4]] !disable_loop_pipelining ![[NUM5]] -// CHECK: ![[NUM4]] = !{} +// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel1"() #0 {{.*}} !disable_loop_pipelining ![[NUM5:[0-9]+]] +// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel2"() #0 {{.*}} +// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel3"() #0 {{.*}} !disable_loop_pipelining ![[NUM5]] // CHECK: ![[NUM5]] = !{i32 1} diff --git a/clang/test/CodeGenSYCL/initiation_interval.cpp b/clang/test/CodeGenSYCL/initiation_interval.cpp index 95ef9ce4cde50..b32cbb037f2ff 100644 --- a/clang/test/CodeGenSYCL/initiation_interval.cpp +++ b/clang/test/CodeGenSYCL/initiation_interval.cpp @@ -39,11 +39,10 @@ int main() { return 0; } -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 !kernel_arg_buffer_location ![[NUM0:[0-9]+]] !initiation_interval ![[NUM1:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 !kernel_arg_buffer_location ![[NUM0]] !initiation_interval ![[NUM42:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 !kernel_arg_buffer_location ![[NUM0]] !initiation_interval ![[NUM2:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 !kernel_arg_buffer_location ![[NUM0]] -// CHECK: ![[NUM0]] = !{} +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !initiation_interval ![[NUM1:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !initiation_interval ![[NUM42:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !initiation_interval ![[NUM2:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} // CHECK: ![[NUM1]] = !{i32 1} // CHECK: ![[NUM42]] = !{i32 42} // CHECK: ![[NUM2]] = !{i32 2} From ca912611a0735bc404db6db80acd55a8edbe7be7 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 6 Apr 2021 16:50:49 -0700 Subject: [PATCH 10/12] Change noalias to be a parameter attribute, not a function metadata Update test --- clang/lib/CodeGen/CGCall.cpp | 3 ++- clang/lib/CodeGen/CodeGenModule.cpp | 15 ++------------- clang/lib/Sema/SemaSYCL.cpp | 6 ++---- .../CodeGenSYCL/accessor_no_alias_property.cpp | 2 +- 4 files changed, 7 insertions(+), 19 deletions(-) diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp index dd10b01aba218..17952404e4ff8 100644 --- a/clang/lib/CodeGen/CGCall.cpp +++ b/clang/lib/CodeGen/CGCall.cpp @@ -2771,7 +2771,8 @@ void CodeGenFunction::EmitFunctionProlog(const CGFunctionInfo &FI, if (Arg->getType().isRestrictQualified() || (CurCodeDecl && CurCodeDecl->hasAttr() && - Arg->getType()->isPointerType())) + Arg->getType()->isPointerType()) || + (Arg->hasAttr() && Arg->getType()->isPointerType())) AI->addAttr(llvm::Attribute::NoAlias); } diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 775f12e86b3e6..be092a0fd97b2 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1580,9 +1580,6 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, // MDNode for the intel_buffer_location attribute. SmallVector argSYCLBufferLocationAttr; - // MDNode for accessor no_alias property. - SmallVector argAccessorNoAliasPropertyAttr; - // MDNode for listing ESIMD kernel pointer arguments originating from // accessors SmallVector argESIMDAccPtrs; @@ -1690,10 +1687,6 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, SYCLBufferLocationAttr->getLocationID())) : llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(-1))); - if (parm->hasAttr()) - argAccessorNoAliasPropertyAttr.push_back( - llvm::ConstantAsMetadata::get(CGF->Builder.getInt1(true))); - if (FD->hasAttr()) argESIMDAccPtrs.push_back(llvm::ConstantAsMetadata::get( CGF->Builder.getInt1(parm->hasAttr()))); @@ -1701,14 +1694,10 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn, bool IsEsimdFunction = FD && FD->hasAttr(); - if (LangOpts.SYCLIsDevice && !IsEsimdFunction) { - - Fn->setMetadata("noalias", llvm::MDNode::get( - VMContext, argAccessorNoAliasPropertyAttr)); - + if (LangOpts.SYCLIsDevice && !IsEsimdFunction) Fn->setMetadata("kernel_arg_buffer_location", llvm::MDNode::get(VMContext, argSYCLBufferLocationAttr)); - } else { + else { Fn->setMetadata("kernel_arg_addr_space", llvm::MDNode::get(VMContext, addressQuals)); Fn->setMetadata("kernel_arg_access_qual", diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 3712cef2a6a25..3e7c04102cff0 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1755,10 +1755,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { void handleNoAliasProperty(ParmVarDecl *Param, QualType PropTy, SourceLocation Loc) { - if (PropTy.isRestrictQualified()) { - ASTContext &Ctx = SemaRef.getASTContext(); - Param->addAttr(RestrictAttr::CreateImplicit(Ctx, Loc)); - } + ASTContext &Ctx = SemaRef.getASTContext(); + Param->addAttr(RestrictAttr::CreateImplicit(Ctx, Loc)); } // Obtain an integer value stored in a template parameter of buffer_location diff --git a/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp b/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp index 54a708c3256db..7f98579286ed8 100644 --- a/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp +++ b/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp @@ -1,6 +1,6 @@ // RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s -// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function{{.*}} !noalias +// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function({{.*}} noalias {{.*}} %_arg_, {{.*}}) #include "Inputs/sycl.hpp" From 8c94561746773c7611cf29b25a22773bf77a1f5e Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Wed, 7 Apr 2021 08:45:17 -0700 Subject: [PATCH 11/12] Undo changes to lit tests --- clang/test/CodeGenSYCL/disable_loop_pipelining.cpp | 7 ++++--- clang/test/CodeGenSYCL/initiation_interval.cpp | 9 +++++---- 2 files changed, 9 insertions(+), 7 deletions(-) diff --git a/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp b/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp index 85c51b090616c..33f83f6b1961f 100644 --- a/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp +++ b/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp @@ -29,7 +29,8 @@ int main() { return 0; } -// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel1"() #0 {{.*}} !disable_loop_pipelining ![[NUM5:[0-9]+]] -// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel2"() #0 {{.*}} -// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel3"() #0 {{.*}} !disable_loop_pipelining ![[NUM5]] +// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel1"() #0 !kernel_arg_buffer_location ![[NUM4:[0-9]+]] !disable_loop_pipelining ![[NUM5:[0-9]+]] +// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel2"() #0 !kernel_arg_buffer_location ![[NUM4]] +// CHECK: define dso_local spir_kernel void @"{{.*}}test_kernel3"() #0 !kernel_arg_buffer_location ![[NUM4]] !disable_loop_pipelining ![[NUM5]] +// CHECK: ![[NUM4]] = !{} // CHECK: ![[NUM5]] = !{i32 1} diff --git a/clang/test/CodeGenSYCL/initiation_interval.cpp b/clang/test/CodeGenSYCL/initiation_interval.cpp index b32cbb037f2ff..95ef9ce4cde50 100644 --- a/clang/test/CodeGenSYCL/initiation_interval.cpp +++ b/clang/test/CodeGenSYCL/initiation_interval.cpp @@ -39,10 +39,11 @@ int main() { return 0; } -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 {{.*}} !initiation_interval ![[NUM1:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !initiation_interval ![[NUM42:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !initiation_interval ![[NUM2:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1"() #0 !kernel_arg_buffer_location ![[NUM0:[0-9]+]] !initiation_interval ![[NUM1:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 !kernel_arg_buffer_location ![[NUM0]] !initiation_interval ![[NUM42:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 !kernel_arg_buffer_location ![[NUM0]] !initiation_interval ![[NUM2:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 !kernel_arg_buffer_location ![[NUM0]] +// CHECK: ![[NUM0]] = !{} // CHECK: ![[NUM1]] = !{i32 1} // CHECK: ![[NUM42]] = !{i32 42} // CHECK: ![[NUM2]] = !{i32 2} From 7dd04ea83ab2b7e0596b1411f80c42d28808c330 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Wed, 7 Apr 2021 11:32:44 -0700 Subject: [PATCH 12/12] Expand test case and add comments --- .../accessor_no_alias_property.cpp | 21 +++++++++++++++++-- 1 file changed, 19 insertions(+), 2 deletions(-) diff --git a/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp b/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp index 7f98579286ed8..682480cf9569f 100644 --- a/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp +++ b/clang/test/CodeGenSYCL/accessor_no_alias_property.cpp @@ -1,6 +1,10 @@ // RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s +// check that noalias parameter attribute is emitted when no_alias accessor property is used +// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE16kernel_function1({{.*}} noalias {{.*}} %_arg_, {{.*}}) -// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function({{.*}} noalias {{.*}} %_arg_, {{.*}}) +// check that noalias parameter attribute is NOT emitted when it is not used +// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE16kernel_function2{{.*}} !kernel_arg_buffer_location +// CHECK-NOT: define {{.*}}spir_kernel void @_ZTSZ4mainE16kernel_function2({{.*}} noalias {{.*}} #include "Inputs/sycl.hpp" @@ -11,9 +15,22 @@ int main() { cl::sycl::ONEAPI::accessor_property_list< cl::sycl::ONEAPI::property::no_alias::instance>> accessorA; - cl::sycl::kernel_single_task( + + cl::sycl::accessor>> + accessorB; + + cl::sycl::kernel_single_task( [=]() { accessorA.use(); }); + + cl::sycl::kernel_single_task( + [=]() { + accessorB.use(); + }); return 0; }