aboutsummaryrefslogtreecommitdiff
path: root/clang/lib/Sema
diff options
context:
space:
mode:
Diffstat (limited to 'clang/lib/Sema')
-rw-r--r--clang/lib/Sema/DeclSpec.cpp2
-rw-r--r--clang/lib/Sema/HLSLBuiltinTypeDeclBuilder.cpp141
-rw-r--r--clang/lib/Sema/HLSLBuiltinTypeDeclBuilder.h7
-rw-r--r--clang/lib/Sema/HLSLExternalSemaSource.cpp3
-rw-r--r--clang/lib/Sema/ParsedAttr.cpp2
-rw-r--r--clang/lib/Sema/SemaAMDGPU.cpp103
-rw-r--r--clang/lib/Sema/SemaDecl.cpp2
-rw-r--r--clang/lib/Sema/SemaExpr.cpp22
-rw-r--r--clang/lib/Sema/SemaHLSL.cpp112
-rw-r--r--clang/lib/Sema/SemaOpenACC.cpp175
-rw-r--r--clang/lib/Sema/SemaOpenACCClause.cpp114
11 files changed, 578 insertions, 105 deletions
diff --git a/clang/lib/Sema/DeclSpec.cpp b/clang/lib/Sema/DeclSpec.cpp
index 8756ce5..184d31e 100644
--- a/clang/lib/Sema/DeclSpec.cpp
+++ b/clang/lib/Sema/DeclSpec.cpp
@@ -197,7 +197,7 @@ DeclaratorChunk DeclaratorChunk::getFunction(bool hasProto,
[&](DeclSpec::TQ TypeQual, StringRef PrintName, SourceLocation SL) {
I.Fun.MethodQualifiers->SetTypeQual(TypeQual, SL);
});
- I.Fun.MethodQualifiers->getAttributes().takeAllFrom(attrs);
+ I.Fun.MethodQualifiers->getAttributes().takeAllPrependingFrom(attrs);
I.Fun.MethodQualifiers->getAttributePool().takeAllFrom(attrs.getPool());
}
diff --git a/clang/lib/Sema/HLSLBuiltinTypeDeclBuilder.cpp b/clang/lib/Sema/HLSLBuiltinTypeDeclBuilder.cpp
index 3c20ccd..40c318a 100644
--- a/clang/lib/Sema/HLSLBuiltinTypeDeclBuilder.cpp
+++ b/clang/lib/Sema/HLSLBuiltinTypeDeclBuilder.cpp
@@ -144,6 +144,7 @@ private:
_2,
_3,
_4,
+ _5,
Handle = 128,
CounterHandle,
LastStmt
@@ -190,6 +191,9 @@ public:
template <typename T>
BuiltinTypeMethodBuilder &
accessCounterHandleFieldOnResource(T ResourceRecord);
+ template <typename ResourceT, typename ValueT>
+ BuiltinTypeMethodBuilder &
+ setCounterHandleFieldOnResource(ResourceT ResourceRecord, ValueT HandleValue);
template <typename T> BuiltinTypeMethodBuilder &returnValue(T ReturnValue);
BuiltinTypeMethodBuilder &returnThis();
BuiltinTypeDeclBuilder &finalize();
@@ -205,6 +209,11 @@ private:
if (!Method)
createDecl();
}
+
+ template <typename ResourceT, typename ValueT>
+ BuiltinTypeMethodBuilder &setFieldOnResource(ResourceT ResourceRecord,
+ ValueT HandleValue,
+ FieldDecl *HandleField);
};
TemplateParameterListBuilder::~TemplateParameterListBuilder() {
@@ -592,13 +601,27 @@ template <typename ResourceT, typename ValueT>
BuiltinTypeMethodBuilder &
BuiltinTypeMethodBuilder::setHandleFieldOnResource(ResourceT ResourceRecord,
ValueT HandleValue) {
+ return setFieldOnResource(ResourceRecord, HandleValue,
+ DeclBuilder.getResourceHandleField());
+}
+
+template <typename ResourceT, typename ValueT>
+BuiltinTypeMethodBuilder &
+BuiltinTypeMethodBuilder::setCounterHandleFieldOnResource(
+ ResourceT ResourceRecord, ValueT HandleValue) {
+ return setFieldOnResource(ResourceRecord, HandleValue,
+ DeclBuilder.getResourceCounterHandleField());
+}
+
+template <typename ResourceT, typename ValueT>
+BuiltinTypeMethodBuilder &BuiltinTypeMethodBuilder::setFieldOnResource(
+ ResourceT ResourceRecord, ValueT HandleValue, FieldDecl *HandleField) {
ensureCompleteDecl();
Expr *ResourceExpr = convertPlaceholder(ResourceRecord);
Expr *HandleValueExpr = convertPlaceholder(HandleValue);
ASTContext &AST = DeclBuilder.SemaRef.getASTContext();
- FieldDecl *HandleField = DeclBuilder.getResourceHandleField();
MemberExpr *HandleMemberExpr = MemberExpr::CreateImplicit(
AST, ResourceExpr, false, HandleField, HandleField->getType(), VK_LValue,
OK_Ordinary);
@@ -829,6 +852,18 @@ BuiltinTypeDeclBuilder &BuiltinTypeDeclBuilder::addDefaultHandleConstructor() {
.finalize();
}
+BuiltinTypeDeclBuilder &
+BuiltinTypeDeclBuilder::addStaticInitializationFunctions(bool HasCounter) {
+ if (HasCounter) {
+ addCreateFromBindingWithImplicitCounter();
+ addCreateFromImplicitBindingWithImplicitCounter();
+ } else {
+ addCreateFromBinding();
+ addCreateFromImplicitBinding();
+ }
+ return *this;
+}
+
// Adds static method that initializes resource from binding:
//
// static Resource<T> __createFromBinding(unsigned registerNo,
@@ -903,6 +938,102 @@ BuiltinTypeDeclBuilder &BuiltinTypeDeclBuilder::addCreateFromImplicitBinding() {
.finalize();
}
+// Adds static method that initializes resource from binding:
+//
+// static Resource<T>
+// __createFromBindingWithImplicitCounter(unsigned registerNo,
+// unsigned spaceNo, int range,
+// unsigned index, const char *name,
+// unsigned counterOrderId) {
+// Resource<T> tmp;
+// tmp.__handle = __builtin_hlsl_resource_handlefrombinding(
+// tmp.__handle, registerNo, spaceNo, range, index, name);
+// tmp.__counter_handle =
+// __builtin_hlsl_resource_counterhandlefromimplicitbinding(
+// tmp.__handle, counterOrderId, spaceNo);
+// return tmp;
+// }
+BuiltinTypeDeclBuilder &
+BuiltinTypeDeclBuilder::addCreateFromBindingWithImplicitCounter() {
+ assert(!Record->isCompleteDefinition() && "record is already complete");
+
+ using PH = BuiltinTypeMethodBuilder::PlaceHolder;
+ ASTContext &AST = SemaRef.getASTContext();
+ QualType HandleType = getResourceHandleField()->getType();
+ QualType RecordType = AST.getTypeDeclType(cast<TypeDecl>(Record));
+ BuiltinTypeMethodBuilder::LocalVar TmpVar("tmp", RecordType);
+
+ return BuiltinTypeMethodBuilder(*this,
+ "__createFromBindingWithImplicitCounter",
+ RecordType, false, false, SC_Static)
+ .addParam("registerNo", AST.UnsignedIntTy)
+ .addParam("spaceNo", AST.UnsignedIntTy)
+ .addParam("range", AST.IntTy)
+ .addParam("index", AST.UnsignedIntTy)
+ .addParam("name", AST.getPointerType(AST.CharTy.withConst()))
+ .addParam("counterOrderId", AST.UnsignedIntTy)
+ .declareLocalVar(TmpVar)
+ .accessHandleFieldOnResource(TmpVar)
+ .callBuiltin("__builtin_hlsl_resource_handlefrombinding", HandleType,
+ PH::LastStmt, PH::_0, PH::_1, PH::_2, PH::_3, PH::_4)
+ .setHandleFieldOnResource(TmpVar, PH::LastStmt)
+ .accessHandleFieldOnResource(TmpVar)
+ .callBuiltin("__builtin_hlsl_resource_counterhandlefromimplicitbinding",
+ HandleType, PH::LastStmt, PH::_5, PH::_1)
+ .setCounterHandleFieldOnResource(TmpVar, PH::LastStmt)
+ .returnValue(TmpVar)
+ .finalize();
+}
+
+// Adds static method that initializes resource from binding:
+//
+// static Resource<T>
+// __createFromImplicitBindingWithImplicitCounter(unsigned orderId,
+// unsigned spaceNo, int range,
+// unsigned index,
+// const char *name,
+// unsigned counterOrderId) {
+// Resource<T> tmp;
+// tmp.__handle = __builtin_hlsl_resource_handlefromimplicitbinding(
+// tmp.__handle, orderId, spaceNo, range, index, name);
+// tmp.__counter_handle =
+// __builtin_hlsl_resource_counterhandlefromimplicitbinding(
+// tmp.__handle, counterOrderId, spaceNo);
+// return tmp;
+// }
+BuiltinTypeDeclBuilder &
+BuiltinTypeDeclBuilder::addCreateFromImplicitBindingWithImplicitCounter() {
+ assert(!Record->isCompleteDefinition() && "record is already complete");
+
+ using PH = BuiltinTypeMethodBuilder::PlaceHolder;
+ ASTContext &AST = SemaRef.getASTContext();
+ QualType HandleType = getResourceHandleField()->getType();
+ QualType RecordType = AST.getTypeDeclType(cast<TypeDecl>(Record));
+ BuiltinTypeMethodBuilder::LocalVar TmpVar("tmp", RecordType);
+
+ return BuiltinTypeMethodBuilder(
+ *this, "__createFromImplicitBindingWithImplicitCounter",
+ RecordType, false, false, SC_Static)
+ .addParam("orderId", AST.UnsignedIntTy)
+ .addParam("spaceNo", AST.UnsignedIntTy)
+ .addParam("range", AST.IntTy)
+ .addParam("index", AST.UnsignedIntTy)
+ .addParam("name", AST.getPointerType(AST.CharTy.withConst()))
+ .addParam("counterOrderId", AST.UnsignedIntTy)
+ .declareLocalVar(TmpVar)
+ .accessHandleFieldOnResource(TmpVar)
+ .callBuiltin("__builtin_hlsl_resource_handlefromimplicitbinding",
+ HandleType, PH::LastStmt, PH::_0, PH::_1, PH::_2, PH::_3,
+ PH::_4)
+ .setHandleFieldOnResource(TmpVar, PH::LastStmt)
+ .accessHandleFieldOnResource(TmpVar)
+ .callBuiltin("__builtin_hlsl_resource_counterhandlefromimplicitbinding",
+ HandleType, PH::LastStmt, PH::_5, PH::_1)
+ .setCounterHandleFieldOnResource(TmpVar, PH::LastStmt)
+ .returnValue(TmpVar)
+ .finalize();
+}
+
BuiltinTypeDeclBuilder &BuiltinTypeDeclBuilder::addCopyConstructor() {
assert(!Record->isCompleteDefinition() && "record is already complete");
@@ -1048,7 +1179,7 @@ BuiltinTypeDeclBuilder &BuiltinTypeDeclBuilder::addIncrementCounterMethod() {
return BuiltinTypeMethodBuilder(*this, "IncrementCounter",
SemaRef.getASTContext().UnsignedIntTy)
.callBuiltin("__builtin_hlsl_buffer_update_counter", QualType(),
- PH::Handle, getConstantIntExpr(1))
+ PH::CounterHandle, getConstantIntExpr(1))
.finalize();
}
@@ -1057,7 +1188,7 @@ BuiltinTypeDeclBuilder &BuiltinTypeDeclBuilder::addDecrementCounterMethod() {
return BuiltinTypeMethodBuilder(*this, "DecrementCounter",
SemaRef.getASTContext().UnsignedIntTy)
.callBuiltin("__builtin_hlsl_buffer_update_counter", QualType(),
- PH::Handle, getConstantIntExpr(-1))
+ PH::CounterHandle, getConstantIntExpr(-1))
.finalize();
}
@@ -1102,7 +1233,7 @@ BuiltinTypeDeclBuilder &BuiltinTypeDeclBuilder::addAppendMethod() {
return BuiltinTypeMethodBuilder(*this, "Append", AST.VoidTy)
.addParam("value", ElemTy)
.callBuiltin("__builtin_hlsl_buffer_update_counter", AST.UnsignedIntTy,
- PH::Handle, getConstantIntExpr(1))
+ PH::CounterHandle, getConstantIntExpr(1))
.callBuiltin("__builtin_hlsl_resource_getpointer",
AST.getPointerType(AddrSpaceElemTy), PH::Handle,
PH::LastStmt)
@@ -1119,7 +1250,7 @@ BuiltinTypeDeclBuilder &BuiltinTypeDeclBuilder::addConsumeMethod() {
AST.getAddrSpaceQualType(ElemTy, LangAS::hlsl_device);
return BuiltinTypeMethodBuilder(*this, "Consume", ElemTy)
.callBuiltin("__builtin_hlsl_buffer_update_counter", AST.UnsignedIntTy,
- PH::Handle, getConstantIntExpr(-1))
+ PH::CounterHandle, getConstantIntExpr(-1))
.callBuiltin("__builtin_hlsl_resource_getpointer",
AST.getPointerType(AddrSpaceElemTy), PH::Handle,
PH::LastStmt)
diff --git a/clang/lib/Sema/HLSLBuiltinTypeDeclBuilder.h b/clang/lib/Sema/HLSLBuiltinTypeDeclBuilder.h
index a981602..86cbd10 100644
--- a/clang/lib/Sema/HLSLBuiltinTypeDeclBuilder.h
+++ b/clang/lib/Sema/HLSLBuiltinTypeDeclBuilder.h
@@ -83,8 +83,7 @@ public:
BuiltinTypeDeclBuilder &addCopyAssignmentOperator();
// Static create methods
- BuiltinTypeDeclBuilder &addCreateFromBinding();
- BuiltinTypeDeclBuilder &addCreateFromImplicitBinding();
+ BuiltinTypeDeclBuilder &addStaticInitializationFunctions(bool HasCounter);
// Builtin types methods
BuiltinTypeDeclBuilder &addLoadMethods();
@@ -96,6 +95,10 @@ public:
BuiltinTypeDeclBuilder &addConsumeMethod();
private:
+ BuiltinTypeDeclBuilder &addCreateFromBinding();
+ BuiltinTypeDeclBuilder &addCreateFromImplicitBinding();
+ BuiltinTypeDeclBuilder &addCreateFromBindingWithImplicitCounter();
+ BuiltinTypeDeclBuilder &addCreateFromImplicitBindingWithImplicitCounter();
BuiltinTypeDeclBuilder &addResourceMember(StringRef MemberName,
ResourceClass RC, bool IsROV,
bool RawBuffer, bool IsCounter,
diff --git a/clang/lib/Sema/HLSLExternalSemaSource.cpp b/clang/lib/Sema/HLSLExternalSemaSource.cpp
index cc43e94..e118dda 100644
--- a/clang/lib/Sema/HLSLExternalSemaSource.cpp
+++ b/clang/lib/Sema/HLSLExternalSemaSource.cpp
@@ -236,8 +236,7 @@ static BuiltinTypeDeclBuilder setupBufferType(CXXRecordDecl *Decl, Sema &S,
.addDefaultHandleConstructor()
.addCopyConstructor()
.addCopyAssignmentOperator()
- .addCreateFromBinding()
- .addCreateFromImplicitBinding();
+ .addStaticInitializationFunctions(HasCounter);
}
// This function is responsible for constructing the constraint expression for
diff --git a/clang/lib/Sema/ParsedAttr.cpp b/clang/lib/Sema/ParsedAttr.cpp
index 294f88e..2b5ad33 100644
--- a/clang/lib/Sema/ParsedAttr.cpp
+++ b/clang/lib/Sema/ParsedAttr.cpp
@@ -304,7 +304,7 @@ bool ParsedAttr::checkAtMostNumArgs(Sema &S, unsigned Num) const {
void clang::takeAndConcatenateAttrs(ParsedAttributes &First,
ParsedAttributes &&Second) {
- First.takeAllAtEndFrom(Second);
+ First.takeAllAppendingFrom(Second);
if (!First.Range.getBegin().isValid())
First.Range.setBegin(Second.Range.getBegin());
diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp
index 3a0c231..e32f437 100644
--- a/clang/lib/Sema/SemaAMDGPU.cpp
+++ b/clang/lib/Sema/SemaAMDGPU.cpp
@@ -11,6 +11,7 @@
//===----------------------------------------------------------------------===//
#include "clang/Sema/SemaAMDGPU.h"
+#include "clang/Basic/DiagnosticFrontend.h"
#include "clang/Basic/DiagnosticSema.h"
#include "clang/Basic/TargetBuiltins.h"
#include "clang/Sema/Ownership.h"
@@ -111,6 +112,108 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_16x8B:
case AMDGPU::BI__builtin_amdgcn_cooperative_atomic_store_8x16B:
return checkCoopAtomicFunctionCall(TheCall, /*IsStore=*/true);
+ case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_1d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_1darray_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_2d_f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_2d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_2darray_f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_2darray_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_3d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_cube_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_mip_1d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_mip_1darray_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_mip_2d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_mip_2darray_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_mip_3d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_load_mip_cube_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_1d_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_1darray_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_2d_f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_2d_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_2darray_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_3d_v4f16_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f32_f32:
+ case AMDGPU::BI__builtin_amdgcn_image_sample_cube_v4f16_f32: {
+ StringRef FeatureList(
+ getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
+ if (!Builtin::evaluateRequiredTargetFeatures(FeatureList,
+ CallerFeatureMap)) {
+ Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)
+ << FD->getDeclName() << FeatureList;
+ return false;
+ }
+
+ unsigned ArgCount = TheCall->getNumArgs() - 1;
+ llvm::APSInt Result;
+
+ return (SemaRef.BuiltinConstantArg(TheCall, 0, Result)) ||
+ (SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) ||
+ (SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result));
+ }
+ case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_1d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_1darray_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_2d_f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_2d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_2darray_f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_2darray_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_3d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_cube_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_1d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_1darray_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_2d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_2darray_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_3d_v4f16_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f32_i32:
+ case AMDGPU::BI__builtin_amdgcn_image_store_mip_cube_v4f16_i32: {
+ StringRef FeatureList(
+ getASTContext().BuiltinInfo.getRequiredFeatures(BuiltinID));
+ if (!Builtin::evaluateRequiredTargetFeatures(FeatureList,
+ CallerFeatureMap)) {
+ Diag(TheCall->getBeginLoc(), diag::err_builtin_needs_feature)
+ << FD->getDeclName() << FeatureList;
+ return false;
+ }
+
+ unsigned ArgCount = TheCall->getNumArgs() - 1;
+ llvm::APSInt Result;
+
+ return (SemaRef.BuiltinConstantArg(TheCall, 1, Result)) ||
+ (SemaRef.BuiltinConstantArg(TheCall, ArgCount, Result)) ||
+ (SemaRef.BuiltinConstantArg(TheCall, (ArgCount - 1), Result));
+ }
default:
return false;
}
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 6eaf7b9..0e83c20 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -14637,7 +14637,7 @@ StmtResult Sema::ActOnCXXForRangeIdentifier(Scope *S, SourceLocation IdentLoc,
Declarator D(DS, ParsedAttributesView::none(), DeclaratorContext::ForInit);
D.SetIdentifier(Ident, IdentLoc);
- D.takeAttributes(Attrs);
+ D.takeAttributesAppending(Attrs);
D.AddTypeInfo(DeclaratorChunk::getReference(0, IdentLoc, /*lvalue*/ false),
IdentLoc);
diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index 4230ea7..01abc1f 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -166,6 +166,11 @@ static void diagnoseUseOfInternalDeclInInlineFunction(Sema &S,
// This is disabled under C++; there are too many ways for this to fire in
// contexts where the warning is a false positive, or where it is technically
// correct but benign.
+ //
+ // WG14 N3622 which removed the constraint entirely in C2y. It is left
+ // enabled in earlier language modes because this is a constraint in those
+ // language modes. But in C2y mode, we still want to issue the "incompatible
+ // with previous standards" diagnostic, too.
if (S.getLangOpts().CPlusPlus)
return;
@@ -190,16 +195,17 @@ static void diagnoseUseOfInternalDeclInInlineFunction(Sema &S,
// This last can give us false negatives, but it's better than warning on
// wrappers for simple C library functions.
const FunctionDecl *UsedFn = dyn_cast<FunctionDecl>(D);
- bool DowngradeWarning = S.getSourceManager().isInMainFile(Loc);
- if (!DowngradeWarning && UsedFn)
- DowngradeWarning = UsedFn->isInlined() || UsedFn->hasAttr<ConstAttr>();
-
- S.Diag(Loc, DowngradeWarning ? diag::ext_internal_in_extern_inline_quiet
- : diag::ext_internal_in_extern_inline)
- << /*IsVar=*/!UsedFn << D;
+ unsigned DiagID;
+ if (S.getLangOpts().C2y)
+ DiagID = diag::warn_c2y_compat_internal_in_extern_inline;
+ else if ((UsedFn && (UsedFn->isInlined() || UsedFn->hasAttr<ConstAttr>())) ||
+ S.getSourceManager().isInMainFile(Loc))
+ DiagID = diag::ext_internal_in_extern_inline_quiet;
+ else
+ DiagID = diag::ext_internal_in_extern_inline;
+ S.Diag(Loc, DiagID) << /*IsVar=*/!UsedFn << D;
S.MaybeSuggestAddingStaticToDecl(Current);
-
S.Diag(D->getCanonicalDecl()->getLocation(), diag::note_entity_declared_at)
<< D;
}
diff --git a/clang/lib/Sema/SemaHLSL.cpp b/clang/lib/Sema/SemaHLSL.cpp
index 09e5d69..17cb1e4 100644
--- a/clang/lib/Sema/SemaHLSL.cpp
+++ b/clang/lib/Sema/SemaHLSL.cpp
@@ -1240,6 +1240,20 @@ static CXXMethodDecl *lookupMethod(Sema &S, CXXRecordDecl *RecordDecl,
} // end anonymous namespace
+static bool hasCounterHandle(const CXXRecordDecl *RD) {
+ if (RD->field_empty())
+ return false;
+ auto It = std::next(RD->field_begin());
+ if (It == RD->field_end())
+ return false;
+ const FieldDecl *SecondField = *It;
+ if (const auto *ResTy =
+ SecondField->getType()->getAs<HLSLAttributedResourceType>()) {
+ return ResTy->getAttrs().IsCounter;
+ }
+ return false;
+}
+
bool SemaHLSL::handleRootSignatureElements(
ArrayRef<hlsl::RootSignatureElement> Elements) {
// Define some common error handling functions
@@ -2973,6 +2987,25 @@ bool SemaHLSL::CheckBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall) {
TheCall->setType(ResourceTy);
break;
}
+ case Builtin::BI__builtin_hlsl_resource_counterhandlefromimplicitbinding: {
+ ASTContext &AST = SemaRef.getASTContext();
+ if (SemaRef.checkArgCount(TheCall, 3) ||
+ CheckResourceHandle(&SemaRef, TheCall, 0) ||
+ CheckArgTypeMatches(&SemaRef, TheCall->getArg(1), AST.UnsignedIntTy) ||
+ CheckArgTypeMatches(&SemaRef, TheCall->getArg(2), AST.UnsignedIntTy))
+ return true;
+
+ QualType MainHandleTy = TheCall->getArg(0)->getType();
+ auto *MainResType = MainHandleTy->getAs<HLSLAttributedResourceType>();
+ auto MainAttrs = MainResType->getAttrs();
+ assert(!MainAttrs.IsCounter && "cannot create a counter from a counter");
+ MainAttrs.IsCounter = true;
+ QualType CounterHandleTy = AST.getHLSLAttributedResourceType(
+ MainResType->getWrappedType(), MainResType->getContainedType(),
+ MainAttrs);
+ TheCall->setType(CounterHandleTy);
+ break;
+ }
case Builtin::BI__builtin_hlsl_and:
case Builtin::BI__builtin_hlsl_or: {
if (SemaRef.checkArgCount(TheCall, 2))
@@ -3780,10 +3813,24 @@ void SemaHLSL::ActOnVariableDeclarator(VarDecl *VD) {
uint32_t OrderID = getNextImplicitBindingOrderID();
if (Binding.hasBinding())
Binding.setImplicitOrderID(OrderID);
- else
+ else {
addImplicitBindingAttrToDecl(
SemaRef, VD, getRegisterType(getResourceArrayHandleType(VD)),
OrderID);
+ // Re-create the binding object to pick up the new attribute.
+ Binding = ResourceBindingAttrs(VD);
+ }
+ }
+
+ // Get to the base type of a potentially multi-dimensional array.
+ QualType Ty = getASTContext().getBaseElementType(VD->getType());
+
+ const CXXRecordDecl *RD = Ty->getAsCXXRecordDecl();
+ if (hasCounterHandle(RD)) {
+ if (!Binding.hasCounterImplicitOrderID()) {
+ uint32_t OrderID = getNextImplicitBindingOrderID();
+ Binding.setCounterImplicitOrderID(OrderID);
+ }
}
}
}
@@ -3808,19 +3855,31 @@ bool SemaHLSL::initGlobalResourceDecl(VarDecl *VD) {
CXXMethodDecl *CreateMethod = nullptr;
llvm::SmallVector<Expr *> Args;
+ bool HasCounter = hasCounterHandle(ResourceDecl);
+ const char *CreateMethodName;
+ if (Binding.isExplicit())
+ CreateMethodName = HasCounter ? "__createFromBindingWithImplicitCounter"
+ : "__createFromBinding";
+ else
+ CreateMethodName = HasCounter
+ ? "__createFromImplicitBindingWithImplicitCounter"
+ : "__createFromImplicitBinding";
+
+ CreateMethod =
+ lookupMethod(SemaRef, ResourceDecl, CreateMethodName, VD->getLocation());
+
+ if (!CreateMethod)
+ // This can happen if someone creates a struct that looks like an HLSL
+ // resource record but does not have the required static create method.
+ // No binding will be generated for it.
+ return false;
+
if (Binding.isExplicit()) {
- // The resource has explicit binding.
- CreateMethod = lookupMethod(SemaRef, ResourceDecl, "__createFromBinding",
- VD->getLocation());
IntegerLiteral *RegSlot =
IntegerLiteral::Create(AST, llvm::APInt(UIntTySize, Binding.getSlot()),
AST.UnsignedIntTy, SourceLocation());
Args.push_back(RegSlot);
} else {
- // The resource has implicit binding.
- CreateMethod =
- lookupMethod(SemaRef, ResourceDecl, "__createFromImplicitBinding",
- VD->getLocation());
uint32_t OrderID = (Binding.hasImplicitOrderID())
? Binding.getImplicitOrderID()
: getNextImplicitBindingOrderID();
@@ -3830,12 +3889,6 @@ bool SemaHLSL::initGlobalResourceDecl(VarDecl *VD) {
Args.push_back(OrderId);
}
- if (!CreateMethod)
- // This can happen if someone creates a struct that looks like an HLSL
- // resource record but does not have the required static create method.
- // No binding will be generated for it.
- return false;
-
IntegerLiteral *Space =
IntegerLiteral::Create(AST, llvm::APInt(UIntTySize, Binding.getSpace()),
AST.UnsignedIntTy, SourceLocation());
@@ -3859,6 +3912,15 @@ bool SemaHLSL::initGlobalResourceDecl(VarDecl *VD) {
Name, nullptr, VK_PRValue, FPOptionsOverride());
Args.push_back(NameCast);
+ if (HasCounter) {
+ // Will this be in the correct order?
+ uint32_t CounterOrderID = getNextImplicitBindingOrderID();
+ IntegerLiteral *CounterId =
+ IntegerLiteral::Create(AST, llvm::APInt(UIntTySize, CounterOrderID),
+ AST.UnsignedIntTy, SourceLocation());
+ Args.push_back(CounterId);
+ }
+
// Make sure the create method template is instantiated and emitted.
if (!CreateMethod->isDefined() && CreateMethod->isTemplateInstantiation())
SemaRef.InstantiateFunctionDefinition(VD->getLocation(), CreateMethod,
@@ -3899,20 +3961,24 @@ bool SemaHLSL::initGlobalResourceArrayDecl(VarDecl *VD) {
ASTContext &AST = SemaRef.getASTContext();
QualType ResElementTy = AST.getBaseElementType(VD->getType());
CXXRecordDecl *ResourceDecl = ResElementTy->getAsCXXRecordDecl();
-
- HLSLResourceBindingAttr *RBA = VD->getAttr<HLSLResourceBindingAttr>();
- HLSLVkBindingAttr *VkBinding = VD->getAttr<HLSLVkBindingAttr>();
CXXMethodDecl *CreateMethod = nullptr;
- if (VkBinding || (RBA && RBA->hasRegisterSlot()))
+ bool HasCounter = hasCounterHandle(ResourceDecl);
+ ResourceBindingAttrs ResourceAttrs(VD);
+ if (ResourceAttrs.isExplicit())
// Resource has explicit binding.
- CreateMethod = lookupMethod(SemaRef, ResourceDecl, "__createFromBinding",
- VD->getLocation());
- else
- // Resource has implicit binding.
CreateMethod =
- lookupMethod(SemaRef, ResourceDecl, "__createFromImplicitBinding",
+ lookupMethod(SemaRef, ResourceDecl,
+ HasCounter ? "__createFromBindingWithImplicitCounter"
+ : "__createFromBinding",
VD->getLocation());
+ else
+ // Resource has implicit binding.
+ CreateMethod = lookupMethod(
+ SemaRef, ResourceDecl,
+ HasCounter ? "__createFromImplicitBindingWithImplicitCounter"
+ : "__createFromImplicitBinding",
+ VD->getLocation());
if (!CreateMethod)
return false;
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 8471f02..779b6e9 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -2759,7 +2759,7 @@ OpenACCPrivateRecipe SemaOpenACC::CreatePrivateInitRecipe(const Expr *VarExpr) {
// Array sections are special, and we have to treat them that way.
if (const auto *ASE =
dyn_cast<ArraySectionExpr>(VarExpr->IgnoreParenImpCasts()))
- VarTy = ArraySectionExpr::getBaseOriginalType(ASE);
+ VarTy = ASE->getElementType();
VarDecl *AllocaDecl = CreateAllocaDecl(
getASTContext(), SemaRef.getCurContext(), VarExpr->getBeginLoc(),
@@ -2795,7 +2795,7 @@ SemaOpenACC::CreateFirstPrivateInitRecipe(const Expr *VarExpr) {
// Array sections are special, and we have to treat them that way.
if (const auto *ASE =
dyn_cast<ArraySectionExpr>(VarExpr->IgnoreParenImpCasts()))
- VarTy = ArraySectionExpr::getBaseOriginalType(ASE);
+ VarTy = ASE->getElementType();
VarDecl *AllocaDecl = CreateAllocaDecl(
getASTContext(), SemaRef.getCurContext(), VarExpr->getBeginLoc(),
@@ -2896,7 +2896,16 @@ OpenACCReductionRecipe SemaOpenACC::CreateReductionInitRecipe(
// Array sections are special, and we have to treat them that way.
if (const auto *ASE =
dyn_cast<ArraySectionExpr>(VarExpr->IgnoreParenImpCasts()))
- VarTy = ArraySectionExpr::getBaseOriginalType(ASE);
+ VarTy = ASE->getElementType();
+
+ llvm::SmallVector<OpenACCReductionRecipe::CombinerRecipe, 1> CombinerRecipes;
+
+ // We use the 'set-ness' of the alloca-decl to determine whether the combiner
+ // is 'set' or not, so we can skip any attempts at it if we're going to fail
+ // at any of the combiners.
+ if (CreateReductionCombinerRecipe(VarExpr->getBeginLoc(), ReductionOperator,
+ VarTy, CombinerRecipes))
+ return OpenACCReductionRecipe::Empty();
VarDecl *AllocaDecl = CreateAllocaDecl(
getASTContext(), SemaRef.getCurContext(), VarExpr->getBeginLoc(),
@@ -2946,5 +2955,163 @@ OpenACCReductionRecipe SemaOpenACC::CreateReductionInitRecipe(
AllocaDecl->setInit(Init.get());
AllocaDecl->setInitStyle(VarDecl::CallInit);
}
- return OpenACCReductionRecipe(AllocaDecl);
+
+ return OpenACCReductionRecipe(AllocaDecl, CombinerRecipes);
+}
+
+bool SemaOpenACC::CreateReductionCombinerRecipe(
+ SourceLocation Loc, OpenACCReductionOperator ReductionOperator,
+ QualType VarTy,
+ llvm::SmallVectorImpl<OpenACCReductionRecipe::CombinerRecipe>
+ &CombinerRecipes) {
+ // Now we can try to generate the 'combiner' recipe. This is a little
+ // complicated in that if the 'VarTy' is an array type, we want to take its
+ // element type so we can generate that. Additionally, if this is a struct,
+ // we have two options: If there is overloaded operators, we want to take
+ // THOSE, else we want to do the individual elements.
+
+ BinaryOperatorKind BinOp;
+ switch (ReductionOperator) {
+ case OpenACCReductionOperator::Invalid:
+ // This can only happen when there is an error, and since these inits
+ // are used for code generation, we can just ignore/not bother doing any
+ // initialization here.
+ CombinerRecipes.push_back({nullptr, nullptr, nullptr});
+ return false;
+ case OpenACCReductionOperator::Addition:
+ BinOp = BinaryOperatorKind::BO_AddAssign;
+ break;
+ case OpenACCReductionOperator::Multiplication:
+ BinOp = BinaryOperatorKind::BO_MulAssign;
+ break;
+ case OpenACCReductionOperator::BitwiseAnd:
+ BinOp = BinaryOperatorKind::BO_AndAssign;
+ break;
+ case OpenACCReductionOperator::BitwiseOr:
+ BinOp = BinaryOperatorKind::BO_OrAssign;
+ break;
+ case OpenACCReductionOperator::BitwiseXOr:
+ BinOp = BinaryOperatorKind::BO_XorAssign;
+ break;
+
+ case OpenACCReductionOperator::Max:
+ case OpenACCReductionOperator::Min:
+ case OpenACCReductionOperator::And:
+ case OpenACCReductionOperator::Or:
+ // We just want a 'NYI' error in the backend, so leave an empty combiner
+ // recipe, and claim success.
+ CombinerRecipes.push_back({nullptr, nullptr, nullptr});
+ return false;
+ }
+
+ // If VarTy is an array type, at the top level only, we want to do our
+ // compares/decomp/etc at the element level.
+ if (auto *AT = getASTContext().getAsArrayType(VarTy))
+ VarTy = AT->getElementType();
+
+ assert(!VarTy->isArrayType() && "Only 1 level of array allowed");
+
+ auto tryCombiner = [&, this](DeclRefExpr *LHSDRE, DeclRefExpr *RHSDRE,
+ bool IncludeTrap) {
+ // TODO: OpenACC: we have to figure out based on the bin-op how to do the
+ // ones that we can't just use compound operators for. So &&, ||, max, and
+ // min aren't really clear what we could do here.
+ if (IncludeTrap) {
+ // Trap all of the errors here, we'll emit our own at the end.
+ Sema::TentativeAnalysisScope Trap{SemaRef};
+
+ return SemaRef.BuildBinOp(SemaRef.getCurScope(), Loc, BinOp, LHSDRE,
+ RHSDRE,
+ /*ForFoldExpr=*/false);
+ } else {
+ return SemaRef.BuildBinOp(SemaRef.getCurScope(), Loc, BinOp, LHSDRE,
+ RHSDRE,
+ /*ForFoldExpr=*/false);
+ }
+ };
+
+ struct CombinerAttemptTy {
+ VarDecl *LHS;
+ DeclRefExpr *LHSDRE;
+ VarDecl *RHS;
+ DeclRefExpr *RHSDRE;
+ Expr *Op;
+ };
+
+ auto formCombiner = [&, this](QualType Ty) -> CombinerAttemptTy {
+ VarDecl *LHSDecl = CreateAllocaDecl(
+ getASTContext(), SemaRef.getCurContext(), Loc,
+ &getASTContext().Idents.get("openacc.reduction.combiner.lhs"), Ty);
+ auto *LHSDRE = DeclRefExpr::Create(
+ getASTContext(), NestedNameSpecifierLoc{}, SourceLocation{}, LHSDecl,
+ /*ReferstoEnclosingVariableOrCapture=*/false,
+ DeclarationNameInfo{DeclarationName{LHSDecl->getDeclName()},
+ LHSDecl->getBeginLoc()},
+ Ty, clang::VK_LValue, LHSDecl, nullptr, NOUR_None);
+ VarDecl *RHSDecl = CreateAllocaDecl(
+ getASTContext(), SemaRef.getCurContext(), Loc,
+ &getASTContext().Idents.get("openacc.reduction.combiner.lhs"), Ty);
+ auto *RHSDRE = DeclRefExpr::Create(
+ getASTContext(), NestedNameSpecifierLoc{}, SourceLocation{}, RHSDecl,
+ /*ReferstoEnclosingVariableOrCapture=*/false,
+ DeclarationNameInfo{DeclarationName{RHSDecl->getDeclName()},
+ RHSDecl->getBeginLoc()},
+ Ty, clang::VK_LValue, RHSDecl, nullptr, NOUR_None);
+
+ ExprResult BinOpResult = tryCombiner(LHSDRE, RHSDRE, /*IncludeTrap=*/true);
+
+ return {LHSDecl, LHSDRE, RHSDecl, RHSDRE, BinOpResult.get()};
+ };
+
+ CombinerAttemptTy TopLevelCombinerInfo = formCombiner(VarTy);
+
+ if (TopLevelCombinerInfo.Op) {
+ if (!TopLevelCombinerInfo.Op->containsErrors() &&
+ TopLevelCombinerInfo.Op->isInstantiationDependent()) {
+ // If this is instantiation dependent, we're just going to 'give up' here
+ // and count on us to get it right during instantaition.
+ CombinerRecipes.push_back({nullptr, nullptr, nullptr});
+ return false;
+ } else if (!TopLevelCombinerInfo.Op->containsErrors()) {
+ // Else, we succeeded, we can just return this combiner.
+ CombinerRecipes.push_back({TopLevelCombinerInfo.LHS,
+ TopLevelCombinerInfo.RHS,
+ TopLevelCombinerInfo.Op});
+ return false;
+ }
+ }
+
+ // Since the 'root' level didn't fail, the only thing that could be successful
+ // is a struct that we decompose on its individual fields.
+
+ RecordDecl *RD = VarTy->getAsRecordDecl();
+ if (!RD) {
+ Diag(Loc, diag::err_acc_reduction_recipe_no_op) << VarTy;
+ tryCombiner(TopLevelCombinerInfo.LHSDRE, TopLevelCombinerInfo.RHSDRE,
+ /*IncludeTrap=*/false);
+ return true;
+ }
+
+ for (const FieldDecl *FD : RD->fields()) {
+ CombinerAttemptTy FieldCombinerInfo = formCombiner(FD->getType());
+
+ if (!FieldCombinerInfo.Op || FieldCombinerInfo.Op->containsErrors()) {
+ Diag(Loc, diag::err_acc_reduction_recipe_no_op) << FD->getType();
+ Diag(FD->getBeginLoc(), diag::note_acc_reduction_recipe_noop_field) << RD;
+ tryCombiner(FieldCombinerInfo.LHSDRE, FieldCombinerInfo.RHSDRE,
+ /*IncludeTrap=*/false);
+ return true;
+ }
+
+ if (FieldCombinerInfo.Op->isInstantiationDependent()) {
+ // If this is instantiation dependent, we're just going to 'give up' here
+ // and count on us to get it right during instantaition.
+ CombinerRecipes.push_back({nullptr, nullptr, nullptr});
+ } else {
+ CombinerRecipes.push_back(
+ {FieldCombinerInfo.LHS, FieldCombinerInfo.RHS, FieldCombinerInfo.Op});
+ }
+ }
+
+ return false;
}
diff --git a/clang/lib/Sema/SemaOpenACCClause.cpp b/clang/lib/Sema/SemaOpenACCClause.cpp
index b086929..881e960 100644
--- a/clang/lib/Sema/SemaOpenACCClause.cpp
+++ b/clang/lib/Sema/SemaOpenACCClause.cpp
@@ -1915,51 +1915,34 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
return Result;
}
-/// OpenACC 3.3 section 2.5.15:
-/// At a mininmum, the supported data types include ... the numerical data types
-/// in C, C++, and Fortran.
-///
-/// If the reduction var is a composite variable, each
-/// member of the composite variable must be a supported datatype for the
-/// reduction operation.
-ExprResult SemaOpenACC::CheckReductionVar(OpenACCDirectiveKind DirectiveKind,
- OpenACCReductionOperator ReductionOp,
- Expr *VarExpr) {
- // For now, we only support 'scalar' types, or composites/arrays of scalar
- // types.
- VarExpr = VarExpr->IgnoreParenCasts();
+bool SemaOpenACC::CheckReductionVarType(Expr *VarExpr) {
SourceLocation VarLoc = VarExpr->getBeginLoc();
SmallVector<PartialDiagnosticAt> Notes;
- QualType CurType = VarExpr->getType();
-
- // For array like things, the expression can either be an array element
- // (subscript expr), array section, or array type. Peel those off, and add
- // notes in case we find an illegal kind. We'll allow scalar or composite of
- // scalars inside of this.
- if (auto *ASE = dyn_cast<ArraySectionExpr>(VarExpr)) {
- QualType BaseType = ArraySectionExpr::getBaseOriginalType(ASE);
+ // The standard isn't clear how many levels of 'array element' or 'subarray'
+ // are permitted, but we can handle as many as we need, so we'll strip them
+ // off here. This will result in CurType being the actual 'type' of the
+ // expression, which is what we are looking to check.
+ QualType CurType = isa<ArraySectionExpr>(VarExpr)
+ ? ArraySectionExpr::getBaseOriginalType(VarExpr)
+ : VarExpr->getType();
+
+ // This can happen when we have a dependent type in an array element that the
+ // above function has tried to 'unwrap'. Since this can only happen with
+ // dependence, just let it go.
+ if (CurType.isNull())
+ return false;
- PartialDiagnostic PD = PDiag(diag::note_acc_reduction_array)
- << diag::OACCReductionArray::Section << BaseType;
- Notes.push_back({ASE->getBeginLoc(), PD});
-
- CurType = getASTContext().getBaseElementType(BaseType);
- } else if (auto *SubExpr = dyn_cast<ArraySubscriptExpr>(VarExpr)) {
- // Array subscript already results in the type of the thing as its type, so
- // there is no type to change here.
- PartialDiagnostic PD =
- PDiag(diag::note_acc_reduction_array)
- << diag::OACCReductionArray::Subscript
- << SubExpr->getBase()->IgnoreParenImpCasts()->getType();
- Notes.push_back({SubExpr->getBeginLoc(), PD});
- } else if (auto *AT = getASTContext().getAsArrayType(CurType)) {
+ // If we are still an array type, we allow 1 level of 'unpeeling' of the
+ // array. The standard isn't clear here whether this is allowed, but
+ // array-of-valid-things makes sense.
+ if (auto *AT = getASTContext().getAsArrayType(CurType)) {
// If we're already the array type, peel off the array and leave the element
// type.
- CurType = getASTContext().getBaseElementType(AT);
PartialDiagnostic PD = PDiag(diag::note_acc_reduction_array)
<< diag::OACCReductionArray::ArrayTy << CurType;
Notes.push_back({VarLoc, PD});
+ CurType = AT->getElementType();
}
auto IsValidMemberOfComposite = [](QualType Ty) {
@@ -1974,31 +1957,26 @@ ExprResult SemaOpenACC::CheckReductionVar(OpenACCDirectiveKind DirectiveKind,
for (auto [Loc, PD] : Notes)
Diag(Loc, PD);
- Diag(VarLoc, diag::note_acc_reduction_type_summary);
+ return Diag(VarLoc, diag::note_acc_reduction_type_summary);
};
// If the type is already scalar, or is dependent, just give up.
if (IsValidMemberOfComposite(CurType)) {
// Nothing to do here, is valid.
} else if (auto *RD = CurType->getAsRecordDecl()) {
- if (!RD->isStruct() && !RD->isClass()) {
- EmitDiags(VarLoc, PDiag(diag::err_acc_reduction_type)
- << RD << diag::OACCReductionTy::NotClassStruct);
- return ExprError();
- }
+ if (!RD->isStruct() && !RD->isClass())
+ return EmitDiags(VarLoc, PDiag(diag::err_acc_reduction_type)
+ << RD
+ << diag::OACCReductionTy::NotClassStruct);
- if (!RD->isCompleteDefinition()) {
- EmitDiags(VarLoc, PDiag(diag::err_acc_reduction_type)
- << RD << diag::OACCReductionTy::NotComplete);
- return ExprError();
- }
+ if (!RD->isCompleteDefinition())
+ return EmitDiags(VarLoc, PDiag(diag::err_acc_reduction_type)
+ << RD << diag::OACCReductionTy::NotComplete);
if (const auto *CXXRD = dyn_cast<CXXRecordDecl>(RD);
- CXXRD && !CXXRD->isAggregate()) {
- EmitDiags(VarLoc, PDiag(diag::err_acc_reduction_type)
- << CXXRD << diag::OACCReductionTy::NotAgg);
- return ExprError();
- }
+ CXXRD && !CXXRD->isAggregate())
+ return EmitDiags(VarLoc, PDiag(diag::err_acc_reduction_type)
+ << CXXRD << diag::OACCReductionTy::NotAgg);
for (FieldDecl *FD : RD->fields()) {
if (!IsValidMemberOfComposite(FD->getType())) {
@@ -2007,17 +1985,37 @@ ExprResult SemaOpenACC::CheckReductionVar(OpenACCDirectiveKind DirectiveKind,
<< FD->getName() << RD->getName();
Notes.push_back({FD->getBeginLoc(), PD});
// TODO: member here.note_acc_reduction_member_of_composite
- EmitDiags(VarLoc, PDiag(diag::err_acc_reduction_type)
- << FD->getType()
- << diag::OACCReductionTy::MemberNotScalar);
- return ExprError();
+ return EmitDiags(VarLoc, PDiag(diag::err_acc_reduction_type)
+ << FD->getType()
+ << diag::OACCReductionTy::MemberNotScalar);
}
}
} else {
- EmitDiags(VarLoc, PDiag(diag::err_acc_reduction_type)
- << CurType << diag::OACCReductionTy::NotScalar);
+ return EmitDiags(VarLoc, PDiag(diag::err_acc_reduction_type)
+ << CurType
+ << diag::OACCReductionTy::NotScalar);
}
+ return false;
+}
+
+/// OpenACC 3.3 section 2.5.15:
+/// At a mininmum, the supported data types include ... the numerical data types
+/// in C, C++, and Fortran.
+///
+/// If the reduction var is a composite variable, each
+/// member of the composite variable must be a supported datatype for the
+/// reduction operation.
+ExprResult SemaOpenACC::CheckReductionVar(OpenACCDirectiveKind DirectiveKind,
+ OpenACCReductionOperator ReductionOp,
+ Expr *VarExpr) {
+ // For now, we only support 'scalar' types, or composites/arrays of scalar
+ // types.
+ VarExpr = VarExpr->IgnoreParenCasts();
+
+ if (CheckReductionVarType(VarExpr))
+ return ExprError();
+
// OpenACC3.3: 2.9.11: Reduction clauses on nested constructs for the same
// reduction 'var' must have the same reduction operator.
if (!VarExpr->isInstantiationDependent()) {
generated by cgit v1.1 at 2025-10-11 00:07:33 +0000