[tint][wgsl] Add unrestricted_pointer_parameters feature
This WGSL feature is currently exposed as experimental.
The feature replaces the chromium_experimental_full_ptr_parameters extension.
Bug: tint:2053
Change-Id: I84cec96496ae16c6ece82716dcf12b0179bb5b6d
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/169264
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/dawn/common/WGSLFeatureMapping.h b/src/dawn/common/WGSLFeatureMapping.h
index 4810cee..3c812dab 100644
--- a/src/dawn/common/WGSLFeatureMapping.h
+++ b/src/dawn/common/WGSLFeatureMapping.h
@@ -30,9 +30,12 @@
#define DAWN_FOREACH_WGSL_FEATURE(X) \
X(kUndefined, Undefined) \
- X(kReadonlyAndReadwriteStorageTextures, ReadonlyAndReadwriteStorageTextures) \
+ /* ------------------------------------------------------------------------ */ \
X(kPacked4X8IntegerDotProduct, Packed4x8IntegerDotProduct) \
X(kPointerCompositeAccess, PointerCompositeAccess) \
+ X(kReadonlyAndReadwriteStorageTextures, ReadonlyAndReadwriteStorageTextures) \
+ X(kUnrestrictedPointerParameters, UnrestrictedPointerParameters) \
+ /* ----------------- entries below are only for testing ------------------ */ \
X(kChromiumTestingUnimplemented, ChromiumTestingUnimplemented) \
X(kChromiumTestingUnsafeExperimental, ChromiumTestingUnsafeExperimental) \
X(kChromiumTestingExperimental, ChromiumTestingExperimental) \
diff --git a/src/dawn/tests/end2end/ShaderTests.cpp b/src/dawn/tests/end2end/ShaderTests.cpp
index e3ba8c0..b2c4762 100644
--- a/src/dawn/tests/end2end/ShaderTests.cpp
+++ b/src/dawn/tests/end2end/ShaderTests.cpp
@@ -38,12 +38,19 @@
class ShaderTests : public DawnTest {
public:
- wgpu::Buffer CreateBuffer(const uint32_t count) {
- std::vector<uint32_t> data(count, 0);
+ wgpu::Buffer CreateBuffer(const std::vector<uint32_t>& data,
+ wgpu::BufferUsage usage = wgpu::BufferUsage::Storage |
+ wgpu::BufferUsage::CopySrc) {
uint64_t bufferSize = static_cast<uint64_t>(data.size() * sizeof(uint32_t));
- return utils::CreateBufferFromData(device, data.data(), bufferSize,
- wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc);
+ return utils::CreateBufferFromData(device, data.data(), bufferSize, usage);
}
+
+ wgpu::Buffer CreateBuffer(const uint32_t count,
+ wgpu::BufferUsage usage = wgpu::BufferUsage::Storage |
+ wgpu::BufferUsage::CopySrc) {
+ return CreateBuffer(std::vector<uint32_t>(count, 0), usage);
+ }
+
wgpu::ComputePipeline CreateComputePipeline(
const std::string& shader,
const char* entryPoint = nullptr,
@@ -2265,6 +2272,94 @@
EXPECT_BUFFER_U32_EQ(42, output, 0);
}
+TEST_P(ShaderTests, UnrestrictedPointerParameters) {
+ // TODO(crbug.com/dawn/2350): Investigate, fix.
+ DAWN_TEST_UNSUPPORTED_IF(IsD3D11());
+
+ wgpu::ComputePipeline pipeline = CreateComputePipeline(R"(
+@binding(0) @group(0) var<uniform> input : array<vec4i, 4>;
+@binding(1) @group(0) var<storage, read_write> output : array<vec4i, 4>;
+
+fn sum(f : ptr<function, i32>,
+ w : ptr<workgroup, atomic<i32>>,
+ p : ptr<private, i32>,
+ u : ptr<uniform, vec4i>) -> vec4i {
+
+ return vec4(*f + atomicLoad(w) + *p) + *u;
+}
+
+struct S {
+ i : i32,
+}
+
+var<private> P0 = S(0);
+var<private> P1 = S(10);
+var<private> P2 = 20;
+var<private> P3 = 30;
+
+struct T {
+ i : atomic<i32>,
+}
+
+var<workgroup> W0 : T;
+var<workgroup> W1 : atomic<i32>;
+var<workgroup> W2 : T;
+var<workgroup> W3 : atomic<i32>;
+
+@compute @workgroup_size(1)
+fn main() {
+ atomicStore(&W0.i, 0);
+ atomicStore(&W1, 100);
+ atomicStore(&W2.i, 200);
+ atomicStore(&W3, 300);
+
+ var F = array(0, 1000, 2000, 3000);
+
+ output[0] = sum(&F[2], &W3, &P1.i, &input[0]); // vec4(2310) + vec4(1, 2, 3, 4)
+ output[1] = sum(&F[1], &W2.i, &P0.i, &input[1]); // vec4(1200) + vec4(4, 3, 2, 1)
+ output[2] = sum(&F[3], &W0.i, &P3, &input[2]); // vec4(3030) + vec4(2, 4, 1, 3)
+ output[3] = sum(&F[2], &W1, &P2, &input[3]); // vec4(2120) + vec4(4, 1, 2, 3)
+}
+)");
+
+ wgpu::Buffer input = CreateBuffer(
+ std::vector<uint32_t>{
+ 1, 2, 3, 4, // [0]
+ 4, 3, 2, 1, // [1]
+ 2, 4, 1, 3, // [2]
+ 4, 1, 2, 3, // [3]
+ },
+ wgpu::BufferUsage::Uniform);
+
+ std::vector<uint32_t> expected{
+ 2311, 2312, 2313, 2314, // [0]
+ 1204, 1203, 1202, 1201, // [1]
+ 3032, 3034, 3031, 3033, // [2]
+ 2124, 2121, 2122, 2123, // [3]
+ };
+
+ wgpu::Buffer output = CreateBuffer(expected.size());
+
+ wgpu::BindGroup bindGroup =
+ utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, input}, {1, output}});
+
+ wgpu::CommandBuffer commands;
+ {
+ wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+ wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+ pass.SetPipeline(pipeline);
+ pass.SetBindGroup(0, bindGroup);
+ pass.DispatchWorkgroups(1);
+ pass.End();
+
+ commands = encoder.Finish();
+ }
+
+ queue.Submit(1, &commands);
+
+ EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), output, 0, expected.size());
+}
+
DAWN_INSTANTIATE_TEST(ShaderTests,
D3D11Backend(),
D3D12Backend(),
diff --git a/src/dawn/tests/unittests/validation/ShaderModuleValidationTests.cpp b/src/dawn/tests/unittests/validation/ShaderModuleValidationTests.cpp
index bf07372..1e6c588 100644
--- a/src/dawn/tests/unittests/validation/ShaderModuleValidationTests.cpp
+++ b/src/dawn/tests/unittests/validation/ShaderModuleValidationTests.cpp
@@ -802,7 +802,6 @@
// Currently the following WGSL extensions are not enabled under any situation.
/*
- {"chromium_experimental_full_ptr_parameters", true, nullptr},
{"chromium_experimental_push_constant", true, nullptr},
{"chromium_internal_relaxed_uniform_layout", true, nullptr},
*/
diff --git a/src/tint/cmd/fuzz/ir/fuzz.cc b/src/tint/cmd/fuzz/ir/fuzz.cc
index 5fa170f..8fbf4cb 100644
--- a/src/tint/cmd/fuzz/ir/fuzz.cc
+++ b/src/tint/cmd/fuzz/ir/fuzz.cc
@@ -46,7 +46,6 @@
bool IsUnsupported(const ast::Enable* enable) {
for (auto ext : enable->extensions) {
switch (ext->name) {
- case tint::wgsl::Extension::kChromiumExperimentalFullPtrParameters:
case tint::wgsl::Extension::kChromiumExperimentalPixelLocal:
case tint::wgsl::Extension::kChromiumExperimentalPushConstant:
case tint::wgsl::Extension::kChromiumInternalDualSourceBlending:
diff --git a/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc b/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc
index 49780b7..f0f2e00 100644
--- a/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc
+++ b/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc
@@ -59,7 +59,9 @@
class DirectVariableAccessTest : public TransformTestBase<testing::Test> {
public:
- std::string Run(std::string in, const DirectVariableAccessOptions& options = {}) {
+ std::string Run(std::string in,
+ const DirectVariableAccessOptions& transform_options = {},
+ const wgsl::writer::ProgramOptions program_options = {}) {
wgsl::reader::Options parser_options;
parser_options.allowed_features = wgsl::AllowedFeatures::Everything();
Source::File file{"test", in};
@@ -73,7 +75,7 @@
return "ProgramToIR() failed:\n" + module.Failure().reason.str();
}
- auto res = DirectVariableAccess(module.Get(), options);
+ auto res = DirectVariableAccess(module.Get(), transform_options);
if (res != Success) {
return "DirectVariableAccess failed:\n" + res.Failure().reason.str();
}
@@ -84,9 +86,6 @@
return "wgsl::writer::Raise failed:\n" + res.Failure().reason.str();
}
- wgsl::writer::ProgramOptions program_options;
- program_options.allowed_features.extensions.insert(
- wgsl::Extension::kChromiumExperimentalFullPtrParameters);
auto program_out = wgsl::writer::IRToProgram(module.Get(), program_options);
if (!program_out.IsValid()) {
return "wgsl::writer::IRToProgram() failed: \n" + program_out.Diagnostics().str() +
@@ -116,8 +115,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_RemoveUncalled, PtrUniform) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> keep_me : i32 = 42i;
fn u(pre : i32, p : ptr<uniform, i32>, post : i32) -> i32 {
@@ -137,8 +134,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_RemoveUncalled, PtrStorage) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> keep_me : i32 = 42i;
fn s(pre : i32, p : ptr<storage, i32>, post : i32) -> i32 {
@@ -157,8 +152,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_RemoveUncalled, PtrWorkgroup) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> keep_me : i32 = 42i;
fn w(pre : i32, p : ptr<workgroup, i32>, post : i32) -> i32 {
@@ -255,8 +248,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_PtrChains, ConstantIndices) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 8>, 8>, 8>;
fn a(pre : i32, p : ptr<uniform, vec4<i32>>, post : i32) -> vec4<i32> {
@@ -312,8 +303,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_PtrChains, DynamicIndices) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 8>, 8>, 8>;
var<private> i : i32;
@@ -401,8 +390,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_PtrChains, DynamicIndicesForLoopInit) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8>, 8>;
var<private> i : i32;
@@ -479,8 +466,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_PtrChains, DynamicIndicesForLoopCond) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8>, 8>;
var<private> i : i32;
@@ -563,8 +548,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_PtrChains, DynamicIndicesForLoopCont) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8>, 8>;
var<private> i : i32;
@@ -647,8 +630,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_PtrChains, DynamicIndicesWhileCond) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8>, 8>;
var<private> i : i32;
@@ -739,8 +720,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_UniformAS, Param_ptr_i32_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : i32;
fn a(pre : i32, p : ptr<uniform, i32>, post : i32) -> i32 {
@@ -771,8 +750,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_UniformAS, Param_ptr_vec4i32_Via_array_DynamicRead) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<vec4<i32>, 8>;
fn a(pre : i32, p : ptr<uniform, vec4<i32>>, post : i32) -> vec4<i32> {
@@ -805,8 +782,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_UniformAS, CallChaining) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
};
@@ -946,8 +921,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_StorageAS, Param_ptr_i32_Via_struct_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
@@ -986,8 +959,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_StorageAS, Param_ptr_arr_i32_Via_struct_write) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
};
@@ -1026,8 +997,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_StorageAS, Param_ptr_vec4i32_Via_array_DynamicWrite) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> S : array<vec4<i32>, 8>;
fn a(pre : i32, p : ptr<storage, vec4<i32>, read_write>, post : i32) {
@@ -1060,8 +1029,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_StorageAS, CallChaining) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
};
@@ -1201,8 +1168,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_WorkgroupAS, Param_ptr_vec4i32_Via_array_StaticRead) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> W : array<vec4<i32>, 8>;
fn a(pre : i32, p : ptr<workgroup, vec4<i32>>, post : i32) -> vec4<i32> {
@@ -1233,8 +1198,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_WorkgroupAS, Param_ptr_vec4i32_Via_array_StaticWrite) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> W : array<vec4<i32>, 8>;
fn a(pre : i32, p : ptr<workgroup, vec4<i32>>, post : i32) {
@@ -1265,8 +1228,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_WorkgroupAS, CallChaining) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
};
@@ -1406,8 +1367,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_PrivateAS, Enabled_Param_ptr_i32_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn a(pre : i32, p : ptr<private, i32>, post : i32) -> i32 {
return *(p);
}
@@ -1438,8 +1397,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_PrivateAS, Enabled_Param_ptr_i32_write) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn a(pre : i32, p : ptr<private, i32>, post : i32) {
*(p) = 42;
}
@@ -1470,8 +1427,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_PrivateAS, Enabled_Param_ptr_i32_Via_struct_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
@@ -1510,8 +1465,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_PrivateAS, Disabled_Param_ptr_i32_Via_struct_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
@@ -1529,15 +1482,16 @@
auto* expect = src;
- auto got = Run(src);
+ wgsl::writer::ProgramOptions program_options;
+ program_options.allowed_features.features.emplace(
+ wgsl::LanguageFeature::kUnrestrictedPointerParameters);
+ auto got = Run(src, /* transform_options */ {}, program_options);
EXPECT_EQ(expect, got);
}
TEST_F(IR_DirectVariableAccessWgslTest_PrivateAS, Enabled_Param_ptr_arr_i32_Via_struct_write) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
};
@@ -1576,8 +1530,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_PrivateAS, Disabled_Param_ptr_arr_i32_Via_struct_write) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4u>,
}
@@ -1595,15 +1547,16 @@
auto* expect = src;
- auto got = Run(src);
+ wgsl::writer::ProgramOptions program_options;
+ program_options.allowed_features.features.emplace(
+ wgsl::LanguageFeature::kUnrestrictedPointerParameters);
+ auto got = Run(src, /* transform_options */ {}, program_options);
EXPECT_EQ(expect, got);
}
TEST_F(IR_DirectVariableAccessWgslTest_PrivateAS, Enabled_Param_ptr_i32_mixed) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
@@ -1660,8 +1613,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_PrivateAS, Disabled_Param_ptr_i32_mixed) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> Pi : i32;
struct str {
@@ -1685,15 +1636,16 @@
auto* expect = src;
- auto got = Run(src);
+ wgsl::writer::ProgramOptions program_options;
+ program_options.allowed_features.features.emplace(
+ wgsl::LanguageFeature::kUnrestrictedPointerParameters);
+ auto got = Run(src, /* transform_options */ {}, program_options);
EXPECT_EQ(expect, got);
}
TEST_F(IR_DirectVariableAccessWgslTest_PrivateAS, Enabled_CallChaining) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
};
@@ -1824,8 +1776,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_PrivateAS, Disabled_CallChaining) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
}
@@ -1877,7 +1827,10 @@
auto* expect = src;
- auto got = Run(src);
+ wgsl::writer::ProgramOptions program_options;
+ program_options.allowed_features.features.emplace(
+ wgsl::LanguageFeature::kUnrestrictedPointerParameters);
+ auto got = Run(src, /* transform_options */ {}, program_options);
EXPECT_EQ(expect, got);
}
@@ -1909,8 +1862,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_FunctionAS, Enabled_Param_ptr_i32_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn a(pre : i32, p : ptr<function, i32>, post : i32) -> i32 {
return *(p);
}
@@ -1939,8 +1890,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_FunctionAS, Enabled_Param_ptr_i32_write) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn a(pre : i32, p : ptr<function, i32>, post : i32) {
*(p) = 42;
}
@@ -1969,8 +1918,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_FunctionAS, Enabled_Param_ptr_i32_Via_struct_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
@@ -2007,8 +1954,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_FunctionAS, Enabled_Param_ptr_arr_i32_Via_struct_write) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
};
@@ -2045,8 +1990,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_FunctionAS, Enabled_Param_ptr_i32_mixed) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
@@ -2100,8 +2043,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_FunctionAS, Disabled_Param_ptr_i32_Via_struct_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn a(pre : i32, p : ptr<function, i32>, post : i32) -> i32 {
return *(p);
}
@@ -2118,15 +2059,16 @@
auto* expect = src;
- auto got = Run(src);
+ wgsl::writer::ProgramOptions program_options;
+ program_options.allowed_features.features.emplace(
+ wgsl::LanguageFeature::kUnrestrictedPointerParameters);
+ auto got = Run(src, /* transform_options */ {}, program_options);
EXPECT_EQ(expect, got);
}
TEST_F(IR_DirectVariableAccessWgslTest_FunctionAS, Disabled_Param_ptr_arr_i32_Via_struct_write) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn a(pre : i32, p : ptr<function, array<i32, 4u>>, post : i32) {
*(p) = array<i32, 4u>();
}
@@ -2143,7 +2085,10 @@
auto* expect = src;
- auto got = Run(src);
+ wgsl::writer::ProgramOptions program_options;
+ program_options.allowed_features.features.emplace(
+ wgsl::LanguageFeature::kUnrestrictedPointerParameters);
+ auto got = Run(src, /* transform_options */ {}, program_options);
EXPECT_EQ(expect, got);
}
@@ -2159,8 +2104,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_BuiltinFn, ArrayLength) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : array<f32>;
fn len(p : ptr<storage, array<f32>>) -> u32 {
@@ -2191,8 +2134,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_BuiltinFn, WorkgroupUniformLoad) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> W : f32;
fn load(p : ptr<workgroup, f32>) -> f32 {
@@ -2232,8 +2173,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_Complex, Param_ptr_mixed_vec4i32_ViaMultiple) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<i32>,
};
@@ -2418,8 +2357,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_Complex, Indexing) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : array<array<array<array<i32, 9>, 9>, 9>, 50>;
fn a(i : i32) -> i32 { return i; }
@@ -2459,8 +2396,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_Complex, IndexingInPtrCall) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : array<array<array<array<i32, 9>, 9>, 9>, 50>;
fn a(pre : i32, i : ptr<storage, i32>, post : i32) -> i32 {
@@ -2504,8 +2439,6 @@
TEST_F(IR_DirectVariableAccessWgslTest_Complex, IndexingDualPointers) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : array<array<array<i32, 9>, 9>, 50>;
@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 9>, 9>, 50>;
diff --git a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
index 9a7c6ea..2026bd2 100644
--- a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
@@ -266,7 +266,6 @@
"GLSL", builder_.AST(), diagnostics_,
Vector{
wgsl::Extension::kChromiumDisableUniformityAnalysis,
- wgsl::Extension::kChromiumExperimentalFullPtrParameters,
wgsl::Extension::kChromiumInternalDualSourceBlending,
wgsl::Extension::kChromiumExperimentalPushConstant,
wgsl::Extension::kF16,
diff --git a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
index 3415d05..d6eec1a 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
@@ -385,7 +385,6 @@
"HLSL", builder_.AST(), diagnostics_,
Vector{
wgsl::Extension::kChromiumDisableUniformityAnalysis,
- wgsl::Extension::kChromiumExperimentalFullPtrParameters,
wgsl::Extension::kChromiumExperimentalPushConstant,
wgsl::Extension::kChromiumExperimentalSubgroups,
wgsl::Extension::kF16,
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
index 3fb75a8..63dcf81 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -277,7 +277,6 @@
"MSL", builder_.AST(), diagnostics_,
Vector{
wgsl::Extension::kChromiumDisableUniformityAnalysis,
- wgsl::Extension::kChromiumExperimentalFullPtrParameters,
wgsl::Extension::kChromiumExperimentalPixelLocal,
wgsl::Extension::kChromiumExperimentalSubgroups,
wgsl::Extension::kChromiumExperimentalFramebufferFetch,
diff --git a/src/tint/lang/msl/writer/ast_raise/module_scope_var_to_entry_point_param.cc b/src/tint/lang/msl/writer/ast_raise/module_scope_var_to_entry_point_param.cc
index 38cad89..8335eac 100644
--- a/src/tint/lang/msl/writer/ast_raise/module_scope_var_to_entry_point_param.cc
+++ b/src/tint/lang/msl/writer/ast_raise/module_scope_var_to_entry_point_param.cc
@@ -375,9 +375,6 @@
if (!private_struct_members.IsEmpty()) {
// Create the private variable struct.
ctx.dst->Structure(PrivateStructName(), std::move(private_struct_members));
- // Passing a pointer to a private variable will now involve passing a pointer to the
- // member of a structure, so enable the extension that allows this.
- ctx.dst->Enable(wgsl::Extension::kChromiumExperimentalFullPtrParameters);
}
// Build a list of `&ident` expressions. We'll use this later to avoid generating
diff --git a/src/tint/lang/msl/writer/ast_raise/module_scope_var_to_entry_point_param_test.cc b/src/tint/lang/msl/writer/ast_raise/module_scope_var_to_entry_point_param_test.cc
index f7630f5..3bbb4fe 100644
--- a/src/tint/lang/msl/writer/ast_raise/module_scope_var_to_entry_point_param_test.cc
+++ b/src/tint/lang/msl/writer/ast_raise/module_scope_var_to_entry_point_param_test.cc
@@ -62,8 +62,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct tint_private_vars_struct {
p : f32,
}
@@ -93,8 +91,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct tint_private_vars_struct {
p : f32,
}
@@ -143,8 +139,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct tint_private_vars_struct {
p : f32,
}
@@ -214,8 +208,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct tint_private_vars_struct {
p : f32,
}
@@ -266,8 +258,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct tint_private_vars_struct {
a : f32,
b : f32,
@@ -299,8 +289,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct tint_private_vars_struct {
a : f32,
b : f32,
@@ -335,8 +323,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct tint_private_vars_struct {
p : f32,
}
@@ -372,8 +358,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct tint_private_vars_struct {
p : f32,
}
@@ -1232,8 +1216,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct tint_private_vars_struct {
p : f32,
p_with_init : f32,
@@ -1278,8 +1260,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
a : f32,
b : f32,
@@ -1339,8 +1319,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
a : f32,
b : f32,
diff --git a/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc b/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc
index 325f528..220dabb 100644
--- a/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc
+++ b/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc
@@ -4652,8 +4652,6 @@
// Test that we can pass a pointer to the vec3 member of the modf return struct to a function
// parameter to which we also pass a pointer to a vec3 member on a host-shareable struct.
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
v : vec3<f32>
}
@@ -4676,7 +4674,6 @@
auto* expect = R"(
enable chromium_internal_relaxed_uniform_layout;
-enable chromium_experimental_full_ptr_parameters;
struct S_tint_packed_vec3 {
@align(16)
@@ -5305,8 +5302,6 @@
TEST_F(PackedVec3Test, VectorPointerParameters) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
v : vec3<f32>,
m : mat3x3<f32>,
@@ -5350,7 +5345,6 @@
auto* expect = R"(
enable chromium_internal_relaxed_uniform_layout;
-enable chromium_experimental_full_ptr_parameters;
struct tint_packed_vec3_f32_array_element {
@align(16)
@@ -5421,8 +5415,6 @@
TEST_F(PackedVec3Test, MatrixPointerParameters) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
m : mat3x3<f32>,
arr_m : array<mat3x3<f32>, 4>,
@@ -5454,7 +5446,6 @@
auto* expect = R"(
enable chromium_internal_relaxed_uniform_layout;
-enable chromium_experimental_full_ptr_parameters;
struct tint_packed_vec3_f32_array_element {
@align(16)
@@ -5523,8 +5514,6 @@
TEST_F(PackedVec3Test, ArrayOfVectorPointerParameters) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
arr_v : array<vec3<f32>, 4>,
}
@@ -5550,7 +5539,6 @@
auto* expect = R"(
enable chromium_internal_relaxed_uniform_layout;
-enable chromium_experimental_full_ptr_parameters;
struct tint_packed_vec3_f32_array_element {
@align(16)
@@ -5610,8 +5598,6 @@
TEST_F(PackedVec3Test, ArrayOfMatrixPointerParameters) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
arr_m : array<mat3x3<f32>, 4>,
}
@@ -5637,7 +5623,6 @@
auto* expect = R"(
enable chromium_internal_relaxed_uniform_layout;
-enable chromium_experimental_full_ptr_parameters;
struct tint_packed_vec3_f32_array_element {
@align(16)
@@ -5713,8 +5698,6 @@
TEST_F(PackedVec3Test, StructPointerParameters) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
v : vec3<f32>,
m : mat3x3<f32>,
@@ -5740,7 +5723,6 @@
auto* expect = R"(
enable chromium_internal_relaxed_uniform_layout;
-enable chromium_experimental_full_ptr_parameters;
struct tint_packed_vec3_f32_array_element {
@align(16)
@@ -6016,8 +5998,6 @@
// Test that we can pass a pointers to a members of both shared and non-shared structs to the
// same function.
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
v : vec3<f32>,
arr : array<vec3<f32>, 4>,
@@ -6047,7 +6027,6 @@
auto* expect = R"(
enable chromium_internal_relaxed_uniform_layout;
-enable chromium_experimental_full_ptr_parameters;
struct tint_packed_vec3_f32_array_element {
@align(16)
@@ -6683,8 +6662,6 @@
TEST_F(PackedVec3Test, MixedAddressSpace_PointerParameters) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
v : vec3<f32>,
m : mat3x3<f32>,
@@ -6715,7 +6692,6 @@
auto* expect = R"(
enable chromium_internal_relaxed_uniform_layout;
-enable chromium_experimental_full_ptr_parameters;
struct tint_packed_vec3_f32_array_element {
@align(16)
@@ -7649,8 +7625,6 @@
// matrix into an array of vec3s in uniform storage.
auto* src = R"(
enable f16;
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> m : mat3x3<f16>;
fn g(p : ptr<uniform, mat3x3<f16>>) -> vec3<f16> {
@@ -7666,7 +7640,6 @@
R"(
enable chromium_internal_relaxed_uniform_layout;
enable f16;
-enable chromium_experimental_full_ptr_parameters;
struct tint_packed_vec3_f16_array_element {
@align(8)
diff --git a/src/tint/lang/spirv/writer/ast_printer/builder.cc b/src/tint/lang/spirv/writer/ast_printer/builder.cc
index b736b28..1895d11 100644
--- a/src/tint/lang/spirv/writer/ast_printer/builder.cc
+++ b/src/tint/lang/spirv/writer/ast_printer/builder.cc
@@ -277,7 +277,6 @@
"SPIR-V", builder_.AST(), builder_.Diagnostics(),
Vector{
wgsl::Extension::kChromiumDisableUniformityAnalysis,
- wgsl::Extension::kChromiumExperimentalFullPtrParameters,
wgsl::Extension::kChromiumExperimentalPushConstant,
wgsl::Extension::kChromiumExperimentalSubgroups,
wgsl::Extension::kF16,
diff --git a/src/tint/lang/wgsl/ast/module_test.cc b/src/tint/lang/wgsl/ast/module_test.cc
index 842a357..8e5f629 100644
--- a/src/tint/lang/wgsl/ast/module_test.cc
+++ b/src/tint/lang/wgsl/ast/module_test.cc
@@ -147,7 +147,7 @@
TEST_F(ModuleTest, Directives) {
auto* enable_1 = Enable(wgsl::Extension::kF16);
auto* diagnostic_1 = DiagnosticDirective(wgsl::DiagnosticSeverity::kWarning, "foo");
- auto* enable_2 = Enable(wgsl::Extension::kChromiumExperimentalFullPtrParameters);
+ auto* enable_2 = Enable(wgsl::Extension::kChromiumExperimentalPixelLocal);
auto* diagnostic_2 = DiagnosticDirective(wgsl::DiagnosticSeverity::kOff, "bar");
Program program(std::move(*this));
diff --git a/src/tint/lang/wgsl/ast/transform/builtin_polyfill.cc b/src/tint/lang/wgsl/ast/transform/builtin_polyfill.cc
index adc8df2..b3a8ece 100644
--- a/src/tint/lang/wgsl/ast/transform/builtin_polyfill.cc
+++ b/src/tint/lang/wgsl/ast/transform/builtin_polyfill.cc
@@ -61,15 +61,7 @@
/// Constructor
/// @param program the source program
/// @param config the transform config
- State(const Program& program, const Config& config) : src(program), cfg(config) {
- has_full_ptr_params = false;
- for (auto* enable : src.AST().Enables()) {
- if (enable->HasExtension(wgsl::Extension::kChromiumExperimentalFullPtrParameters)) {
- has_full_ptr_params = true;
- break;
- }
- }
- }
+ State(const Program& program, const Config& config) : src(program), cfg(config) {}
/// Runs the transform
/// @returns the new program or SkipTransform if the transform is not required
@@ -171,8 +163,6 @@
Hashmap<const sem::BuiltinFn*, Symbol, 8> builtin_polyfills;
/// Polyfill f32 conversion to i32 or u32 (or vectors of)
Hashmap<const core::type::Type*, Symbol, 2> f32_conv_polyfills;
- // Tracks whether the chromium_experimental_full_ptr_parameters extension has been enabled.
- bool has_full_ptr_params = false;
/// True if the transform has made changes (i.e. the program needs cloning)
bool made_changes = false;
@@ -823,10 +813,6 @@
/// @param type the type being loaded
/// @return the polyfill function name
Symbol workgroupUniformLoad(const core::type::Type* type) {
- if (!has_full_ptr_params) {
- b.Enable(wgsl::Extension::kChromiumExperimentalFullPtrParameters);
- has_full_ptr_params = true;
- }
auto name = b.Symbols().New("tint_workgroupUniformLoad");
b.Func(name,
tint::Vector{
diff --git a/src/tint/lang/wgsl/ast/transform/builtin_polyfill_test.cc b/src/tint/lang/wgsl/ast/transform/builtin_polyfill_test.cc
index 4366226..0c64163 100644
--- a/src/tint/lang/wgsl/ast/transform/builtin_polyfill_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/builtin_polyfill_test.cc
@@ -3774,8 +3774,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn tint_workgroupUniformLoad(p : ptr<workgroup, i32>) -> i32 {
workgroupBarrier();
let result = *(p);
@@ -3815,8 +3813,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn tint_workgroupUniformLoad(p : ptr<workgroup, Outer>) -> Outer {
workgroupBarrier();
let result = *(p);
@@ -3848,8 +3844,6 @@
TEST_F(BuiltinPolyfillTest, WorkgroupUniformLoad_AvoidDuplicateEnables) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> a : i32;
var<workgroup> b : u32;
var<workgroup> c : f32;
@@ -3862,8 +3856,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn tint_workgroupUniformLoad(p : ptr<workgroup, i32>) -> i32 {
workgroupBarrier();
let result = *(p);
@@ -3915,8 +3907,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn tint_workgroupUniformLoad_v() -> i32 {
workgroupBarrier();
let result = v;
diff --git a/src/tint/lang/wgsl/ast/transform/direct_variable_access.cc b/src/tint/lang/wgsl/ast/transform/direct_variable_access.cc
index c73dead..4ecea23 100644
--- a/src/tint/lang/wgsl/ast/transform/direct_variable_access.cc
+++ b/src/tint/lang/wgsl/ast/transform/direct_variable_access.cc
@@ -195,10 +195,8 @@
/// The main function for the transform.
/// @returns the ApplyResult
ApplyResult Run() {
- if (!ctx.src->Sem().Module()->Extensions().Contains(
- wgsl::Extension::kChromiumExperimentalFullPtrParameters)) {
- // If the 'chromium_experimental_full_ptr_parameters' extension is not enabled, then
- // there's nothing for this transform to do.
+ // If there are no functions with pointer parameters, then this transform can be skipped.
+ if (!AnyPointerParameters()) {
return SkipTransform;
}
@@ -409,6 +407,18 @@
/// Only valid during the lifetime of the program::CloneContext::Clone().
CloneState* clone_state = nullptr;
+ /// @returns true if any user functions have parameters of a pointer type.
+ bool AnyPointerParameters() const {
+ for (auto* fn : ctx.src->AST().Functions()) {
+ for (auto* param : fn->params) {
+ if (sem.Get(param)->Type()->Is<core::type::Pointer>()) {
+ return true;
+ }
+ }
+ }
+ return false;
+ }
+
/// AppendAccessChain creates or extends an existing AccessChain for the given expression,
/// modifying the #access_chains map.
void AppendAccessChain(const sem::ValueExpression* expr) {
diff --git a/src/tint/lang/wgsl/ast/transform/direct_variable_access.h b/src/tint/lang/wgsl/ast/transform/direct_variable_access.h
index b47f7e3..d1bb9fe 100644
--- a/src/tint/lang/wgsl/ast/transform/direct_variable_access.h
+++ b/src/tint/lang/wgsl/ast/transform/direct_variable_access.h
@@ -34,8 +34,7 @@
/// DirectVariableAccess is a transform that allows usage of pointer parameters in the 'storage',
/// 'uniform' and 'workgroup' address space, and passing of pointers to sub-objects. These pointers
-/// are only allowed by the resolver when the `chromium_experimental_full_ptr_parameters` extension
-/// is enabled.
+/// are allowed with the `unrestricted_pointer_parameters` WGSL feature.
///
/// DirectVariableAccess works by creating specializations of functions that have pointer
/// parameters, one specialization for each pointer argument's unique access chain 'shape' from a
diff --git a/src/tint/lang/wgsl/ast/transform/direct_variable_access_test.cc b/src/tint/lang/wgsl/ast/transform/direct_variable_access_test.cc
index 10f6421..04666b2 100644
--- a/src/tint/lang/wgsl/ast/transform/direct_variable_access_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/direct_variable_access_test.cc
@@ -80,8 +80,6 @@
TEST_F(DirectVariableAccessRemoveUncalledTest, PtrUniform) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> keep_me = 42;
fn u(pre : i32, p : ptr<uniform, i32>, post : i32) -> i32 {
@@ -91,8 +89,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> keep_me = 42;
)";
@@ -103,8 +99,6 @@
TEST_F(DirectVariableAccessRemoveUncalledTest, PtrStorage) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> keep_me = 42;
fn s(pre : i32, p : ptr<storage, i32>, post : i32) -> i32 {
@@ -113,8 +107,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> keep_me = 42;
)";
@@ -125,8 +117,6 @@
TEST_F(DirectVariableAccessRemoveUncalledTest, PtrWorkgroup) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> keep_me = 42;
fn w(pre : i32, p : ptr<workgroup, i32>, post : i32) -> i32 {
@@ -136,8 +126,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> keep_me = 42;
)";
@@ -148,8 +136,6 @@
TEST_F(DirectVariableAccessRemoveUncalledTest, PtrPrivate_Disabled) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> keep_me = 42;
fn f(pre : i32, p : ptr<private, i32>, post : i32) -> i32 {
@@ -166,8 +152,6 @@
TEST_F(DirectVariableAccessRemoveUncalledTest, PtrPrivate_Enabled) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> keep_me = 42;
)";
@@ -180,8 +164,6 @@
TEST_F(DirectVariableAccessRemoveUncalledTest, PtrFunction_Disabled) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> keep_me = 42;
fn f(pre : i32, p : ptr<function, i32>, post : i32) -> i32 {
@@ -198,8 +180,6 @@
TEST_F(DirectVariableAccessRemoveUncalledTest, PtrFunction_Enabled) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> keep_me = 42;
)";
@@ -221,8 +201,6 @@
TEST_F(DirectVariableAccessPtrChainsTest, ConstantIndices) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 8>, 8>, 8>;
fn a(pre : i32, p : ptr<uniform, vec4<i32>>, post : i32) -> vec4<i32> {
@@ -251,8 +229,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 8>, 8>, 8>;
alias U_X_X_X = array<u32, 3u>;
@@ -289,8 +265,6 @@
TEST_F(DirectVariableAccessPtrChainsTest, ConstantIndices_ViaPointerIndex) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 8>, 8>, 8>;
fn a(pre : i32, p : ptr<uniform, vec4<i32>>, post : i32) -> vec4<i32> {
@@ -319,8 +293,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 8>, 8>, 8>;
alias U_X_X_X = array<u32, 3u>;
@@ -357,8 +329,6 @@
TEST_F(DirectVariableAccessPtrChainsTest, HoistIndices) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 8>, 8>, 8>;
var<private> i : i32;
@@ -399,8 +369,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 8>, 8>, 8>;
var<private> i : i32;
@@ -458,8 +426,6 @@
TEST_F(DirectVariableAccessPtrChainsTest, HoistInForLoopInit) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8>, 8>;
var<private> i : i32;
@@ -494,8 +460,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8>, 8>;
var<private> i : i32;
@@ -558,8 +522,6 @@
TEST_F(DirectVariableAccessPtrChainsTest, HoistInForLoopCond) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8>, 8>;
var<private> i : i32;
@@ -596,8 +558,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8>, 8>;
var<private> i : i32;
@@ -648,8 +608,6 @@
TEST_F(DirectVariableAccessPtrChainsTest, HoistInForLoopCont) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8>, 8>;
var<private> i : i32;
@@ -686,8 +644,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8>, 8>;
var<private> i : i32;
@@ -738,8 +694,6 @@
TEST_F(DirectVariableAccessPtrChainsTest, HoistInWhileCond) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8>, 8>;
var<private> i : i32;
@@ -776,8 +730,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8>, 8>;
var<private> i : i32;
@@ -837,8 +789,6 @@
TEST_F(DirectVariableAccessUniformASTest, Param_ptr_i32_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : i32;
fn a(pre : i32, p : ptr<uniform, i32>, post : i32) -> i32 {
@@ -851,8 +801,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : i32;
fn a_U(pre : i32, post : i32) -> i32 {
@@ -871,8 +819,6 @@
TEST_F(DirectVariableAccessUniformASTest, Param_ptr_vec4i32_Via_array_DynamicRead) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<vec4<i32>, 8>;
fn a(pre : i32, p : ptr<uniform, vec4<i32>>, post : i32) -> vec4<i32> {
@@ -886,8 +832,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<vec4<i32>, 8>;
alias U_X = array<u32, 1u>;
@@ -909,8 +853,6 @@
TEST_F(DirectVariableAccessUniformASTest, Param_ptr_vec4i32_Via_array_DynamicRead_ViaPointerDot) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<vec4<i32>, 8>;
fn a(pre : i32, p : ptr<uniform, vec4<i32>>, post : i32) -> vec4<i32> {
@@ -925,8 +867,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> U : array<vec4<i32>, 8>;
alias U_X = array<u32, 1u>;
@@ -949,8 +889,6 @@
TEST_F(DirectVariableAccessUniformASTest, CallChaining) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
};
@@ -1011,8 +949,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
}
@@ -1106,8 +1042,6 @@
TEST_F(DirectVariableAccessUniformASTest, CallChaining_ViaPointerDotOrIndex) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
};
@@ -1168,8 +1102,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
}
@@ -1263,8 +1195,6 @@
TEST_F(DirectVariableAccessUniformASTest, CallChaining2) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
alias T3 = vec4i;
alias T2 = array<T3, 5>;
alias T1 = array<T2, 5>;
@@ -1292,8 +1222,6 @@
auto* expect =
R"(
-enable chromium_experimental_full_ptr_parameters;
-
alias T3 = vec4i;
alias T2 = array<T3, 5>;
@@ -1342,8 +1270,6 @@
TEST_F(DirectVariableAccessStorageASTest, Param_ptr_i32_Via_struct_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
@@ -1360,8 +1286,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
@@ -1384,8 +1308,6 @@
TEST_F(DirectVariableAccessStorageASTest, Param_ptr_arr_i32_Via_struct_write) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
};
@@ -1402,8 +1324,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
}
@@ -1426,8 +1346,6 @@
TEST_F(DirectVariableAccessStorageASTest, Param_ptr_vec4i32_Via_array_DynamicWrite) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> S : array<vec4<i32>, 8>;
fn a(pre : i32, p : ptr<storage, vec4<i32>, read_write>, post : i32) {
@@ -1441,8 +1359,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> S : array<vec4<i32>, 8>;
alias S_X = array<u32, 1u>;
@@ -1464,8 +1380,6 @@
TEST_F(DirectVariableAccessStorageASTest, CallChaining) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
};
@@ -1526,8 +1440,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
}
@@ -1621,8 +1533,6 @@
TEST_F(DirectVariableAccessStorageASTest, CallChaining_ViaPointerDotOrIndex) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
};
@@ -1683,8 +1593,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
}
@@ -1778,8 +1686,6 @@
TEST_F(DirectVariableAccessStorageASTest, CallChaining2) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
alias T3 = vec4i;
alias T2 = array<T3, 5>;
alias T1 = array<T2, 5>;
@@ -1807,8 +1713,6 @@
auto* expect =
R"(
-enable chromium_experimental_full_ptr_parameters;
-
alias T3 = vec4i;
alias T2 = array<T3, 5>;
@@ -1857,8 +1761,6 @@
TEST_F(DirectVariableAccessWorkgroupASTest, Param_ptr_vec4i32_Via_array_StaticRead) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> W : array<vec4<i32>, 8>;
fn a(pre : i32, p : ptr<workgroup, vec4<i32>>, post : i32) -> vec4<i32> {
@@ -1871,8 +1773,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> W : array<vec4<i32>, 8>;
alias W_X = array<u32, 1u>;
@@ -1893,8 +1793,6 @@
TEST_F(DirectVariableAccessWorkgroupASTest, Param_ptr_vec4i32_Via_array_StaticWrite) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> W : array<vec4<i32>, 8>;
fn a(pre : i32, p : ptr<workgroup, vec4<i32>>, post : i32) {
@@ -1907,8 +1805,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> W : array<vec4<i32>, 8>;
alias W_X = array<u32, 1u>;
@@ -1929,8 +1825,6 @@
TEST_F(DirectVariableAccessWorkgroupASTest, CallChaining) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
};
@@ -1991,8 +1885,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
}
@@ -2086,8 +1978,6 @@
TEST_F(DirectVariableAccessWorkgroupASTest, CallChaining_ViaPointerDotOrIndex) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
};
@@ -2148,8 +2038,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
}
@@ -2243,8 +2131,6 @@
TEST_F(DirectVariableAccessWorkgroupASTest, CallChaining2) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
alias T3 = vec4i;
alias T2 = array<T3, 5>;
alias T1 = array<T2, 5>;
@@ -2274,8 +2160,6 @@
auto* expect =
R"(
-enable chromium_experimental_full_ptr_parameters;
-
alias T3 = vec4i;
alias T2 = array<T3, 5>;
@@ -2327,8 +2211,6 @@
TEST_F(DirectVariableAccessPrivateASTest, Enabled_Param_ptr_i32_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn a(pre : i32, p : ptr<private, i32>, post : i32) -> i32 {
return *(p);
}
@@ -2341,8 +2223,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn a_F(pre : i32, p : ptr<private, i32>, post : i32) -> i32 {
return *(p);
}
@@ -2361,8 +2241,6 @@
TEST_F(DirectVariableAccessPrivateASTest, Enabled_Param_ptr_i32_write) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn a(pre : i32, p : ptr<private, i32>, post : i32) {
*(p) = 42;
}
@@ -2375,8 +2253,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn a_F(pre : i32, p : ptr<private, i32>, post : i32) {
*(p) = 42;
}
@@ -2395,8 +2271,6 @@
TEST_F(DirectVariableAccessPrivateASTest, Enabled_Param_ptr_i32_Via_struct_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
@@ -2413,8 +2287,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
@@ -2437,8 +2309,6 @@
TEST_F(DirectVariableAccessPrivateASTest, Disabled_Param_ptr_i32_Via_struct_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
@@ -2463,8 +2333,6 @@
TEST_F(DirectVariableAccessPrivateASTest, Enabled_Param_ptr_arr_i32_Via_struct_write) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
};
@@ -2481,8 +2349,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
}
@@ -2505,8 +2371,6 @@
TEST_F(DirectVariableAccessPrivateASTest, Disabled_Param_ptr_arr_i32_Via_struct_write) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
}
@@ -2531,8 +2395,6 @@
TEST_F(DirectVariableAccessPrivateASTest, Enabled_Param_ptr_i32_mixed) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
@@ -2553,8 +2415,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
@@ -2595,8 +2455,6 @@
TEST_F(DirectVariableAccessPrivateASTest, Disabled_Param_ptr_i32_mixed) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
@@ -2627,8 +2485,6 @@
TEST_F(DirectVariableAccessPrivateASTest, Enabled_CallChaining) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
};
@@ -2689,8 +2545,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
}
@@ -2794,8 +2648,6 @@
TEST_F(DirectVariableAccessPrivateASTest, Enabled_CallChaining2) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
alias T3 = vec4i;
alias T2 = array<T3, 5>;
alias T1 = array<T2, 5>;
@@ -2823,8 +2675,6 @@
auto* expect =
R"(
-enable chromium_experimental_full_ptr_parameters;
-
alias T3 = vec4i;
alias T2 = array<T3, 5>;
@@ -2868,8 +2718,6 @@
TEST_F(DirectVariableAccessPrivateASTest, Disabled_CallChaining) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
}
@@ -2943,8 +2791,6 @@
TEST_F(DirectVariableAccessFunctionASTest, Enabled_LocalPtr) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn f() {
var v : i32;
let p : ptr<function, i32> = &(v);
@@ -2961,8 +2807,6 @@
TEST_F(DirectVariableAccessFunctionASTest, Enabled_Param_ptr_i32_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn a(pre : i32, p : ptr<function, i32>, post : i32) -> i32 {
return *(p);
}
@@ -2974,8 +2818,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn a_F(pre : i32, p : ptr<function, i32>, post : i32) -> i32 {
return *(p);
}
@@ -2993,8 +2835,6 @@
TEST_F(DirectVariableAccessFunctionASTest, Enabled_Param_ptr_i32_write) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn a(pre : i32, p : ptr<function, i32>, post : i32) {
*(p) = 42;
}
@@ -3006,8 +2846,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn a_F(pre : i32, p : ptr<function, i32>, post : i32) {
*(p) = 42;
}
@@ -3025,8 +2863,6 @@
TEST_F(DirectVariableAccessFunctionASTest, Enabled_Param_ptr_i32_Via_struct_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
@@ -3042,8 +2878,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
@@ -3065,8 +2899,6 @@
TEST_F(DirectVariableAccessFunctionASTest, Enabled_Param_ptr_arr_i32_Via_struct_write) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
};
@@ -3082,8 +2914,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
}
@@ -3105,8 +2935,6 @@
TEST_F(DirectVariableAccessFunctionASTest, Enabled_Param_ptr_i32_mixed) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
@@ -3127,8 +2955,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
@@ -3166,8 +2992,6 @@
TEST_F(DirectVariableAccessFunctionASTest, Enabled_CallChaining) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
};
@@ -3219,8 +3043,6 @@
auto* expect =
R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct Inner {
mat : mat3x4<f32>,
}
@@ -3275,8 +3097,6 @@
TEST_F(DirectVariableAccessFunctionASTest, Enabled_CallChaining2) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
alias T3 = vec4i;
alias T2 = array<T3, 5>;
alias T1 = array<T2, 5>;
@@ -3303,8 +3123,6 @@
auto* expect =
R"(
-enable chromium_experimental_full_ptr_parameters;
-
alias T3 = vec4i;
alias T2 = array<T3, 5>;
@@ -3347,8 +3165,6 @@
TEST_F(DirectVariableAccessFunctionASTest, Disabled_Param_ptr_i32_Via_struct_read) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
@@ -3372,8 +3188,6 @@
TEST_F(DirectVariableAccessFunctionASTest, Disabled_Param_ptr_arr_i32_Via_struct_write) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
}
@@ -3406,8 +3220,6 @@
TEST_F(DirectVariableAccessComplexTest, Param_ptr_mixed_vec4i32_ViaMultiple) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<i32>,
};
@@ -3476,8 +3288,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<i32>,
}
@@ -3606,8 +3416,6 @@
TEST_F(DirectVariableAccessComplexTest, Indexing) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : array<array<array<array<i32, 9>, 9>, 9>, 50>;
fn a(i : i32) -> i32 { return i; }
@@ -3624,8 +3432,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : array<array<array<array<i32, 9>, 9>, 9>, 50>;
fn a(i : i32) -> i32 {
@@ -3650,8 +3456,6 @@
TEST_F(DirectVariableAccessComplexTest, IndexingInPtrCall) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : array<array<array<array<i32, 9>, 9>, 9>, 50>;
fn a(pre : i32, i : ptr<storage, i32>, post : i32) -> i32 {
@@ -3670,8 +3474,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : array<array<array<array<i32, 9>, 9>, 9>, 50>;
alias S_X_X_X_X = array<u32, 4u>;
@@ -3698,8 +3500,6 @@
TEST_F(DirectVariableAccessComplexTest, IndexingDualPointers) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : array<array<array<i32, 9>, 9>, 50>;
@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 9>, 9>, 50>;
@@ -3717,8 +3517,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : array<array<array<i32, 9>, 9>, 50>;
@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 9>, 9>, 50>;
diff --git a/src/tint/lang/wgsl/ast/transform/preserve_padding.cc b/src/tint/lang/wgsl/ast/transform/preserve_padding.cc
index 39ebb2d..d6935a9 100644
--- a/src/tint/lang/wgsl/ast/transform/preserve_padding.cc
+++ b/src/tint/lang/wgsl/ast/transform/preserve_padding.cc
@@ -62,32 +62,25 @@
// Gather a list of assignments that need to be transformed.
std::unordered_set<const AssignmentStatement*> assignments_to_transform;
for (auto* node : ctx.src->ASTNodes().Objects()) {
- Switch(
- node, //
- [&](const AssignmentStatement* assign) {
- auto* ty = sem.GetVal(assign->lhs)->Type();
- if (assign->lhs->Is<PhonyExpression>()) {
- // Ignore phony assignment.
- return;
- }
- if (ty->As<core::type::Reference>()->AddressSpace() !=
- core::AddressSpace::kStorage) {
- // We only care about assignments that write to variables in the storage
- // address space, as nothing else is host-visible.
- return;
- }
- if (HasPadding(ty->UnwrapRef())) {
- // The assigned type has padding bytes, so we need to decompose the writes.
- assignments_to_transform.insert(assign);
- }
- },
- [&](const Enable* enable) {
- // Check if the full pointer parameters extension is already enabled.
- if (enable->HasExtension(
- wgsl::Extension::kChromiumExperimentalFullPtrParameters)) {
- ext_enabled = true;
- }
- });
+ Switch(node, //
+ [&](const AssignmentStatement* assign) {
+ auto* ty = sem.GetVal(assign->lhs)->Type();
+ if (assign->lhs->Is<PhonyExpression>()) {
+ // Ignore phony assignment.
+ return;
+ }
+ if (ty->As<core::type::Reference>()->AddressSpace() !=
+ core::AddressSpace::kStorage) {
+ // We only care about assignments that write to variables in the storage
+ // address space, as nothing else is host-visible.
+ return;
+ }
+ if (HasPadding(ty->UnwrapRef())) {
+ // The assigned type has padding bytes, so we need to decompose the
+ // writes.
+ assignments_to_transform.insert(assign);
+ }
+ });
}
if (assignments_to_transform.empty()) {
return SkipTransform;
@@ -127,13 +120,9 @@
// }
// It will be called by passing a pointer to the original LHS:
// assign_helper_T(&lhs, rhs);
- //
- // Since this requires passing pointers to the storage address space, this will also enable
- // the chromium_experimental_full_ptr_parameters extension.
const char* kDestParamName = "dest";
const char* kValueParamName = "value";
auto call_helper = [&](auto&& body) {
- EnableExtension();
auto helper = helpers.GetOrCreate(ty, [&] {
auto helper_name = b.Symbols().New("assign_and_preserve_padding");
tint::Vector<const Parameter*, 2> params = {
@@ -227,14 +216,6 @@
[&](Default) { return false; });
}
- /// Enable the full pointer parameters extension, if we have not already done so.
- void EnableExtension() {
- if (!ext_enabled) {
- b.Enable(wgsl::Extension::kChromiumExperimentalFullPtrParameters);
- ext_enabled = true;
- }
- }
-
private:
/// The program builder
ProgramBuilder b;
@@ -244,8 +225,6 @@
const sem::Info& sem = ctx.src->Sem();
/// Alias to the symbols in ctx.src
const SymbolTable& sym = ctx.src->Symbols();
- /// Flag to track whether we have already enabled the full pointer parameters extension.
- bool ext_enabled = false;
/// Map of semantic types to their assignment helper functions.
Hashmap<const core::type::Type*, Symbol, 8> helpers;
};
diff --git a/src/tint/lang/wgsl/ast/transform/preserve_padding_test.cc b/src/tint/lang/wgsl/ast/transform/preserve_padding_test.cc
index 469448b..be548fc 100644
--- a/src/tint/lang/wgsl/ast/transform/preserve_padding_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/preserve_padding_test.cc
@@ -119,8 +119,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
a : u32,
b : u32,
@@ -175,8 +173,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
a : u32,
b : u32,
@@ -227,8 +223,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
a : u32,
b : vec4<u32>,
@@ -267,8 +261,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
@size(16)
a : u32,
@@ -307,8 +299,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
@size(16)
a : u32,
@@ -362,8 +352,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S1 {
a1 : u32,
b1 : vec3<u32>,
@@ -424,8 +412,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> v : array<vec3<u32>, 4>;
fn assign_and_preserve_padding(dest : ptr<storage, array<vec3<u32>, 4u>, read_write>, value : array<vec3<u32>, 4u>) {
@@ -458,8 +444,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
alias Array = array<array<vec3<u32>, 4>, 3>;
@group(0) @binding(0) var<storage, read_write> v : Array;
@@ -503,8 +487,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
a : u32,
b : array<vec3<u32>, 4>,
@@ -551,8 +533,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> m : mat3x3<f32>;
fn assign_and_preserve_padding(dest : ptr<storage, mat3x3<f32>, read_write>, value : mat3x3<f32>) {
@@ -588,8 +568,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
struct S {
a : u32,
m : mat3x3<f32>,
@@ -631,8 +609,6 @@
)";
auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> arr_m : array<mat3x3<f32>, 4>;
fn assign_and_preserve_padding_1(dest : ptr<storage, mat3x3<f32>, read_write>, value : mat3x3<f32>) {
@@ -676,47 +652,6 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(PreservePaddingTest, AvoidDuplicateEnables) {
- auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
-struct S {
- @size(16) a : u32,
-}
-
-@group(0) @binding(0) var<storage, read_write> v : S;
-
-@compute @workgroup_size(1)
-fn foo() {
- v = S();
-}
-)";
-
- auto* expect = R"(
-enable chromium_experimental_full_ptr_parameters;
-
-struct S {
- @size(16)
- a : u32,
-}
-
-@group(0) @binding(0) var<storage, read_write> v : S;
-
-fn assign_and_preserve_padding(dest : ptr<storage, S, read_write>, value : S) {
- (*(dest)).a = value.a;
-}
-
-@compute @workgroup_size(1)
-fn foo() {
- assign_and_preserve_padding(&(v), S());
-}
-)";
-
- auto got = Run<PreservePadding>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
TEST_F(PreservePaddingTest, NoModify_StructNoPadding) {
auto* src = R"(
struct S {
diff --git a/src/tint/lang/wgsl/ast/transform/robustness_test.cc b/src/tint/lang/wgsl/ast/transform/robustness_test.cc
index 647e669..1cd715f 100644
--- a/src/tint/lang/wgsl/ast/transform/robustness_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/robustness_test.cc
@@ -4614,8 +4614,6 @@
TEST_P(RobustnessTest, Read_PrivatePointerParameter_IndexWithConstant) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> a : array<i32, 4>;
fn x(pre : i32, p : ptr<private, i32>, post : i32) -> i32 {
@@ -4635,8 +4633,6 @@
/* ignore */ src,
/* clamp */ src,
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> a : array<i32, 4>;
fn x(pre : i32, p : ptr<private, i32>, p_predicate : bool, post : i32) -> i32 {
@@ -4663,8 +4659,6 @@
TEST_P(RobustnessTest, Read_WorkgroupPointerParameter_IndexWithConstant) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> a : array<i32, 4>;
fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) -> i32 {
@@ -4684,8 +4678,6 @@
/* ignore */ src,
/* clamp */ src,
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> a : array<i32, 4>;
fn x(pre : i32, p : ptr<workgroup, i32>, p_predicate : bool, post : i32) -> i32 {
@@ -4712,8 +4704,6 @@
TEST_P(RobustnessTest, Read_UniformPointerParameter_IndexWithConstant) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> a : array<vec4i, 4>;
fn x(pre : vec4i, p : ptr<uniform, vec4i>, post : vec4i) -> vec4i {
@@ -4733,8 +4723,6 @@
/* ignore */ src,
/* clamp */ src,
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> a : array<vec4i, 4>;
fn x(pre : vec4i, p : ptr<uniform, vec4i>, p_predicate : bool, post : vec4i) -> vec4i {
@@ -4761,8 +4749,6 @@
TEST_P(RobustnessTest, Read_StoragePointerParameter_IndexWithConstant) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> a : array<vec4i, 4>;
fn x(pre : vec4i, p : ptr<storage, vec4i>, post : vec4i) -> vec4i {
@@ -4782,8 +4768,6 @@
/* ignore */ src,
/* clamp */ src,
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> a : array<vec4i, 4>;
fn x(pre : vec4i, p : ptr<storage, vec4i>, p_predicate : bool, post : vec4i) -> vec4i {
@@ -4810,8 +4794,6 @@
TEST_P(RobustnessTest, Read_FunctionPointerParameter_IndexWithConstant) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn x(pre : i32, p : ptr<function, i32>, post : i32) -> i32 {
return ((pre + *(p)) + post);
}
@@ -4830,8 +4812,6 @@
/* ignore */ src,
/* clamp */ src,
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn x(pre : i32, p : ptr<function, i32>, p_predicate : bool, post : i32) -> i32 {
var predicated_expr : i32;
if (p_predicate) {
@@ -4857,8 +4837,6 @@
TEST_P(RobustnessTest, Read_PrivatePointerParameter_IndexWithLet) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> a : array<i32, 4>;
fn x(pre : i32, p : ptr<private, i32>, post : i32) -> i32 {
@@ -4878,8 +4856,6 @@
auto* expect = Expect(GetParam(),
/* ignore */ src,
/* clamp */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> a : array<i32, 4>;
fn x(pre : i32, p : ptr<private, i32>, post : i32) -> i32 {
@@ -4896,8 +4872,6 @@
}
)",
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> a : array<i32, 4>;
fn x(pre : i32, p : ptr<private, i32>, p_predicate : bool, post : i32) -> i32 {
@@ -4927,8 +4901,6 @@
TEST_P(RobustnessTest, Read_WorkgroupPointerParameter_IndexWithLet) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> a : array<i32, 4>;
fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) -> i32 {
@@ -4948,8 +4920,6 @@
auto* expect = Expect(GetParam(),
/* ignore */ src,
/* clamp */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> a : array<i32, 4>;
fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) -> i32 {
@@ -4966,8 +4936,6 @@
}
)",
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> a : array<i32, 4>;
fn x(pre : i32, p : ptr<workgroup, i32>, p_predicate : bool, post : i32) -> i32 {
@@ -4997,8 +4965,6 @@
TEST_P(RobustnessTest, Read_UniformPointerParameter_IndexWithLet) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> a : array<vec4i, 4>;
fn x(pre : vec4i, p : ptr<uniform, vec4i>, post : vec4i) -> vec4i {
@@ -5018,8 +4984,6 @@
auto* expect = Expect(GetParam(),
/* ignore */ src,
/* clamp */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> a : array<vec4i, 4>;
fn x(pre : vec4i, p : ptr<uniform, vec4i>, post : vec4i) -> vec4i {
@@ -5036,8 +5000,6 @@
}
)",
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> a : array<vec4i, 4>;
fn x(pre : vec4i, p : ptr<uniform, vec4i>, p_predicate : bool, post : vec4i) -> vec4i {
@@ -5067,8 +5029,6 @@
TEST_P(RobustnessTest, Read_StoragePointerParameter_IndexWithLet) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> a : array<vec4i, 4>;
fn x(pre : vec4i, p : ptr<storage, vec4i>, post : vec4i) -> vec4i {
@@ -5088,8 +5048,6 @@
auto* expect = Expect(GetParam(),
/* ignore */ src,
/* clamp */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> a : array<vec4i, 4>;
fn x(pre : vec4i, p : ptr<storage, vec4i>, post : vec4i) -> vec4i {
@@ -5106,8 +5064,6 @@
}
)",
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> a : array<vec4i, 4>;
fn x(pre : vec4i, p : ptr<storage, vec4i>, p_predicate : bool, post : vec4i) -> vec4i {
@@ -5137,8 +5093,6 @@
TEST_P(RobustnessTest, Read_FunctionPointerParameter_IndexWithLet) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn x(pre : i32, p : ptr<function, i32>, post : i32) -> i32 {
return ((pre + *(p)) + post);
}
@@ -5157,8 +5111,6 @@
auto* expect = Expect(GetParam(),
/* ignore */ src,
/* clamp */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn x(pre : i32, p : ptr<function, i32>, post : i32) -> i32 {
return ((pre + *(p)) + post);
}
@@ -5174,8 +5126,6 @@
}
)",
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn x(pre : i32, p : ptr<function, i32>, p_predicate : bool, post : i32) -> i32 {
var predicated_expr : i32;
if (p_predicate) {
@@ -5204,8 +5154,6 @@
TEST_P(RobustnessTest, Write_PrivatePointerParameter_IndexWithConstant) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> a : array<i32, 4>;
fn x(pre : i32, p : ptr<private, i32>, post : i32) {
@@ -5225,8 +5173,6 @@
/* ignore */ src,
/* clamp */ src,
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> a : array<i32, 4>;
fn x(pre : i32, p : ptr<private, i32>, p_predicate : bool, post : i32) {
@@ -5251,8 +5197,6 @@
TEST_P(RobustnessTest, Write_WorkgroupPointerParameter_IndexWithConstant) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> a : array<i32, 4>;
fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) {
@@ -5272,8 +5216,6 @@
/* ignore */ src,
/* clamp */ src,
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> a : array<i32, 4>;
fn x(pre : i32, p : ptr<workgroup, i32>, p_predicate : bool, post : i32) {
@@ -5298,8 +5240,6 @@
TEST_P(RobustnessTest, Write_StoragePointerParameter_IndexWithConstant) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> a : array<i32, 4>;
fn x(pre : i32, p : ptr<storage, i32, read_write>, post : i32) {
@@ -5319,8 +5259,6 @@
/* ignore */ src,
/* clamp */ src,
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> a : array<i32, 4>;
fn x(pre : i32, p : ptr<storage, i32, read_write>, p_predicate : bool, post : i32) {
@@ -5345,8 +5283,6 @@
TEST_P(RobustnessTest, Write_FunctionPointerParameter_IndexWithConstant) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn x(pre : i32, p : ptr<function, i32>, post : i32) {
*(p) = (pre + post);
}
@@ -5365,8 +5301,6 @@
/* ignore */ src,
/* clamp */ src,
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn x(pre : i32, p : ptr<function, i32>, p_predicate : bool, post : i32) {
if (p_predicate) {
*(p) = (pre + post);
@@ -5390,8 +5324,6 @@
TEST_P(RobustnessTest, Write_PrivatePointerParameter_IndexWithLet) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> a : array<i32, 4>;
fn x(pre : i32, p : ptr<private, i32>, post : i32) {
@@ -5411,8 +5343,6 @@
auto* expect = Expect(GetParam(),
/* ignore */ src,
/* clamp */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> a : array<i32, 4>;
fn x(pre : i32, p : ptr<private, i32>, post : i32) {
@@ -5429,8 +5359,6 @@
}
)",
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> a : array<i32, 4>;
fn x(pre : i32, p : ptr<private, i32>, p_predicate : bool, post : i32) {
@@ -5458,8 +5386,6 @@
TEST_P(RobustnessTest, Write_WorkgroupPointerParameter_IndexWithLet) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> a : array<i32, 4>;
fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) {
@@ -5479,8 +5405,6 @@
auto* expect = Expect(GetParam(),
/* ignore */ src,
/* clamp */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> a : array<i32, 4>;
fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) {
@@ -5497,8 +5421,6 @@
}
)",
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> a : array<i32, 4>;
fn x(pre : i32, p : ptr<workgroup, i32>, p_predicate : bool, post : i32) {
@@ -5526,8 +5448,6 @@
TEST_P(RobustnessTest, Write_StoragePointerParameter_IndexWithLet) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> a : array<i32, 4>;
fn x(pre : i32, p : ptr<storage, i32, read_write>, post : i32) {
@@ -5547,8 +5467,6 @@
auto* expect = Expect(GetParam(),
/* ignore */ src,
/* clamp */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> a : array<i32, 4>;
fn x(pre : i32, p : ptr<storage, i32, read_write>, post : i32) {
@@ -5565,8 +5483,6 @@
}
)",
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> a : array<i32, 4>;
fn x(pre : i32, p : ptr<storage, i32, read_write>, p_predicate : bool, post : i32) {
@@ -5594,8 +5510,6 @@
TEST_P(RobustnessTest, Write_FunctionPointerParameter_IndexWithLet) {
auto* src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn x(pre : i32, p : ptr<function, i32>, post : i32) {
*(p) = (pre + post);
}
@@ -5614,8 +5528,6 @@
auto* expect = Expect(GetParam(),
/* ignore */ src,
/* clamp */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn x(pre : i32, p : ptr<function, i32>, post : i32) {
*(p) = (pre + post);
}
@@ -5631,8 +5543,6 @@
}
)",
/* predicate */ R"(
-enable chromium_experimental_full_ptr_parameters;
-
fn x(pre : i32, p : ptr<function, i32>, p_predicate : bool, post : i32) {
if (p_predicate) {
*(p) = (pre + post);
diff --git a/src/tint/lang/wgsl/extension.cc b/src/tint/lang/wgsl/extension.cc
index c2f1762..3069ccf 100644
--- a/src/tint/lang/wgsl/extension.cc
+++ b/src/tint/lang/wgsl/extension.cc
@@ -48,9 +48,6 @@
if (str == "chromium_experimental_framebuffer_fetch") {
return Extension::kChromiumExperimentalFramebufferFetch;
}
- if (str == "chromium_experimental_full_ptr_parameters") {
- return Extension::kChromiumExperimentalFullPtrParameters;
- }
if (str == "chromium_experimental_pixel_local") {
return Extension::kChromiumExperimentalPixelLocal;
}
@@ -80,8 +77,6 @@
return "chromium_disable_uniformity_analysis";
case Extension::kChromiumExperimentalFramebufferFetch:
return "chromium_experimental_framebuffer_fetch";
- case Extension::kChromiumExperimentalFullPtrParameters:
- return "chromium_experimental_full_ptr_parameters";
case Extension::kChromiumExperimentalPixelLocal:
return "chromium_experimental_pixel_local";
case Extension::kChromiumExperimentalPushConstant:
diff --git a/src/tint/lang/wgsl/extension.h b/src/tint/lang/wgsl/extension.h
index 41edcba..7de7e82 100644
--- a/src/tint/lang/wgsl/extension.h
+++ b/src/tint/lang/wgsl/extension.h
@@ -48,7 +48,6 @@
kUndefined,
kChromiumDisableUniformityAnalysis,
kChromiumExperimentalFramebufferFetch,
- kChromiumExperimentalFullPtrParameters,
kChromiumExperimentalPixelLocal,
kChromiumExperimentalPushConstant,
kChromiumExperimentalSubgroups,
@@ -75,22 +74,16 @@
Extension ParseExtension(std::string_view str);
constexpr std::string_view kExtensionStrings[] = {
- "chromium_disable_uniformity_analysis",
- "chromium_experimental_framebuffer_fetch",
- "chromium_experimental_full_ptr_parameters",
- "chromium_experimental_pixel_local",
- "chromium_experimental_push_constant",
- "chromium_experimental_subgroups",
- "chromium_internal_dual_source_blending",
- "chromium_internal_relaxed_uniform_layout",
- "f16",
+ "chromium_disable_uniformity_analysis", "chromium_experimental_framebuffer_fetch",
+ "chromium_experimental_pixel_local", "chromium_experimental_push_constant",
+ "chromium_experimental_subgroups", "chromium_internal_dual_source_blending",
+ "chromium_internal_relaxed_uniform_layout", "f16",
};
/// All extensions
static constexpr Extension kAllExtensions[] = {
Extension::kChromiumDisableUniformityAnalysis,
Extension::kChromiumExperimentalFramebufferFetch,
- Extension::kChromiumExperimentalFullPtrParameters,
Extension::kChromiumExperimentalPixelLocal,
Extension::kChromiumExperimentalPushConstant,
Extension::kChromiumExperimentalSubgroups,
diff --git a/src/tint/lang/wgsl/extension_bench.cc b/src/tint/lang/wgsl/extension_bench.cc
index c2e6bd2..70c0d73 100644
--- a/src/tint/lang/wgsl/extension_bench.cc
+++ b/src/tint/lang/wgsl/extension_bench.cc
@@ -59,55 +59,48 @@
"chromium_experimental_vramebuffeii_fetch",
"chro8WWum_experimental_framebuffer_fetch",
"chromium_eperimenxxMl_framebuffer_fetch",
- "chromium_expeggimeXtal_full_ptr_paraeters",
- "chromium_expVrimental_full_ptr_puraXeer",
- "chromium_experimental_full_ptr3parameters",
- "chromium_experimental_full_ptr_parameters",
- "chromium_experimentalEfull_ptr_parameters",
- "chromium_experimentalfull_ptr_PPaTTameters",
- "chromium_ddxperimental_fullptrxxparameters",
- "chromium_experi44ental_pixel_local",
- "chromium_experimental_VVSixel_local",
- "chroRium_experimental_pix22Rlocal",
+ "chromum_experimental_pixeX_loggal",
+ "chromium_expVrXmntal_ixel_local",
+ "3hromium_experimental_pixel_local",
"chromium_experimental_pixel_local",
- "chromiuF_experiment9lpixel_local",
- "chromium_experimental_pixel_loca",
- "Vhromium_expeOOimentalHpixRRl_lcal",
- "chromiym_experimental_push_contant",
- "nnhro77ium_experimenGal_push_conrrllant",
- "chromium_experimental_push_c4nstan00",
+ "chromium_eEperimental_pixel_local",
+ "chTTomiu_experimentaPP_pixel_local",
+ "cxxromium_expddrimenal_pixel_local",
+ "c44romium_experimental_push_constant",
+ "chromium_experimental_pSSsVV_constant",
+ "chrom22Rm_experimental_pushRonstant",
"chromium_experimental_push_constant",
- "chooomum_experimental_ush_constat",
- "chromium_xperimntal_zzush_constant",
- "chromi11m_experimepptal_psh_ciistant",
- "chromium_experimental_subgroXXps",
- "chromium55eIIperimental_subgnno99ps",
- "chraamiuSS_experimentaHHr_subgrouYs",
+ "chromium_exp9rimFntal_ush_constant",
+ "chrmium_experimental_push_constant",
+ "cOOromium_experiVeHtal_puh_conRRtant",
+ "chromium_eperimental_sybgroups",
+ "chrorri77mGexperimllntal_subgrnnups",
+ "chromium_exp4rimen00al_subgroups",
"chromium_experimental_subgroups",
- "chkkomium_eperimntal_subgroup",
- "jhromium_experRmental_subgogps",
- "chromiubexperiental_subgroups",
- "chromium_internal_dujl_source_blending",
- "chromium_intenal_dual_source_blending",
- "chqomium_internal_dual_source_beding",
+ "chromium_exprimenal_ubgrouoos",
+ "chrmiumexperimenzzal_subgroups",
+ "chrmi11m_experppmeiita_subgroups",
+ "chromium_internal_dual_sourcXX_blending",
+ "chromium_internal_dual_99ou5IInce_blending",
+ "crrroSSium_internYlaadual_source_bHHending",
"chromium_internal_dual_source_blending",
- "chroium_NNnternal_dual_source_blending",
- "chovvium_internal_dual_source_lending",
- "chromium_internQQl_dual_sorce_blending",
- "chromirm_intenal_rfflaxed_unifrm_layout",
- "chromium_internal_jelaxed_uniform_layout",
- "chromium_interna_relNNxed_uwwiform_lay82t",
+ "chromium_internakk_ualsourc_blendHng",
+ "chromium_inRRrnal_dujl_sourceblgnding",
+ "chromiuminternal_duab_source_blendin",
+ "chromium_internal_relaxed_uniform_lajout",
+ "chromium_internal_relxed_uniform_layout",
+ "chroium_inqernal_rlaxed_uniform_layout",
"chromium_internal_relaxed_uniform_layout",
- "chromium_internal_relaxed_uniform_layut",
- "chromium_internal_relaxed_rrniform_layout",
- "chromium_internal_relaxedGuniform_layout",
- "FF16",
- "",
- "rr1",
+ "chromium_internNNl_relaxed_uniform_layou",
+ "chromium_internal_relaxvvd_unifom_laout",
+ "chromium_internalrelaxed_uniQQorm_layout",
+ "ff",
+ "fj6",
+ "wNN2",
"f16",
- "1",
- "DJ1",
- "",
+ "f6",
+ "rr16",
+ "fG6",
};
for (auto _ : state) {
for (auto* str : kStrings) {
diff --git a/src/tint/lang/wgsl/extension_test.cc b/src/tint/lang/wgsl/extension_test.cc
index 08cda9b..8bcbd0c 100644
--- a/src/tint/lang/wgsl/extension_test.cc
+++ b/src/tint/lang/wgsl/extension_test.cc
@@ -59,8 +59,6 @@
static constexpr Case kValidCases[] = {
{"chromium_disable_uniformity_analysis", Extension::kChromiumDisableUniformityAnalysis},
{"chromium_experimental_framebuffer_fetch", Extension::kChromiumExperimentalFramebufferFetch},
- {"chromium_experimental_full_ptr_parameters",
- Extension::kChromiumExperimentalFullPtrParameters},
{"chromium_experimental_pixel_local", Extension::kChromiumExperimentalPixelLocal},
{"chromium_experimental_push_constant", Extension::kChromiumExperimentalPushConstant},
{"chromium_experimental_subgroups", Extension::kChromiumExperimentalSubgroups},
@@ -76,27 +74,24 @@
{"chromium_experimental_framebuf1er_fetch", Extension::kUndefined},
{"chromium_experiqqntal_framebuffer_fetch", Extension::kUndefined},
{"chromium_experimental_framebuffll77_fetch", Extension::kUndefined},
- {"chroium_experimental_full_ptr_paqqppmetHHrs", Extension::kUndefined},
- {"chrium_evperiental_full_ptr_paraceters", Extension::kUndefined},
- {"chromium_expGimental_fullbptr_parameters", Extension::kUndefined},
- {"vhromium_experimental_pixel_liical", Extension::kUndefined},
- {"chromium_experiment8l_pixel_lWWcal", Extension::kUndefined},
- {"chromium_expeimentMl_xxixel_local", Extension::kUndefined},
- {"chrXmium_experimeggtal_ush_constant", Extension::kUndefined},
- {"chromiu_experVmentalpusX_constant", Extension::kUndefined},
- {"chro3ium_experimental_push_constant", Extension::kUndefined},
- {"cEromium_experimental_subgroups", Extension::kUndefined},
- {"TThromium_experiPPental_sugroups", Extension::kUndefined},
- {"chddomium_experimental_subgroxxs", Extension::kUndefined},
- {"chromium_internal_44ual_source_blending", Extension::kUndefined},
- {"chromium_inteSSnal_dual_source_blendinVV", Extension::kUndefined},
- {"chromiuR_interna22dual_source_blenRing", Extension::kUndefined},
- {"chromium_int9rnal_relaxed_Fnifor_layout", Extension::kUndefined},
- {"chrmium_internal_relaxed_uniform_layout", Extension::kUndefined},
- {"VRhHomium_internal_relaxd_uniform_OOayout", Extension::kUndefined},
- {"y1", Extension::kUndefined},
- {"l77rrn6", Extension::kUndefined},
- {"4016", Extension::kUndefined},
+ {"chroqqppum_expermental_pixel_HHocal", Extension::kUndefined},
+ {"chromium_experimenta_piel_occl", Extension::kUndefined},
+ {"chromium_exeGimental_pixel_local", Extension::kUndefined},
+ {"chvomium_experimental_push_constiint", Extension::kUndefined},
+ {"chromiu8WWexperimental_push_constant", Extension::kUndefined},
+ {"chromium_experiMental_push_costanxx", Extension::kUndefined},
+ {"chromiuX_experimentl_sugggroups", Extension::kUndefined},
+ {"chromiu_exuerimntal_XVbgroups", Extension::kUndefined},
+ {"chromium_experimen3al_subgroups", Extension::kUndefined},
+ {"chromium_internal_dual_soErce_blending", Extension::kUndefined},
+ {"chromiuPP_internal_dual_sourceblenTTing", Extension::kUndefined},
+ {"chromim_internadd_dual_sxxurce_blending", Extension::kUndefined},
+ {"chromium_interna44_relaxed_uniform_layout", Extension::kUndefined},
+ {"chromium_internal_relaxed_uniformSSlayouVV", Extension::kUndefined},
+ {"chromiumRnteRnal_re22axed_uniform_layout", Extension::kUndefined},
+ {"96", Extension::kUndefined},
+ {"f1", Extension::kUndefined},
+ {"VOR6", Extension::kUndefined},
};
using ExtensionParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/lang/wgsl/features/language_feature.cc b/src/tint/lang/wgsl/features/language_feature.cc
index 4561eb0..f0391dc 100644
--- a/src/tint/lang/wgsl/features/language_feature.cc
+++ b/src/tint/lang/wgsl/features/language_feature.cc
@@ -66,6 +66,9 @@
if (str == "readonly_and_readwrite_storage_textures") {
return LanguageFeature::kReadonlyAndReadwriteStorageTextures;
}
+ if (str == "unrestricted_pointer_parameters") {
+ return LanguageFeature::kUnrestrictedPointerParameters;
+ }
return LanguageFeature::kUndefined;
}
@@ -89,6 +92,8 @@
return "pointer_composite_access";
case LanguageFeature::kReadonlyAndReadwriteStorageTextures:
return "readonly_and_readwrite_storage_textures";
+ case LanguageFeature::kUnrestrictedPointerParameters:
+ return "unrestricted_pointer_parameters";
}
return "<unknown>";
}
diff --git a/src/tint/lang/wgsl/features/language_feature.h b/src/tint/lang/wgsl/features/language_feature.h
index aa7315b..b7aedbb 100644
--- a/src/tint/lang/wgsl/features/language_feature.h
+++ b/src/tint/lang/wgsl/features/language_feature.h
@@ -54,6 +54,7 @@
kPacked4X8IntegerDotProduct,
kPointerCompositeAccess,
kReadonlyAndReadwriteStorageTextures,
+ kUnrestrictedPointerParameters,
};
/// @param value the enum value
@@ -74,6 +75,7 @@
"packed_4x8_integer_dot_product",
"pointer_composite_access",
"readonly_and_readwrite_storage_textures",
+ "unrestricted_pointer_parameters",
};
/// All features
@@ -86,6 +88,7 @@
LanguageFeature::kPacked4X8IntegerDotProduct,
LanguageFeature::kPointerCompositeAccess,
LanguageFeature::kReadonlyAndReadwriteStorageTextures,
+ LanguageFeature::kUnrestrictedPointerParameters,
};
} // namespace tint::wgsl
diff --git a/src/tint/lang/wgsl/features/status.cc b/src/tint/lang/wgsl/features/status.cc
index f8297ea..2fe2f02 100644
--- a/src/tint/lang/wgsl/features/status.cc
+++ b/src/tint/lang/wgsl/features/status.cc
@@ -33,9 +33,10 @@
FeatureStatus GetLanguageFeatureStatus(LanguageFeature f) {
switch (f) {
- case LanguageFeature::kReadonlyAndReadwriteStorageTextures:
case LanguageFeature::kPacked4X8IntegerDotProduct:
case LanguageFeature::kPointerCompositeAccess:
+ case LanguageFeature::kReadonlyAndReadwriteStorageTextures:
+ case LanguageFeature::kUnrestrictedPointerParameters:
return FeatureStatus::kExperimental;
case LanguageFeature::kUndefined:
return FeatureStatus::kUnknown;
diff --git a/src/tint/lang/wgsl/helpers/append_vector_test.cc b/src/tint/lang/wgsl/helpers/append_vector_test.cc
index f1a95d9..26ebba4 100644
--- a/src/tint/lang/wgsl/helpers/append_vector_test.cc
+++ b/src/tint/lang/wgsl/helpers/append_vector_test.cc
@@ -49,7 +49,7 @@
auto* vec_12 = Call<vec2<i32>>(scalar_1, scalar_2);
WrapInFunction(vec_12, scalar_3);
- resolver::Resolver resolver(this, {});
+ resolver::Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_TRUE(resolver.Resolve());
ASSERT_THAT(resolver.Diagnostics(), testing::IsEmpty());
@@ -90,7 +90,7 @@
auto* vec_12 = Call<vec2<i32>>(scalar_1, scalar_2);
WrapInFunction(vec_12, scalar_3);
- resolver::Resolver resolver(this, {});
+ resolver::Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_TRUE(resolver.Resolve());
ASSERT_THAT(resolver.Diagnostics(), testing::IsEmpty());
@@ -138,7 +138,7 @@
auto* vec_12 = Call<vec2<i32>>(uvec_12);
WrapInFunction(vec_12, scalar_3);
- resolver::Resolver resolver(this, {});
+ resolver::Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_TRUE(resolver.Resolve());
ASSERT_THAT(resolver.Diagnostics(), testing::IsEmpty());
@@ -187,7 +187,7 @@
auto* vec_12 = Call<vec2<i32>>(scalar_1, scalar_2);
WrapInFunction(vec_12, scalar_3);
- resolver::Resolver resolver(this, {});
+ resolver::Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_TRUE(resolver.Resolve());
ASSERT_THAT(resolver.Diagnostics(), testing::IsEmpty());
@@ -233,7 +233,7 @@
auto* vec_123 = Call<vec3<i32>>(scalar_1, scalar_2, scalar_3);
WrapInFunction(vec_123, scalar_4);
- resolver::Resolver resolver(this, {});
+ resolver::Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_TRUE(resolver.Resolve());
ASSERT_THAT(resolver.Diagnostics(), testing::IsEmpty());
@@ -276,7 +276,7 @@
auto* scalar_3 = Expr(3_i);
WrapInFunction(vec_12, scalar_3);
- resolver::Resolver resolver(this, {});
+ resolver::Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_TRUE(resolver.Resolve());
ASSERT_THAT(resolver.Diagnostics(), testing::IsEmpty());
@@ -315,7 +315,7 @@
auto* vec_12 = Call<vec2<i32>>(scalar_1, scalar_2);
WrapInFunction(vec_12, scalar_3);
- resolver::Resolver resolver(this, {});
+ resolver::Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_TRUE(resolver.Resolve());
ASSERT_THAT(resolver.Diagnostics(), testing::IsEmpty());
@@ -356,7 +356,7 @@
auto* scalar_3 = Expr("scalar_3");
WrapInFunction(vec_12, scalar_3);
- resolver::Resolver resolver(this, {});
+ resolver::Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_TRUE(resolver.Resolve());
ASSERT_THAT(resolver.Diagnostics(), testing::IsEmpty());
@@ -394,7 +394,7 @@
auto* scalar_3 = Expr("scalar_3");
WrapInFunction(vec_12, scalar_3);
- resolver::Resolver resolver(this, {});
+ resolver::Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_TRUE(resolver.Resolve());
ASSERT_THAT(resolver.Diagnostics(), testing::IsEmpty());
@@ -436,7 +436,7 @@
auto* scalar_3 = Expr("scalar_3");
WrapInFunction(vec_12, scalar_3);
- resolver::Resolver resolver(this, {});
+ resolver::Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_TRUE(resolver.Resolve());
ASSERT_THAT(resolver.Diagnostics(), testing::IsEmpty());
@@ -472,7 +472,7 @@
auto* vec000 = Call<vec3<i32>>();
WrapInFunction(vec000, scalar);
- resolver::Resolver resolver(this, {});
+ resolver::Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_TRUE(resolver.Resolve());
ASSERT_THAT(resolver.Diagnostics(), testing::IsEmpty());
diff --git a/src/tint/lang/wgsl/language_feature_test.cc b/src/tint/lang/wgsl/language_feature_test.cc
index 991f5db..dc3b564 100644
--- a/src/tint/lang/wgsl/language_feature_test.cc
+++ b/src/tint/lang/wgsl/language_feature_test.cc
@@ -67,6 +67,7 @@
{"pointer_composite_access", LanguageFeature::kPointerCompositeAccess},
{"readonly_and_readwrite_storage_textures",
LanguageFeature::kReadonlyAndReadwriteStorageTextures},
+ {"unrestricted_pointer_parameters", LanguageFeature::kUnrestrictedPointerParameters},
};
static constexpr Case kInvalidCases[] = {
@@ -94,6 +95,9 @@
{"readonlF_and_readwrite_st9rage_textues", LanguageFeature::kUndefined},
{"readonly_and_radwrite_storage_textures", LanguageFeature::kUndefined},
{"readonly_and_readwrite_sOOrage_tVxRRures", LanguageFeature::kUndefined},
+ {"unrestrictd_pointer_payameters", LanguageFeature::kUndefined},
+ {"unrerrt77iGted_poillter_paramenners", LanguageFeature::kUndefined},
+ {"unrestricted4point00r_parameters", LanguageFeature::kUndefined},
};
using LanguageFeatureParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc b/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc
index 0af6015..6f689b3 100644
--- a/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc
@@ -205,7 +205,7 @@
// Error when unknown extension found
EXPECT_TRUE(p->has_error());
EXPECT_EQ(p->error(), R"(1:8: expected extension
-Possible values: 'chromium_disable_uniformity_analysis', 'chromium_experimental_framebuffer_fetch', 'chromium_experimental_full_ptr_parameters', 'chromium_experimental_pixel_local', 'chromium_experimental_push_constant', 'chromium_experimental_subgroups', 'chromium_internal_dual_source_blending', 'chromium_internal_relaxed_uniform_layout', 'f16')");
+Possible values: 'chromium_disable_uniformity_analysis', 'chromium_experimental_framebuffer_fetch', 'chromium_experimental_pixel_local', 'chromium_experimental_push_constant', 'chromium_experimental_subgroups', 'chromium_internal_dual_source_blending', 'chromium_internal_relaxed_uniform_layout', 'f16')");
auto program = p->program();
auto& ast = program.AST();
EXPECT_EQ(ast.Enables().Length(), 0u);
diff --git a/src/tint/lang/wgsl/resolver/builtin_enum_test.cc b/src/tint/lang/wgsl/resolver/builtin_enum_test.cc
index cc7bbeb..14e266d 100644
--- a/src/tint/lang/wgsl/resolver/builtin_enum_test.cc
+++ b/src/tint/lang/wgsl/resolver/builtin_enum_test.cc
@@ -70,7 +70,6 @@
TEST_P(ResolverAddressSpaceUsedWithTemplateArgs, Test) {
// fn f(p : ptr<ADDRESS_SPACE<T>, f32) {}
- Enable(wgsl::Extension::kChromiumExperimentalFullPtrParameters);
auto* tmpl = Ident(Source{{12, 34}}, GetParam(), "T");
Func("f", Vector{Param("p", ty("ptr", tmpl, ty.f32()))}, ty.void_(), tint::Empty);
EXPECT_FALSE(r()->Resolve());
diff --git a/src/tint/lang/wgsl/resolver/builtin_validation_test.cc b/src/tint/lang/wgsl/resolver/builtin_validation_test.cc
index 3bae4eb..638a54c 100644
--- a/src/tint/lang/wgsl/resolver/builtin_validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/builtin_validation_test.cc
@@ -568,7 +568,7 @@
Vector{Expr(1_u), Expr(2_u)})),
});
- auto resolver = Resolver(this, {});
+ Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'dot4I8Packed' requires the "
@@ -597,7 +597,7 @@
Vector{Expr(1_u), Expr(2_u)})),
});
- auto resolver = Resolver(this, {});
+ Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'dot4U8Packed' requires the "
@@ -624,7 +624,7 @@
Return(Call(Source{Source::Location{12, 34}}, "pack4xI8", Call<vec4<i32>>())),
});
- auto resolver = Resolver(this, {});
+ Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'pack4xI8' requires the "
@@ -651,7 +651,7 @@
Return(Call(Source{Source::Location{12, 34}}, "pack4xU8", Call<vec4<u32>>())),
});
- auto resolver = Resolver(this, {});
+ Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'pack4xU8' requires the "
@@ -678,7 +678,7 @@
Return(Call(Source{Source::Location{12, 34}}, "pack4xI8Clamp", Call<vec4<i32>>())),
});
- auto resolver = Resolver(this, {});
+ Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'pack4xI8Clamp' requires the "
@@ -705,7 +705,7 @@
Return(Call(Source{Source::Location{12, 34}}, "pack4xU8Clamp", Call<vec4<u32>>())),
});
- auto resolver = Resolver(this, {});
+ Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'pack4xU8Clamp' requires the "
@@ -732,7 +732,7 @@
Return(Call(Source{Source::Location{12, 34}}, "unpack4xI8", Call<u32>())),
});
- auto resolver = Resolver(this, {});
+ Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'unpack4xI8' requires the "
@@ -759,7 +759,7 @@
Return(Call(Source{Source::Location{12, 34}}, "unpack4xU8", Call<u32>())),
});
- auto resolver = Resolver(this, {});
+ Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'unpack4xU8' requires the "
@@ -988,7 +988,7 @@
CallStmt(Call(Source{Source::Location{12, 34}}, "textureBarrier")),
});
- auto resolver = Resolver(this, {});
+ Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: built-in function 'textureBarrier' requires the "
diff --git a/src/tint/lang/wgsl/resolver/call_validation_test.cc b/src/tint/lang/wgsl/resolver/call_validation_test.cc
index 47da663..9a0d7c1 100644
--- a/src/tint/lang/wgsl/resolver/call_validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/call_validation_test.cc
@@ -146,7 +146,8 @@
EXPECT_EQ(r()->error(), "12:34 error: cannot take the address of expression");
}
-TEST_F(ResolverCallValidationTest, PointerArgument_AddressOfFunctionMember) {
+TEST_F(ResolverCallValidationTest,
+ PointerArgument_AddressOfFunctionMember_WithoutUnrestrictedPointerParameters) {
// struct S { m: i32; };
// fn foo(p: ptr<function, i32>) {}
// fn main() {
@@ -164,22 +165,23 @@
CallStmt(Call("foo", AddressOf(Source{{12, 34}}, MemberAccessor("v", "m")))),
});
- EXPECT_FALSE(r()->Resolve());
- EXPECT_EQ(r()->error(),
- "12:34 error: arguments of pointer type must not point to a subset of the "
- "originating variable");
+ Resolver r{this, wgsl::AllowedFeatures{}};
+
+ EXPECT_FALSE(r.Resolve());
+ EXPECT_EQ(
+ r.error(),
+ R"(12:34 error: arguments of pointer type must not point to a subset of the originating variable)");
}
TEST_F(ResolverCallValidationTest,
- PointerArgument_AddressOfFunctionMember_WithFullPtrParametersExt) {
- // enable chromium_experimental_full_ptr_parameters;
+ PointerArgument_AddressOfFunctionMember_WithUnrestrictedPointerParameters) {
// struct S { m: i32; };
// fn foo(p: ptr<function, i32>) {}
// fn main() {
// var v : S;
// foo(&v.m);
// }
- Enable(wgsl::Extension::kChromiumExperimentalFullPtrParameters);
+
auto* S = Structure("S", Vector{
Member("m", ty.i32()),
});
@@ -322,7 +324,7 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
}
-TEST_F(ResolverCallValidationTest, LetPointer_NotWholeVar) {
+TEST_F(ResolverCallValidationTest, LetPointer_NotWholeVar_WithoutUnrestrictedPointerParameters) {
// fn foo(p : ptr<function, i32>) {}
// @fragment
// fn main() {
@@ -344,14 +346,14 @@
Vector{
Stage(ast::PipelineStage::kFragment),
});
- EXPECT_FALSE(r()->Resolve());
- EXPECT_EQ(r()->error(),
+ Resolver r(this, wgsl::AllowedFeatures{});
+ EXPECT_FALSE(r.Resolve());
+ EXPECT_EQ(r.error(),
"12:34 error: arguments of pointer type must not point to a subset of the "
"originating variable");
}
-TEST_F(ResolverCallValidationTest, LetPointer_NotWholeVar_WithFullPtrParametersExt) {
- // enable chromium_experimental_full_ptr_parameters;
+TEST_F(ResolverCallValidationTest, LetPointer_NotWholeVar_WithUnrestrictedPointerParameters) {
// fn foo(p : ptr<function, i32>) {}
// @fragment
// fn main() {
@@ -359,7 +361,7 @@
// let p: ptr<function, i32> = &(v[0]);
// x(p);
// }
- Enable(wgsl::Extension::kChromiumExperimentalFullPtrParameters);
+
Func("foo",
Vector{
Param("p", ty.ptr<function, i32>()),
@@ -406,7 +408,8 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
}
-TEST_F(ResolverCallValidationTest, ComplexPointerChain_NotWholeVar) {
+TEST_F(ResolverCallValidationTest,
+ ComplexPointerChain_NotWholeVar_WithoutUnrestrictedPointerParameters) {
// fn foo(p : ptr<function, i32>) {}
// @fragment
// fn main() {
@@ -432,14 +435,15 @@
Vector{
Stage(ast::PipelineStage::kFragment),
});
- EXPECT_FALSE(r()->Resolve());
- EXPECT_EQ(r()->error(),
+ Resolver r{this, {}};
+ EXPECT_FALSE(r.Resolve());
+ EXPECT_EQ(r.error(),
"12:34 error: arguments of pointer type must not point to a subset of the "
"originating variable");
}
-TEST_F(ResolverCallValidationTest, ComplexPointerChain_NotWholeVar_WithFullPtrParametersExt) {
- // enable chromium_experimental_full_ptr_parameters;
+TEST_F(ResolverCallValidationTest,
+ ComplexPointerChain_NotWholeVar_WithUnrestrictedPointerParameters) {
// fn foo(p : ptr<function, i32>) {}
// @fragment
// fn main() {
@@ -449,7 +453,7 @@
// let p3 = &(*p2)[0];
// foo(&*p);
// }
- Enable(wgsl::Extension::kChromiumExperimentalFullPtrParameters);
+
Func("foo",
Vector{
Param("p", ty.ptr<function, i32>()),
diff --git a/src/tint/lang/wgsl/resolver/expression_kind_test.cc b/src/tint/lang/wgsl/resolver/expression_kind_test.cc
index ff4d507..94c91c8e 100644
--- a/src/tint/lang/wgsl/resolver/expression_kind_test.cc
+++ b/src/tint/lang/wgsl/resolver/expression_kind_test.cc
@@ -309,7 +309,6 @@
GlobalVar("v", ty("texture_storage_2d", "rgba8unorm", expr), Group(0_u), Binding(0_u));
break;
case Use::kAddressSpace:
- Enable(wgsl::Extension::kChromiumExperimentalFullPtrParameters);
Func(Symbols().New(), Vector{Param("p", ty("ptr", expr, ty.f32()))}, ty.void_(),
tint::Empty);
break;
diff --git a/src/tint/lang/wgsl/resolver/function_validation_test.cc b/src/tint/lang/wgsl/resolver/function_validation_test.cc
index fd015f7..2a01da4 100644
--- a/src/tint/lang/wgsl/resolver/function_validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/function_validation_test.cc
@@ -1047,7 +1047,7 @@
enum class Expectation {
kAlwaysPass,
- kPassWithFullPtrParameterExtension,
+ kPassWithUnrestrictedPointerParameters,
kAlwaysFail,
kInvalid,
};
@@ -1059,7 +1059,11 @@
struct TestWithParams : ResolverTestWithParam<TestParams> {};
using ResolverFunctionParameterValidationTest = TestWithParams;
-TEST_P(ResolverFunctionParameterValidationTest, AddressSpaceNoExtension) {
+TEST_P(ResolverFunctionParameterValidationTest, AddressSpaceWithoutUnrestrictedPointerParameters) {
+ auto features = wgsl::AllowedFeatures::Everything();
+ features.features.erase(wgsl::LanguageFeature::kUnrestrictedPointerParameters);
+ Resolver r(this, features);
+
auto& param = GetParam();
Structure("S", Vector{Member("a", ty.i32())});
auto ptr_type = ty("ptr", Ident(Source{{12, 34}}, param.address_space), ty("S"));
@@ -1071,29 +1075,27 @@
}
if (param.expectation == Expectation::kAlwaysPass) {
- ASSERT_TRUE(r()->Resolve()) << r()->error();
+ ASSERT_TRUE(r.Resolve()) << r.error();
} else {
StringStream ss;
ss << param.address_space;
- EXPECT_FALSE(r()->Resolve());
+ EXPECT_FALSE(r.Resolve());
if (param.expectation == Expectation::kInvalid) {
std::string err = R"(12:34 error: unresolved address space '${addr_space}'
12:34 note: Possible values: 'function', 'pixel_local', 'private', 'push_constant', 'storage', 'uniform', 'workgroup')";
err = tint::ReplaceAll(err, "${addr_space}", tint::ToString(param.address_space));
- EXPECT_EQ(r()->error(), err);
+ EXPECT_EQ(r.error(), err);
} else {
- EXPECT_EQ(r()->error(),
- "12:34 error: function parameter of pointer type cannot be in '" +
- tint::ToString(param.address_space) + "' address space");
+ EXPECT_EQ(r.error(), "12:34 error: function parameter of pointer type cannot be in '" +
+ tint::ToString(param.address_space) + "' address space");
}
}
}
-TEST_P(ResolverFunctionParameterValidationTest, AddressSpaceWithFullPtrParameterExtension) {
+TEST_P(ResolverFunctionParameterValidationTest, AddressSpaceWithUnrestrictedPointerParameters) {
auto& param = GetParam();
Structure("S", Vector{Member("a", ty.i32())});
auto ptr_type = ty("ptr", Ident(Source{{12, 34}}, param.address_space), ty("S"));
auto* arg = Param(Source{{12, 34}}, "p", ptr_type);
- Enable(wgsl::Extension::kChromiumExperimentalFullPtrParameters);
Func("f", Vector{arg}, ty.void_(), tint::Empty);
if (param.address_space == core::AddressSpace::kPixelLocal) {
@@ -1101,7 +1103,7 @@
}
if (param.expectation == Expectation::kAlwaysPass ||
- param.expectation == Expectation::kPassWithFullPtrParameterExtension) {
+ param.expectation == Expectation::kPassWithUnrestrictedPointerParameters) {
ASSERT_TRUE(r()->Resolve()) << r()->error();
} else {
EXPECT_FALSE(r()->Resolve());
@@ -1120,17 +1122,19 @@
INSTANTIATE_TEST_SUITE_P(
ResolverTest,
ResolverFunctionParameterValidationTest,
- testing::Values(
- TestParams{core::AddressSpace::kUndefined, Expectation::kInvalid},
- TestParams{core::AddressSpace::kIn, Expectation::kAlwaysFail},
- TestParams{core::AddressSpace::kOut, Expectation::kAlwaysFail},
- TestParams{core::AddressSpace::kUniform, Expectation::kPassWithFullPtrParameterExtension},
- TestParams{core::AddressSpace::kWorkgroup, Expectation::kPassWithFullPtrParameterExtension},
- TestParams{core::AddressSpace::kHandle, Expectation::kInvalid},
- TestParams{core::AddressSpace::kStorage, Expectation::kPassWithFullPtrParameterExtension},
- TestParams{core::AddressSpace::kPixelLocal, Expectation::kAlwaysFail},
- TestParams{core::AddressSpace::kPrivate, Expectation::kAlwaysPass},
- TestParams{core::AddressSpace::kFunction, Expectation::kAlwaysPass}));
+ testing::Values(TestParams{core::AddressSpace::kUndefined, Expectation::kInvalid},
+ TestParams{core::AddressSpace::kIn, Expectation::kAlwaysFail},
+ TestParams{core::AddressSpace::kOut, Expectation::kAlwaysFail},
+ TestParams{core::AddressSpace::kUniform,
+ Expectation::kPassWithUnrestrictedPointerParameters},
+ TestParams{core::AddressSpace::kWorkgroup,
+ Expectation::kPassWithUnrestrictedPointerParameters},
+ TestParams{core::AddressSpace::kHandle, Expectation::kInvalid},
+ TestParams{core::AddressSpace::kStorage,
+ Expectation::kPassWithUnrestrictedPointerParameters},
+ TestParams{core::AddressSpace::kPixelLocal, Expectation::kAlwaysFail},
+ TestParams{core::AddressSpace::kPrivate, Expectation::kAlwaysPass},
+ TestParams{core::AddressSpace::kFunction, Expectation::kAlwaysPass}));
} // namespace
} // namespace tint::resolver
diff --git a/src/tint/lang/wgsl/resolver/type_validation_test.cc b/src/tint/lang/wgsl/resolver/type_validation_test.cc
index e5fcf1c..0264afc 100644
--- a/src/tint/lang/wgsl/resolver/type_validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/type_validation_test.cc
@@ -1217,7 +1217,7 @@
GlobalVar("a", st, Group(0_a), Binding(0_a));
- auto resolver = Resolver(this, {});
+ auto resolver = Resolver{this, wgsl::AllowedFeatures{}};
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: read-only storage textures require the "
@@ -1246,7 +1246,7 @@
GlobalVar("a", st, Group(0_a), Binding(0_a));
- auto resolver = Resolver(this, {});
+ Resolver resolver{this, wgsl::AllowedFeatures{}};
EXPECT_FALSE(resolver.Resolve());
EXPECT_EQ(resolver.error(),
"12:34 error: read-write storage textures require the "
diff --git a/src/tint/lang/wgsl/resolver/uniformity_test.cc b/src/tint/lang/wgsl/resolver/uniformity_test.cc
index cbab7db..4674237 100644
--- a/src/tint/lang/wgsl/resolver/uniformity_test.cc
+++ b/src/tint/lang/wgsl/resolver/uniformity_test.cc
@@ -3788,8 +3788,6 @@
TEST_F(UniformityAnalysisTest, LoadNonUniformGlobalThroughPointerParameter) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
fn bar(p : ptr<storage, i32, read_write>) {
@@ -3805,15 +3803,15 @@
RunTest(src, false);
EXPECT_EQ(error_,
- R"(test:8:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ R"(test:6:5 error: 'workgroupBarrier' must only be called from uniform control flow
workgroupBarrier();
^^^^^^^^^^^^^^^^
-test:7:3 note: control flow depends on possibly non-uniform value
+test:5:3 note: control flow depends on possibly non-uniform value
if (*p == 0) {
^^
-test:7:8 note: parameter 'p' of 'bar' may be non-uniform
+test:5:8 note: parameter 'p' of 'bar' may be non-uniform
if (*p == 0) {
^
)");
@@ -3821,8 +3819,6 @@
TEST_F(UniformityAnalysisTest, LoadNonUniformGlobalThroughPointerParameter_ViaReturnValue) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
fn bar(p : ptr<storage, i32, read_write>) -> i32 {
@@ -3838,15 +3834,15 @@
RunTest(src, false);
EXPECT_EQ(error_,
- R"(test:12:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ R"(test:10:5 error: 'workgroupBarrier' must only be called from uniform control flow
workgroupBarrier();
^^^^^^^^^^^^^^^^
-test:11:3 note: control flow depends on possibly non-uniform value
+test:9:3 note: control flow depends on possibly non-uniform value
if (0 == bar(&non_uniform)) {
^^
-test:11:12 note: return value of 'bar' may be non-uniform
+test:9:12 note: return value of 'bar' may be non-uniform
if (0 == bar(&non_uniform)) {
^^^^^^^^^^^^^^^^^
)");
@@ -4011,8 +4007,6 @@
TEST_F(UniformityAnalysisTest, LoadUniformThroughNonUniformPointer_ViaParameter) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
fn bar(p : ptr<function, array<i32, 4>>) {
@@ -4033,15 +4027,15 @@
RunTest(src, false);
EXPECT_EQ(error_,
- R"(test:10:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ R"(test:8:5 error: 'workgroupBarrier' must only be called from uniform control flow
workgroupBarrier();
^^^^^^^^^^^^^^^^
-test:9:3 note: control flow depends on possibly non-uniform value
+test:7:3 note: control flow depends on possibly non-uniform value
if (*local_p == 0) {
^^
-test:8:23 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value
+test:6:23 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value
let local_p = &(*p)[non_uniform];
^^^^^^^^^^^
)");
@@ -4049,8 +4043,6 @@
TEST_F(UniformityAnalysisTest, LoadUniformThroughNonUniformPointer_ViaParameterChain) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
fn zoo(p : ptr<function, i32>) {
@@ -4074,31 +4066,31 @@
RunTest(src, false);
EXPECT_EQ(error_,
- R"(test:8:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ R"(test:6:5 error: 'workgroupBarrier' must only be called from uniform control flow
workgroupBarrier();
^^^^^^^^^^^^^^^^
-test:7:3 note: control flow depends on possibly non-uniform value
+test:5:3 note: control flow depends on possibly non-uniform value
if (*p == 0) {
^^
-test:7:8 note: parameter 'p' of 'zoo' may be non-uniform
+test:5:8 note: parameter 'p' of 'zoo' may be non-uniform
if (*p == 0) {
^
-test:13:7 note: possibly non-uniform value passed via pointer here
+test:11:7 note: possibly non-uniform value passed via pointer here
zoo(p);
^
-test:12:8 note: reading from 'p' may result in a non-uniform value
+test:10:8 note: reading from 'p' may result in a non-uniform value
fn bar(p : ptr<function, i32>) {
^
-test:21:7 note: possibly non-uniform value passed via pointer here
+test:19:7 note: possibly non-uniform value passed via pointer here
bar(p);
^
-test:20:14 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value
+test:18:14 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value
let p = &v[non_uniform];
^^^^^^^^^^^
)");
@@ -4138,8 +4130,6 @@
TEST_F(UniformityAnalysisTest, LoadNonUniformThroughUniformPointer_ViaParameter) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
@group(0) @binding(1) var<storage, read> uniform_idx : i32;
@@ -4164,31 +4154,31 @@
RunTest(src, false);
EXPECT_EQ(error_,
- R"(test:9:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ R"(test:7:5 error: 'workgroupBarrier' must only be called from uniform control flow
workgroupBarrier();
^^^^^^^^^^^^^^^^
-test:8:3 note: control flow depends on possibly non-uniform value
+test:6:3 note: control flow depends on possibly non-uniform value
if (*p == 0) {
^^
-test:8:8 note: parameter 'p' of 'zoo' may be non-uniform
+test:6:8 note: parameter 'p' of 'zoo' may be non-uniform
if (*p == 0) {
^
-test:14:7 note: possibly non-uniform value passed via pointer here
+test:12:7 note: possibly non-uniform value passed via pointer here
zoo(p);
^
-test:13:8 note: reading from 'p' may result in a non-uniform value
+test:11:8 note: reading from 'p' may result in a non-uniform value
fn bar(p : ptr<function, i32>) {
^
-test:22:7 note: possibly non-uniform value passed via pointer here
+test:20:7 note: possibly non-uniform value passed via pointer here
bar(p);
^
-test:19:34 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value
+test:17:34 note: reading from read_write storage buffer 'non_uniform' may result in a non-uniform value
var v = array<i32, 4>(0, 0, 0, non_uniform);
^^^^^^^^^^^
)");
@@ -5355,8 +5345,6 @@
TEST_F(UniformityAnalysisTest, AssignUniformToPrivatePointerParameter_StillNonUniform) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> non_uniform : i32;
fn bar(p : ptr<private, i32>) {
@@ -5373,15 +5361,15 @@
RunTest(src, false);
EXPECT_EQ(error_,
- R"(test:9:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ R"(test:7:5 error: 'workgroupBarrier' must only be called from uniform control flow
workgroupBarrier();
^^^^^^^^^^^^^^^^
-test:8:3 note: control flow depends on possibly non-uniform value
+test:6:3 note: control flow depends on possibly non-uniform value
if (*p == 0) {
^^
-test:8:8 note: parameter 'p' of 'bar' may be non-uniform
+test:6:8 note: parameter 'p' of 'bar' may be non-uniform
if (*p == 0) {
^
)");
@@ -5389,8 +5377,6 @@
TEST_F(UniformityAnalysisTest, AssignUniformToWorkgroupPointerParameter_StillNonUniform) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> non_uniform : i32;
fn bar(p : ptr<workgroup, i32>) {
@@ -5407,15 +5393,15 @@
RunTest(src, false);
EXPECT_EQ(error_,
- R"(test:9:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ R"(test:7:5 error: 'workgroupBarrier' must only be called from uniform control flow
workgroupBarrier();
^^^^^^^^^^^^^^^^
-test:8:3 note: control flow depends on possibly non-uniform value
+test:6:3 note: control flow depends on possibly non-uniform value
if (*p == 0) {
^^
-test:8:8 note: parameter 'p' of 'bar' may be non-uniform
+test:6:8 note: parameter 'p' of 'bar' may be non-uniform
if (*p == 0) {
^
)");
@@ -5423,8 +5409,6 @@
TEST_F(UniformityAnalysisTest, AssignUniformToStoragePointerParameter_StillNonUniform) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
fn bar(p : ptr<storage, i32, read_write>) {
@@ -5441,15 +5425,15 @@
RunTest(src, false);
EXPECT_EQ(error_,
- R"(test:9:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ R"(test:7:5 error: 'workgroupBarrier' must only be called from uniform control flow
workgroupBarrier();
^^^^^^^^^^^^^^^^
-test:8:3 note: control flow depends on possibly non-uniform value
+test:6:3 note: control flow depends on possibly non-uniform value
if (*p == 0) {
^^
-test:8:8 note: parameter 'p' of 'bar' may be non-uniform
+test:6:8 note: parameter 'p' of 'bar' may be non-uniform
if (*p == 0) {
^
)");
@@ -5457,8 +5441,6 @@
TEST_F(UniformityAnalysisTest, LoadFromReadOnlyStoragePointerParameter_AlwaysUniform) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read> non_uniform : i32;
fn bar(p : ptr<storage, i32, read>) {
@@ -7799,8 +7781,6 @@
TEST_F(UniformityAnalysisTest,
ArrayElement_AssignUniformToElementWithNonUniformIndex_ViaPartialPointerParameter) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
var<private> non_uniform : i32;
fn bar(p : ptr<function, i32>) {
@@ -7819,15 +7799,15 @@
RunTest(src, false);
EXPECT_EQ(error_,
- R"(test:15:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ R"(test:13:5 error: 'workgroupBarrier' must only be called from uniform control flow
workgroupBarrier();
^^^^^^^^^^^^^^^^
-test:14:3 note: control flow depends on possibly non-uniform value
+test:12:3 note: control flow depends on possibly non-uniform value
if (arr[0] == 0) {
^^
-test:12:17 note: reading from module-scope private variable 'non_uniform' may result in a non-uniform value
+test:10:17 note: reading from module-scope private variable 'non_uniform' may result in a non-uniform value
let p = &(arr[non_uniform]);
^^^^^^^^^^^
)");
@@ -8663,8 +8643,6 @@
TEST_F(UniformityAnalysisTest, ArrayLength_OnPtrArg) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> arr : array<f32>;
fn bar(p : ptr<storage, array<f32>, read_write>) {
@@ -8683,8 +8661,6 @@
TEST_F(UniformityAnalysisTest, ArrayLength_PtrArgRequiredToBeUniformForRetval_Pass) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> arr : array<f32>;
fn length(p : ptr<storage, array<f32>, read_write>) -> u32 {
@@ -8704,8 +8680,6 @@
// TODO(jrprice): This test requires variable pointers.
TEST_F(UniformityAnalysisTest, DISABLED_ArrayLength_PtrArgRequiredToBeUniformForRetval_Fail) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
@group(0) @binding(1) var<storage, read_write> arr1 : array<f32>;
@group(0) @binding(2) var<storage, read_write> arr2 : array<f32>;
@@ -8741,8 +8715,6 @@
TEST_F(UniformityAnalysisTest, ArrayLength_PtrArgRequiredToBeUniformForOtherPtrResult_Pass) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> arr : array<f32>;
fn length(p : ptr<storage, array<f32>, read_write>, out : ptr<function, u32>) {
@@ -8765,8 +8737,6 @@
TEST_F(UniformityAnalysisTest,
DISABLED_ArrayLength_PtrArgRequiredToBeUniformForOtherPtrResult_Fail) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> non_uniform : i32;
@group(0) @binding(1) var<storage, read_write> arr1 : array<f32>;
@group(0) @binding(2) var<storage, read_write> arr2 : array<f32>;
@@ -8805,8 +8775,6 @@
// Test that a single pointer argument can directly require uniformity as well as affecting the
// uniformity of the return value.
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> arr : array<u32>;
fn bar(p : ptr<storage, array<u32>, read_write>) -> u32 {
@@ -8831,15 +8799,15 @@
RunTest(src, false);
EXPECT_EQ(error_,
- R"(test:21:5 error: 'workgroupBarrier' must only be called from uniform control flow
+ R"(test:19:5 error: 'workgroupBarrier' must only be called from uniform control flow
workgroupBarrier();
^^^^^^^^^^^^^^^^
-test:19:3 note: control flow depends on possibly non-uniform value
+test:17:3 note: control flow depends on possibly non-uniform value
if (0 == bar(p)) {
^^
-test:19:12 note: return value of 'bar' may be non-uniform
+test:17:12 note: return value of 'bar' may be non-uniform
if (0 == bar(p)) {
^^^^^^
)");
@@ -8864,8 +8832,6 @@
TEST_F(UniformityAnalysisTest, WorkgroupUniformLoad_ViaPtrArg) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
const wgsize = 4;
var<workgroup> data : array<u32, wgsize>;
@@ -8913,8 +8879,6 @@
TEST_F(UniformityAnalysisTest, WorkgroupUniformLoad_NonUniformPtr_ViaPtrArg) {
std::string src = R"(
-enable chromium_experimental_full_ptr_parameters;
-
const wgsize = 4;
var<workgroup> data : array<u32, wgsize>;
@@ -8933,19 +8897,19 @@
RunTest(src, false);
EXPECT_EQ(error_,
- R"(test:8:31 error: possibly non-uniform value passed here
+ R"(test:6:31 error: possibly non-uniform value passed here
return workgroupUniformLoad(p);
^
-test:8:31 note: parameter 'p' of 'foo' may be non-uniform
+test:6:31 note: parameter 'p' of 'foo' may be non-uniform
return workgroupUniformLoad(p);
^
-test:14:11 note: possibly non-uniform value passed here
+test:12:11 note: possibly non-uniform value passed here
if (foo(&data[idx]) > 0) {
^^^^^^^^^^
-test:14:17 note: builtin 'idx' of 'main' may be non-uniform
+test:12:17 note: builtin 'idx' of 'main' may be non-uniform
if (foo(&data[idx]) > 0) {
^^^
)");
diff --git a/src/tint/lang/wgsl/resolver/validator.cc b/src/tint/lang/wgsl/resolver/validator.cc
index 106d739..e6f7c31 100644
--- a/src/tint/lang/wgsl/resolver/validator.cc
+++ b/src/tint/lang/wgsl/resolver/validator.cc
@@ -845,8 +845,8 @@
case core::AddressSpace::kStorage:
case core::AddressSpace::kUniform:
case core::AddressSpace::kWorkgroup:
- ok = enabled_extensions_.Contains(
- wgsl::Extension::kChromiumExperimentalFullPtrParameters);
+ ok = allowed_features_.features.count(
+ wgsl::LanguageFeature::kUnrestrictedPointerParameters) != 0;
break;
default:
break;
@@ -1912,8 +1912,8 @@
}
if (param_type->Is<core::type::Pointer>() &&
- !enabled_extensions_.Contains(
- wgsl::Extension::kChromiumExperimentalFullPtrParameters)) {
+ !allowed_features_.features.count(
+ wgsl::LanguageFeature::kUnrestrictedPointerParameters)) {
// https://gpuweb.github.io/gpuweb/wgsl/#function-restriction
// Each argument of pointer type to a user-defined function must have the same memory
// view as its root identifier.
diff --git a/src/tint/lang/wgsl/wgsl.def b/src/tint/lang/wgsl/wgsl.def
index 41cd202..e23de8a 100644
--- a/src/tint/lang/wgsl/wgsl.def
+++ b/src/tint/lang/wgsl/wgsl.def
@@ -72,9 +72,6 @@
chromium_disable_uniformity_analysis
// A Chromium-specific extension for push constants
chromium_experimental_push_constant
- // A Chromium-specific extension that enables passing of uniform, storage and workgroup
- // address-spaced pointers as parameters, as well as pointers into sub-objects.
- chromium_experimental_full_ptr_parameters
// A Chromium-specific extension that adds basic subgroup functionality.
chromium_experimental_subgroups
// A Chromium-specific extension that relaxes memory layout requirements for uniform storage.
@@ -91,9 +88,10 @@
// https://gpuweb.github.io/gpuweb/wgsl/#language-extensions-sec
enum language_feature {
- readonly_and_readwrite_storage_textures
packed_4x8_integer_dot_product
pointer_composite_access
+ readonly_and_readwrite_storage_textures
+ unrestricted_pointer_parameters
// Language features used only for testing whose status will never change.
chromium_testing_unimplemented
diff --git a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
index 47b6468..712fab3 100644
--- a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
+++ b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
@@ -267,9 +267,6 @@
attrs.Push(b.Invariant());
}
- if (ParamRequiresFullPtrParameters(param->Type())) {
- Enable(wgsl::Extension::kChromiumExperimentalFullPtrParameters);
- }
return b.Param(name, ty, std::move(attrs));
});
@@ -686,12 +683,6 @@
tint::Switch(
call, //
[&](const core::ir::UserCall* c) {
- for (auto* arg : call->Args()) {
- if (ArgRequiresFullPtrParameters(arg)) {
- Enable(wgsl::Extension::kChromiumExperimentalFullPtrParameters);
- break;
- }
- }
auto* expr = b.Call(NameFor(c->Target()), std::move(args));
if (call->Results().IsEmpty() || !call->Result(0)->IsUsed()) {
Append(b.CallStmt(expr));
@@ -1310,44 +1301,6 @@
return false;
}
}
-
- /// @returns true if a parameter of the type @p ty requires the
- /// kChromiumExperimentalFullPtrParameters extension to be enabled.
- bool ParamRequiresFullPtrParameters(const core::type::Type* ty) {
- if (auto* ptr = ty->As<core::type::Pointer>()) {
- switch (ptr->AddressSpace()) {
- case core::AddressSpace::kUniform:
- case core::AddressSpace::kStorage:
- case core::AddressSpace::kWorkgroup:
- return true;
- default:
- return false;
- }
- }
- return false;
- }
-
- /// @returns true if the argument @p arg requires the kChromiumExperimentalFullPtrParameters
- /// extension to be enabled.
- bool ArgRequiresFullPtrParameters(const core::ir::Value* arg) {
- if (!arg->Type()->Is<core::type::Pointer>()) {
- return false;
- }
-
- auto res = arg->As<core::ir::InstructionResult>();
- while (res) {
- auto* inst = res->Instruction();
- if (inst->Is<core::ir::Access>()) {
- return true; // Passing pointer into sub-object
- }
- if (auto* let = inst->As<core::ir::Let>()) {
- res = let->Value()->As<core::ir::InstructionResult>();
- } else {
- break;
- }
- }
- return false;
- }
};
} // namespace
diff --git a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
index cd77dba..b3119c7 100644
--- a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
+++ b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
@@ -3212,105 +3212,6 @@
}
////////////////////////////////////////////////////////////////////////////////
-// chromium_experimental_full_ptr_parameters
-////////////////////////////////////////////////////////////////////////////////
-TEST_F(IRToProgramTest, Enable_ChromiumExperimentalFullPtrParameters_StoragePtrParameter) {
- auto* fn = b.Function("f", ty.void_());
- fn->SetParams({b.FunctionParam("p", ty.ptr<storage, i32>())});
-
- b.Append(fn->Block(), [&] { b.Return(fn); });
-
- EXPECT_WGSL(R"(
-enable chromium_experimental_full_ptr_parameters;
-
-fn f(p : ptr<storage, i32, read_write>) {
-}
-)");
-}
-
-TEST_F(IRToProgramTest, Enable_ChromiumExperimentalFullPtrParameters_UniformPtrParameter) {
- auto* fn = b.Function("f", ty.void_());
- fn->SetParams({b.FunctionParam("p", ty.ptr<uniform, i32>())});
-
- b.Append(fn->Block(), [&] { b.Return(fn); });
-
- EXPECT_WGSL(R"(
-enable chromium_experimental_full_ptr_parameters;
-
-fn f(p : ptr<uniform, i32>) {
-}
-)");
-}
-
-TEST_F(IRToProgramTest, Enable_ChromiumExperimentalFullPtrParameters_WorkgroupPtrParameter) {
- auto* fn = b.Function("f", ty.void_());
- fn->SetParams({b.FunctionParam("p", ty.ptr<workgroup, i32>())});
-
- b.Append(fn->Block(), [&] { b.Return(fn); });
-
- EXPECT_WGSL(R"(
-enable chromium_experimental_full_ptr_parameters;
-
-fn f(p : ptr<workgroup, i32>) {
-}
-)");
-}
-
-TEST_F(IRToProgramTest, Enable_ChromiumExperimentalFullPtrParameters_SubObjectPtrArg) {
- auto* x = b.Function("x", ty.void_());
- x->SetParams({b.FunctionParam("p", ty.ptr<function, vec3<f32>>())});
- b.Append(x->Block(), [&] { b.Return(x); });
-
- auto* y = b.Function("y", ty.void_());
- b.Append(y->Block(), [&] {
- auto* m = b.Var<function, mat3x3<f32>>();
- auto* v = b.Access(ty.ptr<function, vec3<f32>>(), m, 1_i);
- b.Call(ty.void_(), x, v);
- b.Return(y);
- });
-
- EXPECT_WGSL(R"(
-enable chromium_experimental_full_ptr_parameters;
-
-fn x(p : ptr<function, vec3<f32>>) {
-}
-
-fn y() {
- var v : mat3x3<f32>;
- x(&(v[1i]));
-}
-)");
-}
-
-TEST_F(IRToProgramTest, Enable_ChromiumExperimentalFullPtrParameters_SubObjectPtrArg_ViaLet) {
- auto* x = b.Function("x", ty.void_());
- x->SetParams({b.FunctionParam("p", ty.ptr<function, vec3<f32>>())});
- b.Append(x->Block(), [&] { b.Return(x); });
-
- auto* y = b.Function("y", ty.void_());
- b.Append(y->Block(), [&] {
- auto* m = b.Var<function, mat3x3<f32>>();
- auto* v = b.Access(ty.ptr<function, vec3<f32>>(), m, 1_i);
- auto* l = b.Let("l", v);
- b.Call(ty.void_(), x, l);
- b.Return(y);
- });
-
- EXPECT_WGSL(R"(
-enable chromium_experimental_full_ptr_parameters;
-
-fn x(p : ptr<function, vec3<f32>>) {
-}
-
-fn y() {
- var v : mat3x3<f32>;
- let l = &(v[1i]);
- x(l);
-}
-)");
-}
-
-////////////////////////////////////////////////////////////////////////////////
// chromium_experimental_subgroups
////////////////////////////////////////////////////////////////////////////////
TEST_F(IRToProgramTest, Enable_ChromiumExperimentalSubgroups_SubgroupBallot) {
diff --git a/test/tint/access/ptr.wgsl.expected.spvasm b/test/tint/access/ptr.wgsl.expected.spvasm
index c8f548f..0983663 100644
--- a/test/tint/access/ptr.wgsl.expected.spvasm
+++ b/test/tint/access/ptr.wgsl.expected.spvasm
@@ -17,18 +17,18 @@
OpName %v "v"
OpName %accept_value "accept_value"
OpName %val "val"
- OpName %accept_ptr_deref_call_func "accept_ptr_deref_call_func"
+ OpName %accept_ptr_deref_call_func_F "accept_ptr_deref_call_func_F"
OpName %val_0 "val"
- OpName %accept_ptr_deref_pass_through "accept_ptr_deref_pass_through"
+ OpName %accept_ptr_deref_pass_through_F "accept_ptr_deref_pass_through_F"
OpName %val_1 "val"
OpName %S "S"
OpMemberName %S 0 "a"
OpMemberName %S 1 "b"
- OpName %accept_ptr_to_struct_and_access "accept_ptr_to_struct_and_access"
+ OpName %accept_ptr_to_struct_and_access_F "accept_ptr_to_struct_and_access_F"
OpName %val_2 "val"
- OpName %accept_ptr_to_struct_access_pass_ptr "accept_ptr_to_struct_access_pass_ptr"
+ OpName %accept_ptr_to_struct_access_pass_ptr_F "accept_ptr_to_struct_access_pass_ptr_F"
OpName %val_3 "val"
- OpName %accept_ptr_vec_access_elements "accept_ptr_vec_access_elements"
+ OpName %accept_ptr_vec_access_elements_F "accept_ptr_vec_access_elements_F"
OpName %v1 "v1"
OpName %call_builtin_with_mod_scope_ptr "call_builtin_with_mod_scope_ptr"
OpName %main_inner "main_inner"
@@ -98,7 +98,7 @@
%28 = OpLabel
OpReturnValue %val
OpFunctionEnd
-%accept_ptr_deref_call_func = OpFunction %int None %29
+%accept_ptr_deref_call_func_F = OpFunction %int None %29
%val_0 = OpFunctionParameter %_ptr_Function_int
%33 = OpLabel
%35 = OpLoad %int %val_0
@@ -107,15 +107,15 @@
%39 = OpIAdd %int %35 %36
OpReturnValue %39
OpFunctionEnd
-%accept_ptr_deref_pass_through = OpFunction %int None %29
+%accept_ptr_deref_pass_through_F = OpFunction %int None %29
%val_1 = OpFunctionParameter %_ptr_Function_int
%42 = OpLabel
%44 = OpLoad %int %val_1
- %45 = OpFunctionCall %int %accept_ptr_deref_call_func %val_1
+ %45 = OpFunctionCall %int %accept_ptr_deref_call_func_F %val_1
%46 = OpIAdd %int %44 %45
OpReturnValue %46
OpFunctionEnd
-%accept_ptr_to_struct_and_access = OpFunction %int None %47
+%accept_ptr_to_struct_and_access_F = OpFunction %int None %47
%val_2 = OpFunctionParameter %_ptr_Function_S
%52 = OpLabel
%55 = OpAccessChain %_ptr_Function_int %val_2 %uint_0
@@ -125,7 +125,7 @@
%61 = OpIAdd %int %56 %60
OpReturnValue %61
OpFunctionEnd
-%accept_ptr_to_struct_access_pass_ptr = OpFunction %int None %47
+%accept_ptr_to_struct_access_pass_ptr_F = OpFunction %int None %47
%val_3 = OpFunctionParameter %_ptr_Function_S
%64 = OpLabel
%66 = OpAccessChain %_ptr_Function_int %val_3 %uint_0
@@ -134,7 +134,7 @@
%70 = OpLoad %int %69
OpReturnValue %70
OpFunctionEnd
-%accept_ptr_vec_access_elements = OpFunction %int None %71
+%accept_ptr_vec_access_elements_F = OpFunction %int None %71
%v1 = OpFunctionParameter %_ptr_Function_v3float
%76 = OpLabel
%79 = OpAccessChain %_ptr_Function_float %v1 %uint_0
@@ -165,11 +165,11 @@
OpStore %v2 %108
OpStore %v4 %110
%112 = OpAtomicLoad %int %g1 %uint_2 %uint_0
- %114 = OpFunctionCall %int %accept_ptr_deref_pass_through %v1_0
- %116 = OpFunctionCall %int %accept_ptr_to_struct_and_access %v2
- %118 = OpFunctionCall %int %accept_ptr_to_struct_and_access %v2
- %120 = OpFunctionCall %int %accept_ptr_vec_access_elements %v4
- %122 = OpFunctionCall %int %accept_ptr_to_struct_access_pass_ptr %v2
+ %114 = OpFunctionCall %int %accept_ptr_deref_pass_through_F %v1_0
+ %116 = OpFunctionCall %int %accept_ptr_to_struct_and_access_F %v2
+ %118 = OpFunctionCall %int %accept_ptr_to_struct_and_access_F %v2
+ %120 = OpFunctionCall %int %accept_ptr_vec_access_elements_F %v4
+ %122 = OpFunctionCall %int %accept_ptr_to_struct_access_pass_ptr_F %v2
%124 = OpFunctionCall %int %call_builtin_with_mod_scope_ptr
%126 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0
%127 = OpIAdd %int %114 %116
diff --git a/test/tint/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.spvasm b/test/tint/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.spvasm
index 9a4e1f9..e4e6032 100644
--- a/test/tint/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.spvasm
+++ b/test/tint/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.spvasm
@@ -19,7 +19,7 @@
OpName %result "result"
OpName %S "S"
OpMemberName %S 0 "data"
- OpName %x "x"
+ OpName %x_F "x_F"
OpName %p "p"
OpName %f "f"
OpName %s "s"
@@ -60,7 +60,7 @@
%32 = OpConstantNull %S
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
%int_3 = OpConstant %int 3
- %x = OpFunction %void None %10
+ %x_F = OpFunction %void None %10
%p = OpFunctionParameter %_ptr_Function_S
%19 = OpLabel
%23 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0 %uint_0
@@ -72,7 +72,7 @@
%f = OpFunction %void None %28
%30 = OpLabel
%s = OpVariable %_ptr_Function_S Function %32
- %33 = OpFunctionCall %void %x %s
+ %33 = OpFunctionCall %void %x_F %s
%36 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0 %uint_0
%38 = OpAccessChain %_ptr_Function_int %s %uint_0 %int_3
%39 = OpLoad %int %38
diff --git a/test/tint/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.spvasm b/test/tint/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.spvasm
index d461346..2ade876 100644
--- a/test/tint/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.spvasm
+++ b/test/tint/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.spvasm
@@ -20,7 +20,7 @@
OpName %S "S"
OpMemberName %S 0 "data"
OpName %s "s"
- OpName %x "x"
+ OpName %x_F "x_F"
OpName %p "p"
OpName %f "f"
OpDecorate %ubo_block Block
@@ -61,7 +61,7 @@
%30 = OpTypeFunction %void
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
%int_3 = OpConstant %int 3
- %x = OpFunction %void None %17
+ %x_F = OpFunction %void None %17
%p = OpFunctionParameter %_ptr_Private_S
%21 = OpLabel
%25 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0 %uint_0
@@ -72,7 +72,7 @@
OpFunctionEnd
%f = OpFunction %void None %30
%32 = OpLabel
- %33 = OpFunctionCall %void %x %s
+ %33 = OpFunctionCall %void %x_F %s
%36 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0 %uint_0
%38 = OpAccessChain %_ptr_Private_int %s %uint_0 %int_3
%39 = OpLoad %int %38
diff --git a/test/tint/bug/tint/1086.wgsl.expected.spvasm b/test/tint/bug/tint/1086.wgsl.expected.spvasm
index 2c300ad..9bba4e5 100644
--- a/test/tint/bug/tint/1086.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/1086.wgsl.expected.spvasm
@@ -8,7 +8,7 @@
OpEntryPoint Fragment %f "f"
OpExecutionMode %f OriginUpperLeft
OpName %v "v"
- OpName %x "x"
+ OpName %x_F "x_F"
OpName %p "p"
OpName %g "g"
OpName %f "f"
@@ -19,7 +19,7 @@
%void = OpTypeVoid
%5 = OpTypeFunction %void %_ptr_Private_float
%11 = OpTypeFunction %void
- %x = OpFunction %void None %5
+ %x_F = OpFunction %void None %5
%p = OpFunctionParameter %_ptr_Private_float
%9 = OpLabel
OpStore %p %4
@@ -27,7 +27,7 @@
OpFunctionEnd
%g = OpFunction %void None %11
%13 = OpLabel
- %14 = OpFunctionCall %void %x %v
+ %14 = OpFunctionCall %void %x_F %v
OpReturn
OpFunctionEnd
%f = OpFunction %void None %11
diff --git a/test/tint/bug/tint/2054.wgsl.expected.spvasm b/test/tint/bug/tint/2054.wgsl.expected.spvasm
index 2c6e927..d3100e7 100644
--- a/test/tint/bug/tint/2054.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/2054.wgsl.expected.spvasm
@@ -10,7 +10,7 @@
OpName %out_block "out_block"
OpMemberName %out_block 0 "inner"
OpName %out "out"
- OpName %bar "bar"
+ OpName %bar_F "bar_F"
OpName %p "p"
OpName %foo "foo"
OpName %param "param"
@@ -33,7 +33,7 @@
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
- %bar = OpFunction %void None %5
+ %bar_F = OpFunction %void None %5
%p = OpFunctionParameter %_ptr_Function_float
%10 = OpLabel
%14 = OpFOrdGreaterThanEqual %bool %float_1 %13
@@ -51,7 +51,7 @@
%foo = OpFunction %void None %22
%24 = OpLabel
%param = OpVariable %_ptr_Function_float Function %13
- %26 = OpFunctionCall %void %bar %param
+ %26 = OpFunctionCall %void %bar_F %param
%31 = OpAccessChain %_ptr_StorageBuffer_float %out %uint_0
%32 = OpLoad %float %param
OpStore %31 %32
diff --git a/test/tint/bug/tint/2121.wgsl.expected.spvasm b/test/tint/bug/tint/2121.wgsl.expected.spvasm
index f0651b9..817c5d7 100644
--- a/test/tint/bug/tint/2121.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/2121.wgsl.expected.spvasm
@@ -10,7 +10,7 @@
OpName %vertex_point_size "vertex_point_size"
OpName %VSOut "VSOut"
OpMemberName %VSOut 0 "pos"
- OpName %foo "foo"
+ OpName %foo_F "foo_F"
OpName %out "out"
OpName %pos "pos"
OpName %main_inner "main_inner"
@@ -42,7 +42,7 @@
%28 = OpTypeFunction %VSOut
%32 = OpConstantNull %VSOut
%36 = OpTypeFunction %void
- %foo = OpFunction %void None %9
+ %foo_F = OpFunction %void None %9
%out = OpFunctionParameter %_ptr_Function_VSOut
%15 = OpLabel
%pos = OpVariable %_ptr_Function_v4float Function %5
@@ -55,7 +55,7 @@
%main_inner = OpFunction %VSOut None %28
%30 = OpLabel
%out_0 = OpVariable %_ptr_Function_VSOut Function %32
- %33 = OpFunctionCall %void %foo %out_0
+ %33 = OpFunctionCall %void %foo_F %out_0
%35 = OpLoad %VSOut %out_0
OpReturnValue %35
OpFunctionEnd
diff --git a/test/tint/bug/tint/219.spvasm.expected.spvasm b/test/tint/bug/tint/219.spvasm.expected.spvasm
index 1725762..64431cb 100644
--- a/test/tint/bug/tint/219.spvasm.expected.spvasm
+++ b/test/tint/bug/tint/219.spvasm.expected.spvasm
@@ -7,7 +7,7 @@
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
- OpName %x_200 "x_200"
+ OpName %x_200_F "x_200_F"
OpName %x_201 "x_201"
OpName %main_1 "main_1"
OpName %x_11 "x_11"
@@ -22,7 +22,7 @@
%void = OpTypeVoid
%14 = OpTypeFunction %void
%19 = OpConstantNull %v2float
- %x_200 = OpFunction %float None %1
+ %x_200_F = OpFunction %float None %1
%x_201 = OpFunctionParameter %_ptr_Function_v2float
%7 = OpLabel
%12 = OpAccessChain %_ptr_Function_float %x_201 %uint_0
@@ -32,7 +32,7 @@
%main_1 = OpFunction %void None %14
%17 = OpLabel
%x_11 = OpVariable %_ptr_Function_v2float Function %19
- %20 = OpFunctionCall %float %x_200 %x_11
+ %20 = OpFunctionCall %float %x_200_F %x_11
OpReturn
OpFunctionEnd
%main = OpFunction %void None %14
diff --git a/test/tint/bug/tint/948.wgsl.expected.spvasm b/test/tint/bug/tint/948.wgsl.expected.spvasm
index 55c8b99..207c6fd 100644
--- a/test/tint/bug/tint/948.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/948.wgsl.expected.spvasm
@@ -45,7 +45,7 @@
OpName %stageUnits_1 "stageUnits_1"
OpName %vPosition "vPosition"
OpName %vUV "vUV"
- OpName %getFrameData_f1_ "getFrameData_f1_"
+ OpName %getFrameData_f1__F "getFrameData_f1__F"
OpName %frameID "frameID"
OpName %fX "fX"
OpName %main_1 "main_1"
@@ -203,7 +203,7 @@
%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
%main_out = OpTypeStruct %v4float
%347 = OpTypeFunction %main_out %v2float %v2float %v2float %v2float %v3float %v2float
-%getFrameData_f1_ = OpFunction %mat4v4float None %50
+%getFrameData_f1__F = OpFunction %mat4v4float None %50
%frameID = OpFunctionParameter %_ptr_Function_float
%54 = OpLabel
%fX = OpVariable %_ptr_Function_float Function %38
@@ -407,7 +407,7 @@
%238 = OpLoad %float %frameID_1
%239 = OpFAdd %float %238 %float_0_5
OpStore %param %239
- %240 = OpFunctionCall %mat4v4float %getFrameData_f1_ %param
+ %240 = OpFunctionCall %mat4v4float %getFrameData_f1__F %param
OpStore %frameData %240
%242 = OpAccessChain %_ptr_Function_v4float %frameData %116
%243 = OpLoad %v4float %242
diff --git a/test/tint/bug/tint/949.wgsl.expected.spvasm b/test/tint/bug/tint/949.wgsl.expected.spvasm
index fc44b4f..2f1a75b 100644
--- a/test/tint/bug/tint/949.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/949.wgsl.expected.spvasm
@@ -52,7 +52,7 @@
OpName %glFragColor "glFragColor"
OpName %bumpSamplerSampler "bumpSamplerSampler"
OpName %bumpSamplerTexture "bumpSamplerTexture"
- OpName %cotangent_frame_vf3_vf3_vf2_vf2_ "cotangent_frame_vf3_vf3_vf2_vf2_"
+ OpName %cotangent_frame_vf3_vf3_vf2_vf2__F_F_F_F "cotangent_frame_vf3_vf3_vf2_vf2__F_F_F_F"
OpName %normal_1 "normal_1"
OpName %p "p"
OpName %uv "uv"
@@ -66,17 +66,17 @@
OpName %tangent "tangent"
OpName %bitangent "bitangent"
OpName %invmax "invmax"
- OpName %transposeMat3_mf33_ "transposeMat3_mf33_"
+ OpName %transposeMat3_mf33__F "transposeMat3_mf33__F"
OpName %inMatrix "inMatrix"
OpName %i0 "i0"
OpName %i1 "i1"
OpName %i2 "i2"
OpName %outMatrix "outMatrix"
- OpName %perturbNormalBase_mf33_vf3_f1_ "perturbNormalBase_mf33_vf3_f1_"
+ OpName %perturbNormalBase_mf33_vf3_f1__F_F_F "perturbNormalBase_mf33_vf3_f1__F_F_F"
OpName %cotangentFrame "cotangentFrame"
OpName %normal "normal"
OpName %scale "scale"
- OpName %perturbNormal_mf33_vf3_f1_ "perturbNormal_mf33_vf3_f1_"
+ OpName %perturbNormal_mf33_vf3_f1__F_F_F "perturbNormal_mf33_vf3_f1__F_F_F"
OpName %cotangentFrame_1 "cotangentFrame_1"
OpName %textureSample "textureSample"
OpName %scale_1 "scale_1"
@@ -86,7 +86,7 @@
OpName %lightingInfo "lightingInfo"
OpMemberName %lightingInfo 0 "diffuse"
OpMemberName %lightingInfo 1 "specular"
- OpName %computeHemisphericLighting_vf3_vf3_vf4_vf3_vf3_vf3_f1_ "computeHemisphericLighting_vf3_vf3_vf4_vf3_vf3_vf3_f1_"
+ OpName %computeHemisphericLighting_vf3_vf3_vf4_vf3_vf3_vf3_f1__F_F_F_F_F_F_F "computeHemisphericLighting_vf3_vf3_vf4_vf3_vf3_vf3_f1__F_F_F_F_F_F_F"
OpName %viewDirectionW "viewDirectionW"
OpName %vNormal "vNormal"
OpName %lightData "lightData"
@@ -310,7 +310,7 @@
%uint_3 = OpConstant %uint 3
%main_out = OpTypeStruct %v4float
%646 = OpTypeFunction %main_out %v2float %v4float %bool %v2float %v4float
-%cotangent_frame_vf3_vf3_vf2_vf2_ = OpFunction %mat3v3float None %54
+%cotangent_frame_vf3_vf3_vf2_vf2__F_F_F_F = OpFunction %mat3v3float None %54
%normal_1 = OpFunctionParameter %_ptr_Function_v3float
%p = OpFunctionParameter %_ptr_Function_v3float
%uv = OpFunctionParameter %_ptr_Function_v2float
@@ -406,7 +406,7 @@
%153 = OpCompositeConstruct %mat3v3float %144 %148 %152
OpReturnValue %153
OpFunctionEnd
-%transposeMat3_mf33_ = OpFunction %mat3v3float None %154
+%transposeMat3_mf33__F = OpFunction %mat3v3float None %154
%inMatrix = OpFunctionParameter %_ptr_Function_mat3v3float
%158 = OpLabel
%i0 = OpVariable %_ptr_Function_v3float Function %22
@@ -460,7 +460,7 @@
%212 = OpLoad %mat3v3float %outMatrix
OpReturnValue %212
OpFunctionEnd
-%perturbNormalBase_mf33_vf3_f1_ = OpFunction %v3float None %213
+%perturbNormalBase_mf33_vf3_f1__F_F_F = OpFunction %v3float None %213
%cotangentFrame = OpFunctionParameter %_ptr_Function_mat3v3float
%normal = OpFunctionParameter %_ptr_Function_v3float
%scale = OpFunctionParameter %_ptr_Function_float
@@ -471,7 +471,7 @@
%223 = OpExtInst %v3float %90 Normalize %224
OpReturnValue %223
OpFunctionEnd
-%perturbNormal_mf33_vf3_f1_ = OpFunction %v3float None %213
+%perturbNormal_mf33_vf3_f1__F_F_F = OpFunction %v3float None %213
%cotangentFrame_1 = OpFunctionParameter %_ptr_Function_mat3v3float
%textureSample = OpFunctionParameter %_ptr_Function_v3float
%scale_1 = OpFunctionParameter %_ptr_Function_float
@@ -487,10 +487,10 @@
OpStore %param_1 %241
%243 = OpLoad %float %scale_1
OpStore %param_2 %243
- %244 = OpFunctionCall %v3float %perturbNormalBase_mf33_vf3_f1_ %param %param_1 %param_2
+ %244 = OpFunctionCall %v3float %perturbNormalBase_mf33_vf3_f1__F_F_F %param %param_1 %param_2
OpReturnValue %244
OpFunctionEnd
-%computeHemisphericLighting_vf3_vf3_vf4_vf3_vf3_vf3_f1_ = OpFunction %lightingInfo None %248
+%computeHemisphericLighting_vf3_vf3_vf4_vf3_vf3_vf3_f1__F_F_F_F_F_F_F = OpFunction %lightingInfo None %248
%viewDirectionW = OpFunctionParameter %_ptr_Function_v3float
%vNormal = OpFunctionParameter %_ptr_Function_v3float
%lightData = OpFunctionParameter %_ptr_Function_v4float
@@ -669,11 +669,11 @@
%425 = OpAccessChain %_ptr_Uniform_v2float %x_269 %uint_0 %uint_8
%426 = OpLoad %v2float %425
OpStore %param_6 %426
- %427 = OpFunctionCall %mat3v3float %cotangent_frame_vf3_vf3_vf2_vf2_ %param_3 %param_4 %param_5 %param_6
+ %427 = OpFunctionCall %mat3v3float %cotangent_frame_vf3_vf3_vf2_vf2__F_F_F_F %param_3 %param_4 %param_5 %param_6
OpStore %TBN %427
%432 = OpLoad %mat3v3float %TBN
OpStore %param_7 %432
- %433 = OpFunctionCall %mat3v3float %transposeMat3_mf33_ %param_7
+ %433 = OpFunctionCall %mat3v3float %transposeMat3_mf33__F %param_7
OpStore %invTBN %433
%435 = OpLoad %mat3v3float %invTBN
%436 = OpLoad %v3float %output5
@@ -828,7 +828,7 @@
OpStore %param_9 %550
%551 = OpFDiv %float %float_1 %545
OpStore %param_10 %551
- %552 = OpFunctionCall %v3float %perturbNormal_mf33_vf3_f1_ %param_8 %param_9 %param_10
+ %552 = OpFunctionCall %v3float %perturbNormal_mf33_vf3_f1__F_F_F %param_8 %param_9 %param_10
%556 = OpLoad %v4float %output4
%557 = OpCompositeExtract %float %552 0
%558 = OpCompositeExtract %float %552 1
@@ -900,7 +900,7 @@
OpStore %param_16 %610
%611 = OpLoad %float %glossiness_1
OpStore %param_17 %611
- %612 = OpFunctionCall %lightingInfo %computeHemisphericLighting_vf3_vf3_vf4_vf3_vf3_vf3_f1_ %param_11 %param_12 %param_13 %param_14 %param_15 %param_16 %param_17
+ %612 = OpFunctionCall %lightingInfo %computeHemisphericLighting_vf3_vf3_vf4_vf3_vf3_vf3_f1__F_F_F_F_F_F_F %param_11 %param_12 %param_13 %param_14 %param_15 %param_16 %param_17
OpStore %info %612
OpStore %shadow %float_1
%620 = OpAccessChain %_ptr_Function_v3float %info %uint_0
diff --git a/test/tint/bug/tint/977.spvasm.expected.spvasm b/test/tint/bug/tint/977.spvasm.expected.spvasm
index 5d7f451..3c30bb1 100644
--- a/test/tint/bug/tint/977.spvasm.expected.spvasm
+++ b/test/tint/bug/tint/977.spvasm.expected.spvasm
@@ -26,7 +26,7 @@
OpMemberName %Uniforms 1 "sizeA"
OpMemberName %Uniforms 2 "sizeB"
OpName %x_46 "x_46"
- OpName %binaryOperation_f1_f1_ "binaryOperation_f1_f1_"
+ OpName %binaryOperation_f1_f1__F_F "binaryOperation_f1_f1__F_F"
OpName %a "a"
OpName %b "b"
OpName %tint_return_flag "tint_return_flag"
@@ -107,7 +107,7 @@
%float_n3 = OpConstant %float -3
%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
%104 = OpTypeFunction %void %v3uint
-%binaryOperation_f1_f1_ = OpFunction %float None %24
+%binaryOperation_f1_f1__F_F = OpFunction %float None %24
%a = OpFunctionParameter %_ptr_Function_float
%b = OpFunctionParameter %_ptr_Function_float
%29 = OpLabel
@@ -178,7 +178,7 @@
%96 = OpLoad %int %index
OpStore %param %float_n4
OpStore %param_1 %float_n3
- %99 = OpFunctionCall %float %binaryOperation_f1_f1_ %param %param_1
+ %99 = OpFunctionCall %float %binaryOperation_f1_f1__F_F %param %param_1
%103 = OpAccessChain %_ptr_StorageBuffer_float %resultMatrix %uint_0 %96
OpStore %103 %99
OpReturn
diff --git a/test/tint/builtins/workgroupUniformLoad/via_param.wgsl b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl
index 59f5e8b..f435ed6 100644
--- a/test/tint/builtins/workgroupUniformLoad/via_param.wgsl
+++ b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> v : array<i32, 4>;
fn foo(p : ptr<workgroup, i32>) -> i32 {
diff --git a/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.wgsl b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.wgsl
index 59f5e8b..f435ed6 100644
--- a/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.wgsl
+++ b/test/tint/builtins/workgroupUniformLoad/via_param.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> v : array<i32, 4>;
fn foo(p : ptr<workgroup, i32>) -> i32 {
diff --git a/test/tint/extensions/parsing/multiple.wgsl b/test/tint/extensions/parsing/multiple.wgsl
index 3975c3d..1fb396f 100644
--- a/test/tint/extensions/parsing/multiple.wgsl
+++ b/test/tint/extensions/parsing/multiple.wgsl
@@ -1,5 +1,5 @@
// flags: --hlsl_shader_model 62
-enable chromium_experimental_full_ptr_parameters, f16;
+enable chromium_disable_uniformity_analysis, f16;
@fragment
fn main() -> @location(0) vec4<f32> {
diff --git a/test/tint/extensions/parsing/multiple.wgsl.expected.wgsl b/test/tint/extensions/parsing/multiple.wgsl.expected.wgsl
index 61f22c2..df77540 100644
--- a/test/tint/extensions/parsing/multiple.wgsl.expected.wgsl
+++ b/test/tint/extensions/parsing/multiple.wgsl.expected.wgsl
@@ -1,4 +1,4 @@
-enable chromium_experimental_full_ptr_parameters, f16;
+enable chromium_disable_uniformity_analysis, f16;
@fragment
fn main() -> @location(0) vec4<f32> {
diff --git a/test/tint/ptr_ref/load/param/function/array_in_struct.wgsl b/test/tint/ptr_ref/load/param/function/array_in_struct.wgsl
index 4053c2a..6e2cff6 100644
--- a/test/tint/ptr_ref/load/param/function/array_in_struct.wgsl
+++ b/test/tint/ptr_ref/load/param/function/array_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
};
diff --git a/test/tint/ptr_ref/load/param/function/array_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/function/array_in_struct.wgsl.expected.wgsl
index 8bd5276..4f77626 100644
--- a/test/tint/ptr_ref/load/param/function/array_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/function/array_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
}
diff --git a/test/tint/ptr_ref/load/param/function/i32.wgsl.expected.spvasm b/test/tint/ptr_ref/load/param/function/i32.wgsl.expected.spvasm
index d25f603..9a0d38f 100644
--- a/test/tint/ptr_ref/load/param/function/i32.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/load/param/function/i32.wgsl.expected.spvasm
@@ -7,7 +7,7 @@
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
- OpName %func "func"
+ OpName %func_F "func_F"
OpName %pointer "pointer"
OpName %main "main"
OpName %F "F"
@@ -17,7 +17,7 @@
%void = OpTypeVoid
%9 = OpTypeFunction %void
%14 = OpConstantNull %int
- %func = OpFunction %int None %1
+ %func_F = OpFunction %int None %1
%pointer = OpFunctionParameter %_ptr_Function_int
%6 = OpLabel
%8 = OpLoad %int %pointer
@@ -26,6 +26,6 @@
%main = OpFunction %void None %9
%12 = OpLabel
%F = OpVariable %_ptr_Function_int Function %14
- %15 = OpFunctionCall %int %func %F
+ %15 = OpFunctionCall %int %func_F %F
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/load/param/function/i32_in_struct.wgsl b/test/tint/ptr_ref/load/param/function/i32_in_struct.wgsl
index c3b7d55..6f40512 100644
--- a/test/tint/ptr_ref/load/param/function/i32_in_struct.wgsl
+++ b/test/tint/ptr_ref/load/param/function/i32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/load/param/function/i32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/function/i32_in_struct.wgsl.expected.wgsl
index e5ecd70..04a0bb4 100644
--- a/test/tint/ptr_ref/load/param/function/i32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/function/i32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/load/param/function/struct_in_array.wgsl b/test/tint/ptr_ref/load/param/function/struct_in_array.wgsl
index 103a3dd..7c87787 100644
--- a/test/tint/ptr_ref/load/param/function/struct_in_array.wgsl
+++ b/test/tint/ptr_ref/load/param/function/struct_in_array.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/load/param/function/struct_in_array.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/function/struct_in_array.wgsl.expected.wgsl
index 2d4edb3..74594bf 100644
--- a/test/tint/ptr_ref/load/param/function/struct_in_array.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/function/struct_in_array.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/load/param/function/vec2_f32_in_mat2x2.wgsl b/test/tint/ptr_ref/load/param/function/vec2_f32_in_mat2x2.wgsl
index afe82bc..c91b9da 100644
--- a/test/tint/ptr_ref/load/param/function/vec2_f32_in_mat2x2.wgsl
+++ b/test/tint/ptr_ref/load/param/function/vec2_f32_in_mat2x2.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<function, vec2<f32>>) -> vec2<f32> {
return *pointer;
}
diff --git a/test/tint/ptr_ref/load/param/function/vec2_f32_in_mat2x2.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/function/vec2_f32_in_mat2x2.wgsl.expected.wgsl
index 2c61e24..1aff761 100644
--- a/test/tint/ptr_ref/load/param/function/vec2_f32_in_mat2x2.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/function/vec2_f32_in_mat2x2.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<function, vec2<f32>>) -> vec2<f32> {
return *(pointer);
}
diff --git a/test/tint/ptr_ref/load/param/function/vec4_f32.wgsl.expected.spvasm b/test/tint/ptr_ref/load/param/function/vec4_f32.wgsl.expected.spvasm
index 51c798e..c7d7be4 100644
--- a/test/tint/ptr_ref/load/param/function/vec4_f32.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/load/param/function/vec4_f32.wgsl.expected.spvasm
@@ -7,7 +7,7 @@
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
- OpName %func "func"
+ OpName %func_F "func_F"
OpName %pointer "pointer"
OpName %main "main"
OpName %F "F"
@@ -18,7 +18,7 @@
%void = OpTypeVoid
%10 = OpTypeFunction %void
%15 = OpConstantNull %v4float
- %func = OpFunction %v4float None %1
+ %func_F = OpFunction %v4float None %1
%pointer = OpFunctionParameter %_ptr_Function_v4float
%7 = OpLabel
%9 = OpLoad %v4float %pointer
@@ -27,6 +27,6 @@
%main = OpFunction %void None %10
%13 = OpLabel
%F = OpVariable %_ptr_Function_v4float Function %15
- %16 = OpFunctionCall %v4float %func %F
+ %16 = OpFunctionCall %v4float %func_F %F
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/load/param/function/vec4_f32_in_mat2x4.wgsl b/test/tint/ptr_ref/load/param/function/vec4_f32_in_mat2x4.wgsl
index 8b95569..b65f9ba 100644
--- a/test/tint/ptr_ref/load/param/function/vec4_f32_in_mat2x4.wgsl
+++ b/test/tint/ptr_ref/load/param/function/vec4_f32_in_mat2x4.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<function, vec4<f32>>) -> vec4<f32> {
return *pointer;
}
diff --git a/test/tint/ptr_ref/load/param/function/vec4_f32_in_mat2x4.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/function/vec4_f32_in_mat2x4.wgsl.expected.wgsl
index e5a91d7..9e74cd5 100644
--- a/test/tint/ptr_ref/load/param/function/vec4_f32_in_mat2x4.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/function/vec4_f32_in_mat2x4.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<function, vec4<f32>>) -> vec4<f32> {
return *(pointer);
}
diff --git a/test/tint/ptr_ref/load/param/function/vec4_f32_in_struct.wgsl b/test/tint/ptr_ref/load/param/function/vec4_f32_in_struct.wgsl
index f1167ac..69682cf 100644
--- a/test/tint/ptr_ref/load/param/function/vec4_f32_in_struct.wgsl
+++ b/test/tint/ptr_ref/load/param/function/vec4_f32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
};
diff --git a/test/tint/ptr_ref/load/param/function/vec4_f32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/function/vec4_f32_in_struct.wgsl.expected.wgsl
index f35fc26..1f3a783 100644
--- a/test/tint/ptr_ref/load/param/function/vec4_f32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/function/vec4_f32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
}
diff --git a/test/tint/ptr_ref/load/param/private/array_in_struct.wgsl b/test/tint/ptr_ref/load/param/private/array_in_struct.wgsl
index 867e62b..8ce3524 100644
--- a/test/tint/ptr_ref/load/param/private/array_in_struct.wgsl
+++ b/test/tint/ptr_ref/load/param/private/array_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
};
diff --git a/test/tint/ptr_ref/load/param/private/array_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/private/array_in_struct.wgsl.expected.wgsl
index 1db57bb..deda9b4 100644
--- a/test/tint/ptr_ref/load/param/private/array_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/private/array_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
}
diff --git a/test/tint/ptr_ref/load/param/private/i32.wgsl.expected.spvasm b/test/tint/ptr_ref/load/param/private/i32.wgsl.expected.spvasm
index 66b4e2f..f328a08 100644
--- a/test/tint/ptr_ref/load/param/private/i32.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/load/param/private/i32.wgsl.expected.spvasm
@@ -8,7 +8,7 @@
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpName %P "P"
- OpName %func "func"
+ OpName %func_F "func_F"
OpName %pointer "pointer"
OpName %main "main"
%int = OpTypeInt 32 1
@@ -18,7 +18,7 @@
%5 = OpTypeFunction %int %_ptr_Private_int
%void = OpTypeVoid
%11 = OpTypeFunction %void
- %func = OpFunction %int None %5
+ %func_F = OpFunction %int None %5
%pointer = OpFunctionParameter %_ptr_Private_int
%8 = OpLabel
%10 = OpLoad %int %pointer
@@ -26,6 +26,6 @@
OpFunctionEnd
%main = OpFunction %void None %11
%14 = OpLabel
- %15 = OpFunctionCall %int %func %P
+ %15 = OpFunctionCall %int %func_F %P
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/load/param/private/i32_in_struct.wgsl b/test/tint/ptr_ref/load/param/private/i32_in_struct.wgsl
index 3db15dc..f48aab5 100644
--- a/test/tint/ptr_ref/load/param/private/i32_in_struct.wgsl
+++ b/test/tint/ptr_ref/load/param/private/i32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/load/param/private/i32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/private/i32_in_struct.wgsl.expected.wgsl
index bda048a..5be3925 100644
--- a/test/tint/ptr_ref/load/param/private/i32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/private/i32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/load/param/private/struct_in_array.wgsl b/test/tint/ptr_ref/load/param/private/struct_in_array.wgsl
index b7b8f90..193c2dc 100644
--- a/test/tint/ptr_ref/load/param/private/struct_in_array.wgsl
+++ b/test/tint/ptr_ref/load/param/private/struct_in_array.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/load/param/private/struct_in_array.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/private/struct_in_array.wgsl.expected.wgsl
index 839270e..847d132 100644
--- a/test/tint/ptr_ref/load/param/private/struct_in_array.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/private/struct_in_array.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/load/param/private/vec2_f32_in_mat2x2.wgsl b/test/tint/ptr_ref/load/param/private/vec2_f32_in_mat2x2.wgsl
index 7dc8a0b..92578b9 100644
--- a/test/tint/ptr_ref/load/param/private/vec2_f32_in_mat2x2.wgsl
+++ b/test/tint/ptr_ref/load/param/private/vec2_f32_in_mat2x2.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<private, vec2<f32>>) -> vec2<f32> {
return *pointer;
}
diff --git a/test/tint/ptr_ref/load/param/private/vec2_f32_in_mat2x2.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/private/vec2_f32_in_mat2x2.wgsl.expected.wgsl
index dbc825a..728776c 100644
--- a/test/tint/ptr_ref/load/param/private/vec2_f32_in_mat2x2.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/private/vec2_f32_in_mat2x2.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<private, vec2<f32>>) -> vec2<f32> {
return *(pointer);
}
diff --git a/test/tint/ptr_ref/load/param/private/vec4_f32.wgsl.expected.spvasm b/test/tint/ptr_ref/load/param/private/vec4_f32.wgsl.expected.spvasm
index dfe4930..2c19ccf 100644
--- a/test/tint/ptr_ref/load/param/private/vec4_f32.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/load/param/private/vec4_f32.wgsl.expected.spvasm
@@ -8,7 +8,7 @@
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpName %P "P"
- OpName %func "func"
+ OpName %func_F "func_F"
OpName %pointer "pointer"
OpName %main "main"
%float = OpTypeFloat 32
@@ -19,7 +19,7 @@
%6 = OpTypeFunction %v4float %_ptr_Private_v4float
%void = OpTypeVoid
%12 = OpTypeFunction %void
- %func = OpFunction %v4float None %6
+ %func_F = OpFunction %v4float None %6
%pointer = OpFunctionParameter %_ptr_Private_v4float
%9 = OpLabel
%11 = OpLoad %v4float %pointer
@@ -27,6 +27,6 @@
OpFunctionEnd
%main = OpFunction %void None %12
%15 = OpLabel
- %16 = OpFunctionCall %v4float %func %P
+ %16 = OpFunctionCall %v4float %func_F %P
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/load/param/private/vec4_f32_in_mat2x4.wgsl b/test/tint/ptr_ref/load/param/private/vec4_f32_in_mat2x4.wgsl
index 07dfea9..fc8a503 100644
--- a/test/tint/ptr_ref/load/param/private/vec4_f32_in_mat2x4.wgsl
+++ b/test/tint/ptr_ref/load/param/private/vec4_f32_in_mat2x4.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<private, vec4<f32>>) -> vec4<f32> {
return *pointer;
}
diff --git a/test/tint/ptr_ref/load/param/private/vec4_f32_in_mat2x4.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/private/vec4_f32_in_mat2x4.wgsl.expected.wgsl
index c0cf801..2d03c8e 100644
--- a/test/tint/ptr_ref/load/param/private/vec4_f32_in_mat2x4.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/private/vec4_f32_in_mat2x4.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<private, vec4<f32>>) -> vec4<f32> {
return *(pointer);
}
diff --git a/test/tint/ptr_ref/load/param/private/vec4_f32_in_struct.wgsl b/test/tint/ptr_ref/load/param/private/vec4_f32_in_struct.wgsl
index 90f325e..d38884c 100644
--- a/test/tint/ptr_ref/load/param/private/vec4_f32_in_struct.wgsl
+++ b/test/tint/ptr_ref/load/param/private/vec4_f32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
};
diff --git a/test/tint/ptr_ref/load/param/private/vec4_f32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/private/vec4_f32_in_struct.wgsl.expected.wgsl
index 2c488a9..0365630 100644
--- a/test/tint/ptr_ref/load/param/private/vec4_f32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/private/vec4_f32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
}
diff --git a/test/tint/ptr_ref/load/param/ptr.spvasm.expected.spvasm b/test/tint/ptr_ref/load/param/ptr.spvasm.expected.spvasm
index 47d6500..d0e55b1 100644
--- a/test/tint/ptr_ref/load/param/ptr.spvasm.expected.spvasm
+++ b/test/tint/ptr_ref/load/param/ptr.spvasm.expected.spvasm
@@ -7,7 +7,7 @@
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
- OpName %func "func"
+ OpName %func_F "func_F"
OpName %value "value"
OpName %pointer "pointer"
OpName %main_1 "main_1"
@@ -20,7 +20,7 @@
%11 = OpTypeFunction %void
%15 = OpConstantNull %int
%int_123 = OpConstant %int 123
- %func = OpFunction %int None %1
+ %func_F = OpFunction %int None %1
%value = OpFunctionParameter %int
%pointer = OpFunctionParameter %_ptr_Function_int
%7 = OpLabel
@@ -34,7 +34,7 @@
OpStore %i %15
OpStore %i %int_123
%18 = OpLoad %int %i
- %19 = OpFunctionCall %int %func %18 %i
+ %19 = OpFunctionCall %int %func_F %18 %i
OpReturn
OpFunctionEnd
%main = OpFunction %void None %11
diff --git a/test/tint/ptr_ref/load/param/ptr.wgsl.expected.spvasm b/test/tint/ptr_ref/load/param/ptr.wgsl.expected.spvasm
index bf904d1..7d61a31 100644
--- a/test/tint/ptr_ref/load/param/ptr.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/load/param/ptr.wgsl.expected.spvasm
@@ -7,7 +7,7 @@
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
- OpName %func "func"
+ OpName %func_F "func_F"
OpName %value "value"
OpName %pointer "pointer"
OpName %main "main"
@@ -19,7 +19,7 @@
%11 = OpTypeFunction %void
%int_123 = OpConstant %int 123
%17 = OpConstantNull %int
- %func = OpFunction %int None %1
+ %func_F = OpFunction %int None %1
%value = OpFunctionParameter %int
%pointer = OpFunctionParameter %_ptr_Function_int
%7 = OpLabel
@@ -32,6 +32,6 @@
%i = OpVariable %_ptr_Function_int Function %17
OpStore %i %int_123
%19 = OpLoad %int %i
- %18 = OpFunctionCall %int %func %19 %i
+ %18 = OpFunctionCall %int %func_F %19 %i
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/load/param/storage/array_in_struct.wgsl b/test/tint/ptr_ref/load/param/storage/array_in_struct.wgsl
index eed9668..a393587 100644
--- a/test/tint/ptr_ref/load/param/storage/array_in_struct.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/array_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
};
diff --git a/test/tint/ptr_ref/load/param/storage/array_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/storage/array_in_struct.wgsl.expected.wgsl
index 033d9b1..80789dd 100644
--- a/test/tint/ptr_ref/load/param/storage/array_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/array_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
}
diff --git a/test/tint/ptr_ref/load/param/storage/i32.wgsl b/test/tint/ptr_ref/load/param/storage/i32.wgsl
index ac2a585..341081e 100644
--- a/test/tint/ptr_ref/load/param/storage/i32.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/i32.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : i32;
fn func(pointer : ptr<storage, i32>) -> i32 {
diff --git a/test/tint/ptr_ref/load/param/storage/i32.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/storage/i32.wgsl.expected.wgsl
index b9a422c..955c304 100644
--- a/test/tint/ptr_ref/load/param/storage/i32.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/i32.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : i32;
fn func(pointer : ptr<storage, i32>) -> i32 {
diff --git a/test/tint/ptr_ref/load/param/storage/i32_in_struct.wgsl b/test/tint/ptr_ref/load/param/storage/i32_in_struct.wgsl
index 05631d2..8edbb8f 100644
--- a/test/tint/ptr_ref/load/param/storage/i32_in_struct.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/i32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/load/param/storage/i32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/storage/i32_in_struct.wgsl.expected.wgsl
index 21763e0..7cb41e5 100644
--- a/test/tint/ptr_ref/load/param/storage/i32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/i32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/load/param/storage/struct_in_array.wgsl b/test/tint/ptr_ref/load/param/storage/struct_in_array.wgsl
index 5616c9a..c97e3ab 100644
--- a/test/tint/ptr_ref/load/param/storage/struct_in_array.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/struct_in_array.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/load/param/storage/struct_in_array.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/storage/struct_in_array.wgsl.expected.wgsl
index e16f195..c188052 100644
--- a/test/tint/ptr_ref/load/param/storage/struct_in_array.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/struct_in_array.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/load/param/storage/vec2_f32_in_mat2x2.wgsl b/test/tint/ptr_ref/load/param/storage/vec2_f32_in_mat2x2.wgsl
index 180f5ce..9ce0fa9 100644
--- a/test/tint/ptr_ref/load/param/storage/vec2_f32_in_mat2x2.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/vec2_f32_in_mat2x2.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : mat2x2<f32>;
fn func(pointer : ptr<storage, vec2<f32>>) -> vec2<f32> {
diff --git a/test/tint/ptr_ref/load/param/storage/vec2_f32_in_mat2x2.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/storage/vec2_f32_in_mat2x2.wgsl.expected.wgsl
index 889c31b..06c8846 100644
--- a/test/tint/ptr_ref/load/param/storage/vec2_f32_in_mat2x2.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/vec2_f32_in_mat2x2.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : mat2x2<f32>;
fn func(pointer : ptr<storage, vec2<f32>>) -> vec2<f32> {
diff --git a/test/tint/ptr_ref/load/param/storage/vec4_f32.wgsl b/test/tint/ptr_ref/load/param/storage/vec4_f32.wgsl
index 2a08b41..d13e602 100644
--- a/test/tint/ptr_ref/load/param/storage/vec4_f32.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/vec4_f32.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : vec4<f32>;
fn func(pointer : ptr<storage, vec4<f32>>) -> vec4<f32> {
diff --git a/test/tint/ptr_ref/load/param/storage/vec4_f32.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/storage/vec4_f32.wgsl.expected.wgsl
index a1bcaca..3d74aba 100644
--- a/test/tint/ptr_ref/load/param/storage/vec4_f32.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/vec4_f32.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : vec4<f32>;
fn func(pointer : ptr<storage, vec4<f32>>) -> vec4<f32> {
diff --git a/test/tint/ptr_ref/load/param/storage/vec4_f32_in_mat2x4.wgsl b/test/tint/ptr_ref/load/param/storage/vec4_f32_in_mat2x4.wgsl
index 35ef536..e992bb3 100644
--- a/test/tint/ptr_ref/load/param/storage/vec4_f32_in_mat2x4.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/vec4_f32_in_mat2x4.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : mat2x4<f32>;
fn func(pointer : ptr<storage, vec4<f32>>) -> vec4<f32> {
diff --git a/test/tint/ptr_ref/load/param/storage/vec4_f32_in_mat2x4.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/storage/vec4_f32_in_mat2x4.wgsl.expected.wgsl
index a11814a..174f22c 100644
--- a/test/tint/ptr_ref/load/param/storage/vec4_f32_in_mat2x4.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/vec4_f32_in_mat2x4.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage> S : mat2x4<f32>;
fn func(pointer : ptr<storage, vec4<f32>>) -> vec4<f32> {
diff --git a/test/tint/ptr_ref/load/param/storage/vec4_f32_in_struct.wgsl b/test/tint/ptr_ref/load/param/storage/vec4_f32_in_struct.wgsl
index 5fe26b3..97a8937 100644
--- a/test/tint/ptr_ref/load/param/storage/vec4_f32_in_struct.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/vec4_f32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
};
diff --git a/test/tint/ptr_ref/load/param/storage/vec4_f32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/storage/vec4_f32_in_struct.wgsl.expected.wgsl
index bb2f4cc..08a4e31 100644
--- a/test/tint/ptr_ref/load/param/storage/vec4_f32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/storage/vec4_f32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
}
diff --git a/test/tint/ptr_ref/load/param/uniform/array_in_struct.wgsl b/test/tint/ptr_ref/load/param/uniform/array_in_struct.wgsl
index cccfea0..aebb3dd 100644
--- a/test/tint/ptr_ref/load/param/uniform/array_in_struct.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/array_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<vec4<i32>, 4>,
};
diff --git a/test/tint/ptr_ref/load/param/uniform/array_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/uniform/array_in_struct.wgsl.expected.wgsl
index c2b203e..580d261 100644
--- a/test/tint/ptr_ref/load/param/uniform/array_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/array_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<vec4<i32>, 4>,
}
diff --git a/test/tint/ptr_ref/load/param/uniform/i32.wgsl b/test/tint/ptr_ref/load/param/uniform/i32.wgsl
index 8e37001..0b9a0b7 100644
--- a/test/tint/ptr_ref/load/param/uniform/i32.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/i32.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> S : i32;
fn func(pointer : ptr<uniform, i32>) -> i32 {
diff --git a/test/tint/ptr_ref/load/param/uniform/i32.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/uniform/i32.wgsl.expected.wgsl
index b8a873f..24a9e89 100644
--- a/test/tint/ptr_ref/load/param/uniform/i32.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/i32.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> S : i32;
fn func(pointer : ptr<uniform, i32>) -> i32 {
diff --git a/test/tint/ptr_ref/load/param/uniform/i32_in_struct.wgsl b/test/tint/ptr_ref/load/param/uniform/i32_in_struct.wgsl
index 6db13e0..c3573ab 100644
--- a/test/tint/ptr_ref/load/param/uniform/i32_in_struct.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/i32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/load/param/uniform/i32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/uniform/i32_in_struct.wgsl.expected.wgsl
index c44d00b..56fc937 100644
--- a/test/tint/ptr_ref/load/param/uniform/i32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/i32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/load/param/uniform/struct_in_array.wgsl b/test/tint/ptr_ref/load/param/uniform/struct_in_array.wgsl
index f0c98a4..f602cfb 100644
--- a/test/tint/ptr_ref/load/param/uniform/struct_in_array.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/struct_in_array.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<i32>,
};
diff --git a/test/tint/ptr_ref/load/param/uniform/struct_in_array.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/uniform/struct_in_array.wgsl.expected.wgsl
index be885de..04c15cc 100644
--- a/test/tint/ptr_ref/load/param/uniform/struct_in_array.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/struct_in_array.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<i32>,
}
diff --git a/test/tint/ptr_ref/load/param/uniform/vec2_f32_in_mat2x2.wgsl b/test/tint/ptr_ref/load/param/uniform/vec2_f32_in_mat2x2.wgsl
index e84f1a1..0874bb8 100644
--- a/test/tint/ptr_ref/load/param/uniform/vec2_f32_in_mat2x2.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/vec2_f32_in_mat2x2.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> S : mat2x2<f32>;
fn func(pointer : ptr<uniform, vec2<f32>>) -> vec2<f32> {
diff --git a/test/tint/ptr_ref/load/param/uniform/vec2_f32_in_mat2x2.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/uniform/vec2_f32_in_mat2x2.wgsl.expected.wgsl
index 1891f2a..041bd12 100644
--- a/test/tint/ptr_ref/load/param/uniform/vec2_f32_in_mat2x2.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/vec2_f32_in_mat2x2.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> S : mat2x2<f32>;
fn func(pointer : ptr<uniform, vec2<f32>>) -> vec2<f32> {
diff --git a/test/tint/ptr_ref/load/param/uniform/vec4_f32.wgsl b/test/tint/ptr_ref/load/param/uniform/vec4_f32.wgsl
index 68e7949..cf0470c 100644
--- a/test/tint/ptr_ref/load/param/uniform/vec4_f32.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/vec4_f32.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> S : vec4<f32>;
fn func(pointer : ptr<uniform, vec4<f32>>) -> vec4<f32> {
diff --git a/test/tint/ptr_ref/load/param/uniform/vec4_f32.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/uniform/vec4_f32.wgsl.expected.wgsl
index c5e8e6a..24a9408 100644
--- a/test/tint/ptr_ref/load/param/uniform/vec4_f32.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/vec4_f32.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> S : vec4<f32>;
fn func(pointer : ptr<uniform, vec4<f32>>) -> vec4<f32> {
diff --git a/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_mat2x4.wgsl b/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_mat2x4.wgsl
index a77ca7c..72eb8a4 100644
--- a/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_mat2x4.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_mat2x4.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> S : mat2x4<f32>;
fn func(pointer : ptr<uniform, vec4<f32>>) -> vec4<f32> {
diff --git a/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_mat2x4.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_mat2x4.wgsl.expected.wgsl
index 7f010b8..8b69502 100644
--- a/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_mat2x4.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_mat2x4.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<uniform> S : mat2x4<f32>;
fn func(pointer : ptr<uniform, vec4<f32>>) -> vec4<f32> {
diff --git a/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_struct.wgsl b/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_struct.wgsl
index 7059df4..8c923fa 100644
--- a/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_struct.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
};
diff --git a/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_struct.wgsl.expected.wgsl
index f08290a..8064b35 100644
--- a/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/uniform/vec4_f32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
}
diff --git a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl
index 55a558a..7e1c3a5 100644
--- a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
};
diff --git a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.wgsl
index 1368900c..ebd976a 100644
--- a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
}
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl
index f0b1be9..eb250f8 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : i32;
fn func(pointer : ptr<workgroup, i32>) -> i32 {
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.wgsl
index 8c26859d..65469f1 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : i32;
fn func(pointer : ptr<workgroup, i32>) -> i32 {
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl
index b777a56..e90c5377 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.wgsl
index d335ec2..b168087 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl
index 638ca03..d6f7bed 100644
--- a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.wgsl
index 2a586fb..81bc4ae 100644
--- a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl
index 5d8e9ae..6bb49f9 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : mat2x2<f32>;
fn func(pointer : ptr<workgroup, vec2<f32>>) -> vec2<f32> {
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.wgsl
index 55afec1..687ac5d 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : mat2x2<f32>;
fn func(pointer : ptr<workgroup, vec2<f32>>) -> vec2<f32> {
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl
index 0ccdce7..47157e8 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : vec4<f32>;
fn func(pointer : ptr<workgroup, vec4<f32>>) -> vec4<f32> {
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.wgsl
index 3c9ab35..fe5be6b 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : vec4<f32>;
fn func(pointer : ptr<workgroup, vec4<f32>>) -> vec4<f32> {
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl
index 204d6c2..87524e9 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : mat2x4<f32>;
fn func(pointer : ptr<workgroup, vec4<f32>>) -> vec4<f32> {
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.wgsl
index 5f0f9ae..69a1459 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : mat2x4<f32>;
fn func(pointer : ptr<workgroup, vec4<f32>>) -> vec4<f32> {
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl
index 5251b66..95e977e 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
};
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.wgsl
index 0d3e1d4..3b19185 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
}
diff --git a/test/tint/ptr_ref/store/param/function/array_in_struct.wgsl b/test/tint/ptr_ref/store/param/function/array_in_struct.wgsl
index c39eca4..d3f5cdc 100644
--- a/test/tint/ptr_ref/store/param/function/array_in_struct.wgsl
+++ b/test/tint/ptr_ref/store/param/function/array_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
};
diff --git a/test/tint/ptr_ref/store/param/function/array_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/function/array_in_struct.wgsl.expected.wgsl
index f45fcae..9e26365 100644
--- a/test/tint/ptr_ref/store/param/function/array_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/function/array_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
}
diff --git a/test/tint/ptr_ref/store/param/function/i32.wgsl.expected.spvasm b/test/tint/ptr_ref/store/param/function/i32.wgsl.expected.spvasm
index ed7174b..7fbb39f 100644
--- a/test/tint/ptr_ref/store/param/function/i32.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/store/param/function/i32.wgsl.expected.spvasm
@@ -7,7 +7,7 @@
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
- OpName %func "func"
+ OpName %func_F "func_F"
OpName %pointer "pointer"
OpName %main "main"
OpName %F "F"
@@ -18,7 +18,7 @@
%int_42 = OpConstant %int 42
%10 = OpTypeFunction %void
%14 = OpConstantNull %int
- %func = OpFunction %void None %1
+ %func_F = OpFunction %void None %1
%pointer = OpFunctionParameter %_ptr_Function_int
%7 = OpLabel
OpStore %pointer %int_42
@@ -27,6 +27,6 @@
%main = OpFunction %void None %10
%12 = OpLabel
%F = OpVariable %_ptr_Function_int Function %14
- %15 = OpFunctionCall %void %func %F
+ %15 = OpFunctionCall %void %func_F %F
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/store/param/function/i32_in_struct.wgsl b/test/tint/ptr_ref/store/param/function/i32_in_struct.wgsl
index 971651c..105ca7a 100644
--- a/test/tint/ptr_ref/store/param/function/i32_in_struct.wgsl
+++ b/test/tint/ptr_ref/store/param/function/i32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/store/param/function/i32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/function/i32_in_struct.wgsl.expected.wgsl
index 3db8678..115b428 100644
--- a/test/tint/ptr_ref/store/param/function/i32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/function/i32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/store/param/function/struct_in_array.wgsl b/test/tint/ptr_ref/store/param/function/struct_in_array.wgsl
index 631b70a..0875a4d 100644
--- a/test/tint/ptr_ref/store/param/function/struct_in_array.wgsl
+++ b/test/tint/ptr_ref/store/param/function/struct_in_array.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/store/param/function/struct_in_array.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/function/struct_in_array.wgsl.expected.wgsl
index fe15aa5..18c064b 100644
--- a/test/tint/ptr_ref/store/param/function/struct_in_array.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/function/struct_in_array.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/store/param/function/vec2_f32_in_mat2x2.wgsl b/test/tint/ptr_ref/store/param/function/vec2_f32_in_mat2x2.wgsl
index 3108042..2d7804c 100644
--- a/test/tint/ptr_ref/store/param/function/vec2_f32_in_mat2x2.wgsl
+++ b/test/tint/ptr_ref/store/param/function/vec2_f32_in_mat2x2.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<function, vec2<f32>>) {
*pointer = vec2<f32>();
}
diff --git a/test/tint/ptr_ref/store/param/function/vec2_f32_in_mat2x2.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/function/vec2_f32_in_mat2x2.wgsl.expected.wgsl
index a8dca7e..8ba47d9 100644
--- a/test/tint/ptr_ref/store/param/function/vec2_f32_in_mat2x2.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/function/vec2_f32_in_mat2x2.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<function, vec2<f32>>) {
*(pointer) = vec2<f32>();
}
diff --git a/test/tint/ptr_ref/store/param/function/vec4_f32.wgsl.expected.spvasm b/test/tint/ptr_ref/store/param/function/vec4_f32.wgsl.expected.spvasm
index c6e3a08..ced6948 100644
--- a/test/tint/ptr_ref/store/param/function/vec4_f32.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/store/param/function/vec4_f32.wgsl.expected.spvasm
@@ -7,7 +7,7 @@
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
- OpName %func "func"
+ OpName %func_F "func_F"
OpName %pointer "pointer"
OpName %main "main"
OpName %F "F"
@@ -18,7 +18,7 @@
%1 = OpTypeFunction %void %_ptr_Function_v4float
%10 = OpConstantNull %v4float
%11 = OpTypeFunction %void
- %func = OpFunction %void None %1
+ %func_F = OpFunction %void None %1
%pointer = OpFunctionParameter %_ptr_Function_v4float
%8 = OpLabel
OpStore %pointer %10
@@ -27,6 +27,6 @@
%main = OpFunction %void None %11
%13 = OpLabel
%F = OpVariable %_ptr_Function_v4float Function %10
- %15 = OpFunctionCall %void %func %F
+ %15 = OpFunctionCall %void %func_F %F
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/store/param/function/vec4_f32_in_mat2x4.wgsl b/test/tint/ptr_ref/store/param/function/vec4_f32_in_mat2x4.wgsl
index b71ee41..18f3995 100644
--- a/test/tint/ptr_ref/store/param/function/vec4_f32_in_mat2x4.wgsl
+++ b/test/tint/ptr_ref/store/param/function/vec4_f32_in_mat2x4.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<function, vec4<f32>>) {
*pointer = vec4<f32>();
}
diff --git a/test/tint/ptr_ref/store/param/function/vec4_f32_in_mat2x4.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/function/vec4_f32_in_mat2x4.wgsl.expected.wgsl
index 590aa4f..0b096ba 100644
--- a/test/tint/ptr_ref/store/param/function/vec4_f32_in_mat2x4.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/function/vec4_f32_in_mat2x4.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<function, vec4<f32>>) {
*(pointer) = vec4<f32>();
}
diff --git a/test/tint/ptr_ref/store/param/function/vec4_f32_in_struct.wgsl b/test/tint/ptr_ref/store/param/function/vec4_f32_in_struct.wgsl
index 9351165..8f1d2c5 100644
--- a/test/tint/ptr_ref/store/param/function/vec4_f32_in_struct.wgsl
+++ b/test/tint/ptr_ref/store/param/function/vec4_f32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
};
diff --git a/test/tint/ptr_ref/store/param/function/vec4_f32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/function/vec4_f32_in_struct.wgsl.expected.wgsl
index cf878a0..cc030aa 100644
--- a/test/tint/ptr_ref/store/param/function/vec4_f32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/function/vec4_f32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
}
diff --git a/test/tint/ptr_ref/store/param/private/array_in_struct.wgsl b/test/tint/ptr_ref/store/param/private/array_in_struct.wgsl
index 0d7c2d8..ba11ef2 100644
--- a/test/tint/ptr_ref/store/param/private/array_in_struct.wgsl
+++ b/test/tint/ptr_ref/store/param/private/array_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
};
diff --git a/test/tint/ptr_ref/store/param/private/array_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/private/array_in_struct.wgsl.expected.wgsl
index 99958c7..94ca0b1 100644
--- a/test/tint/ptr_ref/store/param/private/array_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/private/array_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
}
diff --git a/test/tint/ptr_ref/store/param/private/i32.wgsl.expected.spvasm b/test/tint/ptr_ref/store/param/private/i32.wgsl.expected.spvasm
index 867f836..c309afd 100644
--- a/test/tint/ptr_ref/store/param/private/i32.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/store/param/private/i32.wgsl.expected.spvasm
@@ -8,7 +8,7 @@
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpName %P "P"
- OpName %func "func"
+ OpName %func_F "func_F"
OpName %pointer "pointer"
OpName %main "main"
%int = OpTypeInt 32 1
@@ -19,7 +19,7 @@
%5 = OpTypeFunction %void %_ptr_Private_int
%int_42 = OpConstant %int 42
%12 = OpTypeFunction %void
- %func = OpFunction %void None %5
+ %func_F = OpFunction %void None %5
%pointer = OpFunctionParameter %_ptr_Private_int
%9 = OpLabel
OpStore %pointer %int_42
@@ -27,6 +27,6 @@
OpFunctionEnd
%main = OpFunction %void None %12
%14 = OpLabel
- %15 = OpFunctionCall %void %func %P
+ %15 = OpFunctionCall %void %func_F %P
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/store/param/private/i32_in_struct.wgsl b/test/tint/ptr_ref/store/param/private/i32_in_struct.wgsl
index 2238584..fd3ee42 100644
--- a/test/tint/ptr_ref/store/param/private/i32_in_struct.wgsl
+++ b/test/tint/ptr_ref/store/param/private/i32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/store/param/private/i32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/private/i32_in_struct.wgsl.expected.wgsl
index 7052c23..ddb6113 100644
--- a/test/tint/ptr_ref/store/param/private/i32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/private/i32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/store/param/private/struct_in_array.wgsl b/test/tint/ptr_ref/store/param/private/struct_in_array.wgsl
index 2a9d4cd..59bcdcf 100644
--- a/test/tint/ptr_ref/store/param/private/struct_in_array.wgsl
+++ b/test/tint/ptr_ref/store/param/private/struct_in_array.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/store/param/private/struct_in_array.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/private/struct_in_array.wgsl.expected.wgsl
index 3c97892..6ce0314 100644
--- a/test/tint/ptr_ref/store/param/private/struct_in_array.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/private/struct_in_array.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/store/param/private/vec2_f32_in_mat2x2.wgsl b/test/tint/ptr_ref/store/param/private/vec2_f32_in_mat2x2.wgsl
index 012a14d..e3d9154 100644
--- a/test/tint/ptr_ref/store/param/private/vec2_f32_in_mat2x2.wgsl
+++ b/test/tint/ptr_ref/store/param/private/vec2_f32_in_mat2x2.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<private, vec2<f32>>) {
*pointer = vec2<f32>();
}
diff --git a/test/tint/ptr_ref/store/param/private/vec2_f32_in_mat2x2.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/private/vec2_f32_in_mat2x2.wgsl.expected.wgsl
index fe087a1..215e71b 100644
--- a/test/tint/ptr_ref/store/param/private/vec2_f32_in_mat2x2.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/private/vec2_f32_in_mat2x2.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<private, vec2<f32>>) {
*(pointer) = vec2<f32>();
}
diff --git a/test/tint/ptr_ref/store/param/private/vec4_f32.wgsl.expected.spvasm b/test/tint/ptr_ref/store/param/private/vec4_f32.wgsl.expected.spvasm
index 334b687..36a2818 100644
--- a/test/tint/ptr_ref/store/param/private/vec4_f32.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/store/param/private/vec4_f32.wgsl.expected.spvasm
@@ -8,7 +8,7 @@
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpName %P "P"
- OpName %func "func"
+ OpName %func_F "func_F"
OpName %pointer "pointer"
OpName %main "main"
%float = OpTypeFloat 32
@@ -19,7 +19,7 @@
%void = OpTypeVoid
%6 = OpTypeFunction %void %_ptr_Private_v4float
%12 = OpTypeFunction %void
- %func = OpFunction %void None %6
+ %func_F = OpFunction %void None %6
%pointer = OpFunctionParameter %_ptr_Private_v4float
%10 = OpLabel
OpStore %pointer %5
@@ -27,6 +27,6 @@
OpFunctionEnd
%main = OpFunction %void None %12
%14 = OpLabel
- %15 = OpFunctionCall %void %func %P
+ %15 = OpFunctionCall %void %func_F %P
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/store/param/private/vec4_f32_in_mat2x4.wgsl b/test/tint/ptr_ref/store/param/private/vec4_f32_in_mat2x4.wgsl
index 14e4bfa..85dc073 100644
--- a/test/tint/ptr_ref/store/param/private/vec4_f32_in_mat2x4.wgsl
+++ b/test/tint/ptr_ref/store/param/private/vec4_f32_in_mat2x4.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<private, vec4<f32>>) {
*pointer = vec4<f32>();
}
diff --git a/test/tint/ptr_ref/store/param/private/vec4_f32_in_mat2x4.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/private/vec4_f32_in_mat2x4.wgsl.expected.wgsl
index 3ba72c6..1a1175e 100644
--- a/test/tint/ptr_ref/store/param/private/vec4_f32_in_mat2x4.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/private/vec4_f32_in_mat2x4.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
fn func(pointer : ptr<private, vec4<f32>>) {
*(pointer) = vec4<f32>();
}
diff --git a/test/tint/ptr_ref/store/param/private/vec4_f32_in_struct.wgsl b/test/tint/ptr_ref/store/param/private/vec4_f32_in_struct.wgsl
index bdc5cbc..26d4a73 100644
--- a/test/tint/ptr_ref/store/param/private/vec4_f32_in_struct.wgsl
+++ b/test/tint/ptr_ref/store/param/private/vec4_f32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
};
diff --git a/test/tint/ptr_ref/store/param/private/vec4_f32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/private/vec4_f32_in_struct.wgsl.expected.wgsl
index d4c15e1..048d888 100644
--- a/test/tint/ptr_ref/store/param/private/vec4_f32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/private/vec4_f32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
}
diff --git a/test/tint/ptr_ref/store/param/ptr.wgsl.expected.spvasm b/test/tint/ptr_ref/store/param/ptr.wgsl.expected.spvasm
index e750538..1485f53 100644
--- a/test/tint/ptr_ref/store/param/ptr.wgsl.expected.spvasm
+++ b/test/tint/ptr_ref/store/param/ptr.wgsl.expected.spvasm
@@ -7,7 +7,7 @@
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
- OpName %func "func"
+ OpName %func_F "func_F"
OpName %value "value"
OpName %pointer "pointer"
OpName %main "main"
@@ -19,7 +19,7 @@
%10 = OpTypeFunction %void
%int_123 = OpConstant %int 123
%15 = OpConstantNull %int
- %func = OpFunction %void None %1
+ %func_F = OpFunction %void None %1
%value = OpFunctionParameter %int
%pointer = OpFunctionParameter %_ptr_Function_int
%8 = OpLabel
@@ -30,6 +30,6 @@
%12 = OpLabel
%i = OpVariable %_ptr_Function_int Function %15
OpStore %i %int_123
- %16 = OpFunctionCall %void %func %int_123 %i
+ %16 = OpFunctionCall %void %func_F %int_123 %i
OpReturn
OpFunctionEnd
diff --git a/test/tint/ptr_ref/store/param/storage/array_in_struct.wgsl b/test/tint/ptr_ref/store/param/storage/array_in_struct.wgsl
index e0e40eb..ea44a34 100644
--- a/test/tint/ptr_ref/store/param/storage/array_in_struct.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/array_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
};
diff --git a/test/tint/ptr_ref/store/param/storage/array_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/storage/array_in_struct.wgsl.expected.wgsl
index 5545a0d..d962795 100644
--- a/test/tint/ptr_ref/store/param/storage/array_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/array_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
}
diff --git a/test/tint/ptr_ref/store/param/storage/i32.wgsl b/test/tint/ptr_ref/store/param/storage/i32.wgsl
index 8f1c6e9..cfab20c 100644
--- a/test/tint/ptr_ref/store/param/storage/i32.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/i32.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> S : i32;
fn func(pointer : ptr<storage, i32, read_write>) {
diff --git a/test/tint/ptr_ref/store/param/storage/i32.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/storage/i32.wgsl.expected.wgsl
index 6ed869b..531ab7b 100644
--- a/test/tint/ptr_ref/store/param/storage/i32.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/i32.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> S : i32;
fn func(pointer : ptr<storage, i32, read_write>) {
diff --git a/test/tint/ptr_ref/store/param/storage/i32_in_struct.wgsl b/test/tint/ptr_ref/store/param/storage/i32_in_struct.wgsl
index f68d27c..99f5006 100644
--- a/test/tint/ptr_ref/store/param/storage/i32_in_struct.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/i32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/store/param/storage/i32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/storage/i32_in_struct.wgsl.expected.wgsl
index 6a35569..49ce8b8 100644
--- a/test/tint/ptr_ref/store/param/storage/i32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/i32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/store/param/storage/struct_in_array.wgsl b/test/tint/ptr_ref/store/param/storage/struct_in_array.wgsl
index f13db83..196e3c8 100644
--- a/test/tint/ptr_ref/store/param/storage/struct_in_array.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/struct_in_array.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/store/param/storage/struct_in_array.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/storage/struct_in_array.wgsl.expected.wgsl
index 90db206..2feaf3c 100644
--- a/test/tint/ptr_ref/store/param/storage/struct_in_array.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/struct_in_array.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/store/param/storage/vec2_f32_in_mat2x2.wgsl b/test/tint/ptr_ref/store/param/storage/vec2_f32_in_mat2x2.wgsl
index d69899b..8a67978 100644
--- a/test/tint/ptr_ref/store/param/storage/vec2_f32_in_mat2x2.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/vec2_f32_in_mat2x2.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> S : mat2x2<f32>;
fn func(pointer : ptr<storage, vec2<f32>, read_write>) {
diff --git a/test/tint/ptr_ref/store/param/storage/vec2_f32_in_mat2x2.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/storage/vec2_f32_in_mat2x2.wgsl.expected.wgsl
index d31440c..f26fb14 100644
--- a/test/tint/ptr_ref/store/param/storage/vec2_f32_in_mat2x2.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/vec2_f32_in_mat2x2.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> S : mat2x2<f32>;
fn func(pointer : ptr<storage, vec2<f32>, read_write>) {
diff --git a/test/tint/ptr_ref/store/param/storage/vec4_f32.wgsl b/test/tint/ptr_ref/store/param/storage/vec4_f32.wgsl
index b76e544..16d8e24 100644
--- a/test/tint/ptr_ref/store/param/storage/vec4_f32.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/vec4_f32.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> S : vec4<f32>;
fn func(pointer : ptr<storage, vec4<f32>, read_write>) {
diff --git a/test/tint/ptr_ref/store/param/storage/vec4_f32.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/storage/vec4_f32.wgsl.expected.wgsl
index 3a10bfb..3e2c8cc 100644
--- a/test/tint/ptr_ref/store/param/storage/vec4_f32.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/vec4_f32.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> S : vec4<f32>;
fn func(pointer : ptr<storage, vec4<f32>, read_write>) {
diff --git a/test/tint/ptr_ref/store/param/storage/vec4_f32_in_mat2x4.wgsl b/test/tint/ptr_ref/store/param/storage/vec4_f32_in_mat2x4.wgsl
index 82b69e9..1c6afed 100644
--- a/test/tint/ptr_ref/store/param/storage/vec4_f32_in_mat2x4.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/vec4_f32_in_mat2x4.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> S : mat2x4<f32>;
fn func(pointer : ptr<storage, vec4<f32>, read_write>) {
diff --git a/test/tint/ptr_ref/store/param/storage/vec4_f32_in_mat2x4.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/storage/vec4_f32_in_mat2x4.wgsl.expected.wgsl
index 822fba4..7dc050f 100644
--- a/test/tint/ptr_ref/store/param/storage/vec4_f32_in_mat2x4.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/vec4_f32_in_mat2x4.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
@group(0) @binding(0) var<storage, read_write> S : mat2x4<f32>;
fn func(pointer : ptr<storage, vec4<f32>, read_write>) {
diff --git a/test/tint/ptr_ref/store/param/storage/vec4_f32_in_struct.wgsl b/test/tint/ptr_ref/store/param/storage/vec4_f32_in_struct.wgsl
index c49e789..58c5cd7 100644
--- a/test/tint/ptr_ref/store/param/storage/vec4_f32_in_struct.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/vec4_f32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
};
diff --git a/test/tint/ptr_ref/store/param/storage/vec4_f32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/storage/vec4_f32_in_struct.wgsl.expected.wgsl
index 89c0bc3..f3576e7 100644
--- a/test/tint/ptr_ref/store/param/storage/vec4_f32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/storage/vec4_f32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
}
diff --git a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl
index 4bf2e36..fab218b 100644
--- a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
};
diff --git a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.wgsl
index b6f7aeb..263bc27 100644
--- a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
arr : array<i32, 4>,
}
diff --git a/test/tint/ptr_ref/store/param/workgroup/i32.wgsl b/test/tint/ptr_ref/store/param/workgroup/i32.wgsl
index 9d7d66a..cd54b96 100644
--- a/test/tint/ptr_ref/store/param/workgroup/i32.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/i32.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : i32;
fn func(pointer : ptr<workgroup, i32>) {
diff --git a/test/tint/ptr_ref/store/param/workgroup/i32.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/workgroup/i32.wgsl.expected.wgsl
index 3bb4d84..6605ae6 100644
--- a/test/tint/ptr_ref/store/param/workgroup/i32.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/i32.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : i32;
fn func(pointer : ptr<workgroup, i32>) {
diff --git a/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl b/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl
index 3c3209f..edd594a 100644
--- a/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl.expected.wgsl
index 642e2a5..bec39be 100644
--- a/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl
index 5fbef20..1709254 100644
--- a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
};
diff --git a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.wgsl
index 8177527..ebb6b29 100644
--- a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : i32,
}
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl b/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl
index 2c25bbf..afc6510 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : mat2x2<f32>;
fn func(pointer : ptr<workgroup, vec2<f32>>) {
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.wgsl
index 7f66524..faca86f 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : mat2x2<f32>;
fn func(pointer : ptr<workgroup, vec2<f32>>) {
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl b/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl
index 382f211..68a5c91 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : vec4<f32>;
fn func(pointer : ptr<workgroup, vec4<f32>>) {
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl.expected.wgsl
index 1244c4a..469e96d 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : vec4<f32>;
fn func(pointer : ptr<workgroup, vec4<f32>>) {
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl
index 8143ef6..d258f5c 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : mat2x4<f32>;
fn func(pointer : ptr<workgroup, vec4<f32>>) {
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.wgsl
index b58576d..5058fa2 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
var<workgroup> S : mat2x4<f32>;
fn func(pointer : ptr<workgroup, vec4<f32>>) {
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl
index f21f9bf..f2706b2 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
};
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl.expected.wgsl b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl.expected.wgsl
index d6841c1..4d7c0fd 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl.expected.wgsl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl.expected.wgsl
@@ -1,5 +1,3 @@
-enable chromium_experimental_full_ptr_parameters;
-
struct str {
i : vec4<f32>,
}
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.spvasm b/test/tint/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.spvasm
index 1033394..1df16f3 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.spvasm
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.spvasm
@@ -16,7 +16,7 @@
OpMemberName %OuterS 0 "a1"
OpName %InnerS "InnerS"
OpMemberName %InnerS 0 "v"
- OpName %f "f"
+ OpName %f_F "f_F"
OpName %p "p"
OpName %v "v"
OpName %main "main"
@@ -49,7 +49,7 @@
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
%27 = OpTypeFunction %void
%31 = OpConstantNull %OuterS
- %f = OpFunction %void None %6
+ %f_F = OpFunction %void None %6
%p = OpFunctionParameter %_ptr_Function_OuterS
%16 = OpLabel
%v = OpVariable %_ptr_Function_InnerS Function %19
@@ -63,6 +63,6 @@
%main = OpFunction %void None %27
%29 = OpLabel
%s1 = OpVariable %_ptr_Function_OuterS Function %31
- %32 = OpFunctionCall %void %f %s1
+ %32 = OpFunctionCall %void %f_F %s1
OpReturn
OpFunctionEnd
