From f503ede0f45eefa554e7c1d758c44f7e49094792 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 28 Apr 2020 10:41:37 -0700 Subject: [PATCH 01/10] initial code. Will need to be moved and refactored Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaSYCL.cpp | 79 +++++++++++++++++++++++++++++++++++++ 1 file changed, 79 insertions(+) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 72a1551edd9d9..74a2bf7355ec7 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -28,6 +28,9 @@ #include +//CP +#include + using namespace clang; using KernelParamKind = SYCLIntegrationHeader::kernel_param_kind_t; @@ -283,6 +286,17 @@ void Sema::checkSYCLDeviceVarDecl(VarDecl *Var) { checkSYCLVarType(*this, Ty, Loc, Visited); } +// The first four pair with the enumeration in CL/sycl/usm/usm_enums.hpp +enum UsmExpr { + Usm_Host, + Usm_Device, + Usm_Shared, + Unknown, + + Not_Usm +}; + + class MarkDeviceFunction : public RecursiveASTVisitor { public: MarkDeviceFunction(Sema &S) @@ -348,6 +362,71 @@ class MarkDeviceFunction : public RecursiveASTVisitor { return true; } + + + //CP + bool VisitDeclRefExpr(DeclRefExpr *E) { + // Looks like the usm_shared_t thing won't work out. + // Catches static casts of usm allocations, but not templated or single usm::malloc call, or allocators. + + ValueDecl *D = E->getDecl(); + QualType Ty = D->getType(); + if(Ty->isAnyPointerType() && E->refersToEnclosingVariableOrCapture()) { + VarDecl *DVar = dyn_cast(D); + const Expr *Init = DVar->getAnyInitializer(); + if(Init){ + Init = Init->IgnoreCasts(); + QualType InitTy = Init->getType(); + SourceLocation DecLoc = Init->getExprLoc(); + SourceRange RefLoc = E->getSourceRange(); + std::cout << "captured qualtype: " << InitTy.getAsString() << "\n" + << "declared at: " << DecLoc.printToString(SemaRef.getSourceManager()) + << " refereced at: " << RefLoc.printToString(SemaRef.getSourceManager()) << std::endl; + + UsmExpr usesUSM = Unknown; + + //function call + const CallExpr *CE = dyn_cast(Init); + if(CE){ + const FunctionDecl *func = CE->getDirectCallee(); + auto FullName = func->getQualifiedNameAsString(); + //Check to see if this function call is one of the USM allocators. + if((FullName.compare("cl::sycl::malloc_shared") == 0) || (FullName.compare("cl::sycl::aligned_alloc_shared") == 0)) + usesUSM = Usm_Shared; + else if( (FullName.compare("cl::sycl::malloc_device") == 0) || (FullName.compare("cl::sycl::aligned_alloc_device") == 0)) + usesUSM = Usm_Device; + else if((FullName.compare("cl::sycl::malloc_host") == 0) || (FullName.compare("cl::sycl::aligned_alloc_host") == 0)) + usesUSM = Usm_Host; + else if((FullName.compare("cl::sycl::malloc") == 0) || (FullName.compare("cl::sycl::aligned_alloc") == 0)) { + auto LastArgIndex = CE->getNumArgs()-1; + if (LastArgIndex > 1){ + const Expr* LastArgExpr = CE->getArg(LastArgIndex); //DeclRefExpr 0x555823c4a2f0 'cl::sycl::usm::alloc' EnumConstant 0x555823c1cca0 'shared' 'cl::sycl::usm::alloc' + const ValueDecl *LADecl = (dyn_cast(LastArgExpr))->getDecl(); + if(LADecl){ + const EnumConstantDecl *EnumDecl = dyn_cast(LADecl); + if(EnumDecl) + usesUSM = static_cast(EnumDecl->getInitVal().getExtValue()); + } + } + } + + //std::cout << "call expression. callee: " << func->getNameInfo()/*.getName()*/.getAsString() + //<< " " << func->getQualifiedNameAsString() << std::endl; + } + //var usage + + //diagnostics + //if(usesUSM == ) + + + } + } + return true; + } + + + + // The call graph for this translation unit. CallGraph SYCLCG; // The set of functions called by a kernel function. From 027e0305be07859cfcc7f0f14ee5020bb7ba0c6d Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 28 Apr 2020 13:01:05 -0700 Subject: [PATCH 02/10] usm checks Signed-off-by: Chris Perkins --- .../clang/Basic/DiagnosticSemaKinds.td | 4 +++ clang/lib/Sema/SemaSYCL.cpp | 30 ++++++++++++------- 2 files changed, 23 insertions(+), 11 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index b13492fb85152..e9c27b95eea35 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10700,6 +10700,10 @@ def err_sycl_non_trivially_copy_ctor_dtor_type "constructible|destructible}0 class/struct type %1">; def err_sycl_non_std_layout_type : Error< "kernel parameter has non-standard layout class/struct type %0">; +def err_sycl_illegal_memory_reference : Error< + "Illegal memory reference in SYCL device. Use USM (malloc_shared, etc) instead.">; +def note_sycl_capture_declared_here : Note< + "Declared here.">; def err_conflicting_sycl_kernel_attributes : Error< "conflicting attributes applied to a SYCL kernel">; def err_conflicting_sycl_function_attributes : Error< diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 74a2bf7355ec7..9c935c0910d38 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -286,13 +286,10 @@ void Sema::checkSYCLDeviceVarDecl(VarDecl *Var) { checkSYCLVarType(*this, Ty, Loc, Visited); } -// The first four pair with the enumeration in CL/sycl/usm/usm_enums.hpp + enum UsmExpr { - Usm_Host, - Usm_Device, - Usm_Shared, Unknown, - + Uses_Usm, Not_Usm }; @@ -391,6 +388,11 @@ class MarkDeviceFunction : public RecursiveASTVisitor { const FunctionDecl *func = CE->getDirectCallee(); auto FullName = func->getQualifiedNameAsString(); //Check to see if this function call is one of the USM allocators. + if( (FullName.rfind("cl::sycl::malloc", 0) == 0) || (FullName.rfind("cl::sycl::aligned_alloc", 0) == 0)) + usesUSM = Uses_Usm; + else if ( (FullName.compare("malloc")==0) || (FullName.compare("calloc")==0) ) + usesUSM = Not_Usm; +/* if((FullName.compare("cl::sycl::malloc_shared") == 0) || (FullName.compare("cl::sycl::aligned_alloc_shared") == 0)) usesUSM = Usm_Shared; else if( (FullName.compare("cl::sycl::malloc_device") == 0) || (FullName.compare("cl::sycl::aligned_alloc_device") == 0)) @@ -399,24 +401,30 @@ class MarkDeviceFunction : public RecursiveASTVisitor { usesUSM = Usm_Host; else if((FullName.compare("cl::sycl::malloc") == 0) || (FullName.compare("cl::sycl::aligned_alloc") == 0)) { auto LastArgIndex = CE->getNumArgs()-1; - if (LastArgIndex > 1){ - const Expr* LastArgExpr = CE->getArg(LastArgIndex); //DeclRefExpr 0x555823c4a2f0 'cl::sycl::usm::alloc' EnumConstant 0x555823c1cca0 'shared' 'cl::sycl::usm::alloc' + if (LastArgIndex > 0){ + const Expr* LastArgExpr = CE->getArg(LastArgIndex); //e.g. DeclRefExpr 'cl::sycl::usm::alloc' EnumConstant 'shared' 'cl::sycl::usm::alloc' const ValueDecl *LADecl = (dyn_cast(LastArgExpr))->getDecl(); if(LADecl){ const EnumConstantDecl *EnumDecl = dyn_cast(LADecl); - if(EnumDecl) + if(EnumDecl) { usesUSM = static_cast(EnumDecl->getInitVal().getExtValue()); + } } } } +*/ - //std::cout << "call expression. callee: " << func->getNameInfo()/*.getName()*/.getAsString() - //<< " " << func->getQualifiedNameAsString() << std::endl; + std::cout << "call expression. callee: " << func->getNameInfo()/*.getName()*/.getAsString() + << " " << func->getQualifiedNameAsString() << std::endl; } //var usage //diagnostics - //if(usesUSM == ) + std::cout << "usesUSM: " << usesUSM << std::endl; + if(usesUSM == Not_Usm){ + SemaRef.Diag(RefLoc.getBegin(), diag::err_sycl_illegal_memory_reference); + SemaRef.Diag(DecLoc, diag::note_sycl_capture_declared_here); + } } From b59c81228d6e39c05f80e8f1e9f432dfda5516a6 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 29 Apr 2020 08:56:51 -0700 Subject: [PATCH 03/10] note Signed-off-by: Chris Perkins --- .../clang/Basic/DiagnosticSemaKinds.td | 4 +- clang/lib/Sema/SemaSYCL.cpp | 61 ++++++++----------- 2 files changed, 29 insertions(+), 36 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index e9c27b95eea35..891f82f962641 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10701,7 +10701,9 @@ def err_sycl_non_trivially_copy_ctor_dtor_type def err_sycl_non_std_layout_type : Error< "kernel parameter has non-standard layout class/struct type %0">; def err_sycl_illegal_memory_reference : Error< - "Illegal memory reference in SYCL device. Use USM (malloc_shared, etc) instead.">; + "Illegal memory reference in SYCL device kernel. Use USM (malloc_shared, etc) instead.">; +def note_unknown_memory_reference : Note< + "Unknown memory reference in SYCL device kernel. Be sure memory was allocated with USM (malloc_shared, etc).">; def note_sycl_capture_declared_here : Note< "Declared here.">; def err_conflicting_sycl_kernel_attributes : Error< diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 9c935c0910d38..c14ad9531f589 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -286,11 +286,11 @@ void Sema::checkSYCLDeviceVarDecl(VarDecl *Var) { checkSYCLVarType(*this, Ty, Loc, Visited); } - -enum UsmExpr { +// For a DeclRefExpr, determine how it was allocated. +enum ExprAllocation { Unknown, - Uses_Usm, - Not_Usm + USM, + Not_USM }; @@ -366,6 +366,7 @@ class MarkDeviceFunction : public RecursiveASTVisitor { // Looks like the usm_shared_t thing won't work out. // Catches static casts of usm allocations, but not templated or single usm::malloc call, or allocators. + bool speak = false; ValueDecl *D = E->getDecl(); QualType Ty = D->getType(); if(Ty->isAnyPointerType() && E->refersToEnclosingVariableOrCapture()) { @@ -376,11 +377,10 @@ class MarkDeviceFunction : public RecursiveASTVisitor { QualType InitTy = Init->getType(); SourceLocation DecLoc = Init->getExprLoc(); SourceRange RefLoc = E->getSourceRange(); - std::cout << "captured qualtype: " << InitTy.getAsString() << "\n" - << "declared at: " << DecLoc.printToString(SemaRef.getSourceManager()) - << " refereced at: " << RefLoc.printToString(SemaRef.getSourceManager()) << std::endl; - UsmExpr usesUSM = Unknown; + + + ExprAllocation howAllocated = Unknown; //function call const CallExpr *CE = dyn_cast(Init); @@ -389,41 +389,32 @@ class MarkDeviceFunction : public RecursiveASTVisitor { auto FullName = func->getQualifiedNameAsString(); //Check to see if this function call is one of the USM allocators. if( (FullName.rfind("cl::sycl::malloc", 0) == 0) || (FullName.rfind("cl::sycl::aligned_alloc", 0) == 0)) - usesUSM = Uses_Usm; + howAllocated = USM; else if ( (FullName.compare("malloc")==0) || (FullName.compare("calloc")==0) ) - usesUSM = Not_Usm; -/* - if((FullName.compare("cl::sycl::malloc_shared") == 0) || (FullName.compare("cl::sycl::aligned_alloc_shared") == 0)) - usesUSM = Usm_Shared; - else if( (FullName.compare("cl::sycl::malloc_device") == 0) || (FullName.compare("cl::sycl::aligned_alloc_device") == 0)) - usesUSM = Usm_Device; - else if((FullName.compare("cl::sycl::malloc_host") == 0) || (FullName.compare("cl::sycl::aligned_alloc_host") == 0)) - usesUSM = Usm_Host; - else if((FullName.compare("cl::sycl::malloc") == 0) || (FullName.compare("cl::sycl::aligned_alloc") == 0)) { - auto LastArgIndex = CE->getNumArgs()-1; - if (LastArgIndex > 0){ - const Expr* LastArgExpr = CE->getArg(LastArgIndex); //e.g. DeclRefExpr 'cl::sycl::usm::alloc' EnumConstant 'shared' 'cl::sycl::usm::alloc' - const ValueDecl *LADecl = (dyn_cast(LastArgExpr))->getDecl(); - if(LADecl){ - const EnumConstantDecl *EnumDecl = dyn_cast(LADecl); - if(EnumDecl) { - usesUSM = static_cast(EnumDecl->getInitVal().getExtValue()); - } - } - } - } -*/ + howAllocated = Not_USM; - std::cout << "call expression. callee: " << func->getNameInfo()/*.getName()*/.getAsString() - << " " << func->getQualifiedNameAsString() << std::endl; + //std::cout << "call expression. callee: " << func->getNameInfo()/*.getName()*/.getAsString() + //<< " " << func->getQualifiedNameAsString() << std::endl; + } else { + speak = true; } //var usage + + if(speak){ + std::cout << "captured qualtype: " << InitTy.getAsString() << "\n:::: " + << RefLoc.printToString(SemaRef.getSourceManager()) + << "declared at: " << DecLoc.printToString(SemaRef.getSourceManager()) + << std::endl; + } + //diagnostics - std::cout << "usesUSM: " << usesUSM << std::endl; - if(usesUSM == Not_Usm){ + if(howAllocated == Not_USM){ SemaRef.Diag(RefLoc.getBegin(), diag::err_sycl_illegal_memory_reference); SemaRef.Diag(DecLoc, diag::note_sycl_capture_declared_here); + }else if (howAllocated == Unknown){ + SemaRef.Diag(RefLoc.getBegin(), diag::note_unknown_memory_reference); + SemaRef.Diag(DecLoc, diag::note_sycl_capture_declared_here); } From 515fb61ecf70ac654fa3667ad6458c47285c6867 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 29 Apr 2020 11:10:40 -0700 Subject: [PATCH 04/10] call origination in SemaExpr.cpp isVariableCapturable Signed-off-by: Chris Perkins --- clang/include/clang/Sema/Sema.h | 1 + clang/lib/Sema/SemaExpr.cpp | 8 ++++- clang/lib/Sema/SemaSYCL.cpp | 64 ++++++++++++++++++++++++++++++--- 3 files changed, 67 insertions(+), 6 deletions(-) diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index e5ef7433d7a50..b8f241b785186 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -12456,6 +12456,7 @@ class Sema final { bool isKnownGoodSYCLDecl(const Decl *D); void checkSYCLDeviceVarDecl(VarDecl *Var); + void checkSYCLDevicePointerCapture(VarDecl *Var, SourceLocation CaptureLoc); void ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC); void MarkDevice(); diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index bd7db4a9aca18..f2ed245a36596 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -16254,6 +16254,9 @@ static DeclContext *getParentOfCapturingContextOrNull(DeclContext *DC, VarDecl * return nullptr; } +//CP +#include + // Certain capturing entities (lambdas, blocks etc.) are not allowed to capture // certain types of variables (unnamed, variably modified types etc.) // so check for eligibility. @@ -16320,7 +16323,10 @@ static bool isVariableCapturable(CapturingScopeInfo *CSI, VarDecl *Var, S.Diag(Loc, diag::err_opencl_block_ref_block); return false; } - + // SYCL - emit errors for any illegal pointer derefs. + if(Diagnose && S.getLangOpts().SYCLIsDevice && Var->getType()->isAnyPointerType()) + checkSYCLDevicePointerCapture(Var, Loc); + return true; } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c14ad9531f589..7bbff05b430af 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -286,6 +286,60 @@ void Sema::checkSYCLDeviceVarDecl(VarDecl *Var) { checkSYCLVarType(*this, Ty, Loc, Visited); } +void Sema::checkSYCLDevicePointerCapture(VarDecl *Var, SourceLocation CaptureLoc) { + assert(getLangOpts().SYCLIsDevice && + "Should only be called during SYCL compilation"); + bool speak = true; + const Expr *Init = Var->getAnyInitializer(); + if(Init){ + Init = Init->IgnoreCasts(); + QualType InitTy = Init->getType(); + SourceLocation DecLoc = Init->getExprLoc(); + + enum ExprAllocation { + Unknown, + USM, + Not_USM + }; + ExprAllocation howAllocated = Unknown; + + //function call + const CallExpr *CE = dyn_cast(Init); + if(CE){ + const FunctionDecl *func = CE->getDirectCallee(); + auto FullName = func->getQualifiedNameAsString(); + //Check to see if this function call is one of the USM allocators. + if( (FullName.rfind("cl::sycl::malloc", 0) == 0) || (FullName.rfind("cl::sycl::aligned_alloc", 0) == 0)) + howAllocated = USM; + else if ( (FullName.compare("malloc")==0) || (FullName.compare("calloc")==0) ) + howAllocated = Not_USM; + + //std::cout << "call expression. callee: " << func->getNameInfo()/*.getName()*/.getAsString() + //<< " " << func->getQualifiedNameAsString() << std::endl; + } else { + speak = true; + } + //var usage + + if(speak){ + std::cout << "captured qualtype: " << InitTy.getAsString() << "\n:::: " + << CaptureLoc.printToString(getSourceManager()) + << "declared at: " << DecLoc.printToString(getSourceManager()) + << std::endl; + } + + //diagnostics + if(howAllocated == Not_USM){ + SYCLDiagIfDeviceCode(CaptureLoc, diag::err_sycl_illegal_memory_reference); + SYCLDiagIfDeviceCode(DecLoc, diag::note_sycl_capture_declared_here); + }else if (howAllocated == Unknown){ + SYCLDiagIfDeviceCode(CaptureLoc, diag::note_unknown_memory_reference); + SYCLDiagIfDeviceCode(DecLoc, diag::note_sycl_capture_declared_here); + } + + } +} + // For a DeclRefExpr, determine how it was allocated. enum ExprAllocation { Unknown, @@ -396,7 +450,7 @@ class MarkDeviceFunction : public RecursiveASTVisitor { //std::cout << "call expression. callee: " << func->getNameInfo()/*.getName()*/.getAsString() //<< " " << func->getQualifiedNameAsString() << std::endl; } else { - speak = true; + //speak = true; } //var usage @@ -410,11 +464,11 @@ class MarkDeviceFunction : public RecursiveASTVisitor { //diagnostics if(howAllocated == Not_USM){ - SemaRef.Diag(RefLoc.getBegin(), diag::err_sycl_illegal_memory_reference); - SemaRef.Diag(DecLoc, diag::note_sycl_capture_declared_here); + //SemaRef.Diag(RefLoc.getBegin(), diag::err_sycl_illegal_memory_reference); + //SemaRef.Diag(DecLoc, diag::note_sycl_capture_declared_here); }else if (howAllocated == Unknown){ - SemaRef.Diag(RefLoc.getBegin(), diag::note_unknown_memory_reference); - SemaRef.Diag(DecLoc, diag::note_sycl_capture_declared_here); + //SemaRef.Diag(RefLoc.getBegin(), diag::note_unknown_memory_reference); + //SemaRef.Diag(DecLoc, diag::note_sycl_capture_declared_here); } From 0af8083d78b3675fcda02aee12caf62423f27198 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 29 Apr 2020 14:39:29 -0700 Subject: [PATCH 05/10] initial lit test Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaExpr.cpp | 6 +- clang/lib/Sema/SemaSYCL.cpp | 88 +----- .../sycl-pointer-capture-diagnostics.cpp | 260 ++++++++++++++++++ 3 files changed, 268 insertions(+), 86 deletions(-) create mode 100644 clang/test/SemaSYCL/sycl-pointer-capture-diagnostics.cpp diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index f2ed245a36596..57766c734bae4 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -16254,8 +16254,6 @@ static DeclContext *getParentOfCapturingContextOrNull(DeclContext *DC, VarDecl * return nullptr; } -//CP -#include // Certain capturing entities (lambdas, blocks etc.) are not allowed to capture // certain types of variables (unnamed, variably modified types etc.) @@ -16323,9 +16321,9 @@ static bool isVariableCapturable(CapturingScopeInfo *CSI, VarDecl *Var, S.Diag(Loc, diag::err_opencl_block_ref_block); return false; } - // SYCL - emit errors for any illegal pointer derefs. + // SYCL: Emit diagnostics for any illegal pointer derefs. if(Diagnose && S.getLangOpts().SYCLIsDevice && Var->getType()->isAnyPointerType()) - checkSYCLDevicePointerCapture(Var, Loc); + S.checkSYCLDevicePointerCapture(Var, Loc); return true; } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 7bbff05b430af..4769fd0ed2f82 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -289,26 +289,22 @@ void Sema::checkSYCLDeviceVarDecl(VarDecl *Var) { void Sema::checkSYCLDevicePointerCapture(VarDecl *Var, SourceLocation CaptureLoc) { assert(getLangOpts().SYCLIsDevice && "Should only be called during SYCL compilation"); - bool speak = true; + bool speak = false; const Expr *Init = Var->getAnyInitializer(); if(Init){ Init = Init->IgnoreCasts(); QualType InitTy = Init->getType(); SourceLocation DecLoc = Init->getExprLoc(); - enum ExprAllocation { - Unknown, - USM, - Not_USM - }; + enum ExprAllocation { Unknown, USM, Not_USM }; ExprAllocation howAllocated = Unknown; - //function call const CallExpr *CE = dyn_cast(Init); if(CE){ + // Captured variable is result of function call. const FunctionDecl *func = CE->getDirectCallee(); auto FullName = func->getQualifiedNameAsString(); - //Check to see if this function call is one of the USM allocators. + // Check to see if this function call is one of the USM allocators. if( (FullName.rfind("cl::sycl::malloc", 0) == 0) || (FullName.rfind("cl::sycl::aligned_alloc", 0) == 0)) howAllocated = USM; else if ( (FullName.compare("malloc")==0) || (FullName.compare("calloc")==0) ) @@ -317,9 +313,10 @@ void Sema::checkSYCLDevicePointerCapture(VarDecl *Var, SourceLocation CaptureLoc //std::cout << "call expression. callee: " << func->getNameInfo()/*.getName()*/.getAsString() //<< " " << func->getQualifiedNameAsString() << std::endl; } else { + // var usage speak = true; } - //var usage + if(speak){ std::cout << "captured qualtype: " << InitTy.getAsString() << "\n:::: " @@ -340,13 +337,6 @@ void Sema::checkSYCLDevicePointerCapture(VarDecl *Var, SourceLocation CaptureLoc } } -// For a DeclRefExpr, determine how it was allocated. -enum ExprAllocation { - Unknown, - USM, - Not_USM -}; - class MarkDeviceFunction : public RecursiveASTVisitor { public: @@ -414,72 +404,6 @@ class MarkDeviceFunction : public RecursiveASTVisitor { } - - //CP - bool VisitDeclRefExpr(DeclRefExpr *E) { - // Looks like the usm_shared_t thing won't work out. - // Catches static casts of usm allocations, but not templated or single usm::malloc call, or allocators. - - bool speak = false; - ValueDecl *D = E->getDecl(); - QualType Ty = D->getType(); - if(Ty->isAnyPointerType() && E->refersToEnclosingVariableOrCapture()) { - VarDecl *DVar = dyn_cast(D); - const Expr *Init = DVar->getAnyInitializer(); - if(Init){ - Init = Init->IgnoreCasts(); - QualType InitTy = Init->getType(); - SourceLocation DecLoc = Init->getExprLoc(); - SourceRange RefLoc = E->getSourceRange(); - - - - ExprAllocation howAllocated = Unknown; - - //function call - const CallExpr *CE = dyn_cast(Init); - if(CE){ - const FunctionDecl *func = CE->getDirectCallee(); - auto FullName = func->getQualifiedNameAsString(); - //Check to see if this function call is one of the USM allocators. - if( (FullName.rfind("cl::sycl::malloc", 0) == 0) || (FullName.rfind("cl::sycl::aligned_alloc", 0) == 0)) - howAllocated = USM; - else if ( (FullName.compare("malloc")==0) || (FullName.compare("calloc")==0) ) - howAllocated = Not_USM; - - //std::cout << "call expression. callee: " << func->getNameInfo()/*.getName()*/.getAsString() - //<< " " << func->getQualifiedNameAsString() << std::endl; - } else { - //speak = true; - } - //var usage - - - if(speak){ - std::cout << "captured qualtype: " << InitTy.getAsString() << "\n:::: " - << RefLoc.printToString(SemaRef.getSourceManager()) - << "declared at: " << DecLoc.printToString(SemaRef.getSourceManager()) - << std::endl; - } - - //diagnostics - if(howAllocated == Not_USM){ - //SemaRef.Diag(RefLoc.getBegin(), diag::err_sycl_illegal_memory_reference); - //SemaRef.Diag(DecLoc, diag::note_sycl_capture_declared_here); - }else if (howAllocated == Unknown){ - //SemaRef.Diag(RefLoc.getBegin(), diag::note_unknown_memory_reference); - //SemaRef.Diag(DecLoc, diag::note_sycl_capture_declared_here); - } - - - } - } - return true; - } - - - - // The call graph for this translation unit. CallGraph SYCLCG; // The set of functions called by a kernel function. diff --git a/clang/test/SemaSYCL/sycl-pointer-capture-diagnostics.cpp b/clang/test/SemaSYCL/sycl-pointer-capture-diagnostics.cpp new file mode 100644 index 0000000000000..a9f45a4cda662 --- /dev/null +++ b/clang/test/SemaSYCL/sycl-pointer-capture-diagnostics.cpp @@ -0,0 +1,260 @@ +// RUN: %clang_cc1 -fsycl -triple spir64 -fsycl-is-device -verify -fsyntax-only %s +// +// Pointer variables captured by kernel lambda are checked. +// Ensure those diagnostics are working correctly. + +// Mock USM functions trigger warnings, suppress. +#pragma clang diagnostic push +#pragma clang diagnostic ignored "-Wreturn-stack-address" +#pragma clang diagnostic ignored "-Wint-to-pointer-cast" + +namespace std { +class type_info; +typedef __typeof__(sizeof(int)) size_t; +} // namespace std + +inline namespace cl { +namespace sycl { + +template +__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { + kernelFunc(); //#call_kernelFunc +} + +typedef int device; +typedef int context; +typedef double queue; + +// Mock USM memory allocation functions. +namespace usm { +enum class alloc { host, + device, + shared, + unknown }; +} // namespace usm + +void *malloc(std::size_t sz, const device &dev, const context &ctxt, cl::sycl::usm::alloc kind) { + int a = 11; + return (void *)(&a); +} +void *malloc(std::size_t sz, const queue &q, cl::sycl::usm::alloc kind) { + int a = 11; + return (void *)(&a); +} + +void *malloc_device(std::size_t sz, const device &dev, const context &ctxt) { + int a = 11; + return (void *)(&a); +} +void *malloc_device(std::size_t sz, const queue &q) { + int a = 11; + return (void *)(&a); +} + +void *malloc_shared(std::size_t sz, const device &dev, const context &ctxt) { + int a = 11; + return (void *)(&a); +} +void *malloc_shared(std::size_t sz, const queue &q) { + int a = 11; + return (void *)(&a); +} + +void *malloc_host(std::size_t sz, const context &ctxt) { + int a = 12; + return (void *)(&a); +} +void *malloc_host(std::size_t sz, const queue &q) { + int a = 12; + return (void *)(&a); +} + +void *aligned_alloc(std::size_t alignment, std::size_t sz, const device &dev, const context &ctxt, cl::sycl::usm::alloc kind) { + int a = 11; + return (void *)(&a); +} +void *aligned_alloc(std::size_t alignment, std::size_t sz, const queue &q, cl::sycl::usm::alloc kind) { + int a = 11; + return (void *)(&a); +} + +void *aligned_alloc_device(std::size_t alignment, std::size_t sz, const device &dev, const context &ctxt) { + int a = 11; + return (void *)(&a); +} +void *aligned_alloc_device(std::size_t alignment, std::size_t sz, const queue &q) { + int a = 11; + return (void *)(&a); +} + +void *aligned_alloc_shared(std::size_t alignment, std::size_t sz, const device &dev, const context &ctxt) { + int a = 11; + return (void *)(&a); +} +void *aligned_alloc_shared(std::size_t alignment, std::size_t sz, const queue &q) { + int a = 11; + return (void *)(&a); +} + +void *aligned_alloc_host(std::size_t alignment, std::size_t sz, const context &ctxt) { + int a = 12; + return (void *)(&a); +} +void *aligned_alloc_host(std::size_t alignment, std::size_t sz, const queue &q) { + int a = 12; + return (void *)(&a); +} + +//template form +template +T *malloc_shared(std::size_t Count, const device &Dev, const context &Ctxt) { + return static_cast(malloc_shared(Count * sizeof(T), Dev, Ctxt)); +} + +} // namespace sycl +} // namespace cl + +void *malloc(std::size_t sz) { + int a = 11; + return (void *)(&a); +} +void *calloc(std::size_t num, std::size_t sz) { + int a = 11; + return (void *)(&a); +} +// -- END MOCKS + +float calledFromLambda(float *first) { + return first[0]; +} + +int main(int argc, char **argv) { + + int device = 0, context = 0; + double queue = 0; + + //bad pointers + float stackFloat = 20.0; + float *stackFloatP = &stackFloat; //#decl_stackFloatP + + float *frenemy = stackFloatP; //#decl_frenemy + frenemy++; + + float *fromParam = (float *)(argc); //#decl_fromParam + + float *mallocFloatP = static_cast(malloc(sizeof(float) * 2)); //#decl_mallocFloatP + float *mallocFloatP2 = static_cast(malloc(sizeof(float) * 2)); //#decl_mallocFloatP2 + float *callocFloatP = static_cast(calloc(2, sizeof(float))); //#decl_callocFloatP + float *callocFloatP2 = static_cast(calloc(2, sizeof(float))); //#decl_callocFloatP2 + + //usm + float *usmSharedP = static_cast(sycl::malloc_shared(sizeof(float), device, context)); + float *usmSharedP2 = static_cast(sycl::malloc_shared(sizeof(float), queue)); + float *usmSharedP3 = static_cast(sycl::malloc(sizeof(float), device, context, cl::sycl::usm::alloc::shared)); + float *usmSharedP4 = static_cast(sycl::malloc(sizeof(float), queue, cl::sycl::usm::alloc::shared)); + float *usmSharedP5 = sycl::malloc_shared(1, device, context); + + float *usmShAlignP = static_cast(sycl::aligned_alloc_shared(1, sizeof(float), device, context)); + float *usmShAlignP2 = static_cast(sycl::aligned_alloc_shared(1, sizeof(float), queue)); + float *usmShAlignP3 = static_cast(sycl::aligned_alloc(1, sizeof(float), device, context, cl::sycl::usm::alloc::shared)); + float *usmShAlignP4 = static_cast(sycl::aligned_alloc(1, sizeof(float), queue, cl::sycl::usm::alloc::shared)); + + float *usmHostP = static_cast(sycl::malloc_host(sizeof(float), context)); + float *usmHostP2 = static_cast(sycl::malloc_host(sizeof(float), queue)); + float *usmHostP3 = static_cast(sycl::malloc(sizeof(float), device, context, cl::sycl::usm::alloc::host)); + float *usmHostP4 = static_cast(sycl::malloc(sizeof(float), queue, cl::sycl::usm::alloc::host)); + + float *usmHoAlignP = static_cast(sycl::aligned_alloc_host(1, sizeof(float), context)); + float *usmHoAlignP2 = static_cast(sycl::aligned_alloc_host(1, sizeof(float), queue)); + float *usmHoAlignP3 = static_cast(sycl::aligned_alloc(1, sizeof(float), device, context, cl::sycl::usm::alloc::host)); + float *usmHoAlignP4 = static_cast(sycl::aligned_alloc(1, sizeof(float), queue, cl::sycl::usm::alloc::host)); + + float *usmDeviceP = static_cast(sycl::malloc_device(sizeof(float), device, context)); + float *usmDeviceP2 = static_cast(sycl::malloc_device(sizeof(float), queue)); + float *usmDeviceP3 = static_cast(sycl::malloc(sizeof(float), device, context, cl::sycl::usm::alloc::device)); + float *usmDeviceP4 = static_cast(sycl::malloc(sizeof(float), queue, cl::sycl::usm::alloc::device)); + + // --- direct lambda testing --- + cl::sycl::kernel_single_task([=]() { + // --- The following dangerous pointer captures result in errors or notes. + + // expected-note@#call_kernelFunc {{called by 'kernel_single_task Date: Wed, 29 Apr 2020 17:24:06 -0700 Subject: [PATCH 06/10] string literals Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaExpr.cpp | 2 +- clang/lib/Sema/SemaSYCL.cpp | 5 ++++- .../SemaSYCL/sycl-pointer-capture-diagnostics.cpp | 12 ++++++++++++ 3 files changed, 17 insertions(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 57766c734bae4..6f2786662c5ce 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -16322,7 +16322,7 @@ static bool isVariableCapturable(CapturingScopeInfo *CSI, VarDecl *Var, return false; } // SYCL: Emit diagnostics for any illegal pointer derefs. - if(Diagnose && S.getLangOpts().SYCLIsDevice && Var->getType()->isAnyPointerType()) + if(Diagnose && S.getLangOpts().SYCLIsDevice && Var->getType()->isAnyPointerType()) S.checkSYCLDevicePointerCapture(Var, Loc); return true; diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 4769fd0ed2f82..5478f645d01e9 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -301,7 +301,7 @@ void Sema::checkSYCLDevicePointerCapture(VarDecl *Var, SourceLocation CaptureLoc const CallExpr *CE = dyn_cast(Init); if(CE){ - // Captured variable is result of function call. + // Captured pointer is result of function call. const FunctionDecl *func = CE->getDirectCallee(); auto FullName = func->getQualifiedNameAsString(); // Check to see if this function call is one of the USM allocators. @@ -314,6 +314,9 @@ void Sema::checkSYCLDevicePointerCapture(VarDecl *Var, SourceLocation CaptureLoc //<< " " << func->getQualifiedNameAsString() << std::endl; } else { // var usage + const StringLiteral *SL = dyn_cast(Init); + if(SL) + howAllocated = Not_USM; // speak = true; } diff --git a/clang/test/SemaSYCL/sycl-pointer-capture-diagnostics.cpp b/clang/test/SemaSYCL/sycl-pointer-capture-diagnostics.cpp index a9f45a4cda662..3bd350213f70c 100644 --- a/clang/test/SemaSYCL/sycl-pointer-capture-diagnostics.cpp +++ b/clang/test/SemaSYCL/sycl-pointer-capture-diagnostics.cpp @@ -142,6 +142,10 @@ int main(int argc, char **argv) { float *fromParam = (float *)(argc); //#decl_fromParam + // std::string is already caught by 'non-trivially copy constructible' check. + // so we only worry about literal strings. + auto stringLiteral = "omgwtf"; //#decl_stringLiteral + float *mallocFloatP = static_cast(malloc(sizeof(float) * 2)); //#decl_mallocFloatP float *mallocFloatP2 = static_cast(malloc(sizeof(float) * 2)); //#decl_mallocFloatP2 float *callocFloatP = static_cast(calloc(2, sizeof(float))); //#decl_callocFloatP @@ -195,6 +199,10 @@ int main(int argc, char **argv) { // expected-note@+1 {{Unknown memory reference in SYCL device kernel. Be sure memory was allocated with USM (malloc_shared, etc).}} fromParam[0] = 70.0; + // expected-note@#decl_stringLiteral {{Declared here.}} + // expected-error@+1 {{Illegal memory reference in SYCL device kernel. Use USM (malloc_shared, etc) instead.}} + char x = stringLiteral[0]; + // expected-note@#decl_mallocFloatP2 {{Declared here.}} // expected-error@+1 {{Illegal memory reference in SYCL device kernel. Use USM (malloc_shared, etc) instead.}} mallocFloatP2[0] = 80; @@ -207,11 +215,14 @@ int main(int argc, char **argv) { // expected-error@+1 {{Illegal memory reference in SYCL device kernel. Use USM (malloc_shared, etc) instead.}} float someValue = *callocFloatP2; + + // --- Only the first capture of a pointer emits anything. So these violations will NOT emit redundant diagnostics. calledFromLambda(mallocFloatP); stackFloatP[0] = 31.0; frenemy[0] = 41.0; fromParam[0] = 71.0; + char y = stringLiteral[0]; mallocFloatP2[0] = 81; callocFloatP[0] = 81; float someOtherValue = *callocFloatP2; @@ -248,6 +259,7 @@ int main(int argc, char **argv) { stackFloatP[0] = 30.0; frenemy[0] = 40.0; fromParam[0] = 70.0; + char x = stringLiteral[0]; mallocFloatP2[0] = 80; callocFloatP[0] = 80; float someValue = *callocFloatP2; From 81e1b00444e1e16739b92a1c8a147f246856bdff Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 30 Apr 2020 14:23:19 -0700 Subject: [PATCH 07/10] double checking Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaSYCL.cpp | 60 +++--- .../sycl-pointer-capture-diagnostics.cpp | 191 +++++++++--------- 2 files changed, 123 insertions(+), 128 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 5478f645d01e9..1c5a07a104cfb 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -28,8 +28,6 @@ #include -//CP -#include using namespace clang; @@ -289,18 +287,22 @@ void Sema::checkSYCLDeviceVarDecl(VarDecl *Var) { void Sema::checkSYCLDevicePointerCapture(VarDecl *Var, SourceLocation CaptureLoc) { assert(getLangOpts().SYCLIsDevice && "Should only be called during SYCL compilation"); - bool speak = false; + assert(Var->getType()->isAnyPointerType() && + "Should only be called for pointer types being captured."); + + enum ExprAllocation { Unknown, USM, Not_USM }; + ExprAllocation howAllocated = Unknown; + + SourceLocation DecLoc = SourceLocation(); + + // Try to determine provenance of this Var. const Expr *Init = Var->getAnyInitializer(); if(Init){ Init = Init->IgnoreCasts(); - QualType InitTy = Init->getType(); - SourceLocation DecLoc = Init->getExprLoc(); + DecLoc = Init->getExprLoc(); - enum ExprAllocation { Unknown, USM, Not_USM }; - ExprAllocation howAllocated = Unknown; - const CallExpr *CE = dyn_cast(Init); - if(CE){ + if(CE) { // Captured pointer is result of function call. const FunctionDecl *func = CE->getDirectCallee(); auto FullName = func->getQualifiedNameAsString(); @@ -310,34 +312,28 @@ void Sema::checkSYCLDevicePointerCapture(VarDecl *Var, SourceLocation CaptureLoc else if ( (FullName.compare("malloc")==0) || (FullName.compare("calloc")==0) ) howAllocated = Not_USM; - //std::cout << "call expression. callee: " << func->getNameInfo()/*.getName()*/.getAsString() - //<< " " << func->getQualifiedNameAsString() << std::endl; } else { - // var usage + // Var has initialization, but not as return result of a function, disqualify any other obvious bad initialization expressions. const StringLiteral *SL = dyn_cast(Init); if(SL) - howAllocated = Not_USM; // - speak = true; + howAllocated = Not_USM; + + if(Init->isRValue()) // pr-value + howAllocated = Not_USM; } + } + //else: Var does not have local initialization, might be parameter, or initialized via ref, etc. + // Nothing more we can determine at this time. - - if(speak){ - std::cout << "captured qualtype: " << InitTy.getAsString() << "\n:::: " - << CaptureLoc.printToString(getSourceManager()) - << "declared at: " << DecLoc.printToString(getSourceManager()) - << std::endl; - } - - //diagnostics - if(howAllocated == Not_USM){ - SYCLDiagIfDeviceCode(CaptureLoc, diag::err_sycl_illegal_memory_reference); - SYCLDiagIfDeviceCode(DecLoc, diag::note_sycl_capture_declared_here); - }else if (howAllocated == Unknown){ - SYCLDiagIfDeviceCode(CaptureLoc, diag::note_unknown_memory_reference); - SYCLDiagIfDeviceCode(DecLoc, diag::note_sycl_capture_declared_here); - } - - } + //diagnostics + if(howAllocated == Not_USM) + SYCLDiagIfDeviceCode(CaptureLoc, diag::err_sycl_illegal_memory_reference); + else if (howAllocated == Unknown) + SYCLDiagIfDeviceCode(CaptureLoc, diag::note_unknown_memory_reference); + + if(howAllocated != USM && DecLoc.isValid()) + SYCLDiagIfDeviceCode(DecLoc, diag::note_sycl_capture_declared_here); + } diff --git a/clang/test/SemaSYCL/sycl-pointer-capture-diagnostics.cpp b/clang/test/SemaSYCL/sycl-pointer-capture-diagnostics.cpp index 3bd350213f70c..f5bb06b6180f2 100644 --- a/clang/test/SemaSYCL/sycl-pointer-capture-diagnostics.cpp +++ b/clang/test/SemaSYCL/sycl-pointer-capture-diagnostics.cpp @@ -3,11 +3,6 @@ // Pointer variables captured by kernel lambda are checked. // Ensure those diagnostics are working correctly. -// Mock USM functions trigger warnings, suppress. -#pragma clang diagnostic push -#pragma clang diagnostic ignored "-Wreturn-stack-address" -#pragma clang diagnostic ignored "-Wint-to-pointer-cast" - namespace std { class type_info; typedef __typeof__(sizeof(int)) size_t; @@ -25,7 +20,7 @@ typedef int device; typedef int context; typedef double queue; -// Mock USM memory allocation functions. +//-- Mock USM memory allocation functions. namespace usm { enum class alloc { host, device, @@ -33,77 +28,22 @@ enum class alloc { host, unknown }; } // namespace usm -void *malloc(std::size_t sz, const device &dev, const context &ctxt, cl::sycl::usm::alloc kind) { - int a = 11; - return (void *)(&a); -} -void *malloc(std::size_t sz, const queue &q, cl::sycl::usm::alloc kind) { - int a = 11; - return (void *)(&a); -} - -void *malloc_device(std::size_t sz, const device &dev, const context &ctxt) { - int a = 11; - return (void *)(&a); -} -void *malloc_device(std::size_t sz, const queue &q) { - int a = 11; - return (void *)(&a); -} - -void *malloc_shared(std::size_t sz, const device &dev, const context &ctxt) { - int a = 11; - return (void *)(&a); -} -void *malloc_shared(std::size_t sz, const queue &q) { - int a = 11; - return (void *)(&a); -} - -void *malloc_host(std::size_t sz, const context &ctxt) { - int a = 12; - return (void *)(&a); -} -void *malloc_host(std::size_t sz, const queue &q) { - int a = 12; - return (void *)(&a); -} - -void *aligned_alloc(std::size_t alignment, std::size_t sz, const device &dev, const context &ctxt, cl::sycl::usm::alloc kind) { - int a = 11; - return (void *)(&a); -} -void *aligned_alloc(std::size_t alignment, std::size_t sz, const queue &q, cl::sycl::usm::alloc kind) { - int a = 11; - return (void *)(&a); -} - -void *aligned_alloc_device(std::size_t alignment, std::size_t sz, const device &dev, const context &ctxt) { - int a = 11; - return (void *)(&a); -} -void *aligned_alloc_device(std::size_t alignment, std::size_t sz, const queue &q) { - int a = 11; - return (void *)(&a); -} - -void *aligned_alloc_shared(std::size_t alignment, std::size_t sz, const device &dev, const context &ctxt) { - int a = 11; - return (void *)(&a); -} -void *aligned_alloc_shared(std::size_t alignment, std::size_t sz, const queue &q) { - int a = 11; - return (void *)(&a); -} - -void *aligned_alloc_host(std::size_t alignment, std::size_t sz, const context &ctxt) { - int a = 12; - return (void *)(&a); -} -void *aligned_alloc_host(std::size_t alignment, std::size_t sz, const queue &q) { - int a = 12; - return (void *)(&a); -} +void *malloc(std::size_t sz, const device &dev, const context &ctxt, cl::sycl::usm::alloc kind); +void *malloc(std::size_t sz, const queue &q, cl::sycl::usm::alloc kind); +void *malloc_device(std::size_t sz, const device &dev, const context &ctxt); +void *malloc_device(std::size_t sz, const queue &q); +void *malloc_shared(std::size_t sz, const device &dev, const context &ctxt); +void *malloc_shared(std::size_t sz, const queue &q); +void *malloc_host(std::size_t sz, const context &ctxt); +void *malloc_host(std::size_t sz, const queue &q); +void *aligned_alloc(std::size_t alignment, std::size_t sz, const device &dev, const context &ctxt, cl::sycl::usm::alloc kind); +void *aligned_alloc(std::size_t alignment, std::size_t sz, const queue &q, cl::sycl::usm::alloc kind); +void *aligned_alloc_device(std::size_t alignment, std::size_t sz, const device &dev, const context &ctxt); +void *aligned_alloc_device(std::size_t alignment, std::size_t sz, const queue &q); +void *aligned_alloc_shared(std::size_t alignment, std::size_t sz, const device &dev, const context &ctxt); +void *aligned_alloc_shared(std::size_t alignment, std::size_t sz, const queue &q); +void *aligned_alloc_host(std::size_t alignment, std::size_t sz, const context &ctxt); +void *aligned_alloc_host(std::size_t alignment, std::size_t sz, const queue &q); //template form template @@ -114,15 +54,21 @@ T *malloc_shared(std::size_t Count, const device &Dev, const context &Ctxt) { } // namespace sycl } // namespace cl -void *malloc(std::size_t sz) { - int a = 11; - return (void *)(&a); -} -void *calloc(std::size_t num, std::size_t sz) { - int a = 11; - return (void *)(&a); +void *malloc(std::size_t sz); +void *calloc(std::size_t num, std::size_t sz); +//-- End Mocks + +// User functions that might allocate memory in some way unknown to us. +float *unknownFunc(); + +void allocateUSMByHandle(float **pointerHandle, std::size_t sz, sycl::device &dev, sycl::context &ctxt) { + float *mem = sycl::malloc_shared(sz, dev, ctxt); + *pointerHandle = mem; } -// -- END MOCKS + +struct Mesh { + float a; +}; float calledFromLambda(float *first) { return first[0]; @@ -133,19 +79,37 @@ int main(int argc, char **argv) { int device = 0, context = 0; double queue = 0; - //bad pointers + //-- Declarations + + //-- various bad pointers float stackFloat = 20.0; float *stackFloatP = &stackFloat; //#decl_stackFloatP float *frenemy = stackFloatP; //#decl_frenemy frenemy++; + // expected-warning@+1 {{cast to 'float *' from smaller integer type 'int'}} float *fromParam = (float *)(argc); //#decl_fromParam - // std::string is already caught by 'non-trivially copy constructible' check. + // std::string is already caught by 'non-trivially copy constructible' check. // so we only worry about literal strings. auto stringLiteral = "omgwtf"; //#decl_stringLiteral + float *apocryphal = unknownFunc(); //#decl_apocryphal + float *usmByHandle; + allocateUSMByHandle(&usmByHandle, 10, device, context); + + //-- structs + Mesh stackMesh; + stackMesh.a = 31.0; + Mesh *stackMeshP = &stackMesh; //#decl_stackMeshP + Mesh *stackMeshP2; + Mesh *mallocMeshP = static_cast(malloc(sizeof(Mesh))); //#decl_mallocMeshP + Mesh *mallocMeshP2 = static_cast(malloc(sizeof(Mesh))); //#decl_mallocMeshP2 + Mesh *usmMeshP = static_cast(sycl::malloc_shared(sizeof(Mesh), device, context)); //#decl_usmMeshP + Mesh *usmMeshP2 = sycl::malloc_shared(1, device, context); //#decl_usmMeshP2 + + //-- malloc float *mallocFloatP = static_cast(malloc(sizeof(float) * 2)); //#decl_mallocFloatP float *mallocFloatP2 = static_cast(malloc(sizeof(float) * 2)); //#decl_mallocFloatP2 float *callocFloatP = static_cast(calloc(2, sizeof(float))); //#decl_callocFloatP @@ -178,9 +142,10 @@ int main(int argc, char **argv) { float *usmDeviceP3 = static_cast(sycl::malloc(sizeof(float), device, context, cl::sycl::usm::alloc::device)); float *usmDeviceP4 = static_cast(sycl::malloc(sizeof(float), queue, cl::sycl::usm::alloc::device)); - // --- direct lambda testing --- cl::sycl::kernel_single_task([=]() { - // --- The following dangerous pointer captures result in errors or notes. + // --- Captures + + //-- various bad pointers // expected-note@#call_kernelFunc {{called by 'kernel_single_taska; + + // expected-note@+1 {{Unknown memory reference in SYCL device kernel. Be sure memory was allocated with USM (malloc_shared, etc).}} + stackMeshP2->a = 34.0; + + // expected-note@#decl_mallocMeshP {{Declared here.}} + // expected-error@+1 {{Illegal memory reference in SYCL device kernel. Use USM (malloc_shared, etc) instead.}} + float mmpa = mallocMeshP->a; + + // expected-note@#decl_mallocMeshP2 {{Declared here.}} + // expected-error@+1 {{Illegal memory reference in SYCL device kernel. Use USM (malloc_shared, etc) instead.}} + mallocMeshP2->a = 45.0; + + //-- malloc + // expected-note@#decl_mallocFloatP2 {{Declared here.}} // expected-error@+1 {{Illegal memory reference in SYCL device kernel. Use USM (malloc_shared, etc) instead.}} mallocFloatP2[0] = 80; @@ -215,20 +205,25 @@ int main(int argc, char **argv) { // expected-error@+1 {{Illegal memory reference in SYCL device kernel. Use USM (malloc_shared, etc) instead.}} float someValue = *callocFloatP2; - - - // --- Only the first capture of a pointer emits anything. So these violations will NOT emit redundant diagnostics. + // --- Only the first capture of a pointer emits anything. So these accesses will NOT emit redundant diagnostics. calledFromLambda(mallocFloatP); stackFloatP[0] = 31.0; frenemy[0] = 41.0; fromParam[0] = 71.0; char y = stringLiteral[0]; + apocryphal[0] = 88.1; + float smpa2 = stackMeshP->a; + stackMeshP2->a = 34.2; + float mmpa2 = mallocMeshP->a; + mallocMeshP2->a = 45.2; mallocFloatP2[0] = 81; callocFloatP[0] = 81; float someOtherValue = *callocFloatP2; // --- These captures all use USM, and should pass without any notes or errors. calledFromLambda(usmSharedP); + float umpa = usmMeshP->a; + usmMeshP2->a = 61.0; usmSharedP[0] = 1; usmSharedP2[0] = 1; usmSharedP3[0] = 1; @@ -260,6 +255,12 @@ int main(int argc, char **argv) { frenemy[0] = 40.0; fromParam[0] = 70.0; char x = stringLiteral[0]; + apocryphal[0] = 89.0; + usmByHandle[0] = 90.0; + float smpa = stackMeshP->a; + stackMeshP2->a = 34.0; + float mmpa = mallocMeshP->a; + mallocMeshP2->a = 45.0; mallocFloatP2[0] = 80; callocFloatP[0] = 80; float someValue = *callocFloatP2; @@ -267,6 +268,4 @@ int main(int argc, char **argv) { noProblemLambda(); return 0; -} - -#pragma clang diagnostic pop \ No newline at end of file +} \ No newline at end of file From ac0c2eccfc75a704bd7cea907f0afe226dcd4df8 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 30 Apr 2020 14:37:43 -0700 Subject: [PATCH 08/10] clang-format Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaExpr.cpp | 5 ++-- clang/lib/Sema/SemaSYCL.cpp | 54 ++++++++++++++++++++++--------------- 2 files changed, 35 insertions(+), 24 deletions(-) diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 6f2786662c5ce..9d3a7fc1eac41 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -16322,9 +16322,10 @@ static bool isVariableCapturable(CapturingScopeInfo *CSI, VarDecl *Var, return false; } // SYCL: Emit diagnostics for any illegal pointer derefs. - if(Diagnose && S.getLangOpts().SYCLIsDevice && Var->getType()->isAnyPointerType()) + if (Diagnose && S.getLangOpts().SYCLIsDevice && + Var->getType()->isAnyPointerType()) S.checkSYCLDevicePointerCapture(Var, Loc); - + return true; } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 1c5a07a104cfb..3a11da98591f3 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -284,59 +284,69 @@ void Sema::checkSYCLDeviceVarDecl(VarDecl *Var) { checkSYCLVarType(*this, Ty, Loc, Visited); } -void Sema::checkSYCLDevicePointerCapture(VarDecl *Var, SourceLocation CaptureLoc) { +void Sema::checkSYCLDevicePointerCapture(VarDecl *Var, + SourceLocation CaptureLoc) { + // Any pointer captured into the SYCL kernel lambda will fail when + // dereferenced...except USM. If it weren't for USM we could just emit a + // deferred diagnostic for every pointer capture. Instead, we attempt to + // identify which pointers are USM, and which are definitely not. For those + // that are definitely not, we emit an error. For those that are unknown, we + // emit a gentle note suggesting the user ensure they are using USM. For USM + // pointers, we do nothing. assert(getLangOpts().SYCLIsDevice && "Should only be called during SYCL compilation"); - assert(Var->getType()->isAnyPointerType() && - "Should only be called for pointer types being captured."); + assert(Var->getType()->isAnyPointerType() && + "Should only be called for pointer types being captured."); enum ExprAllocation { Unknown, USM, Not_USM }; - ExprAllocation howAllocated = Unknown; + ExprAllocation howAllocated = Unknown; SourceLocation DecLoc = SourceLocation(); // Try to determine provenance of this Var. const Expr *Init = Var->getAnyInitializer(); - if(Init){ + if (Init) { Init = Init->IgnoreCasts(); DecLoc = Init->getExprLoc(); - + const CallExpr *CE = dyn_cast(Init); - if(CE) { + if (CE) { // Captured pointer is result of function call. const FunctionDecl *func = CE->getDirectCallee(); auto FullName = func->getQualifiedNameAsString(); // Check to see if this function call is one of the USM allocators. - if( (FullName.rfind("cl::sycl::malloc", 0) == 0) || (FullName.rfind("cl::sycl::aligned_alloc", 0) == 0)) + if ((FullName.rfind("cl::sycl::malloc", 0) == 0) || + (FullName.rfind("cl::sycl::aligned_alloc", 0) == 0)) howAllocated = USM; - else if ( (FullName.compare("malloc")==0) || (FullName.compare("calloc")==0) ) + else if ((FullName.compare("malloc") == 0) || + (FullName.compare("calloc") == 0)) howAllocated = Not_USM; } else { - // Var has initialization, but not as return result of a function, disqualify any other obvious bad initialization expressions. + // Var has initialization, but not as return result of a function, + // disqualify any other obvious bad initialization expressions. const StringLiteral *SL = dyn_cast(Init); - if(SL) + if (SL) howAllocated = Not_USM; - - if(Init->isRValue()) // pr-value + + if (Init->isRValue()) // pr-value howAllocated = Not_USM; } - } - //else: Var does not have local initialization, might be parameter, or initialized via ref, etc. + } + // else: Var does not have local initialization, might be parameter, or + // initialized via ref, etc. // Nothing more we can determine at this time. - - //diagnostics - if(howAllocated == Not_USM) + + // diagnostics + if (howAllocated == Not_USM) SYCLDiagIfDeviceCode(CaptureLoc, diag::err_sycl_illegal_memory_reference); else if (howAllocated == Unknown) SYCLDiagIfDeviceCode(CaptureLoc, diag::note_unknown_memory_reference); - - if(howAllocated != USM && DecLoc.isValid()) + + if (howAllocated != USM && DecLoc.isValid()) SYCLDiagIfDeviceCode(DecLoc, diag::note_sycl_capture_declared_here); - } - class MarkDeviceFunction : public RecursiveASTVisitor { public: MarkDeviceFunction(Sema &S) From 2ecd89a865bcc0f719e4cb07080a31e8bbeb86fb Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 30 Apr 2020 15:13:53 -0700 Subject: [PATCH 09/10] updated test that new diagnostic was confusing Signed-off-by: Chris Perkins --- clang/test/SemaSYCL/built-in-type-kernel-arg.cpp | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp index 82cd21bf01552..021a318aaef67 100644 --- a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp +++ b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp @@ -21,10 +21,12 @@ void test(const int some_const) { }); } +int *unknownFunc(); + int main() { int data = 5; - int* data_addr = &data; - int* new_data_addr = nullptr; + int* data_addr = unknownFunc(); //#decl_data_addr + int* new_data_addr = unknownFunc(); //#decl_new_data_addr test_struct s; s.data = data; kernel( @@ -38,6 +40,9 @@ int main() { }); kernel( [=]() { + // expected-note@#decl_new_data_addr {{Declared here.}} + // expected-note@#decl_data_addr {{Declared here.}} + // expected-note@+1 2{{Unknown memory reference in SYCL device kernel. Be sure memory was allocated with USM (malloc_shared, etc).}} new_data_addr[0] = data_addr[0]; }); const int some_const = 10; From 1d566d3d188ef63f5d967c04f4b9dbc3bc1e3fce Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 30 Apr 2020 15:29:43 -0700 Subject: [PATCH 10/10] more clang format Signed-off-by: Chris Perkins --- clang/lib/Sema/SemaExpr.cpp | 1 - clang/lib/Sema/SemaSYCL.cpp | 2 -- clang/test/SemaSYCL/built-in-type-kernel-arg.cpp | 4 ++-- 3 files changed, 2 insertions(+), 5 deletions(-) diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 9d3a7fc1eac41..7bcf0911fffac 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -16254,7 +16254,6 @@ static DeclContext *getParentOfCapturingContextOrNull(DeclContext *DC, VarDecl * return nullptr; } - // Certain capturing entities (lambdas, blocks etc.) are not allowed to capture // certain types of variables (unnamed, variably modified types etc.) // so check for eligibility. diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 3a11da98591f3..a979f40258838 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -28,7 +28,6 @@ #include - using namespace clang; using KernelParamKind = SYCLIntegrationHeader::kernel_param_kind_t; @@ -412,7 +411,6 @@ class MarkDeviceFunction : public RecursiveASTVisitor { return true; } - // The call graph for this translation unit. CallGraph SYCLCG; // The set of functions called by a kernel function. diff --git a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp index 021a318aaef67..734830dc531d6 100644 --- a/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp +++ b/clang/test/SemaSYCL/built-in-type-kernel-arg.cpp @@ -25,8 +25,8 @@ int *unknownFunc(); int main() { int data = 5; - int* data_addr = unknownFunc(); //#decl_data_addr - int* new_data_addr = unknownFunc(); //#decl_new_data_addr + int *data_addr = unknownFunc(); //#decl_data_addr + int *new_data_addr = unknownFunc(); //#decl_new_data_addr test_struct s; s.data = data; kernel(