Wave intrinsics. (#3164)

* Wave intrinsics.

* scalar intrinsics.

---------

Co-authored-by: Yong He <yhe@nvidia.com>

10 files changed, 573 insertions, 143 deletions

diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang
index 3c966bb4a..fd668f73a 100644
--- a/source/slang/hlsl.meta.slang
+++ b/source/slang/hlsl.meta.slang
@@ -1353,13 +1353,27 @@ matrix<int,N,M> asint(matrix<int,N,M> x)
 
 // Reinterpret bits of double as a uint (HLSL SM 5.0)
 
-__target_intrinsic(hlsl)
-__target_intrinsic(glsl, "{ uvec2 v = unpackDouble2x32($0); $1 = v.x; $2 = v.y; }")
 __glsl_extension(GL_ARB_gpu_shader5)
-__target_intrinsic(cpp, "$P_asuint($0, $1, $2)")
-__target_intrinsic(cuda, "$P_asuint($0, $1, $2)")
 [__readNone]
-void asuint(double value, out uint lowbits, out uint highbits);
+void asuint(double value, out uint lowbits, out uint highbits)
+{
+    __target_switch
+    {
+    case hlsl: __intrinsic_asm "asuint";
+    case glsl: __intrinsic_asm "{ uvec2 v = unpackDouble2x32($0); $1 = v.x; $2 = v.y; }";
+    case cpp:
+    case cuda:
+        __intrinsic_asm "$P_asuint($0, $1, $2)";
+    case spirv:
+        let uv = spirv_asm
+        {
+            result : $$uint2 = OpBitcast $value;
+        };
+        lowbits = uv.x;
+        highbits = uv.y;
+        return;
+    }
+}
 
 // Reinterpret bits as a uint (HLSL SM 4.0)
 
@@ -2752,20 +2766,30 @@ void InterlockedXor(__ref uint dest, uint value, out uint original_value);
 
 __generic<T : __BuiltinFloatingPointType>
 __target_intrinsic(hlsl)
-__target_intrinsic(cuda, "$P_isfinite($0)")
-__target_intrinsic(cpp, "$P_isfinite($0)")
 [__readNone]
 bool isfinite(T x)
 {
-    return !(isinf(x) || isnan(x));
+    __target_switch
+    {
+    case hlsl: __intrinsic_asm "isfinite";
+    case cuda:
+    case cpp:
+        __intrinsic_asm "$P_isfinite($0)";
+    default:
+        return !(isinf(x) || isnan(x));
+    }
 }
 
 __generic<T : __BuiltinFloatingPointType, let N : int>
-__target_intrinsic(hlsl)
 [__readNone]
 vector<bool, N> isfinite(vector<T, N> x)
 {
-    VECTOR_MAP_UNARY(bool, N, isfinite, x);
+    __target_switch
+    {
+    case hlsl: __intrinsic_asm "isfinite";
+    default:
+        VECTOR_MAP_UNARY(bool, N, isfinite, x);
+    }
 }
 
 __generic<T : __BuiltinFloatingPointType, let N : int, let M : int>
@@ -2778,20 +2802,36 @@ matrix<bool, N, M> isfinite(matrix<T, N, M> x)
 
 // Is floating-point value infinite?
 __generic<T : __BuiltinFloatingPointType>
-__target_intrinsic(hlsl)
-__target_intrinsic(glsl)
-__target_intrinsic(cuda, "$P_isinf($0)")
-__target_intrinsic(cpp, "$P_isinf($0)")
 [__readNone]
-bool isinf(T x);
+bool isinf(T x)
+{
+    __target_switch
+    {
+    case hlsl:
+    case glsl:
+        __intrinsic_asm "isinf";
+    case cuda:
+    case cpp:
+        __intrinsic_asm "$P_isinf($0)";
+    case spirv:
+        return spirv_asm { result:$$bool = OpIsInf $x};
+    }
+}
 
 __generic<T : __BuiltinFloatingPointType, let N : int>
-__target_intrinsic(hlsl)
-__target_intrinsic(glsl)
 [__readNone]
 vector<bool, N> isinf(vector<T, N> x)
 {
-    VECTOR_MAP_UNARY(bool, N, isinf, x);
+    __target_switch
+    {
+    case hlsl:
+    case glsl:
+        __intrinsic_asm "isinf";
+    case spirv:
+        return spirv_asm { result:$$vector<bool,N> = OpIsInf $x};
+    default:
+        VECTOR_MAP_UNARY(bool, N, isinf, x);
+    }
 }
 
 __generic<T : __BuiltinFloatingPointType, let N : int, let M : int>
@@ -2804,20 +2844,36 @@ matrix<bool, N, M> isinf(matrix<T, N, M> x)
 
 // Is floating-point value not-a-number?
 __generic<T : __BuiltinFloatingPointType>
-__target_intrinsic(hlsl)
-__target_intrinsic(glsl)
-__target_intrinsic(cuda, "$P_isnan($0)")
-__target_intrinsic(cpp, "$P_isnan($0)")
 [__readNone]
-bool isnan(T x);
+bool isnan(T x)
+{
+    __target_switch
+    {
+    case hlsl:
+    case glsl:
+        __intrinsic_asm "isnan";
+    case cuda:
+    case cpp:
+        __intrinsic_asm "$P_isnan($0)";
+    case spirv:
+        return spirv_asm { result:$$bool = OpIsNan $x};
+    }
+}
 
 __generic<T : __BuiltinFloatingPointType, let N : int>
-__target_intrinsic(hlsl)
-__target_intrinsic(glsl)
 [__readNone]
 vector<bool, N> isnan(vector<T, N> x)
 {
-    VECTOR_MAP_UNARY(bool, N, isnan, x);
+    __target_switch
+    {
+    case hlsl:
+    case glsl:
+        __intrinsic_asm "isnan";
+    case spirv:
+        return spirv_asm { result:$$vector<bool, N> = OpIsNan $x};
+    default:
+        VECTOR_MAP_UNARY(bool, N, isnan, x);
+    }
 }
 
 __generic<T : __BuiltinFloatingPointType, let N : int, let M : int>
@@ -2832,7 +2888,6 @@ matrix<bool, N, M> isnan(matrix<T, N, M> x)
 
 __generic<T : __BuiltinFloatingPointType>
 __target_intrinsic(hlsl)
-__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Ldexp _0 _1")
 [__readNone]
 T ldexp(T x, T exp)
 {
@@ -2841,7 +2896,6 @@ T ldexp(T x, T exp)
 
 __generic<T : __BuiltinFloatingPointType, let N : int>
 __target_intrinsic(hlsl)
-__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Ldexp _0 _1")
 [__readNone]
 vector<T, N> ldexp(vector<T, N> x, vector<T, N> exp)
 {
@@ -2948,14 +3002,14 @@ __target_intrinsic(hlsl)
 __target_intrinsic(glsl, "(log( $0 ) * $S0( 0.43429448190325182765112891891661) )" )
 __target_intrinsic(cuda, "$P_log10($0)")
 __target_intrinsic(cpp, "$P_log10($0)")
-__target_intrinsic(spirv, "%baseElog = OpExtInst resultType resultId glsl450 Log _0; OpFMul resultType resultId _0 %baseElog const(_p,0.43429448190325182765112891891661)")
+__target_intrinsic(spirv, "%baseElog = OpExtInst resultType resultId glsl450 Log _0; OpFMul resultType resultId %baseElog const(_p,0.43429448190325182765112891891661)")
 [__readNone]
 T log10(T x);
 
 __generic<T : __BuiltinFloatingPointType, let N : int>
 __target_intrinsic(hlsl)
 __target_intrinsic(glsl, "(log( $0 ) * $S0(0.43429448190325182765112891891661) )" )
-__target_intrinsic(spirv, "%baseElog = OpExtInst resultType resultId glsl450 Log _0; OpVectorTimesScalar resultType resultId _0 %baseElog const(_p,0.43429448190325182765112891891661)")
+__target_intrinsic(spirv, "%baseElog = OpExtInst resultType resultId glsl450 Log _0; OpVectorTimesScalar resultType resultId %baseElog const(_p,0.43429448190325182765112891891661)")
 [__readNone]
 vector<T,N> log10(vector<T,N> x)
 {
@@ -3866,24 +3920,58 @@ matrix<T,N,M> saturate(matrix<T,N,M> x)
     MATRIX_MAP_UNARY(T, N, M, saturate, x);
 }
 
+__generic<T:__BuiltinArithmeticType, U:__BuiltinArithmeticType>
+__intrinsic_op($(kIROp_IntCast))
+T __int_cast(U val);
+
+__generic<T:__BuiltinArithmeticType, U:__BuiltinArithmeticType, let N : int>
+__intrinsic_op($(kIROp_IntCast))
+vector<T,N> __int_cast(vector<U,N> val);
+
 // Extract sign of value
 __generic<T : __BuiltinSignedArithmeticType>
-__target_intrinsic(hlsl)
-__target_intrinsic(glsl, "int(sign($0))")
-__target_intrinsic(cuda, "$P_sign($0)")
-__target_intrinsic(cpp, "$P_sign($0)")
-__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fi(FSign, SSign) _0")
 [__readNone]
-int sign(T x);
+int sign(T x)
+{
+    __target_switch
+    {
+    case hlsl: __intrinsic_asm "sign";
+    case glsl: __intrinsic_asm "int(sign($0))";
+    case cuda:
+    case cpp:
+        __intrinsic_asm "$P_sign($0)";
+    case spirv:
+        if (__isFloat<T>())
+            return spirv_asm
+            {
+                %fsign:$$T = OpExtInst glsl450 FSign $x;
+                result:$$int = OpConvertFToS %fsign
+            };
+        else
+            return __int_cast<int>(spirv_asm {OpExtInst $$T result glsl450 SSign $x});
+    }
+}
 
 __generic<T : __BuiltinSignedArithmeticType, let N : int>
-__target_intrinsic(hlsl)
-__target_intrinsic(glsl, "ivec$N0(sign($0))")
-__target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fi(FSign, SSign) _0")
 [__readNone]
 vector<int, N> sign(vector<T, N> x)
 {
-    VECTOR_MAP_UNARY(int, N, sign, x);
+    __target_switch
+    {
+    case hlsl: __intrinsic_asm "sign";
+    case glsl: __intrinsic_asm "ivec$N0(sign($0))";
+    case spirv:
+        if (__isFloat<T>())
+            return spirv_asm
+            {
+                %fsign:$$vector<T, N> = OpExtInst glsl450 FSign $x;
+                result:$$vector<int, N> = OpConvertFToS %fsign
+            };
+        else
+            return __int_cast<int>(spirv_asm {OpExtInst $$vector<T,N> result glsl450 SSign $x});
+    default:
+        VECTOR_MAP_UNARY(int, N, sign, x);
+    }
 }
 
 __generic<T : __BuiltinSignedArithmeticType, let N : int, let M : int>
@@ -4212,10 +4300,9 @@ WaveMask WaveGetConvergedMask()
         __intrinsic_asm "__activemask()";
     case spirv:
         let _true = true;
-        let _scope = 3; // subgroup
         return (spirv_asm
         {
-            OpGroupNonUniformBallot $$uint4 result $_scope $_true
+            OpGroupNonUniformBallot $$uint4 result Subgroup $_true
         }).x;
     }
 }
@@ -4563,8 +4650,11 @@ uint WaveMaskPrefixCountBits(WaveMask mask, bool value)
     case cuda: __intrinsic_asm "__popc(__ballot_sync($0, $1)  & _getLaneLtMask())";
     case hlsl: __intrinsic_asm "WavePrefixCountBits($1)";
     case spirv:
-        let _scope = 3u; // subgroup
-        return spirv_asm {OpGroupNonUniformBallotBitCount $$uint result $_scope 2 $value};
+        return spirv_asm
+        {
+            %mask:$$uint4 = OpGroupNonUniformBallot Subgroup $value;
+            OpGroupNonUniformBallotBitCount $$uint result Subgroup 2 %mask
+        };
     }
 }
 
@@ -5012,9 +5102,9 @@ vector<T,N> WaveMaskPrefixProduct(WaveMask mask, vector<T,N> expr)
             return spirv_asm
             {
                 // TODO: use the correct integer width
-                OpBitcast $$uint %uvalue $expr;
+                OpBitcast $$vector<uint,N> %uvalue $expr;
                 OpGroupNonUniformIMul $$vector<uint,N> %mulResult Subgroup ExclusiveScan %uvalue;
-                OpBitcast $$T result %mulResult
+                OpBitcast $$vector<T,N> result %mulResult
             };
         }
         else if (__isUnsignedInt<T>())
@@ -5074,9 +5164,9 @@ vector<T,N> WaveMaskPrefixSum(WaveMask mask, vector<T,N> expr)
             return spirv_asm
             {
                 // TODO: use the correct integer width
-                %uvalue: $$uint = OpBitcast $expr;
+                %uvalue: $$vector<uint,N> = OpBitcast $expr;
                 %mulResult: $$vector<uint,N> = OpGroupNonUniformIAdd Subgroup ExclusiveScan %uvalue;
-                result: $$T = OpBitcast  %mulResult
+                result: $$vector<T,N> = OpBitcast  %mulResult
             };
         }
         else if (__isUnsignedInt<T>())
@@ -5428,21 +5518,45 @@ matrix<T, N, M> WaveActiveBitXor(matrix<T, N, M> expr)
 __generic<T : __BuiltinArithmeticType>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupMax($0)")
-__target_intrinsic(hlsl)
+__spirv_capability(GroupNonUniformArithmetic)
 T WaveActiveMax(T expr)
 {
-    return WaveMaskMax(WaveGetActiveMask(), expr);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupMax($0)";
+    case hlsl: __intrinsic_asm "WaveActiveMax";
+    case spirv:
+        if (__isFloat<T>())
+            return spirv_asm {OpGroupNonUniformFMax $$T result Subgroup Reduce $expr};
+        else if (__isUnsignedInt<T>())
+            return spirv_asm {OpGroupNonUniformUMax $$T result Subgroup Reduce $expr};
+        else
+            return spirv_asm {OpGroupNonUniformSMax $$T result Subgroup Reduce $expr};
+    default:
+        return WaveMaskMax(WaveGetActiveMask(), expr);
+    }
 }
 
 __generic<T : __BuiltinArithmeticType, let N : int>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupMax($0)")
-__target_intrinsic(hlsl)
+__spirv_capability(GroupNonUniformArithmetic)
 vector<T, N> WaveActiveMax(vector<T, N> expr)
 {
-    return WaveMaskMax(WaveGetActiveMask(), expr);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupMax($0)";
+    case hlsl: __intrinsic_asm "WaveActiveMax";
+    case spirv:
+        if (__isFloat<T>())
+            return spirv_asm {OpGroupNonUniformFMax $$vector<T, N> result Subgroup Reduce $expr};
+        else if (__isUnsignedInt<T>())
+            return spirv_asm {OpGroupNonUniformUMax $$vector<T, N> result Subgroup Reduce $expr};
+        else
+            return spirv_asm {OpGroupNonUniformSMax $$vector<T, N> result Subgroup Reduce $expr};
+    default:
+        return WaveMaskMax(WaveGetActiveMask(), expr);
+    }
 }
 
 __generic<T : __BuiltinArithmeticType, let N : int, let M : int>
@@ -5455,21 +5569,45 @@ matrix<T, N, M> WaveActiveMax(matrix<T, N, M> expr)
 __generic<T : __BuiltinArithmeticType>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupMin($0)")
-__target_intrinsic(hlsl)
+__spirv_capability(GroupNonUniformArithmetic)
 T WaveActiveMin(T expr)
 {
-    return WaveMaskMin(WaveGetActiveMask(), expr);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupMin($0)";
+    case hlsl: __intrinsic_asm "WaveActiveMin";
+    case spirv:
+        if (__isFloat<T>())
+            return spirv_asm {OpGroupNonUniformFMin $$T result Subgroup Reduce $expr};
+        else if (__isUnsignedInt<T>())
+            return spirv_asm {OpGroupNonUniformUMin $$T result Subgroup Reduce $expr};
+        else
+            return spirv_asm {OpGroupNonUniformSMin $$T result Subgroup Reduce $expr};
+    default:
+        return WaveMaskMin(WaveGetActiveMask(), expr);
+    }
 }
 
 __generic<T : __BuiltinArithmeticType, let N : int>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupMin($0)")
-__target_intrinsic(hlsl)
+__spirv_capability(GroupNonUniformArithmetic)
 vector<T, N> WaveActiveMin(vector<T, N> expr)
 {
-    return WaveMaskMin(WaveGetActiveMask(), expr);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupMin($0)";
+    case hlsl: __intrinsic_asm "WaveActiveMinWaveActiveMin";
+    case spirv:
+        if (__isFloat<T>())
+            return spirv_asm {OpGroupNonUniformFMin $$vector<T, N> result Subgroup Reduce $expr};
+        else if (__isUnsignedInt<T>())
+            return spirv_asm {OpGroupNonUniformUMin $$vector<T, N> result Subgroup Reduce $expr};
+        else
+            return spirv_asm {OpGroupNonUniformSMin $$vector<T, N> result Subgroup Reduce $expr};
+    default:
+        return WaveMaskMin(WaveGetActiveMask(), expr);
+    }
 }
 
 __generic<T : __BuiltinArithmeticType, let N : int, let M : int>
@@ -5546,10 +5684,9 @@ bool WaveActiveAllEqual(T value)
     case hlsl:
         __intrinsic_asm "WaveActiveAllEqual";
     case spirv:
-        let _scope = 3u; // subgroup
         return spirv_asm
         {
-            OpGroupNonUniformAllEqual $$bool result $_scope $value
+            OpGroupNonUniformAllEqual $$bool result Subgroup $value
         };
     default:
         return WaveMaskAllEqual(WaveGetActiveMask(), value);
@@ -5569,10 +5706,9 @@ bool WaveActiveAllEqual(vector<T,N> value)
     case hlsl:
         __intrinsic_asm "WaveActiveAllEqual";
     case spirv:
-        let _scope = 3u; // subgroup
         return spirv_asm
         {
-            OpGroupNonUniformAllEqual $$bool result $_scope $value
+            OpGroupNonUniformAllEqual $$bool result Subgroup $value
         };
     default:
         return WaveMaskAllEqual(WaveGetActiveMask(), value);
@@ -5598,10 +5734,9 @@ bool WaveActiveAllTrue(bool condition)
     case hlsl:
         __intrinsic_asm "WaveActiveAllTrue($0)";
     case spirv:
-        let _scope = 3u; // subgroup
         return spirv_asm
         {
-            OpGroupNonUniformAll $$bool result $_scope $condition
+            OpGroupNonUniformAll $$bool result Subgroup $condition
         };
     default:
         return WaveMaskAllTrue(WaveGetActiveMask(), condition);
@@ -5611,8 +5746,6 @@ bool WaveActiveAllTrue(bool condition)
 __glsl_extension(GL_KHR_shader_subgroup_vote)
 __spirv_version(1.3)
 __spirv_capability(GroupNonUniformVote)
-__target_intrinsic(glsl, "subgroupAny($0)")
-__target_intrinsic(hlsl)
 bool WaveActiveAnyTrue(bool condition)
 {
     __target_switch
@@ -5622,10 +5755,9 @@ bool WaveActiveAnyTrue(bool condition)
     case hlsl:
         __intrinsic_asm "WaveActiveAnyTrue($0)";
     case spirv:
-        let _scope = 3u; // subgroup
         return spirv_asm
         {
-            OpGroupNonUniformAny $$bool result $_scope $condition
+            OpGroupNonUniformAny $$bool result Subgroup $condition
         };
     default:
         return WaveMaskAnyTrue(WaveGetActiveMask(), condition);
@@ -5662,14 +5794,25 @@ uint WaveActiveCountBits(bool value)
 
 __glsl_extension(GL_KHR_shader_subgroup_basic)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "(gl_SubgroupSize)")
-__target_intrinsic(cuda, "(warpSize)")
-uint WaveGetLaneCount();
+__spirv_capability(GroupNonUniform)
+uint WaveGetLaneCount()
+{
+    __target_switch
+    {
+    case glsl: __intrinsic_asm  "(gl_SubgroupSize)";
+    case cuda: __intrinsic_asm  "(warpSize)";
+    case hlsl: __intrinsic_asm  "WaveGetLaneCount()";
+    case spirv:
+        return spirv_asm
+        {
+            result:$$uint = OpLoad builtin(SubgroupSize:uint)
+        };
+    }
+}
 
 __glsl_extension(GL_KHR_shader_subgroup_basic)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "(gl_SubgroupInvocationID)")
-__target_intrinsic(cuda, "_getLaneId()")
+__spirv_capability(GroupNonUniform)
 uint WaveGetLaneIndex()
 {
     __target_switch
@@ -5677,14 +5820,11 @@ uint WaveGetLaneIndex()
     case glsl: __intrinsic_asm  "(gl_SubgroupInvocationID)";
     case cuda: __intrinsic_asm  "_getLaneId()";
     case hlsl: __intrinsic_asm  "WaveGetLaneIndex()";
-    /*
     case spirv:
-        let _scope = 3u; // subgroup
         return spirv_asm
         {
-            OpSubgroupLocalInvocationId $$uint result $_scope
+            result:$$uint = OpLoad builtin(SubgroupLocalInvocationId:uint)
         };
-    */
     }
 }
 
@@ -5700,10 +5840,9 @@ bool WaveIsFirstLane()
     case hlsl:
         __intrinsic_asm "WaveIsFirstLane()";
     case spirv:
-        let _scope = 3u; // subgroup
         return spirv_asm
         {
-            OpGroupNonUniformElect $$bool result $_scope
+            OpGroupNonUniformElect $$bool result Subgroup
         };
     default:
         return WaveMaskIsFirstLane(WaveGetActiveMask());
@@ -5718,10 +5857,9 @@ uint _WaveCountBits(uint4 value)
     __target_switch
     {
     case spirv:
-        let _scope = 3u; // Subgroup
         return spirv_asm
         {
-            OpGroupNonUniformBallotBitCount $$uint result $_scope 0 $value
+            OpGroupNonUniformBallotBitCount $$uint result Subgroup Reduce $value
         };
     default:
         // Assume since WaveGetLaneCount should be known at compile time, the branches will hopefully boil away
@@ -5743,22 +5881,62 @@ uint _WaveCountBits(uint4 value)
 __generic<T : __BuiltinArithmeticType>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupExclusiveMul($0)")
-__target_intrinsic(hlsl)
+__spirv_capability(GroupNonUniformArithmetic)
 T WavePrefixProduct(T expr)
 {
-    return WaveMaskPrefixProduct(WaveGetActiveMask(), expr);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupExclusiveMul($0)";
+    case hlsl: __intrinsic_asm "WavePrefixProduct";
+    case spirv:
+        if (__isFloat<T>())
+            return spirv_asm {OpGroupNonUniformFMul $$T result Subgroup ExclusiveScan $expr};
+        else if (__isSignedInt<T>())
+        {
+            return spirv_asm
+            {
+                // TODO: use the correct integer width
+                OpBitcast $$uint %uvalue $expr;
+                OpGroupNonUniformIMul $$uint %mulResult Subgroup ExclusiveScan %uvalue;
+                OpBitcast $$T result %mulResult
+            };
+        }
+        else if (__isUnsignedInt<T>())
+            return spirv_asm {OpGroupNonUniformIMul $$T result Subgroup ExclusiveScan $expr};
+    default:
+        return WaveMaskPrefixProduct(WaveGetActiveMask(), expr);
+    }
 }
 
 
 __generic<T : __BuiltinArithmeticType, let N : int>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupExclusiveMul($0)")
-__target_intrinsic(hlsl)
+__spirv_capability(GroupNonUniformArithmetic)
 vector<T,N> WavePrefixProduct(vector<T,N> expr)
 {
-    return WaveMaskPrefixProduct(WaveGetActiveMask(), expr);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupExclusiveMul($0)";
+    case hlsl: __intrinsic_asm "WavePrefixProduct";
+    case spirv:
+        if (__isFloat<T>())
+            return spirv_asm {OpGroupNonUniformFMul $$vector<T,N> result Subgroup ExclusiveScan $expr};
+        else if (__isSignedInt<T>())
+        {
+            return spirv_asm
+            {
+                // TODO: use the correct integer width
+                OpBitcast $$vector<uint,N> %uvalue $expr;
+                OpGroupNonUniformIMul $$vector<uint,N> %mulResult Subgroup ExclusiveScan %uvalue;
+                OpBitcast $$vector<T,N> result %mulResult
+            };
+        }
+        else if (__isUnsignedInt<T>())
+            return spirv_asm {OpGroupNonUniformIMul $$vector<T,N> result Subgroup ExclusiveScan $expr};
+    default:
+        return WaveMaskPrefixProduct(WaveGetActiveMask(), expr);
+    }
 }
 
 __generic<T : __BuiltinArithmeticType, let N : int, let M : int>
@@ -5771,21 +5949,61 @@ matrix<T, N, M> WavePrefixProduct(matrix<T, N, M> expr)
 __generic<T : __BuiltinArithmeticType>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupExclusiveAdd($0)")
-__target_intrinsic(hlsl)
+__spirv_capability(GroupNonUniformArithmetic)
 T WavePrefixSum(T expr)
 {
-    return WaveMaskPrefixSum(WaveGetActiveMask(), expr);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupExclusiveAdd($0)";
+    case hlsl: __intrinsic_asm "WavePrefixSum";
+    case spirv:
+        if (__isFloat<T>())
+            return spirv_asm {OpGroupNonUniformFAdd $$T result Subgroup ExclusiveScan $expr};
+        else if (__isSignedInt<T>())
+        {
+            return spirv_asm
+            {
+                // TODO: use the correct integer width
+                %uvalue:$$uint = OpBitcast $expr;
+                %mulResult:$$uint = OpGroupNonUniformIAdd Subgroup ExclusiveScan %uvalue;
+                result:$$T = OpBitcast %mulResult
+            };
+        }
+        else if (__isUnsignedInt<T>())
+            return spirv_asm {OpGroupNonUniformIAdd $$T result Subgroup ExclusiveScan $expr};
+    default:
+        return WaveMaskPrefixSum(WaveGetActiveMask(), expr);
+    }
 }
 
 __generic<T : __BuiltinArithmeticType, let N : int>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupExclusiveAdd($0)")
-__target_intrinsic(hlsl)
+__spirv_capability(GroupNonUniformArithmetic)
 vector<T,N> WavePrefixSum(vector<T,N> expr)
 {
-    return WaveMaskPrefixSum(WaveGetActiveMask(), expr);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupExclusiveAdd($0)";
+    case hlsl: __intrinsic_asm "WavePrefixSum";
+    case spirv:
+        if (__isFloat<T>())
+            return spirv_asm {OpGroupNonUniformFAdd $$vector<T,N> result Subgroup ExclusiveScan $expr};
+        else if (__isSignedInt<T>())
+        {
+            return spirv_asm
+            {
+                // TODO: use the correct integer width
+                %uvalue:$$vector<uint,N> = OpBitcast $expr;
+                %mulResult:$$vector<uint,N> = OpGroupNonUniformIAdd Subgroup ExclusiveScan %uvalue;
+                result:$$vector<T,N> = OpBitcast %mulResult
+            };
+        }
+        else if (__isUnsignedInt<T>())
+            return spirv_asm {OpGroupNonUniformIAdd $$vector<T,N> result Subgroup ExclusiveScan $expr};
+    default:
+        return WaveMaskPrefixSum(WaveGetActiveMask(), expr);
+    }
 }
 
 __generic<T : __BuiltinArithmeticType, let N : int, let M : int>
@@ -5798,21 +6016,35 @@ matrix<T,N,M> WavePrefixSum(matrix<T,N,M> expr)
 __generic<T : __BuiltinType>
 __glsl_extension(GL_KHR_shader_subgroup_ballot)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupBroadcastFirst($0)")
-__target_intrinsic(hlsl)
+__spirv_capability(GroupNonUniformBallot)
 T WaveReadLaneFirst(T expr)
 {
-    return WaveMaskReadLaneFirst(WaveGetActiveMask(), expr);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupBroadcastFirst($0)";
+    case hlsl: __intrinsic_asm "WaveReadLaneFirst";
+    case spirv:
+        return spirv_asm {OpGroupNonUniformBroadcastFirst $$T result Subgroup $expr};
+    default:
+        return WaveMaskReadLaneFirst(WaveGetActiveMask(), expr);
+    }
 }
 
 __generic<T : __BuiltinType, let N : int>
 __glsl_extension(GL_KHR_shader_subgroup_ballot)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupBroadcastFirst($0)")
-__target_intrinsic(hlsl)
+__spirv_capability(GroupNonUniformBallot)
 vector<T,N> WaveReadLaneFirst(vector<T,N> expr)
 {
-    return WaveMaskReadLaneFirst(WaveGetActiveMask(), expr);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupBroadcastFirst($0)";
+    case hlsl: __intrinsic_asm "WaveReadLaneFirst";
+    case spirv:
+        return spirv_asm {OpGroupNonUniformBroadcastFirst $$vector<T,N> result Subgroup $expr};
+    default:
+        return WaveMaskReadLaneFirst(WaveGetActiveMask(), expr);
+    }
 }
 
 __generic<T : __BuiltinType, let N : int, let M : int>
@@ -5831,21 +6063,35 @@ matrix<T,N,M> WaveReadLaneFirst(matrix<T,N,M> expr)
 __generic<T : __BuiltinType>
 __glsl_extension(GL_KHR_shader_subgroup_ballot)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupBroadcast($0, $1)")
-__target_intrinsic(hlsl, "WaveReadLaneAt")
+__spirv_capability(GroupNonUniformBallot)
 T WaveBroadcastLaneAt(T value, constexpr int lane)
 {
-    return WaveMaskBroadcastLaneAt(WaveGetActiveMask(), value, lane);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupBroadcast($0, $1)";
+    case hlsl: __intrinsic_asm "WaveReadLaneAt";
+    case spirv:
+        return spirv_asm {OpGroupNonUniformBroadcast $$T result Subgroup $value $lane};
+    default:
+        return WaveMaskBroadcastLaneAt(WaveGetActiveMask(), value, lane);
+    }
 }
 
 __generic<T : __BuiltinType, let N : int>
 __glsl_extension(GL_KHR_shader_subgroup_ballot)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupBroadcast($0, $1)")
-__target_intrinsic(hlsl, "WaveReadLaneAt")
+__spirv_capability(GroupNonUniformBallot)
 vector<T,N> WaveBroadcastLaneAt(vector<T,N> value, constexpr int lane)
 {
-    return WaveMaskBroadcastLaneAt(WaveGetActiveMask(), value, lane);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupBroadcast($0, $1)";
+    case hlsl: __intrinsic_asm "WaveReadLaneAt";
+    case spirv:
+        return spirv_asm {OpGroupNonUniformBroadcast $$vector<T,N> result Subgroup $value $lane};
+    default:
+        return WaveMaskBroadcastLaneAt(WaveGetActiveMask(), value, lane);
+    }
 }
 
 __generic<T : __BuiltinType, let N : int, let M : int>
@@ -5861,21 +6107,35 @@ matrix<T, N, M> WaveBroadcastLaneAt(matrix<T, N, M> value, constexpr int lane)
 __generic<T : __BuiltinType>
 __glsl_extension(GL_KHR_shader_subgroup_shuffle)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupShuffle($0, $1)")
-__target_intrinsic(hlsl)
+__spirv_capability(GroupNonUniformShuffle)
 T WaveReadLaneAt(T value, int lane)
 {
-    return WaveMaskReadLaneAt(WaveGetActiveMask(), value, lane);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupShuffle($0, $1)";
+    case hlsl: __intrinsic_asm "WaveReadLaneAt";
+    case spirv:
+        return spirv_asm {OpGroupNonUniformShuffle $$T result Subgroup $value $lane};
+    default:
+        return WaveMaskReadLaneAt(WaveGetActiveMask(), value, lane);
+    }
 }
 
 __generic<T : __BuiltinType, let N : int>
 __spirv_version(1.3)
 __glsl_extension(GL_KHR_shader_subgroup_shuffle)
-__target_intrinsic(glsl, "subgroupShuffle($0, $1)")
-__target_intrinsic(hlsl)
+__spirv_capability(GroupNonUniformShuffle)
 vector<T,N> WaveReadLaneAt(vector<T,N> value, int lane)
 {
-    return WaveMaskReadLaneAt(WaveGetActiveMask(), value, lane);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupShuffle($0, $1)";
+    case hlsl: __intrinsic_asm "WaveReadLaneAt";
+    case spirv:
+        return spirv_asm {OpGroupNonUniformShuffle $$vector<T,N> result Subgroup $value $lane};
+    default:
+        return WaveMaskReadLaneAt(WaveGetActiveMask(), value, lane);
+    }
 }
 
 __generic<T : __BuiltinType, let N : int, let M : int>
@@ -5892,21 +6152,35 @@ matrix<T, N, M> WaveReadLaneAt(matrix<T, N, M> value, int lane)
 __generic<T : __BuiltinType>
 __glsl_extension(GL_KHR_shader_subgroup_shuffle)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupShuffle($0, $1)")
-__target_intrinsic(hlsl, "WaveReadLaneAt")
+__spirv_capability(GroupNonUniformShuffle)
 T WaveShuffle(T value, int lane)
 {
-    return WaveMaskShuffle(WaveGetActiveMask(), value, lane);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupShuffle($0, $1)";
+    case hlsl: __intrinsic_asm "WaveReadLaneAt";
+    case spirv:
+        return spirv_asm {OpGroupNonUniformShuffle $$T result Subgroup $value $lane};
+    default:
+        return WaveMaskShuffle(WaveGetActiveMask(), value, lane);
+    }
 }
 
 __generic<T : __BuiltinType, let N : int>
 __glsl_extension(GL_KHR_shader_subgroup_shuffle)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupShuffle($0, $1)")
-__target_intrinsic(hlsl, "WaveReadLaneAt")
+__spirv_capability(GroupNonUniformShuffle)
 vector<T,N> WaveShuffle(vector<T,N> value, int lane)
 {
-    return WaveMaskShuffle(WaveGetActiveMask(), value, lane);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupShuffle($0, $1)";
+    case hlsl: __intrinsic_asm "WaveReadLaneAt";
+    case spirv:
+        return spirv_asm {OpGroupNonUniformShuffle $$vector<T,N> result Subgroup $value $lane};
+    default:
+        return WaveMaskShuffle(WaveGetActiveMask(), value, lane);
+    }
 }
 
 __generic<T : __BuiltinType, let N : int, let M : int>
@@ -5918,25 +6192,48 @@ matrix<T, N, M> WaveShuffle(matrix<T, N, M> value, int lane)
 
 __glsl_extension(GL_KHR_shader_subgroup_ballot)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupBallotExclusiveBitCount(subgroupBallot($0))")
-__target_intrinsic(hlsl)
+__spirv_capability(GroupNonUniformBallot)
 uint WavePrefixCountBits(bool value)
 {
-    return WaveMaskPrefixCountBits(WaveGetActiveMask(), value);
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupBallotExclusiveBitCount(subgroupBallot($0))";
+    case hlsl: __intrinsic_asm "WavePrefixCountBits($0)";
+    case spirv:
+        return spirv_asm
+        {
+            %mask:$$uint4 = OpGroupNonUniformBallot Subgroup $value;
+            OpGroupNonUniformBallotBitCount $$uint result Subgroup 2 %mask
+        };
+    default:
+        return WaveMaskPrefixCountBits(WaveGetActiveMask(), value);
+    }
 }
 
 __glsl_extension(GL_KHR_shader_subgroup_ballot)
 __spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupBallot(true)")
-__target_intrinsic(cuda, "make_uint4(__activemask(), 0, 0, 0)")
-__target_intrinsic(hlsl, "WaveActiveBallot(true)")
-uint4 WaveGetConvergedMulti();
+__spirv_capability(GroupNonUniformBallot)
+uint4 WaveGetConvergedMulti()
+{
+    __target_switch
+    {
+    case glsl: __intrinsic_asm "subgroupBallot(true)";
+    case hlsl: __intrinsic_asm "WaveActiveBallot(true)";
+    case cuda: __intrinsic_asm "make_uint4(__activemask(), 0, 0, 0)";
+    case spirv:
+        let _true = true;
+        return spirv_asm
+        {
+            OpGroupNonUniformBallot $$uint4 result Subgroup $_true
+        };
+    }
+}
 
-__glsl_extension(GL_KHR_shader_subgroup_ballot)
-__spirv_version(1.3)
-__target_intrinsic(glsl, "subgroupBallot(true)")
-__target_intrinsic(hlsl, "WaveActiveBallot(true)")
-uint4 WaveGetActiveMulti();
+[ForceInline]
+uint4 WaveGetActiveMulti()
+{
+    return WaveGetConvergedMulti();
+}
 
 // Shader model 6.5 stuff
 // https://github.com/microsoft/DirectX-Specs/blob/master/d3d/HLSL_ShaderModel6_5.md
diff --git a/source/slang/slang-ast-expr.h b/source/slang/slang-ast-expr.h
index 7b09c3334..5603ef2a5 100644
--- a/source/slang/slang-ast-expr.h
+++ b/source/slang/slang-ast-expr.h
@@ -647,6 +647,8 @@ public:
         SlangValue,
         SlangValueAddr,
         SlangType,
+        BuiltinVar,
+        GLSL450Set,
     };
 
     // The flavour and token describes how this was parsed
@@ -672,8 +674,8 @@ public:
     // it as an id created with OpConstant
     bool wrapInId = false;
 
-    // Once we've checked things, the SlangType flavour operands will have this
-    // type populated.
+    // Once we've checked things, the SlangType and BuiltinVar flavour operands
+    // will have this type populated.
     TypeExp type = TypeExp();
 };
 
diff --git a/source/slang/slang-check-expr.cpp b/source/slang/slang-check-expr.cpp
index 43db85e3b..e6d1049a0 100644
--- a/source/slang/slang-check-expr.cpp
+++ b/source/slang/slang-check-expr.cpp
@@ -14,7 +14,7 @@
 #include "slang-ast-natural-layout.h"
 
 #include "slang-lookup.h"
-
+#include "slang-lookup-spirv.h"
 #include "slang-ast-print.h"
 
 namespace Slang
@@ -3969,7 +3969,17 @@ namespace Slang
                         // Then see if it's an opcode (for OpSpecialize)
                         if(!enumValue)
                             enumValue = spirvInfo->opcodes.lookup(operand.token.getContent());
-
+                        if (inst.opcode.knownValue == SpvOpExtInst)
+                        {
+                            if (!enumValue)
+                            {
+                                GLSLstd450 val;
+                                if (lookupGLSLstd450(operand.token.getContent(), val))
+                                {
+                                    enumValue = (SpvWord)val;
+                                }
+                            }
+                        }
                         if(!enumValue)
                         {
                             failed = true;
@@ -3980,6 +3990,19 @@ namespace Slang
                         operand.knownValue = *enumValue;
                         operand.wrapInId = needsIdWrapper;
                     }
+                    else if (operand.flavor == SPIRVAsmOperand::BuiltinVar)
+                    {
+                        operand.type = CheckProperType(operand.type);
+                        auto builtinVarKind = spirvInfo->allEnums.lookup(
+                            SPIRVCoreGrammarInfo::QualifiedEnumName{spirvInfo->operandKinds.lookup(UnownedStringSlice("BuiltIn")).value(), operand.token.getContent()});
+                        if (!builtinVarKind)
+                        {
+                            failed = true;
+                            getSink()->diagnose(operand.token, Diagnostics::spirvUnableToResolveName, operand.token.getContent());
+                            return;
+                        }
+                        operand.knownValue = builtinVarKind.value();
+                    }
                     if(operand.bitwiseOrWith.getCount()
                         && operand.flavor != SPIRVAsmOperand::Literal
                         && operand.flavor != SPIRVAsmOperand::NamedValue)
diff --git a/source/slang/slang-emit-spirv.cpp b/source/slang/slang-emit-spirv.cpp
index 462f408ab..a203f4a80 100644
--- a/source/slang/slang-emit-spirv.cpp
+++ b/source/slang/slang-emit-spirv.cpp
@@ -1454,6 +1454,8 @@ struct SPIRVEmitContext
             return emitGetStringHash(inst);
         default:
             {
+                if (as<IRSPIRVAsmOperand>(inst))
+                    return nullptr;
                 String e = "Unhandled global inst in spirv-emit:\n"
                     + dumpIRToString(inst, {IRDumpOptions::Mode::Detailed, 0});
                 SLANG_UNIMPLEMENTED_X(e.begin());
@@ -1854,6 +1856,8 @@ struct SPIRVEmitContext
         {
         default:
             {
+                if (as<IRSPIRVAsmOperand>(inst))
+                    return nullptr;
                 String e = "Unhandled local inst in spirv-emit:\n"
                     + dumpIRToString(inst, {IRDumpOptions::Mode::Detailed, 0});
                 SLANG_UNIMPLEMENTED_X(e.getBuffer());
@@ -2175,7 +2179,8 @@ struct SPIRVEmitContext
                 auto entryPointDecor = cast<IREntryPointDecoration>(decoration);
                 auto spvStage = mapStageToExecutionModel(entryPointDecor->getProfile().getStage());
                 auto name = entryPointDecor->getName()->getStringSlice();
-                List<IRInst*> params;
+                List<SpvInst*> params;
+                HashSet<SpvInst*> paramsSet;
                 // `interface` part: reference all global variables that are used by this entrypoint.
                 // TODO: we may want to perform more accurate tracking.
                 for (auto globalInst : m_irModule->getModuleInst()->getChildren())
@@ -2184,10 +2189,27 @@ struct SPIRVEmitContext
                     {
                     case kIROp_GlobalVar:
                     case kIROp_GlobalParam:
-                        params.add(globalInst);
+                    {
+                        SpvInst* spvGlobalInst;
+                        if (m_mapIRInstToSpvInst.tryGetValue(globalInst, spvGlobalInst))
+                        {
+                            paramsSet.add(spvGlobalInst);
+                            params.add(spvGlobalInst);
+                        }
+                        break;
+                    }
+                    default:
                         break;
                     }
                 }
+
+                // Add remaining builtin variables that does not have a corresponding IR global var/param.
+                // These variables could be added from SPIRV ASM blocks.
+                for (auto builtinVar : m_builtinGlobalVars)
+                {
+                    if (paramsSet.add(builtinVar.second))
+                        params.add(builtinVar.second);
+                }
                 emitOpEntryPoint(
                     section,
                     decoration,
@@ -3838,6 +3860,18 @@ struct SPIRVEmitContext
                 }
             };
 
+            switch (opcode)
+            {
+            case SpvOpCapability:
+                requireSPIRVCapability((SpvCapability)getIntVal(spvInst->getOperand(1)->getOperand(0)));
+                continue;
+            case SpvOpExtension:
+                ensureExtensionDeclaration(as<IRStringLit>(spvInst->getOperand(1)->getOperand(0))->getStringSlice());
+                continue;
+            default:
+                break;
+            }
+
             last = emitInstCustomOperandFunc(
                 parentForOpCode(opcode, parent),
                 // We want the "result instruction" to refer to the top level
@@ -3922,6 +3956,20 @@ struct SPIRVEmitContext
                             emitOperand(id);
                             break;
                         }
+                        case kIROp_SPIRVAsmOperandBuiltinVar:
+                        {
+                            const auto kind = (SpvBuiltIn)(getIntVal(operand->getOperand(0)));
+                            IRBuilder builder(operand);
+                            builder.setInsertBefore(operand);
+                            auto varInst = getBuiltinGlobalVar(builder.getPtrType(kIROp_PtrType, operand->getDataType(), SpvStorageClassInput), kind);
+                            emitOperand(varInst);
+                            break;
+                        }
+                        case kIROp_SPIRVAsmOperandGLSL450Set:
+                        {
+                            emitOperand(getGLSL450ExtInst());
+                            break;
+                        }
                         default:
                             SLANG_UNREACHABLE("Unhandled case in emitSPIRVAsm");
                         }
diff --git a/source/slang/slang-ir-eliminate-phis.cpp b/source/slang/slang-ir-eliminate-phis.cpp
index c628b937e..a17759fe6 100644
--- a/source/slang/slang-ir-eliminate-phis.cpp
+++ b/source/slang/slang-ir-eliminate-phis.cpp
@@ -531,7 +531,7 @@ struct PhiEliminationContext
                 auto user = use->getUser();
                 m_builder.setInsertBefore(user);
                 auto newVal = m_builder.emitLoad(temp);
-                use->set(newVal);
+                m_builder.replaceOperand(use, newVal);
             }
 
             // Once we've replaced all its uses, there is no need
diff --git a/source/slang/slang-ir-inst-defs.h b/source/slang/slang-ir-inst-defs.h
index e09abcf75..b248012a0 100644
--- a/source/slang/slang-ir-inst-defs.h
+++ b/source/slang/slang-ir-inst-defs.h
@@ -1074,19 +1074,23 @@ INST(SPIRVAsmInst, SPIRVAsmInst, 1, 0)
     // and a reference to a literal integer instruction
     //
     // A literal string or 32-bit integer to be passed as operands
-    INST(SPIRVAsmOperandLiteral, SPIRVAsmOperandLiteral, 1, 0)
+    INST(SPIRVAsmOperandLiteral, SPIRVAsmOperandLiteral, 1, HOISTABLE)
     // A reference to a slang IRInst, either a value or a type
-    INST(SPIRVAsmOperandInst, SPIRVAsmOperandInst, 1, 0)
+    INST(SPIRVAsmOperandInst, SPIRVAsmOperandInst, 1, HOISTABLE)
     // A named enumerator, the value is stored as a constant operand
     // It may have a second operand, which if present is a type with which to
     // construct a constant id to pass, instead of a literal constant
-    INST(SPIRVAsmOperandEnum, SPIRVAsmOperandEnum, 1, 0)
+    INST(SPIRVAsmOperandEnum, SPIRVAsmOperandEnum, 1, HOISTABLE)
+    // A reference to a builtin variable.
+    INST(SPIRVAsmOperandBuiltinVar, SPIRVAsmOperandBuiltinVar, 1, HOISTABLE)
+    // A reference to the glsl450 instruction set.
+    INST(SPIRVAsmOperandGLSL450Set, SPIRVAsmOperandGLSL450Set, 0, HOISTABLE)
     // A string which is given a unique ID in the backend, used to refer to
     // results of other instrucions in the same asm block
-    INST(SPIRVAsmOperandId, SPIRVAsmOperandId, 1, 0)
+    INST(SPIRVAsmOperandId, SPIRVAsmOperandId, 1, HOISTABLE)
     // A special instruction which marks the place to insert the generated
     // result operand
-    INST(SPIRVAsmOperandResult, SPIRVAsmOperandResult, 0, 0)
+    INST(SPIRVAsmOperandResult, SPIRVAsmOperandResult, 0, HOISTABLE)
 INST_RANGE(SPIRVAsmOperand, SPIRVAsmOperandLiteral, SPIRVAsmOperandResult)
 
 
diff --git a/source/slang/slang-ir-insts.h b/source/slang/slang-ir-insts.h
index 1a0a8ed66..9d2f44355 100644
--- a/source/slang/slang-ir-insts.h
+++ b/source/slang/slang-ir-insts.h
@@ -3937,6 +3937,8 @@ public:
     IRSPIRVAsmOperand* emitSPIRVAsmOperandResult();
     IRSPIRVAsmOperand* emitSPIRVAsmOperandEnum(IRInst* inst);
     IRSPIRVAsmOperand* emitSPIRVAsmOperandEnum(IRInst* inst, IRType* constantType);
+    IRSPIRVAsmOperand* emitSPIRVAsmOperandBuiltinVar(IRInst* type, IRInst* builtinKind);
+    IRSPIRVAsmOperand* emitSPIRVAsmOperandGLSL450Set();
     IRSPIRVAsmInst* emitSPIRVAsmInst(IRInst* opcode, List<IRInst*> operands);
     IRSPIRVAsm* emitSPIRVAsm(IRType* type);
     IRInst* emitGenericAsm(UnownedStringSlice asmText);
diff --git a/source/slang/slang-ir.cpp b/source/slang/slang-ir.cpp
index cd49c6df5..d3cfea6e9 100644
--- a/source/slang/slang-ir.cpp
+++ b/source/slang/slang-ir.cpp
@@ -5747,6 +5747,31 @@ namespace Slang
         return i;
     }
 
+    IRSPIRVAsmOperand* IRBuilder::emitSPIRVAsmOperandBuiltinVar(IRInst* type, IRInst* builtinKind)
+    {
+        SLANG_ASSERT(as<IRSPIRVAsm>(m_insertLoc.getParent()));
+        const auto i = createInst<IRSPIRVAsmOperand>(
+            this,
+            kIROp_SPIRVAsmOperandBuiltinVar,
+            (IRType*)type,
+            builtinKind
+        );
+        addInst(i);
+        return i;
+    }
+
+    IRSPIRVAsmOperand* IRBuilder::emitSPIRVAsmOperandGLSL450Set()
+    {
+        SLANG_ASSERT(as<IRSPIRVAsm>(m_insertLoc.getParent()));
+        const auto i = createInst<IRSPIRVAsmOperand>(
+            this,
+            kIROp_SPIRVAsmOperandGLSL450Set,
+            getVoidType()
+        );
+        addInst(i);
+        return i;
+    }
+
     IRSPIRVAsmInst* IRBuilder::emitSPIRVAsmInst(IRInst* opcode, List<IRInst*> operands)
     {
         SLANG_ASSERT(as<IRSPIRVAsm>(m_insertLoc.getParent()));
@@ -7482,6 +7507,9 @@ namespace Slang
         if(as<IRAttr>(this))
             return false;
 
+        if (as<IRSPIRVAsmOperand>(this))
+            return false;
+
         switch(getOp())
         {
         // By default, assume that we might have side effects,
diff --git a/source/slang/slang-lower-to-ir.cpp b/source/slang/slang-lower-to-ir.cpp
index 41347ddef..596d09184 100644
--- a/source/slang/slang-lower-to-ir.cpp
+++ b/source/slang/slang-lower-to-ir.cpp
@@ -3295,6 +3295,17 @@ struct ExprLoweringVisitorBase : ExprVisitor<Derived, LoweredValInfo>
                     else
                         return builder->emitSPIRVAsmOperandEnum(i);
                 }
+            case SPIRVAsmOperand::BuiltinVar:
+                {
+                    const auto kind = operand.knownValue;
+                    auto kindInst = builder->getIntValue(builder->getIntType(), kind);
+                    const auto type = lowerType(context, operand.type.type);
+                    return builder->emitSPIRVAsmOperandBuiltinVar(type, kindInst);
+                }
+            case SPIRVAsmOperand::GLSL450Set:
+                {
+                    return builder->emitSPIRVAsmOperandGLSL450Set();
+                }
             case SPIRVAsmOperand::SlangValue:
                 {
                     IRInst* i;
diff --git a/source/slang/slang-parser.cpp b/source/slang/slang-parser.cpp
index 7e6d7d308..3a1b627bd 100644
--- a/source/slang/slang-parser.cpp
+++ b/source/slang/slang-parser.cpp
@@ -6277,6 +6277,21 @@ namespace Slang
         {
             return SPIRVAsmOperand{SPIRVAsmOperand::ResultMarker, parser->ReadToken()};
         }
+        else if (AdvanceIf(parser, "builtin"))
+        {
+            // reference to a builtin var.
+            parser->ReadToken(TokenType::LParent);
+            auto operand = SPIRVAsmOperand{ SPIRVAsmOperand::BuiltinVar, parser->ReadToken() };
+            parser->ReadToken(TokenType::Colon);
+            AdvanceIf(parser, TokenType::DollarDollar);
+            operand.type = parser->ParseTypeExp();
+            parser->ReadToken(TokenType::RParent);
+            return operand;
+        }
+        else if (parser->LookAheadToken("glsl450"))
+        {
+            return SPIRVAsmOperand{ SPIRVAsmOperand::GLSL450Set, parser->ReadToken() };
+        }
 
         // A regular identifier
         else if(parser->LookAheadToken(TokenType::Identifier))