From 6f43b2698a99cc4f4bb4e905749fb87f24bf391b Mon Sep 17 00:00:00 2001 From: jsmall-nvidia Date: Fri, 27 Mar 2020 18:35:06 -0400 Subject: WaveBroadcastAt/WaveShuffle (#1299) * Support for WaveReadLaneAt with dynamic (but uniform across Wave) on Vk by enabling VK1.4. Fixed wave-lane-at.slang test to test with laneId that is uniform across the Wave. * Added WaveShuffle intrinsic. Test for WaveShuffle intrinsic. * Added some documentation on WaveShuffle * Fix that version required for subgroupBroadcast to be non constexpr is actually 1.5 * Added WaveBroadcastLaneAt Documented WaveShuffle/BroadcastLaneAt/ReadLaneAt * Update docs around WaveBroadcast/Read/Shuffle. Use '_waveShuffle` as name in CUDA prelude to better describe it's more flexible behavior. --- prelude/slang-cuda-prelude.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) (limited to 'prelude') diff --git a/prelude/slang-cuda-prelude.h b/prelude/slang-cuda-prelude.h index dcc585b9c..b5d8b3788 100644 --- a/prelude/slang-cuda-prelude.h +++ b/prelude/slang-cuda-prelude.h @@ -851,7 +851,7 @@ __inline__ __device__ T _waveReadFirstMultiple(T inVal) } template -__inline__ __device__ T _waveReadLaneAtMultiple(T inVal, int lane) +__inline__ __device__ T _waveShuffleMultiple(T inVal, int lane) { typedef typename ElementTypeTrait::Type ElemType; const size_t count = sizeof(T) / sizeof(ElemType); -- cgit v1.2.3