diff options
| -rw-r--r-- | prelude/slang-cuda-prelude.h | 6 | ||||
| -rw-r--r-- | source/slang/hlsl.meta.slang | 50 | ||||
| -rw-r--r-- | source/slang/hlsl.meta.slang.h | 52 |
3 files changed, 103 insertions, 5 deletions
diff --git a/prelude/slang-cuda-prelude.h b/prelude/slang-cuda-prelude.h index 1ca93d9d1..ce8e925a2 100644 --- a/prelude/slang-cuda-prelude.h +++ b/prelude/slang-cuda-prelude.h @@ -441,6 +441,12 @@ __forceinline__ __device__ uint32_t _getLaneId() } #endif +// Return mask of all the lanes less than the current lane +__forceinline__ __device__ int _getLaneLtMask() +{ + return (int(1) << _getLaneId()) - 1; +} + // Note! Note will return true if mask is 0, but thats okay, because there must be one // lane active to execute anything __inline__ __device__ bool _waveIsSingleLane(int mask) diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang index edb678ad6..417f4594d 100644 --- a/source/slang/hlsl.meta.slang +++ b/source/slang/hlsl.meta.slang @@ -1379,6 +1379,9 @@ __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> matrix<T,N,M // Shader model 6.0 stuff +// Information for GLSL wave/subgroup support +// https://github.com/KhronosGroup/GLSL/blob/master/extensions/khr/GL_KHR_shader_subgroup.txt + __generic<T : __BuiltinType> T QuadReadLaneAt(T sourceValue, uint quadLaneID); __generic<T : __BuiltinType, let N : int> vector<T,N> QuadReadLaneAt(vector<T,N> sourceValue, uint quadLaneID); __generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M> QuadReadLaneAt(matrix<T,N,M> sourceValue, uint quadLaneID); @@ -1396,48 +1399,64 @@ __generic<T : __BuiltinType, let N : int> vector<T,N> QuadReadAcrossDiagonal(vec __generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M> QuadReadAcrossDiagonal(matrix<T,N,M> localValue); __generic<T : __BuiltinIntegerType> +__glsl_extension(GL_KHR_shader_subgroup_arithmetic) +__target_intrinsic(glsl, "subgroupAnd($0)") __target_intrinsic(cuda, "_waveAnd(__activemask(), $0)") T WaveActiveBitAnd(T expr); __generic<T : __BuiltinIntegerType, let N : int> vector<T,N> WaveActiveBitAnd(vector<T,N> expr); __generic<T : __BuiltinIntegerType, let N : int, let M : int> matrix<T,N,M> WaveActiveBitAnd(matrix<T,N,M> expr); __generic<T : __BuiltinIntegerType> +__glsl_extension(GL_KHR_shader_subgroup_arithmetic) +__target_intrinsic(glsl, "subgroupOr($0)") __target_intrinsic(cuda, "_waveOr(__activemask(), $0)") T WaveActiveBitOr(T expr); __generic<T : __BuiltinIntegerType, let N : int> vector<T,N> WaveActiveBitOr(vector<T,N> expr); __generic<T : __BuiltinIntegerType, let N : int, let M : int> matrix<T,N,M> WaveActiveBitOr(matrix<T,N,M> expr); __generic<T : __BuiltinIntegerType> +__glsl_extension(GL_KHR_shader_subgroup_arithmetic) +__target_intrinsic(glsl, "subgroupXor($0)") __target_intrinsic(cuda, "_waveXor(__activemask(), $0)") T WaveActiveBitXor(T expr); __generic<T : __BuiltinIntegerType, let N : int> vector<T,N> WaveActiveBitXor(vector<T,N> expr); __generic<T : __BuiltinIntegerType, let N : int, let M : int> matrix<T,N,M> WaveActiveBitXor(matrix<T,N,M> expr); __generic<T : __BuiltinArithmeticType> +__glsl_extension(GL_KHR_shader_subgroup_arithmetic) +__target_intrinsic(glsl, "subgroupMax($0)") __target_intrinsic(cuda, "_waveMax(__activemask(), $0)") T WaveActiveMax(T expr); __generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WaveActiveMax(vector<T,N> expr); __generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WaveActiveMax(matrix<T,N,M> expr); __generic<T : __BuiltinArithmeticType> +__glsl_extension(GL_KHR_shader_subgroup_arithmetic) +__target_intrinsic(glsl, "subgroupMin($0)") __target_intrinsic(cuda, "_waveMin(__activemask(), $0)") T WaveActiveMin(T expr); __generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WaveActiveMin(vector<T,N> expr); __generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WaveActiveMin(matrix<T,N,M> expr); __generic<T : __BuiltinArithmeticType> +__glsl_extension(GL_KHR_shader_subgroup_arithmetic) +__target_intrinsic(glsl, "subgroupMul($0)") __target_intrinsic(cuda, "_waveProduct(__activemask(), $0)") T WaveActiveProduct(T expr); __generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WaveActiveProduct(vector<T,N> expr); __generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WaveActiveProduct(matrix<T,N,M> expr); __generic<T : __BuiltinArithmeticType> +__glsl_extension(GL_KHR_shader_subgroup_arithmetic) +__target_intrinsic(glsl, "subgroupAdd($0)") __target_intrinsic(cuda, "_waveSum(__activemask(), $0)") T WaveActiveSum(T expr); __generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WaveActiveSum(vector<T,N> expr); __generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WaveActiveSum(matrix<T,N,M> expr); __generic<T : __BuiltinType> +__glsl_extension(GL_KHR_shader_subgroup_vote) +__target_intrinsic(glsl, "subgroupAllEqual($0)") __target_intrinsic(cuda, "_waveAllEqual(__activemask(), $0)") bool WaveActiveAllEqual(T value); __generic<T : __BuiltinType, let N : int> vector<bool,N> WaveActiveAllEqual(vector<T,N> value); @@ -1452,24 +1471,40 @@ __generic<T : __BuiltinType, let N : int, let M : int> uint4 WaveMatch(matrix<T, // With the Warp intrinsics there is no mask, and it's just the 'active lanes'. So __activemask() // seems to be appropriate. +__glsl_extension(GL_KHR_shader_subgroup_vote) +__target_intrinsic(glsl, "subgroupAll($0)") __target_intrinsic(cuda, "(__all_sync(__activemask(), $0) != 0)") bool WaveActiveAllTrue(bool condition); + +__glsl_extension(GL_KHR_shader_subgroup_vote) +__target_intrinsic(glsl, "subgroupAny($0)") __target_intrinsic(cuda, "(__any_sync(__activemask(), $0) != 0)") bool WaveActiveAnyTrue(bool condition); +__glsl_extension(GL_KHR_shader_subgroup_ballot) +__target_intrinsic(glsl, "subgroupBallot($0)") __target_intrinsic(cuda, "make_uint4(__ballot_sync(__activemask(), $0), 0, 0, 0)") uint4 WaveActiveBallot(bool condition); +// TODO(JS): +// subgroupBallotBitCount seems to take a uint4 parameter. +__glsl_extension(GL_KHR_shader_subgroup_ballot) +__target_intrinsic(glsl, "subgroupBallotBitCount($0)") __target_intrinsic(cuda, "__popc(__ballot_sync(__activemask(), $0))") uint WaveActiveCountBits(bool value); +__glsl_extension(GL_KHR_shader_subgroup_basic) +__target_intrinsic(glsl, "gl_SubgroupSize") __target_intrinsic(cuda, "(warpSize)") uint WaveGetLaneCount(); +__glsl_extension(GL_KHR_shader_subgroup_basic) +__target_intrinsic(glsl, "gl_SubgroupInvocationID") __target_intrinsic(cuda, "_getLaneId()") uint WaveGetLaneIndex(); -// If there are no *active* lanes less than this one, we must be the lowest lane +__glsl_extension(GL_KHR_shader_subgroup_basic) +__target_intrinsic(glsl, "subgroupElect()") __target_intrinsic(cuda, "_waveIsFirstLane()") bool WaveIsFirstLane(); @@ -1477,11 +1512,15 @@ bool WaveIsFirstLane(); // that would mean different lanes having a different mask, and they all have to have the same mask. __generic<T : __BuiltinArithmeticType> +__glsl_extension(GL_KHR_shader_subgroup_arithmetic) +__target_intrinsic(glsl, "subgroupExcusiveMul($0)") T WavePrefixProduct(T expr); __generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WavePrefixProduct(vector<T,N> expr); __generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WavePrefixProduct(matrix<T,N,M> expr); __generic<T : __BuiltinArithmeticType> +__glsl_extension(GL_KHR_shader_subgroup_arithmetic) +__target_intrinsic(glsl, "subgroupExcusiveAdd($0)") T WavePrefixSum(T expr); __generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WavePrefixSum(vector<T,N> expr); __generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WavePrefixSum(matrix<T,N,M> expr); @@ -1499,7 +1538,10 @@ T WaveMultiPrefixBitXor(T expr); __generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WaveMultiPrefixBitXor(vector<T,N> expr); __generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WaveMultiPrefixBitXor(matrix<T,N,M> expr); -__target_intrinsic(cuda, "__popc(__ballot_sync(__activemask(), $0) & __lanemask_lt())") +// TODO(JS): This takes uvec4 parameter on GLSL +__glsl_extension(GL_KHR_shader_subgroup_ballot) +__target_intrinsic(glsl, "subgroupBallotExclusiveBitCount($0)") +__target_intrinsic(cuda, "__popc(__ballot_sync(__activemask(), $0) & _getLaneLtMask())") uint WavePrefixCountBits(bool value); uint WaveMultiPrefixCountBits(bool value, uint4 mask); @@ -1513,12 +1555,16 @@ __generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WaveMultiPrefixS __generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WaveMultiPrefixSum(matrix<T,N,M> value, uint4 mask); __generic<T : __BuiltinType> +__glsl_extension(GL_KHR_shader_subgroup_ballot) +__target_intrinsic(glsl, "subgoupBroadcastFirst($0)") __target_intrinsic(cuda, "_waveReadFirst($0)") T WaveReadLaneFirst(T expr); __generic<T : __BuiltinType, let N : int> vector<T,N> WaveReadLaneFirst(vector<T,N> expr); __generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M> WaveReadLaneFirst(matrix<T,N,M> expr); __generic<T : __BuiltinType> +__glsl_extension(GL_KHR_shader_subgroup_ballot) +__target_intrinsic(glsl, "subgroupBroadcast($0, $1)") __target_intrinsic(cuda, "__shfl_sync(SLANG_CUDA_WARP_MASK, $0, $1)") T WaveReadLaneAt(T value, int lane); __generic<T : __BuiltinType, let N : int> vector<T,N> WaveReadLaneAt(vector<T,N> value, int lane); diff --git a/source/slang/hlsl.meta.slang.h b/source/slang/hlsl.meta.slang.h index 16a3244ab..34bb15808 100644 --- a/source/slang/hlsl.meta.slang.h +++ b/source/slang/hlsl.meta.slang.h @@ -1455,6 +1455,9 @@ SLANG_RAW("__generic<T : __BuiltinFloatingPointType, let N : int, let M : int> m SLANG_RAW("\n") SLANG_RAW("// Shader model 6.0 stuff\n") SLANG_RAW("\n") +SLANG_RAW("// Information for GLSL wave/subgroup support\n") +SLANG_RAW("// https://github.com/KhronosGroup/GLSL/blob/master/extensions/khr/GL_KHR_shader_subgroup.txt\n") +SLANG_RAW("\n") SLANG_RAW("__generic<T : __BuiltinType> T QuadReadLaneAt(T sourceValue, uint quadLaneID);\n") SLANG_RAW("__generic<T : __BuiltinType, let N : int> vector<T,N> QuadReadLaneAt(vector<T,N> sourceValue, uint quadLaneID);\n") SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M> QuadReadLaneAt(matrix<T,N,M> sourceValue, uint quadLaneID);\n") @@ -1472,48 +1475,64 @@ SLANG_RAW("__generic<T : __BuiltinType, let N : int> vector<T,N> QuadReadAcrossD SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M> QuadReadAcrossDiagonal(matrix<T,N,M> localValue);\n") SLANG_RAW("\n") SLANG_RAW("__generic<T : __BuiltinIntegerType>\n") +SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\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") SLANG_RAW("__generic<T : __BuiltinIntegerType, let N : int> vector<T,N> WaveActiveBitAnd(vector<T,N> expr);\n") SLANG_RAW("__generic<T : __BuiltinIntegerType, let N : int, let M : int> matrix<T,N,M> WaveActiveBitAnd(matrix<T,N,M> expr);\n") SLANG_RAW("\n") SLANG_RAW("__generic<T : __BuiltinIntegerType>\n") +SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\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") SLANG_RAW("__generic<T : __BuiltinIntegerType, let N : int> vector<T,N> WaveActiveBitOr(vector<T,N> expr);\n") SLANG_RAW("__generic<T : __BuiltinIntegerType, let N : int, let M : int> matrix<T,N,M> WaveActiveBitOr(matrix<T,N,M> expr);\n") SLANG_RAW("\n") SLANG_RAW("__generic<T : __BuiltinIntegerType>\n") +SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\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") SLANG_RAW("__generic<T : __BuiltinIntegerType, let N : int> vector<T,N> WaveActiveBitXor(vector<T,N> expr);\n") SLANG_RAW("__generic<T : __BuiltinIntegerType, let N : int, let M : int> matrix<T,N,M> WaveActiveBitXor(matrix<T,N,M> expr);\n") SLANG_RAW("\n") SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n") +SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\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") SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WaveActiveMax(vector<T,N> expr);\n") SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WaveActiveMax(matrix<T,N,M> expr);\n") SLANG_RAW("\n") SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n") +SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\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") SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WaveActiveMin(vector<T,N> expr);\n") SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WaveActiveMin(matrix<T,N,M> expr);\n") SLANG_RAW("\n") SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n") +SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\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") SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WaveActiveProduct(vector<T,N> expr);\n") SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WaveActiveProduct(matrix<T,N,M> expr);\n") SLANG_RAW("\n") SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n") +SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\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") SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WaveActiveSum(vector<T,N> expr);\n") SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WaveActiveSum(matrix<T,N,M> expr);\n") SLANG_RAW("\n") SLANG_RAW("__generic<T : __BuiltinType>\n") +SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_vote)\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") SLANG_RAW("__generic<T : __BuiltinType, let N : int> vector<bool,N> WaveActiveAllEqual(vector<T,N> value);\n") @@ -1528,24 +1547,40 @@ SLANG_RAW("// https://devblogs.nvidia.com/using-cuda-warp-level-primitives/\n") SLANG_RAW("// With the Warp intrinsics there is no mask, and it's just the 'active lanes'. So __activemask()\n") SLANG_RAW("// seems to be appropriate.\n") SLANG_RAW("\n") +SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_vote)\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("__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("__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") 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("__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("__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("__target_intrinsic(glsl, \"gl_SubgroupInvocationID\")\n") SLANG_RAW("__target_intrinsic(cuda, \"_getLaneId()\")\n") SLANG_RAW("uint WaveGetLaneIndex();\n") SLANG_RAW("\n") -SLANG_RAW("// If there are no *active* lanes less than this one, we must be the lowest lane\n") +SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_basic)\n") +SLANG_RAW("__target_intrinsic(glsl, \"subgroupElect()\")\n") SLANG_RAW("__target_intrinsic(cuda, \"_waveIsFirstLane()\")\n") SLANG_RAW("bool WaveIsFirstLane();\n") SLANG_RAW("\n") @@ -1553,11 +1588,15 @@ SLANG_RAW("// TODO(JS): We cannot calculate prefix sums using a mask of __active SLANG_RAW("// that would mean different lanes having a different mask, and they all have to have the same mask.\n") SLANG_RAW("\n") SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n") +SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\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") SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WavePrefixProduct(matrix<T,N,M> expr);\n") SLANG_RAW("\n") SLANG_RAW("__generic<T : __BuiltinArithmeticType>\n") +SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_arithmetic)\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") SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WavePrefixSum(matrix<T,N,M> expr);\n") @@ -1575,7 +1614,10 @@ SLANG_RAW("T WaveMultiPrefixBitXor(T expr);\n") SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WaveMultiPrefixBitXor(vector<T,N> expr);\n") SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WaveMultiPrefixBitXor(matrix<T,N,M> expr);\n") SLANG_RAW("\n") -SLANG_RAW("__target_intrinsic(cuda, \"__popc(__ballot_sync(__activemask(), $0) & __lanemask_lt())\")\n") +SLANG_RAW("// TODO(JS): This takes uvec4 parameter on GLSL\n") +SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_ballot)\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") SLANG_RAW("\n") SLANG_RAW("uint WaveMultiPrefixCountBits(bool value, uint4 mask);\n") @@ -1589,12 +1631,16 @@ SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int> vector<T,N> WaveM SLANG_RAW("__generic<T : __BuiltinArithmeticType, let N : int, let M : int> matrix<T,N,M> WaveMultiPrefixSum(matrix<T,N,M> value, uint4 mask);\n") SLANG_RAW("\n") SLANG_RAW("__generic<T : __BuiltinType>\n") +SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_ballot)\n") +SLANG_RAW("__target_intrinsic(glsl, \"subgoupBroadcastFirst($0)\")\n") SLANG_RAW("__target_intrinsic(cuda, \"_waveReadFirst($0)\")\n") SLANG_RAW("T WaveReadLaneFirst(T expr);\n") SLANG_RAW("__generic<T : __BuiltinType, let N : int> vector<T,N> WaveReadLaneFirst(vector<T,N> expr);\n") SLANG_RAW("__generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M> WaveReadLaneFirst(matrix<T,N,M> expr);\n") SLANG_RAW("\n") SLANG_RAW("__generic<T : __BuiltinType>\n") +SLANG_RAW("__glsl_extension(GL_KHR_shader_subgroup_ballot)\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") SLANG_RAW("__generic<T : __BuiltinType, let N : int> vector<T,N> WaveReadLaneAt(vector<T,N> value, int lane);\n") @@ -1686,7 +1732,7 @@ for (int aa = 0; aa < kBaseBufferAccessLevelCount; ++aa) sb << "};\n"; } -SLANG_RAW("#line 1613 \"hlsl.meta.slang\"") +SLANG_RAW("#line 1659 \"hlsl.meta.slang\"") SLANG_RAW("\n") SLANG_RAW("\n") SLANG_RAW("\n") |