[OpenACC] Implement SubArray Parsing/Sema #90796
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
This implementation takes quite a bit from the OMP implementation of array sections, but only has to enforce the rules as applicable to OpenACC. Additionally, it does its best to create an AST node (with the assistance of RecoveryExprs) with as much checking done as soon as possible in the case of instantiations.
llvmbot
added
clang
|
@llvm/pr-subscribers-clang Author: Erich Keane (erichkeane) ChangesThis implementation takes quite a bit from the OMP implementation of array sections, but only has to enforce the rules as applicable to OpenACC. Additionally, it does its best to create an AST node (with the assistance of RecoveryExprs) with as much checking done as soon as possible in the case of instantiations. Patch is 60.81 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/90796.diff 12 Files Affected:
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 4b074b853bfe65..fbb107679d94ba 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12291,8 +12291,8 @@ def warn_acc_if_self_conflict
"evaluates to true">,
InGroup<DiagGroup<"openacc-self-if-potential-conflict">>;
def err_acc_int_expr_requires_integer
- : Error<"OpenACC %select{clause|directive}0 '%1' requires expression of "
- "integer type (%2 invalid)">;
+ : Error<"OpenACC %select{clause '%1'|directive '%2'|sub-array bound}0 "
+ "requires expression of integer type (%3 invalid)">;
def err_acc_int_expr_incomplete_class_type
: Error<"OpenACC integer expression has incomplete class type %0">;
def err_acc_int_expr_explicit_conversion
@@ -12310,4 +12310,24 @@ def err_acc_num_gangs_num_args
def err_acc_not_a_var_ref
: Error<"OpenACC variable is not a valid variable name, sub-array, array "
"element, or composite variable member">;
+def err_acc_typecheck_subarray_value
+ : Error<"OpenACC sub-array subscripted value is not an array or pointer">;
+def err_acc_subarray_function_type
+ : Error<"OpenACC sub-array cannot be of function type %0">;
+def err_acc_subarray_incomplete_type
+ : Error<"OpenACC sub-array base is of incomplete type %0">;
+def err_acc_subarray_no_length
+ : Error<"OpenACC sub-array length is unspecified and cannot be inferred "
+ "because the subscripted value is %select{not an array|an array of "
+ "unknown bound}0">;
+def err_acc_subarray_negative
+ : Error<"OpenACC sub-array %select{lower bound|length}0 evaluated to "
+ "negative value %1">;
+def err_acc_subarray_out_of_range
+ : Error<"OpenACC sub-array %select{lower bound|length}0 evaluated to a "
+ "value (%1) that would be out of the range of the subscripted "
+ "array size of %2">;
+def err_acc_subarray_base_plus_length_out_of_range
+ : Error<"OpenACC sub-array specified range [%0:%1] would be out of the "
+ "range of the subscripted array size of %2">;
} // end of sema component.
diff --git a/clang/lib/Parse/ParseExpr.cpp b/clang/lib/Parse/ParseExpr.cpp
index 7d6febb04a82c4..5f5f9a79c8c408 100644
--- a/clang/lib/Parse/ParseExpr.cpp
+++ b/clang/lib/Parse/ParseExpr.cpp
@@ -2039,7 +2039,8 @@ Parser::ParsePostfixExpressionSuffix(ExprResult LHS) {
if (Tok.is(tok::colon)) {
// Consume ':'
ColonLocFirst = ConsumeToken();
- Length = Actions.CorrectDelayedTyposInExpr(ParseExpression());
+ if (Tok.isNot(tok::r_square))
+ Length = Actions.CorrectDelayedTyposInExpr(ParseExpression());
}
} else if (ArgExprs.size() <= 1 && getLangOpts().OpenMP) {
ColonProtectionRAIIObject RAII(*this);
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 3ea81e0497c203..a203ece7ff3130 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -16,6 +16,7 @@
#include "clang/Basic/DiagnosticSema.h"
#include "clang/Basic/OpenACCKinds.h"
#include "clang/Sema/Sema.h"
+#include "llvm/ADT/StringExtras.h"
#include "llvm/Support/Casting.h"
using namespace clang;
@@ -367,7 +368,9 @@ ExprResult SemaOpenACC::ActOnIntExpr(OpenACCDirectiveKind DK,
assert(((DK != OpenACCDirectiveKind::Invalid &&
CK == OpenACCClauseKind::Invalid) ||
(DK == OpenACCDirectiveKind::Invalid &&
- CK != OpenACCClauseKind::Invalid)) &&
+ CK != OpenACCClauseKind::Invalid) ||
+ (DK == OpenACCDirectiveKind::Invalid &&
+ CK == OpenACCClauseKind::Invalid)) &&
"Only one of directive or clause kind should be provided");
class IntExprConverter : public Sema::ICEConvertDiagnoser {
@@ -375,6 +378,16 @@ ExprResult SemaOpenACC::ActOnIntExpr(OpenACCDirectiveKind DK,
OpenACCClauseKind ClauseKind;
Expr *IntExpr;
+ // gets the index into the diagnostics so we can use this for clauses,
+ // directives, and sub array.s
+ unsigned getDiagKind() const {
+ if (ClauseKind != OpenACCClauseKind::Invalid)
+ return 0;
+ if (DirectiveKind != OpenACCDirectiveKind::Invalid)
+ return 1;
+ return 2;
+ }
+
public:
IntExprConverter(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
Expr *IntExpr)
@@ -390,12 +403,8 @@ ExprResult SemaOpenACC::ActOnIntExpr(OpenACCDirectiveKind DK,
}
SemaBase::SemaDiagnosticBuilder diagnoseNotInt(Sema &S, SourceLocation Loc,
QualType T) override {
- if (ClauseKind != OpenACCClauseKind::Invalid)
- return S.Diag(Loc, diag::err_acc_int_expr_requires_integer) <<
- /*Clause=*/0 << ClauseKind << T;
-
- return S.Diag(Loc, diag::err_acc_int_expr_requires_integer) <<
- /*Directive=*/1 << DirectiveKind << T;
+ return S.Diag(Loc, diag::err_acc_int_expr_requires_integer)
+ << getDiagKind() << ClauseKind << DirectiveKind << T;
}
SemaBase::SemaDiagnosticBuilder
@@ -503,12 +512,211 @@ ExprResult SemaOpenACC::ActOnArraySectionExpr(Expr *Base, SourceLocation LBLoc,
SourceLocation RBLoc) {
ASTContext &Context = getASTContext();
- // TODO OpenACC: We likely have to reproduce a lot of the same logic from the
- // OMP version of this, but at the moment we don't have a good way to test it,
- // so for now we'll just create the node.
+ // Handle placeholders.
+ if (Base->hasPlaceholderType() &&
+ !Base->hasPlaceholderType(BuiltinType::ArraySection)) {
+ ExprResult Result = SemaRef.CheckPlaceholderExpr(Base);
+ if (Result.isInvalid())
+ return ExprError();
+ Base = Result.get();
+ }
+ if (LowerBound && LowerBound->getType()->isNonOverloadPlaceholderType()) {
+ ExprResult Result = SemaRef.CheckPlaceholderExpr(LowerBound);
+ if (Result.isInvalid())
+ return ExprError();
+ Result = SemaRef.DefaultLvalueConversion(Result.get());
+ if (Result.isInvalid())
+ return ExprError();
+ LowerBound = Result.get();
+ }
+ if (Length && Length->getType()->isNonOverloadPlaceholderType()) {
+ ExprResult Result = SemaRef.CheckPlaceholderExpr(Length);
+ if (Result.isInvalid())
+ return ExprError();
+ Result = SemaRef.DefaultLvalueConversion(Result.get());
+ if (Result.isInvalid())
+ return ExprError();
+ Length = Result.get();
+ }
+
+ // Check the 'base' value, it must be an array or pointer type, and not to/of
+ // a function type.
+ QualType OriginalBaseTy = ArraySectionExpr::getBaseOriginalType(Base);
+ QualType ResultTy;
+ if (!Base->isTypeDependent()) {
+ if (OriginalBaseTy->isAnyPointerType()) {
+ ResultTy = OriginalBaseTy->getPointeeType();
+ } else if (OriginalBaseTy->isArrayType()) {
+ ResultTy = OriginalBaseTy->getAsArrayTypeUnsafe()->getElementType();
+ } else {
+ return ExprError(
+ Diag(Base->getExprLoc(), diag::err_acc_typecheck_subarray_value)
+ << Base->getSourceRange());
+ }
+
+ if (ResultTy->isFunctionType()) {
+ Diag(Base->getExprLoc(), diag::err_acc_subarray_function_type)
+ << ResultTy << Base->getSourceRange();
+ return ExprError();
+ }
+
+ if (SemaRef.RequireCompleteType(Base->getExprLoc(), ResultTy,
+ diag::err_acc_subarray_incomplete_type,
+ Base))
+ return ExprError();
+
+ if (!Base->hasPlaceholderType(BuiltinType::ArraySection)) {
+ ExprResult Result = SemaRef.DefaultFunctionArrayLvalueConversion(Base);
+ if (Result.isInvalid())
+ return ExprError();
+ Base = Result.get();
+ }
+ }
+
+ auto GetRecovery = [&](Expr *E, QualType Ty) {
+ ExprResult Recovery =
+ SemaRef.CreateRecoveryExpr(E->getBeginLoc(), E->getEndLoc(), E, Ty);
+ return Recovery.isUsable() ? Recovery.get() : nullptr;
+ };
+
+ // Ensure both of the expressions are int-exprs.
+ if (LowerBound && !LowerBound->isTypeDependent()) {
+ ExprResult LBRes =
+ ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Invalid,
+ LowerBound->getExprLoc(), LowerBound);
+
+ if (LBRes.isUsable())
+ LBRes = SemaRef.DefaultLvalueConversion(LBRes.get());
+ LowerBound =
+ LBRes.isUsable() ? LBRes.get() : GetRecovery(LowerBound, Context.IntTy);
+ }
+
+ if (Length && !Length->isTypeDependent()) {
+ ExprResult LenRes =
+ ActOnIntExpr(OpenACCDirectiveKind::Invalid, OpenACCClauseKind::Invalid,
+ Length->getExprLoc(), Length);
+
+ if (LenRes.isUsable())
+ LenRes = SemaRef.DefaultLvalueConversion(LenRes.get());
+ Length =
+ LenRes.isUsable() ? LenRes.get() : GetRecovery(Length, Context.IntTy);
+ }
+
+ // Length is required if the base type is not an array of known bounds.
+ if (!Length && (OriginalBaseTy.isNull() ||
+ (!OriginalBaseTy->isDependentType() &&
+ !OriginalBaseTy->isConstantArrayType() &&
+ !OriginalBaseTy->isDependentSizedArrayType()))) {
+ bool IsArray = !OriginalBaseTy.isNull() && OriginalBaseTy->isArrayType();
+ Diag(ColonLoc, diag::err_acc_subarray_no_length) << IsArray;
+ // Fill in a dummy 'length' so that when we instantiate this we don't
+ // double-diagnose here.
+ ExprResult Recovery = SemaRef.CreateRecoveryExpr(
+ ColonLoc, SourceLocation(), ArrayRef<Expr *>{std::nullopt},
+ Context.IntTy);
+ Length = Recovery.isUsable() ? Recovery.get() : nullptr;
+ }
+
+ // Check the values of each of the arguments, they cannot be negative(we
+ // assume), and if the array bound is known, must be within range. As we do
+ // so, do our best to continue with evaluation, we can set the
+ // value/expression to nullptr/nullopt if they are invalid, and treat them as
+ // not present for the rest of evaluation.
+
+ // We don't have to check for dependence, because the dependent size is
+ // represented as a different AST node.
+ std::optional<llvm::APSInt> BaseSize;
+ if (!OriginalBaseTy.isNull() && OriginalBaseTy->isConstantArrayType()) {
+ const auto *ArrayTy = Context.getAsConstantArrayType(OriginalBaseTy);
+ BaseSize = ArrayTy->getSize();
+ }
+
+ auto GetBoundValue = [&](Expr *E) -> std::optional<llvm::APSInt> {
+ if (!E || E->isInstantiationDependent())
+ return std::nullopt;
+
+ Expr::EvalResult Res;
+ if (!E->EvaluateAsInt(Res, Context))
+ return std::nullopt;
+ return Res.Val.getInt();
+ };
+
+ std::optional<llvm::APSInt> LowerBoundValue = GetBoundValue(LowerBound);
+ std::optional<llvm::APSInt> LengthValue = GetBoundValue(Length);
+
+ // Check lower bound for negative or out of range.
+ if (LowerBoundValue.has_value()) {
+ if (LowerBoundValue->isNegative()) {
+ Diag(LowerBound->getExprLoc(), diag::err_acc_subarray_negative)
+ << /*LowerBound=*/0 << toString(*LowerBoundValue, /*Radix=*/10);
+ LowerBoundValue.reset();
+ LowerBound = GetRecovery(LowerBound, LowerBound->getType());
+ } else if (BaseSize.has_value() &&
+ llvm::APSInt::compareValues(*LowerBoundValue, *BaseSize) >= 0) {
+ // Lower bound (start index) must be less than the size of the array.
+ Diag(LowerBound->getExprLoc(), diag::err_acc_subarray_out_of_range)
+ << /*LowerBound=*/0 << toString(*LowerBoundValue, /*Radix=*/10)
+ << toString(*BaseSize, /*Radix=*/10);
+ LowerBoundValue.reset();
+ LowerBound = GetRecovery(LowerBound, LowerBound->getType());
+ }
+ }
+
+ // Check length for negative or out of range.
+ if (LengthValue.has_value()) {
+ if (LengthValue->isNegative()) {
+ Diag(Length->getExprLoc(), diag::err_acc_subarray_negative)
+ << /*Length=*/1 << toString(*LengthValue, /*Radix=*/10);
+ LengthValue.reset();
+ Length = GetRecovery(Length, Length->getType());
+ } else if (BaseSize.has_value() &&
+ llvm::APSInt::compareValues(*LengthValue, *BaseSize) > 0) {
+ // Length must be lessthan or EQUAL to the size of the array.
+ Diag(Length->getExprLoc(), diag::err_acc_subarray_out_of_range)
+ << /*Length=*/1 << toString(*LengthValue, /*Radix=*/10)
+ << toString(*BaseSize, /*Radix=*/10);
+ LengthValue.reset();
+ Length = GetRecovery(Length, Length->getType());
+ }
+ }
+
+ // Adding two APSInts requires matching sign, so extract that here.
+ auto AddAPSInt = [](llvm::APSInt LHS, llvm::APSInt RHS) -> llvm::APSInt {
+ if (LHS.isSigned() == RHS.isSigned())
+ return LHS + RHS;
+
+ unsigned width = std::max(LHS.getBitWidth(), RHS.getBitWidth()) + 1;
+ return llvm::APSInt(LHS.sext(width) + RHS.sext(width), /*Signed=*/true);
+ };
+
+ // If we know all 3 values, we can diagnose that the total value would be out
+ // of range.
+ if (BaseSize.has_value() && LowerBoundValue.has_value() &&
+ LengthValue.has_value() &&
+ llvm::APSInt::compareValues(AddAPSInt(*LowerBoundValue, *LengthValue),
+ *BaseSize) > 0) {
+ Diag(Base->getExprLoc(),
+ diag::err_acc_subarray_base_plus_length_out_of_range)
+ << toString(*LowerBoundValue, /*Radix=*/10)
+ << toString(*LengthValue, /*Radix=*/10)
+ << toString(*BaseSize, /*Radix=*/10);
+
+ LowerBoundValue.reset();
+ LowerBound = GetRecovery(LowerBound, LowerBound->getType());
+ LengthValue.reset();
+ Length = GetRecovery(Length, Length->getType());
+ }
+
+ // If any part of the expression is dependent, return a dependent sub-array.
+ QualType ArrayExprTy = Context.ArraySectionTy;
+ if (Base->isTypeDependent() ||
+ (LowerBound && LowerBound->isInstantiationDependent()) ||
+ (Length && Length->isInstantiationDependent()))
+ ArrayExprTy = Context.DependentTy;
+
return new (Context)
- ArraySectionExpr(Base, LowerBound, Length, Context.ArraySectionTy,
- VK_LValue, OK_Ordinary, ColonLoc, RBLoc);
+ ArraySectionExpr(Base, LowerBound, Length, ArrayExprTy, VK_LValue,
+ OK_Ordinary, ColonLoc, RBLoc);
}
bool SemaOpenACC::ActOnStartStmtDirective(OpenACCDirectiveKind K,
diff --git a/clang/test/ParserOpenACC/parse-cache-construct.c b/clang/test/ParserOpenACC/parse-cache-construct.c
index de26fc2b277a6b..8937aa095d5eaa 100644
--- a/clang/test/ParserOpenACC/parse-cache-construct.c
+++ b/clang/test/ParserOpenACC/parse-cache-construct.c
@@ -134,9 +134,8 @@ void func() {
}
for (int i = 0; i < 10; ++i) {
- // expected-error@+2{{expected expression}}
// expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
- #pragma acc cache(readonly:ArrayPtr[5:])
+ #pragma acc cache(readonly:ArrayPtr[5:1])
}
for (int i = 0; i < 10; ++i) {
diff --git a/clang/test/ParserOpenACC/parse-cache-construct.cpp b/clang/test/ParserOpenACC/parse-cache-construct.cpp
index f1c71e8b584786..374fe2697b63fc 100644
--- a/clang/test/ParserOpenACC/parse-cache-construct.cpp
+++ b/clang/test/ParserOpenACC/parse-cache-construct.cpp
@@ -74,12 +74,12 @@ void use() {
for (int i = 0; i < 10; ++i) {
// expected-error@+2{{OpenACC sub-array is not allowed here}}
// expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
- #pragma acc cache(Arrs.MemArr[3:4].array[1:4])
+ #pragma acc cache(Arrs.MemArr[2:1].array[1:4])
}
for (int i = 0; i < 10; ++i) {
// expected-error@+2{{OpenACC sub-array is not allowed here}}
// expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
- #pragma acc cache(Arrs.MemArr[3:4].array[4])
+ #pragma acc cache(Arrs.MemArr[2:1].array[4])
}
for (int i = 0; i < 10; ++i) {
// expected-error@+3{{expected ']'}}
@@ -88,7 +88,7 @@ void use() {
#pragma acc cache(Arrs.MemArr[3:4:].array[4])
}
for (int i = 0; i < 10; ++i) {
- // expected-error@+2{{expected expression}}
+ // expected-error@+2{{OpenACC sub-array is not allowed here}}
// expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
#pragma acc cache(Arrs.MemArr[:].array[4])
}
@@ -105,7 +105,7 @@ void use() {
#pragma acc cache(Arrs.MemArr[: :].array[4])
}
for (int i = 0; i < 10; ++i) {
- // expected-error@+2{{expected expression}}
+ // expected-error@+2{{OpenACC sub-array is not allowed here}}
// expected-warning@+1{{OpenACC construct 'cache' not yet implemented, pragma ignored}}
#pragma acc cache(Arrs.MemArr[3:].array[4])
}
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 8a439a5ccd4bdc..fa43f42585fc2b 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -499,7 +499,6 @@ void VarListClauses() {
#pragma acc serial copy(HasMem.MemArr[1:3].array[1:2]), seq
for(;;){}
- // expected-error@+3{{expected expression}}
// expected-warning@+2{{OpenACC clause 'copy' not yet implemented, clause ignored}}
// expected-warning@+1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
#pragma acc serial copy(HasMem.MemArr[:]), seq
@@ -519,7 +518,6 @@ void VarListClauses() {
#pragma acc serial copy(HasMem.MemArr[: :]), seq
for(;;){}
- // expected-error@+3{{expected expression}}
// expected-warning@+2{{OpenACC clause 'copy' not yet implemented, clause ignored}}
// expected-warning@+1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
#pragma acc serial copy(HasMem.MemArr[3:]), seq
diff --git a/clang/test/ParserOpenACC/parse-sub-array.cpp b/clang/test/ParserOpenACC/parse-sub-array.cpp
new file mode 100644
index 00000000000000..c0d3f89159e890
--- /dev/null
+++ b/clang/test/ParserOpenACC/parse-sub-array.cpp
@@ -0,0 +1,89 @@
+// RUN: %clang_cc1 %s -verify -fopenacc
+
+void Func(int i, int j) {
+ int array[5];
+#pragma acc parallel private(array[:])
+ while (true);
+#pragma acc parallel private(array[i:])
+ while (true);
+#pragma acc parallel private(array[:j])
+ while (true);
+#pragma acc parallel private(array[i:j])
+ while (true);
+#pragma acc parallel private(array[1:2])
+ while (true);
+
+ // expected-error@+1{{expected unqualified-id}}
+#pragma acc parallel private(array[::])
+ while (true);
+ // expected-error@+2{{expected ']'}}
+ // expected-note@+1{{to match this '['}}
+#pragma acc parallel private(array[1::])
+ while (true);
+ // expected-error@+2{{expected ']'}}
+ // expected-note@+1{{to match this '['}}
+#pragma acc parallel private(array[:2:])
+ while (true);
+ // expected-error@+3{{expected unqualified-id}}
+ // expected-error@+2{{expected ']'}}
+ // expected-note@+1{{to match this '['}}
+#pragma acc parallel private(array[::3])
+ while (true);
+ // expected-error@+2{{expected ']'}}
+ // expected-note@+1{{to match this '['}}
+#pragma acc parallel private(array[1:2:3])
+ while (true);
+}
+
+template<typename T, unsigned I, auto &IPtr>// #IPTR
+void TemplFunc() {
+ T array[I];
+ T array2[2*I];
+ T t; // #tDecl
+#pragma acc parallel private(array[:])
+ while (true);
+#pragma acc parallel private(array[t:])
+ while (true);
+#pragma acc parallel private(array[I-1:])
+ while (true);
+#pragma acc parallel private(array[IPtr:])
+ while (true);
+#pragma acc parallel private(array[:t])
+ while (true);
+#pragma acc parallel private(array[:I])
+ while (true);
+#pragma acc parallel private(array[:IPtr])
+ while (true);
+#pragma acc parallel private(array[t:t])
+ while (true);
+#pragma acc parallel private(array2[I:I])
+ while (true);
+#pragma acc parallel private(array[IPtr:IPtr])
+ while (true);
+
+ // expected-error@+1{{expected unqualified-id}}
+#pragma acc parallel private(array[::])
+ while (true);
+ // expected-error@+3{{'t' is not a class, namespace, or enumeration}}
+ // expected-note@#tDecl{{'t' declared here}}
+ // expected-error@+1{{expected unqualified-id}}
+#pragma acc parallel private(array[t::])
+ while (true);
+ // expected-error@+2{{expected ']'}}
+ // expected-note@+1{{to match this '['}}
+#pragma acc parallel private(array[:I:])
+ while (true);
+ // expected-error@+2{{no member named 'IPtr' in the global namespace}}
+ // expected-note@#IPTR{{'IPtr' declare...
[truncated]
|
alexey-bataev
changed the title
alexey-bataev
approved these changes
May 2, 2024
clang/lib/Sema/SemaOpenACC.cpp
Outdated
| if (LHS.isSigned() == RHS.isSigned()) | ||
| return LHS + RHS; | ||
|
|
||
| unsigned width = std::max(LHS.getBitWidth(), RHS.getBitWidth()) + 1; |
There was a problem hiding this comment.
Suggested change
| unsigned width = std::max(LHS.getBitWidth(), RHS.getBitWidth()) + 1; | |
| unsigned Width = std::max(LHS.getBitWidth(), RHS.getBitWidth()) + 1; |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This implementation takes quite a bit from the OMP implementation of array sections, but only has to enforce the rules as applicable to OpenACC. Additionally, it does its best to create an AST node (with the assistance of RecoveryExprs) with as much checking done as soon as possible in the case of instantiations.