Skip to content

Commit

Permalink
Use Offset instead of ConstOffset for GLSL function textureGatherOffs…
Browse files Browse the repository at this point in the history
…et (#5426)

* 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.
  • Loading branch information
jkwak-work authored Oct 29, 2024
1 parent f65d756 commit 8b3f904
Show file tree
Hide file tree
Showing 3 changed files with 45 additions and 16 deletions.
8 changes: 5 additions & 3 deletions source/slang/hlsl.meta.slang
Original file line number Diff line number Diff line change
Expand Up @@ -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)
{
Expand All @@ -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)
Expand Down Expand Up @@ -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;
};
}
}
Expand Down
27 changes: 27 additions & 0 deletions tests/bugs/gh-5339.slang
Original file line number Diff line number Diff line change
@@ -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);
}
26 changes: 13 additions & 13 deletions tests/glsl-intrinsic/intrinsic-texture.slang
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down Expand Up @@ -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
Expand All @@ -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
Expand All @@ -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
Expand All @@ -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
Expand All @@ -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
Expand All @@ -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
Expand All @@ -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
Expand Down

0 comments on commit 8b3f904

Please sign in to comment.