dawn/native: Add validation on the range of `@subgroup_size` This patch adds validations on the attribute `@subgroup_size` that this value must be between `minExplicitComputeSubgroupSize` and `maxExplicitComputeSubgroupSize`. Bug: 463721943 Change-Id: I445be4eaaaf30d7627da66b8028d16810391f9de Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/285575 Reviewed-by: Corentin Wallez <cwallez@chromium.org> Reviewed-by: Loko Kung <lokokung@google.com> Commit-Queue: Shao, Jiawei <jiawei.shao@intel.com>
diff --git a/src/dawn/native/PhysicalDevice.cpp b/src/dawn/native/PhysicalDevice.cpp index fa828fb..4866c92 100644 --- a/src/dawn/native/PhysicalDevice.cpp +++ b/src/dawn/native/PhysicalDevice.cpp
@@ -221,4 +221,12 @@ wgpu::TextureFormat format, UnpackedPtr<DawnFormatCapabilities>& capabilities) const {} +uint32_t PhysicalDeviceBase::GetMinExplicitComputeSubgroupSize() const { + return mMinExplicitComputeSubgroupSize; +} + +uint32_t PhysicalDeviceBase::GetMaxExplicitComputeSubgroupSize() const { + return mMaxExplicitComputeSubgroupSize; +} + } // namespace dawn::native
diff --git a/src/dawn/native/PhysicalDevice.h b/src/dawn/native/PhysicalDevice.h index 7343a45..898169c 100644 --- a/src/dawn/native/PhysicalDevice.h +++ b/src/dawn/native/PhysicalDevice.h
@@ -133,6 +133,9 @@ InstanceBase* instance, const Surface* surface) const = 0; + uint32_t GetMinExplicitComputeSubgroupSize() const; + uint32_t GetMaxExplicitComputeSubgroupSize() const; + protected: uint32_t mVendorId = 0xFFFFFFFF; std::string mVendorName; @@ -146,6 +149,8 @@ // backend may override this. uint32_t mSubgroupMinSize = kDefaultSubgroupMinSize; uint32_t mSubgroupMaxSize = kDefaultSubgroupMaxSize; + uint32_t mMinExplicitComputeSubgroupSize = kDefaultSubgroupMinSize; + uint32_t mMaxExplicitComputeSubgroupSize = kDefaultSubgroupMaxSize; // Juat a wrapper of ValidateFeatureSupportedWithToggles, return true if a feature is supported // by this adapter AND suitable with given toggles.
diff --git a/src/dawn/native/ShaderModule.cpp b/src/dawn/native/ShaderModule.cpp index 384cdf7..c8a8a95 100644 --- a/src/dawn/native/ShaderModule.cpp +++ b/src/dawn/native/ShaderModule.cpp
@@ -1317,7 +1317,7 @@ bool usesSubgroupMatrix, uint32_t maxSubgroupSize, const LimitsForCompilationRequest& limits, - const LimitsForCompilationRequest& adaterSupportedlimits) { + const LimitsForCompilationRequest& adapterSupportedlimits) { DAWN_INVALID_IF(workgroupInfo.x < 1 || workgroupInfo.y < 1 || workgroupInfo.z < 1, "Entry-point uses workgroup_size(%u, %u, %u) that are below the " "minimum allowed (1, 1, 1).", @@ -1327,11 +1327,11 @@ workgroupInfo.y > limits.maxComputeWorkgroupSizeY || workgroupInfo.z > limits.maxComputeWorkgroupSizeZ) [[unlikely]] { uint32_t maxComputeWorkgroupSizeXAdapterLimit = - adaterSupportedlimits.maxComputeWorkgroupSizeX; + adapterSupportedlimits.maxComputeWorkgroupSizeX; uint32_t maxComputeWorkgroupSizeYAdapterLimit = - adaterSupportedlimits.maxComputeWorkgroupSizeY; + adapterSupportedlimits.maxComputeWorkgroupSizeY; uint32_t maxComputeWorkgroupSizeZAdapterLimit = - adaterSupportedlimits.maxComputeWorkgroupSizeZ; + adapterSupportedlimits.maxComputeWorkgroupSizeZ; std::string increaseLimitAdvice = (workgroupInfo.x <= maxComputeWorkgroupSizeXAdapterLimit && workgroupInfo.y <= maxComputeWorkgroupSizeYAdapterLimit && @@ -1359,7 +1359,7 @@ "The total number of workgroup invocations (%u) exceeds the " "maximum allowed (%u).%s", numInvocations, maxComputeInvocationsPerWorkgroup, - DAWN_INCREASE_LIMIT_MESSAGE(adaterSupportedlimits, + DAWN_INCREASE_LIMIT_MESSAGE(adapterSupportedlimits, maxComputeInvocationsPerWorkgroup, numInvocations)); uint32_t maxComputeWorkgroupStorageSize = limits.maxComputeWorkgroupStorageSize; @@ -1368,7 +1368,7 @@ "The total use of workgroup storage (%u bytes) is larger than " "the maximum allowed (%u bytes).%s", workgroupInfo.storage_size, maxComputeWorkgroupStorageSize, - DAWN_INCREASE_LIMIT_MESSAGE(adaterSupportedlimits, maxComputeWorkgroupStorageSize, + DAWN_INCREASE_LIMIT_MESSAGE(adapterSupportedlimits, maxComputeWorkgroupStorageSize, workgroupInfo.storage_size)); if (usesSubgroupMatrix) { @@ -1393,6 +1393,23 @@ return Extent3D{workgroupInfo.x, workgroupInfo.y, workgroupInfo.z}; } +MaybeError ValidateExplicitComputeSubgroupSize(const tint::WorkgroupInfo& workgroupInfo, + uint32_t minExplicitSubgroupSize, + uint32_t maxExplicitSubgroupSize) { + if (workgroupInfo.subgroup_size.has_value()) { + DAWN_ASSERT(minExplicitSubgroupSize > 0 && maxExplicitSubgroupSize > 0); + const uint32_t explicitSubgroupSize = workgroupInfo.subgroup_size.value(); + DAWN_INVALID_IF( + explicitSubgroupSize < minExplicitSubgroupSize || + explicitSubgroupSize > maxExplicitSubgroupSize, + "The subgroup_size attribute (%u) is not in the allowed range " + "[minExplicitComputeSubgroupSize, maxExplicitComputeSubgroupSize] ([%u, %u]).", + explicitSubgroupSize, minExplicitSubgroupSize, maxExplicitSubgroupSize); + } + + return {}; +} + CachedValidationError::CachedValidationError(std::unique_ptr<ErrorData>&& errorData) { DAWN_ASSERT(errorData->GetType() == InternalErrorType::Validation); message = errorData->GetMessage();
diff --git a/src/dawn/native/ShaderModule.h b/src/dawn/native/ShaderModule.h index 606f23a..37b8e2b 100644 --- a/src/dawn/native/ShaderModule.h +++ b/src/dawn/native/ShaderModule.h
@@ -177,6 +177,10 @@ const LimitsForCompilationRequest& limits, const LimitsForCompilationRequest& adaterSupportedlimits); +MaybeError ValidateExplicitComputeSubgroupSize(const tint::WorkgroupInfo& workgroupInfo, + uint32_t minExplicitSubgroupSize, + uint32_t maxExplicitSubgroupSize); + MaybeError ValidateSubgroupMatrixConfiguration(const tint::SubgroupMatrixInfo& smInfo, const std::vector<SubgroupMatrixConfig>& cfg);
diff --git a/src/dawn/native/d3d/D3DCompilationRequest.h b/src/dawn/native/d3d/D3DCompilationRequest.h index e12b8c8..b4b0a9d 100644 --- a/src/dawn/native/d3d/D3DCompilationRequest.h +++ b/src/dawn/native/d3d/D3DCompilationRequest.h
@@ -74,6 +74,8 @@ X(LimitsForCompilationRequest, limits) \ X(UnsafeUnserializedValue<LimitsForCompilationRequest>, adapterSupportedLimits) \ X(uint32_t, maxSubgroupSize) \ + X(uint32_t, minExplicitComputeSubgroupSize) \ + X(uint32_t, maxExplicitComputeSubgroupSize) \ X(bool, disableSymbolRenaming) \ X(bool, dumpShaders) \ X(bool, dumpShadersOnFailure)
diff --git a/src/dawn/native/d3d/ShaderUtils.cpp b/src/dawn/native/d3d/ShaderUtils.cpp index 7286d04..35ef902 100644 --- a/src/dawn/native/d3d/ShaderUtils.cpp +++ b/src/dawn/native/d3d/ShaderUtils.cpp
@@ -249,6 +249,9 @@ ValidateComputeStageWorkgroupSize( result->workgroup_info, /*usesSubgroupMatrix=*/false, r.maxSubgroupSize, r.limits, r.adapterSupportedLimits.UnsafeGetValue())); + DAWN_TRY(ValidateExplicitComputeSubgroupSize(result->workgroup_info, + r.minExplicitComputeSubgroupSize, + r.maxExplicitComputeSubgroupSize)); } bool usesVertexIndex = false;
diff --git a/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp b/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp index 5c20f50..d07be08 100644 --- a/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp +++ b/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp
@@ -125,6 +125,9 @@ // https://github.com/Microsoft/DirectXShaderCompiler/wiki/Wave-Intrinsics#:~:text=UINT%20WaveLaneCountMax mSubgroupMaxSize = 128u; + mMinExplicitComputeSubgroupSize = mDeviceInfo.waveLaneCountMin; + mMaxExplicitComputeSubgroupSize = mDeviceInfo.waveLaneCountMax; + return {}; } @@ -967,9 +970,9 @@ if (auto* explicitComputeSubgroupSizeConfigs = info.Get<AdapterPropertiesExplicitComputeSubgroupSizeConfigs>()) { explicitComputeSubgroupSizeConfigs->minExplicitComputeSubgroupSize = - mDeviceInfo.waveLaneCountMin; + GetMinExplicitComputeSubgroupSize(); explicitComputeSubgroupSizeConfigs->maxExplicitComputeSubgroupSize = - mDeviceInfo.waveLaneCountMax; + GetMaxExplicitComputeSubgroupSize(); } }
diff --git a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp index d8b3c5b..6d52981 100644 --- a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp +++ b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
@@ -307,6 +307,13 @@ LimitsForCompilationRequest::Create(device->GetAdapter()->GetLimits().v1)); req.hlsl.maxSubgroupSize = device->GetAdapter()->GetPhysicalDevice()->GetSubgroupMaxSize(); + if (device->HasFeature(Feature::ChromiumExperimentalSubgroupSizeControl)) { + req.hlsl.minExplicitComputeSubgroupSize = + device->GetAdapter()->GetPhysicalDevice()->GetMinExplicitComputeSubgroupSize(); + req.hlsl.maxExplicitComputeSubgroupSize = + device->GetAdapter()->GetPhysicalDevice()->GetMaxExplicitComputeSubgroupSize(); + } + CacheResult<d3d::CompiledShader> compiledShader; DAWN_TRY_LOAD_OR_RUN(compiledShader, device, std::move(req), d3d::CompiledShader::FromValidatedBlob, d3d::CompileShader,
diff --git a/src/dawn/native/null/DeviceNull.cpp b/src/dawn/native/null/DeviceNull.cpp index d4c7be0..db3a3da 100644 --- a/src/dawn/native/null/DeviceNull.cpp +++ b/src/dawn/native/null/DeviceNull.cpp
@@ -130,8 +130,10 @@ } if (auto* explicitComputeSubgroupSizeConfigs = info.Get<AdapterPropertiesExplicitComputeSubgroupSizeConfigs>()) { - explicitComputeSubgroupSizeConfigs->minExplicitComputeSubgroupSize = 4; - explicitComputeSubgroupSizeConfigs->maxExplicitComputeSubgroupSize = 128; + explicitComputeSubgroupSizeConfigs->minExplicitComputeSubgroupSize = + GetMinExplicitComputeSubgroupSize(); + explicitComputeSubgroupSizeConfigs->maxExplicitComputeSubgroupSize = + GetMaxExplicitComputeSubgroupSize(); } } @@ -523,6 +525,11 @@ tintResult->workgroup_info, computeStage.metadata->usesSubgroupMatrix, maxSubgroupSize, limits, adapterSupportedLimits)); + DAWN_TRY(ValidateExplicitComputeSubgroupSize( + tintResult->workgroup_info, + GetDevice()->GetAdapter()->GetPhysicalDevice()->GetMinExplicitComputeSubgroupSize(), + GetDevice()->GetAdapter()->GetPhysicalDevice()->GetMaxExplicitComputeSubgroupSize())); + return {}; }
diff --git a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp index 3f8d37c..c9fa553 100644 --- a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp +++ b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
@@ -208,6 +208,9 @@ mSubgroupMinSize = mDeviceInfo.subgroupSizeControlProperties.minSubgroupSize; mSubgroupMaxSize = mDeviceInfo.subgroupSizeControlProperties.maxSubgroupSize; + mMinExplicitComputeSubgroupSize = mDeviceInfo.subgroupSizeControlProperties.minSubgroupSize; + mMaxExplicitComputeSubgroupSize = mDeviceInfo.subgroupSizeControlProperties.maxSubgroupSize; + // Check for essential Vulkan extensions and features // Dawn requires at least Vulkan 1.1 @@ -1531,9 +1534,9 @@ if (auto* explicitComputeSubgroupSizeConfigs = info.Get<AdapterPropertiesExplicitComputeSubgroupSizeConfigs>()) { explicitComputeSubgroupSizeConfigs->minExplicitComputeSubgroupSize = - mDeviceInfo.subgroupSizeControlProperties.minSubgroupSize; + GetMinExplicitComputeSubgroupSize(); explicitComputeSubgroupSizeConfigs->maxExplicitComputeSubgroupSize = - mDeviceInfo.subgroupSizeControlProperties.maxSubgroupSize; + GetMaxExplicitComputeSubgroupSize(); } }
diff --git a/src/dawn/native/vulkan/PhysicalDeviceVk.h b/src/dawn/native/vulkan/PhysicalDeviceVk.h index e6880f8..582b949 100644 --- a/src/dawn/native/vulkan/PhysicalDeviceVk.h +++ b/src/dawn/native/vulkan/PhysicalDeviceVk.h
@@ -25,8 +25,8 @@ // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. -#ifndef SRC_DAWN_NATIVE_VULKAN_ADAPTERVK_H_ -#define SRC_DAWN_NATIVE_VULKAN_ADAPTERVK_H_ +#ifndef SRC_DAWN_NATIVE_VULKAN_PHYSICALDEVICEVK_H_ +#define SRC_DAWN_NATIVE_VULKAN_PHYSICALDEVICEVK_H_ #include <memory> #include <vector> @@ -131,4 +131,4 @@ } // namespace dawn::native::vulkan -#endif // SRC_DAWN_NATIVE_VULKAN_ADAPTERVK_H_ +#endif // SRC_DAWN_NATIVE_VULKAN_PHYSICALDEVICEVK_H_
diff --git a/src/dawn/native/vulkan/ShaderModuleVk.cpp b/src/dawn/native/vulkan/ShaderModuleVk.cpp index 02224df..064875e 100644 --- a/src/dawn/native/vulkan/ShaderModuleVk.cpp +++ b/src/dawn/native/vulkan/ShaderModuleVk.cpp
@@ -112,6 +112,8 @@ X(LimitsForCompilationRequest, limits) \ X(UnsafeUnserializedValue<LimitsForCompilationRequest>, adapterSupportedLimits) \ X(uint32_t, maxSubgroupSize) \ + X(uint32_t, minExplicitComputeSubgroupSize) \ + X(uint32_t, maxExplicitComputeSubgroupSize) \ X(bool, usesSubgroupMatrix) \ X(std::vector<SubgroupMatrixConfig>, subgroupMatrixConfig) \ X(tint::spirv::writer::Options, tintOptions) \ @@ -279,6 +281,12 @@ req.adapterSupportedLimits = UnsafeUnserializedValue( LimitsForCompilationRequest::Create(GetDevice()->GetAdapter()->GetLimits().v1)); req.maxSubgroupSize = GetDevice()->GetAdapter()->GetPhysicalDevice()->GetSubgroupMaxSize(); + if (GetDevice()->HasFeature(Feature::ChromiumExperimentalSubgroupSizeControl)) { + req.minExplicitComputeSubgroupSize = + GetDevice()->GetAdapter()->GetPhysicalDevice()->GetMinExplicitComputeSubgroupSize(); + req.maxExplicitComputeSubgroupSize = + GetDevice()->GetAdapter()->GetPhysicalDevice()->GetMaxExplicitComputeSubgroupSize(); + } CacheResult<CompiledSpirv> compilation; DAWN_TRY_LOAD_OR_RUN( @@ -320,6 +328,9 @@ _, ValidateComputeStageWorkgroupSize( tintResult->workgroup_info, r.usesSubgroupMatrix, r.maxSubgroupSize, r.limits, r.adapterSupportedLimits.UnsafeGetValue())); + DAWN_TRY(ValidateExplicitComputeSubgroupSize(tintResult->workgroup_info, + r.minExplicitComputeSubgroupSize, + r.maxExplicitComputeSubgroupSize)); } DAWN_TRY(ValidateSubgroupMatrixConfiguration(tintResult->subgroup_matrix_info,
diff --git a/src/dawn/tests/end2end/SubgroupsTests.cpp b/src/dawn/tests/end2end/SubgroupsTests.cpp index 019c66d..2f55cd0 100644 --- a/src/dawn/tests/end2end/SubgroupsTests.cpp +++ b/src/dawn/tests/end2end/SubgroupsTests.cpp
@@ -949,30 +949,8 @@ void DoTest(uint32_t subgroupSize) { DAWN_ASSERT(IsPowerOfTwo(subgroupSize)); - std::stringstream code; - code << R"( -enable subgroups; -enable chromium_experimental_subgroup_size_control; - -override kSubgroupSize : u32; - -@group(0) @binding(0) -var<storage, read_write> output: u32; - -@compute @workgroup_size(kSubgroupSize) @subgroup_size(kSubgroupSize) -fn main(@builtin(subgroup_size) sg_size : u32) { - if (subgroupElect()) { - output = sg_size; - } -})"; - wgpu::ShaderModule csModule = utils::CreateShaderModule(device, code.str().c_str()); - - wgpu::ConstantEntry entry = {nullptr, "kSubgroupSize", static_cast<double>(subgroupSize)}; - wgpu::ComputePipelineDescriptor csDesc; - csDesc.compute.module = csModule; - csDesc.compute.constantCount = 1; - csDesc.compute.constants = &entry; - auto pipeline = device.CreateComputePipeline(&csDesc); + wgpu::ComputePipeline pipeline = + CreateComputePipelineWithSubgroupSizeAttribute(subgroupSize, true); uint32_t outputBufferSizeInBytes = sizeof(uint32_t); wgpu::BufferDescriptor outputBufferDesc; @@ -997,6 +975,46 @@ EXPECT_BUFFER_U32_EQ(subgroupSize, outputBuffer, 0); } + wgpu::ComputePipeline CreateComputePipelineWithSubgroupSizeAttribute( + uint32_t subgroupSize, + bool setSubgroupSizeAsOverride) { + std::stringstream code; + + code << R"( +enable subgroups; +enable chromium_experimental_subgroup_size_control;)"; + + if (setSubgroupSizeAsOverride) { + code << "override kSubgroupSize : u32;\n"; + } else { + code << "const kSubgroupSize = " << subgroupSize << ";\n"; + } + + code << R"( +@group(0) @binding(0) +var<storage, read_write> output: u32; + +@compute @workgroup_size(kSubgroupSize) @subgroup_size(kSubgroupSize) +fn main(@builtin(subgroup_size) sg_size : u32) { + if (subgroupElect()) { + output = sg_size; + } +} +)"; + + wgpu::ComputePipelineDescriptor csDesc; + csDesc.compute.module = utils::CreateShaderModule(device, code.str().c_str()); + + wgpu::ConstantEntry entry; + if (setSubgroupSizeAsOverride) { + entry = {nullptr, "kSubgroupSize", static_cast<double>(subgroupSize)}; + csDesc.compute.constantCount = 1; + csDesc.compute.constants = &entry; + } + + return device.CreateComputePipeline(&csDesc); + } + private: bool mSupportsSubgroupSizeControl = false; }; @@ -1020,6 +1038,48 @@ } } +// Test an error occurs when a value that is less than `minExplicitComputeSubgroupSize` is used as +// the attribute `@subgroup_size`. +TEST_P(SubgroupSizeControlTests, LessThanMinExplicitComputeSubgroupSize) { + DAWN_TEST_UNSUPPORTED_IF(!SupportSubgroupSizeControl()); + + wgpu::AdapterInfo info; + wgpu::AdapterPropertiesExplicitComputeSubgroupSizeConfigs subgroupSizeConfigs; + info.nextInChain = &subgroupSizeConfigs; + adapter.GetInfo(&info); + + ASSERT_TRUE(IsPowerOfTwo(subgroupSizeConfigs.minExplicitComputeSubgroupSize)); + + uint32_t invalidSubgroupSize = subgroupSizeConfigs.minExplicitComputeSubgroupSize / 2; + ASSERT_TRUE(invalidSubgroupSize > 0); + + for (bool setSubgroupSizeAsOverride : {true, false}) { + ASSERT_DEVICE_ERROR(CreateComputePipelineWithSubgroupSizeAttribute( + invalidSubgroupSize, setSubgroupSizeAsOverride)); + } +} + +// Test an error occurs when a value that is more than `maxExplicitComputeSubgroupSize` is used as +// the attribute `@subgroup_size`. +TEST_P(SubgroupSizeControlTests, MoreThanMaxExplicitComputeSubgroupSize) { + DAWN_TEST_UNSUPPORTED_IF(!SupportSubgroupSizeControl()); + + wgpu::AdapterInfo info; + wgpu::AdapterPropertiesExplicitComputeSubgroupSizeConfigs subgroupSizeConfigs; + info.nextInChain = &subgroupSizeConfigs; + adapter.GetInfo(&info); + + ASSERT_TRUE(IsPowerOfTwo(subgroupSizeConfigs.maxExplicitComputeSubgroupSize)); + + uint32_t invalidSubgroupSize = subgroupSizeConfigs.maxExplicitComputeSubgroupSize * 2; + ASSERT_TRUE(invalidSubgroupSize > 0); + + for (bool setSubgroupSizeAsOverride : {true, false}) { + ASSERT_DEVICE_ERROR(CreateComputePipelineWithSubgroupSizeAttribute( + invalidSubgroupSize, setSubgroupSizeAsOverride)); + } +} + DAWN_INSTANTIATE_TEST(SubgroupSizeControlTests, D3D12Backend(), MetalBackend(), VulkanBackend()); } // anonymous namespace
diff --git a/src/dawn/tests/unittests/validation/ShaderModuleValidationTests.cpp b/src/dawn/tests/unittests/validation/ShaderModuleValidationTests.cpp index f4bb3a2..1c8e64d 100644 --- a/src/dawn/tests/unittests/validation/ShaderModuleValidationTests.cpp +++ b/src/dawn/tests/unittests/validation/ShaderModuleValidationTests.cpp
@@ -1213,5 +1213,24 @@ TestTotalInvocationsPerWorkgroupAndSubgroupSize({8, 3, 2}, 32, false); } +// Test it is a validation error to use a `@subgroup_size` that is greater than +// `maxExplicitComputeSubgroupSize` or less than `minExplicitComputeSubgroupSize` on current +// adapter. +TEST_F(SubgroupSizeControlValidationTest, ValidateExplicitComputeSubgroupSizes) { + wgpu::AdapterInfo info; + wgpu::AdapterPropertiesExplicitComputeSubgroupSizeConfigs subgroupSizeConfigs; + info.nextInChain = &subgroupSizeConfigs; + adapter.GetInfo(&info); + + for (uint32_t subgroupSize = subgroupSizeConfigs.minExplicitComputeSubgroupSize / 2; + subgroupSize <= subgroupSizeConfigs.maxExplicitComputeSubgroupSize * 2; + subgroupSize *= 2) { + ASSERT_TRUE(IsPowerOfTwo(subgroupSize)); + bool success = subgroupSize >= subgroupSizeConfigs.minExplicitComputeSubgroupSize && + subgroupSize <= subgroupSizeConfigs.maxExplicitComputeSubgroupSize; + TestTotalInvocationsPerWorkgroupAndSubgroupSize({subgroupSize}, subgroupSize, success); + } +} + } // anonymous namespace } // namespace dawn