Skip to content

Commit

Permalink
[WGSL] Texture read needs unsigned coordinates
Browse files Browse the repository at this point in the history 
https://bugs.webkit.org/show_bug.cgi?id=261885
rdar://115843632

Reviewed by Dan Glastonbury.

We convert textureLoad(texture, coordinate) into texture.read(coordinate), but
the problem is that texture load allows either signed or unsigned values in
the coordinate argument, but metal only allows unsigned values.

* Source/WebGPU/WGSL/Metal/MetalFunctionWriter.cpp:
(WGSL::Metal::FunctionDefinitionWriter::visit):

Canonical link: https://commits.webkit.org/268388@main
  • Loading branch information
@tadeuzagallo
tadeuzagallo committed Sep 25, 2023
1 parent 2a66767 commit 88ae5ff
Show file tree
Hide file tree
Showing 2 changed files with 59 additions and 29 deletions.
40 changes: 34 additions & 6 deletions Source/WebGPU/WGSL/Metal/MetalFunctionWriter.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -610,6 +610,7 @@ void FunctionDefinitionWriter::visit(const Type* type)
},
[&](const Texture& texture) {
const char* type;
const char* access = "sample";
switch (texture.kind) {
case Types::Texture::Kind::Texture1d:
type = "texture1d";
Expand All @@ -631,11 +632,12 @@ void FunctionDefinitionWriter::visit(const Type* type)
break;
case Types::Texture::Kind::TextureMultisampled2d:
type = "texture2d_ms";
access = "read";
break;
}
m_stringBuilder.append(type, "<");
visit(texture.element);
m_stringBuilder.append(", access::sample>");
m_stringBuilder.append(", access::", access, ">");
},
[&](const TextureStorage& texture) {
const char* base;
Expand Down Expand Up @@ -879,7 +881,33 @@ void FunctionDefinitionWriter::visit(const Type* type, AST::CallExpression& call
if (!isExternalTexture) {
writer->visit(call.arguments()[0]);
writer->stringBuilder().append(".read");
visitArguments(writer, call, 1);
bool first = true;
writer->stringBuilder().append("(");
const char* cast = "uint";
if (const auto* vector = std::get_if<Types::Vector>(call.arguments()[1].inferredType())) {
switch (vector->size) {
case 2:
cast = "uint2";
break;
case 3:
cast = "uint3";
break;
default:
RELEASE_ASSERT_NOT_REACHED();
}
}
for (unsigned i = 1; i < call.arguments().size(); ++i) {
if (first) {
writer->stringBuilder().append(cast, "(");
writer->visit(call.arguments()[i]);
writer->stringBuilder().append(")");
} else {
writer->stringBuilder().append(", ");
writer->visit(call.arguments()[i]);
}
first = false;
}
writer->stringBuilder().append(")");
return;
}

Expand All @@ -888,16 +916,16 @@ void FunctionDefinitionWriter::visit(const Type* type, AST::CallExpression& call
{
IndentationScope scope(writer->indent());
{
writer->stringBuilder().append(writer->indent(), "auto __coords = (");
writer->stringBuilder().append(writer->indent(), "auto __coords = uint2((");
writer->visit(texture);
writer->stringBuilder().append(".UVRemapMatrix * float3(");
writer->stringBuilder().append(".UVRemapMatrix * float3(float2(");
writer->visit(coordinates);
writer->stringBuilder().append(", 1)).xy;\n");
writer->stringBuilder().append("), 1)).xy);\n");
}
{
writer->stringBuilder().append(writer->indent(), "auto __y = float(");
writer->visit(texture);
writer->stringBuilder().append(".FirstPlane.read(__cords).r);\n");
writer->stringBuilder().append(".FirstPlane.read(__coords).r);\n");
}
{
writer->stringBuilder().append(writer->indent(), "auto __cbcr = float2(");
Expand Down
48 changes: 25 additions & 23 deletions Source/WebGPU/WGSL/tests/valid/overload.wgsl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -2022,36 +2022,36 @@ fn testTrunc()

// 16.7. Texture Built-in Functions (https://gpuweb.github.io/gpuweb/wgsl/#texture-builtin-functions)

var s: sampler;
@group(0) @binding( 0) var s: sampler;

var t1d: texture_1d<f32>;
var t1di: texture_1d<i32>;
var t1du: texture_1d<u32>;
@group(0) @binding( 1) var t1d: texture_1d<f32>;
@group(0) @binding( 2) var t1di: texture_1d<i32>;
@group(0) @binding( 3) var t1du: texture_1d<u32>;

var t2d: texture_2d<f32>;
var t2di: texture_2d<i32>;
var t2du: texture_2d<u32>;
@group(0) @binding( 4) var t2d: texture_2d<f32>;
@group(0) @binding( 5) var t2di: texture_2d<i32>;
@group(0) @binding( 6) var t2du: texture_2d<u32>;

var t2da: texture_2d_array<f32>;
var t2dai: texture_2d_array<i32>;
var t2dau: texture_2d_array<u32>;
@group(0) @binding( 7) var t2da: texture_2d_array<f32>;
@group(0) @binding( 8) var t2dai: texture_2d_array<i32>;
@group(0) @binding( 9) var t2dau: texture_2d_array<u32>;

var t3d: texture_3d<f32>;
var t3di: texture_3d<i32>;
var t3du: texture_3d<u32>;
@group(0) @binding(10) var t3d: texture_3d<f32>;
@group(0) @binding(11) var t3di: texture_3d<i32>;
@group(0) @binding(12) var t3du: texture_3d<u32>;

var tm2d: texture_multisampled_2d<f32>;
var tm2di: texture_multisampled_2d<i32>;
var tm2du: texture_multisampled_2d<u32>;
@group(0) @binding(13) var tm2d: texture_multisampled_2d<f32>;
@group(0) @binding(14) var tm2di: texture_multisampled_2d<i32>;
@group(0) @binding(15) var tm2du: texture_multisampled_2d<u32>;

var te: texture_external;
var tc: texture_cube<f32>;
var tca: texture_cube_array<f32>;
@group(0) @binding(16) var te: texture_external;
@group(0) @binding(17) var tc: texture_cube<f32>;
@group(0) @binding(18) var tca: texture_cube_array<f32>;

var ts1d: texture_storage_2d<rgba8unorm, read>;
var ts2d: texture_storage_2d<rgba16uint, write>;
var ts2da: texture_storage_2d<r32sint, read_write>;
var ts3d: texture_storage_2d<rgba32float, write>;
@group(0) @binding(19) var ts1d: texture_storage_2d<rgba8unorm, read>;
@group(0) @binding(20) var ts2d: texture_storage_2d<rgba16uint, write>;
@group(0) @binding(21) var ts2da: texture_storage_2d<r32sint, read_write>;
@group(0) @binding(22) var ts3d: texture_storage_2d<rgba32float, write>;

var td2d: texture_depth_2d;
var td2da: texture_depth_2d_array;
Expand Down Expand Up @@ -2163,6 +2163,8 @@ fn testTextureGather()
// FIXME: this only applies to texture_depth, implement

// 16.7.4
// RUN: %metal-compile testTextureLoad
@compute @workgroup_size(1)
fn testTextureLoad()
{
// [T < ConcreteInteger, U < ConcreteInteger, S < Concrete32BitNumber].(Texture[S, Texture1d], T, U) => Vector[S, 4],
Expand Down

0 comments on commit 88ae5ff

Please sign in to comment.