Expand error message in ValidateComputeStageWorkgroupSize
This CLs expands error message when validating compute stage workgroup
size to indicate when limits are hit that were not provided in the
device requiredLimits.
Bug: 42240683
Change-Id: I446bb8b89a03b631f30a18b4fb35d7f64e7fdb0d
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/217554
Reviewed-by: Loko Kung <lokokung@google.com>
Commit-Queue: Fr <beaufort.francois@gmail.com>
diff --git a/src/dawn/native/ShaderModule.cpp b/src/dawn/native/ShaderModule.cpp
index 685646e..97b3129 100644
--- a/src/dawn/native/ShaderModule.cpp
+++ b/src/dawn/native/ShaderModule.cpp
@@ -1113,10 +1113,10 @@
}
} // anonymous namespace
-ResultOrError<Extent3D> ValidateComputeStageWorkgroupSize(
- const tint::Program& program,
- const char* entryPointName,
- const LimitsForCompilationRequest& limits) {
+ResultOrError<Extent3D> ValidateComputeStageWorkgroupSize(const tint::Program& program,
+ const char* entryPointName,
+ const LimitsForCompilationRequest& limits,
+ const AdapterBase* adapter) {
tint::inspector::Inspector inspector(program);
// At this point the entry point must exist and must have workgroup size values.
@@ -1125,37 +1125,67 @@
const tint::inspector::WorkgroupSize& workgroup_size = entryPoint.workgroup_size.value();
return ValidateComputeStageWorkgroupSize(workgroup_size.x, workgroup_size.y, workgroup_size.z,
- entryPoint.workgroup_storage_size, limits);
+ entryPoint.workgroup_storage_size, limits, adapter);
}
-ResultOrError<Extent3D> ValidateComputeStageWorkgroupSize(
- uint32_t x,
- uint32_t y,
- uint32_t z,
- size_t workgroupStorageSize,
- const LimitsForCompilationRequest& limits) {
+ResultOrError<Extent3D> ValidateComputeStageWorkgroupSize(uint32_t x,
+ uint32_t y,
+ uint32_t z,
+ size_t workgroupStorageSize,
+ const LimitsForCompilationRequest& limits,
+ const AdapterBase* adapter) {
DAWN_INVALID_IF(x < 1 || y < 1 || z < 1,
"Entry-point uses workgroup_size(%u, %u, %u) that are below the "
"minimum allowed (1, 1, 1).",
x, y, z);
- DAWN_INVALID_IF(x > limits.maxComputeWorkgroupSizeX || y > limits.maxComputeWorkgroupSizeY ||
- z > limits.maxComputeWorkgroupSizeZ,
- "Entry-point uses workgroup_size(%u, %u, %u) that exceeds the "
- "maximum allowed (%u, %u, %u).",
- x, y, z, limits.maxComputeWorkgroupSizeX, limits.maxComputeWorkgroupSizeY,
- limits.maxComputeWorkgroupSizeZ);
+ if (DAWN_UNLIKELY(x > limits.maxComputeWorkgroupSizeX || y > limits.maxComputeWorkgroupSizeY ||
+ z > limits.maxComputeWorkgroupSizeZ)) {
+ SupportedLimits adapterLimits;
+ wgpu::Status status = adapter->APIGetLimits(&adapterLimits);
+ DAWN_ASSERT(status == wgpu::Status::Success);
+
+ uint32_t maxComputeWorkgroupSizeXAdapterLimit =
+ adapterLimits.limits.maxComputeWorkgroupSizeX;
+ uint32_t maxComputeWorkgroupSizeYAdapterLimit =
+ adapterLimits.limits.maxComputeWorkgroupSizeY;
+ uint32_t maxComputeWorkgroupSizeZAdapterLimit =
+ adapterLimits.limits.maxComputeWorkgroupSizeZ;
+ std::string increaseLimitAdvice =
+ (x <= maxComputeWorkgroupSizeXAdapterLimit &&
+ y <= maxComputeWorkgroupSizeYAdapterLimit && z <= maxComputeWorkgroupSizeZAdapterLimit)
+ ? absl::StrFormat(
+ " This adapter supports higher maxComputeWorkgroupSizeX of %u, "
+ "maxComputeWorkgroupSizeY of %u, and maxComputeWorkgroupSizeZ of %u, which "
+ "can be specified in requiredLimits when calling requestDevice(). Limits "
+ "differ by hardware, so always check the adapter limits prior to requesting "
+ "a higher limit.",
+ maxComputeWorkgroupSizeXAdapterLimit, maxComputeWorkgroupSizeYAdapterLimit,
+ maxComputeWorkgroupSizeZAdapterLimit)
+ : "";
+ return DAWN_VALIDATION_ERROR(
+ "Entry-point uses workgroup_size(%u, %u, %u) that exceeds the "
+ "maximum allowed (%u, %u, %u).%s",
+ x, y, z, limits.maxComputeWorkgroupSizeX, limits.maxComputeWorkgroupSizeY,
+ limits.maxComputeWorkgroupSizeZ, increaseLimitAdvice);
+ }
uint64_t numInvocations = static_cast<uint64_t>(x) * y * z;
- DAWN_INVALID_IF(numInvocations > limits.maxComputeInvocationsPerWorkgroup,
- "The total number of workgroup invocations (%u) exceeds the "
- "maximum allowed (%u).",
- numInvocations, limits.maxComputeInvocationsPerWorkgroup);
+ uint32_t maxComputeInvocationsPerWorkgroup = limits.maxComputeInvocationsPerWorkgroup;
+ DAWN_INVALID_IF(
+ numInvocations > maxComputeInvocationsPerWorkgroup,
+ "The total number of workgroup invocations (%u) exceeds the "
+ "maximum allowed (%u).%s",
+ numInvocations, maxComputeInvocationsPerWorkgroup,
+ DAWN_INCREASE_LIMIT_MESSAGE(adapter, maxComputeInvocationsPerWorkgroup, numInvocations));
- DAWN_INVALID_IF(workgroupStorageSize > limits.maxComputeWorkgroupStorageSize,
- "The total use of workgroup storage (%u bytes) is larger than "
- "the maximum allowed (%u bytes).",
- workgroupStorageSize, limits.maxComputeWorkgroupStorageSize);
+ uint32_t maxComputeWorkgroupStorageSize = limits.maxComputeWorkgroupStorageSize;
+ DAWN_INVALID_IF(
+ workgroupStorageSize > maxComputeWorkgroupStorageSize,
+ "The total use of workgroup storage (%u bytes) is larger than "
+ "the maximum allowed (%u bytes).%s",
+ workgroupStorageSize, maxComputeWorkgroupStorageSize,
+ DAWN_INCREASE_LIMIT_MESSAGE(adapter, maxComputeWorkgroupStorageSize, workgroupStorageSize));
return Extent3D{x, y, z};
}
diff --git a/src/dawn/native/ShaderModule.h b/src/dawn/native/ShaderModule.h
index ee688ef..3d371c8 100644
--- a/src/dawn/native/ShaderModule.h
+++ b/src/dawn/native/ShaderModule.h
@@ -145,19 +145,19 @@
// 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);
+ResultOrError<Extent3D> ValidateComputeStageWorkgroupSize(const tint::Program& program,
+ const char* entryPointName,
+ const LimitsForCompilationRequest& limits,
+ const AdapterBase* adapter);
// Return extent3D with workgroup size dimension info if it is valid.
// width = x, height = y, depthOrArrayLength = z.
-ResultOrError<Extent3D> ValidateComputeStageWorkgroupSize(
- uint32_t x,
- uint32_t y,
- uint32_t z,
- size_t workgroupStorageSize,
- const LimitsForCompilationRequest& limits);
+ResultOrError<Extent3D> ValidateComputeStageWorkgroupSize(uint32_t x,
+ uint32_t y,
+ uint32_t z,
+ size_t workgroupStorageSize,
+ const LimitsForCompilationRequest& limits,
+ const AdapterBase* adapter);
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 6c91cd1..d339d7c 100644
--- a/src/dawn/native/d3d/D3DCompilationRequest.h
+++ b/src/dawn/native/d3d/D3DCompilationRequest.h
@@ -32,6 +32,7 @@
#include <string_view>
+#include "dawn/native/Adapter.h"
#include "dawn/native/CacheRequest.h"
#include "dawn/native/Serializable.h"
#include "dawn/native/d3d/d3d_platform.h"
@@ -75,6 +76,7 @@
X(tint::hlsl::writer::Options, tintOptions) \
X(std::optional<tint::ast::transform::SubstituteOverride::Config>, substituteOverrideConfig) \
X(LimitsForCompilationRequest, limits) \
+ X(CacheKey::UnsafeUnkeyedValue<const AdapterBase*>, adapter) \
X(bool, disableSymbolRenaming) \
X(bool, dumpShaders) \
X(bool, useTintIR)
diff --git a/src/dawn/native/d3d/ShaderUtils.cpp b/src/dawn/native/d3d/ShaderUtils.cpp
index 7ddcea0..bf92254 100644
--- a/src/dawn/native/d3d/ShaderUtils.cpp
+++ b/src/dawn/native/d3d/ShaderUtils.cpp
@@ -284,14 +284,15 @@
DAWN_TRY_ASSIGN(
_, ValidateComputeStageWorkgroupSize(
result->workgroup_info.x, result->workgroup_info.y, result->workgroup_info.z,
- result->workgroup_info.storage_size, r.limits));
+ result->workgroup_info.storage_size, r.limits, r.adapter.UnsafeGetValue()));
}
} else {
// Validate workgroup size after program runs transforms.
if (r.stage == SingleShaderStage::Compute) {
Extent3D _;
- DAWN_TRY_ASSIGN(_, ValidateComputeStageWorkgroupSize(
- transformedProgram, kRemappedEntryPointName, r.limits));
+ DAWN_TRY_ASSIGN(
+ _, ValidateComputeStageWorkgroupSize(transformedProgram, kRemappedEntryPointName,
+ r.limits, r.adapter.UnsafeGetValue()));
}
result = tint::hlsl::writer::Generate(transformedProgram, r.tintOptions);
diff --git a/src/dawn/native/d3d11/ShaderModuleD3D11.cpp b/src/dawn/native/d3d11/ShaderModuleD3D11.cpp
index 9bf7b8a..038982a 100644
--- a/src/dawn/native/d3d11/ShaderModuleD3D11.cpp
+++ b/src/dawn/native/d3d11/ShaderModuleD3D11.cpp
@@ -214,6 +214,7 @@
const CombinedLimits& limits = device->GetLimits();
req.hlsl.limits = LimitsForCompilationRequest::Create(limits.v1);
+ req.hlsl.adapter = UnsafeUnkeyedValue(static_cast<const AdapterBase*>(device->GetAdapter()));
req.hlsl.tintOptions.disable_robustness = !device->IsRobustnessEnabled();
req.hlsl.tintOptions.disable_workgroup_init =
diff --git a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
index 7bd281f..c4629bc 100644
--- a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
+++ b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
@@ -374,6 +374,7 @@
const CombinedLimits& limits = device->GetLimits();
req.hlsl.limits = LimitsForCompilationRequest::Create(limits.v1);
+ req.hlsl.adapter = UnsafeUnkeyedValue(static_cast<const AdapterBase*>(device->GetAdapter()));
CacheResult<d3d::CompiledShader> compiledShader;
DAWN_TRY_LOAD_OR_RUN(compiledShader, device, std::move(req), d3d::CompiledShader::FromBlob,
diff --git a/src/dawn/native/metal/ShaderModuleMTL.mm b/src/dawn/native/metal/ShaderModuleMTL.mm
index 1584bf2..3e35318 100644
--- a/src/dawn/native/metal/ShaderModuleMTL.mm
+++ b/src/dawn/native/metal/ShaderModuleMTL.mm
@@ -28,6 +28,7 @@
#include "dawn/native/metal/ShaderModuleMTL.h"
#include "dawn/common/MatchVariant.h"
+#include "dawn/native/Adapter.h"
#include "dawn/native/BindGroupLayout.h"
#include "dawn/native/CacheRequest.h"
#include "dawn/native/Serializable.h"
@@ -63,6 +64,7 @@
X(OptionalVertexPullingTransformConfig, vertexPullingTransformConfig) \
X(std::optional<tint::ast::transform::SubstituteOverride::Config>, substituteOverrideConfig) \
X(LimitsForCompilationRequest, limits) \
+ X(CacheKey::UnsafeUnkeyedValue<const AdapterBase*>, adapter) \
X(std::string, entryPointName) \
X(bool, disableSymbolRenaming) \
X(tint::msl::writer::Options, tintOptions) \
@@ -299,6 +301,7 @@
const CombinedLimits& limits = device->GetLimits();
req.limits = LimitsForCompilationRequest::Create(limits.v1);
+ req.adapter = UnsafeUnkeyedValue(static_cast<const AdapterBase*>(device->GetAdapter()));
CacheResult<MslCompilation> mslCompilation;
DAWN_TRY_LOAD_OR_RUN(
@@ -364,17 +367,19 @@
// overrides to have been substituted.
if (r.stage == SingleShaderStage::Compute) {
// Validate workgroup size and workgroup storage size.
- DAWN_TRY_ASSIGN(localSize,
- ValidateComputeStageWorkgroupSize(
- result->workgroup_info.x, result->workgroup_info.y,
- result->workgroup_info.z,
- result->workgroup_info.storage_size, r.limits));
+ DAWN_TRY_ASSIGN(
+ localSize,
+ ValidateComputeStageWorkgroupSize(
+ result->workgroup_info.x, result->workgroup_info.y,
+ result->workgroup_info.z, result->workgroup_info.storage_size, r.limits,
+ r.adapter.UnsafeGetValue()));
}
} else {
if (r.stage == SingleShaderStage::Compute) {
// Validate workgroup size after program runs transforms.
DAWN_TRY_ASSIGN(localSize, ValidateComputeStageWorkgroupSize(
- program, kRemappedEntryPointName, r.limits));
+ program, kRemappedEntryPointName, r.limits,
+ r.adapter.UnsafeGetValue()));
}
result = tint::msl::writer::Generate(program, r.tintOptions);
diff --git a/src/dawn/native/null/DeviceNull.cpp b/src/dawn/native/null/DeviceNull.cpp
index a9a66ef..1088993 100644
--- a/src/dawn/native/null/DeviceNull.cpp
+++ b/src/dawn/native/null/DeviceNull.cpp
@@ -495,10 +495,10 @@
// Do the workgroup size validation.
const CombinedLimits& limits = GetDevice()->GetLimits();
Extent3D _;
- DAWN_TRY_ASSIGN(
- _, ValidateComputeStageWorkgroupSize(transformedProgram, computeStage.entryPoint.c_str(),
- LimitsForCompilationRequest::Create(limits.v1)));
-
+ DAWN_TRY_ASSIGN(_, ValidateComputeStageWorkgroupSize(
+ transformedProgram, computeStage.entryPoint.c_str(),
+ LimitsForCompilationRequest::Create(limits.v1),
+ static_cast<const AdapterBase*>(GetDevice()->GetAdapter())));
return {};
}
diff --git a/src/dawn/native/opengl/ShaderModuleGL.cpp b/src/dawn/native/opengl/ShaderModuleGL.cpp
index 1b63fee..30cfc1f 100644
--- a/src/dawn/native/opengl/ShaderModuleGL.cpp
+++ b/src/dawn/native/opengl/ShaderModuleGL.cpp
@@ -31,6 +31,7 @@
#include <utility>
#include "absl/container/flat_hash_map.h"
+#include "dawn/native/Adapter.h"
#include "dawn/native/BindGroupLayoutInternal.h"
#include "dawn/native/CacheRequest.h"
#include "dawn/native/Pipeline.h"
@@ -93,6 +94,7 @@
X(SingleShaderStage, stage) \
X(std::optional<tint::ast::transform::SubstituteOverride::Config>, substituteOverrideConfig) \
X(LimitsForCompilationRequest, limits) \
+ X(CacheKey::UnsafeUnkeyedValue<const AdapterBase*>, adapter) \
X(bool, disableSymbolRenaming) \
X(std::vector<InterstageLocationAndName>, interstageVariables) \
X(tint::glsl::writer::Options, tintOptions) \
@@ -450,6 +452,8 @@
req.entryPointName = programmableStage.entryPoint;
req.substituteOverrideConfig = std::move(substituteOverrideConfig);
req.limits = LimitsForCompilationRequest::Create(limits.v1);
+ req.adapter = UnsafeUnkeyedValue(static_cast<const AdapterBase*>(GetDevice()->GetAdapter()));
+
req.platform = UnsafeUnkeyedValue(GetDevice()->GetPlatform());
req.tintOptions.version = tint::glsl::writer::Version(ToTintGLStandard(version.GetStandard()),
@@ -535,10 +539,11 @@
if (r.stage == SingleShaderStage::Compute) {
// Validate workgroup size after program runs transforms.
Extent3D _;
- DAWN_TRY_ASSIGN(_, ValidateComputeStageWorkgroupSize(
- result->workgroup_info.x, result->workgroup_info.y,
- result->workgroup_info.z,
- result->workgroup_info.storage_size, r.limits));
+ DAWN_TRY_ASSIGN(
+ _, ValidateComputeStageWorkgroupSize(
+ result->workgroup_info.x, result->workgroup_info.y,
+ result->workgroup_info.z, result->workgroup_info.storage_size, r.limits,
+ r.adapter.UnsafeGetValue()));
}
return GLSLCompilation{{std::move(result->glsl)}};
diff --git a/src/dawn/native/vulkan/ShaderModuleVk.cpp b/src/dawn/native/vulkan/ShaderModuleVk.cpp
index dcda04d..f28a88f 100644
--- a/src/dawn/native/vulkan/ShaderModuleVk.cpp
+++ b/src/dawn/native/vulkan/ShaderModuleVk.cpp
@@ -36,6 +36,7 @@
#include "absl/container/flat_hash_map.h"
#include "dawn/common/HashUtils.h"
#include "dawn/common/MatchVariant.h"
+#include "dawn/native/Adapter.h"
#include "dawn/native/CacheRequest.h"
#include "dawn/native/PhysicalDevice.h"
#include "dawn/native/Serializable.h"
@@ -197,6 +198,7 @@
X(const tint::Program*, inputProgram) \
X(std::optional<tint::ast::transform::SubstituteOverride::Config>, substituteOverrideConfig) \
X(LimitsForCompilationRequest, limits) \
+ X(CacheKey::UnsafeUnkeyedValue<const AdapterBase*>, adapter) \
X(std::string_view, entryPointName) \
X(tint::spirv::writer::Options, tintOptions) \
X(CacheKey::UnsafeUnkeyedValue<dawn::platform::Platform*>, platform)
@@ -378,6 +380,7 @@
const CombinedLimits& limits = GetDevice()->GetLimits();
req.limits = LimitsForCompilationRequest::Create(limits.v1);
+ req.adapter = UnsafeUnkeyedValue(static_cast<const AdapterBase*>(GetDevice()->GetAdapter()));
CacheResult<CompiledSpirv> compilation;
DAWN_TRY_LOAD_OR_RUN(
@@ -425,10 +428,11 @@
// have been substituted.
if (r.stage == SingleShaderStage::Compute) {
Extent3D _;
- DAWN_TRY_ASSIGN(_, ValidateComputeStageWorkgroupSize(
- tintResult->workgroup_info.x, tintResult->workgroup_info.y,
- tintResult->workgroup_info.z,
- tintResult->workgroup_info.storage_size, r.limits));
+ DAWN_TRY_ASSIGN(
+ _, ValidateComputeStageWorkgroupSize(
+ tintResult->workgroup_info.x, tintResult->workgroup_info.y,
+ tintResult->workgroup_info.z, tintResult->workgroup_info.storage_size,
+ r.limits, r.adapter.UnsafeGetValue()));
}
CompiledSpirv result;
