summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
Diffstat
-rw-r--r--source/slang/hlsl.meta.slang8
-rw-r--r--tests/bugs/gh-5339.slang27
-rw-r--r--tests/glsl-intrinsic/intrinsic-texture.slang26
3 files changed, 45 insertions, 16 deletions
diff --git a/source/slang/hlsl.meta.slang b/source/slang/hlsl.meta.slang
index 2df6c5492..dad634934 100644
--- a/source/slang/hlsl.meta.slang
+++ b/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;
};
}
}
diff --git a/tests/bugs/gh-5339.slang b/tests/bugs/gh-5339.slang
new file mode 100644
index 000000000..163b05bc7
--- /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 c40390c57..c1e73e73d 100644
--- a/tests/glsl-intrinsic/intrinsic-texture.slang
+++ b/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;
@@ -704,8 +704,8 @@ bool textureFuncs( Sampler1D<vector<T,N>> 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<vector<T,N>> 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<vector<T,N>> 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<vector<T,N>> 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<vector<T,N>> 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<vector<T,N>> 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<vector<T,N>> 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
generated by cgit v1.2.3 (git 2.39.1) at 2026-06-01 08:00:20 +0000