diff options
| author | jsmall-nvidia <jsmall@nvidia.com> | 2020-03-06 17:15:58 -0500 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2020-03-06 17:15:58 -0500 |
| commit | cdee13466080b737ed18c148e36af75898285ed6 (patch) | |
| tree | 55fba3a3236732907537c93929baf1e0cf298605 /source | |
| parent | b94a12b91086ea004d9b78fa8a14fd4726af9e76 (diff) | |
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.
Diffstat (limited to 'source')
| -rw-r--r-- | source/slang/core.meta.slang.h | 1864 | ||||
| -rw-r--r-- | source/slang/glsl.meta.slang.h | 202 | ||||
| -rw-r--r-- | source/slang/hlsl.meta.slang.h | 3140 |
3 files changed, 0 insertions, 5206 deletions
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<T:__EnumType> 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<T,U> __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<T> __intrinsic_op(select) T operator?:(bool condition, T ifTrue, T ifFalse);\n") -SLANG_RAW("__generic<T, let N : int> __intrinsic_op(select) vector<T,N> operator?:(vector<bool,N> condition, vector<T,N> ifTrue, vector<T,N> 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<T:__EnumType>\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<T>\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<T>\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<T>\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<T>\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<T = float, let N : int = 4>\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<T,N> 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<T = float, let R : int = 4, let C : int = 4>\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<T>\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<T>\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<T>\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<T> __extension vector<T, " << N << ">\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<T," << M << "> " << 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<T," << M << "> " << kVectorNames[M]; - sb << ", vector<T," << K << "> "; - 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<let N : int> __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<T> __extension matrix<T, " << R << "," << C << ">\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<T," << C << "> 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<U> __init(matrix<U," << R << ", " << C << ">);\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<T," << rr << "," << cc << "> 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<T = float4> "; - - 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> t, "; - sb << "SamplerState s);\n"; - sb << "};\n"; - - sb << "__specialized_for_target(glsl)\n"; - sb << "T texture<T>(Sampler"; - sb << kBaseTextureAccessLevels[accessLevel].name; - sb << name; - if (isMultisample) sb << "MS"; - if (isArray) sb << "Array"; - sb << "<T> 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<T = float4> "; - - 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<T.Scalar, 4>`. - // - static const struct { - char const* genericPrefix; - char const* elementType; - char const* outputType; - } kGatherExtensionCases[] = { - { "__generic<T, let N : int>", "vector<T,N>", "vector<T, 4>" }, - { "", "float", "vector<float, 4>" }, - { "", "int" , "vector<int, 4>"}, - { "", "uint", "vector<uint, 4>"}, - - // 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<float3>` 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<let N : int> "; - 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<let N : int, let M : int> "; - 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<T : " << op.interface << ">\n"; - sb << "__intrinsic_op(" << int(op.opCode) << ") " << resultType << " operator" << op.opName << "(" << qual << "T value);\n"; - - // vector version - sb << "__generic<T : " << op.interface << ", let N : int> "; - sb << fixity; - sb << "__intrinsic_op(" << int(op.opCode) << ") vector<" << resultType << ",N> operator" << op.opName << "(" << qual << "vector<T,N> value);\n"; - - // matrix version - sb << "__generic<T : " << op.interface << ", let N : int, let M : int> "; - sb << fixity; - sb << "__intrinsic_op(" << int(op.opCode) << ") matrix<" << resultType << ",N,M> operator" << op.opName << "(" << qual << "matrix<T,N,M> 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<let N : int> "; - 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<let N : int, let M : int> "; - 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<let N : int> "; - sb << "__intrinsic_op(" << int(op.opCode) << ") vector<" << resultType << ",N> operator" << op.opName << "(" << leftQual << leftType << " left, vector<" << rightType << ",N> right);\n"; - - sb << "__generic<let N : int, let M : int> "; - 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<let N : int> "; - sb << "__intrinsic_op(" << int(op.opCode) << ") vector<" << resultType << ",N> operator" << op.opName << "(" << leftQual << "vector<" << leftType << ",N> left, " << rightType << " right);\n"; - - sb << "__generic<let N : int, let M : int> "; - 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<T : " << op.interface << ">\n"; - sb << "__intrinsic_op(" << int(op.opCode) << ") " << resultType << " operator" << op.opName << "(" << leftQual << leftType << " left, " << rightType << " right);\n"; - - // vector version - sb << "__generic<T : " << op.interface << ", let N : int> "; - 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<T : " << op.interface << ", let N : int, let M : int> "; - 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<T : " << op.interface << ", let N : int> "; - sb << "__intrinsic_op(" << int(op.opCode) << ") vector<" << resultType << ",N> operator" << op.opName << "(" << leftQual << leftType << " left, vector<" << rightType << ",N> right);\n"; - - sb << "__generic<T : " << op.interface << ", let N : int, let M : int> "; - 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<T : " << op.interface << ", let N : int> "; - sb << "__intrinsic_op(" << int(op.opCode) << ") vector<" << resultType << ",N> operator" << op.opName << "(" << leftQual << "vector<" << leftType << ",N> left, " << rightType << " right);\n"; - - sb << "__generic<T : " << op.interface << ", let N : int, let M : int> "; - 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<E : __EnumType>\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<T : __BuiltinArithmeticType, let N : int> __intrinsic_op(mul) vector<T,N> operator*(vector<T,N> x, T y);\n"; -sb << "__generic<T : __BuiltinArithmeticType, let N : int> __intrinsic_op(mul) vector<T,N> operator*(T x, vector<T,N> y);\n"; - -// scalar-matrix and matrix-scalar -sb << "__generic<T : __BuiltinArithmeticType, let N : int, let M :int> __intrinsic_op(mul) matrix<T,N,M> operator*(matrix<T,N,M> x, T y);\n"; -sb << "__generic<T : __BuiltinArithmeticType, let N : int, let M :int> __intrinsic_op(mul) matrix<T,N,M> operator*(T x, matrix<T,N,M> y);\n"; - -// vector-vector (dot product) -sb << "__generic<T : __BuiltinArithmeticType, let N : int> __intrinsic_op(dot) T operator*(vector<T,N> x, vector<T,N> y);\n"; - -// vector-matrix -sb << "__generic<T : __BuiltinArithmeticType, let N : int, let M : int> __intrinsic_op(mul) vector<T,M> operator*(vector<T,N> x, matrix<T,N,M> y);\n"; - -// matrix-vector -sb << "__generic<T : __BuiltinArithmeticType, let N : int, let M : int> __intrinsic_op(mul) vector<T,N> operator*(matrix<T,N,M> x, vector<T,M> y);\n"; - -// matrix-matrix -sb << "__generic<T : __BuiltinArithmeticType, let R : int, let N : int, let C : int> __intrinsic_op(mul) matrix<T,R,C> operator*(matrix<T,R,N> x, matrix<T,N,C> 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<T> "; - sb << "__magic_type(TextureSampler," << int(readFlavor) << ") struct "; - sb << "__sampler" << name; - sb << " {};\n"; - - sb << "__generic<T> "; - sb << "__magic_type(Texture," << int(readFlavor) << ") struct "; - sb << "__texture" << name; - sb << " {};\n"; - - sb << "__generic<T> "; - 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<T> __magic_type(GLSLInputParameterGroupType) struct __GLSLInputParameterGroup {};\n"; -sb << "__generic<T> __magic_type(GLSLOutputParameterGroupType) struct __GLSLOutputParameterGroup {};\n"; -sb << "__generic<T> __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<T>\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<T>\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<T>\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<T, let N : int>\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<T, let N : int>\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<T>\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<T>\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<T>\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<T>\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<TYPE,COUNT> 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<TYPE,ROWS,COLS> 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<TYPE,COUNT> 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<TYPE,ROWS,COLS> 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<TYPE,COUNT> 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<TYPE,ROWS,COLS> 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<T : __BuiltinSignedArithmeticType>\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<T : __BuiltinSignedArithmeticType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> abs(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, abs, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinSignedArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T,N,M> abs(matrix<T,N,M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T acos(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> acos(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, acos, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> acos(matrix<T, N, M> 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<T : __BuiltinType>\n") -SLANG_RAW("__target_intrinsic(glsl, \"bool($0)\")\n") -SLANG_RAW("bool all(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"all(bvec$N0($0))\")\n") -SLANG_RAW("bool all(vector<T,N> 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<T : __BuiltinType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("bool all(matrix<T,N,M> 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<T : __BuiltinType>\n") -SLANG_RAW("__target_intrinsic(glsl, \"bool($0)\")\n") -SLANG_RAW("bool any(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"any(bvec$N0($0))\")\n") -SLANG_RAW("bool any(vector<T, N> 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<T : __BuiltinType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("bool any(matrix<T, N, M> 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<let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"intBitsToFloat\")\n") -SLANG_RAW("vector<float, N> 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<let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"uintBitsToFloat\")\n") -SLANG_RAW("vector<float,N> asfloat(vector<uint,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<let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<float,N,M> 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<let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<float,N,M> asfloat(matrix<uint,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("// 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<let N : int>\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Pos -) -SLANG_RAW(")\n") -SLANG_RAW("vector<float,N> asfloat(vector<float,N> x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<let N : int, let M : int>\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Pos -) -SLANG_RAW(")\n") -SLANG_RAW("matrix<float,N,M> asfloat(matrix<float,N,M> x);\n") -SLANG_RAW("\n") -SLANG_RAW("// Inverse sine (HLSL SM 1.0)\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> asin(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T,N,asin,x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> asin(matrix<T, N, M> 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<let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"floatBitsToInt\")\n") -SLANG_RAW("vector<int, N> asint(vector<float, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(int, N, asint, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"ivec$N0($0)\")\n") -SLANG_RAW("vector<int, N> asint(vector<uint, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(int, N, asint, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<int, N, M> asint(matrix<float, N, M> 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<let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<int, N, M> asint(matrix<uint, N, M> 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<let N : int>\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Pos -) -SLANG_RAW(")\n") -SLANG_RAW("vector<int,N> asint(vector<int,N> x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<let N : int, let M : int>\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Pos -) -SLANG_RAW(")\n") -SLANG_RAW("matrix<int,N,M> asint(matrix<int,N,M> 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<let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"floatBitsToUint\")\n") -SLANG_RAW("vector<uint,N> asuint(vector<float,N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(uint, N, asuint, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"uvec$N0($0)\")\n") -SLANG_RAW("vector<uint, N> asuint(vector<int, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(uint, N, asuint, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<uint,N,M> asuint(matrix<float,N,M> 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<let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<uint, N, M> asuint(matrix<int, N, M> 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<let N : int>\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Pos -) -SLANG_RAW(")\n") -SLANG_RAW("vector<uint,N> asuint(vector<uint,N> x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<let N : int, let M : int>\n") -SLANG_RAW("__intrinsic_op(") -SLANG_SPLICE(kCompoundIntrinsicOp_Pos -) -SLANG_RAW(")\n") -SLANG_RAW("matrix<uint,N,M> asuint(matrix<uint,N,M> x);\n") -SLANG_RAW("\n") -SLANG_RAW("// Inverse tangent (HLSL SM 1.0)\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T atan(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> atan(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, atan, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> atan(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl,\"atan($0,$1)\")\n") -SLANG_RAW("vector<T, N> atan2(vector<T, N> y, vector<T, N> 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<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T,N,M> atan2(matrix<T,N,M> y, matrix<T,N,M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T ceil(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> ceil(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, ceil, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> ceil(matrix<T, N, M> 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<T : __BuiltinArithmeticType>\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<T : __BuiltinArithmeticType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> clamp(vector<T, N> x, vector<T, N> minBound, vector<T, N> maxBound)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return min(max(x, minBound), maxBound);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T,N,M> clamp(matrix<T,N,M> x, matrix<T,N,M> minBound, matrix<T,N,M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("void clip(vector<T,N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" if(any(x < T(0))) discard;\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("void clip(matrix<T,N,M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T cos(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> cos(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T,N, cos, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> cos(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T cosh(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T,N> cosh(vector<T,N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T,N, cosh, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> cosh(matrix<T, N, M> 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<T : __BuiltinArithmeticType>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T,3> cross(vector<T,3> left, vector<T,3> right)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return vector<T,3>(\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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("__target_intrinsic(glsl, dFdx)\n") -SLANG_RAW("T ddx(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, dFdx)\n") -SLANG_RAW("vector<T, N> ddx(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, ddx, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> ddx(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\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<T, N> ddx_coarse(vector<T, N> 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<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> ddx_coarse(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\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<T, N> ddx_fine(vector<T, N> 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<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> ddx_fine(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, dFdy)\n") -SLANG_RAW("vector<T, N> ddy(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, ddy, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> ddy(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\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<T, N> ddy_coarse(vector<T, N> 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<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> ddy_coarse(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\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<T, N> ddy_fine(vector<T, N> 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<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> ddy_fine(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> degrees(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, degrees, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> degrees(matrix<T, N, M> 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 : __BuiltinFloatingPointType, let N : int> T determinant(matrix<T,N,N> 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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("T distance(vector<T, N> x, vector<T, N> 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<T : __BuiltinArithmeticType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("T dot(vector<T, N> x, vector<T, N> 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<T : __BuiltinFloatingPointType> vector<T,4> dst(vector<T,4> x, vector<T,4> 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<T : __BuiltinArithmeticType>\n") -SLANG_RAW("__target_intrinsic(glsl, interpolateAtCentroid)\n") -SLANG_RAW("T EvaluateAttributeAtCentroid(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int>\n") -SLANG_RAW("__target_intrinsic(glsl, interpolateAtCentroid)\n") -SLANG_RAW("vector<T,N> EvaluateAttributeAtCentroid(vector<T,N> x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(glsl, interpolateAtCentroid)\n") -SLANG_RAW("matrix<T,N,M> EvaluateAttributeAtCentroid(matrix<T,N,M> x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType>\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<T : __BuiltinArithmeticType, let N : int>\n") -SLANG_RAW("__target_intrinsic(glsl, \"interpolateAtSample($0, int($1))\")\n") -SLANG_RAW("vector<T,N> EvaluateAttributeAtSample(vector<T,N> x, uint sampleindex);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(glsl, \"interpolateAtSample($0, int($1))\")\n") -SLANG_RAW("matrix<T,N,M> EvaluateAttributeAtSample(matrix<T,N,M> x, uint sampleindex);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType>\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<T : __BuiltinArithmeticType, let N : int>\n") -SLANG_RAW("__target_intrinsic(glsl, \"interpolateAtOffset($0, vec2($1) / 16.0f)\")\n") -SLANG_RAW("vector<T,N> EvaluateAttributeSnapped(vector<T,N> x, int2 offset);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(glsl, \"interpolateAtOffset($0, vec2($1) / 16.0f)\")\n") -SLANG_RAW("matrix<T,N,M> EvaluateAttributeSnapped(matrix<T,N,M> x, int2 offset);\n") -SLANG_RAW("\n") -SLANG_RAW("// Base-e exponent\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T exp(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> exp(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, exp, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> exp(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T exp2(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T,N> exp2(vector<T,N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, exp2, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T,N,M> exp2(matrix<T,N,M> 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<let N : int>\n") -SLANG_RAW("__target_intrinsic(glsl, \"unpackHalf2x16($0).x\")\n") -SLANG_RAW("vector<float,N> f16tof32(vector<uint,N> 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<let N : int>\n") -SLANG_RAW("__target_intrinsic(glsl, \"packHalf2x16(vec2($0,0.0))\")\n") -SLANG_RAW("vector<uint,N> f32tof16(vector<float,N> value);\n") -SLANG_RAW("\n") -SLANG_RAW("// Flip surface normal to face forward, if needed\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T,N> faceforward(vector<T,N> n, vector<T,N> i, vector<T,N> 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<let N : int>\n") -SLANG_RAW("vector<int, N> firstbithigh(vector<int, N> 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<let N : int>\n") -SLANG_RAW("vector<uint,N> firstbithigh(vector<uint,N> 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<let N : int>\n") -SLANG_RAW("vector<int,N> firstbitlow(vector<int,N> 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<let N : int>\n") -SLANG_RAW("vector<uint,N> firstbitlow(vector<uint,N> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T floor(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> floor(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, floor, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> floor(matrix<T, N, M> 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<let N : int> vector<double, N> fma(vector<double, N> a, vector<double, N> b, vector<double, N> c);\n") -SLANG_RAW("__generic<let N : int, let M : int> matrix<double,N,M> fma(matrix<double,N,M> a, matrix<double,N,M> b, matrix<double,N,M> c);\n") -SLANG_RAW("\n") -SLANG_RAW("// Floating point remainder of x/y\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T fmod(T x, T y);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> fmod(vector<T, N> x, vector<T, N> 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<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> fmod(matrix<T, N, M> x, matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("__target_intrinsic(glsl, fract)\n") -SLANG_RAW("T frac(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, fract)\n") -SLANG_RAW("vector<T, N> frac(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, frac, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("matrix<T, N, M> frac(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> frexp(vector<T, N> x, out vector<T, N> 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<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> frexp(matrix<T, N, M> x, out matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T fwidth(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> fwidth(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, fwidth, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> fwidth(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("//__target_intrinsic(glsl, \"(!(isinf($0) || isnan($0)))\")\n") -SLANG_RAW("vector<bool, N> isfinite(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(bool, N, isfinite, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<bool, N, M> isfinite(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("bool isinf(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<bool, N> isinf(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(bool, N, isinf, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<bool, N, M> isinf(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("bool isnan(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<bool, N> isnan(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(bool, N, isnan, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<bool, N, M> isnan(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"($0 * pow(2.0f, $1))\")\n") -SLANG_RAW("vector<T, N> ldexp(vector<T, N> x, vector<T, N> 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<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> ldexp(matrix<T, N, M> x, matrix<T, N, M> 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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("T length(vector<T, N> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, mix)\n") -SLANG_RAW("vector<T, N> lerp(vector<T, N> x, vector<T, N> y, vector<T, N> 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<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T,N,M> lerp(matrix<T,N,M> x, matrix<T,N,M> y, matrix<T,N,M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T log(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> log(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, log, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> log(matrix<T, N, M> 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<T : __BuiltinFloatingPointType> \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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"(log( $0 ) * $S0(0.43429448190325182765112891891661) )\" )\n") -SLANG_RAW("vector<T,N> log10(vector<T,N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, log10, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> \n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T,N,M> log10(matrix<T,N,M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T log2(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T,N> log2(vector<T,N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, log2, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T,N,M> log2(matrix<T,N,M> 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 : __BuiltinArithmeticType> T mad(T mvalue, T avalue, T bvalue);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, fma)\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> mad(vector<T,N> mvalue, vector<T,N> avalue, vector<T,N> bvalue);\n") -SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(glsl, fma)\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> mad(matrix<T,N,M> mvalue, matrix<T,N,M> avalue, matrix<T,N,M> bvalue);\n") -SLANG_RAW("\n") -SLANG_RAW("// maximum\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType>\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<T : __BuiltinArithmeticType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> max(vector<T, N> x, vector<T, N> 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<T : __BuiltinArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> max(matrix<T, N, M> x, matrix<T, N, M> 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<T : __BuiltinArithmeticType>\n") -SLANG_RAW("T min(T x, T y);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T,N> min(vector<T,N> x, vector<T,N> 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<T : __BuiltinArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T,N,M> min(matrix<T,N,M> x, matrix<T,N,M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T modf(T x, out T ip);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T,N> modf(vector<T,N> x, out vector<T,N> 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<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T,N,M> modf(matrix<T,N,M> x, out matrix<T,N,M> 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 : __BuiltinArithmeticType> T mul(T x, T y);\n") -SLANG_RAW("\n") -SLANG_RAW("// scalar-vector and vector-scalar\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> mul(vector<T,N> x, T y);\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> mul(T x, vector<T,N> y);\n") -SLANG_RAW("\n") -SLANG_RAW("// scalar-matrix and matrix-scalar\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M :int> matrix<T,N,M> mul(matrix<T,N,M> x, T y);\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M :int> matrix<T,N,M> mul(T x, matrix<T,N,M> y);\n") -SLANG_RAW("\n") -SLANG_RAW("// vector-vector (dot product)\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> __intrinsic_op(dot) T mul(vector<T,N> x, vector<T,N> y);\n") -SLANG_RAW("\n") -SLANG_RAW("// vector-matrix\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> __intrinsic_op(mulVectorMatrix) vector<T,M> mul(vector<T,N> x, matrix<T,N,M> y);\n") -SLANG_RAW("\n") -SLANG_RAW("// matrix-vector\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> __intrinsic_op(mulMatrixVector) vector<T,N> mul(matrix<T,N,M> x, vector<T,M> y);\n") -SLANG_RAW("\n") -SLANG_RAW("// matrix-matrix\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let R : int, let N : int, let C : int> __intrinsic_op(mulMatrixMatrix) matrix<T,R,C> mul(matrix<T,R,N> x, matrix<T,N,C> 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<let N : int> float noise(vector<float, N> 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<T : __BuiltinFloatingPointType, let N : int> vector<T,N> normalize(vector<T,N> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T pow(T x, T y);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> pow(vector<T, N> x, vector<T, N> 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<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T,N,M> pow(matrix<T,N,M> x, matrix<T,N,M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> radians(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, radians, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> radians(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("vector<T, N> rcp(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, rcp, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\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<T, N, M> rcp(matrix<T, N, M> 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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("vector<T,N> reflect(vector<T,N> i, vector<T,N> 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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("vector<T,N> refract(vector<T,N> i, vector<T,N> 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,N>(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<let N : int> vector<uint,N> reversebits(vector<uint,N> value);\n") -SLANG_RAW("\n") -SLANG_RAW("// Round-to-nearest\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T round(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> round(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, round, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T,N,M> round(matrix<T,N,M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"inversesqrt($0)\")\n") -SLANG_RAW("vector<T, N> rsqrt(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, rsqrt, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> rsqrt(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("T saturate(T x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return clamp<T>(x, T(0), T(1));\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("vector<T,N> saturate(vector<T,N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" return clamp<T,N>(x,\n") -SLANG_RAW(" vector<T,N>(T(0)),\n") -SLANG_RAW(" vector<T,N>(T(1)));\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T,N,M> saturate(matrix<T,N,M> 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<T : __BuiltinSignedArithmeticType>\n") -SLANG_RAW("__target_intrinsic(glsl, \"int(sign($0))\")\n") -SLANG_RAW("int sign(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinSignedArithmeticType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl, \"ivec$N0(sign($0))\")\n") -SLANG_RAW("vector<int, N> sign(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(int, N, sign, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinSignedArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<int, N, M> sign(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T sin(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> sin(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, sin, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> sin(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("void sincos(vector<T,N> x, out vector<T,N> s, out vector<T,N> 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<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("void sincos(matrix<T,N,M> x, out matrix<T,N,M> s, out matrix<T,N,M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T sinh(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> sinh(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, sinh, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> sinh(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T smoothstep(T min, T max, T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> smoothstep(vector<T, N> min, vector<T, N> max, vector<T, N> 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<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> smoothstep(matrix<T, N, M> min, matrix<T, N, M> max, matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T sqrt(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> sqrt(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, sqrt, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> sqrt(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\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<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T,N> step(vector<T,N> y, vector<T,N> 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<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> step(matrix<T, N, M> y, matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T tan(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> tan(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, tan, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> tan(matrix<T, N, M> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T tanh(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T,N> tanh(vector<T,N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, tanh, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T,N,M> tanh(matrix<T,N,M> 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<T : __BuiltinType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("matrix<T, M, N> transpose(matrix<T, N, M> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" matrix<T,M,N> 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<T : __BuiltinFloatingPointType>\n") -SLANG_RAW("T trunc(T x);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("__target_intrinsic(glsl)\n") -SLANG_RAW("vector<T, N> trunc(vector<T, N> x)\n") -SLANG_RAW("{\n") -SLANG_RAW(" VECTOR_MAP_UNARY(T, N, trunc, x);\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n") -SLANG_RAW("__target_intrinsic(hlsl)\n") -SLANG_RAW("matrix<T, N, M> trunc(matrix<T, N, M> 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 : __BuiltinType> T QuadReadLaneAt(T sourceValue, uint quadLaneID);\n") -SLANG_RAW("__generic<T : __BuiltinType, let N : int> vector<T,N> QuadReadLaneAt(vector<T,N> sourceValue, uint quadLaneID);\n") -SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M> QuadReadLaneAt(matrix<T,N,M> sourceValue, uint quadLaneID);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinType> T QuadReadAcrossX(T localValue);\n") -SLANG_RAW("__generic<T : __BuiltinType, let N : int> vector<T,N> QuadReadAcrossX(vector<T,N> localValue);\n") -SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M> QuadReadAcrossX(matrix<T,N,M> localValue);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinType> T QuadReadAcrossY(T localValue);\n") -SLANG_RAW("__generic<T : __BuiltinType, let N : int> vector<T,N> QuadReadAcrossY(vector<T,N> localValue);\n") -SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M> QuadReadAcrossY(matrix<T,N,M> localValue);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinType> T QuadReadAcrossDiagonal(T localValue);\n") -SLANG_RAW("__generic<T : __BuiltinType, let N : int> vector<T,N> QuadReadAcrossDiagonal(vector<T,N> localValue);\n") -SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M> QuadReadAcrossDiagonal(matrix<T,N,M> localValue);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinIntegerType>\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<T : __BuiltinIntegerType, let N : int>\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<T,N> WaveActiveBitAnd(vector<T,N> expr);\n") -SLANG_RAW("__generic<T : __BuiltinIntegerType, let N : int, let M : int>\n") -SLANG_RAW("matrix<T,N,M> WaveActiveBitAnd(matrix<T,N,M> expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinIntegerType>\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<T : __BuiltinIntegerType, let N : int>\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<T,N> WaveActiveBitOr(vector<T,N> expr);\n") -SLANG_RAW("__generic<T : __BuiltinIntegerType, let N : int, let M : int>\n") -SLANG_RAW("matrix<T,N,M> WaveActiveBitOr(matrix<T,N,M> expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinIntegerType>\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<T : __BuiltinIntegerType, let N : int> \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<T,N> WaveActiveBitXor(vector<T,N> expr);\n") -SLANG_RAW("__generic<T : __BuiltinIntegerType, let N : int, let M : int>\n") -SLANG_RAW("matrix<T,N,M> WaveActiveBitXor(matrix<T,N,M> expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType>\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<T : __BuiltinArithmeticType, let N : int>\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<T,N> WaveActiveMax(vector<T,N> expr);\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("matrix<T,N,M> WaveActiveMax(matrix<T,N,M> expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType>\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<T : __BuiltinArithmeticType, let N : int>\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<T,N> WaveActiveMin(vector<T,N> expr);\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("matrix<T,N,M> WaveActiveMin(matrix<T,N,M> expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType>\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<T : __BuiltinArithmeticType, let N : int>\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<T,N> WaveActiveProduct(vector<T,N> expr);\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("matrix<T,N,M> WaveActiveProduct(matrix<T,N,M> expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType>\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<T : __BuiltinArithmeticType, let N : int>\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<T,N> WaveActiveSum(vector<T,N> expr);\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("matrix<T,N,M> WaveActiveSum(matrix<T,N,M> expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinType>\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<T : __BuiltinType, let N : int> \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<bool,N> WaveActiveAllEqual(vector<T,N> value);\n") -SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int>\n") -SLANG_RAW("matrix<bool,N,M> WaveActiveAllEqual(matrix<T,N,M> value);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinType> uint4 WaveMatch(T value);\n") -SLANG_RAW("__generic<T : __BuiltinType, let N : int> uint4 WaveMatch(vector<T,N> value);\n") -SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int> uint4 WaveMatch(matrix<T,N,M> 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<T : __BuiltinArithmeticType>\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<T : __BuiltinArithmeticType, let N : int>\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<T,N> WavePrefixProduct(vector<T,N> expr);\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("matrix<T,N,M> WavePrefixProduct(matrix<T,N,M> expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType>\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<T : __BuiltinArithmeticType, let N : int>\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<T,N> WavePrefixSum(vector<T,N> expr);\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("matrix<T,N,M> WavePrefixSum(matrix<T,N,M> expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType>\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<T : __BuiltinArithmeticType, let N : int>\n") -SLANG_RAW("vector<T,N> WaveMultiPrefixBitAnd(vector<T,N> expr);\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("matrix<T,N,M> WaveMultiPrefixBitAnd(matrix<T,N,M> expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType>\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<T : __BuiltinArithmeticType, let N : int>\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<T,N> WaveMultiPrefixBitOr(vector<T,N> expr);\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("matrix<T,N,M> WaveMultiPrefixBitOr(matrix<T,N,M> expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType>\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<T : __BuiltinArithmeticType, let N : int>\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<T,N> WaveMultiPrefixBitXor(vector<T,N> expr);\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n") -SLANG_RAW("matrix<T,N,M> WaveMultiPrefixBitXor(matrix<T,N,M> 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 : __BuiltinArithmeticType> T WaveMultiPrefixProduct(T value, uint4 mask);\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WaveMultiPrefixProduct(vector<T,N> value, uint4 mask);\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WaveMultiPrefixProduct(matrix<T,N,M> value, uint4 mask);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType> T WaveMultiPrefixSum(T value, uint4 mask);\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WaveMultiPrefixSum(vector<T,N> value, uint4 mask);\n") -SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WaveMultiPrefixSum(matrix<T,N,M> value, uint4 mask);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinType>\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<T : __BuiltinType, let N : int>\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<T,N> WaveReadLaneFirst(vector<T,N> expr);\n") -SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int>\n") -SLANG_RAW("matrix<T,N,M> WaveReadLaneFirst(matrix<T,N,M> expr);\n") -SLANG_RAW("\n") -SLANG_RAW("__generic<T : __BuiltinType>\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<T : __BuiltinType, let N : int>\n") -SLANG_RAW("__spirv_version(1.3)\n") -SLANG_RAW("__target_intrinsic(glsl, \"subgroupBroadcast($0, $1)\")\n") -SLANG_RAW("vector<T,N> WaveReadLaneAt(vector<T,N> value, int lane);\n") -SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int>\n") -SLANG_RAW("matrix<T,N,M> WaveReadLaneAt(matrix<T,N,M> 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<T>\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<Payload>(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<Payload>\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<Payload>\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<payload_t>(\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<Payload>\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<payload_t>\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<A>(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<A>\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<P>(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 <let rayFlags : RAY_FLAG = RAY_FLAG_NONE>\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<T>\n") -SLANG_RAW("{\n") -SLANG_RAW(" T SubpassLoad();\n") -SLANG_RAW("}\n") -SLANG_RAW("\n") -SLANG_RAW("struct VkSubpassInputMS<T>\n") -SLANG_RAW("{\n") -SLANG_RAW(" T SubpassLoad(int sampleIndex);\n") -SLANG_RAW("}\n") |