17 files changed, 483 insertions, 1032 deletions
diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang
index 3d8797fa1..de81cd9f8 100644
--- a/source/slang/hlsl.meta.slang
+++ b/source/slang/hlsl.meta.slang
@@ -2139,6 +2139,7 @@ __generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M> QuadReadAcr
 
 __generic<T : __BuiltinIntegerType>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupAnd($0)")
 __target_intrinsic(cuda, "_waveAnd(__activemask(), $0)")
 T WaveActiveBitAnd(T expr);
@@ -2147,6 +2148,7 @@ __generic<T : __BuiltinIntegerType, let N : int, let M : int> matrix<T,N,M> Wave
 
 __generic<T : __BuiltinIntegerType>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupOr($0)")
 __target_intrinsic(cuda, "_waveOr(__activemask(), $0)")
 T WaveActiveBitOr(T expr);
@@ -2155,6 +2157,7 @@ __generic<T : __BuiltinIntegerType, let N : int, let M : int> matrix<T,N,M> Wave
 
 __generic<T : __BuiltinIntegerType>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupXor($0)")
 __target_intrinsic(cuda, "_waveXor(__activemask(), $0)")
 T WaveActiveBitXor(T expr);
@@ -2163,6 +2166,7 @@ __generic<T : __BuiltinIntegerType, let N : int, let M : int> matrix<T,N,M> Wave
 
 __generic<T : __BuiltinArithmeticType>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupMax($0)")
 __target_intrinsic(cuda, "_waveMax(__activemask(), $0)")
 T WaveActiveMax(T expr);
@@ -2171,6 +2175,7 @@ __generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> W
 
 __generic<T : __BuiltinArithmeticType>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupMin($0)")
 __target_intrinsic(cuda, "_waveMin(__activemask(), $0)")
 T WaveActiveMin(T expr);
@@ -2179,6 +2184,7 @@ __generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> W
 
 __generic<T : __BuiltinArithmeticType>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupMul($0)")
 __target_intrinsic(cuda, "_waveProduct(__activemask(), $0)")
 T WaveActiveProduct(T expr);
@@ -2187,6 +2193,7 @@ __generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> W
 
 __generic<T : __BuiltinArithmeticType>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupAdd($0)")
 __target_intrinsic(cuda, "_waveSum(__activemask(), $0)")
 T WaveActiveSum(T expr);
@@ -2195,6 +2202,7 @@ __generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> W
 
 __generic<T : __BuiltinType>
 __glsl_extension(GL_KHR_shader_subgroup_vote)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupAllEqual($0)")
 __target_intrinsic(cuda, "_waveAllEqual(__activemask(), $0)")
 bool WaveActiveAllEqual(T value);
@@ -2211,16 +2219,19 @@ __generic<T : __BuiltinType, let N : int, let M : int> uint4 WaveMatch(matrix<T,
 // seems to be appropriate.
 
 __glsl_extension(GL_KHR_shader_subgroup_vote)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupAll($0)") 
 __target_intrinsic(cuda, "(__all_sync(__activemask(), $0) != 0)") 
 bool WaveActiveAllTrue(bool condition);
 
 __glsl_extension(GL_KHR_shader_subgroup_vote)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupAny($0)") 
 __target_intrinsic(cuda, "(__any_sync(__activemask(), $0) != 0)")
 bool WaveActiveAnyTrue(bool condition);
 
 __glsl_extension(GL_KHR_shader_subgroup_ballot)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupBallot($0)")
 __target_intrinsic(cuda, "make_uint4(__ballot_sync(__activemask(), $0), 0, 0, 0)")
 uint4 WaveActiveBallot(bool condition);
@@ -2228,21 +2239,25 @@ uint4 WaveActiveBallot(bool condition);
 // TODO(JS): 
 // subgroupBallotBitCount seems to take a uint4 parameter. 
 __glsl_extension(GL_KHR_shader_subgroup_ballot)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupBallotBitCount($0)")
 __target_intrinsic(cuda, "__popc(__ballot_sync(__activemask(), $0))")
 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();
 
 __glsl_extension(GL_KHR_shader_subgroup_basic)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "gl_SubgroupInvocationID")
 __target_intrinsic(cuda, "_getLaneId()")
 uint WaveGetLaneIndex();
 
 __glsl_extension(GL_KHR_shader_subgroup_basic)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupElect()")
 __target_intrinsic(cuda, "_waveIsFirstLane()")
 bool WaveIsFirstLane();
@@ -2252,6 +2267,7 @@ bool WaveIsFirstLane();
 
 __generic<T : __BuiltinArithmeticType>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupExcusiveMul($0)")
 T WavePrefixProduct(T expr);
 __generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WavePrefixProduct(vector<T,N> expr);
@@ -2259,6 +2275,7 @@ __generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> W
 
 __generic<T : __BuiltinArithmeticType>
 __glsl_extension(GL_KHR_shader_subgroup_arithmetic)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupExcusiveAdd($0)")
 T WavePrefixSum(T expr);
 __generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WavePrefixSum(vector<T,N> expr);
@@ -2279,6 +2296,7 @@ __generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> W
 
 // TODO(JS): This takes uvec4 parameter on GLSL
 __glsl_extension(GL_KHR_shader_subgroup_ballot)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupBallotExclusiveBitCount($0)")
 __target_intrinsic(cuda, "__popc(__ballot_sync(__activemask(), $0)  & _getLaneLtMask())")
 uint WavePrefixCountBits(bool value);
@@ -2295,6 +2313,7 @@ __generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> W
 
 __generic<T : __BuiltinType>
 __glsl_extension(GL_KHR_shader_subgroup_ballot)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgoupBroadcastFirst($0)")
 __target_intrinsic(cuda, "_waveReadFirst($0)")
 T WaveReadLaneFirst(T expr);
@@ -2303,6 +2322,7 @@ __generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M> WaveReadLan
 
 __generic<T : __BuiltinType>
 __glsl_extension(GL_KHR_shader_subgroup_ballot)
+__spirv_version(1.3)
 __target_intrinsic(glsl, "subgroupBroadcast($0, $1)")
 __target_intrinsic(cuda, "__shfl_sync(SLANG_CUDA_WARP_MASK, $0, $1)")
 T WaveReadLaneAt(T value, int lane);
diff --git a/source/slang/hlsl.meta.slang.h b/source/slang/hlsl.meta.slang.h
index c382db420..49c968cad 100644
--- a/source/slang/hlsl.meta.slang.h
+++ b/source/slang/hlsl.meta.slang.h
@@ -374,66 +374,20 @@ SLANG_RAW("    __target_intrinsic(glsl, \"EndPrimitive()\")\n")
 SLANG_RAW("    void RestartStrip();\n")
 SLANG_RAW("};\n")
 SLANG_RAW("\n")
-SLANG_RAW("#define VECTOR_MAP_UNARY(TYPE, COUNT, FUNC, VALUE) \\\n")
-SLANG_RAW("    vector<TYPE,COUNT> result; for(int i = 0; i < COUNT; ++i) { result[i] = FUNC(VALUE[i]); } return result\n")
-SLANG_RAW("    \n")
-SLANG_RAW("#define MATRIX_MAP_UNARY(TYPE, ROWS, COLS, FUNC, VALUE) \\\n")
-SLANG_RAW("    matrix<TYPE,ROWS,COLS> result; for(int i = 0; i < ROWS; ++i) { result[i] = FUNC(VALUE[i]); } return result\n")
-SLANG_RAW("\n")
-SLANG_RAW("#define VECTOR_MAP_BINARY(TYPE, COUNT, FUNC, LEFT, RIGHT) \\\n")
-SLANG_RAW("    vector<TYPE,COUNT> result; for(int i = 0; i < COUNT; ++i) { result[i] = FUNC(LEFT[i], RIGHT[i]); } return result\n")
-SLANG_RAW("    \n")
-SLANG_RAW("#define MATRIX_MAP_BINARY(TYPE, ROWS, COLS, FUNC, LEFT, RIGHT) \\\n")
-SLANG_RAW("    matrix<TYPE,ROWS,COLS> result; for(int i = 0; i < ROWS; ++i) { result[i] = FUNC(LEFT[i], RIGHT[i]); } return result\n")
-SLANG_RAW("\n")
-SLANG_RAW("#define VECTOR_MAP_TRINARY(TYPE, COUNT, FUNC, A, B, C) \\\n")
-SLANG_RAW("    vector<TYPE,COUNT> result; for(int i = 0; i < COUNT; ++i) { result[i] = FUNC(A[i], B[i], C[i]); } return result\n")
-SLANG_RAW("    \n")
-SLANG_RAW("#define MATRIX_MAP_TRINARY(TYPE, ROWS, COLS, FUNC, A, B, C) \\\n")
-SLANG_RAW("    matrix<TYPE,ROWS,COLS> result; for(int i = 0; i < ROWS; ++i) { result[i] = FUNC(A[i], B[i], C[i]); } return result\n")
+SLANG_RAW("// Note(tfoley): Trying to systematically add all the HLSL builtins\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Try to terminate the current draw or dispatch call (HLSL SM 4.0)\n")
 SLANG_RAW("void abort();\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Absolute value (HLSL SM 1.0)\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinSignedArithmeticType>\n")
-SLANG_RAW("T abs(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinSignedArithmeticType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> abs(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, abs, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinSignedArithmeticType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T,N,M> abs(matrix<T,N,M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, abs, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinSignedArithmeticType> T abs(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinSignedArithmeticType, let N : int> vector<T,N> abs(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinSignedArithmeticType, let N : int, let M : int> matrix<T,N,M> abs(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Inverse cosine (HLSL SM 1.0)\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T acos(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> acos(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, acos, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> acos(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, acos, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T acos(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> acos(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> acos(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Test if all components are non-zero (HLSL SM 1.0)\n")
 SLANG_RAW("__generic<T : __BuiltinType> bool all(T x);\n")
@@ -457,80 +411,37 @@ SLANG_RAW("__target_intrinsic(glsl, \"bool($0)\")\n")
 SLANG_RAW("bool any(T x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"any(bvec$N0($0))\")\n")
-SLANG_RAW("bool any(vector<T, N> x);\n")
-SLANG_RAW("// TODO: implementation of `any()` in the stdlib is\n")
-SLANG_RAW("// blocked on fixing implementation of `bool` vector\n")
-SLANG_RAW("// `getAt` on the CUDA codegen path.\n")
-SLANG_RAW("/*\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    bool result = false;\n")
-SLANG_RAW("    for(int i = 0; i < N; ++i)\n")
-SLANG_RAW("        result = result || any(x[i]);\n")
-SLANG_RAW("    return result;\n")
-SLANG_RAW("}\n")
-SLANG_RAW("*/\n")
+SLANG_RAW("bool any(vector<T,N> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("bool any(matrix<T, N, M> x);\n")
-SLANG_RAW("/*\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    bool result = false;\n")
-SLANG_RAW("    for(int i = 0; i < N; ++i)\n")
-SLANG_RAW("        result = result || any(x[i]);\n")
-SLANG_RAW("    return result;\n")
-SLANG_RAW("}\n")
-SLANG_RAW("*/\n")
+SLANG_RAW("// TODO: need to define GLSL mapping\n")
+SLANG_RAW("bool any(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Reinterpret bits as a double (HLSL SM 5.0)\n")
 SLANG_RAW("\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"packDouble2x32(uvec2($0, $1))\")\n")
 SLANG_RAW("__glsl_extension(GL_ARB_gpu_shader5)\n")
 SLANG_RAW("double asdouble(uint lowbits, uint highbits);\n")
 SLANG_RAW("\n")
+SLANG_RAW("double asdouble(uint lowbits, uint highbits);\n")
+SLANG_RAW("\n")
 SLANG_RAW("// Reinterpret bits as a float (HLSL SM 4.0)\n")
 SLANG_RAW("\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
+SLANG_RAW("// GLSL Scalar\n")
 SLANG_RAW("__target_intrinsic(glsl, \"intBitsToFloat\")\n")
 SLANG_RAW("float asfloat(int x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"uintBitsToFloat\")\n")
 SLANG_RAW("float asfloat(uint x);\n")
 SLANG_RAW("\n")
+SLANG_RAW("// GLSL Vector\n")
 SLANG_RAW("__generic<let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"intBitsToFloat\")\n")
-SLANG_RAW("vector<float, N> asfloat(vector< int, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(float, N, asfloat, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
+SLANG_RAW("vector<float,N> asfloat(vector< int,N> x);\n")
 SLANG_RAW("__generic<let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"uintBitsToFloat\")\n")
-SLANG_RAW("vector<float,N> asfloat(vector<uint,N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(float, N, asfloat, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<float,N,M> asfloat(matrix< int,N,M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(float, N, M, asfloat, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<float,N,M> asfloat(matrix<uint,N,M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(float, N, M, asfloat, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("vector<float,N> asfloat(vector<uint,N> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// No op\n")
 SLANG_RAW("__intrinsic_op(")
@@ -538,14 +449,12 @@ SLANG_SPLICE(kCompoundIntrinsicOp_Pos
 )
 SLANG_RAW(")\n")
 SLANG_RAW("float asfloat(float x);\n")
-SLANG_RAW("\n")
 SLANG_RAW("__generic<let N : int>\n")
 SLANG_RAW("__intrinsic_op(")
 SLANG_SPLICE(kCompoundIntrinsicOp_Pos
 )
 SLANG_RAW(")\n")
 SLANG_RAW("vector<float,N> asfloat(vector<float,N> x);\n")
-SLANG_RAW("\n")
 SLANG_RAW("__generic<let N : int, let M : int>\n")
 SLANG_RAW("__intrinsic_op(")
 SLANG_SPLICE(kCompoundIntrinsicOp_Pos
@@ -553,66 +462,32 @@ SLANG_SPLICE(kCompoundIntrinsicOp_Pos
 SLANG_RAW(")\n")
 SLANG_RAW("matrix<float,N,M> asfloat(matrix<float,N,M> x);\n")
 SLANG_RAW("\n")
-SLANG_RAW("// Inverse sine (HLSL SM 1.0)\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("T asin(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> asin(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T,N,asin,x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("// Pass thru to HLSL\n")
+SLANG_RAW("float asfloat(uint x);\n")
+SLANG_RAW("float asfloat(int x);\n")
+SLANG_RAW("__generic<let N : int, let M : int> matrix<float,N,M> asfloat(matrix< int,N,M> x);\n")
+SLANG_RAW("__generic<let N : int, let M : int> matrix<float,N,M> asfloat(matrix<uint,N,M> x);\n")
 SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> asin(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T,N,M,asin,x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("// Inverse sine (HLSL SM 1.0)\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T asin(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> asin(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> asin(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Reinterpret bits as an int (HLSL SM 4.0)\n")
 SLANG_RAW("\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
+SLANG_RAW("// GLSL scalar\n")
 SLANG_RAW("__target_intrinsic(glsl, \"floatBitsToInt\")\n")
 SLANG_RAW("int asint(float x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"int($0)\")\n")
 SLANG_RAW("int asint(uint x);\n")
 SLANG_RAW("\n")
+SLANG_RAW("// GLSL Vector\n")
 SLANG_RAW("__generic<let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"floatBitsToInt\")\n")
-SLANG_RAW("vector<int, N> asint(vector<float, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(int, N, asint, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
+SLANG_RAW("vector<int,N> asint(vector<float,N> x);\n")
 SLANG_RAW("__generic<let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"ivec$N0($0)\")\n")
-SLANG_RAW("vector<int, N> asint(vector<uint, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(int, N, asint, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<int, N, M> asint(matrix<float, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(int, N, M, asint, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<int, N, M> asint(matrix<uint, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(int, N, M, asint, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("vector<int,N> asint(vector<uint,N> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// No op\n")
 SLANG_RAW("__intrinsic_op(")
@@ -620,14 +495,12 @@ SLANG_SPLICE(kCompoundIntrinsicOp_Pos
 )
 SLANG_RAW(")\n")
 SLANG_RAW("int asint(int x);\n")
-SLANG_RAW("\n")
 SLANG_RAW("__generic<let N : int>\n")
 SLANG_RAW("__intrinsic_op(")
 SLANG_SPLICE(kCompoundIntrinsicOp_Pos
 )
 SLANG_RAW(")\n")
 SLANG_RAW("vector<int,N> asint(vector<int,N> x);\n")
-SLANG_RAW("\n")
 SLANG_RAW("__generic<let N : int, let M : int>\n")
 SLANG_RAW("__intrinsic_op(")
 SLANG_SPLICE(kCompoundIntrinsicOp_Pos
@@ -635,66 +508,51 @@ SLANG_SPLICE(kCompoundIntrinsicOp_Pos
 SLANG_RAW(")\n")
 SLANG_RAW("matrix<int,N,M> asint(matrix<int,N,M> x);\n")
 SLANG_RAW("\n")
+SLANG_RAW("// Pass thru HLSL\n")
+SLANG_RAW("\n")
+SLANG_RAW("int asint(float x);\n")
+SLANG_RAW("int asint(uint x);\n")
+SLANG_RAW("\n")
+SLANG_RAW("__generic<let N : int> vector<int,N> asint(vector<uint,N> x);\n")
+SLANG_RAW("__generic<let N : int, let M : int> matrix<int,N,M> asint(matrix<float,N,M> x);\n")
+SLANG_RAW("__generic<let N : int, let M : int> matrix<int,N,M> asint(matrix<uint,N,M> x);\n")
+SLANG_RAW("\n")
 SLANG_RAW("// Reinterpret bits of double as a uint (HLSL SM 5.0)\n")
 SLANG_RAW("\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"{ uvec2 v = unpackDouble2x32($0); $1 = v.x; $2 = v.y; }\")\n")
 SLANG_RAW("__glsl_extension(GL_ARB_gpu_shader5)\n")
 SLANG_RAW("void asuint(double value, out uint lowbits, out uint highbits);\n")
 SLANG_RAW("\n")
+SLANG_RAW("void asuint(double value, out uint lowbits, out uint highbits);\n")
+SLANG_RAW("\n")
 SLANG_RAW("// Reinterpret bits as a uint (HLSL SM 4.0)\n")
 SLANG_RAW("\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
+SLANG_RAW("// GLSL Scalar\n")
 SLANG_RAW("__target_intrinsic(glsl, \"floatBitsToUint\")\n")
 SLANG_RAW("uint asuint(float x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"uint($0)\")\n")
 SLANG_RAW("uint asuint(int x);\n")
 SLANG_RAW("\n")
+SLANG_RAW("// GLSL Vector\n")
 SLANG_RAW("__generic<let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"floatBitsToUint\")\n")
-SLANG_RAW("vector<uint,N> asuint(vector<float,N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(uint, N, asuint, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
+SLANG_RAW("vector<uint,N> asuint(vector<float,N> x);\n")
 SLANG_RAW("__generic<let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"uvec$N0($0)\")\n")
-SLANG_RAW("vector<uint, N> asuint(vector<int, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(uint, N, asuint, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<uint,N,M> asuint(matrix<float,N,M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(uint, N, M, asuint, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<uint, N, M> asuint(matrix<int, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(uint, N, M, asuint, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("vector<uint,N> asuint(vector<int,N> x);\n")
 SLANG_RAW("\n")
+SLANG_RAW("// No op\n")
 SLANG_RAW("__intrinsic_op(")
 SLANG_SPLICE(kCompoundIntrinsicOp_Pos
 )
 SLANG_RAW(")\n")
 SLANG_RAW("uint asuint(uint x);\n")
-SLANG_RAW("\n")
 SLANG_RAW("__generic<let N : int>\n")
 SLANG_RAW("__intrinsic_op(")
 SLANG_SPLICE(kCompoundIntrinsicOp_Pos
 )
 SLANG_RAW(")\n")
 SLANG_RAW("vector<uint,N> asuint(vector<uint,N> x);\n")
-SLANG_RAW("\n")
 SLANG_RAW("__generic<let N : int, let M : int>\n")
 SLANG_RAW("__intrinsic_op(")
 SLANG_SPLICE(kCompoundIntrinsicOp_Pos
@@ -702,180 +560,71 @@ SLANG_SPLICE(kCompoundIntrinsicOp_Pos
 SLANG_RAW(")\n")
 SLANG_RAW("matrix<uint,N,M> asuint(matrix<uint,N,M> x);\n")
 SLANG_RAW("\n")
-SLANG_RAW("// Inverse tangent (HLSL SM 1.0)\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T atan(T x);\n")
+SLANG_RAW("// Pass thru HLSL\n")
+SLANG_RAW("uint asuint(float x);\n")
+SLANG_RAW("uint asuint(int x);\n")
 SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> atan(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, atan, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<let N : int> vector<uint,N> asuint(vector<float,N> x);\n")
+SLANG_RAW("__generic<let N : int> vector<uint,N> asuint(vector<int,N> x);\n")
 SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> atan(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, atan, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<let N : int, let M : int> matrix<uint,N,M> asuint(matrix<float,N,M> x);\n")
+SLANG_RAW("__generic<let N : int, let M : int> matrix<uint,N,M> asuint(matrix<int,N,M> x);\n")
+SLANG_RAW("\n")
+SLANG_RAW("// Inverse tangent (HLSL SM 1.0)\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T atan(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> atan(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> atan(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl,\"atan($0,$1)\")\n")
 SLANG_RAW("T atan2(T y, T x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl,\"atan($0,$1)\")\n")
-SLANG_RAW("vector<T, N> atan2(vector<T, N> y, vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_BINARY(T, N, atan2, y, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("vector<T,N> atan2(vector<T,N> y, vector<T,N> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T,N,M> atan2(matrix<T,N,M> y, matrix<T,N,M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_BINARY(T, N, M, atan2, y, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__target_intrinsic(glsl,\"atan($0,$1)\")\n")
+SLANG_RAW("matrix<T,N,M> atan2(matrix<T,N,M> y, matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Ceiling (HLSL SM 1.0)\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T ceil(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> ceil(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, ceil, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> ceil(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, ceil, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T ceil(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> ceil(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> ceil(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Check access status to tiled resource\n")
 SLANG_RAW("bool CheckAccessFullyMapped(uint status);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Clamp (HLSL SM 1.0)\n")
-SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("T clamp(T x, T minBound, T maxBound)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    return min(max(x, minBound), maxBound);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> clamp(vector<T, N> x, vector<T, N> minBound, vector<T, N> maxBound)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    return min(max(x, minBound), maxBound);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T,N,M> clamp(matrix<T,N,M> x, matrix<T,N,M> minBound, matrix<T,N,M> maxBound)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    return min(max(x, minBound), maxBound);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinArithmeticType> T clamp(T x, T min, T max);\n")
+SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> clamp(vector<T,N> x, vector<T,N> min, vector<T,N> max);\n")
+SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> clamp(matrix<T,N,M> x, matrix<T,N,M> min, matrix<T,N,M> max);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Clip (discard) fragment conditionally\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("void clip(T x);\n")
-SLANG_RAW("// TODO: filling this in here requires ability to invoke `operator<(T,T)`\n")
-SLANG_RAW("/*{\n")
-SLANG_RAW("    if(x < T(0)) discard;\n")
-SLANG_RAW("}*/\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("void clip(vector<T,N> x);\n")
-SLANG_RAW("// TODO: filling this in here requires ability to invoke `operator<(T,T)`\n")
-SLANG_RAW("/*{\n")
-SLANG_RAW("    if(any(x < T(0))) discard;\n")
-SLANG_RAW("}*/\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("void clip(matrix<T,N,M> x);\n")
-SLANG_RAW("// TODO: filling this in here requires ability to invoke `operator<(T,T)`\n")
-SLANG_RAW("/*{\n")
-SLANG_RAW("    if(any(x < T(0))) discard;\n")
-SLANG_RAW("}*/\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> void clip(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> void clip(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> void clip(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Cosine\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T cos(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> cos(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T,N, cos, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> cos(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, cos, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T cos(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> cos(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> cos(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Hyperbolic cosine\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T cosh(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T,N> cosh(vector<T,N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T,N, cosh, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> cosh(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, cosh, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T cosh(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> cosh(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> cosh(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Population count\n")
 SLANG_RAW("__target_intrinsic(glsl, \"bitCount\")\n")
 SLANG_RAW("uint countbits(uint value);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Cross product\n")
-SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T,3> cross(vector<T,3> left, vector<T,3> right);\n")
-SLANG_RAW("// TODO: filling this in here requires ability to invoke `operator*(T,T)`, etc.\n")
-SLANG_RAW("/*{\n")
-SLANG_RAW("    return vector<T,3>(\n")
-SLANG_RAW("        left.y * right.z - left.z * right.y,\n")
-SLANG_RAW("        left.z * right.x - left.x * right.z,\n")
-SLANG_RAW("        left.x * right.y - left.y * right.x);\n")
-SLANG_RAW("}*/\n")
-SLANG_RAW("\n")
+SLANG_RAW("__generic<T : __BuiltinArithmeticType> vector<T,3> cross(vector<T,3> x, vector<T,3> y);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Convert encoded color\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("int4 D3DCOLORtoUBYTE4(float4 color)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    let scaled = color.zyxw * 255.001999f;\n")
-SLANG_RAW("    return int4(scaled);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("int4 D3DCOLORtoUBYTE4(float4 x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Partial-difference derivatives\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
@@ -964,25 +713,9 @@ SLANG_RAW("matrix<T,N,M> ddy_fine(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Radians to degrees\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("T degrees(T x);\n")
-SLANG_RAW("// TODO: filling this in here requires ability to invoke `operator*` on T,\n")
-SLANG_RAW("// and convert a constant to `T` for the conversion factor\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("vector<T,N> degrees(vector<T,N> x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> degrees(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, degrees, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T degrees(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> degrees(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> degrees(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Matrix determinant\n")
 SLANG_RAW("\n")
@@ -1057,44 +790,14 @@ SLANG_RAW("__target_intrinsic(glsl, \"interpolateAtOffset($0, vec2($1) / 16.0f)\
 SLANG_RAW("matrix<T,N,M> EvaluateAttributeSnapped(matrix<T,N,M> x, int2 offset);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Base-e exponent\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T exp(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> exp(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, exp, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> exp(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, exp, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T exp(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> exp(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> exp(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Base-2 exponent\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T exp2(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T,N> exp2(vector<T,N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, exp2, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T,N,M> exp2(matrix<T,N,M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, exp2, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T exp2(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> exp2(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> exp2(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Convert 16-bit float stored in low bits of integer\n")
 SLANG_RAW("__target_intrinsic(glsl, \"unpackHalf2x16($0).x\")\n")
@@ -1113,79 +816,38 @@ SLANG_RAW("__target_intrinsic(glsl, \"packHalf2x16(vec2($0,0.0))\")\n")
 SLANG_RAW("vector<uint,N> f32tof16(vector<float,N> value);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Flip surface normal to face forward, if needed\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T,N> faceforward(vector<T,N> n, vector<T,N> i, vector<T,N> ng);\n")
-SLANG_RAW("/*{\n")
-SLANG_RAW("    return dot(ng, i) < T(0.0f) ? n : -n;\n")
-SLANG_RAW("}*/\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> faceforward(vector<T,N> n, vector<T,N> i, vector<T,N> ng);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Find first set bit starting at high bit and working down\n")
 SLANG_RAW("__target_intrinsic(glsl,\"findMSB\")\n")
 SLANG_RAW("int firstbithigh(int value);\n")
 SLANG_RAW("\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl,\"findMSB\")\n")
-SLANG_RAW("__generic<let N : int>\n")
-SLANG_RAW("vector<int, N> firstbithigh(vector<int, N> value)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(int, N, firstbithigh, value);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<let N : int> vector<int,N> firstbithigh(vector<int,N> value);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__target_intrinsic(glsl,\"findMSB\")\n")
 SLANG_RAW("uint firstbithigh(uint value);\n")
 SLANG_RAW("\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl,\"findMSB\")\n")
-SLANG_RAW("__generic<let N : int>\n")
-SLANG_RAW("vector<uint,N> firstbithigh(vector<uint,N> value)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(uint, N, firstbithigh, value);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<let N : int> vector<uint,N> firstbithigh(vector<uint,N> value);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Find first set bit starting at low bit and working up\n")
 SLANG_RAW("__target_intrinsic(glsl,\"findLSB\")\n")
 SLANG_RAW("int firstbitlow(int value);\n")
 SLANG_RAW("\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl,\"findLSB\")\n")
-SLANG_RAW("__generic<let N : int>\n")
-SLANG_RAW("vector<int,N> firstbitlow(vector<int,N> value)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(int, N, firstbitlow, value);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<let N : int> vector<int,N> firstbitlow(vector<int,N> value);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__target_intrinsic(glsl,\"findLSB\")\n")
 SLANG_RAW("uint firstbitlow(uint value);\n")
 SLANG_RAW("\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl,\"findLSB\")\n")
-SLANG_RAW("__generic<let N : int>\n")
-SLANG_RAW("vector<uint,N> firstbitlow(vector<uint,N> value)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(uint, N, firstbitlow, value);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<let N : int> vector<uint,N> firstbitlow(vector<uint,N> value);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Floor (HLSL SM 1.0)\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T floor(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> floor(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, floor, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> floor(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, floor, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T floor(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> floor(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> floor(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Fused multiply-add for doubles\n")
 SLANG_RAW("double fma(double a, double b, double c);\n")
@@ -1193,23 +855,9 @@ SLANG_RAW("__generic<let N : int> vector<double, N> fma(vector<double, N> a, vec
 SLANG_RAW("__generic<let N : int, let M : int> matrix<double,N,M> fma(matrix<double,N,M> a, matrix<double,N,M> b, matrix<double,N,M> c);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Floating point remainder of x/y\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T fmod(T x, T y);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> fmod(vector<T, N> x, vector<T, N> y)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_BINARY(T, N, fmod, x, y);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> fmod(matrix<T, N, M> x, matrix<T, N, M> y)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_BINARY(T, N, M, fmod, x, y);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T fmod(T x, T y);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> fmod(vector<T,N> x, vector<T,N> y);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> fmod(matrix<T,N,M> x, matrix<T,N,M> y);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Fractional part\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
@@ -1217,58 +865,22 @@ SLANG_RAW("__target_intrinsic(glsl, fract)\n")
 SLANG_RAW("T frac(T x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, fract)\n")
-SLANG_RAW("vector<T, N> frac(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, frac, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("vector<T,N> frac(vector<T,N> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("matrix<T, N, M> frac(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, frac, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__target_intrinsic(glsl, fract)\n")
+SLANG_RAW("matrix<T,N,M> frac(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Split float into mantissa and exponent\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("T frexp(T x, out T exp);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> frexp(vector<T, N> x, out vector<T, N> exp)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_BINARY(T, N, frexp, x, exp);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> frexp(matrix<T, N, M> x, out matrix<T, N, M> exp)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_BINARY(T, N, M, frexp, x, exp);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T frexp(T x, out T exp);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> frexp(vector<T,N> x, out vector<T,N> exp);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> frexp(matrix<T,N,M> x, out matrix<T,N,M> exp);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Texture filter width\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T fwidth(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> fwidth(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, fwidth, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> fwidth(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, fwidth, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T fwidth(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> fwidth(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> fwidth(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Get number of samples in render target\n")
 SLANG_RAW("uint GetRenderTargetSampleCount();\n")
@@ -1411,187 +1023,80 @@ SLANG_RAW("\n")
 SLANG_RAW("// Is floating-point value finite?\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(cpu)\n")
-SLANG_RAW("__target_intrinsic(cuda)\n")
-SLANG_RAW("//__target_intrinsic(glsl, \"(!(isinf($0) || isnan($0)))\")\n")
-SLANG_RAW("bool isfinite(T x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    return !(isinf(x) || isnan(x));\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__target_intrinsic(glsl, \"(!(isinf($0) || isnan($0)))\")\n")
+SLANG_RAW("bool isfinite(T x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("//__target_intrinsic(glsl, \"(!(isinf($0) || isnan($0)))\")\n")
-SLANG_RAW("vector<bool, N> isfinite(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(bool, N, isfinite, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__target_intrinsic(glsl, \"(!(isinf($0) || isnan($0)))\")\n")
+SLANG_RAW("vector<bool,N> isfinite(vector<T,N> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<bool, N, M> isfinite(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(bool, N, M, isfinite, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("matrix<bool,N,M> isfinite(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Is floating-point value infinite?\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("bool isinf(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<bool, N> isinf(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(bool, N, isinf, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<bool, N, M> isinf(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(bool, N, M, isinf, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> bool isinf(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<bool,N> isinf(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<bool,N,M> isinf(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Is floating-point value not-a-number?\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("bool isnan(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<bool, N> isnan(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(bool, N, isnan, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<bool, N, M> isnan(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(bool, N, M, isnan, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> bool isnan(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<bool,N> isnan(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<bool,N,M> isnan(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Construct float from mantissa and exponent\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"($0 * pow(2.0f, $1))\")\n")
 SLANG_RAW("T ldexp(T x, T exp);\n")
-SLANG_RAW("/*{\n")
-SLANG_RAW("    return x * exp2(exp);\n")
-SLANG_RAW("}*/\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"($0 * pow(2.0f, $1))\")\n")
-SLANG_RAW("vector<T, N> ldexp(vector<T, N> x, vector<T, N> exp)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_BINARY(T, N, ldexp, x, exp);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("vector<T,N> ldexp(vector<T,N> x, vector<T,N> exp);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> ldexp(matrix<T, N, M> x, matrix<T, N, M> exp)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_BINARY(T, N, M, ldexp, x, exp);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("matrix<T,N,M> ldexp(matrix<T,N,M> x, matrix<T,N,M> exp);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Vector length\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("T length(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    return sqrt(dot(x, x));\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> T length(vector<T,N> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Linear interpolation\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, mix)\n")
 SLANG_RAW("T lerp(T x, T y, T s);\n")
-SLANG_RAW("/*{\n")
-SLANG_RAW("    return x * (1 - s) + y * s;\n")
-SLANG_RAW("}*/\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, mix)\n")
-SLANG_RAW("vector<T, N> lerp(vector<T, N> x, vector<T, N> y, vector<T, N> s)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_TRINARY(T, N, lerp, x, y, s);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("vector<T,N> lerp(vector<T,N> x, vector<T,N> y, vector<T,N> s);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T,N,M> lerp(matrix<T,N,M> x, matrix<T,N,M> y, matrix<T,N,M> s)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_TRINARY(T, N, M, lerp, x, y, s);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__target_intrinsic(glsl, mix)\n")
+SLANG_RAW("matrix<T,N,M> lerp(matrix<T,N,M> x, matrix<T,N,M> y, matrix<T,N,M> s);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Legacy lighting function (obsolete)\n")
 SLANG_RAW("float4 lit(float n_dot_l, float n_dot_h, float m);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Base-e logarithm\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T log(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> log(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, log, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> log(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, log, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T log(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> log(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> log(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Base-10 logarithm\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType> \n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"(log( $0 ) * $S0( 0.43429448190325182765112891891661) )\" )\n")
 SLANG_RAW("T log10(T x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"(log( $0 ) * $S0(0.43429448190325182765112891891661) )\" )\n")
-SLANG_RAW("vector<T,N> log10(vector<T,N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, log10, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("vector<T,N> log10(vector<T,N> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> \n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T,N,M> log10(matrix<T,N,M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, log10, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__target_intrinsic(glsl, \"(log( $0 ) * $S0(0.43429448190325182765112891891661) )\" )\n")
+SLANG_RAW("matrix<T,N,M> log10(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Base-2 logarithm\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T log2(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T,N> log2(vector<T,N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, log2, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T,N,M> log2(matrix<T,N,M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, log2, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T log2(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> log2(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> log2(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// multiply-add\n")
 SLANG_RAW("\n")
@@ -1605,65 +1110,19 @@ SLANG_RAW("__target_intrinsic(glsl, fma)\n")
 SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> mad(matrix<T,N,M> mvalue, matrix<T,N,M> avalue, matrix<T,N,M> bvalue);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// maximum\n")
-SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n")
-SLANG_RAW("T max(T x, T y);\n")
-SLANG_RAW("// Note: a stdlib implementation of `max` (or `min`) will require splitting\n")
-SLANG_RAW("// floating-point and integer cases apart, because the floating-point\n")
-SLANG_RAW("// version needs to correctly handle the case where one of the inputs\n")
-SLANG_RAW("// is not-a-number.\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> max(vector<T, N> x, vector<T, N> y)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_BINARY(T, N, max, x, y);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> max(matrix<T, N, M> x, matrix<T, N, M> y)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_BINARY(T, N, M, max, x, y);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinArithmeticType> T max(T x, T y);\n")
+SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> max(vector<T,N> x, vector<T,N> y);\n")
+SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> max(matrix<T,N,M> x, matrix<T,N,M> y);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// minimum\n")
-SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n")
-SLANG_RAW("T min(T x, T y);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T,N> min(vector<T,N> x, vector<T,N> y)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_BINARY(T, N, min, x, y);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T,N,M> min(matrix<T,N,M> x, matrix<T,N,M> y)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_BINARY(T, N, M, min, x, y);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinArithmeticType> T min(T x, T y);\n")
+SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> min(vector<T,N> x, vector<T,N> y);\n")
+SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> min(matrix<T,N,M> x, matrix<T,N,M> y);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// split into integer and fractional parts (both with same sign)\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T modf(T x, out T ip);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T,N> modf(vector<T,N> x, out vector<T,N> ip)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_BINARY(T, N, modf, x, ip);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T,N,M> modf(matrix<T,N,M> x, out matrix<T,N,M> ip)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_BINARY(T, N, M, modf, x, ip);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T modf(T x, out T ip);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> modf(vector<T,N> x, out vector<T,N> ip);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> modf(matrix<T,N,M> x, out matrix<T,N,M> ip);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// msad4 (whatever that is)\n")
 SLANG_RAW("uint4 msad4(uint reference, uint2 source, uint4 accum);\n")
@@ -1694,16 +1153,8 @@ SLANG_RAW("// matrix-matrix\n")
 SLANG_RAW("__generic<T : __BuiltinArithmeticType, let R : int, let N : int, let C : int> __intrinsic_op(mulMatrixMatrix) matrix<T,R,C> mul(matrix<T,R,N> x, matrix<T,N,C> y);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// noise (deprecated)\n")
-SLANG_RAW("\n")
-SLANG_RAW("float noise(float x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    return 0;\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<let N : int> float noise(vector<float, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    return 0;\n")
-SLANG_RAW("}\n")
+SLANG_RAW("float noise(float x);\n")
+SLANG_RAW("__generic<let N : int> float noise(vector<float, N> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("/// Indicate that an index may be non-uniform at execution time.\n")
 SLANG_RAW("///\n")
@@ -1723,48 +1174,23 @@ SLANG_RAW("/// Note: a future version of Slang may take responsibility for inser
 SLANG_RAW("/// to this function as necessary in output code, rather than make this\n")
 SLANG_RAW("/// the user's responsibility, so that the default behavior of the language\n")
 SLANG_RAW("/// is more semantically \"correct.\"\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, nonuniformEXT)\n")
 SLANG_RAW("__glsl_extension(GL_EXT_nonuniform_qualifier)\n")
 SLANG_RAW("[__readNone]\n")
-SLANG_RAW("uint NonUniformResourceIndex(uint index)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    return index;\n")
-SLANG_RAW("}\n")
+SLANG_RAW("uint NonUniformResourceIndex(uint index);\n")
 SLANG_RAW("\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, nonuniformEXT)\n")
 SLANG_RAW("__glsl_extension(GL_EXT_nonuniform_qualifier)\n")
 SLANG_RAW("[__readNone]\n")
-SLANG_RAW("int NonUniformResourceIndex(int index)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    return index;\n")
-SLANG_RAW("}\n")
+SLANG_RAW("int NonUniformResourceIndex(int index);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Normalize a vector\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> normalize(vector<T,N> x);\n")
-SLANG_RAW("/*{\n")
-SLANG_RAW("    return x / length(x);\n")
-SLANG_RAW("}*/\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Raise to a power\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T pow(T x, T y);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> pow(vector<T, N> x, vector<T, N> y)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_BINARY(T, N, pow, x, y);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T,N,M> pow(matrix<T,N,M> x, matrix<T,N,M> y)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_BINARY(T, N, M, pow, x, y);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T pow(T x, T y);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> pow(vector<T,N> x, vector<T,N> y);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> pow(matrix<T,N,M> x, matrix<T,N,M> y);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Output message\n")
 SLANG_RAW("\n")
@@ -1842,66 +1268,26 @@ SLANG_RAW("    out float RoundedInsideTessFactors,\n")
 SLANG_RAW("    out float UnroundedInsideTessFactors);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Degrees to radians\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("T radians(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> radians(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, radians, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> radians(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, radians, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T radians(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> radians(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> radians(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Approximate reciprocal\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"1.0/($0)\")\n")
 SLANG_RAW("T rcp(T x);\n")
-SLANG_RAW("/*{\n")
-SLANG_RAW("    return T(1) / x;\n")
-SLANG_RAW("}*/\n")
 SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("vector<T, N> rcp(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, rcp, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("// Note: GLSL doesn't define a vector `rcp`, so not intrinsic there\n")
-SLANG_RAW("matrix<T, N, M> rcp(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, rcp, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("// TODO: vector and matrix approx. reciprocals needto be deconstructed for GLSL\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> rcp(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> rcp(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Reflect incident vector across plane with given normal\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
 SLANG_RAW("vector<T,N> reflect(vector<T,N> i, vector<T,N> n);\n")
-SLANG_RAW("/*{\n")
-SLANG_RAW("    return i - T(2) * dot(n,i) * n;\n")
-SLANG_RAW("}*/\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Refract incident vector given surface normal and index of refraction\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
 SLANG_RAW("vector<T,N> refract(vector<T,N> i, vector<T,N> n, float eta);\n")
-SLANG_RAW("/*{\n")
-SLANG_RAW("    let dotNI = dot(n,i);\n")
-SLANG_RAW("    let k = T(1) - eta*eta*(T(1) - dotNI * dotNI);\n")
-SLANG_RAW("    if(k < 0) return vector<T,N>(T(0));\n")
-SLANG_RAW("    return eta * i - (eta * dotNI + sqrt(k)) * n;\n")
-SLANG_RAW("}*/\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Reverse order of bits\n")
 SLANG_RAW("__target_intrinsic(glsl, \"bitfieldReverse\")\n")
@@ -1911,56 +1297,45 @@ SLANG_RAW("__target_intrinsic(glsl, \"bitfieldReverse\")\n")
 SLANG_RAW("__generic<let N : int> vector<uint,N> reversebits(vector<uint,N> value);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Round-to-nearest\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T round(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> round(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, round, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T,N,M> round(matrix<T,N,M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, round, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T round(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> round(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> round(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Reciprocal of square root\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"inversesqrt($0)\")\n")
 SLANG_RAW("T rsqrt(T x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"inversesqrt($0)\")\n")
-SLANG_RAW("vector<T, N> rsqrt(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, rsqrt, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("vector<T,N> rsqrt(vector<T,N> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> rsqrt(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, rsqrt, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__target_intrinsic(glsl, \"inversesqrt($0)\")\n")
+SLANG_RAW("matrix<T,N,M> rsqrt(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Clamp value to [0,1] range\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
+SLANG_RAW("__target_intrinsic(glsl, \"clamp($0, 0, 1)\")\n")
+SLANG_RAW("T saturate(T x);\n")
+SLANG_RAW("\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
+SLANG_RAW("__target_intrinsic(glsl, \"clamp($0, 0, 1)\")\n")
+SLANG_RAW("vector<T,N> saturate(vector<T,N> x);\n")
+SLANG_RAW("\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
+SLANG_RAW("__target_intrinsic(glsl, \"clamp($0, 0, 1)\")\n")
+SLANG_RAW("matrix<T,N,M> saturate(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
+SLANG_RAW("__specialized_for_target(glsl)\n")
 SLANG_RAW("T saturate(T x)\n")
 SLANG_RAW("{\n")
 SLANG_RAW("    return clamp<T>(x, T(0), T(1));\n")
 SLANG_RAW("}\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
+SLANG_RAW("__specialized_for_target(glsl)\n")
 SLANG_RAW("vector<T,N> saturate(vector<T,N> x)\n")
 SLANG_RAW("{\n")
 SLANG_RAW("    return clamp<T,N>(x,\n")
@@ -1968,229 +1343,115 @@ SLANG_RAW("        vector<T,N>(T(0)),\n")
 SLANG_RAW("        vector<T,N>(T(1)));\n")
 SLANG_RAW("}\n")
 SLANG_RAW("\n")
+SLANG_RAW("// HACK: need a helper to turn a scalar into a matrix,\n")
+SLANG_RAW("// because GLSL and HLSL disagree on the semantics of\n")
+SLANG_RAW("// constructing a matrix from a single scalar.\n")
+SLANG_RAW("__generic<T, let N : int, let M : int>\n")
+SLANG_RAW("matrix<T,N,M> __scalarToMatrix(T value);\n")
+SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
+SLANG_RAW("__specialized_for_target(glsl)\n")
 SLANG_RAW("matrix<T,N,M> saturate(matrix<T,N,M> x)\n")
 SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, saturate, x);\n")
+SLANG_RAW("    return clamp<T,N,M>(x,\n")
+SLANG_RAW("        __scalarToMatrix<T,N,M>(T(0)),\n")
+SLANG_RAW("        __scalarToMatrix<T,N,M>(T(1)));\n")
 SLANG_RAW("}\n")
 SLANG_RAW("\n")
+SLANG_RAW("\n")
 SLANG_RAW("// Extract sign of value\n")
 SLANG_RAW("__generic<T : __BuiltinSignedArithmeticType>\n")
 SLANG_RAW("__target_intrinsic(glsl, \"int(sign($0))\")\n")
 SLANG_RAW("int sign(T x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinSignedArithmeticType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"ivec$N0(sign($0))\")\n")
-SLANG_RAW("vector<int, N> sign(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(int, N, sign, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("vector<int,N> sign(vector<T,N> x);\n")
 SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinSignedArithmeticType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<int, N, M> sign(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(int, N, M, sign, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinSignedArithmeticType, let N : int, let M : int> matrix<int,N,M> sign(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Sine\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T sin(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> sin(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, sin, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> sin(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, sin, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T sin(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> sin(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> sin(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Sine and cosine\n")
 SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("void sincos(T x, out T s, out T c)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    s = sin(x);\n")
-SLANG_RAW("    c = cos(x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("void sincos(vector<T,N> x, out vector<T,N> s, out vector<T,N> c)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    s = sin(x);\n")
-SLANG_RAW("    c = cos(x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("void sincos(matrix<T,N,M> x, out matrix<T,N,M> s, out matrix<T,N,M> c)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    s = sin(x);\n")
-SLANG_RAW("    c = cos(x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__target_intrinsic(glsl, \"$1 = sin($0); $2 = cos($0);\")\n")
+SLANG_RAW("void sincos(T x, out T s, out T c);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> void sincos(vector<T,N> x, out vector<T,N> s, out vector<T,N> c);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> void sincos(matrix<T,N,M> x, out matrix<T,N,M> s, out matrix<T,N,M> c);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Hyperbolic Sine\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T sinh(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> sinh(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, sinh, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> sinh(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, sinh, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T sinh(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> sinh(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> sinh(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Smooth step (Hermite interpolation)\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T smoothstep(T min, T max, T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> smoothstep(vector<T, N> min, vector<T, N> max, vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_TRINARY(T, N, smoothstep, min, max, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> smoothstep(matrix<T, N, M> min, matrix<T, N, M> max, matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_TRINARY(T, N, M, smoothstep, min, max, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T smoothstep(T min, T max, T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> smoothstep(vector<T,N> min, vector<T,N> max, vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> smoothstep(matrix<T,N,M> min, matrix<T,N,M> max, matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Square root\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T sqrt(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> sqrt(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, sqrt, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> sqrt(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, sqrt, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T sqrt(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> sqrt(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> sqrt(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Step function\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("T step(T y, T x);\n")
-SLANG_RAW("/*{\n")
-SLANG_RAW("    return x < y ? T(0.0f) : T(1.0f);\n")
-SLANG_RAW("}*/\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T,N> step(vector<T,N> y, vector<T,N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_BINARY(T, N, step, y, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> step(matrix<T, N, M> y, matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_BINARY(T, N, M, step, y, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T step(T y, T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> step(vector<T,N> y, vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> step(matrix<T,N,M> y, matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Tangent\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T tan(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> tan(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, tan, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> tan(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, tan, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T tan(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> tan(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> tan(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Hyperbolic tangent\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T tanh(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T tanh(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> tanh(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> tanh(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T,N> tanh(vector<T,N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, tanh, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("// Legacy texture-fetch operations\n")
 SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T,N,M> tanh(matrix<T,N,M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, tanh, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("/*\n")
+SLANG_RAW("float4 tex1D(sampler1D s, float t);\n")
+SLANG_RAW("float4 tex1D(sampler1D s, float t, float ddx, float ddy);\n")
+SLANG_RAW("float4 tex1Dbias(sampler1D s, float4 t);\n")
+SLANG_RAW("float4 tex1Dgrad(sampler1D s, float t, float ddx, float ddy);\n")
+SLANG_RAW("float4 tex1Dlod(sampler1D s, float4 t);\n")
+SLANG_RAW("float4 tex1Dproj(sampler1D s, float4 t);\n")
+SLANG_RAW("\n")
+SLANG_RAW("float4 tex2D(sampler2D s, float2 t);\n")
+SLANG_RAW("float4 tex2D(sampler2D s, float2 t, float2 ddx, float2 ddy);\n")
+SLANG_RAW("float4 tex2Dbias(sampler2D s, float4 t);\n")
+SLANG_RAW("float4 tex2Dgrad(sampler2D s, float2 t, float2 ddx, float2 ddy);\n")
+SLANG_RAW("float4 tex2Dlod(sampler2D s, float4 t);\n")
+SLANG_RAW("float4 tex2Dproj(sampler2D s, float4 t);\n")
+SLANG_RAW("\n")
+SLANG_RAW("float4 tex3D(sampler3D s, float3 t);\n")
+SLANG_RAW("float4 tex3D(sampler3D s, float3 t, float3 ddx, float3 ddy);\n")
+SLANG_RAW("float4 tex3Dbias(sampler3D s, float4 t);\n")
+SLANG_RAW("float4 tex3Dgrad(sampler3D s, float3 t, float3 ddx, float3 ddy);\n")
+SLANG_RAW("float4 tex3Dlod(sampler3D s, float4 t);\n")
+SLANG_RAW("float4 tex3Dproj(sampler3D s, float4 t);\n")
+SLANG_RAW("\n")
+SLANG_RAW("float4 texCUBE(samplerCUBE s, float3 t);\n")
+SLANG_RAW("float4 texCUBE(samplerCUBE s, float3 t, float3 ddx, float3 ddy);\n")
+SLANG_RAW("float4 texCUBEbias(samplerCUBE s, float4 t);\n")
+SLANG_RAW("float4 texCUBEgrad(samplerCUBE s, float3 t, float3 ddx, float3 ddy);\n")
+SLANG_RAW("float4 texCUBElod(samplerCUBE s, float4 t);\n")
+SLANG_RAW("float4 texCUBEproj(samplerCUBE s, float4 t);\n")
+SLANG_RAW("*/\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Matrix transpose\n")
-SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("matrix<T, M, N> transpose(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    matrix<T,M,N> result;\n")
-SLANG_RAW("    for(int r = 0; r < M; ++r)\n")
-SLANG_RAW("        for(int c = 0; c < N; ++c)\n")
-SLANG_RAW("            result[r][c] = x[c][r];\n")
-SLANG_RAW("    return result;\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int> matrix<T,M,N> transpose(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Truncate to integer\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType>\n")
-SLANG_RAW("T trunc(T x);\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("__target_intrinsic(glsl)\n")
-SLANG_RAW("vector<T, N> trunc(vector<T, N> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    VECTOR_MAP_UNARY(T, N, trunc, x);\n")
-SLANG_RAW("}\n")
-SLANG_RAW("\n")
-SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int>\n")
-SLANG_RAW("__target_intrinsic(hlsl)\n")
-SLANG_RAW("matrix<T, N, M> trunc(matrix<T, N, M> x)\n")
-SLANG_RAW("{\n")
-SLANG_RAW("    MATRIX_MAP_UNARY(T, N, M, trunc, x);\n")
-SLANG_RAW("}\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType> T trunc(T x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int> vector<T,N> trunc(vector<T,N> x);\n")
+SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M> trunc(matrix<T,N,M> x);\n")
 SLANG_RAW("\n")
 SLANG_RAW("// Shader model 6.0 stuff\n")
 SLANG_RAW("\n")
@@ -2215,6 +1476,7 @@ SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M>
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinIntegerType>\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupAnd($0)\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"_waveAnd(__activemask(), $0)\")\n")
 SLANG_RAW("T WaveActiveBitAnd(T expr);\n")
@@ -2223,6 +1485,7 @@ SLANG_RAW("__generic<T : __BuiltinIntegerType, let N : int, let M : int> matrix<
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinIntegerType>\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupOr($0)\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"_waveOr(__activemask(), $0)\")\n")
 SLANG_RAW("T WaveActiveBitOr(T expr);\n")
@@ -2231,6 +1494,7 @@ SLANG_RAW("__generic<T : __BuiltinIntegerType, let N : int, let M : int> matrix<
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinIntegerType>\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupXor($0)\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"_waveXor(__activemask(), $0)\")\n")
 SLANG_RAW("T WaveActiveBitXor(T expr);\n")
@@ -2239,6 +1503,7 @@ SLANG_RAW("__generic<T : __BuiltinIntegerType, let N : int, let M : int> matrix<
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupMax($0)\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"_waveMax(__activemask(), $0)\")\n")
 SLANG_RAW("T WaveActiveMax(T expr);\n")
@@ -2247,6 +1512,7 @@ SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matr
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupMin($0)\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"_waveMin(__activemask(), $0)\")\n")
 SLANG_RAW("T WaveActiveMin(T expr);\n")
@@ -2255,6 +1521,7 @@ SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matr
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupMul($0)\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"_waveProduct(__activemask(), $0)\")\n")
 SLANG_RAW("T WaveActiveProduct(T expr);\n")
@@ -2263,6 +1530,7 @@ SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matr
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupAdd($0)\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"_waveSum(__activemask(), $0)\")\n")
 SLANG_RAW("T WaveActiveSum(T expr);\n")
@@ -2271,6 +1539,7 @@ SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matr
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinType>\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_vote)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupAllEqual($0)\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"_waveAllEqual(__activemask(), $0)\")\n")
 SLANG_RAW("bool WaveActiveAllEqual(T value);\n")
@@ -2287,16 +1556,19 @@ SLANG_RAW("// With the Warp intrinsics there is no mask, and it's just the 'acti
 SLANG_RAW("// seems to be appropriate.\n")
 SLANG_RAW("\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_vote)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupAll($0)\") \n")
 SLANG_RAW("__target_intrinsic(cuda, \"(__all_sync(__activemask(), $0) != 0)\") \n")
 SLANG_RAW("bool WaveActiveAllTrue(bool condition);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_vote)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupAny($0)\") \n")
 SLANG_RAW("__target_intrinsic(cuda, \"(__any_sync(__activemask(), $0) != 0)\")\n")
 SLANG_RAW("bool WaveActiveAnyTrue(bool condition);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_ballot)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupBallot($0)\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"make_uint4(__ballot_sync(__activemask(), $0), 0, 0, 0)\")\n")
 SLANG_RAW("uint4 WaveActiveBallot(bool condition);\n")
@@ -2304,21 +1576,25 @@ SLANG_RAW("\n")
 SLANG_RAW("// TODO(JS): \n")
 SLANG_RAW("// subgroupBallotBitCount seems to take a uint4 parameter. \n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_ballot)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupBallotBitCount($0)\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"__popc(__ballot_sync(__activemask(), $0))\")\n")
 SLANG_RAW("uint WaveActiveCountBits(bool value);\n")
 SLANG_RAW("\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_basic)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"gl_SubgroupSize\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"(warpSize)\")\n")
 SLANG_RAW("uint WaveGetLaneCount();\n")
 SLANG_RAW("\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_basic)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"gl_SubgroupInvocationID\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"_getLaneId()\")\n")
 SLANG_RAW("uint WaveGetLaneIndex();\n")
 SLANG_RAW("\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_basic)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupElect()\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"_waveIsFirstLane()\")\n")
 SLANG_RAW("bool WaveIsFirstLane();\n")
@@ -2328,6 +1604,7 @@ SLANG_RAW("// that would mean different lanes having a different mask, and they
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupExcusiveMul($0)\")\n")
 SLANG_RAW("T WavePrefixProduct(T expr);\n")
 SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WavePrefixProduct(vector<T,N> expr);\n")
@@ -2335,6 +1612,7 @@ SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matr
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupExcusiveAdd($0)\")\n")
 SLANG_RAW("T WavePrefixSum(T expr);\n")
 SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WavePrefixSum(vector<T,N> expr);\n")
@@ -2355,6 +1633,7 @@ SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matr
 SLANG_RAW("\n")
 SLANG_RAW("// TODO(JS): This takes uvec4 parameter on GLSL\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_ballot)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupBallotExclusiveBitCount($0)\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"__popc(__ballot_sync(__activemask(), $0)  & _getLaneLtMask())\")\n")
 SLANG_RAW("uint WavePrefixCountBits(bool value);\n")
@@ -2371,6 +1650,7 @@ SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matr
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinType>\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_ballot)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgoupBroadcastFirst($0)\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"_waveReadFirst($0)\")\n")
 SLANG_RAW("T WaveReadLaneFirst(T expr);\n")
@@ -2379,6 +1659,7 @@ SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M>
 SLANG_RAW("\n")
 SLANG_RAW("__generic<T : __BuiltinType>\n")
 SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_ballot)\n")
+SLANG_RAW("__spirv_version(1.3)\n")
 SLANG_RAW("__target_intrinsic(glsl, \"subgroupBroadcast($0, $1)\")\n")
 SLANG_RAW("__target_intrinsic(cuda, \"__shfl_sync(SLANG_CUDA_WARP_MASK, $0, $1)\")\n")
 SLANG_RAW("T WaveReadLaneAt(T value, int lane);\n")
@@ -2390,6 +1671,37 @@ SLANG_RAW("// `typedef`s to help with the fact that HLSL has been sorta-kinda ca
 SLANG_RAW("typedef Texture2D texture2D;\n")
 SLANG_RAW("\n")
 
+// Component-wise multiplication ops
+for(auto op : binaryOps)
+{
+    switch (op.opCode)
+    {
+    default:
+        continue;
+
+    case kIROp_Mul:
+    case kCompoundIntrinsicOp_MulAssign:
+        break;
+    }
+
+    for (auto type : kBaseTypes)
+    {
+        if ((type.flags & op.flags) == 0)
+            continue;
+
+        char const* leftType = type.name;
+        char const* rightType = leftType;
+        char const* resultType = leftType;
+
+        char const* leftQual = "";
+        if(op.flags & ASSIGNMENT) leftQual = "in out ";
+
+        sb << "__generic<let N : int, let M : int> ";
+        sb << "__intrinsic_op(" << int(op.opCode) << ") matrix<" << resultType << ",N,M> operator" << op.opName << "(" << leftQual << "matrix<" << leftType << ",N,M> left, matrix<" << rightType << ",N,M> right);\n";
+    }
+}
+
+//
 
 // Buffer types
 
@@ -2440,7 +1752,7 @@ for (int aa = 0; aa < kBaseBufferAccessLevelCount; ++aa)
 
     sb << "};\n";
 }
-SLANG_RAW("#line 2367 \"hlsl.meta.slang\"")
+SLANG_RAW("#line 1679 \"hlsl.meta.slang\"")
 SLANG_RAW("\n")
 SLANG_RAW("\n")
 SLANG_RAW("\n")
diff --git a/source/slang/slang-compiler.cpp b/source/slang/slang-compiler.cpp
index 4880cc49d..0da0bce9c 100644
--- a/source/slang/slang-compiler.cpp
+++ b/source/slang/slang-compiler.cpp
@@ -6,6 +6,7 @@
 #include "../core/slang-string-util.h"
 #include "../core/slang-hex-dump-util.h"
 #include "../core/slang-riff.h"
+#include "../core/slang-type-text-util.h"
 
 #include "slang-check.h"
 #include "slang-compiler.h"
@@ -351,49 +352,6 @@ namespace Slang
         return UnownedStringSlice();
     }
 
-    static UnownedStringSlice _getPassThroughAsText(PassThroughMode mode)
-    {
-        switch (mode)
-        {
-            case PassThroughMode::None:     return UnownedStringSlice::fromLiteral("None");
-            case PassThroughMode::Dxc:      return UnownedStringSlice::fromLiteral("Dxc");
-            case PassThroughMode::Fxc:      return UnownedStringSlice::fromLiteral("Fxc");
-            case PassThroughMode::Glslang:  return UnownedStringSlice::fromLiteral("Glslang");
-            case PassThroughMode::Clang:    return UnownedStringSlice::fromLiteral("Clang");
-            case PassThroughMode::VisualStudio: return UnownedStringSlice::fromLiteral("VisualStudio");
-            case PassThroughMode::Gcc:      return UnownedStringSlice::fromLiteral("GCC");
-            case PassThroughMode::GenericCCpp:  return UnownedStringSlice::fromLiteral("Generic C/C++ Compiler");
-            default:                        return UnownedStringSlice::fromLiteral("Unknown");
-        }
-    }
-
-    static const struct
-    {
-        char const*     name;
-        SourceLanguage  sourceLanguage;
-    } kSourceLanguages[] =
-    {
-        { "slang", SourceLanguage::Slang },
-        { "hlsl", SourceLanguage::HLSL },
-        { "glsl", SourceLanguage::GLSL },
-        { "c", SourceLanguage::C },
-        { "cxx", SourceLanguage::CPP },
-        { "cuda", SourceLanguage::CUDA },
-    };
-
-    SourceLanguage findSourceLanguageByName(String const& name)
-    {
-        for (auto entry : kSourceLanguages)
-        {
-            if (name == entry.name)
-            {
-                return entry.sourceLanguage;
-            }
-        }
-
-        return SourceLanguage::Unknown;
-    };
-
     SlangResult checkExternalCompilerSupport(Session* session, PassThroughMode passThrough)
     {
         // Check if the type is supported on this compile
@@ -1209,7 +1167,7 @@ SlangResult dissassembleDXILUsingDXC(
 
         if (!compiler)
         {
-            auto compilerName = _getPassThroughAsText(downstreamCompiler);
+            auto compilerName = TypeTextUtil::getPassThroughAsHumanText((SlangPassThrough)downstreamCompiler);
             if (downstreamCompiler != PassThroughMode::None)
             {
                 sink->diagnose(SourceLoc(), Diagnostics::passThroughCompilerNotFound, compilerName);
diff --git a/source/slang/slang-diagnostic-defs.h b/source/slang/slang-diagnostic-defs.h
index 8bcdb9a09..de3b9dac3 100644
--- a/source/slang/slang-diagnostic-defs.h
+++ b/source/slang/slang-diagnostic-defs.h
@@ -207,6 +207,7 @@ DIAGNOSTIC(20002, Error, syntaxError, "syntax error.");
 DIAGNOSTIC(20004, Error, unexpectedTokenExpectedComponentDefinition, "unexpected token '$0', only component definitions are allowed in a shader scope.")
 DIAGNOSTIC(20008, Error, invalidOperator, "invalid operator '$0'.");
 DIAGNOSTIC(20011, Error, unexpectedColon, "unexpected ':'.")
+DIAGNOSTIC(20012, Error, invalidSPIRVVersion, "Expecting SPIR-V version number as either 'major.minor' or 'spirv_major_minor' (eg for SPIR-V 1.2, '1.2' or 'spirv_1_2')")
 
 //
 // 3xxxx - Semantic analysis
diff --git a/source/slang/slang-emit-glsl-extension-tracker.cpp b/source/slang/slang-emit-glsl-extension-tracker.cpp
index e1bbb3e46..6f2bc678c 100644
--- a/source/slang/slang-emit-glsl-extension-tracker.cpp
+++ b/source/slang/slang-emit-glsl-extension-tracker.cpp
@@ -26,6 +26,14 @@ void GLSLExtensionTracker::requireVersion(ProfileVersion version)
     }
 }
 
+void GLSLExtensionTracker::requireSPIRVVersion(SPIRVVersion version)
+{
+    if (asInteger(version) > asInteger(m_spirvVersion))
+    {
+        m_spirvVersion = version;
+    }
+}
+
 void GLSLExtensionTracker::requireBaseTypeExtension(BaseType baseType)
 {
     uint32_t bit = 1 << int(baseType);
diff --git a/source/slang/slang-emit-glsl-extension-tracker.h b/source/slang/slang-emit-glsl-extension-tracker.h
index 806e430e8..4b1b8b720 100644
--- a/source/slang/slang-emit-glsl-extension-tracker.h
+++ b/source/slang/slang-emit-glsl-extension-tracker.h
@@ -16,6 +16,7 @@ public:
     void requireExtension(const String& name);
     void requireVersion(ProfileVersion version);
     void requireBaseTypeExtension(BaseType baseType);
+    void requireSPIRVVersion(SPIRVVersion version);
 
     ProfileVersion getRequiredProfileVersion() const { return m_profileVersion; }
     const StringBuilder& getExtensionRequireLines() const { return m_extensionRequireLines; }
@@ -25,6 +26,8 @@ protected:
     HashSet<String> m_extensionsRequired;
     StringBuilder m_extensionRequireLines;
 
+    SPIRVVersion m_spirvVersion = makeSPIRVVersion(1, 2);
+
     ProfileVersion m_profileVersion = ProfileVersion::GLSL_110;
 
     static uint32_t _getFlag(BaseType baseType) { return uint32_t(1) << int(baseType); }
diff --git a/source/slang/slang-emit-glsl.cpp b/source/slang/slang-emit-glsl.cpp
index ae7a4130f..131c1882d 100644
--- a/source/slang/slang-emit-glsl.cpp
+++ b/source/slang/slang-emit-glsl.cpp
@@ -25,6 +25,11 @@ void GLSLSourceEmitter::_requireGLSLVersion(ProfileVersion version)
     m_glslExtensionTracker.requireVersion(version);
 }
 
+void GLSLSourceEmitter::_requireSPIRVVersion(SPIRVVersion version)
+{
+    m_glslExtensionTracker.requireSPIRVVersion(version);
+}
+
 void GLSLSourceEmitter::_requireGLSLVersion(int version)
 {
     switch (version)
@@ -1327,12 +1332,21 @@ void GLSLSourceEmitter::handleCallExprDecorationsImpl(IRInst* funcValue)
                 break;
 
             case kIROp_RequireGLSLExtensionDecoration:
+            {
                 _requireGLSLExtension(String(((IRRequireGLSLExtensionDecoration*)decoration)->getExtensionName()));
                 break;
-
+            }
             case kIROp_RequireGLSLVersionDecoration:
+            {
                 _requireGLSLVersion(int(((IRRequireGLSLVersionDecoration*)decoration)->getLanguageVersion()));
                 break;
+            }
+            case kIROp_RequireSPIRVVersionDecoration:
+            {
+                _requireSPIRVVersion(static_cast<IRRequireSPIRVVersionDecoration*>(decoration)->getSPIRVVersion());
+                break;
+            }
+
         }
     }
 }
diff --git a/source/slang/slang-emit-glsl.h b/source/slang/slang-emit-glsl.h
index 37d4054e9..fcd398748 100644
--- a/source/slang/slang-emit-glsl.h
+++ b/source/slang/slang-emit-glsl.h
@@ -60,6 +60,7 @@ protected:
 
     void _requireGLSLVersion(ProfileVersion version);
     void _requireGLSLVersion(int version);
+    void _requireSPIRVVersion(SPIRVVersion version);
 
         // Emit the `flat` qualifier if the underlying type
         // of the variable is an integer type.
diff --git a/source/slang/slang-ir-glsl-legalize.cpp b/source/slang/slang-ir-glsl-legalize.cpp
index 170027405..15a7e3c67 100644
--- a/source/slang/slang-ir-glsl-legalize.cpp
+++ b/source/slang/slang-ir-glsl-legalize.cpp
@@ -188,6 +188,11 @@ struct GLSLLegalizationContext
         glslExtensionTracker->requireExtension(name);
     }
 
+    void requireSPIRVVersion(SPIRVVersion version)
+    {
+        glslExtensionTracker->requireSPIRVVersion(version);
+    }
+
     void requireGLSLVersion(ProfileVersion version)
     {
         glslExtensionTracker->requireVersion(version);
diff --git a/source/slang/slang-ir-inst-defs.h b/source/slang/slang-ir-inst-defs.h
index 53a8ca89d..89fec618c 100644
--- a/source/slang/slang-ir-inst-defs.h
+++ b/source/slang/slang-ir-inst-defs.h
@@ -405,6 +405,7 @@ INST(HighLevelDeclDecoration,               highLevelDecl,          1, 0)
 
     INST(VulkanRayPayloadDecoration,        vulkanRayPayload,       0, 0)
     INST(VulkanHitAttributesDecoration,     vulkanHitAttributes,    0, 0)
+    INST(RequireSPIRVVersionDecoration,     requireSPIRVVersion,    1, 0)
     INST(RequireGLSLVersionDecoration,      requireGLSLVersion,     1, 0)
     INST(RequireGLSLExtensionDecoration,    requireGLSLExtension,   1, 0)
     INST(ReadNoneDecoration,                readNone,               0, 0)
diff --git a/source/slang/slang-ir-insts.h b/source/slang/slang-ir-insts.h
index 26b86107b..bacc3986c 100644
--- a/source/slang/slang-ir-insts.h
+++ b/source/slang/slang-ir-insts.h
@@ -178,6 +178,19 @@ struct IRRequireGLSLVersionDecoration : IRDecoration
     }
 };
 
+struct IRRequireSPIRVVersionDecoration : IRDecoration
+{
+    enum { kOp = kIROp_RequireSPIRVVersionDecoration };
+    IR_LEAF_ISA(RequireGLSLVersionDecoration)
+
+    IRConstant* getSPIRVVersionOperand() { return cast<IRConstant>(getOperand(0)); }
+
+    SPIRVVersion getSPIRVVersion()
+    {
+        return SPIRVVersion(getSPIRVVersionOperand()->value.intVal);
+    }
+};
+
 struct IRRequireGLSLExtensionDecoration : IRDecoration
 {
     enum { kOp = kIROp_RequireGLSLExtensionDecoration };
@@ -2110,6 +2123,11 @@ struct IRBuilder
         addDecoration(value, kIROp_RequireGLSLVersionDecoration, getIntValue(getIntType(), IRIntegerValue(version)));
     }
 
+    void addRequireSPIRVVersionDecoration(IRInst* value, SPIRVVersion version)
+    {
+        addDecoration(value, kIROp_RequireSPIRVVersionDecoration, getIntValue(getIntType(), IRIntegerValue(version)));
+    }
+
     void addPatchConstantFuncDecoration(IRInst* value, IRInst* patchConstantFunc)
     {
         addDecoration(value, kIROp_PatchConstantFuncDecoration, patchConstantFunc);
diff --git a/source/slang/slang-lower-to-ir.cpp b/source/slang/slang-lower-to-ir.cpp
index a4f756b14..57494f83f 100644
--- a/source/slang/slang-lower-to-ir.cpp
+++ b/source/slang/slang-lower-to-ir.cpp
@@ -6292,6 +6292,11 @@ struct DeclLoweringVisitor : DeclVisitor<DeclLoweringVisitor, LoweredValInfo>
         {
             getBuilder()->addRequireGLSLVersionDecoration(irFunc, Int(getIntegerLiteralValue(versionMod->versionNumberToken)));
         }
+        for (auto versionMod : decl->GetModifiersOfType<RequiredSPIRVVersionModifier>())
+        {
+            getBuilder()->addRequireSPIRVVersionDecoration(irFunc, versionMod->spirvVersion);
+        }
+
 
         if (auto attr = decl->FindModifier<InstanceAttribute>())
         {
diff --git a/source/slang/slang-modifier-defs.h b/source/slang/slang-modifier-defs.h
index 6fd660583..40ed945ae 100644
--- a/source/slang/slang-modifier-defs.h
+++ b/source/slang/slang-modifier-defs.h
@@ -82,6 +82,13 @@ FIELD(Token, versionNumberToken)
 END_SYNTAX_CLASS()
 
 
+// A modifier to tag something as an intrinsic that requires
+// a certain SPIRV version to be enabled when used. Specified as "major.minor"
+SYNTAX_CLASS(RequiredSPIRVVersionModifier, Modifier)
+FIELD(Token, token)
+FIELD(SPIRVVersion, spirvVersion)
+END_SYNTAX_CLASS()
+
 SIMPLE_SYNTAX_CLASS(InOutModifier, OutModifier)
 
 // `__ref` modifier for by-reference parameter passing
diff --git a/source/slang/slang-options.cpp b/source/slang/slang-options.cpp
index c97439b29..92ae637b1 100644
--- a/source/slang/slang-options.cpp
+++ b/source/slang/slang-options.cpp
@@ -659,7 +659,7 @@ struct OptionsParser
                     String name;
                     SLANG_RETURN_ON_FAIL(tryReadCommandLineArgument(sink, arg, &argCursor, argEnd, name));
 
-                    SourceLanguage sourceLanguage = findSourceLanguageByName(name);
+                    const SourceLanguage sourceLanguage = (SourceLanguage)TypeTextUtil::findSourceLanguage(name.getUnownedSlice());
 
                     if (sourceLanguage == SourceLanguage::Unknown)
                     {
diff --git a/source/slang/slang-parser.cpp b/source/slang/slang-parser.cpp
index 7a9d8d447..9dfe7ddc6 100644
--- a/source/slang/slang-parser.cpp
+++ b/source/slang/slang-parser.cpp
@@ -4747,6 +4747,88 @@ namespace Slang
 
         return modifier;
     }
+    static RefPtr<RefObject> parseSPIRVVersionModifier(Parser* parser, void* /*userData*/)
+    {
+        auto modifier = new RequiredSPIRVVersionModifier();
+
+        parser->ReadToken(TokenType::LParent);
+
+        UnownedStringSlice typeName;
+        UnownedStringSlice majorVersionText;
+        UnownedStringSlice minorVersionText;
+
+        // There is a legitimate question about how the version should be handled here.
+        switch (parser->tokenReader.peekTokenType())
+        {
+            case TokenType::FloatingPointLiteral:
+            {
+                modifier->token = parser->ReadToken(TokenType::FloatingPointLiteral);
+
+                List<UnownedStringSlice> split;
+                StringUtil::split(modifier->token.Content, '.', split);
+
+                if (split.getCount() != 2)
+                {
+                    parser->sink->diagnose(modifier->token, Diagnostics::invalidSPIRVVersion);
+                    return RefPtr<RefObject>();
+                }
+
+                majorVersionText = split[0];
+                minorVersionText = split[1];
+                break;
+            }
+            case TokenType::Identifier:
+            {
+                modifier->token = parser->ReadToken(TokenType::Identifier);
+
+                List<UnownedStringSlice> split;
+                StringUtil::split(modifier->token.Content, '_', split);
+
+                if (split.getCount() != 3)
+                {
+                    parser->sink->diagnose(modifier->token, Diagnostics::invalidSPIRVVersion);
+                    return RefPtr<RefObject>();
+                }
+
+                typeName = split[0];
+                majorVersionText = split[1];
+                minorVersionText = split[2];
+                break;
+
+            }
+            default:
+            {
+                parser->sink->diagnose(modifier->token, Diagnostics::invalidSPIRVVersion);
+                return RefPtr<RefObject>();
+            }
+        }
+
+        parser->ReadToken(TokenType::RParent);
+
+        // Currently we only support 'spirv' or no typename
+        if (!(typeName == UnownedStringSlice::fromLiteral("spirv") || typeName.getLength() == 0))
+        {
+            parser->sink->diagnose(modifier->token, Diagnostics::invalidSPIRVVersion);
+            return RefPtr<RefObject>();
+        }
+
+        Token intToken = modifier->token;
+        intToken.type = TokenType::IntegerLiteral;
+        intToken.Content = majorVersionText;
+
+        auto majorValue = getIntegerLiteralValue(intToken);
+        intToken.Content = minorVersionText;
+        auto minorValue = getIntegerLiteralValue(intToken);
+
+        if (minorValue < 0 || minorValue > 0xff || majorValue < 0)
+        {
+            parser->sink->diagnose(modifier->token, Diagnostics::invalidSPIRVVersion);
+            return RefPtr<RefObject>();
+        }
+
+        modifier->spirvVersion = makeSPIRVVersion(int(majorValue), int(minorValue));
+        return modifier;
+    }
 
     static RefPtr<RefObject> parseLayoutModifier(Parser* parser, void* /*userData*/)
     {
@@ -5038,6 +5120,7 @@ namespace Slang
         MODIFIER(__specialized_for_target,    parseSpecializedForTargetModifier);
         MODIFIER(__glsl_extension,  parseGLSLExtensionModifier);
         MODIFIER(__glsl_version,    parseGLSLVersionModifier);
+        MODIFIER(__spirv_version,   parseSPIRVVersionModifier);
 
         MODIFIER(__builtin_type,    parseBuiltinTypeModifier);
         MODIFIER(__magic_type,      parseMagicTypeModifier);
diff --git a/source/slang/slang-profile.h b/source/slang/slang-profile.h
index 9bc8197dc..b8d5257bb 100644
--- a/source/slang/slang-profile.h
+++ b/source/slang/slang-profile.h
@@ -104,11 +104,24 @@ namespace Slang
         RawVal raw = Unknown;
     };
 
+
+
     Stage findStageByName(String const& name);
 
     UnownedStringSlice getStageText(Stage stage);
 
-    SourceLanguage findSourceLanguageByName(String const& name);
+
+    // An enum to specify SPIR-V versions.
+    // For the moment they are only differentiated by version number, but it may be necessary to
+    // differentiate specific versions in other ways (such as the VK1.1 / SPIR-V combination in
+    // glslang). The enum is used to encode the representation, and vary as is needed in
+    // an implementation. 
+    enum class SPIRVVersion : uint32_t;
+
+    SLANG_INLINE SPIRVVersion makeSPIRVVersion(int major, int minor) { return SPIRVVersion((uint32_t(major) << 8) | uint32_t(minor)); }
+    SLANG_FORCE_INLINE Index getMajorVersion(SPIRVVersion version) { return Index((uint32_t(version) >> 8) & 0xff); }
+    SLANG_FORCE_INLINE Index getMinorVersion(SPIRVVersion version) { return Index(uint32_t(version) & 0xff); }
+    SLANG_FORCE_INLINE int32_t asInteger(SPIRVVersion version) { return int32_t(version); }
 }
 
 #endif
diff --git a/source/slang/slang-syntax.h b/source/slang/slang-syntax.h
index 660cdb2d5..7a7ea3603 100644
--- a/source/slang/slang-syntax.h
+++ b/source/slang/slang-syntax.h
@@ -32,6 +32,8 @@ namespace Slang
     class Parser;
     class SyntaxNode;
 
+    enum class SPIRVVersion : uint32_t;
+
     typedef RefPtr<RefObject> (*SyntaxParseCallback)(Parser* parser, void* userData);
 
     typedef unsigned int ConversionCost;