Use Offset instead of ConstOffset for GLSL function textureGatherOffset

This commit changes to use `Offset` option on OpImageGather instruction
when translating `textureGatherOffset` to SPIR-V code.

Interestingly GLSL allows the offset value to be a variable not
constant just for the function `textureGatherOffset` while all other
offset variants of texture sampling functions require the offset to be a
constant value.

From a few experiments, I found that the spirv-optimizer changes
`Offset` to `ConstOffset` if the offset is a compile-time value. I also
found that when the offset value is zero, it changes to `None` with no
offset value.

source/slang/hlsl.meta.slang

@@ -2536,7 +2536,7 @@ __generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:
 vector<TElement,4> __texture_gather_offset(
     _Texture<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture,
     SamplerState s,
-    constexpr vector<float, Shape.dimensions+isArray> location,
+    vector<float, Shape.dimensions+isArray> location,
     constexpr vector<int, Shape.planeDimensions> offset,
     int component)
 {
@@ -2557,8 +2557,9 @@ vector<TElement,4> __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<TElement,4> = OpImageGather %sampledImage $location $component ConstOffset $offset;
+            result:$$vector<TElement,4> = OpImageGather %sampledImage $location $component Offset $offset;
         };
     case wgsl:
         if (isShadow == 1)
@@ -2606,7 +2607,8 @@ vector<TElement,4> __texture_gather_offset(
         __intrinsic_asm "textureGatherOffset($0, $1, $2, $3)";
     case spirv:
         return spirv_asm {
-            result:$$vector<TElement,4> = OpImageGather $sampler $location $component ConstOffset $offset;
+            OpCapability ImageGatherExtended;
+            result:$$vector<TElement,4> = OpImageGather $sampler $location $component Offset $offset;
         };
     }
 }

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);
+}

tests/glsl-intrinsic/intrinsic-texture.slang

@@ -167,9 +167,9 @@ bool textureFuncs( Sampler1D<vector<T,N>> 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<vector<T,N>> 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