summaryrefslogtreecommitdiffstats
path: root/source
diff options
context:
space:
mode:
Diffstat (limited to 'source')
-rw-r--r--source/slang-record-replay/record/slang-session.cpp33
-rw-r--r--source/slang-record-replay/record/slang-session.h5
-rw-r--r--source/slang/slang-compiler.h8
-rw-r--r--source/slang/slang-ir-insts-stable-names.cpp61
-rw-r--r--source/slang/slang-ir-insts-stable-names.h10
-rw-r--r--source/slang/slang-ir-insts-stable-names.lua672
-rw-r--r--source/slang/slang-ir-insts.lua82
-rw-r--r--source/slang/slang-ir.h28
-rw-r--r--source/slang/slang-options.cpp83
-rw-r--r--source/slang/slang-serialize-container.cpp3
-rw-r--r--source/slang/slang-serialize-ir.cpp162
-rw-r--r--source/slang/slang-serialize-ir.h8
-rw-r--r--source/slang/slang.cpp46
13 files changed, 1177 insertions, 24 deletions
diff --git a/source/slang-record-replay/record/slang-session.cpp b/source/slang-record-replay/record/slang-session.cpp
index 800d690fa..254e1b091 100644
--- a/source/slang-record-replay/record/slang-session.cpp
+++ b/source/slang-record-replay/record/slang-session.cpp
@@ -92,6 +92,39 @@ SLANG_NO_THROW slang::IModule* SessionRecorder::loadModuleFromIRBlob(
return static_cast<slang::IModule*>(pModuleRecorder);
}
+SLANG_NO_THROW SlangResult SLANG_MCALL SessionRecorder::loadModuleInfoFromIRBlob(
+ slang::IBlob* source,
+ SlangInt& outModuleVersion,
+ const char*& outModuleCompilerVersion,
+ const char*& outModuleName)
+{
+ slangRecordLog(LogLevel::Verbose, "%s\n", __PRETTY_FUNCTION__);
+
+ ParameterRecorder* recorder{};
+ {
+ recorder = m_recordManager->beginMethodRecord(
+ ApiCallId::ISession_loadModuleFromIRBlob,
+ m_sessionHandle);
+ recorder->recordPointer(source);
+ recorder = m_recordManager->endMethodRecord();
+ }
+
+ const auto result = m_actualSession->loadModuleInfoFromIRBlob(
+ source,
+ outModuleVersion,
+ outModuleCompilerVersion,
+ outModuleName);
+
+ {
+ recorder->recordInt64(outModuleVersion);
+ recorder->recordString(outModuleCompilerVersion);
+ recorder->recordString(outModuleName);
+ m_recordManager->apendOutput();
+ }
+
+ return result;
+}
+
SLANG_NO_THROW slang::IModule* SessionRecorder::loadModuleFromSource(
const char* moduleName,
const char* path,
diff --git a/source/slang-record-replay/record/slang-session.h b/source/slang-record-replay/record/slang-session.h
index 9cff7beac..26c591b1c 100644
--- a/source/slang-record-replay/record/slang-session.h
+++ b/source/slang-record-replay/record/slang-session.h
@@ -29,6 +29,11 @@ public:
const char* path,
slang::IBlob* source,
slang::IBlob** outDiagnostics = nullptr) override;
+ SLANG_NO_THROW SlangResult SLANG_MCALL loadModuleInfoFromIRBlob(
+ slang::IBlob* source,
+ SlangInt& outModuleVersion,
+ const char*& outModuleCompilerVersion,
+ const char*& outModuleName) override;
SLANG_NO_THROW slang::IModule* SLANG_MCALL loadModuleFromSource(
const char* moduleName,
const char* path,
diff --git a/source/slang/slang-compiler.h b/source/slang/slang-compiler.h
index 7cdd1614c..6054492bc 100644
--- a/source/slang/slang-compiler.h
+++ b/source/slang/slang-compiler.h
@@ -2251,6 +2251,11 @@ public:
const char* path,
slang::IBlob* source,
slang::IBlob** outDiagnostics = nullptr) override;
+ SLANG_NO_THROW SlangResult SLANG_MCALL loadModuleInfoFromIRBlob(
+ slang::IBlob* source,
+ SlangInt& outModuleVersion,
+ const char*& outModuleCompilerVersion,
+ const char*& outModuleName) override;
SLANG_NO_THROW slang::IModule* SLANG_MCALL loadModuleFromSource(
const char* moduleName,
const char* path,
@@ -2344,6 +2349,9 @@ public:
SourceManager* m_sourceManager = nullptr;
RefPtr<CommandLineContext> m_cmdLineContext;
+ // Used to store strings returned by the api as const char*
+ StringSlicePool m_stringSlicePool;
+
// Name pool for looking up names
NamePool namePool;
diff --git a/source/slang/slang-ir-insts-stable-names.cpp b/source/slang/slang-ir-insts-stable-names.cpp
new file mode 100644
index 000000000..2ac205faf
--- /dev/null
+++ b/source/slang/slang-ir-insts-stable-names.cpp
@@ -0,0 +1,61 @@
+#include "slang-ir-insts-stable-names.h"
+
+namespace Slang
+{
+
+static const UInt kOpcodeToStableName[] = {
+#if 0 // FIDDLE TEMPLATE:
+% local insts = require("source/slang/slang-ir-insts.lua")
+% insts.traverse(function(inst)
+% if inst.is_leaf then
+% if inst.stable_name == nil then
+% error("Instruction is missing stable name: " .. tostring(inst.struct_name))
+% end
+% local stable_name = tostring(inst.stable_name)
+ $stable_name,
+% end
+% end)
+#else // FIDDLE OUTPUT:
+#define FIDDLE_GENERATED_OUTPUT_ID 0
+#include "slang-ir-insts-stable-names.cpp.fiddle"
+#endif // FIDDLE END
+};
+
+static const IROp kStableNameToOpcode[] = {
+#if 0 // FIDDLE TEMPLATE:
+% local insts = require("source/slang/slang-ir-insts.lua")
+% for i = 0, insts.max_stable_name do
+% inst = insts.stable_name_to_inst[i]
+% if inst then
+% local struct_name = inst.struct_name
+ kIROp_$struct_name,
+% else
+ kIROp_Invalid,
+% end
+% end
+#else // FIDDLE OUTPUT:
+#define FIDDLE_GENERATED_OUTPUT_ID 1
+#include "slang-ir-insts-stable-names.cpp.fiddle"
+#endif // FIDDLE END
+};
+
+UInt getOpcodeStableName(IROp op)
+{
+ // Check if the opcode is within valid range
+ if (op >= SLANG_COUNT_OF(kOpcodeToStableName))
+ {
+ return kInvalidStableName;
+ }
+ return kOpcodeToStableName[op];
+}
+
+IROp getStableNameOpcode(UInt stableName)
+{
+ // Check if the stable name is within valid range
+ if (stableName >= SLANG_COUNT_OF(kStableNameToOpcode))
+ {
+ return kIROp_Invalid;
+ }
+ return kStableNameToOpcode[stableName];
+}
+} // namespace Slang
diff --git a/source/slang/slang-ir-insts-stable-names.h b/source/slang/slang-ir-insts-stable-names.h
new file mode 100644
index 000000000..3e364b3bf
--- /dev/null
+++ b/source/slang/slang-ir-insts-stable-names.h
@@ -0,0 +1,10 @@
+#pragma once
+
+#include "core/slang-common.h"
+#include "slang-ir-insts-enum.h"
+namespace Slang
+{
+const UInt kInvalidStableName = ~0u;
+UInt getOpcodeStableName(IROp op);
+IROp getStableNameOpcode(UInt stableName);
+} // namespace Slang
diff --git a/source/slang/slang-ir-insts-stable-names.lua b/source/slang/slang-ir-insts-stable-names.lua
new file mode 100644
index 000000000..24a84b073
--- /dev/null
+++ b/source/slang/slang-ir-insts-stable-names.lua
@@ -0,0 +1,672 @@
+-- This file is machine generated! any entries written below will be preserved,
+-- but things like comments or anything outside the schema won't be preserved
+return {
+ ["nop"] = 0,
+ ["Unrecognized"] = 1,
+ ["Type.BasicType.Void"] = 2,
+ ["Type.BasicType.Bool"] = 3,
+ ["Type.BasicType.Int8"] = 4,
+ ["Type.BasicType.Int16"] = 5,
+ ["Type.BasicType.Int"] = 6,
+ ["Type.BasicType.Int64"] = 7,
+ ["Type.BasicType.UInt8"] = 8,
+ ["Type.BasicType.UInt16"] = 9,
+ ["Type.BasicType.UInt"] = 10,
+ ["Type.BasicType.UInt64"] = 11,
+ ["Type.BasicType.Half"] = 12,
+ ["Type.BasicType.Float"] = 13,
+ ["Type.BasicType.Double"] = 14,
+ ["Type.BasicType.Char"] = 15,
+ ["Type.BasicType.IntPtr"] = 16,
+ ["Type.BasicType.UIntPtr"] = 17,
+ ["Type.AfterBaseType"] = 18,
+ ["Type.StringTypeBase.String"] = 19,
+ ["Type.StringTypeBase.NativeString"] = 20,
+ ["Type.CapabilitySet"] = 21,
+ ["Type.DynamicType"] = 22,
+ ["Type.AnyValueType"] = 23,
+ ["Type.RawPointerTypeBase.RawPointerType"] = 24,
+ ["Type.RawPointerTypeBase.RTTIPointerType"] = 25,
+ ["Type.RawPointerTypeBase.AfterRawPointerTypeBase"] = 26,
+ ["Type.ArrayTypeBase.Array"] = 27,
+ ["Type.ArrayTypeBase.UnsizedArray"] = 28,
+ ["Type.Func"] = 29,
+ ["Type.BasicBlock"] = 30,
+ ["Type.Vec"] = 31,
+ ["Type.Mat"] = 32,
+ ["Type.Conjunction"] = 33,
+ ["Type.Attributed"] = 34,
+ ["Type.Result"] = 35,
+ ["Type.Optional"] = 36,
+ ["Type.Enum"] = 37,
+ ["Type.DifferentialPairTypeBase.DiffPair"] = 38,
+ ["Type.DifferentialPairTypeBase.DiffPairUserCode"] = 39,
+ ["Type.DifferentialPairTypeBase.DiffRefPair"] = 40,
+ ["Type.BwdDiffIntermediateCtxType"] = 41,
+ ["Type.TensorView"] = 42,
+ ["Type.TorchTensor"] = 43,
+ ["Type.ArrayListVector"] = 44,
+ ["Type.Atomic"] = 45,
+ ["Type.BindExistentialsTypeBase.BindExistentials"] = 46,
+ ["Type.BindExistentialsTypeBase.BoundInterface"] = 47,
+ ["Type.Rate.ConstExpr"] = 48,
+ ["Type.Rate.SpecConst"] = 49,
+ ["Type.Rate.GroupShared"] = 50,
+ ["Type.Rate.ActualGlobalRate"] = 51,
+ ["Type.RateQualified"] = 52,
+ ["Type.Kind.Type"] = 53,
+ ["Type.Kind.TypeParameterPack"] = 54,
+ ["Type.Kind.Rate"] = 55,
+ ["Type.Kind.Generic"] = 56,
+ ["Type.PtrTypeBase.Ptr"] = 57,
+ ["Type.PtrTypeBase.Ref"] = 58,
+ ["Type.PtrTypeBase.ConstRef"] = 59,
+ ["Type.PtrTypeBase.PseudoPtr"] = 60,
+ ["Type.PtrTypeBase.OutTypeBase.Out"] = 61,
+ ["Type.PtrTypeBase.OutTypeBase.InOut"] = 62,
+ ["Type.ComPtr"] = 63,
+ ["Type.NativePtr"] = 64,
+ ["Type.DescriptorHandle"] = 65,
+ ["Type.GLSLAtomicUint"] = 66,
+ ["Type.SamplerStateTypeBase.SamplerState"] = 67,
+ ["Type.SamplerStateTypeBase.SamplerComparisonState"] = 68,
+ ["Type.DefaultLayout"] = 69,
+ ["Type.Std140Layout"] = 70,
+ ["Type.Std430Layout"] = 71,
+ ["Type.ScalarLayout"] = 72,
+ ["Type.SubpassInputType"] = 73,
+ ["Type.TextureFootprintType"] = 74,
+ ["Type.TextureShape1DType"] = 75,
+ ["Type.TextureShape2DType"] = 76,
+ ["Type.TextureShape3DType"] = 77,
+ ["Type.TextureShapeCubeDType"] = 78,
+ ["Type.TextureShapeBufferType"] = 79,
+ ["Type.ResourceTypeBase.ResourceType.TextureTypeBase.TextureType"] = 80,
+ ["Type.ResourceTypeBase.ResourceType.TextureTypeBase.GLSLImageType"] = 81,
+ ["Type.UntypedBufferResourceType.ByteAddressBufferTypeBase.ByteAddressBuffer"] = 82,
+ ["Type.UntypedBufferResourceType.ByteAddressBufferTypeBase.RWByteAddressBuffer"] = 83,
+ ["Type.UntypedBufferResourceType.ByteAddressBufferTypeBase.RasterizerOrderedByteAddressBuffer"] = 84,
+ ["Type.UntypedBufferResourceType.RaytracingAccelerationStructure"] = 85,
+ ["Type.HLSLPatchType.InputPatch"] = 86,
+ ["Type.HLSLPatchType.OutputPatch"] = 87,
+ ["Type.GLSLInputAttachment"] = 88,
+ ["Type.BuiltinGenericType.HLSLStreamOutputType.PointStream"] = 89,
+ ["Type.BuiltinGenericType.HLSLStreamOutputType.LineStream"] = 90,
+ ["Type.BuiltinGenericType.HLSLStreamOutputType.TriangleStream"] = 91,
+ ["Type.BuiltinGenericType.MeshOutputType.Vertices"] = 92,
+ ["Type.BuiltinGenericType.MeshOutputType.Indices"] = 93,
+ ["Type.BuiltinGenericType.MeshOutputType.Primitives"] = 94,
+ ["Type.BuiltinGenericType.metal::mesh"] = 95,
+ ["Type.BuiltinGenericType.mesh_grid_properties"] = 96,
+ ["Type.BuiltinGenericType.HLSLStructuredBufferTypeBase.StructuredBuffer"] = 97,
+ ["Type.BuiltinGenericType.HLSLStructuredBufferTypeBase.RWStructuredBuffer"] = 98,
+ ["Type.BuiltinGenericType.HLSLStructuredBufferTypeBase.RasterizerOrderedStructuredBuffer"] = 99,
+ ["Type.BuiltinGenericType.HLSLStructuredBufferTypeBase.AppendStructuredBuffer"] = 100,
+ ["Type.BuiltinGenericType.HLSLStructuredBufferTypeBase.ConsumeStructuredBuffer"] = 101,
+ ["Type.BuiltinGenericType.PointerLikeType.ParameterGroupType.UniformParameterGroupType.ConstantBuffer"] = 102,
+ ["Type.BuiltinGenericType.PointerLikeType.ParameterGroupType.UniformParameterGroupType.TextureBuffer"] = 103,
+ ["Type.BuiltinGenericType.PointerLikeType.ParameterGroupType.UniformParameterGroupType.ParameterBlock"] = 104,
+ ["Type.BuiltinGenericType.PointerLikeType.ParameterGroupType.VaryingParameterGroupType.GLSLInputParameterGroup"] = 105,
+ ["Type.BuiltinGenericType.PointerLikeType.ParameterGroupType.VaryingParameterGroupType.GLSLOutputParameterGroup"] = 106,
+ ["Type.BuiltinGenericType.PointerLikeType.ParameterGroupType.GLSLShaderStorageBuffer"] = 107,
+ ["Type.RayQuery"] = 108,
+ ["Type.HitObject"] = 109,
+ ["Type.CoopVectorType"] = 110,
+ ["Type.CoopMatrixType"] = 111,
+ ["Type.TensorAddressingTensorLayoutType"] = 112,
+ ["Type.TensorAddressingTensorViewType"] = 113,
+ ["Type.MakeTensorAddressingTensorLayout"] = 114,
+ ["Type.MakeTensorAddressingTensorView"] = 115,
+ ["Type.DynamicResource"] = 116,
+ ["Type.struct"] = 117,
+ ["Type.class"] = 118,
+ ["Type.interface"] = 119,
+ ["Type.associated_type"] = 120,
+ ["Type.this_type"] = 121,
+ ["Type.rtti_type"] = 122,
+ ["Type.rtti_handle_type"] = 123,
+ ["Type.TupleTypeBase.tuple_type"] = 124,
+ ["Type.TupleTypeBase.TypePack"] = 125,
+ ["Type.TargetTuple"] = 126,
+ ["Type.ExpandTypeOrVal"] = 127,
+ ["Type.spirvLiteralType"] = 128,
+ ["Type.type_t"] = 129,
+ ["Type.WitnessTableTypeBase.witness_table_t"] = 130,
+ ["Type.WitnessTableTypeBase.witness_table_id_t"] = 131,
+ ["GlobalValueWithCode.GlobalValueWithParams.func"] = 132,
+ ["GlobalValueWithCode.GlobalValueWithParams.generic"] = 133,
+ ["GlobalValueWithCode.global_var"] = 134,
+ ["global_param"] = 135,
+ ["globalConstant"] = 136,
+ ["key"] = 137,
+ ["global_generic_param"] = 138,
+ ["witness_table"] = 139,
+ ["indexedFieldKey"] = 140,
+ ["thisTypeWitness"] = 141,
+ ["TypeEqualityWitness"] = 142,
+ ["global_hashed_string_literals"] = 143,
+ ["module"] = 144,
+ ["block"] = 145,
+ ["Constant.boolConst"] = 146,
+ ["Constant.integer_constant"] = 147,
+ ["Constant.float_constant"] = 148,
+ ["Constant.ptr_constant"] = 149,
+ ["Constant.string_constant"] = 150,
+ ["Constant.blob_constant"] = 151,
+ ["Constant.void_constant"] = 152,
+ ["CapabilitySet.capabilityConjunction"] = 153,
+ ["CapabilitySet.capabilityDisjunction"] = 154,
+ ["undefined"] = 155,
+ ["defaultConstruct"] = 156,
+ ["MakeDifferentialPairBase.MakeDiffPair"] = 157,
+ ["MakeDifferentialPairBase.MakeDiffPairUserCode"] = 158,
+ ["MakeDifferentialPairBase.MakeDiffRefPair"] = 159,
+ ["DifferentialPairGetDifferentialBase.GetDifferential"] = 160,
+ ["DifferentialPairGetDifferentialBase.GetDifferentialUserCode"] = 161,
+ ["DifferentialPairGetDifferentialBase.GetDifferentialPtr"] = 162,
+ ["DifferentialPairGetPrimalBase.GetPrimal"] = 163,
+ ["DifferentialPairGetPrimalBase.GetPrimalUserCode"] = 164,
+ ["DifferentialPairGetPrimalBase.GetPrimalRef"] = 165,
+ ["specialize"] = 166,
+ ["lookupWitness"] = 167,
+ ["GetSequentialID"] = 168,
+ ["bind_global_generic_param"] = 169,
+ ["allocObj"] = 170,
+ ["globalValueRef"] = 171,
+ ["makeUInt64"] = 172,
+ ["makeVector"] = 173,
+ ["makeMatrix"] = 174,
+ ["makeMatrixFromScalar"] = 175,
+ ["matrixReshape"] = 176,
+ ["vectorReshape"] = 177,
+ ["makeArray"] = 178,
+ ["makeArrayFromElement"] = 179,
+ ["makeCoopVector"] = 180,
+ ["makeCoopVectorFromValuePack"] = 181,
+ ["makeStruct"] = 182,
+ ["makeTuple"] = 183,
+ ["makeTargetTuple"] = 184,
+ ["makeValuePack"] = 185,
+ ["getTargetTupleElement"] = 186,
+ ["getTupleElement"] = 187,
+ ["LoadResourceDescriptorFromHeap"] = 188,
+ ["LoadSamplerDescriptorFromHeap"] = 189,
+ ["MakeCombinedTextureSamplerFromHandle"] = 190,
+ ["MakeWitnessPack"] = 191,
+ ["Expand"] = 192,
+ ["Each"] = 193,
+ ["makeResultValue"] = 194,
+ ["makeResultError"] = 195,
+ ["isResultError"] = 196,
+ ["getResultError"] = 197,
+ ["getResultValue"] = 198,
+ ["getOptionalValue"] = 199,
+ ["optionalHasValue"] = 200,
+ ["makeOptionalValue"] = 201,
+ ["makeOptionalNone"] = 202,
+ ["CombinedTextureSamplerGetTexture"] = 203,
+ ["CombinedTextureSamplerGetSampler"] = 204,
+ ["call"] = 205,
+ ["rtti_object"] = 206,
+ ["alloca"] = 207,
+ ["updateElement"] = 208,
+ ["detachDerivative"] = 209,
+ ["bitfieldExtract"] = 210,
+ ["bitfieldInsert"] = 211,
+ ["packAnyValue"] = 212,
+ ["unpackAnyValue"] = 213,
+ ["witness_table_entry"] = 214,
+ ["interface_req_entry"] = 215,
+ ["GetWorkGroupSize"] = 216,
+ ["GetCurrentStage"] = 217,
+ ["param"] = 218,
+ ["field"] = 219,
+ ["var"] = 220,
+ ["load"] = 221,
+ ["store"] = 222,
+ ["AtomicOperation.atomicLoad"] = 223,
+ ["AtomicOperation.atomicStore"] = 224,
+ ["AtomicOperation.atomicExchange"] = 225,
+ ["AtomicOperation.atomicCompareExchange"] = 226,
+ ["AtomicOperation.atomicAdd"] = 227,
+ ["AtomicOperation.atomicSub"] = 228,
+ ["AtomicOperation.atomicAnd"] = 229,
+ ["AtomicOperation.atomicOr"] = 230,
+ ["AtomicOperation.atomicXor"] = 231,
+ ["AtomicOperation.atomicMin"] = 232,
+ ["AtomicOperation.atomicMax"] = 233,
+ ["AtomicOperation.atomicInc"] = 234,
+ ["AtomicOperation.atomicDec"] = 235,
+ ["LoadReverseGradient"] = 236,
+ ["ReverseGradientDiffPairRef"] = 237,
+ ["PrimalParamRef"] = 238,
+ ["DiffParamRef"] = 239,
+ ["IsDifferentialNull"] = 240,
+ ["get_field"] = 241,
+ ["get_field_addr"] = 242,
+ ["getElement"] = 243,
+ ["getElementPtr"] = 244,
+ ["getOffsetPtr"] = 245,
+ ["getAddr"] = 246,
+ ["castDynamicResource"] = 247,
+ ["getNativeStr"] = 248,
+ ["makeString"] = 249,
+ ["getNativePtr"] = 250,
+ ["getManagedPtrWriteRef"] = 251,
+ ["ManagedPtrAttach"] = 252,
+ ["ManagedPtrDetach"] = 253,
+ ["imageSubscript"] = 254,
+ ["imageLoad"] = 255,
+ ["imageStore"] = 256,
+ ["byteAddressBufferLoad"] = 257,
+ ["byteAddressBufferStore"] = 258,
+ ["structuredBufferLoad"] = 259,
+ ["structuredBufferLoadStatus"] = 260,
+ ["rwstructuredBufferLoad"] = 261,
+ ["rwstructuredBufferLoadStatus"] = 262,
+ ["rwstructuredBufferStore"] = 263,
+ ["rwstructuredBufferGetElementPtr"] = 264,
+ ["StructuredBufferAppend"] = 265,
+ ["StructuredBufferConsume"] = 266,
+ ["StructuredBufferGetDimensions"] = 267,
+ ["nonUniformResourceIndex"] = 268,
+ ["getNaturalStride"] = 269,
+ ["meshOutputRef"] = 270,
+ ["meshOutputSet"] = 271,
+ ["metalSetVertex"] = 272,
+ ["metalSetPrimitive"] = 273,
+ ["metalSetIndices"] = 274,
+ ["MetalCastToDepthTexture"] = 275,
+ ["MakeVectorFromScalar"] = 276,
+ ["swizzle"] = 277,
+ ["swizzleSet"] = 278,
+ ["swizzledStore"] = 279,
+ ["TerminatorInst.return_val"] = 280,
+ ["TerminatorInst.yield"] = 281,
+ ["TerminatorInst.UnconditionalBranch.unconditionalBranch"] = 282,
+ ["TerminatorInst.UnconditionalBranch.loop"] = 283,
+ ["TerminatorInst.ConditionalBranch.conditionalBranch"] = 284,
+ ["TerminatorInst.ConditionalBranch.ifElse"] = 285,
+ ["TerminatorInst.throw"] = 286,
+ ["TerminatorInst.tryCall"] = 287,
+ ["TerminatorInst.switch"] = 288,
+ ["TerminatorInst.targetSwitch"] = 289,
+ ["TerminatorInst.GenericAsm"] = 290,
+ ["TerminatorInst.Unreachable.missingReturn"] = 291,
+ ["TerminatorInst.Unreachable.unreachable"] = 292,
+ ["TerminatorInst.defer"] = 293,
+ ["discard"] = 294,
+ ["RequirePrelude"] = 295,
+ ["RequireTargetExtension"] = 296,
+ ["RequireComputeDerivative"] = 297,
+ ["StaticAssert"] = 298,
+ ["Printf"] = 299,
+ ["RequireMaximallyReconverges"] = 300,
+ ["RequireQuadDerivatives"] = 301,
+ ["add"] = 302,
+ ["sub"] = 303,
+ ["mul"] = 304,
+ ["div"] = 305,
+ ["irem"] = 306,
+ ["frem"] = 307,
+ ["shl"] = 308,
+ ["shr"] = 309,
+ ["cmpEQ"] = 310,
+ ["cmpNE"] = 311,
+ ["cmpGT"] = 312,
+ ["cmpLT"] = 313,
+ ["cmpGE"] = 314,
+ ["cmpLE"] = 315,
+ ["and"] = 316,
+ ["xor"] = 317,
+ ["or"] = 318,
+ ["logicalAnd"] = 319,
+ ["logicalOr"] = 320,
+ ["neg"] = 321,
+ ["not"] = 322,
+ ["bitnot"] = 323,
+ ["select"] = 324,
+ ["checkpointObj"] = 325,
+ ["loopExitValue"] = 326,
+ ["getStringHash"] = 327,
+ ["waveGetActiveMask"] = 328,
+ ["waveMaskBallot"] = 329,
+ ["waveMaskMatch"] = 330,
+ ["sample"] = 331,
+ ["sampleGrad"] = 332,
+ ["GroupMemoryBarrierWithGroupSync"] = 333,
+ ["ControlBarrier"] = 334,
+ ["gpuForeach"] = 335,
+ ["getOptiXRayPayloadPtr"] = 336,
+ ["getOptiXHitAttribute"] = 337,
+ ["getOptiXSbtDataPointer"] = 338,
+ ["GetVulkanRayTracingPayloadLocation"] = 339,
+ ["GetLegalizedSPIRVGlobalParamAddr"] = 340,
+ ["GetPerVertexInputArray"] = 341,
+ ["ResolveVaryingInputRef"] = 342,
+ ["ForceVarIntoStructTemporarilyBase.ForceVarIntoStructTemporarily"] = 343,
+ ["ForceVarIntoStructTemporarilyBase.ForceVarIntoRayPayloadStructTemporarily"] = 344,
+ ["MetalAtomicCast"] = 345,
+ ["IsTextureAccess"] = 346,
+ ["IsTextureScalarAccess"] = 347,
+ ["IsTextureArrayAccess"] = 348,
+ ["ExtractTextureFromTextureAccess"] = 349,
+ ["ExtractCoordFromTextureAccess"] = 350,
+ ["ExtractArrayCoordFromTextureAccess"] = 351,
+ ["makeArrayList"] = 352,
+ ["makeTensorView"] = 353,
+ ["allocTorchTensor"] = 354,
+ ["TorchGetCudaStream"] = 355,
+ ["TorchTensorGetView"] = 356,
+ ["CoopMatMapElementIFunc"] = 357,
+ ["allocateOpaqueHandle"] = 358,
+ ["BindingQuery.getRegisterIndex"] = 359,
+ ["BindingQuery.getRegisterSpace"] = 360,
+ ["Decoration.highLevelDecl"] = 361,
+ ["Decoration.layout"] = 362,
+ ["Decoration.branch"] = 363,
+ ["Decoration.flatten"] = 364,
+ ["Decoration.loopControl"] = 365,
+ ["Decoration.loopMaxIters"] = 366,
+ ["Decoration.loopExitPrimalValue"] = 367,
+ ["Decoration.intrinsicOp"] = 368,
+ ["Decoration.TargetSpecificDecoration.TargetSpecificDefinitionDecoration.target"] = 369,
+ ["Decoration.TargetSpecificDecoration.TargetSpecificDefinitionDecoration.targetIntrinsic"] = 370,
+ ["Decoration.TargetSpecificDecoration.requirePrelude"] = 371,
+ ["Decoration.glslOuterArray"] = 372,
+ ["Decoration.TargetSystemValue"] = 373,
+ ["Decoration.interpolationMode"] = 374,
+ ["Decoration.nameHint"] = 375,
+ ["Decoration.PhysicalType"] = 376,
+ ["Decoration.AlignedAddressDecoration"] = 377,
+ ["Decoration.BinaryInterfaceType"] = 378,
+ ["Decoration.transitory"] = 379,
+ ["Decoration.ResultWitness"] = 380,
+ ["Decoration.vulkanRayPayload"] = 381,
+ ["Decoration.vulkanRayPayloadIn"] = 382,
+ ["Decoration.vulkanHitAttributes"] = 383,
+ ["Decoration.vulkanHitObjectAttributes"] = 384,
+ ["Decoration.GlobalVariableShadowingGlobalParameterDecoration"] = 385,
+ ["Decoration.requireSPIRVVersion"] = 386,
+ ["Decoration.requireGLSLVersion"] = 387,
+ ["Decoration.requireGLSLExtension"] = 388,
+ ["Decoration.requireWGSLExtension"] = 389,
+ ["Decoration.requireCUDASMVersion"] = 390,
+ ["Decoration.requireCapabilityAtom"] = 391,
+ ["Decoration.HasExplicitHLSLBinding"] = 392,
+ ["Decoration.DefaultValue"] = 393,
+ ["Decoration.readNone"] = 394,
+ ["Decoration.vulkanCallablePayload"] = 395,
+ ["Decoration.vulkanCallablePayloadIn"] = 396,
+ ["Decoration.earlyDepthStencil"] = 397,
+ ["Decoration.precise"] = 398,
+ ["Decoration.public"] = 399,
+ ["Decoration.hlslExport"] = 400,
+ ["Decoration.downstreamModuleExport"] = 401,
+ ["Decoration.downstreamModuleImport"] = 402,
+ ["Decoration.patchConstantFunc"] = 403,
+ ["Decoration.maxTessFactor"] = 404,
+ ["Decoration.outputControlPoints"] = 405,
+ ["Decoration.outputTopology"] = 406,
+ ["Decoration.partioning"] = 407,
+ ["Decoration.domain"] = 408,
+ ["Decoration.maxVertexCount"] = 409,
+ ["Decoration.instance"] = 410,
+ ["Decoration.numThreads"] = 411,
+ ["Decoration.fpDenormalPreserve"] = 412,
+ ["Decoration.fpDenormalFlushToZero"] = 413,
+ ["Decoration.waveSize"] = 414,
+ ["Decoration.availableInDownstreamIR"] = 415,
+ ["Decoration.GeometryInputPrimitiveTypeDecoration.pointPrimitiveType"] = 416,
+ ["Decoration.GeometryInputPrimitiveTypeDecoration.linePrimitiveType"] = 417,
+ ["Decoration.GeometryInputPrimitiveTypeDecoration.trianglePrimitiveType"] = 418,
+ ["Decoration.GeometryInputPrimitiveTypeDecoration.lineAdjPrimitiveType"] = 419,
+ ["Decoration.GeometryInputPrimitiveTypeDecoration.triangleAdjPrimitiveType"] = 420,
+ ["Decoration.streamOutputTypeDecoration"] = 421,
+ ["Decoration.entryPoint"] = 422,
+ ["Decoration.CudaKernel"] = 423,
+ ["Decoration.CudaHost"] = 424,
+ ["Decoration.TorchEntryPoint"] = 425,
+ ["Decoration.AutoPyBindCUDA"] = 426,
+ ["Decoration.CudaKernelFwdDiffRef"] = 427,
+ ["Decoration.CudaKernelBwdDiffRef"] = 428,
+ ["Decoration.PyBindExportFuncInfo"] = 429,
+ ["Decoration.PyExportDecoration"] = 430,
+ ["Decoration.entryPointParam"] = 431,
+ ["Decoration.dependsOn"] = 432,
+ ["Decoration.keepAlive"] = 433,
+ ["Decoration.noSideEffect"] = 434,
+ ["Decoration.bindExistentialSlots"] = 435,
+ ["Decoration.format"] = 436,
+ ["Decoration.unsafeForceInlineEarly"] = 437,
+ ["Decoration.ForceInline"] = 438,
+ ["Decoration.ForceUnroll"] = 439,
+ ["Decoration.SizeAndAlignment"] = 440,
+ ["Decoration.Offset"] = 441,
+ ["Decoration.LinkageDecoration.import"] = 442,
+ ["Decoration.LinkageDecoration.export"] = 443,
+ ["Decoration.TargetBuiltinVar"] = 444,
+ ["Decoration.UserExtern"] = 445,
+ ["Decoration.externCpp"] = 446,
+ ["Decoration.externC"] = 447,
+ ["Decoration.dllImport"] = 448,
+ ["Decoration.dllExport"] = 449,
+ ["Decoration.cudaDeviceExport"] = 450,
+ ["Decoration.COMInterface"] = 451,
+ ["Decoration.KnownBuiltinDecoration"] = 452,
+ ["Decoration.RTTI_typeSize"] = 453,
+ ["Decoration.AnyValueSize"] = 454,
+ ["Decoration.SpecializeDecoration"] = 455,
+ ["Decoration.SequentialIDDecoration"] = 456,
+ ["Decoration.DynamicDispatchWitnessDecoration"] = 457,
+ ["Decoration.StaticRequirementDecoration"] = 458,
+ ["Decoration.DispatchFuncDecoration"] = 459,
+ ["Decoration.TypeConstraintDecoration"] = 460,
+ ["Decoration.BuiltinDecoration"] = 461,
+ ["Decoration.requiresNVAPI"] = 462,
+ ["Decoration.nvapiMagic"] = 463,
+ ["Decoration.nvapiSlot"] = 464,
+ ["Decoration.noInline"] = 465,
+ ["Decoration.noRefInline"] = 466,
+ ["Decoration.DerivativeGroupQuad"] = 467,
+ ["Decoration.DerivativeGroupLinear"] = 468,
+ ["Decoration.MaximallyReconverges"] = 469,
+ ["Decoration.QuadDerivatives"] = 470,
+ ["Decoration.RequireFullQuads"] = 471,
+ ["Decoration.TempCallArgVar"] = 472,
+ ["Decoration.nonCopyable"] = 473,
+ ["Decoration.DynamicUniform"] = 474,
+ ["Decoration.alwaysFold"] = 475,
+ ["Decoration.output"] = 476,
+ ["Decoration.input"] = 477,
+ ["Decoration.glslLocation"] = 478,
+ ["Decoration.glslOffset"] = 479,
+ ["Decoration.vkStructOffset"] = 480,
+ ["Decoration.raypayload"] = 481,
+ ["Decoration.MeshOutputDecoration.vertices"] = 482,
+ ["Decoration.MeshOutputDecoration.indices"] = 483,
+ ["Decoration.MeshOutputDecoration.primitives"] = 484,
+ ["Decoration.HLSLMeshPayloadDecoration"] = 485,
+ ["Decoration.perprimitive"] = 486,
+ ["Decoration.PositionOutput"] = 487,
+ ["Decoration.PositionInput"] = 488,
+ ["Decoration.PerVertex"] = 489,
+ ["Decoration.StageAccessDecoration.stageReadAccess"] = 490,
+ ["Decoration.StageAccessDecoration.stageWriteAccess"] = 491,
+ ["Decoration.semantic"] = 492,
+ ["Decoration.constructor"] = 493,
+ ["Decoration.method"] = 494,
+ ["Decoration.packoffset"] = 495,
+ ["Decoration.SpecializationConstantDecoration"] = 496,
+ ["Decoration.UserTypeName"] = 497,
+ ["Decoration.CounterBuffer"] = 498,
+ ["Decoration.RequireSPIRVDescriptorIndexingExtensionDecoration"] = 499,
+ ["Decoration.spirvOpDecoration"] = 500,
+ ["Decoration.forwardDifferentiable"] = 501,
+ ["Decoration.AutoDiffOriginalValueDecoration"] = 502,
+ ["Decoration.AutoDiffBuiltinDecoration"] = 503,
+ ["Decoration.fwdDerivative"] = 504,
+ ["Decoration.backwardDifferentiable"] = 505,
+ ["Decoration.primalSubstFunc"] = 506,
+ ["Decoration.backwardDiffPrimalReference"] = 507,
+ ["Decoration.backwardDiffPropagateReference"] = 508,
+ ["Decoration.backwardDiffIntermediateTypeReference"] = 509,
+ ["Decoration.backwardDiffReference"] = 510,
+ ["Decoration.userDefinedBackwardDiffReference"] = 511,
+ ["Decoration.BackwardDerivativePrimalContextDecoration"] = 512,
+ ["Decoration.BackwardDerivativePrimalReturnDecoration"] = 513,
+ ["Decoration.PrimalContextDecoration"] = 514,
+ ["Decoration.loopCounterDecoration"] = 515,
+ ["Decoration.loopCounterUpdateDecoration"] = 516,
+ ["Decoration.AutodiffInstDecoration.primalInstDecoration"] = 517,
+ ["Decoration.AutodiffInstDecoration.diffInstDecoration"] = 518,
+ ["Decoration.AutodiffInstDecoration.mixedDiffInstDecoration"] = 519,
+ ["Decoration.AutodiffInstDecoration.RecomputeBlockDecoration"] = 520,
+ ["Decoration.primalValueKey"] = 521,
+ ["Decoration.primalElementType"] = 522,
+ ["Decoration.IntermediateContextFieldDifferentialTypeDecoration"] = 523,
+ ["Decoration.derivativeMemberDecoration"] = 524,
+ ["Decoration.treatAsDifferentiableDecoration"] = 525,
+ ["Decoration.treatCallAsDifferentiableDecoration"] = 526,
+ ["Decoration.differentiableCallDecoration"] = 527,
+ ["Decoration.optimizableTypeDecoration"] = 528,
+ ["Decoration.ignoreSideEffectsDecoration"] = 529,
+ ["Decoration.CheckpointHintDecoration.PreferCheckpointDecoration"] = 530,
+ ["Decoration.CheckpointHintDecoration.PreferRecomputeDecoration"] = 531,
+ ["Decoration.CheckpointHintDecoration.CheckpointIntermediateDecoration"] = 532,
+ ["Decoration.NonDynamicUniformReturnDecoration"] = 533,
+ ["Decoration.COMWitnessDecoration"] = 534,
+ ["Decoration.DifferentiableTypeDictionaryDecoration"] = 535,
+ ["Decoration.FloatingPointModeOverride"] = 536,
+ ["Decoration.spvBufferBlock"] = 537,
+ ["Decoration.DebugLocation"] = 538,
+ ["Decoration.DebugFunction"] = 539,
+ ["Decoration.spvBlock"] = 540,
+ ["Decoration.NonUniformResource"] = 541,
+ ["Decoration.MemoryQualifierSetDecoration"] = 542,
+ ["Decoration.BitFieldAccessorDecoration"] = 543,
+ ["makeExistential"] = 544,
+ ["makeExistentialWithRTTI"] = 545,
+ ["createExistentialObject"] = 546,
+ ["wrapExistential"] = 547,
+ ["getValueFromBoundInterface"] = 548,
+ ["extractExistentialValue"] = 549,
+ ["extractExistentialType"] = 550,
+ ["extractExistentialWitnessTable"] = 551,
+ ["isNullExistential"] = 552,
+ ["extractTaggedUnionTag"] = 553,
+ ["extractTaggedUnionPayload"] = 554,
+ ["BuiltinCast"] = 555,
+ ["bitCast"] = 556,
+ ["reinterpret"] = 557,
+ ["unmodified"] = 558,
+ ["outImplicitCast"] = 559,
+ ["inOutImplicitCast"] = 560,
+ ["intCast"] = 561,
+ ["floatCast"] = 562,
+ ["castIntToFloat"] = 563,
+ ["castFloatToInt"] = 564,
+ ["CastPtrToBool"] = 565,
+ ["CastPtrToInt"] = 566,
+ ["CastIntToPtr"] = 567,
+ ["castToVoid"] = 568,
+ ["PtrCast"] = 569,
+ ["CastEnumToInt"] = 570,
+ ["CastIntToEnum"] = 571,
+ ["EnumCast"] = 572,
+ ["CastUInt2ToDescriptorHandle"] = 573,
+ ["CastDescriptorHandleToUInt2"] = 574,
+ ["CastDescriptorHandleToResource"] = 575,
+ ["TreatAsDynamicUniform"] = 576,
+ ["sizeOf"] = 577,
+ ["alignOf"] = 578,
+ ["countOf"] = 579,
+ ["GetArrayLength"] = 580,
+ ["IsType"] = 581,
+ ["TypeEquals"] = 582,
+ ["IsInt"] = 583,
+ ["IsBool"] = 584,
+ ["IsFloat"] = 585,
+ ["IsHalf"] = 586,
+ ["IsUnsignedInt"] = 587,
+ ["IsSignedInt"] = 588,
+ ["IsVector"] = 589,
+ ["GetDynamicResourceHeap"] = 590,
+ ["ForwardDifferentiate"] = 591,
+ ["BackwardDifferentiatePrimal"] = 592,
+ ["BackwardDifferentiatePropagate"] = 593,
+ ["BackwardDifferentiate"] = 594,
+ ["PrimalSubstitute"] = 595,
+ ["DispatchKernel"] = 596,
+ ["CudaKernelLaunch"] = 597,
+ ["getEquivalentStructuredBuffer"] = 598,
+ ["getStructuredBufferPtr"] = 599,
+ ["getUntypedBufferPtr"] = 600,
+ ["Layout.varLayout"] = 601,
+ ["Layout.TypeLayout.typeLayout"] = 602,
+ ["Layout.TypeLayout.parameterGroupTypeLayout"] = 603,
+ ["Layout.TypeLayout.arrayTypeLayout"] = 604,
+ ["Layout.TypeLayout.streamOutputTypeLayout"] = 605,
+ ["Layout.TypeLayout.matrixTypeLayout"] = 606,
+ ["Layout.TypeLayout.existentialTypeLayout"] = 607,
+ ["Layout.TypeLayout.structTypeLayout"] = 608,
+ ["Layout.TypeLayout.tupleTypeLayout"] = 609,
+ ["Layout.TypeLayout.structuredBufferTypeLayout"] = 610,
+ ["Layout.TypeLayout.ptrTypeLayout"] = 611,
+ ["Layout.EntryPointLayout"] = 612,
+ ["Attr.pendingLayout"] = 613,
+ ["Attr.stage"] = 614,
+ ["Attr.structFieldLayout"] = 615,
+ ["Attr.tupleFieldLayout"] = 616,
+ ["Attr.caseLayout"] = 617,
+ ["Attr.unorm"] = 618,
+ ["Attr.snorm"] = 619,
+ ["Attr.no_diff"] = 620,
+ ["Attr.nonuniform"] = 621,
+ ["Attr.Aligned"] = 622,
+ ["Attr.SemanticAttr.userSemantic"] = 623,
+ ["Attr.SemanticAttr.systemValueSemantic"] = 624,
+ ["Attr.LayoutResourceInfoAttr.size"] = 625,
+ ["Attr.LayoutResourceInfoAttr.offset"] = 626,
+ ["Attr.FuncThrowType"] = 627,
+ ["LiveRangeMarker.liveRangeStart"] = 628,
+ ["LiveRangeMarker.liveRangeEnd"] = 629,
+ ["SpecializationDictionaryItem"] = 630,
+ ["GenericSpecializationDictionary"] = 631,
+ ["ExistentialFuncSpecializationDictionary"] = 632,
+ ["ExistentialTypeSpecializationDictionary"] = 633,
+ ["DifferentiableTypeDictionaryItem"] = 634,
+ ["DifferentiableTypeAnnotation"] = 635,
+ ["BeginFragmentShaderInterlock"] = 636,
+ ["EndFragmentShaderInterlock"] = 637,
+ ["DebugSource"] = 638,
+ ["DebugLine"] = 639,
+ ["DebugVar"] = 640,
+ ["DebugValue"] = 641,
+ ["DebugInlinedAt"] = 642,
+ ["DebugFunction"] = 643,
+ ["DebugInlinedVariable"] = 644,
+ ["DebugScope"] = 645,
+ ["DebugNoScope"] = 646,
+ ["DebugBuildIdentifier"] = 647,
+ ["EmbeddedDownstreamIR"] = 648,
+ ["SPIRVAsm"] = 649,
+ ["SPIRVAsmInst"] = 650,
+ ["SPIRVAsmOperand.SPIRVAsmOperandLiteral"] = 651,
+ ["SPIRVAsmOperand.SPIRVAsmOperandInst"] = 652,
+ ["SPIRVAsmOperand.SPIRVAsmOperandConvertTexel"] = 653,
+ ["SPIRVAsmOperand.SPIRVAsmOperandRayPayloadFromLocation"] = 654,
+ ["SPIRVAsmOperand.SPIRVAsmOperandRayAttributeFromLocation"] = 655,
+ ["SPIRVAsmOperand.SPIRVAsmOperandRayCallableFromLocation"] = 656,
+ ["SPIRVAsmOperand.SPIRVAsmOperandEnum"] = 657,
+ ["SPIRVAsmOperand.SPIRVAsmOperandBuiltinVar"] = 658,
+ ["SPIRVAsmOperand.SPIRVAsmOperandGLSL450Set"] = 659,
+ ["SPIRVAsmOperand.SPIRVAsmOperandDebugPrintfSet"] = 660,
+ ["SPIRVAsmOperand.SPIRVAsmOperandId"] = 661,
+ ["SPIRVAsmOperand.SPIRVAsmOperandResult"] = 662,
+ ["SPIRVAsmOperand.__truncate"] = 663,
+ ["SPIRVAsmOperand.__entryPoint"] = 664,
+ ["SPIRVAsmOperand.__sampledType"] = 665,
+ ["SPIRVAsmOperand.__imageType"] = 666,
+ ["SPIRVAsmOperand.__sampledImageType"] = 667,
+}
diff --git a/source/slang/slang-ir-insts.lua b/source/slang/slang-ir-insts.lua
index 44dbc1ade..a977e83f0 100644
--- a/source/slang/slang-ir-insts.lua
+++ b/source/slang/slang-ir-insts.lua
@@ -1,6 +1,12 @@
--
-- This file contains the canonical definitions for the instions to the Slang IR.
--- Add new instructions here
+--
+-- Add new instructions here.
+--
+-- !!
+-- !! Please make sure to update the supported module versions in
+-- !! Slang::IRModule accordingly when modifying this file.
+-- !!
--
-- The instructions struct name, i.e. something like "IRVoidType" can be specified with struct_name, otherwise it will be a PascalCase version of the instruction key
--
@@ -13,6 +19,10 @@
local insts = {
{ nop = {} },
+ -- This opcode is used as a placeholder if we were ever to deserialize a
+ -- module which contains an instruction we don't have defined in this file,
+ -- it should never appear except immediately after deserialization.
+ { Unrecognized = {} },
{
Type = {
{
@@ -2171,10 +2181,11 @@ local insts = {
-- A function to calculate some useful properties and put it in the table,
--
--- Returns the tree as well as the flattened tree in preorder
-- Annotates instructions with whether they are a leaf or not
-- Calculates flags from the parent flags
local function process(insts)
+ local stable_names_file = "source/slang/slang-ir-insts-stable-names.lua"
+
local function to_pascal_case(str)
local result = str:gsub("_(.)", function(c)
return c:upper()
@@ -2190,6 +2201,29 @@ local function process(insts)
return true
end
+ -- Load stable names if file is provided
+ local name_to_stable_name = loadfile(stable_names_file)()
+ local stable_name_to_inst = {}
+ local max_stable_name = 0
+
+ -- Build full path for stable name lookup
+ local function build_path(inst, name)
+ local path = { name }
+ local current = inst.parent_inst
+ while current and current.parent_inst do
+ -- Find the name of current in its parent
+ for _, entry in ipairs(current.parent_inst) do
+ local k, v = next(entry)
+ if v == current then
+ table.insert(path, 1, k)
+ break
+ end
+ end
+ current = current.parent_inst
+ end
+ return table.concat(path, ".")
+ end
+
-- Recursively process instructions
local function process_inst(tbl, inherited_flags)
inherited_flags = inherited_flags or {}
@@ -2252,8 +2286,52 @@ local function process(insts)
-- Process the entire tree
process_inst(insts)
+ -- Now add stable names after parent_inst is set
+ local function add_stable_names(tbl)
+ for _, i in ipairs(tbl) do
+ local key, value = next(i)
+
+ -- Build the full path for this instruction
+ local full_path = build_path(value, key)
+
+ -- Look up stable name
+ local stable_id = name_to_stable_name[full_path]
+ if stable_id then
+ value.stable_name = stable_id
+ stable_name_to_inst[stable_id] = value
+ if stable_id > max_stable_name then
+ max_stable_name = stable_id
+ end
+ end
+
+ -- Recursively process children
+ add_stable_names(value)
+ end
+ end
+
+ -- Add stable names to all instructions
+ add_stable_names(insts)
+
+ -- Helper function to traverse the instruction tree
+ local function traverse(callback)
+ local function walk_insts(tbl)
+ for _, i in ipairs(tbl) do
+ local _, value = next(i)
+ callback(value)
+ -- Recursively process nested instructions
+ walk_insts(value)
+ end
+ end
+
+ -- Start walking from the top-level insts
+ walk_insts(insts)
+ end
+
return {
insts = insts,
+ stable_name_to_inst = stable_name_to_inst,
+ max_stable_name = max_stable_name,
+ traverse = traverse,
}
end
diff --git a/source/slang/slang-ir.h b/source/slang/slang-ir.h
index c29a52ff7..712763e86 100644
--- a/source/slang/slang-ir.h
+++ b/source/slang/slang-ir.h
@@ -2357,8 +2357,10 @@ struct IRAnalysis
IRDominatorTree* getDominatorTree();
};
+FIDDLE()
struct IRModule : RefObject
{
+ FIDDLE(...)
public:
enum
{
@@ -2441,9 +2443,28 @@ public:
ContainerPool& getContainerPool() { return m_containerPool; }
+ //
+ // The range of module versions this compiler supports
+ //
+ // This will need to be updated if for example an instruction is removed,
+ // the max supported version should be incremented and the min supported
+ // version set to above the last version an instance of that instruction
+ // could be found
+ //
+ // Additionally this should be updated when new instructions are added,
+ // however only k_maxSupportedModuleVersion needs to be incremented in that
+ // case
+ //
+ // It represents the version of module regarding semantics and doesn't have
+ // anything to do with serialization format
+ //
+ const static UInt k_minSupportedModuleVersion = 0;
+ const static UInt k_maxSupportedModuleVersion = 0;
+
private:
friend struct IRSerialReadContext;
friend struct IRSerialWriteContext;
+ friend struct Fossilized_IRModule;
IRModule() = delete;
@@ -2463,10 +2484,13 @@ private:
/// instructions from an arbitrary IR instruction we expect to find the
/// `IRModuleInst` for the module the instruction belongs to, if any.
///
- IRModuleInst* m_moduleInst = nullptr;
+ FIDDLE() IRModuleInst* m_moduleInst = nullptr;
// The name of the module.
- Name* m_name = nullptr;
+ FIDDLE() Name* m_name = nullptr;
+
+ // The version of the module as it was loaded
+ FIDDLE() UInt m_version = k_maxSupportedModuleVersion;
/// The memory arena from which all IR instructions (and any associated state) in this module
/// are allocated.
diff --git a/source/slang/slang-options.cpp b/source/slang/slang-options.cpp
index af799839c..fde13463e 100644
--- a/source/slang/slang-options.cpp
+++ b/source/slang/slang-options.cpp
@@ -868,7 +868,15 @@ void initCommandOptions(CommandOptions& options)
"-verify-debug-serial-ir",
nullptr,
"Verify IR in the front-end."},
- {OptionKind::DumpModule, "-dump-module", nullptr, "Disassemble and print the module IR."}};
+ {OptionKind::DumpModule, "-dump-module", nullptr, "Disassemble and print the module IR."},
+ {OptionKind::GetModuleInfo,
+ "-get-module-info",
+ nullptr,
+ "Print the name and version of a serialized IR Module"},
+ {OptionKind::GetSupportedModuleVersions,
+ "-get-supported-module-versions",
+ nullptr,
+ "Print the minimum and maximum module versions this compiler supports"}};
_addOptions(makeConstArrayView(debuggingOpts), options);
/* !!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! Experimental !!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! */
@@ -3148,6 +3156,79 @@ SlangResult OptionsParser::_parse(int argc, char const* const* argv)
break;
}
+ case OptionKind::GetModuleInfo:
+ {
+ CommandLineArg fileName;
+ SLANG_RETURN_ON_FAIL(m_reader.expectArg(fileName));
+ auto desc = slang::SessionDesc();
+ ComPtr<slang::ISession> session;
+ m_session->createSession(desc, session.writeRef());
+ ComPtr<slang::IBlob> diagnostics;
+
+ FileStream file;
+ if (SLANG_FAILED(file.init(
+ fileName.value,
+ FileMode::Open,
+ FileAccess::Read,
+ FileShare::None)))
+ {
+ m_sink->diagnose(arg.loc, Diagnostics::cannotOpenFile, fileName.value);
+ return SLANG_FAIL;
+ }
+
+ List<uint8_t> buffer;
+ file.seek(SeekOrigin::End, 0);
+ const Int64 size = file.getPosition();
+ buffer.setCount(size + 1);
+ file.seek(SeekOrigin::Start, 0);
+ SLANG_RETURN_ON_FAIL(file.readExactly(buffer.getBuffer(), (size_t)size));
+ buffer[size] = 0;
+ file.close();
+
+ ComPtr<slang::IModule> module;
+ // Load buffer as an IR blob
+ ComPtr<slang::IBlob> blob;
+ blob = RawBlob::create(buffer.getBuffer(), size);
+
+ const char* moduleName = nullptr;
+ const char* moduleCompilerVersion = nullptr;
+ SlangInt moduleVersion = ~0;
+ SLANG_RETURN_ON_FAIL(session->loadModuleInfoFromIRBlob(
+ blob,
+ moduleVersion,
+ moduleCompilerVersion,
+ moduleName));
+
+ char infoBuffer[512];
+ snprintf(
+ infoBuffer,
+ sizeof(infoBuffer),
+ "Module Name: %s\n"
+ "Module Version: %lld\n"
+ "Compiler Version: %s\n",
+ moduleName ? moduleName : "null",
+ static_cast<long long>(moduleVersion),
+ moduleCompilerVersion ? moduleCompilerVersion : "null");
+
+ m_sink->diagnoseRaw(Severity::Note, infoBuffer);
+
+ break;
+ }
+ case OptionKind::GetSupportedModuleVersions:
+ {
+ char infoBuffer[512];
+ snprintf(
+ infoBuffer,
+ sizeof(infoBuffer),
+ "Minimum supported version: %lu\n"
+ "Maximum supported version: %lu\n",
+ (unsigned long)IRModule::k_minSupportedModuleVersion,
+ (unsigned long)IRModule::k_maxSupportedModuleVersion);
+
+ m_sink->diagnoseRaw(Severity::Note, infoBuffer);
+
+ break;
+ }
case OptionKind::EmitSeparateDebug:
{
// This will emit a separate debug file, containing all debug info in
diff --git a/source/slang/slang-serialize-container.cpp b/source/slang/slang-serialize-container.cpp
index d4d1fc3ad..7325f6e42 100644
--- a/source/slang/slang-serialize-container.cpp
+++ b/source/slang/slang-serialize-container.cpp
@@ -697,7 +697,8 @@ static void calcModuleInstructionList(IRModule* module, List<IRInst*>& instsOut)
return SLANG_FAIL;
}
- readSerializedModuleIR(irChunk, session, sourceLocReader, irReadModule);
+ SLANG_RETURN_ON_FAIL(
+ readSerializedModuleIR(irChunk, session, sourceLocReader, irReadModule));
}
}
diff --git a/source/slang/slang-serialize-ir.cpp b/source/slang/slang-serialize-ir.cpp
index 52a0a6537..55ccca5fc 100644
--- a/source/slang/slang-serialize-ir.cpp
+++ b/source/slang/slang-serialize-ir.cpp
@@ -2,12 +2,19 @@
#include "slang-serialize-ir.h"
#include "core/slang-blob-builder.h"
+#include "slang-ir-insts-stable-names.h"
#include "slang-ir-insts.h"
#include "slang-ir-validate.h"
#include "slang-serialize-fossil.h"
#include "slang-serialize-source-loc.h"
#include "slang-serialize.h"
+#include "slang-tag-version.h"
+#include "slang.h"
+//
+#include "slang-serialize-ir.cpp.fiddle"
+
+FIDDLE()
namespace Slang
{
@@ -16,9 +23,20 @@ namespace Slang
// we want to serialize some sidecar information to help with on-demand loading
// or backwards compat
//
+FIDDLE()
struct IRModuleInfo
{
- RefPtr<IRModule> module;
+ FIDDLE(...)
+ // Include the specific compiler version in serialized output, in case we
+ // ever need to do any version specific workarounds.
+ FIDDLE() String fullVersion = SLANG_TAG_VERSION;
+ // Include this here so that if we need to change the way we serialize
+ // things and maintain backwards compat we can increment this value, for
+ // example if we introduce more instructions with weird payloads like
+ // IRModuleInst or IRConstants.
+ const static UInt kSupportedSerializationVersion = 0;
+ FIDDLE() UInt serializationVersion = kSupportedSerializationVersion;
+ FIDDLE() RefPtr<IRModule> module;
};
//
@@ -68,19 +86,87 @@ struct IRSerialReadContext : IRSerialContext, RefObject
RefPtr<IRModule> _module;
};
-// IROps are serialized as integers
-SLANG_DECLARE_FOSSILIZED_AS(IROp, FossilUInt);
-void serialize(Serializer const& serializer, IROp& value)
+SLANG_DECLARE_FOSSILIZED_AS(Name, String);
+
+/// Fossilized representation of a `IRModule`
+struct Fossilized_IRModule;
+
+SLANG_DECLARE_FOSSILIZED_TYPE(IRModule, Fossilized_IRModule);
+
+struct Fossilized_IRModule : public FossilizedRecordVal
+{
+ Fossilized<decltype(IRModule::m_moduleInst)> m_moduleInst;
+ Fossilized<String> m_name;
+ Fossilized<decltype(IRModule::m_version)> m_version;
+};
+
+//
+// This splice handles any aggregate types, a similar splice is well documented
+// in slang-serialize-ast.cpp
+//
+#if 0 // FIDDLE TEMPLATE:
+% irStructTypes = {
+% Slang.IRModuleInfo,
+% }
+%
+% for _,T in ipairs(irStructTypes) do
+
+/// Fossilized representation of a `$T`
+struct Fossilized_$T;
+
+SLANG_DECLARE_FOSSILIZED_TYPE($T, Fossilized_$T);
+
+/// Serialize a `$T`
+void serialize(IRSerializer const& serializer, $T& value);
+%end
+%for _,T in ipairs(irStructTypes) do
+/// Fossilized representation of a value of type `$T`
+struct Fossilized_$T
+% if T.directSuperClass then
+ : public Fossilized<$(T.directSuperClass)>
+% else
+ : public FossilizedRecordVal
+% end
{
- serializeEnum(serializer, value);
+% for _,f in ipairs(T.directFields) do
+ Fossilized<decltype($T::$f)> $f;
+% end
+};
+
+/// Serialize a `value` of type `$T`
+void serialize(IRSerializer const& serializer, $T& value)
+{
+ SLANG_UNUSED(value);
+ SLANG_SCOPED_SERIALIZER_STRUCT(serializer);
+% if T.directSuperClass then
+ serialize(serializer, static_cast<$(T.directSuperClass)&>(value));
+% end
+% for _,f in ipairs(T.directFields) do
+ serialize(serializer, value.$f);
+% end
}
+% end
+#else // FIDDLE OUTPUT:
+#define FIDDLE_GENERATED_OUTPUT_ID 0
+#include "slang-serialize-ir.cpp.fiddle"
+#endif // FIDDLE END
+
-/// Serialize a `value` of type `IRModuleInfo`, currently no extra information
-/// besides the IRModule
-SLANG_DECLARE_FOSSILIZED_AS_MEMBER(IRModuleInfo, module);
-void serialize(IRSerializer const& serializer, IRModuleInfo& value)
+// IROps are serialized as integers, and given a stable name
+SLANG_DECLARE_FOSSILIZED_AS(IROp, FossilUInt);
+void serialize(Serializer const& serializer, IROp& value)
{
- serialize(serializer, value.module);
+ auto stableName = isWriting(serializer) ? getOpcodeStableName(value) : kInvalidStableName;
+ serializeEnum(serializer, stableName);
+ if (isReading(serializer))
+ {
+ value = getStableNameOpcode(stableName);
+ // It's possible we're reading a module serialized by a future version of
+ // Slang with as-yet unknown instructions.
+ // if this is the case, return IRUnrecognized and we can handle it later
+ if (value == kIROp_Invalid)
+ value = kIROp_Unrecognized;
+ }
}
//
@@ -309,8 +395,9 @@ void serializeObject(IRSerializer const& serializer, IRModule*& value, IRModule*
void IRSerialWriteContext::handleIRModule(IRSerializer const& serializer, IRModule*& value)
{
SLANG_SCOPED_SERIALIZER_STRUCT(serializer);
- serialize(serializer, value->m_name);
serialize(serializer, value->m_moduleInst);
+ serialize(serializer, value->m_name);
+ serialize(serializer, value->m_version);
}
void IRSerialReadContext::handleIRModule(IRSerializer const& serializer, IRModule*& value)
@@ -319,8 +406,9 @@ void IRSerialReadContext::handleIRModule(IRSerializer const& serializer, IRModul
value = new IRModule{_session};
SLANG_ASSERT(!_module);
_module = value;
- serialize(serializer, value->m_name);
serialize(serializer, value->m_moduleInst);
+ serialize(serializer, value->m_name);
+ serialize(serializer, value->m_version);
value->m_moduleInst->module = value;
}
@@ -357,7 +445,9 @@ void writeSerializedModuleIR(
// The flow here is very similar to writeSerializedModuleAST which is very
// well documented.
- IRModuleInfo moduleInfo{.module = irModule};
+ IRModuleInfo moduleInfo;
+ moduleInfo.fullVersion = SLANG_TAG_VERSION;
+ moduleInfo.module = irModule;
BlobBuilder blobBuilder;
{
@@ -375,10 +465,37 @@ void writeSerializedModuleIR(
cursor.addDataChunk(PropertyKeys<IRModule>::IRModule, data, size);
}
+Result readSerializedModuleInfo(
+ RIFF::Chunk const* chunk,
+ String& compilerVersion,
+ UInt& version,
+ String& name)
+{
+ auto dataChunk = as<RIFF::DataChunk>(chunk);
+ if (!dataChunk)
+ {
+ SLANG_UNEXPECTED("invalid format for serialized module IR");
+ }
+
+ Fossil::AnyValPtr rootValPtr =
+ Fossil::getRootValue(dataChunk->getPayload(), dataChunk->getPayloadSize());
+ if (!rootValPtr)
+ {
+ SLANG_UNEXPECTED("invalid format for serialized module IR");
+ }
+
+ Fossilized<IRModuleInfo>* fossilizedModuleInfo = cast<Fossilized<IRModuleInfo>>(rootValPtr);
+ Fossilized<IRModule>* fossilizedModule = fossilizedModuleInfo->module;
+ version = fossilizedModule->m_version;
+ compilerVersion = fossilizedModuleInfo->fullVersion.get();
+ name = fossilizedModuleInfo->module->m_name.get();
+ return SLANG_OK;
+}
+
//
// Read a module, this currently does not do any on-demand loading
//
-void readSerializedModuleIR(
+Result readSerializedModuleIR(
RIFF::Chunk const* chunk,
Session* session,
SerialSourceLocReader* sourceLocReader,
@@ -397,6 +514,13 @@ void readSerializedModuleIR(
SLANG_UNEXPECTED("invalid format for serialized module IR");
}
+ Fossilized<IRModuleInfo>* fossilizedModuleInfo = cast<Fossilized<IRModuleInfo>>(rootValPtr);
+
+ // Only one version supported so far, if we had multiple versions to
+ // support this is where we might branch
+ if (fossilizedModuleInfo->serializationVersion != IRModuleInfo::kSupportedSerializationVersion)
+ return SLANG_FAIL;
+
IRModuleInfo info;
{
auto sharedDecodingContext = RefPtr(new IRSerialReadContext(session, sourceLocReader));
@@ -417,20 +541,28 @@ void readSerializedModuleIR(
// deserialization we didn't necessarily have this information handy at the
// time.
//
- auto go = [](auto&& go, IRInst* parent, IRInst* inst) -> void
+ bool hasUnrecognizedInsts = false;
+ auto go = [&](auto&& go, IRInst* parent, IRInst* inst) -> void
{
+ if (inst->getOp() == kIROp_Unrecognized)
+ hasUnrecognizedInsts = true;
+
inst->parent = parent;
for (const auto child : inst->getDecorationsAndChildren())
go(go, inst, child);
};
go(go, nullptr, info.module->getModuleInst());
+ if (hasUnrecognizedInsts)
+ return SLANG_FAIL;
+
//
// Module is finally valid (or at least as much as it was going it) and
// ready to be used
//
info.module->buildMangledNameToGlobalInstMap();
outIRModule = info.module;
+ return SLANG_OK;
}
} // namespace Slang
diff --git a/source/slang/slang-serialize-ir.h b/source/slang/slang-serialize-ir.h
index 0ac188df7..c2d0380cb 100644
--- a/source/slang/slang-serialize-ir.h
+++ b/source/slang/slang-serialize-ir.h
@@ -15,10 +15,16 @@ void writeSerializedModuleIR(
IRModule* moduleDecl,
SerialSourceLocWriter* sourceLocWriter);
-void readSerializedModuleIR(
+[[nodiscard]] Result readSerializedModuleIR(
RIFF::Chunk const* chunk,
Session* session,
SerialSourceLocReader* sourceLocReader,
RefPtr<IRModule>& outIRModule);
+[[nodiscard]] Result readSerializedModuleInfo(
+ RIFF::Chunk const* chunk,
+ String& compilerVersion,
+ UInt& version,
+ String& name);
+
} // namespace Slang
diff --git a/source/slang/slang.cpp b/source/slang/slang.cpp
index b616eb555..090e34e9a 100644
--- a/source/slang/slang.cpp
+++ b/source/slang/slang.cpp
@@ -763,7 +763,7 @@ SlangResult Session::_readBuiltinModule(
// to deserialize the IR module.
//
RefPtr<IRModule> irModule;
- readSerializedModuleIR(irChunk, this, sourceLocReader, irModule);
+ SLANG_RETURN_ON_FAIL(readSerializedModuleIR(irChunk, this, sourceLocReader, irModule));
irModule->setName(module->getNameObj());
module->setIRModule(irModule);
@@ -1368,6 +1368,7 @@ Linkage::Linkage(Session* session, ASTBuilder* astBuilder, Linkage* builtinLinka
, m_sourceManager(&m_defaultSourceManager)
, m_astBuilder(astBuilder)
, m_cmdLineContext(new CommandLineContext())
+ , m_stringSlicePool(StringSlicePool::Style::Default)
{
getNamePool()->setRootNamePool(session->getRootNamePool());
@@ -1663,6 +1664,47 @@ SLANG_NO_THROW slang::IModule* SLANG_MCALL Linkage::loadModuleFromIRBlob(
return loadModuleFromBlob(moduleName, path, source, ModuleBlobType::IR, outDiagnostics);
}
+SLANG_NO_THROW SlangResult SLANG_MCALL Linkage::loadModuleInfoFromIRBlob(
+ slang::IBlob* source,
+ SlangInt& outModuleVersion,
+ const char*& outModuleCompilerVersion,
+ const char*& outModuleName)
+{
+ // We start by reading the content of the file as
+ // an in-memory RIFF container.
+ //
+ auto rootChunk = RIFF::RootChunk::getFromBlob(source);
+ if (!rootChunk)
+ {
+ return SLANG_FAIL;
+ }
+
+ auto moduleChunk = ModuleChunk::find(rootChunk);
+ if (!moduleChunk)
+ {
+ return SLANG_FAIL;
+ }
+
+ auto irChunk = moduleChunk->findIR();
+ if (!irChunk)
+ {
+ return SLANG_FAIL;
+ }
+
+ RefPtr<IRModule> irModule;
+ String compilerVersion;
+ UInt version;
+ String name;
+ SLANG_RETURN_ON_FAIL(readSerializedModuleInfo(irChunk, compilerVersion, version, name));
+ const auto compilerVersionSlice = m_stringSlicePool.addAndGetSlice(compilerVersion);
+ const auto nameSlice = m_stringSlicePool.addAndGetSlice(name);
+ outModuleCompilerVersion = compilerVersionSlice.begin();
+ outModuleName = nameSlice.begin();
+ outModuleVersion = SlangInt(version);
+
+ return SLANG_OK;
+}
+
SLANG_NO_THROW SlangResult SLANG_MCALL Linkage::createCompositeComponentType(
slang::IComponentType* const* componentTypes,
SlangInt componentTypeCount,
@@ -6801,7 +6843,7 @@ SlangResult Linkage::loadSerializedModuleContents(
module->setModuleDecl(moduleDecl);
RefPtr<IRModule> irModule;
- readSerializedModuleIR(irChunk, session, sourceLocReader, irModule);
+ SLANG_RETURN_ON_FAIL(readSerializedModuleIR(irChunk, session, sourceLocReader, irModule));
module->setIRModule(irModule);
// The handling of file dependencies is complicated, because of
generated by cgit v1.2.3 (git 2.39.1) at 2026-06-01 18:48:15 +0000