diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang index 2df6c5492b..dad6349343 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..163b05bc7d --- /dev/null +++ b/tests/bugs/gh-5339.slang @@ -0,0 +1,27 @@ +//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-NOT:Const + //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..c1e73e73d4 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; @@ -704,8 +704,8 @@ bool textureFuncs( Sampler1D> gsampler1D // GLSL: textureOffset({{.*}}sampler1D // SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler1D - // S-PIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset % - && gvec4(T(0)) == textureOffset(gsampler1D, float(coord), int(0)) + // SPIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset % + && gvec4(T(0)) == textureOffset(gsampler1D, float(coord), int(1)) // GLSL: textureOffset({{.*}}sampler1D // SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler1D @@ -714,7 +714,7 @@ bool textureFuncs( Sampler1D> gsampler1D // GLSL: textureOffset({{.*}}sampler2D // SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2D - // S-PIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset % + // SPIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset % && gvec4(T(0)) == textureOffset(gsampler2D, vec2(coord), offset2D) // GLSL: textureOffset({{.*}}sampler2D @@ -724,7 +724,7 @@ bool textureFuncs( Sampler1D> gsampler1D // GLSL: textureOffset({{.*}}sampler3D // SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler3D - // S-PIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset % + // SPIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset % && gvec4(T(0)) == textureOffset(gsampler3D, vec3(coord), offset3D) // GLSL: textureOffset({{.*}}sampler3D @@ -734,7 +734,7 @@ bool textureFuncs( Sampler1D> gsampler1D // GLSL: textureOffset({{.*}}sampler2DShadow // SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2DShadow - // S-PIR: OpImageSampleDrefImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset % + // SPIR: OpImageSampleDrefImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset % && (float(0) == textureOffset(uniform_sampler2DShadow, vec3(coord), offset2D) || ignoreResultF32) // GLSL: textureOffset({{.*}}sampler2DShadow @@ -744,12 +744,12 @@ bool textureFuncs( Sampler1D> gsampler1D // GLSL: textureOffset({{.*}}sampler2DRect // SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2DRect - // S-PIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset % + // SPIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset % && gvec4(T(0)) == textureOffset(gsampler2DRect, vec2(coord), offset2D) // GLSL: textureOffset({{.*}}sampler2DRectShadow // SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2DRectShadow - // S-PIR: OpImageSampleDrefImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset % + // SPIR: OpImageSampleDrefImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset % && float(0) == textureOffset(uniform_sampler2DRectShadow, vec3(coord), offset2D) // GLSL: textureOffset({{.*}}sampler1DShadow @@ -764,8 +764,8 @@ bool textureFuncs( Sampler1D> gsampler1D // GLSL: textureOffset({{.*}}sampler1DArray // SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler1DArray - // S-PIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset % - && gvec4(T(0)) == textureOffset(gsampler1DArray, vec2(coord), int(0)) + // SPIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset % + && gvec4(T(0)) == textureOffset(gsampler1DArray, vec2(coord), int(1)) // GLSL: textureOffset({{.*}}sampler1DArray // SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler1DArray @@ -774,7 +774,7 @@ bool textureFuncs( Sampler1D> gsampler1D // GLSL: textureOffset({{.*}}sampler2DArray // SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2DArray - // S-PIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset % + // SPIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset % && gvec4(T(0)) == textureOffset(gsampler2DArray, vec3(coord), offset2D) // GLSL: textureOffset({{.*}}sampler2DArray