Skip to content

Commit

Permalink
Merged main:d50d7c37a159 into amd-gfx:1e7a931d181a
Browse files Browse the repository at this point in the history 
Local branch amd-gfx 1e7a931 Merged main:045304701bc6 into amd-gfx:7596d56af686
Remote branch main d50d7c3 [MBP] Prevent rotating a chain contains entry block
  • Loading branch information
Sw authored and Sw committed Dec 14, 2020
2 parents 1e7a931 + d50d7c3 commit 4d7163f
Show file tree
Hide file tree
Showing 38 changed files with 421 additions and 255 deletions.
2 changes: 2 additions & 0 deletions clang/include/clang/Basic/Attr.td
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1079,6 +1079,7 @@ def CUDADeviceBuiltinSurfaceType : InheritableAttr {
let LangOpts = [CUDA];
let Subjects = SubjectList<[CXXRecord]>;
let Documentation = [CUDADeviceBuiltinSurfaceTypeDocs];
let MeaningfulToClassTemplateDefinition = 1;
}

def CUDADeviceBuiltinTextureType : InheritableAttr {
Expand All @@ -1087,6 +1088,7 @@ def CUDADeviceBuiltinTextureType : InheritableAttr {
let LangOpts = [CUDA];
let Subjects = SubjectList<[CXXRecord]>;
let Documentation = [CUDADeviceBuiltinTextureTypeDocs];
let MeaningfulToClassTemplateDefinition = 1;
}

def CUDAGlobal : InheritableAttr {
Expand Down
20 changes: 20 additions & 0 deletions clang/test/SemaCUDA/device-use-host-var.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -158,3 +158,23 @@ void dev_lambda_capture_by_copy(int *out) {
});
}

// Texture references are special. As far as C++ is concerned they are host
// variables that are referenced from device code. However, they are handled
// very differently by the compiler under the hood and such references are
// allowed. Compiler should produce no warning here, but it should diagnose the
// same case without the device_builtin_texture_type attribute.
template <class, int = 1, int = 1>
struct __attribute__((device_builtin_texture_type)) texture {
static texture<int> ref;
__device__ int c() {
auto &x = ref;
}
};

template <class, int = 1, int = 1>
struct not_a_texture {
static not_a_texture<int> ref;
__device__ int c() {
auto &x = ref; // dev-error {{reference to __host__ variable 'ref' in __device__ function}}
}
};
17 changes: 12 additions & 5 deletions flang/include/flang/Common/restorer.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,25 +22,32 @@
namespace Fortran::common {
template <typename A> class Restorer {
public:
explicit Restorer(A &p) : p_{p}, original_{std::move(p)} {}
explicit Restorer(A &p, A original) : p_{p}, original_{std::move(original)} {}
~Restorer() { p_ = std::move(original_); }

// Inhibit any recreation of this restorer that would result in two restorers
// trying to restore the same reference.
Restorer(const Restorer &) = delete;
Restorer(Restorer &&that) = delete;
const Restorer &operator=(const Restorer &) = delete;
const Restorer &operator=(Restorer &&that) = delete;

private:
A &p_;
A original_;
};

template <typename A, typename B>
common::IfNoLvalue<Restorer<A>, B> ScopedSet(A &to, B &&from) {
Restorer<A> result{to};
A original{std::move(to)};
to = std::move(from);
return result;
return Restorer<A>{to, std::move(original)};
}
template <typename A, typename B>
common::IfNoLvalue<Restorer<A>, B> ScopedSet(A &to, const B &from) {
Restorer<A> result{to};
A original{std::move(to)};
to = from;
return result;
return Restorer<A>{to, std::move(original)};
}
} // namespace Fortran::common
#endif // FORTRAN_COMMON_RESTORER_H_
83 changes: 42 additions & 41 deletions flang/include/flang/Optimizer/Dialect/FIROps.td
View file
Original file line number Diff line number Diff line change
Expand Up @@ -249,7 +249,7 @@ class fir_AllocatableOp<string mnemonic, list<OpTrait> traits = []> :
}];

let printer = [{
p << getOperationName() << ' ' << getAttr(inType());
p << getOperationName() << ' ' << (*this)->getAttr(inType());
if (hasLenParams()) {
// print the LEN parameters to a derived type in parens
p << '(' << getLenParams() << " : " << getLenParams().getTypes() << ')';
Expand All @@ -267,7 +267,7 @@ class fir_AllocatableOp<string mnemonic, list<OpTrait> traits = []> :
static constexpr llvm::StringRef lenpName() { return "len_param_count"; }
mlir::Type getAllocatedType();

bool hasLenParams() { return bool{getAttr(lenpName())}; }
bool hasLenParams() { return bool{(*this)->getAttr(lenpName())}; }

unsigned numLenParams() {
if (auto val = (*this)->getAttrOfType<mlir::IntegerAttr>(lenpName()))
Expand Down Expand Up @@ -688,7 +688,7 @@ class fir_IntegralSwitchTerminatorOp<string mnemonic,
p << getOperationName() << ' ';
p.printOperand(getSelector());
p << " : " << getSelector().getType() << " [";
auto cases = getAttrOfType<mlir::ArrayAttr>(getCasesAttr()).getValue();
auto cases = (*this)->getAttrOfType<mlir::ArrayAttr>(getCasesAttr()).getValue();
auto count = getNumConditions();
for (decltype(count) i = 0; i != count; ++i) {
if (i)
Expand All @@ -711,7 +711,7 @@ class fir_IntegralSwitchTerminatorOp<string mnemonic,
getSelector().getType().isa<mlir::IndexType>() ||
getSelector().getType().isa<fir::IntType>()))
return emitOpError("must be an integer");
auto cases = getAttrOfType<mlir::ArrayAttr>(getCasesAttr()).getValue();
auto cases = (*this)->getAttrOfType<mlir::ArrayAttr>(getCasesAttr()).getValue();
auto count = getNumDest();
if (count == 0)
return emitOpError("must have at least one successor");
Expand Down Expand Up @@ -810,7 +810,7 @@ def fir_SelectCaseOp : fir_SwitchTerminatorOp<"select_case"> {
p << getOperationName() << ' ';
p.printOperand(getSelector());
p << " : " << getSelector().getType() << " [";
auto cases = getAttrOfType<mlir::ArrayAttr>(getCasesAttr()).getValue();
auto cases = (*this)->getAttrOfType<mlir::ArrayAttr>(getCasesAttr()).getValue();
auto count = getNumConditions();
for (decltype(count) i = 0; i != count; ++i) {
if (i)
Expand Down Expand Up @@ -839,7 +839,7 @@ def fir_SelectCaseOp : fir_SwitchTerminatorOp<"select_case"> {
getSelector().getType().isa<fir::LogicalType>() ||
getSelector().getType().isa<fir::CharacterType>()))
return emitOpError("must be an integer, character, or logical");
auto cases = getAttrOfType<mlir::ArrayAttr>(getCasesAttr()).getValue();
auto cases = (*this)->getAttrOfType<mlir::ArrayAttr>(getCasesAttr()).getValue();
auto count = getNumDest();
if (count == 0)
return emitOpError("must have at least one successor");
Expand Down Expand Up @@ -925,7 +925,7 @@ def fir_SelectTypeOp : fir_SwitchTerminatorOp<"select_type"> {
p << getOperationName() << ' ';
p.printOperand(getSelector());
p << " : " << getSelector().getType() << " [";
auto cases = getAttrOfType<mlir::ArrayAttr>(getCasesAttr()).getValue();
auto cases = (*this)->getAttrOfType<mlir::ArrayAttr>(getCasesAttr()).getValue();
auto count = getNumConditions();
for (decltype(count) i = 0; i != count; ++i) {
if (i)
Expand All @@ -941,7 +941,7 @@ def fir_SelectTypeOp : fir_SwitchTerminatorOp<"select_type"> {
let verifier = [{
if (!(getSelector().getType().isa<fir::BoxType>()))
return emitOpError("must be a boxed type");
auto cases = getAttrOfType<mlir::ArrayAttr>(getCasesAttr()).getValue();
auto cases = (*this)->getAttrOfType<mlir::ArrayAttr>(getCasesAttr()).getValue();
auto count = getNumDest();
if (count == 0)
return emitOpError("must have at least one successor");
Expand Down Expand Up @@ -1056,7 +1056,7 @@ def fir_EmboxOp : fir_Op<"embox", [NoSideEffect]> {
if (getNumOperands() == 2) {
p << ", ";
p.printOperands(dims());
} else if (auto map = getAttr(layoutName())) {
} else if (auto map = (*this)->getAttr(layoutName())) {
p << " [" << map << ']';
}
p.printOptionalAttrDict(getAttrs(), {layoutName(), lenpName()});
Expand Down Expand Up @@ -1097,9 +1097,9 @@ def fir_EmboxOp : fir_Op<"embox", [NoSideEffect]> {
let extraClassDeclaration = [{
static constexpr llvm::StringRef layoutName() { return "layout_map"; }
static constexpr llvm::StringRef lenpName() { return "len_param_count"; }
bool hasLenParams() { return bool{getAttr(lenpName())}; }
bool hasLenParams() { return bool{(*this)->getAttr(lenpName())}; }
unsigned numLenParams() {
if (auto x = getAttrOfType<mlir::IntegerAttr>(lenpName()))
if (auto x = (*this)->getAttrOfType<mlir::IntegerAttr>(lenpName()))
return x.getInt();
return 0;
}
Expand Down Expand Up @@ -1213,13 +1213,13 @@ def fir_EmboxProcOp : fir_Op<"emboxproc", [NoSideEffect]> {
}];

let printer = [{
p << getOperationName() << ' ' << getAttr("funcname");
p << getOperationName() << ' ' << (*this)->getAttr("funcname");
auto h = host();
if (h) {
p << ", ";
p.printOperand(h);
}
p << " : (" << getAttr("functype");
p << " : (" << (*this)->getAttr("functype");
if (h)
p << ", " << h.getType();
p << ") -> " << getType();
Expand Down Expand Up @@ -1587,7 +1587,7 @@ def fir_CoordinateOp : fir_Op<"coordinate_of", [NoSideEffect]> {
if (!ref().getType().dyn_cast<BoxType>())
return emitOpError("len_param_index must be used on box type");
}
if (auto attr = getAttr(CoordinateOp::baseType())) {
if (auto attr = (*this)->getAttr(CoordinateOp::baseType())) {
if (!attr.isa<mlir::TypeAttr>())
return emitOpError("improperly constructed");
} else {
Expand Down Expand Up @@ -1690,8 +1690,8 @@ def fir_FieldIndexOp : fir_OneResultOp<"field_index", [NoSideEffect]> {

let printer = [{
p << getOperationName() << ' '
<< getAttrOfType<mlir::StringAttr>(fieldAttrName()).getValue() << ", "
<< getAttr(typeAttrName());
<< (*this)->getAttrOfType<mlir::StringAttr>(fieldAttrName()).getValue()
<< ", " << (*this)->getAttr(typeAttrName());
if (getNumOperands()) {
p << '(';
p.printOperands(lenparams());
Expand Down Expand Up @@ -1826,8 +1826,8 @@ def fir_LenParamIndexOp : fir_OneResultOp<"len_param_index", [NoSideEffect]> {

let printer = [{
p << getOperationName() << ' '
<< getAttrOfType<mlir::StringAttr>(fieldAttrName()).getValue() << ", "
<< getAttr(typeAttrName());
<< (*this)->getAttrOfType<mlir::StringAttr>(fieldAttrName()).getValue()
<< ", " << (*this)->getAttr(typeAttrName());
}];

let builders = [
Expand All @@ -1841,7 +1841,7 @@ def fir_LenParamIndexOp : fir_OneResultOp<"len_param_index", [NoSideEffect]> {
static constexpr llvm::StringRef fieldAttrName() { return "field_id"; }
static constexpr llvm::StringRef typeAttrName() { return "on_type"; }
mlir::Type getOnType() {
return getAttrOfType<TypeAttr>(typeAttrName()).getValue();
return (*this)->getAttrOfType<TypeAttr>(typeAttrName()).getValue();
}
}];
}
Expand Down Expand Up @@ -2166,7 +2166,7 @@ def fir_DispatchOp : fir_Op<"dispatch",
}];

let printer = [{
p << getOperationName() << ' ' << getAttr("method") << '(';
p << getOperationName() << ' ' << (*this)->getAttr("method") << '(';
p.printOperand(object());
if (arg_operand_begin() != arg_operand_end()) {
p << ", ";
Expand Down Expand Up @@ -2250,7 +2250,7 @@ def fir_StringLitOp : fir_Op<"string_lit", [NoSideEffect]> {
auto eleTy = getType().cast<fir::SequenceType>().getEleTy();
if (!eleTy.isa<fir::CharacterType>())
return emitOpError("must have !fir.char type");
if (auto xl = getAttr(xlist())) {
if (auto xl = (*this)->getAttr(xlist())) {
auto xList = xl.cast<mlir::ArrayAttr>();
for (auto a : xList)
if (!a.isa<mlir::IntegerAttr>())
Expand All @@ -2265,12 +2265,12 @@ def fir_StringLitOp : fir_Op<"string_lit", [NoSideEffect]> {
static constexpr const char *xlist() { return "xlist"; }

// Get the LEN attribute of this character constant
mlir::Attribute getSize() { return getAttr(size()); }
mlir::Attribute getSize() { return (*this)->getAttr(size()); }
// Get the string value of this character constant
mlir::Attribute getValue() {
if (auto attr = getAttr(value()))
if (auto attr = (*this)->getAttr(value()))
return attr;
return getAttr(xlist());
return (*this)->getAttr(xlist());
}

/// Is this a wide character literal (1 character > 8 bits)
Expand Down Expand Up @@ -2381,7 +2381,7 @@ def fir_CmpfOp : fir_Op<"cmpf",
static CmpFPredicate getPredicateByName(llvm::StringRef name);

CmpFPredicate getPredicate() {
return (CmpFPredicate)getAttrOfType<mlir::IntegerAttr>(
return (CmpFPredicate)(*this)->getAttrOfType<mlir::IntegerAttr>(
getPredicateAttrName()).getInt();
}
}];
Expand Down Expand Up @@ -2415,11 +2415,11 @@ def fir_ConstcOp : fir_Op<"constc", [NoSideEffect]> {

let printer = [{
p << getOperationName() << " (0x";
auto f1 = getAttr(realAttrName()).cast<mlir::FloatAttr>();
auto f1 = (*this)->getAttr(realAttrName()).cast<mlir::FloatAttr>();
auto i1 = f1.getValue().bitcastToAPInt();
p.getStream().write_hex(i1.getZExtValue());
p << ", 0x";
auto f2 = getAttr(imagAttrName()).cast<mlir::FloatAttr>();
auto f2 = (*this)->getAttr(imagAttrName()).cast<mlir::FloatAttr>();
auto i2 = f2.getValue().bitcastToAPInt();
p.getStream().write_hex(i2.getZExtValue());
p << ") : ";
Expand All @@ -2436,8 +2436,8 @@ def fir_ConstcOp : fir_Op<"constc", [NoSideEffect]> {
static constexpr llvm::StringRef realAttrName() { return "real"; }
static constexpr llvm::StringRef imagAttrName() { return "imaginary"; }

mlir::Attribute getReal() { return getAttr(realAttrName()); }
mlir::Attribute getImaginary() { return getAttr(imagAttrName()); }
mlir::Attribute getReal() { return (*this)->getAttr(realAttrName()); }
mlir::Attribute getImaginary() { return (*this)->getAttr(imagAttrName()); }
}];
}

Expand Down Expand Up @@ -2485,7 +2485,7 @@ def fir_CmpcOp : fir_Op<"cmpc",
}

CmpFPredicate getPredicate() {
return (CmpFPredicate)getAttrOfType<mlir::IntegerAttr>(
return (CmpFPredicate)(*this)->getAttrOfType<mlir::IntegerAttr>(
getPredicateAttrName()).getInt();
}
}];
Expand Down Expand Up @@ -2601,7 +2601,7 @@ def fir_GenTypeDescOp : fir_OneResultOp<"gentypedesc", [NoSideEffect]> {
}];

let printer = [{
p << getOperationName() << ' ' << getAttr("in_type");
p << getOperationName() << ' ' << (*this)->getAttr("in_type");
p.printOptionalAttrDict(getAttrs(), {"in_type"});
}];

Expand All @@ -2623,7 +2623,7 @@ def fir_GenTypeDescOp : fir_OneResultOp<"gentypedesc", [NoSideEffect]> {
let extraClassDeclaration = [{
mlir::Type getInType() {
// get the type that the type descriptor describes
return getAttrOfType<mlir::TypeAttr>("in_type").getValue();
return (*this)->getAttrOfType<mlir::TypeAttr>("in_type").getValue();
}
}];
}
Expand Down Expand Up @@ -2697,7 +2697,7 @@ def fir_GlobalOp : fir_Op<"global", [IsolatedFromAbove, Symbol]> {
if (linkName().hasValue())
p << ' ' << linkName().getValue();
p << ' ';
p.printAttributeWithoutType(getAttr(symbolAttrName()));
p.printAttributeWithoutType((*this)->getAttr(symbolAttrName()));
if (auto val = getValueOrNull())
p << '(' << val << ')';
if ((*this)->getAttr(constantAttrName()))
Expand Down Expand Up @@ -2738,7 +2738,7 @@ def fir_GlobalOp : fir_Op<"global", [IsolatedFromAbove, Symbol]> {

/// The printable type of the global
mlir::Type getType() {
return getAttrOfType<TypeAttr>(typeAttrName()).getValue();
return (*this)->getAttrOfType<TypeAttr>(typeAttrName()).getValue();
}

/// The semantic type of the global
Expand Down Expand Up @@ -2768,8 +2768,9 @@ def fir_GlobalOp : fir_Op<"global", [IsolatedFromAbove, Symbol]> {
}

mlir::FlatSymbolRefAttr getSymbol() {
return mlir::FlatSymbolRefAttr::get(getAttrOfType<mlir::StringAttr>(
mlir::SymbolTable::getSymbolAttrName()).getValue(), getContext());
return mlir::FlatSymbolRefAttr::get(
(*this)->getAttrOfType<mlir::StringAttr>(
mlir::SymbolTable::getSymbolAttrName()).getValue(), getContext());
}
}];
}
Expand Down Expand Up @@ -2811,8 +2812,8 @@ def fir_GlobalLenOp : fir_Op<"global_len", []> {
}];

let printer = [{
p << getOperationName() << ' ' << getAttr(lenParamAttrName()) << ", "
<< getAttr(intAttrName());
p << getOperationName() << ' ' << (*this)->getAttr(lenParamAttrName())
<< ", " << (*this)->getAttr(intAttrName());
}];

let extraClassDeclaration = [{
Expand Down Expand Up @@ -2865,7 +2866,7 @@ def fir_DispatchTableOp : fir_Op<"dispatch_table",
}];

let printer = [{
auto tableName = getAttrOfType<StringAttr>(
auto tableName = (*this)->getAttrOfType<StringAttr>(
mlir::SymbolTable::getSymbolAttrName()).getValue();
p << getOperationName() << " @" << tableName;

Expand Down Expand Up @@ -2946,8 +2947,8 @@ def fir_DTEntryOp : fir_Op<"dt_entry", []> {
}];

let printer = [{
p << getOperationName() << ' ' << getAttr(methodAttrName()) << ", "
<< getAttr(procAttrName());
p << getOperationName() << ' ' << (*this)->getAttr(methodAttrName()) << ", "
<< (*this)->getAttr(procAttrName());
}];

let extraClassDeclaration = [{
Expand Down
2 changes: 1 addition & 1 deletion flang/include/flang/Optimizer/Dialect/FIROpsSupport.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@ namespace fir {
/// return true iff the Operation is a non-volatile LoadOp
inline bool nonVolatileLoad(mlir::Operation *op) {
if (auto load = dyn_cast<fir::LoadOp>(op))
return !load.getAttr("volatile");
return !load->getAttr("volatile");
return false;
}

Expand Down
Loading

0 comments on commit 4d7163f

Please sign in to comment.