diff options
| author | ArielG-NV <159081215+ArielG-NV@users.noreply.github.com> | 2024-04-19 23:18:40 -0400 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2024-04-19 20:18:40 -0700 |
| commit | f9bcad35562c1f08638e6d3eb397d370d7d2f8f8 (patch) | |
| tree | 4e2a993689209bd5b597263922af03cb87d07c3d /source/slang/hlsl.meta.slang | |
| parent | 2da28c50d9c3699692eccde4b86d0b8d2323e55c (diff) | |
Initial pass to add capability declarations to stdlib intrinsics. (#3912)
Diffstat (limited to 'source/slang/hlsl.meta.slang')
| -rw-r--r-- | source/slang/hlsl.meta.slang | 1435 |
1 files changed, 986 insertions, 449 deletions
diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang index cdd08b5d7..cb01e9e68 100644 --- a/source/slang/hlsl.meta.slang +++ b/source/slang/hlsl.meta.slang @@ -2,6 +2,9 @@ typedef uint UINT; +__intrinsic_op($(kIROp_RequireGLSLExtension)) +void __requireGLSLExtension(String extensionName); + __intrinsic_op($(kIROp_FloatCast)) T __floatCast<T, U>(U v); @@ -37,24 +40,30 @@ struct GLSLShaderStorageBuffer {} __generic<T,L:IBufferDataLayout> __intrinsic_op($(kIROp_StructuredBufferGetDimensions)) +[require(cpp_cuda_glsl_hlsl_metal_spirv, structuredbuffer_rw)] uint2 __structuredBufferGetDimensions(AppendStructuredBuffer<T,L> buffer); __generic<T,L:IBufferDataLayout> __intrinsic_op($(kIROp_StructuredBufferGetDimensions)) +[require(cpp_cuda_glsl_hlsl_metal_spirv, structuredbuffer_rw)] uint2 __structuredBufferGetDimensions(ConsumeStructuredBuffer<T,L> buffer); __intrinsic_op($(kIROp_StructuredBufferGetDimensions)) +[require(cpp_cuda_glsl_hlsl_metal_spirv, structuredbuffer)] uint2 __structuredBufferGetDimensions<T,L:IBufferDataLayout>(StructuredBuffer<T,L> buffer); __intrinsic_op($(kIROp_StructuredBufferGetDimensions)) +[require(cpp_cuda_glsl_hlsl_metal_spirv, structuredbuffer_rw)] uint2 __structuredBufferGetDimensions<T,L:IBufferDataLayout>(RWStructuredBuffer<T,L> buffer); __intrinsic_op($(kIROp_StructuredBufferGetDimensions)) +[require(cpp_cuda_glsl_hlsl_metal_spirv, structuredbuffer_rw)] uint2 __structuredBufferGetDimensions<T,L:IBufferDataLayout>(RasterizerOrderedStructuredBuffer<T,L> buffer); __generic<T, L:IBufferDataLayout=DefaultDataLayout> __magic_type(HLSLAppendStructuredBufferType) __intrinsic_type($(kIROp_HLSLAppendStructuredBufferType)) +[require(cpp_cuda_glsl_hlsl_spirv, appendstructuredbuffer)] struct AppendStructuredBuffer { __intrinsic_op($(kIROp_StructuredBufferAppend)) @@ -73,6 +82,7 @@ struct AppendStructuredBuffer __magic_type(HLSLByteAddressBufferType) __intrinsic_type($(kIROp_HLSLByteAddressBufferType)) +[require(cpp_cuda_glsl_hlsl_metal_spirv, byteaddressbuffer)] struct ByteAddressBuffer { [__readNone] @@ -213,69 +223,89 @@ struct __TextureImpl<T, Shape: __ITextureShape, let isArray:int, let isMS:int, l // Combined texture sampler specific functions __target_intrinsic(glsl, "texture($0, $1)") +[require(glsl, texture_sm_4_1)] float __glsl_texture<TSampler, TCoord>(TSampler s, TCoord value); +__glsl_extension(GL_EXT_texture_shadow_lod) __target_intrinsic(glsl, "texture($0, $1)") +[require(glsl, texture_shadowlod)] float __glsl_texture_1d_shadow<TSampler, TCoord>(TSampler s, TCoord value); +__glsl_extension(GL_EXT_texture_shadow_lod) __target_intrinsic(glsl, "texture($0, $1, $2)") +[require(glsl, texture_shadowlod)] float __glsl_texture_3d_array_shadow<TSampler, TCoord>(TSampler s, TCoord value, float compare); __glsl_extension(GL_EXT_texture_shadow_lod) - __target_intrinsic(glsl, "textureOffset($0, $1, $2)") +__target_intrinsic(glsl, "textureOffset($0, $1, $2)") +[require(glsl, texture_sm_4_1)] float __glsl_texture_offset<TSampler, TCoord, TOffset>( TSampler s, TCoord value, constexpr TOffset offset); __glsl_extension(GL_EXT_texture_shadow_lod) __target_intrinsic(glsl, "textureOffset($0, $1, $2)") +[require(glsl, texture_shadowlod)] float __glsl_texture_offset_1d_shadow<TSampler, TCoord, TOffset>(TSampler s, TCoord value, constexpr TOffset offset); __glsl_extension(GL_EXT_texture_shadow_lod) __target_intrinsic(glsl, "textureLod($0, $1, 0)") +[require(glsl, texture_sm_4_1)] float __glsl_texture_level_zero<TSampler, TCoord>(TSampler s, TCoord value); __glsl_extension(GL_EXT_texture_shadow_lod) __target_intrinsic(glsl, "textureLod($0, $1, 0)") +[require(glsl, texture_shadowlod)] float __glsl_texture_level_zero_1d_shadow<TSampler, TCoord>(TSampler s, TCoord value); __glsl_extension(GL_EXT_texture_shadow_lod) __target_intrinsic(glsl, "textureLodOffset($0, $1, 0, $2)") +[require(glsl, texture_shadowlod)] float __glsl_texture_offset_level_zero<TSampler, TCoord, TOffset>(TSampler s, TCoord value, constexpr TOffset offset); __glsl_extension(GL_EXT_texture_shadow_lod) __target_intrinsic(glsl, "textureLodOffset($0, $1, 0, $2)") +[require(glsl, texture_shadowlod)] float __glsl_texture_offset_level_zero_1d_shadow<TSampler, TCoord, TOffset>(TSampler s, TCoord value, constexpr TOffset offset); __target_intrinsic(glsl, "texture($p, $2)") +[require(glsl, texture_sm_4_1)] float __glsl_texture<TTexture, TCoord>(TTexture t, SamplerComparisonState s, TCoord value); +__glsl_extension(GL_EXT_texture_shadow_lod) __target_intrinsic(glsl, "texture($p, $2)") +[require(glsl, texture_shadowlod)] float __glsl_texture_1d_shadow<TTexture, TCoord>(TTexture t, SamplerComparisonState s, TCoord value); +__glsl_extension(GL_EXT_texture_shadow_lod) __target_intrinsic(glsl, "texture($p, $2, $3)") +[require(glsl, texture_shadowlod)] float __glsl_texture_3d_array_shadow<TTexture, TCoord>(TTexture t, SamplerComparisonState s, TCoord value, float compare); -__glsl_extension(GL_EXT_texture_shadow_lod) __target_intrinsic(glsl, "textureOffset($p, $2, $3)") +[require(glsl, texture_sm_4_1)] float __glsl_texture_offset<TTexture, TCoord, TOffset>(TTexture t,SamplerComparisonState s, TCoord value, constexpr TOffset offset); __glsl_extension(GL_EXT_texture_shadow_lod) __target_intrinsic(glsl, "textureOffset($p, $2, $3)") +[require(glsl, texture_shadowlod)] float __glsl_texture_offset_1d_shadow<TTexture, TCoord, TOffset>(TTexture t,SamplerComparisonState s, TCoord value, constexpr TOffset offset); -__glsl_extension(GL_EXT_texture_shadow_lod) __target_intrinsic(glsl, "textureLod($p, $2, 0)") +[require(glsl, texture_sm_4_1)] float __glsl_texture_level_zero<TTexture, TCoord>(TTexture t,SamplerComparisonState s, TCoord value); __glsl_extension(GL_EXT_texture_shadow_lod) __target_intrinsic(glsl, "textureLod($p, $2, 0)") +[require(glsl, texture_shadowlod)] float __glsl_texture_level_zero_1d_shadow<TTexture, TCoord>(TTexture t,SamplerComparisonState s, TCoord value); __glsl_extension(GL_EXT_texture_shadow_lod) __target_intrinsic(glsl, "textureLodOffset($p, $2, 0, $3)") +[require(glsl, texture_shadowlod)] float __glsl_texture_offset_level_zero<TTexture, TCoord, TOffset>(TTexture t,SamplerComparisonState s, TCoord value, constexpr TOffset offset); __glsl_extension(GL_EXT_texture_shadow_lod) __target_intrinsic(glsl, "textureLodOffset($p, $2, 0, $3)") +[require(glsl, texture_shadowlod)] float __glsl_texture_offset_level_zero_1d_shadow<TTexture, TCoord, TOffset>(TTexture t,SamplerComparisonState s, TCoord value, constexpr TOffset offset); @@ -297,6 +327,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> [ForceInline] [__readNone] + [require(glsl_hlsl_spirv, texture_querylod)] float CalculateLevelOfDetail(TextureCoord location) { __target_switch @@ -317,6 +348,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> [ForceInline] [__readNone] + [require(glsl_hlsl_spirv, texture_querylod)] float CalculateLevelOfDetailUnclamped(TextureCoord location) { __target_switch @@ -336,6 +368,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> } [__readNone] + [require(cpp_cuda_glsl_hlsl_spirv, texture_sm_4_1_fragment)] T Sample(vector<float, Shape.dimensions+isArray> location) { __target_switch @@ -385,6 +418,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> [__readNone] __glsl_extension(GL_ARB_sparse_texture_clamp) + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_fragment)] T Sample(vector<float, Shape.dimensions+isArray> location, vector<int, Shape.planeDimensions> offset, float clamp) { __target_switch @@ -406,6 +440,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> [__readNone] __target_intrinsic(hlsl) + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_fragment)] T Sample(vector<float, Shape.dimensions+isArray> location, vector<int, Shape.planeDimensions> offset, float clamp, out uint status) { status = 0; @@ -413,6 +448,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> } [__readNone] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_fragment)] T SampleBias(vector<float, Shape.dimensions+isArray> location, float bias) { __target_switch @@ -433,6 +469,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> } [__readNone] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_fragment)] T SampleBias(vector<float, Shape.dimensions+isArray> location, float bias, constexpr vector<int, Shape.planeDimensions> offset) { __target_switch @@ -453,6 +490,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> [__readNone] [ForceInline] + [require(glsl_hlsl_spirv, texture_shadowlod)] float SampleCmp(vector<float, Shape.dimensions+isArray> location, float compareValue) { __target_switch @@ -482,6 +520,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> [__readNone] [ForceInline] + [require(glsl_hlsl_spirv, texture_shadowlod)] float SampleCmpLevelZero(vector<float, Shape.dimensions+isArray> location, float compareValue) { __target_switch @@ -508,6 +547,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> [__readNone] [ForceInline] + [require(glsl_hlsl_spirv, texture_shadowlod)] float SampleCmp(vector<float, Shape.dimensions+isArray> location, float compareValue, constexpr vector<int, Shape.planeDimensions> offset) { __target_switch @@ -533,6 +573,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> [__readNone] [ForceInline] + [require(glsl_hlsl_spirv, texture_shadowlod)] float SampleCmpLevelZero(vector<float, Shape.dimensions+isArray> location, float compareValue, constexpr vector<int, Shape.planeDimensions> offset) { __target_switch @@ -558,6 +599,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> } [__readNone] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1)] T SampleGrad(vector<float, Shape.dimensions+isArray> location, vector<float, Shape.dimensions> gradX, vector<float, Shape.dimensions> gradY) { __target_switch @@ -578,6 +620,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> } [__readNone] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1)] T SampleGrad(vector<float, Shape.dimensions+isArray> location, vector<float, Shape.dimensions> gradX, vector<float, Shape.dimensions> gradY, constexpr vector<int, Shape.dimensions> offset) { __target_switch @@ -598,6 +641,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> __glsl_extension(GL_ARB_sparse_texture_clamp) [__readNone] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1)] T SampleGrad(vector<float, Shape.dimensions+isArray> location, vector<float, Shape.dimensions> gradX, vector<float, Shape.dimensions> gradY, constexpr vector<int, Shape.dimensions> offset, float lodClamp) { __target_switch @@ -619,6 +663,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> [__readNone] [ForceInline] + [require(cpp_cuda_glsl_hlsl_spirv, texture_sm_4_1)] T SampleLevel(vector<float, Shape.dimensions+isArray> location, float level) { __target_switch @@ -670,6 +715,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,1,format> [__readNone] [ForceInline] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1)] T SampleLevel(vector<float, Shape.dimensions+isArray> location, float level, constexpr vector<int, Shape.planeDimensions> offset) { __target_switch @@ -698,6 +744,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,access,isShadow,0,forma [__readNone] [ForceInline] + [require(glsl_hlsl_spirv, texture_querylod)] float CalculateLevelOfDetail(SamplerState s, TextureCoord location) { __target_switch @@ -716,6 +763,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,access,isShadow,0,forma [__readNone] [ForceInline] + [require(glsl_hlsl_spirv, texture_querylod)] float CalculateLevelOfDetailUnclamped(SamplerState s, TextureCoord location) { __target_switch @@ -737,6 +785,7 @@ __generic<T:IFloat, Shape: __ITextureShape, let isArray:int, let isMS:int, let s extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format> { [__readNone] + [require(cpp_cuda_glsl_hlsl_spirv, texture_sm_4_1_fragment)] T Sample(SamplerState s, vector<float, Shape.dimensions+isArray> location) { __target_switch @@ -788,6 +837,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format> } [__readNone] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_fragment)] T Sample(SamplerState s, vector<float, Shape.dimensions+isArray> location, constexpr vector<int, Shape.planeDimensions> offset) { __target_switch @@ -809,6 +859,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format> [__readNone] __glsl_extension(GL_ARB_sparse_texture_clamp) + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_fragment)] T Sample(SamplerState s, vector<float, Shape.dimensions+isArray> location, constexpr vector<int, Shape.planeDimensions> offset, float clamp) { __target_switch @@ -838,6 +889,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format> } [__readNone] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_fragment)] T SampleBias(SamplerState s, vector<float, Shape.dimensions+isArray> location, float bias) { __target_switch @@ -858,6 +910,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format> } [__readNone] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_fragment)] T SampleBias(SamplerState s, vector<float, Shape.dimensions+isArray> location, float bias, constexpr vector<int, Shape.planeDimensions> offset) { __target_switch @@ -878,6 +931,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format> } [__readNone] [ForceInline] + [require(glsl_hlsl_spirv, texture_shadowlod)] float SampleCmp(SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, float compareValue) { __target_switch @@ -907,6 +961,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format> } [__readNone] [ForceInline] + [require(glsl_hlsl_spirv, texture_shadowlod)] float SampleCmpLevelZero(SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, float compareValue) { __target_switch @@ -933,6 +988,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format> } [__readNone] [ForceInline] + [require(glsl_hlsl_spirv, texture_shadowlod)] float SampleCmp(SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, float compareValue, constexpr vector<int, Shape.planeDimensions> offset) { __target_switch @@ -958,6 +1014,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format> } [__readNone] [ForceInline] + [require(glsl_hlsl_spirv, texture_shadowlod)] float SampleCmpLevelZero(SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, float compareValue, constexpr vector<int, Shape.planeDimensions> offset) { __target_switch @@ -984,6 +1041,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format> } [__readNone] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1)] T SampleGrad(SamplerState s, vector<float, Shape.dimensions+isArray> location, vector<float, Shape.dimensions> gradX, vector<float, Shape.dimensions> gradY) { __target_switch @@ -1004,6 +1062,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format> } [__readNone] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1)] T SampleGrad(SamplerState s, vector<float, Shape.dimensions+isArray> location, vector<float, Shape.dimensions> gradX, vector<float, Shape.dimensions> gradY, constexpr vector<int, Shape.dimensions> offset) { __target_switch @@ -1026,6 +1085,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format> __glsl_extension(GL_ARB_sparse_texture_clamp) [__readNone] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1)] T SampleGrad(SamplerState s, vector<float, Shape.dimensions+isArray> location, vector<float, Shape.dimensions> gradX, vector<float, Shape.dimensions> gradY, constexpr vector<int, Shape.dimensions> offset, float lodClamp) { __target_switch @@ -1048,6 +1108,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format> [__readNone] [ForceInline] + [require(cpp_cuda_glsl_hlsl_spirv, texture_sm_4_1)] T SampleLevel(SamplerState s, vector<float, Shape.dimensions+isArray> location, float level) { __target_switch @@ -1100,6 +1161,7 @@ extension __TextureImpl<T,Shape,isArray,isMS,sampleCount,0,isShadow,0,format> [__readNone] [ForceInline] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1)] T SampleLevel(SamplerState s, vector<float, Shape.dimensions+isArray> location, float level, constexpr vector<int, Shape.planeDimensions> offset) { __target_switch @@ -1157,6 +1219,7 @@ ${{{{ __generic<T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let isCombined:int, let format:int> extension __TextureImpl<T,Shape,isArray,1,sampleCount,access,isShadow,isCombined,format> { + [require(cpp_cuda_glsl_hlsl_spirv, texture_sm_4_1_vertex_fragment_geometry)] float2 GetSamplePosition(int s); } @@ -1166,6 +1229,7 @@ Array<T,4> __makeArray<T>(T v0, T v1, T v2, T v3); // Gather for scalar textures. __generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int> [ForceInline] +[require(glsl_spirv, GLSL_400)] vector<TElement,4> __glsl_gather(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerState s, vector<float, Shape.dimensions+isArray> location, int component) { __target_switch @@ -1181,6 +1245,7 @@ vector<TElement,4> __glsl_gather(__TextureImpl<T, Shape, isArray, 0, sampleCount } __generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int> [ForceInline] +[require(glsl_spirv, GLSL_400)] vector<TElement,4> __glsl_gather(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, int component) { __target_switch @@ -1195,6 +1260,7 @@ vector<TElement,4> __glsl_gather(__TextureImpl<T, Shape, isArray, 0, sampleCount } __generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int> [ForceInline] +[require(glsl_spirv, GLSL_400)] vector<TElement,4> __glsl_gather_offset(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerState s, constexpr vector<float, Shape.dimensions+isArray> location, constexpr vector<int, Shape.planeDimensions> offset, int component) { __target_switch @@ -1210,6 +1276,7 @@ vector<TElement,4> __glsl_gather_offset(__TextureImpl<T, Shape, isArray, 0, samp } __generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int> [ForceInline] +[require(glsl_spirv, GLSL_400)] vector<TElement,4> __glsl_gather_offset(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, constexpr vector<int, Shape.planeDimensions> offset, int component) { __target_switch @@ -1224,6 +1291,7 @@ vector<TElement,4> __glsl_gather_offset(__TextureImpl<T, Shape, isArray, 0, samp } __generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int> [ForceInline] +[require(glsl_spirv, GLSL_400)] vector<TElement,4> __glsl_gather_offsets(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerState s, vector<float, Shape.dimensions+isArray> location, constexpr vector<int, Shape.planeDimensions> offset1, constexpr vector<int, Shape.planeDimensions> offset2, @@ -1246,6 +1314,7 @@ vector<TElement,4> __glsl_gather_offsets(__TextureImpl<T, Shape, isArray, 0, sam } __generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int> [ForceInline] +[require(glsl_spirv, GLSL_400)] vector<TElement,4> __glsl_gather_offsets(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, constexpr vector<int, Shape.planeDimensions> offset1, constexpr vector<int, Shape.planeDimensions> offset2, @@ -1267,6 +1336,7 @@ vector<TElement,4> __glsl_gather_offsets(__TextureImpl<T, Shape, isArray, 0, sam } __generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int> [ForceInline] +[require(glsl_spirv, GLSL_400)] vector<TElement,4> __glsl_gatherCmp(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, TElement compareValue) { __target_switch @@ -1282,6 +1352,7 @@ vector<TElement,4> __glsl_gatherCmp(__TextureImpl<T, Shape, isArray, 0, sampleCo } __generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int> [ForceInline] +[require(glsl_spirv, GLSL_400)] vector<TElement,4> __glsl_gatherCmp(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, TElement compareValue) { __target_switch @@ -1296,6 +1367,7 @@ vector<TElement,4> __glsl_gatherCmp(__TextureImpl<T, Shape, isArray, 0, sampleCo } __generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int> [ForceInline] +[require(glsl_spirv, GLSL_400)] vector<TElement,4> __glsl_gatherCmp_offset(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, TElement compareValue, constexpr vector<int, Shape.planeDimensions> offset) { __target_switch @@ -1311,6 +1383,7 @@ vector<TElement,4> __glsl_gatherCmp_offset(__TextureImpl<T, Shape, isArray, 0, s } __generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int> [ForceInline] +[require(glsl_spirv, GLSL_400)] vector<TElement,4> __glsl_gatherCmp_offset(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, TElement compareValue, constexpr vector<int, Shape.planeDimensions> offset) { __target_switch @@ -1325,6 +1398,7 @@ vector<TElement,4> __glsl_gatherCmp_offset(__TextureImpl<T, Shape, isArray, 0, s } __generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int> [ForceInline] +[require(glsl_spirv, GLSL_400)] vector<TElement,4> __glsl_gatherCmp_offsets(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture, SamplerComparisonState s, vector<float, Shape.dimensions+isArray> location, TElement compareValue, vector<int, Shape.planeDimensions> offset1, vector<int, Shape.planeDimensions> offset2, @@ -1346,6 +1420,7 @@ vector<TElement,4> __glsl_gatherCmp_offsets(__TextureImpl<T, Shape, isArray, 0, } __generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:int, let access:int, let isShadow:int, let format:int> [ForceInline] +[require(glsl_spirv, GLSL_400)] vector<TElement,4> __glsl_gatherCmp_offsets(__TextureImpl<T, Shape, isArray, 0, sampleCount, access, isShadow, 1, format> sampler, vector<float, Shape.dimensions+isArray> location, TElement compareValue, vector<int, Shape.planeDimensions> offset1, vector<int, Shape.planeDimensions> offset2, @@ -1396,6 +1471,7 @@ ${{{{ auto glslComponent = (isCmp ? "" :glslComponentNames[componentId == 0 ? 0 : componentId - 1]); }}}} [ForceInline] + [require(glsl_hlsl_spirv, texture_gather)] vector<T,4> Gather$(cmp)$(componentName)($(samplerStateType)$(samplerStateParam) vector<float, Shape.dimensions+isArray> location $(cmpParam)) { __target_switch @@ -1407,6 +1483,7 @@ ${{{{ } } [ForceInline] + [require(glsl_hlsl_spirv, texture_gather)] vector<T,4> Gather$(cmp)$(componentName)($(samplerStateType)$(samplerStateParam) vector<float, Shape.dimensions+isArray> location $(cmpParam), constexpr vector<int, Shape.planeDimensions> offset) { __target_switch @@ -1418,6 +1495,7 @@ ${{{{ } } [ForceInline] + [require(glsl_hlsl_spirv, texture_gather)] vector<T,4> Gather$(cmp)$(componentName)($(samplerStateType)$(samplerStateParam) vector<float, Shape.dimensions+isArray> location $(cmpParam), constexpr vector<int, Shape.planeDimensions> offset1, constexpr vector<int, Shape.planeDimensions> offset2, @@ -1453,6 +1531,7 @@ extension __TextureImpl<T,Shape,isArray,0,sampleCount,0,isShadow,isCombined,form __glsl_extension(GL_EXT_samplerless_texture_functions) [__readNone] + [require(glsl, texture_sm_4_1_samplerless)] T __glsl_load(vector<int, Shape.dimensions+isArray> location) { __intrinsic_asm "$ctexelFetch($0, ($1), 0)$z"; @@ -1461,6 +1540,7 @@ extension __TextureImpl<T,Shape,isArray,0,sampleCount,0,isShadow,isCombined,form __glsl_extension(GL_EXT_samplerless_texture_functions) [__readNone] [ForceInline] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_samplerless)] T Load(vector<int, Shape.dimensions+isArray+1> location) { __target_switch @@ -1497,6 +1577,7 @@ extension __TextureImpl<T,Shape,isArray,0,sampleCount,0,isShadow,isCombined,form __glsl_extension(GL_EXT_samplerless_texture_functions) [__readNone] [ForceInline] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_samplerless)] T Load(vector<int, Shape.dimensions+isArray+1> location, constexpr vector<int, Shape.planeDimensions> offset) { __target_switch @@ -1544,6 +1625,7 @@ extension __TextureImpl<T,Shape,isArray,0,sampleCount,0,isShadow,isCombined,form __glsl_extension(GL_EXT_samplerless_texture_functions) [__readNone] [ForceInline] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_samplerless)] get { __target_switch @@ -1587,6 +1669,7 @@ extension __TextureImpl<T,Shape,isArray,1,sampleCount,0,isShadow,isCombined,form __glsl_extension(GL_EXT_samplerless_texture_functions) [__readNone] [ForceInline] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_samplerless)] T Load(vector<int, Shape.dimensions+isArray> location, int sampleIndex) { __target_switch @@ -1619,6 +1702,7 @@ extension __TextureImpl<T,Shape,isArray,1,sampleCount,0,isShadow,isCombined,form [__readNone] [ForceInline] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1)] T Load(vector<int, Shape.dimensions + isArray + 1> locationAndSampleIndex) { return Load(__vectorReshape<Shape.dimensions + isArray>(locationAndSampleIndex), locationAndSampleIndex[Shape.dimensions + isArray]); @@ -1627,6 +1711,7 @@ extension __TextureImpl<T,Shape,isArray,1,sampleCount,0,isShadow,isCombined,form __glsl_extension(GL_EXT_samplerless_texture_functions) [__readNone] [ForceInline] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_samplerless)] T Load(vector<int, Shape.dimensions+isArray> location, int sampleIndex, constexpr vector<int, Shape.planeDimensions> offset) { __target_switch @@ -1671,6 +1756,7 @@ extension __TextureImpl<T,Shape,isArray,1,sampleCount,0,isShadow,isCombined,form __glsl_extension(GL_EXT_samplerless_texture_functions) [__readNone] [ForceInline] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_samplerless)] get { __target_switch @@ -1680,7 +1766,6 @@ extension __TextureImpl<T,Shape,isArray,1,sampleCount,0,isShadow,isCombined,form __intrinsic_asm "($0).sample[$2][$1]"; case glsl: case spirv: - case cuda: return Load(location, sampleIndex); } } @@ -1700,6 +1785,7 @@ extension __TextureImpl<T,Shape,isArray,0,sampleCount,$(access),isShadow, 0,form { [__readNone] [ForceInline] + [require(cpp_cuda_glsl_hlsl_spirv, texture_sm_4_1)] T Load(vector<int, Shape.dimensions+isArray> location) { __target_switch @@ -1749,6 +1835,7 @@ extension __TextureImpl<T,Shape,isArray,0,sampleCount,$(access),isShadow, 0,form [__readNone] [ForceInline] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1)] T Load(vector<int, Shape.dimensions+isArray> location, vector<int, Shape.dimensions+isArray> offset) { __target_switch @@ -1769,6 +1856,7 @@ extension __TextureImpl<T,Shape,isArray,0,sampleCount,$(access),isShadow, 0,form [__readNone] [ForceInline] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1)] T Load(vector<int, Shape.dimensions+isArray> location, vector<int, Shape.dimensions+isArray> offset, out uint status) { __target_switch @@ -1782,6 +1870,7 @@ extension __TextureImpl<T,Shape,isArray,0,sampleCount,$(access),isShadow, 0,form } } + [require(glsl, texture_sm_4_1)] void __glslImageStore(vector<int, Shape.dimensions+isArray> location, T value) { __intrinsic_asm "imageStore($0, $1, $V2)"; @@ -1791,6 +1880,7 @@ extension __TextureImpl<T,Shape,isArray,0,sampleCount,$(access),isShadow, 0,form { [__readNone] [ForceInline] + [require(cpp_cuda_glsl_hlsl_spirv, texture_sm_4_1)] get { __target_switch @@ -1807,6 +1897,7 @@ extension __TextureImpl<T,Shape,isArray,0,sampleCount,$(access),isShadow, 0,form [nonmutating] [ForceInline] + [require(cpp_cuda_glsl_hlsl_spirv, texture_sm_4_1)] set(T newValue) { __target_switch @@ -1866,6 +1957,7 @@ extension __TextureImpl<T,Shape,isArray,1,sampleCount,$(access),isShadow, 0,form { [__readNone] [ForceInline] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_compute_fragment)] T Load(vector<int, Shape.dimensions+isArray> location, int sampleIndex) { __target_switch @@ -1886,6 +1978,7 @@ extension __TextureImpl<T,Shape,isArray,1,sampleCount,$(access),isShadow, 0,form [__readNone] [ForceInline] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_compute_fragment)] T Load(vector<int, Shape.dimensions+isArray> location, int sampleIndex, vector<int, Shape.dimensions+isArray> offset) { __target_switch @@ -1919,6 +2012,7 @@ extension __TextureImpl<T,Shape,isArray,1,sampleCount,$(access),isShadow, 0,form } } + [require(glsl, texture_sm_4_1_compute_fragment)] void __glslImageStore(vector<int, Shape.dimensions+isArray> location, int sampleIndex, T value) { __intrinsic_asm "imageStore($0, $1, $2, $V3)"; @@ -1928,6 +2022,7 @@ extension __TextureImpl<T,Shape,isArray,1,sampleCount,$(access),isShadow, 0,form { [__readNone] [ForceInline] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_compute_fragment)] get { __target_switch @@ -1937,13 +2032,13 @@ extension __TextureImpl<T,Shape,isArray,1,sampleCount,$(access),isShadow, 0,form __intrinsic_asm "$0.sample[$2][$1]"; case glsl: case spirv: - case cuda: return Load(location, sampleIndex); } } [nonmutating] [ForceInline] + [require(cpp_glsl_hlsl_spirv, texture_sm_4_1_compute_fragment)] set(T newValue) { __target_switch @@ -2007,7 +2102,13 @@ ${{{{ if (shape == kStdlibShapeIndex3D && isArray == 1) continue; const char* textureTypeName = isCombined ? "Sampler" : "Texture"; + + StringBuilder requireStringBuilder; + if (shape == kStdlibShapeIndexCube) + requireStringBuilder << "[require(any_target, texture_cube)]"; + auto requireString = requireStringBuilder.toString(); }}}} +$(requireString) typealias $(accessPrefix[access])$(textureTypeName)$(shapeTypeNames[shape])$(arrayPostFix[isArray])$(msPostFix[isMS])<T=float4, let sampleCount:int=0, let format:int=0> = __TextureImpl<T, __Shape$(shapeTypeNames[shape]), $(isArray), $(isMS), sampleCount, $(access), 0, $(isCombined), format>; ${{{{ } @@ -2021,6 +2122,8 @@ ${{{{ __glsl_version(430) __glsl_extension(GL_EXT_shader_atomic_float) +[ForceInline] +[require(glsl_spirv, atomic_glsl_float1)] float __atomicAdd(__ref float value, float amount) { __target_switch @@ -2038,6 +2141,8 @@ float __atomicAdd(__ref float value, float amount) __glsl_version(430) __glsl_extension(GL_NV_shader_atomic_fp16_vector) +[ForceInline] +[require(glsl_spirv, atomic_glsl_halfvec)] half2 __atomicAdd(__ref half2 value, half2 amount) { __target_switch @@ -2056,18 +2161,23 @@ half2 __atomicAdd(__ref half2 value, half2 amount) // Helper for hlsl, using NVAPI __target_intrinsic(hlsl, "NvInterlockedAddUint64($0, $1, $2)") [__requiresNVAPI] +[require(hlsl, atomic_hlsl_nvapi)] uint2 __atomicAdd(RWByteAddressBuffer buf, uint offset, uint2); // atomic add for hlsl using SM6.6 __target_intrinsic(hlsl, "$0.InterlockedAdd64($1, $2, $3)") +[require(hlsl, atomic_hlsl_sm_6_6)] void __atomicAdd(RWByteAddressBuffer buf, uint offset, int64_t value, out int64_t originalValue); __target_intrinsic(hlsl, "$0.InterlockedAdd64($1, $2, $3)") +[require(hlsl, atomic_hlsl_sm_6_6)] void __atomicAdd(RWByteAddressBuffer buf, uint offset, uint64_t value, out uint64_t originalValue); // Int versions require glsl 4.30 // https://www.khronos.org/registry/OpenGL-Refpages/gl4/html/atomicAdd.xhtml __glsl_version(430) +[ForceInline] +[require(glsl_spirv, atomic_glsl)] int __atomicAdd(__ref int value, int amount) { __target_switch @@ -2082,6 +2192,8 @@ int __atomicAdd(__ref int value, int amount) } __glsl_version(430) +[ForceInline] +[require(glsl_spirv, atomic_glsl)] uint __atomicAdd(__ref uint value, uint amount) { __target_switch @@ -2097,6 +2209,8 @@ uint __atomicAdd(__ref uint value, uint amount) __glsl_version(430) __glsl_extension(GL_EXT_shader_atomic_int64) +[ForceInline] +[require(glsl_spirv, atomic_glsl_int64)] int64_t __atomicAdd(__ref int64_t value, int64_t amount) { __target_switch @@ -2114,6 +2228,8 @@ int64_t __atomicAdd(__ref int64_t value, int64_t amount) __target_intrinsic(glsl, "atomicAdd($0, $1)") __glsl_version(430) __glsl_extension(GL_EXT_shader_atomic_int64) +[ForceInline] +[require(glsl_spirv, atomic_glsl_int64)] uint64_t __atomicAdd(__ref uint64_t value, uint64_t amount) { __target_switch @@ -2134,16 +2250,21 @@ uint64_t __atomicAdd(__ref uint64_t value, uint64_t amount) __target_intrinsic(hlsl, "NvInterlockedCompareExchangeUint64($0, $1, $2, $3)") [__requiresNVAPI] +[require(hlsl, atomic_hlsl_nvapi)] uint2 __cas(RWByteAddressBuffer buf, uint offset, uint2 compareValue, uint2 value); // CAS using SM6.6 __target_intrinsic(hlsl, "$0.InterlockedCompareExchange64($1, $2, $3, $4)") +[require(hlsl, atomic_hlsl_sm_6_6)] void __cas(RWByteAddressBuffer buf, uint offset, in int64_t compare_value, in int64_t value, out int64_t original_value); __target_intrinsic(hlsl, "$0.InterlockedCompareExchange64($1, $2, $3, $4)") +[require(hlsl, atomic_hlsl_sm_6_6)] void __cas(RWByteAddressBuffer buf, uint offset, in uint64_t compare_value, in uint64_t value, out uint64_t original_value); __glsl_version(430) __glsl_extension(GL_EXT_shader_atomic_int64) +[ForceInline] +[require(glsl_spirv, atomic_glsl_int64)] int64_t __cas(__ref int64_t ioValue, int64_t compareValue, int64_t newValue) { __target_switch @@ -2160,6 +2281,8 @@ int64_t __cas(__ref int64_t ioValue, int64_t compareValue, int64_t newValue) __glsl_version(430) __glsl_extension(GL_EXT_shader_atomic_int64) +[ForceInline] +[require(glsl_spirv, atomic_glsl_int64)] uint64_t __cas(__ref uint64_t ioValue, uint64_t compareValue, uint64_t newValue) { __target_switch @@ -2178,10 +2301,13 @@ uint64_t __cas(__ref uint64_t ioValue, uint64_t compareValue, uint64_t newValue) __target_intrinsic(hlsl, "NvInterlockedMaxUint64($0, $1, $2)") [__requiresNVAPI] +[require(hlsl, atomic_hlsl_nvapi)] uint2 __atomicMax(RWByteAddressBuffer buf, uint offset, uint2 value); __glsl_version(430) __glsl_extension(GL_EXT_shader_atomic_int64) +[ForceInline] +[require(glsl_spirv, atomic_glsl_int64)] uint64_t __atomicMax(__ref uint64_t ioValue, uint64_t value) { __target_switch @@ -2197,7 +2323,9 @@ uint64_t __atomicMax(__ref uint64_t ioValue, uint64_t value) } __glsl_version(430) -__glsl_extension(GL_EXT_shader_atomic_float_min_max) +__glsl_extension(GL_EXT_shader_atomic_float2) +[ForceInline] +[require(glsl_spirv, atomic_glsl_float2)] float __atomicMax(__ref float ioValue, float value) { __target_switch @@ -2214,7 +2342,9 @@ float __atomicMax(__ref float ioValue, float value) } __glsl_version(430) -__glsl_extension(GL_EXT_shader_atomic_float_min_max) +__glsl_extension(GL_EXT_shader_atomic_float2) +[ForceInline] +[require(glsl_spirv, atomic_glsl_float2)] half __atomicMax(__ref half ioValue, half value) { __target_switch @@ -2234,10 +2364,13 @@ half __atomicMax(__ref half ioValue, half value) __target_intrinsic(hlsl, "NvInterlockedMinUint64($0, $1, $2)") [__requiresNVAPI] +[require(hlsl, atomic_hlsl_nvapi)] uint2 __atomicMin(RWByteAddressBuffer buf, uint offset, uint2 value); __glsl_version(430) __glsl_extension(GL_EXT_shader_atomic_int64) +[ForceInline] +[require(glsl_spirv, atomic_glsl_int64)] uint64_t __atomicMin(__ref uint64_t ioValue, uint64_t value) { __target_switch @@ -2253,7 +2386,9 @@ uint64_t __atomicMin(__ref uint64_t ioValue, uint64_t value) } __glsl_version(430) -__glsl_extension(GL_EXT_shader_atomic_float_min_max) +__glsl_extension(GL_EXT_shader_atomic_float2) +[ForceInline] +[require(glsl_spirv, atomic_glsl_float2)] float __atomicMin(__ref float ioValue, float value) { __target_switch @@ -2270,7 +2405,9 @@ float __atomicMin(__ref float ioValue, float value) } __glsl_version(430) -__glsl_extension(GL_EXT_shader_atomic_float_min_max) +__glsl_extension(GL_EXT_shader_atomic_float2) +[ForceInline] +[require(glsl_spirv, atomic_glsl_float2)] half __atomicMin(__ref half ioValue, half value) { __target_switch @@ -2290,10 +2427,13 @@ half __atomicMin(__ref half ioValue, half value) __target_intrinsic(hlsl, "NvInterlockedAndUint64($0, $1, $2)") [__requiresNVAPI] +[require(hlsl, atomic_hlsl_nvapi)] uint2 __atomicAnd(RWByteAddressBuffer buf, uint offset, uint2 value); __glsl_version(430) __glsl_extension(GL_EXT_shader_atomic_int64) +[ForceInline] +[require(glsl_spirv, atomic_glsl_int64)] uint64_t __atomicAnd(__ref uint64_t ioValue, uint64_t value) { __target_switch @@ -2312,10 +2452,13 @@ uint64_t __atomicAnd(__ref uint64_t ioValue, uint64_t value) __target_intrinsic(hlsl, "NvInterlockedOrUint64($0, $1, $2)") [__requiresNVAPI] +[require(hlsl, atomic_hlsl_nvapi)] uint2 __atomicOr(RWByteAddressBuffer buf, uint offset, uint2 value); __glsl_version(430) __glsl_extension(GL_EXT_shader_atomic_int64) +[ForceInline] +[require(glsl_spirv, atomic_glsl_int64)] uint64_t __atomicOr(__ref uint64_t ioValue, uint64_t value) { __target_switch @@ -2334,10 +2477,13 @@ uint64_t __atomicOr(__ref uint64_t ioValue, uint64_t value) __target_intrinsic(hlsl, "NvInterlockedXorUint64($0, $1, $2)") [__requiresNVAPI] +[require(hlsl, atomic_hlsl_nvapi)] uint2 __atomicXor(RWByteAddressBuffer buf, uint offset, uint2 value); __glsl_version(430) __glsl_extension(GL_EXT_shader_atomic_int64) +[ForceInline] +[require(glsl_spirv, atomic_glsl_int64)] uint64_t __atomicXor(__ref uint64_t ioValue, uint64_t value) { __target_switch @@ -2356,10 +2502,13 @@ uint64_t __atomicXor(__ref uint64_t ioValue, uint64_t value) __target_intrinsic(hlsl, "NvInterlockedExchangeUint64($0, $1, $2)") [__requiresNVAPI] +[require(hlsl, atomic_hlsl_nvapi)] uint2 __atomicExchange(RWByteAddressBuffer buf, uint offset, uint2 value); __glsl_version(430) __glsl_extension(GL_EXT_shader_atomic_int64) +[ForceInline] +[require(glsl_spirv, atomic_glsl_int64)] uint64_t __atomicExchange(__ref uint64_t ioValue, uint64_t value) { __target_switch @@ -2376,11 +2525,13 @@ uint64_t __atomicExchange(__ref uint64_t ioValue, uint64_t value) // Conversion between uint64_t and uint2 +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] uint2 __asuint2(uint64_t i) { return uint2(uint(i), uint(uint64_t(i) >> 32)); } +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] uint64_t __asuint64(uint2 i) { return (uint64_t(i.y) << 32) | i.x; @@ -2389,18 +2540,23 @@ uint64_t __asuint64(uint2 i) // __intrinsic_op($(kIROp_ByteAddressBufferLoad)) +[require(cpp_cuda_glsl_hlsl_metal_spirv, byteaddressbuffer)] T __byteAddressBufferLoad<T>(ByteAddressBuffer buffer, int offset); __intrinsic_op($(kIROp_ByteAddressBufferLoad)) +[require(cpp_cuda_glsl_hlsl_metal_spirv, byteaddressbuffer_rw)] T __byteAddressBufferLoad<T>(RWByteAddressBuffer buffer, int offset); __intrinsic_op($(kIROp_ByteAddressBufferLoad)) +[require(cpp_cuda_glsl_hlsl_metal_spirv, byteaddressbuffer_rw)] T __byteAddressBufferLoad<T>(RasterizerOrderedByteAddressBuffer buffer, int offset); __intrinsic_op($(kIROp_ByteAddressBufferStore)) +[require(cpp_cuda_glsl_hlsl_metal_spirv, byteaddressbuffer_rw)] void __byteAddressBufferStore<T>(RWByteAddressBuffer buffer, int offset, T value); __intrinsic_op($(kIROp_ByteAddressBufferStore)) +[require(cpp_cuda_glsl_hlsl_metal_spirv, byteaddressbuffer_rw)] void __byteAddressBufferStore<T>(RasterizerOrderedByteAddressBuffer buffer, int offset, T value); __generic<T, L:IBufferDataLayout=DefaultDataLayout> @@ -2423,15 +2579,18 @@ struct StructuredBuffer __target_intrinsic(glsl, "$0._data[$1]") __target_intrinsic(spirv, "%addr = OpAccessChain resultType*StorageBuffer resultId _0 const(int, 0) _1; OpLoad resultType resultId %addr;") [__readNone] + [require(cpp_cuda_glsl_hlsl_spirv, structuredbuffer)] T Load(int location); __intrinsic_op($(kIROp_StructuredBufferLoadStatus)) + [require(cpp_cuda_glsl_hlsl_spirv, structuredbuffer)] T Load(int location, out uint status); __subscript(uint index) -> T { [__readNone] __intrinsic_op($(kIROp_StructuredBufferLoad)) + [require(cpp_cuda_glsl_hlsl_spirv, structuredbuffer)] get; }; }; @@ -2439,6 +2598,7 @@ struct StructuredBuffer __generic<T, L:IBufferDataLayout=DefaultDataLayout> __magic_type(HLSLConsumeStructuredBufferType) __intrinsic_type($(kIROp_HLSLConsumeStructuredBufferType)) +[require(cpp_cuda_glsl_hlsl_spirv, consumestructuredbuffer)] struct ConsumeStructuredBuffer { __intrinsic_op($(kIROp_StructuredBufferConsume)) @@ -2456,6 +2616,7 @@ struct ConsumeStructuredBuffer }; __generic<T, let N : int> +[require(glsl_hlsl_spirv, hull)] __magic_type(HLSLInputPatchType) __intrinsic_type($(kIROp_HLSLInputPatchType)) struct InputPatch @@ -2464,6 +2625,7 @@ struct InputPatch }; __generic<T, let N : int> +[require(glsl_hlsl_spirv, shader_stages_domain_hull)] __magic_type(HLSLOutputPatchType) __intrinsic_type($(kIROp_HLSLOutputPatchType)) struct OutputPatch @@ -2494,11 +2656,13 @@ struct $(item.name) __target_intrinsic(cpp) __target_intrinsic(cuda) [__unsafeForceInlineEarly] + [require(cpp_cuda_glsl_hlsl_spirv, structuredbuffer_rw)] void GetDimensions(out uint dim); [__unsafeForceInlineEarly] __specialized_for_target(spirv) __specialized_for_target(glsl) + [require(cpp_cuda_glsl_hlsl_spirv, structuredbuffer_rw)] void GetDimensions(out uint dim) { dim = __structuredBufferGetDimensions(__getEquivalentStructuredBuffer<uint>(this)).x*4; @@ -2506,6 +2670,7 @@ struct $(item.name) __target_intrinsic(hlsl) [__NoSideEffect] + [require(cpp_cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] uint Load(int location) { return __byteAddressBufferLoad<uint>(this, location); @@ -2516,6 +2681,7 @@ struct $(item.name) __target_intrinsic(hlsl) [__NoSideEffect] + [require(cpp_cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] uint2 Load2(int location) { return __byteAddressBufferLoad<uint2>(this, location); @@ -2526,6 +2692,7 @@ struct $(item.name) __target_intrinsic(hlsl) [__NoSideEffect] + [require(cpp_cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] uint3 Load3(int location) { return __byteAddressBufferLoad<uint3>(this, location); @@ -2536,6 +2703,7 @@ struct $(item.name) __target_intrinsic(hlsl) [__NoSideEffect] + [require(cpp_cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] uint4 Load4(int location) { return __byteAddressBufferLoad<uint4>(this, location); @@ -2545,6 +2713,7 @@ struct $(item.name) uint4 Load4(int location, out uint status); [__NoSideEffect] + [require(cpp_cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] T Load<T>(int location) { return __byteAddressBufferLoad<T>(this, location); @@ -2573,6 +2742,7 @@ ${{{{ __cuda_sm_version(2.0) [__requiresNVAPI] [ForceInline] + [require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda_float1)] void InterlockedAddF32(uint byteAddress, float valueToAdd, out float originalValue) { __target_switch @@ -2592,6 +2762,7 @@ ${{{{ // FP16x2 [__requiresNVAPI] [ForceInline] + [require(hlsl, atomic_hlsl_nvapi)] uint _NvInterlockedAddFp16x2(uint byteAddress, uint fp16x2Value) { __target_switch @@ -2642,6 +2813,8 @@ ${{{{ [__requiresNVAPI] [ForceInline] __cuda_sm_version(2.0) + [ForceInline] + [require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda_float1)] void InterlockedAddF32(uint byteAddress, float valueToAdd) { __target_switch @@ -2661,6 +2834,8 @@ ${{{{ // Int64 Add [ForceInline] __cuda_sm_version(6.0) + [ForceInline] + [require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda6_int64)] void InterlockedAddI64(uint byteAddress, int64_t valueToAdd, out int64_t originalValue) { __target_switch @@ -2680,10 +2855,12 @@ ${{{{ // Without returning original value __cuda_sm_version(6.0) __target_intrinsic(cuda, "atomicAdd($0._getPtrAt<uint64_t>($1), $2)") + [require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda6_int64)] void InterlockedAddI64(uint byteAddress, int64_t valueToAdd); [ForceInline] __specialized_for_target(hlsl) + [ForceInline] void InterlockedAddI64(uint byteAddress, int64_t valueToAdd) { __atomicAdd(this, byteAddress, __asuint2(valueToAdd)); @@ -2692,6 +2869,7 @@ ${{{{ [ForceInline] __specialized_for_target(glsl) __specialized_for_target(spirv) + [ForceInline] void InterlockedAddI64(uint byteAddress, int64_t valueToAdd) { let buf = __getEquivalentStructuredBuffer<int64_t>(this); @@ -2701,9 +2879,11 @@ ${{{{ // Cas uint64_t __target_intrinsic(cuda, "(*$4 = atomicCAS($0._getPtrAt<uint64_t>($1), $2, $3))") + [require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda9_int64)] void InterlockedCompareExchangeU64(uint byteAddress, uint64_t compareValue, uint64_t value, out uint64_t outOriginalValue); __specialized_for_target(hlsl) + [ForceInline] void InterlockedCompareExchangeU64(uint byteAddress, uint64_t compareValue, uint64_t value, out uint64_t outOriginalValue) { outOriginalValue = __asuint64(__cas(this, byteAddress, __asuint2(compareValue), __asuint2(value))); @@ -2711,6 +2891,7 @@ ${{{{ __specialized_for_target(glsl) __specialized_for_target(spirv) + [ForceInline] void InterlockedCompareExchangeU64(uint byteAddress, uint64_t compareValue, uint64_t value, out uint64_t outOriginalValue) { let buf = __getEquivalentStructuredBuffer<uint64_t>(this); @@ -2719,8 +2900,9 @@ ${{{{ // Max - __cuda_sm_version(3.5) + __cuda_sm_version(5.0) __target_intrinsic(cuda, "atomicMax($0._getPtrAt<uint64_t>($1), $2)") + [require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda5_int64)] uint64_t InterlockedMaxU64(uint byteAddress, uint64_t value); __specialized_for_target(hlsl) @@ -2728,14 +2910,15 @@ ${{{{ __specialized_for_target(glsl) __specialized_for_target(spirv) + [ForceInline] uint64_t InterlockedMaxU64(uint byteAddress, uint64_t value) { let buf = __getEquivalentStructuredBuffer<uint64_t>(this); return __atomicMax(buf[byteAddress / 8], value); } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedMax64(uint byteAddress, int64_t value) { __target_switch @@ -2744,8 +2927,8 @@ ${{{{ } } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedMax64(uint byteAddress, int64_t value, out int64_t outOriginalValue) { __target_switch @@ -2754,8 +2937,8 @@ ${{{{ } } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedMax64(uint byteAddress, uint64_t value) { __target_switch @@ -2764,8 +2947,8 @@ ${{{{ } } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedMax64(uint byteAddress, uint64_t value, out uint64_t outOriginalValue) { __target_switch @@ -2776,8 +2959,9 @@ ${{{{ // Min - __cuda_sm_version(3.5) + __cuda_sm_version(5.0) __target_intrinsic(cuda, "atomicMin($0._getPtrAt<uint64_t>($1), $2)") + [require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda5_int64)] uint64_t InterlockedMinU64(uint byteAddress, uint64_t value); __specialized_for_target(hlsl) @@ -2785,14 +2969,15 @@ ${{{{ __specialized_for_target(glsl) __specialized_for_target(spirv) + [ForceInline] uint64_t InterlockedMinU64(uint byteAddress, uint64_t value) { let buf = __getEquivalentStructuredBuffer<uint64_t>(this); return __atomicMin(buf[byteAddress / 8], value); } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedMin64(uint byteAddress, int64_t value) { __target_switch @@ -2801,8 +2986,8 @@ ${{{{ } } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedMin64(uint byteAddress, int64_t value, out int64_t outOriginalValue) { __target_switch @@ -2811,8 +2996,8 @@ ${{{{ } } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedMin64(uint byteAddress, uint64_t value) { __target_switch @@ -2821,8 +3006,8 @@ ${{{{ } } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedMin64(uint byteAddress, uint64_t value, out uint64_t outOriginalValue) { __target_switch @@ -2833,7 +3018,9 @@ ${{{{ // And + __cuda_sm_version(5.0) __target_intrinsic(cuda, "atomicAnd($0._getPtrAt<uint64_t>($1), $2)") + [require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda5_int64)] uint64_t InterlockedAndU64(uint byteAddress, uint64_t value); __specialized_for_target(hlsl) @@ -2841,14 +3028,15 @@ ${{{{ __specialized_for_target(glsl) __specialized_for_target(spirv) + [ForceInline] uint64_t InterlockedAndU64(uint byteAddress, uint64_t value) { let buf = __getEquivalentStructuredBuffer<uint64_t>(this); return __atomicAnd(buf[byteAddress / 8], value); } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedAnd64(uint byteAddress, uint64_t value) { __target_switch @@ -2857,8 +3045,8 @@ ${{{{ } } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedAnd64(uint byteAddress, uint64_t value, out uint64_t outOriginalValue) { __target_switch @@ -2869,7 +3057,9 @@ ${{{{ // Or + __cuda_sm_version(5.0) __target_intrinsic(cuda, "atomicOr($0._getPtrAt<uint64_t>($1), $2)") + [require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda5_int64)] uint64_t InterlockedOrU64(uint byteAddress, uint64_t value); __specialized_for_target(hlsl) @@ -2877,14 +3067,15 @@ ${{{{ __specialized_for_target(glsl) __specialized_for_target(spirv) + [ForceInline] uint64_t InterlockedOrU64(uint byteAddress, uint64_t value) { let buf = __getEquivalentStructuredBuffer<uint64_t>(this); return __atomicOr(buf[byteAddress / 8], value); } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedOr64(uint byteAddress, uint64_t value) { __target_switch @@ -2893,8 +3084,8 @@ ${{{{ } } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedOr64(uint byteAddress, uint64_t value, out uint64_t outOriginalValue) { __target_switch @@ -2905,7 +3096,9 @@ ${{{{ // Xor + __cuda_sm_version(5.0) __target_intrinsic(cuda, "atomicXor($0._getPtrAt<uint64_t>($1), $2)") + [require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda5_int64)] uint64_t InterlockedXorU64(uint byteAddress, uint64_t value); __specialized_for_target(hlsl) @@ -2913,14 +3106,15 @@ ${{{{ __specialized_for_target(glsl) __specialized_for_target(spirv) + [ForceInline] uint64_t InterlockedXorU64(uint byteAddress, uint64_t value) { let buf = __getEquivalentStructuredBuffer<uint64_t>(this); return __atomicXor(buf[byteAddress / 8], value); } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedXor64(uint byteAddress, uint64_t value) { __target_switch @@ -2929,8 +3123,8 @@ ${{{{ } } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedXor64(uint byteAddress, uint64_t value, out uint64_t outOriginalValue) { __target_switch @@ -2942,6 +3136,7 @@ ${{{{ // Exchange __target_intrinsic(cuda, "atomicExch($0._getPtrAt<uint64_t>($1), $2)") + [require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda9_int64)] uint64_t InterlockedExchangeU64(uint byteAddress, uint64_t value); __specialized_for_target(hlsl) @@ -2949,14 +3144,15 @@ ${{{{ __specialized_for_target(glsl) __specialized_for_target(spirv) + [ForceInline] uint64_t InterlockedExchangeU64(uint byteAddress, uint64_t value) { let buf = __getEquivalentStructuredBuffer<uint64_t>(this); return __atomicExchange(buf[byteAddress / 8], value); } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedExchangeFloat(uint byteAddress, float value, out float outOriginalValue) { __target_switch @@ -2965,8 +3161,8 @@ ${{{{ } } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedExchange64(uint byteAddress, int64_t value, out int64_t outOriginalValue) { __target_switch @@ -2975,8 +3171,8 @@ ${{{{ } } - [require(hlsl)] [ForceInline] + [require(hlsl, atomic_hlsl_sm_6_6)] void InterlockedExchange64(uint byteAddress, uint64_t value, out uint64_t outOriginalValue) { __target_switch @@ -2987,6 +3183,7 @@ ${{{{ // SM6.6 6 64bit atomics. [ForceInline] + [require(glsl_hlsl_spirv, atomic_glsl_hlsl_cuda9_int64)] void InterlockedAdd64(uint byteAddress, int64_t valueToAdd) { __target_switch @@ -3000,6 +3197,7 @@ ${{{{ } [ForceInline] + [require(glsl_hlsl_spirv, atomic_glsl_hlsl_cuda9_int64)] void InterlockedAdd64(uint byteAddress, int64_t valueToAdd, out int64_t outOriginalValue) { __target_switch @@ -3009,10 +3207,12 @@ ${{{{ case spirv: let buf = __getEquivalentStructuredBuffer<int64_t>(this); outOriginalValue = __atomicAdd(buf[byteAddress / 8], valueToAdd); + return; } } [ForceInline] + [require(glsl_hlsl_spirv, atomic_glsl_hlsl_cuda9_int64)] void InterlockedAdd64(uint byteAddress, uint64_t valueToAdd) { __target_switch @@ -3026,6 +3226,7 @@ ${{{{ } [ForceInline] + [require(glsl_hlsl_spirv, atomic_glsl_hlsl_cuda9_int64)] void InterlockedAdd64(uint byteAddress, uint64_t valueToAdd, out uint64_t outOriginalValue) { __target_switch @@ -3035,76 +3236,64 @@ ${{{{ case spirv: let buf = __getEquivalentStructuredBuffer<uint64_t>(this); outOriginalValue = __atomicAdd(buf[byteAddress / 8], valueToAdd); + return; } } - __specialized_for_target(hlsl) - void InterlockedCompareExchange64(uint byteAddress, int64_t compareValue, int64_t value, out int64_t outOriginalValue) - { - __cas(this, byteAddress, compareValue, value, outOriginalValue); - } - __specialized_for_target(glsl) - __specialized_for_target(spirv) - void InterlockedCompareExchange64(uint byteAddress, int64_t compareValue, int64_t value, out int64_t outOriginalValue) - { - let buf = __getEquivalentStructuredBuffer<int64_t>(this); - outOriginalValue = __cas(buf[byteAddress / 8], compareValue, value); - } - __specialized_for_target(hlsl) - void InterlockedCompareExchange64(uint byteAddress, uint64_t compareValue, uint64_t value, out uint64_t outOriginalValue) - { - __cas(this, byteAddress, compareValue, value, outOriginalValue); - } - __specialized_for_target(glsl) - __specialized_for_target(spirv) - void InterlockedCompareExchange64(uint byteAddress, uint64_t compareValue, uint64_t value, out uint64_t outOriginalValue) - { - let buf = __getEquivalentStructuredBuffer<uint64_t>(this); - outOriginalValue = __cas(buf[byteAddress / 8], compareValue, value); - } - [require(hlsl)] [ForceInline] - void InterlockedCompareStoreFloatBitwise(uint byteAddress, float compareValue, float value) + void InterlockedCompareExchange64(uint byteAddress, int64_t compareValue, int64_t value, out int64_t outOriginalValue) { __target_switch { - case hlsl: __intrinsic_asm ".InterlockedCompareStoreFloatBitwise"; + case hlsl: + __cas(this, byteAddress, compareValue, value, outOriginalValue); + return; + case glsl: + case spirv: + let buf = __getEquivalentStructuredBuffer<int64_t>(this); + outOriginalValue = __cas(buf[byteAddress / 8], compareValue, value); + return; } } - - [require(hlsl)] [ForceInline] - void InterlockedCompareExchangeFloatBitwise(uint byteAddress, float compareValue, float value, out float outOriginalValue) + void InterlockedCompareExchange64(uint byteAddress, uint64_t compareValue, uint64_t value, out uint64_t outOriginalValue) { __target_switch { - case hlsl: __intrinsic_asm ".InterlockedCompareExchangeFloatBitwise"; + case hlsl: + __cas(this, byteAddress, compareValue, value, outOriginalValue); + return; + case glsl: + case spirv: + let buf = __getEquivalentStructuredBuffer<uint64_t>(this); + outOriginalValue = __cas(buf[byteAddress / 8], compareValue, value); + return; } } - [require(hlsl)] [ForceInline] - void InterlockedCompareStore64(uint byteAddress, int64_t compareValue, int64_t value) + [require(hlsl, atomic_hlsl_sm_6_6)] + void InterlockedCompareStoreFloatBitwise(uint byteAddress, float compareValue, float value) { __target_switch { - case hlsl: __intrinsic_asm ".InterlockedCompareStore64"; + case hlsl: __intrinsic_asm ".InterlockedCompareStoreFloatBitwise"; } } - [require(hlsl)] [ForceInline] - void InterlockedCompareExchange64(uint byteAddress, int64_t compareValue, int64_t value, out int64_t outOriginalValue) + [require(hlsl, atomic_hlsl_sm_6_6)] + void InterlockedCompareExchangeFloatBitwise(uint byteAddress, float compareValue, float value, out float outOriginalValue) { __target_switch { - case hlsl: __intrinsic_asm ".InterlockedCompareExchange64"; + case hlsl: __intrinsic_asm ".InterlockedCompareExchangeFloatBitwise"; } } - [require(hlsl)] [ForceInline] - void InterlockedCompareStore64(uint byteAddress, uint64_t compareValue, uint64_t value) + [require(hlsl, atomic_hlsl_sm_6_6)] + void InterlockedCompareStore64(uint byteAddress, int64_t compareValue, int64_t value) { __target_switch { @@ -3112,13 +3301,13 @@ ${{{{ } } - [require(hlsl)] [ForceInline] - void InterlockedCompareExchange64(uint byteAddress, uint64_t compareValue, uint64_t value, out uint64_t outOriginalValue) + [require(hlsl, atomic_hlsl_sm_6_6)] + void InterlockedCompareStore64(uint byteAddress, uint64_t compareValue, uint64_t value) { __target_switch { - case hlsl: __intrinsic_asm ".InterlockedCompareExchange64"; + case hlsl: __intrinsic_asm ".InterlockedCompareStore64"; } } @@ -3128,6 +3317,7 @@ ${{{{ // Added operations: [ForceInline] + [require(cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] void InterlockedAdd( UINT dest, UINT value, @@ -3145,6 +3335,7 @@ ${{{{ } [ForceInline] + [require(cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] void InterlockedAdd( UINT dest, UINT value) @@ -3160,6 +3351,8 @@ ${{{{ } } + [ForceInline] + [require(cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] void InterlockedAnd( UINT dest, UINT value, @@ -3176,6 +3369,8 @@ ${{{{ } } + [ForceInline] + [require(cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] void InterlockedAnd( UINT dest, UINT value) @@ -3191,6 +3386,8 @@ ${{{{ } } + [ForceInline] + [require(cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] void InterlockedCompareExchange( UINT dest, UINT compare_value, @@ -3208,6 +3405,8 @@ ${{{{ } } + [ForceInline] + [require(cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] void InterlockedCompareStore( UINT dest, UINT compare_value, @@ -3224,6 +3423,8 @@ ${{{{ } } + [ForceInline] + [require(cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] void InterlockedExchange( UINT dest, UINT value, @@ -3240,6 +3441,8 @@ ${{{{ } } + [ForceInline] + [require(cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] void InterlockedMax( UINT dest, UINT value, @@ -3256,6 +3459,8 @@ ${{{{ } } + [ForceInline] + [require(cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] void InterlockedMax( UINT dest, UINT value) @@ -3271,6 +3476,8 @@ ${{{{ } } + [ForceInline] + [require(cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] void InterlockedMin( UINT dest, UINT value, @@ -3287,6 +3494,8 @@ ${{{{ } } + [ForceInline] + [require(cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] void InterlockedMin( UINT dest, UINT value) @@ -3302,6 +3511,8 @@ ${{{{ } } + [ForceInline] + [require(cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] void InterlockedOr( UINT dest, UINT value, @@ -3318,6 +3529,8 @@ ${{{{ } } + [ForceInline] + [require(cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] void InterlockedOr( UINT dest, UINT value) @@ -3333,6 +3546,8 @@ ${{{{ } } + [ForceInline] + [require(cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] void InterlockedXor( UINT dest, UINT value, @@ -3349,6 +3564,8 @@ ${{{{ } } + [ForceInline] + [require(cuda_glsl_hlsl_spirv, byteaddressbuffer_rw)] void InterlockedXor( UINT dest, UINT value) @@ -3423,6 +3640,7 @@ for(auto item : kMutableStructuredBufferCases) { __generic<T, L:IBufferDataLayout=DefaultDataLayout> __magic_type(HLSL$(item.name)Type) __intrinsic_type($(item.op)) +[require(cpp_cuda_glsl_hlsl_metal_spirv, structuredbuffer_rw)] struct $(item.name) { uint DecrementCounter(); @@ -3462,6 +3680,7 @@ ${{{{ }}}} __generic<T> +[require(glsl_hlsl_spirv, geometry)] __magic_type(HLSLPointStreamType) __intrinsic_type($(kIROp_HLSLPointStreamType)) struct PointStream @@ -3490,6 +3709,7 @@ struct PointStream }; __generic<T> +[require(glsl_hlsl_spirv, geometry)] __magic_type(HLSLLineStreamType) __intrinsic_type($(kIROp_HLSLLineStreamType)) struct LineStream @@ -3518,6 +3738,7 @@ struct LineStream }; __generic<T> +[require(glsl_hlsl_spirv, geometry)] __magic_type(HLSLTriangleStreamType) __intrinsic_type($(kIROp_HLSLTriangleStreamType)) struct TriangleStream @@ -3575,6 +3796,7 @@ __target_intrinsic(cuda, "$P_abs($0)") __target_intrinsic(cpp, "$P_abs($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fi(FAbs, SAbs) _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T abs(T x); /*{ // Note: this simple definition may not be appropriate for floating-point inputs @@ -3586,6 +3808,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fi(FAbs, SAbs) _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> abs(vector<T, N> x) { VECTOR_MAP_UNARY(T, N, abs, x); @@ -3594,6 +3817,7 @@ vector<T, N> abs(vector<T, N> x) __generic<T : __BuiltinIntegerType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,N,M> abs(matrix<T,N,M> x) { MATRIX_MAP_UNARY(T, N, M, abs, x); @@ -3606,6 +3830,7 @@ __target_intrinsic(cuda, "$P_abs($0)") __target_intrinsic(cpp, "$P_abs($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fi(FAbs, SAbs) _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T abs(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -3613,6 +3838,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fi(FAbs, SAbs) _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> abs(vector<T, N> x) { VECTOR_MAP_UNARY(T, N, abs, x); @@ -3621,6 +3847,7 @@ vector<T, N> abs(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,N,M> abs(matrix<T,N,M> x) { MATRIX_MAP_UNARY(T, N, M, abs, x); @@ -3635,6 +3862,7 @@ __target_intrinsic(cuda, "$P_acos($0)") __target_intrinsic(cpp, "$P_acos($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Acos _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T acos(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -3642,6 +3870,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Acos _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> acos(vector<T, N> x) { VECTOR_MAP_UNARY(T, N, acos, x); @@ -3650,6 +3879,7 @@ vector<T, N> acos(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> acos(matrix<T, N, M> x) { MATRIX_MAP_UNARY(T, N, M, acos, x); @@ -3740,6 +3970,7 @@ bool all(matrix<T,N,M> x) // Barrier for writes to all memory spaces (HLSL SM 5.0) __glsl_extension(GL_KHR_memory_scope_semantics) +[require(cuda_glsl_hlsl_spirv, memorybarrier_compute)] void AllMemoryBarrier() { __target_switch @@ -3756,6 +3987,7 @@ void AllMemoryBarrier() // Thread-group sync and barrier for writes to all memory spaces (HLSL SM 5.0) __glsl_extension(GL_KHR_memory_scope_semantics) +[require(cuda_glsl_hlsl_spirv, memorybarrier_compute)] void AllMemoryBarrierWithGroupSync() { __target_switch @@ -3868,6 +4100,7 @@ __target_intrinsic(cuda, "$P_asdouble($0, $1)") __target_intrinsic(spirv, "%v = OpCompositeConstruct _type(uint2) resultId _0 _1; OpExtInst resultType resultId glsl450 59 %v") __glsl_extension(GL_ARB_gpu_shader5) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] double asdouble(uint lowbits, uint highbits); // Reinterpret bits as a float (HLSL SM 4.0) @@ -3878,6 +4111,7 @@ __target_intrinsic(cpp, "$P_asfloat($0)") __target_intrinsic(cuda, "$P_asfloat($0)") __target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] float asfloat(int x); __target_intrinsic(hlsl) @@ -3886,6 +4120,7 @@ __target_intrinsic(cpp, "$P_asfloat($0)") __target_intrinsic(cuda, "$P_asfloat($0)") __target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] float asfloat(uint x); __generic<let N : int> @@ -3893,6 +4128,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "intBitsToFloat") __target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] vector<float, N> asfloat(vector< int, N> x) { VECTOR_MAP_UNARY(float, N, asfloat, x); @@ -3903,6 +4139,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "uintBitsToFloat") __target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] vector<float,N> asfloat(vector<uint,N> x) { VECTOR_MAP_UNARY(float, N, asfloat, x); @@ -3911,6 +4148,7 @@ vector<float,N> asfloat(vector<uint,N> x) __generic<let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] matrix<float,N,M> asfloat(matrix< int,N,M> x) { MATRIX_MAP_UNARY(float, N, M, asfloat, x); @@ -3919,6 +4157,7 @@ matrix<float,N,M> asfloat(matrix< int,N,M> x) __generic<let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] matrix<float,N,M> asfloat(matrix<uint,N,M> x) { MATRIX_MAP_UNARY(float, N, M, asfloat, x); @@ -3927,18 +4166,21 @@ matrix<float,N,M> asfloat(matrix<uint,N,M> x) // No op [__unsafeForceInlineEarly] [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] float asfloat(float x) { return x; } __generic<let N : int> [__unsafeForceInlineEarly] [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<float,N> asfloat(vector<float,N> x) { return x; } __generic<let N : int, let M : int> [__unsafeForceInlineEarly] [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<float,N,M> asfloat(matrix<float,N,M> x) { return x; } @@ -3950,6 +4192,7 @@ __target_intrinsic(cuda, "$P_asin($0)") __target_intrinsic(cpp, "$P_asin($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Asin _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T asin(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -3957,6 +4200,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Asin _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> asin(vector<T, N> x) { VECTOR_MAP_UNARY(T,N,asin,x); @@ -3965,6 +4209,7 @@ vector<T, N> asin(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> asin(matrix<T, N, M> x) { MATRIX_MAP_UNARY(T,N,M,asin,x); @@ -3978,6 +4223,7 @@ __target_intrinsic(cpp, "$P_asint($0)") __target_intrinsic(cuda, "$P_asint($0)") __target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] int asint(float x); __target_intrinsic(hlsl) @@ -3986,6 +4232,7 @@ __target_intrinsic(cpp, "$P_asint($0)") __target_intrinsic(cuda, "$P_asint($0)") __target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] int asint(uint x); __generic<let N : int> @@ -3993,6 +4240,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "floatBitsToInt") __target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] vector<int, N> asint(vector<float, N> x) { VECTOR_MAP_UNARY(int, N, asint, x); @@ -4003,6 +4251,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "ivec$N0($0)") __target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] vector<int, N> asint(vector<uint, N> x) { VECTOR_MAP_UNARY(int, N, asint, x); @@ -4011,6 +4260,7 @@ vector<int, N> asint(vector<uint, N> x) __generic<let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] matrix<int, N, M> asint(matrix<float, N, M> x) { MATRIX_MAP_UNARY(int, N, M, asint, x); @@ -4019,6 +4269,7 @@ matrix<int, N, M> asint(matrix<float, N, M> x) __generic<let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] matrix<int, N, M> asint(matrix<uint, N, M> x) { MATRIX_MAP_UNARY(int, N, M, asint, x); @@ -4027,18 +4278,21 @@ matrix<int, N, M> asint(matrix<uint, N, M> x) // No op [__unsafeForceInlineEarly] [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] int asint(int x) { return x; } __generic<let N : int> [__unsafeForceInlineEarly] [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] vector<int,N> asint(vector<int,N> x) { return x; } __generic<let N : int, let M : int> [__unsafeForceInlineEarly] [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] matrix<int,N,M> asint(matrix<int,N,M> x) { return x; } @@ -4046,6 +4300,7 @@ matrix<int,N,M> asint(matrix<int,N,M> x) __glsl_extension(GL_ARB_gpu_shader5) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] void asuint(double value, out uint lowbits, out uint highbits) { __target_switch @@ -4074,6 +4329,7 @@ __target_intrinsic(spirv, "OpBitcast resultType resultId _0") __target_intrinsic(cpp, "$P_asuint($0)") __target_intrinsic(cuda, "$P_asuint($0)") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] uint asuint(float x); __target_intrinsic(hlsl) @@ -4082,6 +4338,7 @@ __target_intrinsic(spirv, "OpBitcast resultType resultId _0") __target_intrinsic(cpp, "$P_asuint($0)") __target_intrinsic(cuda, "$P_asuint($0)") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] uint asuint(int x); __generic<let N : int> @@ -4089,6 +4346,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "floatBitsToUint") __target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] vector<uint,N> asuint(vector<float,N> x) { VECTOR_MAP_UNARY(uint, N, asuint, x); @@ -4099,6 +4357,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "uvec$N0($0)") __target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] vector<uint, N> asuint(vector<int, N> x) { VECTOR_MAP_UNARY(uint, N, asuint, x); @@ -4107,6 +4366,7 @@ vector<uint, N> asuint(vector<int, N> x) __generic<let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] matrix<uint,N,M> asuint(matrix<float,N,M> x) { MATRIX_MAP_UNARY(uint, N, M, asuint, x); @@ -4115,6 +4375,7 @@ matrix<uint,N,M> asuint(matrix<float,N,M> x) __generic<let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_4_0)] matrix<uint, N, M> asuint(matrix<int, N, M> x) { MATRIX_MAP_UNARY(uint, N, M, asuint, x); @@ -4175,13 +4436,16 @@ __target_intrinsic(glsl, "uint16_t(packHalf2x16(vec2($0, 0.0)))") __target_intrinsic(cuda, "__half_as_ushort") __target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] +[require(cuda_glsl_hlsl_spirv, shader5_sm_5_0)] uint16_t asuint16(float16_t value); [__readNone] +[require(cuda_glsl_hlsl_spirv, shader5_sm_5_0)] vector<uint16_t,N> asuint16<let N : int>(vector<float16_t,N> value) { VECTOR_MAP_UNARY(uint16_t, N, asuint16, value); } [__readNone] +[require(cuda_glsl_hlsl_spirv, shader5_sm_5_0)] matrix<uint16_t,R,C> asuint16<let R : int, let C : int>(matrix<float16_t,R,C> value) { MATRIX_MAP_UNARY(uint16_t, R, C, asuint16, value); } @@ -4192,6 +4456,7 @@ __target_intrinsic(glsl, "float16_t(unpackHalf2x16($0).x)") __target_intrinsic(cuda, "__ushort_as_half") __target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] +[require(cuda_glsl_hlsl_spirv, shader5_sm_5_0)] float16_t asfloat16(uint16_t value); [__readNone] @@ -4207,18 +4472,41 @@ matrix<float16_t,R,C> asfloat16<let R : int, let C : int>(matrix<uint16_t,R,C> v __target_intrinsic(hlsl) __target_intrinsic(cuda, "__half_as_short") __target_intrinsic(spirv, "OpBitcast resultType resultId _0") -[__unsafeForceInlineEarly][__readNone] int16_t asint16(float16_t value) { return asuint16(value); } -__target_intrinsic(hlsl) [__unsafeForceInlineEarly][__readNone] vector<int16_t,N> asint16<let N : int>(vector<float16_t,N> value) { return asuint16(value); } -__target_intrinsic(hlsl) [__unsafeForceInlineEarly][__readNone] matrix<int16_t,R,C> asint16<let R : int, let C : int>(matrix<float16_t,R,C> value) { return asuint16(value); } +[__unsafeForceInlineEarly] +[__readNone] +[require(cuda_hlsl_spirv, shader5_sm_5_0)] +int16_t asint16(float16_t value) { return asuint16(value); } + +__target_intrinsic(hlsl) +[__unsafeForceInlineEarly] +[__readNone] +[require(cuda_hlsl_spirv, shader5_sm_5_0)] +vector<int16_t,N> asint16<let N : int>(vector<float16_t,N> value) { return asuint16(value); } + +__target_intrinsic(hlsl) +[__unsafeForceInlineEarly] +[__readNone] +[require(cuda_hlsl_spirv, shader5_sm_5_0)] +matrix<int16_t,R,C> asint16<let R : int, let C : int>(matrix<float16_t,R,C> value) { return asuint16(value); } __target_intrinsic(hlsl) __target_intrinsic(cuda, "__short_as_half") __target_intrinsic(spirv, "OpBitcast resultType resultId _0") [__readNone] -[__unsafeForceInlineEarly] float16_t asfloat16(int16_t value) { return asfloat16(asuint16(value)); } +[__unsafeForceInlineEarly] +[require(cuda_hlsl_spirv, shader5_sm_5_0)] +float16_t asfloat16(int16_t value) { return asfloat16(asuint16(value)); } + +__target_intrinsic(hlsl) +[__unsafeForceInlineEarly] +[__readNone] +vector<float16_t,N> asfloat16<let N : int>(vector<int16_t,N> value) { return asfloat16(asuint16(value)); } -__target_intrinsic(hlsl) [__unsafeForceInlineEarly][__readNone] vector<float16_t,N> asfloat16<let N : int>(vector<int16_t,N> value) { return asfloat16(asuint16(value)); } -__target_intrinsic(hlsl) [__unsafeForceInlineEarly][__readNone] matrix<float16_t,R,C> asfloat16<let R : int, let C : int>(matrix<int16_t,R,C> value) { return asfloat16(asuint16(value)); } +__target_intrinsic(hlsl) +[__unsafeForceInlineEarly] +[__readNone] +[require(cuda_hlsl_spirv, shader5_sm_5_0)] +matrix<float16_t,R,C> asfloat16<let R : int, let C : int>(matrix<int16_t,R,C> value) { return asfloat16(asuint16(value)); } // Inverse tangent (HLSL SM 1.0) __generic<T : __BuiltinFloatingPointType> @@ -4228,6 +4516,7 @@ __target_intrinsic(cuda, "$P_atan($0)") __target_intrinsic(cpp, "$P_atan($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Atan _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T atan(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -4235,6 +4524,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Atan _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> atan(vector<T, N> x) { VECTOR_MAP_UNARY(T, N, atan, x); @@ -4243,6 +4533,7 @@ vector<T, N> atan(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> atan(matrix<T, N, M> x) { MATRIX_MAP_UNARY(T, N, M, atan, x); @@ -4255,6 +4546,7 @@ __target_intrinsic(cuda, "$P_atan2($0, $1)") __target_intrinsic(cpp, "$P_atan2($0, $1)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Atan2 _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T atan2(T y, T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -4262,6 +4554,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl,"atan($0,$1)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Atan2 _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> atan2(vector<T, N> y, vector<T, N> x) { VECTOR_MAP_BINARY(T, N, atan2, y, x); @@ -4270,6 +4563,7 @@ vector<T, N> atan2(vector<T, N> y, vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,N,M> atan2(matrix<T,N,M> y, matrix<T,N,M> x) { MATRIX_MAP_BINARY(T, N, M, atan2, y, x); @@ -4283,6 +4577,7 @@ __target_intrinsic(cuda, "$P_ceil($0)") __target_intrinsic(cpp, "$P_ceil($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Ceil _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T ceil(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -4290,6 +4585,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Ceil _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> ceil(vector<T, N> x) { VECTOR_MAP_UNARY(T, N, ceil, x); @@ -4298,6 +4594,7 @@ vector<T, N> ceil(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> ceil(matrix<T, N, M> x) { MATRIX_MAP_UNARY(T, N, M, ceil, x); @@ -4313,6 +4610,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FClamp, UClamp, SClamp) _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T clamp(T x, T minBound, T maxBound) { return min(max(x, minBound), maxBound); @@ -4323,6 +4621,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FClamp, UClamp, SClamp) _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> clamp(vector<T, N> x, vector<T, N> minBound, vector<T, N> maxBound) { return min(max(x, minBound), maxBound); @@ -4331,6 +4630,7 @@ vector<T, N> clamp(vector<T, N> x, vector<T, N> minBound, vector<T, N> maxBound) __generic<T : __BuiltinIntegerType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,N,M> clamp(matrix<T,N,M> x, matrix<T,N,M> minBound, matrix<T,N,M> maxBound) { return min(max(x, minBound), maxBound); @@ -4341,6 +4641,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FClamp, UClamp, SClamp) _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T clamp(T x, T minBound, T maxBound) { return min(max(x, minBound), maxBound); @@ -4351,6 +4652,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FClamp, UClamp, SClamp) _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> clamp(vector<T, N> x, vector<T, N> minBound, vector<T, N> maxBound) { return min(max(x, minBound), maxBound); @@ -4359,6 +4661,7 @@ vector<T, N> clamp(vector<T, N> x, vector<T, N> minBound, vector<T, N> maxBound) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,N,M> clamp(matrix<T,N,M> x, matrix<T,N,M> minBound, matrix<T,N,M> maxBound) { return min(max(x, minBound), maxBound); @@ -4367,6 +4670,7 @@ matrix<T,N,M> clamp(matrix<T,N,M> x, matrix<T,N,M> minBound, matrix<T,N,M> maxBo // Clip (discard) fragment conditionally __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) +[require(cpp_cuda_glsl_hlsl_spirv, fragment)] void clip(T x) { if(x < T(0)) discard; @@ -4374,6 +4678,7 @@ void clip(T x) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) +[require(cpp_cuda_glsl_hlsl_spirv, fragment)] void clip(vector<T,N> x) { if(any(x < T(0))) discard; @@ -4381,6 +4686,7 @@ void clip(vector<T,N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) +[require(cpp_cuda_glsl_hlsl_spirv, fragment)] void clip(matrix<T,N,M> x) { if(any(x < T(0))) discard; @@ -4394,6 +4700,7 @@ __target_intrinsic(cuda, "$P_cos($0)") __target_intrinsic(cpp, "$P_cos($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Cos _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T cos(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -4444,6 +4751,7 @@ matrix<T, N, M> cosh(matrix<T, N, M> x) // Population count [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] uint countbits(uint value) { __target_switch @@ -4467,6 +4775,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Cross _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T,3> cross(vector<T,3> left, vector<T,3> right) { return vector<T,3>( @@ -4480,6 +4789,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Cross _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, 3> cross(vector<T, 3> left, vector<T, 3> right) { return vector<T, 3>( @@ -4491,6 +4801,7 @@ vector<T, 3> cross(vector<T, 3> left, vector<T, 3> right) // Convert encoded color __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] int4 D3DCOLORtoUBYTE4(float4 color) { let scaled = color.zyxw * 255.001999f; @@ -4504,6 +4815,7 @@ for (auto xOrY : diffDimensions) { }}}} __generic<T : __BuiltinFloatingPointType> [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, fragmentprocessing)] T dd$(xOrY)(T x) { __target_switch @@ -4521,6 +4833,7 @@ T dd$(xOrY)(T x) __generic<T : __BuiltinFloatingPointType, let N : int> [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, fragmentprocessing)] vector<T, N> dd$(xOrY)(vector<T, N> x) { __target_switch @@ -4539,6 +4852,7 @@ vector<T, N> dd$(xOrY)(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, fragmentprocessing)] matrix<T, N, M> dd$(xOrY)(matrix<T, N, M> x) { MATRIX_MAP_UNARY(T, N, M, dd$(xOrY), x); @@ -4547,6 +4861,7 @@ matrix<T, N, M> dd$(xOrY)(matrix<T, N, M> x) __generic<T : __BuiltinFloatingPointType> __glsl_extension(GL_ARB_derivative_control) [__readNone] +[require(glsl_hlsl_spirv, fragmentprocessing_derivativecontrol)] T dd$(xOrY)_coarse(T x) { __target_switch @@ -4560,6 +4875,7 @@ T dd$(xOrY)_coarse(T x) __generic<T : __BuiltinFloatingPointType, let N : int> __glsl_extension(GL_ARB_derivative_control) [__readNone] +[require(glsl_hlsl_spirv, fragmentprocessing_derivativecontrol)] vector<T, N> dd$(xOrY)_coarse(vector<T, N> x) { __target_switch @@ -4581,6 +4897,7 @@ matrix<T, N, M> dd$(xOrY)_coarse(matrix<T, N, M> x) __generic<T : __BuiltinFloatingPointType> __glsl_extension(GL_ARB_derivative_control) [__readNone] +[require(glsl_hlsl_spirv, fragmentprocessing_derivativecontrol)] T dd$(xOrY)_fine(T x) { __target_switch @@ -4594,6 +4911,7 @@ T dd$(xOrY)_fine(T x) __generic<T : __BuiltinFloatingPointType, let N : int> __glsl_extension(GL_ARB_derivative_control) [__readNone] +[require(glsl_hlsl_spirv, fragmentprocessing_derivativecontrol)] vector<T, N> dd$(xOrY)_fine(vector<T, N> x) { __target_switch @@ -4659,6 +4977,7 @@ T determinant(matrix<T,N,N> m); // Barrier for device memory __glsl_extension(GL_KHR_memory_scope_semantics) +[require(cuda_glsl_hlsl_spirv, memorybarrier_compute)] void DeviceMemoryBarrier() { __target_switch @@ -4674,6 +4993,7 @@ void DeviceMemoryBarrier() } __glsl_extension(GL_KHR_memory_scope_semantics) +[require(cuda_glsl_hlsl_spirv, memorybarrier_compute)] void DeviceMemoryBarrierWithGroupSync() { __target_switch @@ -4695,6 +5015,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Distance _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T distance(vector<T, N> x, vector<T, N> y) { return length(x - y); @@ -4702,6 +5023,7 @@ T distance(vector<T, N> x, vector<T, N> y) __generic<T : __BuiltinFloatingPointType> [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T distance(T x, T y) { return length(x - y); @@ -4714,6 +5036,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) [__readNone] [ForceInline] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T dot(T x, T y) { return x * y; @@ -4724,6 +5047,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpDot resultType resultId _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T dot(vector<T, N> x, vector<T, N> y) { T result = T(0); @@ -4735,6 +5059,7 @@ T dot(vector<T, N> x, vector<T, N> y) __generic<T : __BuiltinIntegerType, let N : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T dot(vector<T, N> x, vector<T, N> y) { T result = T(0); @@ -4774,12 +5099,14 @@ __generic<T : __BuiltinArithmeticType> __target_intrinsic(glsl, interpolateAtCentroid) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 InterpolateAtCentroid _0") [__readNone] +[require(glsl_spirv, fragmentprocessing)] T EvaluateAttributeAtCentroid(T x); __generic<T : __BuiltinArithmeticType, let N : int> __target_intrinsic(glsl, interpolateAtCentroid) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 InterpolateAtCentroid _0") [__readNone] +[require(glsl_spirv, fragmentprocessing)] vector<T,N> EvaluateAttributeAtCentroid(vector<T,N> x); __generic<T : __BuiltinArithmeticType, let N : int, let M : int> @@ -4794,17 +5121,20 @@ __generic<T : __BuiltinArithmeticType> __target_intrinsic(glsl, "interpolateAtSample($0, int($1))") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 InterpolateAtSample _0 _1") [__readNone] +[require(glsl_spirv, fragmentprocessing)] T EvaluateAttributeAtSample(T x, uint sampleindex); __generic<T : __BuiltinArithmeticType, let N : int> __target_intrinsic(glsl, "interpolateAtSample($0, int($1))") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 InterpolateAtSample _0 _1") [__readNone] +[require(glsl_spirv, fragmentprocessing)] vector<T,N> EvaluateAttributeAtSample(vector<T,N> x, uint sampleindex); __generic<T : __BuiltinArithmeticType, let N : int, let M : int> __target_intrinsic(glsl, "interpolateAtSample($0, int($1))") [__readNone] +[require(glsl_spirv, fragmentprocessing)] matrix<T,N,M> EvaluateAttributeAtSample(matrix<T,N,M> x, uint sampleindex) { matrix<T,N,M> result; @@ -4819,12 +5149,14 @@ __generic<T : __BuiltinArithmeticType> __target_intrinsic(glsl, "interpolateAtOffset($0, vec2($1) / 16.0f)") __target_intrinsic(spirv, "%foffset = OpConvertSToF _type(float2) resultId _1; %offsetdiv16 = 136 _type(float2) resultId %foffset const(float2, 16.0, 16.0); OpExtInst resultType resultId glsl450 78 _0 %offsetdiv16") [__readNone] +[require(glsl_spirv, fragmentprocessing)] T EvaluateAttributeSnapped(T x, int2 offset); __generic<T : __BuiltinArithmeticType, let N : int> __target_intrinsic(glsl, "interpolateAtOffset($0, vec2($1) / 16.0f)") __target_intrinsic(spirv, "%foffset = OpConvertSToF _type(float2) resultId _1; %offsetdiv16 = 136 _type(float2) resultId %foffset const(float2, 16.0, 16.0); OpExtInst resultType resultId glsl450 78 _0 %offsetdiv16") [__readNone] +[require(glsl_spirv, fragmentprocessing)] vector<T,N> EvaluateAttributeSnapped(vector<T,N> x, int2 offset); __generic<T : __BuiltinArithmeticType, let N : int, let M : int> @@ -4849,6 +5181,7 @@ __target_intrinsic(cuda, "$P_exp($0)") __target_intrinsic(cpp, "$P_exp($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Exp _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T exp(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -4872,10 +5205,8 @@ matrix<T, N, M> exp(matrix<T, N, M> x) // Base-2 exponent __generic<T : __BuiltinFloatingPointType> -__target_intrinsic(hlsl) -__target_intrinsic(cuda, "$P_exp2($0)") -__target_intrinsic(cpp, "$P_exp2($0)") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T exp2(T x) { __target_switch @@ -4896,7 +5227,14 @@ T exp2(T x) result:$$float = OpExtInst glsl450 Exp2 $xf }); } + case hlsl: + __intrinsic_asm "exp2($0)"; + case cpp: + __intrinsic_asm "$P_exp2($0)"; + case cuda: + __intrinsic_asm "$P_exp2($0)"; } + } __generic<T : __BuiltinFloatingPointType, let N : int> @@ -4918,17 +5256,28 @@ matrix<T,N,M> exp2(matrix<T,N,M> x) // Convert 16-bit float stored in low bits of integer -__target_intrinsic(glsl, "unpackHalf2x16($0).x") __glsl_version(420) -__target_intrinsic(hlsl) __cuda_sm_version(6.0) -__target_intrinsic(cuda, "__half2float(__ushort_as_half($0))") -__target_intrinsic(spirv, R"( - %lowBits = OpUConvert _type(uint16_t) resultId _0; - %half = OpBitcast _type(half) resultId %lowBits; - OpFConvert resultType resultId %half)") [__readNone] -float f16tof32(uint value); +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] +float f16tof32(uint value) +{ + __target_switch + { + case glsl: __intrinsic_asm "unpackHalf2x16($0).x"; + case hlsl: __intrinsic_asm "f16tof32($0)"; + case cuda: __intrinsic_asm "__half2float(__ushort_as_half($0))"; + case cpp: __intrinsic_asm "f16tof32($0)"; + case spirv: + { + return spirv_asm { + %lowBits = OpUConvert $$uint16_t $value; + %half = OpBitcast $$half %lowBits; + result:$$float = OpFConvert %half + }; + } + } +} __generic<let N : int> __target_intrinsic(hlsl) @@ -4941,17 +5290,28 @@ vector<float, N> f16tof32(vector<uint, N> value) // Convert to 16-bit float stored in low bits of integer -__target_intrinsic(glsl, "packHalf2x16(vec2($0,0.0))") __glsl_version(420) -__target_intrinsic(hlsl) __cuda_sm_version(6.0) -__target_intrinsic(cuda, "__half_as_ushort(__float2half($0))") -__target_intrinsic(spirv, R"( - %half = OpFConvert _type(half) resultId _0; - %lowBits = OpBitcast _type(uint16_t) resultId %half; - OpUConvert resultType resultId %lowBits)") [__readNone] -uint f32tof16(float value); +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] +uint f32tof16(float value) +{ + __target_switch + { + case glsl: __intrinsic_asm "packHalf2x16(vec2($0,0.0))"; + case hlsl: __intrinsic_asm "f32tof16($0)"; + case cuda: __intrinsic_asm "__half_as_ushort(__float2half($0))"; + case cpp: __intrinsic_asm "f32tof16($0)"; + case spirv: + { + return spirv_asm { + %half = OpFConvert $$half $value; + %lowBits = OpBitcast $$uint16_t %half; + result:$$uint = OpUConvert %lowBits + }; + } + } +} __generic<let N : int> __target_intrinsic(hlsl) @@ -4966,12 +5326,25 @@ vector<uint, N> f32tof16(vector<float, N> value) // It's not clear what happens with float16 time in HLSL -> can the float16 coerce to uint for example? If so that would // give the wrong result -__target_intrinsic(glsl, "unpackHalf2x16($0).x") -__target_intrinsic(cuda, "__half2float") -__target_intrinsic(spirv, "OpFConvert resultType resultId _0") __glsl_version(420) [__readNone] -float f16tof32(float16_t value); +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] +float f16tof32(float16_t value) +{ + __target_switch + { + case glsl: __intrinsic_asm "unpackHalf2x16($0).x"; + case hlsl: __intrinsic_asm "f16tof32($0)"; + case cuda: __intrinsic_asm "__half2float($0)"; + case cpp: __intrinsic_asm "f16tof32($0)"; + case spirv: + { + return spirv_asm { + result:$$float = OpFConvert $value + }; + } + } +} __generic<let N : int> __target_intrinsic(hlsl) @@ -4989,6 +5362,7 @@ __glsl_version(420) __target_intrinsic(cuda, "__float2half") __target_intrinsic(spirv, "OpFConvert resultType resultId _0") [__readNone] +[require(cuda_glsl_spirv, shader5_sm_5_0)] float16_t f32tof16_(float value); __generic<let N : int> @@ -5008,6 +5382,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FaceForward _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_400)] vector<T,N> faceforward(vector<T,N> n, vector<T,N> i, vector<T,N> ng) { return dot(ng, i) < T(0.0f) ? n : -n; @@ -5020,6 +5395,7 @@ __target_intrinsic(cuda, "$P_firstbithigh($0)") __target_intrinsic(cpp, "$P_firstbithigh($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FindSMsb _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] int firstbithigh(int value); __target_intrinsic(hlsl) @@ -5038,6 +5414,7 @@ __target_intrinsic(cuda, "$P_firstbithigh($0)") __target_intrinsic(cpp, "$P_firstbithigh($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FindUMsb _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] uint firstbithigh(uint value); __target_intrinsic(hlsl) @@ -5057,6 +5434,7 @@ __target_intrinsic(cuda, "$P_firstbitlow($0)") __target_intrinsic(cpp, "$P_firstbitlow($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FindILsb _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] int firstbitlow(int value); __target_intrinsic(hlsl) @@ -5075,6 +5453,7 @@ __target_intrinsic(cuda, "$P_firstbitlow($0)") __target_intrinsic(cpp, "$P_firstbitlow($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FindILsb _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] uint firstbitlow(uint value); __target_intrinsic(hlsl) @@ -5096,6 +5475,7 @@ __target_intrinsic(cuda, "$P_floor($0)") __target_intrinsic(cpp, "$P_floor($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Floor _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T floor(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -5123,6 +5503,7 @@ __target_intrinsic(cuda, "$P_fma($0, $1, $2)") __target_intrinsic(cpp, "$P_fma($0, $1, $2)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] T fma(T a, T b, T c) { __target_switch @@ -5132,6 +5513,8 @@ T fma(T a, T b, T c) return mad(a, b, c); else __intrinsic_asm "fma($0, $1, $2)"; + default: + return a*b + c; } } @@ -5140,6 +5523,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] vector<T, N> fma(vector<T, N> a, vector<T, N> b, vector<T, N> c) { VECTOR_MAP_TRINARY(T, N, fma, a, b, c); @@ -5148,6 +5532,7 @@ vector<T, N> fma(vector<T, N> a, vector<T, N> b, vector<T, N> c) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] matrix<T, N, M> fma(matrix<T, N, M> a, matrix<T, N, M> b, matrix<T, N, M> c) { MATRIX_MAP_TRINARY(T, N, M, fma, a, b, c); @@ -5159,6 +5544,7 @@ __target_intrinsic(hlsl) __target_intrinsic(cuda, "$P_fmod($0, $1)") __target_intrinsic(cpp, "$P_fmod($0, $1)") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T fmod(T x, T y) { return x - y * trunc(x/y); @@ -5188,6 +5574,7 @@ __target_intrinsic(cuda, "$P_frac($0)") __target_intrinsic(cpp, "$P_frac($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Fract _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T frac(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -5215,6 +5602,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Frexp _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T frexp(T x, out int exp); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -5241,6 +5629,7 @@ __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpFwidth resultType resultId _0") +[require(glsl_hlsl_spirv, fragmentprocessing)] T fwidth(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -5276,6 +5665,7 @@ matrix<T, N, M> fwidth(matrix<T, N, M> x) __generic<T : __BuiltinType> [__readNone] __glsl_version(450) +[require(glsl_hlsl_spirv, getattributeatvertex)] T GetAttributeAtVertex(T attribute, uint vertexIndex) { __target_switch @@ -5308,6 +5698,7 @@ T GetAttributeAtVertex(T attribute, uint vertexIndex) __generic<T : __BuiltinType, let N : int> [__readNone] __glsl_version(450) +[require(glsl_hlsl_spirv, getattributeatvertex)] vector<T,N> GetAttributeAtVertex(vector<T,N> attribute, uint vertexIndex) { __target_switch @@ -5340,6 +5731,7 @@ vector<T,N> GetAttributeAtVertex(vector<T,N> attribute, uint vertexIndex) __generic<T : __BuiltinType, let N : int, let M : int> [__readNone] __glsl_version(450) +[require(glsl_hlsl_spirv, getattributeatvertex)] matrix<T,N,M> GetAttributeAtVertex(matrix<T,N,M> attribute, uint vertexIndex) { __target_switch @@ -5359,14 +5751,17 @@ matrix<T,N,M> GetAttributeAtVertex(matrix<T,N,M> attribute, uint vertexIndex) // Get number of samples in render target [__readNone] +[require(sm_4_0)] uint GetRenderTargetSampleCount(); // Get position of given sample [__readNone] +[require(sm_4_0)] float2 GetRenderTargetSamplePosition(int Index); // Group memory barrier __glsl_extension(GL_KHR_memory_scope_semantics) +[require(cuda_glsl_hlsl_spirv, memorybarrier_compute)] void GroupMemoryBarrier() { __target_switch @@ -5382,6 +5777,7 @@ void GroupMemoryBarrier() } } +[require(cuda_glsl_hlsl_spirv, memorybarrier_compute)] void __subgroupBarrier() { __target_switch @@ -5398,6 +5794,7 @@ void __subgroupBarrier() } __glsl_extension(GL_KHR_memory_scope_semantics) +[require(cuda_glsl_hlsl_spirv, memorybarrier_compute)] void GroupMemoryBarrierWithGroupSync() { __target_switch @@ -5417,6 +5814,7 @@ void GroupMemoryBarrierWithGroupSync() [ForceInline] __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedAdd(__ref int dest, int value) { __target_switch @@ -5434,6 +5832,7 @@ void InterlockedAdd(__ref int dest, int value) [ForceInline] __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedAdd(__ref uint dest, uint value) { __target_switch @@ -5457,6 +5856,7 @@ void InterlockedAdd(__ref uint dest, int value) [ForceInline] __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedAdd(__ref int dest, int value, out int original_value) { __target_switch @@ -5475,6 +5875,7 @@ void InterlockedAdd(__ref int dest, int value, out int original_value) [ForceInline] __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedAdd(__ref uint dest, uint value, out uint original_value) { __target_switch @@ -5528,6 +5929,7 @@ void InterlockedAdd(__ref uint64_t dest, uint64_t value, out uint64_t origina } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedAnd(__ref int dest, int value) { __target_switch @@ -5544,6 +5946,7 @@ void InterlockedAnd(__ref int dest, int value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedAnd(__ref uint dest, uint value) { __target_switch @@ -5560,6 +5963,7 @@ void InterlockedAnd(__ref uint dest, uint value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedAnd(__ref int dest, int value, out int original_value) { __target_switch @@ -5577,6 +5981,7 @@ void InterlockedAnd(__ref int dest, int value, out int original_value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedAnd(__ref uint dest, uint value, out uint original_value) { __target_switch @@ -5612,6 +6017,7 @@ void InterlockedAnd(__ref uint64_t dest, uint64_t value, out uint64_t origina } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedCompareExchange(__ref int dest, int compare_value, int value, out int original_value) { __target_switch @@ -5629,6 +6035,7 @@ void InterlockedCompareExchange(__ref int dest, int compare_value, int value, } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedCompareExchange(__ref uint dest, uint compare_value, uint value, out uint original_value) { __target_switch @@ -5700,6 +6107,7 @@ void InterlockedCompareExchange(__ref uint64_t dest, uint64_t compare_value, } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedCompareStore(__ref int dest, int compare_value, int value) { __target_switch @@ -5716,6 +6124,7 @@ void InterlockedCompareStore(__ref int dest, int compare_value, int value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedCompareStore(__ref uint dest, uint compare_value, uint value) { __target_switch @@ -5759,6 +6168,7 @@ void InterlockedCompareStore(__ref uint64_t dest, uint64_t compare_value, uint64 } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedExchange(__ref int dest, int value, out int original_value) { __target_switch @@ -5776,6 +6186,7 @@ void InterlockedExchange(__ref int dest, int value, out int original_value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedExchange(__ref uint dest, uint value, out uint original_value) { __target_switch @@ -5847,6 +6258,7 @@ void InterlockedExchange(__ref uint64_t dest, uint64_t value, out uint64_t or } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedMax(__ref int dest, int value) { __target_switch @@ -5863,6 +6275,7 @@ void InterlockedMax(__ref int dest, int value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedMax(__ref uint dest, uint value) { __target_switch @@ -5879,6 +6292,7 @@ void InterlockedMax(__ref uint dest, uint value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedMax(__ref int dest, int value, out int original_value) { __target_switch @@ -5896,6 +6310,7 @@ void InterlockedMax(__ref int dest, int value, out int original_value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedMax(__ref uint dest, uint value, out uint original_value) { __target_switch @@ -5949,6 +6364,7 @@ void InterlockedMax(__ref uint64_t dest, uint64_t value, out uint64_t origina } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedMin(__ref int dest, int value) { __target_switch @@ -5965,6 +6381,7 @@ void InterlockedMin(__ref int dest, int value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedMin(__ref uint dest, uint value) { __target_switch @@ -5981,6 +6398,7 @@ void InterlockedMin(__ref uint dest, uint value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedMin(__ref int dest, int value, out int original_value) { __target_switch @@ -5998,6 +6416,7 @@ void InterlockedMin(__ref int dest, int value, out int original_value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedMin(__ref uint dest, uint value, out uint original_value) { __target_switch @@ -6051,6 +6470,7 @@ void InterlockedMin(__ref uint64_t dest, uint64_t value, out uint64_t origina } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedOr(__ref int dest, int value) { __target_switch @@ -6067,6 +6487,7 @@ void InterlockedOr(__ref int dest, int value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedOr(__ref uint dest, uint value) { __target_switch @@ -6083,6 +6504,7 @@ void InterlockedOr(__ref uint dest, uint value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedOr(__ref int dest, int value, out int original_value) { __target_switch @@ -6100,6 +6522,7 @@ void InterlockedOr(__ref int dest, int value, out int original_value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedOr(__ref uint dest, uint value, out uint original_value) { __target_switch @@ -6135,6 +6558,7 @@ void InterlockedOr(__ref uint64_t dest, uint64_t value, out uint64_t original } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedXor(__ref int dest, int value) { __target_switch @@ -6151,6 +6575,7 @@ void InterlockedXor(__ref int dest, int value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedXor(__ref uint dest, uint value) { __target_switch @@ -6167,6 +6592,7 @@ void InterlockedXor(__ref uint dest, uint value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedXor(__ref int dest, int value, out int original_value) { __target_switch @@ -6184,6 +6610,7 @@ void InterlockedXor(__ref int dest, int value, out int original_value) } __glsl_version(430) +[require(cuda_glsl_hlsl_spirv, atomic_glsl_hlsl_cuda)] void InterlockedXor(__ref uint dest, uint value, out uint original_value) { __target_switch @@ -6224,6 +6651,7 @@ void InterlockedXor(__ref uint64_t dest, uint64_t value, out uint64_t origina __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] bool isfinite(T x) { __target_switch @@ -6239,6 +6667,7 @@ bool isfinite(T x) __generic<T : __BuiltinFloatingPointType, let N : int> [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<bool, N> isfinite(vector<T, N> x) { __target_switch @@ -6252,6 +6681,7 @@ vector<bool, N> isfinite(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<bool, N, M> isfinite(matrix<T, N, M> x) { MATRIX_MAP_UNARY(bool, N, M, isfinite, x); @@ -6260,6 +6690,7 @@ matrix<bool, N, M> isfinite(matrix<T, N, M> x) // Is floating-point value infinite? __generic<T : __BuiltinFloatingPointType> [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] bool isinf(T x) { __target_switch @@ -6277,6 +6708,7 @@ bool isinf(T x) __generic<T : __BuiltinFloatingPointType, let N : int> [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<bool, N> isinf(vector<T, N> x) { __target_switch @@ -6294,6 +6726,7 @@ vector<bool, N> isinf(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<bool, N, M> isinf(matrix<T, N, M> x) { MATRIX_MAP_UNARY(bool, N, M, isinf, x); @@ -6302,6 +6735,7 @@ matrix<bool, N, M> isinf(matrix<T, N, M> x) // Is floating-point value not-a-number? __generic<T : __BuiltinFloatingPointType> [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] bool isnan(T x) { __target_switch @@ -6319,6 +6753,7 @@ bool isnan(T x) __generic<T : __BuiltinFloatingPointType, let N : int> [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<bool, N> isnan(vector<T, N> x) { __target_switch @@ -6336,6 +6771,7 @@ vector<bool, N> isnan(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<bool, N, M> isnan(matrix<T, N, M> x) { MATRIX_MAP_UNARY(bool, N, M, isnan, x); @@ -6346,6 +6782,7 @@ matrix<bool, N, M> isnan(matrix<T, N, M> x) __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T ldexp(T x, T exp) { return x * exp2(exp); @@ -6354,6 +6791,7 @@ T ldexp(T x, T exp) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> ldexp(vector<T, N> x, vector<T, N> exp) { return x * exp2(exp); @@ -6362,6 +6800,7 @@ vector<T, N> ldexp(vector<T, N> x, vector<T, N> exp) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> ldexp(matrix<T, N, M> x, matrix<T, N, M> exp) { MATRIX_MAP_BINARY(T, N, M, ldexp, x, exp); @@ -6373,6 +6812,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Length _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T length(vector<T, N> x) { return sqrt(dot(x, x)); @@ -6380,6 +6820,7 @@ T length(vector<T, N> x) // Scalar float length __generic<T : __BuiltinFloatingPointType> +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T length(T x) { return abs(x); @@ -6391,6 +6832,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, mix) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FMix _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T lerp(T x, T y, T s) { return x * (T(1.0f) - s) + y * s; @@ -6401,6 +6843,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, mix) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 FMix _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> lerp(vector<T, N> x, vector<T, N> y, vector<T, N> s) { return x * (T(1.0f) - s) + y * s; @@ -6409,6 +6852,7 @@ vector<T, N> lerp(vector<T, N> x, vector<T, N> y, vector<T, N> s) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,N,M> lerp(matrix<T,N,M> x, matrix<T,N,M> y, matrix<T,N,M> s) { MATRIX_MAP_TRINARY(T, N, M, lerp, x, y, s); @@ -6417,6 +6861,7 @@ matrix<T,N,M> lerp(matrix<T,N,M> x, matrix<T,N,M> y, matrix<T,N,M> s) // Legacy lighting function (obsolete) __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] float4 lit(float n_dot_l, float n_dot_h, float m) { let ambient = 1.0f; @@ -6433,6 +6878,7 @@ __target_intrinsic(cuda, "$P_log($0)") __target_intrinsic(cpp, "$P_log($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Log _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T log(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -6440,6 +6886,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Log _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> log(vector<T, N> x) { VECTOR_MAP_UNARY(T, N, log, x); @@ -6448,6 +6895,7 @@ vector<T, N> log(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> log(matrix<T, N, M> x) { MATRIX_MAP_UNARY(T, N, M, log, x); @@ -6461,6 +6909,7 @@ __target_intrinsic(cuda, "$P_log10($0)") __target_intrinsic(cpp, "$P_log10($0)") __target_intrinsic(spirv, "%baseElog = OpExtInst resultType resultId glsl450 Log _0; OpFMul resultType resultId %baseElog const(_p,0.43429448190325182765112891891661)") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T log10(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -6468,6 +6917,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "(log( $0 ) * $S0(0.43429448190325182765112891891661) )" ) __target_intrinsic(spirv, "%baseElog = OpExtInst resultType resultId glsl450 Log _0; OpVectorTimesScalar resultType resultId %baseElog const(_p,0.43429448190325182765112891891661)") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T,N> log10(vector<T,N> x) { VECTOR_MAP_UNARY(T, N, log10, x); @@ -6476,6 +6926,7 @@ vector<T,N> log10(vector<T,N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,N,M> log10(matrix<T,N,M> x) { MATRIX_MAP_UNARY(T, N, M, log10, x); @@ -6489,6 +6940,7 @@ __target_intrinsic(cuda, "$P_log2($0)") __target_intrinsic(cpp, "$P_log2($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Log2 _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T log2(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -6496,6 +6948,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Log2 _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T,N> log2(vector<T,N> x) { VECTOR_MAP_UNARY(T, N, log2, x); @@ -6504,6 +6957,7 @@ vector<T,N> log2(vector<T,N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,N,M> log2(matrix<T,N,M> x) { MATRIX_MAP_UNARY(T, N, M, log2, x); @@ -6518,6 +6972,7 @@ __target_intrinsic(cuda, "$P_fma($0, $1, $2)") __target_intrinsic(cpp, "$P_fma($0, $1, $2)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] T mad(T mvalue, T avalue, T bvalue); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -6525,6 +6980,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, fma) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] vector<T, N> mad(vector<T, N> mvalue, vector<T, N> avalue, vector<T, N> bvalue) { VECTOR_MAP_TRINARY(T, N, mad, mvalue, avalue, bvalue); @@ -6533,6 +6989,7 @@ vector<T, N> mad(vector<T, N> mvalue, vector<T, N> avalue, vector<T, N> bvalue) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] matrix<T, N, M> mad(matrix<T, N, M> mvalue, matrix<T, N, M> avalue, matrix<T, N, M> bvalue) { MATRIX_MAP_TRINARY(T, N, M, mad, mvalue, avalue, bvalue); @@ -6545,6 +7002,7 @@ __target_intrinsic(cuda, "$P_fma($0, $1, $2)") __target_intrinsic(cpp, "$P_fma($0, $1, $2)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] T mad(T mvalue, T avalue, T bvalue); __generic<T : __BuiltinIntegerType, let N : int> @@ -6552,6 +7010,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, fma) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Fma _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] vector<T, N> mad(vector<T, N> mvalue, vector<T, N> avalue, vector<T, N> bvalue) { VECTOR_MAP_TRINARY(T, N, mad, mvalue, avalue, bvalue); @@ -6560,6 +7019,7 @@ vector<T, N> mad(vector<T, N> mvalue, vector<T, N> avalue, vector<T, N> bvalue) __generic<T : __BuiltinIntegerType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] matrix<T, N, M> mad(matrix<T, N, M> mvalue, matrix<T, N, M> avalue, matrix<T, N, M> bvalue) { MATRIX_MAP_TRINARY(T, N, M, mad, mvalue, avalue, bvalue); @@ -6574,6 +7034,7 @@ __target_intrinsic(cuda, "$P_max($0, $1)") __target_intrinsic(cpp, "$P_max($0, $1)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMax, UMax, SMax) _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T max(T x, T y); // Note: a stdlib implementation of `max` (or `min`) will require splitting // floating-point and integer cases apart, because the floating-point @@ -6585,6 +7046,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMax, UMax, SMax) _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> max(vector<T, N> x, vector<T, N> y) { VECTOR_MAP_BINARY(T, N, max, x, y); @@ -6593,6 +7055,7 @@ vector<T, N> max(vector<T, N> x, vector<T, N> y) __generic<T : __BuiltinIntegerType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> max(matrix<T, N, M> x, matrix<T, N, M> y) { MATRIX_MAP_BINARY(T, N, M, max, x, y); @@ -6605,6 +7068,7 @@ __target_intrinsic(cuda, "$P_max($0, $1)") __target_intrinsic(cpp, "$P_max($0, $1)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMax, UMax, SMax) _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T max(T x, T y); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -6612,6 +7076,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMax, UMax, SMax) _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> max(vector<T, N> x, vector<T, N> y) { VECTOR_MAP_BINARY(T, N, max, x, y); @@ -6620,6 +7085,7 @@ vector<T, N> max(vector<T, N> x, vector<T, N> y) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> max(matrix<T, N, M> x, matrix<T, N, M> y) { MATRIX_MAP_BINARY(T, N, M, max, x, y); @@ -6633,6 +7099,7 @@ __target_intrinsic(cuda, "$P_min($0, $1)") __target_intrinsic(cpp, "$P_min($0, $1)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMin, UMin, SMin) _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T min(T x, T y); __generic<T : __BuiltinIntegerType, let N : int> @@ -6640,6 +7107,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMin, UMin, SMin) _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T,N> min(vector<T,N> x, vector<T,N> y) { VECTOR_MAP_BINARY(T, N, min, x, y); @@ -6648,6 +7116,7 @@ vector<T,N> min(vector<T,N> x, vector<T,N> y) __generic<T : __BuiltinIntegerType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,N,M> min(matrix<T,N,M> x, matrix<T,N,M> y) { MATRIX_MAP_BINARY(T, N, M, min, x, y); @@ -6660,6 +7129,7 @@ __target_intrinsic(cuda, "$P_min($0, $1)") __target_intrinsic(cpp, "$P_min($0, $1)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMin, UMin, SMin) _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T min(T x, T y); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -6667,6 +7137,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 fus(FMin, UMin, SMin) _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T,N> min(vector<T,N> x, vector<T,N> y) { VECTOR_MAP_BINARY(T, N, min, x, y); @@ -6675,6 +7146,7 @@ vector<T,N> min(vector<T,N> x, vector<T,N> y) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,N,M> min(matrix<T,N,M> x, matrix<T,N,M> y) { MATRIX_MAP_BINARY(T, N, M, min, x, y); @@ -6686,12 +7158,14 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Modf _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T modf(T x, out T ip); __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T,N> modf(vector<T,N> x, out vector<T,N> ip) { VECTOR_MAP_BINARY(T, N, modf, x, ip); @@ -6700,6 +7174,7 @@ vector<T,N> modf(vector<T,N> x, out vector<T,N> ip) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int, let L : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,N,M> modf(matrix<T,N,M> x, out matrix<T,N,M,L> ip) { MATRIX_MAP_BINARY(T, N, M, modf, x, ip); @@ -6708,6 +7183,7 @@ matrix<T,N,M> modf(matrix<T,N,M> x, out matrix<T,N,M,L> ip) // msad4 (whatever that is) __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] uint4 msad4(uint reference, uint2 source, uint4 accum) { int4 bytesRef = (reference >> uint4(24, 16, 8, 0)) & 0xFF; @@ -6730,28 +7206,33 @@ uint4 msad4(uint reference, uint2 source, uint4 accum) __generic<T : __BuiltinArithmeticType> __intrinsic_op($(kIROp_Mul)) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T mul(T x, T y); // scalar-vector and vector-scalar __generic<T : __BuiltinArithmeticType, let N : int> __intrinsic_op($(kIROp_Mul)) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> mul(vector<T, N> x, T y); __generic<T : __BuiltinArithmeticType, let N : int> __intrinsic_op($(kIROp_Mul)) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> mul(T x, vector<T, N> y); // scalar-matrix and matrix-scalar __generic<T : __BuiltinArithmeticType, let N : int, let M :int> __intrinsic_op($(kIROp_Mul)) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> mul(matrix<T, N, M> x, T y); __generic<T : __BuiltinArithmeticType, let N : int, let M :int> __intrinsic_op($(kIROp_Mul)) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> mul(T x, matrix<T, N, M> y); // vector-vector (dot product) @@ -6759,6 +7240,7 @@ __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "dot") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T mul(vector<T, N> x, vector<T, N> y) { return dot(x, y); @@ -6766,6 +7248,7 @@ T mul(vector<T, N> x, vector<T, N> y) __generic<T : __BuiltinIntegerType, let N : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T mul(vector<T, N> x, vector<T, N> y) { return dot(x, y); @@ -6777,6 +7260,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "($1 * $0)") __target_intrinsic(spirv, "OpMatrixTimesVector resultType resultId _1 _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, M> mul(vector<T, N> left, matrix<T, N, M> right) { vector<T,M> result; @@ -6795,6 +7279,7 @@ __generic<T : __BuiltinIntegerType, let N : int, let M : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "($1 * $0)") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, M> mul(vector<T, N> left, matrix<T, N, M> right) { vector<T,M> result; @@ -6813,6 +7298,7 @@ __generic<T : __BuiltinLogicalType, let N : int, let M : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "($1 * $0)") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, M> mul(vector<T, N> left, matrix<T, N, M> right) { vector<T,M> result; @@ -6834,6 +7320,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "($1 * $0)") __target_intrinsic(spirv, "OpVectorTimesMatrix resultType resultId _1 _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T,N> mul(matrix<T,N,M> left, vector<T,M> right) { vector<T,N> result; @@ -6852,6 +7339,7 @@ __generic<T : __BuiltinIntegerType, let N : int, let M : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "($1 * $0)") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T,N> mul(matrix<T,N,M> left, vector<T,M> right) { vector<T,N> result; @@ -6870,6 +7358,7 @@ __generic<T : __BuiltinLogicalType, let N : int, let M : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "($1 * $0)") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T,N> mul(matrix<T,N,M> left, vector<T,M> right) { vector<T,N> result; @@ -6891,6 +7380,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "($1 * $0)") __target_intrinsic(spirv, "OpMatrixTimesMatrix resultType resultId _1 _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,R,C> mul(matrix<T,R,N> left, matrix<T,N,C> right) { matrix<T,R,C> result; @@ -6910,6 +7400,7 @@ __generic<T : __BuiltinIntegerType, let R : int, let N : int, let C : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "($1 * $0)") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,R,C> mul(matrix<T,R,N> left, matrix<T,N,C> right) { matrix<T,R,C> result; @@ -6929,6 +7420,7 @@ __generic<T : __BuiltinLogicalType, let R : int, let N : int, let C : int> __target_intrinsic(hlsl) __target_intrinsic(glsl, "($1 * $0)") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,R,C> mul(matrix<T,R,N> left, matrix<T,N,C> right) { matrix<T,R,C> result; @@ -6980,6 +7472,7 @@ __generic<let N : int> float noise(vector<float, N> x) /// the user's responsibility, so that the default behavior of the language /// is more semantically "correct." [ForceInline] +[require(spirv)] T __copyObject<T>(T v) { __target_switch { @@ -6992,9 +7485,9 @@ T __copyObject<T>(T v) /// `NonUniformResourceIndex` function is used to indicate if the resource index is /// divergent, and ensure scalarization happens correctly for each divergent lane. -[__readNone] __generic<T:__BuiltinArithmeticType> __intrinsic_op($(kIROp_NonUniformResourceIndex)) +[require(cpp_cuda_glsl_hlsl_spirv, nonuniformqualifier)] T NonUniformResourceIndex(T index); /// HLSL allows NonUniformResourceIndex around non int/uint types. @@ -7010,6 +7503,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Normalize _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T,N> normalize(vector<T,N> x) { return x / length(x); @@ -7020,6 +7514,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Normalize _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T normalize(T x) { return x / length(x); @@ -7033,6 +7528,7 @@ __target_intrinsic(cuda, "$P_pow($0, $1)") __target_intrinsic(cpp, "$P_pow($0, $1)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Pow _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T pow(T x, T y); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -7040,6 +7536,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Pow _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> pow(vector<T, N> x, vector<T, N> y) { VECTOR_MAP_BINARY(T, N, pow, x, y); @@ -7048,6 +7545,7 @@ vector<T, N> pow(vector<T, N> x, vector<T, N> y) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,N,M> pow(matrix<T,N,M> x, matrix<T,N,M> y) { MATRIX_MAP_BINARY(T, N, M, pow, x, y); @@ -7082,6 +7580,7 @@ for (int argCount = 0; argCount < 12; argCount++) auto spirvArgs = spirvArgList.toString(); }}}} __glsl_extension(GL_EXT_debug_printf) +[require(cpp_cuda_glsl_hlsl_spirv, printf)] void printf$(genericParamList.toString())(NativeString format $(paramList)) { __target_switch @@ -7105,6 +7604,7 @@ ${{{{ // Tessellation factor fixup routines +[require(hlsl, sm_5_0)] void Process2DQuadTessFactorsAvg( in float4 RawEdgeFactors, in float2 InsideScale, @@ -7112,6 +7612,7 @@ void Process2DQuadTessFactorsAvg( out float2 RoundedInsideTessFactors, out float2 UnroundedInsideTessFactors); +[require(hlsl, sm_5_0)] void Process2DQuadTessFactorsMax( in float4 RawEdgeFactors, in float2 InsideScale, @@ -7119,6 +7620,7 @@ void Process2DQuadTessFactorsMax( out float2 RoundedInsideTessFactors, out float2 UnroundedInsideTessFactors); +[require(hlsl, sm_5_0)] void Process2DQuadTessFactorsMin( in float4 RawEdgeFactors, in float2 InsideScale, @@ -7126,12 +7628,14 @@ void Process2DQuadTessFactorsMin( out float2 RoundedInsideTessFactors, out float2 UnroundedInsideTessFactors); +[require(hlsl, sm_5_0)] void ProcessIsolineTessFactors( in float RawDetailFactor, in float RawDensityFactor, out float RoundedDetailFactor, out float RoundedDensityFactor); +[require(hlsl, sm_5_0)] void ProcessQuadTessFactorsAvg( in float4 RawEdgeFactors, in float InsideScale, @@ -7139,6 +7643,7 @@ void ProcessQuadTessFactorsAvg( out float2 RoundedInsideTessFactors, out float2 UnroundedInsideTessFactors); +[require(hlsl, sm_5_0)] void ProcessQuadTessFactorsMax( in float4 RawEdgeFactors, in float InsideScale, @@ -7146,6 +7651,7 @@ void ProcessQuadTessFactorsMax( out float2 RoundedInsideTessFactors, out float2 UnroundedInsideTessFactors); +[require(hlsl, sm_5_0)] void ProcessQuadTessFactorsMin( in float4 RawEdgeFactors, in float InsideScale, @@ -7153,6 +7659,7 @@ void ProcessQuadTessFactorsMin( out float2 RoundedInsideTessFactors, out float2 UnroundedInsideTessFactors); +[require(hlsl, sm_5_0)] void ProcessTriTessFactorsAvg( in float3 RawEdgeFactors, in float InsideScale, @@ -7160,6 +7667,7 @@ void ProcessTriTessFactorsAvg( out float RoundedInsideTessFactor, out float UnroundedInsideTessFactor); +[require(hlsl, sm_5_0)] void ProcessTriTessFactorsMax( in float3 RawEdgeFactors, in float InsideScale, @@ -7167,6 +7675,7 @@ void ProcessTriTessFactorsMax( out float RoundedInsideTessFactor, out float UnroundedInsideTessFactor); +[require(hlsl, sm_5_0)] void ProcessTriTessFactorsMin( in float3 RawEdgeFactors, in float InsideScale, @@ -7180,6 +7689,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Radians _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T radians(T x) { return x * (T.getPi() / T(180.0f)); @@ -7190,6 +7700,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Radians _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> radians(vector<T, N> x) { return x * (T.getPi() / T(180.0f)); @@ -7198,6 +7709,7 @@ vector<T, N> radians(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> radians(matrix<T, N, M> x) { return x * (T.getPi() / T(180.0f)); @@ -7207,6 +7719,7 @@ matrix<T, N, M> radians(matrix<T, N, M> x) __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T rcp(T x) { return T(1.0) / x; @@ -7215,6 +7728,7 @@ T rcp(T x) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> rcp(vector<T, N> x) { VECTOR_MAP_UNARY(T, N, rcp, x); @@ -7223,6 +7737,7 @@ vector<T, N> rcp(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> rcp(matrix<T, N, M> x) { MATRIX_MAP_UNARY(T, N, M, rcp, x); @@ -7234,6 +7749,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Reflect _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T reflect(T i, T n) { return i - T(2) * dot(n,i) * n; @@ -7244,6 +7760,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Reflect _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T,N> reflect(vector<T,N> i, vector<T,N> n) { return i - T(2) * dot(n,i) * n; @@ -7255,6 +7772,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Refract _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T,N> refract(vector<T,N> i, vector<T,N> n, T eta) { let dotNI = dot(n,i); @@ -7268,6 +7786,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Refract _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T refract(T i, T n, T eta) { let dotNI = dot(n,i); @@ -7278,6 +7797,7 @@ T refract(T i, T n, T eta) // Reverse order of bits [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] uint reversebits(uint value) { __target_switch @@ -7297,6 +7817,7 @@ uint reversebits(uint value) __target_intrinsic(glsl, "bitfieldReverse") __generic<let N : int> [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, shader5_sm_5_0)] vector<uint, N> reversebits(vector<uint, N> value) { __target_switch @@ -7318,6 +7839,7 @@ __target_intrinsic(cuda, "$P_round($0)") __target_intrinsic(cpp, "$P_round($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Round _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T round(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -7325,6 +7847,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Round _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> round(vector<T, N> x) { VECTOR_MAP_UNARY(T, N, round, x); @@ -7333,6 +7856,7 @@ vector<T, N> round(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,N,M> round(matrix<T,N,M> x) { MATRIX_MAP_UNARY(T, N, M, round, x); @@ -7346,6 +7870,7 @@ __target_intrinsic(cuda, "$P_rsqrt($0)") __target_intrinsic(cpp, "$P_rsqrt($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 InverseSqrt _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T rsqrt(T x) { return T(1.0) / sqrt(x); @@ -7356,6 +7881,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl, "inversesqrt($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 InverseSqrt _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> rsqrt(vector<T, N> x) { VECTOR_MAP_UNARY(T, N, rsqrt, x); @@ -7364,6 +7890,7 @@ vector<T, N> rsqrt(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> rsqrt(matrix<T, N, M> x) { MATRIX_MAP_UNARY(T, N, M, rsqrt, x); @@ -7374,6 +7901,7 @@ matrix<T, N, M> rsqrt(matrix<T, N, M> x) __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T saturate(T x) { return clamp<T>(x, T(0), T(1)); @@ -7382,6 +7910,7 @@ T saturate(T x) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T,N> saturate(vector<T,N> x) { return clamp<T,N>(x, @@ -7392,6 +7921,7 @@ vector<T,N> saturate(vector<T,N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,N,M> saturate(matrix<T,N,M> x) { MATRIX_MAP_UNARY(T, N, M, saturate, x); @@ -7469,6 +7999,7 @@ __target_intrinsic(cuda, "$P_sin($0)") __target_intrinsic(cpp, "$P_sin($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Sin _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T sin(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -7476,6 +8007,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Sin _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> sin(vector<T, N> x) { VECTOR_MAP_UNARY(T, N, sin, x); @@ -7484,6 +8016,7 @@ vector<T, N> sin(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> sin(matrix<T, N, M> x) { MATRIX_MAP_UNARY(T, N, M, sin, x); @@ -7494,6 +8027,7 @@ __generic<T : __BuiltinFloatingPointType> __target_intrinsic(hlsl) __target_intrinsic(cuda, "$P_sincos($0, $1, $2)") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] void sincos(T x, out T s, out T c) { s = sin(x); @@ -7503,6 +8037,7 @@ void sincos(T x, out T s, out T c) __generic<T : __BuiltinFloatingPointType, let N : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] void sincos(vector<T,N> x, out vector<T,N> s, out vector<T,N> c) { s = sin(x); @@ -7512,6 +8047,7 @@ void sincos(vector<T,N> x, out vector<T,N> s, out vector<T,N> c) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int, let L1: int, let L2 : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] void sincos(matrix<T,N,M> x, out matrix<T,N,M,L1> s, out matrix<T,N,M,L2> c) { s = sin(x); @@ -7526,6 +8062,7 @@ __target_intrinsic(cuda, "$P_sinh($0)") __target_intrinsic(cpp, "$P_sinh($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Sinh _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T sinh(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -7533,6 +8070,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Sinh _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> sinh(vector<T, N> x) { VECTOR_MAP_UNARY(T, N, sinh, x); @@ -7541,6 +8079,7 @@ vector<T, N> sinh(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> sinh(matrix<T, N, M> x) { MATRIX_MAP_UNARY(T, N, M, sinh, x); @@ -7552,6 +8091,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 SmoothStep _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T smoothstep(T min, T max, T x) { let t = saturate((x - min) / (max - min)); @@ -7563,6 +8103,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 SmoothStep _0 _1 _2") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> smoothstep(vector<T, N> min, vector<T, N> max, vector<T, N> x) { VECTOR_MAP_TRINARY(T, N, smoothstep, min, max, x); @@ -7571,6 +8112,7 @@ vector<T, N> smoothstep(vector<T, N> min, vector<T, N> max, vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> smoothstep(matrix<T, N, M> min, matrix<T, N, M> max, matrix<T, N, M> x) { MATRIX_MAP_TRINARY(T, N, M, smoothstep, min, max, x); @@ -7584,6 +8126,7 @@ __target_intrinsic(cuda, "$P_sqrt($0)") __target_intrinsic(cpp, "$P_sqrt($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Sqrt _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T sqrt(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -7591,6 +8134,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Sqrt _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> sqrt(vector<T, N> x) { VECTOR_MAP_UNARY(T, N, sqrt, x); @@ -7599,6 +8143,7 @@ vector<T, N> sqrt(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> sqrt(matrix<T, N, M> x) { MATRIX_MAP_UNARY(T, N, M, sqrt, x); @@ -7610,6 +8155,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Step _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T step(T y, T x) { return x < y ? T(0.0f) : T(1.0f); @@ -7620,6 +8166,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Step _0 _1") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T,N> step(vector<T,N> y, vector<T,N> x) { VECTOR_MAP_BINARY(T, N, step, y, x); @@ -7628,6 +8175,7 @@ vector<T,N> step(vector<T,N> y, vector<T,N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> step(matrix<T, N, M> y, matrix<T, N, M> x) { MATRIX_MAP_BINARY(T, N, M, step, y, x); @@ -7641,6 +8189,7 @@ __target_intrinsic(cuda, "$P_tan($0)") __target_intrinsic(cpp, "$P_tan($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Tan _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T tan(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -7648,6 +8197,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Tan _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> tan(vector<T, N> x) { VECTOR_MAP_UNARY(T, N, tan, x); @@ -7656,6 +8206,7 @@ vector<T, N> tan(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> tan(matrix<T, N, M> x) { MATRIX_MAP_UNARY(T, N, M, tan, x); @@ -7669,6 +8220,7 @@ __target_intrinsic(cuda, "$P_tanh($0)") __target_intrinsic(cpp, "$P_tanh($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Tanh _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T tanh(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -7676,6 +8228,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Tanh _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T,N> tanh(vector<T,N> x) { VECTOR_MAP_UNARY(T, N, tanh, x); @@ -7684,6 +8237,7 @@ vector<T,N> tanh(vector<T,N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T,N,M> tanh(matrix<T,N,M> x) { MATRIX_MAP_UNARY(T, N, M, tanh, x); @@ -7695,6 +8249,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpTranspose resultType resultId _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] [PreferRecompute] matrix<T, M, N> transpose(matrix<T, N, M> x) { @@ -7709,6 +8264,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpTranspose resultType resultId _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] [PreferRecompute] matrix<T, M, N> transpose(matrix<T, N, M> x) { @@ -7723,6 +8279,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpTranspose resultType resultId _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] [PreferRecompute] [OverloadRank(-1)] matrix<T, M, N> transpose(matrix<T, N, M> x) @@ -7742,6 +8299,7 @@ __target_intrinsic(cuda, "$P_trunc($0)") __target_intrinsic(cpp, "$P_trunc($0)") __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Trunc _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] T trunc(T x); __generic<T : __BuiltinFloatingPointType, let N : int> @@ -7749,6 +8307,7 @@ __target_intrinsic(hlsl) __target_intrinsic(glsl) __target_intrinsic(spirv, "OpExtInst resultType resultId glsl450 Trunc _0") [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] vector<T, N> trunc(vector<T, N> x) { VECTOR_MAP_UNARY(T, N, trunc, x); @@ -7757,6 +8316,7 @@ vector<T, N> trunc(vector<T, N> x) __generic<T : __BuiltinFloatingPointType, let N : int, let M : int> __target_intrinsic(hlsl) [__readNone] +[require(cpp_cuda_glsl_hlsl_spirv, sm_2_0_GLSL_140)] matrix<T, N, M> trunc(matrix<T, N, M> x) { MATRIX_MAP_UNARY(T, N, M, trunc, x); @@ -7768,6 +8328,7 @@ typedef uint WaveMask; __glsl_extension(GL_KHR_shader_subgroup_ballot) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] WaveMask WaveGetConvergedMask() { __target_switch @@ -7793,6 +8354,7 @@ WaveMask __WaveGetActiveMask(); __glsl_extension(GL_KHR_shader_subgroup_ballot) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] WaveMask WaveGetActiveMask() { __target_switch @@ -7808,13 +8370,14 @@ WaveMask WaveGetActiveMask() OpCapability GroupNonUniformBallot; OpGroupNonUniformBallot $$uint4 result Subgroup $_true }).x; - default: + case cuda: return __WaveGetActiveMask(); } } __glsl_extension(GL_KHR_shader_subgroup_basic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_basic)] bool WaveMaskIsFirstLane(WaveMask mask) { __target_switch @@ -7831,13 +8394,12 @@ bool WaveMaskIsFirstLane(WaveMask mask) OpCapability GroupNonUniformBallot; OpGroupNonUniformElect $$bool result Subgroup }; - default: - return false; } } __glsl_extension(GL_KHR_shader_subgroup_vote) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_vote)] bool WaveMaskAllTrue(WaveMask mask, bool condition) { __target_switch @@ -7854,13 +8416,12 @@ bool WaveMaskAllTrue(WaveMask mask, bool condition) OpCapability GroupNonUniformBallot; OpGroupNonUniformAll $$bool result Subgroup $condition }; - default: - return false; } } __glsl_extension(GL_KHR_shader_subgroup_vote) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_vote)] bool WaveMaskAnyTrue(WaveMask mask, bool condition) { __target_switch @@ -7877,13 +8438,12 @@ bool WaveMaskAnyTrue(WaveMask mask, bool condition) OpCapability GroupNonUniformBallot; OpGroupNonUniformAny $$bool result Subgroup $condition }; - default: - return false; } } __glsl_extension(GL_KHR_shader_subgroup_ballot) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] WaveMask WaveMaskBallot(WaveMask mask, bool condition) { __target_switch @@ -7900,11 +8460,10 @@ WaveMask WaveMaskBallot(WaveMask mask, bool condition) OpCapability GroupNonUniformBallot; OpGroupNonUniformBallot $$uint4 result Subgroup $condition }).x; - default: - return 0; } } +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] uint WaveMaskCountBits(WaveMask mask, bool value) { __target_switch @@ -7933,6 +8492,7 @@ uint WaveMaskCountBits(WaveMask mask, bool value) // It seems this can only mean the active threads are the "threads the program flow would lead to". This implies a lockstep // "straight SIMD" style interpretation. That being the case this op on HLSL is just a memory barrier without any Sync. +[require(cuda_glsl_hlsl_spirv, memorybarrier_compute)] void AllMemoryBarrierWithWaveMaskSync(WaveMask mask) { __target_switch @@ -7964,6 +8524,7 @@ void AllMemoryBarrierWithWaveMaskSync(WaveMask mask) // aspect of HLSL seems to make everything in lock step - but that's not quite so, it only has to apparently be that way as far as the programmers // model appears - divergence could perhaps potentially still happen. +[require(cuda_glsl_hlsl_spirv, memorybarrier_compute)] void GroupMemoryBarrierWithWaveMaskSync(WaveMask mask) { __target_switch @@ -7979,6 +8540,7 @@ void GroupMemoryBarrierWithWaveMaskSync(WaveMask mask) } } +[require(cuda_glsl_hlsl_spirv, memorybarrier_compute)] void AllMemoryBarrierWithWaveSync() { __target_switch @@ -7994,6 +8556,7 @@ void AllMemoryBarrierWithWaveSync() } } +[require(cuda_glsl_hlsl_spirv, memorybarrier_compute)] void GroupMemoryBarrierWithWaveSync() { __target_switch @@ -8019,6 +8582,7 @@ void GroupMemoryBarrierWithWaveSync() __generic<T : __BuiltinType> __glsl_extension(GL_KHR_shader_subgroup_ballot) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] T WaveMaskBroadcastLaneAt(WaveMask mask, T value, constexpr int lane) { __target_switch @@ -8038,6 +8602,7 @@ T WaveMaskBroadcastLaneAt(WaveMask mask, T value, constexpr int lane) __generic<T : __BuiltinType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_ballot) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] vector<T,N> WaveMaskBroadcastLaneAt(WaveMask mask, vector<T,N> value, constexpr int lane) { __target_switch @@ -8063,6 +8628,7 @@ matrix<T,N,M> WaveMaskBroadcastLaneAt(WaveMask mask, matrix<T,N,M> value, conste __generic<T : __BuiltinType> __glsl_extension(GL_KHR_shader_subgroup_shuffle) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_shuffle)] T WaveMaskReadLaneAt(WaveMask mask, T value, int lane) { __target_switch @@ -8079,8 +8645,9 @@ T WaveMaskReadLaneAt(WaveMask mask, T value, int lane) } } __generic<T : __BuiltinType, let N : int> -__spirv_version(1.3)__glsl_extension(GL_KHR_shader_subgroup_shuffle) +__glsl_extension(GL_KHR_shader_subgroup_shuffle) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_shuffle)] vector<T,N> WaveMaskReadLaneAt(WaveMask mask, vector<T,N> value, int lane) { __target_switch @@ -8125,6 +8692,7 @@ matrix<T,N,M> WaveMaskShuffle(WaveMask mask, matrix<T,N,M> value, int lane) __glsl_extension(GL_KHR_shader_subgroup_ballot) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] uint WaveMaskPrefixCountBits(WaveMask mask, bool value) { __target_switch @@ -8147,6 +8715,7 @@ uint WaveMaskPrefixCountBits(WaveMask mask, bool value) __generic<T : __BuiltinIntegerType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WaveMaskBitAnd(WaveMask mask, T expr) { __target_switch @@ -8165,6 +8734,7 @@ T WaveMaskBitAnd(WaveMask mask, T expr) __generic<T : __BuiltinIntegerType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T,N> WaveMaskBitAnd(WaveMask mask, vector<T,N> expr) { __target_switch @@ -8187,6 +8757,7 @@ matrix<T,N,M> WaveMaskBitAnd(WaveMask mask, matrix<T,N,M> expr); __generic<T : __BuiltinIntegerType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WaveMaskBitOr(WaveMask mask, T expr) { __target_switch @@ -8204,6 +8775,7 @@ T WaveMaskBitOr(WaveMask mask, T expr) __generic<T : __BuiltinIntegerType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T,N> WaveMaskBitOr(WaveMask mask, vector<T,N> expr) { __target_switch @@ -8226,6 +8798,7 @@ matrix<T,N,M> WaveMaskBitOr(WaveMask mask, matrix<T,N,M> expr); __generic<T : __BuiltinIntegerType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WaveMaskBitXor(WaveMask mask, T expr) { __target_switch @@ -8243,6 +8816,7 @@ T WaveMaskBitXor(WaveMask mask, T expr) __generic<T : __BuiltinIntegerType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T,N> WaveMaskBitXor(WaveMask mask, vector<T,N> expr) { __target_switch @@ -8265,6 +8839,7 @@ matrix<T,N,M> WaveMaskBitXor(WaveMask mask, matrix<T,N,M> expr); __generic<T : __BuiltinArithmeticType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WaveMaskMax(WaveMask mask, T expr) { __target_switch @@ -8285,6 +8860,7 @@ T WaveMaskMax(WaveMask mask, T expr) __generic<T : __BuiltinArithmeticType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T,N> WaveMaskMax(WaveMask mask, vector<T,N> expr) { __target_switch @@ -8311,6 +8887,7 @@ matrix<T,N,M> WaveMaskMax(WaveMask mask, matrix<T,N,M> expr); __generic<T : __BuiltinArithmeticType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WaveMaskMin(WaveMask mask, T expr) { __target_switch @@ -8332,6 +8909,7 @@ T WaveMaskMin(WaveMask mask, T expr) __generic<T : __BuiltinArithmeticType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T,N> WaveMaskMin(WaveMask mask, vector<T,N> expr) { __target_switch @@ -8358,6 +8936,7 @@ matrix<T,N,M> WaveMaskMin(WaveMask mask, matrix<T,N,M> expr); __generic<T : __BuiltinArithmeticType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WaveMaskProduct(WaveMask mask, T expr) { __target_switch @@ -8383,6 +8962,7 @@ T WaveMaskProduct(WaveMask mask, T expr) __generic<T : __BuiltinArithmeticType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T,N> WaveMaskProduct(WaveMask mask, vector<T,N> expr) { __target_switch @@ -8410,12 +8990,10 @@ __target_intrinsic(cuda, "_waveProductMultiple($0, $1)") __target_intrinsic(hlsl, "WaveActiveProduct($1)") matrix<T,N,M> WaveMaskProduct(WaveMask mask, matrix<T,N,M> expr); -__intrinsic_op($(kIROp_RequireGLSLExtension)) -void __requireGLSLExtension(String extensionName); - __generic<T : __BuiltinArithmeticType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WaveMaskSum(WaveMask mask, T expr) { __target_switch @@ -8442,6 +9020,7 @@ T WaveMaskSum(WaveMask mask, T expr) __generic<T : __BuiltinArithmeticType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T,N> WaveMaskSum(WaveMask mask, vector<T,N> expr) { __target_switch @@ -8474,6 +9053,7 @@ __generic<T : __BuiltinType> __glsl_extension(GL_KHR_shader_subgroup_vote) __spirv_version(1.3) __cuda_sm_version(7.0) +[require(cuda_glsl_hlsl_spirv, subgroup_vote)] bool WaveMaskAllEqual(WaveMask mask, T value) { __target_switch @@ -8491,14 +9071,13 @@ bool WaveMaskAllEqual(WaveMask mask, T value) OpCapability GroupNonUniformVote; OpGroupNonUniformAllEqual $$bool result Subgroup $value }; - default: - return false; } } __generic<T : __BuiltinType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_vote) __spirv_version(1.3) __cuda_sm_version(7.0) +[require(cuda_glsl_hlsl_spirv, subgroup_vote)] bool WaveMaskAllEqual(WaveMask mask, vector<T,N> value) { __target_switch @@ -8516,8 +9095,6 @@ bool WaveMaskAllEqual(WaveMask mask, vector<T,N> value) OpCapability GroupNonUniformVote; OpGroupNonUniformAllEqual $$bool result Subgroup $value }; - default: - return false; } } __generic<T : __BuiltinType, let N : int, let M : int> @@ -8531,6 +9108,7 @@ bool WaveMaskAllEqual(WaveMask mask, matrix<T,N,M> value); __generic<T : __BuiltinArithmeticType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WaveMaskPrefixProduct(WaveMask mask, T expr) { __target_switch @@ -8557,6 +9135,7 @@ T WaveMaskPrefixProduct(WaveMask mask, T expr) __generic<T : __BuiltinArithmeticType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T,N> WaveMaskPrefixProduct(WaveMask mask, vector<T,N> expr) { __target_switch @@ -8588,6 +9167,7 @@ matrix<T,N,M> WaveMaskPrefixProduct(WaveMask mask, matrix<T,N,M> expr); __generic<T : __BuiltinArithmeticType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WaveMaskPrefixSum(WaveMask mask, T expr) { __target_switch @@ -8615,6 +9195,7 @@ T WaveMaskPrefixSum(WaveMask mask, T expr) __generic<T : __BuiltinArithmeticType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T,N> WaveMaskPrefixSum(WaveMask mask, vector<T,N> expr) { __target_switch @@ -8646,6 +9227,7 @@ matrix<T,N,M> WaveMaskPrefixSum(WaveMask mask, matrix<T,N,M> expr); __generic<T : __BuiltinType> __glsl_extension(GL_KHR_shader_subgroup_ballot) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] T WaveMaskReadLaneFirst(WaveMask mask, T expr) { __target_switch @@ -8660,6 +9242,7 @@ T WaveMaskReadLaneFirst(WaveMask mask, T expr) __generic<T : __BuiltinType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_ballot) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] vector<T,N> WaveMaskReadLaneFirst(WaveMask mask, vector<T,N> expr) { __target_switch @@ -8684,6 +9267,7 @@ __generic<T : __BuiltinType> __glsl_extension(GL_NV_shader_subgroup_partitioned) __spirv_version(1.1) __cuda_sm_version(7.0) +[require(cuda_glsl_hlsl_spirv, subgroup_partitioned)] WaveMask WaveMaskMatch(WaveMask mask, T value) { __target_switch @@ -8704,6 +9288,7 @@ __generic<T : __BuiltinType, let N : int> __glsl_extension(GL_NV_shader_subgroup_partitioned) __spirv_version(1.1) __cuda_sm_version(7.0) +[require(cuda_glsl_hlsl_spirv, subgroup_partitioned)] WaveMask WaveMaskMatch(WaveMask mask, vector<T,N> value) { __target_switch @@ -8733,6 +9318,7 @@ WaveMask WaveMaskMatch(WaveMask mask, matrix<T,N,M> value); __generic<T : __BuiltinArithmeticType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WaveMaskPrefixBitAnd(WaveMask mask, T expr) { __target_switch @@ -8748,6 +9334,7 @@ T WaveMaskPrefixBitAnd(WaveMask mask, T expr) __generic<T : __BuiltinArithmeticType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T,N> WaveMaskPrefixBitAnd(WaveMask mask, vector<T,N> expr) { __target_switch @@ -8768,6 +9355,7 @@ matrix<T,N,M> WaveMaskPrefixBitAnd(WaveMask mask, matrix<T,N,M> expr); __generic<T : __BuiltinArithmeticType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WaveMaskPrefixBitOr(WaveMask mask, T expr) { __target_switch @@ -8783,6 +9371,7 @@ T WaveMaskPrefixBitOr(WaveMask mask, T expr) __generic<T : __BuiltinArithmeticType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T,N> WaveMaskPrefixBitOr(WaveMask mask, vector<T,N> expr) { __target_switch @@ -8803,6 +9392,7 @@ matrix<T,N,M> WaveMaskPrefixBitOr(WaveMask mask, matrix<T,N,M> expr); __generic<T : __BuiltinArithmeticType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WaveMaskPrefixBitXor(WaveMask mask, T expr) { __target_switch @@ -8818,6 +9408,7 @@ T WaveMaskPrefixBitXor(WaveMask mask, T expr) __generic<T : __BuiltinArithmeticType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T,N> WaveMaskPrefixBitXor(WaveMask mask, vector<T,N> expr) { __target_switch @@ -8843,6 +9434,7 @@ matrix<T,N,M> WaveMaskPrefixBitXor(WaveMask mask, matrix<T,N,M> expr); __generic<T : __BuiltinType> __glsl_extension(GL_KHR_shader_subgroup_quad) __spirv_version(1.3) +[require(glsl_hlsl_spirv, subgroup_quad)] T QuadReadLaneAt(T sourceValue, uint quadLaneID) { __target_switch @@ -8861,6 +9453,7 @@ T QuadReadLaneAt(T sourceValue, uint quadLaneID) __generic<T : __BuiltinType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_quad) __spirv_version(1.3) +[require(glsl_hlsl_spirv, subgroup_quad)] vector<T,N> QuadReadLaneAt(vector<T,N> sourceValue, uint quadLaneID) { __target_switch @@ -8882,6 +9475,7 @@ __generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M> QuadReadLan __generic<T : __BuiltinType> __glsl_extension(GL_KHR_shader_subgroup_quad) __spirv_version(1.3) +[require(glsl_hlsl_spirv, subgroup_quad)] T QuadReadAcrossX(T localValue) { __target_switch @@ -8902,6 +9496,7 @@ T QuadReadAcrossX(T localValue) __generic<T : __BuiltinType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_quad) __spirv_version(1.3) +[require(glsl_hlsl_spirv, subgroup_quad)] vector<T,N> QuadReadAcrossX(vector<T,N> localValue) { __target_switch @@ -8923,6 +9518,7 @@ __generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M> QuadReadAcr __generic<T : __BuiltinType> __glsl_extension(GL_KHR_shader_subgroup_quad) __spirv_version(1.3) +[require(glsl_hlsl_spirv, subgroup_quad)] T QuadReadAcrossY(T localValue) { __target_switch @@ -8942,6 +9538,7 @@ T QuadReadAcrossY(T localValue) __generic<T : __BuiltinType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_quad) __spirv_version(1.3) +[require(glsl_hlsl_spirv, subgroup_quad)] vector<T,N> QuadReadAcrossY(vector<T,N> localValue) { __target_switch @@ -8964,6 +9561,7 @@ __generic<T : __BuiltinType, let N : int, let M : int> matrix<T,N,M> QuadReadAcr __generic<T : __BuiltinType> __glsl_extension(GL_KHR_shader_subgroup_quad) __spirv_version(1.3) +[require(glsl_hlsl_spirv, subgroup_quad)] T QuadReadAcrossDiagonal(T localValue) { __target_switch @@ -8983,6 +9581,7 @@ T QuadReadAcrossDiagonal(T localValue) __generic<T : __BuiltinType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_quad) __spirv_version(1.3) +[require(glsl_hlsl_spirv, subgroup_quad)] vector<T,N> QuadReadAcrossDiagonal(vector<T,N> localValue) { __target_switch @@ -9011,6 +9610,7 @@ for (auto opName : kWaveActiveBitOpEntries) { __generic<T : __BuiltinIntegerType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WaveActive$(opName.hlslName)(T expr) { __target_switch @@ -9027,6 +9627,7 @@ T WaveActive$(opName.hlslName)(T expr) __generic<T : __BuiltinIntegerType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T, N> WaveActive$(opName.hlslName)(vector<T, N> expr) { __target_switch @@ -9059,6 +9660,7 @@ for (const char* opName : kWaveActiveMinMaxNames) { __generic<T : __BuiltinArithmeticType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WaveActive$(opName)(T expr) { __target_switch @@ -9080,6 +9682,7 @@ T WaveActive$(opName)(T expr) __generic<T : __BuiltinArithmeticType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T, N> WaveActive$(opName)(vector<T, N> expr) { __target_switch @@ -9119,6 +9722,7 @@ for (auto opName : kWaveActivProductSumNames) { __generic<T : __BuiltinArithmeticType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WaveActive$(opName.hlslName)(T expr) { __target_switch @@ -9151,6 +9755,7 @@ __generic<T : __BuiltinArithmeticType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) __target_intrinsic(hlsl) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T,N> WaveActive$(opName.hlslName)(vector<T,N> expr) { __target_switch @@ -9192,6 +9797,7 @@ ${{{{ __generic<T : __BuiltinType> __glsl_extension(GL_KHR_shader_subgroup_vote) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_vote)] bool WaveActiveAllEqual(T value) { __target_switch @@ -9214,6 +9820,7 @@ bool WaveActiveAllEqual(T value) __generic<T : __BuiltinType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_vote) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_vote)] bool WaveActiveAllEqual(vector<T,N> value) { __target_switch @@ -9242,6 +9849,7 @@ bool WaveActiveAllEqual(matrix<T, N, M> value) __glsl_extension(GL_KHR_shader_subgroup_vote) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_vote)] bool WaveActiveAllTrue(bool condition) { __target_switch @@ -9263,6 +9871,7 @@ bool WaveActiveAllTrue(bool condition) __glsl_extension(GL_KHR_shader_subgroup_vote) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_vote)] bool WaveActiveAnyTrue(bool condition) { __target_switch @@ -9285,6 +9894,7 @@ bool WaveActiveAnyTrue(bool condition) __glsl_extension(GL_KHR_shader_subgroup_ballot) __spirv_version(1.3) [NonUniformReturn] +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] uint4 WaveActiveBallot(bool condition) { __target_switch @@ -9313,6 +9923,7 @@ uint WaveActiveCountBits(bool value) __glsl_extension(GL_KHR_shader_subgroup_basic) __spirv_version(1.3) [NonUniformReturn] +[require(cuda_glsl_hlsl_spirv, subgroup_basic)] uint WaveGetLaneCount() { __target_switch @@ -9332,6 +9943,7 @@ uint WaveGetLaneCount() __glsl_extension(GL_KHR_shader_subgroup_basic) __spirv_version(1.3) [NonUniformReturn] +[require(cuda_glsl_hlsl_spirv, subgroup_basic)] uint WaveGetLaneIndex() { __target_switch @@ -9351,6 +9963,7 @@ uint WaveGetLaneIndex() __glsl_extension(GL_KHR_shader_subgroup_basic) __spirv_version(1.3) [NonUniformReturn] +[require(cuda_glsl_hlsl_spirv, subgroup_basic)] bool WaveIsFirstLane() { __target_switch @@ -9372,6 +9985,8 @@ bool WaveIsFirstLane() // It's useful to have a wave uint4 version of countbits, because some wave functions return uint4. // This implementation tries to limit the amount of work required by the actual lane count. +__spirv_version(1.3) +[require(cpp_cuda_glsl_hlsl_spirv, subgroup_ballot)] uint _WaveCountBits(uint4 value) { __target_switch @@ -9402,6 +10017,7 @@ uint _WaveCountBits(uint4 value) __generic<T : __BuiltinArithmeticType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WavePrefixProduct(T expr) { __target_switch @@ -9434,6 +10050,7 @@ T WavePrefixProduct(T expr) __generic<T : __BuiltinArithmeticType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T,N> WavePrefixProduct(vector<T,N> expr) { __target_switch @@ -9461,6 +10078,7 @@ vector<T,N> WavePrefixProduct(vector<T,N> expr) __generic<T : __BuiltinArithmeticType, let N : int, let M : int> __target_intrinsic(hlsl) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] matrix<T, N, M> WavePrefixProduct(matrix<T, N, M> expr) { return WaveMaskPrefixProduct(WaveGetActiveMask(), expr); @@ -9469,6 +10087,7 @@ matrix<T, N, M> WavePrefixProduct(matrix<T, N, M> expr) __generic<T : __BuiltinArithmeticType> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] T WavePrefixSum(T expr) { __target_switch @@ -9497,6 +10116,7 @@ T WavePrefixSum(T expr) __generic<T : __BuiltinArithmeticType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_arithmetic)] vector<T,N> WavePrefixSum(vector<T,N> expr) { __target_switch @@ -9532,6 +10152,7 @@ matrix<T,N,M> WavePrefixSum(matrix<T,N,M> expr) __generic<T : __BuiltinType> __glsl_extension(GL_KHR_shader_subgroup_ballot) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] T WaveReadLaneFirst(T expr) { __target_switch @@ -9550,6 +10171,7 @@ T WaveReadLaneFirst(T expr) __generic<T : __BuiltinType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_ballot) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] vector<T,N> WaveReadLaneFirst(vector<T,N> expr) { __target_switch @@ -9581,6 +10203,7 @@ matrix<T,N,M> WaveReadLaneFirst(matrix<T,N,M> expr) __generic<T : __BuiltinType> __glsl_extension(GL_KHR_shader_subgroup_ballot) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] T WaveBroadcastLaneAt(T value, constexpr int lane) { __target_switch @@ -9600,6 +10223,7 @@ T WaveBroadcastLaneAt(T value, constexpr int lane) __generic<T : __BuiltinType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_ballot) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] vector<T,N> WaveBroadcastLaneAt(vector<T,N> value, constexpr int lane) { __target_switch @@ -9629,6 +10253,7 @@ matrix<T, N, M> WaveBroadcastLaneAt(matrix<T, N, M> value, constexpr int lane) __generic<T : __BuiltinType> __glsl_extension(GL_KHR_shader_subgroup_shuffle) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_shuffle)] T WaveReadLaneAt(T value, int lane) { __target_switch @@ -9648,6 +10273,7 @@ T WaveReadLaneAt(T value, int lane) __generic<T : __BuiltinType, let N : int> __spirv_version(1.3) __glsl_extension(GL_KHR_shader_subgroup_shuffle) +[require(cuda_glsl_hlsl_spirv, subgroup_shuffle)] vector<T,N> WaveReadLaneAt(vector<T,N> value, int lane) { __target_switch @@ -9678,6 +10304,7 @@ matrix<T, N, M> WaveReadLaneAt(matrix<T, N, M> value, int lane) __generic<T : __BuiltinType> __glsl_extension(GL_KHR_shader_subgroup_shuffle) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_shuffle)] T WaveShuffle(T value, int lane) { __target_switch @@ -9697,6 +10324,7 @@ T WaveShuffle(T value, int lane) __generic<T : __BuiltinType, let N : int> __glsl_extension(GL_KHR_shader_subgroup_shuffle) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_shuffle)] vector<T,N> WaveShuffle(vector<T,N> value, int lane) { __target_switch @@ -9722,6 +10350,7 @@ matrix<T, N, M> WaveShuffle(matrix<T, N, M> value, int lane) __glsl_extension(GL_KHR_shader_subgroup_ballot) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] uint WavePrefixCountBits(bool value) { __target_switch @@ -9743,6 +10372,7 @@ uint WavePrefixCountBits(bool value) __glsl_extension(GL_KHR_shader_subgroup_ballot) __spirv_version(1.3) +[require(cuda_glsl_hlsl_spirv, subgroup_ballot)] uint4 WaveGetConvergedMulti() { __target_switch @@ -9793,6 +10423,7 @@ uint4 WaveMatch(matrix<T,N,M> value) __target_intrinsic(hlsl) __target_intrinsic(cuda, "_popc(__ballot_sync(($1).x, $0) & _getLaneLtMask())") +[require(cuda_hlsl, waveprefix)] uint WaveMultiPrefixCountBits(bool value, uint4 mask); __generic<T : __BuiltinArithmeticType> @@ -9801,6 +10432,7 @@ __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) __target_intrinsic(glsl, "subgroupExclusiveAnd($0)") __target_intrinsic(cuda, "_wavePrefixAnd(_getMultiPrefixMask(($1).x), $0)") +[require(cuda_glsl_hlsl, waveprefix)] T WaveMultiPrefixBitAnd(T expr, uint4 mask); __target_intrinsic(hlsl) @@ -9809,32 +10441,37 @@ __spirv_version(1.3) __target_intrinsic(glsl, "subgroupExclusiveAnd($0)") __target_intrinsic(cuda, "_wavePrefixAndMultiple(_getMultiPrefixMask(($1).x), $0)") __generic<T : __BuiltinArithmeticType, let N : int> +[require(cuda_glsl_hlsl, waveprefix)] vector<T,N> WaveMultiPrefixBitAnd(vector<T,N> expr, uint4 mask); __generic<T : __BuiltinArithmeticType, let N : int, let M : int> __target_intrinsic(hlsl) __target_intrinsic(cuda, "_wavePrefixAndMultiple(_getMultiPrefixMask(($1).x), $0)") +[require(cuda_hlsl, waveprefix)] matrix<T,N,M> WaveMultiPrefixBitAnd(matrix<T,N,M> expr, uint4 mask); __generic<T : __BuiltinArithmeticType> __target_intrinsic(hlsl) __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) -//__target_intrinsic(glsl, "subgroupExclusiveOr($0)") +__target_intrinsic(glsl, "subgroupExclusiveOr($0)") __target_intrinsic(cuda, "_wavePrefixOr(, _getMultiPrefixMask(($1).x), $0)") +[require(cuda_glsl_hlsl, waveprefix)] T WaveMultiPrefixBitOr(T expr, uint4 mask); __generic<T : __BuiltinArithmeticType, let N : int> __target_intrinsic(hlsl) __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) -//__target_intrinsic(glsl, "subgroupExclusiveOr($0)") +__target_intrinsic(glsl, "subgroupExclusiveOr($0)") __target_intrinsic(cuda, "_wavePrefixOrMultiple(_getMultiPrefixMask(($1).x), $0)") +[require(cuda_glsl_hlsl, waveprefix)] vector<T,N> WaveMultiPrefixBitOr(vector<T,N> expr, uint4 mask); __generic<T : __BuiltinArithmeticType, let N : int, let M : int> __target_intrinsic(hlsl) __target_intrinsic(cuda, "_wavePrefixOrMultiple(_getMultiPrefixMask(($1).x), $0)") +[require(cuda_hlsl, waveprefix)] matrix<T,N,M> WaveMultiPrefixBitOr(matrix<T,N,M> expr, uint4 mask); __generic<T : __BuiltinArithmeticType> @@ -9843,6 +10480,7 @@ __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) __target_intrinsic(glsl, "subgroupExclusiveXor($0)") __target_intrinsic(cuda, "_wavePrefixXor(_getMultiPrefixMask(($1).x), $0)") +[require(cuda_glsl_hlsl, waveprefix)] T WaveMultiPrefixBitXor(T expr, uint4 mask); __generic<T : __BuiltinArithmeticType, let N : int> @@ -9851,41 +10489,49 @@ __glsl_extension(GL_KHR_shader_subgroup_arithmetic) __spirv_version(1.3) __target_intrinsic(glsl, "subgroupExclusiveXor($0)") __target_intrinsic(cuda, "_wavePrefixXorMultiple(_getMultiPrefixMask(($1).x), $0)") +[require(cuda_glsl_hlsl, waveprefix)] vector<T,N> WaveMultiPrefixBitXor(vector<T,N> expr, uint4 mask); __generic<T : __BuiltinArithmeticType, let N : int, let M : int> __target_intrinsic(hlsl) __target_intrinsic(cuda, "_wavePrefixXorMultiple(_getMultiPrefixMask(($1).x), $0)") +[require(cuda_hlsl, waveprefix)] matrix<T,N,M> WaveMultiPrefixBitXor(matrix<T,N,M> expr, uint4 mask); __generic<T : __BuiltinArithmeticType> __target_intrinsic(hlsl) __target_intrinsic(cuda, "_wavePrefixProduct(_getMultiPrefixMask(($1).x), $0)") +[require(cuda_hlsl, waveprefix)] T WaveMultiPrefixProduct(T value, uint4 mask); __generic<T : __BuiltinArithmeticType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(cuda, "_wavePrefixProductMultiple(_getMultiPrefixMask(($1).x), $0)") +[require(cuda_hlsl, waveprefix)] vector<T,N> WaveMultiPrefixProduct(vector<T,N> value, uint4 mask); __generic<T : __BuiltinArithmeticType, let N : int, let M : int> __target_intrinsic(hlsl) __target_intrinsic(cuda, "_wavePrefixProductMultiple(_getMultiPrefixMask(($1).x), $0)") +[require(cuda_hlsl, waveprefix)] matrix<T,N,M> WaveMultiPrefixProduct(matrix<T,N,M> value, uint4 mask); __generic<T : __BuiltinArithmeticType> __target_intrinsic(hlsl) __target_intrinsic(cuda, "_wavePrefixSum(_getMultiPrefixMask(($1).x), $0)") +[require(cuda_hlsl, waveprefix)] T WaveMultiPrefixSum(T value, uint4 mask); __generic<T : __BuiltinArithmeticType, let N : int> __target_intrinsic(hlsl) __target_intrinsic(cuda, "_wavePrefixSumMultiple(_getMultiPrefixMask(($1).x), $0 )") +[require(cuda_hlsl, waveprefix)] vector<T,N> WaveMultiPrefixSum(vector<T,N> value, uint4 mask); __generic<T : __BuiltinArithmeticType, let N : int, let M : int> __target_intrinsic(hlsl) __target_intrinsic(cuda, "_wavePrefixSumMultiple(_getMultiPrefixMask(($1).x), $0)") +[require(cuda_hlsl, waveprefix)] matrix<T,N,M> WaveMultiPrefixSum(matrix<T,N,M> value, uint4 mask); // `typedef`s to help with the fact that HLSL has been sorta-kinda case insensitive at various points @@ -9918,12 +10564,15 @@ for (int aa = 0; aa < kBaseBufferAccessLevelCount; ++aa) char const* glslTextureSizeFunc = (isReadOnly) ? "textureSize" : "imageSize"; char const* glslLoadFuncName = (isReadOnly) ? "texelFetch" : "imageLoad"; char const* spvLoadInstName = (isReadOnly) ? "OpImageFetch" : "OpImageRead"; + char const* requireToSetQuery = (isReadOnly) ? "[require(glsl_hlsl_spirv, texture_size)]" : "[require(glsl_hlsl_spirv, image_size)]"; + char const* requireToSet = (isReadOnly) ? "[require(glsl_hlsl_spirv, texture_sm_4_1)]" : "[require(glsl_hlsl_spirv, texture_sm_4_1_compute_fragment)]"; }}}} __generic<T, let format:int> extension __TextureImpl<T, __ShapeBuffer, 0, 0, 0, $(aa), 0, 0, format> { [__readNone] + $(requireToSetQuery) void GetDimensions(out uint dim) { __target_switch @@ -9940,6 +10589,7 @@ extension __TextureImpl<T, __ShapeBuffer, 0, 0, 0, $(aa), 0, 0, format> __glsl_extension(GL_EXT_samplerless_texture_functions) $(isReadOnly?"[__readNone] ":"") + $(requireToSet) T Load(int location) { __target_switch @@ -9954,17 +10604,21 @@ extension __TextureImpl<T, __ShapeBuffer, 0, 0, 0, $(aa), 0, 0, format> } $(isReadOnly?"[__readNone] ":"") + $(requireToSet) T Load(int location, out uint status); __subscript(uint index) -> T { $(isReadOnly?"[__readNone] ":"") [ForceInline] + $(requireToSet) get { return Load((int)index); } ${{{{ if (access != SLANG_RESOURCE_ACCESS_READ) { }}}} - [nonmutating] set + [nonmutating] + $(requireToSet) + set { __target_switch { @@ -10054,6 +10708,7 @@ struct RaytracingAccelerationStructure {}; // 10.1.5 - Intersection Attributes Structure __target_intrinsic(hlsl, BuiltInTriangleIntersectionAttributes) +[require(cpp_cuda_glsl_hlsl_spirv, rayobject)] struct BuiltInTriangleIntersectionAttributes { __target_intrinsic(hlsl, barycentrics) @@ -10072,8 +10727,8 @@ struct BuiltInTriangleIntersectionAttributes // `executeCallableNV` is the GLSL intrinsic that will be used to implement // `CallShader()` for GLSL-based targets. // -__target_intrinsic(_GL_EXT_ray_tracing, "executeCallableEXT") -[require(glsl, raytracing)] +__target_intrinsic(glsl, "executeCallableEXT") +[require(glsl, raytracing_raygen_closesthit_miss_callable)] void __executeCallable(uint shaderIndex, int payloadLocation); // Next is the custom intrinsic that will compute the payload location @@ -10089,9 +10744,7 @@ int __callablePayloadLocation(__ref Payload payload); // GLSL equivalent. // __generic<Payload> -[require(glsl, raytracing)] -[require(spirv, raytracing)] -[require(hlsl, raytracing)] +[require(glsl_hlsl_spirv, raytracing_raygen_closesthit_miss_callable)] void CallShader(uint shaderIndex, inout Payload payload) { __target_switch @@ -10132,8 +10785,8 @@ __intrinsic_op($(kIROp_ForceVarIntoStructTemporarily)) Ref<T> __forceVarIntoStructTemporarily(inout T maybeStruct); __target_intrinsic(hlsl, "TraceRay") -[require(hlsl, raytracing)] __generic<payload_t> +[require(hlsl, raytracing)] void __traceRayHLSL( RaytracingAccelerationStructure AccelerationStructure, uint RayFlags, @@ -10144,8 +10797,8 @@ void __traceRayHLSL( RayDesc Ray, inout payload_t Payload); -__target_intrinsic(_GL_EXT_ray_tracing, "traceRayEXT") -[require(glsl, raytracing)] +__target_intrinsic(glsl, "traceRayEXT") +[require(glsl, raytracing_raygen_closesthit_miss)] void __traceRay( RaytracingAccelerationStructure AccelerationStructure, uint RayFlags, @@ -10171,10 +10824,8 @@ __intrinsic_op($(kIROp_GetVulkanRayTracingPayloadLocation)) int __rayPayloadLocation(__ref Payload payload); [ForceInline] -[require(glsl, raytracing)] -[require(spirv, raytracing)] -[require(hlsl, raytracing)] __generic<payload_t> +[require(cuda_glsl_hlsl_spirv, raytracing_raygen_closesthit_miss)] void TraceRay( RaytracingAccelerationStructure AccelerationStructure, uint RayFlags, @@ -10256,8 +10907,8 @@ void TraceRay( // https://github.com/KhronosGroup/GLSL/blob/master/extensions/nv/GLSL_NV_ray_tracing_motion_blur.txt __target_intrinsic(hlsl, "TraceMotionRay") -[require(hlsl, raytracing_motionblur)] __generic<payload_t> +[require(hlsl, raytracing_motionblur)] void __traceMotionRayHLSL( RaytracingAccelerationStructure AccelerationStructure, uint RayFlags, @@ -10271,7 +10922,7 @@ void __traceMotionRayHLSL( __glsl_extension(GL_NV_ray_tracing_motion_blur) __target_intrinsic(glsl, "traceRayMotionNV") -[require(glsl, raytracing_motionblur)] +[require(glsl, raytracing_motionblur_raygen_closesthit_miss)] void __traceMotionRay( RaytracingAccelerationStructure AccelerationStructure, uint RayFlags, @@ -10287,9 +10938,7 @@ void __traceMotionRay( int PayloadLocation); [ForceInline] -[require(glsl, raytracing_motionblur)] -[require(spirv, raytracing_motionblur)] -[require(hlsl, raytracing_motionblur)] +[require(glsl_hlsl_spirv, raytracing_motionblur_raygen_closesthit_miss)] __generic<payload_t> void TraceMotionRay( RaytracingAccelerationStructure AccelerationStructure, @@ -10373,17 +11022,13 @@ void TraceMotionRay( } // 10.3.3 -__target_intrinsic(hlsl) -[require(hlsl, raytracing)] -bool ReportHit<A>(float tHit, uint hitKind, A attributes); -[require(spirv, raytracing)] -[require(glsl, raytracing)] +[require(glsl_spirv, raytracing_intersection)] bool __reportIntersection(float tHit, uint hitKind) { __target_switch { - case _GL_EXT_ray_tracing: __intrinsic_asm "reportIntersectionEXT"; + case glsl: __intrinsic_asm "reportIntersectionEXT"; case spirv: return spirv_asm { @@ -10393,27 +11038,31 @@ bool __reportIntersection(float tHit, uint hitKind) } __generic<A> -__specialized_for_target(glsl) -__specialized_for_target(spirv) +[ForceInline] +[require(glsl_hlsl_spirv, raytracing_intersection)] bool ReportHit(float tHit, uint hitKind, A attributes) { - [__vulkanHitAttributes] - static A a; - - a = attributes; - return __reportIntersection(tHit, hitKind); + __target_switch + { + case hlsl: + __intrinsic_asm "ReportHit($0, $1, $2)"; + case glsl: + case spirv: + [__vulkanHitAttributes] + static A a; + a = attributes; + return __reportIntersection(tHit, hitKind); + } } // 10.3.4 -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(cuda_glsl_hlsl_spirv, raytracing_anyhit)] void IgnoreHit() { __target_switch { case hlsl: __intrinsic_asm "IgnoreHit"; - case _GL_EXT_ray_tracing: __intrinsic_asm "ignoreIntersectionEXT;"; + case glsl: __intrinsic_asm "ignoreIntersectionEXT;"; case cuda: __intrinsic_asm "optixIgnoreIntersection"; case spirv: spirv_asm @@ -10424,15 +11073,13 @@ void IgnoreHit() } // 10.3.5 -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(cuda_glsl_hlsl_spirv, raytracing_anyhit)] void AcceptHitAndEndSearch() { __target_switch { case hlsl: __intrinsic_asm "AcceptHitAndEndSearch"; - case _GL_EXT_ray_tracing: __intrinsic_asm "terminateRayEXT;"; + case glsl: __intrinsic_asm "terminateRayEXT;"; case cuda: __intrinsic_asm "optixTerminateRay"; case spirv: spirv_asm @@ -10450,15 +11097,13 @@ void AcceptHitAndEndSearch() // 10.4.1 - Ray Dispatch System Values [NonUniformReturn] -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(cuda_glsl_hlsl_spirv, raytracing_allstages)] uint3 DispatchRaysIndex() { __target_switch { case hlsl: __intrinsic_asm "DispatchRaysIndex"; - case _GL_EXT_ray_tracing: __intrinsic_asm "(gl_LaunchIDEXT)"; + case glsl: __intrinsic_asm "(gl_LaunchIDEXT)"; case cuda: __intrinsic_asm "optixGetLaunchIndex"; case spirv: return spirv_asm @@ -10468,15 +11113,13 @@ uint3 DispatchRaysIndex() } } -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(cuda_glsl_hlsl_spirv, raytracing_allstages)] uint3 DispatchRaysDimensions() { __target_switch { case hlsl: __intrinsic_asm "DispatchRaysDimensions"; - case _GL_EXT_ray_tracing: __intrinsic_asm "(gl_LaunchSizeEXT)"; + case glsl: __intrinsic_asm "(gl_LaunchSizeEXT)"; case cuda: __intrinsic_asm "optixGetLaunchDimensions"; case spirv: return spirv_asm @@ -10489,15 +11132,13 @@ uint3 DispatchRaysDimensions() // 10.4.2 - Ray System Values [NonUniformReturn] -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(cuda_glsl_hlsl_spirv, raytracing_anyhit_closesthit_intersection_miss)] float3 WorldRayOrigin() { __target_switch { case hlsl: __intrinsic_asm "WorldRayOrigin"; - case _GL_EXT_ray_tracing: __intrinsic_asm "(gl_WorldRayOriginEXT)"; + case glsl: __intrinsic_asm "(gl_WorldRayOriginEXT)"; case cuda: __intrinsic_asm "optixGetWorldRayOrigin"; case spirv: return spirv_asm @@ -10508,15 +11149,13 @@ float3 WorldRayOrigin() } [NonUniformReturn] -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(cuda_glsl_hlsl_spirv, raytracing_anyhit_closesthit_intersection_miss)] float3 WorldRayDirection() { __target_switch { case hlsl: __intrinsic_asm "WorldRayDirection"; - case _GL_EXT_ray_tracing: __intrinsic_asm "(gl_WorldRayDirectionEXT)"; + case glsl: __intrinsic_asm "(gl_WorldRayDirectionEXT)"; case cuda: __intrinsic_asm "optixGetWorldRayDirection"; case spirv: return spirv_asm @@ -10527,15 +11166,13 @@ float3 WorldRayDirection() } [NonUniformReturn] -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(cuda_glsl_hlsl_spirv, raytracing_anyhit_closesthit_intersection_miss)] float RayTMin() { __target_switch { case hlsl: __intrinsic_asm "RayTMin"; - case _GL_EXT_ray_tracing: __intrinsic_asm "(gl_RayTminEXT)"; + case glsl: __intrinsic_asm "(gl_RayTminEXT)"; case cuda: __intrinsic_asm "optixGetRayTmin"; case spirv: return spirv_asm @@ -10556,15 +11193,13 @@ float RayTMin() // to the appropriate Vulkan stages. // [NonUniformReturn] -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(cuda_glsl_hlsl_spirv, raytracing_anyhit_closesthit_intersection_miss)] float RayTCurrent() { __target_switch { case hlsl: __intrinsic_asm "RayTCurrent"; - case _GL_EXT_ray_tracing: __intrinsic_asm "(gl_RayTmaxEXT)"; + case glsl: __intrinsic_asm "(gl_RayTmaxEXT)"; case cuda: __intrinsic_asm "optixGetRayTmax"; case spirv: return spirv_asm @@ -10574,15 +11209,13 @@ float RayTCurrent() } } -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(cuda_glsl_hlsl_spirv, raytracing_anyhit_closesthit_intersection_miss)] uint RayFlags() { __target_switch { case hlsl: __intrinsic_asm "RayFlags"; - case _GL_EXT_ray_tracing: __intrinsic_asm "(gl_IncomingRayFlagsEXT)"; + case glsl: __intrinsic_asm "(gl_IncomingRayFlagsEXT)"; case cuda: __intrinsic_asm "optixGetRayFlags"; case spirv: return spirv_asm @@ -10595,15 +11228,13 @@ uint RayFlags() // 10.4.3 - Primitive/Object Space System Values [NonUniformReturn] -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(cuda_glsl_hlsl_spirv, raytracing_anyhit_closesthit_intersection)] uint InstanceIndex() { __target_switch { case hlsl: __intrinsic_asm "InstanceIndex"; - case _GL_EXT_ray_tracing: __intrinsic_asm "(gl_InstanceID)"; + case glsl: __intrinsic_asm "(gl_InstanceID)"; case cuda: __intrinsic_asm "optixGetInstanceIndex"; case spirv: return spirv_asm @@ -10614,15 +11245,13 @@ uint InstanceIndex() } [NonUniformReturn] -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(cuda_glsl_hlsl_spirv, raytracing_anyhit_closesthit_intersection)] uint InstanceID() { __target_switch { case hlsl: __intrinsic_asm "InstanceID"; - case _GL_EXT_ray_tracing: __intrinsic_asm "(gl_InstanceCustomIndexEXT)"; + case glsl: __intrinsic_asm "(gl_InstanceCustomIndexEXT)"; case cuda: __intrinsic_asm "optixGetInstanceId"; case spirv: return spirv_asm @@ -10633,15 +11262,13 @@ uint InstanceID() } [NonUniformReturn] -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(cuda_glsl_hlsl_spirv, raytracing_anyhit_closesthit_intersection)] uint PrimitiveIndex() { __target_switch { case hlsl: __intrinsic_asm "PrimitiveIndex"; - case _GL_EXT_ray_tracing: __intrinsic_asm "(gl_PrimitiveID)"; + case glsl: __intrinsic_asm "(gl_PrimitiveID)"; case cuda: __intrinsic_asm "optixGetPrimitiveIndex"; case spirv: return spirv_asm @@ -10652,15 +11279,13 @@ uint PrimitiveIndex() } [NonUniformReturn] -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(cuda_glsl_hlsl_spirv, raytracing_anyhit_closesthit_intersection)] float3 ObjectRayOrigin() { __target_switch { case hlsl: __intrinsic_asm "ObjectRayOrigin"; - case _GL_EXT_ray_tracing: __intrinsic_asm "(gl_ObjectRayOriginEXT)"; + case glsl: __intrinsic_asm "(gl_ObjectRayOriginEXT)"; case cuda: __intrinsic_asm "optixGetObjectRayOrigin"; case spirv: return spirv_asm @@ -10671,15 +11296,13 @@ float3 ObjectRayOrigin() } [NonUniformReturn] -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(cuda_glsl_hlsl_spirv, raytracing_anyhit_closesthit_intersection)] float3 ObjectRayDirection() { __target_switch { case hlsl: __intrinsic_asm "ObjectRayDirection"; - case _GL_EXT_ray_tracing: __intrinsic_asm "(gl_ObjectRayDirectionEXT)"; + case glsl: __intrinsic_asm "(gl_ObjectRayDirectionEXT)"; case cuda: __intrinsic_asm "optixGetObjectRayDirection"; case spirv: return spirv_asm @@ -10692,15 +11315,13 @@ float3 ObjectRayDirection() // TODO: optix has an optixGetObjectToWorldTransformMatrix function that returns 12 // floats by reference. [NonUniformReturn] -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(glsl_hlsl_spirv, raytracing_anyhit_closesthit_intersection)] float3x4 ObjectToWorld3x4() { __target_switch { case hlsl: __intrinsic_asm "ObjectToWorld3x4"; - case _GL_EXT_ray_tracing: __intrinsic_asm "transpose(gl_ObjectToWorldEXT)"; + case glsl: __intrinsic_asm "transpose(gl_ObjectToWorldEXT)"; case spirv: return spirv_asm { @@ -10711,15 +11332,13 @@ float3x4 ObjectToWorld3x4() } [NonUniformReturn] -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(glsl_hlsl_spirv, raytracing_anyhit_closesthit_intersection)] float3x4 WorldToObject3x4() { __target_switch { case hlsl: __intrinsic_asm "WorldToObject3x4"; - case _GL_EXT_ray_tracing: __intrinsic_asm "transpose(gl_WorldToObjectEXT)"; + case glsl: __intrinsic_asm "transpose(gl_WorldToObjectEXT)"; case spirv: return spirv_asm { @@ -10730,15 +11349,13 @@ float3x4 WorldToObject3x4() } [NonUniformReturn] -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(glsl_hlsl_spirv, raytracing_anyhit_closesthit_intersection)] float4x3 ObjectToWorld4x3() { __target_switch { case hlsl: __intrinsic_asm "ObjectToWorld4x3"; - case _GL_EXT_ray_tracing: __intrinsic_asm "(gl_ObjectToWorldEXT)"; + case glsl: __intrinsic_asm "(gl_ObjectToWorldEXT)"; case spirv: return spirv_asm { @@ -10748,15 +11365,13 @@ float4x3 ObjectToWorld4x3() } [NonUniformReturn] -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(glsl_hlsl_spirv, raytracing_anyhit_closesthit_intersection)] float4x3 WorldToObject4x3() { __target_switch { case hlsl: __intrinsic_asm "WorldToObject4x3"; - case _GL_EXT_ray_tracing: __intrinsic_asm "(gl_WorldToObjectEXT)"; + case glsl: __intrinsic_asm "(gl_WorldToObjectEXT)"; case spirv: return spirv_asm { @@ -10772,9 +11387,7 @@ float4x3 WorldToObject4x3() __glsl_extension(GL_NV_ray_tracing_motion_blur) __glsl_extension(GL_EXT_ray_tracing) [NonUniformReturn] -[require(hlsl, raytracing_motionblur)] -[require(glsl, raytracing_motionblur)] -[require(spirv, raytracing_motionblur)] +[require(glsl_hlsl_spirv, raytracing_motionblur_anyhit_closesthit_intersection_miss)] float RayCurrentTime() { __target_switch @@ -10803,15 +11416,13 @@ float RayCurrentTime() // 10.4.4 - Hit Specific System values [NonUniformReturn] -[require(hlsl, raytracing)] -[require(glsl, raytracing)] -[require(spirv, raytracing)] +[require(cuda_glsl_hlsl_spirv, raytracing_anyhit_closesthit)] uint HitKind() { __target_switch { case hlsl: __intrinsic_asm "HitKind"; - case _GL_EXT_ray_tracing: __intrinsic_asm "(gl_HitKindEXT)"; + case glsl: __intrinsic_asm "(gl_HitKindEXT)"; case cuda: __intrinsic_asm "optixGetHitKind"; case spirv: return spirv_asm @@ -10851,6 +11462,7 @@ float dot2add(float2 left, float2 right, float acc); // Set the number of output vertices and primitives for a mesh shader invocation. __glsl_extension(GL_EXT_mesh_shader) __glsl_version(450) +[require(glsl_hlsl_spirv, meshshading)] void SetMeshOutputCounts(uint vertexCount, uint primitiveCount) { __target_switch @@ -10875,6 +11487,7 @@ void SetMeshOutputCounts(uint vertexCount, uint primitiveCount) // This function doesn't return. // [KnownBuiltin("DispatchMesh")] +[require(glsl_hlsl_spirv, meshshading)] void DispatchMesh<P>(uint threadGroupCountX, uint threadGroupCountY, uint threadGroupCountZ, __ref P meshPayload) { __target_switch @@ -11004,10 +11617,8 @@ extension __TextureImpl<T,__Shape2D, 1, 0, 0, $(kStdlibResourceAccessFeedback), // Get the index of the geometry that was hit in an intersection, any-hit, or closest-hit shader __glsl_extension(GL_EXT_ray_tracing) -[require(glsl, raytracing)] -[require(spirv, raytracing)] -[require(hlsl, raytracing)] [NonUniformReturn] +[require(glsl_hlsl_spirv, raytracing_anyhit_closesthit_intersection)] uint GeometryIndex() { __target_switch @@ -11024,9 +11635,8 @@ uint GeometryIndex() // https://github.com/KhronosGroup/GLSL/blob/master/extensions/ext/GLSL_EXT_ray_tracing_position_fetch.txt __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_EXT_ray_tracing_position_fetch) -[require(glsl, raytracing_pos)] -[require(spirv, raytracing_pos)] [ForceInline] +[require(glsl_spirv, raytracing_position)] float3 HitTriangleVertexPosition(uint index) { __target_switch @@ -11124,8 +11734,7 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> __target_intrinsic(glsl, "rayQueryInitializeEXT($0, $1, $2, $3, $4, $5, $6, $7)") __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] + [require(glsl_spirv, rayquery)] [mutating] void __rayQueryInitializeEXT( RaytracingAccelerationStructure accelerationStructure, @@ -11158,11 +11767,9 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> // `RayQuery` to get the effective ray flags, which // must obey any API-imposed restrictions. // - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [__unsafeForceInlineEarly] [mutating] + [require(glsl_hlsl_spirv, rayquery)] void TraceRayInline( RaytracingAccelerationStructure accelerationStructure, RAY_FLAG rayFlags, @@ -11200,10 +11807,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> // that was found. // __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [mutating] + [require(glsl_hlsl_spirv, rayquery)] bool Proceed() { __target_switch @@ -11224,10 +11829,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> // `Proceed()` calls will return `false`. // __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [mutating] + [require(glsl_hlsl_spirv, rayquery)] void Abort() { __target_switch @@ -11240,11 +11843,9 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> // Commit the current non-opaque triangle hit. __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [__NoSideEffect] [mutating] + [require(glsl_hlsl_spirv, rayquery)] void CommitNonOpaqueTriangleHit() { __target_switch @@ -11257,11 +11858,9 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> // Commit the current procedural primitive hit, with hit time `t`. __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [__NoSideEffect] [mutating] + [require(glsl_hlsl_spirv, rayquery)] void CommitProceduralPrimitiveHit(float t) { __target_switch @@ -11282,11 +11881,9 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> // user code. // __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [__NoSideEffect] [NonUniformReturn] + [require(glsl_hlsl_spirv, rayquery)] CANDIDATE_TYPE CandidateType() { __target_switch @@ -11303,11 +11900,9 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> // Get the status of the committed (closest) hit, if any. __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [__NoSideEffect] [NonUniformReturn] + [require(glsl_hlsl_spirv, rayquery)] COMMITTED_STATUS CommittedStatus() { __target_switch @@ -11324,11 +11919,9 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [__NoSideEffect] [NonUniformReturn] + [require(glsl_hlsl_spirv, rayquery)] bool CandidateProceduralPrimitiveNonOpaque() { __target_switch @@ -11346,11 +11939,9 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [__NoSideEffect] [NonUniformReturn] + [require(glsl_hlsl_spirv, rayquery)] float CandidateTriangleRayT() { __target_switch @@ -11366,11 +11957,9 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [__NoSideEffect] [NonUniformReturn] + [require(glsl_hlsl_spirv, rayquery)] float CommittedRayT() { __target_switch @@ -11388,9 +11977,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> ///missing hlsl equivlent; only implemented for glsl & spirv __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] int CandidateRayInstanceCustomIndex() { __target_switch @@ -11406,9 +11994,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] int CommittedRayInstanceCustomIndex() { __target_switch @@ -11424,9 +12011,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] int CandidateRayInstanceId() { __target_switch @@ -11442,9 +12028,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] int CommittedRayInstanceId() { __target_switch @@ -11460,9 +12045,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] uint CandidateRayInstanceShaderBindingTableRecordOffset() { __target_switch @@ -11478,9 +12062,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] uint CommittedRayInstanceShaderBindingTableRecordOffset() { __target_switch @@ -11496,9 +12079,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] int CandidateRayGeometryIndex() { __target_switch @@ -11514,9 +12096,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] int CommittedRayGeometryIndex() { __target_switch @@ -11532,9 +12113,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] int CandidateRayPrimitiveIndex() { __target_switch @@ -11550,9 +12130,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] int CommittedRayPrimitiveIndex() { __target_switch @@ -11568,9 +12147,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] float2 CandidateRayBarycentrics() { __target_switch @@ -11586,9 +12164,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] float2 CommittedRayBarycentrics() { __target_switch @@ -11604,9 +12181,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] bool CandidateRayFrontFace() { __target_switch @@ -11622,9 +12198,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] bool CommittedRayFrontFace() { __target_switch @@ -11640,9 +12215,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] float3 CandidateRayObjectRayDirection() { __target_switch @@ -11658,9 +12232,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] float3 CommittedRayObjectRayDirection() { __target_switch @@ -11676,9 +12249,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] float3 CandidateRayObjectRayOrigin() { __target_switch @@ -11694,9 +12266,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] float3 CommittedRayObjectRayOrigin() { __target_switch @@ -11712,9 +12283,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] float4x3 CandidateRayObjectToWorld() { __target_switch @@ -11730,9 +12300,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] float4x3 CommittedRayObjectToWorld() { __target_switch @@ -11748,9 +12317,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] float4x3 CandidateRayWorldToObject() { __target_switch @@ -11766,9 +12334,8 @@ struct RayQuery <let rayFlagsGeneric : RAY_FLAG = RAY_FLAG_NONE> } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] [__NoSideEffect] + [require(glsl_spirv, rayquery)] float4x3 CommittedRayWorldToObject() { __target_switch @@ -11796,7 +12363,7 @@ ${{{{ __glsl_extension(GL_EXT_ray_query) __glsl_extension(GL_EXT_ray_tracing_position_fetch) - [require(glsl, rayquery_pos)] + [require(glsl, rayquery_position)] [__NoSideEffect] void __glslGetIntersectionTriangleVertexPositions$(ccName)(out float3 arr[3]) { @@ -11804,8 +12371,8 @@ ${{{{ }; __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery_pos)] - [require(spirv, rayquery_pos)] + [require(glsl, rayquery_position)] + [require(spirv, rayquery_position)] [__NoSideEffect] float3[3] $(ccName)GetIntersectionTriangleVertexPositions() { @@ -11838,11 +12405,9 @@ ${{{{ }}}} __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [__NoSideEffect] [NonUniformReturn] + [require(glsl_hlsl_spirv, rayquery)] float3x4 $(ccName)$(matName)3x4() { __target_switch @@ -11860,11 +12425,9 @@ ${{{{ } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [__readNone] [NonUniformReturn] + [require(glsl_hlsl_spirv, rayquery)] float4x3 $(ccName)$(matName)4x3() { __target_switch @@ -11905,11 +12468,9 @@ ${{{{ }}}} __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [__NoSideEffect] [NonUniformReturn] + [require(glsl_hlsl_spirv, rayquery)] $(method.type) $(ccName)$(method.hlslName)() { __target_switch @@ -11932,11 +12493,9 @@ ${{{{ // Access properties of the ray being traced. __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [__NoSideEffect] [NonUniformReturn] + [require(glsl_hlsl_spirv, rayquery)] uint RayFlags() { __target_switch @@ -11952,11 +12511,9 @@ ${{{{ } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [__NoSideEffect] [NonUniformReturn] + [require(glsl_hlsl_spirv, rayquery)] float3 WorldRayOrigin() { __target_switch @@ -11972,11 +12529,9 @@ ${{{{ } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [__NoSideEffect] [NonUniformReturn] + [require(glsl_hlsl_spirv, rayquery)] float3 WorldRayDirection() { __target_switch @@ -11992,11 +12547,9 @@ ${{{{ } __glsl_extension(GL_EXT_ray_query) - [require(glsl, rayquery)] - [require(spirv, rayquery)] - [require(hlsl, rayquery)] [__NoSideEffect] [NonUniformReturn] + [require(glsl_hlsl_spirv, rayquery)] float RayTMin() { __target_switch @@ -12018,6 +12571,7 @@ ${{{{ __magic_type(SubpassInputType) __intrinsic_type($(kIROp_SubpassInputType)) +[require(glsl_hlsl_spirv, subpass)] struct __SubpassImpl<T, let isMS:int> { } @@ -12029,6 +12583,7 @@ __generic<T> extension __SubpassImpl<T, 0> { [ForceInline] + [require(hlsl_spirv, subpass)] T SubpassLoad() { __target_switch @@ -12054,6 +12609,7 @@ __generic<T> extension __SubpassImpl<T, 1> { [ForceInline] + [require(hlsl_spirv, subpass)] T SubpassLoad(int sample) { __target_switch @@ -12136,10 +12692,8 @@ struct HitObject /// Executes ray traversal (including anyhit and intersection shaders) like TraceRay, but returns the /// resulting hit information as a HitObject and does not trigger closesthit or miss shaders. - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] static HitObject TraceRay<payload_t>( RaytracingAccelerationStructure AccelerationStructure, uint RayFlags, @@ -12232,9 +12786,7 @@ struct HitObject /// Executes motion ray traversal (including anyhit and intersection shaders) like TraceRay, but returns the /// resulting hit information as a HitObject and does not trigger closesthit or miss shaders. [ForceInline] - [require(glsl, ser_motion)] - [require(spirv, ser_motion)] - [require(hlsl, ser_motion)] + [require(glsl_hlsl_spirv, ser_motion_raygen_closesthit_miss)] static HitObject TraceMotionRay<payload_t>( RaytracingAccelerationStructure AccelerationStructure, uint RayFlags, @@ -12332,10 +12884,8 @@ struct HitObject /// TraceRay. The computed index must reference a valid hit group record in the shader table. The /// Attributes parameter must either be an attribute struct, such as /// BuiltInTriangleIntersectionAttributes, or another HitObject to copy the attributes from. - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] static HitObject MakeHit<attr_t>( RaytracingAccelerationStructure AccelerationStructure, uint InstanceIndex, @@ -12419,10 +12969,8 @@ struct HitObject /// See MakeHit but handles Motion /// Currently only supported on VK - [require(glsl, ser_motion)] - [require(spirv, ser_motion)] - [require(hlsl, ser_motion)] - [ForceInline] + [ForceInline] + [require(glsl_hlsl_spirv, ser_motion_raygen_closesthit_miss)] static HitObject MakeMotionHit<attr_t>( RaytracingAccelerationStructure AccelerationStructure, uint InstanceIndex, @@ -12503,10 +13051,8 @@ struct HitObject /// reference a valid hit group record in the shader table. The Attributes parameter must either be an /// attribute struct, such as BuiltInTriangleIntersectionAttributes, or another HitObject to copy the /// attributes from. - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] static HitObject MakeHit<attr_t>( uint HitGroupRecordIndex, RaytracingAccelerationStructure AccelerationStructure, @@ -12583,9 +13129,8 @@ struct HitObject } /// See MakeHit but handles Motion /// Currently only supported on VK - [require(glsl, ser_motion)] - [require(spirv, ser_motion)] [ForceInline] + [require(glsl_spirv, ser_motion_raygen_closesthit_miss)] static HitObject MakeMotionHit<attr_t>( uint HitGroupRecordIndex, RaytracingAccelerationStructure AccelerationStructure, @@ -12655,10 +13200,8 @@ struct HitObject /// tracing a ray. The provided shader table index must reference a valid miss record in the shader /// table. [__requiresNVAPI] - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] static HitObject MakeMiss( uint MissShaderIndex, RayDesc Ray) @@ -12692,10 +13235,8 @@ struct HitObject /// See MakeMiss but handles Motion /// Currently only supported on VK - [require(glsl, ser_motion)] - [require(spirv, ser_motion)] - [require(hlsl, ser_motion)] [ForceInline] + [require(glsl_hlsl_spirv, ser_motion_raygen_closesthit_miss)] static HitObject MakeMotionMiss( uint MissShaderIndex, RayDesc Ray, @@ -12737,10 +13278,8 @@ struct HitObject /// scenarios where future control flow for some threads is known to process neither a hit nor a /// miss. [__requiresNVAPI] - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] static HitObject MakeNop() { __target_switch @@ -12771,10 +13310,8 @@ struct HitObject /// Invokes closesthit or miss shading for the specified hit object. In case of a NOP HitObject, no /// shader is invoked. [__requiresNVAPI] - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] static void Invoke<payload_t>( RaytracingAccelerationStructure AccelerationStructure, HitObject HitOrMiss, @@ -12825,10 +13362,8 @@ struct HitObject /// Returns true if the HitObject encodes a miss, otherwise returns false. [__requiresNVAPI] - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] bool IsMiss() { __target_switch @@ -12847,10 +13382,8 @@ struct HitObject /// Returns true if the HitObject encodes a hit, otherwise returns false. [__requiresNVAPI] - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] bool IsHit() { __target_switch @@ -12869,10 +13402,8 @@ struct HitObject /// Returns true if the HitObject encodes a nop, otherwise returns false. [__requiresNVAPI] - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] bool IsNop() { __target_switch @@ -12891,10 +13422,8 @@ struct HitObject /// Queries ray properties from HitObject. Valid if the hit object represents a hit or a miss. [__requiresNVAPI] - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] RayDesc GetRayDesc() { __target_switch @@ -12923,10 +13452,8 @@ struct HitObject /// Queries shader table index from HitObject. Valid if the hit object represents a hit or a miss. [__requiresNVAPI] __glsl_extension(GL_EXT_ray_tracing) - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] uint GetShaderTableIndex() { __target_switch @@ -12946,10 +13473,8 @@ struct HitObject /// Returns the instance index of a hit. Valid if the hit object represents a hit. [__requiresNVAPI] __glsl_extension(GL_EXT_ray_tracing) - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] uint GetInstanceIndex() { __target_switch @@ -12969,10 +13494,8 @@ struct HitObject /// Returns the instance ID of a hit. Valid if the hit object represents a hit. [__requiresNVAPI] __glsl_extension(GL_EXT_ray_tracing) - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] uint GetInstanceID() { __target_switch @@ -12992,10 +13515,8 @@ struct HitObject /// Returns the geometry index of a hit. Valid if the hit object represents a hit. [__requiresNVAPI] __glsl_extension(GL_EXT_ray_tracing) - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] uint GetGeometryIndex() { __target_switch @@ -13015,10 +13536,8 @@ struct HitObject /// Returns the primitive index of a hit. Valid if the hit object represents a hit. [__requiresNVAPI] __glsl_extension(GL_EXT_ray_tracing) - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] uint GetPrimitiveIndex() { __target_switch @@ -13038,10 +13557,8 @@ struct HitObject /// Returns the hit kind. Valid if the hit object represents a hit. [__requiresNVAPI] __glsl_extension(GL_EXT_ray_tracing) - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] uint GetHitKind() { __target_switch @@ -13060,10 +13577,8 @@ struct HitObject [__requiresNVAPI] __glsl_extension(GL_EXT_ray_tracing) - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] float4x3 GetWorldToObject() { __target_switch @@ -13082,10 +13597,8 @@ struct HitObject [__requiresNVAPI] __glsl_extension(GL_EXT_ray_tracing) - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_hlsl_spirv, ser_raygen_closesthit_miss)] float4x3 GetObjectToWorld() { __target_switch @@ -13102,9 +13615,8 @@ struct HitObject } } - [require(glsl, ser)] - [require(spirv, ser)] [ForceInline] + [require(glsl_spirv, ser_raygen_closesthit_miss)] float GetCurrentTime() { __target_switch { @@ -13120,9 +13632,8 @@ struct HitObject } } - [require(glsl, ser)] - [require(spirv, ser)] [ForceInline] + [require(glsl_spirv, ser_raygen_closesthit_miss)] float3 GetObjectRayOrigin() { __target_switch { @@ -13138,9 +13649,8 @@ struct HitObject } } - [require(glsl, ser)] - [require(spirv, ser)] [ForceInline] + [require(glsl_spirv, ser_raygen_closesthit_miss)] float3 GetObjectRayDirection() { __target_switch { @@ -13156,9 +13666,8 @@ struct HitObject } } - [require(glsl, ser)] - [require(spirv, ser)] [ForceInline] + [require(glsl_spirv, ser_raygen_closesthit_miss)] uint2 GetShaderRecordBufferHandle() { __target_switch { @@ -13175,10 +13684,8 @@ struct HitObject } /// Returns the attributes of a hit. Valid if the hit object represents a hit or a miss. - [require(glsl, ser)] - [require(spirv, ser)] - [require(hlsl, ser)] [ForceInline] + [require(glsl_spirv, ser_raygen_closesthit_miss)] attr_t GetAttributes<attr_t>() { __target_switch @@ -13226,12 +13733,12 @@ struct HitObject __target_intrinsic(hlsl, "NvGetAttributesFromHitObject($0, $1)") [__requiresNVAPI] - [require(hlsl, ser)] + [require(hlsl, ser_raygen_closesthit_miss)] void __hlslGetAttributesFromHitObject<T>(out T t); __target_intrinsic(hlsl, "NvMakeHitWithRecordIndex") [__requiresNVAPI] - [require(hlsl, ser)] + [require(hlsl, ser_raygen_closesthit_miss)] static void __hlslMakeHitWithRecordIndex<attr_t>( uint HitGroupRecordIndex, RaytracingAccelerationStructure AccelerationStructure, @@ -13245,7 +13752,7 @@ struct HitObject __target_intrinsic(hlsl, "NvMakeHit") [__requiresNVAPI] - [require(hlsl, ser)] + [require(hlsl, ser_raygen_closesthit_miss)] static void __hlslMakeHit<attr_t>(RaytracingAccelerationStructure AccelerationStructure, uint InstanceIndex, uint GeometryIndex, @@ -13259,7 +13766,7 @@ struct HitObject __target_intrinsic(hlsl, "NvTraceRayHitObject") [__requiresNVAPI] - [require(hlsl, ser)] + [require(hlsl, ser_raygen_closesthit_miss)] static void __hlslTraceRay<payload_t>( RaytracingAccelerationStructure AccelerationStructure, uint RayFlags, @@ -13277,8 +13784,8 @@ struct HitObject __glsl_extension(GL_NV_shader_invocation_reorder) __glsl_extension(GL_EXT_ray_tracing) - [require(glsl, ser)] __target_intrinsic(glsl, "hitObjectRecordMissNV") + [require(glsl, ser_raygen_closesthit_miss)] static void __glslMakeMiss( out HitObject hitObj, uint MissShaderIndex, @@ -13291,8 +13798,8 @@ struct HitObject __glsl_extension(GL_NV_shader_invocation_reorder) __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_ray_tracing_motion_blur) - [require(glsl, ser_motion)] __target_intrinsic(glsl, "hitObjectRecordMissMotionNV") + [require(glsl, ser_motion_raygen_closesthit_miss)] static void __glslMakeMotionMiss( out HitObject hitObj, uint MissShaderIndex, @@ -13304,45 +13811,45 @@ struct HitObject __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) - [require(glsl, ser)] __target_intrinsic(glsl, "hitObjectRecordEmptyNV") + [require(glsl, ser_raygen_closesthit_miss)] static void __glslMakeNop(out HitObject hitObj); __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) - [require(glsl, ser)] __target_intrinsic(glsl, "hitObjectGetObjectRayDirectionNV($0)") + [require(glsl, ser_raygen_closesthit_miss)] float3 __glslGetRayDirection(); __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) - [require(glsl, ser)] __target_intrinsic(glsl, "hitObjectGetWorldRayDirectionNV($0)") + [require(glsl, ser_raygen_closesthit_miss)] float3 __glslGetRayWorldDirection(); __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) - [require(glsl, ser)] __target_intrinsic(glsl, "hitObjectGetWorldRayOriginNV($0)") + [require(glsl, ser_raygen_closesthit_miss)] float3 __glslGetRayWorldOrigin(); __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) - [require(glsl, ser)] __target_intrinsic(glsl, "hitObjectGetRayTMaxNV($0)") + [require(glsl, ser_raygen_closesthit_miss)] float __glslGetTMax(); __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) - [require(glsl, ser)] __target_intrinsic(glsl, "hitObjectGetRayTMinNV($0)") + [require(glsl, ser_raygen_closesthit_miss)] float __glslGetTMin(); // "void hitObjectRecordHitWithIndexNV(hitObjectNV, accelerationStructureEXT,int,int,int,uint,uint,vec3,float,vec3,float,int);" __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) - [require(glsl, ser)] __target_intrinsic(glsl, "hitObjectRecordHitWithIndexNV") + [require(glsl, ser_raygen_closesthit_miss)] static void __glslMakeHitWithIndex( out HitObject hitObj, RaytracingAccelerationStructure accelerationStructure, @@ -13361,8 +13868,8 @@ struct HitObject __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) __glsl_extension(GL_NV_ray_tracing_motion_blur) - [require(glsl, ser_motion)] __target_intrinsic(glsl, "hitObjectRecordHitWithIndexMotionNV") + [require(glsl, ser_motion_raygen_closesthit_miss)] static void __glslMakeMotionHitWithIndex( out HitObject hitObj, RaytracingAccelerationStructure accelerationStructure, @@ -13381,8 +13888,8 @@ struct HitObject // "void hitObjectRecordHitNV(hitObjectNV,accelerationStructureEXT,int,int,int,uint,uint,uint,vec3,float,vec3,float,int);" __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) - [require(glsl, ser)] __target_intrinsic(glsl, "hitObjectRecordHitNV") + [require(glsl, ser_raygen_closesthit_miss)] static void __glslMakeHit( out HitObject hitObj, RaytracingAccelerationStructure accelerationStructure, @@ -13402,8 +13909,8 @@ struct HitObject __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) __glsl_extension(GL_NV_ray_tracing_motion_blur) - [require(glsl, ser_motion)] __target_intrinsic(glsl, "hitObjectRecordHitMotionNV") + [require(glsl, ser_motion_raygen_closesthit_miss)] static void __glslMakeMotionHit( out HitObject hitObj, RaytracingAccelerationStructure accelerationStructure, @@ -13423,14 +13930,14 @@ struct HitObject __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) - [require(glsl, ser)] __target_intrinsic(glsl, "hitObjectGetAttributesNV($0, $1)") + [require(glsl, ser_raygen_closesthit_miss)] void __glslGetAttributes(int attributeLocation); __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) - [require(glsl, ser)] __target_intrinsic(glsl, "hitObjectTraceRayNV") + [require(glsl, ser_raygen_closesthit_miss)] static void __glslTraceRay( out HitObject hitObject, RaytracingAccelerationStructure accelerationStructure, @@ -13448,8 +13955,8 @@ struct HitObject __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) __glsl_extension(GL_NV_ray_tracing_motion_blur) - [require(glsl, ser_motion)] __target_intrinsic(glsl, "hitObjectTraceRayMotionNV") + [require(glsl, ser_motion_raygen_closesthit_miss)] static void __glslTraceMotionRay( out HitObject hitObject, RaytracingAccelerationStructure accelerationStructure, @@ -13467,8 +13974,8 @@ struct HitObject __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) - [require(glsl, ser)] __target_intrinsic(glsl, "hitObjectExecuteShaderNV") + [require(glsl, ser_raygen_closesthit_miss)] static void __glslInvoke( HitObject hitObj, int payload); @@ -13484,9 +13991,8 @@ struct HitObject [__requiresNVAPI] __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) -[require(glsl, ser)] -[require(spirv, ser)] -[require(hlsl, ser)] +[ForceInline] +[require(glsl_hlsl_spirv, ser_raygen)] void ReorderThread( uint CoherenceHint, uint NumCoherenceHintBitsFromLSB ) { __target_switch @@ -13522,9 +14028,8 @@ void ReorderThread( uint CoherenceHint, uint NumCoherenceHintBitsFromLSB ) [__requiresNVAPI] __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) -[require(glsl, ser)] -[require(spirv, ser)] -[require(hlsl, ser)] +[ForceInline] +[require(glsl_hlsl_spirv, ser_raygen)] void ReorderThread( HitObject HitOrMiss, uint CoherenceHint, uint NumCoherenceHintBitsFromLSB ) { __target_switch @@ -13550,9 +14055,8 @@ void ReorderThread( HitObject HitOrMiss, uint CoherenceHint, uint NumCoherenceHi [__requiresNVAPI] __glsl_extension(GL_EXT_ray_tracing) __glsl_extension(GL_NV_shader_invocation_reorder) -[require(glsl, ser)] -[require(spirv, ser)] -[require(hlsl, ser)] +[ForceInline] +[require(glsl_hlsl_spirv, ser_raygen)] void ReorderThread( HitObject HitOrMiss ) { __target_switch @@ -13574,14 +14078,26 @@ void ReorderThread( HitObject HitOrMiss ) /// /// There doesn't appear to be an equivalent for debugBreak for HLSL -__target_intrinsic(hlsl, "/* debugBreak() not currently supported for HLSL */") -__target_intrinsic(cuda,"__brkpt()") -__target_intrinsic(cpp, "SLANG_BREAKPOINT(0)") -void debugBreak(); __specialized_for_target(glsl) [[vk::spirv_instruction(1, "NonSemantic.DebugBreak")]] -void debugBreak(); +void __glslDebugBreak(); + +[ForceInline] +[require(cpp_cuda_glsl_hlsl, breakpoint)] +void debugBreak() +{ + __target_switch + { + case hlsl: __intrinsic_asm "/* debugBreak() not currently supported for HLSL */"; + case cuda: __intrinsic_asm "__brkpt()"; + case cpp: __intrinsic_asm "SLANG_BREAKPOINT(0)"; + case glsl: + __glslDebugBreak(); + return; + } +} + // // Realtime Clock support @@ -13591,8 +14107,8 @@ void debugBreak(); [__requiresNVAPI] __glsl_extension(GL_EXT_shader_realtime_clock) -[require(shaderclock)] [NonUniformReturn] +[require(cpp_cuda_glsl_hlsl_spirv, shaderclock)] uint getRealtimeClockLow() { __target_switch @@ -13611,14 +14127,15 @@ uint getRealtimeClockLow() } __target_intrinsic(cpp, "std::chrono::high_resolution_clock::now().time_since_epoch().count()") - __target_intrinsic(cuda, "clock64") +__target_intrinsic(cuda, "clock64") [NonUniformReturn] +[require(cpp_cuda, shaderclock)] int64_t __cudaCppGetRealtimeClock(); [__requiresNVAPI] __glsl_extension(GL_EXT_shader_realtime_clock) -[require(shaderclock)] [NonUniformReturn] +[require(cpp_cuda_glsl_hlsl_spirv, shaderclock)] uint2 getRealtimeClock() { __target_switch @@ -13648,16 +14165,19 @@ uint2 getRealtimeClock() __target_intrinsic(cuda, "(threadIdx)") [__readNone] [NonUniformReturn] +[require(cuda)] uint3 cudaThreadIdx(); __target_intrinsic(cuda, "(blockIdx)") [__readNone] [NonUniformReturn] +[require(cuda)] uint3 cudaBlockIdx(); __target_intrinsic(cuda, "(blockDim)") [__readNone] [NonUniformReturn] +[require(cuda)] uint3 cudaBlockDim(); // @@ -13832,6 +14352,7 @@ ${{{{ [__NoSideEffect] [__requiresNVAPI] +[require(glsl_hlsl_spirv, texturefootprint)] vector<uint, ND> __textureFootprintGetAnchor<let ND:int>(__TextureFootprintData<ND> data, int nd) { __target_switch @@ -13849,6 +14370,7 @@ vector<uint, ND> __textureFootprintGetAnchor<let ND:int>(__TextureFootprintData< [__NoSideEffect] [__requiresNVAPI] +[require(glsl_hlsl_spirv, texturefootprint)] vector<uint, ND> __textureFootprintGetOffset<let ND:int>(__TextureFootprintData<ND> data, int nd) { __target_switch @@ -13865,6 +14387,7 @@ vector<uint, ND> __textureFootprintGetOffset<let ND:int>(__TextureFootprintData< } __intrinsic_type($(kIROp_TextureFootprintType)) +[require(glsl_hlsl_spirv, texturefootprint)] struct __TextureFootprintData<let ND:int> { typealias Anchor = vector<uint, ND>; @@ -14011,6 +14534,7 @@ ${ [__NoSideEffect] __glsl_version(450) __glsl_extension(GL_NV_shader_texture_footprint) + [require(glsl_spirv, texturefootprint)] bool __queryFootprintGLSL( SamplerState sampler, Coords coords, @@ -14037,6 +14561,7 @@ ${ [__NoSideEffect] __glsl_version(450) __glsl_extension(GL_NV_shader_texture_footprint) + [require(glsl_spirv, texturefootprint)] bool __queryFootprintGLSL( SamplerState sampler, Coords coords, @@ -14065,8 +14590,7 @@ ${ __glsl_version(450) __glsl_extension(GL_NV_shader_texture_footprint) __glsl_extension(GL_ARB_sparse_texture_clamp) - __target_intrinsic(glsl, - "textureFootprintClampNV($p, $*2)") + [require(glsl_spirv, texturefootprintclamp)] bool __queryFootprintClampGLSL( SamplerState sampler, Coords coords, @@ -14096,6 +14620,7 @@ ${ __glsl_version(450) __glsl_extension(GL_NV_shader_texture_footprint) __glsl_extension(GL_ARB_sparse_texture_clamp) + [require(glsl_spirv, texturefootprintclamp)] bool __queryFootprintClampGLSL( SamplerState sampler, Coords coords, @@ -14126,6 +14651,7 @@ ${ __glsl_version(450) __glsl_extension(GL_NV_shader_texture_footprint) [__requiresNVAPI] + [require(glsl_spirv, texturefootprint)] bool __queryFootprintLodGLSL( SamplerState sampler, Coords coords, @@ -14158,6 +14684,7 @@ ${{{ __glsl_version(450) __glsl_extension(GL_NV_shader_texture_footprint) [__requiresNVAPI] + [require(glsl_spirv, texturefootprint)] bool __queryFootprintGradGLSL( SamplerState sampler, Coords coords, @@ -14187,6 +14714,7 @@ ${{{ __glsl_version(450) __glsl_extension(GL_NV_shader_texture_footprint) __glsl_extension(GL_ARB_sparse_texture_clamp) + [require(glsl_spirv, texturefootprintclamp)] bool __queryFootprintGradClampGLSL( SamplerState sampler, Coords coords, @@ -14256,6 +14784,7 @@ for(auto levelChoice : kLevelChoices) [__requiresNVAPI] __target_intrinsic(hlsl, "NvFootprint$(CoarseOrFine)($1, $2, $3, $4, NV_EXTN_TEXTURE_$!0D, $*5)") + [require(hlsl, texturefootprint)] static __FootprintData __queryFootprint$(CoarseOrFine)NVAPI( int nd, uint textureSpace, @@ -14270,6 +14799,7 @@ for(auto levelChoice : kLevelChoices) [__requiresNVAPI] __target_intrinsic(hlsl, "NvFootprint$(CoarseOrFine)Bias($1, $2, $3, $4, NV_EXTN_TEXTURE_$!0D, $*5)") + [require(hlsl, texturefootprint)] static __FootprintData __queryFootprint$(CoarseOrFine)BiasNVAPI( int nd, uint textureSpace, @@ -14285,6 +14815,7 @@ for(auto levelChoice : kLevelChoices) [__requiresNVAPI] __target_intrinsic(hlsl, "NvFootprint$(CoarseOrFine)Level($1, $2, $3, $4, NV_EXTN_TEXTURE_$!0D, $*5)") + [require(hlsl, texturefootprint)] static __FootprintData __queryFootprint$(CoarseOrFine)LevelNVAPI( int nd, uint textureSpace, @@ -14300,6 +14831,7 @@ for(auto levelChoice : kLevelChoices) [__requiresNVAPI] __target_intrinsic(hlsl, "NvFootprint$(CoarseOrFine)Grad($1, $2, $3, $4, NV_EXTN_TEXTURE_$!0D, $*5)") + [require(hlsl, texturefootprint)] static __FootprintData __queryFootprint$(CoarseOrFine)GradNVAPI( int nd, uint textureSpace, @@ -14564,6 +15096,7 @@ extension __TextureImpl<float, Shape, 0, 0, 0, $(kStdlibResourceAccessReadWrite) [__requiresNVAPI] [ForceInline] __glsl_extension(GL_EXT_shader_atomic_float) + [require(glsl_hlsl_spirv, atomic_glsl_hlsl_cuda_float1)] void InterlockedAddF32(vector<uint, Shape.dimensions> coord, float value, out float originalValue) { __target_switch @@ -14628,6 +15161,7 @@ __generic<T, let Alignment : int = 16> __intrinsic_type($(kIROp_HLSLConstBufferPointerType)) __glsl_extension(GL_EXT_buffer_reference) __magic_type(ConstBufferPointerType) +[require(glsl_spirv, bufferreference)] struct ConstBufferPointer { __glsl_version(450) @@ -14657,6 +15191,7 @@ struct ConstBufferPointer __glsl_version(450) __glsl_extension(GL_EXT_shader_explicit_arithmetic_types_int64) __glsl_extension(GL_EXT_buffer_reference) + [require(glsl_spirv, bufferreference_int64)] static ConstBufferPointer<T> fromUInt(uint64_t val) { __target_switch @@ -14673,6 +15208,7 @@ struct ConstBufferPointer __glsl_version(450) __glsl_extension(GL_EXT_shader_explicit_arithmetic_types_int64) __glsl_extension(GL_EXT_buffer_reference) + [require(glsl_spirv, bufferreference_int64)] uint64_t toUInt() { __target_switch @@ -14691,6 +15227,7 @@ struct ConstBufferPointer __glsl_extension(GL_EXT_buffer_reference) [__NoSideEffect] [ForceInline] + [require(glsl_spirv, bufferreference_int64)] bool isValid() { __target_switch |
