Support for WaveReadLaneAt with dynamic (but uniform across Wave) on Vk by enabling VK1.4. (#1297)

Fixed wave-lane-at.slang test to test with laneId that is uniform across the Wave.

1 files changed, 5 insertions, 2 deletions

diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang
index fe348c60d..739b8579d 100644
--- a/source/slang/hlsl.meta.slang
+++ b/source/slang/hlsl.meta.slang
@@ -2722,14 +2722,17 @@ matrix<T,N,M> WaveReadLaneFirst(matrix<T,N,M> expr);
 
 // NOTE! On GLSL based targets the lane index *must* be a compile time expression!
 // See https://github.com/KhronosGroup/GLSL/blob/master/extensions/khr/GL_KHR_shader_subgroup.txt
+// It is allowed to be 'dynamically uniform within the subgroup' if it's SPIR-V 1.4.
+// TODO(JS): For now we'll use 1.4, but aim for the future for the compiler to determine
+// if the line the is compile constant, and reduce requirement to 1.3
 __generic<T : __BuiltinType>
 __glsl_extension(GL_KHR_shader_subgroup_ballot)
-__spirv_version(1.3)
+__spirv_version(1.4)
 __target_intrinsic(glsl, "subgroupBroadcast($0, $1)")
 __target_intrinsic(cuda, "__shfl_sync(__activemask(), $0, $1)")
 T WaveReadLaneAt(T value, int lane);
 __generic<T : __BuiltinType, let N : int>
-__spirv_version(1.3)
+__spirv_version(1.4)
 __target_intrinsic(glsl, "subgroupBroadcast($0, $1)")
 __target_intrinsic(cuda, "_waveReadLaneAtMultiple($0, $1)")
 vector<T,N> WaveReadLaneAt(vector<T,N> value, int lane);