From adf91fa8e8b9f56a6ab826de60fec906698ab2d6 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Wed, 18 Aug 2021 13:40:19 -0700 Subject: [PATCH 1/9] [SYCL] Avoid nullptr dereferencing of YDimExpr and ZDimExpr Fix the Klocworks exposed bug for non sycl:: or intel:: usage by refactoring and adding checks to avoid nullptr dereferencing --- clang/lib/Sema/SemaDeclAttr.cpp | 19 +++++++++++++------ 1 file changed, 13 insertions(+), 6 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index bcbac73fd300c..110053825b215 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3177,25 +3177,32 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { ASTContext &Ctx = S.getASTContext(); - if (!XDimExpr->isValueDependent() && !YDimExpr->isValueDependent() && - !ZDimExpr->isValueDependent()) { - llvm::APSInt XDimVal, YDimVal, ZDimVal; - ExprResult XDim = S.VerifyIntegerConstantExpression(XDimExpr, &XDimVal); - ExprResult YDim = S.VerifyIntegerConstantExpression(YDimExpr, &YDimVal); - ExprResult ZDim = S.VerifyIntegerConstantExpression(ZDimExpr, &ZDimVal); + bool YDimExprIsDereferencable = YDimExpr && (!YDimExpr->isValueDependent()); + bool ZDimExprIsDereferencable = ZDimExpr && (!ZDimExpr->isValueDependent()); + llvm::APSInt XDimVal, YDimVal, ZDimVal; + if (!XDimExpr->isValueDependent()) { + ExprResult XDim = S.VerifyIntegerConstantExpression(XDimExpr, &XDimVal); if (XDim.isInvalid()) return; XDimExpr = XDim.get(); + } + if (YDimExprIsDereferencable) { + ExprResult YDim = S.VerifyIntegerConstantExpression(YDimExpr, &YDimVal); if (YDim.isInvalid()) return; YDimExpr = YDim.get(); + } + if (ZDimExprIsDereferencable) { + ExprResult ZDim = S.VerifyIntegerConstantExpression(ZDimExpr, &ZDimVal); if (ZDim.isInvalid()) return; ZDimExpr = ZDim.get(); + } + if ((XDimVal >= 1) && (YDimVal >= 1) && (ZDimVal >= 1)) { // If the num_simd_work_items attribute is specified on a declaration it // must evenly divide the index that increments fastest in the // reqd_work_group_size attribute. In OpenCL, the first argument increments From e9f8624b342238dd9fda95e63bd70b661c545b92 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 24 Aug 2021 11:07:16 -0700 Subject: [PATCH 2/9] Assert that YDimExpr and ZDimExpr can't practically be nullptr --- clang/lib/Sema/SemaDeclAttr.cpp | 34 +++++++++++++-------------------- 1 file changed, 13 insertions(+), 21 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 110053825b215..4c3235e67ac58 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3148,13 +3148,13 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { // to the default value 1, but only if the sycl:: or intel:: // reqd_work_group_size spelling was used. auto SetDefaultValue = [](Sema &S, const ParsedAttr &AL, SourceLocation loc) { - Expr *E = - (AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && AL.hasScope() && - (AL.getScopeName()->isStr("sycl") || - AL.getScopeName()->isStr("intel"))) - ? IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), - S.Context.IntTy, AL.getLoc()) - : nullptr; + assert((AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && AL.hasScope() && + (AL.getScopeName()->isStr("sycl") || + AL.getScopeName()->isStr("intel"))) && + "Attribute does not exist in sycl:: or intel:: scope"); + + Expr *E = IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), + S.Context.IntTy, AL.getLoc()); return E; }; @@ -3176,33 +3176,25 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { } ASTContext &Ctx = S.getASTContext(); - - bool YDimExprIsDereferencable = YDimExpr && (!YDimExpr->isValueDependent()); - bool ZDimExprIsDereferencable = ZDimExpr && (!ZDimExpr->isValueDependent()); - - llvm::APSInt XDimVal, YDimVal, ZDimVal; - if (!XDimExpr->isValueDependent()) { + if (!XDimExpr->isValueDependent() && !YDimExpr->isValueDependent() && + !ZDimExpr->isValueDependent()) { + llvm::APSInt XDimVal, YDimVal, ZDimVal; ExprResult XDim = S.VerifyIntegerConstantExpression(XDimExpr, &XDimVal); + ExprResult YDim = S.VerifyIntegerConstantExpression(YDimExpr, &YDimVal); + ExprResult ZDim = S.VerifyIntegerConstantExpression(ZDimExpr, &ZDimVal); + if (XDim.isInvalid()) return; XDimExpr = XDim.get(); - } - if (YDimExprIsDereferencable) { - ExprResult YDim = S.VerifyIntegerConstantExpression(YDimExpr, &YDimVal); if (YDim.isInvalid()) return; YDimExpr = YDim.get(); - } - if (ZDimExprIsDereferencable) { - ExprResult ZDim = S.VerifyIntegerConstantExpression(ZDimExpr, &ZDimVal); if (ZDim.isInvalid()) return; ZDimExpr = ZDim.get(); - } - if ((XDimVal >= 1) && (YDimVal >= 1) && (ZDimVal >= 1)) { // If the num_simd_work_items attribute is specified on a declaration it // must evenly divide the index that increments fastest in the // reqd_work_group_size attribute. In OpenCL, the first argument increments From 5dbf9d5f5286447c1799f31a2ecfddb0f5fbb2e6 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 24 Aug 2021 11:16:52 -0700 Subject: [PATCH 3/9] Remove local variable E; return literal directly Fix whitespace --- clang/lib/Sema/SemaDeclAttr.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 4c3235e67ac58..eb0144c59fad8 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3153,9 +3153,8 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { AL.getScopeName()->isStr("intel"))) && "Attribute does not exist in sycl:: or intel:: scope"); - Expr *E = IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), - S.Context.IntTy, AL.getLoc()); - return E; + return (IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), + S.Context.IntTy, AL.getLoc())); }; Expr *YDimExpr = AL.isArgExpr(1) ? AL.getArgAsExpr(1) @@ -3176,6 +3175,7 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { } ASTContext &Ctx = S.getASTContext(); + if (!XDimExpr->isValueDependent() && !YDimExpr->isValueDependent() && !ZDimExpr->isValueDependent()) { llvm::APSInt XDimVal, YDimVal, ZDimVal; From 690d6cec01d32e42319ba828a0f77e85a656bfa8 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 24 Aug 2021 11:23:46 -0700 Subject: [PATCH 4/9] Fix parenthesis --- clang/lib/Sema/SemaDeclAttr.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index eb0144c59fad8..65de2b143508f 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3153,8 +3153,8 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { AL.getScopeName()->isStr("intel"))) && "Attribute does not exist in sycl:: or intel:: scope"); - return (IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), - S.Context.IntTy, AL.getLoc())); + return IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), + S.Context.IntTy, AL.getLoc()); }; Expr *YDimExpr = AL.isArgExpr(1) ? AL.getArgAsExpr(1) From 11f0a79a0a3d66603439fd8a5a8cfd1dc6b77522 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 2 Sep 2021 15:11:11 -0700 Subject: [PATCH 5/9] Assert that YDimExpr or ZDimExpr cannot be null at a different place Do it after checking for 3 argument requirement for OpenCL and cl:: spellings --- clang/lib/Sema/SemaDeclAttr.cpp | 16 ++++++++++------ 1 file changed, 10 insertions(+), 6 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 65de2b143508f..f4f336c867c9c 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3148,13 +3148,15 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { // to the default value 1, but only if the sycl:: or intel:: // reqd_work_group_size spelling was used. auto SetDefaultValue = [](Sema &S, const ParsedAttr &AL, SourceLocation loc) { - assert((AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && AL.hasScope() && - (AL.getScopeName()->isStr("sycl") || - AL.getScopeName()->isStr("intel"))) && - "Attribute does not exist in sycl:: or intel:: scope"); + Expr *E = + (AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && AL.hasScope() && + (AL.getScopeName()->isStr("sycl") || + AL.getScopeName()->isStr("intel"))) + ? IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), + S.Context.IntTy, AL.getLoc()) + : nullptr; - return IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), - S.Context.IntTy, AL.getLoc()); + return E; }; Expr *YDimExpr = AL.isArgExpr(1) ? AL.getArgAsExpr(1) @@ -3176,6 +3178,8 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { ASTContext &Ctx = S.getASTContext(); + assert(YDimExpr && "YDimExpr cannot be NULL"); + assert(ZDimExpr && "ZDimExpr cannot be NULL"); if (!XDimExpr->isValueDependent() && !YDimExpr->isValueDependent() && !ZDimExpr->isValueDependent()) { llvm::APSInt XDimVal, YDimVal, ZDimVal; From 3d4000a4b483dcad6a0b808de14e2f32e13c34c5 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Thu, 2 Sep 2021 15:15:30 -0700 Subject: [PATCH 6/9] Remove whitespace --- clang/lib/Sema/SemaDeclAttr.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index f4f336c867c9c..e5ea254c05bb5 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3155,7 +3155,6 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { ? IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), S.Context.IntTy, AL.getLoc()) : nullptr; - return E; }; From 7f49a73eb6bbac3d0cdeb50a7462a9ef04592ba0 Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 7 Sep 2021 14:16:33 -0700 Subject: [PATCH 7/9] Remove the ability of SetDefaultValue lambda to return nullptr Assert in SetDefaultValue that it is not possible for sycl:: or intel:: cases to have NULL arguments Move error checking of three arguments for OpenCL and cl:: cases ahead of the assert in SetDefaultValue --- clang/lib/Sema/SemaDeclAttr.cpp | 45 +++++++++++++++------------------ 1 file changed, 20 insertions(+), 25 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index e5ea254c05bb5..6bf5e0ce27419 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3141,29 +3141,6 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { return; S.CheckDeprecatedSYCLAttributeSpelling(AL); - - Expr *XDimExpr = AL.getArgAsExpr(0); - - // If no attribute argument is specified, set the second and third argument - // to the default value 1, but only if the sycl:: or intel:: - // reqd_work_group_size spelling was used. - auto SetDefaultValue = [](Sema &S, const ParsedAttr &AL, SourceLocation loc) { - Expr *E = - (AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && AL.hasScope() && - (AL.getScopeName()->isStr("sycl") || - AL.getScopeName()->isStr("intel"))) - ? IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), - S.Context.IntTy, AL.getLoc()) - : nullptr; - return E; - }; - - Expr *YDimExpr = AL.isArgExpr(1) ? AL.getArgAsExpr(1) - : SetDefaultValue(S, AL, AL.getLoc()); - - Expr *ZDimExpr = AL.isArgExpr(2) ? AL.getArgAsExpr(2) - : SetDefaultValue(S, AL, AL.getLoc()); - // __attribute__((reqd_work_group_size)), [[cl::reqd_work_group_size]], and // [[intel::max_work_group_size]] all require exactly three arguments. if ((AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && @@ -3175,10 +3152,28 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { return; } + Expr *XDimExpr = AL.getArgAsExpr(0); + + // If no attribute argument is specified, set the second and third argument + // to the default value 1, but only if the sycl:: or intel:: + // reqd_work_group_size spelling was used. + // + auto SetDefaultValue = [](Sema &S, const ParsedAttr &AL) { + assert(AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && AL.hasScope() && + (AL.getScopeName()->isStr("sycl") || + AL.getScopeName()->isStr("intel"))); + return IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), + S.Context.IntTy, AL.getLoc()); + }; + + Expr *YDimExpr = + AL.isArgExpr(1) ? AL.getArgAsExpr(1) : SetDefaultValue(S, AL); + + Expr *ZDimExpr = + AL.isArgExpr(2) ? AL.getArgAsExpr(2) : SetDefaultValue(S, AL); + ASTContext &Ctx = S.getASTContext(); - assert(YDimExpr && "YDimExpr cannot be NULL"); - assert(ZDimExpr && "ZDimExpr cannot be NULL"); if (!XDimExpr->isValueDependent() && !YDimExpr->isValueDependent() && !ZDimExpr->isValueDependent()) { llvm::APSInt XDimVal, YDimVal, ZDimVal; From 5ae220f326f76afba481d5b4ee58e9beb6f24ffc Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Tue, 7 Sep 2021 15:18:04 -0700 Subject: [PATCH 8/9] Rebase and remove checks for intel::reqd_work_group_size spelling --- clang/lib/Sema/SemaDeclAttr.cpp | 7 +++---- 1 file changed, 3 insertions(+), 4 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index dd80ef7d1b6d6..39ba640ab7416 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3155,13 +3155,12 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { Expr *XDimExpr = AL.getArgAsExpr(0); // If no attribute argument is specified, set the second and third argument - // to the default value 1, but only if the sycl:: or intel:: - // reqd_work_group_size spelling was used. + // to the default value 1, but only if the sycl::reqd_work_group_size + // spelling was used. // auto SetDefaultValue = [](Sema &S, const ParsedAttr &AL) { assert(AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && AL.hasScope() && - (AL.getScopeName()->isStr("sycl") || - AL.getScopeName()->isStr("intel"))); + AL.getScopeName()->isStr("sycl")); return IntegerLiteral::Create(S.Context, llvm::APInt(32, 1), S.Context.IntTy, AL.getLoc()); }; From e114ef6f1d5c8cf4f05dbe2bed549f51406b09bd Mon Sep 17 00:00:00 2001 From: Sindhu Chittireddy Date: Wed, 8 Sep 2021 09:19:31 -0700 Subject: [PATCH 9/9] Remove extra `//` --- clang/lib/Sema/SemaDeclAttr.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 39ba640ab7416..384ef396e4ca8 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3157,7 +3157,6 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { // If no attribute argument is specified, set the second and third argument // to the default value 1, but only if the sycl::reqd_work_group_size // spelling was used. - // auto SetDefaultValue = [](Sema &S, const ParsedAttr &AL) { assert(AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && AL.hasScope() && AL.getScopeName()->isStr("sycl"));