From ef3552d9cbde0c06ddc77150b8ab68674d8422d3 Mon Sep 17 00:00:00 2001 From: Jay Kwak <82421531+jkwak-work@users.noreply.github.com> Date: Mon, 23 Sep 2024 19:46:32 -0700 Subject: [PATCH] Feature/wgsl intrinsic texture gather (#5141) This PR implements the texture gather functions for WGSL. The pattern was very similar to how Metal was implemented. Before copy and paste from the Metal implementation, I had to clean up the Metal implementation to make it more readable and maintainable. Gather functions are available only for 2D and 3D textures. Their `array` and `depth` variants may or may not be supported depending on the target. `static_assert` ensures that Gather functions are available only for 2D and 3D textures. Removed incorrect use of "$p" argument for targeting GLSL. --- source/slang/hlsl.meta.slang | 475 +++++++++++++++---------- source/slang/slang-capabilities.capdef | 4 + tests/wgsl/texture-sampler-less.slang | 8 +- tests/wgsl/texture.slang | 8 +- 4 files changed, 303 insertions(+), 192 deletions(-) diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang index 8d1b3202e5..3ea5481b19 100644 --- a/source/slang/hlsl.meta.slang +++ b/source/slang/hlsl.meta.slang @@ -2188,60 +2188,68 @@ extension __TextureImpl __makeArray(T v0, T v1, T v2, T v3); -// Gather for scalar textures. + +// Beginning of Texture Gather __generic [ForceInline] -[require(glsl_metal_spirv, texture_gather)] -vector __texture_gather(__TextureImpl texture, SamplerState s, vector location, int component) +[require(glsl_metal_spirv_wgsl, texture_gather)] +vector __texture_gather( + __TextureImpl texture, + SamplerState s, + vector location, + int component) { __target_switch { case glsl: - __intrinsic_asm "textureGather($p, $2, $3)"; + __intrinsic_asm "textureGather($0, $1, $2, $3)"; case metal: - if (isShadow == 0) + if (isArray == 1) { switch (Shape.flavor) { case $(SLANG_TEXTURE_2D): - if (isArray == 1) - { - // Tv gather(sampler s, float2 coord, uint array, int2 offset = int2(0), component c = component::x) const - __intrinsic_asm "$0.gather($1, ($2).xy, uint(($2).z), int2(0), metal::component($3))"; - } - else - { - // Tv gather(sampler s, float2 coord, int2 offset = int2(0), component c = component::x) const - __intrinsic_asm "$0.gather($1, $2, int2(0), metal::component($3))"; - } - break; + // Tv gather(sampler s, float2 coord, uint array, int2 offset = int2(0), component c = component::x) const + __intrinsic_asm "$0.gather($1, ($2).xy, uint(($2).z), int2(0), metal::component($3))"; case $(SLANG_TEXTURE_CUBE): - if (isArray == 1) - { - // Tv gather(sampler s, float3 coord, uint array, component c = component::x) const - __intrinsic_asm "$0.gather($1, ($2).xyz, uint(($2).w), metal::component($3))"; - } - else - { - // Tv gather(sampler s, float3 coord, component c = component::x) const - __intrinsic_asm "$0.gather($1, $2, metal::component($3))"; - } - break; + // Tv gather(sampler s, float3 coord, uint array, component c = component::x) const + __intrinsic_asm "$0.gather($1, ($2).xyz, uint(($2).w), metal::component($3))"; } } - // TODO: This needs to be handled by the capability system - __intrinsic_asm ""; + if (Shape.flavor == $(SLANG_TEXTURE_CUBE)) + { + // Tv gather(sampler s, float3 coord, component c = component::x) const + __intrinsic_asm "$0.gather($1, $2, metal::component($3))"; + } + // Tv gather(sampler s, float2 coord, int2 offset = int2(0), component c = component::x) const + __intrinsic_asm "$0.gather($1, $2, int2(0), metal::component($3))"; case spirv: return spirv_asm { %sampledImage : __sampledImageType(texture) = OpSampledImage $texture $s; result:$$vector = OpImageGather %sampledImage $location $component; }; + case wgsl: + if (isArray == 1) + { + switch (Shape.flavor) + { + case $(SLANG_TEXTURE_2D): + __intrinsic_asm "textureGather($3, $0, $1, ($2).xy, u32(($2).z))"; + case $(SLANG_TEXTURE_CUBE): + __intrinsic_asm "textureGather($3, $0, $1, ($2).xyz, u32(($2).w))"; + } + } + __intrinsic_asm "textureGather($3, $0, $1, $2)"; } } + __generic [ForceInline] [require(glsl_spirv, texture_gather)] -vector __texture_gather(__TextureImpl sampler, vector location, int component) +vector __texture_gather( + __TextureImpl sampler, + vector location, + int component) { __target_switch { @@ -2253,45 +2261,60 @@ vector __texture_gather(__TextureImpl [ForceInline] -[require(glsl_metal_spirv, texture_gather)] -vector __texture_gather_offset(__TextureImpl texture, SamplerState s, constexpr vector location, constexpr vector offset, int component) +[require(glsl_metal_spirv_wgsl, texture_gather)] +vector __texture_gather_offset( + __TextureImpl texture, + SamplerState s, + constexpr vector location, + constexpr vector offset, + int component) { __target_switch { case glsl: - __intrinsic_asm "textureGatherOffset($p, $2, $3, $4)"; + __intrinsic_asm "textureGatherOffset($0, $1, $2, $3, $4)"; case metal: - if (Shape.flavor == $(SLANG_TEXTURE_2D)) + static_assert(Shape.flavor == $(SLANG_TEXTURE_2D), + "Metal supports offset variant of Gather only for 2D textures"); + + if (isArray == 1) { - if (isShadow == 0) - { - if (isArray == 1) - { - // Tv gather(sampler s, float2 coord, uint array, int2 offset = int2(0), component c = component::x) const - __intrinsic_asm "$0.gather($1, ($2).xy, uint(($2).z), $3, metal::component($4))"; - } - else - { - // Tv gather(sampler s, float2 coord, int2 offset = int2(0), component c = component::x) const - __intrinsic_asm "$0.gather($1, $2, $3, metal::component($4))"; - } - } + // Tv gather(sampler s, float2 coord, uint array, int2 offset = int2(0), component c = component::x) const + __intrinsic_asm "$0.gather($1, ($2).xy, uint(($2).z), $3, metal::component($4))"; } - // TODO: This needs to be handled by the capability system - __intrinsic_asm ""; + // Tv gather(sampler s, float2 coord, int2 offset = int2(0), component c = component::x) const + __intrinsic_asm "$0.gather($1, $2, $3, metal::component($4))"; case spirv: return spirv_asm { %sampledImage : __sampledImageType(texture) = OpSampledImage $texture $s; result:$$vector = OpImageGather %sampledImage $location $component ConstOffset $offset; }; + case wgsl: + if (isArray == 1) + { + switch (Shape.flavor) + { + case $(SLANG_TEXTURE_2D): + __intrinsic_asm "textureGather($4, $0, $1, ($2).xy, u32(($2).z), $3)"; + case $(SLANG_TEXTURE_CUBE): + __intrinsic_asm "textureGather($4, $0, $1, ($2).xyz, u32(($2).w), $3)"; + } + } + __intrinsic_asm "textureGather($4, $0, $1, $2, $3)"; } } + __generic [ForceInline] [require(glsl_spirv, texture_gather)] -vector __texture_gather_offset(__TextureImpl sampler, vector location, constexpr vector offset, int component) +vector __texture_gather_offset( + __TextureImpl sampler, + vector location, + constexpr vector offset, + int component) { __target_switch { @@ -2303,10 +2326,14 @@ vector __texture_gather_offset(__TextureImpl [ForceInline] [require(glsl_spirv, texture_gather)] -vector __texture_gather_offsets(__TextureImpl texture, SamplerState s, vector location, +vector __texture_gather_offsets( + __TextureImpl texture, + SamplerState s, + vector location, constexpr vector offset1, constexpr vector offset2, constexpr vector offset3, @@ -2316,7 +2343,7 @@ vector __texture_gather_offsets(__TextureImpl __texture_gather_offsets(__TextureImpl [ForceInline] [require(glsl_spirv, texture_gather)] -vector __texture_gather_offsets(__TextureImpl sampler, vector location, - +vector __texture_gather_offsets( + __TextureImpl sampler, + vector location, constexpr vector offset1, constexpr vector offset2, constexpr vector offset3, @@ -2349,58 +2378,62 @@ vector __texture_gather_offsets(__TextureImpl [ForceInline] -[require(glsl_metal_spirv, texture_gather)] -vector __texture_gatherCmp(__TextureImpl texture, SamplerComparisonState s, vector location, TElement compareValue) +[require(glsl_metal_spirv_wgsl, texture_gather)] +vector __texture_gatherCmp( + __TextureImpl texture, + SamplerComparisonState s, + vector location, + TElement compareValue) { __target_switch { case glsl: - __intrinsic_asm "textureGather($p, $2, $3)"; + __intrinsic_asm "textureGather($0, $1, $2, $3)"; case metal: - if (isShadow == 1) + if (isArray == 1) { - if (Shape.flavor == $(SLANG_TEXTURE_2D)) - { - if (isArray == 1) - { - // Tv gather_compare(sampler s, float2 coord, uint array, float compare_value, int2 offset = int2(0)) const - __intrinsic_asm "$0.gather_compare($1, ($2).xy, uint(($2).z), $3)"; - } - else - { - // Tv gather_compare(sampler s, float2 coord, float compare_value, int2 offset = int2(0)) const - __intrinsic_asm "$0.gather_compare($1, $2, $3)"; - } - } - else if (Shape.flavor == $(SLANG_TEXTURE_CUBE)) + switch (Shape.flavor) { - if (isArray == 1) - { - // Tv gather_compare(sampler s, float3 coord, uint array, float compare_value) const - __intrinsic_asm "$0.gather_compare($1, ($2).xyz, uint(($2).w), $3)"; - } - else - { - // Tv gather_compare(sampler s, float3 coord, float compare_value) const - __intrinsic_asm "$0.gather_compare($1, $2, $3)"; - } + case $(SLANG_TEXTURE_2D): + // Tv gather_compare(sampler s, float2 coord, uint array, float compare_value, int2 offset = int2(0)) const + __intrinsic_asm "$0.gather_compare($1, ($2).xy, uint(($2).z), $3)"; + case $(SLANG_TEXTURE_CUBE): + // Tv gather_compare(sampler s, float3 coord, uint array, float compare_value) const + __intrinsic_asm "$0.gather_compare($1, ($2).xyz, uint(($2).w), $3)"; } } - // TODO: This needs to be handled by the capability system - __intrinsic_asm ""; + // Tv gather_compare(sampler s, float2 coord, float compare_value, int2 offset = int2(0)) const + __intrinsic_asm "$0.gather_compare($1, $2, $3)"; case spirv: return spirv_asm { %sampledImage : __sampledImageType(texture) = OpSampledImage $texture $s; result:$$vector = OpImageDrefGather %sampledImage $location $compareValue; }; + case wgsl: + if (isArray == 1) + { + switch (Shape.flavor) + { + case $(SLANG_TEXTURE_2D): + __intrinsic_asm "textureGatherCompare($0, $1, ($2).xy, u32(($2).z), $3)"; + case $(SLANG_TEXTURE_CUBE): + __intrinsic_asm "textureGatherCompare($0, $1, ($2).xyz, u32(($2).w), $3)"; + } + } + __intrinsic_asm "textureGatherCompare($0, $1, $2, $3)"; } } + __generic [ForceInline] [require(glsl_spirv, texture_gather)] -vector __texture_gatherCmp(__TextureImpl sampler, vector location, TElement compareValue) +vector __texture_gatherCmp( + __TextureImpl sampler, + vector location, + TElement compareValue) { __target_switch { @@ -2412,45 +2445,60 @@ vector __texture_gatherCmp(__TextureImpl [ForceInline] -[require(glsl_metal_spirv, texture_gather)] -vector __texture_gatherCmp_offset(__TextureImpl texture, SamplerComparisonState s, vector location, TElement compareValue, constexpr vector offset) +[require(glsl_metal_spirv_wgsl, texture_gather)] +vector __texture_gatherCmp_offset( + __TextureImpl texture, + SamplerComparisonState s, + vector location, + TElement compareValue, + constexpr vector offset) { __target_switch { case glsl: - __intrinsic_asm "textureGatherOffset($p, $2, $3, $4)"; + __intrinsic_asm "textureGatherOffset($0, $1, $2, $3, $4)"; case metal: - if (isShadow == 1) + static_assert(Shape.flavor == $(SLANG_TEXTURE_2D), + "Metal supports depth compare Gather only for 2D texture"); + + if (isArray == 1) { - if (Shape.flavor == $(SLANG_TEXTURE_2D)) - { - if (isArray == 1) - { - // Tv gather_compare(sampler s, float2 coord, uint array, float compare_value, int2 offset = int2(0)) const - __intrinsic_asm "$0.gather_compare($1, ($2).xy, uint(($2).z), $3, $4)"; - } - else - { - // Tv gather_compare(sampler s, float2 coord, float compare_value, int2 offset = int2(0)) const - __intrinsic_asm "$0.gather_compare($1, $2, $3, $4)"; - } - } + // Tv gather_compare(sampler s, float2 coord, uint array, float compare_value, int2 offset = int2(0)) const + __intrinsic_asm "$0.gather_compare($1, ($2).xy, uint(($2).z), $3, $4)"; } - // TODO: This needs to be handled by the capability system - __intrinsic_asm ""; + // Tv gather_compare(sampler s, float2 coord, float compare_value, int2 offset = int2(0)) const + __intrinsic_asm "$0.gather_compare($1, $2, $3, $4)"; case spirv: return spirv_asm { %sampledImage : __sampledImageType(texture) = OpSampledImage $texture $s; result:$$vector = OpImageDrefGather %sampledImage $location $compareValue ConstOffset $offset; }; + case wgsl: + if (isArray == 1) + { + switch (Shape.flavor) + { + case $(SLANG_TEXTURE_2D): + __intrinsic_asm "textureGatherCompare($0, $1, ($2).xy, u32(($2).z), $3, $4)"; + case $(SLANG_TEXTURE_CUBE): + __intrinsic_asm "textureGatherCompare($0, $1, ($2).xyz, u32(($2).w), $3, $4)"; + } + } + __intrinsic_asm "textureGatherCompare($0, $1, $2, $3, $4)"; } } + __generic [ForceInline] [require(glsl_spirv, texture_gather)] -vector __texture_gatherCmp_offset(__TextureImpl sampler, vector location, TElement compareValue, constexpr vector offset) +vector __texture_gatherCmp_offset( + __TextureImpl sampler, + vector location, + TElement compareValue, + constexpr vector offset) { __target_switch { @@ -2462,10 +2510,14 @@ vector __texture_gatherCmp_offset(__TextureImpl [ForceInline] [require(glsl_spirv, texture_gather)] -vector __texture_gatherCmp_offsets(__TextureImpl texture, SamplerComparisonState s, vector location, TElement compareValue, +vector __texture_gatherCmp_offsets( + __TextureImpl texture, + SamplerComparisonState s, vector location, + TElement compareValue, vector offset1, vector offset2, vector offset3, @@ -2474,7 +2526,7 @@ vector __texture_gatherCmp_offsets(__TextureImpl __texture_gatherCmp_offsets(__TextureImpl [ForceInline] [require(glsl_spirv, texture_gather)] -vector __texture_gatherCmp_offsets(__TextureImpl sampler, vector location, TElement compareValue, +vector __texture_gatherCmp_offsets( + __TextureImpl sampler, + vector location, + TElement compareValue, vector offset1, vector offset2, vector offset3, @@ -2510,124 +2566,167 @@ ${{{{ for (int isCombined = 0; isCombined < 2; isCombined++) for (int isScalarTexture = 0; isScalarTexture < 2; isScalarTexture++) { - if (isScalarTexture == 0) - { - sb << "__generic\n"; - sb << "extension __TextureImpl\n"; - } - else - { - sb << "__generic\n"; - sb << "extension __TextureImpl,Shape,isArray,0,sampleCount,0,isShadow," << isCombined << ",format>\n"; - } + const char* extSizeParam = isScalarTexture ? "" : ", let N:int"; + const char* extTexType = isScalarTexture ? "T" : "vector"; + }}}} -{ // begin extension for gather +// Gather for [TextureType = $(extTexType), isCombined = $(isCombined)] +__generic +extension __TextureImpl<$(extTexType),Shape,isArray,0,sampleCount,0,isShadow,$(isCombined),format> +{ ${{{{ - - // Gather component - const char* samplerStateParam = isCombined ? "" : " s,"; - const char* getTexture = isCombined ? "__getTexture()" : "this"; - for (int isCmp = 0; isCmp < 2; ++isCmp) - { - const char* cmp = isCmp ? "Cmp" : ""; - const char* cmpParam = isCmp ? ", T compareValue" : ""; - const char* compareArg = isCmp ? ", compareValue" : ""; - const char* samplerStateType = isCombined ? "" : (isCmp ? "SamplerComparisonState" : "SamplerState"); - const char* getSampler = isCombined ? (isCmp ? " __getComparisonSampler()," : " __getSampler(),") : samplerStateParam; - const char* componentNames[] = { "", "Red", "Green", "Blue", "Alpha"}; - const char* glslComponentNames[] = { ", 0", ", 1", ", 2", ", 3" }; - - for (auto componentId = 0; componentId < 5; componentId++) - { - auto componentName = componentNames[componentId]; - auto glslComponent = (isCmp ? "" :glslComponentNames[componentId == 0 ? 0 : componentId - 1]); - - for (bool isStatus : { false, true }) - { - const char* statusDecl = isStatus ? ", out uint status" : ""; - const char* statusInit = isStatus ? " status = 0;\n" : ""; - const char* statusCapWithMetal = isStatus ? "hlsl" : "glsl_hlsl_metal_spirv"; - const char* statusCapWithoutMetal = isStatus ? "hlsl" : "glsl_hlsl_spirv"; + for (int isShadow = 0; isShadow < 2; isShadow++) + for (auto componentId = 0; componentId < 5; componentId++) + { + const char* compareFunc = isShadow ? "Cmp" : ""; + const char* compareParam = isShadow ? ", T compareValue" : ""; + const char* compareArg = isShadow ? ", compareValue" : ""; + + // Some targets support the combined texture natively + const char* samplerParam = isCombined ? "" : (isShadow ? "SamplerComparisonState s," : "SamplerState s,"); + const char* samplerArg = isCombined ? "" : ", s"; + const char* getTexture = isCombined ? "__getTexture()" : "this"; + const char* getSampler = isCombined ? (isShadow ? ", __getComparisonSampler()" : ", __getSampler()") : samplerArg; + + const char* componentFuncString[] = { "", "Red", "Green", "Blue", "Alpha"}; + const char* componentArgString[] = { ", 0", ", 0", ", 1", ", 2", ", 3" }; + const char* componentFunc = componentFuncString[componentId]; + const char* componentArg = (isShadow ? "" : componentArgString[componentId]); }}}} [ForceInline] - [require($(statusCapWithMetal), texture_gather)] - vector Gather$(cmp)$(componentName)($(samplerStateType)$(samplerStateParam) vector location $(cmpParam) $(statusDecl)) + [require(glsl_hlsl_metal_spirv_wgsl, texture_gather)] + vector Gather$(compareFunc)$(componentFunc)( + $(samplerParam) + vector location + $(compareParam)) { - $(statusInit) + static_assert(Shape.flavor == $(SLANG_TEXTURE_2D) || Shape.flavor == $(SLANG_TEXTURE_CUBE), + "Gather is supported only for 2D and 3D textures"); + __target_switch { - case hlsl: __intrinsic_asm ".Gather$(cmp)$(componentName)"; -${{{{ - if (!isStatus) - { -}}}} + case hlsl: __intrinsic_asm ".Gather$(compareFunc)$(componentFunc)"; case metal: - return __texture_gather$(cmp)($(getTexture),$(getSampler) location $(compareArg) $(glslComponent)); + case wgsl: + return __texture_gather$(compareFunc)($(getTexture) $(getSampler), location $(compareArg) $(componentArg)); case glsl: case spirv: - return __texture_gather$(cmp)(this,$(samplerStateParam) location $(compareArg) $(glslComponent)); -${{{{ - } // if(!isStatus) -}}}} + return __texture_gather$(compareFunc)(this $(samplerArg), location $(compareArg) $(componentArg)); } } + [ForceInline] - [require($(statusCapWithMetal), texture_gather)] - vector Gather$(cmp)$(componentName)($(samplerStateType)$(samplerStateParam) vector location $(cmpParam), constexpr vector offset $(statusDecl)) + [require(hlsl, texture_gather)] + vector Gather$(compareFunc)$(componentFunc)( + $(samplerParam) + vector location + $(compareParam), + out uint status) { - $(statusInit) + static_assert(Shape.flavor == $(SLANG_TEXTURE_2D) || Shape.flavor == $(SLANG_TEXTURE_CUBE), + "Gather is supported only for 2D and 3D textures"); + __target_switch { - case hlsl: __intrinsic_asm ".Gather$(cmp)$(componentName)"; -${{{{ - if (!isStatus) - { -}}}} + case hlsl: __intrinsic_asm ".Gather$(compareFunc)$(componentFunc)"; + } + } + + [ForceInline] + [require(glsl_hlsl_metal_spirv_wgsl, texture_gather)] + vector Gather$(compareFunc)$(componentFunc)( + $(samplerParam) + vector location + $(compareParam), + constexpr vector offset) + { + static_assert(Shape.flavor == $(SLANG_TEXTURE_2D) || Shape.flavor == $(SLANG_TEXTURE_CUBE), + "Gather is supported only for 2D and 3D textures"); + + __target_switch + { + case hlsl: __intrinsic_asm ".Gather$(compareFunc)$(componentFunc)"; case metal: - return __texture_gather$(cmp)_offset($(getTexture),$(getSampler) location $(compareArg), offset $(glslComponent)); + case wgsl: + return __texture_gather$(compareFunc)_offset($(getTexture) $(getSampler), location $(compareArg), offset $(componentArg)); case glsl: case spirv: - return __texture_gather$(cmp)_offset(this,$(samplerStateParam) location $(compareArg), offset $(glslComponent)); -${{{{ - } // if(!isStatus) -}}}} + return __texture_gather$(compareFunc)_offset(this $(samplerArg), location $(compareArg), offset $(componentArg)); } } + [ForceInline] - [require($(statusCapWithoutMetal), texture_gather)] - vector Gather$(cmp)$(componentName)($(samplerStateType)$(samplerStateParam) vector location $(cmpParam), + [require(hlsl, texture_gather)] + vector Gather$(compareFunc)$(componentFunc)( + $(samplerParam) + vector location + $(compareParam), + constexpr vector offset, + out uint status) + { + static_assert(Shape.flavor == $(SLANG_TEXTURE_2D) || Shape.flavor == $(SLANG_TEXTURE_CUBE), + "Gather is supported only for 2D and 3D textures"); + + __target_switch + { + case hlsl: __intrinsic_asm ".Gather$(compareFunc)$(componentFunc)"; + } + } + + [ForceInline] + [require(glsl_hlsl_spirv, texture_gather)] + vector Gather$(compareFunc)$(componentFunc)( + $(samplerParam) + vector location + $(compareParam), constexpr vector offset1, constexpr vector offset2, constexpr vector offset3, - constexpr vector offset4 - $(statusDecl)) + constexpr vector offset4) { - $(statusInit) + static_assert(Shape.flavor == $(SLANG_TEXTURE_2D) || Shape.flavor == $(SLANG_TEXTURE_CUBE), + "Gather is supported only for 2D and 3D textures"); + __target_switch { - case hlsl: __intrinsic_asm ".Gather$(cmp)$(componentName)"; -${{{{ - if (!isStatus) - { -}}}} + case hlsl: __intrinsic_asm ".Gather$(compareFunc)$(componentFunc)"; case glsl: case spirv: - return __texture_gather$(cmp)_offsets(this,$(samplerStateParam) location $(compareArg), offset1,offset2,offset3,offset4 $(glslComponent)); -${{{{ - } // if(!isStatus) -}}}} + return __texture_gather$(compareFunc)_offsets(this $(samplerArg), location $(compareArg), offset1,offset2,offset3,offset4 $(componentArg)); + } + } + + [ForceInline] + [require(hlsl, texture_gather)] + vector Gather$(compareFunc)$(componentFunc)( + $(samplerParam) + vector location + $(compareParam), + constexpr vector offset1, + constexpr vector offset2, + constexpr vector offset3, + constexpr vector offset4, + out uint status) + { + static_assert(Shape.flavor == $(SLANG_TEXTURE_2D) || Shape.flavor == $(SLANG_TEXTURE_CUBE), + "Gather is supported only for 2D and 3D textures"); + + __target_switch + { + case hlsl: __intrinsic_asm ".Gather$(compareFunc)$(componentFunc)"; } } + ${{{{ - } // for (isStatus) - } // for (componentId) - } // for (isCmp) + } // for (componentId) }}}} -} // end extension for gather +} // End of: Gather for [TextureType = $(extTexType), isCombined = $(isCombined)] ${{{{ } // for (isScalarTexture) }}}} +// End of all Texture Gather + // Load/Subscript for readonly, no MS textures diff --git a/source/slang/slang-capabilities.capdef b/source/slang/slang-capabilities.capdef index 102671ffbc..801d54eca8 100644 --- a/source/slang/slang-capabilities.capdef +++ b/source/slang/slang-capabilities.capdef @@ -374,6 +374,10 @@ alias glsl_hlsl_metal_spirv_wgsl = glsl | hlsl | metal | spirv | wgsl; /// [Compound] alias glsl_metal_spirv = glsl | metal | spirv; +/// GLSL, Metal, SPIRV and WGSL code-gen targets +/// [Compound] +alias glsl_metal_spirv_wgsl = glsl | metal | spirv | wgsl; + /// GLSL, and SPIRV code-gen targets /// [Compound] alias glsl_spirv = glsl | spirv; diff --git a/tests/wgsl/texture-sampler-less.slang b/tests/wgsl/texture-sampler-less.slang index 1a6cb53412..893e867b7e 100644 --- a/tests/wgsl/texture-sampler-less.slang +++ b/tests/wgsl/texture-sampler-less.slang @@ -311,23 +311,27 @@ bool TEST_texture( // https://www.w3.org/TR/WGSL/#texturegather // ================================== -#if 0 + // WGSL: textureGather({{.*}}t2D && all(Tv4(T(0)) == t2D.Gather(float2(u, u))) + // WGSL: textureGather({{.*}}tCube && all(Tv4(T(0)) == tCube.Gather(normalize(float3(u, 1 - u, u)))) + // WGSL: textureGather({{.*}}t2DArray && all(Tv4(T(0)) == t2DArray.Gather(float3(u, u, 0))) + // WGSL: textureGather({{.*}}tCubeArray && all(Tv4(T(0)) == tCubeArray.Gather(float4(normalize(float3(u, 1 - u, u)), 0))) #if TEST_WHEN_CONSTEXPR_WORKS_FOR_OFFSET // Offset variant + // W-GSL: textureGather({{.*}}t2D && all(Tv4(T(0)) == t2D.Gather(float2(u2, u), int2(0, 0))) + // W-GSL: textureGather({{.*}}t2DArray && all(Tv4(T(0)) == t2DArray.Gather(float3(u2, u, 0), int2(0, 0))) #endif // #if TEST_WHEN_CONSTEXPR_WORKS_FOR_OFFSET -#endif // ===================================== // T SampleGrad() diff --git a/tests/wgsl/texture.slang b/tests/wgsl/texture.slang index 999555a55c..af39cf52ab 100644 --- a/tests/wgsl/texture.slang +++ b/tests/wgsl/texture.slang @@ -346,21 +346,25 @@ bool TEST_texture( // https://www.w3.org/TR/WGSL/#texturegather // ================================== -#if 0 + // WGSL: textureGather({{.*}}t2D && all(Tv4(T(0)) == t2D.Gather(samplerState, float2(u, u))) + // WGSL: textureGather({{.*}}tCube && all(Tv4(T(0)) == tCube.Gather(samplerState, normalize(float3(u, 1 - u, u)))) + // WGSL: textureGather({{.*}}t2DArray && all(Tv4(T(0)) == t2DArray.Gather(samplerState, float3(u, u, 0))) + // WGSL: textureGather({{.*}}tCubeArray && all(Tv4(T(0)) == tCubeArray.Gather(samplerState, float4(normalize(float3(u, 1 - u, u)), 0))) // Offset variant + // WGSL: textureGather({{.*}}t2D && all(Tv4(T(0)) == t2D.Gather(samplerState, float2(u2, u), int2(0, 0))) + // WGSL: textureGather({{.*}}t2DArray && all(Tv4(T(0)) == t2DArray.Gather(samplerState, float3(u2, u, 0), int2(0, 0))) -#endif // ===================================== // T SampleGrad()