Stable names and backwards compat for serialized IR modules (#7644)

* stable names

* tests, options and ci for stable names

* Add back compat design document

* fix warnings

* formatting

* comment

* neaten

* regenerate command line reference

* consolidate ci scripts

* faster ci

* remove libreadline

* Move new function to end of interface

---------

Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>

11 files changed, 1139 insertions, 24 deletions

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