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>

20 files changed, 1728 insertions, 25 deletions

diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml
index b18bc7bad..1f831415f 100644
--- a/.github/workflows/ci.yml
+++ b/.github/workflows/ci.yml
@@ -123,7 +123,7 @@ jobs:
         run: |
           if [[ "${{ matrix.os }}" = "linux" ]]; then
               sudo apt-get update
-              sudo apt-get install -y libx11-dev
+              sudo apt-get install -y libx11-dev 
           fi
       - name: Setup Node.js
         if: matrix.os == 'linux'
@@ -139,6 +139,12 @@ jobs:
           platform: ${{matrix.platform}}
           config: ${{matrix.config}}
           build-llvm: ${{ matrix.build-llvm }}
+
+      # Don't need to check this on every config
+      - name: Check Stable Names Table
+        if: ${{ matrix.os == 'linux' && matrix.config == 'debug' }}
+        run: ./extras/check-ir-stable-names-gh-actions.sh
+
       - name: Build Slang
         if: steps.filter.outputs.should-run == 'true'
         run: |
diff --git a/docs/command-line-slangc-reference.md b/docs/command-line-slangc-reference.md
index d8d58ec9b..c7d216a6e 100644
--- a/docs/command-line-slangc-reference.md
+++ b/docs/command-line-slangc-reference.md
@@ -681,6 +681,16 @@ Verify IR in the front-end.
 Disassemble and print the module IR. 
 
 
+<a id="get-module-info"></a>
+### -get-module-info
+Print the name and version of a serialized IR Module 
+
+
+<a id="get-supported-module-versions"></a>
+### -get-supported-module-versions
+Print the minimum and maximum module versions this compiler supports 
+
+
 
 <a id="Repro"></a>
 ## Repro
diff --git a/docs/design/backwards-compat-for-ir-modules.md b/docs/design/backwards-compat-for-ir-modules.md
new file mode 100644
index 000000000..4f859c66f
--- /dev/null
+++ b/docs/design/backwards-compat-for-ir-modules.md
@@ -0,0 +1,199 @@
+# Design Document: Slang IR Module Backwards Compatibility
+
+## Overview
+
+This document describes the design and implementation of backwards compatibility support for serialized Slang IR modules. The feature enables Slang to load IR modules compiled with different versions of the compiler, providing version information and graceful handling of incompatible modules.
+
+## Motivation
+
+As Slang evolves, the intermediate representation (IR) may change with new instructions being added or existing ones being modified. Without backwards compatibility:
+
+- Users cannot load modules compiled with older versions of Slang
+- There's no way to detect version mismatches between modules
+- Module compatibility issues are opaque to users
+
+This feature addresses these issues by introducing versioning and stable instruction naming.
+
+## User-Facing Changes
+
+### New Command Line Options
+
+1. **`-get-module-info <module-file>`**
+
+   - Prints information about a serialized IR module without loading it
+   - Output includes:
+     - Module name
+     - Module version
+     - Compiler version that created the module
+   - Example usage: `slangc -get-module-info mymodule.slang-module`
+
+2. **`-get-supported-module-versions`**
+   - Prints the range of module versions this compiler supports
+   - Output includes minimum and maximum supported versions
+   - Example usage: `slangc -get-supported-module-versions`
+
+### API Changes
+
+New method in `ISession` interface:
+
+```cpp
+SlangResult loadModuleInfoFromIRBlob(
+    slang::IBlob* source,
+    SlangInt& outModuleVersion,
+    const char*& outModuleCompilerVersion,
+    const char*& outModuleName);
+```
+
+This allows programmatic inspection of module metadata without full deserialization.
+
+## Technical Design
+
+### Stable Instruction Names
+
+The core mechanism for backwards compatibility is the introduction of stable names for IR instructions:
+
+1. **Stable Name Table** (`slang-ir-insts-stable-names.lua`)
+
+   - Maps instruction names to unique integer IDs
+   - IDs are permanent once assigned
+   - New instructions get new IDs, never reusing old ones
+
+2. **Runtime Mapping**
+   - `getOpcodeStableName(IROp)`: Convert runtime opcode to stable ID
+   - `getStableNameOpcode(UInt)`: Convert stable ID back to runtime opcode
+   - Unknown stable IDs map to `kIROp_Unrecognized`
+
+### Module Versioning
+
+Two types of versions are tracked:
+
+1. **Module Version** (`IRModule::m_version`)
+
+   - Semantic version of the IR instruction set
+   - Range: `k_minSupportedModuleVersion` to `k_maxSupportedModuleVersion`
+   - Stored in each serialized module
+
+2. **Serialization Version** (`IRModuleInfo::serializationVersion`)
+   - Version of the serialization format itself
+   - Currently version 0
+   - Allows future changes to serialization structure
+
+### Compiler Version Tracking
+
+Each module stores the exact compiler version (`SLANG_TAG_VERSION`) that created it. This enables version-specific workarounds if needed in the future.
+
+### Validation System
+
+A GitHub Actions workflow (`check-ir-stable-names.yml`) ensures consistency:
+
+1. **Check Mode**: Validates that:
+
+   - All IR instructions have stable names
+   - No duplicate stable IDs exist
+   - The stable name table is a bijection with current instructions
+
+2. **Update Mode**: Automatically assigns stable IDs to new instructions
+
+The validation is implemented in `check-ir-stable-names.lua` which:
+
+- Loads instruction definitions from `slang-ir-insts.lua`
+- Compares against `slang-ir-insts-stable-names.lua`
+- Reports missing entries or inconsistencies
+
+## Breaking Changes and Version Management
+
+### When to Update Module Version
+
+The module version must be updated when:
+
+1. **Adding Instructions** (Minor Version Bump)
+
+   - Increment `k_maxSupportedModuleVersion`
+   - Older compilers can still load modules that don't use new instructions
+
+2. **Removing Instructions** (Major Version Bump)
+
+   - Increment `k_maxSupportedModuleVersion`
+   - Update `k_minSupportedModuleVersion` to exclude versions with removed instructions
+   - This breaks compatibility with older modules using removed instructions
+
+3. **Changing Instruction Semantics**
+   - Even if the instruction name remains the same
+   - Requires version bump to prevent incorrect behavior
+   - To avoid bumping the minimum supported version, one may instead introduce
+     a new instruction and just bump `k_maxSupportedModuleVersion`
+
+### Serialization Format Changes
+
+Changes to how data is serialized (not what data) require updating `serializationVersion`:
+
+- Changes to the RIFF container structure
+- Different encoding for instruction payloads
+- Reordering of serialized data
+
+## Implementation Details
+
+### Module Loading Flow
+
+1. **Version Check**
+
+   ```cpp
+   if (fossilizedModuleInfo->serializationVersion != IRModuleInfo::kSupportedSerializationVersion)
+       return SLANG_FAIL;
+   ```
+
+2. **Instruction Deserialization**
+
+   - Stable IDs are converted to runtime opcodes
+   - Unknown IDs become `kIROp_Unrecognized`
+
+3. **Validation Pass**
+   - After deserialization, check for any `kIROp_Unrecognized` instructions
+   - Fail loading if any are found
+
+### Error Handling
+
+- Incompatible serialization versions: Immediate failure
+- Unknown instructions: Mark as unrecognized, fail after full deserialization
+  (this should be caught by the next check)
+- Module version out of range: Fail after deserialization
+
+## Future Considerations
+
+### Potential Enhancements
+
+1. **Graceful Degradation**
+
+   - Skip unrecognized instructions if they're not critical
+   - Provide compatibility shims for removed instructions
+
+2. **Module Migration Tools**
+
+   - Utility to upgrade old modules to new formats
+   - Batch processing for large codebases
+
+### Maintenance Guidelines
+
+1. **Regular CI Validation**
+
+   - The GitHub Action ensures stable names stay synchronized
+   - Catches missing entries before merge
+
+2. **Version Documentation**
+
+   - Maintain changelog of what changed in each module version
+   - Document any version-specific workarounds
+
+3. **Testing**
+   - Test loading of modules from previous versions
+   - Verify error messages for incompatible modules
+
+## Conclusion
+
+This backwards compatibility system provides a robust foundation for Slang IR evolution while maintaining compatibility where possible. The combination of stable instruction naming, comprehensive versioning, and automated validation ensures that:
+
+- Users can reliably use modules across Slang versions
+- Developers can evolve the IR with clear compatibility boundaries
+- Version mismatches are detected and reported clearly
+
+The system is designed to be maintainable and extensible, with clear guidelines for when and how to make breaking changes.
diff --git a/extras/check-ir-stable-names-gh-actions.sh b/extras/check-ir-stable-names-gh-actions.sh
new file mode 100755
index 000000000..78c9945a6
--- /dev/null
+++ b/extras/check-ir-stable-names-gh-actions.sh
@@ -0,0 +1,29 @@
+#!/usr/bin/env bash
+set -e
+
+# The github action runners don't have libreadline and we don't need it
+make -C external/lua MYCFLAGS="-DLUA_USE_POSIX" MYLIBS=""
+external/lua/lua -v
+
+# Run the check
+if ! ./external/lua/lua extras/check-ir-stable-names.lua check; then
+  echo "Check failed. Running update..."
+  ./external/lua/lua extras/check-ir-stable-names.lua update
+
+  echo -e "\n=== Diff of changes made ==="
+  git diff --no-index --color=always source/slang/slang-ir-insts-stable-names.lua || true
+
+  # Also create a summary for GitHub Actions
+  if [ -n "$GITHUB_STEP_SUMMARY" ]; then
+    {
+      echo "## IR Stable Names Table Update Required"
+      echo "The following changes need to be made to \`source/slang/slang-ir-insts-stable-names.lua\`:"
+      echo '```diff'
+      git diff --no-index source/slang/slang-ir-insts-stable-names.lua
+      echo '```'
+    } >>"$GITHUB_STEP_SUMMARY"
+  fi
+
+  # Fail the job since the check failed
+  exit 1
+fi
diff --git a/extras/check-ir-stable-names.lua b/extras/check-ir-stable-names.lua
new file mode 100644
index 000000000..90cff4683
--- /dev/null
+++ b/extras/check-ir-stable-names.lua
@@ -0,0 +1,280 @@
+-- Helper function to flatten the instruction hierarchy
+local function flatten_instructions(insts, prefix, result)
+	prefix = prefix or ""
+	result = result or {}
+
+	for _, entry in ipairs(insts) do
+		for name, data in pairs(entry) do
+			local full_name = prefix == "" and name or (prefix .. "." .. name)
+
+			-- If it's a table with numeric indices, it has children
+			if type(data) == "table" and #data > 0 then
+				flatten_instructions(data, full_name, result)
+			else
+				-- Add the current instruction
+				table.insert(result, full_name)
+			end
+		end
+	end
+
+	return result
+end
+
+-- Load instruction definitions
+local function load_instructions(filename)
+	local chunk, err = loadfile(filename)
+	if not chunk then
+		error("Failed to load instruction file: " .. filename .. " - " .. (err or "unknown error"))
+	end
+
+	-- Just execute it normally
+	local result = chunk()
+
+	-- If the file sets a global 'insts', use that
+	if result.insts then
+		return result.insts
+	end
+
+	error("Instruction file must return a table with 'insts' entry")
+end
+
+-- Load stable names table
+local function load_stable_names(filename)
+	local file = io.open(filename, "r")
+	if not file then
+		-- File doesn't exist, return empty table
+		return {}
+	end
+	file:close()
+
+	local chunk, err = loadfile(filename)
+	if not chunk then
+		error("Failed to load stable names file: " .. filename .. " - " .. (err or "unknown error"))
+	end
+
+	local result = chunk()
+
+	-- Validate structure
+	if type(result) ~= "table" then
+		error("Stable names file must return a table")
+	end
+
+	for name, id in pairs(result) do
+		if type(name) ~= "string" then
+			error(string.format("Invalid key: expected string, got %s", type(name)))
+		end
+		if type(id) ~= "number" then
+			error(string.format("Invalid value for '%s': expected number, got %s", name, type(id)))
+		end
+	end
+
+	return result
+end
+
+-- Save stable names table
+local function save_stable_names(filename, stable_names)
+	local file, err = io.open(filename, "w")
+	if not file then
+		error("Failed to open file for writing: " .. filename .. " - " .. (err or "unknown error"))
+	end
+
+	file:write("-- This file is machine generated! any entries written below will be preserved,\n")
+	file:write("-- but things like comments or anything outside the schema won't be preserved\n")
+	file:write("return {\n")
+
+	-- Sort by ID for consistent output
+	local sorted_entries = {}
+	for name, id in pairs(stable_names) do
+		table.insert(sorted_entries, { name = name, id = id })
+	end
+	table.sort(sorted_entries, function(a, b)
+		return a.id < b.id
+	end)
+
+	for _, entry in ipairs(sorted_entries) do
+		-- Escape quotes in name
+		local escaped_name = entry.name:gsub('"', '\\"')
+		file:write(string.format('\t["%s"] = %d,\n', escaped_name, entry.id))
+	end
+	file:write("}\n")
+	file:close()
+end
+
+-- Check for unique IDs
+local function check_unique_ids(stable_names)
+	local seen_ids = {}
+	local duplicates = {}
+
+	for name, id in pairs(stable_names) do
+		if seen_ids[id] then
+			if not duplicates[id] then
+				duplicates[id] = { seen_ids[id] }
+			end
+			table.insert(duplicates[id], name)
+		else
+			seen_ids[id] = name
+		end
+	end
+
+	return duplicates
+end
+
+-- Check bijection
+local function check_bijection(inst_names, stable_names)
+	local missing_from_stable = {}
+	local extra_in_stable = {}
+
+	-- Check for instructions missing from stable names
+	for _, name in ipairs(inst_names) do
+		if stable_names[name] == nil then
+			table.insert(missing_from_stable, name)
+		end
+	end
+
+	-- Check for stable names not in instructions
+	local inst_name_set = {}
+	for _, name in ipairs(inst_names) do
+		inst_name_set[name] = true
+	end
+
+	for name, _ in pairs(stable_names) do
+		if not inst_name_set[name] then
+			table.insert(extra_in_stable, name)
+		end
+	end
+
+	return missing_from_stable, extra_in_stable
+end
+
+-- Get next available ID
+local function get_next_id(stable_names)
+	local max_id = -1
+	for _, id in pairs(stable_names) do
+		if id > max_id then
+			max_id = id
+		end
+	end
+	return max_id + 1
+end
+
+-- Print usage
+local function print_usage()
+	print("Usage: lua check_instructions.lua check|update [inst_file] [stable_file]")
+	print("Commands:")
+	print("  check  - Check bijection and uniqueness (default)")
+	print("  update - Add missing instructions to stable names")
+end
+
+-- Main program
+local function main(args)
+	local command = args[1] or "check"
+	local inst_file = args[2] or "source/slang/slang-ir-insts.lua"
+	local stable_file = args[3] or "source/slang/slang-ir-insts-stable-names.lua"
+
+	-- Validate command
+	local valid_commands = { check = true, update = true }
+	if not valid_commands[command] then
+		print("ERROR: Invalid command: " .. command)
+		print_usage()
+		return 1
+	end
+
+	-- Load data with error handling
+	local ok, insts_or_err = pcall(load_instructions, inst_file)
+	if not ok then
+		print("ERROR: " .. insts_or_err)
+		return 1
+	end
+	local insts = insts_or_err
+
+	ok, stable_names = pcall(load_stable_names, stable_file)
+	if not ok then
+		print("ERROR: " .. stable_names)
+		return 1
+	end
+
+	-- Flatten instruction hierarchy
+	local inst_names = flatten_instructions(insts)
+
+	local has_errors = false
+
+	if command == "check" or command == "all" then
+		print("=== Checking stable names ===")
+
+		-- Check unique IDs
+		local duplicate_ids = check_unique_ids(stable_names)
+		if next(duplicate_ids) then
+			has_errors = true
+			print("ERROR: Duplicate IDs found:")
+			for id, names in pairs(duplicate_ids) do
+				print(string.format("  - ID %d used by: %s", id, table.concat(names, ", ")))
+			end
+		else
+			print("✓ All IDs are unique")
+		end
+
+		-- Check bijection
+		local missing, extra = check_bijection(inst_names, stable_names)
+
+		if #missing > 0 then
+			has_errors = true
+			print(string.format("ERROR: %d instructions missing from stable names:", #missing))
+			for _, name in ipairs(missing) do
+				print("  - " .. name)
+			end
+		else
+			print("✓ All instructions have stable names")
+		end
+
+		if #extra > 0 then
+			print(string.format("WARNING: %d extra entries in stable names (not in instructions):", #extra))
+			for _, name in ipairs(extra) do
+				print("  - " .. name)
+			end
+		else
+			print("✓ No extra entries in stable names")
+		end
+
+		if not has_errors and #extra == 0 then
+			print("✓ Is a bijection")
+		end
+	end
+
+	if command == "update" or command == "all" then
+		print("=== Updating stable names ===")
+
+		-- Don't update if there are errors
+		if has_errors then
+			print("ERROR: Cannot update due to errors in existing stable names")
+			return 1
+		end
+
+		local missing, _ = check_bijection(inst_names, stable_names)
+
+		if #missing > 0 then
+			-- Add missing instructions
+			local next_id = get_next_id(stable_names)
+
+			for _, name in ipairs(missing) do
+				stable_names[name] = next_id
+				next_id = next_id + 1
+			end
+
+			-- Save updated file
+			local ok, err = pcall(save_stable_names, stable_file, stable_names)
+			if not ok then
+				print("ERROR: Failed to save: " .. err)
+				return 1
+			end
+
+			print(string.format("Added %d new instructions to %s", #missing, stable_file))
+		else
+			print("No missing instructions to add")
+		end
+	end
+
+	return has_errors and 1 or 0
+end
+
+-- Run the program
+os.exit(main(arg) or 0)
diff --git a/include/slang.h b/include/slang.h
index 83009d2e1..9d7fe2d49 100644
--- a/include/slang.h
+++ b/include/slang.h
@@ -1034,6 +1034,9 @@ typedef uint32_t SlangSizeT;
         SkipDownstreamLinking, // bool, experimental
         DumpModule,
 
+        GetModuleInfo,              // Print serialized module version and name
+        GetSupportedModuleVersions, // Print the min and max module versions this compiler supports
+
         EmitSeparateDebug, // bool
 
         // Floating point denormal handling modes
@@ -4167,6 +4170,16 @@ struct ISession : public ISlangUnknown
         slang::TypeReflection* interfaceType,
         uint32_t* outRTTIDataBuffer,
         uint32_t bufferSizeInBytes) = 0;
+
+    /** Read module info (name and version) from a module blob
+     *
+     * The returned pointers are valid for as long as the session.
+     */
+    virtual SLANG_NO_THROW SlangResult SLANG_MCALL loadModuleInfoFromIRBlob(
+        slang::IBlob* source,
+        SlangInt& outModuleVersion,
+        const char*& outModuleCompilerVersion,
+        const char*& outModuleName) = 0;
 };
 
     #define SLANG_UUID_ISession ISession::getTypeGuid()
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
diff --git a/tests/ir/dump-module-info.slang b/tests/ir/dump-module-info.slang
new file mode 100644
index 000000000..9ce421161
--- /dev/null
+++ b/tests/ir/dump-module-info.slang
@@ -0,0 +1,13 @@
+//TEST:COMPILE: tests/ir/dump-module-info.slang -o tests/ir/dump-module-info.slang-module 
+
+//TEST:SIMPLE(filecheck=CHECK): -get-module-info tests/ir/dump-module-info.slang-module
+
+module "foo";
+
+// CHECK: Module Name: foo
+// This will need bumping whenever we bump the ir module version
+// CHECK: Module Version: 0
+// Just check that this is in the output with some string
+// CHECK: Compiler Version: {{.+}}
+
+