From cdee13466080b737ed18c148e36af75898285ed6 Mon Sep 17 00:00:00 2001 From: jsmall-nvidia Date: Fri, 6 Mar 2020 17:15:58 -0500 Subject: Remove generated header files (#1264) * Update slang-binaries to verison with SPIR-V version support. * Support vec and matrix Wave intrinsics on vk. Added wave-vector.slang test Add wave-diverge.slang test Add support for more wave intrinsics to vk. * Test out Wave intrinsic support for matrices. * Remove matrix glsl intrinsics -> not available. Fix some typo. * Remove generated slang generated headers. --- source/slang/core.meta.slang.h | 1864 ------------------------ source/slang/glsl.meta.slang.h | 202 --- source/slang/hlsl.meta.slang.h | 3140 ---------------------------------------- 3 files changed, 5206 deletions(-) delete mode 100644 source/slang/core.meta.slang.h delete mode 100644 source/slang/glsl.meta.slang.h delete mode 100644 source/slang/hlsl.meta.slang.h (limited to 'source') diff --git a/source/slang/core.meta.slang.h b/source/slang/core.meta.slang.h deleted file mode 100644 index f793d0c1b..000000000 --- a/source/slang/core.meta.slang.h +++ /dev/null @@ -1,1864 +0,0 @@ -SLANG_RAW("// Slang `core` library\n") -SLANG_RAW("\n") -SLANG_RAW("// Aliases for base types\n") -SLANG_RAW("typedef half float16_t;\n") -SLANG_RAW("typedef float float32_t;\n") -SLANG_RAW("typedef double float64_t;\n") -SLANG_RAW("\n") -SLANG_RAW("typedef int int32_t;\n") -SLANG_RAW("typedef uint uint32_t;\n") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("// Modifier for variables that must resolve to compile-time constants\n") -SLANG_RAW("// as part of translation.\n") -SLANG_RAW("syntax constexpr : ConstExprModifier;\n") -SLANG_RAW("\n") -SLANG_RAW("// Modifier for variables that should have writes be made\n") -SLANG_RAW("// visible at the global-memory scope\n") -SLANG_RAW("syntax globallycoherent : GloballyCoherentModifier;\n") -SLANG_RAW("\n") -SLANG_RAW("// A type that can be used as an operand for builtins\n") -SLANG_RAW("interface __BuiltinType {}\n") -SLANG_RAW("\n") -SLANG_RAW("// A type that can be used for arithmetic operations\n") -SLANG_RAW("interface __BuiltinArithmeticType : __BuiltinType\n") -SLANG_RAW("{\n") -SLANG_RAW(" /// Initialize from a 32-bit signed integer value.\n") -SLANG_RAW(" __init(int value);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW(" /// A type that can be used for logical/bitwsie operations\n") -SLANG_RAW("interface __BuiltinLogicalType : __BuiltinType {}\n") -SLANG_RAW("\n") -SLANG_RAW("// A type that logically has a sign (positive/negative/zero)\n") -SLANG_RAW("interface __BuiltinSignedArithmeticType : __BuiltinArithmeticType {}\n") -SLANG_RAW("\n") -SLANG_RAW("// A type that can represent integers\n") -SLANG_RAW("interface __BuiltinIntegerType : __BuiltinArithmeticType\n") -SLANG_RAW("{}\n") -SLANG_RAW("\n") -SLANG_RAW("// A type that can represent non-integers\n") -SLANG_RAW("interface __BuiltinRealType : __BuiltinSignedArithmeticType {}\n") -SLANG_RAW("\n") -SLANG_RAW("// A type that uses a floating-point representation\n") -SLANG_RAW("interface __BuiltinFloatingPointType : __BuiltinRealType\n") -SLANG_RAW("{\n") -SLANG_RAW(" /// Initialize from a 32-bit floating-point value.\n") -SLANG_RAW(" __init(float value);\n") -SLANG_RAW("\n") -SLANG_RAW(" /// Get the value of the mathematical constant pi in this type.\n") -SLANG_RAW(" static This getPi();\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// A type resulting from an `enum` declaration.\n") -SLANG_RAW("__magic_type(EnumTypeType)\n") -SLANG_RAW("interface __EnumType\n") -SLANG_RAW("{\n") -SLANG_RAW(" // The type of tags for this `enum`\n") -SLANG_RAW(" //\n") -SLANG_RAW(" // Note: using `__Tag` instead of `Tag` to avoid any\n") -SLANG_RAW(" // conflict if a user had an `enum` case called `Tag`\n") -SLANG_RAW(" associatedtype __Tag : __BuiltinIntegerType;\n") -SLANG_RAW("};\n") -SLANG_RAW("\n") -SLANG_RAW("// Use an extension to declare that every `enum` type\n") -SLANG_RAW("// inherits an initializer based on the tag type.\n") -SLANG_RAW("//\n") -SLANG_RAW("// Note: there is an important and subtle point here.\n") -SLANG_RAW("// If we declared these initializers inside the `interface`\n") -SLANG_RAW("// declaration above, then they would implicitly be\n") -SLANG_RAW("// *requirements* of the `__EnumType` interface, and any\n") -SLANG_RAW("// type that declares conformance to it would need to\n") -SLANG_RAW("// provide implementations. That would put the onus on\n") -SLANG_RAW("// the semantic checker to synthesize such initializers\n") -SLANG_RAW("// when conforming an `enum` type to `__EnumType` (just\n") -SLANG_RAW("// as it currently synthesizes the `__Tag` requirement.\n") -SLANG_RAW("// Putting the declaration in an `extension` makes them\n") -SLANG_RAW("// concrete declerations rather than interface requirements.\n") -SLANG_RAW("// (Admittedly, they are \"concrete\" declarations with\n") -SLANG_RAW("// no bodies, because currently all initializers are\n") -SLANG_RAW("// assumed to be intrinsics).\n") -SLANG_RAW("//\n") -SLANG_RAW("// TODO: It might be more accurate to express this as:\n") -SLANG_RAW("//\n") -SLANG_RAW("// __generic extension T { ... }\n") -SLANG_RAW("//\n") -SLANG_RAW("// That alternative would express an extension of every\n") -SLANG_RAW("// type that conforms to `__EnumType`, rather than an\n") -SLANG_RAW("// extension of `__EnumType` itself. The distinction\n") -SLANG_RAW("// is subtle, and unfortunately not one the Slang type\n") -SLANG_RAW("// checker is equiped to handle right now. For now we\n") -SLANG_RAW("// will stick with the syntax that actually works, even\n") -SLANG_RAW("// if it might be the less technically correct one.\n") -SLANG_RAW("//\n") -SLANG_RAW("//\n") -SLANG_RAW("extension __EnumType\n") -SLANG_RAW("{\n") -SLANG_RAW(" // TODO: this should be a single initializer using\n") -SLANG_RAW(" // the `__Tag` associated type from the `__EnumType`\n") -SLANG_RAW(" // interface, but right now the scoping for looking\n") -SLANG_RAW(" // up that type isn't working right.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" __init(int value);\n") -SLANG_RAW(" __init(uint value);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// A type resulting from an `enum` declaration\n") -SLANG_RAW("// with the `[flags]` attribute.\n") -SLANG_RAW("interface __FlagsEnumType : __EnumType\n") -SLANG_RAW("{\n") -SLANG_RAW("};\n") -SLANG_RAW("\n") -SLANG_RAW("// The \"comma operator\" is effectively just a generic function that returns its second\n") -SLANG_RAW("// argument. The left-to-right evaluation order guaranteed by Slang then ensures that\n") -SLANG_RAW("// `left` is evaluated before `right`.\n") -SLANG_RAW("//\n") -SLANG_RAW("__generic __intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Sequence -) -SLANG_RAW(") U operator,(T left, U right);\n") -SLANG_RAW("\n") -SLANG_RAW("// The ternary `?:` operator does not short-circuit in HLSL, and Slang continues to\n") -SLANG_RAW("// follow that definition, so that this operator is effectively just an ordinary\n") -SLANG_RAW("// function, rather than a special-case piece of syntax.\n") -SLANG_RAW("//\n") -SLANG_RAW("__generic __intrinsic_op(select) T operator?:(bool condition, T ifTrue, T ifFalse);\n") -SLANG_RAW("__generic __intrinsic_op(select) vector operator?:(vector condition, vector ifTrue, vector ifFalse);\n") -SLANG_RAW("\n") - -// We are going to use code generation to produce the -// declarations for all of our base types. -static const int kBaseTypeCount = sizeof(kBaseTypes) / sizeof(kBaseTypes[0]); -for (int tt = 0; tt < kBaseTypeCount; ++tt) -{ -SLANG_RAW("#line 131 \"core.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("__builtin_type(") -SLANG_SPLICE(int(kBaseTypes[tt].tag) -) -SLANG_RAW(")\n") -SLANG_RAW("struct ") -SLANG_SPLICE(kBaseTypes[tt].name -) -SLANG_RAW("\n") -SLANG_RAW(" : __BuiltinType\n") -SLANG_RAW("\n") - - switch (kBaseTypes[tt].tag) - { - case BaseType::Half: - case BaseType::Float: - case BaseType::Double: -SLANG_RAW("#line 143 \"core.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW(" , __BuiltinFloatingPointType\n") -SLANG_RAW(" , __BuiltinRealType\n") -SLANG_RAW(" , __BuiltinSignedArithmeticType\n") -SLANG_RAW(" , __BuiltinArithmeticType\n") - - break; - case BaseType::Int8: - case BaseType::Int16: - case BaseType::Int: - case BaseType::Int64: -SLANG_RAW("#line 154 \"core.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW(" , __BuiltinSignedArithmeticType\n") - - ; // fall through to: - case BaseType::UInt8: - case BaseType::UInt16: - case BaseType::UInt: - case BaseType::UInt64: -SLANG_RAW("#line 162 \"core.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW(" , __BuiltinArithmeticType\n") -SLANG_RAW(" , __BuiltinIntegerType\n") - - ; // fall through to: - case BaseType::Bool: -SLANG_RAW("#line 168 \"core.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW(" , __BuiltinLogicalType\n") - - break; - - default: - break; - } -SLANG_RAW("#line 176 \"core.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW("{\n") -SLANG_RAW("\n") - - // Declare initializers to convert from various other types - for (int ss = 0; ss < kBaseTypeCount; ++ss) - { - // Don't allow conversion from `void` - if (kBaseTypes[ss].tag == BaseType::Void) - continue; - - // We need to emit a modifier so that the semantic-checking - // layer will know it can use these operations for implicit - // conversion. - ConversionCost conversionCost = getBaseTypeConversionCost( - kBaseTypes[tt], - kBaseTypes[ss]); -SLANG_RAW("#line 193 \"core.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW(" __implicit_conversion(") -SLANG_SPLICE(conversionCost -) -SLANG_RAW(")\n") -SLANG_RAW(" __init(") -SLANG_SPLICE(kBaseTypes[ss].name -) -SLANG_RAW(" value);\n") -SLANG_RAW("\n") - - } - - // If this is a basic integer type, then define explicit - // initializers that take a value of an `enum` type. - // - // TODO: This should actually be restricted, so that this - // only applies `where T.__Tag == Self`, but we don't have - // the needed features in our type system to implement - // that constraint right now. - // - switch (kBaseTypes[tt].tag) - { - // TODO: should this cover the full gamut of integer types? - case BaseType::Int: - case BaseType::UInt: -SLANG_RAW("#line 214 \"core.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW(" __generic\n") -SLANG_RAW(" __init(T value);\n") - - break; - - default: - break; - } - - // If this is a floating-point type, then we need to - // define the basic `getPi()` function that is used - // to implement generic versions of `degrees()` and - // `radians()`. - // - switch (kBaseTypes[tt].tag) - { - default: - break; - case BaseType::Half: - case BaseType::Float: - case BaseType::Double: -SLANG_RAW("#line 236 \"core.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW(" static ") -SLANG_SPLICE(kBaseTypes[tt].name -) -SLANG_RAW(" getPi() { return ") -SLANG_SPLICE(kBaseTypes[tt].name -) -SLANG_RAW("(3.14159265358979323846264338328); }\n") - - break; - } -SLANG_RAW("#line 241 \"core.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") - -} - -// Declare built-in pointer type -// (eventually we can have the traditional syntax sugar for this) -SLANG_RAW("#line 250 \"core.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__magic_type(PtrType)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_PtrType -) -SLANG_RAW(")\n") -SLANG_RAW("struct Ptr\n") -SLANG_RAW("{};\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__magic_type(OutType)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_OutType -) -SLANG_RAW(")\n") -SLANG_RAW("struct Out\n") -SLANG_RAW("{};\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__magic_type(InOutType)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_InOutType -) -SLANG_RAW(")\n") -SLANG_RAW("struct InOut\n") -SLANG_RAW("{};\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__magic_type(RefType)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_RefType -) -SLANG_RAW(")\n") -SLANG_RAW("struct Ref\n") -SLANG_RAW("{};\n") -SLANG_RAW("\n") -SLANG_RAW("__magic_type(StringType)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_StringType -) -SLANG_RAW(")\n") -SLANG_RAW("struct String\n") -SLANG_RAW("{};\n") -SLANG_RAW("\n") -SLANG_RAW(" /// An `N` component vector with elements of type `T`.\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__magic_type(Vector)\n") -SLANG_RAW("struct vector\n") -SLANG_RAW("{\n") -SLANG_RAW(" /// The element type of the vector\n") -SLANG_RAW(" typedef T Element;\n") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW(" /// Initialize a vector where all elements have the same scalar `value`.\n") -SLANG_RAW(" __implicit_conversion(") -SLANG_SPLICE(kConversionCost_ScalarToVector -) -SLANG_RAW(")\n") -SLANG_RAW(" __intrinsic_op(") -SLANG_SPLICE(kIROp_constructVectorFromScalar -) -SLANG_RAW(")\n") -SLANG_RAW(" __init(T value);\n") -SLANG_RAW("\n") -SLANG_RAW(" /// Initialize a vector from a value of the same type\n") -SLANG_RAW(" // TODO: we should revise semantic checking so this kind of \"identity\" conversion is not required\n") -SLANG_RAW(" __init(vector value);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW(" /// A matrix with `R` rows and `C` columns, with elements of type `T`.\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__magic_type(Matrix)\n") -SLANG_RAW("struct matrix\n") -SLANG_RAW("{\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") - -static const struct { - char const* name; - char const* glslPrefix; -} kTypes[] = -{ - {"half", "f16"}, - {"float", ""}, - {"double", "d"}, - - {"float16_t", "f16"}, - {"float32_t", "f32"}, - {"float64_t", "f64"}, - - {"int8_t", "i8"}, - {"int16_t", "i16"}, - {"int32_t", "i32"}, - {"int", "i"}, - {"int64_t", "i64"}, - - {"uint8_t", "u8"}, - {"uint16_t", "u16"}, - {"uint32_t", "u32"}, - {"uint", "u"}, - {"uint64_t", "u64"}, - - {"bool", "b"}, -}; -static const int kTypeCount = sizeof(kTypes) / sizeof(kTypes[0]); - -for (int tt = 0; tt < kTypeCount; ++tt) -{ - // Declare HLSL vector types - for (int ii = 1; ii <= 4; ++ii) - { - sb << "typedef vector<" << kTypes[tt].name << "," << ii << "> " << kTypes[tt].name << ii << ";\n"; - } - - // Declare HLSL matrix types - for (int rr = 2; rr <= 4; ++rr) - for (int cc = 2; cc <= 4; ++cc) - { - sb << "typedef matrix<" << kTypes[tt].name << "," << rr << "," << cc << "> " << kTypes[tt].name << rr << "x" << cc << ";\n"; - } -} - -// Declare additional built-in generic types -SLANG_RAW("#line 353 \"core.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_ConstantBufferType -) -SLANG_RAW(")\n") -SLANG_RAW("__magic_type(ConstantBuffer)\n") -SLANG_RAW("struct ConstantBuffer {}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_TextureBufferType -) -SLANG_RAW(")\n") -SLANG_RAW("__magic_type(TextureBuffer)\n") -SLANG_RAW("struct TextureBuffer {}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_ParameterBlockType -) -SLANG_RAW(")\n") -SLANG_RAW("__magic_type(ParameterBlockType)\n") -SLANG_RAW("struct ParameterBlock {}\n") -SLANG_RAW("\n") - - -static const char* kComponentNames[]{ "x", "y", "z", "w" }; -static const char* kVectorNames[]{ "", "x", "xy", "xyz", "xyzw" }; - -// Need to add constructors to the types above -for (int N = 2; N <= 4; ++N) -{ - sb << "__generic __extension vector\n{\n"; - - // initialize from N scalars - sb << "__init("; - for (int ii = 0; ii < N; ++ii) - { - if (ii != 0) sb << ", "; - sb << "T " << kComponentNames[ii]; - } - sb << ");\n"; - - // Initialize from an M-vector and then scalars - for (int M = 2; M < N; ++M) - { - sb << "__init(vector " << kVectorNames[M]; - for (int ii = M; ii < N; ++ii) - { - sb << ", T " << kComponentNames[ii]; - } - sb << ");\n"; - } - - // Initialize from two vectors, of size M and N-M - for(int M = 2; M <= (N-2); ++M) - { - int K = N - M; - SLANG_ASSERT(K >= 2); - - sb << "__init(vector " << kVectorNames[M]; - sb << ", vector "; - for (int ii = 0; ii < K; ++ii) - { - // The component names for the second parameter - // must start at `M`, so that we get signatures like: - // - // __init(float2 xy, float2 zw) - // - sb << kComponentNames[M + ii]; - } - sb << ");\n"; - } - - sb << "}\n"; -} - -// The above extension was generic in the *type* of the vector, -// but explicit in the *size*. We will now declare an extension -// for each builtin type that is generic in the size. -// -for (int tt = 0; tt < kBaseTypeCount; ++tt) -{ - if(kBaseTypes[tt].tag == BaseType::Void) continue; - - sb << "__generic __extension vector<" - << kBaseTypes[tt].name << ",N>\n{\n"; - - for (int ff = 0; ff < kBaseTypeCount; ++ff) - { - if(kBaseTypes[ff].tag == BaseType::Void) continue; - - - if( tt != ff ) - { - auto cost = getBaseTypeConversionCost( - kBaseTypes[tt], - kBaseTypes[ff]); - - // Implicit conversion from a vector of the same - // size, but different element type. - sb << " __implicit_conversion(" << cost << ")\n"; - sb << " __init(vector<" << kBaseTypes[ff].name << ",N> value);\n"; - - // Constructor to make a vector from a scalar of another type. - cost += kConversionCost_ScalarToVector; - sb << " __implicit_conversion(" << cost << ")\n"; - sb << " __init(" << kBaseTypes[ff].name << " value);\n"; - } - } - - sb << "}\n"; -} - -for( int R = 2; R <= 4; ++R ) -for( int C = 2; C <= 4; ++C ) -{ - sb << "__generic __extension matrix\n{\n"; - - // initialize from R*C scalars - sb << "__init("; - for( int ii = 0; ii < R; ++ii ) - for( int jj = 0; jj < C; ++jj ) - { - if ((ii+jj) != 0) sb << ", "; - sb << "T m" << ii << jj; - } - sb << ");\n"; - - // Initialize from R C-vectors - sb << "__init("; - for (int ii = 0; ii < R; ++ii) - { - if(ii != 0) sb << ", "; - sb << "vector row" << ii; - } - sb << ");\n"; - - - // initialize from another matrix of the same size - // - // TODO(tfoley): See comment about how this overlaps - // with implicit conversion, in the `vector` case above - sb << "__generic __init(matrix);\n"; - - // initialize from a matrix of larger size - for(int rr = R; rr <= 4; ++rr) - for( int cc = C; cc <= 4; ++cc ) - { - if(rr == R && cc == C) continue; - sb << "__init(matrix value);\n"; - } - - sb << "}\n"; -} -SLANG_RAW("#line 501 \"core.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW(" /// Sampling state for filtered texture fetches.\n") -SLANG_RAW("__magic_type(SamplerState, ") -SLANG_SPLICE(int(SamplerStateFlavor::SamplerState) -) -SLANG_RAW(")\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_SamplerStateType -) -SLANG_RAW(")\n") -SLANG_RAW("struct SamplerState\n") -SLANG_RAW("{\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW(" /// Sampling state for filtered texture fetches that include a comparison operation before filtering.\n") -SLANG_RAW("__magic_type(SamplerState, ") -SLANG_SPLICE(int(SamplerStateFlavor::SamplerComparisonState) -) -SLANG_RAW(")\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_SamplerComparisonStateType -) -SLANG_RAW(")\n") -SLANG_RAW("struct SamplerComparisonState\n") -SLANG_RAW("{\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") - - -static const struct { - char const* name; - TextureFlavor::Shape baseShape; - int coordCount; -} kBaseTextureTypes[] = { - { "Texture1D", TextureFlavor::Shape::Shape1D, 1 }, - { "Texture2D", TextureFlavor::Shape::Shape2D, 2 }, - { "Texture3D", TextureFlavor::Shape::Shape3D, 3 }, - { "TextureCube", TextureFlavor::Shape::ShapeCube, 3 }, -}; -static const int kBaseTextureTypeCount = sizeof(kBaseTextureTypes) / sizeof(kBaseTextureTypes[0]); - - -static const struct { - char const* name; - SlangResourceAccess access; -} kBaseTextureAccessLevels[] = { - { "", SLANG_RESOURCE_ACCESS_READ }, - { "RW", SLANG_RESOURCE_ACCESS_READ_WRITE }, - { "RasterizerOrdered", SLANG_RESOURCE_ACCESS_RASTER_ORDERED }, -}; -static const int kBaseTextureAccessLevelCount = sizeof(kBaseTextureAccessLevels) / sizeof(kBaseTextureAccessLevels[0]); - -// Declare the GLSL types here for compatibility... -// -// TODO: The stdlib should include a module that declares the GLSL types, to keep -// them separate... -for (int tt = 0; tt < kBaseTextureTypeCount; ++tt) -{ - char const* name = kBaseTextureTypes[tt].name; - TextureFlavor::Shape baseShape = kBaseTextureTypes[tt].baseShape; - - for (int isArray = 0; isArray < 2; ++isArray) - { - // Arrays of 3D textures aren't allowed - if (isArray && baseShape == TextureFlavor::Shape::Shape3D) continue; - - for (int isMultisample = 0; isMultisample < 2; ++isMultisample) - for (int accessLevel = 0; accessLevel < kBaseTextureAccessLevelCount; ++accessLevel) - { - auto access = kBaseTextureAccessLevels[accessLevel].access; - - // TODO: any constraints to enforce on what gets to be multisampled? - - unsigned flavor = baseShape; - if (isArray) flavor |= TextureFlavor::ArrayFlag; - if (isMultisample) flavor |= TextureFlavor::MultisampleFlag; -// if (isShadow) flavor |= TextureFlavor::ShadowFlag; - - flavor |= (access << 8); - - // emit a generic signature - // TODO: allow for multisample count to come in as well... - sb << "__generic "; - - sb << "__magic_type(TextureSampler," << int(flavor) << ")\n"; - sb << "__intrinsic_type(" << (kIROp_TextureSamplerType + (int(flavor) << kIROpMeta_OtherShift)) << ")\n"; - sb << "struct Sampler"; - sb << kBaseTextureAccessLevels[accessLevel].name; - sb << name; - if (isMultisample) sb << "MS"; - if (isArray) sb << "Array"; -// if (isShadow) sb << "Shadow"; - sb << "\n{\n"; - sb << "__specialized_for_target(glsl)\n"; - sb << "__init("; - sb << kBaseTextureAccessLevels[accessLevel].name; - sb << name; - if (isMultisample) sb << "MS"; - if (isArray) sb << "Array"; - sb << " t, "; - sb << "SamplerState s);\n"; - sb << "};\n"; - - sb << "__specialized_for_target(glsl)\n"; - sb << "T texture(Sampler"; - sb << kBaseTextureAccessLevels[accessLevel].name; - sb << name; - if (isMultisample) sb << "MS"; - if (isArray) sb << "Array"; - sb << " t, float" << kBaseTextureTypes[tt].coordCount + isArray << " location);\n"; - } - } -} - -for (int tt = 0; tt < kBaseTextureTypeCount; ++tt) -{ - char const* name = kBaseTextureTypes[tt].name; - TextureFlavor::Shape baseShape = kBaseTextureTypes[tt].baseShape; - - for (int isArray = 0; isArray < 2; ++isArray) - { - // Arrays of 3D textures aren't allowed - if (isArray && baseShape == TextureFlavor::Shape::Shape3D) continue; - - for (int isMultisample = 0; isMultisample < 2; ++isMultisample) - for (int accessLevel = 0; accessLevel < kBaseTextureAccessLevelCount; ++accessLevel) - { - auto access = kBaseTextureAccessLevels[accessLevel].access; - - // TODO: any constraints to enforce on what gets to be multisampled? - - unsigned flavor = baseShape; - if (isArray) flavor |= TextureFlavor::ArrayFlag; - if (isMultisample) flavor |= TextureFlavor::MultisampleFlag; -// if (isShadow) flavor |= TextureFlavor::ShadowFlag; - - flavor |= (access << 8); - - // emit a generic signature - // TODO: allow for multisample count to come in as well... - sb << "__generic "; - - sb << "__magic_type(Texture," << int(flavor) << ")\n"; - sb << "__intrinsic_type(" << (kIROp_TextureType + (int(flavor) << kIROpMeta_OtherShift)) << ")\n"; - sb << "struct "; - sb << kBaseTextureAccessLevels[accessLevel].name; - sb << name; - if (isMultisample) sb << "MS"; - if (isArray) sb << "Array"; -// if (isShadow) sb << "Shadow"; - sb << "\n{"; - - if( !isMultisample ) - { - sb << "float CalculateLevelOfDetail(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount << " location);\n"; - - sb << "float CalculateLevelOfDetailUnclamped(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount << " location);\n"; - } - - // `GetDimensions` - - for(int isFloat = 0; isFloat < 2; ++isFloat) - for(int includeMipInfo = 0; includeMipInfo < 2; ++includeMipInfo) - { - { - sb << "__glsl_version(450)\n"; - sb << "__glsl_extension(GL_EXT_samplerless_texture_functions)"; - sb << "__target_intrinsic(glsl, \"("; - - int aa = 1; - String lodStr = ", 0"; - if (includeMipInfo) - { - int mipLevelArg = aa++; - lodStr = ", int($"; - lodStr.append(mipLevelArg); - lodStr.append(")"); - } - - String opStr = " = textureSize($0" + lodStr; - switch( access ) - { - case SLANG_RESOURCE_ACCESS_READ_WRITE: - case SLANG_RESOURCE_ACCESS_RASTER_ORDERED: - opStr = " = imageSize($0"; - break; - - default: - break; - } - - - int cc = 0; - switch(baseShape) - { - case TextureFlavor::Shape::Shape1D: - sb << "($" << aa++ << opStr << "))"; - cc = 1; - break; - - case TextureFlavor::Shape::Shape2D: - case TextureFlavor::Shape::ShapeCube: - sb << "($" << aa++ << opStr << ").x)"; - sb << ", ($" << aa++ << opStr << ").y)"; - cc = 2; - break; - - case TextureFlavor::Shape::Shape3D: - sb << "($" << aa++ << opStr << ").x)"; - sb << ", ($" << aa++ << opStr << ").y)"; - sb << ", ($" << aa++ << opStr << ").z)"; - cc = 3; - break; - - default: - SLANG_UNEXPECTED("unhandled resource shape"); - break; - } - - if(isArray) - { - sb << ", ($" << aa++ << opStr << ")." << kComponentNames[cc] << ")"; - } - - if(isMultisample) - { - sb << ", ($" << aa++ << " = textureSamples($0))"; - } - - if (includeMipInfo) - { - sb << ", ($" << aa++ << " = textureQueryLevels($0))"; - } - - - sb << ")\")\n"; - } - - char const* t = isFloat ? "out float " : "out uint "; - - sb << "void GetDimensions("; - if(includeMipInfo) - sb << "uint mipLevel, "; - - switch(baseShape) - { - case TextureFlavor::Shape::Shape1D: - sb << t << "width"; - break; - - case TextureFlavor::Shape::Shape2D: - case TextureFlavor::Shape::ShapeCube: - sb << t << "width,"; - sb << t << "height"; - break; - - case TextureFlavor::Shape::Shape3D: - sb << t << "width,"; - sb << t << "height,"; - sb << t << "depth"; - break; - - default: - assert(!"unexpected"); - break; - } - - if(isArray) - { - sb << ", " << t << "elements"; - } - - if(isMultisample) - { - sb << ", " << t << "sampleCount"; - } - - if(includeMipInfo) - sb << ", " << t << "numberOfLevels"; - - sb << ");\n"; - } - - // `GetSamplePosition()` - if( isMultisample ) - { - sb << "float2 GetSamplePosition(int s);\n"; - } - - // `Load()` - - if( kBaseTextureTypes[tt].coordCount + isArray < 4 ) - { - // The `Load()` operation on an ordinary `Texture2D` takes - // an `int3` for the location, where `.xy` holds the texel - // coordinates, and `.z` holds the mip level to use. - // - // The third coordinate for mip level is absent in - // `Texure2DMS.Load()` and `RWTexture2D.Load`. This pattern - // is repreated for all the other texture shapes. - // - bool needsMipLevel = !isMultisample && (access == SLANG_RESOURCE_ACCESS_READ); - - int loadCoordCount = kBaseTextureTypes[tt].coordCount + isArray + (needsMipLevel?1:0); - - char const* glslFuncName = (access == SLANG_RESOURCE_ACCESS_READ) ? "texelFetch" : "imageLoad"; - - // When translating to GLSL, we need to break apart the `location` argument. - // - // TODO: this should realy be handled by having this member actually get lowered! - static const char* kGLSLLoadCoordsSwizzle[] = { "", "", "x", "xy", "xyz", "xyzw" }; - static const char* kGLSLLoadLODSwizzle[] = { "", "", "y", "z", "w", "error" }; - - // TODO: The GLSL translations here only handle the read-only texture - // cases (stuff that lowers to `texture*` in GLSL) and not the stuff - // that lowers to `image*`. - // - // At some point it may make sense to separate the read-only and - // `RW`/`RasterizerOrdered` cases here rather than try to share code. - - if (isMultisample) - { - sb << "__glsl_extension(GL_EXT_samplerless_texture_functions)"; - sb << "__target_intrinsic(glsl, \"$c" << glslFuncName << "($0, $1, $2)$z\")\n"; - } - else - { - sb << "__glsl_extension(GL_EXT_samplerless_texture_functions)"; - sb << "__target_intrinsic(glsl, \"$c" << glslFuncName << "($0, "; - if( needsMipLevel ) - { - sb << "($1)." << kGLSLLoadCoordsSwizzle[loadCoordCount] << ", ($1)." << kGLSLLoadLODSwizzle[loadCoordCount]; - } - else - { - sb << "$1"; - } - sb << ")$z\")\n"; - - } - - // CUDA - if (isMultisample) - { - } - else - { - if (access == SLANG_RESOURCE_ACCESS_READ_WRITE) - { - const int coordCount = kBaseTextureTypes[tt].coordCount; - const int vecCount = coordCount + int(isArray); - - if( baseShape != TextureFlavor::Shape::ShapeCube ) - { - sb << "__target_intrinsic(cuda, \"surf" << coordCount << "D"; - if (isArray) - { - sb << "Layered"; - } - sb << "read"; - sb << "<$T0>($0"; - for (int i = 0; i < coordCount; ++i) - { - sb << ", ($1)"; - if (vecCount > 1) - { - sb << '.' << char(i + 'x'); - } - } - if (isArray) - { - sb << ", int(($1)." << char(coordCount + 'x') << ")"; - } - sb << ", SLANG_CUDA_BOUNDARY_MODE)\")\n"; - } - else - { - sb << "__target_intrinsic(cuda, \"surfCubemap"; - if (isArray) - { - sb << "Layered"; - } - sb << "read"; - sb << "<$T0>($0, ($1).x, ($1).y, ($1).z"; - if (isArray) - { - sb << ", int(($1).w)"; - } - sb << ", SLANG_CUDA_BOUNDARY_MODE)\")\n"; - } - } - else if (access == SLANG_RESOURCE_ACCESS_READ) - { - // We can allow this on Texture1D - if( baseShape == TextureFlavor::Shape::Shape1D && isArray == false) - { - sb << "__target_intrinsic(cuda, \"tex1Dfetch<$T0>($0, ($1).x)\")\n"; - } - } - } - - sb << "T Load("; - sb << "int" << loadCoordCount << " location"; - if(isMultisample) - { - sb << ", int sampleIndex"; - } - sb << ");\n"; - - // GLSL - if (isMultisample) - { - sb << "__glsl_extension(GL_EXT_samplerless_texture_functions)"; - sb << "__target_intrinsic(glsl, \"$c" << glslFuncName << "($0, $0, $1, $2)$z\")\n"; - } - else - { - sb << "__glsl_extension(GL_EXT_samplerless_texture_functions)"; - sb << "__target_intrinsic(glsl, \"$c" << glslFuncName << "($0, "; - if( needsMipLevel ) - { - sb << "($1)." << kGLSLLoadCoordsSwizzle[loadCoordCount] << ", ($1)." << kGLSLLoadLODSwizzle[loadCoordCount]; - } - else - { - sb << "$1, 0"; - } - sb << ", $2)$z\")\n"; - } - - - - sb << "T Load("; - sb << "int" << loadCoordCount << " location"; - if(isMultisample) - { - sb << ", int sampleIndex"; - } - sb << ", constexpr int" << kBaseTextureTypes[tt].coordCount << " offset"; - sb << ");\n"; - - - sb << "T Load("; - sb << "int" << loadCoordCount << " location"; - if(isMultisample) - { - sb << ", int sampleIndex"; - } - sb << ", constexpr int" << kBaseTextureTypes[tt].coordCount << " offset"; - sb << ", out uint status"; - sb << ");\n"; - } - - if(baseShape != TextureFlavor::Shape::ShapeCube) - { - int N = kBaseTextureTypes[tt].coordCount + isArray; - - char const* uintNs[] = { "", "uint", "uint2", "uint3", "uint4" }; - char const* ivecNs[] = { "", "int", "ivec2", "ivec3", "ivec4" }; - - auto uintN = uintNs[N]; - auto ivecN = ivecNs[N]; - - // subscript operator - sb << "__subscript(" << uintN << " location) -> T {\n"; - - // !!!!!!!!!!!!!!!!!!!! get !!!!!!!!!!!!!!!!!!!!!!! - - // GLSL/SPIR-V distinguished sampled vs. non-sampled images - { - switch( access ) - { - case SLANG_RESOURCE_ACCESS_NONE: - case SLANG_RESOURCE_ACCESS_READ: - sb << "__glsl_extension(GL_EXT_samplerless_texture_functions)"; - sb << "__target_intrinsic(glsl, \"$ctexelFetch($0, " << ivecN << "($1)"; - if( !isMultisample ) - { - sb << ", 0"; - } - else - { - // TODO: how to handle passing through sample index? - sb << ", 0"; - } - break; - - default: - sb << "__target_intrinsic(glsl, \"$cimageLoad($0, " << ivecN << "($1)"; - if( isMultisample ) - { - // TODO: how to handle passing through sample index? - sb << ", 0"; - } - break; - } - sb << ")$z\")\n"; - } - - // CUDA - { - if (access == SLANG_RESOURCE_ACCESS_READ_WRITE) - { - const int coordCount = kBaseTextureTypes[tt].coordCount; - const int vecCount = coordCount + int(isArray); - - sb << "__target_intrinsic(cuda, \"surf"; - if( baseShape != TextureFlavor::Shape::ShapeCube ) - { - sb << coordCount << "D"; - } - else - { - sb << "Cubemap"; - } - - sb << (isArray ? "Layered" : ""); - sb << "read<$T0>($0"; - - for (int i = 0; i < vecCount; ++i) - { - sb << ", ($1)"; - if (vecCount > 1) - { - sb << '.' << char(i + 'x'); - } - } - - sb << ", SLANG_CUDA_BOUNDARY_MODE)\")\n"; - } - else if (access == SLANG_RESOURCE_ACCESS_READ) - { - // We can allow this on Texture1D - if( baseShape == TextureFlavor::Shape::Shape1D && isArray == false) - { - sb << "__target_intrinsic(cuda, \"tex1Dfetch<$T0>($0, $1)\")\n"; - } - } - } - - // Output that has get - sb << " get;\n"; - - // !!!!!!!!!!!!!!!!!!!! set !!!!!!!!!!!!!!!!!!!!!!! - - if (!(access == SLANG_RESOURCE_ACCESS_NONE || access == SLANG_RESOURCE_ACCESS_READ)) - { - // GLSL - sb << "__target_intrinsic(glsl, \"imageStore($0, " << ivecN << "($1), $V2)\")\n"; - - // CUDA - { - const int coordCount = kBaseTextureTypes[tt].coordCount; - const int vecCount = coordCount + int(isArray); - - sb << "__target_intrinsic(cuda, \"surf"; - if( baseShape != TextureFlavor::Shape::ShapeCube ) - { - sb << coordCount << "D"; - } - else - { - sb << "Cubemap"; - } - - sb << (isArray ? "Layered" : ""); - sb << "write<$T0>($2, $0"; - for (int i = 0; i < vecCount; ++i) - { - sb << ", ($1)"; - if (vecCount > 1) - { - sb << '.' << char(i + 'x'); - } - } - - sb << ", SLANG_CUDA_BOUNDARY_MODE)\")\n"; - } - - // Set - sb << " set;\n"; - } - - // !!!!!!!!!!!!!!!!!! ref !!!!!!!!!!!!!!!!!!!!!!!!! - - // Depending on the access level of the texture type, - // we either have just a getter (the default), or both - // a getter and setter. - switch( access ) - { - case SLANG_RESOURCE_ACCESS_NONE: - case SLANG_RESOURCE_ACCESS_READ: - break; - default: - sb << "__intrinsic_op(" << int(kIROp_ImageSubscript) << ") ref;\n"; - break; - } - - sb << "}\n"; - } - - if( !isMultisample ) - { - // `Sample()` - - sb << "__target_intrinsic(glsl, \"$ctexture($p, $2)$z\")\n"; - - // CUDA - { - const int coordCount = kBaseTextureTypes[tt].coordCount; - const int vecCount = coordCount + int(isArray); - - if( baseShape != TextureFlavor::Shape::ShapeCube ) - { - sb << "__target_intrinsic(cuda, \"tex" << coordCount << "D"; - if (isArray) - { - sb << "Layered"; - } - sb << "<$T0>($0"; - for (int i = 0; i < coordCount; ++i) - { - sb << ", ($2)"; - if (vecCount > 1) - { - sb << '.' << char(i + 'x'); - } - } - if (isArray) - { - sb << ", int(($2)." << char(coordCount + 'x') << ")"; - } - sb << ")\")\n"; - } - else - { - sb << "__target_intrinsic(cuda, \"texCubemap"; - if (isArray) - { - sb << "Layered"; - } - sb << "<$T0>($0, ($2).x, ($2).y, ($2).z"; - if (isArray) - { - sb << ", int(($2).w)"; - } - sb << ")\")\n"; - } - } - - sb << "T Sample(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount + isArray << " location);\n"; - - if( baseShape != TextureFlavor::Shape::ShapeCube ) - { - sb << "__target_intrinsic(glsl, \"$ctextureOffset($p, $2, $3)$z\")\n"; - sb << "T Sample(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount + isArray << " location, "; - sb << "constexpr int" << kBaseTextureTypes[tt].coordCount << " offset);\n"; - } - - sb << "T Sample(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount + isArray << " location, "; - if( baseShape != TextureFlavor::Shape::ShapeCube ) - { - sb << "constexpr int" << kBaseTextureTypes[tt].coordCount << " offset, "; - } - sb << "float clamp);\n"; - - sb << "T Sample(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount + isArray << " location, "; - if( baseShape != TextureFlavor::Shape::ShapeCube ) - { - sb << "constexpr int" << kBaseTextureTypes[tt].coordCount << " offset, "; - } - sb << "float clamp, out uint status);\n"; - - // `SampleBias()` - sb << "__target_intrinsic(glsl, \"$ctexture($p, $2, $3)$z\")\n"; - sb << "T SampleBias(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount + isArray << " location, float bias);\n"; - - if( baseShape != TextureFlavor::Shape::ShapeCube ) - { - sb << "__target_intrinsic(glsl, \"$ctextureOffset($p, $2, $3, $4)$z\")\n"; - sb << "T SampleBias(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount + isArray << " location, float bias, "; - sb << "constexpr int" << kBaseTextureTypes[tt].coordCount << " offset);\n"; - } - int baseCoordCount = kBaseTextureTypes[tt].coordCount; - int arrCoordCount = baseCoordCount + isArray; - if (arrCoordCount <= 3) - { - // `SampleCmp()` and `SampleCmpLevelZero` - sb << "__target_intrinsic(glsl, \"texture($p, vec" << arrCoordCount + 1 << "($2, $3))\")"; - sb << "float SampleCmp(SamplerComparisonState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount + isArray << " location, "; - sb << "float compareValue"; - sb << ");\n"; - sb << "__target_intrinsic(glsl, \"texture($p, vec" << arrCoordCount + 1 << "($2, $3))\")"; - sb << "float SampleCmpLevelZero(SamplerComparisonState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount + isArray << " location, "; - sb << "float compareValue"; - sb << ");\n"; - } - if (arrCoordCount < 3) - { - int extCoordCount = arrCoordCount + 1; - - if (extCoordCount < 3) - extCoordCount = 3; - - sb << "__target_intrinsic(glsl, \"$ctextureLod($p, "; - - sb << "vec" << extCoordCount << "($2,"; - for (int ii = arrCoordCount; ii < extCoordCount - 1; ++ii) - { - sb << " 0.0,"; - } - sb << "$3)"; - - sb << ", 0.0)$z\")\n"; - } - else if(arrCoordCount <= 3) - { - int extCoordCount = arrCoordCount + 1; - - if (extCoordCount < 3) - extCoordCount = 3; - - sb << "__target_intrinsic(glsl, \"$ctextureGrad($p, "; - - sb << "vec" << extCoordCount << "($2,"; - for (int ii = arrCoordCount; ii < extCoordCount - 1; ++ii) - { - sb << " 0.0,"; - } - sb << "$3)"; - - // Construct gradients - sb << ", vec" << baseCoordCount << "(0.0)"; - sb << ", vec" << baseCoordCount << "(0.0)"; - sb << ")$z\")\n"; - } - - - if( baseShape != TextureFlavor::Shape::ShapeCube ) - { - // Note(tfoley): MSDN seems confused, and claims that the `offset` - // parameter for `SampleCmp` is available for everything but 3D - // textures, while `Sample` and `SampleBias` are consistent in - // saying they only exclude `offset` for cube maps (which makes - // sense). I'm going to assume the documentation for `SampleCmp` - // is just wrong. - - sb << "float SampleCmp(SamplerComparisonState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount + isArray << " location, "; - sb << "float compareValue, "; - sb << "constexpr int" << kBaseTextureTypes[tt].coordCount << " offset);\n"; - - sb << "float SampleCmpLevelZero(SamplerComparisonState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount + isArray << " location, "; - sb << "float compareValue, "; - sb << "constexpr int" << kBaseTextureTypes[tt].coordCount << " offset);\n"; - } - - // TODO(JS): Not clear how to map this to CUDA, because in HLSL, the gradient is a vector based on - // the dimension. On CUDA there is texNDGrad, but it always just takes ddx, ddy. - // I could just assume 0 for elements not supplied, and ignore z. For now will just leave - sb << "__target_intrinsic(glsl, \"$ctextureGrad($p, $2, $3, $4)$z\")\n"; - sb << "T SampleGrad(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount + isArray << " location, "; - sb << "float" << kBaseTextureTypes[tt].coordCount << " gradX, "; - sb << "float" << kBaseTextureTypes[tt].coordCount << " gradY"; - sb << ");\n"; - - if( baseShape != TextureFlavor::Shape::ShapeCube ) - { - sb << "__target_intrinsic(glsl, \"$ctextureGradOffset($p, $2, $3, $4, $5)$z\")\n"; - sb << "T SampleGrad(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount + isArray << " location, "; - sb << "float" << kBaseTextureTypes[tt].coordCount << " gradX, "; - sb << "float" << kBaseTextureTypes[tt].coordCount << " gradY, "; - sb << "constexpr int" << kBaseTextureTypes[tt].coordCount << " offset);\n"; - } - - // `SampleLevel` - - sb << "__target_intrinsic(glsl, \"$ctextureLod($p, $2, $3)$z\")\n"; - - // CUDA - { - const int coordCount = kBaseTextureTypes[tt].coordCount; - const int vecCount = coordCount + int(isArray); - - if( baseShape != TextureFlavor::Shape::ShapeCube ) - { - sb << "__target_intrinsic(cuda, \"tex" << coordCount << "D"; - if (isArray) - { - sb << "Layered"; - } - sb << "Lod<$T0>($0"; - for (int i = 0; i < coordCount; ++i) - { - sb << ", ($2)"; - if (vecCount > 1) - { - sb << '.' << char(i + 'x'); - } - } - if (isArray) - { - sb << ", int(($2)." << char(coordCount + 'x') << ")"; - } - sb << ", $3)\")\n"; - } - else - { - sb << "__target_intrinsic(cuda, \"texCubemap"; - if (isArray) - { - sb << "Layered"; - } - sb << "Lod<$T0>($0, ($2).x, ($2).y, ($2).z"; - if (isArray) - { - sb << ", int(($2).w)"; - } - sb << ", $3)\")\n"; - } - } - - sb << "T SampleLevel(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount + isArray << " location, "; - sb << "float level);\n"; - - if( baseShape != TextureFlavor::Shape::ShapeCube ) - { - sb << "__target_intrinsic(glsl, \"$ctextureLodOffset($p, $2, $3, $4)$z\")\n"; - sb << "T SampleLevel(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount + isArray << " location, "; - sb << "float level, "; - sb << "constexpr int" << kBaseTextureTypes[tt].coordCount << " offset);\n"; - } - } - - sb << "\n};\n"; - - // `Gather*()` operations are handled via an `extension` declaration, - // because this lets us capture the element type of the texture. - // - // TODO: longer-term there should be something like a `TextureElementType` - // interface, that both scalars and vectors implement, that then exposes - // a `Scalar` associated type, and `Gather` can return `vector`. - // - static const struct { - char const* genericPrefix; - char const* elementType; - char const* outputType; - } kGatherExtensionCases[] = { - { "__generic", "vector", "vector" }, - { "", "float", "vector" }, - { "", "int" , "vector"}, - { "", "uint", "vector"}, - - // TODO: need a case here for scalars `T`, but also - // need to ensure that case doesn't accidentally match - // for `T = vector<...>`, which requires actual checking - // of constraints on generic parameters. - }; - for(auto cc : kGatherExtensionCases) - { - // TODO: this should really be an `if` around the entire `Gather` logic - if (isMultisample) break; - - EMIT_LINE_DIRECTIVE(); - sb << cc.genericPrefix << " __extension "; - sb << kBaseTextureAccessLevels[accessLevel].name; - sb << name; - if (isArray) sb << "Array"; - sb << "<" << cc.elementType << " >"; - sb << "\n{\n"; - - // `Gather` - // (tricky because it returns a 4-vector of the element type - // of the texture components...) - // - // TODO: is it actually correct to restrict these so that, e.g., - // `GatherAlpha()` isn't allowed on `Texture2D` because - // it nominally doesn't have an alpha component? - static const struct { - int componentIndex; - char const* componentName; - } kGatherComponets[] = { - { 0, "" }, - { 0, "Red" }, - { 1, "Green" }, - { 2, "Blue" }, - { 3, "Alpha" }, - }; - - for(auto kk : kGatherComponets) - { - auto componentIndex = kk.componentIndex; - auto componentName = kk.componentName; - - auto outputType = cc.outputType; - - EMIT_LINE_DIRECTIVE(); - - sb << "__target_intrinsic(glsl, \"textureGather($p, $2, " << componentIndex << ")\")\n"; - if (kBaseTextureTypes[tt].coordCount == 2) - { - // Gather only works on 2D in CUDA - // "It is based on the base type of DataType except when readMode is equal to cudaReadModeNormalizedFloat (see Texture Reference API), in which case it is always float4." - sb << "__target_intrinsic(cuda, \"tex2Dgather<$T0>($0, ($2).x, ($2).y, " << componentIndex << ")\")\n"; - } - sb << outputType << " Gather" << componentName << "(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount << " location);\n"; - - EMIT_LINE_DIRECTIVE(); - sb << "__target_intrinsic(glsl, \"textureGatherOffset($p, $2, $3, " << componentIndex << ")\")\n"; - sb << outputType << " Gather" << componentName << "(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount << " location, "; - sb << "constexpr int" << kBaseTextureTypes[tt].coordCount << " offset);\n"; - - EMIT_LINE_DIRECTIVE(); - sb << outputType << " Gather" << componentName << "(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount << " location, "; - sb << "constexpr int" << kBaseTextureTypes[tt].coordCount << " offset, "; - sb << "out uint status);\n"; - - EMIT_LINE_DIRECTIVE(); - sb << "__target_intrinsic(glsl, \"textureGatherOffsets($p, $2, int" << kBaseTextureTypes[tt].coordCount << "[]($3, $4, $5, $6), " << componentIndex << ")\")\n"; - sb << outputType << " Gather" << componentName << "(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount << " location, "; - sb << "int" << kBaseTextureTypes[tt].coordCount << " offset1, "; - sb << "int" << kBaseTextureTypes[tt].coordCount << " offset2, "; - sb << "int" << kBaseTextureTypes[tt].coordCount << " offset3, "; - sb << "int" << kBaseTextureTypes[tt].coordCount << " offset4);\n"; - - EMIT_LINE_DIRECTIVE(); - sb << outputType << " Gather" << componentName << "(SamplerState s, "; - sb << "float" << kBaseTextureTypes[tt].coordCount << " location, "; - sb << "int" << kBaseTextureTypes[tt].coordCount << " offset1, "; - sb << "int" << kBaseTextureTypes[tt].coordCount << " offset2, "; - sb << "int" << kBaseTextureTypes[tt].coordCount << " offset3, "; - sb << "int" << kBaseTextureTypes[tt].coordCount << " offset4, "; - sb << "out uint status);\n"; - } - - EMIT_LINE_DIRECTIVE(); - sb << "\n}\n"; - } - } - } -} - - -for (auto op : unaryOps) -{ - char const* fixity = (op.flags & POSTFIX) != 0 ? "__postfix " : "__prefix "; - char const* qual = (op.flags & ASSIGNMENT) != 0 ? "in out " : ""; - - for (auto type : kBaseTypes) - { - if ((type.flags & op.flags) == 0) - continue; - - char const* resultType = type.name; - if (op.flags & BOOL_RESULT) resultType = "bool"; - - // scalar version - sb << fixity; - sb << "__intrinsic_op(" << int(op.opCode) << ") " << resultType << " operator" << op.opName << "(" << qual << type.name << " value);\n"; - - // vector version - sb << "__generic "; - sb << fixity; - sb << "__intrinsic_op(" << int(op.opCode) << ") vector<" << resultType << ",N> operator" << op.opName << "(" << qual << "vector<" << type.name << ",N> value);\n"; - - // matrix version - sb << "__generic "; - sb << fixity; - sb << "__intrinsic_op(" << int(op.opCode) << ") matrix<" << resultType << ",N,M> operator" << op.opName << "(" << qual << "matrix<" << type.name << ",N,M> value);\n"; - } - - // Synthesize generic versions - if(op.interface) - { - char const* resultType = "T"; - if (op.flags & BOOL_RESULT) resultType = "bool"; - - // scalar version - sb << fixity; - sb << "__generic\n"; - sb << "__intrinsic_op(" << int(op.opCode) << ") " << resultType << " operator" << op.opName << "(" << qual << "T value);\n"; - - // vector version - sb << "__generic "; - sb << fixity; - sb << "__intrinsic_op(" << int(op.opCode) << ") vector<" << resultType << ",N> operator" << op.opName << "(" << qual << "vector value);\n"; - - // matrix version - sb << "__generic "; - sb << fixity; - sb << "__intrinsic_op(" << int(op.opCode) << ") matrix<" << resultType << ",N,M> operator" << op.opName << "(" << qual << "matrix value);\n"; - } -} - -for (auto op : binaryOps) -{ - char const* leftQual = ""; - if(op.flags & ASSIGNMENT) leftQual = "in out "; - - for (auto type : kBaseTypes) - { - if ((type.flags & op.flags) == 0) - continue; - - char const* leftType = type.name; - char const* rightType = leftType; - char const* resultType = leftType; - - if (op.flags & BOOL_RESULT) resultType = "bool"; - - // TODO: We should handle a `SHIFT` flag on the op - // by changing `rightType` to `int` in order to - // account for the fact that the shift amount should - // always have a fixed type independent of the LHS. - // - // (It is unclear why this change hadn't been made - // already, so it is possible that such a change - // breaks overload resolution or other parts of - // the compiler) - - // scalar version - sb << "__intrinsic_op(" << int(op.opCode) << ") " << resultType << " operator" << op.opName << "(" << leftQual << leftType << " left, " << rightType << " right);\n"; - - // vector version - sb << "__generic "; - sb << "__intrinsic_op(" << int(op.opCode) << ") vector<" << resultType << ",N> operator" << op.opName << "(" << leftQual << "vector<" << leftType << ",N> left, vector<" << rightType << ",N> right);\n"; - - // matrix version - sb << "__generic "; - sb << "__intrinsic_op(" << int(op.opCode) << ") matrix<" << resultType << ",N,M> operator" << op.opName << "(" << leftQual << "matrix<" << leftType << ",N,M> left, matrix<" << rightType << ",N,M> right);\n"; - - // We currently synthesize addiitonal overloads - // for the case where one or the other operand - // is a scalar. This choice serves a few purposes: - // - // 1. It avoids introducing scalar-to-vector or - // scalar-to-matrix promotions before the operator, - // which might allow some back ends to produce - // more optimal code. - // - // 2. It avoids concerns about making overload resolution - // and the inference rules for `N` and `M` able to - // handle the mixed vector/scalar or matrix/scalar case. - // - // 3. Having explicit overloads for the matrix/scalar cases - // here means that we do *not* need to support a general - // implicit conversion from scalars to matrices, unless - // we decide we want to. - // - // Note: Case (2) of the motivation shouldn't really apply - // any more, because we end up having to support similar - // inteference for built-in binary math functions where - // vectors and scalars might be combined (and where defining - // additional overloads to cover all the combinations doesn't - // seem practical or desirable). - // - // TODO: We should consider whether dropping these extra - // overloads is possible and worth it. The optimization - // concern (1) could possibly be addressed in specific - // back-ends. The issue (3) about not wanting to support - // implicit scalar-to-matrix conversion may be moot if - // we end up needing to support mixed scalar/matrix input - // for builtin in non-operator functions anyway. - - // scalar-vector and scalar-matrix - if (!(op.flags & ASSIGNMENT)) - { - sb << "__generic "; - sb << "__intrinsic_op(" << int(op.opCode) << ") vector<" << resultType << ",N> operator" << op.opName << "(" << leftQual << leftType << " left, vector<" << rightType << ",N> right);\n"; - - sb << "__generic "; - sb << "__intrinsic_op(" << int(op.opCode) << ") matrix<" << resultType << ",N,M> operator" << op.opName << "(" << leftQual << leftType << " left, matrix<" << rightType << ",N,M> right);\n"; - } - - // vector-scalar and matrix-scalar - sb << "__generic "; - sb << "__intrinsic_op(" << int(op.opCode) << ") vector<" << resultType << ",N> operator" << op.opName << "(" << leftQual << "vector<" << leftType << ",N> left, " << rightType << " right);\n"; - - sb << "__generic "; - sb << "__intrinsic_op(" << int(op.opCode) << ") matrix<" << resultType << ",N,M> operator" << op.opName << "(" << leftQual << "matrix<" << leftType << ",N,M> left, " << rightType << " right);\n"; - } - - // Synthesize generic versions - if(op.interface) - { - char const* leftType = "T"; - char const* rightType = leftType; - char const* resultType = leftType; - - if (op.flags & BOOL_RESULT) resultType = "bool"; - // TODO: handle `SHIFT` - - // scalar version - sb << "__generic\n"; - sb << "__intrinsic_op(" << int(op.opCode) << ") " << resultType << " operator" << op.opName << "(" << leftQual << leftType << " left, " << rightType << " right);\n"; - - // vector version - sb << "__generic "; - sb << "__intrinsic_op(" << int(op.opCode) << ") vector<" << resultType << ",N> operator" << op.opName << "(" << leftQual << "vector<" << leftType << ",N> left, vector<" << rightType << ",N> right);\n"; - - // matrix version - sb << "__generic "; - sb << "__intrinsic_op(" << int(op.opCode) << ") matrix<" << resultType << ",N,M> operator" << op.opName << "(" << leftQual << "matrix<" << leftType << ",N,M> left, matrix<" << rightType << ",N,M> right);\n"; - - // scalar-vector and scalar-matrix - if (!(op.flags & ASSIGNMENT)) - { - sb << "__generic "; - sb << "__intrinsic_op(" << int(op.opCode) << ") vector<" << resultType << ",N> operator" << op.opName << "(" << leftQual << leftType << " left, vector<" << rightType << ",N> right);\n"; - - sb << "__generic "; - sb << "__intrinsic_op(" << int(op.opCode) << ") matrix<" << resultType << ",N,M> operator" << op.opName << "(" << leftQual << leftType << " left, matrix<" << rightType << ",N,M> right);\n"; - } - - // vector-scalar and matrix-scalar - sb << "__generic "; - sb << "__intrinsic_op(" << int(op.opCode) << ") vector<" << resultType << ",N> operator" << op.opName << "(" << leftQual << "vector<" << leftType << ",N> left, " << rightType << " right);\n"; - - sb << "__generic "; - sb << "__intrinsic_op(" << int(op.opCode) << ") matrix<" << resultType << ",N,M> operator" << op.opName << "(" << leftQual << "matrix<" << leftType << ",N,M> left, " << rightType << " right);\n"; - } -} -SLANG_RAW("#line 1632 \"core.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("// Specialized function\n") -SLANG_RAW("\n") -SLANG_RAW("__intrinsic_op\n") -SLANG_RAW("int getStringHash(String string);\n") -SLANG_RAW("\n") -SLANG_RAW("// Operators to apply to `enum` types\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kIROp_Eql -) -SLANG_RAW(")\n") -SLANG_RAW("bool operator==(E left, E right);\n") -SLANG_RAW("\n") -SLANG_RAW("// Binding Attributes\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(DeclBase)\n") -SLANG_RAW("attribute_syntax [vk_binding(binding: int, set: int = 0)]\t\t\t: GLSLBindingAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(DeclBase)\n") -SLANG_RAW("attribute_syntax [gl_binding(binding: int, set: int = 0)]\t\t\t: GLSLBindingAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(VarDeclBase)\n") -SLANG_RAW("attribute_syntax [vk_shader_record]\t\t\t : ShaderRecordAttribute;\n") -SLANG_RAW("__attributeTarget(VarDeclBase)\n") -SLANG_RAW("attribute_syntax [shader_record]\t\t\t : ShaderRecordAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(DeclBase)\n") -SLANG_RAW("attribute_syntax [vk_push_constant]\t\t\t\t\t\t\t\t\t: PushConstantAttribute;\n") -SLANG_RAW("__attributeTarget(DeclBase)\n") -SLANG_RAW("attribute_syntax [push_constant]\t\t\t\t\t\t\t\t\t: PushConstantAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(VarDeclBase)\n") -SLANG_RAW("attribute_syntax [vk_location(locaiton : int)] : GLSLLocationAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(VarDeclBase)\n") -SLANG_RAW("attribute_syntax [vk_index(index : int)] : GLSLIndexAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("// Statement Attributes\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(LoopStmt)\n") -SLANG_RAW("attribute_syntax [unroll(count: int = 0)] : UnrollAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(LoopStmt)\n") -SLANG_RAW("attribute_syntax [loop] : LoopAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(LoopStmt)\n") -SLANG_RAW("attribute_syntax [fastopt] : FastOptAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(LoopStmt)\n") -SLANG_RAW("attribute_syntax [allow_uav_condition] : AllowUAVConditionAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(IfStmt)\n") -SLANG_RAW("attribute_syntax [flatten] : FlattenAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(IfStmt)\n") -SLANG_RAW("__attributeTarget(SwitchStmt)\n") -SLANG_RAW("attribute_syntax [branch] : BranchAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(SwitchStmt)\n") -SLANG_RAW("attribute_syntax [forcecase] : ForceCaseAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(SwitchStmt)\n") -SLANG_RAW("attribute_syntax [call] : CallAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("// Entry-point Attributes\n") -SLANG_RAW("\n") -SLANG_RAW("// All Stages\n") -SLANG_RAW("__attributeTarget(FuncDecl)\n") -SLANG_RAW("attribute_syntax [shader(stage)] : EntryPointAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("// Hull Shader\n") -SLANG_RAW("__attributeTarget(FuncDecl)\n") -SLANG_RAW("attribute_syntax [maxtessfactor(factor: float)] : MaxTessFactorAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(FuncDecl)\n") -SLANG_RAW("attribute_syntax [outputcontrolpoints(count: int)] : OutputControlPointsAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(FuncDecl)\n") -SLANG_RAW("attribute_syntax [outputtopology(topology)] : OutputTopologyAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(FuncDecl)\n") -SLANG_RAW("attribute_syntax [partitioning(mode)] : PartitioningAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(FuncDecl)\n") -SLANG_RAW("attribute_syntax [patchconstantfunc(name)] : PatchConstantFuncAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("// Hull/Domain Shader\n") -SLANG_RAW("__attributeTarget(FuncDecl)\n") -SLANG_RAW("attribute_syntax [domain(domain)] : DomainAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("// Geometry Shader\n") -SLANG_RAW("__attributeTarget(FuncDecl)\n") -SLANG_RAW("attribute_syntax [maxvertexcount(count: int)] : MaxVertexCountAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(FuncDecl)\n") -SLANG_RAW("attribute_syntax [instance(count: int)] : InstanceAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("// Fragment (\"Pixel\") Shader\n") -SLANG_RAW("__attributeTarget(FuncDecl)\n") -SLANG_RAW("attribute_syntax [earlydepthstencil] : EarlyDepthStencilAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("// Compute Shader\n") -SLANG_RAW("__attributeTarget(FuncDecl)\n") -SLANG_RAW("attribute_syntax [numthreads(x: int, y: int = 1, z: int = 1)] : NumThreadsAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("//\n") -SLANG_RAW("__attributeTarget(VarDeclBase)\n") -SLANG_RAW("attribute_syntax [__vulkanRayPayload] : VulkanRayPayloadAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(VarDeclBase)\n") -SLANG_RAW("attribute_syntax [__vulkanCallablePayload] : VulkanCallablePayloadAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(VarDeclBase)\n") -SLANG_RAW("attribute_syntax [__vulkanHitAttributes] : VulkanHitAttributesAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(FunctionDeclBase)\n") -SLANG_RAW("attribute_syntax [mutating] : MutatingAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW(" /// Indicates that a function computes its result as a function of its arguments without loading/storing any memory or other state.\n") -SLANG_RAW(" ///\n") -SLANG_RAW(" /// This is equivalent to the LLVM `readnone` function attribute.\n") -SLANG_RAW("__attributeTarget(FunctionDeclBase)\n") -SLANG_RAW("attribute_syntax [__readNone] : ReadNoneAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("enum _AttributeTargets\n") -SLANG_RAW("{\n") -SLANG_RAW(" Struct = ") -SLANG_SPLICE( (int) UserDefinedAttributeTargets::Struct -) -SLANG_RAW(",\n") -SLANG_RAW(" Var = ") -SLANG_SPLICE( (int) UserDefinedAttributeTargets::Var -) -SLANG_RAW(",\n") -SLANG_RAW(" Function = ") -SLANG_SPLICE( (int) UserDefinedAttributeTargets::Function -) -SLANG_RAW(",\n") -SLANG_RAW("};\n") -SLANG_RAW("__attributeTarget(StructDecl)\n") -SLANG_RAW("attribute_syntax [__AttributeUsage(target : _AttributeTargets)] : AttributeUsageAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(VarDeclBase)\n") -SLANG_RAW("attribute_syntax [format(format : String)] : FormatAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("__attributeTarget(Decl)\n") -SLANG_RAW("attribute_syntax [allow(diagnostic: String)] : AllowAttribute;\n") -SLANG_RAW("\n") -SLANG_RAW("// Linking\n") -SLANG_RAW("__attributeTarget(Decl)\n") -SLANG_RAW("attribute_syntax [__extern] : ExternAttribute;\n") -SLANG_RAW("\n") diff --git a/source/slang/glsl.meta.slang.h b/source/slang/glsl.meta.slang.h deleted file mode 100644 index f7634db31..000000000 --- a/source/slang/glsl.meta.slang.h +++ /dev/null @@ -1,202 +0,0 @@ -SLANG_RAW("// Slang GLSL compatibility library\n") -SLANG_RAW("\n") - -static const struct { - char const* name; - char const* glslPrefix; -} kTypes[] = -{ - {"float", ""}, - {"int", "i"}, - {"uint", "u"}, - {"bool", "b"}, -}; -static const int kTypeCount = sizeof(kTypes) / sizeof(kTypes[0]); - -for( int tt = 0; tt < kTypeCount; ++tt ) -{ - // Declare GLSL aliases for HLSL types - for (int vv = 2; vv <= 4; ++vv) - { - sb << "typedef vector<" << kTypes[tt].name << "," << vv << "> " << kTypes[tt].glslPrefix << "vec" << vv << ";\n"; - sb << "typedef matrix<" << kTypes[tt].name << "," << vv << "," << vv << "> " << kTypes[tt].glslPrefix << "mat" << vv << ";\n"; - } - for (int rr = 2; rr <= 4; ++rr) - for (int cc = 2; cc <= 4; ++cc) - { - sb << "typedef matrix<" << kTypes[tt].name << "," << rr << "," << cc << "> " << kTypes[tt].glslPrefix << "mat" << rr << "x" << cc << ";\n"; - } -} - -// Multiplication operations for vectors + matrices - -// scalar-vector and vector-scalar -sb << "__generic __intrinsic_op(mul) vector operator*(vector x, T y);\n"; -sb << "__generic __intrinsic_op(mul) vector operator*(T x, vector y);\n"; - -// scalar-matrix and matrix-scalar -sb << "__generic __intrinsic_op(mul) matrix operator*(matrix x, T y);\n"; -sb << "__generic __intrinsic_op(mul) matrix operator*(T x, matrix y);\n"; - -// vector-vector (dot product) -sb << "__generic __intrinsic_op(dot) T operator*(vector x, vector y);\n"; - -// vector-matrix -sb << "__generic __intrinsic_op(mul) vector operator*(vector x, matrix y);\n"; - -// matrix-vector -sb << "__generic __intrinsic_op(mul) vector operator*(matrix x, vector y);\n"; - -// matrix-matrix -sb << "__generic __intrinsic_op(mul) matrix operator*(matrix x, matrix y);\n"; - - - -// - -// TODO(tfoley): Need to handle `RW*` variants of texture types as well... -static const struct { - char const* name; - TextureFlavor::Shape baseShape; - int coordCount; -} kBaseTextureTypes[] = { - { "1D", TextureFlavor::Shape::Shape1D, 1 }, - { "2D", TextureFlavor::Shape::Shape2D, 2 }, - { "3D", TextureFlavor::Shape::Shape3D, 3 }, - { "Cube", TextureFlavor::Shape::ShapeCube, 3 }, - { "Buffer", TextureFlavor::Shape::ShapeBuffer, 1 }, -}; -static const int kBaseTextureTypeCount = sizeof(kBaseTextureTypes) / sizeof(kBaseTextureTypes[0]); - - -static const struct { - char const* name; - SlangResourceAccess access; -} kBaseTextureAccessLevels[] = { - { "", SLANG_RESOURCE_ACCESS_READ }, - { "RW", SLANG_RESOURCE_ACCESS_READ_WRITE }, - { "RasterizerOrdered", SLANG_RESOURCE_ACCESS_RASTER_ORDERED }, -}; -static const int kBaseTextureAccessLevelCount = sizeof(kBaseTextureAccessLevels) / sizeof(kBaseTextureAccessLevels[0]); - -for (int tt = 0; tt < kBaseTextureTypeCount; ++tt) -{ - char const* shapeName = kBaseTextureTypes[tt].name; - TextureFlavor::Shape baseShape = kBaseTextureTypes[tt].baseShape; - - for (int isArray = 0; isArray < 2; ++isArray) - { - // Arrays of 3D textures aren't allowed - if (isArray && baseShape == TextureFlavor::Shape::Shape3D) continue; - - for (int isMultisample = 0; isMultisample < 2; ++isMultisample) - { - auto readAccess = SLANG_RESOURCE_ACCESS_READ; - auto readWriteAccess = SLANG_RESOURCE_ACCESS_READ_WRITE; - - // TODO: any constraints to enforce on what gets to be multisampled? - - - unsigned flavor = baseShape; - if (isArray) flavor |= TextureFlavor::ArrayFlag; - if (isMultisample) flavor |= TextureFlavor::MultisampleFlag; -// if (isShadow) flavor |= TextureFlavor::ShadowFlag; - - - - unsigned readFlavor = flavor | (readAccess << 8); - unsigned readWriteFlavor = flavor | (readWriteAccess << 8); - - StringBuilder nameBuilder; - nameBuilder << shapeName; - if (isMultisample) nameBuilder << "MS"; - if (isArray) nameBuilder << "Array"; - auto name = nameBuilder.ProduceString(); - - sb << "__generic "; - sb << "__magic_type(TextureSampler," << int(readFlavor) << ") struct "; - sb << "__sampler" << name; - sb << " {};\n"; - - sb << "__generic "; - sb << "__magic_type(Texture," << int(readFlavor) << ") struct "; - sb << "__texture" << name; - sb << " {};\n"; - - sb << "__generic "; - sb << "__magic_type(GLSLImageType," << int(readWriteFlavor) << ") struct "; - sb << "__image" << name; - sb << " {};\n"; - - // TODO(tfoley): flesh this out for all the available prefixes - static const struct - { - char const* prefix; - char const* elementType; - } kTextureElementTypes[] = { - { "", "vec4" }, - { "i", "ivec4" }, - { "u", "uvec4" }, - { nullptr, nullptr }, - }; - for( auto ee = kTextureElementTypes; ee->prefix; ++ee ) - { - sb << "typedef __sampler" << name << "<" << ee->elementType << "> " << ee->prefix << "sampler" << name << ";\n"; - sb << "typedef __texture" << name << "<" << ee->elementType << "> " << ee->prefix << "texture" << name << ";\n"; - sb << "typedef __image" << name << "<" << ee->elementType << "> " << ee->prefix << "image" << name << ";\n"; - } - } - } -} - -sb << "__generic __magic_type(GLSLInputParameterGroupType) struct __GLSLInputParameterGroup {};\n"; -sb << "__generic __magic_type(GLSLOutputParameterGroupType) struct __GLSLOutputParameterGroup {};\n"; -sb << "__generic __magic_type(GLSLShaderStorageBufferType) struct __GLSLShaderStorageBuffer {};\n"; - -sb << "__magic_type(SamplerState," << int(SamplerStateFlavor::SamplerState) << ") struct sampler {};"; - -sb << "__magic_type(GLSLInputAttachmentType) struct subpassInput {};"; - -// Define additional keywords - -sb << "syntax buffer : GLSLBufferModifier;\n"; - -// [GLSL 4.3] Storage Qualifiers - -// TODO: need to support `shared` here with its GLSL meaning - -sb << "syntax patch : GLSLPatchModifier;\n"; -// `centroid` and `sample` handled centrally - -// [GLSL 4.5] Interpolation Qualifiers -sb << "syntax smooth : SimpleModifier;\n"; -sb << "syntax flat : SimpleModifier;\n"; -sb << "syntax noperspective : SimpleModifier;\n"; - - -// [GLSL 4.3.2] Constant Qualifier - -// We need to handle GLSL `const` separately from HLSL `const`, -// since they mean such different things. - -// [GLSL 4.7.2] Precision Qualifiers -sb << "syntax highp : SimpleModifier;\n"; -sb << "syntax mediump : SimpleModifier;\n"; -sb << "syntax lowp : SimpleModifier;\n"; - -// [GLSL 4.8.1] The Invariant Qualifier - -sb << "syntax invariant : SimpleModifier;\n"; - -// [GLSL 4.10] Memory Qualifiers - -sb << "syntax coherent : SimpleModifier;\n"; -sb << "syntax volatile : SimpleModifier;\n"; -sb << "syntax restrict : SimpleModifier;\n"; -sb << "syntax readonly : GLSLReadOnlyModifier;\n"; -sb << "syntax writeonly : GLSLWriteOnlyModifier;\n"; - -// We will treat `subroutine` as a qualifier for now -sb << "syntax subroutine : SimpleModifier;\n"; -SLANG_RAW("\n") -SLANG_RAW("\n") diff --git a/source/slang/hlsl.meta.slang.h b/source/slang/hlsl.meta.slang.h deleted file mode 100644 index 216f45bbe..000000000 --- a/source/slang/hlsl.meta.slang.h +++ /dev/null @@ -1,3140 +0,0 @@ -SLANG_RAW("// Slang HLSL compatibility library\n") -SLANG_RAW("\n") -SLANG_RAW("typedef uint UINT;\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__magic_type(HLSLAppendStructuredBufferType)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_HLSLAppendStructuredBufferType -) -SLANG_RAW(")\n") -SLANG_RAW("struct AppendStructuredBuffer\n") -SLANG_RAW("{\n") -SLANG_RAW(" void Append(T value);\n") -SLANG_RAW("\n") -SLANG_RAW(" void GetDimensions(\n") -SLANG_RAW(" out uint numStructs,\n") -SLANG_RAW(" out uint stride);\n") -SLANG_RAW("};\n") -SLANG_RAW("\n") -SLANG_RAW("__magic_type(HLSLByteAddressBufferType)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_HLSLByteAddressBufferType -) -SLANG_RAW(")\n") -SLANG_RAW("struct ByteAddressBuffer\n") -SLANG_RAW("{\n") -SLANG_RAW(" __target_intrinsic(glsl, \"$1 = $0._data.length() * 4\")\n") -SLANG_RAW(" void GetDimensions(\n") -SLANG_RAW(" out uint dim);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"$0._data[$1/4]\")\n") -SLANG_RAW(" uint Load(int location);\n") -SLANG_RAW("\n") -SLANG_RAW(" uint Load(int location, out uint status);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"uvec2($0._data[$1/4], $0._data[$1/4+1])\")\n") -SLANG_RAW(" uint2 Load2(int location);\n") -SLANG_RAW("\n") -SLANG_RAW(" uint2 Load2(int location, out uint status);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"uvec3($0._data[$1/4], $0._data[$1/4+1], $0._data[$1/4+2])\")\n") -SLANG_RAW(" uint3 Load3(int location);\n") -SLANG_RAW("\n") -SLANG_RAW(" uint3 Load3(int location, out uint status);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"uvec4($0._data[$1/4], $0._data[$1/4+1], $0._data[$1/4+2], $0._data[$1/4+3])\")\n") -SLANG_RAW(" uint4 Load4(int location);\n") -SLANG_RAW("\n") -SLANG_RAW(" uint4 Load4(int location, out uint status);\n") -SLANG_RAW("};\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__magic_type(HLSLStructuredBufferType)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_HLSLStructuredBufferType -) -SLANG_RAW(")\n") -SLANG_RAW("struct StructuredBuffer\n") -SLANG_RAW("{\n") -SLANG_RAW(" __target_intrinsic(glsl, \"$1 = $0._data.length(); $2 = 0\")\n") -SLANG_RAW(" void GetDimensions(\n") -SLANG_RAW(" out uint numStructs,\n") -SLANG_RAW(" out uint stride);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"$0._data[$1]\") \n") -SLANG_RAW(" T Load(int location);\n") -SLANG_RAW(" T Load(int location, out uint status);\n") -SLANG_RAW("\n") -SLANG_RAW(" __subscript(uint index) -> T\n") -SLANG_RAW(" {\n") -SLANG_RAW(" __target_intrinsic(glsl, \"$0._data[$1]\")\n") -SLANG_RAW(" get;\n") -SLANG_RAW(" };\n") -SLANG_RAW("};\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__magic_type(HLSLConsumeStructuredBufferType)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_HLSLConsumeStructuredBufferType -) -SLANG_RAW(")\n") -SLANG_RAW("struct ConsumeStructuredBuffer\n") -SLANG_RAW("{\n") -SLANG_RAW(" T Consume();\n") -SLANG_RAW("\n") -SLANG_RAW(" void GetDimensions(\n") -SLANG_RAW(" out uint numStructs,\n") -SLANG_RAW(" out uint stride);\n") -SLANG_RAW("};\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__magic_type(HLSLInputPatchType)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_HLSLInputPatchType -) -SLANG_RAW(")\n") -SLANG_RAW("struct InputPatch\n") -SLANG_RAW("{\n") -SLANG_RAW(" __subscript(uint index) -> T;\n") -SLANG_RAW("};\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__magic_type(HLSLOutputPatchType)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_HLSLOutputPatchType -) -SLANG_RAW(")\n") -SLANG_RAW("struct OutputPatch\n") -SLANG_RAW("{\n") -SLANG_RAW(" __subscript(uint index) -> T;\n") -SLANG_RAW("};\n") -SLANG_RAW("\n") - -static const struct { - IROp op; - char const* name; -} kMutableByteAddressBufferCases[] = -{ - { kIROp_HLSLRWByteAddressBufferType, "RWByteAddressBuffer" }, - { kIROp_HLSLRasterizerOrderedByteAddressBufferType, "RasterizerOrderedByteAddressBuffer" }, -}; -for(auto item : kMutableByteAddressBufferCases) { -SLANG_RAW("#line 105 \"hlsl.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("__magic_type(HLSL") -SLANG_SPLICE(item.name -) -SLANG_RAW("Type)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(item.op -) -SLANG_RAW(")\n") -SLANG_RAW("struct ") -SLANG_SPLICE(item.name -) -SLANG_RAW("\n") -SLANG_RAW("{\n") -SLANG_RAW(" // Note(tfoley): supports all operations from `ByteAddressBuffer`\n") -SLANG_RAW(" // TODO(tfoley): can this be made a sub-type?\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"$1 = $0._data.length() * 4\")\n") -SLANG_RAW(" void GetDimensions(\n") -SLANG_RAW(" out uint dim);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"$0._data[$1/4]\")\n") -SLANG_RAW(" uint Load(int location);\n") -SLANG_RAW("\n") -SLANG_RAW(" uint Load(int location, out uint status);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"uvec2($0._data[$1/4], $0._data[$1/4+1])\")\n") -SLANG_RAW(" uint2 Load2(int location);\n") -SLANG_RAW("\n") -SLANG_RAW(" uint2 Load2(int location, out uint status);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"uvec3($0._data[$1/4], $0._data[$1/4+1], $0._data[$1/4+2])\")\n") -SLANG_RAW(" uint3 Load3(int location);\n") -SLANG_RAW("\n") -SLANG_RAW(" uint3 Load3(int location, out uint status);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"uvec4($0._data[$1/4], $0._data[$1/4+1], $0._data[$1/4+2], $0._data[$1/4+3])\")\n") -SLANG_RAW(" uint4 Load4(int location);\n") -SLANG_RAW("\n") -SLANG_RAW(" uint4 Load4(int location, out uint status);\n") -SLANG_RAW("\n") -SLANG_RAW(" // Added operations:\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"($3 = atomicAdd($0._data[$1/4], $2))\")\n") -SLANG_RAW(" void InterlockedAdd(\n") -SLANG_RAW(" UINT dest,\n") -SLANG_RAW(" UINT value,\n") -SLANG_RAW(" out UINT original_value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"atomicAdd($0._data[$1/4], $2)\")\n") -SLANG_RAW(" void InterlockedAdd(\n") -SLANG_RAW(" UINT dest,\n") -SLANG_RAW(" UINT value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"($3 = atomicAnd($0._data[$1/4], $2))\")\n") -SLANG_RAW(" void InterlockedAnd(\n") -SLANG_RAW(" UINT dest,\n") -SLANG_RAW(" UINT value,\n") -SLANG_RAW(" out UINT original_value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"atomicAnd($0._data[$1/4], $2)\")\n") -SLANG_RAW(" void InterlockedAnd(\n") -SLANG_RAW(" UINT dest,\n") -SLANG_RAW(" UINT value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"($4 = atomicCompSwap($0._data[$1/4], $2, $3))\")\n") -SLANG_RAW(" void InterlockedCompareExchange(\n") -SLANG_RAW(" UINT dest,\n") -SLANG_RAW(" UINT compare_value,\n") -SLANG_RAW(" UINT value,\n") -SLANG_RAW(" out UINT original_value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"atomicCompSwap($0._data[$1/4], $2, $3)\")\n") -SLANG_RAW(" void InterlockedCompareStore(\n") -SLANG_RAW(" UINT dest,\n") -SLANG_RAW(" UINT compare_value,\n") -SLANG_RAW(" UINT value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"($3 = atomicExchange($0._data[$1/4], $2))\")\n") -SLANG_RAW(" void InterlockedExchange(\n") -SLANG_RAW(" UINT dest,\n") -SLANG_RAW(" UINT value,\n") -SLANG_RAW(" out UINT original_value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"($3 = atomicMax($0._data[$1/4], $2))\")\n") -SLANG_RAW(" void InterlockedMax(\n") -SLANG_RAW(" UINT dest,\n") -SLANG_RAW(" UINT value,\n") -SLANG_RAW(" out UINT original_value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"atomicMax($0._data[$1/4], $2)\")\n") -SLANG_RAW(" void InterlockedMax(\n") -SLANG_RAW(" UINT dest,\n") -SLANG_RAW(" UINT value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"($3 = atomicMin($0._data[$1/4], $2))\")\n") -SLANG_RAW(" void InterlockedMin(\n") -SLANG_RAW(" UINT dest,\n") -SLANG_RAW(" UINT value,\n") -SLANG_RAW(" out UINT original_value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"atomicMin($0._data[$1/4], $2)\")\n") -SLANG_RAW(" void InterlockedMin(\n") -SLANG_RAW(" UINT dest,\n") -SLANG_RAW(" UINT value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"($3 = atomicOr($0._data[$1/4], $2))\")\n") -SLANG_RAW(" void InterlockedOr(\n") -SLANG_RAW(" UINT dest,\n") -SLANG_RAW(" UINT value,\n") -SLANG_RAW(" out UINT original_value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"atomicOr($0._data[$1/4], $2)\")\n") -SLANG_RAW(" void InterlockedOr(\n") -SLANG_RAW(" UINT dest,\n") -SLANG_RAW(" UINT value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"($3 = atomicXor($0._data[$1/4], $2))\")\n") -SLANG_RAW(" void InterlockedXor(\n") -SLANG_RAW(" UINT dest,\n") -SLANG_RAW(" UINT value,\n") -SLANG_RAW(" out UINT original_value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"atomicXor($0._data[$1/4], $2)\")\n") -SLANG_RAW(" void InterlockedXor(\n") -SLANG_RAW(" UINT dest,\n") -SLANG_RAW(" UINT value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"$0._data[$1/4] = $2\")\n") -SLANG_RAW(" void Store(\n") -SLANG_RAW(" uint address,\n") -SLANG_RAW(" uint value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"$0._data[$1/4] = $2.x, $0._data[$1/4+1] = $2.y\")\n") -SLANG_RAW(" void Store2(\n") -SLANG_RAW(" uint address,\n") -SLANG_RAW(" uint2 value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"$0._data[$1/4] = $2.x, $0._data[$1/4+1] = $2.y, $0._data[$1/4+2] = $2.z\")\n") -SLANG_RAW(" void Store3(\n") -SLANG_RAW(" uint address,\n") -SLANG_RAW(" uint3 value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"$0._data[$1/4] = $2.x, $0._data[$1/4+1] = $2.y, $0._data[$1/4+2] = $2.z, $0._data[$1/4+3] = $2.w\")\n") -SLANG_RAW(" void Store4(\n") -SLANG_RAW(" uint address,\n") -SLANG_RAW(" uint4 value);\n") -SLANG_RAW("};\n") -SLANG_RAW("\n") - -} -SLANG_RAW("#line 248 \"hlsl.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW("\n") - -static const struct { - IROp op; - char const* name; -} kMutableStructuredBufferCases[] = -{ - { kIROp_HLSLRWStructuredBufferType, "RWStructuredBuffer" }, - { kIROp_HLSLRasterizerOrderedStructuredBufferType, "RasterizerOrderedStructuredBuffer" }, -}; -for(auto item : kMutableStructuredBufferCases) { -SLANG_RAW("#line 260 \"hlsl.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__magic_type(HLSL") -SLANG_SPLICE(item.name -) -SLANG_RAW("Type)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(item.op -) -SLANG_RAW(")\n") -SLANG_RAW("struct ") -SLANG_SPLICE(item.name -) -SLANG_RAW("\n") -SLANG_RAW("{\n") -SLANG_RAW(" uint DecrementCounter();\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"$1 = $0._data.length(); $2 = 0\")\n") -SLANG_RAW(" void GetDimensions(\n") -SLANG_RAW(" out uint numStructs,\n") -SLANG_RAW(" out uint stride);\n") -SLANG_RAW("\n") -SLANG_RAW(" uint IncrementCounter();\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"$0._data[$1]\")\n") -SLANG_RAW(" T Load(int location);\n") -SLANG_RAW(" T Load(int location, out uint status);\n") -SLANG_RAW("\n") -SLANG_RAW(" __subscript(uint index) -> T\n") -SLANG_RAW(" {\n") -SLANG_RAW(" __target_intrinsic(glsl, \"$0._data[$1]\")\n") -SLANG_RAW(" ref;\n") -SLANG_RAW(" }\n") -SLANG_RAW("};\n") -SLANG_RAW("\n") - -} -SLANG_RAW("#line 290 \"hlsl.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__magic_type(HLSLPointStreamType)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_HLSLPointStreamType -) -SLANG_RAW(")\n") -SLANG_RAW("struct PointStream\n") -SLANG_RAW("{\n") -SLANG_RAW(" __target_intrinsic(glsl, \"EmitVertex()\")\n") -SLANG_RAW(" void Append(T value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"EndPrimitive()\")\n") -SLANG_RAW(" void RestartStrip();\n") -SLANG_RAW("};\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__magic_type(HLSLLineStreamType)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_HLSLLineStreamType -) -SLANG_RAW(")\n") -SLANG_RAW("struct LineStream\n") -SLANG_RAW("{\n") -SLANG_RAW(" __target_intrinsic(glsl, \"EmitVertex()\")\n") -SLANG_RAW(" void Append(T value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"EndPrimitive()\")\n") -SLANG_RAW(" void RestartStrip();\n") -SLANG_RAW("};\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__magic_type(HLSLTriangleStreamType)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_HLSLTriangleStreamType -) -SLANG_RAW(")\n") -SLANG_RAW("struct TriangleStream\n") -SLANG_RAW("{\n") -SLANG_RAW(" __target_intrinsic(glsl, \"EmitVertex()\")\n") -SLANG_RAW(" void Append(T value);\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(glsl, \"EndPrimitive()\")\n") -SLANG_RAW(" void RestartStrip();\n") -SLANG_RAW("};\n") -SLANG_RAW("\n") -SLANG_RAW("#define VECTOR_MAP_UNARY(TYPE, COUNT, FUNC, VALUE) \\\n") -SLANG_RAW(" vector result; for(int i = 0; i < COUNT; ++i) { result[i] = FUNC(VALUE[i]); } return result\n") -SLANG_RAW(" \n") -SLANG_RAW("#define MATRIX_MAP_UNARY(TYPE, ROWS, COLS, FUNC, VALUE) \\\n") -SLANG_RAW(" matrix result; for(int i = 0; i < ROWS; ++i) { result[i] = FUNC(VALUE[i]); } return result\n") -SLANG_RAW("\n") -SLANG_RAW("#define VECTOR_MAP_BINARY(TYPE, COUNT, FUNC, LEFT, RIGHT) \\\n") -SLANG_RAW(" vector result; for(int i = 0; i < COUNT; ++i) { result[i] = FUNC(LEFT[i], RIGHT[i]); } return result\n") -SLANG_RAW(" \n") -SLANG_RAW("#define MATRIX_MAP_BINARY(TYPE, ROWS, COLS, FUNC, LEFT, RIGHT) \\\n") -SLANG_RAW(" matrix result; for(int i = 0; i < ROWS; ++i) { result[i] = FUNC(LEFT[i], RIGHT[i]); } return result\n") -SLANG_RAW("\n") -SLANG_RAW("#define VECTOR_MAP_TRINARY(TYPE, COUNT, FUNC, A, B, C) \\\n") -SLANG_RAW(" vector result; for(int i = 0; i < COUNT; ++i) { result[i] = FUNC(A[i], B[i], C[i]); } return result\n") -SLANG_RAW(" \n") -SLANG_RAW("#define MATRIX_MAP_TRINARY(TYPE, ROWS, COLS, FUNC, A, B, C) \\\n") -SLANG_RAW(" matrix result; for(int i = 0; i < ROWS; ++i) { result[i] = FUNC(A[i], B[i], C[i]); } return result\n") -SLANG_RAW("\n") -SLANG_RAW("// Try to terminate the current draw or dispatch call (HLSL SM 4.0)\n") -SLANG_RAW("void abort();\n") -SLANG_RAW("\n") -SLANG_RAW("// Absolute value (HLSL SM 1.0)\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T abs(T x);\n") -SLANG_RAW("/*{\n") -SLANG_RAW(" return x < 0 ? -x : x;\n") -SLANG_RAW("}*/\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector abs(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, abs, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix abs(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, abs, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Inverse cosine (HLSL SM 1.0)\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T acos(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector acos(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, acos, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix acos(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, acos, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Test if all components are non-zero (HLSL SM 1.0)\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, \"bool($0)\")\n") -SLANG_RAW("bool all(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"all(bvec$N0($0))\")\n") -SLANG_RAW("bool all(vector x);\n") -SLANG_RAW("// TODO: implementation of `all()` in the stdlib is\n") -SLANG_RAW("// blocked on fixing implementation of `bool` vector\n") -SLANG_RAW("// `getAt` on the CUDA codegen path.\n") -SLANG_RAW("/*\n") -SLANG_RAW("{\n") -SLANG_RAW(" bool result = true;\n") -SLANG_RAW(" for(int i = 0; i < N; ++i)\n") -SLANG_RAW(" result = result && all(x[i]);\n") -SLANG_RAW(" return result;\n") -SLANG_RAW("}\n") -SLANG_RAW("*/\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("bool all(matrix x);\n") -SLANG_RAW("/*\n") -SLANG_RAW("{\n") -SLANG_RAW(" bool result = true;\n") -SLANG_RAW(" for(int i = 0; i < N; ++i)\n") -SLANG_RAW(" result = result && all(x[i]);\n") -SLANG_RAW(" return result;\n") -SLANG_RAW("}\n") -SLANG_RAW("*/\n") -SLANG_RAW("\n") -SLANG_RAW("// Barrier for writes to all memory spaces (HLSL SM 5.0)\n") -SLANG_RAW("__target_intrinsic(glsl, \"memoryBarrier(), groupMemoryBarrier(), memoryBarrierImage(), memoryBarrierBuffer()\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"__threadfence()\")\n") -SLANG_RAW("void AllMemoryBarrier();\n") -SLANG_RAW("\n") -SLANG_RAW("// Thread-group sync and barrier for writes to all memory spaces (HLSL SM 5.0)\n") -SLANG_RAW("__target_intrinsic(glsl, \"memoryBarrier(), groupMemoryBarrier(), memoryBarrierImage(), memoryBarrierBuffer(), barrier()\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"__syncthreads()\")\n") -SLANG_RAW("void AllMemoryBarrierWithGroupSync();\n") -SLANG_RAW("\n") -SLANG_RAW("// Test if any components is non-zero (HLSL SM 1.0)\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, \"bool($0)\")\n") -SLANG_RAW("bool any(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"any(bvec$N0($0))\")\n") -SLANG_RAW("bool any(vector x);\n") -SLANG_RAW("// TODO: implementation of `any()` in the stdlib is\n") -SLANG_RAW("// blocked on fixing implementation of `bool` vector\n") -SLANG_RAW("// `getAt` on the CUDA codegen path.\n") -SLANG_RAW("/*\n") -SLANG_RAW("{\n") -SLANG_RAW(" bool result = false;\n") -SLANG_RAW(" for(int i = 0; i < N; ++i)\n") -SLANG_RAW(" result = result || any(x[i]);\n") -SLANG_RAW(" return result;\n") -SLANG_RAW("}\n") -SLANG_RAW("*/\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("bool any(matrix x);\n") -SLANG_RAW("/*\n") -SLANG_RAW("{\n") -SLANG_RAW(" bool result = false;\n") -SLANG_RAW(" for(int i = 0; i < N; ++i)\n") -SLANG_RAW(" result = result || any(x[i]);\n") -SLANG_RAW(" return result;\n") -SLANG_RAW("}\n") -SLANG_RAW("*/\n") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("// Reinterpret bits as a double (HLSL SM 5.0)\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"packDouble2x32(uvec2($0, $1))\")\n") -SLANG_RAW("__glsl_extension(GL_ARB_gpu_shader5)\n") -SLANG_RAW("double asdouble(uint lowbits, uint highbits);\n") -SLANG_RAW("\n") -SLANG_RAW("// Reinterpret bits as a float (HLSL SM 4.0)\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"intBitsToFloat\")\n") -SLANG_RAW("float asfloat(int x);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"uintBitsToFloat\")\n") -SLANG_RAW("float asfloat(uint x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"intBitsToFloat\")\n") -SLANG_RAW("vector asfloat(vector< int, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(float, N, asfloat, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"uintBitsToFloat\")\n") -SLANG_RAW("vector asfloat(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(float, N, asfloat, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix asfloat(matrix< int,N,M> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(float, N, M, asfloat, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix asfloat(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(float, N, M, asfloat, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// No op\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Pos -) -SLANG_RAW(")\n") -SLANG_RAW("float asfloat(float x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Pos -) -SLANG_RAW(")\n") -SLANG_RAW("vector asfloat(vector x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Pos -) -SLANG_RAW(")\n") -SLANG_RAW("matrix asfloat(matrix x);\n") -SLANG_RAW("\n") -SLANG_RAW("// Inverse sine (HLSL SM 1.0)\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("T asin(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector asin(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T,N,asin,x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix asin(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T,N,M,asin,x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Reinterpret bits as an int (HLSL SM 4.0)\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"floatBitsToInt\")\n") -SLANG_RAW("int asint(float x);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"int($0)\")\n") -SLANG_RAW("int asint(uint x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"floatBitsToInt\")\n") -SLANG_RAW("vector asint(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(int, N, asint, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"ivec$N0($0)\")\n") -SLANG_RAW("vector asint(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(int, N, asint, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix asint(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(int, N, M, asint, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix asint(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(int, N, M, asint, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// No op\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Pos -) -SLANG_RAW(")\n") -SLANG_RAW("int asint(int x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Pos -) -SLANG_RAW(")\n") -SLANG_RAW("vector asint(vector x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Pos -) -SLANG_RAW(")\n") -SLANG_RAW("matrix asint(matrix x);\n") -SLANG_RAW("\n") -SLANG_RAW("// Reinterpret bits of double as a uint (HLSL SM 5.0)\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"{ uvec2 v = unpackDouble2x32($0); $1 = v.x; $2 = v.y; }\")\n") -SLANG_RAW("__glsl_extension(GL_ARB_gpu_shader5)\n") -SLANG_RAW("void asuint(double value, out uint lowbits, out uint highbits);\n") -SLANG_RAW("\n") -SLANG_RAW("// Reinterpret bits as a uint (HLSL SM 4.0)\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"floatBitsToUint\")\n") -SLANG_RAW("uint asuint(float x);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"uint($0)\")\n") -SLANG_RAW("uint asuint(int x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"floatBitsToUint\")\n") -SLANG_RAW("vector asuint(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(uint, N, asuint, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"uvec$N0($0)\")\n") -SLANG_RAW("vector asuint(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(uint, N, asuint, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix asuint(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(uint, N, M, asuint, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix asuint(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(uint, N, M, asuint, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Pos -) -SLANG_RAW(")\n") -SLANG_RAW("uint asuint(uint x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Pos -) -SLANG_RAW(")\n") -SLANG_RAW("vector asuint(vector x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Pos -) -SLANG_RAW(")\n") -SLANG_RAW("matrix asuint(matrix x);\n") -SLANG_RAW("\n") -SLANG_RAW("// Inverse tangent (HLSL SM 1.0)\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T atan(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector atan(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, atan, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix atan(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, atan, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl,\"atan($0,$1)\")\n") -SLANG_RAW("T atan2(T y, T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl,\"atan($0,$1)\")\n") -SLANG_RAW("vector atan2(vector y, vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_BINARY(T, N, atan2, y, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix atan2(matrix y, matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_BINARY(T, N, M, atan2, y, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Ceiling (HLSL SM 1.0)\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T ceil(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector ceil(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, ceil, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix ceil(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, ceil, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("// Check access status to tiled resource\n") -SLANG_RAW("bool CheckAccessFullyMapped(uint status);\n") -SLANG_RAW("\n") -SLANG_RAW("// Clamp (HLSL SM 1.0)\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("T clamp(T x, T minBound, T maxBound)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return min(max(x, minBound), maxBound);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector clamp(vector x, vector minBound, vector maxBound)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return min(max(x, minBound), maxBound);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix clamp(matrix x, matrix minBound, matrix maxBound)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return min(max(x, minBound), maxBound);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Clip (discard) fragment conditionally\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("void clip(T x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" if(x < T(0)) discard;\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("void clip(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" if(any(x < T(0))) discard;\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("void clip(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" if(any(x < T(0))) discard;\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Cosine\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T cos(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector cos(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T,N, cos, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix cos(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, cos, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Hyperbolic cosine\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T cosh(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector cosh(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T,N, cosh, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix cosh(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, cosh, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Population count\n") -SLANG_RAW("__target_intrinsic(glsl, \"bitCount\")\n") -SLANG_RAW("uint countbits(uint value);\n") -SLANG_RAW("\n") -SLANG_RAW("// Cross product\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector cross(vector left, vector right)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return vector(\n") -SLANG_RAW(" left.y * right.z - left.z * right.y,\n") -SLANG_RAW(" left.z * right.x - left.x * right.z,\n") -SLANG_RAW(" left.x * right.y - left.y * right.x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Convert encoded color\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("int4 D3DCOLORtoUBYTE4(float4 color)\n") -SLANG_RAW("{\n") -SLANG_RAW(" let scaled = color.zyxw * 255.001999f;\n") -SLANG_RAW(" return int4(scaled);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Partial-difference derivatives\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, dFdx)\n") -SLANG_RAW("T ddx(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, dFdx)\n") -SLANG_RAW("vector ddx(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, ddx, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix ddx(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, ddx, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__glsl_extension(GL_ARB_derivative_control)\n") -SLANG_RAW("__target_intrinsic(glsl, dFdxCoarse)\n") -SLANG_RAW("T ddx_coarse(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__glsl_extension(GL_ARB_derivative_control)\n") -SLANG_RAW("__target_intrinsic(glsl, dFdxCoarse)\n") -SLANG_RAW("vector ddx_coarse(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, ddx_coarse, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix ddx_coarse(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, ddx_coarse, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__glsl_extension(GL_ARB_derivative_control)\n") -SLANG_RAW("__target_intrinsic(glsl, dFdxFine)\n") -SLANG_RAW("T ddx_fine(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__glsl_extension(GL_ARB_derivative_control)\n") -SLANG_RAW("__target_intrinsic(glsl, dFdxFine)\n") -SLANG_RAW("vector ddx_fine(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, ddx_fine, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix ddx_fine(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, ddx_fine, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, dFdy)\n") -SLANG_RAW("T ddy(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, dFdy)\n") -SLANG_RAW("vector ddy(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, ddy, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix ddy(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, ddy, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_ARB_derivative_control)\n") -SLANG_RAW("__target_intrinsic(glsl, dFdyCoarse)\n") -SLANG_RAW("T ddy_coarse(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__glsl_extension(GL_ARB_derivative_control)\n") -SLANG_RAW("__target_intrinsic(glsl, dFdyCoarse)\n") -SLANG_RAW("vector ddy_coarse(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, ddy_coarse, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix ddy_coarse(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, ddy_coarse, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__glsl_extension(GL_ARB_derivative_control)\n") -SLANG_RAW("__target_intrinsic(glsl, dFdyFine)\n") -SLANG_RAW("T ddy_fine(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__glsl_extension(GL_ARB_derivative_control)\n") -SLANG_RAW("__target_intrinsic(glsl, dFdyFine)\n") -SLANG_RAW("vector ddy_fine(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, ddy_fine, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix ddy_fine(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, ddy_fine, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("// Radians to degrees\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("T degrees(T x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return x * (T(180) / T.getPi());\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector degrees(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, degrees, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix degrees(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, degrees, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Matrix determinant\n") -SLANG_RAW("\n") -SLANG_RAW("__generic T determinant(matrix m);\n") -SLANG_RAW("\n") -SLANG_RAW("// Barrier for device memory\n") -SLANG_RAW("__target_intrinsic(glsl, \"memoryBarrier(), memoryBarrierImage(), memoryBarrierBuffer()\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"__threadfence()\")\n") -SLANG_RAW("void DeviceMemoryBarrier();\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"memoryBarrier(), memoryBarrierImage(), memoryBarrierBuffer(), barrier()\")\n") -SLANG_RAW("__target_intrinsic(glsl, \"__syncthreads()\")\n") -SLANG_RAW("void DeviceMemoryBarrierWithGroupSync();\n") -SLANG_RAW("\n") -SLANG_RAW("// Vector distance\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("T distance(vector x, vector y)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return length(x - y);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Vector dot product\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("T dot(vector x, vector y)\n") -SLANG_RAW("{\n") -SLANG_RAW(" T result = T(0);\n") -SLANG_RAW(" for(int i = 0; i < N; ++i)\n") -SLANG_RAW(" result += x[i] * y[i];\n") -SLANG_RAW(" return result;\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Helper for computing distance terms for lighting (obsolete)\n") -SLANG_RAW("\n") -SLANG_RAW("__generic vector dst(vector x, vector y);\n") -SLANG_RAW("\n") -SLANG_RAW("// Error message\n") -SLANG_RAW("\n") -SLANG_RAW("// void errorf( string format, ... );\n") -SLANG_RAW("\n") -SLANG_RAW("// Attribute evaluation\n") -SLANG_RAW("\n") -SLANG_RAW("// TODO: The matrix cases of these functions won't actuall work\n") -SLANG_RAW("// when compiled to GLSL, since they only support scalar/vector\n") -SLANG_RAW("\n") -SLANG_RAW("// TODO: Should these be constrains to `__BuiltinFloatingPointType`?\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, interpolateAtCentroid)\n") -SLANG_RAW("T EvaluateAttributeAtCentroid(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, interpolateAtCentroid)\n") -SLANG_RAW("vector EvaluateAttributeAtCentroid(vector x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, interpolateAtCentroid)\n") -SLANG_RAW("matrix EvaluateAttributeAtCentroid(matrix x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, \"interpolateAtSample($0, int($1))\")\n") -SLANG_RAW("T EvaluateAttributeAtSample(T x, uint sampleindex);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, \"interpolateAtSample($0, int($1))\")\n") -SLANG_RAW("vector EvaluateAttributeAtSample(vector x, uint sampleindex);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, \"interpolateAtSample($0, int($1))\")\n") -SLANG_RAW("matrix EvaluateAttributeAtSample(matrix x, uint sampleindex);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, \"interpolateAtOffset($0, vec2($1) / 16.0f)\")\n") -SLANG_RAW("T EvaluateAttributeSnapped(T x, int2 offset);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, \"interpolateAtOffset($0, vec2($1) / 16.0f)\")\n") -SLANG_RAW("vector EvaluateAttributeSnapped(vector x, int2 offset);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, \"interpolateAtOffset($0, vec2($1) / 16.0f)\")\n") -SLANG_RAW("matrix EvaluateAttributeSnapped(matrix x, int2 offset);\n") -SLANG_RAW("\n") -SLANG_RAW("// Base-e exponent\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T exp(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector exp(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, exp, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix exp(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, exp, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Base-2 exponent\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T exp2(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector exp2(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, exp2, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix exp2(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, exp2, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Convert 16-bit float stored in low bits of integer\n") -SLANG_RAW("__target_intrinsic(glsl, \"unpackHalf2x16($0).x\")\n") -SLANG_RAW("float f16tof32(uint value);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, \"unpackHalf2x16($0).x\")\n") -SLANG_RAW("vector f16tof32(vector value);\n") -SLANG_RAW("\n") -SLANG_RAW("// Convert to 16-bit float stored in low bits of integer\n") -SLANG_RAW("__target_intrinsic(glsl, \"packHalf2x16(vec2($0,0.0))\")\n") -SLANG_RAW("uint f32tof16(float value);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, \"packHalf2x16(vec2($0,0.0))\")\n") -SLANG_RAW("vector f32tof16(vector value);\n") -SLANG_RAW("\n") -SLANG_RAW("// Flip surface normal to face forward, if needed\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector faceforward(vector n, vector i, vector ng)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return dot(ng, i) < T(0.0f) ? n : -n;\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Find first set bit starting at high bit and working down\n") -SLANG_RAW("__target_intrinsic(glsl,\"findMSB\")\n") -SLANG_RAW("int firstbithigh(int value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl,\"findMSB\")\n") -SLANG_RAW("__generic\n") -SLANG_RAW("vector firstbithigh(vector value)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(int, N, firstbithigh, value);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl,\"findMSB\")\n") -SLANG_RAW("uint firstbithigh(uint value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl,\"findMSB\")\n") -SLANG_RAW("__generic\n") -SLANG_RAW("vector firstbithigh(vector value)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(uint, N, firstbithigh, value);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Find first set bit starting at low bit and working up\n") -SLANG_RAW("__target_intrinsic(glsl,\"findLSB\")\n") -SLANG_RAW("int firstbitlow(int value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl,\"findLSB\")\n") -SLANG_RAW("__generic\n") -SLANG_RAW("vector firstbitlow(vector value)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(int, N, firstbitlow, value);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl,\"findLSB\")\n") -SLANG_RAW("uint firstbitlow(uint value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl,\"findLSB\")\n") -SLANG_RAW("__generic\n") -SLANG_RAW("vector firstbitlow(vector value)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(uint, N, firstbitlow, value);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Floor (HLSL SM 1.0)\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T floor(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector floor(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, floor, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix floor(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, floor, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Fused multiply-add for doubles\n") -SLANG_RAW("double fma(double a, double b, double c);\n") -SLANG_RAW("__generic vector fma(vector a, vector b, vector c);\n") -SLANG_RAW("__generic matrix fma(matrix a, matrix b, matrix c);\n") -SLANG_RAW("\n") -SLANG_RAW("// Floating point remainder of x/y\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T fmod(T x, T y);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector fmod(vector x, vector y)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_BINARY(T, N, fmod, x, y);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix fmod(matrix x, matrix y)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_BINARY(T, N, M, fmod, x, y);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Fractional part\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, fract)\n") -SLANG_RAW("T frac(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, fract)\n") -SLANG_RAW("vector frac(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, frac, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix frac(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, frac, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Split float into mantissa and exponent\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("T frexp(T x, out T exp);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector frexp(vector x, out vector exp)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_BINARY(T, N, frexp, x, exp);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix frexp(matrix x, out matrix exp)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_BINARY(T, N, M, frexp, x, exp);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Texture filter width\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T fwidth(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector fwidth(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, fwidth, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix fwidth(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, fwidth, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Get number of samples in render target\n") -SLANG_RAW("uint GetRenderTargetSampleCount();\n") -SLANG_RAW("\n") -SLANG_RAW("// Get position of given sample\n") -SLANG_RAW("float2 GetRenderTargetSamplePosition(int Index);\n") -SLANG_RAW("\n") -SLANG_RAW("// Group memory barrier\n") -SLANG_RAW("__target_intrinsic(glsl, \"groupMemoryBarrier\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"__threadfence_block\")\n") -SLANG_RAW("void GroupMemoryBarrier();\n") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"groupMemoryBarrier(), barrier()\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"__syncthreads()\")\n") -SLANG_RAW("void GroupMemoryBarrierWithGroupSync();\n") -SLANG_RAW("\n") -SLANG_RAW("// Atomics\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"$atomicAdd($A, $1)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"atomicAdd(&$0, $1)\")\n") -SLANG_RAW("void InterlockedAdd(__ref int dest, int value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"$atomicAdd($A, $1)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"atomicAdd((uint*)&$0, $1)\")\n") -SLANG_RAW("void InterlockedAdd(__ref uint dest, uint value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicAdd($A, $1))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicAdd(&$0, $1))\")\n") -SLANG_RAW("void InterlockedAdd(__ref int dest, int value, out int original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicAdd($A, $1))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($2 = (uint)atomicAdd((uint*)&$0, $1))\")\n") -SLANG_RAW("void InterlockedAdd(__ref uint dest, uint value, out uint original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"$atomicAnd($A, $1)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"atomicAnd(&$0, $1)\")\n") -SLANG_RAW("void InterlockedAnd(__ref int dest, int value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"$atomicAnd($A, $1)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"atomicAnd((int*)&$0, $1)\")\n") -SLANG_RAW("void InterlockedAnd(__ref uint dest, uint value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicAnd($A, $1))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicAnd(&$0, $1))\")\n") -SLANG_RAW("void InterlockedAnd(__ref int dest, int value, out int original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicAnd($A, $1))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicAnd((int*)&$0, $1))\")\n") -SLANG_RAW("void InterlockedAnd(__ref uint dest, uint value, out uint original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($3 = $atomicCompSwap($A, $1, $2))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($3 = atomicCAS(&$0, $1, $2))\")\n") -SLANG_RAW("void InterlockedCompareExchange(__ref int dest, int compare_value, int value, out int original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($3 = $atomicCompSwap($A, $1, $2))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($3 = (uint)atomicCAS((int*)&$0, $1, $2))\")\n") -SLANG_RAW("void InterlockedCompareExchange(__ref uint dest, uint compare_value, uint value, out uint original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"$atomicCompSwap($A, $1, $2)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"atomicCAS(&$0, $1, $2)\")\n") -SLANG_RAW("void InterlockedCompareStore(__ref int dest, int compare_value, int value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"$atomicCompSwap($A, $1, $2)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"atomicCAS((int*)&$0, $1, $2)\")\n") -SLANG_RAW("void InterlockedCompareStore(__ref uint dest, uint compare_value, uint value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicExchange($A, $1))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicExch(&$0, $1))\")\n") -SLANG_RAW("void InterlockedExchange(__ref int dest, int value, out int original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicExchange($A, $1))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($2 = (uint)atomicExch((int*)&$0, $1))\")\n") -SLANG_RAW("void InterlockedExchange(__ref uint dest, uint value, out uint original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"$atomicMax($A, $1)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"atomicMax(&$0, $1)\")\n") -SLANG_RAW("void InterlockedMax(__ref int dest, int value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"$atomicMax($A, $1)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"atomicMax((int*)&$0, $1)\")\n") -SLANG_RAW("void InterlockedMax(__ref uint dest, uint value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicMax($A, $1))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicMax(&$0, $1))\")\n") -SLANG_RAW("void InterlockedMax(__ref int dest, int value, out int original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicMax($A, $1))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($2 = (uint)atomicMax((int*)&$0, $1))\")\n") -SLANG_RAW("void InterlockedMax(__ref uint dest, uint value, out uint original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"$atomicMin($A, $1)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"atomicMin(&$0, $1)\")\n") -SLANG_RAW("void InterlockedMin(__ref int dest, int value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"$atomicMin($A, $1)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"atomicMin((int*)&$0, $1)\")\n") -SLANG_RAW("void InterlockedMin(__ref uint dest, uint value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicMin($A, $1))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicMin(&$0, $1))\")\n") -SLANG_RAW("void InterlockedMin(__ref int dest, int value, out int original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicMin($A, $1))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($2 = (uint)atomicMin((int*)&$0, $1))\")\n") -SLANG_RAW("void InterlockedMin(__ref uint dest, uint value, out uint original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"$atomicOr($A, $1)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"atomicOr(&$0, $1)\")\n") -SLANG_RAW("void InterlockedOr(__ref int dest, int value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"$atomicOr($A, $1)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"atomicOr((int*)&$0, $1)\")\n") -SLANG_RAW("void InterlockedOr(__ref uint dest, uint value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicOr($A, $1))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicOr(&$0, $1))\")\n") -SLANG_RAW("void InterlockedOr(__ref int dest, int value, out int original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicOr($A, $1))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($2 = (uint)atomicOr((int*)&$0, $1))\")\n") -SLANG_RAW("void InterlockedOr(__ref uint dest, uint value, out uint original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"$atomicXor($A, $1)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"atomicXor(&$0, $1)\")\n") -SLANG_RAW("void InterlockedXor(__ref int dest, int value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"$atomicXor($A, $1)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"atomicXor((int*)&$0, $1)\")\n") -SLANG_RAW("void InterlockedXor(__ref uint dest, uint value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicXor($A, $1))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($2 = atomicXor(&$0, $1))\")\n") -SLANG_RAW("void InterlockedXor(__ref int dest, int value, out int original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"($2 = $atomicXor($A, $1))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"($2 = (uint)atomicXor((int*)&$0, $1))\")\n") -SLANG_RAW("void InterlockedXor(__ref uint dest, uint value, out uint original_value);\n") -SLANG_RAW("\n") -SLANG_RAW("// Is floating-point value finite?\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(cpu)\n") -SLANG_RAW("__target_intrinsic(cuda)\n") -SLANG_RAW("//__target_intrinsic(glsl, \"(!(isinf($0) || isnan($0)))\")\n") -SLANG_RAW("bool isfinite(T x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return !(isinf(x) || isnan(x));\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("//__target_intrinsic(glsl, \"(!(isinf($0) || isnan($0)))\")\n") -SLANG_RAW("vector isfinite(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(bool, N, isfinite, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix isfinite(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(bool, N, M, isfinite, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Is floating-point value infinite?\n") -SLANG_RAW("__generic\n") -SLANG_RAW("bool isinf(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector isinf(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(bool, N, isinf, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix isinf(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(bool, N, M, isinf, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Is floating-point value not-a-number?\n") -SLANG_RAW("__generic\n") -SLANG_RAW("bool isnan(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector isnan(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(bool, N, isnan, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix isnan(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(bool, N, M, isnan, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Construct float from mantissa and exponent\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"($0 * pow(2.0f, $1))\")\n") -SLANG_RAW("T ldexp(T x, T exp);\n") -SLANG_RAW("/*{\n") -SLANG_RAW(" return x * exp2(exp);\n") -SLANG_RAW("}*/\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"($0 * pow(2.0f, $1))\")\n") -SLANG_RAW("vector ldexp(vector x, vector exp)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_BINARY(T, N, ldexp, x, exp);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix ldexp(matrix x, matrix exp)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_BINARY(T, N, M, ldexp, x, exp);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Vector length\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("T length(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return sqrt(dot(x, x));\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Linear interpolation\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, mix)\n") -SLANG_RAW("T lerp(T x, T y, T s);\n") -SLANG_RAW("/*{\n") -SLANG_RAW(" return x * (1 - s) + y * s;\n") -SLANG_RAW("}*/\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, mix)\n") -SLANG_RAW("vector lerp(vector x, vector y, vector s)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_TRINARY(T, N, lerp, x, y, s);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix lerp(matrix x, matrix y, matrix s)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_TRINARY(T, N, M, lerp, x, y, s);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Legacy lighting function (obsolete)\n") -SLANG_RAW("float4 lit(float n_dot_l, float n_dot_h, float m);\n") -SLANG_RAW("\n") -SLANG_RAW("// Base-e logarithm\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T log(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector log(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, log, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix log(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, log, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Base-10 logarithm\n") -SLANG_RAW("__generic \n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"(log( $0 ) * $S0( 0.43429448190325182765112891891661) )\" )\n") -SLANG_RAW("T log10(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"(log( $0 ) * $S0(0.43429448190325182765112891891661) )\" )\n") -SLANG_RAW("vector log10(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, log10, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic \n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix log10(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, log10, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Base-2 logarithm\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T log2(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector log2(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, log2, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix log2(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, log2, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// multiply-add\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, fma)\n") -SLANG_RAW("__generic T mad(T mvalue, T avalue, T bvalue);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, fma)\n") -SLANG_RAW("__generic vector mad(vector mvalue, vector avalue, vector bvalue);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, fma)\n") -SLANG_RAW("__generic matrix mad(matrix mvalue, matrix avalue, matrix bvalue);\n") -SLANG_RAW("\n") -SLANG_RAW("// maximum\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T max(T x, T y);\n") -SLANG_RAW("// Note: a stdlib implementation of `max` (or `min`) will require splitting\n") -SLANG_RAW("// floating-point and integer cases apart, because the floating-point\n") -SLANG_RAW("// version needs to correctly handle the case where one of the inputs\n") -SLANG_RAW("// is not-a-number.\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector max(vector x, vector y)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_BINARY(T, N, max, x, y);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix max(matrix x, matrix y)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_BINARY(T, N, M, max, x, y);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// minimum\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T min(T x, T y);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector min(vector x, vector y)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_BINARY(T, N, min, x, y);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix min(matrix x, matrix y)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_BINARY(T, N, M, min, x, y);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// split into integer and fractional parts (both with same sign)\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T modf(T x, out T ip);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector modf(vector x, out vector ip)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_BINARY(T, N, modf, x, ip);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix modf(matrix x, out matrix ip)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_BINARY(T, N, M, modf, x, ip);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// msad4 (whatever that is)\n") -SLANG_RAW("uint4 msad4(uint reference, uint2 source, uint4 accum);\n") -SLANG_RAW("\n") -SLANG_RAW("// General inner products\n") -SLANG_RAW("\n") -SLANG_RAW("// scalar-scalar\n") -SLANG_RAW("__generic T mul(T x, T y);\n") -SLANG_RAW("\n") -SLANG_RAW("// scalar-vector and vector-scalar\n") -SLANG_RAW("__generic vector mul(vector x, T y);\n") -SLANG_RAW("__generic vector mul(T x, vector y);\n") -SLANG_RAW("\n") -SLANG_RAW("// scalar-matrix and matrix-scalar\n") -SLANG_RAW("__generic matrix mul(matrix x, T y);\n") -SLANG_RAW("__generic matrix mul(T x, matrix y);\n") -SLANG_RAW("\n") -SLANG_RAW("// vector-vector (dot product)\n") -SLANG_RAW("__generic __intrinsic_op(dot) T mul(vector x, vector y);\n") -SLANG_RAW("\n") -SLANG_RAW("// vector-matrix\n") -SLANG_RAW("__generic __intrinsic_op(mulVectorMatrix) vector mul(vector x, matrix y);\n") -SLANG_RAW("\n") -SLANG_RAW("// matrix-vector\n") -SLANG_RAW("__generic __intrinsic_op(mulMatrixVector) vector mul(matrix x, vector y);\n") -SLANG_RAW("\n") -SLANG_RAW("// matrix-matrix\n") -SLANG_RAW("__generic __intrinsic_op(mulMatrixMatrix) matrix mul(matrix x, matrix y);\n") -SLANG_RAW("\n") -SLANG_RAW("// noise (deprecated)\n") -SLANG_RAW("\n") -SLANG_RAW("float noise(float x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return 0;\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic float noise(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return 0;\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("/// Indicate that an index may be non-uniform at execution time.\n") -SLANG_RAW("///\n") -SLANG_RAW("/// Shader Model 5.1 and 6.x introduce support for dynamic indexing\n") -SLANG_RAW("/// of arrays of resources, but place the restriction that *by default*\n") -SLANG_RAW("/// the implementation can assume that any value used as an index into\n") -SLANG_RAW("/// such arrays will be dynamically uniform across an entire `Draw` or `Dispatch`\n") -SLANG_RAW("/// (when using instancing, the value must be uniform across all instances;\n") -SLANG_RAW("/// it does not seem that the restriction extends to draws within a multi-draw).\n") -SLANG_RAW("///\n") -SLANG_RAW("/// In order to indicate to the implementation that it cannot make the\n") -SLANG_RAW("/// uniformity assumption, a shader programmer is required to pass the index\n") -SLANG_RAW("/// to the `NonUniformResourceIndex` function before using it as an index.\n") -SLANG_RAW("/// The function superficially acts like an identity function.\n") -SLANG_RAW("///\n") -SLANG_RAW("/// Note: a future version of Slang may take responsibility for inserting calls\n") -SLANG_RAW("/// to this function as necessary in output code, rather than make this\n") -SLANG_RAW("/// the user's responsibility, so that the default behavior of the language\n") -SLANG_RAW("/// is more semantically \"correct.\"\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, nonuniformEXT)\n") -SLANG_RAW("__glsl_extension(GL_EXT_nonuniform_qualifier)\n") -SLANG_RAW("[__readNone]\n") -SLANG_RAW("uint NonUniformResourceIndex(uint index)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return index;\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, nonuniformEXT)\n") -SLANG_RAW("__glsl_extension(GL_EXT_nonuniform_qualifier)\n") -SLANG_RAW("[__readNone]\n") -SLANG_RAW("int NonUniformResourceIndex(int index)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return index;\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Normalize a vector\n") -SLANG_RAW("__generic vector normalize(vector x);\n") -SLANG_RAW("/*{\n") -SLANG_RAW(" return x / length(x);\n") -SLANG_RAW("}*/\n") -SLANG_RAW("\n") -SLANG_RAW("// Raise to a power\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T pow(T x, T y);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector pow(vector x, vector y)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_BINARY(T, N, pow, x, y);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix pow(matrix x, matrix y)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_BINARY(T, N, M, pow, x, y);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Output message\n") -SLANG_RAW("\n") -SLANG_RAW("// void printf( string format, ... );\n") -SLANG_RAW("\n") -SLANG_RAW("// Tessellation factor fixup routines\n") -SLANG_RAW("\n") -SLANG_RAW("void Process2DQuadTessFactorsAvg(\n") -SLANG_RAW(" in float4 RawEdgeFactors,\n") -SLANG_RAW(" in float2 InsideScale,\n") -SLANG_RAW(" out float4 RoundedEdgeTessFactors,\n") -SLANG_RAW(" out float2 RoundedInsideTessFactors,\n") -SLANG_RAW(" out float2 UnroundedInsideTessFactors);\n") -SLANG_RAW("\n") -SLANG_RAW("void Process2DQuadTessFactorsMax(\n") -SLANG_RAW(" in float4 RawEdgeFactors,\n") -SLANG_RAW(" in float2 InsideScale,\n") -SLANG_RAW(" out float4 RoundedEdgeTessFactors,\n") -SLANG_RAW(" out float2 RoundedInsideTessFactors,\n") -SLANG_RAW(" out float2 UnroundedInsideTessFactors);\n") -SLANG_RAW("\n") -SLANG_RAW("void Process2DQuadTessFactorsMin(\n") -SLANG_RAW(" in float4 RawEdgeFactors,\n") -SLANG_RAW(" in float2 InsideScale,\n") -SLANG_RAW(" out float4 RoundedEdgeTessFactors,\n") -SLANG_RAW(" out float2 RoundedInsideTessFactors,\n") -SLANG_RAW(" out float2 UnroundedInsideTessFactors);\n") -SLANG_RAW("\n") -SLANG_RAW("void ProcessIsolineTessFactors(\n") -SLANG_RAW(" in float RawDetailFactor,\n") -SLANG_RAW(" in float RawDensityFactor,\n") -SLANG_RAW(" out float RoundedDetailFactor,\n") -SLANG_RAW(" out float RoundedDensityFactor);\n") -SLANG_RAW("\n") -SLANG_RAW("void ProcessQuadTessFactorsAvg(\n") -SLANG_RAW(" in float4 RawEdgeFactors,\n") -SLANG_RAW(" in float InsideScale,\n") -SLANG_RAW(" out float4 RoundedEdgeTessFactors,\n") -SLANG_RAW(" out float2 RoundedInsideTessFactors,\n") -SLANG_RAW(" out float2 UnroundedInsideTessFactors);\n") -SLANG_RAW("\n") -SLANG_RAW("void ProcessQuadTessFactorsMax(\n") -SLANG_RAW(" in float4 RawEdgeFactors,\n") -SLANG_RAW(" in float InsideScale,\n") -SLANG_RAW(" out float4 RoundedEdgeTessFactors,\n") -SLANG_RAW(" out float2 RoundedInsideTessFactors,\n") -SLANG_RAW(" out float2 UnroundedInsideTessFactors);\n") -SLANG_RAW("\n") -SLANG_RAW("void ProcessQuadTessFactorsMin(\n") -SLANG_RAW(" in float4 RawEdgeFactors,\n") -SLANG_RAW(" in float InsideScale,\n") -SLANG_RAW(" out float4 RoundedEdgeTessFactors,\n") -SLANG_RAW(" out float2 RoundedInsideTessFactors,\n") -SLANG_RAW(" out float2 UnroundedInsideTessFactors);\n") -SLANG_RAW("\n") -SLANG_RAW("void ProcessTriTessFactorsAvg(\n") -SLANG_RAW(" in float3 RawEdgeFactors,\n") -SLANG_RAW(" in float InsideScale,\n") -SLANG_RAW(" out float3 RoundedEdgeTessFactors,\n") -SLANG_RAW(" out float RoundedInsideTessFactor,\n") -SLANG_RAW(" out float UnroundedInsideTessFactor);\n") -SLANG_RAW("\n") -SLANG_RAW("void ProcessTriTessFactorsMax(\n") -SLANG_RAW(" in float3 RawEdgeFactors,\n") -SLANG_RAW(" in float InsideScale,\n") -SLANG_RAW(" out float3 RoundedEdgeTessFactors,\n") -SLANG_RAW(" out float RoundedInsideTessFactor,\n") -SLANG_RAW(" out float UnroundedInsideTessFactor);\n") -SLANG_RAW("\n") -SLANG_RAW("void ProcessTriTessFactorsMin(\n") -SLANG_RAW(" in float3 RawEdgeFactors,\n") -SLANG_RAW(" in float InsideScale,\n") -SLANG_RAW(" out float3 RoundedEdgeTessFactors,\n") -SLANG_RAW(" out float RoundedInsideTessFactors,\n") -SLANG_RAW(" out float UnroundedInsideTessFactors);\n") -SLANG_RAW("\n") -SLANG_RAW("// Degrees to radians\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("T radians(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector radians(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, radians, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix radians(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, radians, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Approximate reciprocal\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"1.0/($0)\")\n") -SLANG_RAW("T rcp(T x);\n") -SLANG_RAW("/*{\n") -SLANG_RAW(" return T(1) / x;\n") -SLANG_RAW("}*/\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("vector rcp(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, rcp, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("// Note: GLSL doesn't define a vector `rcp`, so not intrinsic there\n") -SLANG_RAW("matrix rcp(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, rcp, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Reflect incident vector across plane with given normal\n") -SLANG_RAW("__generic\n") -SLANG_RAW("vector reflect(vector i, vector n);\n") -SLANG_RAW("/*{\n") -SLANG_RAW(" return i - T(2) * dot(n,i) * n;\n") -SLANG_RAW("}*/\n") -SLANG_RAW("\n") -SLANG_RAW("// Refract incident vector given surface normal and index of refraction\n") -SLANG_RAW("__generic\n") -SLANG_RAW("vector refract(vector i, vector n, float eta);\n") -SLANG_RAW("/*{\n") -SLANG_RAW(" let dotNI = dot(n,i);\n") -SLANG_RAW(" let k = T(1) - eta*eta*(T(1) - dotNI * dotNI);\n") -SLANG_RAW(" if(k < 0) return vector(T(0));\n") -SLANG_RAW(" return eta * i - (eta * dotNI + sqrt(k)) * n;\n") -SLANG_RAW("}*/\n") -SLANG_RAW("\n") -SLANG_RAW("// Reverse order of bits\n") -SLANG_RAW("__target_intrinsic(glsl, \"bitfieldReverse\")\n") -SLANG_RAW("uint reversebits(uint value);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"bitfieldReverse\")\n") -SLANG_RAW("__generic vector reversebits(vector value);\n") -SLANG_RAW("\n") -SLANG_RAW("// Round-to-nearest\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T round(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector round(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, round, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix round(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, round, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Reciprocal of square root\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"inversesqrt($0)\")\n") -SLANG_RAW("T rsqrt(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"inversesqrt($0)\")\n") -SLANG_RAW("vector rsqrt(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, rsqrt, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix rsqrt(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, rsqrt, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Clamp value to [0,1] range\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("T saturate(T x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return clamp(x, T(0), T(1));\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("vector saturate(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return clamp(x,\n") -SLANG_RAW(" vector(T(0)),\n") -SLANG_RAW(" vector(T(1)));\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix saturate(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, saturate, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Extract sign of value\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, \"int(sign($0))\")\n") -SLANG_RAW("int sign(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"ivec$N0(sign($0))\")\n") -SLANG_RAW("vector sign(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(int, N, sign, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix sign(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(int, N, M, sign, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("// Sine\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T sin(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector sin(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, sin, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix sin(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, sin, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Sine and cosine\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("void sincos(T x, out T s, out T c)\n") -SLANG_RAW("{\n") -SLANG_RAW(" s = sin(x);\n") -SLANG_RAW(" c = cos(x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("void sincos(vector x, out vector s, out vector c)\n") -SLANG_RAW("{\n") -SLANG_RAW(" s = sin(x);\n") -SLANG_RAW(" c = cos(x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("void sincos(matrix x, out matrix s, out matrix c)\n") -SLANG_RAW("{\n") -SLANG_RAW(" s = sin(x);\n") -SLANG_RAW(" c = cos(x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Hyperbolic Sine\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T sinh(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector sinh(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, sinh, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix sinh(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, sinh, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Smooth step (Hermite interpolation)\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T smoothstep(T min, T max, T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector smoothstep(vector min, vector max, vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_TRINARY(T, N, smoothstep, min, max, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix smoothstep(matrix min, matrix max, matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_TRINARY(T, N, M, smoothstep, min, max, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Square root\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T sqrt(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector sqrt(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, sqrt, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix sqrt(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, sqrt, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Step function\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("T step(T y, T x);\n") -SLANG_RAW("/*{\n") -SLANG_RAW(" return x < y ? T(0.0f) : T(1.0f);\n") -SLANG_RAW("}*/\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector step(vector y, vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_BINARY(T, N, step, y, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix step(matrix y, matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_BINARY(T, N, M, step, y, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Tangent\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T tan(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector tan(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, tan, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix tan(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, tan, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Hyperbolic tangent\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T tanh(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector tanh(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, tanh, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix tanh(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, tanh, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Matrix transpose\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("matrix transpose(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" matrix result;\n") -SLANG_RAW(" for(int r = 0; r < M; ++r)\n") -SLANG_RAW(" for(int c = 0; c < N; ++c)\n") -SLANG_RAW(" result[r][c] = x[c][r];\n") -SLANG_RAW(" return result;\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Truncate to integer\n") -SLANG_RAW("__generic\n") -SLANG_RAW("T trunc(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector trunc(vector x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, trunc, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix trunc(matrix x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" MATRIX_MAP_UNARY(T, N, M, trunc, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// Shader model 6.0 stuff\n") -SLANG_RAW("\n") -SLANG_RAW("// Information for GLSL wave/subgroup support\n") -SLANG_RAW("// https://github.com/KhronosGroup/GLSL/blob/master/extensions/khr/GL_KHR_shader_subgroup.txt\n") -SLANG_RAW("\n") -SLANG_RAW("__generic T QuadReadLaneAt(T sourceValue, uint quadLaneID);\n") -SLANG_RAW("__generic vector QuadReadLaneAt(vector sourceValue, uint quadLaneID);\n") -SLANG_RAW("__generic matrix QuadReadLaneAt(matrix sourceValue, uint quadLaneID);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic T QuadReadAcrossX(T localValue);\n") -SLANG_RAW("__generic vector QuadReadAcrossX(vector localValue);\n") -SLANG_RAW("__generic matrix QuadReadAcrossX(matrix localValue);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic T QuadReadAcrossY(T localValue);\n") -SLANG_RAW("__generic vector QuadReadAcrossY(vector localValue);\n") -SLANG_RAW("__generic matrix QuadReadAcrossY(matrix localValue);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic T QuadReadAcrossDiagonal(T localValue);\n") -SLANG_RAW("__generic vector QuadReadAcrossDiagonal(vector localValue);\n") -SLANG_RAW("__generic matrix QuadReadAcrossDiagonal(matrix localValue);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupAnd($0)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"_waveAnd(__activemask(), $0)\")\n") -SLANG_RAW("T WaveActiveBitAnd(T expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupAnd($0)\")\n") -SLANG_RAW("vector WaveActiveBitAnd(vector expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix WaveActiveBitAnd(matrix expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupOr($0)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"_waveOr(__activemask(), $0)\")\n") -SLANG_RAW("T WaveActiveBitOr(T expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupOr($0)\")\n") -SLANG_RAW("vector WaveActiveBitOr(vector expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix WaveActiveBitOr(matrix expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupXor($0)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"_waveXor(__activemask(), $0)\")\n") -SLANG_RAW("T WaveActiveBitXor(T expr);\n") -SLANG_RAW("__generic \n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupXor($0)\")\n") -SLANG_RAW("vector WaveActiveBitXor(vector expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix WaveActiveBitXor(matrix expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupMax($0)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"_waveMax(__activemask(), $0)\")\n") -SLANG_RAW("T WaveActiveMax(T expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupMax($0)\")\n") -SLANG_RAW("vector WaveActiveMax(vector expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix WaveActiveMax(matrix expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupMin($0)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"_waveMin(__activemask(), $0)\")\n") -SLANG_RAW("T WaveActiveMin(T expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupMin($0)\")\n") -SLANG_RAW("vector WaveActiveMin(vector expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix WaveActiveMin(matrix expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupMul($0)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"_waveProduct(__activemask(), $0)\")\n") -SLANG_RAW("T WaveActiveProduct(T expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupMul($0)\")\n") -SLANG_RAW("vector WaveActiveProduct(vector expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix WaveActiveProduct(matrix expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupAdd($0)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"_waveSum(__activemask(), $0)\")\n") -SLANG_RAW("T WaveActiveSum(T expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupAdd($0)\")\n") -SLANG_RAW("vector WaveActiveSum(vector expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix WaveActiveSum(matrix expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_vote)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupAllEqual($0)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"_waveAllEqual(__activemask(), $0)\")\n") -SLANG_RAW("bool WaveActiveAllEqual(T value);\n") -SLANG_RAW("__generic \n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_vote)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupAllEqual($0)\")\n") -SLANG_RAW("vector WaveActiveAllEqual(vector value);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix WaveActiveAllEqual(matrix value);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic uint4 WaveMatch(T value);\n") -SLANG_RAW("__generic uint4 WaveMatch(vector value);\n") -SLANG_RAW("__generic uint4 WaveMatch(matrix value);\n") -SLANG_RAW("\n") -SLANG_RAW("// TODO(JS): For CUDA the article claims mask has to be used carefully\n") -SLANG_RAW("// https://devblogs.nvidia.com/using-cuda-warp-level-primitives/\n") -SLANG_RAW("// With the Warp intrinsics there is no mask, and it's just the 'active lanes'. So __activemask()\n") -SLANG_RAW("// seems to be appropriate.\n") -SLANG_RAW("\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_vote)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupAll($0)\") \n") -SLANG_RAW("__target_intrinsic(cuda, \"(__all_sync(__activemask(), $0) != 0)\") \n") -SLANG_RAW("bool WaveActiveAllTrue(bool condition);\n") -SLANG_RAW("\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_vote)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupAny($0)\") \n") -SLANG_RAW("__target_intrinsic(cuda, \"(__any_sync(__activemask(), $0) != 0)\")\n") -SLANG_RAW("bool WaveActiveAnyTrue(bool condition);\n") -SLANG_RAW("\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_ballot)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupBallot($0)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"make_uint4(__ballot_sync(__activemask(), $0), 0, 0, 0)\")\n") -SLANG_RAW("uint4 WaveActiveBallot(bool condition);\n") -SLANG_RAW("\n") -SLANG_RAW("// TODO(JS): \n") -SLANG_RAW("// subgroupBallotBitCount seems to take a uint4 parameter. \n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_ballot)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupBallotBitCount($0)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"__popc(__ballot_sync(__activemask(), $0))\")\n") -SLANG_RAW("uint WaveActiveCountBits(bool value);\n") -SLANG_RAW("\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_basic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"gl_SubgroupSize\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"(warpSize)\")\n") -SLANG_RAW("uint WaveGetLaneCount();\n") -SLANG_RAW("\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_basic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"gl_SubgroupInvocationID\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"_getLaneId()\")\n") -SLANG_RAW("uint WaveGetLaneIndex();\n") -SLANG_RAW("\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_basic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupElect()\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"_waveIsFirstLane()\")\n") -SLANG_RAW("bool WaveIsFirstLane();\n") -SLANG_RAW("\n") -SLANG_RAW("// TODO(JS): We cannot calculate prefix sums using a mask of __activemask() & __lanemask_lt(), because (amongst other reasons)\n") -SLANG_RAW("// that would mean different lanes having a different mask, and they all have to have the same mask.\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupExcusiveMul($0)\")\n") -SLANG_RAW("T WavePrefixProduct(T expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupExcusiveMul($0)\")\n") -SLANG_RAW("vector WavePrefixProduct(vector expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix WavePrefixProduct(matrix expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupExcusiveAdd($0)\")\n") -SLANG_RAW("T WavePrefixSum(T expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupExcusiveAdd($0)\")\n") -SLANG_RAW("vector WavePrefixSum(vector expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix WavePrefixSum(matrix expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupExcusiveAnd($0)\")\n") -SLANG_RAW("T WaveMultiPrefixBitAnd(T expr);\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupExcusiveAnd($0)\")\n") -SLANG_RAW("__generic\n") -SLANG_RAW("vector WaveMultiPrefixBitAnd(vector expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix WaveMultiPrefixBitAnd(matrix expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupExcusiveOr($0)\")\n") -SLANG_RAW("T WaveMultiPrefixBitOr(T expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupExcusiveOr($0)\")\n") -SLANG_RAW("vector WaveMultiPrefixBitOr(vector expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix WaveMultiPrefixBitOr(matrix expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupExcusiveXor($0)\")\n") -SLANG_RAW("T WaveMultiPrefixBitXor(T expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupExcusiveXor($0)\")\n") -SLANG_RAW("vector WaveMultiPrefixBitXor(vector expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix WaveMultiPrefixBitXor(matrix expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_ballot)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupBallotExclusiveBitCount(subgroupBallot($0))\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"__popc(__ballot_sync(__activemask(), $0) & _getLaneLtMask())\")\n") -SLANG_RAW("uint WavePrefixCountBits(bool value);\n") -SLANG_RAW("\n") -SLANG_RAW("uint WaveMultiPrefixCountBits(bool value, uint4 mask);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic T WaveMultiPrefixProduct(T value, uint4 mask);\n") -SLANG_RAW("__generic vector WaveMultiPrefixProduct(vector value, uint4 mask);\n") -SLANG_RAW("__generic matrix WaveMultiPrefixProduct(matrix value, uint4 mask);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic T WaveMultiPrefixSum(T value, uint4 mask);\n") -SLANG_RAW("__generic vector WaveMultiPrefixSum(vector value, uint4 mask);\n") -SLANG_RAW("__generic matrix WaveMultiPrefixSum(matrix value, uint4 mask);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_ballot)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupBroadcastFirst($0)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"_waveReadFirst($0)\")\n") -SLANG_RAW("T WaveReadLaneFirst(T expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_ballot)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupBroadcastFirst($0)\")\n") -SLANG_RAW("vector WaveReadLaneFirst(vector expr);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix WaveReadLaneFirst(matrix expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_ballot)\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupBroadcast($0, $1)\")\n") -SLANG_RAW("__target_intrinsic(cuda, \"__shfl_sync(SLANG_CUDA_WARP_MASK, $0, $1)\")\n") -SLANG_RAW("T WaveReadLaneAt(T value, int lane);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupBroadcast($0, $1)\")\n") -SLANG_RAW("vector WaveReadLaneAt(vector value, int lane);\n") -SLANG_RAW("__generic\n") -SLANG_RAW("matrix WaveReadLaneAt(matrix value, int lane);\n") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("// `typedef`s to help with the fact that HLSL has been sorta-kinda case insensitive at various points\n") -SLANG_RAW("typedef Texture2D texture2D;\n") -SLANG_RAW("\n") - - -// Buffer types - -static const struct { - char const* name; - SlangResourceAccess access; -} kBaseBufferAccessLevels[] = { - { "", SLANG_RESOURCE_ACCESS_READ }, - { "RW", SLANG_RESOURCE_ACCESS_READ_WRITE }, - { "RasterizerOrdered", SLANG_RESOURCE_ACCESS_RASTER_ORDERED }, -}; -static const int kBaseBufferAccessLevelCount = sizeof(kBaseBufferAccessLevels) / sizeof(kBaseBufferAccessLevels[0]); - -for (int aa = 0; aa < kBaseBufferAccessLevelCount; ++aa) -{ - auto access = kBaseBufferAccessLevels[aa].access; - auto flavor = TextureFlavor::create(TextureFlavor::Shape::ShapeBuffer, access).flavor; - sb << "__generic\n"; - sb << "__magic_type(Texture," << int(flavor) << ")\n"; - sb << "__intrinsic_type(" << (kIROp_TextureType + (int(flavor) << kIROpMeta_OtherShift)) << ")\n"; - sb << "struct "; - sb << kBaseBufferAccessLevels[aa].name; - sb << "Buffer {\n"; - - sb << "void GetDimensions(out uint dim);\n"; - - char const* glslLoadFuncName = (access == SLANG_RESOURCE_ACCESS_READ) ? "texelFetch" : "imageLoad"; - - sb << "__glsl_extension(GL_EXT_samplerless_texture_functions)"; - sb << "__target_intrinsic(glsl, \"" << glslLoadFuncName << "($0, $1)$z\")\n"; - sb << "T Load(int location);\n"; - - sb << "T Load(int location, out uint status);\n"; - - sb << "__subscript(uint index) -> T {\n"; - - sb << "__glsl_extension(GL_EXT_samplerless_texture_functions)"; - sb << "__target_intrinsic(glsl, \"" << glslLoadFuncName << "($0, int($1))$z\") get;\n"; - - if (access != SLANG_RESOURCE_ACCESS_READ) - { - sb << "__target_intrinsic(glsl, \"imageStore($0, int($1), $V2)\") set;\n"; - - sb << "__intrinsic_op(" << int(kIROp_ImageSubscript) << ") ref;\n"; - } - - sb << "}\n"; - - sb << "};\n"; -} -SLANG_RAW("#line 2558 \"hlsl.meta.slang\"") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("\n") -SLANG_RAW("// DirectX Raytracing (DXR) Support\n") -SLANG_RAW("//\n") -SLANG_RAW("// The following is based on the experimental DXR SDK v0.09.01.\n") -SLANG_RAW("//\n") -SLANG_RAW("// Numbering follows the sections in the \"D3D12 Raytracing Functional Spec\" v0.09 (2018-03-12)\n") -SLANG_RAW("//\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.1.1 - Ray Flags\n") -SLANG_RAW("\n") -SLANG_RAW("typedef uint RAY_FLAG;\n") -SLANG_RAW("\n") -SLANG_RAW("static const RAY_FLAG RAY_FLAG_NONE = 0x00;\n") -SLANG_RAW("static const RAY_FLAG RAY_FLAG_FORCE_OPAQUE = 0x01;\n") -SLANG_RAW("static const RAY_FLAG RAY_FLAG_FORCE_NON_OPAQUE = 0x02;\n") -SLANG_RAW("static const RAY_FLAG RAY_FLAG_ACCEPT_FIRST_HIT_AND_END_SEARCH = 0x04;\n") -SLANG_RAW("static const RAY_FLAG RAY_FLAG_SKIP_CLOSEST_HIT_SHADER = 0x08;\n") -SLANG_RAW("static const RAY_FLAG RAY_FLAG_CULL_BACK_FACING_TRIANGLES = 0x10;\n") -SLANG_RAW("static const RAY_FLAG RAY_FLAG_CULL_FRONT_FACING_TRIANGLES = 0x20;\n") -SLANG_RAW("static const RAY_FLAG RAY_FLAG_CULL_OPAQUE = 0x40;\n") -SLANG_RAW("static const RAY_FLAG RAY_FLAG_CULL_NON_OPAQUE = 0x80;\n") -SLANG_RAW("static const RAY_FLAG RAY_FLAG_SKIP_TRIANGLES = 0x100;\n") -SLANG_RAW("static const RAY_FLAG RAY_FLAG_SKIP_PROCEDURAL_PRIMITIVES = 0x200;\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.1.2 - Ray Description Structure\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(hlsl, RayDesc)\n") -SLANG_RAW("struct RayDesc\n") -SLANG_RAW("{\n") -SLANG_RAW(" __target_intrinsic(hlsl, Origin)\n") -SLANG_RAW(" float3 Origin;\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(hlsl, TMin)\n") -SLANG_RAW(" float TMin;\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(hlsl, Direction)\n") -SLANG_RAW(" float3 Direction;\n") -SLANG_RAW("\n") -SLANG_RAW(" __target_intrinsic(hlsl, TMax)\n") -SLANG_RAW(" float TMax;\n") -SLANG_RAW("};\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.1.3 - Ray Acceleration Structure\n") -SLANG_RAW("\n") -SLANG_RAW("__builtin\n") -SLANG_RAW("__magic_type(RaytracingAccelerationStructureType)\n") -SLANG_RAW("__intrinsic_type(") -SLANG_SPLICE(kIROp_RaytracingAccelerationStructureType -) -SLANG_RAW(")\n") -SLANG_RAW("struct RaytracingAccelerationStructure {};\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.1.4 - Subobject Definitions\n") -SLANG_RAW("\n") -SLANG_RAW("// TODO: We may decide to support these, but their reliance on C++ implicit\n") -SLANG_RAW("// constructor call syntax (`SomeType someVar(arg0, arg1);`) makes them\n") -SLANG_RAW("// annoying for the current Slang parsing strategy, and using global variables\n") -SLANG_RAW("// for this stuff comes across as a kludge rather than the best possible design.\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.1.5 - Intersection Attributes Structure\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(hlsl, BuiltInTriangleIntersectionAttributes)\n") -SLANG_RAW("struct BuiltInTriangleIntersectionAttributes\n") -SLANG_RAW("{\n") -SLANG_RAW(" __target_intrinsic(hlsl, barycentrics)\n") -SLANG_RAW(" float2 barycentrics;\n") -SLANG_RAW("};\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.2 Shaders\n") -SLANG_RAW("\n") -SLANG_RAW("// Right now new shader stages need to be added directly to the compiler\n") -SLANG_RAW("// implementation, rather than being something that can be declared in the stdlib.\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.3 - Intrinsics\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.3.1\n") -SLANG_RAW("\n") -SLANG_RAW("void CallShader(uint shaderIndex, inout Payload payload);\n") -SLANG_RAW("\n") -SLANG_RAW("// `executeCallableNV` is the GLSL intrinsic that will be used to implement\n") -SLANG_RAW("// `CallShader()` for GLSL-based targets.\n") -SLANG_RAW("//\n") -SLANG_RAW("__target_intrinsic(glsl, \"executeCallableNV\")\n") -SLANG_RAW("void __executeCallableNV(uint shaderIndex, int payloadLocation);\n") -SLANG_RAW("\n") -SLANG_RAW("// Next is the custom intrinsic that will compute the payload location\n") -SLANG_RAW("// for a type being used in a `CallShader()` call for GLSL-based targets.\n") -SLANG_RAW("//\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, \"$XC\")\n") -SLANG_RAW("[__readNone]\n") -SLANG_RAW("int __callablePayloadLocation(Payload payload);\n") -SLANG_RAW("\n") -SLANG_RAW("// Now we provide a hard-coded definition of `CallShader()` for GLSL-based\n") -SLANG_RAW("// targets, which maps the generic HLSL operation into the non-generic\n") -SLANG_RAW("// GLSL equivalent.\n") -SLANG_RAW("//\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__specialized_for_target(glsl)\n") -SLANG_RAW("void CallShader(uint shaderIndex, inout Payload payload)\n") -SLANG_RAW("{\n") -SLANG_RAW(" [__vulkanRayPayload]\n") -SLANG_RAW(" static Payload p;\n") -SLANG_RAW("\n") -SLANG_RAW(" p = payload;\n") -SLANG_RAW(" __executeCallableNV(shaderIndex, __callablePayloadLocation(p));\n") -SLANG_RAW(" payload = p;\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.3.2\n") -SLANG_RAW("void TraceRay(\n") -SLANG_RAW(" RaytracingAccelerationStructure AccelerationStructure,\n") -SLANG_RAW(" uint RayFlags,\n") -SLANG_RAW(" uint InstanceInclusionMask,\n") -SLANG_RAW(" uint RayContributionToHitGroupIndex,\n") -SLANG_RAW(" uint MultiplierForGeometryContributionToHitGroupIndex,\n") -SLANG_RAW(" uint MissShaderIndex,\n") -SLANG_RAW(" RayDesc Ray,\n") -SLANG_RAW(" inout payload_t Payload);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"traceNV\")\n") -SLANG_RAW("void __traceNV(\n") -SLANG_RAW(" RaytracingAccelerationStructure AccelerationStructure,\n") -SLANG_RAW(" uint RayFlags,\n") -SLANG_RAW(" uint InstanceInclusionMask,\n") -SLANG_RAW(" uint RayContributionToHitGroupIndex,\n") -SLANG_RAW(" uint MultiplierForGeometryContributionToHitGroupIndex,\n") -SLANG_RAW(" uint MissShaderIndex,\n") -SLANG_RAW(" float3 Origin,\n") -SLANG_RAW(" float TMin,\n") -SLANG_RAW(" float3 Direction,\n") -SLANG_RAW(" float TMax,\n") -SLANG_RAW(" int PayloadLocation);\n") -SLANG_RAW("\n") -SLANG_RAW("// TODO: Slang's parsing logic currently puts modifiers on\n") -SLANG_RAW("// the `GenericDecl` rather than the inner decl when\n") -SLANG_RAW("// using our default syntax, which seems wrong. We need\n") -SLANG_RAW("// to fix this, but for now using the expanded `__generic`\n") -SLANG_RAW("// syntax works in a pinch.\n") -SLANG_RAW("//\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__target_intrinsic(glsl, \"$XP\")\n") -SLANG_RAW("[__readNone]\n") -SLANG_RAW("int __rayPayloadLocation(Payload payload);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__specialized_for_target(glsl)\n") -SLANG_RAW("void TraceRay(\n") -SLANG_RAW(" RaytracingAccelerationStructure AccelerationStructure,\n") -SLANG_RAW(" uint RayFlags,\n") -SLANG_RAW(" uint InstanceInclusionMask,\n") -SLANG_RAW(" uint RayContributionToHitGroupIndex,\n") -SLANG_RAW(" uint MultiplierForGeometryContributionToHitGroupIndex,\n") -SLANG_RAW(" uint MissShaderIndex,\n") -SLANG_RAW(" RayDesc Ray,\n") -SLANG_RAW(" inout payload_t Payload)\n") -SLANG_RAW("{\n") -SLANG_RAW(" [__vulkanRayPayload]\n") -SLANG_RAW(" static payload_t p;\n") -SLANG_RAW("\n") -SLANG_RAW(" p = Payload;\n") -SLANG_RAW(" __traceNV(\n") -SLANG_RAW(" AccelerationStructure,\n") -SLANG_RAW(" RayFlags,\n") -SLANG_RAW(" InstanceInclusionMask,\n") -SLANG_RAW(" RayContributionToHitGroupIndex,\n") -SLANG_RAW(" MultiplierForGeometryContributionToHitGroupIndex,\n") -SLANG_RAW(" MissShaderIndex,\n") -SLANG_RAW(" Ray.Origin,\n") -SLANG_RAW(" Ray.TMin,\n") -SLANG_RAW(" Ray.Direction,\n") -SLANG_RAW(" Ray.TMax,\n") -SLANG_RAW(" __rayPayloadLocation(p));\n") -SLANG_RAW(" Payload = p;\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.3.3\n") -SLANG_RAW("bool ReportHit(float tHit, uint hitKind, A attributes);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"reportIntersectionNV\")\n") -SLANG_RAW("bool __reportIntersectionNV(float tHit, uint hitKind);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic\n") -SLANG_RAW("__specialized_for_target(glsl)\n") -SLANG_RAW("bool ReportHit(float tHit, uint hitKind, A attributes)\n") -SLANG_RAW("{\n") -SLANG_RAW(" [__vulkanHitAttributes]\n") -SLANG_RAW(" static A a;\n") -SLANG_RAW("\n") -SLANG_RAW(" a = attributes;\n") -SLANG_RAW(" return __reportIntersectionNV(tHit, hitKind);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.3.4\n") -SLANG_RAW("__target_intrinsic(glsl, ignoreIntersectionNV)\n") -SLANG_RAW("void IgnoreHit();\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.3.5\n") -SLANG_RAW("__target_intrinsic(glsl, terminateRayNV)\n") -SLANG_RAW("void AcceptHitAndEndSearch();\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.4 - System Values and Special Semantics\n") -SLANG_RAW("\n") -SLANG_RAW("// TODO: Many of these functions need to be restricted so that\n") -SLANG_RAW("// they can only be accessed from specific stages.\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.4.1 - Ray Dispatch System Values\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"(gl_LaunchIDNV)\")\n") -SLANG_RAW("uint3 DispatchRaysIndex();\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"(gl_LaunchSizeNV)\")\n") -SLANG_RAW("uint3 DispatchRaysDimensions();\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.4.2 - Ray System Values\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"(gl_WorldRayOriginNV)\")\n") -SLANG_RAW("float3 WorldRayOrigin();\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"(gl_WorldRayDirectionNV)\")\n") -SLANG_RAW("float3 WorldRayDirection();\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"(gl_RayTminNV)\")\n") -SLANG_RAW("float RayTMin();\n") -SLANG_RAW("\n") -SLANG_RAW("// Note: The `RayTCurrent()` intrinsic should translate to\n") -SLANG_RAW("// either `gl_HitTNV` (for hit shaders) or `gl_RayTmaxNV`\n") -SLANG_RAW("// (for intersection shaders). Right now we are handling this\n") -SLANG_RAW("// during code emission, for simplicity.\n") -SLANG_RAW("//\n") -SLANG_RAW("// TODO: Once the compiler supports a more refined concept\n") -SLANG_RAW("// of profiles/capabilities and overloading based on them,\n") -SLANG_RAW("// we should simply provide two overloads here, specialized\n") -SLANG_RAW("// to the appropriate Vulkan stages.\n") -SLANG_RAW("//\n") -SLANG_RAW("__target_intrinsic(glsl, \"$XT\")\n") -SLANG_RAW("float RayTCurrent();\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"(gl_IncomingRayFlagsNV)\")\n") -SLANG_RAW("uint RayFlags();\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.4.3 - Primitive/Object Space System Values\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"(gl_InstanceCustomIndexNV)\")\n") -SLANG_RAW("uint InstanceIndex();\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"(gl_InstanceID)\")\n") -SLANG_RAW("uint InstanceID();\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"(gl_PrimitiveID)\")\n") -SLANG_RAW("uint PrimitiveIndex();\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"(gl_ObjectRayOriginNV)\")\n") -SLANG_RAW("float3 ObjectRayOrigin();\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"(gl_ObjectRayDirectionNV)\")\n") -SLANG_RAW("float3 ObjectRayDirection();\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"transpose(gl_ObjectToWorldNV)\")\n") -SLANG_RAW("float3x4 ObjectToWorld3x4();\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"transpose(gl_WorldToObjectNV)\")\n") -SLANG_RAW("float3x4 WorldToObject3x4();\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"(gl_ObjectToWorldNV)\")\n") -SLANG_RAW("float4x3 ObjectToWorld4x3();\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, \"(gl_WorldToObjectNV)\")\n") -SLANG_RAW("float4x3 WorldToObject4x3();\n") -SLANG_RAW("\n") -SLANG_RAW("// Note: The provisional DXR spec included these unadorned\n") -SLANG_RAW("// `ObjectToWorld()` and `WorldToObject()` functions, so\n") -SLANG_RAW("// we will forward them to the new names as a convience\n") -SLANG_RAW("// for users who are porting their code.\n") -SLANG_RAW("//\n") -SLANG_RAW("// TODO: Should we provide a deprecation warning on these\n") -SLANG_RAW("// declarations, so that users can know they aren't coding\n") -SLANG_RAW("// against the final spec?\n") -SLANG_RAW("//\n") -SLANG_RAW("float3x4 ObjectToWorld() { return ObjectToWorld3x4(); }\n") -SLANG_RAW("float3x4 WorldToObject() { return WorldToObject3x4(); }\n") -SLANG_RAW("\n") -SLANG_RAW("// 10.4.4 - Hit Specific System values\n") -SLANG_RAW("__target_intrinsic(glsl, \"(gl_HitKindNV)\")\n") -SLANG_RAW("uint HitKind();\n") -SLANG_RAW("\n") -SLANG_RAW("// Pre-defined hit kinds (not documented explicitly)\n") -SLANG_RAW("static const uint HIT_KIND_TRIANGLE_FRONT_FACE = 254;\n") -SLANG_RAW("static const uint HIT_KIND_TRIANGLE_BACK_FACE = 255;\n") -SLANG_RAW("\n") -SLANG_RAW("//\n") -SLANG_RAW("// Shader Model 6.4\n") -SLANG_RAW("//\n") -SLANG_RAW("\n") -SLANG_RAW("// Treats `left` and `right` as 4-component vectors of `UInt8` and computes `dot(left, right) + acc`\n") -SLANG_RAW("uint dot4add_u8packed(uint left, uint right, uint acc);\n") -SLANG_RAW("\n") -SLANG_RAW("// Treats `left` and `right` as 4-component vectors of `Int8` and computes `dot(left, right) + acc`\n") -SLANG_RAW("int dot4add_i8packed(uint left, uint right, int acc);\n") -SLANG_RAW("\n") -SLANG_RAW("// Computes `dot(left, right) + acc`.\n") -SLANG_RAW("//\n") -SLANG_RAW("// May not produce infinities or NaNs for intermediate results that overflow the range of `half`\n") -SLANG_RAW("float dot2add(float2 left, float2 right, float acc);\n") -SLANG_RAW("\n") -SLANG_RAW("//\n") -SLANG_RAW("// Shader Model 6.5\n") -SLANG_RAW("//\n") -SLANG_RAW("\n") -SLANG_RAW("//\n") -SLANG_RAW("// Mesh Shaders\n") -SLANG_RAW("//\n") -SLANG_RAW("\n") -SLANG_RAW("// Set the number of output vertices and primitives for a mesh shader invocation.\n") -SLANG_RAW("void SetMeshOutputCounts(uint vertexCount, uint primitiveCount);\n") -SLANG_RAW("\n") -SLANG_RAW("// Specify the number of downstream mesh shader thread groups to invoke from an amplification shader,\n") -SLANG_RAW("// and provide the values for per-mesh payload parameters.\n") -SLANG_RAW("//\n") -SLANG_RAW("void DispatchMesh

(uint threadGroupCountX, uint threadGroupCountY, uint threadGroupCountZ, P meshPayload);\n") -SLANG_RAW("\n") -SLANG_RAW("//\n") -SLANG_RAW("// TODO: \"Sampler feedback\" types `FeedbackTexture2D` and `FeedbackTexture2DArray`.\n") -SLANG_RAW("//\n") -SLANG_RAW("\n") -SLANG_RAW("//\n") -SLANG_RAW("// DXR 1.1 and `TraceRayInline` support\n") -SLANG_RAW("//\n") -SLANG_RAW("\n") -SLANG_RAW("// Get the index of the geometry that was hit in an intersection, any-hit, or closest-hit shader\n") -SLANG_RAW("uint GeometryIndex();\n") -SLANG_RAW("\n") -SLANG_RAW("// Status of whether a (closest) hit has been committed in a `RayQuery`.\n") -SLANG_RAW("typedef uint COMMITTED_STATUS;\n") -SLANG_RAW("\n") -SLANG_RAW("// No hit committed.\n") -SLANG_RAW("static const COMMITTED_STATUS COMMITTED_NOTHING = 0;\n") -SLANG_RAW("\n") -SLANG_RAW("// Closest hit is a triangle.\n") -SLANG_RAW("//\n") -SLANG_RAW("// This could be an opaque triangle hit found by the fixed-function\n") -SLANG_RAW("// traversal and intersection implementation, or a non-opaque\n") -SLANG_RAW("// triangle hit committed by user code with `RayQuery.CommitNonOpaqueTriangleHit`\n") -SLANG_RAW("//\n") -SLANG_RAW("static const COMMITTED_STATUS COMMITTED_TRIANGLE_HIT = 1;\n") -SLANG_RAW("\n") -SLANG_RAW("// Closest hit is a procedural primitive.\n") -SLANG_RAW("//\n") -SLANG_RAW("// A procedural hit primitive is committed using `RayQuery.CommitProceduralPrimitiveHit`.\n") -SLANG_RAW("static const COMMITTED_STATUS COMMITTED_PROCEDURAL_PRIMITIVE_HIT = 2;\n") -SLANG_RAW("\n") -SLANG_RAW("// Type of candidate hit that a `RayQuery` is pausing at.\n") -SLANG_RAW("//\n") -SLANG_RAW("// A `RayQuery` can automatically commit hits with opaque triangles,\n") -SLANG_RAW("// but yields to user code for other hits to allow them to be\n") -SLANG_RAW("// dismissed or committed.\n") -SLANG_RAW("//\n") -SLANG_RAW("typedef uint CANDIDATE_TYPE;\n") -SLANG_RAW("\n") -SLANG_RAW("// Candidate hit is a non-opaque triangle.\n") -SLANG_RAW("static const CANDIDATE_TYPE CANDIDATE_NON_OPAQUE_TRIANGLE = 0;\n") -SLANG_RAW("\n") -SLANG_RAW("// Candidate hit is a procedural primitive.\n") -SLANG_RAW("static const CANDIDATE_TYPE CANDIDATE_PROCEDURAL_PRIMITIVE = 1;\n") -SLANG_RAW("\n") -SLANG_RAW("// Handle to state of an in-progress ray-tracing query.\n") -SLANG_RAW("//\n") -SLANG_RAW("// The ray query is effectively a coroutine that user shader\n") -SLANG_RAW("// code can resume to continue tracing the ray, and which yields\n") -SLANG_RAW("// back to the user code at interesting events along the ray.\n") -SLANG_RAW("//\n") -SLANG_RAW("__target_intrinsic(hlsl, RayQuery)\n") -SLANG_RAW("struct RayQuery \n") -SLANG_RAW("{\n") -SLANG_RAW(" // Initialize a ray-tracing query.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" // This method may be called on a \"fresh\" ray query, or\n") -SLANG_RAW(" // on one that is already tracing a ray. In the latter\n") -SLANG_RAW(" // case any state related to the ray previously being\n") -SLANG_RAW(" // traced is overwritten.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" // The `rayFlags` here will be bitwise ORed with\n") -SLANG_RAW(" // the `rayFlags` passed as a generic argument to\n") -SLANG_RAW(" // `RayQuery` to get the effective ray flags, which\n") -SLANG_RAW(" // must obey any API-imposed restrictions.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" void TraceRayInline(RaytracingAccelerationStructure accelerationStructure, RAY_FLAG rayFlags, uint instanceInclusionMask, RayDesc ray);\n") -SLANG_RAW("\n") -SLANG_RAW(" // Resume the ray query coroutine.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" // If the coroutine suspends because of encountering\n") -SLANG_RAW(" // a candidate hit that cannot be resolved with fixed-funciton\n") -SLANG_RAW(" // logic, this function returns `true`, and the `Candidate*()`\n") -SLANG_RAW(" // functions should be used by application code to resolve\n") -SLANG_RAW(" // the candidate hit (by either committing or ignoring it).\n") -SLANG_RAW(" //\n") -SLANG_RAW(" // If the coroutine terminates because traversal is\n") -SLANG_RAW(" // complete (or has been aborted), this function returns\n") -SLANG_RAW(" // `false`, and application code should use the `Committed*()`\n") -SLANG_RAW(" // functions to appropriately handle the closest hit (it any)\n") -SLANG_RAW(" // that was found.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" bool Proceed();\n") -SLANG_RAW("\n") -SLANG_RAW(" // Causes the ray query to terminate.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" // This function cases the ray query to act as if\n") -SLANG_RAW(" // traversal has terminated, so that subsequent\n") -SLANG_RAW(" // `Proceed()` calls will return `false`.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" void Abort();\n") -SLANG_RAW("\n") -SLANG_RAW(" // Get the type of candidate hit being considered.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" // The ray query coroutine will suspend when it encounters\n") -SLANG_RAW(" // a hit that cannot be resolved with fixed-function logic\n") -SLANG_RAW(" // (either a non-opaque triangle or a procedural primitive).\n") -SLANG_RAW(" // In either of those cases, `CandidateType()` will return\n") -SLANG_RAW(" // the kind of candidate hit that must be resolved by\n") -SLANG_RAW(" // user code.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" CANDIDATE_TYPE CandidateType();\n") -SLANG_RAW("\n") -SLANG_RAW(" // Access properties of a candidate hit.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" float3x4 CandidateObjectToWorld3x4();\n") -SLANG_RAW(" float4x3 CandidateObjectToWorld4x3();\n") -SLANG_RAW(" float3x4 CandidateWorldToObject3x4();\n") -SLANG_RAW(" float4x3 CandidateWorldToObject4x3();\n") -SLANG_RAW(" uint CandidateInstanceIndex();\n") -SLANG_RAW(" uint CandidateInstanceID();\n") -SLANG_RAW(" uint CandidateGeometryIndex();\n") -SLANG_RAW(" uint CandidatePrimitiveIndex();\n") -SLANG_RAW(" uint CandidateInstanceContributionToHitGroupIndex();\n") -SLANG_RAW("\n") -SLANG_RAW(" // Access properties of the ray being traced\n") -SLANG_RAW(" // in the object space of a candidate hit.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" float3 CandidateObjectRayOrigin();\n") -SLANG_RAW(" float3 CandidateObjectRayDirection();\n") -SLANG_RAW("\n") -SLANG_RAW(" // Access properties of a candidate procedural primitive hit.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" bool CandidateProceduralPrimitiveNonOpaque();\n") -SLANG_RAW("\n") -SLANG_RAW(" // Access properties of a candidate no-opaque triangle hit.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" bool CandidateTriangleFrontFace();\n") -SLANG_RAW(" float2 CandidateTriangleBarycentrics();\n") -SLANG_RAW(" float CandidateTriangleRayT();\n") -SLANG_RAW("\n") -SLANG_RAW(" // Commit the current non-opaque triangle hit.\n") -SLANG_RAW(" void CommitNonOpaqueTriangleHit();\n") -SLANG_RAW("\n") -SLANG_RAW(" // Commit the current procedural primitive hit, with hit time `t`.\n") -SLANG_RAW(" void CommitProceduralPrimitiveHit(float t);\n") -SLANG_RAW("\n") -SLANG_RAW(" // Get the status of the committed (closest) hit, if any.\n") -SLANG_RAW(" COMMITTED_STATUS CommittedStatus();\n") -SLANG_RAW("\n") -SLANG_RAW(" // Access properties of the committed hit.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" float3x4 CommittedObjectToWorld3x4();\n") -SLANG_RAW(" float4x3 CommittedObjectToWorld4x3();\n") -SLANG_RAW(" float3x4 CommittedWorldToObject3x4();\n") -SLANG_RAW(" float4x3 CommittedWorldToObject4x3();\n") -SLANG_RAW(" float CommittedRayT();\n") -SLANG_RAW(" uint CommittedInstanceIndex();\n") -SLANG_RAW(" uint CommittedInstanceID();\n") -SLANG_RAW(" uint CommittedGeometryIndex();\n") -SLANG_RAW(" uint CommittedPrimitiveIndex();\n") -SLANG_RAW(" uint CommittedInstanceContributionToHitGroupIndex();\n") -SLANG_RAW("\n") -SLANG_RAW(" // Access properties of the ray being traced\n") -SLANG_RAW(" // in the object space of a committed hit.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" float3 CommittedObjectRayOrigin();\n") -SLANG_RAW(" float3 CommittedObjectRayDirection();\n") -SLANG_RAW("\n") -SLANG_RAW(" // Access properties of a committed triangle hit.\n") -SLANG_RAW(" //\n") -SLANG_RAW(" bool CommittedTriangleFrontFace();\n") -SLANG_RAW(" float2 CommittedTriangleBarycentrics();\n") -SLANG_RAW("\n") -SLANG_RAW(" // Access properties of the ray being traced.\n") -SLANG_RAW(" uint RayFlags();\n") -SLANG_RAW(" float3 WorldRayOrigin();\n") -SLANG_RAW(" float3 WorldRayDirection();\n") -SLANG_RAW(" float RayTMin();\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("//\n") -SLANG_RAW("// Vulkan/SPIR-V specific features\n") -SLANG_RAW("//\n") -SLANG_RAW("\n") -SLANG_RAW("struct VkSubpassInput\n") -SLANG_RAW("{\n") -SLANG_RAW(" T SubpassLoad();\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("struct VkSubpassInputMS\n") -SLANG_RAW("{\n") -SLANG_RAW(" T SubpassLoad(int sampleIndex);\n") -SLANG_RAW("}\n") -- cgit v1.2.3