From 5b0b8436123aa2faa9b682ed45efe2bd7edbf01b Mon Sep 17 00:00:00 2001 From: jsmall-nvidia Date: Fri, 27 Mar 2020 14:49:41 -0400 Subject: 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. --- source/slang/hlsl.meta.slang | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) (limited to 'source') 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 WaveReadLaneFirst(matrix 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 __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 -__spirv_version(1.3) +__spirv_version(1.4) __target_intrinsic(glsl, "subgroupBroadcast($0, $1)") __target_intrinsic(cuda, "_waveReadLaneAtMultiple($0, $1)") vector WaveReadLaneAt(vector value, int lane); -- cgit v1.2.3