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 | |
| 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.
| -rw-r--r-- | docs/building.md | 11 | ||||
| -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 |
4 files changed, 11 insertions, 5206 deletions
diff --git a/docs/building.md b/docs/building.md index 8e44b9040..316a48e75 100644 --- a/docs/building.md +++ b/docs/building.md @@ -23,6 +23,17 @@ For Linux and other targets the section below on `premake` describes the process Some targets below are described as 'unofficial'. In practice this means that they are not tested as part of contiguous integration. Thus unfortunately it is quite possible from time to time for them to break on a merge of a PR. That said, if broken it is likely only very minor changes are needed to make them work again. +### Generated Files + +Slang as part of it's build process generates header files, which are then used to compile the main Slang project. If you use `premake` to create your project, it will automatically generate these files before compiling the rest of the Slang. These are the current header generations which are created via the `slang-generate` tool... + +* core.meta.slang -> core.meta.slang.h +* hlsl.meta.slang -> hlsl.meta.slang.h + +It may be necessary or desirable to create a build of Slang without using `premake`. + +One way to do this would be to first compile slang-generate and then invoke it directly or as a dependency in your build. Another perhaps simpler way would be to first compile the same Slang source on another system that does support `premake`, or using a preexisting build mechanism (such as Visual Studio projects on Windows). Then copy the generated header files to your target system. This is appropriate because the generated files are indentical across platforms. It does of course mean that if `core.meta.slang` or `hlsl.meta.slang` files change the headers will need to be regenerated. + ## Premake Slang uses the tool [`premake5`](https://premake.github.io/) in order to generate projects that can be built on different targets. On Linux premake will generate Makefile/s and on windows it will generate a Visual Studio solution. Information on invoking premake for different kinds of targets can be found [here](https://github.com/premake/premake-core/wiki/Using-Premake). You can also run with `--help` to see available command line options 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") |