From f3d637ba4d90bc2e23db07f1a9df5a6be7533f08 Mon Sep 17 00:00:00 2001 From: Tim Foley Date: Thu, 4 Jun 2020 11:53:13 -0700 Subject: First steps toward inheritance for struct types (#1366) * First steps toward inheritance for struct types This change adds the ability for a `struct` type to declare a base type that is another `struct`: ```hlsl struct Base { int baseMember; } struct Derived : Base { int derivedMember; } ``` The semantics of the feature are that code like the above desugars into code like: ```hlsl struct Base { int baseMember; } struct Derived { Base _base; int derivedMember; } ``` At points where a member from the base type is being projected out, or the value is being implicitly cast to the base type, the compiler transforms the code to reference the implicitly-generated `_base` member. That means code like this: ```hlsl void f(Base b); ... Derived d = ...; int x = d.baseMember; f(d); ``` gets transformed into a form like this: ```hlsl void f(Base b); ... Derived d = ...; int x = d._base.baseMember; f(d._base); ``` Note that as a result of this choice, the behavior when passing a `Derived` value to a function that expects a `Base` (including to inherited member functions) is that of "object shearing" from the C++ world: the called function can only "see" the `Base` part of the argument, and any operations performed on it will behave as if the value was indeed a `Base`. There is no polymorphism going on because Slang doesn't currently have `virtual` methods. In an attempt to work toward inheritance being a robust feature, this change adds a bunch of more detailed logic for checking the bases of various declarations: * An `interface` declaration is only allowed to inherit from other `interface`s * An `extension` declaration can only introduce inheritance from `interface`s * A `struct` declaration can only inherit from at most one other `struct`, and that `struct` must be the first entry in the list of bases This change also adds a mechanism to control whether a `struct` or `interface` in one module can inherit from a `struct` or `interface` declared in another module: * If the base declaration is marked `[open]`, then the inheritance is allowed * If the base declaration is marked `[sealed]`, then the inheritance is allowed * If it is not marked otherwise, a `struct` is implicitly `[sealed]` * If it is not marked otherwise, an `interface` is implicitly `[open]` These seem like reasonable defaults. In order to safeguard the standard library a bit, the interfaces for builtin types have been marked `[sealed]` to make sure that a user cannot declare a `struct` and then mark it as a `BuiltinFloatingPointType`. This step should bring us a bit closer to being able to document and expose these interfaces for built-in types so that users can write code that is generic over them. There are some big caveats with this work, such that it really only represents a stepping-stone toward a usable inheritance feature. The most important caveats are: * If a `Derived` type tries to conform to an interface, such that one or more interface requirements are satisfied with members inherited from the `Base` type, that is likely to cause a crash or incorrect code generation. * If a `Derived` type tries to inherit from a `Base` type that conforms to one or more interfaces, the witness table generated for the conformance of `Derived` to that interface is likely to lead to a crash or incorrect code generation. It is clear that solving both of those issues will be necessary before we can really promote `struct` inheritance as a feature for users to try out. * fixup: trying to appease clang error * fixups: review feedback --- source/slang/core.meta.slang | 14 ++ source/slang/slang-ast-expr.h | 16 +- source/slang/slang-ast-modifier.h | 10 + source/slang/slang-check-conformance.cpp | 71 +++--- source/slang/slang-check-constraint.cpp | 4 +- source/slang/slang-check-conversion.cpp | 10 +- source/slang/slang-check-decl.cpp | 377 ++++++++++++++++++++++++++----- source/slang/slang-check-expr.cpp | 9 +- source/slang/slang-check-impl.h | 60 +++-- source/slang/slang-diagnostic-defs.h | 16 +- source/slang/slang-lower-to-ir.cpp | 168 +++++++++++--- 11 files changed, 614 insertions(+), 141 deletions(-) (limited to 'source') diff --git a/source/slang/core.meta.slang b/source/slang/core.meta.slang index 2b43206a3..0c39f2c2f 100644 --- a/source/slang/core.meta.slang +++ b/source/slang/core.meta.slang @@ -18,9 +18,11 @@ syntax constexpr : ConstExprModifier; syntax globallycoherent : GloballyCoherentModifier; // A type that can be used as an operand for builtins +[sealed] interface __BuiltinType {} // A type that can be used for arithmetic operations +[sealed] interface __BuiltinArithmeticType : __BuiltinType { /// Initialize from a 32-bit signed integer value. @@ -28,19 +30,24 @@ interface __BuiltinArithmeticType : __BuiltinType } /// A type that can be used for logical/bitwsie operations +[sealed] interface __BuiltinLogicalType : __BuiltinType {} // A type that logically has a sign (positive/negative/zero) +[sealed] interface __BuiltinSignedArithmeticType : __BuiltinArithmeticType {} // A type that can represent integers +[sealed] interface __BuiltinIntegerType : __BuiltinArithmeticType {} // A type that can represent non-integers +[sealed] interface __BuiltinRealType : __BuiltinSignedArithmeticType {} // A type that uses a floating-point representation +[sealed] interface __BuiltinFloatingPointType : __BuiltinRealType { /// Initialize from a 32-bit floating-point value. @@ -1917,3 +1924,10 @@ attribute_syntax [__extern] : ExternAttribute; __attributeTarget(FunctionDeclBase) attribute_syntax [__unsafeForceInlineEarly] : UnsafeForceInlineEarlyAttribute; + +// Inheritance Control +__attributeTarget(AggTypeDecl) +attribute_syntax [sealed] : SealedAttribute; + +__attributeTarget(AggTypeDecl) +attribute_syntax [open] : OpenAttribute; diff --git a/source/slang/slang-ast-expr.h b/source/slang/slang-ast-expr.h index e7111e631..96e90ae02 100644 --- a/source/slang/slang-ast-expr.h +++ b/source/slang/slang-ast-expr.h @@ -228,15 +228,21 @@ class ImplicitCastExpr : public TypeCastExpr SLANG_CLASS(ImplicitCastExpr) }; - /// A cast from a value to an interface ("existential") type. -class CastToInterfaceExpr: public Expr + /// A cast of a value to a super-type of its type. + /// + /// The type being cast to is stored as this expression's `type`. + /// +class CastToSuperTypeExpr: public Expr { - SLANG_CLASS(CastToInterfaceExpr) + SLANG_CLASS(CastToSuperTypeExpr) - /// The value being cast to an interface type + /// The value being cast to a super type + /// + /// The type being case from is `valueArg->type`. + /// RefPtr valueArg; - /// A witness showing that `valueArg` conforms to the chosen interface + /// A witness showing that `valueArg`'s type is a sub-type of this expression's `type` RefPtr witnessArg; }; diff --git a/source/slang/slang-ast-modifier.h b/source/slang/slang-ast-modifier.h index 065118c7a..4badbd02a 100644 --- a/source/slang/slang-ast-modifier.h +++ b/source/slang/slang-ast-modifier.h @@ -873,4 +873,14 @@ class UnsafeForceInlineEarlyAttribute : public Attribute SLANG_CLASS(UnsafeForceInlineEarlyAttribute) }; + /// An attribute that marks a type declaration as either allowing or + /// disallowing the type to be inherited from in other modules. +class InheritanceControlAttribute : public Attribute { SLANG_CLASS(InheritanceControlAttribute) }; + + /// An attribute that marks a type declaration as allowing the type to be inherited from in other modules. +class OpenAttribute : public InheritanceControlAttribute { SLANG_CLASS(OpenAttribute) }; + + /// An attribute that marks a type declaration as disallowing the type to be inherited from in other modules. +class SealedAttribute : public InheritanceControlAttribute { SLANG_CLASS(SealedAttribute) }; + } // namespace Slang diff --git a/source/slang/slang-check-conformance.cpp b/source/slang/slang-check-conformance.cpp index 4d54a93fb..f57e7cb38 100644 --- a/source/slang/slang-check-conformance.cpp +++ b/source/slang/slang-check-conformance.cpp @@ -18,15 +18,15 @@ namespace Slang } RefPtr SemanticsVisitor::createTypeWitness( - RefPtr type, - DeclRef interfaceDeclRef, + RefPtr subType, + DeclRef superTypeDeclRef, TypeWitnessBreadcrumb* inBreadcrumbs) { if(!inBreadcrumbs) { // We need to construct a witness to the fact - // that `type` has been proven to be *equal* - // to `interfaceDeclRef`. + // that `subType` has been proven to be *equal* + // to `superTypeDeclRef`. // SLANG_UNEXPECTED("reflexive type witness"); UNREACHABLE_RETURN(nullptr); @@ -145,15 +145,15 @@ namespace Slang } } - bool SemanticsVisitor::doesTypeConformToInterfaceImpl( - RefPtr originalType, - RefPtr type, - DeclRef interfaceDeclRef, + bool SemanticsVisitor::_isDeclaredSubtype( + RefPtr originalSubType, + RefPtr subType, + DeclRef superTypeDeclRef, RefPtr* outWitness, TypeWitnessBreadcrumb* inBreadcrumbs) { // for now look up a conformance member... - if(auto declRefType = as(type)) + if(auto declRefType = as(subType)) { auto declRef = declRefType->declRef; @@ -162,11 +162,11 @@ namespace Slang // TODO: This is actually a bit more complicated, as // the interface needs to be "object-safe" for us to // really make this determination... - if(declRef == interfaceDeclRef) + if(declRef == superTypeDeclRef) { if(outWitness) { - *outWitness = createTypeWitness(originalType, interfaceDeclRef, inBreadcrumbs); + *outWitness = createTypeWitness(originalSubType, superTypeDeclRef, inBreadcrumbs); } return true; } @@ -198,11 +198,11 @@ namespace Slang TypeWitnessBreadcrumb breadcrumb; breadcrumb.prev = inBreadcrumbs; - breadcrumb.sub = type; + breadcrumb.sub = subType; breadcrumb.sup = inheritedType; breadcrumb.declRef = inheritanceDeclRef; - if(doesTypeConformToInterfaceImpl(originalType, inheritedType, interfaceDeclRef, outWitness, &breadcrumb)) + if(_isDeclaredSubtype(originalSubType, inheritedType, superTypeDeclRef, outWitness, &breadcrumb)) { return true; } @@ -214,10 +214,10 @@ namespace Slang auto inheritedType = getSup(m_astBuilder, genConstraintDeclRef); TypeWitnessBreadcrumb breadcrumb; breadcrumb.prev = inBreadcrumbs; - breadcrumb.sub = type; + breadcrumb.sub = subType; breadcrumb.sup = inheritedType; breadcrumb.declRef = genConstraintDeclRef; - if (doesTypeConformToInterfaceImpl(originalType, inheritedType, interfaceDeclRef, outWitness, &breadcrumb)) + if (_isDeclaredSubtype(originalSubType, inheritedType, superTypeDeclRef, outWitness, &breadcrumb)) { return true; } @@ -252,14 +252,14 @@ namespace Slang breadcrumb.sup = sup; breadcrumb.declRef = constraintDeclRef; - if(doesTypeConformToInterfaceImpl(originalType, sup, interfaceDeclRef, outWitness, &breadcrumb)) + if(_isDeclaredSubtype(originalSubType, sup, superTypeDeclRef, outWitness, &breadcrumb)) { return true; } } } } - else if(auto taggedUnionType = as(type)) + else if(auto taggedUnionType = as(subType)) { // A tagged union type conforms to an interface if all of // the constituent types in the tagged union conform. @@ -276,10 +276,10 @@ namespace Slang { RefPtr caseWitness; - if(!doesTypeConformToInterfaceImpl( + if(!_isDeclaredSubtype( caseType, caseType, - interfaceDeclRef, + superTypeDeclRef, outWitness ? &caseWitness : nullptr, nullptr)) { @@ -304,8 +304,11 @@ namespace Slang // We will start out being conservative about what we accept // here, just to keep things simple. // - if(!isInterfaceSafeForTaggedUnion(interfaceDeclRef)) - return false; + if( auto superInterfaceDeclRef = superTypeDeclRef.as() ) + { + if(!isInterfaceSafeForTaggedUnion(superInterfaceDeclRef)) + return false; + } // If we reach this point then we have a concrete // witness for each of the case types, and that is @@ -315,7 +318,7 @@ namespace Slang { RefPtr taggedUnionWitness = m_astBuilder->create(); taggedUnionWitness->sub = taggedUnionType; - taggedUnionWitness->sup = DeclRefType::create(m_astBuilder, interfaceDeclRef); + taggedUnionWitness->sup = DeclRefType::create(m_astBuilder, superTypeDeclRef); taggedUnionWitness->caseWitnesses.swapWith(caseWitnesses); *outWitness = taggedUnionWitness; @@ -327,22 +330,30 @@ namespace Slang return false; } - bool SemanticsVisitor::DoesTypeConformToInterface( - RefPtr type, - DeclRef interfaceDeclRef) + bool SemanticsVisitor::isDeclaredSubtype( + RefPtr subType, + DeclRef superTypeDeclRef) { - return doesTypeConformToInterfaceImpl(type, type, interfaceDeclRef, nullptr, nullptr); + return _isDeclaredSubtype(subType, subType, superTypeDeclRef, nullptr, nullptr); } - RefPtr SemanticsVisitor::tryGetInterfaceConformanceWitness( - RefPtr type, - DeclRef interfaceDeclRef) + RefPtr SemanticsVisitor::tryGetSubtypeWitness( + RefPtr subType, + DeclRef superTypeDeclRef) { RefPtr result; - doesTypeConformToInterfaceImpl(type, type, interfaceDeclRef, &result, nullptr); + _isDeclaredSubtype(subType, subType, superTypeDeclRef, &result, nullptr); return result; } + + RefPtr SemanticsVisitor::tryGetInterfaceConformanceWitness( + RefPtr type, + DeclRef interfaceDeclRef) + { + return tryGetSubtypeWitness(type, interfaceDeclRef); + } + RefPtr SemanticsVisitor::createTypeEqualityWitness( Type* type) { diff --git a/source/slang/slang-check-constraint.cpp b/source/slang/slang-check-constraint.cpp index 427ba9ec2..c37af8892 100644 --- a/source/slang/slang-check-constraint.cpp +++ b/source/slang/slang-check-constraint.cpp @@ -80,7 +80,7 @@ namespace Slang DeclRef interfaceDeclRef) { // The most basic test here should be: does the type declare conformance to the trait. - if(DoesTypeConformToInterface(type, interfaceDeclRef)) + if(isDeclaredSubtype(type, interfaceDeclRef)) return type; // Just because `type` doesn't conform to the given `interfaceDeclRef`, that @@ -119,7 +119,7 @@ namespace Slang continue; // We only want to consider types that implement the target interface. - if(!DoesTypeConformToInterface(candidateType, interfaceDeclRef)) + if(!isDeclaredSubtype(candidateType, interfaceDeclRef)) continue; // We only want to consider types where we can implicitly convert from `type` diff --git a/source/slang/slang-check-conversion.cpp b/source/slang/slang-check-conversion.cpp index 1f34e240d..c4e809025 100644 --- a/source/slang/slang-check-conversion.cpp +++ b/source/slang/slang-check-conversion.cpp @@ -544,12 +544,12 @@ namespace Slang if (auto toDeclRefType = as(toType)) { auto toTypeDeclRef = toDeclRefType->declRef; - if (auto interfaceDeclRef = toTypeDeclRef.as()) + if (auto toAggTypeDeclRef = toTypeDeclRef.as()) { - if(auto witness = tryGetInterfaceConformanceWitness(fromType, interfaceDeclRef)) + if(auto witness = tryGetSubtypeWitness(fromType, toAggTypeDeclRef)) { if (outToExpr) - *outToExpr = createCastToInterfaceExpr(toType, fromExpr, witness); + *outToExpr = createCastToSuperTypeExpr(toType, fromExpr, witness); if (outCost) *outCost = kConversionCost_CastToInterface; return true; @@ -835,12 +835,12 @@ namespace Slang return castExpr; } - RefPtr SemanticsVisitor::createCastToInterfaceExpr( + RefPtr SemanticsVisitor::createCastToSuperTypeExpr( RefPtr toType, RefPtr fromExpr, RefPtr witness) { - RefPtr expr = m_astBuilder->create(); + RefPtr expr = m_astBuilder->create(); expr->loc = fromExpr->loc; expr->type = QualType(toType); expr->valueArg = fromExpr; diff --git a/source/slang/slang-check-decl.cpp b/source/slang/slang-check-decl.cpp index 66a8ceaf8..8486bf107 100644 --- a/source/slang/slang-check-decl.cpp +++ b/source/slang/slang-check-decl.cpp @@ -117,10 +117,24 @@ namespace Slang void visitInheritanceDecl(InheritanceDecl* inheritanceDecl); - void visitAggTypeDecl(AggTypeDecl* decl); + /// Validate that `decl` isn't illegally inheriting from a type in another module. + /// + /// This call checks a single `inheritanceDecl` to make sure that it either + /// * names a base type from the same module as `decl`, or + /// * names a type that allows cross-module inheritance + void _validateCrossModuleInheritance( + AggTypeDeclBase* decl, + InheritanceDecl* inheritanceDecl); + + void visitInterfaceDecl(InterfaceDecl* decl); + + void visitStructDecl(StructDecl* decl); void visitEnumDecl(EnumDecl* decl); + /// Validate that the target type of an extension `decl` is valid. + void _validateExtensionDeclTargetType(ExtensionDecl* decl); + void visitExtensionDecl(ExtensionDecl* decl); }; @@ -984,25 +998,10 @@ namespace Slang base = TranslateTypeNode(base); inheritanceDecl->base = base; - // For now we only allow inheritance from interfaces, so - // we will validate that the type expression names an interface - - if(auto declRefType = as(base.type)) - { - if(auto interfaceDeclRef = declRefType->declRef.as()) - { - return; - } - } - else if(base.type.is()) - { - // If an error was already produced, don't emit a cascading error. - return; - } - - // If type expression didn't name an interface, we'll emit an error here - // TODO: deal with the case of an error in the type expression (don't cascade) - getSink()->diagnose( base.exp, Diagnostics::expectedAnInterfaceGot, base.type); + // Note: we do not check whether the type being inherited from + // is valid to use for inheritance here, because there could + // be contextual factors that need to be taken into account + // based on the declaration that is doing the inheriting. } // Concretize interface conformances so that we have witnesses as required for lookup. @@ -1660,6 +1659,12 @@ namespace Slang inheritanceDecl, baseInterfaceDeclRef); } + else if( auto structDeclRef = baseTypeDeclRef.as() ) + { + // The type is saying it inherits from a `struct`, + // which doesn't require any checking at present + return nullptr; + } } getSink()->diagnose(inheritanceDecl, Diagnostics::unimplemented, "type not supported for inheritance"); @@ -1761,15 +1766,200 @@ namespace Slang } } - void SemanticsDeclBasesVisitor::visitAggTypeDecl(AggTypeDecl* decl) + void SemanticsDeclBasesVisitor::_validateCrossModuleInheritance( + AggTypeDeclBase* decl, + InheritanceDecl* inheritanceDecl) { - // TODO: We need to enumerate the bases here, - // and ideally form a "class precedence list" - // from them. + // Within a single module, users should be allowed to inherit + // one type from another more or less freely, so long as they + // don't violate fundamental validity conditions around + // inheritance. + // + // When an inheritance relationship is declared in one module, + // and the base type is in another module, we may want to + // enforce more restrictions. As a strong example, we probably + // don't want people to declare their own subtype of `int` + // or `Texture2D`. + // + // We start by checking if the type being inherited from is + // a decl-ref type, since that means it refers to a declaration + // that can be localized to its original module. + // + auto baseType = inheritanceDecl->base.type; + auto baseDeclRefType = as(baseType); + if( !baseDeclRefType ) + { + return; + } + auto baseDecl = baseDeclRefType->declRef.decl; + + // Using the parent/child hierarchy baked into `Decl`s we + // can find the modules that contain both the `decl` doing + // the inheriting, and the `baseDeclRefType` that is being + // inherited from. + // + // If those modules are the same, then we aren't seeing any + // kind of cross-module inheritance here, and there is nothing + // that needs enforcing. + // + auto moduleWithInheritance = getModule(decl); + auto moduleWithBaseType = getModule(baseDecl); + if( moduleWithInheritance == moduleWithBaseType ) + { + return; + } + + if( baseDecl->hasModifier() ) + { + // If the original declaration had the `[sealed]` attribute on it, + // then it explicitly does *not* allow inheritance from other + // modules. + // + getSink()->diagnose(inheritanceDecl, Diagnostics::cannotInheritFromExplicitlySealedDeclarationInAnotherModule, baseType, moduleWithBaseType->getModuleDecl()->getName()); + return; + } + else if( baseDecl->hasModifier() ) + { + // Conversely, if the original declaration had the `[open]` attribute + // on it, then it explicit *does* allow inheritance from other + // modules. + // + // In this case we don't need to check anything: the inheritance + // is allowed. + } + else if( as(baseDecl) ) + { + // If an interface isn't explicitly marked `[open]` or `[sealed]`, + // then the default behavior is to treat it as `[open]`, since + // interfaces are most often used to define protocols that + // users of a module can opt into. + } + else + { + // For any non-interface type, if the declaration didn't specify + // `[open]` or `[sealed]` then we assume `[sealed]` is the default. + // + getSink()->diagnose(inheritanceDecl, Diagnostics::cannotInheritFromImplicitlySealedDeclarationInAnotherModule, baseType, moduleWithBaseType->getModuleDecl()->getName()); + return; + } + } + void SemanticsDeclBasesVisitor::visitInterfaceDecl(InterfaceDecl* decl) + { for( auto inheritanceDecl : decl->getMembersOfType() ) { ensureDecl(inheritanceDecl, DeclCheckState::CanUseBaseOfInheritanceDecl); + auto baseType = inheritanceDecl->base.type; + + // It is possible that there was an error in checking the base type + // expression, and in such a case we shouldn't emit a cascading error. + // + if( auto baseErrorType = as(baseType) ) + { + continue; + } + + // An `interface` type can only inherit from other `interface` types. + // + // TODO: In the long run it might make sense for an interface to support + // an inheritance clause naming a non-interface type, with the meaning + // that any type that implements the interface must be a sub-type of the + // type named in the inheritance clause. + // + auto baseDeclRefType = as(baseType); + if( !baseDeclRefType ) + { + getSink()->diagnose(inheritanceDecl, Diagnostics::baseOfInterfaceMustBeInterface, decl, baseType); + continue; + } + + auto baseDeclRef = baseDeclRefType->declRef; + auto baseInterfaceDeclRef = baseDeclRef.as(); + if( !baseInterfaceDeclRef ) + { + getSink()->diagnose(inheritanceDecl, Diagnostics::baseOfInterfaceMustBeInterface, decl, baseType); + continue; + } + + // TODO: At this point we have the `baseInterfaceDeclRef` + // and could use it to perform further validity checks, + // and/or to build up a more refined representation of + // the inheritance graph for this type (e.g., a "class + // precedence list"). + // + // E.g., we can/should check that we aren't introducing + // a circular inheritance relationship. + + _validateCrossModuleInheritance(decl, inheritanceDecl); + } + } + + void SemanticsDeclBasesVisitor::visitStructDecl(StructDecl* decl) + { + // A `struct` type can only inherit from `struct` or `interface` types. + // + // Furthermore, only the first inheritance clause (in source + // order) is allowed to declare a base `struct` type. + // + Index inheritanceClauseCounter = 0; + for( auto inheritanceDecl : decl->getMembersOfType() ) + { + Index inheritanceClauseIndex = inheritanceClauseCounter++; + + ensureDecl(inheritanceDecl, DeclCheckState::CanUseBaseOfInheritanceDecl); + auto baseType = inheritanceDecl->base.type; + + // It is possible that there was an error in checking the base type + // expression, and in such a case we shouldn't emit a cascading error. + // + if( auto baseErrorType = as(baseType) ) + { + continue; + } + + auto baseDeclRefType = as(baseType); + if( !baseDeclRefType ) + { + getSink()->diagnose(inheritanceDecl, Diagnostics::baseOfStructMustBeStructOrInterface, decl, baseType); + continue; + } + + auto baseDeclRef = baseDeclRefType->declRef; + if( auto baseInterfaceDeclRef = baseDeclRef.as() ) + { + } + else if( auto baseStructDeclRef = baseDeclRef.as() ) + { + // To simplify the task of reading and maintaining code, + // we require that when a `struct` inherits from another + // `struct`, the base `struct` is the first item in + // the list of bases (before any interfaces). + // + // This constraint also has the secondary effect of restricting + // it so that a `struct` cannot multiply inherit from other + // `struct` types. + // + if( inheritanceClauseIndex != 0 ) + { + getSink()->diagnose(inheritanceDecl, Diagnostics::baseStructMustBeListedFirst, decl, baseType); + } + } + else + { + getSink()->diagnose(inheritanceDecl, Diagnostics::baseOfStructMustBeStructOrInterface, decl, baseType); + continue; + } + + // TODO: At this point we have the `baseDeclRef` + // and could use it to perform further validity checks, + // and/or to build up a more refined representation of + // the inheritance graph for this type (e.g., a "class + // precedence list"). + // + // E.g., we can/should check that we aren't introducing + // a circular inheritance relationship. + + _validateCrossModuleInheritance(decl, inheritanceDecl); } } @@ -1795,46 +1985,74 @@ namespace Slang void SemanticsDeclBasesVisitor::visitEnumDecl(EnumDecl* decl) { - // Look at inheritance clauses, and - // see if one of them is making the enum - // "inherit" from a concrete type. - // This will become the "tag" type - // of the enum. + // An `enum` type can inherit from interfaces, and also + // from a single "tag" type that must: + // + // * be a built-in integer type + // * come first in the list of base types + // + Index inheritanceClauseCounter = 0; RefPtr tagType; InheritanceDecl* tagTypeInheritanceDecl = nullptr; for(auto inheritanceDecl : decl->getMembersOfType()) { + Index inheritanceClauseIndex = inheritanceClauseCounter++; + ensureDecl(inheritanceDecl, DeclCheckState::CanUseBaseOfInheritanceDecl); + auto baseType = inheritanceDecl->base.type; - // Look at the type being inherited from. - auto superType = inheritanceDecl->base.type; + // It is possible that there was an error in checking the base type + // expression, and in such a case we shouldn't emit a cascading error. + // + if( auto baseErrorType = as(baseType) ) + { + continue; + } - if(auto errorType = as(superType)) + auto baseDeclRefType = as(baseType); + if( !baseDeclRefType ) { - // Ignore any erroneous inheritance clauses. + getSink()->diagnose(inheritanceDecl, Diagnostics::baseOfEnumMustBeIntegerOrInterface, decl, baseType); continue; } - else if(auto declRefType = as(superType)) + + auto baseDeclRef = baseDeclRefType->declRef; + if( auto baseInterfaceDeclRef = baseDeclRef.as() ) + { + _validateCrossModuleInheritance(decl, inheritanceDecl); + } + else if( auto baseStructDeclRef = baseDeclRef.as() ) { - if(auto interfaceDeclRef = declRefType->declRef.as()) + // To simplify the task of reading and maintaining code, + // we require that when an `enum` declares an explicit + // underlying tag type using an inheritance clause, that + // type must be the first item in the list of bases. + // + // This constraint also has the secondary effect of restricting + // it so that an `enum` can't possibly have multiple tag + // types declared. + // + if( inheritanceClauseIndex != 0 ) { - // Don't consider interface bases as candidates for - // the tag type. - continue; + getSink()->diagnose(inheritanceDecl, Diagnostics::tagTypeMustBeListedFirst, decl, baseType); + } + else + { + tagType = baseType; + tagTypeInheritanceDecl = inheritanceDecl; } - } - if(tagType) - { - // We already found a tag type. - getSink()->diagnose(inheritanceDecl, Diagnostics::enumTypeAlreadyHasTagType); - getSink()->diagnose(tagTypeInheritanceDecl, Diagnostics::seePreviousTagType); - break; + // Note: we do *not* apply the code that validates + // cross-module inheritance to a base that represnts + // a tag type, because declaring a tag type for an + // `enum` doesn't actually make it into a subtype + // of the tag type, and thus doesn't violate the + // rules when the tag type is `sealed`. } else { - tagType = superType; - tagTypeInheritanceDecl = inheritanceDecl; + getSink()->diagnose(inheritanceDecl, Diagnostics::baseOfEnumMustBeIntegerOrInterface, decl, baseType); + continue; } } @@ -1845,6 +2063,7 @@ namespace Slang // `enum` types that have a "raw representation" like this from // ones that are purely abstract and don't expose their // type of their tag. + // if(!tagType) { tagType = m_astBuilder->getIntType(); @@ -2852,10 +3071,8 @@ namespace Slang } } - void SemanticsDeclBasesVisitor::visitExtensionDecl(ExtensionDecl* decl) + void SemanticsDeclBasesVisitor::_validateExtensionDeclTargetType(ExtensionDecl* decl) { - decl->targetType = CheckProperType(decl->targetType); - if (auto targetDeclRefType = as(decl->targetType)) { // Attach our extension to that type as a candidate... @@ -2867,7 +3084,65 @@ namespace Slang return; } } - getSink()->diagnose(decl->targetType.exp, Diagnostics::unimplemented, "expected a nominal type here"); + getSink()->diagnose(decl->targetType.exp, Diagnostics::unimplemented, "an 'extension' can only extend a nominal type"); + } + + void SemanticsDeclBasesVisitor::visitExtensionDecl(ExtensionDecl* decl) + { + // We check the target type expression, and then validate + // that the type it names is one that it makes sense + // to extend. + // + decl->targetType = CheckProperType(decl->targetType); + _validateExtensionDeclTargetType(decl); + + for( auto inheritanceDecl : decl->getMembersOfType() ) + { + ensureDecl(inheritanceDecl, DeclCheckState::CanUseBaseOfInheritanceDecl); + auto baseType = inheritanceDecl->base.type; + + // It is possible that there was an error in checking the base type + // expression, and in such a case we shouldn't emit a cascading error. + // + if( auto baseErrorType = as(baseType) ) + { + continue; + } + + // An `extension` can only introduce inheritance from `interface` types. + // + // TODO: It might in theory make sense to allow an `extension` to + // introduce a non-`interface` base if we decide that an `extension` + // within the same module as the type it extends counts as just + // a continuation of the type's body (like a `partial class` in C#). + // + auto baseDeclRefType = as(baseType); + if( !baseDeclRefType ) + { + getSink()->diagnose(inheritanceDecl, Diagnostics::baseOfExtensionMustBeInterface, decl, baseType); + continue; + } + + auto baseDeclRef = baseDeclRefType->declRef; + auto baseInterfaceDeclRef = baseDeclRef.as(); + if( !baseInterfaceDeclRef ) + { + getSink()->diagnose(inheritanceDecl, Diagnostics::baseOfExtensionMustBeInterface, decl, baseType); + continue; + } + + // TODO: At this point we have the `baseInterfaceDeclRef` + // and could use it to perform further validity checks, + // and/or to build up a more refined representation of + // the inheritance graph for this extension (e.g., a "class + // precedence list"). + // + // E.g., we can/should check that we aren't introducing + // an inheritance relationship that already existed + // on the type as originally declared. + + _validateCrossModuleInheritance(decl, inheritanceDecl); + } } RefPtr SemanticsVisitor::calcThisType(DeclRef declRef) diff --git a/source/slang/slang-check-expr.cpp b/source/slang/slang-check-expr.cpp index d8eef571e..c0beb8262 100644 --- a/source/slang/slang-check-expr.cpp +++ b/source/slang/slang-check-expr.cpp @@ -325,10 +325,17 @@ namespace Slang { // TODO: do we need to make something more // explicit here? - bb = ConstructDeclRefExpr( + auto expr = ConstructDeclRefExpr( breadcrumb->declRef, bb, loc); + + if(bb && bb->type.isLeftValue) + { + expr->type.isLeftValue = true; + } + + bb = expr; } break; diff --git a/source/slang/slang-check-impl.h b/source/slang/slang-check-impl.h index 42edb5df7..c316dd820 100644 --- a/source/slang/slang-check-impl.h +++ b/source/slang/slang-check-impl.h @@ -645,7 +645,7 @@ namespace Slang /// which packages up the value, its type, and the witness /// of its conformance to the interface. /// - RefPtr createCastToInterfaceExpr( + RefPtr createCastToSuperTypeExpr( RefPtr toType, RefPtr fromExpr, RefPtr witness); @@ -924,9 +924,15 @@ namespace Slang RefPtr createSimpleSubtypeWitness( TypeWitnessBreadcrumb* breadcrumb); + /// Create a withness that `subType` is a sub-type of `superTypeDeclRef`. + /// + /// The `inBreadcrumbs` parameter represents a linked list of steps + /// in the process that validated the sub-type relationship, which + /// will be used to inform the construction of the witness. + /// RefPtr createTypeWitness( - RefPtr type, - DeclRef interfaceDeclRef, + RefPtr subType, + DeclRef superTypeDeclRef, TypeWitnessBreadcrumb* inBreadcrumbs); /// Is the given interface one that a tagged-union type can conform to? @@ -950,20 +956,48 @@ namespace Slang DeclRef interfaceDeclRef, DeclRef requirementDeclRef); - bool doesTypeConformToInterfaceImpl( - RefPtr originalType, - RefPtr type, - DeclRef interfaceDeclRef, + /// Check whether `subType` is declared a sub-type of `superTypeDeclRef` + /// + /// If this function returns `true` (because the subtype relationship holds), + /// then `outWitness` will be set to a value that serves as a witness + /// to the subtype relationship. + /// + /// This function may be used to validate a transitive subtype relationship + /// where, e.g., `A : C` becase `A : B` and `B : C`. In such a case, a recursive + /// call to `_isDeclaredSubtype` may occur where `originalSubType` is `A`, + /// `subType` is `C`, and `superTypeDeclRef` is `C`. The `inBreadcrumbs` in that + /// case would include information for the `A : B` relationship, which can be + /// used to construct a witness for `A : C` from the `A : B` and `B : C` witnesses. + /// + bool _isDeclaredSubtype( + RefPtr originalSubType, + RefPtr subType, + DeclRef superTypeDeclRef, RefPtr* outWitness, TypeWitnessBreadcrumb* inBreadcrumbs); - bool DoesTypeConformToInterface( - RefPtr type, - DeclRef interfaceDeclRef); + /// Check whether `subType` is a sub-type of `superTypeDeclRef`. + bool isDeclaredSubtype( + RefPtr subType, + DeclRef superTypeDeclRef); + + /// Check whether `subType` is a sub-type of `superTypeDeclRef`, + /// and return a witness to the sub-type relationship if it holds + /// (return null otherwise). + /// + RefPtr tryGetSubtypeWitness( + RefPtr subType, + DeclRef superTypeDeclRef); + /// Check whether `type` conforms to `interfaceDeclRef`, + /// and return a witness to the conformance if it holds + /// (return null otherwise). + /// + /// This function is equivalent to `tryGetSubtypeWitness()`. + /// RefPtr tryGetInterfaceConformanceWitness( - RefPtr type, - DeclRef interfaceDeclRef); + RefPtr type, + DeclRef interfaceDeclRef); /// Does there exist an implicit conversion from `fromType` to `toType`? bool canConvertImplicitly( @@ -1381,7 +1415,7 @@ namespace Slang CASE(OverloadedExpr) CASE(OverloadedExpr2) CASE(AggTypeCtorExpr) - CASE(CastToInterfaceExpr) + CASE(CastToSuperTypeExpr) CASE(LetExpr) CASE(ExtractExistentialValueExpr) diff --git a/source/slang/slang-diagnostic-defs.h b/source/slang/slang-diagnostic-defs.h index 6e77aa45d..41c674ddb 100644 --- a/source/slang/slang-diagnostic-defs.h +++ b/source/slang/slang-diagnostic-defs.h @@ -289,8 +289,6 @@ DIAGNOSTIC(31120, Error, invalidAttributeTarget, "invalid syntax target for user // Enums DIAGNOSTIC(32000, Error, invalidEnumTagType, "invalid tag type for 'enum': '$0'") -DIAGNOSTIC(32001, Error, enumTypeAlreadyHasTagType, "'enum' type has already declared a tag type") -DIAGNOSTIC(32002, Note, seePreviousTagType, "see previous tag type declaration") DIAGNOSTIC(32003, Error, unexpectedEnumTagExpr, "unexpected form for 'enum' tag value expression") @@ -320,6 +318,18 @@ DIAGNOSTIC(30610, Error, ambiguousDefaultInitializerForType, "more than one defa // 307xx: parameters DIAGNOSTIC(30700, Error, outputParameterCannotHaveDefaultValue, "an 'out' or 'inout' parameter cannot have a default-value expression"); +// 308xx: inheritance +DIAGNOSTIC(30810, Error, baseOfInterfaceMustBeInterface, "interface '$0' cannot inherit from non-interface type '$1'") +DIAGNOSTIC(30811, Error, baseOfStructMustBeStructOrInterface, "struct '$0' cannot inherit from type '$1' that is neither a struct nor an interface") +DIAGNOSTIC(30812, Error, baseOfEnumMustBeIntegerOrInterface, "enum '$0' cannot inherit from type '$1' that is neither an interface not a builtin integer type") +DIAGNOSTIC(30810, Error, baseOfExtensionMustBeInterface, "extension cannot inherit from non-interface type '$1'") + +DIAGNOSTIC(30820, Error, baseStructMustBeListedFirst, "a struct type may only inherit from one other struct type, and that type must appear first in the list of bases") +DIAGNOSTIC(30821, Error, tagTypeMustBeListedFirst, "an unum type may only have a single tag type, and that type must be listed first in the list of bases") + +DIAGNOSTIC(30820, Error, cannotInheritFromExplicitlySealedDeclarationInAnotherModule, "cannot inherit from type '$0' marked 'sealed' in module '$1'") +DIAGNOSTIC(30821, Error, cannotInheritFromImplicitlySealedDeclarationInAnotherModule, "cannot inherit from type '$0' in module '$1' because it is implicitly 'sealed'; mark the base type 'open' to allow inheritance across modules") + // 39999 waiting to be placed in the right range DIAGNOSTIC(39999, Error, expectedIntegerConstantWrongType, "expected integer constant (found: '$0')") @@ -343,8 +353,6 @@ DIAGNOSTIC(39999, Error, expectedAGeneric, "expected a generic when using '<...> DIAGNOSTIC(39999, Error, genericArgumentInferenceFailed, "could not specialize generic for arguments of type $0") DIAGNOSTIC(39999, Note, genericSignatureTried, "see declaration of $0") -DIAGNOSTIC(39999, Error, expectedAnInterfaceGot, "expected an interface, got '$0'") - DIAGNOSTIC(39999, Error, ambiguousReference, "ambiguous reference to '$0'"); DIAGNOSTIC(39999, Error, ambiguousExpression, "ambiguous reference"); diff --git a/source/slang/slang-lower-to-ir.cpp b/source/slang/slang-lower-to-ir.cpp index f369729d2..c52025244 100644 --- a/source/slang/slang-lower-to-ir.cpp +++ b/source/slang/slang-lower-to-ir.cpp @@ -701,17 +701,17 @@ LoweredValInfo emitCallToDeclRef( } IRInst* getFieldKey( - IRGenContext* context, - DeclRef field) + IRGenContext* context, + DeclRef field) { return getSimpleVal(context, emitDeclRef(context, field, context->irBuilder->getKeyType())); } LoweredValInfo extractField( - IRGenContext* context, - IRType* fieldType, - LoweredValInfo base, - DeclRef field) + IRGenContext* context, + IRType* fieldType, + LoweredValInfo base, + DeclRef field) { IRBuilder* builder = context->irBuilder; @@ -2054,6 +2054,21 @@ struct ExprLoweringVisitorBase : ExprVisitor } else if(auto constraintDeclRef = declRef.as()) { + auto superType = getSup(getASTBuilder(), constraintDeclRef); + if(auto superDeclRefType = as(superType)) + { + if(auto superStructDeclRef = superDeclRefType->declRef.template as()) + { + // The constraint is saying that the given type inherits + // from a concrete `struct` type, which means it should + // be satisfied by a witness that represents a field + // (TODO: or a chain of fields) to fetch to get the + // final value. + // + return extractField(loweredType, loweredBase, constraintDeclRef); + } + } + // The code is making use of a "witness" that a value of // some generic type conforms to an interface. // @@ -2748,30 +2763,87 @@ struct ExprLoweringVisitorBase : ExprVisitor UNREACHABLE_RETURN(LoweredValInfo()); } - LoweredValInfo visitCastToInterfaceExpr( - CastToInterfaceExpr* expr) + /// Emit code to cast `value` to a concrete `superType` (e.g., a `struct`). + /// + /// The `subTypeWitness` is expected to witness the sub-type relationship + /// by naming a field (or chain of fields) that leads from the type of + /// `value` to the field that stores its members for `superType`. + /// + LoweredValInfo emitCastToConcreteSuperTypeRec( + LoweredValInfo const& value, + IRType* superType, + Val* subTypeWitness) { - // We have an expression that is "up-casting" some concrete value - // to an existential type (aka interface type), using a subtype witness - // (which will lower as a witness table) to show that the conversion - // is valid. - // - // At the IR level, this will become a `makeExistential` instruction, - // which collects the above information into a single IR-level value. - // A dynamic CPU implementation of Slang might encode an existential - // as a "fat pointer" representation, which includes a pointer to - // data for the concrete value, plus a pointer to the witness table. + if( auto declaredSubtypeWitness = as(subTypeWitness) ) + { + return extractField(superType, value, declaredSubtypeWitness->declRef); + } + else + { + SLANG_ASSERT(!"unhandled"); + return nullptr; + } + } + + LoweredValInfo visitCastToSuperTypeExpr( + CastToSuperTypeExpr* expr) + { + auto superType = lowerType(context, expr->type); + auto value = lowerRValueExpr(context, expr->valueArg); + + // The actual operation that we need to perform here + // depends on the kind of subtype relationship we + // are making use of. // - // Note: if/when Slang supports more general existential types, such - // as compositions of interface (e.g., `IReadable & IWritable`), then - // we should probably extend the AST and IR mechanism here to accept - // a sequence of witness tables. + // The first important case is when the super type is + // an interface type, such that casting from a concrete + // value to that type creates a value of existential + // type that binds together the concrete value and the + // witness table that represents the subtype relationship. // - auto existentialType = lowerType(context, expr->type); - auto concreteValue = getSimpleVal(context, lowerRValueExpr(context, expr->valueArg)); - auto witnessTable = lowerSimpleVal(context, expr->witnessArg); - auto existentialValue = getBuilder()->emitMakeExistential(existentialType, concreteValue, witnessTable); - return LoweredValInfo::simple(existentialValue); + if( auto declRefType = as(expr->type) ) + { + auto declRef = declRefType->declRef; + if( auto interfaceDeclRef = declRef.as() ) + { + // We have an expression that is "up-casting" some concrete value + // to an existential type (aka interface type), using a subtype witness + // (which will lower as a witness table) to show that the conversion + // is valid. + // + auto witnessTable = lowerSimpleVal(context, expr->witnessArg); + + // At the IR level, this will become a `makeExistential` instruction, + // which collects the above information into a single IR-level value. + // A dynamic CPU implementation of Slang might encode an existential + // as a "fat pointer" representation, which includes a pointer to + // data for the concrete value, plus a pointer to the witness table. + // + // Note: if/when Slang supports more general existential types, such + // as compositions of interface (e.g., `IReadable & IWritable`), then + // we should probably extend the AST and IR mechanism here to accept + // a sequence of witness tables. + // + auto concreteValue = getSimpleVal(context, value); + auto existentialValue = getBuilder()->emitMakeExistential( + superType, + concreteValue, + witnessTable); + return LoweredValInfo::simple(existentialValue); + } + else if( auto structDeclRef = declRef.as() ) + { + // We are up-casting to a concrete `struct` super-type, + // such that the witness will represent a field of the super-type + // that is stored in instances of the sub-type (or a chain + // of such fields for a transitive witness). + // + return emitCastToConcreteSuperTypeRec(value, superType, expr->witnessArg); + } + } + + SLANG_UNEXPECTED("unexpected case of subtype relationship"); + UNREACHABLE_RETURN(LoweredValInfo()); } LoweredValInfo subscriptValue( @@ -2815,9 +2887,9 @@ struct ExprLoweringVisitorBase : ExprVisitor } LoweredValInfo extractField( - IRType* fieldType, - LoweredValInfo base, - DeclRef field) + IRType* fieldType, + LoweredValInfo base, + DeclRef field) { return Slang::extractField(context, fieldType, base, field); } @@ -4524,6 +4596,21 @@ struct DeclLoweringVisitor : DeclVisitor // What is the super-type that we have declared we inherit from? RefPtr superType = inheritanceDecl->base.type; + if(auto superDeclRefType = as(superType)) + { + if( auto superStructDeclRef = superDeclRefType->declRef.as() ) + { + // TODO: the witness that a type inherits from a `struct` + // type should probably be a key that will be used for + // a field that holds the base type... + // + auto irKey = getBuilder()->createStructKey(); + auto keyVal = LoweredValInfo::simple(irKey); + setGlobalValue(context, inheritanceDecl, keyVal); + return keyVal; + } + } + // Construct the mangled name for the witness table, which depends // on the type that is conforming, and the type that it conforms to. // @@ -5270,6 +5357,27 @@ struct DeclLoweringVisitor : DeclVisitor subBuilder->setInsertInto(irStruct); + // A `struct` that inherits from another `struct` must start + // with a member for the direct base type. + // + for( auto inheritanceDecl : decl->getMembersOfType() ) + { + auto superType = inheritanceDecl->base; + if(auto superDeclRefType = as(superType)) + { + if(auto superStructDeclRef = superDeclRefType->declRef.as()) + { + auto superKey = (IRStructKey*) getSimpleVal(context, ensureDecl(context, inheritanceDecl)); + auto irSuperType = lowerType(context, superType.type); + subBuilder->createStructField( + irStruct, + superKey, + irSuperType); + } + } + } + + for (auto fieldDecl : decl->getMembersOfType()) { if (fieldDecl->hasModifier()) -- cgit v1.2.3 From 4e2f2771eb2f8991014d957848a6a25aa49c0aaf Mon Sep 17 00:00:00 2001 From: Yong He Date: Thu, 4 Jun 2020 12:42:17 -0700 Subject: Emit [loop] attribute to output HLSL. --- source/slang/slang-emit-hlsl.cpp | 9 +++- source/slang/slang-ir-insts.h | 1 + source/slang/slang-lower-to-ir.cpp | 4 ++ tests/cross-compile/loop-attribs.slang | 19 ++++++++ tests/cross-compile/loop-attribs.slang.expected | 61 +++++++++++++++++++++++++ 5 files changed, 93 insertions(+), 1 deletion(-) create mode 100644 tests/cross-compile/loop-attribs.slang create mode 100644 tests/cross-compile/loop-attribs.slang.expected (limited to 'source') diff --git a/source/slang/slang-emit-hlsl.cpp b/source/slang/slang-emit-hlsl.cpp index 5ebd7e9fc..b15680f29 100644 --- a/source/slang/slang-emit-hlsl.cpp +++ b/source/slang/slang-emit-hlsl.cpp @@ -673,9 +673,16 @@ void HLSLSourceEmitter::emitVectorTypeNameImpl(IRType* elementType, IRIntegerVal void HLSLSourceEmitter::emitLoopControlDecorationImpl(IRLoopControlDecoration* decl) { - if (decl->getMode() == kIRLoopControl_Unroll) + switch (decl->getMode()) { + case kIRLoopControl_Unroll: m_writer->emit("[unroll]\n"); + break; + case kIRLoopControl_Loop: + m_writer->emit("[loop]\n"); + break; + default: + SLANG_UNREACHABLE("emit loop control decoration"); } } diff --git a/source/slang/slang-ir-insts.h b/source/slang/slang-ir-insts.h index 06a5c1412..0ae05311b 100644 --- a/source/slang/slang-ir-insts.h +++ b/source/slang/slang-ir-insts.h @@ -42,6 +42,7 @@ struct IRHighLevelDeclDecoration : IRDecoration enum IRLoopControl { kIRLoopControl_Unroll, + kIRLoopControl_Loop, }; struct IRLoopControlDecoration : IRDecoration diff --git a/source/slang/slang-lower-to-ir.cpp b/source/slang/slang-lower-to-ir.cpp index f369729d2..030c5fc74 100644 --- a/source/slang/slang-lower-to-ir.cpp +++ b/source/slang/slang-lower-to-ir.cpp @@ -3281,6 +3281,10 @@ struct StmtLoweringVisitor : StmtVisitor { getBuilder()->addLoopControlDecoration(inst, kIRLoopControl_Unroll); } + else if( stmt->findModifier() ) + { + getBuilder()->addLoopControlDecoration(inst, kIRLoopControl_Loop); + } // TODO: handle other cases here } diff --git a/tests/cross-compile/loop-attribs.slang b/tests/cross-compile/loop-attribs.slang new file mode 100644 index 000000000..ddedb67df --- /dev/null +++ b/tests/cross-compile/loop-attribs.slang @@ -0,0 +1,19 @@ +// loop-attribs.slang +// Test that loop attributes are correctly emitted to the resulting HLSL. + +//TEST:SIMPLE:-target hlsl -entry main -stage fragment -profile sm_6_0 + +float4 main() : SV_Target +{ + float sum = 0.0f; + + [loop] + for (int i = 0; i < 100; i++) + sum += float(i); + + [unroll(10)] + for (int j = 0; j < 100; j++) + sum += float(j); + + return float4(sum, 0, 0, 0); +} \ No newline at end of file diff --git a/tests/cross-compile/loop-attribs.slang.expected b/tests/cross-compile/loop-attribs.slang.expected new file mode 100644 index 000000000..0ecf0194f --- /dev/null +++ b/tests/cross-compile/loop-attribs.slang.expected @@ -0,0 +1,61 @@ +result code = 0 +standard error = { +} +standard output = { +#pragma pack_matrix(column_major) + +#line 6 "tests/cross-compile/loop-attribs.slang" +vector main() : SV_TARGET +{ + int i_0; + float sum_0; + int j_0; + float sum_1; + i_0 = int(0); + sum_0 = 0.00000000000000000000; + [loop] + for(;;) + { + +#line 11 + if(i_0 < int(100)) + { + } + else + { + break; + } + float _S1 = sum_0 + (float) i_0; + +#line 11 + int _S2 = i_0 + (int) int(1); + i_0 = _S2; + sum_0 = _S1; + } + j_0 = int(0); + sum_1 = sum_0; + [unroll] + for(;;) + { + +#line 15 + if(j_0 < int(100)) + { + } + else + { + break; + } + float _S3 = sum_1 + (float) j_0; + +#line 15 + int _S4 = j_0 + (int) int(1); + j_0 = _S4; + sum_1 = _S3; + } + +#line 18 + return vector(sum_1, (float) int(0), (float) int(0), (float) int(0)); +} + +} -- cgit v1.2.3 From 2c30d36c3df77d832443087f83940bf8f4f039f0 Mon Sep 17 00:00:00 2001 From: Yong He Date: Thu, 4 Jun 2020 14:06:10 -0700 Subject: Remove aborting in emitLoopControlDecoration default case. --- source/slang/slang-emit-hlsl.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'source') diff --git a/source/slang/slang-emit-hlsl.cpp b/source/slang/slang-emit-hlsl.cpp index b15680f29..83bfb8f2a 100644 --- a/source/slang/slang-emit-hlsl.cpp +++ b/source/slang/slang-emit-hlsl.cpp @@ -682,7 +682,7 @@ void HLSLSourceEmitter::emitLoopControlDecorationImpl(IRLoopControlDecoration* d m_writer->emit("[loop]\n"); break; default: - SLANG_UNREACHABLE("emit loop control decoration"); + break; } } -- cgit v1.2.3 From ecac0c76a93bd123d5a5a86473716ad166f8c5d6 Mon Sep 17 00:00:00 2001 From: jsmall-nvidia Date: Thu, 4 Jun 2020 17:15:38 -0400 Subject: Add a ASTBuilder to a Module (#1369) Only construct on valid ASTBuilder (was being called on nullptr on occassion) --- source/slang/slang-ast-builder.cpp | 12 +++++++----- source/slang/slang-compiler.h | 7 +++++++ source/slang/slang.cpp | 24 +++++++++++++++++------- 3 files changed, 31 insertions(+), 12 deletions(-) (limited to 'source') diff --git a/source/slang/slang-ast-builder.cpp b/source/slang/slang-ast-builder.cpp index 316390353..67ce80120 100644 --- a/source/slang/slang-ast-builder.cpp +++ b/source/slang/slang-ast-builder.cpp @@ -21,18 +21,20 @@ void SharedASTBuilder::init(Session* session) m_session = session; // We just want as a place to store allocations of shared types - RefPtr astBuilder(new ASTBuilder); - - astBuilder->m_sharedASTBuilder = this; + { + RefPtr astBuilder(new ASTBuilder); + astBuilder->m_sharedASTBuilder = this; + m_astBuilder = astBuilder.detach(); + } + // Clear the built in types memset(m_builtinTypes, 0, sizeof(m_builtinTypes)); + // Create common shared types m_errorType = m_astBuilder->create(); m_initializerListType = m_astBuilder->create(); m_overloadedType = m_astBuilder->create(); - m_astBuilder = astBuilder.detach(); - // We can just iterate over the class pointers. // NOTE! That this adds the names of the abstract classes too(!) for (Index i = 0; i < Index(ASTNodeType::CountOf); ++i) diff --git a/source/slang/slang-compiler.h b/source/slang/slang-compiler.h index 2c23a89bc..d5de51ac4 100644 --- a/source/slang/slang-compiler.h +++ b/source/slang/slang-compiler.h @@ -970,6 +970,9 @@ namespace Slang List const& getModuleDependencies() SLANG_OVERRIDE { return m_moduleDependencyList.getModuleList(); } List const& getFilePathDependencies() SLANG_OVERRIDE { return m_filePathDependencyList.getFilePathList(); } + /// Get the ASTBuilder + ASTBuilder* getASTBuilder() { return &m_astBuilder; } + /// Collect information on the shader parameters of the module. /// /// This method should only be called once, after the core @@ -1044,6 +1047,10 @@ namespace Slang // that was defined as part of the module. // List> m_entryPoints; + + // The builder that owns all of the AST nodes from parsing the source of + // this module. + ASTBuilder m_astBuilder; }; typedef Module LoadedModule; diff --git a/source/slang/slang.cpp b/source/slang/slang.cpp index 793b81ebf..331277bf3 100644 --- a/source/slang/slang.cpp +++ b/source/slang/slang.cpp @@ -124,13 +124,17 @@ void Session::init() // Use to create a ASTBuilder RefPtr builtinAstBuilder(new ASTBuilder(m_sharedASTBuilder)); - + + // And the global ASTBuilder + globalAstBuilder = new ASTBuilder(m_sharedASTBuilder); + // Make sure our source manager is initialized builtinSourceManager.initialize(nullptr, nullptr); - + // Built in linkage uses the built in builder m_builtinLinkage = new Linkage(this, builtinAstBuilder); + // Because the `Session` retains the builtin `Linkage`, // we need to make sure that the parent pointer inside // `Linkage` doesn't create a retain cycle. @@ -952,7 +956,10 @@ void FrontEndCompileRequest::parseTranslationUnit( combinedPreprocessorDefinitions.Add(def.Key, def.Value); auto module = translationUnit->getModule(); - RefPtr translationUnitSyntax = linkage->getASTBuilder()->create(); + + ASTBuilder* astBuilder = module->getASTBuilder(); + + RefPtr translationUnitSyntax = astBuilder->create(); translationUnitSyntax->nameAndLoc.name = translationUnit->moduleName; translationUnitSyntax->module = module; module->setModuleDecl(translationUnitSyntax); @@ -966,13 +973,13 @@ void FrontEndCompileRequest::parseTranslationUnit( // // We are adding the marker here, before we even parse the // code in the module, in case the subsequent steps would - // like to treat the standard library differently. Alternatiely + // like to treat the standard library differently. Alternatively // we could pass down the `m_isStandardLibraryCode` flag to // these passes. // if( m_isStandardLibraryCode ) { - translationUnitSyntax->modifiers.first = linkage->getASTBuilder()->create(); + translationUnitSyntax->modifiers.first = astBuilder->create(); } for (auto sourceFile : translationUnit->getSourceFiles()) @@ -986,7 +993,7 @@ void FrontEndCompileRequest::parseTranslationUnit( module); parseSourceFile( - linkage->getASTBuilder(), + astBuilder, translationUnit, tokens, getSink(), @@ -1054,7 +1061,9 @@ void FrontEndCompileRequest::generateIR() // * it can dump ir // * it can generate diagnostics - /// Generate IR for translation unit + /// Generate IR for translation unit. + /// TODO(JS): Use the linkage ASTBuilder, because it seems possible that cross module constructs are possible in + /// ir lowering. RefPtr irModule(generateIRForTranslationUnit(getLinkage()->getASTBuilder(), translationUnit)); if (verifyDebugSerialization) @@ -1771,6 +1780,7 @@ void FilePathDependencyList::addDependency(Module* module) Module::Module(Linkage* linkage) : ComponentType(linkage) + , m_astBuilder(linkage->getASTBuilder()->getSharedASTBuilder()) { addModuleDependency(this); } -- cgit v1.2.3 From 3bb780724830ae830657a47e4eba008a4c0f4ff7 Mon Sep 17 00:00:00 2001 From: Tim Foley Date: Fri, 5 Jun 2020 10:00:15 -0700 Subject: Fixes for active mask synthesis + tests (#1370) * Fixes for active mask synthesis + tests There are two fixes here: * The code generation that follows active mask synthesis was requiring CUDA SM architecture version 7.0 for one of the introduced instructions, but not all of them. This change centralizes the handling of upgrading the required CUDA SM architecture version, and makes sure that the instructions introduced by active mask synthesis request version 7.0. * The tests for active mask synthesis were not flagged as requiring the `cuda_sm_7_0` feature when invoking `render-test-tool`, which meant they run but produce unexpected results when invoked on a GPU without the required semantics for functions like `__ballot_sync()`. This change adds the missing `-render-feature cuda_sm_7_0` to those tests. * fixup: mark more tests that rely on implicit active mask --- source/slang/slang-emit-cuda.cpp | 23 +++++++++++----------- source/slang/slang-emit-cuda.h | 3 +++ tests/hlsl-intrinsic/active-mask/for-break.slang | 2 +- .../active-mask/for-continue-ext.slang | 2 +- .../hlsl-intrinsic/active-mask/for-continue.slang | 2 +- tests/hlsl-intrinsic/active-mask/for.slang | 2 +- .../active-mask/if-conditional-exit.slang | 2 +- .../hlsl-intrinsic/active-mask/if-early-exit.slang | 2 +- .../hlsl-intrinsic/active-mask/if-one-sided.slang | 2 +- tests/hlsl-intrinsic/active-mask/if.slang | 2 +- .../active-mask/switch-no-default.slang | 2 +- .../active-mask/switch-trivial-fallthrough.slang | 2 +- tests/hlsl-intrinsic/active-mask/switch.slang | 7 ++++++- tests/hlsl-intrinsic/wave-active-product.slang | 2 +- tests/hlsl-intrinsic/wave-broadcast-lane-at.slang | 2 +- tests/hlsl-intrinsic/wave-diverge.slang | 2 +- tests/hlsl-intrinsic/wave-is-first-lane.slang | 2 +- tests/hlsl-intrinsic/wave-matrix.slang | 2 +- tests/hlsl-intrinsic/wave-prefix-product.slang | 2 +- tests/hlsl-intrinsic/wave-prefix-sum.slang | 2 +- tests/hlsl-intrinsic/wave-read-lane-at.slang | 2 +- tests/hlsl-intrinsic/wave-shuffle-vk.slang | 2 +- tests/hlsl-intrinsic/wave-shuffle.slang | 2 +- tests/hlsl-intrinsic/wave-vector.slang | 2 +- tests/hlsl-intrinsic/wave.slang | 2 +- 25 files changed, 43 insertions(+), 34 deletions(-) (limited to 'source') diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp index 639e7f737..25b06027d 100644 --- a/source/slang/slang-emit-cuda.cpp +++ b/source/slang/slang-emit-cuda.cpp @@ -520,6 +520,8 @@ bool CUDASourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu } case kIROp_WaveMaskBallot: { + _requireCUDASMVersion(SemanticVersion(7, 0)); + m_writer->emit("__ballot_sync("); emitOperand(inst->getOperand(0), getInfo(EmitOp::General)); m_writer->emit(", "); @@ -529,12 +531,7 @@ bool CUDASourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu } case kIROp_WaveMaskMatch: { - SemanticVersion version; - version.set(7, 0); - if (version > m_extensionTracker->m_smVersion) - { - m_extensionTracker->m_smVersion = version; - } + _requireCUDASMVersion(SemanticVersion(7, 0)); m_writer->emit("__match_any_sync("); emitOperand(inst->getOperand(0), getInfo(EmitOp::General)); @@ -549,6 +546,14 @@ bool CUDASourceEmitter::tryEmitInstExprImpl(IRInst* inst, const EmitOpInfo& inOu return Super::tryEmitInstExprImpl(inst, inOuterPrec); } +void CUDASourceEmitter::_requireCUDASMVersion(SemanticVersion const& version) +{ + if (version > m_extensionTracker->m_smVersion) + { + m_extensionTracker->m_smVersion = version; + } +} + void CUDASourceEmitter::handleCallExprDecorationsImpl(IRInst* funcValue) { // Does this function declare any requirements on GLSL version or @@ -566,11 +571,7 @@ void CUDASourceEmitter::handleCallExprDecorationsImpl(IRInst* funcValue) { SemanticVersion version; version.setFromInteger(SemanticVersion::IntegerType(smDecoration->getCUDASMVersion())); - - if (version > m_extensionTracker->m_smVersion) - { - m_extensionTracker->m_smVersion = version; - } + _requireCUDASMVersion(version); } } } diff --git a/source/slang/slang-emit-cuda.h b/source/slang/slang-emit-cuda.h index 669bf2d20..c0a5ac5dc 100644 --- a/source/slang/slang-emit-cuda.h +++ b/source/slang/slang-emit-cuda.h @@ -79,6 +79,9 @@ protected: void _emitInitializerList(IRType* elementType, IRUse* operands, Index operandCount); void _emitInitializerListValue(IRType* elementType, IRInst* value); + /// Ensure that the generated code is compiled for at least CUDA SM `version` + void _requireCUDASMVersion(SemanticVersion const& version); + RefPtr m_extensionTracker; }; diff --git a/tests/hlsl-intrinsic/active-mask/for-break.slang b/tests/hlsl-intrinsic/active-mask/for-break.slang index d0c35a017..440e7481c 100644 --- a/tests/hlsl-intrinsic/active-mask/for-break.slang +++ b/tests/hlsl-intrinsic/active-mask/for-break.slang @@ -8,7 +8,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 -xslang -DHACK //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -xslang -DHACK -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer buffer; diff --git a/tests/hlsl-intrinsic/active-mask/for-continue-ext.slang b/tests/hlsl-intrinsic/active-mask/for-continue-ext.slang index 67ee536eb..881c0389e 100644 --- a/tests/hlsl-intrinsic/active-mask/for-continue-ext.slang +++ b/tests/hlsl-intrinsic/active-mask/for-continue-ext.slang @@ -9,7 +9,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 -xslang -DHACK //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -xslang -DHACK -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer buffer; diff --git a/tests/hlsl-intrinsic/active-mask/for-continue.slang b/tests/hlsl-intrinsic/active-mask/for-continue.slang index 6bb0f23fb..7700a937a 100644 --- a/tests/hlsl-intrinsic/active-mask/for-continue.slang +++ b/tests/hlsl-intrinsic/active-mask/for-continue.slang @@ -9,7 +9,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 -xslang -DHACK //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -xslang -DHACK -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer buffer; diff --git a/tests/hlsl-intrinsic/active-mask/for.slang b/tests/hlsl-intrinsic/active-mask/for.slang index aac66f3c9..c6fef5e58 100644 --- a/tests/hlsl-intrinsic/active-mask/for.slang +++ b/tests/hlsl-intrinsic/active-mask/for.slang @@ -8,7 +8,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer buffer; diff --git a/tests/hlsl-intrinsic/active-mask/if-conditional-exit.slang b/tests/hlsl-intrinsic/active-mask/if-conditional-exit.slang index 78815e942..9f90b9380 100644 --- a/tests/hlsl-intrinsic/active-mask/if-conditional-exit.slang +++ b/tests/hlsl-intrinsic/active-mask/if-conditional-exit.slang @@ -9,7 +9,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer diff --git a/tests/hlsl-intrinsic/active-mask/if-early-exit.slang b/tests/hlsl-intrinsic/active-mask/if-early-exit.slang index 8532035a3..35efb9a6d 100644 --- a/tests/hlsl-intrinsic/active-mask/if-early-exit.slang +++ b/tests/hlsl-intrinsic/active-mask/if-early-exit.slang @@ -8,7 +8,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer buffer; diff --git a/tests/hlsl-intrinsic/active-mask/if-one-sided.slang b/tests/hlsl-intrinsic/active-mask/if-one-sided.slang index 0662fc1e7..9b3259f3b 100644 --- a/tests/hlsl-intrinsic/active-mask/if-one-sided.slang +++ b/tests/hlsl-intrinsic/active-mask/if-one-sided.slang @@ -6,7 +6,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer buffer; diff --git a/tests/hlsl-intrinsic/active-mask/if.slang b/tests/hlsl-intrinsic/active-mask/if.slang index 897ba6d1b..077bb4b40 100644 --- a/tests/hlsl-intrinsic/active-mask/if.slang +++ b/tests/hlsl-intrinsic/active-mask/if.slang @@ -6,7 +6,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer buffer; diff --git a/tests/hlsl-intrinsic/active-mask/switch-no-default.slang b/tests/hlsl-intrinsic/active-mask/switch-no-default.slang index d103bc4c7..f3060631b 100644 --- a/tests/hlsl-intrinsic/active-mask/switch-no-default.slang +++ b/tests/hlsl-intrinsic/active-mask/switch-no-default.slang @@ -9,7 +9,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 -xslang -DHACK //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -xslang -DHACK -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer buffer; diff --git a/tests/hlsl-intrinsic/active-mask/switch-trivial-fallthrough.slang b/tests/hlsl-intrinsic/active-mask/switch-trivial-fallthrough.slang index ceb7d236d..b5674c8d1 100644 --- a/tests/hlsl-intrinsic/active-mask/switch-trivial-fallthrough.slang +++ b/tests/hlsl-intrinsic/active-mask/switch-trivial-fallthrough.slang @@ -13,7 +13,7 @@ // target because we do not synthesize the active // mask value we want/expect to see. // -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer buffer; diff --git a/tests/hlsl-intrinsic/active-mask/switch.slang b/tests/hlsl-intrinsic/active-mask/switch.slang index 41f379f3d..0ab0c7173 100644 --- a/tests/hlsl-intrinsic/active-mask/switch.slang +++ b/tests/hlsl-intrinsic/active-mask/switch.slang @@ -6,7 +6,12 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 -xslang -DHACK //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -xslang -DHACK -//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute + +// Note: this test is currently disabled on the CUDA +// target because we do not synthesize the active +// mask value we want/expect to see. +// +//DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name buffer RWStructuredBuffer buffer; diff --git a/tests/hlsl-intrinsic/wave-active-product.slang b/tests/hlsl-intrinsic/wave-active-product.slang index ca3fdcb77..67b9ca587 100644 --- a/tests/hlsl-intrinsic/wave-active-product.slang +++ b/tests/hlsl-intrinsic/wave-active-product.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-broadcast-lane-at.slang b/tests/hlsl-intrinsic/wave-broadcast-lane-at.slang index b6f5d3847..bb1523219 100644 --- a/tests/hlsl-intrinsic/wave-broadcast-lane-at.slang +++ b/tests/hlsl-intrinsic/wave-broadcast-lane-at.slang @@ -3,7 +3,7 @@ //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 // Disabled on VK because glsl can't do WaveReadLaneAt on matrix. //DISABLE_TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-diverge.slang b/tests/hlsl-intrinsic/wave-diverge.slang index ab83a1553..2d500d0c8 100644 --- a/tests/hlsl-intrinsic/wave-diverge.slang +++ b/tests/hlsl-intrinsic/wave-diverge.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-is-first-lane.slang b/tests/hlsl-intrinsic/wave-is-first-lane.slang index 39e7dab5d..e1f0faef3 100644 --- a/tests/hlsl-intrinsic/wave-is-first-lane.slang +++ b/tests/hlsl-intrinsic/wave-is-first-lane.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute, vulkan):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute, vulkan):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-matrix.slang b/tests/hlsl-intrinsic/wave-matrix.slang index b5af69f5d..50acc4a71 100644 --- a/tests/hlsl-intrinsic/wave-matrix.slang +++ b/tests/hlsl-intrinsic/wave-matrix.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //DISABLE_TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-prefix-product.slang b/tests/hlsl-intrinsic/wave-prefix-product.slang index a56912616..436ce06f7 100644 --- a/tests/hlsl-intrinsic/wave-prefix-product.slang +++ b/tests/hlsl-intrinsic/wave-prefix-product.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-prefix-sum.slang b/tests/hlsl-intrinsic/wave-prefix-sum.slang index 343a7afbd..358debedc 100644 --- a/tests/hlsl-intrinsic/wave-prefix-sum.slang +++ b/tests/hlsl-intrinsic/wave-prefix-sum.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-read-lane-at.slang b/tests/hlsl-intrinsic/wave-read-lane-at.slang index c3caaa4e8..4d840ad0f 100644 --- a/tests/hlsl-intrinsic/wave-read-lane-at.slang +++ b/tests/hlsl-intrinsic/wave-read-lane-at.slang @@ -3,7 +3,7 @@ //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 // Disabled on VK because glsl can't do WaveReadLaneAt on matrix. //DISABLE_TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-shuffle-vk.slang b/tests/hlsl-intrinsic/wave-shuffle-vk.slang index 75aa392ea..820a69e89 100644 --- a/tests/hlsl-intrinsic/wave-shuffle-vk.slang +++ b/tests/hlsl-intrinsic/wave-shuffle-vk.slang @@ -3,7 +3,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-shuffle.slang b/tests/hlsl-intrinsic/wave-shuffle.slang index 093babcce..608e57ffd 100644 --- a/tests/hlsl-intrinsic/wave-shuffle.slang +++ b/tests/hlsl-intrinsic/wave-shuffle.slang @@ -4,7 +4,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 // Disabled because vk doesn't currently support matrix types. See wave-shuffle-vk.slang //DISABLE_TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer outputBuffer; diff --git a/tests/hlsl-intrinsic/wave-vector.slang b/tests/hlsl-intrinsic/wave-vector.slang index 8d2868600..784f16c16 100644 --- a/tests/hlsl-intrinsic/wave-vector.slang +++ b/tests/hlsl-intrinsic/wave-vector.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer outputBuffer; diff --git a/tests/hlsl-intrinsic/wave.slang b/tests/hlsl-intrinsic/wave.slang index d8273080c..eb647d0fe 100644 --- a/tests/hlsl-intrinsic/wave.slang +++ b/tests/hlsl-intrinsic/wave.slang @@ -2,7 +2,7 @@ //DISABLE_TEST(compute):COMPARE_COMPUTE_EX:-slang -compute //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -use-dxil -profile cs_6_0 //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute +//TEST(compute):COMPARE_COMPUTE_EX:-cuda -compute -render-features cuda_sm_7_0 //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer RWStructuredBuffer outputBuffer; -- cgit v1.2.3