Skip to content

Commit 621b8bb

Browse files
csyonghecsyonghe
authored
Merge branch 'master' into cheneym2/slang_issue_4124
2 parents ece22d0 + 8b3f904 commit 621b8bb

File tree

3 files changed

+45
-16
lines changed

3 files changed

+45
-16
lines changed

‎source/slang/hlsl.meta.slang

+5-3
Original file line numberDiff line numberDiff line change
@@ -2536,7 +2536,7 @@ __generic<TElement, T, Shape: __ITextureShape, let isArray:int, let sampleCount:
25362536
vector<TElement,4> __texture_gather_offset(
25372537
_Texture<T, Shape, isArray, 0, sampleCount, access, isShadow, 0, format> texture,
25382538
SamplerState s,
2539-
constexpr vector<float, Shape.dimensions+isArray> location,
2539+
vector<float, Shape.dimensions+isArray> location,
25402540
constexpr vector<int, Shape.planeDimensions> offset,
25412541
int component)
25422542
{
@@ -2557,8 +2557,9 @@ vector<TElement,4> __texture_gather_offset(
25572557
__intrinsic_asm "$0.gather($1, $2, $3, metal::component($4))";
25582558
case spirv:
25592559
return spirv_asm {
2560+
OpCapability ImageGatherExtended;
25602561
%sampledImage : __sampledImageType(texture) = OpSampledImage $texture $s;
2561-
result:$$vector<TElement,4> = OpImageGather %sampledImage $location $component ConstOffset $offset;
2562+
result:$$vector<TElement,4> = OpImageGather %sampledImage $location $component Offset $offset;
25622563
};
25632564
case wgsl:
25642565
if (isShadow == 1)
@@ -2606,7 +2607,8 @@ vector<TElement,4> __texture_gather_offset(
26062607
__intrinsic_asm "textureGatherOffset($0, $1, $2, $3)";
26072608
case spirv:
26082609
return spirv_asm {
2609-
result:$$vector<TElement,4> = OpImageGather $sampler $location $component ConstOffset $offset;
2610+
OpCapability ImageGatherExtended;
2611+
result:$$vector<TElement,4> = OpImageGather $sampler $location $component Offset $offset;
26102612
};
26112613
}
26122614
}

‎tests/bugs/gh-5339.slang

+27
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,27 @@
1+
//TEST:SIMPLE(filecheck=SPV): -allow-glsl -target spirv-asm -entry computeMain -stage compute
2+
3+
// Test if we are correctly using `Offset` option instead of `ConstOffset`
4+
// when the offset value is not a compile-time constant.
5+
6+
//SPV:OpCapability ImageGatherExtended
7+
8+
#extension GL_EXT_gpu_shader5 : require
9+
10+
layout (location = 0) in highp vec2 v_texCoord;
11+
12+
layout (binding = 0) uniform highp sampler2D u_sampler;
13+
layout (binding = 1) uniform offset { highp ivec2 u_offset; };
14+
15+
//TEST_INPUT:ubuffer(data=[0], stride=4):out,name=outputBuffer
16+
buffer MyBlockName
17+
{
18+
vec4 result;
19+
} outputBuffer;
20+
21+
void computeMain(uint3 dispatchThreadID : SV_DispatchThreadID)
22+
{
23+
//SPV:OpImageGather %
24+
//SPV-NOT:Const
25+
//SPV-SAME: Offset %
26+
outputBuffer.result = textureGatherOffset(u_sampler, v_texCoord, u_offset);
27+
}

‎tests/glsl-intrinsic/intrinsic-texture.slang

+13-13
Original file line numberDiff line numberDiff line change
@@ -167,9 +167,9 @@ bool textureFuncs( Sampler1D<vector<T,N>> gsampler1D
167167

168168
constexpr float coord = 0.5;
169169

170-
constexpr ivec2 offset2D = ivec2(0);
171-
constexpr ivec3 offset3D = ivec3(0);
172-
constexpr ivec2 offsets[4] = { ivec2(0), ivec2(0), ivec2(0), ivec2(0) };
170+
constexpr ivec2 offset2D = ivec2(2);
171+
constexpr ivec3 offset3D = ivec3(3);
172+
constexpr ivec2 offsets[4] = { ivec2(1), ivec2(2), ivec2(3), ivec2(4) };
173173

174174
bool ignoreResultF32 = ignoreResult && T is float;
175175
bool ignoreResultI32 = ignoreResult && T is int32_t;
@@ -704,8 +704,8 @@ bool textureFuncs( Sampler1D<vector<T,N>> gsampler1D
704704

705705
// GLSL: textureOffset({{.*}}sampler1D
706706
// SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler1D
707-
// S-PIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset %
708-
&& gvec4(T(0)) == textureOffset(gsampler1D, float(coord), int(0))
707+
// SPIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset %
708+
&& gvec4(T(0)) == textureOffset(gsampler1D, float(coord), int(1))
709709

710710
// GLSL: textureOffset({{.*}}sampler1D
711711
// SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler1D
@@ -714,7 +714,7 @@ bool textureFuncs( Sampler1D<vector<T,N>> gsampler1D
714714

715715
// GLSL: textureOffset({{.*}}sampler2D
716716
// SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2D
717-
// S-PIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset %
717+
// SPIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset %
718718
&& gvec4(T(0)) == textureOffset(gsampler2D, vec2(coord), offset2D)
719719

720720
// GLSL: textureOffset({{.*}}sampler2D
@@ -724,7 +724,7 @@ bool textureFuncs( Sampler1D<vector<T,N>> gsampler1D
724724

725725
// GLSL: textureOffset({{.*}}sampler3D
726726
// SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler3D
727-
// S-PIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset %
727+
// SPIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset %
728728
&& gvec4(T(0)) == textureOffset(gsampler3D, vec3(coord), offset3D)
729729

730730
// GLSL: textureOffset({{.*}}sampler3D
@@ -734,7 +734,7 @@ bool textureFuncs( Sampler1D<vector<T,N>> gsampler1D
734734

735735
// GLSL: textureOffset({{.*}}sampler2DShadow
736736
// SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2DShadow
737-
// S-PIR: OpImageSampleDrefImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset %
737+
// SPIR: OpImageSampleDrefImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset %
738738
&& (float(0) == textureOffset(uniform_sampler2DShadow, vec3(coord), offset2D) || ignoreResultF32)
739739

740740
// GLSL: textureOffset({{.*}}sampler2DShadow
@@ -744,12 +744,12 @@ bool textureFuncs( Sampler1D<vector<T,N>> gsampler1D
744744

745745
// GLSL: textureOffset({{.*}}sampler2DRect
746746
// SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2DRect
747-
// S-PIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset %
747+
// SPIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset %
748748
&& gvec4(T(0)) == textureOffset(gsampler2DRect, vec2(coord), offset2D)
749749

750750
// GLSL: textureOffset({{.*}}sampler2DRectShadow
751751
// SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2DRectShadow
752-
// S-PIR: OpImageSampleDrefImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset %
752+
// SPIR: OpImageSampleDrefImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset %
753753
&& float(0) == textureOffset(uniform_sampler2DRectShadow, vec3(coord), offset2D)
754754

755755
// GLSL: textureOffset({{.*}}sampler1DShadow
@@ -764,8 +764,8 @@ bool textureFuncs( Sampler1D<vector<T,N>> gsampler1D
764764

765765
// GLSL: textureOffset({{.*}}sampler1DArray
766766
// SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler1DArray
767-
// S-PIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset %
768-
&& gvec4(T(0)) == textureOffset(gsampler1DArray, vec2(coord), int(0))
767+
// SPIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset %
768+
&& gvec4(T(0)) == textureOffset(gsampler1DArray, vec2(coord), int(1))
769769

770770
// GLSL: textureOffset({{.*}}sampler1DArray
771771
// SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler1DArray
@@ -774,7 +774,7 @@ bool textureFuncs( Sampler1D<vector<T,N>> gsampler1D
774774

775775
// GLSL: textureOffset({{.*}}sampler2DArray
776776
// SPIR: [[LOAD:%[1-9][0-9]*]] = OpLoad{{.*}}sampler2DArray
777-
// S-PIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}} ConstOffset %
777+
// SPIR: OpImageSampleImplicitLod {{.*}}[[LOAD]]{{.*}}ConstOffset %
778778
&& gvec4(T(0)) == textureOffset(gsampler2DArray, vec3(coord), offset2D)
779779

780780
// GLSL: textureOffset({{.*}}sampler2DArray

0 commit comments

Comments
 (0)