diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang index fb73496c9f..315024a9bf 100644 --- a/source/slang/hlsl.meta.slang +++ b/source/slang/hlsl.meta.slang @@ -2536,7 +2536,7 @@ __generic __texture_gather_offset( _Texture texture, SamplerState s, - constexpr vector location, + vector location, constexpr vector offset, int component) { @@ -2557,8 +2557,9 @@ vector __texture_gather_offset( __intrinsic_asm "$0.gather($1, $2, $3, metal::component($4))"; case spirv: return spirv_asm { + OpCapability ImageGatherExtended; %sampledImage : __sampledImageType(texture) = OpSampledImage $texture $s; - result:$$vector = OpImageGather %sampledImage $location $component ConstOffset $offset; + result:$$vector = OpImageGather %sampledImage $location $component Offset $offset; }; case wgsl: if (isShadow == 1) @@ -2606,7 +2607,8 @@ vector __texture_gather_offset( __intrinsic_asm "textureGatherOffset($0, $1, $2, $3)"; case spirv: return spirv_asm { - result:$$vector = OpImageGather $sampler $location $component ConstOffset $offset; + OpCapability ImageGatherExtended; + result:$$vector = OpImageGather $sampler $location $component Offset $offset; }; } } diff --git a/tests/bugs/gh-5339.slang b/tests/bugs/gh-5339.slang new file mode 100644 index 0000000000..59b9e5e574 --- /dev/null +++ b/tests/bugs/gh-5339.slang @@ -0,0 +1,26 @@ +//TEST:SIMPLE(filecheck=SPV): -allow-glsl -target spirv-asm -entry computeMain -stage compute + +// Test if we are correctly using `Offset` option instead of `ConstOffset` +// when the offset value is not a compile-time constant. + +//SPV:OpCapability ImageGatherExtended + +#extension GL_EXT_gpu_shader5 : require + +layout (location = 0) in highp vec2 v_texCoord; + +layout (binding = 0) uniform highp sampler2D u_sampler; +layout (binding = 1) uniform offset { highp ivec2 u_offset; }; + +//TEST_INPUT:ubuffer(data=[0], stride=4):out,name=outputBuffer +buffer MyBlockName +{ + vec4 result; +} outputBuffer; + +void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID) +{ + //SPV:OpImageGather % + //SPV-SAME: Offset % + outputBuffer.result = textureGatherOffset(u_sampler, v_texCoord, u_offset); +} diff --git a/tests/glsl-intrinsic/intrinsic-texture.slang b/tests/glsl-intrinsic/intrinsic-texture.slang index c40390c57c..0d6bd151ac 100644 --- a/tests/glsl-intrinsic/intrinsic-texture.slang +++ b/tests/glsl-intrinsic/intrinsic-texture.slang @@ -167,9 +167,9 @@ bool textureFuncs( Sampler1D> gsampler1D constexpr float coord = 0.5; - constexpr ivec2 offset2D = ivec2(0); - constexpr ivec3 offset3D = ivec3(0); - constexpr ivec2 offsets[4] = { ivec2(0), ivec2(0), ivec2(0), ivec2(0) }; + constexpr ivec2 offset2D = ivec2(2); + constexpr ivec3 offset3D = ivec3(3); + constexpr ivec2 offsets[4] = { ivec2(1), ivec2(2), ivec2(3), ivec2(4) }; bool ignoreResultF32 = ignoreResult && T is float; bool ignoreResultI32 = ignoreResult && T is int32_t; @@ -1321,35 +1321,35 @@ bool textureFuncs( Sampler1D> gsampler1D // GLSL-COUNT-2: textureGatherOffset({{.*}}sampler2D // SPIR-COUNT-2: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2D - // SPIR: OpImageGather {{.*}}[[LOAD]]{{.*}} ConstOffset % + // SPIR: OpImageGather {{.*}}[[LOAD]]{{.*}}Offset % && gvec4(T(0)) == textureGatherOffset(gsampler2D, vec2(coord), offset2D) && gvec4(T(0)) == textureGatherOffset(gsampler2D, vec2(coord), offset2D, int(0)) // GLSL-COUNT-2: textureGatherOffset({{.*}}sampler2DArray // SPIR-COUNT-2: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2DArray - // SPIR: OpImageGather {{.*}}[[LOAD]]{{.*}} ConstOffset % + // SPIR: OpImageGather {{.*}}[[LOAD]]{{.*}}Offset % && gvec4(T(0)) == textureGatherOffset(gsampler2DArray, vec3(coord), offset2D) && gvec4(T(0)) == textureGatherOffset(gsampler2DArray, vec3(coord), offset2D, int(0)) // GLSL: textureGatherOffset({{.*}}sampler2DShadow // SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2DShadow - // SPIR: OpImageDrefGather {{.*}}[[LOAD]]{{.*}} ConstOffset % + // SPIR: OpImageDrefGather {{.*}}[[LOAD]]{{.*}}Offset % && (vec4(0) == textureGatherOffset(uniform_sampler2DShadow, vec2(coord), float(0), offset2D) || ignoreResultF32) // GLSL: textureGatherOffset({{.*}}sampler2DArrayShadow // SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2DArrayShadow - // SPIR: OpImageDrefGather {{.*}}[[LOAD]]{{.*}} ConstOffset % + // SPIR: OpImageDrefGather {{.*}}[[LOAD]]{{.*}}Offset % && vec4(0) == textureGatherOffset(uniform_sampler2DArrayShadow, vec3(coord), float(0), offset2D) // GLSL-COUNT-2: textureGatherOffset({{.*}}sampler2DRect // SPIR-COUNT-2: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2DRect - // SPIR: OpImageGather {{.*}}[[LOAD]]{{.*}} ConstOffset % + // SPIR: OpImageGather {{.*}}[[LOAD]]{{.*}}Offset % && gvec4(T(0)) == textureGatherOffset(gsampler2DRect, vec2(coord), offset2D) && gvec4(T(0)) == textureGatherOffset(gsampler2DRect, vec2(coord), offset2D, int(0)) // GLSL: textureGatherOffset({{.*}}sampler2DRectShadow // SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2DRectShadow - // SPIR: OpImageDrefGather {{.*}}[[LOAD]]{{.*}} ConstOffset % + // SPIR: OpImageDrefGather {{.*}}[[LOAD]]{{.*}}Offset % && vec4(0) == textureGatherOffset(uniform_sampler2DRectShadow, vec2(coord), float(0), offset2D) // GLSL-COUNT-2: textureGatherOffsets({{.*}}sampler2D