Skip to content

Commit 6e9fa89

Browse files
authored
[SYCL][NFC] Avoid nullptr dereferencing of YDimExpr and ZDimExpr (#4367)
Add assert that arguments of work_group_size attributes cannot be nullptr Klocwork exposed a (false positive) bug that nullptr derferencing of YDimExpr and ZDimExpr is possible for non sycl:: usage. This is not practically possible because such (OpenCL and cl::) spellings of the attribute require three arguments. This patch 1. Removes the ability of SetDefaultValue lambda to return nullptr 2. Asserts in SetDefaultValue that it is not possible for sycl:: case to have NULL arguments 3. Moves error checking for three arguments for OpenCL and cl:: cases ahead of the assert 4. Removes check for intel::reqd_work_group_size spelling
1 parent d6bf2b6 commit 6e9fa89

File tree

1 file changed

+18
-23
lines changed

1 file changed

+18
-23
lines changed

clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 18 additions & 23 deletions
Original file line numberDiff line numberDiff line change
@@ -3141,29 +3141,6 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
31413141
return;
31423142

31433143
S.CheckDeprecatedSYCLAttributeSpelling(AL);
3144-
3145-
Expr *XDimExpr = AL.getArgAsExpr(0);
3146-
3147-
// If no attribute argument is specified, set the second and third argument
3148-
// to the default value 1, but only if the sycl:: or intel::
3149-
// reqd_work_group_size spelling was used.
3150-
auto SetDefaultValue = [](Sema &S, const ParsedAttr &AL, SourceLocation loc) {
3151-
Expr *E =
3152-
(AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && AL.hasScope() &&
3153-
(AL.getScopeName()->isStr("sycl") ||
3154-
AL.getScopeName()->isStr("intel")))
3155-
? IntegerLiteral::Create(S.Context, llvm::APInt(32, 1),
3156-
S.Context.IntTy, AL.getLoc())
3157-
: nullptr;
3158-
return E;
3159-
};
3160-
3161-
Expr *YDimExpr = AL.isArgExpr(1) ? AL.getArgAsExpr(1)
3162-
: SetDefaultValue(S, AL, AL.getLoc());
3163-
3164-
Expr *ZDimExpr = AL.isArgExpr(2) ? AL.getArgAsExpr(2)
3165-
: SetDefaultValue(S, AL, AL.getLoc());
3166-
31673144
// __attribute__((reqd_work_group_size)), [[cl::reqd_work_group_size]], and
31683145
// [[intel::max_work_group_size]] all require exactly three arguments.
31693146
if ((AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize &&
@@ -3175,6 +3152,24 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
31753152
return;
31763153
}
31773154

3155+
Expr *XDimExpr = AL.getArgAsExpr(0);
3156+
3157+
// If no attribute argument is specified, set the second and third argument
3158+
// to the default value 1, but only if the sycl::reqd_work_group_size
3159+
// spelling was used.
3160+
auto SetDefaultValue = [](Sema &S, const ParsedAttr &AL) {
3161+
assert(AL.getKind() == ParsedAttr::AT_ReqdWorkGroupSize && AL.hasScope() &&
3162+
AL.getScopeName()->isStr("sycl"));
3163+
return IntegerLiteral::Create(S.Context, llvm::APInt(32, 1),
3164+
S.Context.IntTy, AL.getLoc());
3165+
};
3166+
3167+
Expr *YDimExpr =
3168+
AL.isArgExpr(1) ? AL.getArgAsExpr(1) : SetDefaultValue(S, AL);
3169+
3170+
Expr *ZDimExpr =
3171+
AL.isArgExpr(2) ? AL.getArgAsExpr(2) : SetDefaultValue(S, AL);
3172+
31783173
ASTContext &Ctx = S.getASTContext();
31793174

31803175
if (!XDimExpr->isValueDependent() && !YDimExpr->isValueDependent() &&

0 commit comments

Comments
 (0)