Remove deprecated ChromiumExperimentalSubgroups feature
Bug: 349125474, 377868468
Change-Id: I0d569c9efe5c4fd4071ddfa15459ae20a345de0c
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/215174
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: David Neto <dneto@google.com>
diff --git a/docs/dawn/features/shader_features.md b/docs/dawn/features/shader_features.md
index 3fdc484..68185f9 100644
--- a/docs/dawn/features/shader_features.md
+++ b/docs/dawn/features/shader_features.md
@@ -3,15 +3,6 @@
New shader features require the extensions to be enabled at device creation in addition to adding a `enable` directive in the WGSL.
This file regroups the documentation for all the Dawn-specific shader-only features.
-## ChromiumExperimentalSubgroups
-
-This adds support for the [`chromium_experimental_subgroups`](Link to do) WGSL `enable`.
-Currently used to investigate subgroup functionality and not for general use.
-
-It also provides the `wgpu::DawnExperimentalSubgroupLimits` structure used to gather data about the subgroup minimum and maximum size on the device.
-(the limit cannot be changed when requesting a device)
-`wgpu::DawnExperimentalSubgroupLimits` is populated by chaining it to the `wgpu::SupportedLimits` in the calls to `wgpu::Adapter::GetLimits` and `wgpu::Device::GetLimits`.
-
## ChromiumExperimentalSubgroupsUniformControlFlow
Used to expose that the device supports `VK_KHR_shader_subgroup_uniform_control_flow`.
diff --git a/src/dawn/dawn.json b/src/dawn/dawn.json
index 00b3fd9..ddf0f03 100644
--- a/src/dawn/dawn.json
+++ b/src/dawn/dawn.json
@@ -1185,15 +1185,6 @@
{"name": "compute", "type": "programmable stage descriptor"}
]
},
- "dawn compute pipeline full subgroups": {
- "category": "structure",
- "chained": "in",
- "chain roots": ["compute pipeline descriptor"],
- "tags": ["dawn"],
- "members": [
- {"name": "requires full subgroups", "type": "bool", "default": "false"}
- ]
- },
"alpha mode": {
"category": "enum",
"tags": ["dawn"],
@@ -2479,54 +2470,53 @@
{"value": 8, "name": "dual source blending", "tags": ["dawn"]},
{"value": 9, "name": "D3D11 multithread protected", "tags": ["dawn", "native"]},
{"value": 10, "name": "ANGLE texture sharing", "tags": ["dawn", "native"]},
- {"value": 11, "name": "chromium experimental subgroups", "tags": ["dawn"]},
- {"value": 12, "name": "chromium experimental subgroup uniform control flow", "tags": ["dawn"]},
- {"value": 13, "name": "pixel local storage coherent", "tags": ["dawn"]},
- {"value": 14, "name": "pixel local storage non coherent", "tags": ["dawn"]},
- {"value": 15, "name": "unorm16 texture formats", "tags": ["dawn"]},
- {"value": 16, "name": "snorm16 texture formats", "tags": ["dawn"]},
- {"value": 17, "name": "multi planar format extended usages", "tags": ["dawn"]},
- {"value": 18, "name": "multi planar format p010", "tags": ["dawn"]},
- {"value": 19, "name": "host mapped pointer", "tags": ["dawn"]},
- {"value": 20, "name": "multi planar render targets", "tags": ["dawn"]},
- {"value": 21, "name": "multi planar format nv12a", "tags": ["dawn"]},
- {"value": 22, "name": "framebuffer fetch", "tags": ["dawn"]},
- {"value": 23, "name": "buffer map extended usages", "tags": ["dawn"]},
- {"value": 24, "name": "adapter properties memory heaps", "tags": ["dawn"]},
- {"value": 25, "name": "adapter properties D3D", "tags": ["dawn"]},
- {"value": 26, "name": "adapter properties vk", "tags": ["dawn"]},
- {"value": 27, "name": "r8 unorm storage", "tags": ["dawn"]},
- {"value": 28, "name": "format capabilities", "tags": ["dawn"]},
- {"value": 29, "name": "drm format capabilities", "tags": ["dawn"]},
- {"value": 30, "name": "norm16 texture formats", "tags": ["dawn"]},
- {"value": 31, "name": "multi planar format nv16", "tags": ["dawn"]},
- {"value": 32, "name": "multi planar format nv24", "tags": ["dawn"]},
- {"value": 33, "name": "multi planar format p210", "tags": ["dawn"]},
- {"value": 34, "name": "multi planar format p410", "tags": ["dawn"]},
+ {"value": 11, "name": "chromium experimental subgroup uniform control flow", "tags": ["dawn"]},
+ {"value": 12, "name": "pixel local storage coherent", "tags": ["dawn"]},
+ {"value": 13, "name": "pixel local storage non coherent", "tags": ["dawn"]},
+ {"value": 14, "name": "unorm16 texture formats", "tags": ["dawn"]},
+ {"value": 15, "name": "snorm16 texture formats", "tags": ["dawn"]},
+ {"value": 16, "name": "multi planar format extended usages", "tags": ["dawn"]},
+ {"value": 17, "name": "multi planar format p010", "tags": ["dawn"]},
+ {"value": 18, "name": "host mapped pointer", "tags": ["dawn"]},
+ {"value": 19, "name": "multi planar render targets", "tags": ["dawn"]},
+ {"value": 20, "name": "multi planar format nv12a", "tags": ["dawn"]},
+ {"value": 21, "name": "framebuffer fetch", "tags": ["dawn"]},
+ {"value": 22, "name": "buffer map extended usages", "tags": ["dawn"]},
+ {"value": 23, "name": "adapter properties memory heaps", "tags": ["dawn"]},
+ {"value": 24, "name": "adapter properties D3D", "tags": ["dawn"]},
+ {"value": 25, "name": "adapter properties vk", "tags": ["dawn"]},
+ {"value": 26, "name": "r8 unorm storage", "tags": ["dawn"]},
+ {"value": 27, "name": "format capabilities", "tags": ["dawn"]},
+ {"value": 28, "name": "drm format capabilities", "tags": ["dawn"]},
+ {"value": 29, "name": "norm16 texture formats", "tags": ["dawn"]},
+ {"value": 30, "name": "multi planar format nv16", "tags": ["dawn"]},
+ {"value": 31, "name": "multi planar format nv24", "tags": ["dawn"]},
+ {"value": 32, "name": "multi planar format p210", "tags": ["dawn"]},
+ {"value": 33, "name": "multi planar format p410", "tags": ["dawn"]},
- {"value": 35, "name": "shared texture memory vk dedicated allocation", "tags": ["dawn", "native"]},
- {"value": 36, "name": "shared texture memory a hardware buffer", "tags": ["dawn", "native"]},
- {"value": 37, "name": "shared texture memory dma buf", "tags": ["dawn", "native"]},
- {"value": 38, "name": "shared texture memory opaque FD", "tags": ["dawn", "native"]},
- {"value": 39, "name": "shared texture memory zircon handle", "tags": ["dawn", "native"]},
- {"value": 40, "name": "shared texture memory DXGI shared handle", "tags": ["dawn", "native"]},
- {"value": 41, "name": "shared texture memory D3D11 texture 2D", "tags": ["dawn", "native"]},
- {"value": 42, "name": "shared texture memory IO surface", "tags": ["dawn", "native"]},
- {"value": 43, "name": "shared texture memory EGL image", "tags": ["dawn", "native"]},
- {"value": 44, "name": "shared fence vk semaphore opaque FD", "tags": ["dawn", "native"]},
- {"value": 45, "name": "shared fence sync FD", "tags": ["dawn", "native"]},
- {"value": 46, "name": "shared fence vk semaphore zircon handle", "tags": ["dawn", "native"]},
- {"value": 47, "name": "shared fence DXGI shared handle", "tags": ["dawn", "native"]},
- {"value": 48, "name": "shared fence MTL shared event", "tags": ["dawn", "native"]},
- {"value": 49, "name": "shared buffer memory D3D12 resource", "tags": ["dawn", "native"]},
- {"value": 50, "name": "static samplers", "tags": ["dawn"]},
- {"value": 51, "name": "y cb cr vulkan samplers", "tags": ["dawn"]},
- {"value": 52, "name": "shader module compilation options", "tags": ["dawn"]},
+ {"value": 34, "name": "shared texture memory vk dedicated allocation", "tags": ["dawn", "native"]},
+ {"value": 35, "name": "shared texture memory a hardware buffer", "tags": ["dawn", "native"]},
+ {"value": 36, "name": "shared texture memory dma buf", "tags": ["dawn", "native"]},
+ {"value": 37, "name": "shared texture memory opaque FD", "tags": ["dawn", "native"]},
+ {"value": 38, "name": "shared texture memory zircon handle", "tags": ["dawn", "native"]},
+ {"value": 39, "name": "shared texture memory DXGI shared handle", "tags": ["dawn", "native"]},
+ {"value": 40, "name": "shared texture memory D3D11 texture 2D", "tags": ["dawn", "native"]},
+ {"value": 41, "name": "shared texture memory IO surface", "tags": ["dawn", "native"]},
+ {"value": 42, "name": "shared texture memory EGL image", "tags": ["dawn", "native"]},
+ {"value": 43, "name": "shared fence vk semaphore opaque FD", "tags": ["dawn", "native"]},
+ {"value": 44, "name": "shared fence sync FD", "tags": ["dawn", "native"]},
+ {"value": 45, "name": "shared fence vk semaphore zircon handle", "tags": ["dawn", "native"]},
+ {"value": 46, "name": "shared fence DXGI shared handle", "tags": ["dawn", "native"]},
+ {"value": 47, "name": "shared fence MTL shared event", "tags": ["dawn", "native"]},
+ {"value": 48, "name": "shared buffer memory D3D12 resource", "tags": ["dawn", "native"]},
+ {"value": 49, "name": "static samplers", "tags": ["dawn"]},
+ {"value": 50, "name": "y cb cr vulkan samplers", "tags": ["dawn"]},
+ {"value": 51, "name": "shader module compilation options", "tags": ["dawn"]},
- {"value": 53, "name": "dawn load resolve texture", "tags": ["dawn"]},
- {"value": 54, "name": "dawn partial load resolve texture", "tags": ["dawn"]},
- {"value": 55, "name": "multi draw indirect", "tags": ["dawn"]},
- {"value": 56, "name": "clip distances", "tags": ["dawn"]}
+ {"value": 52, "name": "dawn load resolve texture", "tags": ["dawn"]},
+ {"value": 53, "name": "dawn partial load resolve texture", "tags": ["dawn"]},
+ {"value": 54, "name": "multi draw indirect", "tags": ["dawn"]},
+ {"value": 55, "name": "clip distances", "tags": ["dawn"]}
]
},
"filter mode": {
@@ -4035,44 +4025,43 @@
{"value": 20, "name": "adapter properties memory heaps", "tags": ["dawn"]},
{"value": 21, "name": "adapter properties D3D", "tags": ["dawn"]},
{"value": 22, "name": "adapter properties vk", "tags": ["dawn"]},
- {"value": 23, "name": "dawn compute pipeline full subgroups", "tags": ["dawn"]},
- {"value": 24, "name": "dawn wire WGSL control", "tags": ["dawn"]},
- {"value": 25, "name": "dawn WGSL blocklist", "tags": ["dawn", "native"]},
- {"value": 26, "name": "drm format capabilities", "tags": ["dawn"]},
- {"value": 27, "name": "shader module compilation options", "tags": ["dawn"]},
- {"value": 28, "name": "color target state expand resolve texture dawn", "tags": ["dawn"]},
- {"value": 29, "name": "render pass descriptor expand resolve rect", "tags": ["dawn"]},
+ {"value": 23, "name": "dawn wire WGSL control", "tags": ["dawn"]},
+ {"value": 24, "name": "dawn WGSL blocklist", "tags": ["dawn", "native"]},
+ {"value": 25, "name": "drm format capabilities", "tags": ["dawn"]},
+ {"value": 26, "name": "shader module compilation options", "tags": ["dawn"]},
+ {"value": 27, "name": "color target state expand resolve texture dawn", "tags": ["dawn"]},
+ {"value": 28, "name": "render pass descriptor expand resolve rect", "tags": ["dawn"]},
- {"value": 30, "name": "shared texture memory vk dedicated allocation descriptor", "tags": ["dawn", "native"]},
- {"value": 31, "name": "shared texture memory a hardware buffer descriptor", "tags": ["dawn", "native"]},
- {"value": 32, "name": "shared texture memory dma buf descriptor", "tags": ["dawn", "native"]},
- {"value": 33, "name": "shared texture memory opaque FD descriptor", "tags": ["dawn", "native"]},
- {"value": 34, "name": "shared texture memory zircon handle descriptor", "tags": ["dawn", "native"]},
- {"value": 35, "name": "shared texture memory DXGI shared handle descriptor", "tags": ["dawn", "native"]},
- {"value": 36, "name": "shared texture memory D3D11 texture 2D descriptor", "tags": ["dawn", "native"]},
- {"value": 37, "name": "shared texture memory IO surface descriptor", "tags": ["dawn", "native"]},
- {"value": 38, "name": "shared texture memory EGL image descriptor", "tags": ["dawn", "native"]},
- {"value": 39, "name": "shared texture memory initialized begin state", "tags": ["dawn", "native"]},
- {"value": 40, "name": "shared texture memory initialized end state", "tags": ["dawn", "native"]},
- {"value": 41, "name": "shared texture memory vk image layout begin state", "tags": ["dawn", "native"]},
- {"value": 42, "name": "shared texture memory vk image layout end state", "tags": ["dawn", "native"]},
- {"value": 43, "name": "shared texture memory D3D swapchain begin state", "tags": ["dawn", "native"]},
- {"value": 44, "name": "shared fence vk semaphore opaque FD descriptor", "tags": ["dawn", "native"]},
- {"value": 45, "name": "shared fence vk semaphore opaque FD export info", "tags": ["dawn", "native"]},
- {"value": 46, "name": "shared fence sync FD descriptor", "tags": ["dawn", "native"]},
- {"value": 47, "name": "shared fence sync FD export info", "tags": ["dawn", "native"]},
- {"value": 48, "name": "shared fence vk semaphore zircon handle descriptor", "tags": ["dawn", "native"]},
- {"value": 49, "name": "shared fence vk semaphore zircon handle export info", "tags": ["dawn", "native"]},
- {"value": 50, "name": "shared fence DXGI shared handle descriptor", "tags": ["dawn", "native"]},
- {"value": 51, "name": "shared fence DXGI shared handle export info", "tags": ["dawn", "native"]},
- {"value": 52, "name": "shared fence MTL shared event descriptor", "tags": ["dawn", "native"]},
- {"value": 53, "name": "shared fence MTL shared event export info", "tags": ["dawn", "native"]},
- {"value": 54, "name": "shared buffer memory D3D12 resource descriptor", "tags": ["dawn", "native"]},
- {"value": 55, "name": "static sampler binding layout", "tags": ["dawn"]},
- {"value": 56, "name": "y cb cr vk descriptor", "tags": ["dawn"]},
- {"value": 57, "name": "shared texture memory a hardware buffer properties", "tags": ["dawn", "native"]},
- {"value": 58, "name": "a hardware buffer properties", "tags": ["dawn", "native"]},
- {"value": 59, "name": "dawn experimental immediate data limits", "tags": ["dawn"]}
+ {"value": 29, "name": "shared texture memory vk dedicated allocation descriptor", "tags": ["dawn", "native"]},
+ {"value": 30, "name": "shared texture memory a hardware buffer descriptor", "tags": ["dawn", "native"]},
+ {"value": 31, "name": "shared texture memory dma buf descriptor", "tags": ["dawn", "native"]},
+ {"value": 32, "name": "shared texture memory opaque FD descriptor", "tags": ["dawn", "native"]},
+ {"value": 33, "name": "shared texture memory zircon handle descriptor", "tags": ["dawn", "native"]},
+ {"value": 34, "name": "shared texture memory DXGI shared handle descriptor", "tags": ["dawn", "native"]},
+ {"value": 35, "name": "shared texture memory D3D11 texture 2D descriptor", "tags": ["dawn", "native"]},
+ {"value": 36, "name": "shared texture memory IO surface descriptor", "tags": ["dawn", "native"]},
+ {"value": 37, "name": "shared texture memory EGL image descriptor", "tags": ["dawn", "native"]},
+ {"value": 38, "name": "shared texture memory initialized begin state", "tags": ["dawn", "native"]},
+ {"value": 39, "name": "shared texture memory initialized end state", "tags": ["dawn", "native"]},
+ {"value": 40, "name": "shared texture memory vk image layout begin state", "tags": ["dawn", "native"]},
+ {"value": 41, "name": "shared texture memory vk image layout end state", "tags": ["dawn", "native"]},
+ {"value": 42, "name": "shared texture memory D3D swapchain begin state", "tags": ["dawn", "native"]},
+ {"value": 43, "name": "shared fence vk semaphore opaque FD descriptor", "tags": ["dawn", "native"]},
+ {"value": 44, "name": "shared fence vk semaphore opaque FD export info", "tags": ["dawn", "native"]},
+ {"value": 45, "name": "shared fence sync FD descriptor", "tags": ["dawn", "native"]},
+ {"value": 46, "name": "shared fence sync FD export info", "tags": ["dawn", "native"]},
+ {"value": 47, "name": "shared fence vk semaphore zircon handle descriptor", "tags": ["dawn", "native"]},
+ {"value": 48, "name": "shared fence vk semaphore zircon handle export info", "tags": ["dawn", "native"]},
+ {"value": 49, "name": "shared fence DXGI shared handle descriptor", "tags": ["dawn", "native"]},
+ {"value": 50, "name": "shared fence DXGI shared handle export info", "tags": ["dawn", "native"]},
+ {"value": 51, "name": "shared fence MTL shared event descriptor", "tags": ["dawn", "native"]},
+ {"value": 52, "name": "shared fence MTL shared event export info", "tags": ["dawn", "native"]},
+ {"value": 53, "name": "shared buffer memory D3D12 resource descriptor", "tags": ["dawn", "native"]},
+ {"value": 54, "name": "static sampler binding layout", "tags": ["dawn"]},
+ {"value": 55, "name": "y cb cr vk descriptor", "tags": ["dawn"]},
+ {"value": 56, "name": "shared texture memory a hardware buffer properties", "tags": ["dawn", "native"]},
+ {"value": 57, "name": "a hardware buffer properties", "tags": ["dawn", "native"]},
+ {"value": 58, "name": "dawn experimental immediate data limits", "tags": ["dawn"]}
]
},
diff --git a/src/dawn/native/Adapter.cpp b/src/dawn/native/Adapter.cpp
index c2f0e70..2ea98ed 100644
--- a/src/dawn/native/Adapter.cpp
+++ b/src/dawn/native/Adapter.cpp
@@ -287,12 +287,6 @@
// creating a device for internal usage with AllowUnsafeAPI enabled from an adapter that
// disabled AllowUnsafeAPIS.
for (wgpu::FeatureName requiredFeature : requiredFeatureSet) {
- // TODO(349125474): Remove deprecated ChromiumExperimentalSubgroups.
- if (requiredFeature == wgpu::FeatureName::ChromiumExperimentalSubgroups) {
- GetInstance()->EmitDeprecationWarning(
- "Feature chromium-experimental-subgroups is deprecated. Use features subgroups and "
- "subgroups-f16 instead.");
- }
FeatureValidationResult result =
mPhysicalDevice->ValidateFeatureSupportedWithToggles(requiredFeature, deviceToggles);
DAWN_INVALID_IF(!result.success, "Invalid feature required: %s",
@@ -302,12 +296,9 @@
// TODO(349125474): Decide if this validation is needed, see
// https://github.com/gpuweb/gpuweb/issues/4734 for detail.
if (requiredFeatureSet.count(wgpu::FeatureName::SubgroupsF16) > 0) {
- // TODO(349125474): Remove deprecated ChromiumExperimentalSubgroups.
- DAWN_INVALID_IF(
- (requiredFeatureSet.count(wgpu::FeatureName::Subgroups) == 0) &&
- (requiredFeatureSet.count(wgpu::FeatureName::ChromiumExperimentalSubgroups) == 0),
- "Feature %s must be required together with feature %s.",
- wgpu::FeatureName::SubgroupsF16, wgpu::FeatureName::Subgroups);
+ DAWN_INVALID_IF((requiredFeatureSet.count(wgpu::FeatureName::Subgroups) == 0),
+ "Feature %s must be required together with feature %s.",
+ wgpu::FeatureName::SubgroupsF16, wgpu::FeatureName::Subgroups);
DAWN_INVALID_IF(requiredFeatureSet.count(wgpu::FeatureName::ShaderF16) == 0,
"Feature %s must be required together with feature %s.",
wgpu::FeatureName::SubgroupsF16, wgpu::FeatureName::ShaderF16);
diff --git a/src/dawn/native/ComputePipeline.cpp b/src/dawn/native/ComputePipeline.cpp
index 8aeefad..c693829 100644
--- a/src/dawn/native/ComputePipeline.cpp
+++ b/src/dawn/native/ComputePipeline.cpp
@@ -36,16 +36,6 @@
MaybeError ValidateComputePipelineDescriptor(DeviceBase* device,
const ComputePipelineDescriptor* descriptor) {
- UnpackedPtr<ComputePipelineDescriptor> unpacked;
- DAWN_TRY_ASSIGN(unpacked, ValidateAndUnpack(descriptor));
- auto* fullSubgroupsOption = unpacked.Get<DawnComputePipelineFullSubgroups>();
- // TODO(349125474): Decide what to do with fullSubgroupsOption before removing deprecated
- // ChromiumExperimentalSubgroups.
- DAWN_INVALID_IF(
- (fullSubgroupsOption && !device->HasFeature(Feature::ChromiumExperimentalSubgroups)),
- "DawnComputePipelineFullSubgroups is used without %s enabled.",
- ToAPI(Feature::ChromiumExperimentalSubgroups));
-
if (descriptor->layout != nullptr) {
DAWN_TRY(device->ValidateObject(descriptor->layout));
}
@@ -70,15 +60,10 @@
descriptor->layout,
descriptor->label,
{{SingleShaderStage::Compute, descriptor->compute.module, descriptor->compute.entryPoint,
- descriptor->compute.constantCount, descriptor->compute.constants}}),
- mRequiresFullSubgroups(false) {
+ descriptor->compute.constantCount, descriptor->compute.constants}}) {
SetContentHash(ComputeContentHash());
GetObjectTrackingList()->Track(this);
- if (auto* fullSubgroupsOption = descriptor.Get<DawnComputePipelineFullSubgroups>()) {
- mRequiresFullSubgroups = fullSubgroupsOption->requiresFullSubgroups;
- }
-
// Initialize the cache key to include the cache type and device information.
StreamIn(&mCacheKey, CacheKey::Type::ComputePipeline, device->GetCacheKey());
}
@@ -94,10 +79,6 @@
Uncache();
}
-bool ComputePipelineBase::IsFullSubgroupsRequired() const {
- return mRequiresFullSubgroups;
-}
-
// static
Ref<ComputePipelineBase> ComputePipelineBase::MakeError(DeviceBase* device, StringView label) {
class ErrorComputePipeline final : public ComputePipelineBase {
@@ -120,8 +101,7 @@
bool ComputePipelineBase::EqualityFunc::operator()(const ComputePipelineBase* a,
const ComputePipelineBase* b) const {
- return PipelineBase::EqualForCache(a, b) &&
- (a->IsFullSubgroupsRequired() == b->IsFullSubgroupsRequired());
+ return PipelineBase::EqualForCache(a, b);
}
} // namespace dawn::native
diff --git a/src/dawn/native/ComputePipeline.h b/src/dawn/native/ComputePipeline.h
index a3641d2..0827283 100644
--- a/src/dawn/native/ComputePipeline.h
+++ b/src/dawn/native/ComputePipeline.h
@@ -57,15 +57,11 @@
bool operator()(const ComputePipelineBase* a, const ComputePipelineBase* b) const;
};
- bool IsFullSubgroupsRequired() const;
-
protected:
void DestroyImpl() override;
private:
ComputePipelineBase(DeviceBase* device, ObjectBase::ErrorTag tag, StringView label);
-
- bool mRequiresFullSubgroups;
};
} // namespace dawn::native
diff --git a/src/dawn/native/Device.cpp b/src/dawn/native/Device.cpp
index 232f7d81..0e06f9a 100644
--- a/src/dawn/native/Device.cpp
+++ b/src/dawn/native/Device.cpp
@@ -1821,11 +1821,6 @@
if (mEnabledFeatures.IsEnabled(Feature::ShaderF16)) {
mWGSLAllowedFeatures.extensions.insert(tint::wgsl::Extension::kF16);
}
- // TODO(349125474): Remove deprecated ChromiumExperimentalSubgroups.
- if (mEnabledFeatures.IsEnabled(Feature::ChromiumExperimentalSubgroups)) {
- mWGSLAllowedFeatures.extensions.insert(
- tint::wgsl::Extension::kChromiumExperimentalSubgroups);
- }
if (mEnabledFeatures.IsEnabled(Feature::Subgroups)) {
mWGSLAllowedFeatures.extensions.insert(tint::wgsl::Extension::kSubgroups);
}
@@ -1973,9 +1968,7 @@
if (auto* subgroupLimits = unpacked.Get<DawnExperimentalSubgroupLimits>()) {
wgpu::ChainedStructOut* originalChain = subgroupLimits->nextInChain;
- // TODO(349125474): Remove deprecated ChromiumExperimentalSubgroups.
- if (!(HasFeature(Feature::Subgroups) ||
- HasFeature(Feature::ChromiumExperimentalSubgroups))) {
+ if (!HasFeature(Feature::Subgroups)) {
// If subgroups feature is not enabled, return the default-initialized
// DawnExperimentalSubgroupLimits object, where minSubgroupSize and
// maxSubgroupSize are WGPU_LIMIT_U32_UNDEFINED.
diff --git a/src/dawn/native/Features.cpp b/src/dawn/native/Features.cpp
index 52d2964..63cbdab 100644
--- a/src/dawn/native/Features.cpp
+++ b/src/dawn/native/Features.cpp
@@ -98,15 +98,6 @@
{"Allows textures with formats \"r32float\" \"rg32float\" and \"rgba32float\" to be "
"blendable.",
"https://gpuweb.github.io/gpuweb/#float32-blendable", FeatureInfo::FeatureState::Stable}},
- {Feature::ChromiumExperimentalSubgroups,
- {"DEPRECATED, use subgroups and subgroups-f16 features instead. "
- "Experimental, allows using subgroup and supports the \"enable "
- "chromium_experimental_subgroups\" directive in WGSL. Only used to investigate the semantic "
- "of subgroups and should not be relied upon. Note that currently \"enable "
- "chromium_experimental_subgroups\" feature allows using subgroups functions with f16 types "
- "within WGSL, but doesn't ensure that backend supports it.",
- "https://dawn.googlesource.com/dawn/+/refs/heads/main/docs/dawn/features/shader_features.md",
- FeatureInfo::FeatureState::Experimental}},
{Feature::ChromiumExperimentalSubgroupUniformControlFlow,
{"Experimental, supports VK_KHR_shader_subgroup_uniform_control_flow on Vulkan devices. Only "
"used to investigate the semantic of subgroups and should not be relied upon.",
diff --git a/src/dawn/native/ShaderModule.cpp b/src/dawn/native/ShaderModule.cpp
index 70c4c38..1d2b193 100644
--- a/src/dawn/native/ShaderModule.cpp
+++ b/src/dawn/native/ShaderModule.cpp
@@ -1069,8 +1069,7 @@
ResultOrError<Extent3D> ValidateComputeStageWorkgroupSize(
const tint::Program& program,
const char* entryPointName,
- const LimitsForCompilationRequest& limits,
- std::optional<uint32_t> maxSubgroupSizeForFullSubgroups) {
+ const LimitsForCompilationRequest& limits) {
tint::inspector::Inspector inspector(program);
// At this point the entry point must exist and must have workgroup size values.
tint::inspector::EntryPoint entryPoint = inspector.GetEntryPoint(entryPointName);
@@ -1104,14 +1103,6 @@
"the maximum allowed (%u bytes).",
workgroupStorageSize, limits.maxComputeWorkgroupStorageSize);
- // Validate workgroup_size.x is a multiple of maxSubgroupSizeForFullSubgroups if
- // it holds a value.
- DAWN_INVALID_IF(maxSubgroupSizeForFullSubgroups &&
- (workgroup_size.x % *maxSubgroupSizeForFullSubgroups != 0),
- "the X dimension of the workgroup size (%d) must be a multiple of "
- "maxSubgroupSize (%d) if full subgroups required in compute pipeline",
- workgroup_size.x, *maxSubgroupSizeForFullSubgroups);
-
return Extent3D{workgroup_size.x, workgroup_size.y, workgroup_size.z};
}
diff --git a/src/dawn/native/ShaderModule.h b/src/dawn/native/ShaderModule.h
index 6ba8db7..c421433 100644
--- a/src/dawn/native/ShaderModule.h
+++ b/src/dawn/native/ShaderModule.h
@@ -143,14 +143,12 @@
const EntryPointMetadata& entryPoint,
const PipelineLayoutBase* layout);
-// Return extent3D with workgroup size dimension info if it is valid. Also validate workgroup_size.x
-// is a multiple of maxSubgroupSizeForFullSubgroups if it holds a value.
+// Return extent3D with workgroup size dimension info if it is valid.
// width = x, height = y, depthOrArrayLength = z.
ResultOrError<Extent3D> ValidateComputeStageWorkgroupSize(
const tint::Program& program,
const char* entryPointName,
- const LimitsForCompilationRequest& limits,
- std::optional<uint32_t> maxSubgroupSizeForFullSubgroups);
+ const LimitsForCompilationRequest& limits);
RequiredBufferSizes ComputeRequiredBufferSizesForLayout(const EntryPointMetadata& entryPoint,
const PipelineLayoutBase* layout);
diff --git a/src/dawn/native/d3d/D3DCompilationRequest.h b/src/dawn/native/d3d/D3DCompilationRequest.h
index 475b113..6c91cd1 100644
--- a/src/dawn/native/d3d/D3DCompilationRequest.h
+++ b/src/dawn/native/d3d/D3DCompilationRequest.h
@@ -77,8 +77,7 @@
X(LimitsForCompilationRequest, limits) \
X(bool, disableSymbolRenaming) \
X(bool, dumpShaders) \
- X(bool, useTintIR) \
- X(std::optional<uint32_t>, maxSubgroupSizeForFullSubgroups)
+ X(bool, useTintIR)
#define D3D_BYTECODE_COMPILATION_REQUEST_MEMBERS(X) \
X(bool, hasShaderF16Feature) \
diff --git a/src/dawn/native/d3d/ShaderUtils.cpp b/src/dawn/native/d3d/ShaderUtils.cpp
index 4889a66..ad21be2 100644
--- a/src/dawn/native/d3d/ShaderUtils.cpp
+++ b/src/dawn/native/d3d/ShaderUtils.cpp
@@ -253,9 +253,8 @@
// Validate workgroup size after program runs transforms.
if (r.stage == SingleShaderStage::Compute) {
Extent3D _;
- DAWN_TRY_ASSIGN(
- _, ValidateComputeStageWorkgroupSize(transformedProgram, remappedEntryPointName->data(),
- r.limits, r.maxSubgroupSizeForFullSubgroups));
+ DAWN_TRY_ASSIGN(_, ValidateComputeStageWorkgroupSize(
+ transformedProgram, remappedEntryPointName->data(), r.limits));
}
bool usesVertexIndex = false;
diff --git a/src/dawn/native/d3d12/ComputePipelineD3D12.cpp b/src/dawn/native/d3d12/ComputePipelineD3D12.cpp
index b818190..92ee191 100644
--- a/src/dawn/native/d3d12/ComputePipelineD3D12.cpp
+++ b/src/dawn/native/d3d12/ComputePipelineD3D12.cpp
@@ -84,15 +84,9 @@
d3dDesc.pRootSignature = ToBackend(GetLayout())->GetRootSignature();
d3d::CompiledShader compiledShader;
- DAWN_TRY_ASSIGN(
- compiledShader,
- module->Compile(
- computeStage, SingleShaderStage::Compute, ToBackend(GetLayout()), compileFlags,
- /* usedInterstageVariables */ {},
- /* maxSubgroupSizeForFullSubgroups */
- IsFullSubgroupsRequired()
- ? std::make_optional(device->GetLimits().experimentalSubgroupLimits.maxSubgroupSize)
- : std::nullopt));
+ DAWN_TRY_ASSIGN(compiledShader, module->Compile(computeStage, SingleShaderStage::Compute,
+ ToBackend(GetLayout()), compileFlags,
+ /* usedInterstageVariables */ {}));
d3dDesc.CS = {compiledShader.shaderBlob.Data(), compiledShader.shaderBlob.Size()};
StreamIn(&mCacheKey, d3dDesc, ToBackend(GetLayout())->GetRootSignatureBlob());
diff --git a/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp b/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp
index cdaf7e5..b8fc14d 100644
--- a/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp
+++ b/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp
@@ -163,10 +163,8 @@
shaderF16Enabled = true;
}
- // ChromiumExperimentalSubgroups requires SM >= 6.0 and capabilities flags.
+ // Subgroups feature requires SM >= 6.0 and capabilities flags.
if (GetBackend()->IsDXCAvailable() && mDeviceInfo.supportsWaveOps) {
- // TODO(349125474): Remove deprecated ChromiumExperimentalSubgroups.
- EnableFeature(Feature::ChromiumExperimentalSubgroups);
EnableFeature(Feature::Subgroups);
// D3D12 devices that support both native f16 and wave ops can support subgroups-f16.
if (shaderF16Enabled) {
@@ -399,8 +397,6 @@
case wgpu::FeatureName::ShaderF16:
case wgpu::FeatureName::Subgroups:
case wgpu::FeatureName::SubgroupsF16:
- // TODO(349125474): Remove deprecated ChromiumExperimentalSubgroups.
- case wgpu::FeatureName::ChromiumExperimentalSubgroups:
return FeatureValidationResult(
absl::StrFormat("Feature %s requires DXC for D3D12.", feature));
default:
diff --git a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
index 4fd8126..a2d8aa6 100644
--- a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
+++ b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
@@ -130,8 +130,8 @@
SingleShaderStage stage,
const PipelineLayout* layout,
uint32_t compileFlags,
- const std::optional<dawn::native::d3d::InterStageShaderVariablesMask>& usedInterstageVariables,
- std::optional<uint32_t> maxSubgroupSizeForFullSubgroups) {
+ const std::optional<dawn::native::d3d::InterStageShaderVariablesMask>&
+ usedInterstageVariables) {
Device* device = ToBackend(GetDevice());
TRACE_EVENT0(device->GetPlatform(), General, "ShaderModuleD3D12::Compile");
DAWN_ASSERT(!IsError());
@@ -146,7 +146,6 @@
req.hlsl.disableSymbolRenaming = device->IsToggleEnabled(Toggle::DisableSymbolRenaming);
req.hlsl.dumpShaders = device->IsToggleEnabled(Toggle::DumpShaders);
req.hlsl.useTintIR = device->IsToggleEnabled(Toggle::UseTintIR);
- req.hlsl.maxSubgroupSizeForFullSubgroups = maxSubgroupSizeForFullSubgroups;
req.bytecode.hasShaderF16Feature = device->HasFeature(Feature::ShaderF16);
req.bytecode.compileFlags = compileFlags;
diff --git a/src/dawn/native/d3d12/ShaderModuleD3D12.h b/src/dawn/native/d3d12/ShaderModuleD3D12.h
index f5e143d..cab1dd0 100644
--- a/src/dawn/native/d3d12/ShaderModuleD3D12.h
+++ b/src/dawn/native/d3d12/ShaderModuleD3D12.h
@@ -62,8 +62,7 @@
const PipelineLayout* layout,
uint32_t compileFlags,
const std::optional<dawn::native::d3d::InterStageShaderVariablesMask>&
- usedInterstageVariables = {},
- std::optional<uint32_t> maxSubgroupSizeForFullSubgroups = std::nullopt);
+ usedInterstageVariables = {});
private:
ShaderModule(Device* device,
diff --git a/src/dawn/native/metal/ComputePipelineMTL.mm b/src/dawn/native/metal/ComputePipelineMTL.mm
index 759f628..b7d784f 100644
--- a/src/dawn/native/metal/ComputePipelineMTL.mm
+++ b/src/dawn/native/metal/ComputePipelineMTL.mm
@@ -59,15 +59,10 @@
ShaderModule::MetalFunctionData computeData;
DAWN_TRY(ToBackend(computeStage.module.Get())
- ->CreateFunction(
- SingleShaderStage::Compute, computeStage, ToBackend(GetLayout()), &computeData,
- /* sampleMask */ 0xFFFFFFFF,
- /* renderPipeline */ nullptr,
- /* maxSubgroupSizeForFullSubgroups */
- IsFullSubgroupsRequired()
- ? std::make_optional(
- GetDevice()->GetLimits().experimentalSubgroupLimits.maxSubgroupSize)
- : std::nullopt));
+ ->CreateFunction(SingleShaderStage::Compute, computeStage, ToBackend(GetLayout()),
+ &computeData,
+ /* sampleMask */ 0xFFFFFFFF,
+ /* renderPipeline */ nullptr));
NSError* error = nullptr;
NSRef<NSString> label = MakeDebugName(GetDevice(), "Dawn_ComputePipeline", GetLabel());
@@ -78,10 +73,6 @@
descriptor.computeFunction = computeData.function.Get();
descriptor.label = label.Get();
- if (IsFullSubgroupsRequired()) {
- descriptor.threadGroupSizeIsMultipleOfThreadExecutionWidth = true;
- }
-
platform::metrics::DawnHistogramTimer timer(GetDevice()->GetPlatform());
mMtlComputePipelineState.Acquire([mtlDevice
newComputePipelineStateWithDescriptor:descriptor
diff --git a/src/dawn/native/metal/PhysicalDeviceMTL.mm b/src/dawn/native/metal/PhysicalDeviceMTL.mm
index dd63da6..0454bf8 100644
--- a/src/dawn/native/metal/PhysicalDeviceMTL.mm
+++ b/src/dawn/native/metal/PhysicalDeviceMTL.mm
@@ -737,8 +737,6 @@
[*mDevice supportsFamily:MTLGPUFamilyMac2]) {
EnableFeature(Feature::Subgroups);
EnableFeature(Feature::SubgroupsF16);
- // TODO(349125474): Remove deprecated ChromiumExperimentalSubgroups.
- EnableFeature(Feature::ChromiumExperimentalSubgroups);
}
}
diff --git a/src/dawn/native/metal/ShaderModuleMTL.h b/src/dawn/native/metal/ShaderModuleMTL.h
index 9f34270..fbb0ebf 100644
--- a/src/dawn/native/metal/ShaderModuleMTL.h
+++ b/src/dawn/native/metal/ShaderModuleMTL.h
@@ -64,14 +64,12 @@
MTLSize localWorkgroupSize;
};
- MaybeError CreateFunction(
- SingleShaderStage stage,
- const ProgrammableStage& programmableStage,
- const PipelineLayout* layout,
- MetalFunctionData* out,
- uint32_t sampleMask = 0xFFFFFFFF,
- const RenderPipeline* renderPipeline = nullptr,
- std::optional<uint32_t> maxSubgroupSizeForFullSubgroups = std::nullopt);
+ MaybeError CreateFunction(SingleShaderStage stage,
+ const ProgrammableStage& programmableStage,
+ const PipelineLayout* layout,
+ MetalFunctionData* out,
+ uint32_t sampleMask = 0xFFFFFFFF,
+ const RenderPipeline* renderPipeline = nullptr);
private:
ShaderModule(Device* device,
diff --git a/src/dawn/native/metal/ShaderModuleMTL.mm b/src/dawn/native/metal/ShaderModuleMTL.mm
index 173489a..04e93cd 100644
--- a/src/dawn/native/metal/ShaderModuleMTL.mm
+++ b/src/dawn/native/metal/ShaderModuleMTL.mm
@@ -64,8 +64,7 @@
X(bool, disableSymbolRenaming) \
X(tint::msl::writer::Options, tintOptions) \
X(bool, use_tint_ir) \
- X(CacheKey::UnsafeUnkeyedValue<dawn::platform::Platform*>, platform) \
- X(std::optional<uint32_t>, maxSubgroupSizeForFullSubgroups)
+ X(CacheKey::UnsafeUnkeyedValue<dawn::platform::Platform*>, platform)
DAWN_MAKE_CACHE_REQUEST(MslCompilationRequest, MSL_COMPILATION_REQUEST_MEMBERS);
#undef MSL_COMPILATION_REQUEST_MEMBERS
@@ -210,8 +209,7 @@
ShaderModule::MetalFunctionData* out,
uint32_t sampleMask,
const RenderPipeline* renderPipeline,
- const BindingInfoArray& moduleBindingInfo,
- std::optional<uint32_t> maxSubgroupSizeForFullSubgroups) {
+ const BindingInfoArray& moduleBindingInfo) {
ScopedTintICEHandler scopedICEHandler(device);
std::ostringstream errorStream;
@@ -275,7 +273,6 @@
req.entryPointName = programmableStage.entryPoint.c_str();
req.disableSymbolRenaming = device->IsToggleEnabled(Toggle::DisableSymbolRenaming);
req.platform = UnsafeUnkeyedValue(device->GetPlatform());
- req.maxSubgroupSizeForFullSubgroups = maxSubgroupSizeForFullSubgroups;
req.tintOptions.disable_robustness = !device->IsRobustnessEnabled();
req.tintOptions.buffer_size_ubo_index = kBufferLengthBufferSlot;
@@ -358,8 +355,7 @@
if (r.stage == SingleShaderStage::Compute) {
// Validate workgroup size after program runs transforms.
DAWN_TRY_ASSIGN(localSize, ValidateComputeStageWorkgroupSize(
- program, remappedEntryPointName.data(), r.limits,
- r.maxSubgroupSizeForFullSubgroups));
+ program, remappedEntryPointName.data(), r.limits));
}
TRACE_EVENT0(r.platform.UnsafeGetValue(), General, "tint::msl::writer::Generate");
@@ -419,8 +415,7 @@
const PipelineLayout* layout,
ShaderModule::MetalFunctionData* out,
uint32_t sampleMask,
- const RenderPipeline* renderPipeline,
- std::optional<uint32_t> maxSubgroupSizeForFullSubgroups) {
+ const RenderPipeline* renderPipeline) {
TRACE_EVENT1(GetDevice()->GetPlatform(), General, "metal::ShaderModule::CreateFunction",
"label", utils::GetLabelForTrace(GetLabel()));
@@ -437,8 +432,7 @@
CacheResult<MslCompilation> mslCompilation;
DAWN_TRY_ASSIGN(mslCompilation,
TranslateToMSL(GetDevice(), programmableStage, stage, layout, out, sampleMask,
- renderPipeline, GetEntryPoint(entryPointName).bindings,
- maxSubgroupSizeForFullSubgroups));
+ renderPipeline, GetEntryPoint(entryPointName).bindings));
out->needsStorageBufferLength = mslCompilation->needsStorageBufferLength;
out->workgroupAllocations = std::move(mslCompilation->workgroupAllocations);
diff --git a/src/dawn/native/null/DeviceNull.cpp b/src/dawn/native/null/DeviceNull.cpp
index ab4889c..3a2a120 100644
--- a/src/dawn/native/null/DeviceNull.cpp
+++ b/src/dawn/native/null/DeviceNull.cpp
@@ -488,17 +488,12 @@
DAWN_TRY_ASSIGN(transformedProgram, RunTransforms(&transformManager, &(tintProgram->program),
transformInputs, nullptr, nullptr));
- // Do the workgroup size validation, although different backend will have different
- // fullSubgroups parameter.
+ // Do the workgroup size validation.
const CombinedLimits& limits = GetDevice()->GetLimits();
Extent3D _;
DAWN_TRY_ASSIGN(
- _, ValidateComputeStageWorkgroupSize(
- transformedProgram, computeStage.entryPoint.c_str(),
- LimitsForCompilationRequest::Create(limits.v1), /* maxSubgroupSizeForFullSubgroups */
- IsFullSubgroupsRequired()
- ? std::make_optional(limits.experimentalSubgroupLimits.maxSubgroupSize)
- : std::nullopt));
+ _, ValidateComputeStageWorkgroupSize(transformedProgram, computeStage.entryPoint.c_str(),
+ LimitsForCompilationRequest::Create(limits.v1)));
return {};
}
diff --git a/src/dawn/native/opengl/ShaderModuleGL.cpp b/src/dawn/native/opengl/ShaderModuleGL.cpp
index 2adfbf6..8b54ba2 100644
--- a/src/dawn/native/opengl/ShaderModuleGL.cpp
+++ b/src/dawn/native/opengl/ShaderModuleGL.cpp
@@ -569,8 +569,7 @@
// Validate workgroup size after program runs transforms.
Extent3D _;
DAWN_TRY_ASSIGN(_, ValidateComputeStageWorkgroupSize(
- program, remappedEntryPoint.c_str(), r.limits,
- /* fullSubgroups */ {}));
+ program, remappedEntryPoint.c_str(), r.limits));
}
// Intentionally assign entry point to empty to avoid a redundant 'SingleEntryPoint'
diff --git a/src/dawn/native/vulkan/ComputePipelineVk.cpp b/src/dawn/native/vulkan/ComputePipelineVk.cpp
index 013989b..e5b6a55 100644
--- a/src/dawn/native/vulkan/ComputePipelineVk.cpp
+++ b/src/dawn/native/vulkan/ComputePipelineVk.cpp
@@ -78,38 +78,20 @@
ShaderModule* module = ToBackend(computeStage.module.Get());
ShaderModule::ModuleAndSpirv moduleAndSpirv;
- DAWN_TRY_ASSIGN(
- moduleAndSpirv,
- module->GetHandleAndSpirv(
- SingleShaderStage::Compute, computeStage, layout,
- /*clampFragDepth*/ false,
- /*emitPointSize*/ false,
- /* maxSubgroupSizeForFullSubgroups */
- IsFullSubgroupsRequired()
- ? std::make_optional(device->GetLimits().experimentalSubgroupLimits.maxSubgroupSize)
- : std::nullopt));
+ DAWN_TRY_ASSIGN(moduleAndSpirv,
+ module->GetHandleAndSpirv(SingleShaderStage::Compute, computeStage, layout,
+ /*clampFragDepth*/ false,
+ /*emitPointSize*/ false));
createInfo.stage.module = moduleAndSpirv.module;
createInfo.stage.pName = moduleAndSpirv.remappedEntryPoint.c_str();
-
- if (IsFullSubgroupsRequired()) {
- // Workgroup size validation is handled in ValidateComputeStageWorkgroupSize when compiling
- // shader module. Vulkan device that support Subgroups feature must support
- // computeFullSubgroups.
- DAWN_ASSERT(device->GetDeviceInfo().subgroupSizeControlFeatures.computeFullSubgroups);
- createInfo.stage.flags |= VK_PIPELINE_SHADER_STAGE_CREATE_REQUIRE_FULL_SUBGROUPS_BIT |
- VK_PIPELINE_SHADER_STAGE_CREATE_ALLOW_VARYING_SUBGROUP_SIZE_BIT;
- }
-
createInfo.stage.pSpecializationInfo = nullptr;
VkPipelineShaderStageRequiredSubgroupSizeCreateInfoEXT subgroupSizeInfo = {};
PNextChainBuilder stageExtChain(&createInfo.stage);
uint32_t computeSubgroupSize = device->GetComputeSubgroupSize();
- // If experimental full subgroups is required, pipeline is created with varying subgroup size
- // enabled, and thus do not use explicit subgroup size control.
- if (computeSubgroupSize != 0u && !IsFullSubgroupsRequired()) {
+ if (computeSubgroupSize != 0u) {
DAWN_ASSERT(device->GetDeviceInfo().HasExt(DeviceExt::SubgroupSizeControl));
subgroupSizeInfo.requiredSubgroupSize = computeSubgroupSize;
stageExtChain.Add(
diff --git a/src/dawn/native/vulkan/DeviceVk.cpp b/src/dawn/native/vulkan/DeviceVk.cpp
index 5c3df42..22c080c 100644
--- a/src/dawn/native/vulkan/DeviceVk.cpp
+++ b/src/dawn/native/vulkan/DeviceVk.cpp
@@ -497,24 +497,14 @@
}
// Set device feature for subgroups with f16 types.
- // TODO(349125474): Remove deprecated ChromiumExperimentalSubgroups.
- if (HasFeature(Feature::SubgroupsF16) || HasFeature(Feature::ChromiumExperimentalSubgroups)) {
- // If ChromiumExperimentalSubgroups feature is required, set the shaderSubgroupExtendedTypes
- // as-is, so that subgroups functions with f16 can be used if supported by backend.
- if (HasFeature(Feature::ChromiumExperimentalSubgroups)) {
- if (usedKnobs.HasExt(DeviceExt::ShaderSubgroupExtendedTypes)) {
- usedKnobs.shaderSubgroupExtendedTypes = mDeviceInfo.shaderSubgroupExtendedTypes;
- featuresChain.Add(&usedKnobs.shaderSubgroupExtendedTypes);
- }
- } else {
- DAWN_ASSERT(usedKnobs.HasExt(DeviceExt::ShaderSubgroupExtendedTypes) &&
- mDeviceInfo.shaderSubgroupExtendedTypes.shaderSubgroupExtendedTypes ==
- VK_TRUE &&
- HasFeature(Feature::ShaderF16) && HasFeature(Feature::Subgroups));
+ if (HasFeature(Feature::SubgroupsF16)) {
+ DAWN_ASSERT(usedKnobs.HasExt(DeviceExt::ShaderSubgroupExtendedTypes) &&
+ mDeviceInfo.shaderSubgroupExtendedTypes.shaderSubgroupExtendedTypes ==
+ VK_TRUE &&
+ HasFeature(Feature::ShaderF16) && HasFeature(Feature::Subgroups));
- usedKnobs.shaderSubgroupExtendedTypes = mDeviceInfo.shaderSubgroupExtendedTypes;
- featuresChain.Add(&usedKnobs.shaderSubgroupExtendedTypes);
- }
+ usedKnobs.shaderSubgroupExtendedTypes = mDeviceInfo.shaderSubgroupExtendedTypes;
+ featuresChain.Add(&usedKnobs.shaderSubgroupExtendedTypes);
}
if (HasFeature(Feature::DualSourceBlending)) {
diff --git a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
index 84499d3..7e3702f 100644
--- a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
+++ b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
@@ -386,26 +386,6 @@
EnableFeature(Feature::AdapterPropertiesVk);
EnableFeature(Feature::DawnLoadResolveTexture);
- // TODO(349125474): Remove deprecated ChromiumExperimentalSubgroups.
- // Enable ChromiumExperimentalSubgroups feature if:
- // 1. Vulkan API version is 1.1 or later, and
- // 2. subgroupSupportedStages includes compute stage bit, and
- // 3. subgroupSupportedOperations includes basic and ballot bits, and
- // 4. VK_EXT_subgroup_size_control extension is valid, and both subgroupSizeControl
- // and computeFullSubgroups is TRUE in VkPhysicalDeviceSubgroupSizeControlFeaturesEXT.
- // Notes that these requirement doesn't ensure all subgroups features are supported by the
- // Vulkan backend. For example, currently ChromiumExperimentalSubgroups feature allows using
- // subgroups functions with f16 types in WGSL, but doesn't ensure that backend supports it.
- if ((mDeviceInfo.properties.apiVersion >= VK_API_VERSION_1_1) &&
- (mDeviceInfo.subgroupProperties.supportedStages & VK_SHADER_STAGE_COMPUTE_BIT) &&
- (mDeviceInfo.subgroupProperties.supportedOperations & VK_SUBGROUP_FEATURE_BASIC_BIT) &&
- (mDeviceInfo.subgroupProperties.supportedOperations & VK_SUBGROUP_FEATURE_BALLOT_BIT) &&
- (mDeviceInfo.HasExt(DeviceExt::SubgroupSizeControl)) &&
- (mDeviceInfo.subgroupSizeControlFeatures.subgroupSizeControl == VK_TRUE) &&
- (mDeviceInfo.subgroupSizeControlFeatures.computeFullSubgroups == VK_TRUE)) {
- EnableFeature(Feature::ChromiumExperimentalSubgroups);
- }
-
// Enable Subgroups feature if:
// 1. Vulkan API version is 1.1 or later, and
// 2. subgroupSupportedStages includes compute and fragment stage bit, and
diff --git a/src/dawn/native/vulkan/RenderPipelineVk.cpp b/src/dawn/native/vulkan/RenderPipelineVk.cpp
index 0b7e3c9..e8a0868 100644
--- a/src/dawn/native/vulkan/RenderPipelineVk.cpp
+++ b/src/dawn/native/vulkan/RenderPipelineVk.cpp
@@ -367,8 +367,7 @@
ShaderModule::ModuleAndSpirv moduleAndSpirv;
DAWN_TRY_ASSIGN(moduleAndSpirv, ToBackend(programmableStage.module)
->GetHandleAndSpirv(stage, programmableStage, layout,
- clampFragDepth, emitPointSize,
- /* fullSubgroups */ {}));
+ clampFragDepth, emitPointSize));
mHasInputAttachment = mHasInputAttachment || moduleAndSpirv.hasInputAttachment;
// Record cache key for each shader since it will become inaccessible later on.
StreamIn(&mCacheKey, stream::Iterable(moduleAndSpirv.spirv, moduleAndSpirv.wordCount));
diff --git a/src/dawn/native/vulkan/ShaderModuleVk.cpp b/src/dawn/native/vulkan/ShaderModuleVk.cpp
index 3ff33cd..6aeacac 100644
--- a/src/dawn/native/vulkan/ShaderModuleVk.cpp
+++ b/src/dawn/native/vulkan/ShaderModuleVk.cpp
@@ -76,9 +76,6 @@
if (!std::equal(constants.begin(), constants.end(), other.constants.begin())) {
return false;
}
- if (maxSubgroupSizeForFullSubgroups != other.maxSubgroupSizeForFullSubgroups) {
- return false;
- }
if (emitPointSize != other.emitPointSize) {
return false;
}
@@ -205,8 +202,7 @@
X(std::string_view, entryPointName) \
X(bool, disableSymbolRenaming) \
X(tint::spirv::writer::Options, tintOptions) \
- X(CacheKey::UnsafeUnkeyedValue<dawn::platform::Platform*>, platform) \
- X(std::optional<uint32_t>, maxSubgroupSizeForFullSubgroups)
+ X(CacheKey::UnsafeUnkeyedValue<dawn::platform::Platform*>, platform)
DAWN_MAKE_CACHE_REQUEST(SpirvCompilationRequest, SPIRV_COMPILATION_REQUEST_MEMBERS);
#undef SPIRV_COMPILATION_REQUEST_MEMBERS
@@ -218,8 +214,7 @@
const ProgrammableStage& programmableStage,
const PipelineLayout* layout,
bool clampFragDepth,
- bool emitPointSize,
- std::optional<uint32_t> maxSubgroupSizeForFullSubgroups) {
+ bool emitPointSize) {
TRACE_EVENT0(GetDevice()->GetPlatform(), General, "ShaderModuleVk::GetHandleAndSpirv");
ScopedTintICEHandler scopedICEHandler(GetDevice());
@@ -228,9 +223,9 @@
// TODO(chromium:345359083): Improve the computation of the cache key. For example, it isn't
// ideal to use `reinterpret_cast<uintptr_t>(layout)` as the layout may be freed and
// reallocated during the runtime.
- auto cacheKey = TransformedShaderModuleCacheKey{
- reinterpret_cast<uintptr_t>(layout), programmableStage.entryPoint.c_str(),
- programmableStage.constants, maxSubgroupSizeForFullSubgroups, emitPointSize};
+ auto cacheKey = TransformedShaderModuleCacheKey{reinterpret_cast<uintptr_t>(layout),
+ programmableStage.entryPoint.c_str(),
+ programmableStage.constants, emitPointSize};
auto handleAndSpirv = mTransformedShaderModuleCache->Find(cacheKey);
if (handleAndSpirv.has_value()) {
return std::move(*handleAndSpirv);
@@ -346,7 +341,6 @@
req.disableSymbolRenaming = GetDevice()->IsToggleEnabled(Toggle::DisableSymbolRenaming);
req.platform = UnsafeUnkeyedValue(GetDevice()->GetPlatform());
req.substituteOverrideConfig = std::move(substituteOverrideConfig);
- req.maxSubgroupSizeForFullSubgroups = maxSubgroupSizeForFullSubgroups;
req.tintOptions.statically_paired_texture_binding_points =
std::move(statically_paired_texture_binding_points);
req.tintOptions.clamp_frag_depth = clampFragDepth;
@@ -440,8 +434,7 @@
if (r.stage == SingleShaderStage::Compute) {
Extent3D _;
DAWN_TRY_ASSIGN(_, ValidateComputeStageWorkgroupSize(
- program, remappedEntryPoint.c_str(), r.limits,
- r.maxSubgroupSizeForFullSubgroups));
+ program, remappedEntryPoint.c_str(), r.limits));
}
TRACE_EVENT0(r.platform.UnsafeGetValue(), General, "tint::spirv::writer::Generate()");
diff --git a/src/dawn/native/vulkan/ShaderModuleVk.h b/src/dawn/native/vulkan/ShaderModuleVk.h
index aad21e5..a4acf8d 100644
--- a/src/dawn/native/vulkan/ShaderModuleVk.h
+++ b/src/dawn/native/vulkan/ShaderModuleVk.h
@@ -51,7 +51,6 @@
uintptr_t layoutPtr;
std::string entryPoint;
PipelineConstantEntries constants;
- std::optional<uint32_t> maxSubgroupSizeForFullSubgroups;
bool emitPointSize;
bool operator==(const TransformedShaderModuleCacheKey& other) const;
@@ -81,13 +80,11 @@
ShaderModuleParseResult* parseResult,
OwnedCompilationMessages* compilationMessages);
- ResultOrError<ModuleAndSpirv> GetHandleAndSpirv(
- SingleShaderStage stage,
- const ProgrammableStage& programmableStage,
- const PipelineLayout* layout,
- bool clampFragDepth,
- bool emitPointSize,
- std::optional<uint32_t> maxSubgroupSizeForFullSubgroups);
+ ResultOrError<ModuleAndSpirv> GetHandleAndSpirv(SingleShaderStage stage,
+ const ProgrammableStage& programmableStage,
+ const PipelineLayout* layout,
+ bool clampFragDepth,
+ bool emitPointSize);
private:
ShaderModule(Device* device,
diff --git a/src/dawn/node/binding/Converter.cpp b/src/dawn/node/binding/Converter.cpp
index d4034fe..1f2e38f 100644
--- a/src/dawn/node/binding/Converter.cpp
+++ b/src/dawn/node/binding/Converter.cpp
@@ -1526,9 +1526,6 @@
case interop::GPUFeatureName::kDualSourceBlending:
out = wgpu::FeatureName::DualSourceBlending;
return true;
- case interop::GPUFeatureName::kChromiumExperimentalSubgroups:
- out = wgpu::FeatureName::ChromiumExperimentalSubgroups;
- return true;
case interop::GPUFeatureName::kChromiumExperimentalSubgroupUniformControlFlow:
out = wgpu::FeatureName::ChromiumExperimentalSubgroupUniformControlFlow;
return true;
@@ -1553,7 +1550,6 @@
switch (in) {
CASE(BGRA8UnormStorage, kBgra8UnormStorage);
- CASE(ChromiumExperimentalSubgroups, kChromiumExperimentalSubgroups);
CASE(ChromiumExperimentalSubgroupUniformControlFlow,
kChromiumExperimentalSubgroupUniformControlFlow);
CASE(Depth32FloatStencil8, kDepth32FloatStencil8);
diff --git a/src/dawn/node/binding/GPUAdapter.cpp b/src/dawn/node/binding/GPUAdapter.cpp
index 2b38434..8b4cab9 100644
--- a/src/dawn/node/binding/GPUAdapter.cpp
+++ b/src/dawn/node/binding/GPUAdapter.cpp
@@ -109,9 +109,7 @@
wgpu::ChainedStructOut** limitsListTail = &limits.nextInChain;
// Query the subgroup limits only if subgroups feature is available on the adapter.
- // TODO(349125474): Remove deprecated ChromiumExperimentalSubgroups.
- if (wgpuAdapter.HasFeature(FeatureName::Subgroups) ||
- wgpuAdapter.HasFeature(FeatureName::ChromiumExperimentalSubgroups)) {
+ if (wgpuAdapter.HasFeature(FeatureName::Subgroups)) {
InsertInChain(&subgroupLimits);
}
diff --git a/src/dawn/node/binding/GPUDevice.cpp b/src/dawn/node/binding/GPUDevice.cpp
index 473bc18..b64475e 100644
--- a/src/dawn/node/binding/GPUDevice.cpp
+++ b/src/dawn/node/binding/GPUDevice.cpp
@@ -193,9 +193,7 @@
};
// Query the subgroup limits only if subgroups feature is enabled on the device.
- // TODO(349125474): Remove deprecated ChromiumExperimentalSubgroups.
- if (device_.HasFeature(wgpu::FeatureName::Subgroups) ||
- device_.HasFeature(wgpu::FeatureName::ChromiumExperimentalSubgroups)) {
+ if (device_.HasFeature(wgpu::FeatureName::Subgroups)) {
InsertInChain(&subgroupLimits);
}
diff --git a/src/dawn/node/interop/DawnExtensions.idl b/src/dawn/node/interop/DawnExtensions.idl
index 583f204..43e2f06 100644
--- a/src/dawn/node/interop/DawnExtensions.idl
+++ b/src/dawn/node/interop/DawnExtensions.idl
@@ -29,7 +29,6 @@
// upstream webgpu.idl.
enum GPUFeatureName {
- "chromium-experimental-subgroups",
"chromium-experimental-subgroup-uniform-control-flow",
// subgroups, subgroups-f16 and multi-draw-indirect features are not in webgpu.idl yet.
"subgroups",
diff --git a/src/dawn/tests/DawnTest.cpp b/src/dawn/tests/DawnTest.cpp
index f9c926f9..6bd0220 100644
--- a/src/dawn/tests/DawnTest.cpp
+++ b/src/dawn/tests/DawnTest.cpp
@@ -1107,11 +1107,6 @@
for (uint32_t i = 0; i < descriptor.requiredFeatureCount; ++i) {
requiredFeatureSet.insert(descriptor.requiredFeatures[i]);
}
- // ChromiumExperimentalSubgroups feature is deprecated.
- // TODO(349125474): Remove deprecated ChromiumExperimentalSubgroups.
- if (requiredFeatureSet.count(wgpu::FeatureName::ChromiumExperimentalSubgroups)) {
- expectedDeprecatedCount++;
- }
return expectedDeprecatedCount;
}
diff --git a/src/dawn/tests/end2end/SubgroupsTests.cpp b/src/dawn/tests/end2end/SubgroupsTests.cpp
index 8142acf..eb7c4ba 100644
--- a/src/dawn/tests/end2end/SubgroupsTests.cpp
+++ b/src/dawn/tests/end2end/SubgroupsTests.cpp
@@ -53,26 +53,16 @@
mRequiredShaderF16Feature = true;
requiredFeatures.push_back(wgpu::FeatureName::ShaderF16);
}
-
- // Require either ChromiumExperimentalSubgroups or Subgroups/F16, but not both of them, so
- // that we can test the code path not involving ChromiumExperimentalSubgroups.
- if (GetParam().mUseChromiumExperimentalSubgroups) {
- if (SupportsFeatures({wgpu::FeatureName::ChromiumExperimentalSubgroups})) {
- mRequiredChromiumExperimentalSubgroups = true;
- requiredFeatures.push_back(wgpu::FeatureName::ChromiumExperimentalSubgroups);
- }
- } else {
- if (SupportsFeatures({wgpu::FeatureName::Subgroups})) {
- mRequiredSubgroupsFeature = true;
- requiredFeatures.push_back(wgpu::FeatureName::Subgroups);
- }
- if (SupportsFeatures({wgpu::FeatureName::SubgroupsF16})) {
- // SubgroupsF16 feature could be supported only if ShaderF16 and Subgroups features
- // are also supported.
- DAWN_ASSERT(mRequiredShaderF16Feature && mRequiredSubgroupsFeature);
- mRequiredSubgroupsF16Feature = true;
- requiredFeatures.push_back(wgpu::FeatureName::SubgroupsF16);
- }
+ if (SupportsFeatures({wgpu::FeatureName::Subgroups})) {
+ mRequiredSubgroupsFeature = true;
+ requiredFeatures.push_back(wgpu::FeatureName::Subgroups);
+ }
+ if (SupportsFeatures({wgpu::FeatureName::SubgroupsF16})) {
+ // SubgroupsF16 feature could be supported only if ShaderF16 and Subgroups features
+ // are also supported.
+ DAWN_ASSERT(mRequiredShaderF16Feature && mRequiredSubgroupsFeature);
+ mRequiredSubgroupsF16Feature = true;
+ requiredFeatures.push_back(wgpu::FeatureName::SubgroupsF16);
}
mSubgroupsF16SupportedByBackend = SupportsFeatures({wgpu::FeatureName::SubgroupsF16});
@@ -85,47 +75,29 @@
if (mRequiredShaderF16Feature) {
code << "enable f16;";
}
- if (GetParam().mUseChromiumExperimentalSubgroups) {
- code << "enable chromium_experimental_subgroups;";
- } else {
- if (mRequiredSubgroupsFeature) {
- code << "enable subgroups;";
- }
- if (mRequiredSubgroupsF16Feature) {
- code << "enable subgroups_f16;";
- }
+ if (mRequiredSubgroupsFeature) {
+ code << "enable subgroups;";
+ }
+ if (mRequiredSubgroupsF16Feature) {
+ code << "enable subgroups_f16;";
}
return code;
}
bool IsShaderF16EnabledInWGSL() const { return mRequiredShaderF16Feature; }
- bool IsSubgroupsEnabledInWGSL() const {
- return mRequiredSubgroupsFeature || mRequiredChromiumExperimentalSubgroups;
- }
- bool IsSubgroupsF16EnabledInWGSL() const {
- return mRequiredSubgroupsF16Feature || mRequiredChromiumExperimentalSubgroups;
- }
- bool IsChromiumExperimentalSubgroupsRequired() const {
- return mRequiredChromiumExperimentalSubgroups;
- }
+ bool IsSubgroupsEnabledInWGSL() const { return mRequiredSubgroupsFeature; }
+ bool IsSubgroupsF16EnabledInWGSL() const { return mRequiredSubgroupsF16Feature; }
bool IsSubgroupsF16SupportedByBackend() const { return mSubgroupsF16SupportedByBackend; }
private:
bool mRequiredShaderF16Feature = false;
bool mRequiredSubgroupsFeature = false;
bool mRequiredSubgroupsF16Feature = false;
- bool mRequiredChromiumExperimentalSubgroups = false;
- // Indicates that backend actually supports using subgroups functions with f16 types. Note that
- // using ChromiumExperimentalSubgroups allows subgroups_f16 extension in WGSL, but does not
- // ensure that backend supports using it.
+ // Indicates that backend actually supports using subgroups functions with f16 types.
bool mSubgroupsF16SupportedByBackend = false;
};
-using UseChromiumExperimentalSubgroups = bool;
-DAWN_TEST_PARAM_STRUCT(SubgroupsShaderTestsParams, UseChromiumExperimentalSubgroups);
-
-class SubgroupsShaderTests
- : public SubgroupsTestsBase<SubgroupsShaderTestsParams> {
+class SubgroupsShaderTests : public SubgroupsTestsBase<AdapterTestParam> {
protected:
// Testing reading subgroup_size. The shader declares a workgroup size of [workgroupSize, 1, 1],
// in which each invocation read the workgroup_size built-in value and write back to output
@@ -242,13 +214,13 @@
}
// DawnTestBase::CreateDeviceImpl always enables allow_unsafe_apis toggle.
-DAWN_INSTANTIATE_TEST_P(SubgroupsShaderTests,
- {D3D12Backend(), D3D12Backend({}, {"use_dxc"}), MetalBackend(),
- VulkanBackend()},
- {false, true} // UseChromiumExperimentalSubgroups
-);
+DAWN_INSTANTIATE_TEST(SubgroupsShaderTests,
+ D3D12Backend(),
+ D3D12Backend({}, {"use_dxc"}),
+ MetalBackend(),
+ VulkanBackend());
-class SubgroupsShaderTestsFragment : public SubgroupsTestsBase<SubgroupsShaderTestsParams> {
+class SubgroupsShaderTestsFragment : public SubgroupsTestsBase<AdapterTestParam> {
protected:
// Testing reading subgroup_size in fragment shader. There is no workgroup size here and
// subgroup_size is varying.
@@ -354,11 +326,11 @@
}
// DawnTestBase::CreateDeviceImpl always enables allow_unsafe_apis toggle.
-DAWN_INSTANTIATE_TEST_P(SubgroupsShaderTestsFragment,
- {D3D12Backend(), D3D12Backend({}, {"use_dxc"}), MetalBackend(),
- VulkanBackend()},
- {false, true} // UseChromiumExperimentalSubgroups
-);
+DAWN_INSTANTIATE_TEST(SubgroupsShaderTestsFragment,
+ D3D12Backend(),
+ D3D12Backend({}, {"use_dxc"}),
+ MetalBackend(),
+ VulkanBackend());
enum class BroadcastType {
I32,
@@ -406,9 +378,7 @@
return o;
}
-using UseChromiumExperimentalSubgroups = bool;
DAWN_TEST_PARAM_STRUCT(SubgroupsBroadcastTestsParams,
- UseChromiumExperimentalSubgroups,
BroadcastType,
SubgroupBroadcastValueOfInvocation0);
@@ -614,15 +584,6 @@
DAWN_TEST_UNSUPPORTED_IF(!IsSubgroupsEnabledInWGSL());
}
- if (IsChromiumExperimentalSubgroupsRequired()) {
- // Adreno 640 does not support subgroups in the fragment stage and therefore will not
- // actually be supported for the subgroup feature in WGSL. In addition to missing the
- // fragment stage subgroups also appear to have implementation issues in compute for this
- // device. See crbug/351745820
- DAWN_SUPPRESS_TEST_IF(gpu_info::IsQualcomm_PCIAdreno6xx(
- GetParam().adapterProperties.vendorID, GetParam().adapterProperties.deviceID));
- }
-
for (uint32_t workgroupSize : {1, 2, 3, 4, 7, 8, 15, 16, 31, 32, 63, 64, 127, 128, 255, 256}) {
TestBroadcastSubgroupSize(workgroupSize);
}
@@ -632,7 +593,6 @@
DAWN_INSTANTIATE_TEST_P(SubgroupsBroadcastTests,
{D3D12Backend(), D3D12Backend({}, {"use_dxc"}), MetalBackend(),
VulkanBackend()},
- {false, true}, // UseChromiumExperimentalSubgroups
{
BroadcastType::I32,
BroadcastType::U32,
@@ -644,169 +604,6 @@
// SubgroupBroadcastValueOfInvocation0
);
-using UseChromiumExperimentalSubgroups = bool;
-DAWN_TEST_PARAM_STRUCT(SubgroupsFullSubgroupsTestsParams,
- UseChromiumExperimentalSubgroups);
-
-class SubgroupsFullSubgroupsTests
- : public SubgroupsTestsBase<SubgroupsFullSubgroupsTestsParams> {
- protected:
- // Helper function that create shader module with subgroups extension required and a empty
- // compute entry point, named main, of given workgroup size
- wgpu::ShaderModule CreateShaderModuleWithSubgroupsRequired(WGPUExtent3D workgroupSize = {1, 1,
- 1}) {
- std::stringstream code;
-
- EnableExtensions(code) << R"(
- @compute @workgroup_size()"
- << workgroupSize.width << ", " << workgroupSize.height << ", "
- << workgroupSize.depthOrArrayLayers << R"()
- fn main() {}
-)";
- return utils::CreateShaderModule(device, code.str().c_str());
- }
-
- // Helper function that create shader module with subgroups extension required and a empty
- // compute entry point, named main, of workgroup size that are override constants.
- wgpu::ShaderModule CreateShaderModuleWithOverrideWorkgroupSize() {
- std::stringstream code;
- EnableExtensions(code) << R"(
- override wgs_x: u32;
- override wgs_y: u32;
- override wgs_z: u32;
-
- @compute @workgroup_size(wgs_x, wgs_y, wgs_z)
- fn main() {}
-)";
- return utils::CreateShaderModule(device, code.str().c_str());
- }
-
- struct TestCase {
- WGPUExtent3D workgroupSize;
- bool isFullSubgroups;
- };
-
- // Helper function that generate workgroup size cases for full subgroups test, based on device
- // reported max subgroup size.
- std::vector<TestCase> GenerateFullSubgroupsWorkgroupSizeCases() {
- wgpu::SupportedLimits limits{};
- wgpu::DawnExperimentalSubgroupLimits subgroupLimits{};
- limits.nextInChain = &subgroupLimits;
- EXPECT_EQ(device.GetLimits(&limits), wgpu::Status::Success);
- uint32_t maxSubgroupSize = subgroupLimits.maxSubgroupSize;
- EXPECT_TRUE(1 <= maxSubgroupSize && maxSubgroupSize <= 128);
- // maxSubgroupSize should be a power of 2.
- EXPECT_TRUE(IsPowerOfTwo(maxSubgroupSize));
-
- std::vector<TestCase> cases;
-
- // workgroup_size.x = maxSubgroupSize, is a multiple of maxSubgroupSize.
- cases.push_back({{maxSubgroupSize, 1, 1}, true});
- // Note that maxSubgroupSize is no larger than 128, so threads in the wrokgroups below is no
- // more than 256, fits in the maxComputeInvocationsPerWorkgroup limit which is at least 256.
- cases.push_back({{maxSubgroupSize * 2, 1, 1}, true});
- cases.push_back({{maxSubgroupSize, 2, 1}, true});
- cases.push_back({{maxSubgroupSize, 1, 2}, true});
-
- EXPECT_TRUE(maxSubgroupSize >= 4);
- // workgroup_size.x = maxSubgroupSize / 2, not a multiple of maxSubgroupSize.
- cases.push_back({{maxSubgroupSize / 2, 1, 1}, false});
- cases.push_back({{maxSubgroupSize / 2, 2, 1}, false});
- // workgroup_size.x = maxSubgroupSize - 1, not a multiple of maxSubgroupSize.
- cases.push_back({{maxSubgroupSize - 1, 1, 1}, false});
- // workgroup_size.x = maxSubgroupSize * 2 - 1, not a multiple of maxSubgroupSize if
- // maxSubgroupSize > 1.
- cases.push_back({{maxSubgroupSize * 2 - 1, 1, 1}, false});
- // workgroup_size.x = 1, not a multiple of maxSubgroupSize. Test that validation
- // checks the x dimension of workgroup size instead of others.
- cases.push_back({{1, maxSubgroupSize, 1}, false});
-
- return cases;
- }
-};
-
-// Test that creating compute pipeline with full subgroups required will validate the workgroup size
-// as expected, when using compute shader with literal workgroup size.
-TEST_P(SubgroupsFullSubgroupsTests,
- ComputePipelineRequiringFullSubgroupsWithLiteralWorkgroupSize) {
- // Currently DawnComputePipelineFullSubgroups only supported with ChromiumExperimentalSubgroups
- // enabled.
- DAWN_TEST_UNSUPPORTED_IF(!IsChromiumExperimentalSubgroupsRequired());
-
- // Keep all success compute pipeline alive, so that we can test the compute pipeline cache.
- std::vector<wgpu::ComputePipeline> computePipelines;
-
- for (const TestCase& c : GenerateFullSubgroupsWorkgroupSizeCases()) {
- // Reuse the shader module for both not requiring and requiring full subgroups cases, to
- // test that cached compute pipeline will not be used unexpectedly.
- auto shaderModule = CreateShaderModuleWithSubgroupsRequired(c.workgroupSize);
- for (bool requiresFullSubgroups : {false, true}) {
- wgpu::ComputePipelineDescriptor csDesc;
- csDesc.compute.module = shaderModule;
-
- wgpu::DawnComputePipelineFullSubgroups fullSubgroupsOption;
- fullSubgroupsOption.requiresFullSubgroups = requiresFullSubgroups;
- csDesc.nextInChain = &fullSubgroupsOption;
-
- // It should be a validation error if full subgroups is required but given workgroup
- // size does not fit.
- if (requiresFullSubgroups && !c.isFullSubgroups) {
- ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&csDesc));
- } else {
- // Otherwise, creating compute pipeline should succeed.
- computePipelines.push_back(device.CreateComputePipeline(&csDesc));
- }
- }
- }
-}
-
-// Test that creating compute pipeline with full subgroups required will validate the workgroup size
-// as expected, when using compute shader with override constants workgroup size.
-TEST_P(SubgroupsFullSubgroupsTests,
- ComputePipelineRequiringFullSubgroupsWithOverrideWorkgroupSize) {
- // Currently DawnComputePipelineFullSubgroups only supported with ChromiumExperimentalSubgroups
- // enabled.
- DAWN_TEST_UNSUPPORTED_IF(!IsChromiumExperimentalSubgroupsRequired());
- // Reuse the same shader module for all case to test the validation happened as expected.
- auto shaderModule = CreateShaderModuleWithOverrideWorkgroupSize();
- // Keep all success compute pipeline alive, so that we can test the compute pipeline cache.
- std::vector<wgpu::ComputePipeline> computePipelines;
-
- for (const TestCase& c : GenerateFullSubgroupsWorkgroupSizeCases()) {
- for (bool requiresFullSubgroups : {false, true}) {
- std::vector<wgpu::ConstantEntry> constants{
- {nullptr, "wgs_x", static_cast<double>(c.workgroupSize.width)},
- {nullptr, "wgs_y", static_cast<double>(c.workgroupSize.height)},
- {nullptr, "wgs_z", static_cast<double>(c.workgroupSize.depthOrArrayLayers)},
- };
-
- wgpu::ComputePipelineDescriptor csDesc;
- csDesc.compute.module = shaderModule;
- csDesc.compute.constants = constants.data();
- csDesc.compute.constantCount = constants.size();
-
- wgpu::DawnComputePipelineFullSubgroups fullSubgroupsOption;
- fullSubgroupsOption.requiresFullSubgroups = requiresFullSubgroups;
- csDesc.nextInChain = &fullSubgroupsOption;
-
- // It should be a validation error if full subgroups is required but given workgroup
- // size does not fit.
- if (requiresFullSubgroups && !c.isFullSubgroups) {
- ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&csDesc));
- } else {
- // Otherwise, creating compute pipeline should succeed.
- computePipelines.push_back(device.CreateComputePipeline(&csDesc));
- }
- }
- }
-}
-
-// DawnTestBase::CreateDeviceImpl always enables allow_unsafe_apis toggle.
-DAWN_INSTANTIATE_TEST_P(SubgroupsFullSubgroupsTests,
- {D3D12Backend(), D3D12Backend({}, {"use_dxc"}), MetalBackend(),
- VulkanBackend()},
- {false, true} // UseChromiumExperimentalSubgroups
-);
// Core functions that may be polyfilled
enum class SubgroupIntrinsicOp : uint8_t {
@@ -852,7 +649,6 @@
}
DAWN_TEST_PARAM_STRUCT(SubgroupsShaderInclusiveTestsParams,
- UseChromiumExperimentalSubgroups,
SubgroupIntrinsicOp,
SubgroupOpDataType);
@@ -993,15 +789,6 @@
DAWN_TEST_UNSUPPORTED_IF(!IsSubgroupsEnabledInWGSL());
}
- if (IsChromiumExperimentalSubgroupsRequired()) {
- // Adreno 640 does not support subgroups in the fragment stage and therefore will not
- // actually be supported for the subgroup feature in WGSL. In addition to missing the
- // fragment stage subgroups also appear to have implementation issues in compute for this
- // device. See crbug/351745820
- DAWN_SUPPRESS_TEST_IF(gpu_info::IsQualcomm_PCIAdreno6xx(
- GetParam().adapterProperties.vendorID, GetParam().adapterProperties.deviceID));
- }
-
for (uint32_t workgroupSize : {1, 2, 3, 4, 7, 8, 15, 16, 31, 32, 63, 64, 127, 128, 255, 256}) {
TestReadSubgroupSize(workgroupSize);
}
@@ -1010,7 +797,6 @@
DAWN_INSTANTIATE_TEST_P(SubgroupsShaderInclusiveTest,
{D3D12Backend(), D3D12Backend({}, {"use_dxc"}), MetalBackend(),
VulkanBackend()},
- {false, true}, // UseChromiumExperimentalSubgroups
{SubgroupIntrinsicOp::Add, SubgroupIntrinsicOp::Mul},
{
SubgroupOpDataType::F32,
diff --git a/src/dawn/tests/unittests/validation/ComputeValidationTests.cpp b/src/dawn/tests/unittests/validation/ComputeValidationTests.cpp
index 6221aa1..bc54e0c 100644
--- a/src/dawn/tests/unittests/validation/ComputeValidationTests.cpp
+++ b/src/dawn/tests/unittests/validation/ComputeValidationTests.cpp
@@ -69,158 +69,6 @@
ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&csDesc));
}
-// Test that creating a compute pipeline with chained DawnComputePipelineFullSubgroups on a device
-// that don't enable ChromiumExperimentalSubgroups feature fails.
-TEST_F(ComputePipelineValidationTest, UnexpectedDawnComputePipelineFullSubgroups) {
- auto computeModule = CreateShaderModule();
-
- wgpu::ComputePipelineDescriptor csDesc;
- csDesc.compute.module = computeModule;
-
- wgpu::DawnComputePipelineFullSubgroups subgroupOptions;
- subgroupOptions.requiresFullSubgroups = false;
- csDesc.nextInChain = &subgroupOptions;
-
- ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&csDesc));
-}
-
-// Tests that requiring ChromiumExperimentalSubgroups feature, for DawnComputePipelineFullSubgroups
-// testing.
-// TODO(349125474): Revisit these tests when removing deprecated ChromiumExperimentalSubgroups.
-class ComputePipelineValidationTestWithChromiumExperimentalSubgroupsFeatureEnabled
- : public ComputePipelineValidationTest {
- protected:
- std::vector<wgpu::FeatureName> GetRequiredFeatures() override {
- return {wgpu::FeatureName::ChromiumExperimentalSubgroups};
- }
-
- // Helper function that create a shader module with compute entry point named main and
- // workgroup size with override constants (wgs_x, wgs_y, wgs_z).
- wgpu::ShaderModule CreateShaderModuleWithOverrideWorkgroupSize() {
- // Note that we don't need to require subgroups WGSL extension in the shader since we don't
- // use subgroup built-in in this empty entry point.
- return utils::CreateShaderModule(device, R"(
- override wgs_x: u32;
- override wgs_y: u32;
- override wgs_z: u32;
-
- @compute @workgroup_size(wgs_x, wgs_y, wgs_z)
- fn main() {
- })");
- }
-};
-
-// Test that creating a compute pipeline with basic shader module and chained
-// DawnComputePipelineFullSubgroups not requiring fullSubgroups succeeds.
-TEST_F(ComputePipelineValidationTestWithChromiumExperimentalSubgroupsFeatureEnabled,
- DawnComputePipelineFullSubgroupsNotRequired) {
- auto computeModule = CreateShaderModule();
-
- wgpu::PipelineLayout pl = utils::MakeBasicPipelineLayout(device, nullptr);
-
- wgpu::ComputePipelineDescriptor csDesc;
- csDesc.layout = pl;
- csDesc.compute.module = computeModule;
-
- wgpu::DawnComputePipelineFullSubgroups subgroupOptions;
- subgroupOptions.requiresFullSubgroups = false;
- csDesc.nextInChain = &subgroupOptions;
-
- device.CreateComputePipeline(&csDesc);
-}
-
-// Test that creating a compute pipeline with basic shader module and chained
-// DawnComputePipelineFullSubgroups requiring fullSubgroups fails if x dimension of workgroup size
-// is not a multiple of maxSubgroupSize. Note that ValidationTest use Null backend, which assume a
-// maxSubgroupSize of 128.
-TEST_F(ComputePipelineValidationTestWithChromiumExperimentalSubgroupsFeatureEnabled,
- DawnComputePipelineFullSubgroupsRequired_WorkgroupSizeInvalid) {
- // Can not require full subgroups with workgroup size {127, 1, 1}
- auto computeModule = CreateShaderModule(127);
-
- wgpu::ComputePipelineDescriptor csDesc;
- csDesc.compute.module = computeModule;
-
- wgpu::DawnComputePipelineFullSubgroups subgroupOptions;
- subgroupOptions.requiresFullSubgroups = true;
- csDesc.nextInChain = &subgroupOptions;
-
- ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&csDesc));
-}
-
-// Test that creating a compute pipeline with basic shader module and chained
-// DawnComputePipelineFullSubgroups requiring fullSubgroups succeeds if x dimension of workgroup
-// size is a multiple of maxSubgroupSize. Note that ValidationTest use Null backend, which assume a
-// maxSubgroupSize of 128.
-TEST_F(ComputePipelineValidationTestWithChromiumExperimentalSubgroupsFeatureEnabled,
- DawnComputePipelineFullSubgroupsRequired_WorkgroupSizeValid) {
- // Can require full subgroups with workgroup size {128, 1, 1}
- auto computeModule = CreateShaderModule(128);
-
- wgpu::ComputePipelineDescriptor csDesc;
- csDesc.compute.module = computeModule;
-
- wgpu::DawnComputePipelineFullSubgroups subgroupOptions;
- subgroupOptions.requiresFullSubgroups = true;
- csDesc.nextInChain = &subgroupOptions;
-
- device.CreateComputePipeline(&csDesc);
-}
-
-// Test that creating a compute pipeline with override workgroup size shader module and chained
-// DawnComputePipelineFullSubgroups requiring fullSubgroups fails if x dimension of workgroup size
-// is not a multiple of maxSubgroupSize. Note that ValidationTest use Null backend, which assume a
-// maxSubgroupSize of 128.
-TEST_F(ComputePipelineValidationTestWithChromiumExperimentalSubgroupsFeatureEnabled,
- DawnComputePipelineFullSubgroupsRequired_OverrideWorkgroupSizeInvalid) {
- auto computeModule = CreateShaderModuleWithOverrideWorkgroupSize();
-
- // Can not require full subgroups with workgroup size {127, 1, 1}
- std::vector<wgpu::ConstantEntry> constants{
- {nullptr, "wgs_x", 127},
- {nullptr, "wgs_y", 1},
- {nullptr, "wgs_z", 1},
- };
-
- wgpu::ComputePipelineDescriptor csDesc;
- csDesc.compute.module = computeModule;
- csDesc.compute.constants = constants.data();
- csDesc.compute.constantCount = constants.size();
-
- wgpu::DawnComputePipelineFullSubgroups subgroupOptions;
- subgroupOptions.requiresFullSubgroups = true;
- csDesc.nextInChain = &subgroupOptions;
-
- ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&csDesc));
-}
-
-// Test that creating a compute pipeline with override workgroup size shader module and chained
-// DawnComputePipelineFullSubgroups requiring fullSubgroups succeeds if x dimension of workgroup
-// size is a multiple of maxSubgroupSize. Note that ValidationTest use Null backend, which assume a
-// maxSubgroupSize of 128.
-TEST_F(ComputePipelineValidationTestWithChromiumExperimentalSubgroupsFeatureEnabled,
- DawnComputePipelineFullSubgroupsRequired_OverrideWorkgroupSizeValid) {
- auto computeModule = CreateShaderModuleWithOverrideWorkgroupSize();
-
- // Can require full subgroups with workgroup size {128, 1, 1}
- std::vector<wgpu::ConstantEntry> constants{
- {nullptr, "wgs_x", 128},
- {nullptr, "wgs_y", 1},
- {nullptr, "wgs_z", 1},
- };
-
- wgpu::ComputePipelineDescriptor csDesc;
- csDesc.compute.module = computeModule;
- csDesc.compute.constants = constants.data();
- csDesc.compute.constantCount = constants.size();
-
- wgpu::DawnComputePipelineFullSubgroups subgroupOptions;
- subgroupOptions.requiresFullSubgroups = true;
- csDesc.nextInChain = &subgroupOptions;
-
- device.CreateComputePipeline(&csDesc);
-}
-
// TODO(cwallez@chromium.org): Add a regression test for Disptach validation trying to access the
// input state.
diff --git a/src/dawn/tests/unittests/validation/DeviceValidationTests.cpp b/src/dawn/tests/unittests/validation/DeviceValidationTests.cpp
index ebbfee9..5f165f9 100644
--- a/src/dawn/tests/unittests/validation/DeviceValidationTests.cpp
+++ b/src/dawn/tests/unittests/validation/DeviceValidationTests.cpp
@@ -247,44 +247,37 @@
TEST_F(RequestDeviceValidationTest, SubgroupsF16FeatureDependency) {
for (bool requireShaderF16 : {false, true}) {
for (bool requireSubgroups : {false, true}) {
- // TODO(349125474): Remove deprecated ChromiumExperimentalSubgroups.
- for (bool requireChromiumExperimentalSubgroups : {false, true}) {
- std::vector<wgpu::FeatureName> features;
- if (requireShaderF16) {
- features.push_back(wgpu::FeatureName::ShaderF16);
- }
- if (requireSubgroups) {
- features.push_back(wgpu::FeatureName::Subgroups);
- }
- if (requireChromiumExperimentalSubgroups) {
- features.push_back(wgpu::FeatureName::ChromiumExperimentalSubgroups);
- }
- features.push_back(wgpu::FeatureName::SubgroupsF16);
-
- wgpu::DeviceDescriptor descriptor;
- descriptor.requiredFeatureCount = features.size();
- descriptor.requiredFeatures = features.data();
-
- // Device request with subgroups-f16 feature can only success if shader-f16 feature
- // and subgroups or chromium-experimental-subgroups features are required as well.
- const bool isSuccess =
- (requireSubgroups || requireChromiumExperimentalSubgroups) && requireShaderF16;
-
- if (isSuccess) {
- EXPECT_CALL(mRequestDeviceCallback, Call(wgpu::RequestDeviceStatus::Success,
- NotNull(), EmptySizedString()))
- .Times(1);
- } else {
- EXPECT_CALL(mRequestDeviceCallback, Call(wgpu::RequestDeviceStatus::Error,
- IsNull(), NonEmptySizedString()))
- .Times(1);
- }
-
- EXPECT_DEPRECATION_WARNINGS(
- adapter.RequestDevice(&descriptor, wgpu::CallbackMode::AllowSpontaneous,
- mRequestDeviceCallback.Callback()),
- GetDeviceCreationDeprecationWarningExpectation(descriptor));
+ std::vector<wgpu::FeatureName> features;
+ if (requireShaderF16) {
+ features.push_back(wgpu::FeatureName::ShaderF16);
}
+ if (requireSubgroups) {
+ features.push_back(wgpu::FeatureName::Subgroups);
+ }
+ features.push_back(wgpu::FeatureName::SubgroupsF16);
+
+ wgpu::DeviceDescriptor descriptor;
+ descriptor.requiredFeatureCount = features.size();
+ descriptor.requiredFeatures = features.data();
+
+ // Device request with subgroups-f16 feature can only success if shader-f16 feature
+ // and subgroups features are required as well.
+ const bool isSuccess = requireSubgroups && requireShaderF16;
+
+ if (isSuccess) {
+ EXPECT_CALL(mRequestDeviceCallback,
+ Call(wgpu::RequestDeviceStatus::Success, NotNull(), EmptySizedString()))
+ .Times(1);
+ } else {
+ EXPECT_CALL(mRequestDeviceCallback,
+ Call(wgpu::RequestDeviceStatus::Error, IsNull(), NonEmptySizedString()))
+ .Times(1);
+ }
+
+ EXPECT_DEPRECATION_WARNINGS(
+ adapter.RequestDevice(&descriptor, wgpu::CallbackMode::AllowSpontaneous,
+ mRequestDeviceCallback.Callback()),
+ GetDeviceCreationDeprecationWarningExpectation(descriptor));
}
}
}
diff --git a/src/dawn/tests/unittests/validation/ShaderModuleValidationTests.cpp b/src/dawn/tests/unittests/validation/ShaderModuleValidationTests.cpp
index 3b59c80..6a2fd93 100644
--- a/src/dawn/tests/unittests/validation/ShaderModuleValidationTests.cpp
+++ b/src/dawn/tests/unittests/validation/ShaderModuleValidationTests.cpp
@@ -757,7 +757,6 @@
{"f16", false, {"shader-f16"}, {}},
{"clip_distances", false, {"clip-distances"}, {}},
{"dual_source_blending", false, {"dual-source-blending"}, {}},
- {"chromium_experimental_subgroups", true, {"chromium-experimental-subgroups"}, {}},
{"subgroups", false, {"subgroups"}, {}},
{"subgroups_f16", false, {"shader-f16", "subgroups", "subgroups-f16"}, {"f16", "subgroups"}},
{"chromium_experimental_pixel_local", true, {"pixel-local-storage-coherent"}, {}},
diff --git a/src/dawn/tests/unittests/validation/ValidationTest.cpp b/src/dawn/tests/unittests/validation/ValidationTest.cpp
index a21e080..78b8998 100644
--- a/src/dawn/tests/unittests/validation/ValidationTest.cpp
+++ b/src/dawn/tests/unittests/validation/ValidationTest.cpp
@@ -319,11 +319,6 @@
for (uint32_t i = 0; i < descriptor.requiredFeatureCount; ++i) {
requiredFeatureSet.insert(descriptor.requiredFeatures[i]);
}
- // ChromiumExperimentalSubgroups feature is deprecated.
- // TODO(349125474): Remove deprecated ChromiumExperimentalSubgroups.
- if (requiredFeatureSet.count(wgpu::FeatureName::ChromiumExperimentalSubgroups)) {
- expectedDeprecatedCount++;
- }
return expectedDeprecatedCount;
}
diff --git a/src/dawn/wire/SupportedFeatures.cpp b/src/dawn/wire/SupportedFeatures.cpp
index ff26aa2..02c2554 100644
--- a/src/dawn/wire/SupportedFeatures.cpp
+++ b/src/dawn/wire/SupportedFeatures.cpp
@@ -93,7 +93,6 @@
case WGPUFeatureName_MSAARenderToSingleSampled:
case WGPUFeatureName_DualSourceBlending:
case WGPUFeatureName_ANGLETextureSharing:
- case WGPUFeatureName_ChromiumExperimentalSubgroups:
case WGPUFeatureName_ChromiumExperimentalSubgroupUniformControlFlow:
case WGPUFeatureName_PixelLocalStorageCoherent:
case WGPUFeatureName_PixelLocalStorageNonCoherent:
diff --git a/src/tint/cmd/fuzz/wgsl/dictionary.txt b/src/tint/cmd/fuzz/wgsl/dictionary.txt
index b5e1a79..97c9fc4 100644
--- a/src/tint/cmd/fuzz/wgsl/dictionary.txt
+++ b/src/tint/cmd/fuzz/wgsl/dictionary.txt
@@ -162,7 +162,6 @@
"chromium_experimental_pixel_local"
"chromium_experimental_push_constant"
"chromium_experimental_subgroup_matrix"
-"chromium_experimental_subgroups"
"chromium_internal_graphite"
"chromium_internal_input_attachments"
"chromium_internal_relaxed_uniform_layout"
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 51d6955..43edbff 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
@@ -395,7 +395,6 @@
wgsl::Extension::kChromiumDisableUniformityAnalysis,
wgsl::Extension::kChromiumExperimentalPixelLocal,
wgsl::Extension::kChromiumExperimentalPushConstant,
- wgsl::Extension::kChromiumExperimentalSubgroups,
wgsl::Extension::kChromiumInternalGraphite,
wgsl::Extension::kClipDistances,
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 37ac38f..eafb0e6 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -287,7 +287,6 @@
wgsl::Extension::kChromiumDisableUniformityAnalysis,
wgsl::Extension::kChromiumExperimentalFramebufferFetch,
wgsl::Extension::kChromiumExperimentalPixelLocal,
- wgsl::Extension::kChromiumExperimentalSubgroups,
wgsl::Extension::kChromiumInternalGraphite,
wgsl::Extension::kChromiumInternalRelaxedUniformLayout,
wgsl::Extension::kClipDistances,
diff --git a/src/tint/lang/msl/writer/ast_raise/subgroup_ballot_test.cc b/src/tint/lang/msl/writer/ast_raise/subgroup_ballot_test.cc
index 2ab7d19..d5d6ab5 100644
--- a/src/tint/lang/msl/writer/ast_raise/subgroup_ballot_test.cc
+++ b/src/tint/lang/msl/writer/ast_raise/subgroup_ballot_test.cc
@@ -42,7 +42,7 @@
TEST_F(SubgroupBallotTest, DirectUse) {
auto* src = R"(
-enable chromium_experimental_subgroups;
+enable subgroups;
@compute @workgroup_size(64)
fn foo() {
@@ -53,7 +53,7 @@
auto* expect =
R"(
-enable chromium_experimental_subgroups;
+enable subgroups;
@internal(simd_ballot) @internal(disable_validation__function_has_no_body)
fn tint_msl_simd_ballot(pred : bool) -> vec2<u32>
@@ -84,7 +84,7 @@
TEST_F(SubgroupBallotTest, IndirectUse) {
auto* src = R"(
-enable chromium_experimental_subgroups;
+enable subgroups;
fn bar() -> vec4u {
let pred = true;
@@ -99,7 +99,7 @@
auto* expect =
R"(
-enable chromium_experimental_subgroups;
+enable subgroups;
@internal(simd_ballot) @internal(disable_validation__function_has_no_body)
fn tint_msl_simd_ballot(pred : bool) -> vec2<u32>
@@ -134,7 +134,7 @@
TEST_F(SubgroupBallotTest, PreexistingSubgroupSizeBuiltin) {
auto* src = R"(
-enable chromium_experimental_subgroups;
+enable subgroups;
@compute @workgroup_size(64)
fn foo(@builtin(workgroup_id) group_id: vec3u,
@@ -148,7 +148,7 @@
auto* expect =
R"(
-enable chromium_experimental_subgroups;
+enable subgroups;
@internal(simd_ballot) @internal(disable_validation__function_has_no_body)
fn tint_msl_simd_ballot(pred : bool) -> vec2<u32>
diff --git a/src/tint/lang/wgsl/ast/transform/canonicalize_entry_point_io_test.cc b/src/tint/lang/wgsl/ast/transform/canonicalize_entry_point_io_test.cc
index 1bd2d64..90a3bc1 100644
--- a/src/tint/lang/wgsl/ast/transform/canonicalize_entry_point_io_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/canonicalize_entry_point_io_test.cc
@@ -4339,7 +4339,7 @@
TEST_F(CanonicalizeEntryPointIOTest, SubgroupBuiltins_Hlsl) {
auto* src = R"(
-enable chromium_experimental_subgroups;
+enable subgroups;
@compute @workgroup_size(64)
fn frag_main(@builtin(subgroup_invocation_id) id : u32,
@@ -4349,7 +4349,7 @@
)";
auto* expect = R"(
-enable chromium_experimental_subgroups;
+enable subgroups;
@internal(intrinsic_wave_get_lane_index) @internal(disable_validation__function_has_no_body)
fn __WaveGetLaneIndex() -> u32
@@ -4376,7 +4376,7 @@
TEST_F(CanonicalizeEntryPointIOTest, SubgroupBuiltinsStruct_Hlsl) {
auto* src = R"(
-enable chromium_experimental_subgroups;
+enable subgroups;
struct Inputs {
@builtin(subgroup_invocation_id) id : u32,
@@ -4390,7 +4390,7 @@
)";
auto* expect = R"(
-enable chromium_experimental_subgroups;
+enable subgroups;
@internal(intrinsic_wave_get_lane_index) @internal(disable_validation__function_has_no_body)
fn __WaveGetLaneIndex() -> u32
diff --git a/src/tint/lang/wgsl/extension.cc b/src/tint/lang/wgsl/extension.cc
index 118d908..2510fec 100644
--- a/src/tint/lang/wgsl/extension.cc
+++ b/src/tint/lang/wgsl/extension.cc
@@ -57,9 +57,6 @@
if (str == "chromium_experimental_subgroup_matrix") {
return Extension::kChromiumExperimentalSubgroupMatrix;
}
- if (str == "chromium_experimental_subgroups") {
- return Extension::kChromiumExperimentalSubgroups;
- }
if (str == "chromium_internal_graphite") {
return Extension::kChromiumInternalGraphite;
}
@@ -101,8 +98,6 @@
return "chromium_experimental_push_constant";
case Extension::kChromiumExperimentalSubgroupMatrix:
return "chromium_experimental_subgroup_matrix";
- case Extension::kChromiumExperimentalSubgroups:
- return "chromium_experimental_subgroups";
case Extension::kChromiumInternalGraphite:
return "chromium_internal_graphite";
case Extension::kChromiumInternalInputAttachments:
diff --git a/src/tint/lang/wgsl/extension.h b/src/tint/lang/wgsl/extension.h
index 9c11792..cd480fb 100644
--- a/src/tint/lang/wgsl/extension.h
+++ b/src/tint/lang/wgsl/extension.h
@@ -51,7 +51,6 @@
kChromiumExperimentalPixelLocal,
kChromiumExperimentalPushConstant,
kChromiumExperimentalSubgroupMatrix,
- kChromiumExperimentalSubgroups,
kChromiumInternalGraphite,
kChromiumInternalInputAttachments,
kChromiumInternalRelaxedUniformLayout,
@@ -85,7 +84,6 @@
"chromium_experimental_pixel_local",
"chromium_experimental_push_constant",
"chromium_experimental_subgroup_matrix",
- "chromium_experimental_subgroups",
"chromium_internal_graphite",
"chromium_internal_input_attachments",
"chromium_internal_relaxed_uniform_layout",
@@ -103,7 +101,6 @@
Extension::kChromiumExperimentalPixelLocal,
Extension::kChromiumExperimentalPushConstant,
Extension::kChromiumExperimentalSubgroupMatrix,
- Extension::kChromiumExperimentalSubgroups,
Extension::kChromiumInternalGraphite,
Extension::kChromiumInternalInputAttachments,
Extension::kChromiumInternalRelaxedUniformLayout,
diff --git a/src/tint/lang/wgsl/extension_bench.cc b/src/tint/lang/wgsl/extension_bench.cc
index 8883e4d..2f8861d 100644
--- a/src/tint/lang/wgsl/extension_bench.cc
+++ b/src/tint/lang/wgsl/extension_bench.cc
@@ -80,69 +80,62 @@
"chromium_exeimoontal_subgroup_atrix",
"chromium_experimnal_subgroup_mazzrix",
"chro11ium_experienial_subgrppup_matrix",
- "chromium_experimental_subgroXXps",
- "chromium55eIIperimental_subgnno99ps",
- "chraamiuSS_experimentaHHr_subgrouYs",
- "chromium_experimental_subgroups",
- "chkkomium_eperimntal_subgroup",
- "jhromium_experRmental_subgogps",
- "chromiubexperiental_subgroups",
- "chromiumjinternal_graphite",
- "chromium_inernal_graphite",
- "cromiu_internaq_graphite",
+ "chXXomium_internal_graphite",
+ "chromi55m_internnal_gra99hiIIe",
+ "chSSomiuY_internal_aarHHphrrte",
"chromium_internal_graphite",
- "chromium_intenalNNgraphite",
- "chromiuminternal_gvaphite",
- "chromium_internal_grphitQQ",
- "chromirm_ffnternalinpt_attachments",
- "chromium_internal_input_attachmenjs",
- "chwwomiu2_interNNal_inpu_att8chments",
+ "kkhromium_nternal_rahHte",
+ "chromium_nRegnaj_graphite",
+ "chromium_ntebnal_gaphite",
+ "chromium_internal_input_atjachments",
+ "chromium_internal_inpt_attachments",
+ "chromium_nteral_iqput_attachments",
"chromium_internal_input_attachments",
- "chromium_internalinput_attachments",
- "crrromium_internal_input_attachments",
- "Ghromium_internal_input_attachments",
- "chromium_internalFFrelaxed_uniform_layout",
- "chromEum_internal_relaxed_unifrmlyout",
- "chromium_internalrrrelaxd_uniform_layout",
+ "chromium_internal_input_aNNtachents",
+ "chromium_internalinpt_attavvhments",
+ "chromium_internal_inut_attacQQments",
+ "chromirm_intenal_rfflaxed_unifrm_layout",
+ "chromium_internal_jelaxed_uniform_layout",
+ "chromium_interna_relNNxed_uwwiform_lay82t",
"chromium_internal_relaxed_uniform_layout",
- "chromiuminternal_relaxed_uniform_layut",
- "cXroDium_internal_rJJlaed_uniform_layout",
- "chromium_int8nal_relaed_uniform_layut",
- "clip_ksa11ce",
- "cli_distances",
- "Jlp_distances",
+ "chromium_internal_relaxed_uniform_layut",
+ "chromium_internal_relaxed_rrniform_layout",
+ "chromium_internal_relaxedGuniform_layout",
+ "clip_distanceFF",
+ "cEipdtances",
+ "cli_rristances",
"clip_distances",
- "clipcdistances",
- "clip_dOstances",
- "clip_disKK__vvttes",
- "d5l_so8rce_blendxxng",
- "dual_so__Fcebleqqing",
- "dual_source_bqqnding",
+ "lip_distanes",
+ "DXp_diJJtances",
+ "cl8pdistane",
+ "dul_okrc_blen11ing",
+ "dua_source_blending",
+ "duJl_source_blendig",
"dual_source_blending",
- "dual_so6633ce_blnding",
- "d9al_soource_bl6Qttding",
- "dual_source66blendin",
- "zzO6",
- "fyy6",
- "HZ",
+ "dual_source_clending",
+ "dual_sOurce_blending",
+ "dualKKs__urce_blttvnding",
+ "xx8",
+ "__F",
+ "f1q",
"f16",
- "4qW6",
- "fOO",
- "oY",
- "ubgrou",
- "subgrupF",
- "subgrwps",
+ "331O",
+ "ftt6QQ",
+ "666",
+ "zzxbO6rops",
+ "subgyyoups",
+ "HHugroZs",
"subgroups",
- "GbrKupff",
- "KKuqgroups",
- "subgmmou3s",
- "subgroupsf16",
- "subgroupsf1q",
- "sbgroupsbbf16",
+ "sWW44roupq",
+ "sOObgoups",
+ "sbgroYps",
+ "subroups_f",
+ "suFgoups_f16",
+ "subgowps_f16",
"subgroups_f16",
- "sbiiroups_f1",
- "sbgroqOps_f16",
- "TTubgroups_fvv6",
+ "suffgKups_f6",
+ "KKubgroqps_f16",
+ "subFroup3mmf16",
};
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 a8e92b2..518e08c 100644
--- a/src/tint/lang/wgsl/extension_test.cc
+++ b/src/tint/lang/wgsl/extension_test.cc
@@ -62,7 +62,6 @@
{"chromium_experimental_pixel_local", Extension::kChromiumExperimentalPixelLocal},
{"chromium_experimental_push_constant", Extension::kChromiumExperimentalPushConstant},
{"chromium_experimental_subgroup_matrix", Extension::kChromiumExperimentalSubgroupMatrix},
- {"chromium_experimental_subgroups", Extension::kChromiumExperimentalSubgroups},
{"chromium_internal_graphite", Extension::kChromiumInternalGraphite},
{"chromium_internal_input_attachments", Extension::kChromiumInternalInputAttachments},
{"chromium_internal_relaxed_uniform_layout", Extension::kChromiumInternalRelaxedUniformLayout},
@@ -89,33 +88,30 @@
{"cXromium_experimeggtal_subgrup_matrix", Extension::kUndefined},
{"chrXmiuu_expeimVtal_subgroup_matrix", Extension::kUndefined},
{"chromium_e3perimental_subgroup_matrix", Extension::kUndefined},
- {"cEromium_experimental_subgroups", Extension::kUndefined},
- {"TThromium_experiPPental_sugroups", Extension::kUndefined},
- {"chddomium_experimental_subgroxxs", Extension::kUndefined},
- {"chromi44m_internal_graphite", Extension::kUndefined},
- {"chromSSuVV_internal_graphite", Extension::kUndefined},
- {"cRromium_nternR22_graphite", Extension::kUndefined},
- {"chromium_int9rnaF_inpu_attachments", Extension::kUndefined},
- {"chrmium_internal_input_attachments", Extension::kUndefined},
- {"cOOromium_internVlHinput_ttachRRents", Extension::kUndefined},
- {"chromium_internl_relaxyd_uniform_layout", Extension::kUndefined},
- {"chromnnum_internrr77_Gelaxell_uniform_layout", Extension::kUndefined},
- {"chromium_intern4l_relaxe00_uniform_layout", Extension::kUndefined},
- {"cli_iootanes", Extension::kUndefined},
- {"czip_dstances", Extension::kUndefined},
- {"lipp_distan11es", Extension::kUndefined},
- {"XXual_source_blending", Extension::kUndefined},
- {"dual_sou5IIcennblend99ng", Extension::kUndefined},
- {"aSSual_soYrce_blerrdinHH", Extension::kUndefined},
- {"U", Extension::kUndefined},
- {"jV3", Extension::kUndefined},
- {"", Extension::kUndefined},
- {"uGGSrops", Extension::kUndefined},
- {"subgous", Extension::kUndefined},
- {"bgrups", Extension::kUndefined},
- {"subgrokp_jj6", Extension::kUndefined},
- {"ubgroupsffr6", Extension::kUndefined},
- {"subgroujs_f16", Extension::kUndefined},
+ {"chromium_internal_graphitE", Extension::kUndefined},
+ {"chromium_nternal_gTTPaphite", Extension::kUndefined},
+ {"chxxdomium_interal_graphite", Extension::kUndefined},
+ {"c44romium_internal_input_attachments", Extension::kUndefined},
+ {"chromium_internal_inputSSaVVtachments", Extension::kUndefined},
+ {"chrom22Rm_internal_input_atRchments", Extension::kUndefined},
+ {"chromium_int9rnal_relaxed_Fnifor_layout", Extension::kUndefined},
+ {"chrmium_internal_relaxed_uniform_layout", Extension::kUndefined},
+ {"VRhHomium_internal_relaxd_uniform_OOayout", Extension::kUndefined},
+ {"clip_distyces", Extension::kUndefined},
+ {"clipGlr77inntances", Extension::kUndefined},
+ {"clip_d04stances", Extension::kUndefined},
+ {"dua_ource_bledoong", Extension::kUndefined},
+ {"dualsorce_blendzzng", Extension::kUndefined},
+ {"ua_sopiirce_bl11nding", Extension::kUndefined},
+ {"f1XX", Extension::kUndefined},
+ {"55199II", Extension::kUndefined},
+ {"frSSHHa", Extension::kUndefined},
+ {"kkbHups", Extension::kUndefined},
+ {"jggoupRR", Extension::kUndefined},
+ {"sugoupb", Extension::kUndefined},
+ {"subgroupj_f16", Extension::kUndefined},
+ {"subgroups_f6", Extension::kUndefined},
+ {"sgroupq_f16", Extension::kUndefined},
};
using ExtensionParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/lang/wgsl/helpers/check_supported_extensions_test.cc b/src/tint/lang/wgsl/helpers/check_supported_extensions_test.cc
index 6d8252d..e17683f 100644
--- a/src/tint/lang/wgsl/helpers/check_supported_extensions_test.cc
+++ b/src/tint/lang/wgsl/helpers/check_supported_extensions_test.cc
@@ -42,7 +42,7 @@
ASSERT_TRUE(CheckSupportedExtensions("writer", AST(), Diagnostics(),
Vector{
wgsl::Extension::kF16,
- wgsl::Extension::kChromiumExperimentalSubgroups,
+ wgsl::Extension::kChromiumExperimentalSubgroupMatrix,
}));
}
@@ -51,7 +51,7 @@
ASSERT_FALSE(CheckSupportedExtensions("writer", AST(), Diagnostics(),
Vector{
- wgsl::Extension::kChromiumExperimentalSubgroups,
+ wgsl::Extension::kChromiumExperimentalSubgroupMatrix,
}));
EXPECT_EQ(Diagnostics().Str(), "12:34 error: writer backend does not support extension 'f16'");
}
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 60290b4..9d1ece2 100644
--- a/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc
+++ b/src/tint/lang/wgsl/reader/parser/enable_directive_test.cc
@@ -82,8 +82,7 @@
// Test a valid enable directive with multiple extensions.
TEST_F(EnableDirectiveTest, Multiple) {
- auto p = parser(
- "enable f16, chromium_disable_uniformity_analysis, chromium_experimental_subgroups;");
+ auto p = parser("enable f16, chromium_disable_uniformity_analysis, subgroups;");
p->enable_directive();
EXPECT_FALSE(p->has_error()) << p->error();
auto program = p->program();
@@ -101,19 +100,18 @@
EXPECT_EQ(enable->extensions[1]->source.range.begin.column, 13u);
EXPECT_EQ(enable->extensions[1]->source.range.end.line, 1u);
EXPECT_EQ(enable->extensions[1]->source.range.end.column, 49u);
- EXPECT_EQ(enable->extensions[2]->name, wgsl::Extension::kChromiumExperimentalSubgroups);
+ EXPECT_EQ(enable->extensions[2]->name, wgsl::Extension::kSubgroups);
EXPECT_EQ(enable->extensions[2]->source.range.begin.line, 1u);
EXPECT_EQ(enable->extensions[2]->source.range.begin.column, 51u);
EXPECT_EQ(enable->extensions[2]->source.range.end.line, 1u);
- EXPECT_EQ(enable->extensions[2]->source.range.end.column, 82u);
+ EXPECT_EQ(enable->extensions[2]->source.range.end.column, 60u);
ASSERT_EQ(ast.GlobalDeclarations().Length(), 1u);
EXPECT_EQ(ast.GlobalDeclarations()[0], enable);
}
// Test a valid enable directive with multiple extensions.
TEST_F(EnableDirectiveTest, MultipleTrailingComma) {
- auto p = parser(
- "enable f16, chromium_disable_uniformity_analysis, chromium_experimental_subgroups,;");
+ auto p = parser("enable f16, chromium_disable_uniformity_analysis, subgroups,;");
p->enable_directive();
EXPECT_FALSE(p->has_error()) << p->error();
auto program = p->program();
@@ -131,11 +129,11 @@
EXPECT_EQ(enable->extensions[1]->source.range.begin.column, 13u);
EXPECT_EQ(enable->extensions[1]->source.range.end.line, 1u);
EXPECT_EQ(enable->extensions[1]->source.range.end.column, 49u);
- EXPECT_EQ(enable->extensions[2]->name, wgsl::Extension::kChromiumExperimentalSubgroups);
+ EXPECT_EQ(enable->extensions[2]->name, wgsl::Extension::kSubgroups);
EXPECT_EQ(enable->extensions[2]->source.range.begin.line, 1u);
EXPECT_EQ(enable->extensions[2]->source.range.begin.column, 51u);
EXPECT_EQ(enable->extensions[2]->source.range.end.line, 1u);
- EXPECT_EQ(enable->extensions[2]->source.range.end.column, 82u);
+ EXPECT_EQ(enable->extensions[2]->source.range.end.column, 60u);
ASSERT_EQ(ast.GlobalDeclarations().Length(), 1u);
EXPECT_EQ(ast.GlobalDeclarations()[0], enable);
}
@@ -205,7 +203,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_pixel_local', 'chromium_experimental_push_constant', 'chromium_experimental_subgroup_matrix', 'chromium_experimental_subgroups', 'chromium_internal_graphite', 'chromium_internal_input_attachments', 'chromium_internal_relaxed_uniform_layout', 'clip_distances', 'dual_source_blending', 'f16', 'subgroups', 'subgroups_f16')");
+Possible values: 'chromium_disable_uniformity_analysis', 'chromium_experimental_framebuffer_fetch', 'chromium_experimental_pixel_local', 'chromium_experimental_push_constant', 'chromium_experimental_subgroup_matrix', 'chromium_internal_graphite', 'chromium_internal_input_attachments', 'chromium_internal_relaxed_uniform_layout', 'clip_distances', 'dual_source_blending', 'f16', 'subgroups', 'subgroups_f16')");
auto program = p->program();
auto& ast = program.AST();
EXPECT_EQ(ast.Enables().Length(), 0u);
diff --git a/src/tint/lang/wgsl/resolver/builtin_validation_test.cc b/src/tint/lang/wgsl/resolver/builtin_validation_test.cc
index eaa3afe..7c6e138 100644
--- a/src/tint/lang/wgsl/resolver/builtin_validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/builtin_validation_test.cc
@@ -879,19 +879,6 @@
)");
}
-TEST_F(ResolverBuiltinValidationTest, SubgroupBallotWithExperimentalExtension) {
- // enable chromium_experimental_subgroups;
- // fn func -> vec4<u32> { return subgroupBallot(true); }
- Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
-
- Func("func", tint::Empty, ty.vec4<u32>(),
- Vector{
- Return(Call("subgroupBallot", true)),
- });
-
- EXPECT_TRUE(r()->Resolve());
-}
-
TEST_F(ResolverBuiltinValidationTest, SubgroupBroadcastWithoutExtension) {
// fn func -> i32 { return subgroupBroadcast(1,0); }
Func("func", tint::Empty, ty.i32(),
@@ -918,19 +905,6 @@
EXPECT_TRUE(r()->Resolve());
}
-TEST_F(ResolverBuiltinValidationTest, SubgroupBroadcastWithExperimentalExtension) {
- // enable chromium_experimental_subgroups;
- // fn func -> i32 { return subgroupBroadcast(1,0); }
- Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
-
- Func("func", tint::Empty, ty.i32(),
- Vector{
- Return(Call("subgroupBroadcast", 1_i, 0_u)),
- });
-
- EXPECT_TRUE(r()->Resolve());
-}
-
TEST_F(ResolverBuiltinValidationTest, SubgroupBroadcastWithoutExtension_F16) {
// enable f16;
// enable subgroups;
@@ -965,21 +939,6 @@
EXPECT_TRUE(r()->Resolve());
}
-TEST_F(ResolverBuiltinValidationTest, SubgroupBroadcastWithExperimentalExtension_F16) {
- // enable f16;
- // enable chromium_experimental_subgroups;
- // fn func -> f16 { return subgroupBroadcast(1.h,0); }
- Enable(wgsl::Extension::kF16);
- Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
-
- Func("func", tint::Empty, ty.f16(),
- Vector{
- Return(Call("subgroupBroadcast", 1_h, 0_u)),
- });
-
- EXPECT_TRUE(r()->Resolve());
-}
-
TEST_F(ResolverBuiltinValidationTest, SubgroupBroadcastWithoutExtension_VecF16) {
// enable f16;
// enable subgroups;
@@ -1014,21 +973,6 @@
EXPECT_TRUE(r()->Resolve());
}
-TEST_F(ResolverBuiltinValidationTest, SubgroupBroadcastWithExperimentalExtension_VecF16) {
- // enable f16;
- // enable chromium_experimental_subgroups;
- // fn func -> vec4<f16> { return subgroupBroadcast(vec4(1.h),0); }
- Enable(wgsl::Extension::kF16);
- Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
-
- Func("func", tint::Empty, ty.vec4<f16>(),
- Vector{
- Return(Call("subgroupBroadcast", Call(ty.vec4<f16>(), 1_h), 0_u)),
- });
-
- EXPECT_TRUE(r()->Resolve());
-}
-
TEST_F(ResolverBuiltinValidationTest, SubroupBroadcastInComputeStage) {
// @vertex fn func { dpdx(1.0); }
diff --git a/src/tint/lang/wgsl/resolver/subgroups_extension_test.cc b/src/tint/lang/wgsl/resolver/subgroups_extension_test.cc
index 2b29b52..3fa8e4d 100644
--- a/src/tint/lang/wgsl/resolver/subgroups_extension_test.cc
+++ b/src/tint/lang/wgsl/resolver/subgroups_extension_test.cc
@@ -95,17 +95,6 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
}
-// Using a subgroup_size builtin attribute with chromium_experimental_subgroups enabled should pass.
-TEST_F(ResolverSubgroupsExtensionTest, UseSubgroupSizeAttribWithExperimentalExtension) {
- Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
- Structure("Inputs",
- Vector{
- Member("a", ty.u32(), Vector{Builtin(core::BuiltinValue::kSubgroupSize)}),
- });
-
- EXPECT_TRUE(r()->Resolve()) << r()->error();
-}
-
// Using a subgroup_invocation_id builtin attribute with subgroups enabled should pass.
TEST_F(ResolverSubgroupsExtensionTest, UseSubgroupInvocationIdAttribWithExtension) {
Enable(wgsl::Extension::kSubgroups);
@@ -117,18 +106,6 @@
EXPECT_TRUE(r()->Resolve()) << r()->error();
}
-// Using a subgroup_invocation_id builtin attribute with chromium_experimental_subgroups enabled
-// should pass.
-TEST_F(ResolverSubgroupsExtensionTest, UseSubgroupInvocationIdAttribWithExperimentalExtension) {
- Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
- Structure("Inputs",
- Vector{
- Member("a", ty.u32(), Vector{Builtin(core::BuiltinValue::kSubgroupInvocationId)}),
- });
-
- EXPECT_TRUE(r()->Resolve()) << r()->error();
-}
-
// Using an i32 for a subgroup_size builtin input should fail.
TEST_F(ResolverSubgroupsExtensionTest, SubgroupSizeI32Error) {
Enable(wgsl::Extension::kSubgroups);
diff --git a/src/tint/lang/wgsl/resolver/uniformity_test.cc b/src/tint/lang/wgsl/resolver/uniformity_test.cc
index 576ff69..28b1648 100644
--- a/src/tint/lang/wgsl/resolver/uniformity_test.cc
+++ b/src/tint/lang/wgsl/resolver/uniformity_test.cc
@@ -617,10 +617,9 @@
class ComputeBuiltin : public UniformityAnalysisTestBase,
public ::testing::TestWithParam<BuiltinEntry> {};
TEST_P(ComputeBuiltin, AsParam) {
- std::string src = std::string((GetParam().name == "subgroup_size")
- ? R"(enable chromium_experimental_subgroups;
+ std::string src = std::string((GetParam().name == "subgroup_size") ? R"(enable subgroups;
)"
- : "") +
+ : "") +
R"(
@compute @workgroup_size(64)
fn main(@builtin()" + GetParam().name +
@@ -652,10 +651,9 @@
}
TEST_P(ComputeBuiltin, InStruct) {
- std::string src = std::string((GetParam().name == "subgroup_size")
- ? R"(enable chromium_experimental_subgroups;
+ std::string src = std::string((GetParam().name == "subgroup_size") ? R"(enable subgroups;
)"
- : "") +
+ : "") +
R"(
struct S {
@builtin()" + GetParam().name +
@@ -738,10 +736,9 @@
class FragmentBuiltin : public UniformityAnalysisTestBase,
public ::testing::TestWithParam<BuiltinEntry> {};
TEST_P(FragmentBuiltin, AsParam) {
- std::string src = std::string((GetParam().name == "subgroup_size")
- ? R"(enable chromium_experimental_subgroups;
+ std::string src = std::string((GetParam().name == "subgroup_size") ? R"(enable subgroups;
)"
- : R"(
+ : R"(
)") +
R"(
@fragment
@@ -773,10 +770,9 @@
}
TEST_P(FragmentBuiltin, InStruct) {
- std::string src = std::string((GetParam().name == "subgroup_size")
- ? R"(enable chromium_experimental_subgroups;
+ std::string src = std::string((GetParam().name == "subgroup_size") ? R"(enable subgroups;
)"
- : R"(
+ : R"(
)") +
R"(
struct S {
diff --git a/src/tint/lang/wgsl/resolver/validator.cc b/src/tint/lang/wgsl/resolver/validator.cc
index 9e806a6..21df623 100644
--- a/src/tint/lang/wgsl/resolver/validator.cc
+++ b/src/tint/lang/wgsl/resolver/validator.cc
@@ -1080,8 +1080,7 @@
break;
case core::BuiltinValue::kSubgroupInvocationId:
case core::BuiltinValue::kSubgroupSize:
- if (!(enabled_extensions_.Contains(wgsl::Extension::kChromiumExperimentalSubgroups) ||
- enabled_extensions_.Contains(wgsl::Extension::kSubgroups))) {
+ if (!enabled_extensions_.Contains(wgsl::Extension::kSubgroups)) {
AddError(attr->source)
<< "use of " << style::Attribute("@builtin")
<< style::Code("(", style::Enum(builtin), ")")
@@ -2016,19 +2015,15 @@
}
if (builtin->IsSubgroup()) {
- // The `chromium_experimental_subgroups` extension enables all subgroup features. Otherwise,
- // we need `subgroups`, or `subgroups_f16` for f16 functions.
- if (!enabled_extensions_.Contains(wgsl::Extension::kChromiumExperimentalSubgroups)) {
- auto ext = wgsl::Extension::kSubgroups;
- if (builtin->ReturnType()->DeepestElement()->Is<core::type::F16>()) {
- ext = wgsl::Extension::kSubgroupsF16;
- }
- if (!enabled_extensions_.Contains(ext)) {
- AddError(call->Declaration()->source)
- << "cannot call built-in function " << style::Function(builtin->Fn())
- << " without extension " << style::Code(wgsl::ToString(ext));
- return false;
- }
+ auto ext = wgsl::Extension::kSubgroups;
+ if (builtin->ReturnType()->DeepestElement()->Is<core::type::F16>()) {
+ ext = wgsl::Extension::kSubgroupsF16;
+ }
+ if (!enabled_extensions_.Contains(ext)) {
+ AddError(call->Declaration()->source)
+ << "cannot call built-in function " << style::Function(builtin->Fn())
+ << " without extension " << style::Code(wgsl::ToString(ext));
+ return false;
}
}
diff --git a/src/tint/lang/wgsl/wgsl.def b/src/tint/lang/wgsl/wgsl.def
index 4f29985..57d8f5e 100644
--- a/src/tint/lang/wgsl/wgsl.def
+++ b/src/tint/lang/wgsl/wgsl.def
@@ -84,8 +84,6 @@
chromium_disable_uniformity_analysis
// A Chromium-specific extension for push constants
chromium_experimental_push_constant
- // A Chromium-specific extension that adds basic subgroup functionality.
- chromium_experimental_subgroups
// A Chromium-specific extension that enables features for graphite
chromium_internal_graphite
// A Chromium-specific extension that enables features for input
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 05a4711..8ebaf97 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
@@ -248,11 +248,11 @@
attrs.Push(b.Builtin(core::BuiltinValue::kSampleMask));
break;
case core::BuiltinValue::kSubgroupInvocationId:
- Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
+ Enable(wgsl::Extension::kSubgroups);
attrs.Push(b.Builtin(core::BuiltinValue::kSubgroupInvocationId));
break;
case core::BuiltinValue::kSubgroupSize:
- Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
+ Enable(wgsl::Extension::kSubgroups);
attrs.Push(b.Builtin(core::BuiltinValue::kSubgroupSize));
break;
case core::BuiltinValue::kClipDistances:
@@ -681,7 +681,9 @@
case wgsl::BuiltinFn::kQuadSwapX:
case wgsl::BuiltinFn::kQuadSwapY:
case wgsl::BuiltinFn::kQuadSwapDiagonal:
- Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
+ Enable(wgsl::Extension::kF16);
+ Enable(wgsl::Extension::kSubgroups);
+ Enable(wgsl::Extension::kSubgroupsF16);
break;
default:
break;
@@ -1081,7 +1083,7 @@
}
if (auto builtin = ir_attrs.builtin) {
if (RequiresSubgroups(*builtin)) {
- Enable(wgsl::Extension::kChromiumExperimentalSubgroups);
+ Enable(wgsl::Extension::kSubgroups);
} else if (*builtin == core::BuiltinValue::kClipDistances) {
Enable(wgsl::Extension::kClipDistances);
}
@@ -1258,8 +1260,7 @@
}
}
- /// @returns true if the builtin value requires the kChromiumExperimentalSubgroups extension to
- /// be enabled.
+ /// @returns true if the builtin value requires the kSubgroups extension to be enabled.
bool RequiresSubgroups(core::BuiltinValue builtin) {
switch (builtin) {
case core::BuiltinValue::kSubgroupInvocationId:
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 578b0cb..4d38ae4 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
@@ -291,7 +291,7 @@
fn->Block()->Append(b.Return(fn));
EXPECT_WGSL(R"(
-enable chromium_experimental_subgroups;
+enable subgroups;
@compute @workgroup_size(3u, 4u, 5u)
fn f(@builtin(local_invocation_id) v : vec3<u32>, @builtin(local_invocation_index) v_1 : u32, @builtin(global_invocation_id) v_2 : vec3<u32>, @builtin(workgroup_id) v_3 : vec3<u32>, @builtin(num_workgroups) v_4 : vec3<u32>, @builtin(subgroup_invocation_id) v_5 : u32, @builtin(subgroup_size) v_6 : u32) {
@@ -311,7 +311,7 @@
fn->Block()->Append(b.Return(fn));
EXPECT_WGSL(R"(
-enable chromium_experimental_subgroups;
+enable subgroups;
@fragment
fn f(@builtin(front_facing) v : bool, @builtin(sample_index) v_1 : u32, @builtin(sample_mask) v_2 : u32, @builtin(subgroup_size) v_3 : u32) {
@@ -2662,96 +2662,6 @@
)");
}
-////////////////////////////////////////////////////////////////////////////////
-// chromium_experimental_subgroups
-////////////////////////////////////////////////////////////////////////////////
-TEST_F(IRToProgramTest, Enable_ChromiumExperimentalSubgroups_SubgroupBallot) {
- auto* fn = b.Function("f", ty.void_());
- b.Append(fn->Block(), [&] {
- auto* call = b.CallWithResult<wgsl::ir::BuiltinCall>(
- b.InstructionResult(ty.vec4<u32>()), wgsl::BuiltinFn::kSubgroupBallot, true);
- b.Let("v", call);
- b.Return(fn);
- });
-
- EXPECT_WGSL(R"(
-enable chromium_experimental_subgroups;
-
-fn f() {
- let v = subgroupBallot(true);
-}
-)");
-}
-
-TEST_F(IRToProgramTest, Enable_ChromiumExperimentalSubgroups_SubgroupBroadcast) {
- auto* fn = b.Function("f", ty.void_());
- b.Append(fn->Block(), [&] {
- auto* one = b.Value(1_u);
- auto* call = b.CallWithResult<wgsl::ir::BuiltinCall>(
- b.InstructionResult(ty.u32()), wgsl::BuiltinFn::kSubgroupBroadcast, Vector{one, one});
- b.Let("v", call);
- b.Return(fn);
- });
-
- EXPECT_WGSL(R"(
-enable chromium_experimental_subgroups;
-
-fn f() {
- let v = subgroupBroadcast(1u, 1u);
-}
-)");
-}
-
-TEST_F(IRToProgramTest, Enable_ChromiumExperimentalSubgroups_StructBuiltin_SubgroupInvocationId) {
- core::type::Manager::StructMemberDesc member;
- member.name = mod.symbols.New("a");
- member.type = ty.u32();
- member.attributes.builtin = core::BuiltinValue::kSubgroupInvocationId;
-
- auto* S = ty.Struct(mod.symbols.New("S"), {member});
-
- auto* fn = b.Function("f", ty.void_());
- fn->SetParams({b.FunctionParam(S)});
- b.Append(fn->Block(), [&] { b.Return(fn); });
-
- EXPECT_WGSL(R"(
-enable chromium_experimental_subgroups;
-
-struct S {
- @builtin(subgroup_invocation_id)
- a : u32,
-}
-
-fn f(v : S) {
-}
-)");
-}
-
-TEST_F(IRToProgramTest, Enable_ChromiumExperimentalSubgroups_StructBuiltin_SubgroupSize) {
- core::type::Manager::StructMemberDesc member;
- member.name = mod.symbols.New("a");
- member.type = ty.u32();
- member.attributes.builtin = core::BuiltinValue::kSubgroupSize;
-
- auto* S = ty.Struct(mod.symbols.New("S"), {member});
-
- auto* fn = b.Function("f", ty.void_());
- fn->SetParams({b.FunctionParam(S)});
- b.Append(fn->Block(), [&] { b.Return(fn); });
-
- EXPECT_WGSL(R"(
-enable chromium_experimental_subgroups;
-
-struct S {
- @builtin(subgroup_size)
- a : u32,
-}
-
-fn f(v : S) {
-}
-)");
-}
-
TEST_F(IRToProgramTest, Enable_ChromiumExperimentalFramebufferFetch_StructColor) {
core::type::Manager::StructMemberDesc member;
member.name = mod.symbols.New("a");
diff --git a/test/tint/builtins/compute_subgroup_inclusive.wgsl b/test/tint/builtins/compute_subgroup_inclusive.wgsl
index ffdfa1d..27c317f 100644
--- a/test/tint/builtins/compute_subgroup_inclusive.wgsl
+++ b/test/tint/builtins/compute_subgroup_inclusive.wgsl
@@ -1,4 +1,4 @@
-enable chromium_experimental_subgroups;
+enable subgroups;
@compute @workgroup_size(1)
fn main() {
diff --git a/test/tint/builtins/compute_subgroup_inclusive.wgsl.expected.wgsl b/test/tint/builtins/compute_subgroup_inclusive.wgsl.expected.wgsl
index f5bbe9d..17158b4 100644
--- a/test/tint/builtins/compute_subgroup_inclusive.wgsl.expected.wgsl
+++ b/test/tint/builtins/compute_subgroup_inclusive.wgsl.expected.wgsl
@@ -1,4 +1,4 @@
-enable chromium_experimental_subgroups;
+enable subgroups;
@compute @workgroup_size(1)
fn main() {
diff --git a/test/tint/types/functions/shader_io/compute_subgroup_builtins.wgsl b/test/tint/types/functions/shader_io/compute_subgroup_builtins.wgsl
index f821795..0766830 100644
--- a/test/tint/types/functions/shader_io/compute_subgroup_builtins.wgsl
+++ b/test/tint/types/functions/shader_io/compute_subgroup_builtins.wgsl
@@ -1,4 +1,4 @@
-enable chromium_experimental_subgroups;
+enable subgroups;
@group(0) @binding(0)
var<storage, read_write> output: array<u32>;
diff --git a/test/tint/types/functions/shader_io/compute_subgroup_builtins.wgsl.expected.wgsl b/test/tint/types/functions/shader_io/compute_subgroup_builtins.wgsl.expected.wgsl
index f237c19..b417e25 100644
--- a/test/tint/types/functions/shader_io/compute_subgroup_builtins.wgsl.expected.wgsl
+++ b/test/tint/types/functions/shader_io/compute_subgroup_builtins.wgsl.expected.wgsl
@@ -1,4 +1,4 @@
-enable chromium_experimental_subgroups;
+enable subgroups;
@group(0) @binding(0) var<storage, read_write> output : array<u32>;
diff --git a/test/tint/types/functions/shader_io/compute_subgroup_builtins_struct.wgsl b/test/tint/types/functions/shader_io/compute_subgroup_builtins_struct.wgsl
index 6f425ba..9d43000 100644
--- a/test/tint/types/functions/shader_io/compute_subgroup_builtins_struct.wgsl
+++ b/test/tint/types/functions/shader_io/compute_subgroup_builtins_struct.wgsl
@@ -1,4 +1,4 @@
-enable chromium_experimental_subgroups;
+enable subgroups;
@group(0) @binding(0)
var<storage, read_write> output: array<u32>;
diff --git a/test/tint/types/functions/shader_io/compute_subgroup_builtins_struct.wgsl.expected.wgsl b/test/tint/types/functions/shader_io/compute_subgroup_builtins_struct.wgsl.expected.wgsl
index dacd4b7..e8f7f06 100644
--- a/test/tint/types/functions/shader_io/compute_subgroup_builtins_struct.wgsl.expected.wgsl
+++ b/test/tint/types/functions/shader_io/compute_subgroup_builtins_struct.wgsl.expected.wgsl
@@ -1,4 +1,4 @@
-enable chromium_experimental_subgroups;
+enable subgroups;
@group(0) @binding(0) var<storage, read_write> output : array<u32>;
diff --git a/test/tint/types/functions/shader_io/fragment_subgroup_builtins.wgsl b/test/tint/types/functions/shader_io/fragment_subgroup_builtins.wgsl
index 372412a..abe6eb9 100644
--- a/test/tint/types/functions/shader_io/fragment_subgroup_builtins.wgsl
+++ b/test/tint/types/functions/shader_io/fragment_subgroup_builtins.wgsl
@@ -1,4 +1,4 @@
-enable chromium_experimental_subgroups;
+enable subgroups;
@group(0) @binding(0)
var<storage, read_write> output: array<u32>;
diff --git a/test/tint/types/functions/shader_io/fragment_subgroup_builtins.wgsl.expected.wgsl b/test/tint/types/functions/shader_io/fragment_subgroup_builtins.wgsl.expected.wgsl
index 0a06e13..11d1b4b 100644
--- a/test/tint/types/functions/shader_io/fragment_subgroup_builtins.wgsl.expected.wgsl
+++ b/test/tint/types/functions/shader_io/fragment_subgroup_builtins.wgsl.expected.wgsl
@@ -1,4 +1,4 @@
-enable chromium_experimental_subgroups;
+enable subgroups;
@group(0) @binding(0) var<storage, read_write> output : array<u32>;
diff --git a/test/tint/types/functions/shader_io/fragment_subgroup_builtins_struct.wgsl b/test/tint/types/functions/shader_io/fragment_subgroup_builtins_struct.wgsl
index 4b10630..b8f9efe 100644
--- a/test/tint/types/functions/shader_io/fragment_subgroup_builtins_struct.wgsl
+++ b/test/tint/types/functions/shader_io/fragment_subgroup_builtins_struct.wgsl
@@ -1,4 +1,4 @@
-enable chromium_experimental_subgroups;
+enable subgroups;
@group(0) @binding(0)
var<storage, read_write> output: array<u32>;
diff --git a/test/tint/types/functions/shader_io/fragment_subgroup_builtins_struct.wgsl.expected.wgsl b/test/tint/types/functions/shader_io/fragment_subgroup_builtins_struct.wgsl.expected.wgsl
index b7e71f3..6ab0528 100644
--- a/test/tint/types/functions/shader_io/fragment_subgroup_builtins_struct.wgsl.expected.wgsl
+++ b/test/tint/types/functions/shader_io/fragment_subgroup_builtins_struct.wgsl.expected.wgsl
@@ -1,4 +1,4 @@
-enable chromium_experimental_subgroups;
+enable subgroups;
@group(0) @binding(0) var<storage, read_write> output : array<u32>;
