diff --git a/source/slang/diff.meta.slang b/source/slang/diff.meta.slang index b9cc0b103b..d6d9a7600e 100644 --- a/source/slang/diff.meta.slang +++ b/source/slang/diff.meta.slang @@ -1023,9 +1023,14 @@ struct DiffTensorView } }; +// Note: we need `TorchTensorType` so Slang is aware +// of when we use `TorchTensorType` in a synthisized +// function. This is important since `[CudaHost]` is +// required for any function using `TorchTensorType` /// Represents the handle of a Torch tensor object. __generic __intrinsic_type($(kIROp_TorchTensorType)) +__magic_type(TorchTensorType) struct TorchTensor { __intrinsic_op($(kIROp_TorchTensorGetView)) diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang index 7b3113fceb..58923d838e 100644 --- a/source/slang/hlsl.meta.slang +++ b/source/slang/hlsl.meta.slang @@ -19766,6 +19766,11 @@ struct TextureFootprint : __TextureFootprintData { bool _isSingleLevel; + __init(__TextureFootprintData data, bool _isSingleLevelIn) + { + _isSingleLevel = _isSingleLevelIn; + } + property isSingleLevel : bool { [__NoSideEffect] diff --git a/source/slang/slang-ast-builder.h b/source/slang/slang-ast-builder.h index 08951513dc..75c90e27a5 100644 --- a/source/slang/slang-ast-builder.h +++ b/source/slang/slang-ast-builder.h @@ -80,6 +80,15 @@ class SharedASTBuilder : public RefObject } return m_thisTypeName; } + + Name* getCtorName() + { + if (!m_ctorName) + { + m_ctorName = getNamePool()->getName("$init"); + } + return m_ctorName; + } protected: // State shared between ASTBuilders @@ -114,6 +123,7 @@ class SharedASTBuilder : public RefObject NamePool* m_namePool = nullptr; Name* m_thisTypeName = nullptr; + Name* m_ctorName = nullptr; // This is a private builder used for these shared types ASTBuilder* m_astBuilder = nullptr; diff --git a/source/slang/slang-ast-decl.h b/source/slang/slang-ast-decl.h index 36aa3a313c..226b2d66f8 100644 --- a/source/slang/slang-ast-decl.h +++ b/source/slang/slang-ast-decl.h @@ -367,14 +367,30 @@ class FunctionDeclBase : public CallableDecl Stmt* body = nullptr; }; +enum class ConstructorTags : int +{ + None = 0, + + /// Indicates whether the declaration was synthesized by + /// Slang and not actually provided by the user + Synthesized = 1 << 0, + + /// Derived classes will call this ctor if they need a memberwise ctor for public members. + MemberwiseCtorForPublicVisibility = 1 << 1, + + /// Derived classes will call this ctor if they need a memberwise ctor for public and internal members. + /// A ctor with `MemberwiseCtorForInternalVisibility` may also contain 'isMemberwiseCtorForPublicVisibility' + MemberwiseCtorForInternalVisibility = 1 << 2, +}; + // A constructor/initializer to create instances of a type class ConstructorDecl : public FunctionDeclBase { SLANG_AST_CLASS(ConstructorDecl) - // Indicates whether the declaration was synthesized by - // slang and not actually provided by the user - bool isSynthesized = false; + int options = (int)ConstructorTags::None; + void addOption(ConstructorTags option) { options |= (int)option; } + bool containsOption(ConstructorTags option) { return options & (int)option; } }; // A subscript operation used to index instances of a type diff --git a/source/slang/slang-ast-modifier.h b/source/slang/slang-ast-modifier.h index 69f86a43d2..8d9263bd45 100644 --- a/source/slang/slang-ast-modifier.h +++ b/source/slang/slang-ast-modifier.h @@ -38,6 +38,9 @@ class ToBeSynthesizedModifier : public Modifier {SLANG_AST_CLASS(ToBeSynthesized // Marks that the definition of a decl is synthesized. class SynthesizedModifier : public Modifier { SLANG_AST_CLASS(SynthesizedModifier) }; +// Marks that the definition of a decl is our $DefaultInit function. +class DefaultInitModifier : public Modifier { SLANG_AST_CLASS(DefaultInitModifier) }; + // Marks a synthesized variable as local temporary variable. class LocalTempVarModifier : public Modifier { SLANG_AST_CLASS(LocalTempVarModifier) }; diff --git a/source/slang/slang-ast-support-types.h b/source/slang/slang-ast-support-types.h index 56101bb919..665c4cb2c6 100644 --- a/source/slang/slang-ast-support-types.h +++ b/source/slang/slang-ast-support-types.h @@ -470,7 +470,6 @@ namespace Slang /// functions, so it belongs in the last phase of checking. /// DefinitionChecked, - DefaultConstructorReadyForUse = DefinitionChecked, /// The capabilities required by the decl is infered and validated. /// @@ -491,6 +490,8 @@ namespace Slang CanUseFuncSignature = ReadyForReference, CanSpecializeGeneric = ReadyForReference, CanReadInterfaceRequirements = ReadyForLookup, + DefaultConstructorReadyForUse = ReadyForConformances, + VarInitExprAreChecked = DefinitionChecked, }; /// A `DeclCheckState` plus a bit to track whether a declaration is currently being checked. diff --git a/source/slang/slang-ast-type.h b/source/slang/slang-ast-type.h index 46ea3ea559..081d136ef1 100644 --- a/source/slang/slang-ast-type.h +++ b/source/slang/slang-ast-type.h @@ -523,6 +523,11 @@ class TensorViewType : public BuiltinType Type* getElementType(); }; +class TorchTensorType : public BuiltinType +{ + SLANG_AST_CLASS(TorchTensorType) +}; + // Base class for built in string types class StringTypeBase : public BuiltinType { diff --git a/source/slang/slang-check-conversion.cpp b/source/slang/slang-check-conversion.cpp index c0d7feaff3..5a762cb181 100644 --- a/source/slang/slang-check-conversion.cpp +++ b/source/slang/slang-check-conversion.cpp @@ -15,7 +15,7 @@ namespace Slang ConversionCost SemanticsVisitor::getImplicitConversionCost( Decl* decl) { - if(auto modifier = decl->findModifier()) + if (auto modifier = decl->findModifier()) { return modifier->cost; } @@ -35,46 +35,55 @@ namespace Slang } bool SemanticsVisitor::isEffectivelyScalarForInitializerLists( - Type* type) + Type* type) { - if(as(type)) return false; - if(as(type)) return false; - if(as(type)) return false; + if (as(type)) return false; + if (as(type)) return false; + if (as(type)) return false; - if(as(type)) + if (as(type)) { return true; } - if(as(type)) + if (as(type)) { return true; } - if(as(type)) + if (as(type)) { return true; } - if(as(type)) + if (as(type)) { return true; } - if(auto declRefType = as(type)) + if (auto declRefType = as(type)) { - if(as(declRefType->getDeclRef())) + if (as(declRefType->getDeclRef())) return false; } return true; } + bool isSameSimpleExprType(Type* type1, Type* type2) + { + if (as(type1) && as(type2)) return true; + if (as(type1) && as(type2)) return true; + if (as(type1) && as(type2)) return true; + if (as(type1) && as(type2)) return true; + return false; + } + bool SemanticsVisitor::shouldUseInitializerDirectly( - Type* toType, - Expr* fromExpr) + Type* toType, + Expr* fromExpr) { // A nested initializer list should always be used directly. // - if(as(fromExpr)) + if (as(fromExpr)) { return true; } @@ -82,34 +91,45 @@ namespace Slang // If the desired type is a scalar, then we should always initialize // directly, since it isn't an aggregate. // - if(isEffectivelyScalarForInitializerLists(toType)) + if (isEffectivelyScalarForInitializerLists(toType)) return true; // If the type we are initializing isn't effectively scalar, // but the initialization expression *is*, then it doesn't // seem like direct initialization is intended. // - if(isEffectivelyScalarForInitializerLists(fromExpr->type)) + if (isEffectivelyScalarForInitializerLists(fromExpr->type)) return false; - // Once the above cases are handled, the main thing - // we want to check for is whether a direct initialization - // is possible (a type conversion exists). - // - return canCoerce(toType, fromExpr->type, fromExpr); + // If 2 types are "broadly equal" we know the types can be coerced directly + // since otherwise trying to solve `float3 tmp[2] = {float3(val), float3(val)};` will break down into searching for `float` + // (which we don't have in our init-list argument list, this will error). + // + // This logic should be completly cleaned up at some point by only using constructors for `{...}` and attempting to use + // constructors before fallback coerce logic is tried. + if (toType->equals(fromExpr->type)) + return true; + + // `vector` and `vector` may not be equal types, but, we have no choice in this senario but to assume coerce logic + // can resolve such situations, otherwise we need to fail + if (isSameSimpleExprType(toType, fromExpr->type)) + return true; + + return false; } + template bool SemanticsVisitor::_readValueFromInitializerList( - Type* toType, - Expr** outToExpr, + Type* toType, + Expr** outToExpr, InitializerListExpr* fromInitializerListExpr, - UInt &ioInitArgIndex) + UInt& ioInitArgIndex) { // First, we will check if we have run out of arguments // on the initializer list. // UInt initArgCount = fromInitializerListExpr->args.getCount(); - if(ioInitArgIndex >= initArgCount) + if (ioInitArgIndex >= initArgCount) { // If we are at the end of the initializer list, // then our ability to read an argument depends @@ -128,7 +148,7 @@ namespace Slang // for aggregate initialization. // auto firstInitExpr = fromInitializerListExpr->args[ioInitArgIndex]; - if(shouldUseInitializerDirectly(toType, firstInitExpr)) + if (shouldUseInitializerDirectly(toType, firstInitExpr) && canCoerce(toType, firstInitExpr->type, firstInitExpr)) { ioInitArgIndex++; return _coerce( @@ -144,7 +164,7 @@ namespace Slang // expressions, then everything could be thrown off and we // shouldn't keep trying to read arguments. // - if( IsErrorExpr(firstInitExpr) ) + if (IsErrorExpr(firstInitExpr)) { // Stop reading arguments, as if we'd reached // the end of the list. @@ -156,7 +176,7 @@ namespace Slang // The fallback case is to recursively read the // type from the same list as an aggregate. // - return _readAggregateValueFromInitializerList( + return _readAggregateValueFromInitializerList( toType, outToExpr, fromInitializerListExpr, @@ -166,17 +186,17 @@ namespace Slang DeclRefType* findBaseStructType(ASTBuilder* astBuilder, DeclRef structTypeDeclRef) { auto inheritanceDecl = getMembersOfType(astBuilder, structTypeDeclRef).getFirstOrNull(); - if(!inheritanceDecl) + if (!inheritanceDecl) return nullptr; auto baseType = getBaseType(astBuilder, inheritanceDecl); auto baseDeclRefType = as(baseType); - if(!baseDeclRefType) + if (!baseDeclRefType) return nullptr; auto baseDeclRef = baseDeclRefType->getDeclRef(); auto baseStructDeclRef = baseDeclRef.as(); - if(!baseStructDeclRef) + if (!baseStructDeclRef) return nullptr; return baseDeclRefType; @@ -201,11 +221,24 @@ namespace Slang return baseStructDeclRef; } + GenericAppDeclRef* _getGenericAppDeclRefType(Type* argType) + { + auto argDeclRefType = as(argType); + if (argDeclRefType) + { + if (auto genericType = as(argDeclRefType->getDeclRefBase())) + return genericType; + } + + return nullptr; + } + + template bool SemanticsVisitor::_readAggregateValueFromInitializerList( - Type* inToType, - Expr** outToExpr, + Type* inToType, + Expr** outToExpr, InitializerListExpr* fromInitializerListExpr, - UInt &ioArgIndex) + UInt& ioArgIndex) { auto toType = inToType; UInt argCount = fromInitializerListExpr->args.getCount(); @@ -214,14 +247,22 @@ namespace Slang // we will collect the new arguments here List coercedArgs; - if(isEffectivelyScalarForInitializerLists(toType)) + if (isEffectivelyScalarForInitializerLists(toType)) { // For any type that is effectively a non-aggregate, // we expect to read a single value from the initializer list // - if(ioArgIndex < argCount) + if (ioArgIndex < argCount) { auto arg = fromInitializerListExpr->args[ioArgIndex++]; + + if (!canCoerce(toType, arg->type, nullptr)) + { + if (outToExpr) + *outToExpr = arg; + return true; + } + return _coerce( CoercionSite::Initializer, toType, @@ -250,46 +291,46 @@ namespace Slang UInt elementCount = 0; if (auto constElementCount = as(toElementCount)) { - elementCount = (UInt) constElementCount->getValue(); + elementCount = (UInt)constElementCount->getValue(); } else { // We don't know the element count statically, // so what are we supposed to be doing? // - if(outToExpr) + if (outToExpr) { getSink()->diagnose(fromInitializerListExpr, Diagnostics::cannotUseInitializerListForVectorOfUnknownSize, toElementCount); } return false; } - for(UInt ee = 0; ee < elementCount; ++ee) + for (UInt ee = 0; ee < elementCount; ++ee) { Expr* coercedArg = nullptr; - bool argResult = _readValueFromInitializerList( + bool argResult = _readValueFromInitializerList( toElementType, outToExpr ? &coercedArg : nullptr, fromInitializerListExpr, ioArgIndex); // No point in trying further if any argument fails - if(!argResult) + if (!argResult) return false; - if( coercedArg ) + if (coercedArg) { coercedArgs.add(coercedArg); } } } - else if(auto toArrayType = as(toType)) + else if (auto toArrayType = as(toType)) { // TODO(tfoley): If we can compute the size of the array statically, // then we want to check that there aren't too many initializers present auto toElementType = toArrayType->getElementType(); - if(!toArrayType->isUnsized()) + if (!toArrayType->isUnsized()) { auto toElementCount = toArrayType->getElementCount(); @@ -299,34 +340,34 @@ namespace Slang UInt elementCount = 0; if (auto constElementCount = as(toElementCount)) { - elementCount = (UInt) constElementCount->getValue(); + elementCount = (UInt)constElementCount->getValue(); } else { // We don't know the element count statically, // so what are we supposed to be doing? // - if(outToExpr) + if (outToExpr) { getSink()->diagnose(fromInitializerListExpr, Diagnostics::cannotUseInitializerListForArrayOfUnknownSize, toElementCount); } return false; } - for(UInt ee = 0; ee < elementCount; ++ee) + for (UInt ee = 0; ee < elementCount; ++ee) { Expr* coercedArg = nullptr; - bool argResult = _readValueFromInitializerList( + bool argResult = _readValueFromInitializerList( toElementType, outToExpr ? &coercedArg : nullptr, fromInitializerListExpr, ioArgIndex); // No point in trying further if any argument fails - if(!argResult) + if (!argResult) return false; - if( coercedArg ) + if (coercedArg) { coercedArgs.add(coercedArg); } @@ -339,22 +380,22 @@ namespace Slang // the element count. // UInt elementCount = 0; - while(ioArgIndex < argCount) + while (ioArgIndex < argCount) { Expr* coercedArg = nullptr; - bool argResult = _readValueFromInitializerList( + bool argResult = _readValueFromInitializerList( toElementType, outToExpr ? &coercedArg : nullptr, fromInitializerListExpr, ioArgIndex); // No point in trying further if any argument fails - if(!argResult) + if (!argResult) return false; elementCount++; - if( coercedArg ) + if (coercedArg) { coercedArgs.add(coercedArg); } @@ -366,7 +407,7 @@ namespace Slang m_astBuilder->getIntVal(m_astBuilder->getIntType(), elementCount)); } } - else if(auto toMatrixType = as(toType)) + else if (auto toMatrixType = as(toType)) { // In the general case, the initializer list might comprise // both vectors and scalars. @@ -388,88 +429,246 @@ namespace Slang if (auto constRowCount = as(toMatrixType->getRowCount())) { - rowCount = (UInt) constRowCount->getValue(); + rowCount = (UInt)constRowCount->getValue(); } else { // We don't know the element count statically, // so what are we supposed to be doing? // - if(outToExpr) + if (outToExpr) { getSink()->diagnose(fromInitializerListExpr, Diagnostics::cannotUseInitializerListForMatrixOfUnknownSize, toMatrixType->getRowCount()); } return false; } - for(UInt rr = 0; rr < rowCount; ++rr) + for (UInt rr = 0; rr < rowCount; ++rr) { Expr* coercedArg = nullptr; - bool argResult = _readValueFromInitializerList( + bool argResult = _readValueFromInitializerList( toRowType, outToExpr ? &coercedArg : nullptr, fromInitializerListExpr, ioArgIndex); // No point in trying further if any argument fails - if(!argResult) + if (!argResult) return false; - if( coercedArg ) + if (coercedArg) { coercedArgs.add(coercedArg); } } } - else if(auto toDeclRefType = as(toType)) + else if (auto toDeclRefType = as(toType)) { - auto toTypeDeclRef = toDeclRefType->getDeclRef(); - if(auto toStructDeclRef = toTypeDeclRef.as()) + // We allow using legacy "init list" logic as a backup if something fails + if constexpr (UseLegacyLogic) { - // Trying to initialize a `struct` type given an initializer list. - // - // Before we iterate over the fields, we want to check if this struct - // inherits from another `struct` type. If so, we want to read - // an initializer for that base type first. - // - if (auto baseStructType = findBaseStructType(m_astBuilder, toStructDeclRef)) + auto toTypeDeclRef = toDeclRefType->getDeclRef(); + if (auto toStructDeclRef = toTypeDeclRef.as()) { - Expr* coercedArg = nullptr; - bool argResult = _readValueFromInitializerList( - baseStructType, - outToExpr ? &coercedArg : nullptr, - fromInitializerListExpr, - ioArgIndex); - - // No point in trying further if any argument fails - if (!argResult) - return false; + // Trying to initialize a `struct` type given an initializer list. + // + // Before we iterate over the fields, we want to check if this struct + // inherits from another `struct` type. If so, we want to read + // an initializer for that base type first. + // + if (auto baseStructType = findBaseStructType(m_astBuilder, toStructDeclRef)) + { + Expr* coercedArg = nullptr; + bool argResult = _readValueFromInitializerList( + baseStructType, + outToExpr ? &coercedArg : nullptr, + fromInitializerListExpr, + ioArgIndex); + + // No point in trying further if any argument fails + if (!argResult) + return false; + + if (coercedArg) + { + coercedArgs.add(coercedArg); + } + } - if (coercedArg) + // We will go through the fields in order and try to match them + // up with initializer arguments. + // + for (auto fieldDeclRef : getMembersOfType(m_astBuilder, toStructDeclRef, MemberFilterStyle::Instance)) { - coercedArgs.add(coercedArg); + Expr* coercedArg = nullptr; + bool argResult = _readValueFromInitializerList( + getType(m_astBuilder, fieldDeclRef), + outToExpr ? &coercedArg : nullptr, + fromInitializerListExpr, + ioArgIndex); + + // No point in trying further if any argument fails + if (!argResult) + return false; + + if (coercedArg) + { + coercedArgs.add(coercedArg); + } } } - - // We will go through the fields in order and try to match them - // up with initializer arguments. - // - for(auto fieldDeclRef : getMembersOfType(m_astBuilder, toStructDeclRef, MemberFilterStyle::Instance)) + } + else + { + auto toTypeDeclRef = toDeclRefType->getDeclRef(); + if (auto toStructDeclRef = toTypeDeclRef.as()) { - Expr* coercedArg = nullptr; - bool argResult = _readValueFromInitializerList( - getType(m_astBuilder, fieldDeclRef), - outToExpr ? &coercedArg : nullptr, - fromInitializerListExpr, - ioArgIndex); + // We will try to coerce the initializer list into a struct following these steps: + // 1. Ensure our `StructDecl` has declaratations created for all auto-generated constructors + // 2. We have a `{}` (0 arguments), try to coerce + // 3. We have a `{arg1, arg2...}`, try to coerce - // No point in trying further if any argument fails - if(!argResult) - return false; + auto toStructDecl = toStructDeclRef.getDecl(); + // 1. Ensure our `StructDecl` has declaratations created for all auto-generated constructors + ensureDecl(toStructDecl, DeclCheckState::DefaultConstructorReadyForUse); - if( coercedArg ) + // 2. We have a `{}` (0 arguments), try to coerce + if (argCount == 0) { - coercedArgs.add(coercedArg); + if (outToExpr) + { + *outToExpr = constructDefaultInitListFunc(this, toStructDecl, toType, ConstructDefaultInitListOptions::CheckToAvoidRecursion); + (*outToExpr)->loc = fromInitializerListExpr->loc; + } + return true; + } + + // 3. We have a `{arg1, arg2...}`, try to coerce + // + // For this situation we have more steps: + // a. Find constructor candidate for our particular init-list + // b. Collect and coerce init-list arguments into constructor parameters + // c. Create InvokeExpr for our valid ConstructorDecl + List ctorList = _getCtorList(this->getASTBuilder(), this, toStructDecl, nullptr); + bool allowCStyleInitList = isCStyleStructDecl(this, toStructDecl, ctorList); + + List maybeArgList; + maybeArgList.reserve(argCount); + Index ioArgIndexMirror = ioArgIndex; + UInt ioArgIndexCandidate = 0; + + for (auto& ctor : ctorList) + { + // a. Find constructor candidate for our particular init-list + + // Don't try to init default ctor with this logic + auto ctorParamCount = ctor->getParameters().getCount(); + if (ctorParamCount == 0) + continue; + + ioArgIndexCandidate = ioArgIndexMirror; + ioArgIndex = ctorParamCount; + + // if allowCStyleInitList, process any ctor which comes next using the most members possible. ioArgIndex may not be 0. + // if !allowCStyleInitList, process any ctor that exactly matched our argument count. ioArgIndex must start at 0. + if (!allowCStyleInitList && ctorParamCount != Index(argCount)) + continue; + + // b. Collect and coerce init-list arguments into constructor parameters + List maybeCandidate; + auto parameters = getParameters(m_astBuilder, ctor); + auto parametersCount = parameters.getCount(); + + for (auto index = coercedArgs.getCount(); index < parametersCount; index++) + { + // If we ran out of elements and allow using a c-style-partial-initialization-list, end early. + if ((Index)ioArgIndexCandidate > fromInitializerListExpr->args.getCount()) + break; + + auto ctorParam = parameters[index]; + auto paramType = ctorParam.getDecl()->type.type; + for (auto i : getMembersOfType(m_astBuilder, toStructDeclRef, MemberFilterStyle::Instance)) + if (i.getDecl()->type.type == paramType) + paramType = getType(m_astBuilder, i); + + Expr* coercedArg = nullptr; + _readValueFromInitializerList( + paramType, + &coercedArg, + fromInitializerListExpr, + ioArgIndexCandidate); + + if (coercedArg) + maybeArgList.add(coercedArg); + else + break; + } + + if (!allowCStyleInitList && maybeArgList.getCount() != ctorParamCount) + continue; + + // Skip non-visible constructors. + if (!isDeclVisible(ctor)) + { + // if an exact argument match with our init-list we know the user meant to use a + // member-wise constructor, error + if (ctorParamCount == Index(argCount)) + getSink()->diagnose(fromInitializerListExpr, Diagnostics::declIsNotVisible, ctor); + continue; + } + + // c. Create InvokeExpr for our valid ConstructorDecl + // We cannot fail anymore, set `ioArgIndex` to the 'used up arg count'. + if (outToExpr) + { + ioArgIndex = ioArgIndexCandidate; + + for (auto i : maybeArgList) + coercedArgs.add(i); + + List argTypes; + for (auto i : coercedArgs) + argTypes.add(i->type); + + auto* varExpr = getASTBuilder()->create(); + varExpr->type = (QualType)getASTBuilder()->getTypeType(toType); + varExpr->declRef = isDeclRefTypeOf(toType); + auto* constructorExpr = getASTBuilder()->create(); + constructorExpr->functionExpr = varExpr; + constructorExpr->arguments.addRange(coercedArgs); + auto resolvedConstructorExpr = CheckExpr(constructorExpr); + + *outToExpr = resolvedConstructorExpr; + } + return true; + } + + // If we have a generic being compared to another generic (with different generic arguments) + // coerce logic will fail regardless of if generics are valid together. Example is below: + // + // MyStruct tmp = {MyStructBase(), 1}; // Assume 'U' is unresolved at this point in time but equal to T + // + // To handle this since this is not verifiable coerce logic: + // 1. We need to ensure we don't have any matching constructors + // 2. if '1.' is true we can assign the possibly compatible generics and let generic resolution + // diagnose an error down the line if types are in-compatible + + if (auto toGenericType = _getGenericAppDeclRefType(toType)) + { + auto arg = fromInitializerListExpr->args[ioArgIndexMirror]; + if (auto fromGenericType = _getGenericAppDeclRefType(getType(m_astBuilder, arg))) + { + for (;;) + { + if (toGenericType->getBase() != fromGenericType->getBase()) + break; + + ioArgIndex = ioArgIndexMirror; + *outToExpr = arg; + ioArgIndex++; + return true; + } + } } } } @@ -481,7 +680,7 @@ namespace Slang // list invalid if we are trying to read something // off of it that wasn't handled by the cases above. // - if(outToExpr) + if (outToExpr) { getSink()->diagnose(fromInitializerListExpr, Diagnostics::cannotUseInitializerListForType, inToType); } @@ -491,7 +690,7 @@ namespace Slang // We were able to coerce all the arguments given, and so // we need to construct a suitable expression to remember the result // - if(outToExpr) + if (outToExpr) { auto toInitializerListExpr = m_astBuilder->create(); toInitializerListExpr->loc = fromInitializerListExpr->loc; @@ -503,10 +702,10 @@ namespace Slang // if (auto func = getParentFuncOfVisitor()) { - if (func->findModifier() && + if (func->findModifier() && !isTypeDifferentiable(toType)) { - for (auto &arg : toInitializerListExpr->args) + for (auto& arg : toInitializerListExpr->args) { if (isTypeDifferentiable(arg->type.type)) { @@ -525,9 +724,32 @@ namespace Slang return true; } + void _maybeRunLegacyLogic( + SemanticsVisitor* visitor, + Type* toType, + Expr** outToExpr, + InitializerListExpr* fromInitializerListExpr, + bool& result, + UInt& argIndex, + bool& ranLegacyLogic + ) + { + // Try to use legacy struct logic if we have a 'struct' toType for Slang backwards compatability + // Note: if we remove this entire `if` block, "LegacyLogic" will be no-longer be enabled + if (auto declRefType = as(toType)) + { + if (declRefType->getDeclRef().as()) + { + argIndex = 0; + ranLegacyLogic = true; + result = visitor->_readAggregateValueFromInitializerList(toType, outToExpr, fromInitializerListExpr, argIndex); + } + } + } + bool SemanticsVisitor::_coerceInitializerList( - Type* toType, - Expr** outToExpr, + Type* toType, + Expr** outToExpr, InitializerListExpr* fromInitializerListExpr) { UInt argCount = fromInitializerListExpr->args.getCount(); @@ -543,21 +765,38 @@ namespace Slang // If this isn't prohibited, then we can proceed to try and coerce from // the initializer list itself; assuming that coercion is closed under // composition this shouldn't fail. - if(!as(fromInitializerListExpr->type) && - !canCoerce(toType, fromInitializerListExpr->type, nullptr)) + if (!as(fromInitializerListExpr->type) && + !canCoerce(toType, fromInitializerListExpr->type, nullptr)) return _failedCoercion(toType, outToExpr, fromInitializerListExpr); - if(!_readAggregateValueFromInitializerList(toType, outToExpr, fromInitializerListExpr, argIndex)) - return false; + // If we fail, try to run legacy logic. + bool usedLegacy = false; + bool result = _readAggregateValueFromInitializerList(toType, outToExpr, fromInitializerListExpr, argIndex); + // Note: if we remove this entire `if` block for a `return false`, "LegacyLogic" will be no-longer be enabled for "failed" case + if(!result) + { + _maybeRunLegacyLogic(this, toType, outToExpr, fromInitializerListExpr, result, argIndex, usedLegacy); + if (!result) + return false; + } if(argIndex != argCount) { if( outToExpr ) { - getSink()->diagnose(fromInitializerListExpr, Diagnostics::tooManyInitializers, argIndex, argCount); + // Note: if we remove this entire `if` block, "LegacyLogic" will be no-longer be enabled for "too-many-arg" case + if (!usedLegacy) + _maybeRunLegacyLogic(this, toType, outToExpr, fromInitializerListExpr, result, argIndex, usedLegacy); + // If we fail without/with legacy logic, error as usual + if (argIndex != argCount) + getSink()->diagnose(fromInitializerListExpr, Diagnostics::cannotFindMatchingInitListToConstructorCount, argCount, toType); } } + // Warn that legacy logic is the reason for code to compile + if (usedLegacy) + getSink()->diagnose(fromInitializerListExpr, Diagnostics::usingLegacyInitListStructLogic); + return true; } @@ -925,9 +1164,7 @@ namespace Slang } if (outToExpr) { - auto* defaultExpr = getASTBuilder()->create(); - defaultExpr->type = QualType(toType); - *outToExpr = defaultExpr; + *outToExpr = createDefaultConstructExprForType(getASTBuilder(), QualType(toType), {}); } return true; } diff --git a/source/slang/slang-check-decl.cpp b/source/slang/slang-check-decl.cpp index 2307c16696..7c7e1101cb 100644 --- a/source/slang/slang-check-decl.cpp +++ b/source/slang/slang-check-decl.cpp @@ -1799,10 +1799,52 @@ namespace Slang checkVisibility(varDecl); } - static ConstructorDecl* _createCtor(SemanticsDeclVisitorBase* visitor, ASTBuilder* m_astBuilder, AggTypeDecl* decl) + void addVisibilityModifier(ASTBuilder* builder, Decl* decl, DeclVisibility vis) + { + switch (vis) + { + case DeclVisibility::Public: + addModifier(decl, builder->create()); + break; + case DeclVisibility::Internal: + addModifier(decl, builder->create()); + break; + case DeclVisibility::Private: + addModifier(decl, builder->create()); + break; + default: + break; + } + } + + static void addAutoDiffModifiersToFunc( + SemanticsDeclVisitorBase* visitor, + ASTBuilder* m_astBuilder, + FunctionDeclBase* func) + { + if (visitor->isTypeDifferentiable(func->returnType.type)) + { + addModifier(func, m_astBuilder->create()); + addModifier(func, m_astBuilder->create()); + } + else + addModifier(func, m_astBuilder->create()); + } + + + struct CreateCtorArg + { + VarDeclBase* m_argToCopy; + Expr* m_initExpr = nullptr; + }; + static ConstructorDecl* _createCtor( + SemanticsDeclVisitorBase* visitor, + ASTBuilder* m_astBuilder, + AggTypeDecl* decl, + List&& argList, + DeclVisibility visibility) { auto ctor = m_astBuilder->create(); - addModifier(ctor, m_astBuilder->create()); auto ctorName = visitor->getName("$init"); ctor->ownedScope = m_astBuilder->create(); ctor->ownedScope->containerDecl = ctor; @@ -1823,61 +1865,27 @@ namespace Slang body->closingSourceLoc = ctor->closingSourceLoc; ctor->body = body; body->body = m_astBuilder->create(); - ctor->isSynthesized = true; - decl->addMember(ctor); - return ctor; - } - - static ConstructorDecl* _getDefaultCtor(StructDecl* structDecl) - { - for (auto ctor : structDecl->getMembersOfType()) - { - if (!ctor->body || ctor->members.getCount() != 0) - continue; - return ctor; - } - return nullptr; - } - - - static List _getCtorList(ASTBuilder* m_astBuilder, SemanticsVisitor* visitor, StructDecl* structDecl, ConstructorDecl** defaultCtorOut) - { - List ctorList; - - auto ctorLookupResult = lookUpMember( - m_astBuilder, - visitor, - visitor->getName("$init"), - DeclRefType::create(m_astBuilder, structDecl), - structDecl->ownedScope, - LookupMask::Function, - (LookupOptions)((Index)LookupOptions::IgnoreInheritance | (Index)LookupOptions::NoDeref)); + addModifier(ctor, m_astBuilder->create()); + ctor->addOption(ConstructorTags::Synthesized); - if (!ctorLookupResult.isValid()) - return ctorList; + // kIROp_TorchTensorType can only be used in host-callable functions + if(auto intrinsicType = decl->findModifier()) + if(intrinsicType->irOp == kIROp_TorchTensorType) + addModifier(ctor, m_astBuilder->create()); - auto lookupResultHandle = [&](LookupResultItem& item) + for (auto arg : argList) { - auto ctor = as(item.declRef.getDecl()); - if (!ctor || !ctor->body) - return; - ctorList.add(ctor); - if (ctor->members.getCount() != 0) - return; - *defaultCtorOut = ctor; - }; - if (ctorLookupResult.items.getCount() == 0) - { - lookupResultHandle(ctorLookupResult.item); - return ctorList; - } - - for (auto m : ctorLookupResult.items) - { - lookupResultHandle(m); - } - - return ctorList; + auto param = m_astBuilder->create(); + param->type = (TypeExp)arg.m_argToCopy->type; + param->initExpr = arg.m_initExpr; + param->parentDecl = ctor; + param->loc = ctor->loc; + ctor->members.add(param); + } + addVisibilityModifier(m_astBuilder, ctor, visibility); + addAutoDiffModifiersToFunc(visitor, m_astBuilder, ctor); + decl->addMember(ctor); + return ctor; } void SemanticsDeclHeaderVisitor::visitStructDecl(StructDecl* structDecl) @@ -1942,72 +1950,6 @@ namespace Slang return true; } - bool isDefaultInitializable(VarDeclBase* varDecl) - { - if (!DiagnoseIsAllowedInitExpr(varDecl, nullptr)) - return false; - - // Find struct and modifiers associated with varDecl - StructDecl* structDecl = as(varDecl); - if (auto declRefType = as(varDecl->getType())) - { - if (auto genericAppRefDecl = as(declRefType->getDeclRefBase())) - { - auto baseGenericRefType = genericAppRefDecl->getBase()->getDecl(); - if (auto baseTypeStruct = as(baseGenericRefType)) - { - structDecl = baseTypeStruct; - } - else if (auto genericDecl = as(baseGenericRefType)) - { - if(auto innerTypeStruct = as(genericDecl->inner)) - structDecl = innerTypeStruct; - } - } - } - if (structDecl) - { - // find if a type is non-copyable - if (structDecl->findModifier()) - return false; - } - - return true; - } - - static Expr* constructDefaultInitExprForVar(SemanticsVisitor* visitor, VarDeclBase* varDecl) - { - if (!varDecl->type || !varDecl->type.type) - return nullptr; - - if (!isDefaultInitializable(varDecl)) - return nullptr; - - ConstructorDecl* defaultCtor = nullptr; - auto declRefType = as(varDecl->type.type); - if (declRefType) - { - if (auto structDecl = as(declRefType->getDeclRef().getDecl())) - { - defaultCtor = _getDefaultCtor(structDecl); - } - } - - if (defaultCtor) - { - auto* invoke = visitor->getASTBuilder()->create(); - auto member = visitor->getASTBuilder()->getMemberDeclRef(declRefType->getDeclRef(), defaultCtor); - invoke->functionExpr = visitor->ConstructDeclRefExpr(member, nullptr, defaultCtor->getName(), defaultCtor->loc, nullptr); - return invoke; - } - else - { - auto* defaultCall = visitor->getASTBuilder()->create(); - defaultCall->type = QualType(varDecl->type); - return defaultCall; - } - } - void SemanticsDeclBodyVisitor::checkVarDeclCommon(VarDeclBase* varDecl) { DiagnoseIsAllowedInitExpr(varDecl, getSink()); @@ -2018,7 +1960,7 @@ namespace Slang && as(varDecl) ) { - varDecl->initExpr = constructDefaultInitExprForVar(this, varDecl); + varDecl->initExpr = constructDefaultInitExprForVar(this, varDecl->type, varDecl); } if (auto initExpr = varDecl->initExpr) @@ -2215,24 +2157,6 @@ namespace Slang } } - void addVisibilityModifier(ASTBuilder* builder, Decl* decl, DeclVisibility vis) - { - switch (vis) - { - case DeclVisibility::Public: - addModifier(decl, builder->create()); - break; - case DeclVisibility::Internal: - addModifier(decl, builder->create()); - break; - case DeclVisibility::Private: - addModifier(decl, builder->create()); - break; - default: - break; - } - } - bool SemanticsVisitor::trySynthesizeDifferentialAssociatedTypeRequirementWitness( ConformanceCheckingContext* context, DeclRef requirementDeclRef, @@ -2661,6 +2585,95 @@ namespace Slang } } + // Annotate ctor as a memberwise ctor. A non synthisized function may be annotated as a memberwise ctor + // if it has a compatible parameter list with a memberwise ctor. + void _annotateMemberwiseCtorWithVisibility(ConstructorDecl* ctor, DeclVisibility visibility) + { + switch (visibility) + { + case DeclVisibility::Public: + ctor->addOption(ConstructorTags::MemberwiseCtorForPublicVisibility); + break; + case DeclVisibility::Internal: + ctor->addOption(ConstructorTags::MemberwiseCtorForInternalVisibility); + break; + default: + break; + } + } + + static ConstructorDecl* _tryToGenerateCtorWithArgList( + SemanticsDeclVisitorBase* visitor, + ASTBuilder* astBuilder, + List&& args, + List& existingCtorList, + StructDecl* structDecl, + DeclVisibility visibility) + { + // Find any existing ctor type conflicts + for (auto ctor : existingCtorList) + { + auto existingCtorArgs = ctor->getParameters(); + auto existingCtorArgsLength = ctor->getParameters().getCount(); + auto newCtorArgsLength = args.getCount(); + + // Different arg count, ctor are not conflicting + if (existingCtorArgsLength != newCtorArgsLength) + continue; + + // if both arg lists are empty, return, we cannot have 2 default ctor's + if (!existingCtorArgsLength && !newCtorArgsLength) + return nullptr; + + // Check if every newCtorArg[i] is castable to existingCtorArg[i] (if so this is a conflicting ctor) + bool equalCtor = true; + for (Index i = 0; i < newCtorArgsLength; i++) + { + auto newCtorArg = args[i].m_argToCopy; + auto existingCtorArg = existingCtorArgs[i]; + if (visitor->getConversionCost(newCtorArg->getType(), existingCtorArg->getType()) == kConversionCost_Impossible) + { + equalCtor = false; + break; + } + } + if (equalCtor) + { + _annotateMemberwiseCtorWithVisibility(ctor, visibility); + return nullptr; + } + } + + // We did not have any ctor conflicts, auto-generate a ctor + auto generatedCtor = _createCtor(visitor, astBuilder, structDecl, std::move(args), visibility); + existingCtorList.add(generatedCtor); + _annotateMemberwiseCtorWithVisibility(generatedCtor, visibility); + return generatedCtor; + } + + /// TODO: cache result + /// Since we require to generate all declarations for ctor's in the conformance pass + /// Slang needs to be 100% certain that the ctor we are generating is viable, else + /// the wittness tables we will setup with our ctor-declarations will be invalid. + bool _structHasMemberWithValue(ASTBuilder* m_astBuilder, StructDecl* structDecl) + { + if (getMembersOfType(m_astBuilder, structDecl, MemberFilterStyle::Instance).getFirstOrNull()) + return true; + + for (auto inheritanceMember : getMembersOfType(m_astBuilder, structDecl, MemberFilterStyle::Instance)) + { + auto declRefType = as(inheritanceMember.getDecl()->base.type); + if (!declRefType) + continue; + auto baseStruct = as(declRefType->getDeclRef().getDecl()); + if (!baseStruct) + continue; + if (_structHasMemberWithValue(m_astBuilder, baseStruct)) + return true; + } + return false; + } + // Concretize interface conformances so that we have witnesses as required for lookup. // for lookup. struct SemanticsDeclConformancesVisitor @@ -2679,6 +2692,209 @@ namespace Slang // void visitAggTypeDecl(AggTypeDecl* aggTypeDecl) { + if (auto structDecl = as(aggTypeDecl)) + { + // First part of auto-generating constructors/functions is inside `SemanticsDeclConformancesVisitor::visitAggTypeDecl` + // * Create declarations for constructors/functions. This is important so Slang may use the declarations before + // we have the chance to synthisize function bodies + // Second part of auto-generating constructors/functions is inside `SemanticsDeclBodyVisitor::visitAggTypeDecl` + // * Add bodies to our declarations. + // + // These are split up so we can assign declarations before we finish definitions + // + // + // Steps: + // 1. Generate $DefaultInit + // 2. Add an declaration for default-ctor if missing a real default-ctor. + // 3. Generate 'member-wise' constructors + + if (_structHasMemberWithValue(getASTBuilder(), structDecl)) + { + // 1. Generate $DefaultInit + // + // $DefaultInit is a synthisized static-function only used In 2 cases: + // 1. if `{}` is used inside a `__init()` + // 2. if `{}` is used and a user has a 'synthisized __init()` + // + // Use of $DefaultInit is only for functionality of `{}` to avoid hacks. + { + auto defaultInitFunc = m_astBuilder->create(); + auto ctorName = getName("$DefaultInit"); + defaultInitFunc->ownedScope = m_astBuilder->create(); + defaultInitFunc->ownedScope->containerDecl = defaultInitFunc; + defaultInitFunc->ownedScope->parent = getScope(structDecl); + defaultInitFunc->parentDecl = structDecl; + defaultInitFunc->loc = structDecl->loc; + defaultInitFunc->closingSourceLoc = defaultInitFunc->loc; + defaultInitFunc->nameAndLoc.name = ctorName; + defaultInitFunc->nameAndLoc.loc = defaultInitFunc->loc; + defaultInitFunc->returnType.type = calcThisType(makeDeclRef(structDecl)); + auto body = m_astBuilder->create(); + body->scopeDecl = m_astBuilder->create(); + body->scopeDecl->ownedScope = m_astBuilder->create(); + body->scopeDecl->ownedScope->parent = getScope(defaultInitFunc); + body->scopeDecl->parentDecl = defaultInitFunc; + body->scopeDecl->loc = defaultInitFunc->loc; + body->scopeDecl->closingSourceLoc = defaultInitFunc->loc; + body->closingSourceLoc = defaultInitFunc->closingSourceLoc; + defaultInitFunc->body = body; + body->body = m_astBuilder->create(); + + addAutoDiffModifiersToFunc(this, m_astBuilder, defaultInitFunc); + addModifier(defaultInitFunc, m_astBuilder->create()); + + addModifier(defaultInitFunc, m_astBuilder->create()); + addVisibilityModifier(m_astBuilder, defaultInitFunc, getDeclVisibility(structDecl)); + addModifier(defaultInitFunc, m_astBuilder->create()); + structDecl->addMember(defaultInitFunc); + } + + // 2. Add an declaration for default-ctor if missing a real default-ctor. + ConstructorDecl* defaultCtor = nullptr; + List ctorList = _getCtorList(this->getASTBuilder(), this, structDecl, &defaultCtor); + if (!defaultCtor) + { + defaultCtor = _createCtor(this, m_astBuilder, structDecl, {}, getDeclVisibility(structDecl)); + ctorList.add(defaultCtor); + } + + + // 3. Generate 'member-wise' constructors + // + // Add an empty constructor for all combinations of visibility and access + // which is possible: + // public constructor - usable *outside class scope* in a *different module* + List publicCtorArgs; + // public-internal constructor - usable *outside class scope* in the *same module* + List publicInternalCtorArgs; + // public-private-internal constructor - usable *inside class scope* in the *same module* + List publicInternalPrivateCtorArgs; + + // Harvest parameters which map to the base type ctor. + // Note: assumes 1 structDecl, N number inheritance decl + for (auto inheritanceMember : structDecl->getMembersOfType()) + { + auto declRefType = as(inheritanceMember->base.type); + if (!declRefType) + continue; + auto baseStruct = as(declRefType->getDeclRef().getDecl()); + if (!baseStruct) + continue; + + DeclVisibility baseVisibilityToDerived = (isVisibilityOfDeclVisibleInScope(baseStruct, DeclVisibility::Internal, structDecl->ownedScope)) ? DeclVisibility::Internal : DeclVisibility::Public; + + ConstructorDecl* ctorForPublic = nullptr; + ConstructorDecl* ctorForInternal = nullptr; + + // From our synthisized ctor's find the publicMemberCtor and internalMemberCtor. + List baseCtorList = _getCtorList(this->getASTBuilder(), this, baseStruct, nullptr); + for (auto i : baseCtorList) + { + if (i->containsOption(ConstructorTags::MemberwiseCtorForPublicVisibility)) + ctorForPublic = i; + if (i->containsOption(ConstructorTags::MemberwiseCtorForInternalVisibility)) + ctorForInternal = i; + } + + // If base is not defined in the same module (or if the internal ctor is missing) + // set the parameter list to a base-type member-wise ctor. + if (baseVisibilityToDerived == DeclVisibility::Public + || !ctorForInternal) + { + ctorForInternal = ctorForPublic; + } + + if (ctorForPublic) + { + for (auto i : ctorForPublic->getParameters()) + { + publicCtorArgs.add({ i,nullptr }); + } + for (auto i : ctorForInternal->getParameters()) + { + publicInternalCtorArgs.add({ i,nullptr }); + publicInternalPrivateCtorArgs.add({ i,nullptr }); + } + } + } + + // If we have an internal field which lacks an initExpr we cannot allow + // a memberwise ctor to be synthesized for public members. + // This principal also applies for private members and internal/public member-wise ctor synthisis. + DeclVisibility maxVisibilityCtorToGenerate = getDeclVisibility(structDecl); + + // Here we collect all variables which will be apart of our member-wise ctor's. + for (auto varDeclRef : getMembersOfType(getASTBuilder(), structDecl, MemberFilterStyle::Instance)) + { + auto varDecl = varDeclRef.getDecl(); + ensureDecl(varDecl, DeclCheckState::TypesFullyResolved); + + // Do not map a read-only variable for default-init + if (!getTypeForDeclRef(m_astBuilder, varDeclRef, varDeclRef.getLoc()).isLeftValue) + continue; + + auto declVisibility = getDeclVisibility(varDecl); + switch (declVisibility) + { + case DeclVisibility::Private: + publicInternalPrivateCtorArgs.add({ varDecl, nullptr }); + if (!varDecl->initExpr) + maxVisibilityCtorToGenerate = DeclVisibility::Private; + break; + case DeclVisibility::Internal: + publicInternalPrivateCtorArgs.add({ varDecl, nullptr }); + publicInternalCtorArgs.add({ varDecl, nullptr }); + if (!varDecl->initExpr) + maxVisibilityCtorToGenerate = DeclVisibility::Internal; + break; + case DeclVisibility::Public: + publicInternalPrivateCtorArgs.add({ varDecl, nullptr }); + publicInternalCtorArgs.add({ varDecl, nullptr }); + publicCtorArgs.add({ varDecl, nullptr }); + break; + default: + // Unknown visibility + SLANG_ASSERT(false); + break; + } + } + + List generatedMemberwiseCtors; + if (maxVisibilityCtorToGenerate >= DeclVisibility::Public) + if (auto generatedCtor = _tryToGenerateCtorWithArgList(this, this->getASTBuilder(), std::move(publicCtorArgs), ctorList, structDecl, DeclVisibility::Public)) + generatedMemberwiseCtors.add(generatedCtor); + if (maxVisibilityCtorToGenerate >= DeclVisibility::Internal) + if (auto generatedCtor = _tryToGenerateCtorWithArgList(this, this->getASTBuilder(), std::move(publicInternalCtorArgs), ctorList, structDecl, DeclVisibility::Internal)) + generatedMemberwiseCtors.add(generatedCtor); + if (maxVisibilityCtorToGenerate >= DeclVisibility::Private) + if (auto generatedCtor = _tryToGenerateCtorWithArgList(this, this->getASTBuilder(), std::move(publicInternalPrivateCtorArgs), ctorList, structDecl, DeclVisibility::Private)) + generatedMemberwiseCtors.add(generatedCtor); + + // `checkIfCStyleStruct` must check after we add all possible Ctors. + // If we are a CStyleStruct add DefaultConstructExpr to all params (excluding arg 0) + bool isCStyleStruct = isCStyleStructDecl(this, structDecl, ctorList); + if (isCStyleStruct && generatedMemberwiseCtors.getCount() > 0) + { + // We know the user provided 0 non-default ctor's, we only had a chance to generate non default Ctors above in this AST pass. + SLANG_ASSERT(generatedMemberwiseCtors.getCount() == 1); + for (auto generatedCtor : generatedMemberwiseCtors) + { + Index paramIndex = 0; + for (ParamDecl* i : generatedCtor->getParameters()) + { + // Never annotate the first parameter to prevent conflict with "DefaultCtor" + if (paramIndex == 0) + { + paramIndex++; + continue; + } + paramIndex++; + i->initExpr = createDefaultConstructExprForType(this->getASTBuilder(), (QualType)i->type, i->loc); + } + } + } + } + } checkAggTypeConformance(aggTypeDecl); } @@ -6108,12 +6324,17 @@ namespace Slang // requests will be handled further down. For now we include // lookup results that might be usable, but not as-is. // - LookupResult lookupResult; - if (!isWrapperTypeDecl(context->parentDecl)) - { - lookupResult = lookUpMember(m_astBuilder, this, name, subType, nullptr, LookupMask::Default, LookupOptions::IgnoreBaseInterfaces); + LookupResult lookupResult = lookUpMember(m_astBuilder, this, name, subType, nullptr, LookupMask::Default, LookupOptions((uint8_t)LookupOptions::IgnoreBaseInterfaces)); - if (!lookupResult.isValid()) + if (!lookupResult.isValid()) + { + // If we failed to look up a member with the name of the + // requirement, it may be possible that we can still synthesis the + // implementation if this is one of the known builtin requirements. + // Otherwise, report diagnostic now. + if (!requiredMemberDeclRef.getDecl()->hasModifier() && + !(requiredMemberDeclRef.as() && + getInner(requiredMemberDeclRef.as())->hasModifier())) { // If we failed to look up a member with the name of the // requirement, it may be possible that we can still synthesis the @@ -8327,6 +8548,73 @@ namespace Slang return as(stmt->body); } + static bool _isSynthisizedMemberwiseCtor(ConstructorDecl* ctor) + { + return ctor->containsOption(ConstructorTags::Synthesized) + && ctor->getParameters().getCount() != 0 + && !allParamHaveInitExpr(ctor); + } + + template + bool _containsTargetType(ASTBuilder* m_astBuilder, NodeBase* target, HashSet& visitedNodes) + { + if (!target) + return false; + + if (as(target)) + return true; + + if (visitedNodes.contains(target)) + return false; + visitedNodes.add(target); + + if (auto aggTypeDecl = as(target)) + { + for (auto i : getMembersOfType(m_astBuilder, aggTypeDecl, MemberFilterStyle::Instance)) + if (_containsTargetType(m_astBuilder, i.getDecl()->type, visitedNodes)) + return true; + } + else if (auto genericDecl = as(target)) + { + for (auto i : genericDecl->getArgs()) + if (_containsTargetType(m_astBuilder, as(i), visitedNodes)) + return true; + return _containsTargetType(m_astBuilder, genericDecl->getBase(), visitedNodes); + } + else if (auto vectorExpr = as(target)) + { + return _containsTargetType(m_astBuilder, vectorExpr->getElementType(), visitedNodes); + } + else if (auto matrixExpr = as(target)) + { + return _containsTargetType(m_astBuilder, matrixExpr->getElementType(), visitedNodes); + } + else if (auto arrayExpr = as(target)) + { + return _containsTargetType(m_astBuilder, arrayExpr->getElementType(), visitedNodes); + } + else if (auto basicExpr = as(target)) + { + return _containsTargetType(m_astBuilder, basicExpr->getDeclRef(), visitedNodes); + } + else if (auto declRefType = as(target)) + { + return _containsTargetType(m_astBuilder, declRefType->getDeclRef(), visitedNodes); + } + else if (auto directDeclRef = as(target)) + { + return _containsTargetType(m_astBuilder, directDeclRef->getDecl(), visitedNodes); + } + return false; + } + + template + bool containsTargetType(ASTBuilder* m_astBuilder, NodeBase* target) + { + HashSet visitedNodes; + return _containsTargetType(m_astBuilder, target, visitedNodes); + } + void SemanticsDeclBodyVisitor::visitAggTypeDecl(AggTypeDecl* aggTypeDecl) { if (aggTypeDecl->hasTag(TypeTag::Incomplete) && aggTypeDecl->hasModifier()) @@ -8337,68 +8625,114 @@ namespace Slang auto structDecl = as(aggTypeDecl); if (!structDecl) return; + + // First part of auto-generating constructors/functions is inside `SemanticsDeclConformancesVisitor::visitAggTypeDecl` + // * Create declarations for constructors/functions. This is important so Slang may use the declarations before + // we have the chance to synthisize function bodies + // Second part of auto-generating constructors/functions is inside `SemanticsDeclBodyVisitor::visitAggTypeDecl` + // * Add bodies to our declarations. + // + // These are split up so we can assign declarations before we finish definitions + // + // + + auto structDeclType = DeclRefType::create(m_astBuilder, structDecl); + // Collect ctor info struct DeclAndCtorInfo { - StructDecl* parent = nullptr; - ConstructorDecl* defaultCtor = nullptr; - List ctorList; + AggTypeDecl* m_inheritanceBaseDecl; + InheritanceDecl* m_inheritanceDecl; + ConstructorDecl* m_defaultCtor = nullptr; + + struct CtorAndInsertOffset + { + ConstructorDecl* m_ctor; + Index m_insertOffset; + }; + + List m_ctorInfoList; DeclAndCtorInfo() { } - DeclAndCtorInfo(ASTBuilder* m_astBuilder, SemanticsVisitor* visitor, StructDecl* parent, const bool getOnlyDefault) + DeclAndCtorInfo(ASTBuilder* m_astBuilder, SemanticsVisitor* visitor, AggTypeDecl* inheritanceBaseDecl, InheritanceDecl* inheritanceDecl, const bool getOnlyDefault) : m_inheritanceBaseDecl(inheritanceBaseDecl), m_inheritanceDecl(inheritanceDecl) { + // only get ctor info from structDecl + StructDecl* structDecl = as(m_inheritanceBaseDecl); + if (!structDecl) + return; + if (getOnlyDefault) - defaultCtor = _getDefaultCtor(parent); + m_defaultCtor = _getDefaultCtor(structDecl); else - ctorList = _getCtorList(m_astBuilder, visitor, parent, &defaultCtor); + { + auto ctorList = _getCtorList(m_astBuilder, visitor, structDecl, &m_defaultCtor); + m_ctorInfoList.reserve(ctorList.getCount()); + for(auto i : ctorList) + m_ctorInfoList.add({i, 0}); + } } }; - List inheritanceDefaultCtorList{}; + + List inheritanceInfoList{}; for (auto inheritanceMember : structDecl->getMembersOfType()) { auto declRefType = as(inheritanceMember->base.type); if (!declRefType) continue; - auto structOfInheritance = as(declRefType->getDeclRef().getDecl()); - if (!structOfInheritance) + auto typeOfInheritance = as(declRefType->getDeclRef().getDecl()); + if (!typeOfInheritance) continue; - inheritanceDefaultCtorList.add(DeclAndCtorInfo(m_astBuilder, this, structOfInheritance, true)); + inheritanceInfoList.add(DeclAndCtorInfo(m_astBuilder, this, typeOfInheritance, inheritanceMember, false)); } - DeclAndCtorInfo structDeclInfo = DeclAndCtorInfo(m_astBuilder, this, structDecl, false); + DeclAndCtorInfo structDeclInfo = DeclAndCtorInfo(m_astBuilder, this, structDecl, nullptr, false); // ensure all varDecl members are processed up to SemanticsBodyVisitor so we can be sure that if init expressions // of members are to be synthisised, they are. - bool isDefaultInitializableType = isSubtype(DeclRefType::create(m_astBuilder, structDecl), m_astBuilder->getDefaultInitializableType(), IsSubTypeOptions::None); + bool isDefaultInitializableType = isSubtype(structDeclType, m_astBuilder->getDefaultInitializableType(), IsSubTypeOptions::None); for (auto m : structDecl->members) { auto varDeclBase = as(m); if (!varDeclBase) continue; - ensureDecl(m->getDefaultDeclRef(), DeclCheckState::DefaultConstructorReadyForUse); + ensureDecl(m->getDefaultDeclRef(), DeclCheckState::VarInitExprAreChecked); if (!isDefaultInitializableType || varDeclBase->initExpr) continue; - varDeclBase->initExpr = constructDefaultInitExprForVar(this, varDeclBase); + varDeclBase->initExpr = constructDefaultInitExprForVar(this, varDeclBase->type, varDeclBase); } - Index insertOffset = 0; Dictionary cachedDeclToCheckedVar; - for (auto ctor : structDeclInfo.ctorList) + + // Insert 'this->base->__init()' into 'this->__init()'. + /* + __init() { - auto seqStmt = _ensureCtorBodyIsSeqStmt(m_astBuilder, ctor); + super::__init(); // we add this call here + } + */ + for (auto& ctorInfo : structDeclInfo.m_ctorInfoList) + { + auto ctor = ctorInfo.m_ctor; + + // Synthesized constructors should not inject a defaultCtor of base-struct-type + // unless ctor has 0 parameters. This is to allow a ctor which desires to be + // init-list initialized to control which members initialize. + if(_isSynthisizedMemberwiseCtor(ctor)) + continue; + auto seqStmtChild = m_astBuilder->create(); - seqStmtChild->stmts.reserve(inheritanceDefaultCtorList.getCount()); - for (auto& declInfo : inheritanceDefaultCtorList) + seqStmtChild->stmts.reserve(inheritanceInfoList.getCount()); + for (auto& declInfo : inheritanceInfoList) { - if (!declInfo.defaultCtor) + if (!declInfo.m_defaultCtor) continue; auto ctorToInvoke = m_astBuilder->create(); - ctorToInvoke->declRef = declInfo.defaultCtor->getDefaultDeclRef(); - ctorToInvoke->name = declInfo.defaultCtor->getName(); - ctorToInvoke->loc = declInfo.defaultCtor->loc; - ctorToInvoke->type = m_astBuilder->getFuncType(ArrayView(), structDeclInfo.defaultCtor->returnType.type); + ctorToInvoke->declRef = declInfo.m_defaultCtor->getDefaultDeclRef(); + ctorToInvoke->name = declInfo.m_defaultCtor->getName(); + ctorToInvoke->loc = declInfo.m_defaultCtor->loc; + ctorToInvoke->type = m_astBuilder->getFuncType(ArrayView(), structDeclInfo.m_defaultCtor->returnType.type); auto invoke = m_astBuilder->create(); invoke->functionExpr = ctorToInvoke; @@ -8407,8 +8741,13 @@ namespace Slang thisExpr->scope = ctor->ownedScope; thisExpr->type = ctor->returnType.type; + // A base may not have any value, if this is the case do not insert a `__init()` + // since coercing is impossible. + if (!canCoerce(declInfo.m_defaultCtor->returnType.type, ctor->returnType.type, thisExpr)) + continue; + auto assign = m_astBuilder->create(); - assign->left = coerce(CoercionSite::Initializer, declInfo.defaultCtor->returnType.type, thisExpr); + assign->left = coerce(CoercionSite::Initializer, declInfo.m_defaultCtor->returnType.type, thisExpr); assign->right = invoke; auto stmt = m_astBuilder->create(); stmt->expression = assign; @@ -8420,53 +8759,68 @@ namespace Slang if (seqStmtChild->stmts.getCount() == 0) continue; - seqStmt->stmts.insert(0, seqStmtChild); - insertOffset = 1; + auto seqStmt = _ensureCtorBodyIsSeqStmt(m_astBuilder, ctor); + seqStmt->stmts.insert(ctorInfo.m_insertOffset++, seqStmtChild); } - for (auto ctor : structDeclInfo.ctorList) + // Collect all 'per instance' members of our struct + List> membersOfStructDeclInstance; + for (auto i : getMembersOfType(m_astBuilder, structDecl, MemberFilterStyle::Instance)) + membersOfStructDeclInstance.add(i); + + // Assign member variable init expressions + /* + struct Foo { + int a = 5; + __init() + { + a = 5; // we add this initExpr logic here. + } + } + */ + for (auto& ctorInfo : structDeclInfo.m_ctorInfoList) + { + auto ctor = ctorInfo.m_ctor; + ThisExpr* thisExpr = m_astBuilder->create(); thisExpr->scope = ctor->ownedScope; thisExpr->type = ctor->returnType.type; - auto seqStmt = _ensureCtorBodyIsSeqStmt(m_astBuilder, ctor); auto seqStmtChild = m_astBuilder->create(); seqStmtChild->stmts.reserve(structDecl->members.getCount()); - for (auto& m : structDecl->members) - { - auto varDeclBase = as(m); - // Static variables are initialized at start of runtime, not inside a constructor - if (!varDeclBase - || !varDeclBase->initExpr - || varDeclBase->hasModifier()) + for (auto varDeclBaseRef : membersOfStructDeclInstance) + { + auto varDeclBase = varDeclBaseRef.getDecl(); + Expr* initExpr = varDeclBase->initExpr; + if (!initExpr) continue; MemberExpr* memberExpr = m_astBuilder->create(); memberExpr->baseExpression = thisExpr; - memberExpr->declRef = m->getDefaultDeclRef(); + memberExpr->declRef = varDeclBase->getDefaultDeclRef(); memberExpr->scope = ctor->ownedScope; - memberExpr->loc = m->loc; - memberExpr->name = m->getName(); - memberExpr->type = DeclRefType::create(getASTBuilder(), m->getDefaultDeclRef()); + memberExpr->loc = varDeclBase->loc; + memberExpr->name = varDeclBase->getName(); + memberExpr->type = DeclRefType::create(getASTBuilder(), varDeclBase->getDefaultDeclRef()); auto assign = m_astBuilder->create(); assign->left = memberExpr; - assign->right = varDeclBase->initExpr; - assign->loc = m->loc; + assign->right = initExpr; + assign->loc = varDeclBase->loc; auto stmt = m_astBuilder->create(); stmt->expression = assign; - stmt->loc = m->loc; + stmt->loc = varDeclBase->loc; Expr* checkedMemberVarExpr; - if (cachedDeclToCheckedVar.containsKey(m)) - checkedMemberVarExpr = cachedDeclToCheckedVar[m]; + if (cachedDeclToCheckedVar.containsKey(varDeclBase)) + checkedMemberVarExpr = cachedDeclToCheckedVar[varDeclBase]; else { checkedMemberVarExpr = CheckTerm(memberExpr); - cachedDeclToCheckedVar.add({ m, checkedMemberVarExpr }); + cachedDeclToCheckedVar.add({ varDeclBase, checkedMemberVarExpr }); } if (!checkedMemberVarExpr->type.isLeftValue) continue; @@ -8475,18 +8829,306 @@ namespace Slang } if (seqStmtChild->stmts.getCount() == 0) continue; - seqStmt->stmts.insert(insertOffset, seqStmtChild); + + auto seqStmt = _ensureCtorBodyIsSeqStmt(m_astBuilder, ctor); + seqStmt->stmts.insert(ctorInfo.m_insertOffset++, seqStmtChild); } - if (structDeclInfo.defaultCtor) + // Note: we assume only 1 inheritance decl currently. + // Pre-calculate if we have a base-type and its associated ctor-list + // We also must ensure inheritance-decl is checked, else we may not have up-to-date 'DerivativeMemberAttribute' modifiers + DeclAndCtorInfo* baseStructInfo = nullptr; + for(auto& i : inheritanceInfoList) { - auto seqStmt = as(as(structDeclInfo.defaultCtor->body)->body); - if (seqStmt && seqStmt->stmts.getCount() == 0) + ensureDecl(i.m_inheritanceDecl, DeclCheckState::VarInitExprAreChecked); + if(as(i.m_inheritanceBaseDecl)) + baseStructInfo = &i; + } + + // pre-calculate any requirements of a CudaHostAttribute + HashSet requiresCudaHostModifier; + for (auto member : membersOfStructDeclInstance) + if (containsTargetType(m_astBuilder, member.getDecl()->type.type)) + requiresCudaHostModifier.add(member.getDecl()); + auto addCudaHostModifierIfRequired = [&](FunctionDeclBase* func, VarDeclBase* member, bool& foundCudaHostModifier) { - structDecl->members.remove(structDeclInfo.defaultCtor); - structDecl->invalidateMemberDictionary(); - structDecl->buildMemberDictionary(); + // Manage CUDA host modifier based on inheritance + if (!foundCudaHostModifier && requiresCudaHostModifier.contains(member)) + { + foundCudaHostModifier = true; + addModifier(func, m_astBuilder->create()); + } + }; + + + + // Insert parameters as values for member-wise init expression. + for (auto& ctorInfo : structDeclInfo.m_ctorInfoList) + { + auto ctor = ctorInfo.m_ctor; + + if(!_isSynthisizedMemberwiseCtor(ctor)) + continue; + + auto ctorVisibility = getDeclVisibility(ctor); + + Index paramIndex = 0; + auto paramList = ctor->getParameters(); + + Index memberIndex = 0; + auto members = membersOfStructDeclInstance; + + auto seqStmt = _ensureCtorBodyIsSeqStmt(m_astBuilder, ctor); + auto seqStmtChild = m_astBuilder->create(); + + ThisExpr* thisExpr = m_astBuilder->create(); + thisExpr->scope = ctor->ownedScope; + thisExpr->type = ctor->returnType.type; + + auto paramCount = paramList.getCount(); + + // ensure synth'ed ctor has correct modifier to tag as a "CUDA host function" + bool foundCudaHostModifier = false; + while (paramIndex < paramCount) + { + auto param = paramList[paramIndex]; + /* + 2 parts: + + 1. Insert super::__init inside this->__init() + 2. Assign *parameters* to related *varDecl member of our struct* + */ + + // Part 1 + // If we have a base type, the first arg is a 'base->__init(...)'. We need to find this 'base->__init(...)' + // and assign the parameters needed to call '__init(...)' + if (paramIndex == 0 && baseStructInfo && baseStructInfo->m_ctorInfoList.getCount() > 0) + { + auto baseStruct = baseStructInfo->m_inheritanceBaseDecl; + + // First find a member-wise 'base->__init()' which maps to this current ctor being filled + ConstructorDecl* baseCtor = nullptr; + ConstructorTags memberwiseCtorToCall = ConstructorTags::MemberwiseCtorForPublicVisibility; + if (isVisibilityOfDeclVisibleInScope(baseStruct, DeclVisibility::Internal, structDecl->ownedScope)) + memberwiseCtorToCall = ConstructorTags::MemberwiseCtorForInternalVisibility; + for (auto& i : baseStructInfo->m_ctorInfoList) + { + if (i.m_ctor->containsOption(memberwiseCtorToCall)) + { + baseCtor = i.m_ctor; + break; + } + } + if (baseCtor) + { + // Manage CUDA host modifier based on inheritance + if (baseCtor->findModifier()) + { + foundCudaHostModifier = true; + addModifier(ctor, m_astBuilder->create()); + } + + auto baseCtorParamCount = baseCtor->getParameters().getCount(); + + // Now assign the found ctor and add it to our auto-synthisized ctor + auto ctorToInvokeExpr = m_astBuilder->create(); + ctorToInvokeExpr->declRef = baseCtor->getDefaultDeclRef(); + ctorToInvokeExpr->name = baseCtor->getName(); + ctorToInvokeExpr->loc = baseCtor->loc; + ctorToInvokeExpr->type = m_astBuilder->getFuncType(ArrayView(), baseCtor->returnType.type); + + auto invoke = m_astBuilder->create(); + invoke->functionExpr = ctorToInvokeExpr; + for (; paramIndex < baseCtorParamCount; paramIndex++) + { + auto paramToAdd = paramList[paramIndex]; + auto paramType = paramToAdd->getType(); + auto paramExpr = m_astBuilder->create(); + paramExpr->scope = ctor->ownedScope; + paramExpr->declRef = paramToAdd; + paramExpr->type = paramType; + paramExpr->loc = paramToAdd->loc; + + invoke->arguments.add(paramExpr); + } + + auto assign = m_astBuilder->create(); + assign->left = coerce(CoercionSite::Initializer, baseCtor->returnType.type, thisExpr); + assign->right = invoke; + auto stmt = m_astBuilder->create(); + stmt->expression = assign; + stmt->loc = ctor->loc; + seqStmtChild->stmts.add(stmt); + + continue; + } + } + + // Part 2 + // Regular assignment logic + paramIndex++; + auto paramType = param->getType(); + auto paramExpr = m_astBuilder->create(); + paramExpr->scope = ctor->ownedScope; + paramExpr->declRef = param; + paramExpr->type = paramType; + paramExpr->loc = param->loc; + + // Do not map a variable which is not visible + // Do not map a read-only variable for default-init + // Do not map compiler generated 'internal processing' variable + VarDeclBase* member = members[memberIndex++].getDecl(); + while (getDeclVisibility(member) < ctorVisibility + || !getTypeForDeclRef(m_astBuilder, member, member->loc).isLeftValue + || member->getName() && member->getName()->text[0] == '$' + ) + { + // Should note be possible to be out of range unless compiler generated a synth-ctor wrong. + SLANG_ASSERT(memberIndex < members.getCount()); + member = members[memberIndex++].getDecl(); + } + + // We need to ensure member is `no_diff` if it cannot be differentiated, `ctor` modifiers do not matter + // in this case since member-wise ctor is always differentiable or "treat as differentiable". + if (!isTypeDifferentiable(member->getType()) || member->hasModifier()) + { + auto noDiffMod = m_astBuilder->create(); + noDiffMod->loc = param->loc; + addModifier(param, noDiffMod); + } + + addCudaHostModifierIfRequired(ctor, member, foundCudaHostModifier); + + MemberExpr* memberExpr = m_astBuilder->create(); + memberExpr->baseExpression = thisExpr; + memberExpr->declRef = member->getDefaultDeclRef(); + memberExpr->scope = ctor->ownedScope; + memberExpr->loc = member->loc; + memberExpr->name = member->getName(); + memberExpr->type = DeclRefType::create(getASTBuilder(), memberExpr->declRef); + + auto assign = m_astBuilder->create(); + assign->left = memberExpr; + assign->right = paramExpr; + assign->loc = paramExpr->loc; + + auto stmt = m_astBuilder->create(); + stmt->expression = assign; + stmt->loc = assign->loc; + + seqStmtChild->stmts.add(stmt); } + seqStmt->stmts.insert(ctorInfo.m_insertOffset++, seqStmtChild); + } + + // Fill in our $DefaultInit + if (auto defaultInitListFunc = findDefaultInitListFunc(structDecl)) + { + SLANG_ASSERT(defaultInitListFunc->getParameters().getCount() == 0); + SLANG_ASSERT(defaultInitListFunc->findModifier()); + + SLANG_ASSERT(as(defaultInitListFunc->body)); + auto block = as(defaultInitListFunc->body); + SLANG_ASSERT(as(block->body)); + auto seqStmt = as(block->body); + + /* + 3 steps: + 1. Assign DefaultConstructExpr to `this` + 2. Insert `super::$DefaultInit()` into `this->$DefaultInit` + 3. Assign initExpr to all members + */ + + bool foundCudaHostModifier = false; + + auto defaultConstructExpr = createDefaultConstructExprForType(m_astBuilder, structDeclType, seqStmt->loc); + + // Part 1 + // Assign DefaultConstructExpr to `this` + auto structVarDecl = m_astBuilder->create(); + addModifier(structVarDecl, m_astBuilder->create()); + structVarDecl->type = TypeExp(structDeclType); + structVarDecl->initExpr = defaultConstructExpr; + structVarDecl->parentDecl = defaultInitListFunc; + structVarDecl->loc = seqStmt->loc; + ensureDecl(structVarDecl, DeclCheckState::VarInitExprAreChecked); + + auto structVarDeclExpr = m_astBuilder->create(); + structVarDeclExpr->declRef = structVarDecl->getDefaultDeclRef(); + structVarDeclExpr->name = structVarDecl->getName(); + structVarDeclExpr->loc = structVarDecl->loc; + structVarDeclExpr->type = GetTypeForDeclRef(structVarDecl, structVarDecl->loc); + structVarDeclExpr->type.isLeftValue = true; + structVarDeclExpr->scope = block->scopeDecl->ownedScope; + + auto structAssignExpr = m_astBuilder->create(); + structAssignExpr->left = structVarDeclExpr; + structAssignExpr->right = defaultConstructExpr; + auto structAssignExprStmt = m_astBuilder->create(); + structAssignExprStmt->expression = structAssignExpr; + structAssignExprStmt->loc = seqStmt->loc; + seqStmt->stmts.add(structAssignExprStmt); + + // Part 2 + // Insert `super::$DefaultInit()` into `this->$DefaultInit` + // Only insert if we have values to insert for + if (baseStructInfo) + { + if (auto baseStructDecl = as(baseStructInfo->m_inheritanceBaseDecl)) + { + if (_structHasMemberWithValue(getASTBuilder(), structDecl)) + { + Expr* defaultInitFunc = constructDefaultInitListFunc(this, baseStructDecl, baseStructInfo->m_inheritanceDecl->base.type, ConstructDefaultInitListOptions::PreferDefaultInitFunc); + auto assign = m_astBuilder->create(); + assign->left = coerce(CoercionSite::Initializer, defaultInitFunc->type, structVarDeclExpr); + assign->right = defaultInitFunc; + auto stmt = m_astBuilder->create(); + stmt->expression = assign; + stmt->loc = seqStmt->loc; + seqStmt->stmts.add(stmt); + } + } + } + + // Part 3 + // Assign initExpr to all members of structDecl + for (auto memberRef : membersOfStructDeclInstance) + { + + auto member = memberRef.getDecl(); + addCudaHostModifierIfRequired(defaultInitListFunc, member, foundCudaHostModifier); + + auto initExpr = member->initExpr; + if (!initExpr) + continue; + + MemberExpr* memberExpr = m_astBuilder->create(); + memberExpr->baseExpression = structVarDeclExpr; + memberExpr->declRef = member->getDefaultDeclRef(); + memberExpr->scope = defaultInitListFunc->ownedScope; + memberExpr->loc = member->loc; + memberExpr->name = member->getName(); + memberExpr->type = GetTypeForDeclRef(member, member->loc); + auto checkedMemberVarExpr = CheckTerm(memberExpr); + if (!checkedMemberVarExpr->type.isLeftValue) + continue; + + auto assign = m_astBuilder->create(); + assign->left = checkedMemberVarExpr; + assign->right = initExpr; + assign->loc = checkedMemberVarExpr->loc; + + auto stmt = m_astBuilder->create(); + stmt->expression = assign; + stmt->loc = assign->loc; + seqStmt->stmts.add(stmt); + } + + auto returnStmt = m_astBuilder->create(); + returnStmt->loc = structVarDeclExpr->loc; + returnStmt->expression = structVarDeclExpr; + + seqStmt->stmts.add(returnStmt); } } @@ -10732,12 +11374,6 @@ namespace Slang void SemanticsDeclAttributesVisitor::visitStructDecl(StructDecl* structDecl) { - // add a empty deault CTor if missing; checking in attributes - // to avoid circular checking logic - auto defaultCtor = _getDefaultCtor(structDecl); - if (!defaultCtor) - _createCtor(this, m_astBuilder, structDecl); - int backingWidth = 0; [[maybe_unused]] int totalWidth = 0; diff --git a/source/slang/slang-check-expr.cpp b/source/slang/slang-check-expr.cpp index 5233008fda..28ca8c98cc 100644 --- a/source/slang/slang-check-expr.cpp +++ b/source/slang/slang-check-expr.cpp @@ -887,9 +887,8 @@ namespace Slang return DeclVisibility::Public; } - bool SemanticsVisitor::isDeclVisibleFromScope(DeclRef declRef, Scope* scope) + bool SemanticsVisitor::isVisibilityOfDeclVisibleInScope(DeclRef declRef, DeclVisibility visibility, Scope* scope) { - auto visibility = getDeclVisibility(declRef.getDecl()); if (visibility == DeclVisibility::Public) return true; if (visibility == DeclVisibility::Internal) @@ -918,7 +917,17 @@ namespace Slang } return false; } - return false; + return false; + } + + bool SemanticsVisitor::isDeclVisibleFromScope(DeclRef declRef, Scope* scope) + { + return isVisibilityOfDeclVisibleInScope(declRef, getDeclVisibility(declRef.getDecl()), scope); + } + + bool SemanticsVisitor::isDeclVisible(DeclRef declRef) + { + return isDeclVisibleFromScope(declRef, m_outerScope); } LookupResult SemanticsVisitor::filterLookupResultByVisibility(const LookupResult& lookupResult) @@ -2652,7 +2661,7 @@ namespace Slang if (auto varExpr = as(expr->functionExpr)) { - if ((varExpr->name->text == "&&") || (varExpr->name->text == "||")) + if (varExpr->name && (varExpr->name->text == "&&" || varExpr->name->text == "||")) { // We only use short-circuiting in scalar input, will fall back // to non-short-circuiting in vector input. @@ -2802,9 +2811,13 @@ namespace Slang } else { - getSink()->diagnose( - m_treatAsDifferentiableExpr, - Diagnostics::useOfNoDiffOnDifferentiableFunc); + // We should only error if our TreatAsDifferentiableExpr::Flavor is NoDiff. + if (m_treatAsDifferentiableExpr->flavor == TreatAsDifferentiableExpr::Flavor::NoDiff) + { + getSink()->diagnose( + m_treatAsDifferentiableExpr, + Diagnostics::useOfNoDiffOnDifferentiableFunc); + } } } } diff --git a/source/slang/slang-check-impl.h b/source/slang/slang-check-impl.h index 63c16e6d23..b1a38c9eca 100644 --- a/source/slang/slang-check-impl.h +++ b/source/slang/slang-check-impl.h @@ -740,6 +740,15 @@ namespace Slang m_mapTypePairToImplicitCastMethod[key] = candidate; } + void cacheIsCStyleStruct(StructDecl* structDecl, bool isCStyleStruct) + { + m_isCStyleStruct[structDecl] = isCStyleStruct; + } + bool* tryGetIsCStyleStructFromCache(StructDecl* structDecl) + { + return m_isCStyleStruct.tryGetValue(structDecl); + } + // Get the inner most generic decl that a decl-ref is dependent on. // For example, `Foo` depends on the generic decl that defines `T`. // @@ -865,6 +874,7 @@ namespace Slang Dictionary, InheritanceInfo> m_mapDeclRefToInheritanceInfo; Dictionary m_mapTypePairToSubtypeWitness; Dictionary m_mapTypePairToImplicitCastMethod; + Dictionary m_isCStyleStruct; }; /// Local/scoped state of the semantic-checking system @@ -1256,7 +1266,9 @@ namespace Slang Expr* originalExpr); DeclVisibility getTypeVisibility(Type* type); + bool isVisibilityOfDeclVisibleInScope(DeclRef declRef, DeclVisibility visibility, Scope* scope); bool isDeclVisibleFromScope(DeclRef declRef, Scope* scope); + bool isDeclVisible(DeclRef declRef); LookupResult filterLookupResultByVisibility(const LookupResult& lookupResult); LookupResult filterLookupResultByVisibilityAndDiagnose(const LookupResult& lookupResult, SourceLoc loc, bool& outDiagnosed); @@ -1495,11 +1507,12 @@ namespace Slang /// If the routine fails and `outToExpr` is non-null, /// then a suitable diagnostic will be emitted. /// + template bool _readValueFromInitializerList( Type* toType, Expr** outToExpr, InitializerListExpr* fromInitializerListExpr, - UInt &ioInitArgIndex); + UInt &ioInitArgIndex); /// Read an aggregate value from an initializer list expression. /// @@ -1520,6 +1533,7 @@ namespace Slang /// If the routine fails and `outToExpr` is non-null, /// then a suitable diagnostic will be emitted. /// + template bool _readAggregateValueFromInitializerList( Type* inToType, Expr** outToExpr, @@ -2827,12 +2841,12 @@ namespace Slang // deal with this cases here, even if they are no-ops. // + // Do not error and just return since a term may + // be checked so it can be resolved into a valid + // term (argument of an 'Invoke' for example) #define CASE(NAME) \ Expr* visit##NAME(NAME* expr) \ { \ - if (!getShared()->isInLanguageServer()) \ - SLANG_DIAGNOSE_UNEXPECTED(getSink(), expr, "should not appear in input syntax"); \ - expr->type = m_astBuilder->getErrorType(); \ return expr; \ } @@ -3005,4 +3019,36 @@ namespace Slang FrontEndEntryPointRequest* entryPointReq); bool resolveStageOfProfileWithEntryPoint(Profile& entryPointProfile, CompilerOptionSet& optionSet, const List>& targets, FuncDecl* entryPointFuncDecl, DiagnosticSink* sink); + + LookupResult lookUpMember( + ASTBuilder* astBuilder, + SemanticsVisitor* semantics, + Name* name, + Type* type, + Scope* sourceScope, + LookupMask mask, + LookupOptions options); + + ConstructorDecl* _getDefaultCtor(StructDecl* structDecl); + bool allParamHaveInitExpr(ConstructorDecl* ctor); + List _getCtorList(ASTBuilder* m_astBuilder, SemanticsVisitor* visitor, StructDecl* structDecl, ConstructorDecl** defaultCtorOut); + bool DiagnoseIsAllowedInitExpr(VarDeclBase* varDecl, DiagnosticSink* sink); + bool isDefaultInitializable(Type* varDeclType, VarDeclBase* associatedDecl); + Expr* constructDefaultInitExprForVar(SemanticsVisitor* visitor, TypeExp varDeclType, VarDeclBase* decl); + + enum class ConstructDefaultInitListOptions : UInt + { + None = 0 << 0, + PreferDefaultInitFunc = 1 << 0, + CheckToAvoidRecursion = 1 << 2, + }; + Expr* constructDefaultInitListFunc(SemanticsVisitor* visitor, StructDecl* structDecl, Type* structDeclType, ConstructDefaultInitListOptions options); + FuncDecl* findDefaultInitListFunc(StructDecl* structDecl); + + bool isCStyleStructDecl(SemanticsVisitor* visitor, StructDecl* decl, List const& ctorList); + + DefaultConstructExpr* createDefaultConstructExprForType(ASTBuilder* m_astBuilder, QualType type, SourceLoc loc); + + DeclRefBase* _getDeclRefFromVal(Val* val); + } diff --git a/source/slang/slang-check-overload.cpp b/source/slang/slang-check-overload.cpp index c738f5b49e..13dcb80de8 100644 --- a/source/slang/slang-check-overload.cpp +++ b/source/slang/slang-check-overload.cpp @@ -1508,6 +1508,11 @@ namespace Slang OverloadCandidate* left, OverloadCandidate* right) { + // If candidates are equal (which is possible if we have an overload from `AssocType` and `StructDecl`) + // We need to pick 1 to keep. Overlapping lookups with `__init()` are common with auto-diff. + if (left->item.declRef == right->item.declRef) + return -1; + // If one candidate got further along in validation, pick it if (left->status != right->status) return int(right->status) - int(left->status); diff --git a/source/slang/slang-constructor-utility.cpp b/source/slang/slang-constructor-utility.cpp new file mode 100644 index 0000000000..e4ce2a8591 --- /dev/null +++ b/source/slang/slang-constructor-utility.cpp @@ -0,0 +1,276 @@ +// slang-constructor-utility.cpp +#include "slang-check-impl.h" + +namespace Slang +{ + + DefaultConstructExpr* createDefaultConstructExprForType(ASTBuilder* m_astBuilder, QualType type, SourceLoc loc) + { + auto defaultConstructExpr = m_astBuilder->create(); + defaultConstructExpr->type = type; + defaultConstructExpr->loc = loc; + return defaultConstructExpr; + } + + bool allParamHaveInitExpr(ConstructorDecl* ctor) + { + for (auto i : ctor->getParameters()) + if (!i->initExpr) + return false; + return true; + } + + static inline bool _isDefaultCtor(ConstructorDecl* ctor) + { + // 1. default ctor must have definition + // 2. either 2.1 or 2.2 is safisfied + // 2.1. default ctor must have no parameters + // 2.2. default ctor can have parameters, but all parameters have init expr (Because we won't differentiate this case from 2.) + if (!ctor->body || (ctor->members.getCount() != 0 && !allParamHaveInitExpr(ctor))) + return false; + + return true; + } + + ConstructorDecl* _getDefaultCtor(StructDecl* structDecl) + { + for (auto ctor : structDecl->getMembersOfType()) + { + if (_isDefaultCtor(ctor)) + return ctor; + } + return nullptr; + } + + List _getCtorList(ASTBuilder* m_astBuilder, SemanticsVisitor* visitor, StructDecl* structDecl, ConstructorDecl** defaultCtorOut) + { + List ctorList; + + auto ctorLookupResult = lookUpMember( + m_astBuilder, + visitor, + m_astBuilder->getSharedASTBuilder()->getCtorName(), + DeclRefType::create(m_astBuilder, structDecl), + structDecl->ownedScope, + LookupMask::Function, + (LookupOptions)((uint8_t)LookupOptions::IgnoreInheritance | (uint8_t)LookupOptions::IgnoreBaseInterfaces | (uint8_t)LookupOptions::NoDeref)); + + if (!ctorLookupResult.isValid()) + return ctorList; + + auto lookupResultHandle = [&](LookupResultItem& item) + { + auto ctor = as(item.declRef.getDecl()); + if (!ctor) + return; + ctorList.add(ctor); + if (ctor->members.getCount() != 0 + && !allParamHaveInitExpr(ctor) + || !defaultCtorOut) + return; + *defaultCtorOut = ctor; + }; + if (ctorLookupResult.items.getCount() == 0) + { + lookupResultHandle(ctorLookupResult.item); + return ctorList; + } + + for (auto m : ctorLookupResult) + { + lookupResultHandle(m); + } + + return ctorList; + } + + bool isDefaultInitializable(Type* varDeclType, VarDeclBase* associatedDecl) + { + if (!DiagnoseIsAllowedInitExpr(associatedDecl, nullptr)) + return false; + + // Find struct and modifiers associated with varDecl + StructDecl* structDecl = nullptr; + if (auto declRefType = as(varDeclType)) + { + if (auto genericAppRefDecl = as(declRefType->getDeclRefBase())) + { + auto baseGenericRefType = genericAppRefDecl->getBase()->getDecl(); + if (auto baseTypeStruct = as(baseGenericRefType)) + { + structDecl = baseTypeStruct; + } + else if (auto genericDecl = as(baseGenericRefType)) + { + if (auto innerTypeStruct = as(genericDecl->inner)) + structDecl = innerTypeStruct; + } + } + else + { + structDecl = as(declRefType->getDeclRef().getDecl()); + } + } + if (structDecl) + { + // find if a type is non-copyable + if (structDecl->findModifier()) + return false; + } + + return true; + } + + Expr* constructDefaultInitExprForVar(SemanticsVisitor* visitor, TypeExp varDeclType, VarDeclBase* decl) + { + if (!varDeclType || !varDeclType.type) + return nullptr; + + if (!isDefaultInitializable(varDeclType.type, decl)) + return nullptr; + + ConstructorDecl* defaultCtor = nullptr; + auto declRefType = as(varDeclType.type); + if (declRefType) + { + if (auto structDecl = as(declRefType->getDeclRef().getDecl())) + { + defaultCtor = _getDefaultCtor(structDecl); + } + } + + if (defaultCtor) + { + auto* invoke = visitor->getASTBuilder()->create(); + auto member = visitor->getASTBuilder()->getMemberDeclRef(declRefType->getDeclRef(), defaultCtor); + invoke->functionExpr = visitor->ConstructDeclRefExpr(member, nullptr, defaultCtor->getName(), defaultCtor->loc, nullptr); + invoke->type = varDeclType.type; + return invoke; + } + else + { + return createDefaultConstructExprForType(visitor->getASTBuilder(), QualType(varDeclType.type), {}); + } + } + + FuncDecl* findDefaultInitListFunc(StructDecl* structDecl) + { + for (auto funcDecl : structDecl->getMembersOfType()) + { + if (!funcDecl->findModifier()) + continue; + return funcDecl; + } + return nullptr; + } + + Expr* constructDefaultCtorInvocationExpr(SemanticsVisitor* visitor, StructDecl* structDecl, Type* structDeclType, ConstructorDecl* defaultCtor) + { + auto* invoke = visitor->getASTBuilder()->create(); + auto member = visitor->getASTBuilder()->getMemberDeclRef(structDecl->getDefaultDeclRef(), defaultCtor); + invoke->functionExpr = visitor->ConstructDeclRefExpr(member, nullptr, defaultCtor->getName(), defaultCtor->loc, nullptr); + invoke->type = structDeclType; + return invoke; + } + Expr* constructDefaultInitListFunc(SemanticsVisitor* visitor, StructDecl* structDecl, Type* structDeclType, ConstructDefaultInitListOptions options) + { + SLANG_ASSERT(structDecl); + + // 1. Prefer non-synth default-ctor + // * Skip this option if `ConstructDefaultInitListOptions::PreferDefaultInitFunc` is true + // * Skip this option if `ConstructDefaultInitListOptions::CheckToAvoidRecursion` detects recursion + // * Only user-defined ctor will try and have recursion of `{}` + // 2. Prefer $DefaultInit + // 3. Prefer any default-ctor + // 4. Use `DefaultConstructExpr` + + // 1. + auto defaultCtor = _getDefaultCtor(structDecl); + if(defaultCtor + && !defaultCtor->containsOption(ConstructorTags::Synthesized) + && !((UInt)options & (UInt)ConstructDefaultInitListOptions::PreferDefaultInitFunc)) + { + bool canCreateCtor = true; + if(((UInt)options & (UInt)ConstructDefaultInitListOptions::CheckToAvoidRecursion)) + { + auto callingScope = visitor->getOuterScope(); + while (callingScope) + { + if (callingScope->containerDecl == defaultCtor) + { + canCreateCtor = false; + break; + } + callingScope = callingScope->parent; + } + } + if(canCreateCtor) + return constructDefaultCtorInvocationExpr(visitor, structDecl, structDeclType, defaultCtor); + } + + // 2. + if (auto defaultInitFunc = findDefaultInitListFunc(structDecl)) + { + auto* invoke = visitor->getASTBuilder()->create(); + DeclRef member; + auto declRefType = as(structDeclType); + if(declRefType && as(declRefType->getDeclRefBase())) + member = visitor->getASTBuilder()->getMemberDeclRef(as(declRefType->getDeclRefBase()), defaultInitFunc); + else + member = visitor->getASTBuilder()->getMemberDeclRef(structDecl, defaultInitFunc); + + invoke->functionExpr = visitor->ConstructDeclRefExpr(member, nullptr, defaultInitFunc->getName(), defaultInitFunc->loc, nullptr); + invoke->type = structDeclType; + return invoke; + } + + // 3. + if (defaultCtor) + return constructDefaultCtorInvocationExpr(visitor, structDecl, structDeclType, defaultCtor); + + // 4. + return createDefaultConstructExprForType(visitor->getASTBuilder(), QualType(structDeclType), {}); + } + + bool isCStyleStructDecl(SemanticsVisitor* visitor, StructDecl* structDecl, List const& ctorList) + { + // CStyleStruct follows the following rules: + // 1. Does not contain a non 'Synthesized' Ctor (excluding 'DefaultCtor') + // + // 2. Only contains 1 'non-default' ctor regardless of synthisis or not, else + // `__init(int, int)` and `__init(int)` would have ambiguity for + // c-style-initialization of `MyStruct[3] tmp = {1,2, 1,2, 1,2};` + // + // 3. Every `VarDeclBase*` member has the same visibility + + auto isCStyleStruct = visitor->getShared()->tryGetIsCStyleStructFromCache(structDecl); + + if (isCStyleStruct) + return *isCStyleStruct; + + // Add to IsCStyleStruct cache + int nonDefaultInitCount = 0; + for (auto i : ctorList) + { + // Default ctor is always fine + if (i->getParameters().getCount() == 0) + continue; + + // Cannot contain user defined ctor which is a non default ctor + if (!i->containsOption(ConstructorTags::Synthesized)) + { + visitor->getShared()->cacheIsCStyleStruct(structDecl, false); + return false; + } + + // Cannot contain 2+ non-default init's + nonDefaultInitCount++; + if (nonDefaultInitCount > 1) + { + visitor->getShared()->cacheIsCStyleStruct(structDecl, false); + return false; + } + } + return true; + } +} diff --git a/source/slang/slang-diagnostic-defs.h b/source/slang/slang-diagnostic-defs.h index 3a14441178..28ff0c69ed 100644 --- a/source/slang/slang-diagnostic-defs.h +++ b/source/slang/slang-diagnostic-defs.h @@ -522,11 +522,12 @@ DIAGNOSTIC(30401, Error, invalidTypeForConstraint, "type '$0' cannot be used as DIAGNOSTIC(30402, Error, invalidConstraintSubType, "type '$0' is not a valid left hand side of a type constraint.") // 305xx: initializer lists -DIAGNOSTIC(30500, Error, tooManyInitializers, "too many initializers (expected $0, got $1)") +DIAGNOSTIC(30500, Error, cannotFindMatchingInitListToConstructorCount, "cannot find matching constructor to call with arguments count of '$0' for '$1'") DIAGNOSTIC(30501, Error, cannotUseInitializerListForArrayOfUnknownSize, "cannot use initializer list for array of statically unknown size '$0'") DIAGNOSTIC(30502, Error, cannotUseInitializerListForVectorOfUnknownSize, "cannot use initializer list for vector of statically unknown size '$0'") DIAGNOSTIC(30503, Error, cannotUseInitializerListForMatrixOfUnknownSize, "cannot use initializer list for matrix of statically unknown size '$0' rows") DIAGNOSTIC(30504, Error, cannotUseInitializerListForType, "cannot use initializer list for type '$0'") +DIAGNOSTIC(30505, Warning, usingLegacyInitListStructLogic, "detected deprecated legacy initializer-list (`{...}`) resolution logic, we advise following proper initializer-list rules as per Slang's user-guide") // 3062x: variables DIAGNOSTIC(30620, Error, varWithoutTypeMustHaveInitializer, "a variable declaration without an initial-value expression must be given an explicit type") @@ -662,7 +663,7 @@ DIAGNOSTIC(38027, Error, mismatchExistentialSlotArgCount, "expected $0 existenti DIAGNOSTIC(38029, Error, typeArgumentDoesNotConformToInterface, "type argument '$0' does not conform to the required interface '$1'") DIAGNOSTIC(38031, Error, invalidUseOfNoDiff, "'no_diff' can only be used to decorate a call or a subscript operation") -DIAGNOSTIC(38032, Error, useOfNoDiffOnDifferentiableFunc, "use 'no_diff' on a call to a differentiable function has no meaning.") +DIAGNOSTIC(38032, Error, useOfNoDiffOnDifferentiableFunc, "using 'no_diff' on a call to a differentiable function has no meaning.") DIAGNOSTIC(38033, Error, cannotUseNoDiffInNonDifferentiableFunc, "cannot use 'no_diff' in a non-differentiable function.") DIAGNOSTIC(38034, Error, cannotUseConstRefOnDifferentiableParameter, "cannot use '__constref' on a differentiable parameter.") DIAGNOSTIC(38034, Error, cannotUseConstRefOnDifferentiableMemberMethod, "cannot use '[constref]' on a differentiable member method of a differentiable type.") diff --git a/source/slang/slang-ir-insts.h b/source/slang/slang-ir-insts.h index f31a56673d..19386d4dde 100644 --- a/source/slang/slang-ir-insts.h +++ b/source/slang/slang-ir-insts.h @@ -3922,6 +3922,7 @@ struct IRBuilder /// If `fallback` is true, will emit `DefaultConstruct` inst on unknown types. /// Otherwise, returns nullptr if we can't materialize the inst. IRInst* emitDefaultConstruct(IRType* type, bool fallback = true); + IRInst* _emitDefaultConstruct(IRType* type, bool fallback, HashSet visitedTypes); /// Emits a raw `DefaultConstruct` opcode without attempting to fold/materialize /// the inst. diff --git a/source/slang/slang-ir-use-uninitialized-values.cpp b/source/slang/slang-ir-use-uninitialized-values.cpp index 56b13aa099..744bedf615 100644 --- a/source/slang/slang-ir-use-uninitialized-values.cpp +++ b/source/slang/slang-ir-use-uninitialized-values.cpp @@ -154,6 +154,10 @@ namespace Slang if (!type) return true; + // In case we have a resource type (which can be assigned values in many unpredictable ways) + if (isResourceType(type)) + return true; + if (as(type)) return true; diff --git a/source/slang/slang-ir.cpp b/source/slang/slang-ir.cpp index b89929f550..7cc8e0697a 100644 --- a/source/slang/slang-ir.cpp +++ b/source/slang/slang-ir.cpp @@ -3841,8 +3841,19 @@ namespace Slang return emitIntrinsicInst(type, kIROp_DefaultConstruct, 0, nullptr); } - IRInst* IRBuilder::emitDefaultConstruct(IRType* type, bool fallback) + IRInst* IRBuilder::_emitDefaultConstruct(IRType* type, bool fallback, HashSet visitedTypes) { + // Slang generally detects recursive type-uses in IR, + // This means that DefaultConstruct may crash unless we + // track visited types with `visitedTypes.contains(type)` + // to avoid infinite looping of type-checks. + // + // Slang may be asked to default init a `RWTexture2D`, + // if so, adding `isResourceType(type)` ensures we don't + // generate garbage for resource types. + if (visitedTypes.contains(type) || isResourceType(type)) + return emitUndefined(type); + visitedTypes.add(type); IRType* actualType = type; for (;;) { @@ -3889,7 +3900,7 @@ namespace Slang return getNullPtrValue(type); case kIROp_OptionalType: { - auto inner = emitDefaultConstruct(as(actualType)->getValueType(), fallback); + auto inner = _emitDefaultConstruct(as(actualType)->getValueType(), fallback, visitedTypes); if (!inner) return nullptr; return emitMakeOptionalNone(type, inner); @@ -3903,7 +3914,7 @@ namespace Slang auto operand = tupleType->getOperand(i); if (as(operand)) break; - auto inner = emitDefaultConstruct((IRType*)operand, fallback); + auto inner = _emitDefaultConstruct((IRType*)operand, fallback, visitedTypes); if (!inner) return nullptr; elements.add(inner); @@ -3917,7 +3928,7 @@ namespace Slang for (auto field : structType->getFields()) { auto fieldType = field->getFieldType(); - auto inner = emitDefaultConstruct(fieldType, fallback); + auto inner = _emitDefaultConstruct(fieldType, fallback, visitedTypes); if (!inner) return nullptr; elements.add(inner); @@ -3929,7 +3940,7 @@ namespace Slang auto arrayType = as(actualType); if (auto count = as(arrayType->getElementCount())) { - auto element = emitDefaultConstruct(arrayType->getElementType(), fallback); + auto element = _emitDefaultConstruct(arrayType->getElementType(), fallback, visitedTypes); if (!element) return nullptr; List elements; @@ -3946,14 +3957,14 @@ namespace Slang } case kIROp_VectorType: { - auto inner = emitDefaultConstruct(as(actualType)->getElementType(), fallback); + auto inner = _emitDefaultConstruct(as(actualType)->getElementType(), fallback, visitedTypes); if (!inner) return nullptr; return emitIntrinsicInst(type, kIROp_MakeVectorFromScalar, 1, &inner); } case kIROp_MatrixType: { - auto inner = emitDefaultConstruct(as(actualType)->getElementType(), fallback); + auto inner = _emitDefaultConstruct(as(actualType)->getElementType(), fallback, visitedTypes); if (!inner) return nullptr; return emitIntrinsicInst(type, kIROp_MakeMatrixFromScalar, 1, &inner); @@ -3967,6 +3978,12 @@ namespace Slang } return nullptr; } + + IRInst* IRBuilder::emitDefaultConstruct(IRType* type, bool fallback) + { + return _emitDefaultConstruct(type, fallback, {}); + } + IRInst* IRBuilder::emitEmbeddedDownstreamIR(CodeGenTarget target, ISlangBlob *blob) { diff --git a/source/slang/slang-lower-to-ir.cpp b/source/slang/slang-lower-to-ir.cpp index 8134677437..08aa77587d 100644 --- a/source/slang/slang-lower-to-ir.cpp +++ b/source/slang/slang-lower-to-ir.cpp @@ -10002,7 +10002,7 @@ struct DeclLoweringVisitor : DeclVisitor } // Used for diagnostics - getBuilder()->addConstructorDecoration(irFunc, constructorDecl->isSynthesized); + getBuilder()->addConstructorDecoration(irFunc, constructorDecl->containsOption(ConstructorTags::Synthesized)); } // We lower whatever statement was stored on the declaration diff --git a/source/slang/slang.natvis b/source/slang/slang.natvis index 70b57e3397..d93ec09e7a 100644 --- a/source/slang/slang.natvis +++ b/source/slang/slang.natvis @@ -286,6 +286,7 @@ (Slang::ContinueStmt*)&astNodeType (Slang::ReturnStmt*)&astNodeType (Slang::ExpressionStmt*)&astNodeType + (Slang::TargetSwitchStmt*)&astNodeType (Slang::Stmt*)this,! diff --git a/tests/autodiff/bsdf/bsdf-auto-rev.slang b/tests/autodiff/bsdf/bsdf-auto-rev.slang index 1bb1989bfb..5e4242ee3c 100644 --- a/tests/autodiff/bsdf/bsdf-auto-rev.slang +++ b/tests/autodiff/bsdf/bsdf-auto-rev.slang @@ -3,10 +3,10 @@ implementing "bsdf-sample"; struct ShadingData { - float3 V; - float3 N; - float3 T; - float3 B; + float3 V = {}; + float3 N = {}; + float3 T = {}; + float3 B = {}; float3 fromLocal(float3 v) { diff --git a/tests/autodiff/cuda-kernel-export.slang b/tests/autodiff/cuda-kernel-export.slang index 39815ede85..f454113219 100644 --- a/tests/autodiff/cuda-kernel-export.slang +++ b/tests/autodiff/cuda-kernel-export.slang @@ -10,8 +10,8 @@ struct MySubType struct MyType { - float2 v; - MySubType sub[2]; + float2 v = {}; + MySubType sub[2] = {}; } struct MyInput diff --git a/tests/autodiff/generic-autodiff-1.slang b/tests/autodiff/generic-autodiff-1.slang index 9ab0d5fef1..be5fc99bf4 100644 --- a/tests/autodiff/generic-autodiff-1.slang +++ b/tests/autodiff/generic-autodiff-1.slang @@ -15,7 +15,7 @@ interface IInterface struct A : IInterface { - float z; + float z = {}; [mutating] float sample() { z = z + 1.0; return 1.0; } }; diff --git a/tests/autodiff/material/DiffuseMaterialInstance.slang b/tests/autodiff/material/DiffuseMaterialInstance.slang index c6f0e315b3..3a3401b69d 100644 --- a/tests/autodiff/material/DiffuseMaterialInstance.slang +++ b/tests/autodiff/material/DiffuseMaterialInstance.slang @@ -5,7 +5,7 @@ __exported import IBSDF; public struct DiffuseBSDF : IBSDF { - public float3 albedo; + public float3 albedo = {}; [TreatAsDifferentiable] public float3 eval(const float3 wiLocal, const float3 woLocal) diff --git a/tests/autodiff/material/GlossyMaterialInstance.slang b/tests/autodiff/material/GlossyMaterialInstance.slang index fb1ce55c0a..4e0a0f0bde 100644 --- a/tests/autodiff/material/GlossyMaterialInstance.slang +++ b/tests/autodiff/material/GlossyMaterialInstance.slang @@ -5,7 +5,7 @@ __exported import IBSDF; public struct GlossyBSDF : IBSDF { - public float3 albedo; + public float3 albedo = {}; public float3 getAlbedo() { diff --git a/tests/autodiff/modify-vector-param.slang b/tests/autodiff/modify-vector-param.slang index d6ddd73868..2084f00647 100644 --- a/tests/autodiff/modify-vector-param.slang +++ b/tests/autodiff/modify-vector-param.slang @@ -8,7 +8,7 @@ typedef float.Differential dfloat; struct Params { - float vv; + float vv = {}; } [Differentiable] diff --git a/tests/autodiff/reverse-while-loop-2.slang b/tests/autodiff/reverse-while-loop-2.slang index 9ad9ac466b..ffcb6c83a1 100644 --- a/tests/autodiff/reverse-while-loop-2.slang +++ b/tests/autodiff/reverse-while-loop-2.slang @@ -10,7 +10,7 @@ typedef float.Differential dfloat; struct P { - bool terminated; + bool terminated = {}; bool isTerminated() { return terminated; } bool isHit() { return !terminated; } } diff --git a/tests/autodiff/reverse-while-loop-3.slang b/tests/autodiff/reverse-while-loop-3.slang index 0a27437e97..d00e9bda94 100644 --- a/tests/autodiff/reverse-while-loop-3.slang +++ b/tests/autodiff/reverse-while-loop-3.slang @@ -14,7 +14,7 @@ typedef float.Differential dfloat; struct P { - bool terminated; + bool terminated = {}; bool isTerminated() { return terminated; } bool isHit() { return !terminated; } }; @@ -25,7 +25,7 @@ struct Data { this.t = dataIn; } - float t; + float t = {}; }; void updateData(Data data) diff --git a/tests/bindings/nested-parameter-block-2.slang b/tests/bindings/nested-parameter-block-2.slang index 790e3d060a..1107739126 100644 --- a/tests/bindings/nested-parameter-block-2.slang +++ b/tests/bindings/nested-parameter-block-2.slang @@ -6,7 +6,7 @@ struct CB { - uint4 value; + uint4 value = {}; } struct MaterialSystem diff --git a/tests/bindings/nested-parameter-block-3.slang b/tests/bindings/nested-parameter-block-3.slang index a9cfcfdaab..3fd07960fd 100644 --- a/tests/bindings/nested-parameter-block-3.slang +++ b/tests/bindings/nested-parameter-block-3.slang @@ -11,7 +11,7 @@ struct CB { - uint4 value; + uint4 value = {}; } struct MaterialSystem diff --git a/tests/bugs/c-style-cast-coerce.slang b/tests/bugs/c-style-cast-coerce.slang index e01903e2cc..b123638415 100644 --- a/tests/bugs/c-style-cast-coerce.slang +++ b/tests/bugs/c-style-cast-coerce.slang @@ -1,4 +1,4 @@ -//TEST:SIMPLE: +//TEST:SIMPLE(filecheck=CHECK): -target hlsl // It used to be the case that coercions of already-coerced initializer lists // didn't take into account the more specific type. One way of triggering this @@ -6,6 +6,8 @@ // syntax which creates the equivalent of `{} : S`. In the example below this // then proceeds to be coerced to type T. +//CHECK: c-style-cast-coerce.slang(17): error 30019: expected an expression of type 'T', got 'S' + struct S {}; struct T {}; diff --git a/tests/bugs/c-style-cast-coerce.slang.expected b/tests/bugs/c-style-cast-coerce.slang.expected deleted file mode 100644 index a09c6842fe..0000000000 --- a/tests/bugs/c-style-cast-coerce.slang.expected +++ /dev/null @@ -1,8 +0,0 @@ -result code = -1 -standard error = { -tests/bugs/c-style-cast-coerce.slang(15): error 30019: expected an expression of type 'T', got 'S' - T t = (S)0; - ^ -} -standard output = { -} diff --git a/tests/bugs/gh-3601.slang b/tests/bugs/gh-3601.slang index d12b480acf..8852cb9658 100644 --- a/tests/bugs/gh-3601.slang +++ b/tests/bugs/gh-3601.slang @@ -42,7 +42,7 @@ void main(int id : SV_DispatchThreadID) *pData1 = 3; *(int2*)pData = int2(1, 2); pData1[-1] = 2; - buffer[0].pNext[1] = {5}; + buffer[0].pNext[1] = {5, 0}; // CHECK: OpConvertPtrToU // CHECK: OpINotEqual if (pData1) diff --git a/tests/bugs/gh-3845-2.slang b/tests/bugs/gh-3845-2.slang index 06a3d1c999..6073ffdad2 100644 --- a/tests/bugs/gh-3845-2.slang +++ b/tests/bugs/gh-3845-2.slang @@ -1,7 +1,7 @@ //TEST(compute):COMPARE_COMPUTE(filecheck-buffer=CHECK):-shaderobj -output-using-type //TEST(compute):COMPARE_COMPUTE(filecheck-buffer=CHECK):-shaderobj -vk -output-using-type -struct Foo {float v;}; +struct Foo {float v = {};}; namespace foo { typedef ::Foo Foo; // unexpected '::', expected identifier; works in dxc/hlsl diff --git a/tests/bugs/gh-3845.slang b/tests/bugs/gh-3845.slang index fe8da5acf9..858a8bfbed 100644 --- a/tests/bugs/gh-3845.slang +++ b/tests/bugs/gh-3845.slang @@ -2,7 +2,7 @@ //TEST(compute):COMPARE_COMPUTE(filecheck-buffer=CHECK):-shaderobj -vk -output-using-type -struct Foo {float v;}; +struct Foo {float v = {};}; namespace foo { typedef Foo Foo; // cyclic reference 'Foo'; allowed in dxc/hlsl diff --git a/tests/bugs/optional.slang b/tests/bugs/optional.slang index 7a220b23d5..14932d8f25 100644 --- a/tests/bugs/optional.slang +++ b/tests/bugs/optional.slang @@ -6,7 +6,7 @@ interface IFoo void foo(); } -struct S : IFoo { int x; void foo(); } +struct S : IFoo { int x = {}; void foo(); } struct P { diff --git a/tests/bugs/specialize-function-array-args.slang b/tests/bugs/specialize-function-array-args.slang index 8a572eda4c..b1f1f13c2d 100644 --- a/tests/bugs/specialize-function-array-args.slang +++ b/tests/bugs/specialize-function-array-args.slang @@ -13,7 +13,7 @@ struct Params { RWStructuredBuffer buffer; - int a[2]; + int a[2] = {}; }; //TEST_INPUT:ubuffer(data=[9 9 9 9], stride=4):out,name=outputBuffer @@ -38,4 +38,4 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) int tid = int(dispatchThreadID.x); outputBuffer[tid] = doSomething(params, tid); -} \ No newline at end of file +} diff --git a/tests/bugs/stk-chk.slang b/tests/bugs/stk-chk.slang index c745ecc72a..165e4b0972 100644 --- a/tests/bugs/stk-chk.slang +++ b/tests/bugs/stk-chk.slang @@ -5,7 +5,7 @@ struct LargeStruct { - int values[4096]; + int values[4096] = {}; }; //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name=outputBuffer @@ -52,4 +52,4 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) int index = int(dispatchThreadID.x); outputBuffer[index] = _calc1(index) + 1; -} \ No newline at end of file +} diff --git a/tests/compute/assoctype-complex.slang b/tests/compute/assoctype-complex.slang index bd53be4857..7fc43735da 100644 --- a/tests/compute/assoctype-complex.slang +++ b/tests/compute/assoctype-complex.slang @@ -18,7 +18,7 @@ interface ISimple struct Val : IBase { typedef int V; - int base; + int base = {}; V sub(V a0, V a1) { return a0 - a1 + base; @@ -49,4 +49,4 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) v1.base = 2; int outVal = test(s, v0, v1); // == 4.0 outputBuffer[dispatchThreadID.x] = outVal; -} \ No newline at end of file +} diff --git a/tests/compute/assoctype-lookup.slang b/tests/compute/assoctype-lookup.slang index 348391e217..edfdfedfb8 100644 --- a/tests/compute/assoctype-lookup.slang +++ b/tests/compute/assoctype-lookup.slang @@ -16,8 +16,8 @@ struct StandardBoneWeightSet : IBoneWeightSet { struct PackedType { - uint boneIds : BONEIDS; - uint boneWeights : BONEWEIGHTS; + uint boneIds : BONEIDS = {}; + uint boneWeights : BONEWEIGHTS = {}; }; PackedType field; }; diff --git a/tests/compute/buffer-layout.slang b/tests/compute/buffer-layout.slang index a456fa1428..5ba4b9f456 100644 --- a/tests/compute/buffer-layout.slang +++ b/tests/compute/buffer-layout.slang @@ -14,8 +14,8 @@ RWStructuredBuffer outputBuffer; struct A { - float x; - float y; + float x = {}; + float y = {}; } struct S @@ -29,7 +29,7 @@ struct S // // 0 4 4 // - float z; + float z = {}; // The `std140` and D3D constant buffer ruless both // ensure a minimum of 16-byte alignment on `struct` @@ -62,7 +62,7 @@ struct S // struct 12 4 4 // std430 12 4 4 // - int b; + int b = {}; // As our next stress-test case, we will insert an // array with elements that aren't a multiple of @@ -89,7 +89,7 @@ struct S // struct 16 16 4 // std430 16 16 8 // - float2 c[2]; + float2 c[2] = {}; // Now we put in one more ordinary `int` field // just to probe the offset computed so far. diff --git a/tests/compute/cbuffer-legalize.slang b/tests/compute/cbuffer-legalize.slang index d0c0a87774..b892a74b9b 100644 --- a/tests/compute/cbuffer-legalize.slang +++ b/tests/compute/cbuffer-legalize.slang @@ -8,7 +8,7 @@ struct P { - uint4 c; + uint4 c = {}; Texture2D t; SamplerState s; }; diff --git a/tests/compute/dynamic-dispatch-1.slang b/tests/compute/dynamic-dispatch-1.slang index 10d3552ab8..57ee654656 100644 --- a/tests/compute/dynamic-dispatch-1.slang +++ b/tests/compute/dynamic-dispatch-1.slang @@ -18,7 +18,7 @@ int GenericCompute(T obj, int inVal) struct Impl : IInterface { - int base; + int base = {}; int Compute(int inVal) { return base + inVal * inVal; } }; diff --git a/tests/compute/dynamic-dispatch-10.slang b/tests/compute/dynamic-dispatch-10.slang index ddd51c4b92..62230a1c5d 100644 --- a/tests/compute/dynamic-dispatch-10.slang +++ b/tests/compute/dynamic-dispatch-10.slang @@ -25,7 +25,7 @@ int GenericCompute1(T obj, IInterface obj2, int inVal) struct Impl : IInterface { - int base; + int base = {}; int Compute(int inVal) { return base + inVal * inVal; } }; diff --git a/tests/compute/dynamic-dispatch-14.slang b/tests/compute/dynamic-dispatch-14.slang index 641fe938c8..f19a7efa48 100644 --- a/tests/compute/dynamic-dispatch-14.slang +++ b/tests/compute/dynamic-dispatch-14.slang @@ -47,10 +47,10 @@ void computeMain(int3 dispatchThreadID : SV_DispatchThreadID) // Type must be marked `public` to ensure it is visible in the generated DLL. export struct MyImpl : IInterface { - int val; + int val = {}; export struct TAssoc : IAssoc { - int val; + int val = {}; int eval() { return val; } } TAssoc run(int input) @@ -63,10 +63,10 @@ export struct MyImpl : IInterface export struct MyImpl2 : IInterface { - int val; + int val = {}; export struct TAssoc : IAssoc { - int val; + int val = {}; int eval() { return val; } } TAssoc run(int input) diff --git a/tests/compute/dynamic-dispatch-15.slang b/tests/compute/dynamic-dispatch-15.slang index 2ab169281c..6be4e6505e 100644 --- a/tests/compute/dynamic-dispatch-15.slang +++ b/tests/compute/dynamic-dispatch-15.slang @@ -18,7 +18,7 @@ RWStructuredBuffer gOutputBuffer; RWStructuredBuffer gObj; [numthreads(1, 1, 1)] -void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) { // Test unpacking. float result = 0.0; @@ -48,7 +48,7 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) // Type must be marked `public` to ensure it is visible in the generated DLL. export struct FloatVal : IInterface { - float val; + float val = {}; float run() { return val; @@ -58,7 +58,7 @@ interface ISomething{void g();} struct Float4Struct : ISomething { float4 val; void g() {} } export struct Float4Val : IInterface { - Float4Struct val; + Float4Struct val = {}; float run() { return val.val.x; @@ -66,7 +66,7 @@ export struct Float4Val : IInterface }; export struct IntVal : IInterface { - int val; + int val = {}; float run() { return val; @@ -74,7 +74,7 @@ export struct IntVal : IInterface }; export struct Int4Val : IInterface { - int4 val; + int4 val = {}; float run() { return val.x; diff --git a/tests/compute/dynamic-dispatch-16.slang b/tests/compute/dynamic-dispatch-16.slang index b9fecb9664..9ce00a3d7c 100644 --- a/tests/compute/dynamic-dispatch-16.slang +++ b/tests/compute/dynamic-dispatch-16.slang @@ -12,8 +12,8 @@ interface IInterface struct UserDefinedPackedType { - float3 val; - uint flags; + float3 val = {}; + uint flags = {}; }; //TEST_INPUT:ubuffer(data=[0], stride=4):out,name=gOutputBuffer @@ -40,14 +40,14 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) struct FloatVal : IInterface { - float val; + float val = {}; float run() { return val; } }; interface ISomething{void g();} -struct Float4Struct : ISomething { float4 val; void g() {} } +struct Float4Struct : ISomething { float4 val = {}; void g() {} } struct Float4Val : IInterface { Float4Struct val; diff --git a/tests/compute/dynamic-dispatch-17.slang b/tests/compute/dynamic-dispatch-17.slang index e89482a5b7..ab99de2f23 100644 --- a/tests/compute/dynamic-dispatch-17.slang +++ b/tests/compute/dynamic-dispatch-17.slang @@ -17,8 +17,8 @@ interface IInterface struct UserDefinedPackedType { - float3 val; - uint flags; + float3 val = {}; + uint flags = {}; }; //TEST_INPUT:ubuffer(data=[0], stride=4):out,name=gOutputBuffer @@ -49,14 +49,14 @@ struct ReturnsZero : IReturnsZero } struct FloatVal : IInterface { - float val; + float val = {}; float run() { Z z = Z(); return val + z.get(); } }; -struct Float4Struct { float4 val; } +struct Float4Struct { float4 val = {}; } struct Float4Val : IInterface { Float4Struct val; diff --git a/tests/compute/dynamic-dispatch-18.slang b/tests/compute/dynamic-dispatch-18.slang index 23cdabbbdc..79dedc82f5 100644 --- a/tests/compute/dynamic-dispatch-18.slang +++ b/tests/compute/dynamic-dispatch-18.slang @@ -19,10 +19,10 @@ interface IInterface struct Impl : IInterface { - int data; + int data = {}; struct Getter : IReturnsZero { - bool data; + bool data = {}; float get() { if (data) return 0.0; else return 1.0;} } Getter createGetter() { Getter g; g.data = true; return g; } diff --git a/tests/compute/dynamic-dispatch-2.slang b/tests/compute/dynamic-dispatch-2.slang index 06aed6e71c..ec306c9bea 100644 --- a/tests/compute/dynamic-dispatch-2.slang +++ b/tests/compute/dynamic-dispatch-2.slang @@ -29,11 +29,11 @@ struct Impl : IInterface { struct Assoc : IAssoc { - int val; + int val = {}; int get() { return val; } static int getBase() { return -1; } }; - int base; + int base = {}; int Compute(int inVal) { return base + inVal * inVal; } }; diff --git a/tests/compute/dynamic-dispatch-3.slang b/tests/compute/dynamic-dispatch-3.slang index 94430e638a..21f52aa15b 100644 --- a/tests/compute/dynamic-dispatch-3.slang +++ b/tests/compute/dynamic-dispatch-3.slang @@ -38,11 +38,11 @@ struct Impl : IInterface { struct Assoc : IAssoc { - int val; + int val = {}; int get() { return val; } static int getBase(T t) { return t.getVal(); } }; - int base; + int base = {}; int Compute(int inVal) { return base + inVal * inVal; } }; diff --git a/tests/compute/dynamic-dispatch-5.slang b/tests/compute/dynamic-dispatch-5.slang index 61d4024af5..dbe462323e 100644 --- a/tests/compute/dynamic-dispatch-5.slang +++ b/tests/compute/dynamic-dispatch-5.slang @@ -17,7 +17,7 @@ int GenericCompute(int inVal, T obj, T other) struct Impl : IInterface { - int base; + int base = {}; int Compute(int inVal, This other) { return other.base + base + inVal; } }; int test(int inVal) diff --git a/tests/compute/dynamic-dispatch-7.slang b/tests/compute/dynamic-dispatch-7.slang index 7f516b0d55..c8dada852a 100644 --- a/tests/compute/dynamic-dispatch-7.slang +++ b/tests/compute/dynamic-dispatch-7.slang @@ -45,7 +45,7 @@ struct Impl : IInterface { struct TAssoc : IAssoc { - int base; + int base = {}; int Compute() { return base; diff --git a/tests/compute/dynamic-dispatch-8.slang b/tests/compute/dynamic-dispatch-8.slang index 62d3f0a9d4..d5a9cfc8ed 100644 --- a/tests/compute/dynamic-dispatch-8.slang +++ b/tests/compute/dynamic-dispatch-8.slang @@ -18,7 +18,7 @@ int GenericCompute(IInterface obj, int inVal) struct Impl : IInterface { - int base; + int base = {}; int Compute(int inVal) { return base + inVal * inVal; } }; diff --git a/tests/compute/dynamic-dispatch-9.slang b/tests/compute/dynamic-dispatch-9.slang index 77112a318b..0672181076 100644 --- a/tests/compute/dynamic-dispatch-9.slang +++ b/tests/compute/dynamic-dispatch-9.slang @@ -28,7 +28,7 @@ int GenericCompute1(T obj, int inVal) struct Impl : IInterface { - int base; + int base = {}; int Compute(int inVal) { return base + inVal * inVal; } }; diff --git a/tests/compute/explicit-this-expr.slang b/tests/compute/explicit-this-expr.slang index baf718be0d..d73584a246 100644 --- a/tests/compute/explicit-this-expr.slang +++ b/tests/compute/explicit-this-expr.slang @@ -7,7 +7,7 @@ struct A { - float x; + float x = {}; float addWith(float y) { @@ -33,4 +33,4 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) float inVal = float(tid); float outVal = test(inVal); outputBuffer[tid] = outVal; -} \ No newline at end of file +} diff --git a/tests/compute/extension-multi-interface.slang b/tests/compute/extension-multi-interface.slang index 23c9174110..94dc168fab 100644 --- a/tests/compute/extension-multi-interface.slang +++ b/tests/compute/extension-multi-interface.slang @@ -22,7 +22,7 @@ interface IAddAndSub struct Simple : IAdd { - float base; + float base = {}; float addf(float u, float v) { return u+v; @@ -49,4 +49,4 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) s.base = 0.0; float outVal = testAddSub(s); outputBuffer[dispatchThreadID.x] = outVal; -} \ No newline at end of file +} diff --git a/tests/compute/generic-interface-method-simple.slang b/tests/compute/generic-interface-method-simple.slang index 0e20947d16..dd8d265427 100644 --- a/tests/compute/generic-interface-method-simple.slang +++ b/tests/compute/generic-interface-method-simple.slang @@ -22,7 +22,7 @@ struct StandardVertexInterpolant : IVertexInterpolant struct MaterialPattern1 : IBRDFPattern { - float base; + float base = {}; __generic float evalPattern(TVertexInterpolant interpolants) { @@ -46,4 +46,4 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) mp1.base = 0.5; float outVal = test(mp1, vertInterp); outputBuffer[dispatchThreadID.x] = outVal; -} \ No newline at end of file +} diff --git a/tests/compute/generic-interface-method.slang b/tests/compute/generic-interface-method.slang index fade14d412..56d12141a0 100644 --- a/tests/compute/generic-interface-method.slang +++ b/tests/compute/generic-interface-method.slang @@ -6,25 +6,27 @@ RWStructuredBuffer outputBuffer; struct DisneyBRDFPattern { - float3 baseColor; - float3 normal; - float specular, metallic, roughness; - float opacity; - float3 emmissive; - float ambientOcclusion; + float3 baseColor = {}; + float3 normal = {}; + float specular = {}; + float metallic = {}; + float roughness = {}; + float opacity = {}; + float3 emmissive = {}; + float ambientOcclusion = {}; }; struct VertexPosition { - float3 pos; - float3 normal; - float2 uv; + float3 pos = {}; + float3 normal = {}; + float2 uv = {}; }; struct CameraView { - float3 camPos; - float3 camDir; + float3 camPos = {}; + float3 camDir = {}; }; interface IVertexInterpolant @@ -84,4 +86,4 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) MaterialPattern1 mp1; float outVal = test(vertInterp, mp1); outputBuffer[dispatchThreadID.x] = outVal; -} \ No newline at end of file +} diff --git a/tests/compute/generics-constrained.slang b/tests/compute/generics-constrained.slang index 435b840d1f..d166c18bce 100644 --- a/tests/compute/generics-constrained.slang +++ b/tests/compute/generics-constrained.slang @@ -10,7 +10,7 @@ interface Helper struct A : Helper { - float a; + float a = {}; float getHelp() { @@ -42,4 +42,4 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) float outVal = testHelp(a); outputBuffer[tid] = outVal; -} \ No newline at end of file +} diff --git a/tests/compute/half-structured-buffer.slang b/tests/compute/half-structured-buffer.slang index 973e84e7b1..7329f17e13 100644 --- a/tests/compute/half-structured-buffer.slang +++ b/tests/compute/half-structured-buffer.slang @@ -6,9 +6,9 @@ struct Thing { - uint pos; - float radius; - half4 color; + uint pos = {}; + float radius = {}; + half4 color = {}; }; //TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=16):out,name outputBuffer @@ -29,4 +29,4 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) thing.radius = v; outputBuffer[tid] = thing; -} \ No newline at end of file +} diff --git a/tests/compute/implicit-generic-app.slang b/tests/compute/implicit-generic-app.slang index c41b5e31d8..460e1b02bf 100644 --- a/tests/compute/implicit-generic-app.slang +++ b/tests/compute/implicit-generic-app.slang @@ -12,7 +12,7 @@ interface ISimple struct Simple : ISimple { - int val; + int val = {}; int getVal() { return val; } }; @@ -38,4 +38,4 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) int tid = (int) dispatchThreadID.x; int outVal = test(tid); outputBuffer[tid] = outVal; -} \ No newline at end of file +} diff --git a/tests/compute/implicit-this-expr.slang b/tests/compute/implicit-this-expr.slang index e8b7200938..1cd7711c33 100644 --- a/tests/compute/implicit-this-expr.slang +++ b/tests/compute/implicit-this-expr.slang @@ -6,7 +6,7 @@ struct A { - float x; + float x = {}; float addWith(float y) { @@ -31,4 +31,4 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) float inVal = float(tid); float outVal = test(inVal); outputBuffer[tid] = outVal; -} \ No newline at end of file +} diff --git a/tests/compute/init-list-defaults.slang.expected.txt b/tests/compute/init-list-defaults.slang.expected.txt deleted file mode 100644 index fa9fe4c8a4..0000000000 --- a/tests/compute/init-list-defaults.slang.expected.txt +++ /dev/null @@ -1,4 +0,0 @@ -1142 -2453 -3060 -4000 diff --git a/tests/compute/initializer-list.slang.expected.txt b/tests/compute/initializer-list.slang.expected.txt deleted file mode 100644 index a0d427709c..0000000000 --- a/tests/compute/initializer-list.slang.expected.txt +++ /dev/null @@ -1,4 +0,0 @@ -10 -11 -12 -13 diff --git a/tests/compute/interface-assoc-type-param.slang b/tests/compute/interface-assoc-type-param.slang index a234d457f0..acdc38a98a 100644 --- a/tests/compute/interface-assoc-type-param.slang +++ b/tests/compute/interface-assoc-type-param.slang @@ -18,10 +18,10 @@ interface IEval export struct Impl : IInterface { - uint val; + uint val = {}; struct TEval : IEval { - uint val; + uint val = {}; uint eval() { return val; diff --git a/tests/compute/mutating-methods.slang b/tests/compute/mutating-methods.slang index 6f561e9a66..7f5e23043e 100644 --- a/tests/compute/mutating-methods.slang +++ b/tests/compute/mutating-methods.slang @@ -11,7 +11,7 @@ interface IAccumulator struct Accumulator : IAccumulator { - int state; + int state = {}; [mutating] void accumulate(int v) { diff --git a/tests/compute/pack-any-value-16bit.slang b/tests/compute/pack-any-value-16bit.slang index 07ad0590e0..ff01c68029 100644 --- a/tests/compute/pack-any-value-16bit.slang +++ b/tests/compute/pack-any-value-16bit.slang @@ -12,12 +12,12 @@ interface IInterface struct Val : IInterface { - int16_t v0; - float f0; - uint16_t v1; - float16_t v2; - half v3; - uint16_t v4; + int16_t v0 = {}; + float f0 = {}; + uint16_t v1 = {}; + float16_t v2 = {}; + half v3 = {}; + uint16_t v4 = {}; float run() { return v0 + f0 + v1 + v2 + v3 + v4; @@ -26,7 +26,7 @@ struct Val : IInterface struct UserDefinedPackedType { - uint4 values[2]; + uint4 values[2] = {}; }; //TEST_INPUT:ubuffer(data=[0 0 0], stride=4):out,name=gOutputBuffer diff --git a/tests/compute/pointer-emit.slang b/tests/compute/pointer-emit.slang index 0516e2dfa1..3d72c9a660 100644 --- a/tests/compute/pointer-emit.slang +++ b/tests/compute/pointer-emit.slang @@ -4,13 +4,13 @@ struct Something { - int4 field0; - int field1; + int4 field0 = {}; + int field1 = {}; }; struct ArrayType { - Something data[2]; + Something data[2] = {}; }; ArrayType getArray() diff --git a/tests/compute/scope-operator.slang b/tests/compute/scope-operator.slang index 0622adc4a2..2d90702a36 100644 --- a/tests/compute/scope-operator.slang +++ b/tests/compute/scope-operator.slang @@ -17,10 +17,10 @@ enum Color struct Thing { - int a; + int a = {}; struct Another { - int b; + int b = {}; } }; @@ -75,4 +75,4 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) val = test(val); outputBuffer[tid] = val; -} \ No newline at end of file +} diff --git a/tests/compute/struct-default-init.slang b/tests/compute/struct-default-init.slang index dc0e0218af..5a0d0e6930 100644 --- a/tests/compute/struct-default-init.slang +++ b/tests/compute/struct-default-init.slang @@ -7,6 +7,21 @@ struct Test int b = 1; int c = 0; int d = 1 + 1; + __init(int a_in) + { + a = a_in; + } + __init(int a_in, int b_in) + { + a = a_in; + b = b_in; + } + __init(int a_in, int b_in, int c_in) + { + a = a_in; + b = b_in; + c = c_in; + } } int test(int inVal) diff --git a/tests/compute/struct-make.slang b/tests/compute/struct-make.slang index b7adf4e838..e5340d596d 100644 --- a/tests/compute/struct-make.slang +++ b/tests/compute/struct-make.slang @@ -7,7 +7,7 @@ struct Thing { - int a; + int a = {}; static Thing make(int v) { Thing thing; diff --git a/tests/compute/texture-sample-grad-offset-clamp.slang b/tests/compute/texture-sample-grad-offset-clamp.slang index b83cc12647..0cee3fedd7 100644 --- a/tests/compute/texture-sample-grad-offset-clamp.slang +++ b/tests/compute/texture-sample-grad-offset-clamp.slang @@ -16,20 +16,20 @@ cbuffer Uniforms struct AssembledVertex { - float3 position; - float3 color; - float2 uv; + float3 position = {}; + float3 color = {}; + float2 uv = {}; }; struct CoarseVertex { - float3 color; - float2 uv; + float3 color = {}; + float2 uv = {}; }; struct Fragment { - float4 color; + float4 color = {}; }; @@ -43,7 +43,7 @@ struct VertexStageInput struct VertexStageOutput { CoarseVertex coarseVertex : CoarseVertex; - float4 sv_position : SV_Position; + float4 sv_position : SV_Position = {}; }; VertexStageOutput vertexMain(VertexStageInput input) diff --git a/tests/compute/texture-sampling.slang b/tests/compute/texture-sampling.slang index 0e319680a3..30c9b45682 100644 --- a/tests/compute/texture-sampling.slang +++ b/tests/compute/texture-sampling.slang @@ -30,20 +30,20 @@ cbuffer Uniforms struct AssembledVertex { - float3 position; - float3 color; - float2 uv; + float3 position = {}; + float3 color = {}; + float2 uv = {}; }; struct CoarseVertex { - float3 color; - float2 uv; + float3 color = {}; + float2 uv = {}; }; struct Fragment { - float4 color; + float4 color = {}; }; @@ -57,7 +57,7 @@ struct VertexStageInput struct VertexStageOutput { CoarseVertex coarseVertex : CoarseVertex; - float4 sv_position : SV_Position; + float4 sv_position : SV_Position = {}; }; VertexStageOutput vertexMain(VertexStageInput input) diff --git a/tests/compute/type-legalize-global-with-init.slang b/tests/compute/type-legalize-global-with-init.slang index 573ac98499..a3a1429aae 100644 --- a/tests/compute/type-legalize-global-with-init.slang +++ b/tests/compute/type-legalize-global-with-init.slang @@ -14,19 +14,15 @@ RWStructuredBuffer inputBuffer; static const RWStructuredBuffer gBuffer = inputBuffer; -struct Stuff -{ - RWStructuredBuffer a; - RWStructuredBuffer b; -} - -static const Stuff gStuff = { inputBuffer, inputBuffer }; +// Note: due to #4874 we cannot assign functions to globals +// which contain resource types. +static const RWStructuredBuffer gStuff = inputBuffer; uint test(uint x) { return gBuffer[x] - + gStuff.a[x + 1] * 16 - + gStuff.b[x + 2] * 256; + + gStuff[x + 1] * 16 + + gStuff[x + 2] * 256; } [numthreads(4, 1, 1)] diff --git a/tests/diagnostics/autodiff.slang b/tests/diagnostics/autodiff.slang index 7905b48b6d..36ac69c7cf 100644 --- a/tests/diagnostics/autodiff.slang +++ b/tests/diagnostics/autodiff.slang @@ -1,4 +1,13 @@ -//DIAGNOSTIC_TEST:SIMPLE: +//DIAGNOSTIC_TEST:SIMPLE(filecheck=CHECK): + +//CHECK: tests/diagnostics/autodiff.slang({{..}}): error 38031: +//CHECK-NEXT: float x1 = no_diff x; + +//CHECK: tests/diagnostics/autodiff.slang({{..}}): error 38032: +//CHECK-NEXT: return no_diff f(x); + +//CHECK: tests/diagnostics/autodiff.slang({{..}}): error 38033: +//CHECK-NEXT: return no_diff nonDiff(x); float nonDiff(float x) { diff --git a/tests/diagnostics/autodiff.slang.expected b/tests/diagnostics/autodiff.slang.expected deleted file mode 100644 index 95da2a6d3b..0000000000 --- a/tests/diagnostics/autodiff.slang.expected +++ /dev/null @@ -1,14 +0,0 @@ -result code = -1 -standard error = { -tests/diagnostics/autodiff.slang(30): error 38031: 'no_diff' can only be used to decorate a call or a subscript operation - float x1 = no_diff x; // invalid use of no_diff here. - ^~~~~~~ -tests/diagnostics/autodiff.slang(31): error 38032: use 'no_diff' on a call to a differentiable function has no meaning. - return no_diff f(x); // no_diff on a differentiable call has no meaning. - ^~~~~~~ -tests/diagnostics/autodiff.slang(36): error 38033: cannot use 'no_diff' in a non-differentiable function. - return no_diff nonDiff(x); // no_diff in a non-differentiable function - ^~~~~~~ -} -standard output = { -} diff --git a/tests/diagnostics/deprecation.slang b/tests/diagnostics/deprecation.slang index 75dfd51d80..7fad5c8fe1 100644 --- a/tests/diagnostics/deprecation.slang +++ b/tests/diagnostics/deprecation.slang @@ -20,7 +20,7 @@ const int pi = 3; struct Billiards { [deprecated("puns")] - bool bs[22]; + bool bs[22] = {}; } int main() diff --git a/tests/diagnostics/mismatching-types.slang b/tests/diagnostics/mismatching-types.slang index 15fc1d0e3f..edd4e33750 100644 --- a/tests/diagnostics/mismatching-types.slang +++ b/tests/diagnostics/mismatching-types.slang @@ -1,5 +1,5 @@ // mismatching-types.slang -//DIAGNOSTIC_TEST:SIMPLE:-target hlsl +//TEST:SIMPLE(filecheck=CHECK): -target hlsl Texture1D tex; @@ -47,21 +47,21 @@ void main(uint3 dispatchThreadID : SV_DispatchThreadID) NonGenericOuter c; NonGenericOuter.GenericInner d; - // expected an expression of type 'GenericOuter', got 'int' +//CHECK: tests/diagnostics/mismatching-types.slang(51): error 30019: expected an expression of type 'GenericOuter', got 'int' a = 0; - // expected an expression of type 'GenericOuter.GenericInner', got 'int' +//CHECK: tests/diagnostics/mismatching-types.slang(53): error 30019: expected an expression of type 'GenericOuter.GenericInner', got 'int' a.g = 0; - // expected an expression of type 'GenericOuter.NonGenericInner', got 'int' +//CHECK: tests/diagnostics/mismatching-types.slang(55): error 30019: expected an expression of type 'GenericOuter.NonGenericInner', got 'int' a.ng = 0; - // expected an expression of type 'GenericOuter.GenericInner', got 'GenericOuter.GenericInner' +//CHECK: tests/diagnostics/mismatching-types.slang(57): error 30019: expected an expression of type 'GenericOuter.GenericInner', got 'GenericOuter.GenericInner' a.g = b.g; - // expected an expression of type 'GenericOuter.NonGenericInner', got 'GenericOuter.NonGenericInner' +//CHECK: tests/diagnostics/mismatching-types.slang(59): error 30019: expected an expression of type 'GenericOuter.NonGenericInner', got 'GenericOuter.NonGenericInner' a.ng = b.ng; - // expected an expression of type 'NonGenericOuter.GenericInner', got 'int' +//CHECK: tests/diagnostics/mismatching-types.slang(61): error 30019: expected an expression of type 'GenericInner', got 'int' c.i = 0; - // expected an expression of type 'NonGenericOuter.GenericInner', got 'NonGenericOuter.GenericInner' +//CHECK: tests/diagnostics/mismatching-types.slang(63): error 30019: expected an expression of type 'GenericInner', got 'GenericInner' c.i = c.f; - // expected an expression of type 'NonGenericOuter.GenericInner.ReallyNested', got 'int' +//CHECK: tests/diagnostics/mismatching-types.slang(65): error 30019: expected an expression of type 'GenericInner.ReallyNested', got 'int' c.i.n = 0; // OK c.i.n.val = 0; @@ -70,8 +70,8 @@ void main(uint3 dispatchThreadID : SV_DispatchThreadID) // OK c.i = d; - // expected an expression of type 'Texture1D', got 'Texture1D' +//CHECK: tests/diagnostics/mismatching-types.slang(74): error 30019: expected an expression of type 'Texture1D', got 'Texture1D' Texture1D t1 = tex; - // expected an expression of type 'Texture2D', got 'Texture1D' +//CHECK: tests/diagnostics/mismatching-types.slang(76): error 30019: expected an expression of type 'Texture2D', got 'Texture1D' Texture2D t2 = tex; } \ No newline at end of file diff --git a/tests/diagnostics/variable-redeclaration.slang b/tests/diagnostics/variable-redeclaration.slang index bbd6a07c01..d2d9f90d70 100644 --- a/tests/diagnostics/variable-redeclaration.slang +++ b/tests/diagnostics/variable-redeclaration.slang @@ -1,6 +1,6 @@ // variable-redeclaration.slang -//DIAGNOSTIC_TEST:SIMPLE: +//TEST:SIMPLE(filecheck=CHECK): -target hlsl // This test confirms that the compiler produces @@ -52,3 +52,15 @@ int testParameterRedeclaration( { return size; } + +//CHECK: tests/diagnostics/variable-redeclaration.slang(14): error 30200: declaration of 'gA' conflicts with existing declaration +//CHECK: tests/diagnostics/variable-redeclaration.slang(12): note: see previous declaration of 'gA' +//CHECK: tests/diagnostics/variable-redeclaration.slang(44): error 30200: declaration of 'f' conflicts with existing declaration +//CHECK: tests/diagnostics/variable-redeclaration.slang(43): note: see previous declaration of 'f' +//CHECK: tests/diagnostics/variable-redeclaration.slang(51): error 30200: declaration of 'size' conflicts with existing declaration +//CHECK: tests/diagnostics/variable-redeclaration.slang(50): note: see previous declaration of 'size' +//CHECK: tests/diagnostics/variable-redeclaration.slang(21): error 30200: declaration of 'y' conflicts with existing declaration +//CHECK: tests/diagnostics/variable-redeclaration.slang(20): note: see previous declaration of 'y' +//CHECK: tests/diagnostics/variable-redeclaration.slang(53): error 39999: ambiguous reference to 'size' +//CHECK: tests/diagnostics/variable-redeclaration.slang(51): note 39999: candidate: float size +//CHECK: tests/diagnostics/variable-redeclaration.slang(50): note 39999: candidate: int size \ No newline at end of file diff --git a/tests/diagnostics/variable-redeclaration.slang.expected b/tests/diagnostics/variable-redeclaration.slang.expected deleted file mode 100644 index 944037b1cd..0000000000 --- a/tests/diagnostics/variable-redeclaration.slang.expected +++ /dev/null @@ -1,38 +0,0 @@ -result code = -1 -standard error = { -tests/diagnostics/variable-redeclaration.slang(14): error 30200: declaration of 'gA' conflicts with existing declaration -static Texture2D gA; - ^~ -tests/diagnostics/variable-redeclaration.slang(12): note: see previous declaration of 'gA' -static int gA; - ^~ -tests/diagnostics/variable-redeclaration.slang(44): error 30200: declaration of 'f' conflicts with existing declaration - float f; - ^ -tests/diagnostics/variable-redeclaration.slang(43): note: see previous declaration of 'f' - int f; - ^ -tests/diagnostics/variable-redeclaration.slang(51): error 30200: declaration of 'size' conflicts with existing declaration - float size) - ^~~~ -tests/diagnostics/variable-redeclaration.slang(50): note: see previous declaration of 'size' - int size, - ^~~~ -tests/diagnostics/variable-redeclaration.slang(21): error 30200: declaration of 'y' conflicts with existing declaration - int y = x; - ^ -tests/diagnostics/variable-redeclaration.slang(20): note: see previous declaration of 'y' - int y = x; - ^ -tests/diagnostics/variable-redeclaration.slang(53): error 39999: ambiguous reference to 'size' - return size; - ^~~~ -tests/diagnostics/variable-redeclaration.slang(51): note 39999: candidate: float size - float size) - ^~~~ -tests/diagnostics/variable-redeclaration.slang(50): note 39999: candidate: int size - int size, - ^~~~ -} -standard output = { -} diff --git a/tests/front-end/struct.slang b/tests/front-end/struct.slang index 16b1bebc7f..075b39cb1a 100644 --- a/tests/front-end/struct.slang +++ b/tests/front-end/struct.slang @@ -7,8 +7,8 @@ // struct declaration struct Foo { - float3 a; - float3 b; + float3 a = {}; + float3 b = {}; }; // function on a struct @@ -46,4 +46,4 @@ template shader Test() out @Fragment vec4 colorTarget = vec4(result,1); } -*/ \ No newline at end of file +*/ diff --git a/tests/glsl-intrinsic/intrinsic-texture.slang b/tests/glsl-intrinsic/intrinsic-texture.slang index c40390c57c..e7ee87b59f 100644 --- a/tests/glsl-intrinsic/intrinsic-texture.slang +++ b/tests/glsl-intrinsic/intrinsic-texture.slang @@ -16,7 +16,7 @@ //TEST_INPUT:ubuffer(data=[0], stride=4):out,name=outputBuffer buffer MyBlockName { - int result; + int result = 0; } outputBuffer; // float diff --git a/tests/hlsl-intrinsic/size-of/align-of.slang b/tests/hlsl-intrinsic/size-of/align-of.slang index c5be345c8c..61b9cbfffb 100644 --- a/tests/hlsl-intrinsic/size-of/align-of.slang +++ b/tests/hlsl-intrinsic/size-of/align-of.slang @@ -15,14 +15,14 @@ enum Enum : int8_t struct Hey { - double b; - Enum e; + double b = {}; + Enum e = {}; }; struct Thing : Hey { - int a; - float b; + int a = {}; + float b = {}; }; [numthreads(4, 1, 1)] diff --git a/tests/hlsl-intrinsic/size-of/size-of.slang b/tests/hlsl-intrinsic/size-of/size-of.slang index fc531bc4ef..f6db0e1a52 100644 --- a/tests/hlsl-intrinsic/size-of/size-of.slang +++ b/tests/hlsl-intrinsic/size-of/size-of.slang @@ -15,14 +15,14 @@ enum Enum : int8_t struct Hey { - double b; - Enum e; + double b = {}; + Enum e = {}; }; struct Thing : Hey { - int a; - float b; + int a = {}; + float b = {}; }; [numthreads(4, 1, 1)] diff --git a/tests/language-feature/extensions/this-in-extension.slang b/tests/language-feature/extensions/this-in-extension.slang index 374eabe6fb..e3b5eea276 100644 --- a/tests/language-feature/extensions/this-in-extension.slang +++ b/tests/language-feature/extensions/this-in-extension.slang @@ -6,22 +6,28 @@ interface IFoo { - static const This identity; + static const This identityZero; + static const This identityFive; } __generic extension T { - This getIdentity() + This getIdentityZero() { - return identity; + return identityZero; + } + This getIdentityFive() + { + return identityFive; } } struct FooImpl : IFoo { int v = 1; - static const This identity = This(); + static const This identityZero = This(); + static const This identityFive = This(5); } @@ -32,7 +38,10 @@ RWStructuredBuffer outputBuffer; void computeMain(int3 dispatchThreadID : SV_DispatchThreadID) { FooImpl foo; - var ident = foo.getIdentity(); + var i0 = foo.getIdentityZero(); + var i5 = foo.getIdentityFive(); // CHECK: 1 - outputBuffer[0] = ident.v; + outputBuffer[0] = i0.v; + // CHECK: 5 + outputBuffer[1] = i5.v; } diff --git a/tests/language-feature/generics/generic-interface-1.slang b/tests/language-feature/generics/generic-interface-1.slang index 217e7f06f6..b51c4320f2 100644 --- a/tests/language-feature/generics/generic-interface-1.slang +++ b/tests/language-feature/generics/generic-interface-1.slang @@ -13,7 +13,7 @@ bool test(IEqlTestable v0, T v1) struct MyType : IEqlTestable { - int val; + int val = 0; bool testEql(MyType v1) { return val == v1.val; diff --git a/tests/language-feature/generics/where-4.slang b/tests/language-feature/generics/where-4.slang index 9f86ca7dc8..0e6523d59c 100644 --- a/tests/language-feature/generics/where-4.slang +++ b/tests/language-feature/generics/where-4.slang @@ -19,7 +19,7 @@ interface IFoo struct BarImpl : IBar { typealias TB = int; - int x; + int x = 0; int get() { return x; } } diff --git a/tests/language-feature/ifunc/functor.slang b/tests/language-feature/ifunc/functor.slang index 73987cbbf2..dc2661704f 100644 --- a/tests/language-feature/ifunc/functor.slang +++ b/tests/language-feature/ifunc/functor.slang @@ -5,7 +5,7 @@ struct Functor : IMutatingFunc { - int context; + int context = 0; [mutating] float operator()(float p) diff --git a/tests/language-feature/initializer-lists/constructor-inheritance.slang b/tests/language-feature/initializer-lists/constructor-inheritance.slang new file mode 100644 index 0000000000..90cb6bf4e5 --- /dev/null +++ b/tests/language-feature/initializer-lists/constructor-inheritance.slang @@ -0,0 +1,168 @@ +//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=BUF): -vk -compute -shaderobj + +// Test senario where we have constructors (partially defined through inheritance) +// and our init-list syntax calls into these constructors. + +public struct TestDerived1 +{ + private int val1 = 0; + public int val2 = 0; +} +public struct Test1 : TestDerived1 +{ + private int val3 = 0; + public int val4 = 0; +} + +public struct TestDerived2 +{ + private int val1 = 0; + public int val2 = 0; + + // This call overrides the memberwise constructor + // of the base type. + __init(int val) + { + this.val2 = val + 5; + } +} +public struct Test2 : TestDerived2 +{ + private int val3 = 0; + public int val4 = 0; + + // Does not override memberwise constructor + __init(int val) + { + val4 = val + 5; + } +} + +public struct TestDerived3 +{ + private int val1 = 0; + public int val2 = 0; + + // Does not get called since the derived type + // memberwise constructor is overloaded. + __init(int val) + { + this.val2 = val + 5; + } +} +public struct Test3 : TestDerived3 +{ + private int val3 = 0; + public int val4 = 0; + // This call overrides memberwise constructor, + // we ignore any inherited memberwise constructor. + __init(int val1, int val2) + { + this.val2 = val1 + 5; + this.val4 = val2 + 10; + } +} + +public struct TestDerived4 +{ + private int val1 = 0; + public int val2 = 0; +} +public struct Test4 : TestDerived4 +{ + private int val3 = 0; + public int val4 = 0; + __init(int val4, int val2) + { + this.val2 = val2 + 10; + this.val4 = val4; + } +} + +public struct TestDerived5 +{ + private int val1 = 0; + public int val2 = 0; +} +public struct Test5 : TestDerived5 +{ + private int val3 = 0; + public int val4 = 0; + + int getVal3() + { + return val3; + } + + static Test5 callPrivateMemberwiseCtor() + { + return { 1, 2, 3 }; + } +} + +public struct TestDerived6 +{ + private int val1 = 0; + public int val2 = 0; +} +public struct Test6 : TestDerived6 +{ + private int val3 = 0; + public int val4 = 0; + + __init(int val2, int val3, int val4) + { + this.val2 = val2*2; + this.val3 = val3*2; + this.val4 = val4*2; + } + + int getVal3() + { + return val3; + } + + static Test6 callPrivateMemberwiseCtor() + { + return { 1, 2, 3 }; + } +} + +//TEST_INPUT: ubuffer(data=[0], stride=4):out,name=outputBuffer +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test1 val1 = { 1, 2 }; + Test2 val2 = { 1, 2 }; + Test3 val3 = { 1, 2 }; + Test4 val4 = { 1, 2 }; + Test5 val5 = Test5::callPrivateMemberwiseCtor(); + Test6 val6 = Test6::callPrivateMemberwiseCtor(); + + outputBuffer[0] = (true + && val1.val2 == 1 + && val1.val4 == 2 + + && val2.val2 == 6 + && val2.val4 == 2 + + && val3.val2 == 6 + && val3.val4 == 12 + + && val4.val2 == 12 + && val4.val4 == 1 + + //&& val5.val2 == 1 + //&& val5.getVal3() == 2 + //&& val5.val4 == 3 + + //&& val5.val2 == 2 + //&& val5.getVal3() == 4 + //&& val5.val4 == 6 + + ) ? 1 : 0; + +// BUF: 1 +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/cstyle-init-list-2.slang b/tests/language-feature/initializer-lists/cstyle-init-list-2.slang new file mode 100644 index 0000000000..a5554299ae --- /dev/null +++ b/tests/language-feature/initializer-lists/cstyle-init-list-2.slang @@ -0,0 +1,41 @@ +//TEST:SIMPLE(filecheck=CHECK): -target dxil -entry computeMain -stage compute -profile sm_6_5 +//TEST(compute):COMPARE_COMPUTE(filecheck-buffer=BUF): -shaderobj +//TEST(compute, vulkan):COMPARE_COMPUTE(filecheck-buffer=BUF): -vk -shaderobj + +struct TestInner +{ + float4 a = {}; +}; + +TestInner makeTestInner(float v) +{ + TestInner val; + val.a = (float4)v; + return val; +} + +struct Test +{ + float3 a; + float3 b; + TestInner c; +}; + + +//CHECK: computeMain + +//TEST_INPUT: ubuffer(data=[0], stride=4):out,name=outputBuffer +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test test = { float3(1, 2, 3), (float3)2, makeTestInner(3) }; + +//BUF: 1 + outputBuffer[0] = true + && all(test.a == float3(1, 2, 3)) + && all(test.b == float3(2, 2, 2)) + && all(test.c.a == float4(3, 3, 3, 3)) + ; +} diff --git a/tests/language-feature/initializer-lists/cstyle-init-list.slang b/tests/language-feature/initializer-lists/cstyle-init-list.slang new file mode 100644 index 0000000000..228d7d4e05 --- /dev/null +++ b/tests/language-feature/initializer-lists/cstyle-init-list.slang @@ -0,0 +1,29 @@ +//TEST:SIMPLE(filecheck=CHECK): -target hlsl -entry computeMain -stage compute + +struct Test +{ + uint a; + uint b; +}; + + +//CHECK: computeMain + +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test test[3] = {1, 2, 3, 4, 5, 6}; + + outputBuffer[0] = true + && test[0].a == 1 + && test[0].b == 2 + + && test[1].a == 3 + && test[1].b == 4 + + && test[2].a == 5 + && test[2].b == 6 + ; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/default-init-16bit-types.slang b/tests/language-feature/initializer-lists/default-init-16bit-types.slang index 9620534e25..c5d0518fa7 100644 --- a/tests/language-feature/initializer-lists/default-init-16bit-types.slang +++ b/tests/language-feature/initializer-lists/default-init-16bit-types.slang @@ -1,6 +1,6 @@ // simple-namespace.slang -//TEST(compute):COMPARE_COMPUTE:-vk -render-feature int16 +//TEST(compute):COMPARE_COMPUTE(filecheck-buffer=BUF):-vk -render-feature int16 // Test that default initialization works with 16-bit types under Vulkan. @@ -31,5 +31,10 @@ void computeMain(int3 dispatchThreadID : SV_DispatchThreadID) int tid = dispatchThreadID.x; int inVal = tid; int outVal = test(inVal); + +//BUF: 0 +//BUF: 1111 +//BUF: 2222 +//BUF: 3333 outputBuffer[tid] = outVal; } diff --git a/tests/language-feature/initializer-lists/default-init-16bit-types.slang.expected.txt b/tests/language-feature/initializer-lists/default-init-16bit-types.slang.expected.txt deleted file mode 100644 index 6dbd5c939b..0000000000 --- a/tests/language-feature/initializer-lists/default-init-16bit-types.slang.expected.txt +++ /dev/null @@ -1,4 +0,0 @@ -0 -1111 -2222 -3333 diff --git a/tests/language-feature/initializer-lists/default-init-list-2.slang b/tests/language-feature/initializer-lists/default-init-list-2.slang new file mode 100644 index 0000000000..b45e1d64ec --- /dev/null +++ b/tests/language-feature/initializer-lists/default-init-list-2.slang @@ -0,0 +1,40 @@ +//TEST:SIMPLE(filecheck=CHECK): -target hlsl -entry computeMain -stage compute +//TEST(compute):COMPARE_COMPUTE(filecheck-buffer=BUF): -shaderobj + +//TEST_INPUT: ubuffer(data=[0], stride=4):out,name=outputBuffer +RWStructuredBuffer outputBuffer; + +public struct Test +{ +//CHECK: = 5 + public uint a; +//CHECK: = 0 + public uint b; + __init() + { + // This code will do 2 things: + // 1. call `$ZeroInit()` + // 2. assign `a = 5` + this = {}; + a = 5; + } +}; + +Test getDefault() +{ + Test test = {}; + return test; +} + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test test = getDefault(); + +//BUF: 1 + outputBuffer[0] = (true + && test.a == 5 + && test.b == 0 + ) ? 1 : 0 + ; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/default-init-list-inheritance.slang b/tests/language-feature/initializer-lists/default-init-list-inheritance.slang new file mode 100644 index 0000000000..d4feab9d6a --- /dev/null +++ b/tests/language-feature/initializer-lists/default-init-list-inheritance.slang @@ -0,0 +1,42 @@ +//TEST:SIMPLE(filecheck=CHECK): -target hlsl -entry computeMain -stage compute +//TEST(compute, vulkan):COMPARE_COMPUTE_EX(filecheck-buffer=BUF):-vk -compute -shaderobj -xslang -Wno-41021 + +//TEST_INPUT: ubuffer(data=[0], stride=4):out,name=outputBuffer +RWStructuredBuffer outputBuffer; + +public struct TestDerived +{ +//CHECK: = 5 + public uint a = 5; +//CHECK: = 0 + public uint b; +} + +public struct Test : TestDerived +{ +//CHECK: = 0 + public uint c; +//CHECK: = 6 + public uint d = 6; +}; + +Test getDefault() +{ + Test test = {}; + return test; +} + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test test = getDefault(); + +//BUF: 1 + outputBuffer[0] = (true + && test.a == 5 + && test.b == 0 + && test.c == 0 + && test.d == 6 + ) ? 1 : 0 + ; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/default-init-list.slang b/tests/language-feature/initializer-lists/default-init-list.slang new file mode 100644 index 0000000000..70cee85c41 --- /dev/null +++ b/tests/language-feature/initializer-lists/default-init-list.slang @@ -0,0 +1,32 @@ +//TEST:SIMPLE(filecheck=CHECK): -target hlsl -entry computeMain -stage compute +//TEST(compute):COMPARE_COMPUTE(filecheck-buffer=BUF): -shaderobj + +//TEST_INPUT: ubuffer(data=[0], stride=4):out,name=outputBuffer +RWStructuredBuffer outputBuffer; + +public struct Test +{ +//CHECK: = 5 + public uint a = 5; +//CHECK: = 0 + public uint b; +}; + +Test getDefault() +{ + Test test = {}; + return test; +} + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test test = getDefault(); + +//BUF: 1 + outputBuffer[0] = (true + && test.a == 5 + && test.b == 0 + ) ? 1 : 0 + ; +} \ No newline at end of file diff --git a/tests/compute/init-list-defaults.slang b/tests/language-feature/initializer-lists/defaults.slang similarity index 78% rename from tests/compute/init-list-defaults.slang rename to tests/language-feature/initializer-lists/defaults.slang index 0494501098..379ccb75d0 100644 --- a/tests/compute/init-list-defaults.slang +++ b/tests/language-feature/initializer-lists/defaults.slang @@ -1,6 +1,6 @@ // init-list-defaults.slang -//TEST(compute):COMPARE_COMPUTE: -shaderobj -//TEST(compute):COMPARE_COMPUTE:-cpu -shaderobj +//TEST(compute):COMPARE_COMPUTE(filecheck-buffer=BUF): -shaderobj +//TEST(compute):COMPARE_COMPUTE(filecheck-buffer=BUF):-cpu -shaderobj // Confirm that initializer lists correctly default-initialize elements past those specified. @@ -34,5 +34,9 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) int inVal = int(tid); int outVal = test(inVal); +//BUF: 1142 +//BUF: 2453 +//BUF: 3060 +//BUF: 4000 outputBuffer[tid] = outVal; } \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/deprecated-init-list-behavior-1.slang b/tests/language-feature/initializer-lists/deprecated-init-list-behavior-1.slang new file mode 100644 index 0000000000..0dbf2da46f --- /dev/null +++ b/tests/language-feature/initializer-lists/deprecated-init-list-behavior-1.slang @@ -0,0 +1,38 @@ +//TEST:SIMPLE(filecheck=CHECK): -target hlsl -entry computeMain -stage compute + +struct InnerTest +{ + uint a; +}; + +struct Test +{ + InnerTest a; + uint b; + float2 c; +}; + +struct TestOuter : Test +{ + __init(float val) + { + } +}; + +//CHECK: warning 30505 +//CHECK: computeMain + +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + TestOuter test = {1, 1, float3(1).x, float3(1).z}; + + outputBuffer[0] = true + && test.a.a == 1 + && test.b == 1 + && test.c[0] == 1 + && test.c[1] == 1 + ; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/deprecated-init-list-behavior-2.slang b/tests/language-feature/initializer-lists/deprecated-init-list-behavior-2.slang new file mode 100644 index 0000000000..1a4a80faa9 --- /dev/null +++ b/tests/language-feature/initializer-lists/deprecated-init-list-behavior-2.slang @@ -0,0 +1,35 @@ +//TEST:SIMPLE(filecheck=CHECK): -target hlsl -entry computeMain -stage compute + +struct Test +{ + uint a; + uint b; + __init(int a_in, int b_in) + { + a_in = 0; + b_in = 0; + } +}; + + +//CHECK: warning 30505 +//CHECK: computeMain + +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test test[3] = {1, 2, 3, 4, 5, 6}; + + outputBuffer[0] = true + && test[0].a == 1 + && test[0].b == 2 + + && test[1].a == 1 + && test[1].b == 2 + + && test[2].a == 1 + && test[2].b == 2 + ; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/dont-synth-constructor-conflicting-default.slang b/tests/language-feature/initializer-lists/dont-synth-constructor-conflicting-default.slang new file mode 100644 index 0000000000..73dd55409e --- /dev/null +++ b/tests/language-feature/initializer-lists/dont-synth-constructor-conflicting-default.slang @@ -0,0 +1,23 @@ +//TEST:SIMPLE(filecheck=CHECK): -target hlsl -entry computeMain -stage compute + +struct Test +{ + uint a; + uint b; + __init(int b = 16) + { + this.b = b; + } +}; + +RWStructuredBuffer outputBuffer; + +//CHECK: computeMain + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test test = Test(); + + outputBuffer[0] = test.a; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/dont-synth-internal-constructor.slang b/tests/language-feature/initializer-lists/dont-synth-internal-constructor.slang new file mode 100644 index 0000000000..261c84a403 --- /dev/null +++ b/tests/language-feature/initializer-lists/dont-synth-internal-constructor.slang @@ -0,0 +1,38 @@ +//TEST:SIMPLE(filecheck=ERROR): -target spirv -entry computeMain -stage compute + +// Test senario where an 'private' member is not default-initialized +// for a struct but we call into a "internal" memberwise ctor. +// The called "internal" memberwise ctor should not be synth'ed. + +//ERROR: error 30500 +public struct TestNested +{ + public int val1; + public int val2; +} + +public struct Test +{ + private int val1; + public int val2; + public TestNested nested1; + + int getVal1() + { + return val1; + } +}; + +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test val1 = { 1, { 2, 3 } }; + outputBuffer[0] = (true + && val1.getVal1() == 2 + && val1.val2 == 2 + && val1.nested1.val1 == 2 + && val1.nested1.val2 == 2 + ) ? 1 : 0; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/dont-synth-public-constructor-helper.slang b/tests/language-feature/initializer-lists/dont-synth-public-constructor-helper.slang new file mode 100644 index 0000000000..8b128415c0 --- /dev/null +++ b/tests/language-feature/initializer-lists/dont-synth-public-constructor-helper.slang @@ -0,0 +1,23 @@ +// Have a seperate file to ensure we are not trying to call an 'internal' synth'ed constructor. + +public struct TestNested +{ + public int val1; + public int val2; +} + +public struct Test +{ +#ifdef TEST_INTERNAL + internal int val1; +#else + private int val1; +#endif + public int val2; + public TestNested nested1; + + public int getVal1() + { + return val1; + } +}; \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/dont-synth-public-constructor.slang b/tests/language-feature/initializer-lists/dont-synth-public-constructor.slang new file mode 100644 index 0000000000..646e522cd8 --- /dev/null +++ b/tests/language-feature/initializer-lists/dont-synth-public-constructor.slang @@ -0,0 +1,23 @@ +//TEST:SIMPLE(filecheck=ERROR): -target spirv -entry computeMain -stage compute -DTEST_INTERNAL +//TEST:SIMPLE(filecheck=ERROR): -target spirv -entry computeMain -stage compute + +// Test senario where a 'private'/'internal' member is not default-initialized +// for a struct but we call into a "public" memberwise ctor. +// The called "public" memberwise ctor should not be synth'ed. + +// ERROR: error 30500 +import dont_synth_public_constructor_helper; + +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test val1 = { 1, { 2, 3 } }; + outputBuffer[0] = (true + && val1.getVal1() == 2 + && val1.val2 == 2 + && val1.nested1.val1 == 2 + && val1.nested1.val2 == 2 + ) ? 1 : 0; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/init-list-with-autodiff-attributes-1.slang b/tests/language-feature/initializer-lists/init-list-with-autodiff-attributes-1.slang new file mode 100644 index 0000000000..a747686988 --- /dev/null +++ b/tests/language-feature/initializer-lists/init-list-with-autodiff-attributes-1.slang @@ -0,0 +1,41 @@ +//TEST:SIMPLE(filecheck=CHECK): -target hlsl -entry computeMain -stage compute + +struct TestInner : IDifferentiable +{ + typealias Differential = TestInner; + int a; +} +struct Test : IDifferentiable +{ + uint a; + uint b; + TestInner test; + + [Differentiable] + __init(uint data1, uint data2, TestInner data) + { + a = data1; + b = data2; + test = {}; + } +}; + +struct Test2 : IDifferentiable +{ + typealias Differential = Test2; + TestInner a; +} + +//CHECK: computeMain + +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test.Differential test = { {} }; + + outputBuffer[0] = true + && test.test.a == 0 + ; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/init-list-with-autodiff-attributes-2.slang b/tests/language-feature/initializer-lists/init-list-with-autodiff-attributes-2.slang new file mode 100644 index 0000000000..aa35c4d412 --- /dev/null +++ b/tests/language-feature/initializer-lists/init-list-with-autodiff-attributes-2.slang @@ -0,0 +1,30 @@ +//TEST:SIMPLE(filecheck=CHECK): -target hlsl -entry computeMain -stage compute + +struct Test : IDifferentiable +{ + float3 a; + no_diff float3 b; +}; + +//CHECK: computeMain + +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test.Differential test1 = {}; + Test test2 = {}; + outputBuffer[0] = true + && test1.a[0] == 0 + && test1.a[1] == 0 + && test1.a[2] == 0 + + && test1.a[0] == 0 + && test1.a[1] == 0 + && test1.a[2] == 0 + && test1.b[0] == 0 + && test1.b[1] == 0 + && test1.b[2] == 0 + ; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/init-list-with-autodiff-attributes-3.slang b/tests/language-feature/initializer-lists/init-list-with-autodiff-attributes-3.slang new file mode 100644 index 0000000000..17df592a46 --- /dev/null +++ b/tests/language-feature/initializer-lists/init-list-with-autodiff-attributes-3.slang @@ -0,0 +1,31 @@ +//TEST:SIMPLE(filecheck=CHECK): -target hlsl -entry computeMain -stage compute + +struct Test1 : IDifferentiable +{ + float3 a; +}; +struct TestWrapper +{ + Test1 data1; +} + +//CHECK: computeMain + +[Differentiable] +TestWrapper getWrapper() +{ + return { float3(1) }; +} + +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + TestWrapper test1 = getWrapper(); + outputBuffer[0] = true + && test1.data1.a[0] == 1 + && test1.data1.a[1] == 1 + && test1.data1.a[2] == 1 + ; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/init-list-with-autodiff-attributes-4.slang b/tests/language-feature/initializer-lists/init-list-with-autodiff-attributes-4.slang new file mode 100644 index 0000000000..9791bbe8ca --- /dev/null +++ b/tests/language-feature/initializer-lists/init-list-with-autodiff-attributes-4.slang @@ -0,0 +1,29 @@ +//TEST:SIMPLE(filecheck=CHECK): -target hlsl -entry computeMain -stage compute + +struct Test1 : IDifferentiable +{ + float3 a; + no_diff float3 b; +}; + +//CHECK: computeMain + +[Differentiable] +Test1 getWrapper() +{ + return { float3(1), float3(1) }; +} + +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test1 test1 = getWrapper(); + outputBuffer[0] = true + && test1.a[0] == 1 + && test1.a[1] == 1 + && test1.a[2] == 1 + && test1.b[0] == 1 + ; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/partial-init-list.slang b/tests/language-feature/initializer-lists/partial-init-list.slang new file mode 100644 index 0000000000..a63508d498 --- /dev/null +++ b/tests/language-feature/initializer-lists/partial-init-list.slang @@ -0,0 +1,41 @@ +//TEST:SIMPLE(filecheck=CHECK): -target hlsl -entry computeMain -stage compute +//TEST(compute):COMPARE_COMPUTE(filecheck-buffer=BUF): -shaderobj + +//CHECK: computeMain + +//TEST_INPUT: ubuffer(data=[0], stride=4):out,name=outputBuffer +RWStructuredBuffer outputBuffer; + +public struct Test +{ + public uint a; + public uint b; +}; + +Test getTest1() +{ + Test test = {1, 2}; + return test; +} +Test getTest2() +{ + Test test = {3}; + return test; +} + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test test1 = getTest1(); + Test test2 = getTest2(); + +//BUF: 1 + outputBuffer[0] = (true + && test1.a == 1 + && test1.b == 2 + + && test2.a == 3 + && test2.b == 0 + ) ? 1 : 0 + ; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/readonly-member.slang b/tests/language-feature/initializer-lists/readonly-member.slang new file mode 100644 index 0000000000..25f620cb06 --- /dev/null +++ b/tests/language-feature/initializer-lists/readonly-member.slang @@ -0,0 +1,21 @@ +//TEST:SIMPLE(filecheck=CHECK): -target hlsl -allow-glsl -entry computeMain -stage compute + +struct Test +{ + readonly uint a; + writeonly uint b; +}; + +//CHECK: computeMain + +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test test = {2}; + + outputBuffer[0] = true + && test.b == 2; + ; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/struct-swizzle-initializer-list.slang b/tests/language-feature/initializer-lists/struct-swizzle-initializer-list.slang new file mode 100644 index 0000000000..15ec428023 --- /dev/null +++ b/tests/language-feature/initializer-lists/struct-swizzle-initializer-list.slang @@ -0,0 +1,29 @@ +//TEST:SIMPLE(filecheck=CHECK): -target dxil -entry computeMain -stage compute -profile sm_6_5 + +//CHECK: computeMain + +RWStructuredBuffer outputBuffer; + +struct Test +{ + float3 a; + float3 b; + float2 c; +}; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + float3 position = {1, 2, 3}; + const Test test[3] = + { + mul(float4x4(1), float4(position, 1.f)).xyz, mul(float4x4(1), float4(position, 1.f)).xyz - mul(float4x4(1), float4(position, 1.f)).xyz, float2(2), + mul(float4x4(1), float4(position, 1.f)).xyz, mul(float4x4(1), float4(position, 1.f)).xyz - mul(float4x4(1), float4(position, 1.f)).xyz, float2(2), + mul(float4x4(1), float4(position, 1.f)).xyz, mul(float4x4(1), float4(position, 1.f)).xyz - mul(float4x4(1), float4(position, 1.f)).xyz, float2(2) + }; + outputBuffer[0] = true + && test[0].a[0] != 99 + && test[0].b[0] != 99 + && test[1].c[0] != 99 + ; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/synth-constructor-private-internal-helper.slang b/tests/language-feature/initializer-lists/synth-constructor-private-internal-helper.slang new file mode 100644 index 0000000000..876e180345 --- /dev/null +++ b/tests/language-feature/initializer-lists/synth-constructor-private-internal-helper.slang @@ -0,0 +1,48 @@ +public struct TestNested +{ + private int val1 = 0; + internal int val2; + + public int getVal1() + { + return val1; + } + public int getVal2() + { + return val2; + } +} + +public struct Test +{ + private int val1 = 0; + internal int val2; + private TestNested nested1 = { 0 }; + private TestNested nested2 = { 0 }; + + public static Test getTest() + { + return { 1, 2, { 3 }, { 4 }}; + } + public int getVal1() + { + return val1; + } + public int getVal2() + { + return val2; + } + public TestNested getNested1() + { + return nested1; + } + public TestNested getNested2() + { + return nested2; + } +}; + +public Test getTest() +{ + return { 1 }; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/synth-constructor-private-internal.slang b/tests/language-feature/initializer-lists/synth-constructor-private-internal.slang new file mode 100644 index 0000000000..64c33b9811 --- /dev/null +++ b/tests/language-feature/initializer-lists/synth-constructor-private-internal.slang @@ -0,0 +1,27 @@ +//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=BUF): -vk -compute -shaderobj + +import synth_constructor_private_internal_helper; + +//TEST_INPUT: ubuffer(data=[0], stride=4):out,name=outputBuffer +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test val1 = getTest(); + Test val2 = Test::getTest(); + outputBuffer[0] = (true + && val1.getVal2() == 1 + && val1.getNested1().getVal1() == 0 + && val1.getNested1().getVal2() == 0 + && val1.getNested2().getVal1() == 0 + && val1.getNested2().getVal2() == 0 + + && val2.getVal1() == 1 + && val2.getVal2() == 2 + && val2.getNested1().getVal2() == 3 + && val2.getNested2().getVal2() == 4 + ) ? 1 : 0; + +// BUF: 1 +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/synth-constructor-private.slang b/tests/language-feature/initializer-lists/synth-constructor-private.slang new file mode 100644 index 0000000000..e79114fca5 --- /dev/null +++ b/tests/language-feature/initializer-lists/synth-constructor-private.slang @@ -0,0 +1,62 @@ +//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=BUF): -vk -compute -shaderobj + +public struct TestNested +{ + private int val1 = {}; + private int val2 = {}; + + int getVal1() + { + return val1; + } + int getVal2() + { + return val2; + } +} + +public struct Test +{ + private int val1; + private int val2; + private TestNested nested1; + + static Test getTest() + { + return { 1, 2, {}}; + } + + int getVal1() + { + return val1; + } + int getVal2() + { + return val2; + } + TestNested getNested1() + { + return nested1; + } +}; + +//TEST_INPUT: ubuffer(data=[0], stride=4):out,name=outputBuffer +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test val1 = {}; + Test val2 = Test::getTest(); + outputBuffer[0] = (true + && val1.getVal1() == 0 + && val1.getNested1().getVal1() == 0 + + && val2.getVal1() == 1 + && val2.getVal2() == 2 + && val2.getNested1().getVal1() == 0 + && val2.getNested1().getVal2() == 0 + ) ? 1 : 0; + +// BUF: 1 +} diff --git a/tests/language-feature/initializer-lists/synth-constructor-public-private-internal-helper.slang b/tests/language-feature/initializer-lists/synth-constructor-public-private-internal-helper.slang new file mode 100644 index 0000000000..81c80a0db8 --- /dev/null +++ b/tests/language-feature/initializer-lists/synth-constructor-public-private-internal-helper.slang @@ -0,0 +1,28 @@ +public struct NestedTest +{ + internal int val1 = 0; + private int val2 = 0; + public int val3; + + public int getVal1() + { + return val1; + } +} +public struct Test +{ + internal int val1 = 0; + private int val2 = 0; + public int val3; + public NestedTest nested1 = { 0, 0 }; + + public int getVal1() + { + return val1; + } +} + +public Test makeInternal() +{ + return { 1, 2, { 3, 4 } }; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/synth-constructor-public-private-internal.slang b/tests/language-feature/initializer-lists/synth-constructor-public-private-internal.slang new file mode 100644 index 0000000000..53355bd869 --- /dev/null +++ b/tests/language-feature/initializer-lists/synth-constructor-public-private-internal.slang @@ -0,0 +1,24 @@ +//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=BUF): -vk -compute -shaderobj + +import synth_constructor_public_private_internal_helper; + +//TEST_INPUT: ubuffer(data=[0], stride=4):out,name=outputBuffer +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test val1 = { 1, { 2 } }; + Test val2 = makeInternal(); + outputBuffer[0] = (true + && val1.val3 == 1 + && val1.nested1.val3 == 2 + + && val2.getVal1() == 1 + && val2.val3 == 2 + && val2.nested1.getVal1() == 3 + && val2.nested1.val3 == 4 + ) ? 1 : 0; + +// BUF: 1 +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/synth-constructor-public-private.slang b/tests/language-feature/initializer-lists/synth-constructor-public-private.slang new file mode 100644 index 0000000000..181ccffe94 --- /dev/null +++ b/tests/language-feature/initializer-lists/synth-constructor-public-private.slang @@ -0,0 +1,50 @@ +//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=BUF): -vk -compute -shaderobj + +public struct TestNested +{ + private int val1 = 0; + public int val2; +} + +public struct Test +{ + private int val1 = 0; + public int val2; + public TestNested nested1 = { 0 }; + private TestNested nested2 = { 0 }; + + static Test getTest() + { + return { 1, 2, { 3 }, { 4 } }; + } + + int getVal1() + { + return val1; + } + TestNested getNested2() + { + return nested2; + } +}; + +//TEST_INPUT: ubuffer(data=[0], stride=4):out,name=outputBuffer +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test val1 = { 1, { 2 } }; + Test val2 = Test::getTest(); + outputBuffer[0] = (true + && val1.val2 == 1 + && val1.nested1.val2 == 2 + + && val2.getVal1() == 1 + && val2.val2 == 2 + && val2.nested1.val2 == 3 + && val2.getNested2().val2 == 4 + ) ? 1 : 0; + +// BUF: 1 +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/synth-constructor-public.slang b/tests/language-feature/initializer-lists/synth-constructor-public.slang new file mode 100644 index 0000000000..9e942cf54c --- /dev/null +++ b/tests/language-feature/initializer-lists/synth-constructor-public.slang @@ -0,0 +1,31 @@ +//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=BUF): -vk -compute -shaderobj + +public struct TestNested +{ + public int val1 = {}; + public int val2 = {}; +} + +public struct Test +{ + public int val1; + public int val2; + public TestNested nested1; +}; + +//TEST_INPUT: ubuffer(data=[0], stride=4):out,name=outputBuffer +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test val = { 1, 2, { 3, 4 } }; + outputBuffer[0] = (true + && val.val1 == 1 + && val.val2 == 2 + && val.nested1.val1 == 3 + && val.nested1.val2 == 4 + ) ? 1 : 0; + +// BUF: 1 +} diff --git a/tests/language-feature/initializer-lists/synth-constructor-visibility-error.slang b/tests/language-feature/initializer-lists/synth-constructor-visibility-error.slang new file mode 100644 index 0000000000..6895c9f625 --- /dev/null +++ b/tests/language-feature/initializer-lists/synth-constructor-visibility-error.slang @@ -0,0 +1,37 @@ +//TEST:SIMPLE(filecheck=ERROR): -target spirv -entry computeMain -stage compute + +// Test senario where user calls into a synth'ed memberwise constructor +// of "private" visibility from an "internal" visible scope. + +//ERROR: error 30600 +public struct TestNested +{ + public int val1; + public int val2; +} + +public struct Test +{ + private int val1; + public int val2; + public TestNested nested1; + + int getVal1() + { + return val1; + } +}; + +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test val1 = { 2, 1, { 1, 1 } }; + outputBuffer[0] = (true + && val1.getVal1() == 2 + && val1.val2 == 2 + && val1.nested1.val1 == 2 + && val1.nested1.val2 == 2 + ) ? 1 : 0; +} \ No newline at end of file diff --git a/tests/compute/initializer-list.slang b/tests/language-feature/initializer-lists/synth-constructor.slang similarity index 62% rename from tests/compute/initializer-list.slang rename to tests/language-feature/initializer-lists/synth-constructor.slang index 85ad0d8cfb..ac683ae79f 100644 --- a/tests/compute/initializer-list.slang +++ b/tests/language-feature/initializer-lists/synth-constructor.slang @@ -1,5 +1,5 @@ -//TEST(compute):COMPARE_COMPUTE: -shaderobj -//TEST(compute):COMPARE_COMPUTE:-cpu -shaderobj +//TEST(compute):COMPARE_COMPUTE(filecheck-buffer=BUF): -shaderobj +//TEST(compute):COMPARE_COMPUTE(filecheck-buffer=BUF):-cpu -shaderobj struct Test { @@ -11,7 +11,7 @@ struct Test uint test(uint val) { - Test t = { float4(1.0f), 16, 99.0f }; + Test t = { { 1.0f, 1.0f, 1.0f, 1.0f }, 16, 99.0f }; return val + t.b; } @@ -26,5 +26,9 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) uint inVal = tid; uint outVal = test(inVal); +//BUF: 10 +//BUF-NEXT: 11 +//BUF-NEXT: 12 +//BUF-NEXT: 13 outputBuffer[tid] = outVal; } \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/torch-tensor.slang b/tests/language-feature/initializer-lists/torch-tensor.slang new file mode 100644 index 0000000000..f5ea1ac493 --- /dev/null +++ b/tests/language-feature/initializer-lists/torch-tensor.slang @@ -0,0 +1,71 @@ +//TEST:SIMPLE(filecheck=CUDA): -target cuda -line-directive-mode none +//TEST:SIMPLE(filecheck=TORCH): -target torch -line-directive-mode none + + +struct MySubType +{ + TorchTensor array[2]; +} + +struct MyType +{ + float2 v; + MySubType sub[2]; +} + +struct MyType2 +{ + float2 v; + TorchTensor tensor; +} + +struct ReturnType +{ + MyType t1; + MyType2 t2; +} + +struct MyInput +{ + TorchTensor inValues; + float normalVal; +} + +// CUDA: __global__ void myKernel(TensorView inValues_[[#]], TensorView outValues_[[#]]) +[CudaKernel] +void myKernel(TensorView inValues, TensorView outValues) +{ + if (cudaThreadIdx().x > 0) + return; + outValues.store(cudaThreadIdx().x, sin(inValues.load(cudaThreadIdx().x))); +} + +// TORCH: {{^SLANG_PRELUDE_EXPORT$}} +// TORCH-NEXT: void myKernel(TensorView {{[[:alnum:]_]+}}, TensorView {{[[:alnum:]_]+}}); +// +// TORCH: {{^SLANG_PRELUDE_EXPORT$}} +// TORCH-NEXT: runCompute(std::tuple input_[[#]]) +[TorchEntryPoint] +export __extern_cpp ReturnType runCompute(MyInput input) +{ + MyType rs; + var outValues = TorchTensor.alloc(1); + let inValues = input.inValues; + + __dispatch_kernel(myKernel, uint3(1, 1, 1), uint3(32, 1, 1))(inValues, outValues); + + rs.v = float2(1.0, 2.0); + rs.sub[0].array[0] = outValues; + rs.sub[0].array[1] = inValues; + + rs.sub[1].array[0] = inValues; + rs.sub[1].array[1] = outValues; + + MyType2 rs2; + rs2.tensor = outValues; + + ReturnType returnVal; + returnVal.t1 = rs; + returnVal.t2 = rs2; + return returnVal; +} diff --git a/tests/language-feature/initializer-lists/user-constructor-visibility-error.slang b/tests/language-feature/initializer-lists/user-constructor-visibility-error.slang new file mode 100644 index 0000000000..802f03fb9e --- /dev/null +++ b/tests/language-feature/initializer-lists/user-constructor-visibility-error.slang @@ -0,0 +1,45 @@ +//TEST:SIMPLE(filecheck=ERROR): -target spirv -entry computeMain -stage compute + +// Test senario where user calls into a user defined constructor +// of "private" visibility from an "internal" visible scope. + +//ERROR: error 30600 +public struct TestNested +{ + public int val1; + public int val2; +} + +public struct Test +{ + private int val1; + public int val2; + public TestNested nested1; + + int getVal1() + { + return val1; + } + + private __init(int in1) + { + this.val1 = in1; + this.val2 = in1; + this.nested1.val1 = in1; + this.nested1.val2 = in1; + } +}; + +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test val1 = { 2 }; + outputBuffer[0] = (true + && val1.getVal1() == 2 + && val1.val2 == 2 + && val1.nested1.val1 == 2 + && val1.nested1.val2 == 2 + ) ? 1 : 0; +} \ No newline at end of file diff --git a/tests/language-feature/initializer-lists/user-constructor.slang b/tests/language-feature/initializer-lists/user-constructor.slang new file mode 100644 index 0000000000..69074b6874 --- /dev/null +++ b/tests/language-feature/initializer-lists/user-constructor.slang @@ -0,0 +1,75 @@ +//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=BUF): -vk -compute -shaderobj + +// Test senario where user defines constructors +// and our init-list syntax calls into those constructors + +public struct TestNested +{ + public int val1 = {}; + public int val2 = {}; +} + +public struct Test +{ + private int val1; + public int val2; + public TestNested nested1; + + int getVal1() + { + return val1; + } + + __init(int in1, TestNested in2) + { + this.val1 = in1; + this.val2 = in1; + this.nested1 = in2; + this.nested1.val1 = in1; + } + + __init(int in1) + { + this.val1 = in1; + this.val2 = in1; + this.nested1.val1 = in1; + this.nested1.val2 = in1; + } + + __init() + { + this.val1 = 5; + this.val2 = 5; + this.nested1.val1 = 5; + this.nested1.val2 = 5; + } +}; + +//TEST_INPUT: ubuffer(data=[0], stride=4):out,name=outputBuffer +RWStructuredBuffer outputBuffer; + +[numthreads(1, 1, 1)] +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + Test val1 = { 1, { 2, 3 } }; + Test val2 = { 2 }; + Test val3 = {}; + outputBuffer[0] = (true + && val1.getVal1() == 1 + && val1.val2 == 1 + && val1.nested1.val1 == 1 + && val1.nested1.val2 == 3 + + && val2.getVal1() == 2 + && val2.val2 == 2 + && val2.nested1.val1 == 2 + && val2.nested1.val2 == 2 + + && val3.getVal1() == 5 + && val3.val2 == 5 + && val3.nested1.val1 == 5 + && val3.nested1.val2 == 5 + ) ? 1 : 0; + +// BUF: 1 +} diff --git a/tests/language-feature/struct-in-func.slang b/tests/language-feature/struct-in-func.slang index d3bd046f08..21bdda11b7 100644 --- a/tests/language-feature/struct-in-func.slang +++ b/tests/language-feature/struct-in-func.slang @@ -5,8 +5,8 @@ int test(int r) { struct Val { - int a; - int b; + int a = 0; + int b = 0; } Val a; a.a = 1; diff --git a/tests/language-feature/swizzles/matrix-swizzle-write-array.slang b/tests/language-feature/swizzles/matrix-swizzle-write-array.slang index 616a19b198..907d5f51c1 100644 --- a/tests/language-feature/swizzles/matrix-swizzle-write-array.slang +++ b/tests/language-feature/swizzles/matrix-swizzle-write-array.slang @@ -11,7 +11,7 @@ RWStructuredBuffer outputBuffer; struct MySubscriptable { - float2x2 m; + float2x2 m = {}; __subscript() -> float2x2 { get { return m; } diff --git a/tests/language-feature/types/optional.slang b/tests/language-feature/types/optional.slang index ae1711cd7a..a945f0a624 100644 --- a/tests/language-feature/types/optional.slang +++ b/tests/language-feature/types/optional.slang @@ -18,7 +18,7 @@ interface IFoo //TEST_INPUT: type_conformance Impl1:IFoo = 0 struct Impl1 : IFoo { - int data; + int data = 0; int method() { return data; } } diff --git a/tests/legalization/vec1.slang b/tests/legalization/vec1.slang index f3de085b00..5ec79e0847 100644 --- a/tests/legalization/vec1.slang +++ b/tests/legalization/vec1.slang @@ -15,13 +15,13 @@ RWStructuredBuffer outputBuffer; struct V { // 1-vector of 1-vector - vector, 1> oo; + vector, 1> oo = {}; // 1-vector of n-vector - vector, 1> on; + vector, 1> on = {}; // n-vector of 1-vector - vector, 4> no; + vector, 4> no = {}; }; vector get1Vec(int x) diff --git a/tests/pipeline/rasterization/mesh/hlsl-syntax.slang b/tests/pipeline/rasterization/mesh/hlsl-syntax.slang index 76ac3a598e..9e4614660e 100644 --- a/tests/pipeline/rasterization/mesh/hlsl-syntax.slang +++ b/tests/pipeline/rasterization/mesh/hlsl-syntax.slang @@ -2,7 +2,21 @@ // Test that we can ingest hlsl mesh output syntax -//TEST:CROSS_COMPILE:-target spirv -profile sm_6_5 -entry main -stage mesh +//TEST:SIMPLE(filecheck=CHECK):-target spirv -profile sm_6_5 -entry main -stage mesh + +//CHECK: OpCapability MeshShadingEXT +//CHECK: OpExtension "SPV_EXT_mesh_shader" +//CHECK-DAG: OpEntryPoint MeshEXT +//CHECK-DAG: OpExecutionMode %{{.*}} OutputPrimitivesEXT 1 +//CHECK-DAG: OpExecutionMode %{{.*}} OutputTrianglesEXT + +//CHECK: OpDecorate %gl_LocalInvocationIndex BuiltIn LocalInvocationIndex +//CHECK-NOT: OpDecorate %gl_PrimitiveTriangleIndicesEXT BuiltIn PrimitiveTriangleIndicesEXT + +//CHECK: %gl_LocalInvocationIndex = OpVariable +//CHECK-NOT: %gl_PrimitiveTriangleIndicesEXT = OpVariable + +//CHECK: OpSetMeshOutputsEXT %uint_3 %uint_1 const static float2 positions[3] = { float2(0.0, -0.5), diff --git a/tests/pipeline/rasterization/mesh/nested-component-write.slang b/tests/pipeline/rasterization/mesh/nested-component-write.slang index 681331a51a..273eae1037 100644 --- a/tests/pipeline/rasterization/mesh/nested-component-write.slang +++ b/tests/pipeline/rasterization/mesh/nested-component-write.slang @@ -7,7 +7,7 @@ struct Foo { - float4 pos : SV_Position; + float4 pos : SV_Position = {}; Bar bar; }; @@ -18,7 +18,7 @@ struct Bar struct Baz { - float3 color; + float3 color = {}; }; struct Vertex diff --git a/tests/pipeline/rasterization/mesh/passing-outputs.slang b/tests/pipeline/rasterization/mesh/passing-outputs.slang index 3fd6fc4a0d..45d9405ab1 100644 --- a/tests/pipeline/rasterization/mesh/passing-outputs.slang +++ b/tests/pipeline/rasterization/mesh/passing-outputs.slang @@ -2,7 +2,7 @@ // This tests that writing to individual components of the output struct works -//TEST:SIMPLE(filecheck=SPIRV): -target spirv-assembly -entry main -stage mesh -profile glsl_450+spirv_1_4 +//TEST:SIMPLE(filecheck=SPIRV): -target spirv-assembly -entry main -stage mesh -profile glsl_450+GL_EXT_mesh_shader // DXC is stricter than we are about passing references to individual mesh shader outputs // We could get around this by doing what we do for GLSL, i.e. use a temporary diff --git a/tests/pipeline/ray-tracing/trace-ray-inline.slang b/tests/pipeline/ray-tracing/trace-ray-inline.slang index b8688c2f5c..09b470a91b 100644 --- a/tests/pipeline/ray-tracing/trace-ray-inline.slang +++ b/tests/pipeline/ray-tracing/trace-ray-inline.slang @@ -1,14 +1,24 @@ // trace-ray-inline.slang -//TEST:CROSS_COMPILE:-target dxil-asm -stage compute -profile sm_6_5 -entry main -line-directive-mode none -//TEST:SIMPLE(filecheck=CHECK):-target spirv-asm -stage compute -profile glsl_460+GL_EXT_ray_query -entry main -line-directive-mode none - -// CHECK: OpCapability RayQueryKHR -// CHECK: OpExtension "SPV_KHR_ray_query" -// CHECK: OpRayQueryInitializeKHR -// CHECK: OpRayQueryProceedKHR -// CHECK: OpRayQueryGetIntersectionTypeKHR -// CHECK: OpRayQueryConfirmIntersectionKHR +//TEST:SIMPLE(filecheck=HLSL):-target hlsl -stage compute -profile sm_6_5 -entry main +//TEST:SIMPLE(filecheck=DXIL):-target dxil -stage compute -profile sm_6_5 -entry main +//TEST:SIMPLE(filecheck=SPV):-target spirv-asm -stage compute -profile glsl_460+GL_EXT_ray_query -entry main + +// HLSL: {{.*}}.TraceRayInline +// HLSL: {{.*}}.Proceed +// HLSL: {{.*}}.CandidateType +// HLSL: {{.*}}.CommitProceduralPrimitiveHit +// HLSL: {{.*}}.CommitNonOpaqueTriangleHit + +/// ensure the compiled HLSL is valid +// DXIL: @main + +// SPV: OpCapability RayQueryKHR +// SPV: OpExtension "SPV_KHR_ray_query" +// SPV: OpRayQueryInitializeKHR +// SPV: OpRayQueryProceedKHR +// SPV: OpRayQueryGetIntersectionTypeKHR +// SPV: OpRayQueryConfirmIntersectionKHR // The goal of this shader is to use all the main pieces // of functionality in DXR 1.1's `TraceRayInline` feature, diff --git a/tests/spirv/array-uniform-param.slang b/tests/spirv/array-uniform-param.slang index 235e85bbd0..5379e71161 100644 --- a/tests/spirv/array-uniform-param.slang +++ b/tests/spirv/array-uniform-param.slang @@ -10,12 +10,12 @@ RWStructuredBuffer resultBuffer; struct Data { - uint arr[4]; + uint arr[4] = {}; } struct Data2 { - uint arr[4u]; + uint arr[4u] = {}; } //TEST_INPUT: set Uniforms.d = new Data{{1,2,3,4}}; diff --git a/tests/spirv/bit-field.slang b/tests/spirv/bit-field.slang index 878b4483dc..412ea064a8 100644 --- a/tests/spirv/bit-field.slang +++ b/tests/spirv/bit-field.slang @@ -15,7 +15,7 @@ enum Flag } struct Data { - uint flags; + uint flags = {}; [mutating] void setFlag(Flag g, bool v) { const uint bit = uint(g) << 8; diff --git a/tests/spirv/debug-info.slang b/tests/spirv/debug-info.slang index dc225f8ccb..c12b3d5813 100644 --- a/tests/spirv/debug-info.slang +++ b/tests/spirv/debug-info.slang @@ -2,8 +2,8 @@ struct TestType { - float memberA; - float3 memberB; + float memberA = {}; + float3 memberB = {}; RWStructuredBuffer memberC; float getValue() { diff --git a/tests/spirv/debug-type-pointer.slang b/tests/spirv/debug-type-pointer.slang index fd3248c592..72ad671d5c 100644 --- a/tests/spirv/debug-type-pointer.slang +++ b/tests/spirv/debug-type-pointer.slang @@ -14,18 +14,18 @@ RWStructuredBuffer outputBuffer; struct LinkedNode { - float value; + float value = {}; //SPV: [[TYPE_float:%[1-9][0-9]*]] = OpExtInst %void %{{[0-9]*}} DebugTypeBasic [[STRING_float]] //SPV: [[TYPE_float_ptr:%[1-9][0-9]*]] = OpExtInst %void %{{[0-9]*}} DebugTypePointer [[TYPE_float]] //SPV: DebugTypeMember [[STRING_pValue]] [[TYPE_float_ptr]] - float *pValue; + float *pValue = {}; // "MEMBER_pNext" hasn't been declared yet, but OpExtInstWithForwardRefs requires at least one of these "forward declared ID". //SPV: [[FORWARD_LinkedNode:%[1-9][0-9]*]] = OpExtInstWithForwardRefsKHR %void %{{[0-9]*}} DebugTypeComposite [[STRING_LinkedNode]] {{.*}} [[MEMBER_pNext:%[1-9][0-9]*]] //SPV: [[TYPE_pLinkedNode:%[1-9][0-9]*]] = OpExtInst %void %{{[0-9]*}} DebugTypePointer [[FORWARD_LinkedNode]] //SPV: [[MEMBER_pNext]] = {{.*}} DebugTypeMember [[STRING_pNext]] [[TYPE_pLinkedNode]] - LinkedNode *pNext; + LinkedNode *pNext = {}; }; //SPV: OpExtInst %void %{{[0-9]*}} DebugTypeComposite [[STRING_LinkedNode]] {{.*}} [[MEMBER_pNext]] diff --git a/tests/spirv/debug-value-dynamic-index.slang b/tests/spirv/debug-value-dynamic-index.slang index 591b1f6dad..4a5e63c15a 100644 --- a/tests/spirv/debug-value-dynamic-index.slang +++ b/tests/spirv/debug-value-dynamic-index.slang @@ -2,9 +2,9 @@ struct TestType { - float memberA; - float3 memberB; - float arrayVal[10]; + float memberA = {}; + float3 memberB = {}; + float arrayVal[10] = {}; RWStructuredBuffer memberC; float getValue() { diff --git a/tests/spirv/large-struct-pack.slang b/tests/spirv/large-struct-pack.slang index df15ac67c0..3e7d2c3f35 100644 --- a/tests/spirv/large-struct-pack.slang +++ b/tests/spirv/large-struct-pack.slang @@ -8,7 +8,7 @@ struct WorkData { - int B[1024]; + int B[1024] = {}; }; //TEST_INPUT:set resultBuffer = out ubuffer(data=[0 0 0 0], stride=4, count=1024) diff --git a/tests/spirv/mutating-method-syn.slang b/tests/spirv/mutating-method-syn.slang index 14adc001f7..fb3bb7c9cb 100644 --- a/tests/spirv/mutating-method-syn.slang +++ b/tests/spirv/mutating-method-syn.slang @@ -10,7 +10,7 @@ interface IFoo struct Val : IFoo { - int x; + int x = {}; int bar(int y) { diff --git a/tools/slang-unit-test/unit-test-decl-tree-reflection.cpp b/tools/slang-unit-test/unit-test-decl-tree-reflection.cpp index 89579c5850..8c5d01186e 100644 --- a/tools/slang-unit-test/unit-test-decl-tree-reflection.cpp +++ b/tools/slang-unit-test/unit-test-decl-tree-reflection.cpp @@ -42,7 +42,7 @@ SLANG_UNIT_TEST(declTreeReflection) // Source for a module that contains an undecorated entrypoint. const char* userSourceBody = R"( [__AttributeUsage(_AttributeTargets.Function)] - struct MyFuncPropertyAttribute {int v;} + struct MyFuncPropertyAttribute {int v = {};} [MyFuncProperty(1024)] [Differentiable] @@ -57,13 +57,13 @@ SLANG_UNIT_TEST(declTreeReflection) struct MyType { - int x; + int x = {}; float f(float x) { return x; } } struct MyGenericType { - T z; + T z = {}; __init(T _z) { z = _z; } @@ -79,7 +79,7 @@ SLANG_UNIT_TEST(declTreeReflection) { struct MyStruct { - int x; + int x = {}; } } @@ -118,10 +118,10 @@ SLANG_UNIT_TEST(declTreeReflection) SLANG_CHECK(moduleDeclReflection->getKind() == slang::DeclReflection::Kind::Module); SLANG_CHECK(moduleDeclReflection->getChildrenCount() == 9); - // First declaration should be a struct with 1 variable + // First declaration should be a struct with 1 variable, 1 constructor (memberwise ctor), 1 funcDecl ($ZeroInit), 1 default ctor auto firstDecl = moduleDeclReflection->getChild(0); SLANG_CHECK(firstDecl->getKind() == slang::DeclReflection::Kind::Struct); - SLANG_CHECK(firstDecl->getChildrenCount() == 1); + SLANG_CHECK(firstDecl->getChildrenCount() == 4); { slang::TypeReflection* type = firstDecl->getType();