diff options
62 files changed, 3950 insertions, 871 deletions
diff --git a/.github/workflows/linux-arm64.yml b/.github/workflows/linux-arm64.yml index 8a2f958a2..bc1991444 100644 --- a/.github/workflows/linux-arm64.yml +++ b/.github/workflows/linux-arm64.yml @@ -53,4 +53,5 @@ jobs: CONFIGURATION=${{matrix.configuration}} CC=${{matrix.compiler}} ARCH=${{matrix.platform}} + PATH="${PATH:+${PATH}:}$(pwd)/external/slang-binaries/spirv-tools/$(uname -m)-linux/bin" source ./github_test.sh diff --git a/.github/workflows/linux.yml b/.github/workflows/linux.yml index fa4e4f446..435346837 100644 --- a/.github/workflows/linux.yml +++ b/.github/workflows/linux.yml @@ -69,4 +69,5 @@ jobs: CONFIGURATION=${{matrix.configuration}} CC=${{matrix.compiler}} ARCH=${{matrix.platform}} + PATH="${PATH:+${PATH}:}$(pwd)/external/slang-binaries/spirv-tools/$(uname -m)-linux/bin" source ./github_test.sh diff --git a/.github/workflows/windows-selfhosted.yml b/.github/workflows/windows-selfhosted.yml index 62370feb4..414f7515c 100644 --- a/.github/workflows/windows-selfhosted.yml +++ b/.github/workflows/windows-selfhosted.yml @@ -39,6 +39,6 @@ jobs: - name: test run: | $slangTestBinDir = ".\bin\windows-${{matrix.testPlatform}}\${{matrix.configuration}}\"; - $env:Path += ";$slangTestBinDir"; + $spirvToolsBinDir = ".\external\slang-binaries\spirv-tools\windows-${{matrix.testPlatform}}\bin\"; + $env:Path += ";$slangTestBinDir;$spirvToolsBinDir"; & "$slangTestBinDir\slang-test.exe" -appveyor -bindir "$slangTestBinDir\" -platform ${{matrix.testPlatform}} -configuration ${{matrix.configuration}} -category ${{matrix.testCategory}} -api all-cpu 2>&1; -
\ No newline at end of file diff --git a/.github/workflows/windows.yml b/.github/workflows/windows.yml index 3d8fe3ea8..f3ac5fc04 100644 --- a/.github/workflows/windows.yml +++ b/.github/workflows/windows.yml @@ -63,7 +63,8 @@ jobs: - name: test run: | $slangTestBinDir = ".\bin\windows-${{matrix.testPlatform}}\${{matrix.configuration}}\"; - $env:Path += ";$slangTestBinDir"; + $spirvToolsBinDir = ".\external\slang-binaries\spirv-tools\windows-${{matrix.testPlatform}}\bin\"; + $env:Path += ";$slangTestBinDir;$spirvToolsBinDir"; Expand-Archive "vk_swiftshader_windows_${{matrix.testPlatform}}.zip" -DestinationPath $slangTestBinDir; & "$slangTestBinDir\slang-test.exe" -api all-dx12 -appveyor -bindir "$slangTestBinDir\" -platform ${{matrix.testPlatform}} -configuration ${{matrix.configuration}} -category ${{matrix.testCategory}} 2>&1; -
\ No newline at end of file + diff --git a/build/visual-studio/slang/slang.vcxproj b/build/visual-studio/slang/slang.vcxproj index 7d655f0c5..e9b34374c 100644 --- a/build/visual-studio/slang/slang.vcxproj +++ b/build/visual-studio/slang/slang.vcxproj @@ -344,6 +344,7 @@ IF EXIST ..\..\..\external\slang-glslang\bin\windows-aarch64\release\slang-glsla <ClInclude Include="..\..\..\source\slang\slang-emit-hlsl.h" />
<ClInclude Include="..\..\..\source\slang\slang-emit-precedence.h" />
<ClInclude Include="..\..\..\source\slang\slang-emit-source-writer.h" />
+ <ClInclude Include="..\..\..\source\slang\slang-emit-spirv-ops.h" />
<ClInclude Include="..\..\..\source\slang\slang-emit-torch.h" />
<ClInclude Include="..\..\..\source\slang\slang-glsl-extension-tracker.h" />
<ClInclude Include="..\..\..\source\slang\slang-hlsl-to-vulkan-layout-options.h" />
@@ -396,6 +397,7 @@ IF EXIST ..\..\..\external\slang-glslang\bin\windows-aarch64\release\slang-glsla <ClInclude Include="..\..\..\source\slang\slang-ir-inst-defs.h" />
<ClInclude Include="..\..\..\source\slang\slang-ir-inst-pass-base.h" />
<ClInclude Include="..\..\..\source\slang\slang-ir-insts.h" />
+ <ClInclude Include="..\..\..\source\slang\slang-ir-layout-on-types.h" />
<ClInclude Include="..\..\..\source\slang\slang-ir-layout.h" />
<ClInclude Include="..\..\..\source\slang\slang-ir-legalize-array-return-type.h" />
<ClInclude Include="..\..\..\source\slang\slang-ir-legalize-mesh-outputs.h" />
@@ -603,6 +605,7 @@ IF EXIST ..\..\..\external\slang-glslang\bin\windows-aarch64\release\slang-glsla <ClCompile Include="..\..\..\source\slang\slang-ir-glsl-liveness.cpp" />
<ClCompile Include="..\..\..\source\slang\slang-ir-init-local-var.cpp" />
<ClCompile Include="..\..\..\source\slang\slang-ir-inline.cpp" />
+ <ClCompile Include="..\..\..\source\slang\slang-ir-layout-on-types.cpp" />
<ClCompile Include="..\..\..\source\slang\slang-ir-layout.cpp" />
<ClCompile Include="..\..\..\source\slang\slang-ir-legalize-array-return-type.cpp" />
<ClCompile Include="..\..\..\source\slang\slang-ir-legalize-mesh-outputs.cpp" />
diff --git a/build/visual-studio/slang/slang.vcxproj.filters b/build/visual-studio/slang/slang.vcxproj.filters index 8e8023520..b95f1894d 100644 --- a/build/visual-studio/slang/slang.vcxproj.filters +++ b/build/visual-studio/slang/slang.vcxproj.filters @@ -120,6 +120,9 @@ <ClInclude Include="..\..\..\source\slang\slang-emit-source-writer.h">
<Filter>Header Files</Filter>
</ClInclude>
+ <ClInclude Include="..\..\..\source\slang\slang-emit-spirv-ops.h">
+ <Filter>Header Files</Filter>
+ </ClInclude>
<ClInclude Include="..\..\..\source\slang\slang-emit-torch.h">
<Filter>Header Files</Filter>
</ClInclude>
@@ -276,6 +279,9 @@ <ClInclude Include="..\..\..\source\slang\slang-ir-insts.h">
<Filter>Header Files</Filter>
</ClInclude>
+ <ClInclude Include="..\..\..\source\slang\slang-ir-layout-on-types.h">
+ <Filter>Header Files</Filter>
+ </ClInclude>
<ClInclude Include="..\..\..\source\slang\slang-ir-layout.h">
<Filter>Header Files</Filter>
</ClInclude>
@@ -893,6 +899,9 @@ <ClCompile Include="..\..\..\source\slang\slang-ir-inline.cpp">
<Filter>Source Files</Filter>
</ClCompile>
+ <ClCompile Include="..\..\..\source\slang\slang-ir-layout-on-types.cpp">
+ <Filter>Source Files</Filter>
+ </ClCompile>
<ClCompile Include="..\..\..\source\slang\slang-ir-layout.cpp">
<Filter>Source Files</Filter>
</ClCompile>
diff --git a/docs/command-line-slangc-reference.md b/docs/command-line-slangc-reference.md index 8bd15bbee..862ed9022 100644 --- a/docs/command-line-slangc-reference.md +++ b/docs/command-line-slangc-reference.md @@ -550,9 +550,9 @@ Verify IR in the front-end. Experimental options (use at your own risk) -<a id="emit-spirv-directly"></a> -## -emit-spirv-directly -Generate SPIR-V output directly (otherwise through GLSL and using the glslang compiler) +<a id="emit-spirv-via-glsl"></a> +## -emit-spirv-via-glsl +Generate SPIR-V output by generating and compiling GLSL using glslang <a id="file-system"></a> @@ -852,7 +852,7 @@ A capability describes an optional feature that a target may or may not support. * `c` * `cpp` * `cuda` -* `spirv_direct` +* `spirv` * `GL_NV_ray_tracing` : enables the GL_NV_ray_tracing extension * `GL_EXT_ray_tracing` : enables the GL_EXT_ray_tracing extension * `GL_NV_fragment_shader_barycentric` : enables the GL_NV_fragment_shader_barycentric extension diff --git a/docs/design/stdlib-intrinsics.md b/docs/design/stdlib-intrinsics.md index 0aaa99057..76614beba 100644 --- a/docs/design/stdlib-intrinsics.md +++ b/docs/design/stdlib-intrinsics.md @@ -77,7 +77,7 @@ This is a widely used and somewhat complicated intrinsic. Placed on a declaratio * cuda - CUDA * cpp - C++ output (used for exe, shared-library or host-callable) -* spirv_direct - Used for slangs SPIR-V direct mechanism +* spirv - Used for slangs SPIR-V direct mechanism A function definition can have a `target_intrinsic` *and* a body. In that case, the body will be used for targets where the `target_intrinsic` isn't defined. diff --git a/external/slang-binaries b/external/slang-binaries -Subproject dede65e9fabaea0ad59c38eda8730b02f043052 +Subproject a68cd451448c1a480dd3df339b64457d7e46c34 diff --git a/premake5.lua b/premake5.lua index 404016b65..daffde45d 100644 --- a/premake5.lua +++ b/premake5.lua @@ -228,6 +228,14 @@ newoption { allowed = { { "true", "True"}, { "false", "False" } } } +newoption { + trigger = "default-spirv-direct", + description = "(Optional) Development flag to make the default SPIR-V path generate directly rather than via GLSL", + value = "bool", + default = "false", + allowed = { { "true", "True"}, { "false", "False" } } +} + buildLocation = _OPTIONS["build-location"] executeBinary = (_OPTIONS["execute-binary"] == "true") buildGlslang = (_OPTIONS["build-glslang"] == "true") @@ -245,6 +253,7 @@ fullDebugValidation = (_OPTIONS["full-debug-validation"] == "true") enableAsan = (_OPTIONS["enable-asan"] == "true") dxOnVk = (_OPTIONS["dx-on-vk"] == "true") enableAftermath = (_OPTIONS["enable-aftermath"] == "true") +defaultSPIRVDirect = (_OPTIONS["default-spirv-direct"] == "true") -- If stdlib embedding is enabled, disable stdlib source embedding by default disableStdlibSource = enableEmbedStdLib @@ -435,6 +444,10 @@ workspace "slang" defines { "SLANG_CONFIG_DX_ON_VK" } end + if defaultSPIRVDirect then + defines { "SLANG_CONFIG_DEFAULT_SPIRV_DIRECT" } + end + function dump(o) if type(o) == 'table' then local s = '{ ' diff --git a/slang-gfx.h b/slang-gfx.h index 577c8666e..82d8b768e 100644 --- a/slang-gfx.h +++ b/slang-gfx.h @@ -2204,7 +2204,7 @@ public: const char* targetProfile = nullptr; // (optional) Target shader profile. If null this will be set to platform dependent default. SlangFloatingPointMode floatingPointMode = SLANG_FLOATING_POINT_MODE_DEFAULT; SlangOptimizationLevel optimizationLevel = SLANG_OPTIMIZATION_LEVEL_DEFAULT; - SlangTargetFlags targetFlags = 0; + SlangTargetFlags targetFlags = kDefaultTargetFlags; SlangLineDirectiveMode lineDirectiveMode = SLANG_LINE_DIRECTIVE_MODE_DEFAULT; }; @@ -699,9 +699,14 @@ extern "C" /* When set, will dump out the IR between intermediate compilation steps.*/ SLANG_TARGET_FLAG_DUMP_IR = 1 << 9, - /* When set, will generate SPIRV directly instead of going through glslang. */ + /* When set, will generate SPIRV directly rather than via glslang. */ SLANG_TARGET_FLAG_GENERATE_SPIRV_DIRECTLY = 1 << 10, }; +#if defined(SLANG_CONFIG_DEFAULT_SPIRV_DIRECT) + constexpr static SlangTargetFlags kDefaultTargetFlags = SLANG_TARGET_FLAG_GENERATE_SPIRV_DIRECTLY; +#else + constexpr static SlangTargetFlags kDefaultTargetFlags = 0; +#endif /*! @brief Options to control floating-point precision guarantees for a target. @@ -737,6 +742,7 @@ extern "C" SLANG_SOURCE_LANGUAGE_C, SLANG_SOURCE_LANGUAGE_CPP, SLANG_SOURCE_LANGUAGE_CUDA, + SLANG_SOURCE_LANGUAGE_SPIRV, SLANG_SOURCE_LANGUAGE_COUNT_OF, }; @@ -4152,7 +4158,7 @@ namespace slang SlangProfileID profile = SLANG_PROFILE_UNKNOWN; /** Flags for the code generation target. Currently unused. */ - SlangTargetFlags flags = 0; + SlangTargetFlags flags = kDefaultTargetFlags; /** Default mode to use for floating-point operations on the target. */ diff --git a/source/compiler-core/slang-spirv-dis-compiler.cpp b/source/compiler-core/slang-spirv-dis-compiler.cpp index 04e5c8e4a..0e484c7c5 100644 --- a/source/compiler-core/slang-spirv-dis-compiler.cpp +++ b/source/compiler-core/slang-spirv-dis-compiler.cpp @@ -1,8 +1,12 @@ #include "slang-spirv-dis-compiler.h" #include "../core/slang-common.h" +#include "../core/slang-string-util.h" +#include "../core/slang-string.h" +#include "slang-artifact-desc-util.h" #include "slang-artifact-representation.h" #include "slang-artifact-util.h" +#include "slang-artifact-representation-impl.h" namespace Slang { @@ -39,19 +43,23 @@ SlangResult SLANG_MCALL SPIRVDisDownstreamCompiler::convert( ISlangBlob* fromBlob; SLANG_RETURN_ON_FAIL(from->loadBlob(ArtifactKeep::No, &fromBlob)); + ComPtr<IOSFileArtifactRepresentation> fromFile; + SLANG_RETURN_ON_FAIL(from->requireFile(ArtifactKeep::No, fromFile.writeRef())); + + String toFile; + File::generateTemporary(UnownedStringSlice("spv-asm"), toFile); + // Set up our process CommandLine commandLine; commandLine.m_executableLocation.setName("spirv-dis"); + commandLine.addArg("--comment"); + commandLine.addArg(fromFile->getPath()); + commandLine.addArg("-o"); + commandLine.addArg(toFile); RefPtr<Process> p; SLANG_RETURN_ON_FAIL(Process::create(commandLine, 0, p)); - const auto in = p->getStream(StdStreamType::In); - const auto out = p->getStream(StdStreamType::Out); const auto err = p->getStream(StdStreamType::ErrorOut); - // Write the assembly - SLANG_RETURN_ON_FAIL(in->write(fromBlob->getBufferPointer(), fromBlob->getBufferSize())); - in->close(); - // Wait for it to finish if(!p->waitForTermination(1000)) return SLANG_FAIL; @@ -61,18 +69,26 @@ SlangResult SLANG_MCALL SPIRVDisDownstreamCompiler::convert( SLANG_RETURN_ON_FAIL(StreamUtil::readAll(err, 0, errData)); fwrite(errData.getBuffer(), errData.getCount(), 1, stderr); + // If spirv-dis failed, we fail const auto ret = p->getReturnValue(); if(ret != 0) return SLANG_FAIL; - // Read the disassembly - List<Byte> outData; - SLANG_RETURN_ON_FAIL(StreamUtil::readAll(out, 0, outData)); - - // Wobble it into an artifact - ComPtr<ISlangBlob> outBlob = RawBlob::create(outData.getBuffer(), outData.getCount()); + // Normalize line endings + String outContents; + SLANG_RETURN_ON_FAIL(File::readAllText(toFile, outContents)); + StringBuilder outBuilder; + StringUtil::appendStandardLines(outContents.getUnownedSlice(), outBuilder); + SLANG_RETURN_ON_FAIL(File::writeAllBytes(toFile, outBuilder.getBuffer(), outBuilder.getLength())); + + // Return as a file artifact + auto fileRep = OSFileArtifactRepresentation::create( + IOSFileArtifactRepresentation::Kind::Owned, + toFile.getUnownedSlice(), + nullptr + ); auto artifact = ArtifactUtil::createArtifact(to); - artifact->addRepresentationUnknown(outBlob.detach()); + artifact->addRepresentation(fileRep.detach()); *outArtifact = artifact.detach(); return SLANG_OK; diff --git a/source/core/slang-array.h b/source/core/slang-array.h index 86fdd301f..f1ed4fe0d 100644 --- a/source/core/slang-array.h +++ b/source/core/slang-array.h @@ -101,6 +101,50 @@ namespace Slang Index m_count = 0; }; + template<typename T> + class Array<T, 0> + { + public: + T* begin() { return nullptr; } + const T* begin() const { return nullptr; } + + const T* end() const { return nullptr; } + T* end() { return nullptr; } + + inline Index getCapacity() const { return 0; } + inline Index getCount() const { return 0; } + inline void setCount(Index newCount) + { + SLANG_ASSERT(newCount == 0); + } + inline const T* getBuffer() const { return nullptr; } + inline T* getBuffer() { return nullptr; } + inline void clear() {} + + template<typename T2> + Index indexOf(const T2& val) const { return getView().indexOf(val); } + template<typename T2> + Index lastIndexOf(const T2& val) const { return getView().lastIndexOf(val); } + template<typename Func> + Index findFirstIndex(const Func& predicate) const { return getView().findFirstIndex(predicate); } + template<typename Func> + Index findLastIndex(const Func& predicate) const { return getView().findLastIndex(predicate); } + + inline ConstArrayView<T> getView() const { return ConstArrayView<T>(nullptr, 0); } + inline ConstArrayView<T> getView(Index start, Index count) const + { + SLANG_ASSERT(start == 0 && count == 0); + return ConstArrayView<T>(nullptr, 0); + } + + inline ArrayView<T> getView() { return ArrayView<T>(nullptr, 0); } + inline ArrayView<T> getView(Index start, Index count) + { + SLANG_ASSERT(start == 0 && count == 0); + return ArrayView<T>(nullptr, 0); + } + }; + template<typename T, typename ...TArgs> struct FirstType { diff --git a/source/slang/core.meta.slang b/source/slang/core.meta.slang index ebf4b5d76..157b83653 100644 --- a/source/slang/core.meta.slang +++ b/source/slang/core.meta.slang @@ -2078,6 +2078,16 @@ struct TextureTypeInfo sb << "__target_intrinsic(glsl, \"$ctextureLod($p, $2, $3)$z\")\n"; + // SPIR-V + { + // TODO: + // Need to: + // - Construct sampled image type OpTypeSampledImage of image type + // - Construct OpSampledImage from image and sampler + // - Call OpImageSampleExplicitLod + // test ./tests/compute/texture-simpler.slang + } + // CUDA { const int coordCount = base.coordCount; diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang index 015b1c4ad..d721bec25 100644 --- a/source/slang/hlsl.meta.slang +++ b/source/slang/hlsl.meta.slang @@ -235,9 +235,10 @@ struct StructuredBuffer out uint numStructs, out uint stride); - //__target_intrinsic(glsl, "$0._data[$1]") - //__target_intrinsic(spirv_direct, "%addr = OpAccessChain resultType*StorageBuffer resultId _0 const(int, 0) _1; OpLoad resultType resultId %addr;") __intrinsic_op($(kIROp_StructuredBufferLoad)) + __target_intrinsic(glsl, "$0._data[$1]") + __target_intrinsic(spirv, "%addr = OpAccessChain resultType*StorageBuffer resultId _0 const(int, 0) _1; OpLoad resultType resultId %addr;") + [__readNone] T Load(int location); __intrinsic_op($(kIROp_StructuredBufferLoadStatus)) @@ -245,8 +246,6 @@ struct StructuredBuffer __subscript(uint index) -> T { - //__target_intrinsic(glsl, "$0._data[$1]") - //__target_intrinsic(spirv_direct, "%addr = OpAccessChain resultType*StorageBuffer resultId _0 const(int, 0) _1; OpLoad resultType resultId %addr;") [__readNone] __intrinsic_op($(kIROp_StructuredBufferLoad)) get; @@ -724,8 +723,6 @@ struct $(item.name) uint IncrementCounter(); - //__target_intrinsic(glsl, "$0._data[$1]") - //__target_intrinsic(spirv_direct, "%addr = OpAccessChain resultType*StorageBuffer resultId _0 const(int, 0) _1; OpLoad resultType resultId %addr;") [__NoSideEffect] __intrinsic_op($(kIROp_RWStructuredBufferLoad)) T Load(int location); @@ -736,8 +733,6 @@ struct $(item.name) __subscript(uint index) -> T { - //__target_intrinsic(glsl, "$0._data[$1]") - //__target_intrinsic(spirv_direct, "*StorageBuffer OpAccessChain resultType resultId _0 const(int, 0) _1") [__NoSideEffect] __intrinsic_op($(kIROp_RWStructuredBufferGetElementPtr)) ref; @@ -812,7 +807,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_abs($0)") __target_intrinsic(cpp, "$P_abs($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fi(FAbs, SAbs) _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fi(FAbs, SAbs) _0") [__readNone] T abs(T x); /*{ @@ -823,7 +818,7 @@ T abs(T x); __generic<T : __BuiltinIntegerType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fi(FAbs, SAbs) _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fi(FAbs, SAbs) _0") [__readNone] vector<T, N> abs(vector<T, N> x) { @@ -843,14 +838,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_abs($0)") __target_intrinsic(cpp, "$P_abs($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fi(FAbs, SAbs) _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fi(FAbs, SAbs) _0") [__readNone] T abs(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fi(FAbs, SAbs) _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fi(FAbs, SAbs) _0") [__readNone] vector<T, N> abs(vector<T, N> x) { @@ -872,14 +867,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_acos($0)") __target_intrinsic(cpp, "$P_acos($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Acos _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Acos _0") [__readNone] T acos(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Acos _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Acos _0") [__readNone] vector<T, N> acos(vector<T, N> x) { @@ -899,12 +894,18 @@ __generic<T : __BuiltinType> __target_intrinsic(cpp, "bool($0)") __target_intrinsic(cuda, "bool($0)") __target_intrinsic(glsl, "bool($0)") +__target_intrinsic(spirv, boolean(T), "OpCopyObject resultType resultId _0") +__target_intrinsic(spirv, integral(T), "OpINotEqual resultType resultId _0 const(,0)") +__target_intrinsic(spirv, floating(T), "OpFUnordNotEqual resultType resultId _0 const(,0)") [__readNone] bool all(T x); __generic<T : __BuiltinType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "all(bvec$N0($0))") +__target_intrinsic(spirv, boolean(T), "OpAll resultType resultId _0") +__target_intrinsic(spirv, integral(T), "%c = OpINotEqual resultType resultId _0 const(,0); OpAll resultType resultId %c") +__target_intrinsic(spirv, floating(T), "%c = OpFUnordNotEqual resultType resultId _0 const(,0); OpAll resultType resultId %c") [__readNone] bool all(vector<T,N> x) { @@ -974,7 +975,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "packDouble2x32(uvec2($0, $1))") __target_intrinsic(cpp, "$P_asdouble($0, $1)") __target_intrinsic(cuda, "$P_asdouble($0, $1)") -__target_intrinsic(spirv_direct, "%v = OpCompositeConstruct _type(uint2) resultId _0 _1; OpExtInst resultType resultId glsl450 59 %v") +__target_intrinsic(spirv, "%v = OpCompositeConstruct _type(uint2) resultId _0 _1; OpExtInst resultType resultId glsl450 59 %v") __glsl_extension(GL_ARB_gpu_shader5) [__readNone] double asdouble(uint lowbits, uint highbits); @@ -985,7 +986,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "intBitsToFloat") __target_intrinsic(cpp, "$P_asfloat($0)") __target_intrinsic(cuda, "$P_asfloat($0)") -__target_intrinsic(spirv_direct, "OpBitcast resultType resultId _0") +__target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] float asfloat(int x); @@ -993,14 +994,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "uintBitsToFloat") __target_intrinsic(cpp, "$P_asfloat($0)") __target_intrinsic(cuda, "$P_asfloat($0)") -__target_intrinsic(spirv_direct, "OpBitcast resultType resultId _0") +__target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] float asfloat(uint x); __generic<let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "intBitsToFloat") -__target_intrinsic(spirv_direct, "OpBitcast resultType resultId _0") +__target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] vector<float, N> asfloat(vector< int, N> x) { @@ -1010,7 +1011,7 @@ vector<float, N> asfloat(vector< int, N> x) __generic<let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "uintBitsToFloat") -__target_intrinsic(spirv_direct, "OpBitcast resultType resultId _0") +__target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] vector<float,N> asfloat(vector<uint,N> x) { @@ -1057,14 +1058,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_asin($0)") __target_intrinsic(cpp, "$P_asin($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Asin _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Asin _0") [__readNone] T asin(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Asin _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Asin _0") [__readNone] vector<T, N> asin(vector<T, N> x) { @@ -1085,7 +1086,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "floatBitsToInt") __target_intrinsic(cpp, "$P_asint($0)") __target_intrinsic(cuda, "$P_asint($0)") -__target_intrinsic(spirv_direct, "OpBitcast resultType resultId _0") +__target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] int asint(float x); @@ -1093,14 +1094,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "int($0)") __target_intrinsic(cpp, "$P_asint($0)") __target_intrinsic(cuda, "$P_asint($0)") -__target_intrinsic(spirv_direct, "OpBitcast resultType resultId _0") +__target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] int asint(uint x); __generic<let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "floatBitsToInt") -__target_intrinsic(spirv_direct, "OpBitcast resultType resultId _0") +__target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] vector<int, N> asint(vector<float, N> x) { @@ -1110,7 +1111,7 @@ vector<int, N> asint(vector<float, N> x) __generic<let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "ivec$N0($0)") -__target_intrinsic(spirv_direct, "OpBitcast resultType resultId _0") +__target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] vector<int, N> asint(vector<uint, N> x) { @@ -1165,7 +1166,7 @@ void asuint(double value, out uint lowbits, out uint highbits); __target_intrinsic(hlsl) __target_intrinsic(glsl, "floatBitsToUint") -__target_intrinsic(spirv_direct, "OpBitcast resultType resultId _0") +__target_intrinsic(spirv, "OpBitcast resultType resultId _0") __target_intrinsic(cpp, "$P_asuint($0)") __target_intrinsic(cuda, "$P_asuint($0)") [__readNone] @@ -1173,7 +1174,7 @@ uint asuint(float x); __target_intrinsic(hlsl) __target_intrinsic(glsl, "uint($0)") -__target_intrinsic(spirv_direct, "OpBitcast resultType resultId _0") +__target_intrinsic(spirv, "OpBitcast resultType resultId _0") __target_intrinsic(cpp, "$P_asuint($0)") __target_intrinsic(cuda, "$P_asuint($0)") [__readNone] @@ -1182,7 +1183,7 @@ uint asuint(int x); __generic<let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "floatBitsToUint") -__target_intrinsic(spirv_direct, "OpBitcast resultType resultId _0") +__target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] vector<uint,N> asuint(vector<float,N> x) { @@ -1192,7 +1193,7 @@ vector<uint,N> asuint(vector<float,N> x) __generic<let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "uvec$N0($0)") -__target_intrinsic(spirv_direct, "OpBitcast resultType resultId _0") +__target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] vector<uint, N> asuint(vector<int, N> x) { @@ -1317,14 +1318,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_atan($0)") __target_intrinsic(cpp, "$P_atan($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Atan _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Atan _0") [__readNone] T atan(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Atan _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Atan _0") [__readNone] vector<T, N> atan(vector<T, N> x) { @@ -1344,14 +1345,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl,"atan($0,$1)") __target_intrinsic(cuda, "$P_atan2($0, $1)") __target_intrinsic(cpp, "$P_atan2($0, $1)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Atan2 _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Atan2 _0 _1") [__readNone] T atan2(T y, T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl,"atan($0,$1)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Atan2 _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Atan2 _0 _1") [__readNone] vector<T, N> atan2(vector<T, N> y, vector<T, N> x) { @@ -1372,14 +1373,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_ceil($0)") __target_intrinsic(cpp, "$P_ceil($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Ceil _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Ceil _0") [__readNone] T ceil(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Ceil _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Ceil _0") [__readNone] vector<T, N> ceil(vector<T, N> x) { @@ -1402,7 +1403,7 @@ bool CheckAccessFullyMapped(uint status); __generic<T : __BuiltinIntegerType> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fus(FClamp, UClamp, SClamp) _0 _1 _2") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FClamp, UClamp, SClamp) _0 _1 _2") [__readNone] T clamp(T x, T minBound, T maxBound) { @@ -1412,7 +1413,7 @@ T clamp(T x, T minBound, T maxBound) __generic<T : __BuiltinIntegerType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fus(FClamp, UClamp, SClamp) _0 _1 _2") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FClamp, UClamp, SClamp) _0 _1 _2") [__readNone] vector<T, N> clamp(vector<T, N> x, vector<T, N> minBound, vector<T, N> maxBound) { @@ -1430,7 +1431,7 @@ matrix<T,N,M> clamp(matrix<T,N,M> x, matrix<T,N,M> minBound, matrix<T,N,M> maxBo __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fus(FClamp, UClamp, SClamp) _0 _1 _2") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FClamp, UClamp, SClamp) _0 _1 _2") [__readNone] T clamp(T x, T minBound, T maxBound) { @@ -1440,7 +1441,7 @@ T clamp(T x, T minBound, T maxBound) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fus(FClamp, UClamp, SClamp) _0 _1 _2") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FClamp, UClamp, SClamp) _0 _1 _2") [__readNone] vector<T, N> clamp(vector<T, N> x, vector<T, N> minBound, vector<T, N> maxBound) { @@ -1483,14 +1484,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_cos($0)") __target_intrinsic(cpp, "$P_cos($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Cos _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Cos _0") [__readNone] T cos(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Cos _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Cos _0") [__readNone] vector<T, N> cos(vector<T, N> x) { @@ -1511,14 +1512,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_cosh($0)") __target_intrinsic(cpp, "$P_cosh($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Cosh _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Cosh _0") [__readNone] T cosh(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Cosh _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Cosh _0") [__readNone] vector<T,N> cosh(vector<T,N> x) { @@ -1546,7 +1547,7 @@ uint countbits(uint value); __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Cross _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Cross _0 _1") [__readNone] vector<T,3> cross(vector<T,3> left, vector<T,3> right) { @@ -1559,7 +1560,7 @@ vector<T,3> cross(vector<T,3> left, vector<T,3> right) __generic<T : __BuiltinIntegerType> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Cross _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Cross _0 _1") [__readNone] vector<T, 3> cross(vector<T, 3> left, vector<T, 3> right) { @@ -1729,7 +1730,7 @@ matrix<T, N, M> ddy_fine(matrix<T, N, M> x) __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Degrees _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Degrees _0") [__readNone] T degrees(T x) { @@ -1739,7 +1740,7 @@ T degrees(T x) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Degrees _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Degrees _0") [__readNone] vector<T, N> degrees(vector<T, N> x) { @@ -1759,7 +1760,7 @@ matrix<T, N, M> degrees(matrix<T, N, M> x) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Determinant _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Determinant _0") [__readNone] [PreferCheckpoint] T determinant(matrix<T,N,N> m); @@ -1778,7 +1779,7 @@ void DeviceMemoryBarrierWithGroupSync(); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Distance _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Distance _0 _1") [__readNone] T distance(vector<T, N> x, vector<T, N> y) { @@ -1840,13 +1841,13 @@ RWStructuredBuffer<T> __getEquivalentStructuredBuffer<T>(RWByteAddressBuffer b); __generic<T : __BuiltinArithmeticType> __target_intrinsic(glsl, interpolateAtCentroid) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 InterpolateAtCentroid _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 InterpolateAtCentroid _0") [__readNone] T EvaluateAttributeAtCentroid(T x); __generic<T : __BuiltinArithmeticType, let N : int> __target_intrinsic(glsl, interpolateAtCentroid) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 InterpolateAtCentroid _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 InterpolateAtCentroid _0") [__readNone] vector<T,N> EvaluateAttributeAtCentroid(vector<T,N> x); @@ -1860,13 +1861,13 @@ matrix<T,N,M> EvaluateAttributeAtCentroid(matrix<T,N,M> x) __generic<T : __BuiltinArithmeticType> __target_intrinsic(glsl, "interpolateAtSample($0, int($1))") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 InterpolateAtSample _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 InterpolateAtSample _0 _1") [__readNone] T EvaluateAttributeAtSample(T x, uint sampleindex); __generic<T : __BuiltinArithmeticType, let N : int> __target_intrinsic(glsl, "interpolateAtSample($0, int($1))") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 InterpolateAtSample _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 InterpolateAtSample _0 _1") [__readNone] vector<T,N> EvaluateAttributeAtSample(vector<T,N> x, uint sampleindex); @@ -1885,13 +1886,13 @@ matrix<T,N,M> EvaluateAttributeAtSample(matrix<T,N,M> x, uint sampleindex) __generic<T : __BuiltinArithmeticType> __target_intrinsic(glsl, "interpolateAtOffset($0, vec2($1) / 16.0f)") -__target_intrinsic(spirv_direct, "%foffset = OpConvertSToF _type(float2) resultId _1; %offsetdiv16 = 136 _type(float2) resultId %foffset const(float2, 16.0, 16.0); OpExtInst resultType resultId glsl450 78 _0 %offsetdiv16") +__target_intrinsic(spirv, "%foffset = OpConvertSToF _type(float2) resultId _1; %offsetdiv16 = 136 _type(float2) resultId %foffset const(float2, 16.0, 16.0); OpExtInst resultType resultId glsl450 78 _0 %offsetdiv16") [__readNone] T EvaluateAttributeSnapped(T x, int2 offset); __generic<T : __BuiltinArithmeticType, let N : int> __target_intrinsic(glsl, "interpolateAtOffset($0, vec2($1) / 16.0f)") -__target_intrinsic(spirv_direct, "%foffset = OpConvertSToF _type(float2) resultId _1; %offsetdiv16 = 136 _type(float2) resultId %foffset const(float2, 16.0, 16.0); OpExtInst resultType resultId glsl450 78 _0 %offsetdiv16") +__target_intrinsic(spirv, "%foffset = OpConvertSToF _type(float2) resultId _1; %offsetdiv16 = 136 _type(float2) resultId %foffset const(float2, 16.0, 16.0); OpExtInst resultType resultId glsl450 78 _0 %offsetdiv16") [__readNone] vector<T,N> EvaluateAttributeSnapped(vector<T,N> x, int2 offset); @@ -1915,14 +1916,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_exp($0)") __target_intrinsic(cpp, "$P_exp($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Exp _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Exp _0") [__readNone] T exp(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Exp _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Exp _0") [__readNone] vector<T, N> exp(vector<T, N> x) { @@ -1944,14 +1945,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_exp2($0)") __target_intrinsic(cpp, "$P_exp2($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Exp2 _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Exp2 _0") [__readNone] T exp2(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Exp2 _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Exp2 _0") [__readNone] vector<T,N> exp2(vector<T,N> x) { @@ -2055,13 +2056,13 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl,"findMSB") __target_intrinsic(cuda, "$P_firstbithigh($0)") __target_intrinsic(cpp, "$P_firstbithigh($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 FindSMsb _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FindSMsb _0") [__readNone] int firstbithigh(int value); __target_intrinsic(hlsl) __target_intrinsic(glsl,"findMSB") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 FindSMsb _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FindSMsb _0") __generic<let N : int> [__readNone] vector<int, N> firstbithigh(vector<int, N> value) @@ -2073,13 +2074,13 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl,"findMSB") __target_intrinsic(cuda, "$P_firstbithigh($0)") __target_intrinsic(cpp, "$P_firstbithigh($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 FindUMsb _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FindUMsb _0") [__readNone] uint firstbithigh(uint value); __target_intrinsic(hlsl) __target_intrinsic(glsl,"findMSB") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 FindUMsb _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FindUMsb _0") __generic<let N : int> [__readNone] vector<uint,N> firstbithigh(vector<uint,N> value) @@ -2092,13 +2093,13 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl,"findLSB") __target_intrinsic(cuda, "$P_firstbitlow($0)") __target_intrinsic(cpp, "$P_firstbitlow($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 FindILsb _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FindILsb _0") [__readNone] int firstbitlow(int value); __target_intrinsic(hlsl) __target_intrinsic(glsl,"findLSB") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 FindILsb _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FindILsb _0") __generic<let N : int> [__readNone] vector<int,N> firstbitlow(vector<int,N> value) @@ -2110,14 +2111,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl,"findLSB") __target_intrinsic(cuda, "$P_firstbitlow($0)") __target_intrinsic(cpp, "$P_firstbitlow($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 FindILsb _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FindILsb _0") [__readNone] uint firstbitlow(uint value); __target_intrinsic(hlsl) __target_intrinsic(glsl,"findLSB") __generic<let N : int> -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 FindILsb _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FindILsb _0") [__readNone] vector<uint,N> firstbitlow(vector<uint,N> value) { @@ -2131,14 +2132,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_floor($0)") __target_intrinsic(cpp, "$P_floor($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Floor _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Floor _0") [__readNone] T floor(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Floor _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Floor _0") [__readNone] vector<T, N> floor(vector<T, N> x) { @@ -2158,14 +2159,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_fma($0, $1, $2)") __target_intrinsic(cpp, "$P_fma($0, $1, $2)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") [__readNone] double fma(double a, double b, double c); __generic<let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") [__readNone] vector<double, N> fma(vector<double, N> a, vector<double, N> b, vector<double, N> c) { @@ -2213,14 +2214,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, fract) __target_intrinsic(cuda, "$P_frac($0)") __target_intrinsic(cpp, "$P_frac($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Fract _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Fract _0") [__readNone] T frac(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, fract) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Fract _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Fract _0") [__readNone] vector<T, N> frac(vector<T, N> x) { @@ -2238,14 +2239,14 @@ matrix<T, N, M> frac(matrix<T, N, M> x) __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Frexp _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Frexp _0 _1") [__readNone] T frexp(T x, out T exp); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Frexp _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Frexp _0 _1") [__readNone] vector<T, N> frexp(vector<T, N> x, out vector<T, N> exp) { @@ -2360,11 +2361,23 @@ float2 GetRenderTargetSamplePosition(int Index); // Group memory barrier __target_intrinsic(glsl, "groupMemoryBarrier") __target_intrinsic(cuda, "__threadfence_block") +__target_intrinsic(spirv, "OpMemoryBarrier const(int,ScopeWorkgroup)" + "const(int, MemorySemanticsAcquireReleaseMask" + "| MemorySemanticsUniformMemoryMask" + "| MemorySemanticsImageMemoryMask" + "| MemorySemanticsAtomicCounterMemoryMask" + "| MemorySemanticsWorkgroupMemoryMask)") void GroupMemoryBarrier(); __target_intrinsic(glsl, "groupMemoryBarrier(), barrier()") __target_intrinsic(cuda, "__syncthreads()") +__target_intrinsic(spirv, "OpControlBarrier const(int,ScopeWorkgroup) const(int, ScopeWorkgroup)" + "const(int, MemorySemanticsAcquireReleaseMask" + "| MemorySemanticsUniformMemoryMask" + "| MemorySemanticsImageMemoryMask" + "| MemorySemanticsAtomicCounterMemoryMask" + "| MemorySemanticsWorkgroupMemoryMask)") void GroupMemoryBarrierWithGroupSync(); // Atomics @@ -2387,6 +2400,9 @@ void InterlockedAdd(__ref int dest, int value, out int original_value); __target_intrinsic(glsl, "($2 = $atomicAdd($A, $1))") __target_intrinsic(cuda, "(*$2 = (uint)atomicAdd((uint*)$0, $1))") +__target_intrinsic(spirv, "%old = OpAtomicIAdd _type(uint) resultId _0" + "const(int, ScopeDevice) const(int, MemorySemanticsMaskNone) _1;" + "OpStore _2 %old;") void InterlockedAdd(__ref uint dest, uint value, out uint original_value); __target_intrinsic(glsl, "$atomicAnd($A, $1)") @@ -2577,7 +2593,7 @@ matrix<bool, N, M> isnan(matrix<T, N, M> x) __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Ldexp _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Ldexp _0 _1") [__readNone] T ldexp(T x, T exp) { @@ -2586,7 +2602,7 @@ T ldexp(T x, T exp) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Ldexp _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Ldexp _0 _1") [__readNone] vector<T, N> ldexp(vector<T, N> x, vector<T, N> exp) { @@ -2605,7 +2621,7 @@ matrix<T, N, M> ldexp(matrix<T, N, M> x, matrix<T, N, M> exp) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Length _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Length _0") [__readNone] T length(vector<T, N> x) { @@ -2614,16 +2630,16 @@ T length(vector<T, N> x) // Scalar float length __generic<T : __BuiltinFloatingPointType> -T length(T x) -{ - return abs(x); +T length(T x) +{ + return abs(x); } // Linear interpolation __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) __target_intrinsic(glsl, mix) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 FMix _0 _1 _2") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FMix _0 _1 _2") [__readNone] T lerp(T x, T y, T s) { @@ -2633,7 +2649,7 @@ T lerp(T x, T y, T s) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, mix) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 FMix _0 _1 _2") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FMix _0 _1 _2") [__readNone] vector<T, N> lerp(vector<T, N> x, vector<T, N> y, vector<T, N> s) { @@ -2665,14 +2681,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_log($0)") __target_intrinsic(cpp, "$P_log($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Log _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Log _0") [__readNone] T log(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Log _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Log _0") [__readNone] vector<T, N> log(vector<T, N> x) { @@ -2693,14 +2709,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "(log( $0 ) * $S0( 0.43429448190325182765112891891661) )" ) __target_intrinsic(cuda, "$P_log10($0)") __target_intrinsic(cpp, "$P_log10($0)") -__target_intrinsic(spirv_direct, "%baseElog = OpExtInst resultType resultId glsl450 Log _0; OpFMul resultType resultId _0 %baseElog const(_p,0.43429448190325182765112891891661)") +__target_intrinsic(spirv, "%baseElog = OpExtInst resultType resultId glsl450 Log _0; OpFMul resultType resultId _0 %baseElog const(_p,0.43429448190325182765112891891661)") [__readNone] T log10(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "(log( $0 ) * $S0(0.43429448190325182765112891891661) )" ) -__target_intrinsic(spirv_direct, "%baseElog = OpExtInst resultType resultId glsl450 Log _0; OpVectorTimesScalar resultType resultId _0 %baseElog const(_p,0.43429448190325182765112891891661)") +__target_intrinsic(spirv, "%baseElog = OpExtInst resultType resultId glsl450 Log _0; OpVectorTimesScalar resultType resultId _0 %baseElog const(_p,0.43429448190325182765112891891661)") [__readNone] vector<T,N> log10(vector<T,N> x) { @@ -2721,14 +2737,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_log2($0)") __target_intrinsic(cpp, "$P_log2($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Log2 _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Log2 _0") [__readNone] T log2(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Log2 _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Log2 _0") [__readNone] vector<T,N> log2(vector<T,N> x) { @@ -2750,14 +2766,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, fma) __target_intrinsic(cuda, "$P_fma($0, $1, $2)") __target_intrinsic(cpp, "$P_fma($0, $1, $2)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") [__readNone] T mad(T mvalue, T avalue, T bvalue); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, fma) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") [__readNone] vector<T, N> mad(vector<T, N> mvalue, vector<T, N> avalue, vector<T, N> bvalue) { @@ -2777,14 +2793,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, fma) __target_intrinsic(cuda, "$P_fma($0, $1, $2)") __target_intrinsic(cpp, "$P_fma($0, $1, $2)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") [__readNone] T mad(T mvalue, T avalue, T bvalue); __generic<T : __BuiltinIntegerType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, fma) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") [__readNone] vector<T, N> mad(vector<T, N> mvalue, vector<T, N> avalue, vector<T, N> bvalue) { @@ -2806,7 +2822,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_max($0, $1)") __target_intrinsic(cpp, "$P_max($0, $1)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fus(FMax, UMax, SMax) _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMax, UMax, SMax) _0") [__readNone] T max(T x, T y); // Note: a stdlib implementation of `max` (or `min`) will require splitting @@ -2817,7 +2833,7 @@ T max(T x, T y); __generic<T : __BuiltinIntegerType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fus(FMax, UMax, SMax) _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMax, UMax, SMax) _0") [__readNone] vector<T, N> max(vector<T, N> x, vector<T, N> y) { @@ -2837,14 +2853,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_max($0, $1)") __target_intrinsic(cpp, "$P_max($0, $1)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fus(FMax, UMax, SMax) _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMax, UMax, SMax) _0") [__readNone] T max(T x, T y); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fus(FMax, UMax, SMax) _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMax, UMax, SMax) _0") [__readNone] vector<T, N> max(vector<T, N> x, vector<T, N> y) { @@ -2865,14 +2881,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_min($0, $1)") __target_intrinsic(cpp, "$P_min($0, $1)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fus(FMin, UMin, SMin) _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMin, UMin, SMin) _0") [__readNone] T min(T x, T y); __generic<T : __BuiltinIntegerType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fus(FMin, UMin, SMin) _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMin, UMin, SMin) _0") [__readNone] vector<T,N> min(vector<T,N> x, vector<T,N> y) { @@ -2892,14 +2908,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_min($0, $1)") __target_intrinsic(cpp, "$P_min($0, $1)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fus(FMin, UMin, SMin) _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMin, UMin, SMin) _0") [__readNone] T min(T x, T y); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fus(FMin, UMin, SMin) _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMin, UMin, SMin) _0") [__readNone] vector<T,N> min(vector<T,N> x, vector<T,N> y) { @@ -3225,10 +3241,10 @@ int NonUniformResourceIndex(int index) return index; } -/// HLSL allows NonUniformResourceIndex around non int/uint types. +/// HLSL allows NonUniformResourceIndex around non int/uint types. /// It's effect is presumably to ignore it, which the following implementation does. /// We should also look to add a warning for this scenario. -[__unsafeForceInlineEarly] +[__unsafeForceInlineEarly] [deprecated("NonUniformResourceIndex on a type other than uint/int is depreciated and has no effect")] T NonUniformResourceIndex<T>(T value) { return value; } @@ -3236,7 +3252,7 @@ T NonUniformResourceIndex<T>(T value) { return value; } __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Normalize _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Normalize _0") [__readNone] vector<T,N> normalize(vector<T,N> x) { @@ -3249,14 +3265,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_pow($0, $1)") __target_intrinsic(cpp, "$P_pow($0, $1)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Pow _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Pow _0 _1") [__readNone] T pow(T x, T y); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Pow _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Pow _0 _1") [__readNone] vector<T, N> pow(vector<T, N> x, vector<T, N> y) { @@ -3412,7 +3428,7 @@ void ProcessTriTessFactorsMin( __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Radians _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Radians _0") [__readNone] T radians(T x) { @@ -3422,7 +3438,7 @@ T radians(T x) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Radians _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Radians _0") [__readNone] vector<T, N> radians(vector<T, N> x) { @@ -3466,7 +3482,7 @@ matrix<T, N, M> rcp(matrix<T, N, M> x) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Reflect _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Reflect _0 _1") [__readNone] vector<T,N> reflect(vector<T,N> i, vector<T,N> n) { @@ -3477,7 +3493,7 @@ vector<T,N> reflect(vector<T,N> i, vector<T,N> n) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Refract _0 _1 _2") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Refract _0 _1 _2") [__readNone] vector<T,N> refract(vector<T,N> i, vector<T,N> n, T eta) { @@ -3509,14 +3525,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_round($0)") __target_intrinsic(cpp, "$P_round($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Round _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Round _0") [__readNone] T round(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Round _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Round _0") [__readNone] vector<T, N> round(vector<T, N> x) { @@ -3537,7 +3553,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "inversesqrt($0)") __target_intrinsic(cuda, "$P_rsqrt($0)") __target_intrinsic(cpp, "$P_rsqrt($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 InverseSqrt _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 InverseSqrt _0") [__readNone] T rsqrt(T x) { @@ -3547,7 +3563,7 @@ T rsqrt(T x) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "inversesqrt($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 InverseSqrt _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 InverseSqrt _0") [__readNone] vector<T, N> rsqrt(vector<T, N> x) { @@ -3596,14 +3612,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "int(sign($0))") __target_intrinsic(cuda, "$P_sign($0)") __target_intrinsic(cpp, "$P_sign($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fi(FSign, SSign) _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fi(FSign, SSign) _0") [__readNone] int sign(T x); __generic<T : __BuiltinSignedArithmeticType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "ivec$N0(sign($0))") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 fi(FSign, SSign) _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fi(FSign, SSign) _0") [__readNone] vector<int, N> sign(vector<T, N> x) { @@ -3626,14 +3642,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_sin($0)") __target_intrinsic(cpp, "$P_sin($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Sin _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Sin _0") [__readNone] T sin(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Sin _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Sin _0") [__readNone] vector<T, N> sin(vector<T, N> x) { @@ -3683,14 +3699,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_sinh($0)") __target_intrinsic(cpp, "$P_sinh($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Sinh _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Sinh _0") [__readNone] T sinh(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Sinh _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Sinh _0") [__readNone] vector<T, N> sinh(vector<T, N> x) { @@ -3709,7 +3725,7 @@ matrix<T, N, M> sinh(matrix<T, N, M> x) __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 SmoothStep _0 _1 _2") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 SmoothStep _0 _1 _2") [__readNone] T smoothstep(T min, T max, T x) { @@ -3720,7 +3736,7 @@ T smoothstep(T min, T max, T x) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 SmoothStep _0 _1 _2") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 SmoothStep _0 _1 _2") [__readNone] vector<T, N> smoothstep(vector<T, N> min, vector<T, N> max, vector<T, N> x) { @@ -3741,14 +3757,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_sqrt($0)") __target_intrinsic(cpp, "$P_sqrt($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Sqrt _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Sqrt _0") [__readNone] T sqrt(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Sqrt _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Sqrt _0") [__readNone] vector<T, N> sqrt(vector<T, N> x) { @@ -3767,7 +3783,7 @@ matrix<T, N, M> sqrt(matrix<T, N, M> x) __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Step _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Step _0 _1") [__readNone] T step(T y, T x) { @@ -3777,7 +3793,7 @@ T step(T y, T x) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Step _0 _1") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Step _0 _1") [__readNone] vector<T,N> step(vector<T,N> y, vector<T,N> x) { @@ -3798,14 +3814,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_tan($0)") __target_intrinsic(cpp, "$P_tan($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Tan _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Tan _0") [__readNone] T tan(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Tan _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Tan _0") [__readNone] vector<T, N> tan(vector<T, N> x) { @@ -3826,14 +3842,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_tanh($0)") __target_intrinsic(cpp, "$P_tanh($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Tanh _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Tanh _0") [__readNone] T tanh(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Tanh _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Tanh _0") [__readNone] vector<T,N> tanh(vector<T,N> x) { @@ -3895,14 +3911,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(cuda, "$P_trunc($0)") __target_intrinsic(cpp, "$P_trunc($0)") -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Trunc _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Trunc _0") [__readNone] T trunc(T x); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) -__target_intrinsic(spirv_direct, "OpExtInst resultType resultId glsl450 Trunc _0") +__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Trunc _0") [__readNone] vector<T, N> trunc(vector<T, N> x) { diff --git a/source/slang/slang-ast-modifier.h b/source/slang/slang-ast-modifier.h index 01cf3a24a..7b772e645 100644 --- a/source/slang/slang-ast-modifier.h +++ b/source/slang/slang-ast-modifier.h @@ -79,8 +79,16 @@ class TargetIntrinsicModifier : public Modifier // is an intrisic for. Token targetToken; - // A custom definition for the operation - Token definitionToken; + // A custom definition for the operation, one of either an ident or a + // string (the concatenation of several string literals) + Token definitionIdent; + String definitionString; + bool isString; + + // A predicate to be used on an identifier to guard this intrinsic + Token predicateToken; + NameLoc scrutinee; + DeclRef<Decl> scrutineeDeclRef; }; // A modifier that marks a declaration as representing a diff --git a/source/slang/slang-capability-defs.h b/source/slang/slang-capability-defs.h index e9deb3fd6..c523d9cda 100644 --- a/source/slang/slang-capability-defs.h +++ b/source/slang/slang-capability-defs.h @@ -57,7 +57,7 @@ SLANG_CAPABILITY_ATOM1(C, c, Concrete,TargetFormat,0,TEXTUAL_SOUR SLANG_CAPABILITY_ATOM1(CPP, cpp, Concrete,TargetFormat,0,TEXTUAL_SOURCE) SLANG_CAPABILITY_ATOM1(CUDA, cuda, Concrete,TargetFormat,0,TEXTUAL_SOURCE) -SLANG_CAPABILITY_ATOM0(SPIRV_DIRECT, spirv_direct, Concrete, TargetFormat, 0) +SLANG_CAPABILITY_ATOM0(SPIRV_DIRECT, spirv, Concrete, TargetFormat, 0) // We have multiple capabilities for the various SPIR-V versions, // arranged so that they inherit from one another to represent which versions diff --git a/source/slang/slang-check-modifier.cpp b/source/slang/slang-check-modifier.cpp index ab34c83dd..21d7669ce 100644 --- a/source/slang/slang-check-modifier.cpp +++ b/source/slang/slang-check-modifier.cpp @@ -984,6 +984,37 @@ namespace Slang return packOffsetModifier; } + if(auto targetIntrinsic = as<TargetIntrinsicModifier>(m)) + { + // TODO: verify that the predicate is one we understand + if(targetIntrinsic->scrutinee.name) + { + if(auto genDecl = as<ContainerDecl>(syntaxNode)) + { + auto scrutineeResults = lookUp( + m_astBuilder, + this, + targetIntrinsic->scrutinee.name, + genDecl->ownedScope); + if(!scrutineeResults.isValid()) + { + getSink()->diagnose( + targetIntrinsic->scrutinee.loc, + Diagnostics::undefinedIdentifier2, + targetIntrinsic->scrutinee.name); + } + if(scrutineeResults.isOverloaded()) + { + getSink()->diagnose( + targetIntrinsic->scrutinee.loc, + Diagnostics::ambiguousReference, + targetIntrinsic->scrutinee.name); + } + targetIntrinsic->scrutineeDeclRef = scrutineeResults.item.declRef; + } + } + } + // Default behavior is to leave things as they are, // and assume that modifiers are mostly already checked. // diff --git a/source/slang/slang-compiler.cpp b/source/slang/slang-compiler.cpp index afc7899a8..f786512d1 100644 --- a/source/slang/slang-compiler.cpp +++ b/source/slang/slang-compiler.cpp @@ -456,6 +456,10 @@ namespace Slang { return SourceLanguage::CUDA; } + case PassThroughMode::SpirvDis: + { + return SourceLanguage::SPIRV; + } default: break; } @@ -485,7 +489,7 @@ namespace Slang case CodeGenTarget::SPIRVAssembly: case CodeGenTarget::SPIRV: { - return PassThroughMode::Glslang; + return PassThroughMode::SpirvDis; } case CodeGenTarget::DXBytecode: case CodeGenTarget::DXBytecodeAssembly: diff --git a/source/slang/slang-compiler.h b/source/slang/slang-compiler.h index e42e97b89..040773546 100755 --- a/source/slang/slang-compiler.h +++ b/source/slang/slang-compiler.h @@ -1597,7 +1597,7 @@ namespace Slang bool shouldEmitSPIRVDirectly() { - return (targetFlags & SLANG_TARGET_FLAG_GENERATE_SPIRV_DIRECTLY) != 0; + return targetFlags & SLANG_TARGET_FLAG_GENERATE_SPIRV_DIRECTLY; } bool isWholeProgramRequest() @@ -1636,7 +1636,7 @@ namespace Slang private: Linkage* linkage = nullptr; CodeGenTarget format = CodeGenTarget::Unknown; - SlangTargetFlags targetFlags = 0; + SlangTargetFlags targetFlags = kDefaultTargetFlags; Slang::Profile targetProfile = Slang::Profile(); FloatingPointMode floatingPointMode = FloatingPointMode::Default; List<CapabilityAtom> rawCapabilities; diff --git a/source/slang/slang-doc-markdown-writer.cpp b/source/slang/slang-doc-markdown-writer.cpp index c13dc9668..0e7856f1f 100644 --- a/source/slang/slang-doc-markdown-writer.cpp +++ b/source/slang/slang-doc-markdown-writer.cpp @@ -427,7 +427,7 @@ static DocMarkdownWriter::Requirement _getRequirementFromTargetToken(const Token } auto targetName = tok.getContent(); - if (targetName == "spirv_direct") + if (targetName == "spirv") { return Requirement{CodeGenTarget::SPIRV, UnownedStringSlice("")}; } diff --git a/source/slang/slang-emit-spirv-ops.h b/source/slang/slang-emit-spirv-ops.h new file mode 100644 index 000000000..445f6f7d0 --- /dev/null +++ b/source/slang/slang-emit-spirv-ops.h @@ -0,0 +1,2183 @@ +#ifdef SLANG_IN_SPIRV_EMIT_CONTEXT +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpUndef +template<typename T> +SpvInst* emitOpUndef(SpvInstParent* parent, IRInst* inst, const T& idResultType) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpUndef, idResultType, kResultID); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpName +template<typename T> +SpvInst* emitOpName( + SpvInstParent* parent, + IRInst* inst, + const T& target, + const UnownedStringSlice& name +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpName, target, name); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpExtension +SpvInst* emitOpExtension(SpvInstParent* parent, IRInst* inst, const UnownedStringSlice& name) +{ + return emitInst(parent, inst, SpvOpExtension, name); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpExtInstImport +SpvInst* emitOpExtInstImport(SpvInstParent* parent, IRInst* inst, const UnownedStringSlice& name) +{ + return emitInst(parent, inst, SpvOpExtInstImport, kResultID, name); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemoryModel +SpvInst* emitOpMemoryModel( + SpvInstParent* parent, + IRInst* inst, + SpvAddressingModel addressingModel, + SpvMemoryModel memoryModel +) +{ + return emitInst(parent, inst, SpvOpMemoryModel, addressingModel, memoryModel); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpEntryPoint +template<typename T, typename Ts> +SpvInst* emitOpEntryPoint( + SpvInstParent* parent, + IRInst* inst, + SpvExecutionModel executionModel, + const T& entryPoint, + const UnownedStringSlice& name, + const Ts& interfaces +) +{ + static_assert(isSingular<T>); + static_assert(isPlural<Ts>); + return emitInst(parent, inst, SpvOpEntryPoint, executionModel, entryPoint, name, interfaces); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpExecutionMode +template<typename T> +SpvInst* emitOpExecutionMode( + SpvInstParent* parent, + IRInst* inst, + const T& entryPoint, + SpvExecutionMode mode +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpExecutionMode, entryPoint, mode); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpExecutionMode +template<typename T> +SpvInst* emitOpExecutionModeLocalSize( + SpvInstParent* parent, + IRInst* inst, + const T& entryPoint, + const SpvLiteralInteger& xSize, + const SpvLiteralInteger& ySize, + const SpvLiteralInteger& zSize +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, inst, SpvOpExecutionMode, entryPoint, SpvExecutionModeLocalSize, xSize, ySize, zSize + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpExecutionMode +template<typename T1, typename T2, typename T3, typename T4> +SpvInst* emitOpExecutionModeLocalSizeId( + SpvInstParent* parent, + IRInst* inst, + const T1& entryPoint, + const T2& xSize, + const T3& ySize, + const T4& zSize +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + static_assert(isSingular<T4>); + return emitInst( + parent, inst, SpvOpExecutionMode, entryPoint, SpvExecutionModeLocalSizeId, xSize, ySize, zSize + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpCapability +SpvInst* emitOpCapability(SpvInstParent* parent, IRInst* inst, SpvCapability capability) +{ + return emitInst(parent, inst, SpvOpCapability, capability); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeVoid +SpvInst* emitOpTypeVoid(IRInst* inst) +{ + return emitInstMemoized( + getSection(SpvLogicalSectionID::ConstantsAndTypes), inst, SpvOpTypeVoid, kResultID + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeBool +SpvInst* emitOpTypeBool(IRInst* inst) +{ + return emitInstMemoized( + getSection(SpvLogicalSectionID::ConstantsAndTypes), inst, SpvOpTypeBool, kResultID + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeInt +SpvInst* emitOpTypeInt( + IRInst* inst, + const SpvLiteralInteger& width, + const SpvLiteralInteger& signedness +) +{ + return emitInstMemoized( + getSection(SpvLogicalSectionID::ConstantsAndTypes), + inst, + SpvOpTypeInt, + kResultID, + width, + signedness + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeFloat +SpvInst* emitOpTypeFloat(IRInst* inst, const SpvLiteralInteger& width) +{ + return emitInstMemoized( + getSection(SpvLogicalSectionID::ConstantsAndTypes), inst, SpvOpTypeFloat, kResultID, width + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeVector +template<typename T> +SpvInst* emitOpTypeVector( + IRInst* inst, + const T& componentType, + const SpvLiteralInteger& componentCount +) +{ + static_assert(isSingular<T>); + return emitInstMemoized( + getSection(SpvLogicalSectionID::ConstantsAndTypes), + inst, + SpvOpTypeVector, + kResultID, + componentType, + componentCount + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeMatrix +template<typename T> +SpvInst* emitOpTypeMatrix(IRInst* inst, const T& columnType, const SpvLiteralInteger& columnCount) +{ + static_assert(isSingular<T>); + return emitInstMemoized( + getSection(SpvLogicalSectionID::ConstantsAndTypes), + inst, + SpvOpTypeMatrix, + kResultID, + columnType, + columnCount + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeImage +template<typename T> +SpvInst* emitOpTypeImage( + IRInst* inst, + const T& sampledType, + SpvDim dim, + const SpvLiteralInteger& depth, + const SpvLiteralInteger& arrayed, + const SpvLiteralInteger& mS, + const SpvLiteralInteger& sampled, + SpvImageFormat imageFormat, + OptionalOperand<SpvAccessQualifier> accessQualifier = SkipThisOptionalOperand{} +) +{ + static_assert(isSingular<T>); + return emitInstMemoized( + getSection(SpvLogicalSectionID::ConstantsAndTypes), + inst, + SpvOpTypeImage, + kResultID, + sampledType, + dim, + depth, + arrayed, + mS, + sampled, + imageFormat, + accessQualifier + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeSampler +SpvInst* emitOpTypeSampler(IRInst* inst) +{ + return emitInstMemoized( + getSection(SpvLogicalSectionID::ConstantsAndTypes), inst, SpvOpTypeSampler, kResultID + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeArray +template<typename T1, typename T2> +SpvInst* emitOpTypeArray(IRInst* inst, const T1& elementType, const T2& length) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInstMemoized( + getSection(SpvLogicalSectionID::ConstantsAndTypes), + inst, + SpvOpTypeArray, + kResultID, + elementType, + length + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeRuntimeArray +template<typename T> +SpvInst* emitOpTypeRuntimeArray(IRInst* inst, const T& elementType) +{ + static_assert(isSingular<T>); + return emitInstMemoized( + getSection(SpvLogicalSectionID::ConstantsAndTypes), + inst, + SpvOpTypeRuntimeArray, + kResultID, + elementType + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeStruct +template<typename Ts> +SpvInst* emitOpTypeStruct(IRInst* inst, const Ts& member0TypeMember1TypeEtc) +{ + static_assert(isPlural<Ts>); + return emitInstMemoized( + getSection(SpvLogicalSectionID::ConstantsAndTypes), + inst, + SpvOpTypeStruct, + kResultID, + member0TypeMember1TypeEtc + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypePointer +template<typename T> +SpvInst* emitOpTypePointer(IRInst* inst, SpvStorageClass storageClass, const T& type) +{ + static_assert(isSingular<T>); + return emitInstMemoized( + getSection(SpvLogicalSectionID::ConstantsAndTypes), + inst, + SpvOpTypePointer, + kResultID, + storageClass, + type + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpTypeFunction +template<typename T, typename Ts> +SpvInst* emitOpTypeFunction( + IRInst* inst, + const T& returnType, + const Ts& parameter0TypeParameter1TypeEtc +) +{ + static_assert(isSingular<T>); + static_assert(isPlural<Ts>); + return emitInstMemoized( + getSection(SpvLogicalSectionID::ConstantsAndTypes), + inst, + SpvOpTypeFunction, + kResultID, + returnType, + parameter0TypeParameter1TypeEtc + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpConstantTrue +template<typename T> +SpvInst* emitOpConstantTrue(IRInst* inst, const T& idResultType) +{ + static_assert(isSingular<T>); + return emitInst( + getSection(SpvLogicalSectionID::ConstantsAndTypes), + inst, + SpvOpConstantTrue, + idResultType, + kResultID + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpConstantFalse +template<typename T> +SpvInst* emitOpConstantFalse(IRInst* inst, const T& idResultType) +{ + static_assert(isSingular<T>); + return emitInst( + getSection(SpvLogicalSectionID::ConstantsAndTypes), + inst, + SpvOpConstantFalse, + idResultType, + kResultID + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpConstant +template<typename T> +SpvInst* emitOpConstant(IRInst* inst, const T& idResultType, const SpvLiteralBits& value) +{ + static_assert(isSingular<T>); + return emitInst( + getSection(SpvLogicalSectionID::ConstantsAndTypes), + inst, + SpvOpConstant, + idResultType, + kResultID, + value + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpConstantComposite +template<typename T, typename Ts> +SpvInst* emitOpConstantComposite(IRInst* inst, const T& idResultType, const Ts& constituents) +{ + static_assert(isSingular<T>); + static_assert(isPlural<Ts>); + return emitInst( + getSection(SpvLogicalSectionID::ConstantsAndTypes), + inst, + SpvOpConstantComposite, + idResultType, + kResultID, + constituents + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpConstantNull +template<typename T> +SpvInst* emitOpConstantNull(IRInst* inst, const T& idResultType) +{ + static_assert(isSingular<T>); + return emitInst( + getSection(SpvLogicalSectionID::ConstantsAndTypes), + inst, + SpvOpConstantNull, + idResultType, + kResultID + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFunction +template<typename T1, typename T2> +SpvInst* emitOpFunction( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + SpvFunctionControlMask functionControl, + const T2& functionType +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst( + parent, inst, SpvOpFunction, idResultType, kResultID, functionControl, functionType + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFunctionParameter +template<typename T> +SpvInst* emitOpFunctionParameter(SpvInstParent* parent, IRInst* inst, const T& idResultType) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpFunctionParameter, idResultType, kResultID); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFunctionEnd +SpvInst* emitOpFunctionEnd(SpvInstParent* parent, IRInst* inst) +{ + return emitInst(parent, inst, SpvOpFunctionEnd); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFunctionCall +template<typename T1, typename T2, typename Ts> +SpvInst* emitOpFunctionCall( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& function, + const Ts& argument0Argument1Etc +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isPlural<Ts>); + return emitInst( + parent, inst, SpvOpFunctionCall, idResultType, kResultID, function, argument0Argument1Etc + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpVariable +template<typename T, typename Opt = SkipThisOptionalOperand> +SpvInst* emitOpVariable( + SpvInstParent* parent, + IRInst* inst, + const T& idResultType, + SpvStorageClass storageClass, + const Opt& initializer = SkipThisOptionalOperand{} +) +{ + static_assert(isSingular<T>); + static_assert(isSingular<Opt>); + return emitInst( + parent, inst, SpvOpVariable, idResultType, kResultID, storageClass, initializer + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpLoad +template<typename T1, typename T2> +SpvInst* emitOpLoad( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& pointer, + OptionalOperand<SpvMemoryAccessMask> memoryAccess = SkipThisOptionalOperand{} +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpLoad, idResultType, kResultID, pointer, memoryAccess); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpLoad +template<typename T1, typename T2> +SpvInst* emitOpLoadAligned( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& pointer, + const SpvLiteralInteger& literalInteger +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst( + parent, + inst, + SpvOpLoad, + idResultType, + kResultID, + pointer, + SpvMemoryAccessAlignedMask, + literalInteger + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpStore +template<typename T1, typename T2> +SpvInst* emitOpStore( + SpvInstParent* parent, + IRInst* inst, + const T1& pointer, + const T2& object, + OptionalOperand<SpvMemoryAccessMask> memoryAccess = SkipThisOptionalOperand{} +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpStore, pointer, object, memoryAccess); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpStore +template<typename T1, typename T2> +SpvInst* emitOpStoreAligned( + SpvInstParent* parent, + IRInst* inst, + const T1& pointer, + const T2& object, + const SpvLiteralInteger& literalInteger +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst( + parent, inst, SpvOpStore, pointer, object, SpvMemoryAccessAlignedMask, literalInteger + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpAccessChain +template<typename T1, typename T2, typename Ts> +SpvInst* emitOpAccessChain( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& base, + const Ts& indexes +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isPlural<Ts>); + return emitInst(parent, inst, SpvOpAccessChain, idResultType, kResultID, base, indexes); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpDecorate +template<typename T> +SpvInst* emitOpDecorate( + SpvInstParent* parent, + IRInst* inst, + const T& target, + SpvDecoration decoration +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpDecorate, target, decoration); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpDecorate +template<typename T> +SpvInst* emitOpDecorateSpecId( + SpvInstParent* parent, + IRInst* inst, + const T& target, + const SpvLiteralInteger& specializationConstantID +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, inst, SpvOpDecorate, target, SpvDecorationSpecId, specializationConstantID + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpDecorate +template<typename T> +SpvInst* emitOpDecorateArrayStride( + SpvInstParent* parent, + IRInst* inst, + const T& target, + const SpvLiteralInteger& arrayStride +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpDecorate, target, SpvDecorationArrayStride, arrayStride); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpDecorate +template<typename T> +SpvInst* emitOpDecorateMatrixStride( + SpvInstParent* parent, + IRInst* inst, + const T& target, + const SpvLiteralInteger& matrixStride +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpDecorate, target, SpvDecorationMatrixStride, matrixStride); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpDecorate +template<typename T> +SpvInst* emitOpDecorateBuiltIn( + SpvInstParent* parent, + IRInst* inst, + const T& target, + SpvBuiltIn builtIn +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpDecorate, target, SpvDecorationBuiltIn, builtIn); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpDecorate +template<typename T1, typename T2> +SpvInst* emitOpDecorateUniformId( + SpvInstParent* parent, + IRInst* inst, + const T1& target, + const T2& execution +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpDecorate, target, SpvDecorationUniformId, execution); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpDecorate +template<typename T> +SpvInst* emitOpDecorateLocation( + SpvInstParent* parent, + IRInst* inst, + const T& target, + const SpvLiteralInteger& location +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpDecorate, target, SpvDecorationLocation, location); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpDecorate +template<typename T> +SpvInst* emitOpDecorateComponent( + SpvInstParent* parent, + IRInst* inst, + const T& target, + const SpvLiteralInteger& component +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpDecorate, target, SpvDecorationComponent, component); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpDecorate +template<typename T> +SpvInst* emitOpDecorateIndex( + SpvInstParent* parent, + IRInst* inst, + const T& target, + const SpvLiteralInteger& index +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpDecorate, target, SpvDecorationIndex, index); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpDecorate +template<typename T> +SpvInst* emitOpDecorateBinding( + SpvInstParent* parent, + IRInst* inst, + const T& target, + const SpvLiteralInteger& bindingPoint +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpDecorate, target, SpvDecorationBinding, bindingPoint); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpDecorate +template<typename T> +SpvInst* emitOpDecorateDescriptorSet( + SpvInstParent* parent, + IRInst* inst, + const T& target, + const SpvLiteralInteger& descriptorSet +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpDecorate, target, SpvDecorationDescriptorSet, descriptorSet); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpDecorate +template<typename T> +SpvInst* emitOpDecorateOffset( + SpvInstParent* parent, + IRInst* inst, + const T& target, + const SpvLiteralInteger& byteOffset +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpDecorate, target, SpvDecorationOffset, byteOffset); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpDecorate +template<typename T> +SpvInst* emitOpDecorateFPRoundingMode( + SpvInstParent* parent, + IRInst* inst, + const T& target, + SpvFPRoundingMode floatingPointRoundingMode +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, inst, SpvOpDecorate, target, SpvDecorationFPRoundingMode, floatingPointRoundingMode + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpDecorate +template<typename T1, typename T2> +SpvInst* emitOpDecorateCounterBuffer( + SpvInstParent* parent, + IRInst* inst, + const T1& target, + const T2& counterBuffer +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpDecorate, target, SpvDecorationCounterBuffer, counterBuffer); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpDecorate +template<typename T> +SpvInst* emitOpDecorateUserSemantic( + SpvInstParent* parent, + IRInst* inst, + const T& target, + const UnownedStringSlice& semantic +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpDecorate, target, SpvDecorationUserSemantic, semantic); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemberDecorate +template<typename T> +SpvInst* emitOpMemberDecorate( + SpvInstParent* parent, + IRInst* inst, + const T& structureType, + const SpvLiteralInteger& member, + SpvDecoration decoration +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpMemberDecorate, structureType, member, decoration); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemberDecorate +template<typename T> +SpvInst* emitOpMemberDecorateSpecId( + SpvInstParent* parent, + IRInst* inst, + const T& structureType, + const SpvLiteralInteger& member, + const SpvLiteralInteger& specializationConstantID +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, + inst, + SpvOpMemberDecorate, + structureType, + member, + SpvDecorationSpecId, + specializationConstantID + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemberDecorate +template<typename T> +SpvInst* emitOpMemberDecorateArrayStride( + SpvInstParent* parent, + IRInst* inst, + const T& structureType, + const SpvLiteralInteger& member, + const SpvLiteralInteger& arrayStride +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, + inst, + SpvOpMemberDecorate, + structureType, + member, + SpvDecorationArrayStride, + arrayStride + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemberDecorate +template<typename T> +SpvInst* emitOpMemberDecorateMatrixStride( + SpvInstParent* parent, + IRInst* inst, + const T& structureType, + const SpvLiteralInteger& member, + const SpvLiteralInteger& matrixStride +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, + inst, + SpvOpMemberDecorate, + structureType, + member, + SpvDecorationMatrixStride, + matrixStride + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemberDecorate +template<typename T> +SpvInst* emitOpMemberDecorateBuiltIn( + SpvInstParent* parent, + IRInst* inst, + const T& structureType, + const SpvLiteralInteger& member, + SpvBuiltIn builtIn +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, inst, SpvOpMemberDecorate, structureType, member, SpvDecorationBuiltIn, builtIn + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemberDecorate +template<typename T1, typename T2> +SpvInst* emitOpMemberDecorateUniformId( + SpvInstParent* parent, + IRInst* inst, + const T1& structureType, + const SpvLiteralInteger& member, + const T2& execution +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst( + parent, inst, SpvOpMemberDecorate, structureType, member, SpvDecorationUniformId, execution + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemberDecorate +template<typename T> +SpvInst* emitOpMemberDecorateLocation( + SpvInstParent* parent, + IRInst* inst, + const T& structureType, + const SpvLiteralInteger& member, + const SpvLiteralInteger& location +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, inst, SpvOpMemberDecorate, structureType, member, SpvDecorationLocation, location + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemberDecorate +template<typename T> +SpvInst* emitOpMemberDecorateComponent( + SpvInstParent* parent, + IRInst* inst, + const T& structureType, + const SpvLiteralInteger& member, + const SpvLiteralInteger& component +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, inst, SpvOpMemberDecorate, structureType, member, SpvDecorationComponent, component + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemberDecorate +template<typename T> +SpvInst* emitOpMemberDecorateIndex( + SpvInstParent* parent, + IRInst* inst, + const T& structureType, + const SpvLiteralInteger& member, + const SpvLiteralInteger& index +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, inst, SpvOpMemberDecorate, structureType, member, SpvDecorationIndex, index + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemberDecorate +template<typename T> +SpvInst* emitOpMemberDecorateBinding( + SpvInstParent* parent, + IRInst* inst, + const T& structureType, + const SpvLiteralInteger& member, + const SpvLiteralInteger& bindingPoint +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, inst, SpvOpMemberDecorate, structureType, member, SpvDecorationBinding, bindingPoint + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemberDecorate +template<typename T> +SpvInst* emitOpMemberDecorateDescriptorSet( + SpvInstParent* parent, + IRInst* inst, + const T& structureType, + const SpvLiteralInteger& member, + const SpvLiteralInteger& descriptorSet +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, + inst, + SpvOpMemberDecorate, + structureType, + member, + SpvDecorationDescriptorSet, + descriptorSet + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemberDecorate +template<typename T> +SpvInst* emitOpMemberDecorateOffset( + SpvInstParent* parent, + IRInst* inst, + const T& structureType, + const SpvLiteralInteger& member, + const SpvLiteralInteger& byteOffset +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, inst, SpvOpMemberDecorate, structureType, member, SpvDecorationOffset, byteOffset + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemberDecorate +template<typename T> +SpvInst* emitOpMemberDecorateFPRoundingMode( + SpvInstParent* parent, + IRInst* inst, + const T& structureType, + const SpvLiteralInteger& member, + SpvFPRoundingMode floatingPointRoundingMode +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, + inst, + SpvOpMemberDecorate, + structureType, + member, + SpvDecorationFPRoundingMode, + floatingPointRoundingMode + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemberDecorate +template<typename T1, typename T2> +SpvInst* emitOpMemberDecorateCounterBuffer( + SpvInstParent* parent, + IRInst* inst, + const T1& structureType, + const SpvLiteralInteger& member, + const T2& counterBuffer +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst( + parent, + inst, + SpvOpMemberDecorate, + structureType, + member, + SpvDecorationCounterBuffer, + counterBuffer + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpMemberDecorate +template<typename T> +SpvInst* emitOpMemberDecorateUserSemantic( + SpvInstParent* parent, + IRInst* inst, + const T& structureType, + const SpvLiteralInteger& member, + const UnownedStringSlice& semantic +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, inst, SpvOpMemberDecorate, structureType, member, SpvDecorationUserSemantic, semantic + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpVectorShuffle +template<typename T1, typename T2, typename T3, Index N> +SpvInst* emitOpVectorShuffle( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& vector1, + const T3& vector2, + const Array<SpvLiteralInteger, N>& components +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst( + parent, inst, SpvOpVectorShuffle, idResultType, kResultID, vector1, vector2, components + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpCompositeConstruct +template<typename T, typename Ts> +SpvInst* emitOpCompositeConstruct( + SpvInstParent* parent, + IRInst* inst, + const T& idResultType, + const Ts& constituents +) +{ + static_assert(isSingular<T>); + static_assert(isPlural<Ts>); + return emitInst(parent, inst, SpvOpCompositeConstruct, idResultType, kResultID, constituents); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpCompositeExtract +template<typename T1, typename T2, Index N> +SpvInst* emitOpCompositeExtract( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& composite, + const Array<SpvLiteralInteger, N>& indexes +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst( + parent, inst, SpvOpCompositeExtract, idResultType, kResultID, composite, indexes + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpCopyObject +template<typename T1, typename T2> +SpvInst* emitOpCopyObject( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpCopyObject, idResultType, kResultID, operand); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpConvertFToU +template<typename T1, typename T2> +SpvInst* emitOpConvertFToU( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& floatValue +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpConvertFToU, idResultType, kResultID, floatValue); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpConvertFToS +template<typename T1, typename T2> +SpvInst* emitOpConvertFToS( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& floatValue +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpConvertFToS, idResultType, kResultID, floatValue); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpConvertSToF +template<typename T1, typename T2> +SpvInst* emitOpConvertSToF( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& signedValue +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpConvertSToF, idResultType, kResultID, signedValue); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpConvertUToF +template<typename T1, typename T2> +SpvInst* emitOpConvertUToF( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& unsignedValue +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpConvertUToF, idResultType, kResultID, unsignedValue); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpUConvert +template<typename T1, typename T2> +SpvInst* emitOpUConvert( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& unsignedValue +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpUConvert, idResultType, kResultID, unsignedValue); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpSConvert +template<typename T1, typename T2> +SpvInst* emitOpSConvert( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& signedValue +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpSConvert, idResultType, kResultID, signedValue); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFConvert +template<typename T1, typename T2> +SpvInst* emitOpFConvert( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& floatValue +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpFConvert, idResultType, kResultID, floatValue); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpBitcast +template<typename T1, typename T2> +SpvInst* emitOpBitcast( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpBitcast, idResultType, kResultID, operand); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpSNegate +template<typename T1, typename T2> +SpvInst* emitOpSNegate( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpSNegate, idResultType, kResultID, operand); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFNegate +template<typename T1, typename T2> +SpvInst* emitOpFNegate( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpFNegate, idResultType, kResultID, operand); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpIAdd +template<typename T1, typename T2, typename T3> +SpvInst* emitOpIAdd( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpIAdd, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFAdd +template<typename T1, typename T2, typename T3> +SpvInst* emitOpFAdd( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpFAdd, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpISub +template<typename T1, typename T2, typename T3> +SpvInst* emitOpISub( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpISub, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFSub +template<typename T1, typename T2, typename T3> +SpvInst* emitOpFSub( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpFSub, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpIMul +template<typename T1, typename T2, typename T3> +SpvInst* emitOpIMul( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpIMul, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFMul +template<typename T1, typename T2, typename T3> +SpvInst* emitOpFMul( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpFMul, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpUDiv +template<typename T1, typename T2, typename T3> +SpvInst* emitOpUDiv( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpUDiv, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpSDiv +template<typename T1, typename T2, typename T3> +SpvInst* emitOpSDiv( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpSDiv, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFDiv +template<typename T1, typename T2, typename T3> +SpvInst* emitOpFDiv( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpFDiv, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpUMod +template<typename T1, typename T2, typename T3> +SpvInst* emitOpUMod( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpUMod, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpSRem +template<typename T1, typename T2, typename T3> +SpvInst* emitOpSRem( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpSRem, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFRem +template<typename T1, typename T2, typename T3> +SpvInst* emitOpFRem( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpFRem, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpIAddCarry +template<typename T1, typename T2, typename T3> +SpvInst* emitOpIAddCarry( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpIAddCarry, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpISubBorrow +template<typename T1, typename T2, typename T3> +SpvInst* emitOpISubBorrow( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpISubBorrow, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpLogicalEqual +template<typename T1, typename T2, typename T3> +SpvInst* emitOpLogicalEqual( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpLogicalEqual, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpLogicalNotEqual +template<typename T1, typename T2, typename T3> +SpvInst* emitOpLogicalNotEqual( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst( + parent, inst, SpvOpLogicalNotEqual, idResultType, kResultID, operand1, operand2 + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpLogicalOr +template<typename T1, typename T2, typename T3> +SpvInst* emitOpLogicalOr( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpLogicalOr, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpLogicalAnd +template<typename T1, typename T2, typename T3> +SpvInst* emitOpLogicalAnd( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpLogicalAnd, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpLogicalNot +template<typename T1, typename T2> +SpvInst* emitOpLogicalNot( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpLogicalNot, idResultType, kResultID, operand); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpIEqual +template<typename T1, typename T2, typename T3> +SpvInst* emitOpIEqual( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpIEqual, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpINotEqual +template<typename T1, typename T2, typename T3> +SpvInst* emitOpINotEqual( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpINotEqual, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpUGreaterThan +template<typename T1, typename T2, typename T3> +SpvInst* emitOpUGreaterThan( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpUGreaterThan, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpSGreaterThan +template<typename T1, typename T2, typename T3> +SpvInst* emitOpSGreaterThan( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpSGreaterThan, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpUGreaterThanEqual +template<typename T1, typename T2, typename T3> +SpvInst* emitOpUGreaterThanEqual( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst( + parent, inst, SpvOpUGreaterThanEqual, idResultType, kResultID, operand1, operand2 + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpSGreaterThanEqual +template<typename T1, typename T2, typename T3> +SpvInst* emitOpSGreaterThanEqual( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst( + parent, inst, SpvOpSGreaterThanEqual, idResultType, kResultID, operand1, operand2 + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpULessThan +template<typename T1, typename T2, typename T3> +SpvInst* emitOpULessThan( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpULessThan, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpSLessThan +template<typename T1, typename T2, typename T3> +SpvInst* emitOpSLessThan( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpSLessThan, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpULessThanEqual +template<typename T1, typename T2, typename T3> +SpvInst* emitOpULessThanEqual( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpULessThanEqual, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpSLessThanEqual +template<typename T1, typename T2, typename T3> +SpvInst* emitOpSLessThanEqual( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpSLessThanEqual, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFOrdEqual +template<typename T1, typename T2, typename T3> +SpvInst* emitOpFOrdEqual( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpFOrdEqual, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFOrdNotEqual +template<typename T1, typename T2, typename T3> +SpvInst* emitOpFOrdNotEqual( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpFOrdNotEqual, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFOrdLessThan +template<typename T1, typename T2, typename T3> +SpvInst* emitOpFOrdLessThan( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpFOrdLessThan, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFOrdGreaterThan +template<typename T1, typename T2, typename T3> +SpvInst* emitOpFOrdGreaterThan( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst( + parent, inst, SpvOpFOrdGreaterThan, idResultType, kResultID, operand1, operand2 + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFOrdLessThanEqual +template<typename T1, typename T2, typename T3> +SpvInst* emitOpFOrdLessThanEqual( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst( + parent, inst, SpvOpFOrdLessThanEqual, idResultType, kResultID, operand1, operand2 + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpFOrdGreaterThanEqual +template<typename T1, typename T2, typename T3> +SpvInst* emitOpFOrdGreaterThanEqual( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst( + parent, inst, SpvOpFOrdGreaterThanEqual, idResultType, kResultID, operand1, operand2 + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpShiftRightLogical +template<typename T1, typename T2, typename T3> +SpvInst* emitOpShiftRightLogical( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& base, + const T3& shift +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpShiftRightLogical, idResultType, kResultID, base, shift); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpShiftRightArithmetic +template<typename T1, typename T2, typename T3> +SpvInst* emitOpShiftRightArithmetic( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& base, + const T3& shift +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpShiftRightArithmetic, idResultType, kResultID, base, shift); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpShiftLeftLogical +template<typename T1, typename T2, typename T3> +SpvInst* emitOpShiftLeftLogical( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& base, + const T3& shift +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpShiftLeftLogical, idResultType, kResultID, base, shift); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpBitwiseOr +template<typename T1, typename T2, typename T3> +SpvInst* emitOpBitwiseOr( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpBitwiseOr, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpBitwiseXor +template<typename T1, typename T2, typename T3> +SpvInst* emitOpBitwiseXor( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpBitwiseXor, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpBitwiseAnd +template<typename T1, typename T2, typename T3> +SpvInst* emitOpBitwiseAnd( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& operand1, + const T3& operand2 +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst(parent, inst, SpvOpBitwiseAnd, idResultType, kResultID, operand1, operand2); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpBitReverse +template<typename T1, typename T2> +SpvInst* emitOpBitReverse( + SpvInstParent* parent, + IRInst* inst, + const T1& idResultType, + const T2& base +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpBitReverse, idResultType, kResultID, base); +} + +// OpPhi elided, please use emitInst directly + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpLoopMerge +template<typename T1, typename T2> +SpvInst* emitOpLoopMerge( + SpvInstParent* parent, + IRInst* inst, + const T1& mergeBlock, + const T2& continueTarget, + SpvLoopControlMask loopControl +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst(parent, inst, SpvOpLoopMerge, mergeBlock, continueTarget, loopControl); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpLoopMerge +template<typename T1, typename T2> +SpvInst* emitOpLoopMergeDependencyLength( + SpvInstParent* parent, + IRInst* inst, + const T1& mergeBlock, + const T2& continueTarget, + const SpvLiteralInteger& literalInteger +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst( + parent, + inst, + SpvOpLoopMerge, + mergeBlock, + continueTarget, + SpvLoopControlDependencyLengthMask, + literalInteger + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpLoopMerge +template<typename T1, typename T2> +SpvInst* emitOpLoopMergeMinIterations( + SpvInstParent* parent, + IRInst* inst, + const T1& mergeBlock, + const T2& continueTarget, + const SpvLiteralInteger& literalInteger +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst( + parent, + inst, + SpvOpLoopMerge, + mergeBlock, + continueTarget, + SpvLoopControlMinIterationsMask, + literalInteger + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpLoopMerge +template<typename T1, typename T2> +SpvInst* emitOpLoopMergeMaxIterations( + SpvInstParent* parent, + IRInst* inst, + const T1& mergeBlock, + const T2& continueTarget, + const SpvLiteralInteger& literalInteger +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst( + parent, + inst, + SpvOpLoopMerge, + mergeBlock, + continueTarget, + SpvLoopControlMaxIterationsMask, + literalInteger + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpLoopMerge +template<typename T1, typename T2> +SpvInst* emitOpLoopMergeIterationMultiple( + SpvInstParent* parent, + IRInst* inst, + const T1& mergeBlock, + const T2& continueTarget, + const SpvLiteralInteger& literalInteger +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst( + parent, + inst, + SpvOpLoopMerge, + mergeBlock, + continueTarget, + SpvLoopControlIterationMultipleMask, + literalInteger + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpLoopMerge +template<typename T1, typename T2> +SpvInst* emitOpLoopMergePeelCount( + SpvInstParent* parent, + IRInst* inst, + const T1& mergeBlock, + const T2& continueTarget, + const SpvLiteralInteger& literalInteger +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst( + parent, + inst, + SpvOpLoopMerge, + mergeBlock, + continueTarget, + SpvLoopControlPeelCountMask, + literalInteger + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpLoopMerge +template<typename T1, typename T2> +SpvInst* emitOpLoopMergePartialCount( + SpvInstParent* parent, + IRInst* inst, + const T1& mergeBlock, + const T2& continueTarget, + const SpvLiteralInteger& literalInteger +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + return emitInst( + parent, + inst, + SpvOpLoopMerge, + mergeBlock, + continueTarget, + SpvLoopControlPartialCountMask, + literalInteger + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpSelectionMerge +template<typename T> +SpvInst* emitOpSelectionMerge( + SpvInstParent* parent, + IRInst* inst, + const T& mergeBlock, + SpvSelectionControlMask selectionControl +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpSelectionMerge, mergeBlock, selectionControl); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpLabel +SpvInst* emitOpLabel(SpvInstParent* parent, IRInst* inst) +{ + return emitInst(parent, inst, SpvOpLabel, kResultID); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpBranch +template<typename T> +SpvInst* emitOpBranch(SpvInstParent* parent, IRInst* inst, const T& targetLabel) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpBranch, targetLabel); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpBranchConditional +template<typename T1, typename T2, typename T3, Index N> +SpvInst* emitOpBranchConditional( + SpvInstParent* parent, + IRInst* inst, + const T1& condition, + const T2& trueLabel, + const T3& falseLabel, + const Array<SpvLiteralInteger, N>& branchWeights +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + return emitInst( + parent, inst, SpvOpBranchConditional, condition, trueLabel, falseLabel, branchWeights + ); +} + +// OpSwitch elided, please use emitInst directly + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpKill +SpvInst* emitOpKill(SpvInstParent* parent, IRInst* inst) +{ + return emitInst(parent, inst, SpvOpKill); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpReturn +SpvInst* emitOpReturn(SpvInstParent* parent, IRInst* inst) +{ + return emitInst(parent, inst, SpvOpReturn); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpReturnValue +template<typename T> +SpvInst* emitOpReturnValue(SpvInstParent* parent, IRInst* inst, const T& value) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpReturnValue, value); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpUnreachable +SpvInst* emitOpUnreachable(SpvInstParent* parent, IRInst* inst) +{ + return emitInst(parent, inst, SpvOpUnreachable); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpExecutionModeId +template<typename T> +SpvInst* emitOpExecutionModeId( + SpvInstParent* parent, + IRInst* inst, + const T& entryPoint, + SpvExecutionMode mode +) +{ + static_assert(isSingular<T>); + return emitInst(parent, inst, SpvOpExecutionModeId, entryPoint, mode); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpExecutionModeId +template<typename T> +SpvInst* emitOpExecutionModeIdLocalSize( + SpvInstParent* parent, + IRInst* inst, + const T& entryPoint, + const SpvLiteralInteger& xSize, + const SpvLiteralInteger& ySize, + const SpvLiteralInteger& zSize +) +{ + static_assert(isSingular<T>); + return emitInst( + parent, inst, SpvOpExecutionModeId, entryPoint, SpvExecutionModeLocalSize, xSize, ySize, zSize + ); +} + +// https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpExecutionModeId +template<typename T1, typename T2, typename T3, typename T4> +SpvInst* emitOpExecutionModeIdLocalSizeId( + SpvInstParent* parent, + IRInst* inst, + const T1& entryPoint, + const T2& xSize, + const T3& ySize, + const T4& zSize +) +{ + static_assert(isSingular<T1>); + static_assert(isSingular<T2>); + static_assert(isSingular<T3>); + static_assert(isSingular<T4>); + return emitInst( + parent, + inst, + SpvOpExecutionModeId, + entryPoint, + SpvExecutionModeLocalSizeId, + xSize, + ySize, + zSize + ); +} +#endif // SLANG_IN_SPIRV_EMIT_CONTEXT diff --git a/source/slang/slang-emit-spirv.cpp b/source/slang/slang-emit-spirv.cpp index bafbb79f1..3947e8468 100644 --- a/source/slang/slang-emit-spirv.cpp +++ b/source/slang/slang-emit-spirv.cpp @@ -12,6 +12,7 @@ #include "slang-spirv-val.h" #include "spirv/unified1/spirv.h" #include "../core/slang-memory-arena.h" +#include <type_traits> namespace Slang { @@ -275,6 +276,80 @@ struct SpvSnippetEmitContext List<SpvWord> argumentIds; }; +// A structure which can hold an integer literal, either one word or several +struct SpvLiteralInteger +{ + static SpvLiteralInteger from32(int32_t value) { return from32(uint32_t(value)); } + static SpvLiteralInteger from32(uint32_t value) { return SpvLiteralInteger{{value}}; } + static SpvLiteralInteger from64(int64_t value) { return from64(uint64_t(value)); } + static SpvLiteralInteger from64(uint64_t value) { return SpvLiteralInteger{{SpvWord(value), SpvWord(value >> 32)}}; } + List<SpvWord> value; // Words, stored low words to high (TODO, SmallArray or something here) +}; + +// A structure which can hold bitwise literal, either one word or several +struct SpvLiteralBits +{ + static SpvLiteralBits from32(uint32_t value) { return SpvLiteralBits{{value}}; } + static SpvLiteralBits from64(uint64_t value) { return SpvLiteralBits{{SpvWord(value), SpvWord(value >> 32)}}; } + List<SpvWord> value; // Words, stored low words to high (TODO, SmallArray or something here) +}; + +// As a convenience, there are often cases where +// we will want to emit all of the operands of some +// IR instruction as <id> operands of a SPIR-V +// instruction. This is handy in cases where the +// Slang IR and SPIR-V instructions agree on the +// number, order, and meaning of their operands. +/// Helper type for emitting all the operands of the current IR instruction +struct OperandsOf +{ + OperandsOf(IRInst* irInst) + : irInst(irInst) + {} + + IRInst* irInst = nullptr; +}; + +/// Helper type for not emitting an operand in this position +struct SkipThisOptionalOperand {}; + +template<typename T> +struct OptionalOperand +{ + static_assert(std::is_trivial_v<T>); + OptionalOperand(SkipThisOptionalOperand) : present(false) {} + OptionalOperand(T value) : present(true), value(value) {} + bool present; + T value; +}; + +template<typename T> +OptionalOperand<T> nullOptionOperand() +{ + return OptionalOperand<T>{false}; +} + +template<typename T> +OptionalOperand<T> someOptionOperand(T t) +{ + return OptionalOperand<T>{true, t}; +} + +template<typename T> +constexpr bool isPlural = false; +template<typename T> +constexpr bool isPlural<List<T>> = true; +template<typename T> +constexpr bool isPlural<IROperandList<T>> = true; +template<typename T, Index N> +constexpr bool isPlural<Array<T, N>> = true; +template<> +constexpr bool isPlural<OperandsOf> = true; +template<> +constexpr bool isPlural<IRUse*> = true; +template<typename T> +constexpr bool isSingular = !isPlural<T>; + // Now that we've defined the intermediate data structures we will // use to represent SPIR-V code during emission, we will move on // to defining the main context type that will drive SPIR-V @@ -471,10 +546,12 @@ struct SPIRVEmitContext // Holds a stack of instructions operands *BEFORE* they added to the instruction. List<SpvWord> m_operandStack; - // The current instruction being constructed. Cannot add operands unless it is set. + // The current instruction being constructed. Cannot add operands unless it + // is set, or we are peeking at some operands to see if we have them memoized SpvInst* m_currentInst = nullptr; + bool m_peekingOperands = false; - // Operands can only be added when inside of a InstConstructScope + // Operands can only be added when inside of a InstConstructScope or... struct InstConstructScope { SLANG_FORCE_INLINE operator SpvInst*() const { return m_inst; } @@ -495,6 +572,28 @@ struct SPIRVEmitContext Index m_operandsStartIndex; ///< The start index for operands of m_inst }; + // ...If we're speculatively adding them to see if we have a memoized results + struct OperandMemoizeScope + { + OperandMemoizeScope(SPIRVEmitContext* context) : m_context(context) + { + m_tmpOperandStack.swapWith(m_context->m_operandStack); + std::swap(m_tmpPeeking, m_context->m_peekingOperands); + std::swap(m_tmpInst, m_context->m_currentInst); + } + ~OperandMemoizeScope() + { + std::swap(m_tmpInst, m_context->m_currentInst); + std::swap(m_tmpPeeking, m_context->m_peekingOperands); + m_tmpOperandStack.swapWith(m_context->m_operandStack); + } + + SPIRVEmitContext* m_context; + List<SpvWord> m_tmpOperandStack; + bool m_tmpPeeking = true; + SpvInst* m_tmpInst = nullptr; + }; + /// Holds memory for instructions and operands. MemoryArena m_memoryArena; @@ -586,7 +685,7 @@ struct SPIRVEmitContext void emitOperand(SpvWord word) { // Can only add operands if we are constructing an instruction (ie in _beginInst/_endInst) - SLANG_ASSERT(m_currentInst); + SLANG_ASSERT(m_currentInst || m_peekingOperands); m_operandStack.add(word); } @@ -623,7 +722,7 @@ struct SPIRVEmitContext void emitOperand(UnownedStringSlice const& text) { // Can only emitOperands if we are in an instruction - SLANG_ASSERT(m_currentInst); + SLANG_ASSERT(m_currentInst || m_peekingOperands); SLANG_COMPILE_TIME_ASSERT(sizeof(SpvWord) == 4); // Assert that `text` doesn't contain any embedded nul bytes, since they @@ -670,16 +769,53 @@ struct SPIRVEmitContext void emitOperand(ResultIDToken) { + // This is the one case we shouldn't be peeking at operands, as it + // depends on having an instruction under construction SLANG_ASSERT(m_currentInst); // A result <id> operand uses the <id> of the instruction itself (which is m_currentInst) emitOperand(getID(m_currentInst)); } - void emitOperand(SpvDecoration decoration) { emitOperand((SpvWord)decoration); } + void emitOperand(const SpvLiteralBits& bits) + { + for(const auto v : bits.value) + emitOperand(v); + } - void emitOperand(SpvBuiltIn builtin) { emitOperand((SpvWord)builtin); } - void emitOperand(SpvStorageClass val) { emitOperand((SpvWord)val); } + void emitOperand(const SpvLiteralInteger& integer) + { + for(const auto v : integer.value) + emitOperand(v); + } + + template<typename T> + void emitOperand(const List<T>& os) + { + for(const auto& o : os) + emitOperand(o); + } + + template<typename T> + void emitOperand(const IROperandList<T>& os) + { + for(const auto& o : os) + emitOperand(o); + } + + template<typename T, Index N> + void emitOperand(const Array<T, N>& os) + { + for(const auto& o : os) + emitOperand(o); + } + + template<typename T> + void emitOperand(const ArrayView<T>& os) + { + for(const auto& o : os) + emitOperand(o); + } template<typename TConstant> struct ConstantValueKey @@ -697,7 +833,7 @@ struct SPIRVEmitContext }; Dictionary<ConstantValueKey<IRIntegerValue>, SpvInst*> m_spvIntConstants; Dictionary<ConstantValueKey<IRFloatingPointValue>, SpvInst*> m_spvFloatConstants; - SpvInst* emitIntConstant(IRIntegerValue val, IRType* type) + SpvInst* emitIntConstant(IRIntegerValue val, IRType* type, IRInst* inst = nullptr) { ConstantValueKey<IRIntegerValue> key; key.value = val; @@ -705,8 +841,6 @@ struct SPIRVEmitContext SpvInst* result = nullptr; if (m_spvIntConstants.tryGetValue(key, result)) return result; - SpvWord valWord; - memcpy(&valWord, &val, sizeof(SpvWord)); switch (type->getOp()) { case kIROp_Int64Type: @@ -716,34 +850,27 @@ struct SPIRVEmitContext case kIROp_UIntPtrType: #endif { - SpvWord valHighWord; - memcpy(&valHighWord, (char*)(&val) + 4, sizeof(SpvWord)); - result = emitInst( - getSection(SpvLogicalSectionID::ConstantsAndTypes), - nullptr, - SpvOpConstant, + result = emitOpConstant( + inst, type, - kResultID, - valWord, - valHighWord); + SpvLiteralBits::from64(uint64_t(val)) + ); break; } default: { - result = emitInst( - getSection(SpvLogicalSectionID::ConstantsAndTypes), - nullptr, - SpvOpConstant, + result = emitOpConstant( + inst, type, - kResultID, - valWord); + SpvLiteralBits::from32(uint32_t(val)) + ); break; } } m_spvIntConstants[key] = result; return result; } - SpvInst* emitFloatConstant(IRFloatingPointValue val, IRType* type) + SpvInst* emitFloatConstant(IRFloatingPointValue val, IRType* type, IRInst* inst = nullptr) { ConstantValueKey<IRFloatingPointValue> key; key.value = val; @@ -751,50 +878,34 @@ struct SPIRVEmitContext SpvInst* result = nullptr; if (m_spvFloatConstants.tryGetValue(key, result)) return result; - SpvWord valWord; - memcpy(&valWord, &val, sizeof(SpvWord)); if (type->getOp() == kIROp_DoubleType) { - SpvWord valHighWord; - memcpy(&valHighWord, (char*)(&val) + 4, sizeof(SpvWord)); - result = emitInst( - getSection(SpvLogicalSectionID::ConstantsAndTypes), - nullptr, - SpvOpConstant, + result = emitOpConstant( + inst, type, - kResultID, - valWord, - valHighWord); + SpvLiteralBits::from64(uint64_t(DoubleAsInt64(val)))); } - else + else if(type->getOp() == kIROp_FloatType) { - result = emitInst( - getSection(SpvLogicalSectionID::ConstantsAndTypes), - nullptr, - SpvOpConstant, + result = emitOpConstant( + inst, + type, + SpvLiteralBits::from32(uint32_t(FloatAsInt(float(val))))); + } + else if(type->getOp() == kIROp_HalfType) + { + result = emitOpConstant( + inst, type, - kResultID, - valWord); + SpvLiteralBits::from32(uint32_t(FloatToHalf(float(val))))); + } + else + { + SLANG_UNEXPECTED("missing case in SPIR-V emitFloatConstant"); } m_spvFloatConstants[key] = result; return result; } - // As another convenience, there are often cases where - // we will want to emit all of the operands of some - // IR instruction as <id> operands of a SPIR-V - // instruction. This is handy in cases where the - // Slang IR and SPIR-V instructions agree on the - // number, order, and meaning of their operands. - - /// Helper type for emitting all the operands of the current IR instruction - struct OperandsOf - { - OperandsOf(IRInst* irInst) - : irInst(irInst) - {} - - IRInst* irInst = nullptr; - }; /// Emit operand words for all the operands of a given IR instruction void emitOperand(OperandsOf const& other) @@ -807,101 +918,104 @@ struct SPIRVEmitContext } } + /// Do nothing + void emitOperand(SkipThisOptionalOperand) { } + + template<typename T> + void emitOperand(OptionalOperand<T> o) + { + if(o.present) + emitOperand(o.value); + } + // With the above routines, code can easily construct a SPIR-V // instruction with arbitrary operands over multiple lines of code. // - // In many cases, however, it is desirable to be able to emit - // an instruction more compactly, and for that we will introduce - // a number of `emitInst()` helpers that handle creating an - // instruction, filling in its operands, and adding it to a parent. + // The safe way to call these routines is encoded in the below `emitInst` + // function. // - // These routines are overloaded on the number of operands, and - // also templates to work with any of the types for which - // `emitOperand()` works. + // This allows one to generically output a SPIR-V instruction with any + // desired operands. // - // In all of these cases, the caller takes responsibility for - // correctly matching the SPIR-V encoding rules for the chosen - // opcode, including whether a type <id> or result <id> is - // required. - - SpvInst* emitInst(SpvInstParent* parent, IRInst* irInst, SpvOp opcode) + // This function performs no checks that it is actually being used + // correctly with respect to the SPIR-V rules for each opcode. As such, a + // more type safe function for each opcode is included in + // 'slang-emit-spirv-ops.h', and available in this class. You are + // encouraged to use these instead. + // + template<typename... Operands> + SpvInst* emitInst(SpvInstParent* parent, IRInst* irInst, SpvOp opcode, const Operands& ...ops) { InstConstructScope scopeInst(this, opcode, irInst); SpvInst* spvInst = scopeInst; + (emitOperand(ops), ...); parent->addInst(spvInst); return spvInst; } - template<typename A> - SpvInst* emitInst(SpvInstParent* parent, IRInst* irInst, SpvOp opcode, A const& a) + template<typename OperandEmitFunc> + SpvInst* emitInstCustomOperandFunc(SpvInstParent* parent, IRInst* irInst, SpvOp opcode, const OperandEmitFunc& f) { InstConstructScope scopeInst(this, opcode, irInst); SpvInst* spvInst = scopeInst; - emitOperand(a); + f(); parent->addInst(spvInst); return spvInst; } - template<typename A, typename B> - SpvInst* emitInst(SpvInstParent* parent, IRInst* irInst, SpvOp opcode, A const& a, B const& b) + // Emits a SPV Inst with deduplication + // This is used where our IR doesn't guarantee uniqueness but SPIR-V + // requires it + template<typename... Operands> + SpvInst* emitInstMemoized( + SpvInstParent* parent, + IRInst* irInst, + SpvOp opcode, + // We take the resultId here explicitly here to make sure we don't try + // and memoize its value. + ResultIDToken resultId, + const Operands& ...ops + ) { - InstConstructScope scopeInst(this, opcode, irInst); - SpvInst* spvInst = scopeInst; - emitOperand(a); - emitOperand(b); - parent->addInst(spvInst); - return spvInst; - } + List<SpvWord> ourOperands; + { + auto scopePeek = OperandMemoizeScope(this); + (emitOperand(ops), ...); + // Steal our operands back, so we don't have to calculate them + // again + ourOperands = std::move(m_operandStack); + } - template<typename A, typename B, typename C> - SpvInst* emitInst(SpvInstParent* parent, IRInst* irInst, SpvOp opcode, A const& a, B const& b, C const& c) - { - InstConstructScope scopeInst(this, opcode, irInst); - SpvInst* spvInst = scopeInst; - emitOperand(a); - emitOperand(b); - emitOperand(c); - parent->addInst(spvInst); - return spvInst; - } + // Hash the whole global stack and opcode + SpvTypeInstKey key; + key.words.add(opcode); + key.words.addRange(ourOperands); - template<typename A, typename B, typename C, typename D> - SpvInst* emitInst(SpvInstParent* parent, IRInst* irInst, SpvOp opcode, A const& a, B const& b, C const& c, D const& d) - { - InstConstructScope scopeInst(this, opcode, irInst); - SpvInst* spvInst = scopeInst; - emitOperand(a); - emitOperand(b); - emitOperand(c); - emitOperand(d); - parent->addInst(spvInst); - return spvInst; - } + // If we have seen this before, return the memoized instruction + if (SpvInst** memoized = m_spvTypeInsts.tryGetValue(key)) + return *memoized; - template<typename A, typename B, typename C, typename D, typename E> - SpvInst* emitInst(SpvInstParent* parent, IRInst* irInst, SpvOp opcode, A const& a, B const& b, C const& c, D const& d, E const& e) - { + // Otherwise, we can construct our instruction and record the result InstConstructScope scopeInst(this, opcode, irInst); SpvInst* spvInst = scopeInst; - emitOperand(a); - emitOperand(b); - emitOperand(c); - emitOperand(d); - emitOperand(e); - parent->addInst(spvInst); - return spvInst; - } + m_spvTypeInsts[key] = spvInst; + + // Emit our operands, this time with the resultId too + emitOperand(resultId); + m_operandStack.addRange(ourOperands); - template<typename OperandEmitFunc> - SpvInst* emitInstCustomOperandFunc(SpvInstParent* parent, IRInst* irInst, SpvOp opcode, const OperandEmitFunc& f) - { - InstConstructScope scopeInst(this, opcode, irInst); - SpvInst* spvInst = scopeInst; - f(); parent->addInst(spvInst); return spvInst; } + // + // Specific emit funcs + // + +# define SLANG_IN_SPIRV_EMIT_CONTEXT +# include "slang-emit-spirv-ops.h" +# undef SLANG_IN_SPIRV_EMIT_CONTEXT + /// The SPIRV OpExtInstImport inst that represents the GLSL450 /// extended instruction set. SpvInst* m_glsl450ExtInst = nullptr; @@ -910,11 +1024,9 @@ struct SPIRVEmitContext { if (m_glsl450ExtInst) return m_glsl450ExtInst; - m_glsl450ExtInst = emitInst( + m_glsl450ExtInst = emitOpExtInstImport( getSection(SpvLogicalSectionID::ExtIntInstImports), nullptr, - SpvOpExtInstImport, - kResultID, UnownedStringSlice("GLSL.std.450")); return m_glsl450ExtInst; } @@ -937,7 +1049,11 @@ struct SPIRVEmitContext // For now we will always emit the `Shader` capability, // since every Vulkan shader module will use it. // - emitInst(getSection(SpvLogicalSectionID::Capabilities), nullptr, SpvOpCapability, SpvCapabilityShader); + emitOpCapability( + getSection(SpvLogicalSectionID::Capabilities), + nullptr, + SpvCapabilityShader + ); // [2.4: Logical Layout of a Module] // @@ -953,7 +1069,12 @@ struct SPIRVEmitContext // a requirement, but it is what glslang produces, // so we will use it for now. // - emitInst(getSection(SpvLogicalSectionID::MemoryModel), nullptr, SpvOpMemoryModel, SpvAddressingModelLogical, SpvMemoryModelGLSL450); + emitOpMemoryModel( + getSection(SpvLogicalSectionID::MemoryModel), + nullptr, + SpvAddressingModelLogical, + SpvMemoryModelGLSL450 + ); } Dictionary<UnownedStringSlice, SpvInst*> m_extensionInsts; @@ -962,8 +1083,11 @@ struct SPIRVEmitContext SpvInst* result = nullptr; if (m_extensionInsts.tryGetValue(name, result)) return result; - result = - emitInst(getSection(SpvLogicalSectionID::Extensions), nullptr, SpvOpExtension, name); + result = emitOpExtension( + getSection(SpvLogicalSectionID::Extensions), + nullptr, + name + ); m_extensionInsts[name] = result; return result; } @@ -983,31 +1107,6 @@ struct SPIRVEmitContext Dictionary<SpvTypeInstKey, SpvInst*> m_spvTypeInsts; - // Emits a SPV Inst that represents a type, with deduplications since - // our IR doesn't currently guarantee types are unique in generated SPV. - SpvInst* emitTypeInst(IRInst* typeInst, SpvOp opcode, ArrayView<SpvWord> operands) - { - SpvTypeInstKey key; - key.words.add((SpvWord)opcode); - for (auto op : operands) - key.words.add(op); - SpvInst* result = nullptr; - if (m_spvTypeInsts.tryGetValue(key, result)) - { - return result; - } - result = emitInstCustomOperandFunc( - getSection(SpvLogicalSectionID::ConstantsAndTypes), typeInst, opcode, [&]() { - emitOperand(kResultID); - for (auto op : operands) - { - emitOperand(op); - } - }); - m_spvTypeInsts[key] = result; - return result; - } - // Next, let's look at emitting some of the instructions // that can occur at global scope. @@ -1017,21 +1116,13 @@ struct SPIRVEmitContext /// SpvInst* emitGlobalInst(IRInst* inst) { - switch( inst->getOp() ) + switch( inst->getOp() & kIROpMask_OpMask ) { // [3.32.6: Type-Declaration Instructions] // -#define CASE(IROP, SPVOP) \ - case IROP: return emitTypeInst(inst, SPVOP, ArrayView<SpvWord>()); - - // > OpTypeVoid - CASE(kIROp_VoidType, SpvOpTypeVoid); - - // > OpTypeBool - CASE(kIROp_BoolType, SpvOpTypeBool); - -#undef CASE + case kIROp_VoidType: return emitOpTypeVoid(inst); + case kIROp_BoolType: return emitOpTypeBool(inst); // > OpTypeInt @@ -1045,10 +1136,11 @@ struct SPIRVEmitContext case kIROp_Int64Type: { const IntInfo i = getIntTypeInfo(as<IRType>(inst)); - return emitTypeInst( + return emitOpTypeInt( inst, - SpvOpTypeInt, - makeArray(static_cast<SpvWord>(i.width), SpvWord{i.isSigned}).getView()); + SpvLiteralInteger::from32(int32_t(i.width)), + SpvLiteralInteger::from32(i.isSigned) + ); } // > OpTypeFloat @@ -1058,7 +1150,7 @@ struct SPIRVEmitContext case kIROp_DoubleType: { const FloatInfo i = getFloatingTypeInfo(as<IRType>(inst)); - return emitTypeInst(inst, SpvOpTypeFloat, makeArray(static_cast<SpvWord>(i.width)).getView()); + return emitOpTypeFloat(inst, SpvLiteralInteger::from32(int32_t(i.width))); } case kIROp_PtrType: @@ -1073,22 +1165,24 @@ struct SPIRVEmitContext storageClass = (SpvStorageClass)ptrType->getAddressSpace(); if (storageClass == SpvStorageClassStorageBuffer) ensureExtensionDeclaration(UnownedStringSlice("SPV_KHR_storage_buffer_storage_class")); - auto operands = makeArray<SpvWord>( - (SpvWord)storageClass, getID(ensureInst(inst->getOperand(0)))); - return emitTypeInst( - inst, SpvOpTypePointer, operands.getView()); + return emitOpTypePointer( + inst, + storageClass, + inst->getOperand(0) + ); } + case kIROp_ConstantBufferType: + SLANG_UNEXPECTED("Constant buffer type remaining in spirv emit"); case kIROp_StructType: { - auto spvStructType = emitInstCustomOperandFunc( - getSection(SpvLogicalSectionID::ConstantsAndTypes), inst, SpvOpTypeStruct, [&]() { - emitOperand(kResultID); - for (auto field : static_cast<IRStructType*>(inst)->getFields()) - { - emitOperand(field->getFieldType()); - // TODO: decorate offset - } - }); + List<IRType*> types; + // TODO: decorate offset + for (auto field : static_cast<IRStructType*>(inst)->getFields()) + types.add(field->getFieldType()); + auto spvStructType = emitOpTypeStruct( + inst, + types + ); emitDecorations(inst, getID(spvStructType)); return spvStructType; } @@ -1107,15 +1201,13 @@ struct SPIRVEmitContext static_cast<IRBasicType*>(matrixType->getElementType())->getBaseType(), static_cast<IRIntLit*>(matrixType->getRowCount())->getValue(), nullptr); - auto matrixSPVType = emitInst( - getSection(SpvLogicalSectionID::ConstantsAndTypes), + const auto columnCount = static_cast<IRIntLit*>(matrixType->getColumnCount())->getValue(); + auto matrixSPVType = emitOpTypeMatrix( inst, - SpvOpTypeMatrix, - kResultID, vectorSpvType, - (SpvWord)static_cast<IRIntLit*>(matrixType->getColumnCount())->getValue()); + SpvLiteralInteger::from32(int32_t(columnCount)) + ); // TODO: properly compute matrix stride. - auto columnCount = static_cast<IRIntLit*>(matrixType->getRowCount())->getValue(); uint32_t stride = 0; switch (columnCount) { @@ -1132,39 +1224,85 @@ struct SPIRVEmitContext default: break; } - emitInst( + // TODO: This decoration is not legal here. It must be placed + // on a struct member (which may entail wrapping matrices) + emitOpDecorate( + getSection(SpvLogicalSectionID::Annotations), + nullptr, + matrixSPVType, + SpvDecorationRowMajor); + emitOpDecorateMatrixStride( getSection(SpvLogicalSectionID::Annotations), nullptr, - SpvOpDecorate, matrixSPVType, - SpvDecorationRowMajor, - SpvDecorationMatrixStride, - stride); + SpvLiteralInteger::from32(stride)); return matrixSPVType; } + case kIROp_ArrayType: case kIROp_UnsizedArrayType: { - auto elementType = static_cast<IRUnsizedArrayType*>(inst)->getElementType(); - auto runtimeArrayType = emitInst( - getSection(SpvLogicalSectionID::ConstantsAndTypes), - nullptr, - SpvOpTypeRuntimeArray, - kResultID, - elementType); + const auto elementType = static_cast<IRArrayTypeBase*>(inst)->getElementType(); + const auto arrayType = inst->getOp() == kIROp_ArrayType + ? emitOpTypeArray(inst, elementType, static_cast<IRArrayTypeBase*>(inst)->getElementCount()) + : emitOpTypeRuntimeArray(inst, elementType); // TODO: properly decorate stride. + // TODO: don't do this more than once IRSizeAndAlignment sizeAndAlignment; getNaturalSizeAndAlignment(this->m_targetRequest, elementType, &sizeAndAlignment); - emitInst( + emitOpDecorateArrayStride( getSection(SpvLogicalSectionID::Annotations), nullptr, - SpvOpDecorate, - runtimeArrayType, - SpvDecorationArrayStride, - (SpvWord)sizeAndAlignment.getStride()); - return runtimeArrayType; + arrayType, + SpvLiteralInteger::from32(int32_t(sizeAndAlignment.getStride()))); + return arrayType; + } + + case kIROp_TextureType: + { + const auto texTypeInst = as<IRTextureType>(inst); + const auto sampledType = texTypeInst->getElementType(); + SpvDim dim = SpvDim1D; // Silence uninitialized warnings from msvc... + switch(texTypeInst->GetBaseShape()) + { + case TextureFlavor::Shape1D: + case TextureFlavor::Shape1DArray: + dim = SpvDim1D; + break; + case TextureFlavor::Shape2D: + case TextureFlavor::Shape2DArray: + dim = SpvDim2D; + break; + case TextureFlavor::Shape3D: + dim = SpvDim3D; + break; + case TextureFlavor::ShapeCube: + case TextureFlavor::ShapeCubeArray: + dim = SpvDimCube; + break; + case TextureFlavor::ShapeBuffer: + dim = SpvDimBuffer; + break; + } + bool arrayed = texTypeInst->isArray(); + SpvWord depth = 2; // No knowledge of if this is a depth image + bool ms = texTypeInst->isMultisample(); + // TODO: can we do better here? + SpvWord sampled = 0; // Only known at run time + // TODO: can we do better? + SpvImageFormat format = SpvImageFormatUnknown; + return emitOpTypeImage( + inst, + sampledType, + dim, + SpvLiteralInteger::from32(depth), + SpvLiteralInteger::from32(arrayed), + SpvLiteralInteger::from32(ms), + SpvLiteralInteger::from32(sampled), + format + ); } - // > OpTypeImage - // > OpTypeSampler + case kIROp_SamplerStateType: + return emitOpTypeSampler(inst); // > OpTypeArray // > OpTypeRuntimeArray // > OpTypeOpaque @@ -1177,7 +1315,11 @@ struct SPIRVEmitContext // with the result-type operand coming first, // followed by operand sfor all the parameter types. // - return emitInst(getSection(SpvLogicalSectionID::ConstantsAndTypes), inst, SpvOpTypeFunction, kResultID, OperandsOf(inst)); + return emitOpTypeFunction( + inst, + static_cast<IRFuncType*>(inst)->getResultType(), + static_cast<IRFuncType*>(inst)->getParamTypes() + ); case kIROp_RateQualifiedType: { @@ -1208,10 +1350,21 @@ struct SPIRVEmitContext return emitGlobalVar(as<IRGlobalVar>(inst)); // ... + case kIROp_Specialize: + { + const auto s = as<IRSpecialize>(inst); + const auto g = s->getBase(); + const auto e = + "Specialize instruction remains in IR for SPIR-V emit, is something undefined?\n" + + dumpIRToString(g); + SLANG_UNEXPECTED(e.getBuffer()); + } default: - String e = "Unhandled global inst in spirv-emit: " - + dumpIRToString(inst, {IRDumpOptions::Mode::Detailed, 0}); - SLANG_UNIMPLEMENTED_X(e.begin()); + { + String e = "Unhandled global inst in spirv-emit:\n" + + dumpIRToString(inst, {IRDumpOptions::Mode::Detailed, 0}); + SLANG_UNIMPLEMENTED_X(e.begin()); + } } } @@ -1228,9 +1381,11 @@ struct SPIRVEmitContext builder.getBasicType(baseType), builder.getIntValue(builder.getIntType(), elementCount)); } - auto operands = - makeArray<SpvWord>(getID(ensureInst(inst->getElementType())), (SpvWord)elementCount); - auto result = emitTypeInst(inst, SpvOpTypeVector, operands.getView()); + auto result = emitOpTypeVector( + inst, + inst->getElementType(), + SpvLiteralInteger::from32(int32_t(elementCount)) + ); return result; } @@ -1246,49 +1401,44 @@ struct SPIRVEmitContext break; case LayoutResourceKind::VaryingInput: - emitInst( + emitOpDecorateLocation( getSection(SpvLogicalSectionID::Annotations), nullptr, - SpvOpDecorate, varInst, - SpvDecorationLocation, - (SpvWord)index); - emitInst( + SpvLiteralInteger::from32(int32_t(index)) + ); + emitOpDecorateIndex( getSection(SpvLogicalSectionID::Annotations), nullptr, - SpvOpDecorate, varInst, - SpvDecorationIndex, - (SpvWord)space); + SpvLiteralInteger::from32(int32_t(space)) + ); break; case LayoutResourceKind::VaryingOutput: - emitInst( + emitOpDecorateLocation( getSection(SpvLogicalSectionID::Annotations), nullptr, - SpvOpDecorate, varInst, - SpvDecorationLocation, - (SpvWord)index); + SpvLiteralInteger::from32(int32_t(index)) + ); if (space) { - emitInst( + emitOpDecorateIndex( getSection(SpvLogicalSectionID::Annotations), nullptr, - SpvOpDecorate, varInst, - SpvDecorationIndex, - (SpvWord)space); + SpvLiteralInteger::from32(int32_t(space)) + ); } break; case LayoutResourceKind::SpecializationConstant: - emitInst( + emitOpDecorateSpecId( getSection(SpvLogicalSectionID::Annotations), nullptr, - SpvOpDecorate, varInst, - SpvDecorationSpecId, - (SpvWord)index); + SpvLiteralInteger::from32(int32_t(index)) + ); break; case LayoutResourceKind::ConstantBuffer: @@ -1296,20 +1446,18 @@ struct SPIRVEmitContext case LayoutResourceKind::UnorderedAccess: case LayoutResourceKind::SamplerState: case LayoutResourceKind::DescriptorTableSlot: - emitInst( + emitOpDecorateBinding( getSection(SpvLogicalSectionID::Annotations), nullptr, - SpvOpDecorate, varInst, - SpvDecorationBinding, - (SpvWord)index); - emitInst( + SpvLiteralInteger::from32(int32_t(index)) + ); + emitOpDecorateDescriptorSet( getSection(SpvLogicalSectionID::Annotations), nullptr, - SpvOpDecorate, varInst, - SpvDecorationDescriptorSet, - (SpvWord)space); + SpvLiteralInteger::from32(int32_t(space)) + ); break; default: break; @@ -1331,13 +1479,12 @@ struct SPIRVEmitContext registerInst(param, systemValInst); return systemValInst; } - auto varInst = emitInst( + auto varInst = emitOpVariable( getSection(SpvLogicalSectionID::GlobalVariables), param, - SpvOpVariable, param->getDataType(), - kResultID, - storageClass); + storageClass + ); emitVarLayout(varInst, layout); return varInst; } @@ -1346,21 +1493,20 @@ struct SPIRVEmitContext SpvInst* emitGlobalVar(IRGlobalVar* globalVar) { auto layout = getVarLayout(globalVar); - SLANG_ASSERT(layout); auto storageClass = SpvStorageClassUniform; if (auto ptrType = as<IRPtrTypeBase>(globalVar->getDataType())) { if (ptrType->hasAddressSpace()) storageClass = (SpvStorageClass)ptrType->getAddressSpace(); } - auto varInst = emitInst( + auto varInst = emitOpVariable( getSection(SpvLogicalSectionID::GlobalVariables), globalVar, - SpvOpVariable, globalVar->getDataType(), - kResultID, - storageClass); - emitVarLayout(varInst, layout); + storageClass + ); + if(layout) + emitVarLayout(varInst, layout); return varInst; } @@ -1438,11 +1584,13 @@ struct SPIRVEmitContext // type is given as a later operand. Slan IR instead uses // the type of a function instruction store, you know, its *type*. // - SpvInst* spvFunc = emitInst(section, irFunc, SpvOpFunction, + SpvInst* spvFunc = emitOpFunction( + section, + irFunc, irFunc->getDataType()->getResultType(), - kResultID, spvFunctionControl, - irFunc->getDataType()); + irFunc->getDataType() + ); // > OpFunctionParameter // @@ -1473,7 +1621,7 @@ struct SPIRVEmitContext // for( auto irBlock : irFunc->getBlocks() ) { - auto spvBlock = emitInst(spvFunc, irBlock, SpvOpLabel, kResultID); + auto spvBlock = emitOpLabel(spvFunc, irBlock); if (irBlock == irFunc->getFirstBlock()) { // OpVariable @@ -1494,7 +1642,7 @@ struct SPIRVEmitContext { if (irInst->getOp() == kIROp_loop) { - emitInst(spvFunc, irInst, SpvOpLabel, kResultID); + emitOpLabel(spvFunc, irInst); } } } @@ -1557,7 +1705,7 @@ struct SPIRVEmitContext // structure we will make the `OpFunctionEnd` be the last child of // the `OpFunction`. // - emitInst(spvFunc, nullptr, SpvOpFunctionEnd); + emitOpFunctionEnd(spvFunc, nullptr); // We will emit any decorations pertinent to the function to the // appropriate section of the module. @@ -1597,7 +1745,7 @@ struct SPIRVEmitContext { default: { - String e = "Unhandled local inst in spirv-emit: " + String e = "Unhandled local inst in spirv-emit:\n" + dumpIRToString(inst, {IRDumpOptions::Mode::Detailed, 0}); SLANG_UNIMPLEMENTED_X(e.getBuffer()); } @@ -1606,7 +1754,7 @@ struct SPIRVEmitContext case kIROp_Var: return emitVar(parent, inst); case kIROp_Call: - return emitCall(parent, inst); + return emitCall(parent, static_cast<IRCall*>(inst)); case kIROp_FieldAddress: return emitFieldAddress(parent, as<IRFieldAddress>(inst)); case kIROp_FieldExtract: @@ -1615,6 +1763,8 @@ struct SPIRVEmitContext return emitGetElementPtr(parent, as<IRGetElementPtr>(inst)); case kIROp_GetElement: return emitGetElement(parent, as<IRGetElement>(inst)); + case kIROp_MakeStruct: + return emitCompositeConstruct(parent, inst); case kIROp_Load: return emitLoad(parent, as<IRLoad>(inst)); case kIROp_Store: @@ -1643,8 +1793,12 @@ struct SPIRVEmitContext // TODO: break emitConstruct into separate functions for each opcode. return emitConstruct(parent, inst); case kIROp_BitCast: - return emitInst( - parent, inst, SpvOpBitcast, inst->getDataType(), kResultID, inst->getOperand(0)); + return emitOpBitcast( + parent, + inst, + inst->getDataType(), + inst->getOperand(0) + ); case kIROp_Add: case kIROp_Sub: case kIROp_Mul: @@ -1670,16 +1824,11 @@ struct SPIRVEmitContext return emitArithmetic(parent, inst); case kIROp_Return: if (as<IRReturn>(inst)->getVal()->getOp() == kIROp_VoidLit) - { - return emitInst(parent, inst, SpvOpReturn); - } + return emitOpReturn(parent, inst); else - { - return emitInst( - parent, inst, SpvOpReturnValue, as<IRReturn>(inst)->getVal()); - } + return emitOpReturnValue(parent, inst, as<IRReturn>(inst)->getVal()); case kIROp_discard: - return emitInst(parent, inst, SpvOpKill); + return emitOpKill(parent, inst); case kIROp_unconditionalBranch: { // If we are jumping to the main block of a loop, @@ -1688,15 +1837,9 @@ struct SPIRVEmitContext auto targetBlock = as<IRUnconditionalBranch>(inst)->getTargetBlock(); IRInst* loopInst = nullptr; if (isLoopTargetBlock(targetBlock, loopInst)) - { - return emitInst(parent, inst, SpvOpBranch, getIRInstSpvID(loopInst)); - } + return emitOpBranch(parent, inst, getIRInstSpvID(loopInst)); // Otherwise, emit a normal branch inst into the target block. - return emitInst( - parent, - inst, - SpvOpBranch, - getIRInstSpvID(targetBlock)); + return emitOpBranch(parent, inst, getIRInstSpvID(targetBlock)); } case kIROp_loop: { @@ -1710,7 +1853,7 @@ struct SPIRVEmitContext // Note: the body of the loop header block is emitted // after everything else to ensure Phi instructions (which come // from the actual loop target block) are emitted first. - emitInst(parent, nullptr, SpvOpBranch, blockId); + emitOpBranch(parent, nullptr, blockId); return block; } @@ -1718,26 +1861,22 @@ struct SPIRVEmitContext { auto ifelseInst = as<IRIfElse>(inst); auto afterBlockID = getIRInstSpvID(ifelseInst->getAfterBlock()); - emitInst( - parent, - nullptr, - SpvOpSelectionMerge, - afterBlockID, - 0); + emitOpSelectionMerge(parent, nullptr, afterBlockID, SpvSelectionControlMaskNone); auto falseLabel = ifelseInst->getFalseBlock(); - return emitInst( + return emitOpBranchConditional( parent, inst, - SpvOpBranchConditional, ifelseInst->getCondition(), ifelseInst->getTrueBlock(), - falseLabel ? getID(ensureInst(falseLabel)) : afterBlockID); + falseLabel ? getID(ensureInst(falseLabel)) : afterBlockID, + makeArray<SpvLiteralInteger>() + ); } case kIROp_Switch: { auto switchInst = as<IRSwitch>(inst); auto mergeBlockID = getIRInstSpvID(switchInst->getBreakLabel()); - emitInst(parent, nullptr, SpvOpSelectionMerge, mergeBlockID, 0); + emitOpSelectionMerge(parent, nullptr, mergeBlockID, SpvSelectionControlMaskNone); return emitInstCustomOperandFunc(parent, inst, SpvOpSwitch, [&]() { emitOperand(switchInst->getCondition()); auto defaultLabel = switchInst->getDefaultLabel(); @@ -1754,9 +1893,22 @@ struct SPIRVEmitContext }); } case kIROp_Unreachable: - return emitInst(parent, inst, SpvOpUnreachable); + return emitOpUnreachable(parent, inst); case kIROp_conditionalBranch: SLANG_UNEXPECTED("Unstructured branching is not supported by SPIRV."); + case kIROp_MakeVector: + return emitConstruct(parent, inst); + case kIROp_MakeVectorFromScalar: + { + const auto scalar = inst->getOperand(0); + const auto vecTy = as<IRVectorType>(inst->getDataType()); + SLANG_ASSERT(vecTy); + const auto numElems = as<IRIntLit>(vecTy->getElementCount()); + SLANG_ASSERT(numElems); + return emitSplat(parent, inst, scalar, numElems->getValue()); + } + case kIROp_MakeArray: + return emitConstruct(parent, inst); } } @@ -1767,86 +1919,29 @@ struct SPIRVEmitContext case kIROp_IntLit: { auto value = as<IRIntLit>(inst)->getValue(); - switch (as<IRBasicType>(inst->getDataType())->getBaseType()) - { - case BaseType::Int64: - case BaseType::UInt64: - case BaseType::IntPtr: - case BaseType::UIntPtr: - return emitInst( - getSection(SpvLogicalSectionID::ConstantsAndTypes), - inst, - SpvOpConstant, - inst->getDataType(), - kResultID, - (SpvWord)(value & 0xFFFFFFFF), - (SpvWord)((value >> 32) & 0xFFFFFFFF)); - default: - return emitInst( - getSection(SpvLogicalSectionID::ConstantsAndTypes), - inst, - SpvOpConstant, - inst->getDataType(), - kResultID, - (SpvWord)value); - } + return emitIntConstant(value, inst->getDataType(), inst); } case kIROp_FloatLit: { - auto value = as<IRConstant>(inst)->value.floatVal; - switch (as<IRBasicType>(inst->getDataType())->getBaseType()) - { - case BaseType::Half: - return emitInst( - getSection(SpvLogicalSectionID::ConstantsAndTypes), - inst, - SpvOpConstant, - inst->getDataType(), - kResultID, - (SpvWord)(FloatToHalf((float)value))); - case BaseType::Float: - return emitInst( - getSection(SpvLogicalSectionID::ConstantsAndTypes), - inst, - SpvOpConstant, - inst->getDataType(), - kResultID, - (SpvWord)(FloatAsInt((float)value))); - case BaseType::Double: - { - auto ival = DoubleAsInt64(value); - return emitInst( - getSection(SpvLogicalSectionID::ConstantsAndTypes), - inst, - SpvOpConstant, - inst->getDataType(), - kResultID, - (SpvWord)(ival&0xFFFFFFFF), - (SpvWord)(ival>>32)); - } - default: - return nullptr; - } + const auto value = as<IRConstant>(inst)->value.floatVal; + const auto type = inst->getDataType(); + return emitFloatConstant(value, type, inst); } case kIROp_BoolLit: { - if (as<IRBoolLit>(inst)->getValue()) + if (cast<IRBoolLit>(inst)->getValue()) { - return emitInst( - getSection(SpvLogicalSectionID::ConstantsAndTypes), + return emitOpConstantTrue( inst, - SpvOpConstantTrue, - inst->getDataType(), - kResultID); + inst->getDataType() + ); } else { - return emitInst( - getSection(SpvLogicalSectionID::ConstantsAndTypes), + return emitOpConstantFalse( inst, - SpvOpConstantFalse, - inst->getDataType(), - kResultID); + inst->getDataType() + ); } } default: @@ -1897,6 +1992,39 @@ struct SPIRVEmitContext default: break; + case kIROp_LayoutDecoration: + { + // Basic offsets for structs used in buffers + if(const auto typeLayout = as<IRTypeLayout>(as<IRLayoutDecoration>(decoration)->getLayout())) + { + if(const auto structTypeLayout = as<IRStructTypeLayout>(typeLayout)) + { + auto section = getSection(SpvLogicalSectionID::Annotations); + SpvWord i = 0; + for(const auto fieldLayoutAttr : structTypeLayout->getFieldLayoutAttrs()) + { + if(const auto structFieldLayoutAttr = as<IRStructFieldLayoutAttr>(fieldLayoutAttr)) + { + const auto varLayout = structFieldLayoutAttr->getLayout(); + if(const auto varOffsetAttr = varLayout->findOffsetAttr(LayoutResourceKind::Uniform)) + { + const auto offset = static_cast<SpvWord>(varOffsetAttr->getOffset()); + emitOpMemberDecorateOffset( + section, + fieldLayoutAttr, + dstID, + SpvLiteralInteger::from32(i), + SpvLiteralInteger::from32(offset) + ); + } + } + ++i; + } + } + } + } + break; + // [3.32.2. Debug Instructions] // // > OpName @@ -1905,7 +2033,11 @@ struct SPIRVEmitContext { auto section = getSection(SpvLogicalSectionID::DebugNames); auto nameHint = cast<IRNameHintDecoration>(decoration); - emitInst(section, decoration, SpvOpName, dstID, nameHint->getName()); + // We can't associate this spirv instruction with our + // irInstruction, our instruction may be a hint on several + // values, however this decoration is specific to a single + // dstID. + emitOpName(section, nullptr, dstID, nameHint->getName()); } break; @@ -1932,23 +2064,27 @@ struct SPIRVEmitContext auto entryPointDecor = cast<IREntryPointDecoration>(decoration); auto spvStage = mapStageToExecutionModel(entryPointDecor->getProfile().getStage()); auto name = entryPointDecor->getName()->getStringSlice(); - emitInstCustomOperandFunc(section, decoration, SpvOpEntryPoint, [&]() { - emitOperand(spvStage); - emitOperand(dstID); - emitOperand(name); - // `interface` part: reference all global variables that are used by this entrypoint. - // TODO: we may want to perform more accurate tracking. - for (auto globalInst : m_irModule->getModuleInst()->getChildren()) + List<IRInst*> params; + // `interface` part: reference all global variables that are used by this entrypoint. + // TODO: we may want to perform more accurate tracking. + for (auto globalInst : m_irModule->getModuleInst()->getChildren()) + { + switch (globalInst->getOp()) { - switch (globalInst->getOp()) - { - case kIROp_GlobalVar: - case kIROp_GlobalParam: - emitOperand(getIRInstSpvID(globalInst)); - break; - } + case kIROp_GlobalVar: + case kIROp_GlobalParam: + params.add(globalInst); + break; } - }); + } + emitOpEntryPoint( + section, + decoration, + spvStage, + dstID, + name, + params + ); } break; @@ -1969,29 +2105,25 @@ struct SPIRVEmitContext // in those positions in the Slang IR). // auto numThreads = cast<IRNumThreadsDecoration>(decoration); - emitInst(section, decoration, SpvOpExecutionMode, dstID, SpvExecutionModeLocalSize, - SpvWord(numThreads->getX()->getValue()), - SpvWord(numThreads->getY()->getValue()), - SpvWord(numThreads->getZ()->getValue())); + emitOpExecutionModeLocalSize( + section, + decoration, + dstID, + SpvLiteralInteger::from32(int32_t(numThreads->getX()->getValue())), + SpvLiteralInteger::from32(int32_t(numThreads->getY()->getValue())), + SpvLiteralInteger::from32(int32_t(numThreads->getZ()->getValue())) + ); } break; case kIROp_SPIRVBufferBlockDecoration: { - emitInst( + emitOpDecorate( getSection(SpvLogicalSectionID::Annotations), decoration, - SpvOpDecorate, - dstID, - SpvDecorationBlock); - emitInst( - getSection(SpvLogicalSectionID::Annotations), - nullptr, - SpvOpMemberDecorate, dstID, - 0, - SpvDecorationOffset, - 0); + SpvDecorationBlock + ); } break; // ... @@ -2035,20 +2167,18 @@ struct SPIRVEmitContext builder.setInsertBefore(type); auto ptrType = as<IRPtrTypeBase>(type); SLANG_ASSERT(ptrType && "`getBuiltinGlobalVar`: `type` must be ptr type."); - auto varInst = emitInst( + auto varInst = emitOpVariable( getSection(SpvLogicalSectionID::GlobalVariables), nullptr, - SpvOpVariable, type, - kResultID, - (SpvStorageClass)ptrType->getAddressSpace()); - emitInst( + static_cast<SpvStorageClass>(ptrType->getAddressSpace()) + ); + emitOpDecorateBuiltIn( getSection(SpvLogicalSectionID::Annotations), nullptr, - SpvOpDecorate, varInst, - SpvDecorationBuiltIn, - builtinVal); + builtinVal + ); m_builtinGlobalVars[builtinVal] = varInst; return varInst; } @@ -2078,7 +2208,7 @@ struct SPIRVEmitContext SpvInst* emitParam(SpvInstParent* parent, IRInst* inst) { - return emitInst(parent, inst, SpvOpFunctionParameter, inst->getFullType(), kResultID); + return emitOpFunctionParameter(parent, inst, inst->getFullType()); } SpvInst* emitVar(SpvInstParent* parent, IRInst* inst) @@ -2090,7 +2220,7 @@ struct SPIRVEmitContext { storageClass = (SpvStorageClass)ptrType->getAddressSpace(); } - return emitInst(parent, inst, SpvOpVariable, inst->getFullType(), kResultID, storageClass); + return emitOpVariable(parent, inst, inst->getFullType(), storageClass); } /// Cached `IRParam` indices in an `IRBlock`. For use in `getParamIndexInBlock`. @@ -2143,29 +2273,29 @@ struct SPIRVEmitContext void emitLoopHeaderBlock(IRLoop* loopInst, SpvInst* loopHeaderBlock) { - SpvWord loopControl = 0; + SpvLoopControlMask loopControl = SpvLoopControlMaskNone; if (auto loopControlDecoration = loopInst->findDecoration<IRLoopControlDecoration>()) { switch (loopControlDecoration->getMode()) { case IRLoopControl::kIRLoopControl_Unroll: - loopControl = 0x1; + loopControl = SpvLoopControlUnrollMask; break; case IRLoopControl::kIRLoopControl_Loop: - loopControl = 0x2; + loopControl = SpvLoopControlDontUnrollMask; break; default: break; } } - emitInst( + emitOpLoopMerge( loopHeaderBlock, nullptr, - SpvOpLoopMerge, getIRInstSpvID(loopInst->getBreakBlock()), getIRInstSpvID(loopInst->getContinueBlock()), - loopControl); - emitInst(loopHeaderBlock, nullptr, SpvOpBranch, loopInst->getTargetBlock()); + loopControl + ); + emitOpBranch(loopHeaderBlock, nullptr, loopInst->getTargetBlock()); } SpvInst* emitPhi(SpvInstParent* parent, IRParam* inst) @@ -2225,9 +2355,9 @@ struct SPIRVEmitContext }); } - SpvInst* emitCall(SpvInstParent* parent, IRInst* inst) + SpvInst* emitCall(SpvInstParent* parent, IRCall* inst) { - auto funcValue = inst->getOperand(0); + auto funcValue = inst->getCallee(); // Does this function declare any requirements. handleRequiredCapabilities(funcValue); @@ -2241,8 +2371,13 @@ struct SPIRVEmitContext } else { - return emitInst( - parent, inst, SpvOpFunctionCall, inst->getFullType(), kResultID, OperandsOf(inst)); + return emitOpFunctionCall( + parent, + inst, + inst->getFullType(), + funcValue, + inst->getArgsList() + ); } } @@ -2274,14 +2409,16 @@ struct SPIRVEmitContext // different storage-class qualifier. We need to pre-create these // storage-class-qualified result pointer types so they can be used // during inlining of the snippet. - if (auto oldPtrType = as<IRPtrTypeBase>(inst->getDataType())) { - for (auto storageClass : snippet->usedResultTypeStorageClasses) + IRBuilder builder(m_irModule); + builder.setInsertBefore(inst); + for (auto storageClass : snippet->usedPtrResultTypeStorageClasses) { - IRBuilder builder(m_irModule); - builder.setInsertBefore(inst); auto newPtrType = builder.getPtrType( - oldPtrType->getOp(), oldPtrType->getValueType(), storageClass); + kIROp_PtrType, + inst->getDataType(), + storageClass + ); context.qualifiedResultTypes[storageClass] = newPtrType; } } @@ -2309,14 +2446,11 @@ struct SPIRVEmitContext auto floatType = builder.getType(kIROp_FloatType); auto element1 = emitFloatConstant(constant.floatValues[0], floatType); auto element2 = emitFloatConstant(constant.floatValues[1], floatType); - result = emitInst( - getSection(SpvLogicalSectionID::ConstantsAndTypes), + result = emitOpConstantComposite( nullptr, - SpvOpConstantComposite, builder.getVectorType(floatType, builder.getIntValue(builder.getIntType(), 2)), - kResultID, - element1, - element2); + makeArray(element1, element2) + ); } break; case SpvSnippet::ASMType::Int: @@ -2327,14 +2461,11 @@ struct SPIRVEmitContext auto uintType = builder.getType(kIROp_UIntType); auto element1 = emitIntConstant((IRIntegerValue)constant.intValues[0], uintType); auto element2 = emitIntConstant((IRIntegerValue)constant.intValues[1], uintType); - result = emitInst( - getSection(SpvLogicalSectionID::ConstantsAndTypes), + result = emitOpConstantComposite( nullptr, - SpvOpConstantComposite, builder.getVectorType(uintType, builder.getIntValue(builder.getIntType(), 2)), - kResultID, - element1, - element2); + makeArray(element1, element2) + ); } break; } @@ -2356,6 +2487,9 @@ struct SPIRVEmitContext case SpvSnippet::ASMType::Int: irType = builder.getIntType(); break; + case SpvSnippet::ASMType::UInt: + irType = builder.getUIntType(); + break; case SpvSnippet::ASMType::Float2: irType = builder.getVectorType( builder.getType(kIROp_FloatType), builder.getIntValue(builder.getIntType(), 2)); @@ -2533,23 +2667,27 @@ struct SPIRVEmitContext baseStructType = as<IRStructType>(base->getDataType()); auto structPtrType = builder.getPtrType(baseStructType); - auto varInst = emitInst( - parent, nullptr, SpvOpVariable, structPtrType, kResultID, SpvStorageClassFunction); - emitInst(parent, nullptr, SpvOpStore, varInst, base); + auto varInst = emitOpVariable( + parent, + nullptr, + structPtrType, + SpvStorageClassFunction + ); + emitOpStore(parent, nullptr, varInst, base); baseId = getID(varInst); } - SLANG_ASSERT(baseStructType && "field_address require base to be a struct."); + SLANG_ASSERT(baseStructType && "field_address requires base to be a struct."); auto fieldId = emitIntConstant( getStructFieldId(baseStructType, as<IRStructKey>(fieldAddress->getField())), builder.getIntType()); - return emitInst( + SLANG_ASSERT(as<IRPtrTypeBase>(fieldAddress->getFullType())); + return emitOpAccessChain( parent, fieldAddress, - SpvOpAccessChain, fieldAddress->getFullType(), - kResultID, baseId, - fieldId); + makeArray(fieldId) + ); } SpvInst* emitFieldExtract(SpvInstParent* parent, IRFieldExtract* inst) @@ -2558,100 +2696,84 @@ struct SPIRVEmitContext builder.setInsertBefore(inst); IRStructType* baseStructType = as<IRStructType>(inst->getBase()->getDataType()); - SLANG_ASSERT(baseStructType && "field_extract require base to be a struct."); - auto fieldId = emitIntConstant( - getStructFieldId(baseStructType, as<IRStructKey>(inst->getField())), - builder.getIntType()); + SLANG_ASSERT(baseStructType && "field_extract requires base to be a struct."); + auto fieldId = static_cast<SpvWord>(getStructFieldId( + baseStructType, + as<IRStructKey>(inst->getField()))); - return emitInst( + return emitOpCompositeExtract( parent, inst, - SpvOpCompositeExtract, inst->getDataType(), - kResultID, inst->getBase(), - fieldId); + makeArray(SpvLiteralInteger::from32(fieldId)) + ); } SpvInst* emitGetElementPtr(SpvInstParent* parent, IRGetElementPtr* inst) { auto base = inst->getBase(); SpvWord baseId = 0; - IRArrayType* baseArrayType = nullptr; // Only used in debug build, but we don't want a warning/error for an unused initialized variable - SLANG_UNUSED(baseArrayType); if (auto ptrLikeType = as<IRPointerLikeType>(base->getDataType())) { - baseArrayType = as<IRArrayType>(ptrLikeType->getElementType()); baseId = getID(ensureInst(base)); } else if (auto ptrType = as<IRPtrTypeBase>(base->getDataType())) { - baseArrayType = as<IRArrayType>(ptrType->getValueType()); baseId = getID(ensureInst(base)); } else { SLANG_ASSERT(!"invalid IR: base of getElementPtr must be a pointer."); } - SLANG_ASSERT(baseArrayType && "getElementPtr require base to be an array."); - return emitInst( + SLANG_ASSERT(as<IRPtrTypeBase>(inst->getFullType())); + return emitOpAccessChain( parent, inst, - SpvOpAccessChain, inst->getFullType(), - kResultID, baseId, - inst->getIndex()); + makeArray(inst->getIndex()) + ); } SpvInst* emitGetElement(SpvInstParent* parent, IRGetElement* inst) { auto base = inst->getBase(); - SpvWord baseId = 0; - IRArrayType* baseArrayType = nullptr; - // Only used in debug build, but we don't want a warning/error for an unused initialized variable - SLANG_UNUSED(baseArrayType); - - if (auto ptrLikeType = as<IRPointerLikeType>(base->getDataType())) - { - baseArrayType = as<IRArrayType>(ptrLikeType->getElementType()); - baseId = getID(ensureInst(base)); - } - else if (auto ptrType = as<IRPtrTypeBase>(base->getDataType())) - { - baseArrayType = as<IRArrayType>(ptrType->getValueType()); - baseId = getID(ensureInst(base)); - } - else - { - SLANG_ASSERT(!"invalid IR: base of getElement must be a pointer."); - } - SLANG_ASSERT(baseArrayType && "getElement require base to be an array."); + const auto baseTy = base->getDataType(); + SLANG_ASSERT( + as<IRPointerLikeType>(baseTy) || + as<IRArrayType>(baseTy) || + as<IRVectorType>(baseTy) || + as<IRMatrixType>(baseTy)); IRBuilder builder(m_irModule); builder.setInsertBefore(inst); - auto ptr = emitInst( + auto ptr = emitOpAccessChain( parent, nullptr, - SpvOpAccessChain, builder.getPtrType(inst->getFullType()), - kResultID, - baseId, - inst->getIndex()); - return emitInst(parent, inst, SpvOpLoad, inst->getFullType(), kResultID, ptr); + inst->getBase(), + makeArray(inst->getIndex()) + ); + return emitOpLoad( + parent, + inst, + inst->getFullType(), + ptr + ); } SpvInst* emitLoad(SpvInstParent* parent, IRLoad* inst) { - return emitInst(parent, inst, SpvOpLoad, inst->getDataType(), kResultID, inst->getPtr()); + return emitOpLoad(parent, inst, inst->getDataType(), inst->getPtr()); } SpvInst* emitStore(SpvInstParent* parent, IRStore* inst) { - return emitInst(parent, inst, SpvOpStore, inst->getPtr(), inst->getVal()); + return emitOpStore(parent, inst, inst->getPtr(), inst->getVal()); } SpvInst* emitStructuredBufferLoad(SpvInstParent* parent, IRInst* inst) @@ -2682,14 +2804,14 @@ struct SPIRVEmitContext { if (inst->getElementCount() == 1) { - return emitInst( + const auto index = as<IRIntLit>(inst->getElementIndex(0))->getValue(); + return emitOpCompositeExtract( parent, inst, - SpvOpCompositeExtract, inst->getDataType(), - kResultID, inst->getBase(), - (SpvWord)as<IRIntLit>(inst->getElementIndex(0))->getValue()); + makeArray(SpvLiteralInteger::from32(int32_t(index))) + ); } else { @@ -2727,25 +2849,22 @@ struct SPIRVEmitContext const auto fromInfo = getIntTypeInfo(fromType); const auto toInfo = getIntTypeInfo(toType); - const auto convertWith = [&](auto op){ - return emitInst(parent, inst, op, toTypeV, kResultID, inst->getOperand(0)); - }; if(fromInfo == toInfo) - return convertWith(SpvOpCopyObject); + return emitOpCopyObject(parent, inst, toTypeV, inst->getOperand(0)); else if(fromInfo.width == toInfo.width) - return convertWith(SpvOpBitcast); + return emitOpBitcast(parent, inst, toTypeV, inst->getOperand(0)); else if(!fromInfo.isSigned && !toInfo.isSigned) // unsigned to unsigned, don't sign extend - return convertWith(SpvOpUConvert); + return emitOpUConvert(parent, inst, toTypeV, inst->getOperand(0)); else if(toInfo.isSigned) // unsigned to signed, sign extend - return convertWith(SpvOpSConvert); + return emitOpSConvert(parent, inst, toTypeV, inst->getOperand(0)); else if(fromInfo.isSigned) // signed to unsigned, sign extend - return convertWith(SpvOpSConvert); + return emitOpSConvert(parent, inst, toTypeV, inst->getOperand(0)); else if(fromInfo.isSigned && toInfo.isSigned) // signed to signed, sign extend - return convertWith(SpvOpSConvert); + return emitOpSConvert(parent, inst, toTypeV, inst->getOperand(0)); SLANG_UNREACHABLE(__func__); } @@ -2761,7 +2880,7 @@ struct SPIRVEmitContext SLANG_ASSERT(isFloatingType(toType)); SLANG_ASSERT(!isTypeEqual(fromType, toType)); - return emitInst(parent, inst, SpvOpFConvert, toTypeV, kResultID, inst->getOperand(0)); + return emitOpFConvert(parent, inst, toTypeV, inst->getOperand(0)); } SpvInst* emitIntToFloatCast(SpvInstParent* parent, IRCastIntToFloat* inst) @@ -2776,11 +2895,9 @@ struct SPIRVEmitContext const auto fromInfo = getIntTypeInfo(fromType); - const auto convertWith = [&](auto op){ - return emitInst(parent, inst, op, toTypeV, kResultID, inst->getOperand(0)); - }; - - return convertWith(fromInfo.isSigned ? SpvOpConvertSToF : SpvOpConvertUToF); + return fromInfo.isSigned + ? emitOpConvertSToF(parent, inst, toTypeV, inst->getOperand(0)) + : emitOpConvertUToF(parent, inst, toTypeV, inst->getOperand(0)); } SpvInst* emitFloatToIntCast(SpvInstParent* parent, IRCastFloatToInt* inst) @@ -2795,11 +2912,14 @@ struct SPIRVEmitContext const auto toInfo = getIntTypeInfo(toType); - const auto convertWith = [&](auto op){ - return emitInst(parent, inst, op, toTypeV, kResultID, inst->getOperand(0)); - }; + return toInfo.isSigned + ? emitOpConvertFToS(parent, inst, toTypeV, inst->getOperand(0)) + : emitOpConvertFToU(parent, inst, toTypeV, inst->getOperand(0)); + } - return convertWith(toInfo.isSigned ? SpvOpConvertFToS : SpvOpConvertFToU); + SpvInst* emitCompositeConstruct(SpvInstParent* parent, IRInst* inst) + { + return emitOpCompositeConstruct(parent, inst, inst->getDataType(), OperandsOf(inst)); } SpvInst* emitConstruct(SpvInstParent* parent, IRInst* inst) @@ -2809,20 +2929,16 @@ struct SPIRVEmitContext if (inst->getOperandCount() == 1) { if (inst->getDataType() == inst->getOperand(0)->getDataType()) - return emitInst( + return emitOpCopyObject( parent, inst, - SpvOpCopyObject, inst->getFullType(), - kResultID, inst->getOperand(0)); else - return emitInst( + return emitOpBitcast( parent, inst, - SpvOpBitcast, inst->getFullType(), - kResultID, inst->getOperand(0)); } else @@ -2833,33 +2949,24 @@ struct SPIRVEmitContext } else { - return emitInst( - parent, - inst, - SpvOpCompositeConstruct, - inst->getDataType(), - kResultID, - OperandsOf(inst)); + return emitCompositeConstruct(parent, inst); } } - SpvInst* emitSplat(SpvInstParent* parent, IRInst* scalar, IRIntegerValue numElems) + SpvInst* emitSplat(SpvInstParent* parent, IRInst* inst, IRInst* scalar, IRIntegerValue numElems) { const auto scalarTy = as<IRBasicType>(scalar->getDataType()); + SLANG_ASSERT(scalarTy); const auto spvVecTy = ensureVectorType( scalarTy->getBaseType(), numElems, nullptr); - return emitInstCustomOperandFunc( + return emitOpCompositeConstruct( parent, - nullptr, - SpvOpCompositeConstruct, - [&](){ - emitOperand(spvVecTy); - emitOperand(kResultID); - for(Int i = 0; i < numElems; ++i) - emitOperand(scalar); - }); + inst, + spvVecTy, + List<IRInst*>::makeRepeated(scalar, Index(numElems)) + ); } bool isSignedType(IRType* type) @@ -3019,13 +3126,13 @@ struct SPIRVEmitContext { const auto len = as<IRIntLit>(lVec->getElementCount()); SLANG_ASSERT(len); - return go(l, emitSplat(parent, r, len->getValue())); + return go(l, emitSplat(parent, nullptr, r, len->getValue())); } else if (!lVec && rVec) { const auto len = as<IRIntLit>(rVec->getElementCount()); SLANG_ASSERT(len); - return go(emitSplat(parent, l, len->getValue()), r); + return go(emitSplat(parent, nullptr, l, len->getValue()), r); } return go(l, r); } @@ -3038,11 +3145,11 @@ struct SPIRVEmitContext { if (m_capabilities.add(capability)) { - emitInst( + emitOpCapability( getSection(SpvLogicalSectionID::Capabilities), nullptr, - SpvOpCapability, - capability); + capability + ); } } @@ -3073,12 +3180,6 @@ struct SPIRVEmitContext } } - void diagnoseUnhandledInst(IRInst* inst) - { - m_sink->diagnose( - inst, Diagnostics::unimplemented, "unexpected IR opcode during code emit"); - } - SPIRVEmitContext(IRModule* module, TargetRequest* target, DiagnosticSink* sink) : SPIRVEmitSharedContext(module, target, sink) , m_irModule(module) @@ -3101,6 +3202,16 @@ SlangResult emitSPIRVFromIR( SPIRVEmitContext context(irModule, targetRequest, sink); legalizeIRForSPIRV(&context, irModule, irEntryPoints, codeGenContext); +#if 0 + DiagnosticSinkWriter writer(codeGenContext->getSink()); + dumpIR( + irModule, + {IRDumpOptions::Mode::Simplified, 0}, + "BEFORE SPIR-V EMIT", + codeGenContext->getSourceManager(), + &writer); +#endif + context.emitFrontMatter(); for (auto irEntryPoint : irEntryPoints) { @@ -3110,7 +3221,7 @@ SlangResult emitSPIRVFromIR( spirvOut.addRange( (uint8_t const*) context.m_words.getBuffer(), - context.m_words.getCount() * sizeof(context.m_words[0])); + context.m_words.getCount() * Index(sizeof(context.m_words[0]))); const auto validationResult = debugValidateSPIRV(spirvOut); // If validation isn't available, don't say it failed, it's just a debug diff --git a/source/slang/slang-ir-inst-defs.h b/source/slang/slang-ir-inst-defs.h index c3c92f9ba..623fca32f 100644 --- a/source/slang/slang-ir-inst-defs.h +++ b/source/slang/slang-ir-inst-defs.h @@ -997,6 +997,8 @@ INST(GetEquivalentStructuredBuffer, getEquivalentStructuredBuffer, 1, 0) INST(MatrixTypeLayout, matrixTypeLayout, 1, HOISTABLE) INST(ExistentialTypeLayout, existentialTypeLayout, 0, HOISTABLE) INST(StructTypeLayout, structTypeLayout, 0, HOISTABLE) + INST(TupleTypeLayout, tupleTypeLayout, 0, HOISTABLE) + INST(StructuredBufferTypeLayout, structuredBufferTypeLayout, 1, HOISTABLE) // TODO(JS): Ideally we'd have the layout to the pointed to value type (ie 1 instead of 0 here). But to avoid infinite recursion we don't. INST(PointerTypeLayout, ptrTypeLayout, 0, HOISTABLE) INST_RANGE(TypeLayout, TypeLayoutBase, PointerTypeLayout) @@ -1008,6 +1010,7 @@ INST_RANGE(Layout, VarLayout, EntryPointLayout) INST(PendingLayoutAttr, pendingLayout, 1, HOISTABLE) INST(StageAttr, stage, 1, HOISTABLE) INST(StructFieldLayoutAttr, fieldLayout, 2, HOISTABLE) + INST(TupleFieldLayoutAttr, fieldLayout, 1, HOISTABLE) INST(CaseTypeLayoutAttr, caseLayout, 1, HOISTABLE) INST(UNormAttr, unorm, 0, HOISTABLE) INST(SNormAttr, snorm, 0, HOISTABLE) diff --git a/source/slang/slang-ir-insts.h b/source/slang/slang-ir-insts.h index dfd662be3..4af21c50a 100644 --- a/source/slang/slang-ir-insts.h +++ b/source/slang/slang-ir-insts.h @@ -82,6 +82,27 @@ struct IRTargetSpecificDecoration : IRDecoration IRCapabilitySet* getTargetCapsOperand() { return cast<IRCapabilitySet>(getOperand(0)); } CapabilitySet getTargetCaps() { return getTargetCapsOperand()->getCaps(); } + + bool hasPredicate() + { + return getOperandCount() >= 4; + } + + UnownedStringSlice getTypePredicate() + { + SLANG_ASSERT(getOperandCount() == 4); + const auto lit = as<IRStringLit>(getOperand(2)); + SLANG_ASSERT(lit); + return lit->getStringSlice(); + } + + IRType* getTypeScrutinee() + { + SLANG_ASSERT(getOperandCount() == 4); + const auto t = as<IRType>(getOperand(3)); + SLANG_ASSERT(t); + return t; + } }; struct IRTargetDecoration : IRTargetSpecificDecoration @@ -1546,6 +1567,38 @@ struct IRArrayTypeLayout : IRTypeLayout }; }; + /// Specialized layout information for structured buffer types +struct IRStructuredBufferTypeLayout : IRTypeLayout +{ + typedef IRTypeLayout Super; + + IR_LEAF_ISA(StructuredBufferTypeLayout) + + IRTypeLayout* getElementTypeLayout() + { + return cast<IRTypeLayout>(getOperand(0)); + } + + struct Builder : Super::Builder + { + Builder(IRBuilder* irBuilder, IRTypeLayout* elementTypeLayout) + : Super::Builder(irBuilder) + , m_elementTypeLayout(elementTypeLayout) + {} + + IRStructuredBufferTypeLayout* build() + { + return cast<IRStructuredBufferTypeLayout>(Super::Builder::build()); + } + + protected: + IROp getOp() SLANG_OVERRIDE { return kIROp_StructuredBufferTypeLayout; } + void addOperandsImpl(List<IRInst*>& ioOperands) override; + + IRTypeLayout* m_elementTypeLayout; + }; +}; + /* TODO(JS): It would arguably be "more correct" if the IRPointerTypeLayout, contained a refence to the value/target @@ -1728,6 +1781,74 @@ struct IRStructTypeLayout : IRTypeLayout }; }; + /// Attribute that specifies the layout for one field of a structure type. +struct IRTupleFieldLayoutAttr : IRAttr +{ + IR_LEAF_ISA(TupleFieldLayoutAttr) + + IRTypeLayout* getLayout() + { + return cast<IRTypeLayout>(getOperand(1)); + } +}; + + /// Specialized layout information for tuple types. +struct IRTupleTypeLayout : IRTypeLayout +{ + IR_LEAF_ISA(TupleTypeLayout) + + typedef IRTypeLayout Super; + + /// Get all of the attributes that represent field layouts. + IROperandList<IRTupleFieldLayoutAttr> getFieldLayoutAttrs() + { + return findAttrs<IRTupleFieldLayoutAttr>(); + } + + /// Get the number of fields for which layout information is stored. + UInt getFieldCount() + { + return getFieldLayoutAttrs().getCount(); + } + + /// Get the layout information for a field by `index` + IRTypeLayout* getFieldLayout(UInt index) + { + return getFieldLayoutAttrs()[index]->getLayout(); + } + + /// Specialized builder for tuple type layouts. + struct Builder : Super::Builder + { + Builder(IRBuilder* irBuilder) + : Super::Builder(irBuilder) + {} + + void addField(IRTypeLayout* layout) + { + FieldInfo info; + info.layout = layout; + m_fields.add(info); + } + + IRTupleTypeLayout* build() + { + return cast<IRTupleTypeLayout>(Super::Builder::build()); + } + + protected: + IROp getOp() SLANG_OVERRIDE { return kIROp_TupleTypeLayout; } + void addAttrsImpl(List<IRInst*>& ioOperands) override; + + struct FieldInfo + { + IRTypeLayout* layout; + }; + + List<FieldInfo> m_fields; + }; +}; + /// Attribute that represents the layout for one case of a union type struct IRCaseTypeLayoutAttr : IRAttr { @@ -1958,6 +2079,10 @@ struct IRCall : IRInst UInt getArgCount() { return getOperandCount() - 1; } IRUse* getArgs() { return getOperands() + 1; } + IROperandList<IRInst> getArgsList() + { + return IROperandList<IRInst>(getOperands() + 1, getOperands() + getOperandCount()); + } IRInst* getArg(UInt index) { return getOperand(index + 1); } }; @@ -3435,6 +3560,9 @@ public: IRVar* emitVar( IRType* type); + IRVar* emitVar( + IRType* type, + IRIntegerValue addressSpace); IRInst* emitLoad( IRType* type, @@ -3722,6 +3850,18 @@ public: return addDecoration(value, op, operands, SLANG_COUNT_OF(operands)); } + IRDecoration* addDecoration(IRInst* value, IROp op, IRInst* operand0, IRInst* operand1, IRInst* operand2) + { + IRInst* operands[] = { operand0, operand1, operand2 }; + return addDecoration(value, op, operands, SLANG_COUNT_OF(operands)); + } + + IRDecoration* addDecoration(IRInst* value, IROp op, IRInst* operand0, IRInst* operand1, IRInst* operand2, IRInst* operand3) + { + IRInst* operands[] = { operand0, operand1, operand2, operand3 }; + return addDecoration(value, op, operands, SLANG_COUNT_OF(operands)); + } + template<typename T> void addSimpleDecoration(IRInst* value) { @@ -3747,6 +3887,8 @@ public: IRStructFieldLayoutAttr* getFieldLayoutAttr( IRInst* key, IRVarLayout* layout); + IRTupleFieldLayoutAttr* getTupleFieldLayoutAttr( + IRTypeLayout* layout); IRCaseTypeLayoutAttr* getCaseTypeLayoutAttr( IRTypeLayout* layout); @@ -3844,14 +3986,26 @@ public: addDecoration(value, kIROp_SemanticDecoration, getStringValue(text), getIntValue(getIntType(), index)); } - void addTargetIntrinsicDecoration(IRInst* value, IRInst* caps, UnownedStringSlice const& definition) + void addTargetIntrinsicDecoration(IRInst* value, IRInst* caps, UnownedStringSlice const& definition, UnownedStringSlice const& predicate, IRInst* typeScrutinee) { - addDecoration(value, kIROp_TargetIntrinsicDecoration, caps, getStringValue(definition)); + typeScrutinee + ? addDecoration( + value, + kIROp_TargetIntrinsicDecoration, + caps, + getStringValue(definition), + getStringValue(predicate), + typeScrutinee) + : addDecoration( + value, + kIROp_TargetIntrinsicDecoration, + caps, + getStringValue(definition)); } - void addTargetIntrinsicDecoration(IRInst* value, CapabilitySet const& caps, UnownedStringSlice const& definition) + void addTargetIntrinsicDecoration(IRInst* value, CapabilitySet const& caps, UnownedStringSlice const& definition, UnownedStringSlice const& predicate = UnownedStringSlice{}, IRInst* typeScrutinee = nullptr) { - addTargetIntrinsicDecoration(value, getCapabilityValue(caps), definition); + addTargetIntrinsicDecoration(value, getCapabilityValue(caps), definition, predicate, typeScrutinee); } void addTargetDecoration(IRInst* value, IRInst* caps) diff --git a/source/slang/slang-ir-layout-on-types.cpp b/source/slang/slang-ir-layout-on-types.cpp new file mode 100644 index 000000000..a1e480067 --- /dev/null +++ b/source/slang/slang-ir-layout-on-types.cpp @@ -0,0 +1,153 @@ +#include "slang-ir-layout-on-types.h" + +#include "slang-ir-insts.h" +#include "slang-ir.h" + +namespace Slang +{ + +struct TypeTypeLayout +{ + IRType* type; + IRTypeLayout* layout; +}; + +template<typename Struct, typename StructuredBuffer, typename Array, typename Base> +static void zipPreorderTypeAndTypeLayout( + Struct struct_, + StructuredBuffer structuredBuffer, + Array array, + Base base, + IRType* type, + IRTypeLayout* layout) +{ + auto go = [&](auto& go, IRType* type, IRLayout* layout) -> void{ + if(const auto structTypeLayout = as<IRStructTypeLayout>(layout)) + { + const auto structType = as<IRStructType>(type); + SLANG_ASSERT(structType); + struct_(structType, structTypeLayout); + + Index i = 0; + for(const auto field : structType->getFields()) + { + const auto fieldVarLayout = structTypeLayout->getFieldLayout(i); + const auto fieldTypeLayout = fieldVarLayout->getTypeLayout(); + const auto fieldType = field->getFieldType(); + go(go, fieldType, fieldTypeLayout); + ++i; + } + } + else if(const auto structuredBufferTypeLayout = as<IRStructuredBufferTypeLayout>(layout)) + { + const auto structuredBufferType = as<IRHLSLStructuredBufferTypeBase>(type); + SLANG_ASSERT(structuredBufferType); + structuredBuffer(structuredBufferType, structuredBufferTypeLayout); + + go(go, structuredBufferType->getElementType(), structuredBufferTypeLayout->getElementTypeLayout()); + } + else if(const auto arrayTypeLayout = as<IRArrayTypeLayout>(layout)) + { + const auto arrayType = as<IRArrayTypeBase>(type); + SLANG_ASSERT(arrayType); + array(arrayType , arrayTypeLayout); + + go(go, arrayType->getElementType(), arrayTypeLayout->getElementTypeLayout()); + } + else if(const auto existentialTypeLayout = as<IRExistentialTypeLayout>(layout)) + { + // TODO: To resolve this we should + // - Record the interface width in each IRExistentialTypeLayout + // - To allow us to replace them with IRTupleTypeLayout when we lower existentials + // - To allow us to replace them with IRStructTypeLayout when we lower tuples + SLANG_UNEXPECTED("Existentials (or their layouts) have not been erased before SPIR-V emit"); + } + else + { + base((IRType*)type, (IRTypeLayout*)layout); + } + }; + go(go, type, layout); +} + +static Dictionary<IRType*, IRTypeLayout*> associateTypesWithLayouts(IRModule* module) +{ + List<TypeTypeLayout> worklist; + for(auto globalInst : module->getGlobalInsts()) + { + if(const auto varLayout = as<IRVarLayout>(globalInst)) + { + auto typeLayout = varLayout->getTypeLayout(); + List<IRInst*> vars; + traverseUsers(varLayout, [&](IRInst* varLayoutUser){ + if(const auto dec = as<IRLayoutDecoration>(varLayoutUser)) + { + if(const auto globalParam = as<IRGlobalParam>(dec->getParent())) + vars.add(globalParam); + else + { + // todo + } + } + else if(as<IREntryPointLayout>(varLayoutUser)) + { + // todo + } + else if(as<IRStructFieldLayoutAttr>(varLayoutUser)) + { + // todo + } + else if(as<IRParameterGroupTypeLayout>(varLayoutUser)) + { + // todo + } + else + SLANG_UNEXPECTED("Var layout was used somewhere unexpected"); + }); + if(vars.getCount() == 1) + { + auto type = vars[0]->getDataType(); + worklist.add({type, typeLayout}); + } + else if(vars.getCount() > 2) + { + SLANG_UNIMPLEMENTED_X("vars with different layouts"); + } + } + } + + Dictionary<IRType*, IRTypeLayout*> ret; + while(!worklist.getCount() == 0) + { + const auto ttl = worklist.getLast(); + worklist.removeLast(); + const auto add = [&](IRType* type, IRTypeLayout* typeLayout){ + const auto* otherTypeLayout = ret.tryGetValueOrAdd(type, typeLayout); + if(otherTypeLayout) + SLANG_ASSERT(*otherTypeLayout == typeLayout); + }; + zipPreorderTypeAndTypeLayout(add, add, add, add, ttl.type, ttl.layout); + } + + return ret; +} + +static void decorateTypesWithLayouts(IRModule* module, const Dictionary<IRType*, IRTypeLayout*>& assocs) +{ + IRBuilder builder(module); + for(const auto& [type, layout] : assocs) + { + builder.setInsertBefore(type); + // TODO: types with more than one decoration (needs deduplicating) + builder.addLayoutDecoration(type, layout); + } +} + +void placeTypeLayoutsOnTypes(IRModule* module, CodeGenContext* codeGenContext) +{ + SLANG_ASSERT(module); + SLANG_ASSERT(codeGenContext); + const auto assocs = associateTypesWithLayouts(module); + decorateTypesWithLayouts(module, assocs); +} +} diff --git a/source/slang/slang-ir-layout-on-types.h b/source/slang/slang-ir-layout-on-types.h new file mode 100644 index 000000000..4cf6f887b --- /dev/null +++ b/source/slang/slang-ir-layout-on-types.h @@ -0,0 +1,8 @@ +#pragma once + +namespace Slang +{ + struct IRModule; + struct CodeGenContext; + void placeTypeLayoutsOnTypes(IRModule* module, CodeGenContext* codeGenContext); +} diff --git a/source/slang/slang-ir-specialize.cpp b/source/slang/slang-ir-specialize.cpp index f1de6e408..83e2cad70 100644 --- a/source/slang/slang-ir-specialize.cpp +++ b/source/slang/slang-ir-specialize.cpp @@ -288,17 +288,10 @@ struct SpecializationContext IRGeneric* g = generic; for (;;) { - // We can't specialize a generic if it is marked as - // being imported from an external module (in which - // case its definition is not available to us). - // - if (!isDefinition(g)) - return false; - // Given the generic `g`, we will find the value // it appears to return in its body. // - auto val = findGenericReturnVal(g); + const auto val = findGenericReturnVal(g); if (!val) return false; @@ -311,6 +304,29 @@ struct SpecializationContext continue; } + // HACK: there are conflicting features in our target intrinsic decorations + // Some reference generic parameter with the `$G0` syntax (the + // first generic parameter) while other have a predicate + // `boolean(T)` where the T is resolved properly in IR lowering. We + // need to specialize the latter and not the former. + // + // The solution is to remove the `$G` syntax and replace that with + // resolved types too. + bool intrinsicNeedsSpecialization = false; + for(const auto dec : val->getDecorations()) + { + // TODO: We should probably take into account our target to see if the intrinsic applies + if(const auto intrinsicDec = as<IRTargetIntrinsicDecoration>(dec)) + intrinsicNeedsSpecialization = intrinsicNeedsSpecialization || intrinsicDec->hasPredicate(); + } + + // We can't specialize a generic if it is marked as + // being imported from an external module (in which + // case its definition is not available to us). + // + if (!isDefinition(g) && !intrinsicNeedsSpecialization) + return false; + // We should never specialize intrinsic types. // // TODO: This logic assumes that having *any* target diff --git a/source/slang/slang-ir-spirv-legalize.cpp b/source/slang/slang-ir-spirv-legalize.cpp index cc9bf4164..0fb472df7 100644 --- a/source/slang/slang-ir-spirv-legalize.cpp +++ b/source/slang/slang-ir-spirv-legalize.cpp @@ -3,10 +3,12 @@ #include "slang-ir-glsl-legalize.h" +#include "slang-ir-layout-on-types.h" #include "slang-ir.h" #include "slang-ir-insts.h" #include "slang-emit-base.h" #include "slang-glsl-extension-tracker.h" +#include "slang-type-layout.h" namespace Slang { @@ -72,11 +74,24 @@ struct SPIRVLegalizationContext : public SourceEmitterBase storageClass = SpvStorageClassInput; } } + else if(const auto parameterGroupTypeLayout = + as<IRParameterGroupTypeLayout>(layout->getTypeLayout())) + { + storageClass = SpvStorageClassUniform; + } + } + + // Strip any HLSL wrappers + auto innerType = inst->getFullType(); + if(const auto constantBufferType = as<IRConstantBufferType>(innerType)) + { + innerType = constantBufferType->getElementType(); } + // Make a pointer type of storageClass. IRBuilder builder(m_sharedContext->m_irModule); builder.setInsertBefore(inst); - ptrType = builder.getPtrType(kIROp_PtrType, inst->getFullType(), storageClass); + ptrType = builder.getPtrType(kIROp_PtrType, innerType, storageClass); inst->setFullType(ptrType); // Insert an explicit load at each use site. List<IRUse*> uses; @@ -109,41 +124,40 @@ struct SPIRVLegalizationContext : public SourceEmitterBase return; } - auto varLayout = getVarLayout(inst); - if (!varLayout) - return; - SpvStorageClass storageClass = SpvStorageClassPrivate; - for (auto rr : varLayout->getOffsetAttrs()) + if (as<IRGroupSharedRate>(inst->getRate())) { - switch (rr->getResourceKind()) - { - case LayoutResourceKind::Uniform: - case LayoutResourceKind::ShaderResource: - case LayoutResourceKind::DescriptorTableSlot: - storageClass = SpvStorageClassUniform; - break; - case LayoutResourceKind::VaryingInput: - storageClass = SpvStorageClassInput; - break; - case LayoutResourceKind::VaryingOutput: - storageClass = SpvStorageClassOutput; - break; - case LayoutResourceKind::UnorderedAccess: - storageClass = SpvStorageClassStorageBuffer; - break; - case LayoutResourceKind::PushConstantBuffer: - storageClass = SpvStorageClassPushConstant; - break; - default: - break; - } + storageClass = SpvStorageClassWorkgroup; } - auto rate = inst->getRate(); - if (as<IRGroupSharedRate>(rate)) + else if (const auto varLayout = getVarLayout(inst)) { - storageClass = SpvStorageClassWorkgroup; + for (auto rr : varLayout->getOffsetAttrs()) + { + switch (rr->getResourceKind()) + { + case LayoutResourceKind::Uniform: + case LayoutResourceKind::ShaderResource: + case LayoutResourceKind::DescriptorTableSlot: + storageClass = SpvStorageClassUniform; + break; + case LayoutResourceKind::VaryingInput: + storageClass = SpvStorageClassInput; + break; + case LayoutResourceKind::VaryingOutput: + storageClass = SpvStorageClassOutput; + break; + case LayoutResourceKind::UnorderedAccess: + storageClass = SpvStorageClassStorageBuffer; + break; + case LayoutResourceKind::PushConstantBuffer: + storageClass = SpvStorageClassPushConstant; + break; + default: + break; + } + } } + IRBuilder builder(m_sharedContext->m_irModule); builder.setInsertBefore(inst); auto newPtrType = @@ -182,6 +196,30 @@ struct SPIRVLegalizationContext : public SourceEmitterBase } } + // Replace getElement(x, i) with, y = store(x); p = getElementPtr(y, i); load(p) + // SPIR-V has no support for dynamic indexing into values like we do. + // It may be advantageous however to do this further up the pipeline + void processGetElement(IRGetElement* inst) + { + const auto x = inst->getBase(); + List<IRInst*> indices; + IRGetElement* c = inst; + do + { + indices.add(c->getIndex()); + } while(c = as<IRGetElement>(c->getBase()), c); + IRBuilder builder(m_sharedContext->m_irModule); + builder.setInsertBefore(inst); + IRInst* y = builder.emitVar(x->getDataType(), SpvStorageClassFunction); + builder.emitStore(y, x); + for(Index i = indices.getCount() - 1; i >= 0; --i) + y = builder.emitElementAddress(y, indices[i]); + const auto newInst = builder.emitLoad(y); + inst->replaceUsesWith(newInst); + inst->removeAndDeallocate(); + addUsersToWorkList(newInst); + } + void processGetElementPtrImpl(IRInst* gepInst, IRInst* base, IRInst* index) { if (auto ptrType = as<IRPtrTypeBase>(base->getDataType())) @@ -243,13 +281,34 @@ struct SPIRVLegalizationContext : public SourceEmitterBase void processStructuredBufferType(IRHLSLStructuredBufferTypeBase* inst) { + // const auto elementTypeLayout = inst->getElementType()->findDecoration<IRTypeLayout>(); + // SLANG_ASSERT(elementTypeLayout); + const auto typeLayoutDecoration = inst->findDecoration<IRLayoutDecoration>(); + SLANG_ASSERT(typeLayoutDecoration); + const auto typeLayout = as<IRStructuredBufferTypeLayout>(typeLayoutDecoration->getLayout()); + SLANG_ASSERT(typeLayout); + const auto elemTypeLayout = typeLayout->getElementTypeLayout(); + SLANG_ASSERT(elemTypeLayout); + IRBuilder builder(m_sharedContext->m_irModule); + builder.setInsertBefore(inst); - auto arrayType = builder.getUnsizedArrayType(inst->getElementType()); - auto structType = builder.createStructType(); - auto arrayKey = builder.createStructKey(); + const auto arrayType = builder.getUnsizedArrayType(inst->getElementType()); + IRArrayTypeLayout::Builder arrayTypeLayoutBuilder(&builder, elemTypeLayout); + const auto arrayTypeLayout = arrayTypeLayoutBuilder.build(); + IRVarLayout::Builder varLayoutBuilder(&builder, arrayTypeLayout); + varLayoutBuilder.findOrAddResourceInfo(LayoutResourceKind::Uniform); + const auto arrayFieldVarLayout = varLayoutBuilder.build(); + + const auto structType = builder.createStructType(); + IRStructTypeLayout::Builder structTypeLayoutBuilder(&builder); + const auto arrayKey = builder.createStructKey(); builder.createStructField(structType, arrayKey, arrayType); - auto ptrType = builder.getPtrType(kIROp_PtrType, structType, SpvStorageClassStorageBuffer); + structTypeLayoutBuilder.addField(arrayKey, arrayFieldVarLayout); + const auto structTypeLayout = structTypeLayoutBuilder.build(); + + const auto ptrType = builder.getPtrType(kIROp_PtrType, structType, SpvStorageClassStorageBuffer); + StringBuilder nameSb; switch (inst->getOp()) { @@ -268,6 +327,8 @@ struct SPIRVLegalizationContext : public SourceEmitterBase } builder.addNameHintDecoration(structType, nameSb.getUnownedSlice()); builder.addDecoration(structType, kIROp_SPIRVBufferBlockDecoration); + SLANG_ASSERT(!structType->findDecoration<IRLayoutDecoration>()); + builder.addLayoutDecoration(structType, structTypeLayout); inst->replaceUsesWith(ptrType); inst->removeAndDeallocate(); addUsersToWorkList(ptrType); @@ -291,6 +352,9 @@ struct SPIRVLegalizationContext : public SourceEmitterBase case kIROp_Call: processCall(as<IRCall>(inst)); break; + case kIROp_GetElement: + processGetElement(as<IRGetElement>(inst)); + break; case kIROp_GetElementPtr: processGetElementPtr(as<IRGetElementPtr>(inst)); break; @@ -344,6 +408,7 @@ void legalizeIRForSPIRV( const List<IRFunc*>& entryPoints, CodeGenContext* codeGenContext) { + placeTypeLayoutsOnTypes(module, codeGenContext); GLSLExtensionTracker extensionTracker; legalizeEntryPointsForGLSL(module->getSession(), module, entryPoints, codeGenContext, &extensionTracker); legalizeSPIRV(context, module); diff --git a/source/slang/slang-ir-spirv-snippet.cpp b/source/slang/slang-ir-spirv-snippet.cpp index a263ad7f7..40a506d64 100644 --- a/source/slang/slang-ir-spirv-snippet.cpp +++ b/source/slang/slang-ir-spirv-snippet.cpp @@ -32,6 +32,8 @@ SpvSnippet::ASMType parseASMType(Slang::Misc::TokenReader& tokenReader) return SpvSnippet::ASMType::Float2; else if (word == "int") return SpvSnippet::ASMType::Int; + else if (word == "uint") + return SpvSnippet::ASMType::UInt; else if (word == "_p") return SpvSnippet::ASMType::FloatOrDouble; return SpvSnippet::ASMType::None; @@ -190,12 +192,13 @@ RefPtr<SpvSnippet> SpvSnippet::parse(UnownedStringSlice definition) operand.content = (SpvWord)0xFFFFFFFF; if (tokenReader.AdvanceIf("*")) { - // A "*" at operand qualifies the use of `resultType` with - // a storage class, but does not modify `resultType` itself. + // A "*" at operand qualifies the use of `resultType` as + // `ptr(resultType, storage class), but does + // not modify `resultType` itself. auto storageClass = tokenReader.ReadWord(); auto spvStorageClass = translateStorageClass(storageClass); operand.content = spvStorageClass; - snippet->usedResultTypeStorageClasses.add(spvStorageClass); + snippet->usedPtrResultTypeStorageClasses.add(spvStorageClass); } inst.operands.add(operand); } diff --git a/source/slang/slang-ir-spirv-snippet.h b/source/slang/slang-ir-spirv-snippet.h index e524abe3a..75edf0b38 100644 --- a/source/slang/slang-ir-spirv-snippet.h +++ b/source/slang/slang-ir-spirv-snippet.h @@ -63,6 +63,7 @@ struct SpvSnippet : public RefObject { None, Int, + UInt, Float, Double, FloatOrDouble, // Float or double type, depending on the result type of the intrinsic. @@ -110,6 +111,8 @@ struct SpvSnippet : public RefObject floatValues[1] == other.floatValues[1]; case ASMType::Int: return intValues[0] == other.intValues[0]; + case ASMType::UInt: + return intValues[0] == other.intValues[0]; case ASMType::UInt2: return intValues[0] == other.intValues[0] && intValues[1] == other.intValues[1]; default: @@ -125,7 +128,7 @@ struct SpvSnippet : public RefObject }; List<ASMInst> instructions; - List<SpvStorageClass> usedResultTypeStorageClasses; + HashSet<SpvStorageClass> usedPtrResultTypeStorageClasses; List<ASMConstant> constants; SpvStorageClass resultStorageClass = SpvStorageClassMax; diff --git a/source/slang/slang-ir.cpp b/source/slang/slang-ir.cpp index 49c91dc22..131c8d407 100644 --- a/source/slang/slang-ir.cpp +++ b/source/slang/slang-ir.cpp @@ -959,6 +959,19 @@ namespace Slang } // + // IRTupleTypeLayout + // + + void IRTupleTypeLayout::Builder::addAttrsImpl(List<IRInst*>& ioOperands) + { + auto irBuilder = getIRBuilder(); + for(auto field : m_fields) + { + ioOperands.add(irBuilder->getTupleFieldLayoutAttr(field.layout)); + } + } + + // // IRArrayTypeLayout // @@ -967,6 +980,15 @@ namespace Slang ioOperands.add(m_elementTypeLayout); } + // + // IRStructuredBufferTypeLayout + // + + void IRStructuredBufferTypeLayout::Builder::addOperandsImpl(List<IRInst*>& ioOperands) + { + ioOperands.add(m_elementTypeLayout); + } + // // IRPointerTypeLayout // @@ -4414,6 +4436,19 @@ namespace Slang return inst; } + IRVar* IRBuilder::emitVar( + IRType* type, + IRIntegerValue addressSpace) + { + auto allocatedType = getPtrType(kIROp_PtrType, type, addressSpace); + auto inst = createInst<IRVar>( + this, + kIROp_Var, + allocatedType); + addInst(inst); + return inst; + } + IRInst* IRBuilder::emitLoadReverseGradient(IRType* type, IRInst* diffValue) { auto inst = createInst<IRLoadReverseGradient>( @@ -5669,6 +5704,18 @@ namespace Slang operands)); } + IRTupleFieldLayoutAttr* IRBuilder::getTupleFieldLayoutAttr( + IRTypeLayout* layout) + { + IRInst* operands[] = { layout }; + + return cast<IRTupleFieldLayoutAttr>(createIntrinsicInst( + getVoidType(), + kIROp_TupleFieldLayoutAttr, + SLANG_COUNT_OF(operands), + operands)); + } + IRCaseTypeLayoutAttr* IRBuilder::getCaseTypeLayoutAttr( IRTypeLayout* layout) { @@ -7457,6 +7504,21 @@ namespace Slang if (decorationCaps.isIncompatibleWith(targetCaps)) continue; + if(decoration->hasPredicate()) + { + const auto scrutinee = decoration->getTypeScrutinee(); + const auto predicate = decoration->getTypePredicate(); + const auto predicateFun = + predicate == "boolean" ? [](auto t){ return t->getOp() == kIROp_BoolType; } + : predicate == "integral" ? isIntegralType + : predicate == "floating" ? isFloatingType + : nullptr; + + SLANG_ASSERT(predicateFun); + if(!predicateFun(scrutinee)) + continue; + } + if(!bestDecoration || decorationCaps.isBetterForTarget(bestCaps, targetCaps)) { bestDecoration = decoration; diff --git a/source/slang/slang-ir.h b/source/slang/slang-ir.h index c5778004b..68ca403bd 100644 --- a/source/slang/slang-ir.h +++ b/source/slang/slang-ir.h @@ -1672,6 +1672,10 @@ struct IRFuncType : IRType IRType* getResultType() { return (IRType*) getOperand(0); } UInt getParamCount() { return getOperandCount() - 1; } IRType* getParamType(UInt index) { return (IRType*)getOperand(1 + index); } + IROperandList<IRType> getParamTypes() + { + return IROperandList<IRType>(getOperands() + 1, getOperands() + getOperandCount()); + } IR_LEAF_ISA(FuncType) }; diff --git a/source/slang/slang-lower-to-ir.cpp b/source/slang/slang-lower-to-ir.cpp index d67d57001..f89f92711 100644 --- a/source/slang/slang-lower-to-ir.cpp +++ b/source/slang/slang-lower-to-ir.cpp @@ -6790,7 +6790,7 @@ struct DeclLoweringVisitor : DeclVisitor<DeclLoweringVisitor, LoweredValInfo> builder->addHighLevelDeclDecoration(irParam, decl); } - addTargetIntrinsicDecorations(irParam, decl); + addTargetIntrinsicDecorations(nullptr, irParam, decl); // A global variable's SSA value is a *pointer* to // the underlying storage. @@ -7401,7 +7401,7 @@ struct DeclLoweringVisitor : DeclVisitor<DeclLoweringVisitor, LoweredValInfo> irInterface->moveToEnd(); - addTargetIntrinsicDecorations(irInterface, decl); + addTargetIntrinsicDecorations(subContext, irInterface, decl); auto finalVal = finishOuterGenerics(subBuilder, irInterface, outerGeneric); return LoweredValInfo::simple(finalVal); @@ -7602,7 +7602,7 @@ struct DeclLoweringVisitor : DeclVisitor<DeclLoweringVisitor, LoweredValInfo> // type declarations later, from the top-level emit logic. irAggType->moveToEnd(); - addTargetIntrinsicDecorations(irAggType, decl); + addTargetIntrinsicDecorations(subContext, irAggType, decl); for (auto modifier : decl->modifiers) { if (as<NonCopyableTypeAttribute>(modifier)) @@ -7747,7 +7747,7 @@ struct DeclLoweringVisitor : DeclVisitor<DeclLoweringVisitor, LoweredValInfo> // We allow a field to be marked as a target intrinsic, // so that we can override its mangled name in the // output for the chosen target. - addTargetIntrinsicDecorations(irFieldKey, fieldDecl); + addTargetIntrinsicDecorations(nullptr, irFieldKey, fieldDecl); return LoweredValInfo::simple(irFieldKey); } @@ -8152,6 +8152,7 @@ struct DeclLoweringVisitor : DeclVisitor<DeclLoweringVisitor, LoweredValInfo> // Attach target-intrinsic decorations to an instruction, // based on modifiers on an AST declaration. void addTargetIntrinsicDecorations( + IRGenContext* subContext, IRInst* irInst, Decl* decl) { @@ -8160,14 +8161,13 @@ struct DeclLoweringVisitor : DeclVisitor<DeclLoweringVisitor, LoweredValInfo> for (auto targetMod : decl->getModifiersOfType<TargetIntrinsicModifier>()) { String definition; - auto definitionToken = targetMod->definitionToken; - if (definitionToken.type == TokenType::StringLiteral) + if(targetMod->isString) { - definition = getStringLiteralTokenValue(definitionToken); + definition = targetMod->definitionString; } - else if(definitionToken.type == TokenType::Identifier) + else if (targetMod->definitionIdent.type == TokenType::Identifier) { - definition = definitionToken.getContent(); + definition = targetMod->definitionIdent.getContent(); } else { @@ -8201,7 +8201,19 @@ struct DeclLoweringVisitor : DeclVisitor<DeclLoweringVisitor, LoweredValInfo> targetCaps = CapabilitySet(targetCap); } - builder->addTargetIntrinsicDecoration(irInst, targetCaps, definition.getUnownedSlice()); + IRInst* scrutinee = nullptr; + UnownedStringSlice predicate; + if(targetMod->scrutineeDeclRef) + { + const auto s = subContext->findLoweredDecl(targetMod->scrutineeDeclRef.getDecl()); + if(s && s->flavor == LoweredValInfo::Flavor::Simple) + { + scrutinee = s->val; + predicate = targetMod->predicateToken.getContent(); + } + } + + builder->addTargetIntrinsicDecoration(irInst, targetCaps, definition.getUnownedSlice(), predicate, scrutinee); } if(const auto nvapiMod = decl->findModifier<NVAPIMagicModifier>()) @@ -8685,7 +8697,7 @@ struct DeclLoweringVisitor : DeclVisitor<DeclLoweringVisitor, LoweredValInfo> // If this declaration was marked as having a target-specific lowering // for a particular target, then handle that here. - addTargetIntrinsicDecorations(irFunc, decl); + addTargetIntrinsicDecorations(subContext, irFunc, decl); addCatchAllIntrinsicDecorationIfNeeded(irFunc, decl); @@ -10019,6 +10031,12 @@ IRTypeLayout* lowerTypeLayout( return _lowerTypeLayoutCommon(context, &builder, paramGroupTypeLayout); } + else if( auto structuredBufferTypeLayout = as<StructuredBufferTypeLayout>(typeLayout)) + { + auto irElementTypeLayout = lowerTypeLayout(context, structuredBufferTypeLayout->elementTypeLayout); + IRStructuredBufferTypeLayout::Builder builder(context->irBuilder, irElementTypeLayout); + return _lowerTypeLayoutCommon(context, &builder, structuredBufferTypeLayout); + } else if( auto structTypeLayout = as<StructTypeLayout>(typeLayout) ) { IRStructTypeLayout::Builder builder(context->irBuilder); diff --git a/source/slang/slang-options.cpp b/source/slang/slang-options.cpp index 85602b744..216040f98 100644 --- a/source/slang/slang-options.cpp +++ b/source/slang/slang-options.cpp @@ -99,6 +99,9 @@ enum class OptionKind GLSLForceScalarLayout, EnableEffectAnnotations, + + EmitSpirvViaGLSL, + EmitSpirvDirectly, // Downstream @@ -131,7 +134,6 @@ enum class OptionKind // Experimental - EmitSpirvDirectly, FileSystem, Heterogeneous, NoMangle, @@ -506,6 +508,17 @@ void initCommandOptions(CommandOptions& options) { OptionKind::EnableEffectAnnotations, "-enable-effect-annotations", nullptr, "Enables support for legacy effect annotation syntax."}, +#if defined(SLANG_CONFIG_DEFAULT_SPIRV_DIRECT) + { OptionKind::EmitSpirvViaGLSL, "-emit-spirv-via-glsl", nullptr, + "Generate SPIR-V output by compiling generated GLSL with glslang" }, + { OptionKind::EmitSpirvDirectly, "-emit-spirv-directly", nullptr, + "Generate SPIR-V output direclty (default)" }, +#else + { OptionKind::EmitSpirvViaGLSL, "-emit-spirv-via-glsl", nullptr, + "Generate SPIR-V output by compiling generated GLSL with glslang (default)" }, + { OptionKind::EmitSpirvDirectly, "-emit-spirv-directly", nullptr, + "Generate SPIR-V output direclty rather than by compiling generated GLSL with glslang" }, +#endif }; _addOptions(makeConstArrayView(targetOpts), options); @@ -583,9 +596,6 @@ void initCommandOptions(CommandOptions& options) const Option experimentalOpts[] = { - { OptionKind::EmitSpirvDirectly, "-emit-spirv-directly", nullptr, - "Generate SPIR-V output directly (otherwise through " - "GLSL and using the glslang compiler)"}, { OptionKind::FileSystem, "-file-system", "-file-system <file-system-type>", "Set the filesystem hook to use for a compile request."}, { OptionKind::Heterogeneous, "-heterogeneous", nullptr, "Output heterogeneity-related code." }, @@ -726,7 +736,7 @@ struct OptionsParser { CodeGenTarget format = CodeGenTarget::Unknown; ProfileVersion profileVersion = ProfileVersion::Unknown; - SlangTargetFlags targetFlags = 0; + SlangTargetFlags targetFlags = kDefaultTargetFlags; int targetID = -1; FloatingPointMode floatingPointMode = FloatingPointMode::Default; bool forceGLSLScalarLayout = false; @@ -2280,7 +2290,16 @@ SlangResult OptionsParser::_parse( // We retun an error so after this has successfully passed, we quit return SLANG_FAIL; } - case OptionKind::EmitSpirvDirectly: getCurrentTarget()->targetFlags |= SLANG_TARGET_FLAG_GENERATE_SPIRV_DIRECTLY; break; + case OptionKind::EmitSpirvViaGLSL: + { + getCurrentTarget()->targetFlags &= ~SLANG_TARGET_FLAG_GENERATE_SPIRV_DIRECTLY; + } + break; + case OptionKind::EmitSpirvDirectly: + { + getCurrentTarget()->targetFlags |= SLANG_TARGET_FLAG_GENERATE_SPIRV_DIRECTLY; + } + break; case OptionKind::DefaultDownstreamCompiler: { diff --git a/source/slang/slang-parser.cpp b/source/slang/slang-parser.cpp index 80d14795c..98c227159 100644 --- a/source/slang/slang-parser.cpp +++ b/source/slang/slang-parser.cpp @@ -6412,6 +6412,7 @@ namespace Slang static NodeBase* parseTargetIntrinsicModifier(Parser* parser, void* /*userData*/) { auto modifier = parser->astBuilder->create<TargetIntrinsicModifier>(); + modifier->isString = false; if (AdvanceIf(parser, TokenType::LParent)) { @@ -6419,13 +6420,31 @@ namespace Slang if( AdvanceIf(parser, TokenType::Comma) ) { + if(parser->LookAheadToken(TokenType::LParent, 1)) + { + modifier->predicateToken = parser->ReadToken(TokenType::Identifier); + parser->ReadToken(); + modifier->scrutinee = NameLoc(parser->ReadToken(TokenType::Identifier)); + parser->ReadToken(TokenType::RParent); + parser->ReadToken(TokenType::Comma); + } if( parser->LookAheadToken(TokenType::StringLiteral) ) { - modifier->definitionToken = parser->ReadToken(); + bool first = true; + do + { + const auto t = parser->ReadToken(); + first + ? void(first = false) + : modifier->definitionString.append(" "); + modifier->definitionString.append(getStringLiteralTokenValue(t)); + modifier->isString = true; + } + while(parser->LookAheadToken(TokenType::StringLiteral)); } else { - modifier->definitionToken = parser->ReadToken(TokenType::Identifier); + modifier->definitionIdent = parser->ReadToken(TokenType::Identifier); } } diff --git a/source/slang/slang-profile.h b/source/slang/slang-profile.h index 0adb9dee9..a0284215f 100644 --- a/source/slang/slang-profile.h +++ b/source/slang/slang-profile.h @@ -16,6 +16,7 @@ namespace Slang C = SLANG_SOURCE_LANGUAGE_C, CPP = SLANG_SOURCE_LANGUAGE_CPP, CUDA = SLANG_SOURCE_LANGUAGE_CUDA, + SPIRV = SLANG_SOURCE_LANGUAGE_SPIRV, CountOf = SLANG_SOURCE_LANGUAGE_COUNT_OF, }; diff --git a/source/slang/slang-serialize-container.h b/source/slang/slang-serialize-container.h index 9ee0625bf..7cbc97aa2 100644 --- a/source/slang/slang-serialize-container.h +++ b/source/slang/slang-serialize-container.h @@ -36,7 +36,7 @@ struct SerialContainerData struct Target { CodeGenTarget codeGenTarget = CodeGenTarget::Unknown; - SlangTargetFlags flags = 0; + SlangTargetFlags flags = kDefaultTargetFlags; Profile profile; FloatingPointMode floatingPointMode = FloatingPointMode::Default; }; diff --git a/source/slang/slang-spirv-val.cpp b/source/slang/slang-spirv-val.cpp index 54bf5348b..c2381fa0e 100644 --- a/source/slang/slang-spirv-val.cpp +++ b/source/slang/slang-spirv-val.cpp @@ -3,27 +3,44 @@ namespace Slang { -SlangResult debugDisassembleSPIRV(const List<uint8_t>& spirv, String& outDis) +static SlangResult disassembleSPIRV(const List<uint8_t>& spirv, String& outErr, String& outDis) { + // Set up our process CommandLine commandLine; commandLine.m_executableLocation.setName("spirv-dis"); RefPtr<Process> p; - const auto createResult = Process::create(commandLine, 0, p); + // If we failed to even start the process, then validation isn't available - SLANG_RETURN_ON_FAIL(createResult); + SLANG_RETURN_ON_FAIL(Process::create(commandLine, 0, p)); const auto in = p->getStream(StdStreamType::In); const auto out = p->getStream(StdStreamType::Out); + const auto err = p->getStream(StdStreamType::ErrorOut); + // Write the assembly SLANG_RETURN_ON_FAIL(in->write(spirv.getBuffer(), spirv.getCount())); in->close(); + // Wait for it to finish - if (!p->waitForTermination(1000)) + if(!p->waitForTermination(1000)) return SLANG_FAIL; + // TODO: allow inheriting stderr in Process List<Byte> outData; SLANG_RETURN_ON_FAIL(StreamUtil::readAll(out, 0, outData)); - outDis = String((const char*)outData.getBuffer()); - return SLANG_OK; + outErr = String( + reinterpret_cast<const char*>(outData.begin()), + reinterpret_cast<const char*>(outData.end()) + ); + + outData.clear(); + SLANG_RETURN_ON_FAIL(StreamUtil::readAll(err, 0, outData)); + outDis = String( + reinterpret_cast<const char*>(outData.begin()), + reinterpret_cast<const char*>(outData.end()) + ); + + const auto ret = p->getReturnValue(); + return ret == 0 ? SLANG_OK : SLANG_FAIL; } SlangResult debugValidateSPIRV(const List<uint8_t>& spirv) @@ -57,10 +74,13 @@ SlangResult debugValidateSPIRV(const List<uint8_t>& spirv) SLANG_RETURN_ON_FAIL(StreamUtil::readAll(err, 0, outData)); fwrite(outData.getBuffer(), outData.getCount(), 1, stderr); const auto ret = p->getReturnValue(); - if (SLANG_FAILED(ret)) + + if(ret != 0) { + String spirvDisErr; String spirvDis; - debugDisassembleSPIRV(spirv, spirvDis); + disassembleSPIRV(spirv, spirvDisErr, spirvDis); + fwrite(spirvDisErr.getBuffer(), spirvDisErr.getLength(), 1, stderr); fwrite(spirvDis.getBuffer(), spirvDis.getLength(), 1, stderr); } diff --git a/tests/autodiff/custom-intrinsic.slang b/tests/autodiff/custom-intrinsic.slang index dd122a674..1954c6f8d 100644 --- a/tests/autodiff/custom-intrinsic.slang +++ b/tests/autodiff/custom-intrinsic.slang @@ -15,7 +15,7 @@ namespace myintrinsiclib __target_intrinsic(glsl, "exp($0)") __target_intrinsic(cuda, "$P_exp($0)") __target_intrinsic(cpp, "$P_exp($0)") - __target_intrinsic(spirv_direct, "12 resultType resultId glsl450 27 _0") + __target_intrinsic(spirv, "12 resultType resultId glsl450 27 _0") [ForwardDerivative(d_myexp<T>)] T myexp(T x); @@ -34,7 +34,7 @@ namespace myintrinsiclib __target_intrinsic(glsl, "sin($0)") __target_intrinsic(cuda, "$P_sin($0)") __target_intrinsic(cpp, "$P_sin($0)") - __target_intrinsic(spirv_direct, "12 resultType resultId glsl450 13 _0") + __target_intrinsic(spirv, "12 resultType resultId glsl450 13 _0") [ForwardDerivative(d_mysin<T>)] T mysin(T x); @@ -52,7 +52,7 @@ namespace myintrinsiclib __target_intrinsic(glsl, "cos($0)") __target_intrinsic(cuda, "$P_cos($0)") __target_intrinsic(cpp, "$P_cos($0)") - __target_intrinsic(spirv_direct, "12 resultType resultId glsl450 14 _0") + __target_intrinsic(spirv, "12 resultType resultId glsl450 14 _0") [ForwardDerivative(d_mycos<T>)] T mycos(T x); diff --git a/tests/bugs/gh-3085.slang b/tests/bugs/gh-3085.slang index 9250d501d..86992aec6 100644 --- a/tests/bugs/gh-3085.slang +++ b/tests/bugs/gh-3085.slang @@ -1,5 +1,5 @@ //TEST:SIMPLE(filecheck=CHECK): -entry MainCs -stage compute -profile glsl_450 -target spirv -//CHECK: EntryPoint +//CHECK: OpEntryPoint RWTexture2D<float4> g_Test; @@ -26,4 +26,4 @@ void MainCs(uint3 groupID: SV_GroupID, uint3 groupThreadID: SV_GroupThreadID) // color+= LoadSourceColor ( pixelPos , int2 ( i%3 , i/3 ) , msaaSampleIndex ) . rgb ; //} g_Test[int2(pixelPos.x / 2, pixelPos.y + 0)] = float4(color, 1.0); -}
\ No newline at end of file +} diff --git a/tests/bugs/gh-841.slang b/tests/bugs/gh-841.slang index 8dc687e77..0608023e5 100644 --- a/tests/bugs/gh-841.slang +++ b/tests/bugs/gh-841.slang @@ -12,9 +12,9 @@ struct RasterVertex float4 c : COLOR; // Make sure that the input value in location 1 is decorated as Flat - // SPV-DAG: [[#VAL:]]{{.*}}:{{.*}} Variable Input - // SPV-DAG: Decorate [[#VAL]]{{.*}} Location 1 - // SPV-DAG: Decorate [[#VAL]]{{.*}} Flat + // SPV-DAG: [[VAL:%[_A-Za-z0-9]+]] = OpVariable {{.*}} Input + // SPV-DAG: OpDecorate [[VAL]] Location 1 + // SPV-DAG: OpDecorate [[VAL]] Flat // // Likewise for GLSL // GLSL: flat layout(location = 1) diff --git a/tests/bugs/interlocked-add-uint-int.slang b/tests/bugs/interlocked-add-uint-int.slang index 10baa7f17..4396a3dc5 100644 --- a/tests/bugs/interlocked-add-uint-int.slang +++ b/tests/bugs/interlocked-add-uint-int.slang @@ -1,5 +1,5 @@ //TEST:SIMPLE(filecheck=CHECK): -target spirv -profile glsl_450 -stage compute -entry MainCs -line-directive-mode none -//CHECK: {{.*}} AtomicIAdd +//CHECK: {{.*}} OpAtomicIAdd RWBuffer<uint> g_InterlockTest; [numthreads(1,1,1)] @@ -16,4 +16,4 @@ void MainCs( uint2 dispatchThreadID : SV_DispatchThreadID, uint2 groupThreadID : { InterlockedAdd( g_InterlockTest[ 0 ], -int(nVertexCount) ); } -}
\ No newline at end of file +} diff --git a/tests/bugs/texture-array-samplecmplevelzero.slang b/tests/bugs/texture-array-samplecmplevelzero.slang index 6176f74dc..c8afba603 100644 --- a/tests/bugs/texture-array-samplecmplevelzero.slang +++ b/tests/bugs/texture-array-samplecmplevelzero.slang @@ -1,5 +1,5 @@ //TEST_DISABLED:SIMPLE(filecheck=CHECK): -entry MainPs -stage fragment -profile glsl_450 -target spirv -//CHECK: EntryPoint +//CHECK: OpEntryPoint SamplerComparisonState g_tSFMShadowDepthTexture_sampler; Texture2DArray g_tSFMShadowDepthTexture; diff --git a/tests/compute/simple.slang b/tests/compute/simple.slang index 0d4efc2e2..96cccc090 100644 --- a/tests/compute/simple.slang +++ b/tests/compute/simple.slang @@ -1,6 +1,7 @@ //TEST(smoke,compute):COMPARE_COMPUTE:-shaderobj //TEST(smoke,compute):COMPARE_COMPUTE:-dx12 -use-dxil -shaderobj //TEST(smoke,compute):COMPARE_COMPUTE:-cpu -shaderobj +//TEST(smoke,compute):COMPARE_COMPUTE:-vk -shaderobj //TEST(smoke,compute):COMPARE_COMPUTE(filecheck-buffer=CHECK):-cpu -shaderobj // CHECK: 0 diff --git a/tests/compute/structured-buffer-of-struct.slang b/tests/compute/structured-buffer-of-struct.slang new file mode 100644 index 000000000..7bfb38df8 --- /dev/null +++ b/tests/compute/structured-buffer-of-struct.slang @@ -0,0 +1,37 @@ +//TEST(compute, vulkan):COMPARE_COMPUTE(filecheck-buffer=CHECK):-vk -shaderobj + +// CHECK: 2 +// CHECK-NEXT: 1 +// CHECK-NEXT: 4 +// CHECK-NEXT: 3 +// CHECK-NEXT: 6 +// CHECK-NEXT: 5 +// CHECK-NEXT: 8 +// CHECK-NEXT: 7 + +struct S +{ + int a; + int b; +}; + +//TEST_INPUT:ubuffer(data=[1 2 3 4 5 6 7 8], stride=8):out,name=buf +RWStructuredBuffer<S> buf; + +S test(S s) +{ + s.a = s.a ^ s.b; + s.b = s.a ^ s.b; + s.a = s.a ^ s.b; + return s; +} + +[numthreads(4, 1, 1)] +// void computeMain(uint i : SV_GroupIndex) +// { +// buf[i] = test(buf[i]); +// } +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + buf[dispatchThreadID.x] = test(buf[dispatchThreadID.x]); +} diff --git a/tests/compute/texture-simpler.slang b/tests/compute/texture-simpler.slang new file mode 100644 index 000000000..ce4b76f77 --- /dev/null +++ b/tests/compute/texture-simpler.slang @@ -0,0 +1,16 @@ +//TEST(compute, vulkan):COMPARE_COMPUTE:-vk -compute -shaderobj -output-using-type + +//TEST_INPUT: Texture2D(size=4, content = one):name t2D +Texture2D<float> t2D; +//TEST_INPUT: Sampler:name samplerState +SamplerState samplerState; + +//TEST_INPUT: ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer +RWStructuredBuffer<float> outputBuffer; + +[numthreads(4, 1, 1)] +void computeMain(uint i : SV_GroupIndex) +{ + float u = i * 0.25; + outputBuffer[i] = t2D.SampleLevel(samplerState, float2(u, u), 0); +} diff --git a/tests/compute/texture-simpler.slang.expected.txt b/tests/compute/texture-simpler.slang.expected.txt new file mode 100644 index 000000000..334dd67a8 --- /dev/null +++ b/tests/compute/texture-simpler.slang.expected.txt @@ -0,0 +1,5 @@ +type: float +1.000000 +1.000000 +1.000000 +1.000000 diff --git a/tests/cross-compile/dxc-error.hlsl b/tests/cross-compile/dxc-error.hlsl index 0e6b6e489..8d73d7a1c 100644 --- a/tests/cross-compile/dxc-error.hlsl +++ b/tests/cross-compile/dxc-error.hlsl @@ -1,4 +1,4 @@ -//DIAGNOSTIC_TEST(dxc):SIMPLE:-pass-through dxc -target dxil -entry computeMain -stage compute -profile sm_6_1 +//DIAGNOSTIC_TEST(dxc):SIMPLE(filecheck=CHECK):-pass-through dxc -target dxil -entry computeMain -stage compute -profile sm_6_1 [numthreads(4, 1, 1)] void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) @@ -6,4 +6,7 @@ void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) uint tid = dispatchThreadID.x; // Error should be here... as gOutputBuffer is not defined... gOutputBuffer[tid] = dispatchThreadID.x * 0.5f; -}
\ No newline at end of file + // CHECK: : tests/cross-compile/dxc-error.hlsl([[#@LINE-1]]): error : use of undeclared identifier 'gOutputBuffer' + // CHECK-NEXT: : note : gOutputBuffer[tid] = dispatchThreadID.x * 0.5f; + // CHECK-NEXT: : note : ^ +} diff --git a/tests/cross-compile/dxc-error.hlsl.expected b/tests/cross-compile/dxc-error.hlsl.expected deleted file mode 100644 index c47ca80a6..000000000 --- a/tests/cross-compile/dxc-error.hlsl.expected +++ /dev/null @@ -1,8 +0,0 @@ -result code = -1 -standard error = { -dxc 1.7: tests/cross-compile/dxc-error.hlsl(8): error : use of undeclared identifier 'gOutputBuffer' -dxc 1.7: note : gOutputBuffer[tid] = dispatchThreadID.x * 0.5f; -dxc 1.7: note : ^ -} -standard output = { -} diff --git a/tests/gpu-feature/texture/query/footprint/nv-shader-texture-footprint-simple.slang b/tests/gpu-feature/texture/query/footprint/nv-shader-texture-footprint-simple.slang index b219ca829..782c7e7e5 100644 --- a/tests/gpu-feature/texture/query/footprint/nv-shader-texture-footprint-simple.slang +++ b/tests/gpu-feature/texture/query/footprint/nv-shader-texture-footprint-simple.slang @@ -1,7 +1,9 @@ -//TEST:SIMPLE(filecheck=SPIRV):-target spirv-assembly -entry fragmentMain -stage fragment +//DISABLED_TEST:SIMPLE(filecheck=SPIRV):-target spirv-assembly -entry fragmentMain -stage fragment //DISABLED_TEST:SIMPLE(filecheck=DXIL):-target dxil-assembly -entry fragmentMain -stage fragment //TEST:SIMPLE(filecheck=HLSL):-target hlsl -entry fragmentMain -stage fragment +// TODO: The SPIRV test is fine locally, but seems to lead to an impossible to debug hang in CI + //DISABLED_TEST:SIMPLE:-target spirv-assembly -entry fragmentMain -stage fragment //DISABLED_TEST:SIMPLE:-target dxil-assembly -entry fragmentMain -stage fragment //DISABLED_TEST:SIMPLE:-target hlsl -entry fragmentMain -stage fragment @@ -62,8 +64,8 @@ void fragmentMain( accumulate(r, texture.queryFootprintCoarse(granularity, sampler, coords)); -// SPIRV: Extension "SPV_NV_shader_image_footprint" -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpExtension "SPV_NV_shader_image_footprint" +// SPIRV: OpImageSampleFootprintNV // HLSL: NvFootprintCoarse diff --git a/tests/gpu-feature/texture/query/footprint/nv-shader-texture-footprint.slang b/tests/gpu-feature/texture/query/footprint/nv-shader-texture-footprint.slang index 12ccd59fc..d5e406615 100644 --- a/tests/gpu-feature/texture/query/footprint/nv-shader-texture-footprint.slang +++ b/tests/gpu-feature/texture/query/footprint/nv-shader-texture-footprint.slang @@ -1,7 +1,9 @@ -//TEST:SIMPLE(filecheck=SPIRV):-target spirv-assembly -entry fragmentMain -stage fragment -DENABLE_CLAMP +//DISABLED_TEST:SIMPLE(filecheck=SPIRV):-target spirv-assembly -entry fragmentMain -stage fragment -DENABLE_CLAMP //DISABLED_TEST:SIMPLE(filecheck=DXIL):-target dxil-assembly -entry fragmentMain -stage fragment //TEST:SIMPLE(filecheck=HLSL):-target hlsl -entry fragmentMain -stage fragment +// TODO: The SPIRV test is fine locally, but seems to lead to an impossible to debug hang in CI + //DISABLED_TEST:SIMPLE:-target spirv-assembly -entry fragmentMain -stage fragment -DENABLE_CLAMP //DISABLED_TEST:SIMPLE:-target dxil-assembly -entry fragmentMain -stage fragment //DISABLED_TEST:SIMPLE:-target hlsl -entry fragmentMain -stage fragment @@ -74,19 +76,19 @@ void accumulate_3D(inout uint r, TextureFootprint3D f) accumulate(r, f.isSingleLevel); } -// SPRIV: Extension "SPV_NV_shader_image_footprint" +// SPRIV: OpExtension "SPV_NV_shader_image_footprint" void fragmentMain( float v : VARYING, uniform Texture2D texture_2D, -// SPIRV: TypeImage +// SPIRV: OpTypeImage // SPIRV-SAME: 2D // HLSL-DAG: Texture2D{{.*}}register(t0) uniform Texture3D texture_3D, -// SPIRV: TypeImage +// SPIRV: OpTypeImage // SPIRV-SAME: 3D // HLSL-DAG: Texture3D{{.*}}register(t1) @@ -135,8 +137,8 @@ CLAMP( accumulate##ND(r, texture##ND.queryFootprint##COARSE_OR_FINE##GradClamp TEST_COMMON(_3D, COARSE_OR_FINE) \ /* end */ -// SPIRV-DAG: [[TRUE:[0-9]+]]:{{.*}}ConstantTrue -// SPIRV-DAG: [[FALSE:[0-9]+]]:{{.*}}ConstantFalse +// SPIRV-DAG: [[TRUE:%[a-zA-Z0-9_]+]] = OpConstantTrue +// SPIRV-DAG: [[FALSE:%[a-zA-Z0-9_]+]] = OpConstantFalse { let coords = c.xy; @@ -145,57 +147,57 @@ CLAMP( accumulate##ND(r, texture##ND.queryFootprint##COARSE_OR_FINE##GradClamp TEST_2D(Coarse); -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[TRUE]] // HLSL-DAG: NvFootprintCoarse{{.*}}NV_EXTN_TEXTURE_2D -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[TRUE]] Bias // HLSL-DAG: NvFootprintCoarseBias{{.*}}NV_EXTN_TEXTURE_2D -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[TRUE]] MinLod -// SPIRV: ImageSampleFootprintNV -// SPIRV-SAME: [[TRUE]] Bias MinLod +// SPIRV: OpImageSampleFootprintNV +// SPIRV-SAME: [[TRUE]] Bias|MinLod -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[TRUE]] Lod // HLSL-DAG: NvFootprintCoarseLevel{{.*}}NV_EXTN_TEXTURE_2D -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[TRUE]] Grad // HLSL-DAG: NvFootprintCoarseGrad{{.*}}NV_EXTN_TEXTURE_2D -// SPIRV: ImageSampleFootprintNV -// SPIRV-SAME: [[TRUE]] Grad MinLod +// SPIRV: OpImageSampleFootprintNV +// SPIRV-SAME: [[TRUE]] Grad|MinLod TEST_2D(Fine); -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[FALSE]] // HLSL-DAG: NvFootprintFine{{.*}}NV_EXTN_TEXTURE_2D -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[FALSE]] Bias // HLSL-DAG: NvFootprintFineBias{{.*}}NV_EXTN_TEXTURE_2D -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[FALSE]] MinLod -// SPIRV: ImageSampleFootprintNV -// SPIRV-SAME: [[FALSE]] Bias MinLod +// SPIRV: OpImageSampleFootprintNV +// SPIRV-SAME: [[FALSE]] Bias|MinLod -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[FALSE]] Lod // HLSL-DAG: NvFootprintFineLevel{{.*}}NV_EXTN_TEXTURE_2D -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[FALSE]] Grad // HLSL-DAG: NvFootprintFineGrad{{.*}}NV_EXTN_TEXTURE_2D -// SPIRV: ImageSampleFootprintNV -// SPIRV-SAME: [[FALSE]] Grad MinLod +// SPIRV: OpImageSampleFootprintNV +// SPIRV-SAME: [[FALSE]] Grad|MinLod } @@ -206,41 +208,41 @@ CLAMP( accumulate##ND(r, texture##ND.queryFootprint##COARSE_OR_FINE##GradClamp TEST_3D(Coarse); -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[TRUE]] // HLSL-DAG: NvFootprintCoarse{{.*}}NV_EXTN_TEXTURE_3D -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[TRUE]] Bias // HLSL-DAG: NvFootprintCoarseBias{{.*}}NV_EXTN_TEXTURE_3D -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[TRUE]] MinLod -// SPIRV: ImageSampleFootprintNV -// SPIRV-SAME: [[TRUE]] Bias MinLod +// SPIRV: OpImageSampleFootprintNV +// SPIRV-SAME: [[TRUE]] Bias|MinLod -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[TRUE]] Lod // HLSL-DAG: NvFootprintCoarseLevel{{.*}}NV_EXTN_TEXTURE_3D TEST_3D(Fine); -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[FALSE]] // HLSL-DAG: NvFootprintFine{{.*}}NV_EXTN_TEXTURE_3D -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[FALSE]] Bias // HLSL-DAG: NvFootprintFineBias{{.*}}NV_EXTN_TEXTURE_3D -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[FALSE]] MinLod -// SPIRV: ImageSampleFootprintNV -// SPIRV-SAME: [[FALSE]] Bias MinLod +// SPIRV: OpImageSampleFootprintNV +// SPIRV-SAME: [[FALSE]] Bias|MinLod -// SPIRV: ImageSampleFootprintNV +// SPIRV: OpImageSampleFootprintNV // SPIRV-SAME: [[FALSE]] Lod // HLSL-DAG: NvFootprintFineLevel{{.*}}NV_EXTN_TEXTURE_3D diff --git a/tests/hlsl-intrinsic/const-buffer-pointer.slang b/tests/hlsl-intrinsic/const-buffer-pointer.slang index 33b538a60..eddfc3be0 100644 --- a/tests/hlsl-intrinsic/const-buffer-pointer.slang +++ b/tests/hlsl-intrinsic/const-buffer-pointer.slang @@ -1,7 +1,7 @@ //TEST:SIMPLE(filecheck=CHECK):-target glsl -profile glsl_450 -entry main -stage compute //TEST:SIMPLE(filecheck=SPV):-target spirv -profile glsl_450 -entry main -stage compute -// SPV: EntryPoint GLCompute {{.*}} "main" {{.*}} +// SPV: OpEntryPoint GLCompute {{.*}} "main" {{.*}} struct MyStruct { @@ -36,4 +36,4 @@ void main(int3 tid: SV_DispatchThreadID) { MyStruct s = gGlobals.pStruct.get(); outputBuffer[tid.x] = uint(s.position.x); -}
\ No newline at end of file +} diff --git a/tests/hlsl-intrinsic/scalar-float.slang b/tests/hlsl-intrinsic/scalar-float.slang index bc9d24a0e..4623d4835 100644 --- a/tests/hlsl-intrinsic/scalar-float.slang +++ b/tests/hlsl-intrinsic/scalar-float.slang @@ -2,7 +2,7 @@ //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -shaderobj //TEST(compute):COMPARE_COMPUTE_EX:-slang -compute -dx12 -shaderobj //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -shaderobj -//DISABLED_TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -shaderobj -emit-spirv-directly +//DISABLED_TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -shaderobj //TEST(compute, vulkan):COMPARE_COMPUTE_EX:-cuda -compute -shaderobj //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name outputBuffer diff --git a/tests/spirv/direct-spirv-compute-simple.slang b/tests/spirv/direct-spirv-compute-simple.slang index 2fa2798e5..4c33d8a72 100644 --- a/tests/spirv/direct-spirv-compute-simple.slang +++ b/tests/spirv/direct-spirv-compute-simple.slang @@ -1,6 +1,6 @@ // direct-spirv-compute-simple.slang -//TESTD:SIMPLE:-target spirv -entry computeMain -stage compute -emit-spirv-directly -//TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -emit-spirv-directly +//TESTD:SIMPLE:-target spirv -entry computeMain -stage compute +//TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute // Test runinng a shader generated from direct SPIR-V emit. diff --git a/tests/spirv/direct-spirv-control-flow-2.slang b/tests/spirv/direct-spirv-control-flow-2.slang index 7dd829bda..697125500 100644 --- a/tests/spirv/direct-spirv-control-flow-2.slang +++ b/tests/spirv/direct-spirv-control-flow-2.slang @@ -1,7 +1,7 @@ // direct-spirv-control-flow-2.slang -//TESTD:SIMPLE:-target spirv -entry computeMain -stage compute -emit-spirv-directly -//TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -output-using-type -emit-spirv-directly +//TESTD:SIMPLE:-target spirv -entry computeMain -stage compute +//TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -output-using-type // Test direct SPIR-V emit on control flows. diff --git a/tests/spirv/direct-spirv-control-flow.slang b/tests/spirv/direct-spirv-control-flow.slang index 10bee2522..744519438 100644 --- a/tests/spirv/direct-spirv-control-flow.slang +++ b/tests/spirv/direct-spirv-control-flow.slang @@ -1,6 +1,6 @@ // direct-spirv-control-flow.slang -//TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute -emit-spirv-directly +//TEST(compute, vulkan):COMPARE_COMPUTE_EX:-vk -compute // Test direct SPIRV emit on control fl. diff --git a/tests/spirv/direct-spirv-emit.slang.expected b/tests/spirv/direct-spirv-emit.slang.expected index aff2198be..55cc718a8 100644 --- a/tests/spirv/direct-spirv-emit.slang.expected +++ b/tests/spirv/direct-spirv-emit.slang.expected @@ -2,19 +2,26 @@ result code = 0 standard error = { } standard output = { -// Module Version 10500 -// Generated by (magic number): 0 -// Id's are bound by 5 +; SPIR-V +; Version: 1.5 +; Generator: Khronos; 0 +; Bound: 5 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %computeMain "computeMain" + OpExecutionMode %computeMain LocalSize 4 1 1 - Capability Shader - MemoryModel Logical GLSL450 - EntryPoint GLCompute 2 "computeMain" - ExecutionMode 2 LocalSize 4 1 1 - Name 2 "computeMain" - 1: TypeVoid - 3: TypeFunction 1 - 2(computeMain): 1 Function None 3 - 4: Label - Return - FunctionEnd + ; Debug Information + OpName %computeMain "computeMain" ; id %2 + + ; Types, variables and constants + %void = OpTypeVoid + %3 = OpTypeFunction %void + + ; Function computeMain +%computeMain = OpFunction %void None %3 + %4 = OpLabel + OpReturn + OpFunctionEnd } diff --git a/tools/gfx/vulkan/vk-shader-program.cpp b/tools/gfx/vulkan/vk-shader-program.cpp index ee307ed04..caa219c99 100644 --- a/tools/gfx/vulkan/vk-shader-program.cpp +++ b/tools/gfx/vulkan/vk-shader-program.cpp @@ -74,9 +74,10 @@ Result ShaderProgramImpl::createShaderModule( // to control the entry point naming behavior in SPIRV-direct path // so we can remove the ad-hoc logic here. auto realEntryPointName = entryPointInfo->getNameOverride(); - const char* spirvBinaryEntryPointName = "main"; - if (m_device->m_desc.slang.targetFlags & SLANG_TARGET_FLAG_GENERATE_SPIRV_DIRECTLY) - spirvBinaryEntryPointName = realEntryPointName; + const char* spirvBinaryEntryPointName = + m_device->m_desc.slang.targetFlags & SLANG_TARGET_FLAG_GENERATE_SPIRV_DIRECTLY + ? realEntryPointName + : "main"; m_stageCreateInfos.add(compileEntryPoint( spirvBinaryEntryPointName, kernelCode, diff --git a/tools/render-test/options.cpp b/tools/render-test/options.cpp index 3b2c07ce6..2f304b7c0 100644 --- a/tools/render-test/options.cpp +++ b/tools/render-test/options.cpp @@ -143,7 +143,11 @@ static gfx::DeviceType _toRenderType(Slang::RenderApiType apiType) } else if (argValue == "-emit-spirv-directly") { - outOptions.generateSPIRVDirectly = true; + outOptions.generateSPIRVDirectly= true; + } + else if (argValue == "-emit-spirv-via-glsl") + { + outOptions.generateSPIRVDirectly= false; } else if (argValue == "-only-startup") { diff --git a/tools/render-test/options.h b/tools/render-test/options.h index 22b14c862..c2466f671 100644 --- a/tools/render-test/options.h +++ b/tools/render-test/options.h @@ -71,7 +71,11 @@ struct Options Slang::DownstreamArgs downstreamArgs; ///< Args to downstream tools. Here it's just slang +#if defined(SLANG_CONFIG_DEFAULT_SPIRV_DIRECT) + bool generateSPIRVDirectly = true; +#else bool generateSPIRVDirectly = false; +#endif Options() { downstreamArgs.addName("slang"); } |