D3D12: Support feature chromium_experimental_dp4a
Bug: tint:1497
Test: dawn_end2end_tests
Change-Id: I57d5c06c15c0c366c7cc239426e5eee3a7237101
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/90028
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Commit-Queue: Corentin Wallez <cwallez@chromium.org>
diff --git a/dawn.json b/dawn.json
index 2e82ff3..4fa8cb9 100644
--- a/dawn.json
+++ b/dawn.json
@@ -1201,6 +1201,7 @@
{"name": "depth clamping", "type": "bool", "default": "false"},
{"name": "depth24 unorm stencil8", "type": "bool", "default": "false"},
{"name": "depth32 float stencil8", "type": "bool", "default": "false"},
+ {"name": "chromium experimental dp4a", "type": "bool", "default": "false"},
{"name": "invalid feature", "type": "bool", "default": "false"},
{"name": "dawn internal usages", "type": "bool", "default": "false"},
{"name": "dawn native", "type": "bool", "default": "false"},
@@ -1361,7 +1362,8 @@
{"value": 1001, "name": "dawn shader float 16", "tags": ["dawn"]},
{"value": 1002, "name": "dawn internal usages", "tags": ["dawn"]},
{"value": 1003, "name": "dawn multi planar formats", "tags": ["dawn"]},
- {"value": 1004, "name": "dawn native", "tags": ["dawn", "native"]}
+ {"value": 1004, "name": "dawn native", "tags": ["dawn", "native"]},
+ {"value": 1005, "name": "chromium experimental dp4a", "tags": ["dawn"]}
]
},
"filter mode": {
diff --git a/src/dawn/native/Device.cpp b/src/dawn/native/Device.cpp
index 1f11514..fe7b00a 100644
--- a/src/dawn/native/Device.cpp
+++ b/src/dawn/native/Device.cpp
@@ -182,6 +182,8 @@
if (togglesDesc != nullptr) {
ApplyToggleOverrides(togglesDesc);
}
+
+ SetDefaultToggles();
ApplyFeatures(descriptor);
DawnCacheDeviceDescriptor defaultCacheDesc = {};
@@ -198,7 +200,6 @@
}
mFormatTable = BuildFormatTable(this);
- SetDefaultToggles();
SetWGSLExtensionAllowList();
@@ -1234,13 +1235,25 @@
}
bool DeviceBase::IsFeatureEnabled(Feature feature) const {
- return mEnabledFeatures.IsEnabled(feature);
+ if (mEnabledFeatures.IsEnabled(feature)) {
+ // Currently we can only use DXC to compile HLSL shaders using float16, and
+ // ChromiumExperimentalDp4a is an experimental feature which can only be enabled with toggle
+ // "use_dxc".
+ if (feature == Feature::ChromiumExperimentalDp4a || feature == Feature::ShaderFloat16) {
+ return IsToggleEnabled(Toggle::UseDXC);
+ }
+ return true;
+ }
+ return false;
}
void DeviceBase::SetWGSLExtensionAllowList() {
// Set the WGSL extensions allow list based on device's enabled features and other
// propority. For example:
// mWGSLExtensionAllowList.insert("InternalExtensionForTesting");
+ if (IsFeatureEnabled(Feature::ChromiumExperimentalDp4a)) {
+ mWGSLExtensionAllowList.insert("chromium_experimental_dp4a");
+ }
}
WGSLExtensionSet DeviceBase::GetWGSLExtensionAllowList() const {
diff --git a/src/dawn/native/Device.h b/src/dawn/native/Device.h
index c6822de..6c66f6b 100644
--- a/src/dawn/native/Device.h
+++ b/src/dawn/native/Device.h
@@ -265,7 +265,11 @@
QueueBase* APIGetQueue();
bool APIGetLimits(SupportedLimits* limits) const;
+ // Note that we should not use this function to query the features which can only be enabled
+ // behind toggles (use IsFeatureEnabled() instead).
bool APIHasFeature(wgpu::FeatureName feature) const;
+ // Note that we should not use this function to query the features which can only be enabled
+ // behind toggles (use IsFeatureEnabled() instead).
size_t APIEnumerateFeatures(wgpu::FeatureName* features) const;
void APIInjectError(wgpu::ErrorType type, const char* message);
bool APITick();
diff --git a/src/dawn/native/Features.cpp b/src/dawn/native/Features.cpp
index c33b4af..e2b82ee 100644
--- a/src/dawn/native/Features.cpp
+++ b/src/dawn/native/Features.cpp
@@ -32,64 +32,69 @@
using FeatureEnumAndInfoList =
std::array<FeatureEnumAndInfo, static_cast<size_t>(Feature::EnumCount)>;
-static constexpr FeatureEnumAndInfoList kFeatureNameAndInfoList = {
- {{Feature::TextureCompressionBC,
- {"texture-compression-bc", "Support Block Compressed (BC) texture formats",
- "https://bugs.chromium.org/p/dawn/issues/detail?id=42"},
- &WGPUDeviceProperties::textureCompressionBC},
- {Feature::TextureCompressionETC2,
- {"texture-compression-etc2",
- "Support Ericsson Texture Compressed (ETC2/EAC) texture "
- "formats",
- "https://bugs.chromium.org/p/dawn/issues/detail?id=955"},
- &WGPUDeviceProperties::textureCompressionETC2},
- {Feature::TextureCompressionASTC,
- {"texture-compression-astc",
- "Support Adaptable Scalable Texture Compressed (ASTC) "
- "texture formats",
- "https://bugs.chromium.org/p/dawn/issues/detail?id=955"},
- &WGPUDeviceProperties::textureCompressionASTC},
- {Feature::ShaderFloat16,
- {"shader-float16",
- "Support 16bit float arithmetic and declarations in uniform and storage buffers",
- "https://bugs.chromium.org/p/dawn/issues/detail?id=426"},
- &WGPUDeviceProperties::shaderFloat16},
- {Feature::PipelineStatisticsQuery,
- {"pipeline-statistics-query", "Support Pipeline Statistics Query",
- "https://bugs.chromium.org/p/dawn/issues/detail?id=434"},
- &WGPUDeviceProperties::pipelineStatisticsQuery},
- {Feature::TimestampQuery,
- {"timestamp-query", "Support Timestamp Query",
- "https://bugs.chromium.org/p/dawn/issues/detail?id=434"},
- &WGPUDeviceProperties::timestampQuery},
- {Feature::DepthClamping,
- {"depth-clamping", "Clamp depth to [0, 1] in NDC space instead of clipping",
- "https://bugs.chromium.org/p/dawn/issues/detail?id=716"},
- &WGPUDeviceProperties::depthClamping},
- {Feature::Depth24UnormStencil8,
- {"depth24unorm-stencil8", "Support depth24unorm-stencil8 texture format",
- "https://bugs.chromium.org/p/dawn/issues/detail?id=690"},
- &WGPUDeviceProperties::depth24UnormStencil8},
- {Feature::Depth32FloatStencil8,
- {"depth32float-stencil8", "Support depth32float-stencil8 texture format",
- "https://bugs.chromium.org/p/dawn/issues/detail?id=690"},
- &WGPUDeviceProperties::depth32FloatStencil8},
- {Feature::DawnInternalUsages,
- {"dawn-internal-usages",
- "Add internal usages to resources to affect how the texture is allocated, but not "
- "frontend validation. Other internal commands may access this usage.",
- "https://dawn.googlesource.com/dawn/+/refs/heads/main/docs/dawn/features/"
- "dawn_internal_usages.md"},
- &WGPUDeviceProperties::dawnInternalUsages},
- {Feature::MultiPlanarFormats,
- {"multiplanar-formats", "Import and use multi-planar texture formats with per plane views",
- "https://bugs.chromium.org/p/dawn/issues/detail?id=551"},
- &WGPUDeviceProperties::multiPlanarFormats},
- {Feature::DawnNative,
- {"dawn-native", "WebGPU is running on top of dawn_native.",
- "https://dawn.googlesource.com/dawn/+/refs/heads/main/docs/dawn/features/"
- "dawn_native.md"},
- &WGPUDeviceProperties::dawnNative}}};
+static constexpr FeatureEnumAndInfoList kFeatureNameAndInfoList = {{
+ {Feature::TextureCompressionBC,
+ {"texture-compression-bc", "Support Block Compressed (BC) texture formats",
+ "https://bugs.chromium.org/p/dawn/issues/detail?id=42"},
+ &WGPUDeviceProperties::textureCompressionBC},
+ {Feature::TextureCompressionETC2,
+ {"texture-compression-etc2",
+ "Support Ericsson Texture Compressed (ETC2/EAC) texture "
+ "formats",
+ "https://bugs.chromium.org/p/dawn/issues/detail?id=955"},
+ &WGPUDeviceProperties::textureCompressionETC2},
+ {Feature::TextureCompressionASTC,
+ {"texture-compression-astc",
+ "Support Adaptable Scalable Texture Compressed (ASTC) "
+ "texture formats",
+ "https://bugs.chromium.org/p/dawn/issues/detail?id=955"},
+ &WGPUDeviceProperties::textureCompressionASTC},
+ {Feature::ShaderFloat16,
+ {"shader-float16",
+ "Support 16bit float arithmetic and declarations in uniform and storage buffers",
+ "https://bugs.chromium.org/p/dawn/issues/detail?id=426"},
+ &WGPUDeviceProperties::shaderFloat16},
+ {Feature::PipelineStatisticsQuery,
+ {"pipeline-statistics-query", "Support Pipeline Statistics Query",
+ "https://bugs.chromium.org/p/dawn/issues/detail?id=434"},
+ &WGPUDeviceProperties::pipelineStatisticsQuery},
+ {Feature::TimestampQuery,
+ {"timestamp-query", "Support Timestamp Query",
+ "https://bugs.chromium.org/p/dawn/issues/detail?id=434"},
+ &WGPUDeviceProperties::timestampQuery},
+ {Feature::DepthClamping,
+ {"depth-clamping", "Clamp depth to [0, 1] in NDC space instead of clipping",
+ "https://bugs.chromium.org/p/dawn/issues/detail?id=716"},
+ &WGPUDeviceProperties::depthClamping},
+ {Feature::Depth24UnormStencil8,
+ {"depth24unorm-stencil8", "Support depth24unorm-stencil8 texture format",
+ "https://bugs.chromium.org/p/dawn/issues/detail?id=690"},
+ &WGPUDeviceProperties::depth24UnormStencil8},
+ {Feature::Depth32FloatStencil8,
+ {"depth32float-stencil8", "Support depth32float-stencil8 texture format",
+ "https://bugs.chromium.org/p/dawn/issues/detail?id=690"},
+ &WGPUDeviceProperties::depth32FloatStencil8},
+ {Feature::ChromiumExperimentalDp4a,
+ {"chromium-experimental-dp4a", "Support experimental DP4a instructions in WGSL",
+ "https://bugs.chromium.org/p/tint/issues/detail?id=1497"},
+ &WGPUDeviceProperties::chromiumExperimentalDp4a},
+ {Feature::DawnInternalUsages,
+ {"dawn-internal-usages",
+ "Add internal usages to resources to affect how the texture is allocated, but not "
+ "frontend validation. Other internal commands may access this usage.",
+ "https://dawn.googlesource.com/dawn/+/refs/heads/main/docs/dawn/features/"
+ "dawn_internal_usages.md"},
+ &WGPUDeviceProperties::dawnInternalUsages},
+ {Feature::MultiPlanarFormats,
+ {"multiplanar-formats", "Import and use multi-planar texture formats with per plane views",
+ "https://bugs.chromium.org/p/dawn/issues/detail?id=551"},
+ &WGPUDeviceProperties::multiPlanarFormats},
+ {Feature::DawnNative,
+ {"dawn-native", "WebGPU is running on top of dawn_native.",
+ "https://dawn.googlesource.com/dawn/+/refs/heads/main/docs/dawn/features/"
+ "dawn_native.md"},
+ &WGPUDeviceProperties::dawnNative},
+}};
Feature FromAPIFeature(wgpu::FeatureName feature) {
switch (feature) {
@@ -120,6 +125,8 @@
return Feature::MultiPlanarFormats;
case wgpu::FeatureName::DawnNative:
return Feature::DawnNative;
+ case wgpu::FeatureName::ChromiumExperimentalDp4a:
+ return Feature::ChromiumExperimentalDp4a;
case wgpu::FeatureName::IndirectFirstInstance:
return Feature::InvalidEnum;
@@ -153,6 +160,8 @@
return wgpu::FeatureName::DawnMultiPlanarFormats;
case Feature::DawnNative:
return wgpu::FeatureName::DawnNative;
+ case Feature::ChromiumExperimentalDp4a:
+ return wgpu::FeatureName::ChromiumExperimentalDp4a;
case Feature::EnumCount:
break;
diff --git a/src/dawn/native/Features.h b/src/dawn/native/Features.h
index c2439b2..a87eaea 100644
--- a/src/dawn/native/Features.h
+++ b/src/dawn/native/Features.h
@@ -36,6 +36,7 @@
DepthClamping,
Depth24UnormStencil8,
Depth32FloatStencil8,
+ ChromiumExperimentalDp4a,
// Dawn-specific
DawnInternalUsages,
diff --git a/src/dawn/native/d3d12/AdapterD3D12.cpp b/src/dawn/native/d3d12/AdapterD3D12.cpp
index 5f48014..4ac4134 100644
--- a/src/dawn/native/d3d12/AdapterD3D12.cpp
+++ b/src/dawn/native/d3d12/AdapterD3D12.cpp
@@ -23,6 +23,7 @@
#include "dawn/native/d3d12/D3D12Error.h"
#include "dawn/native/d3d12/DeviceD3D12.h"
#include "dawn/native/d3d12/PlatformFunctions.h"
+#include "dawn/native/d3d12/UtilsD3D12.h"
namespace dawn::native::d3d12 {
@@ -137,6 +138,17 @@
mSupportedFeatures.EnableFeature(Feature::Depth24UnormStencil8);
mSupportedFeatures.EnableFeature(Feature::Depth32FloatStencil8);
+ if (GetBackend()->GetFunctions()->IsDXCAvailable()) {
+ uint64_t dxcVersion = 0;
+ DAWN_TRY_ASSIGN(dxcVersion, GetBackend()->GetDXCompilerVersion());
+ constexpr uint64_t kLeastMajorVersionForDP4a = 1;
+ constexpr uint64_t kLeastMinorVersionForDP4a = 4;
+ if (mDeviceInfo.supportsDP4a &&
+ dxcVersion >= MakeDXCVersion(kLeastMajorVersionForDP4a, kLeastMinorVersionForDP4a)) {
+ mSupportedFeatures.EnableFeature(Feature::ChromiumExperimentalDp4a);
+ }
+ }
+
return {};
}
diff --git a/src/dawn/native/d3d12/BackendD3D12.cpp b/src/dawn/native/d3d12/BackendD3D12.cpp
index 1232dbf..3d0b13f 100644
--- a/src/dawn/native/d3d12/BackendD3D12.cpp
+++ b/src/dawn/native/d3d12/BackendD3D12.cpp
@@ -21,6 +21,7 @@
#include "dawn/native/d3d12/AdapterD3D12.h"
#include "dawn/native/d3d12/D3D12Error.h"
#include "dawn/native/d3d12/PlatformFunctions.h"
+#include "dawn/native/d3d12/UtilsD3D12.h"
namespace dawn::native::d3d12 {
@@ -141,6 +142,21 @@
return mDxcValidator;
}
+ResultOrError<uint64_t> Backend::GetDXCompilerVersion() {
+ DAWN_TRY(EnsureDxcValidator());
+
+ ComPtr<IDxcVersionInfo> versionInfo;
+ DAWN_TRY(CheckHRESULT(mDxcValidator.As(&versionInfo),
+ "D3D12 QueryInterface IDxcValidator to IDxcVersionInfo"));
+
+ uint32_t compilerMajor, compilerMinor;
+ DAWN_TRY(CheckHRESULT(versionInfo->GetVersion(&compilerMajor, &compilerMinor),
+ "IDxcVersionInfo::GetVersion"));
+
+ // Pack both into a single version number.
+ return MakeDXCVersion(compilerMajor, compilerMinor);
+}
+
const PlatformFunctions* Backend::GetFunctions() const {
return mFunctions.get();
}
diff --git a/src/dawn/native/d3d12/BackendD3D12.h b/src/dawn/native/d3d12/BackendD3D12.h
index 9a983f2..1bf1ead 100644
--- a/src/dawn/native/d3d12/BackendD3D12.h
+++ b/src/dawn/native/d3d12/BackendD3D12.h
@@ -40,6 +40,7 @@
ComPtr<IDxcLibrary> GetDxcLibrary() const;
ComPtr<IDxcCompiler> GetDxcCompiler() const;
ComPtr<IDxcValidator> GetDxcValidator() const;
+ ResultOrError<uint64_t> GetDXCompilerVersion();
const PlatformFunctions* GetFunctions() const;
diff --git a/src/dawn/native/d3d12/D3D12Info.cpp b/src/dawn/native/d3d12/D3D12Info.cpp
index b8fc896..6a4c435 100644
--- a/src/dawn/native/d3d12/D3D12Info.cpp
+++ b/src/dawn/native/d3d12/D3D12Info.cpp
@@ -75,10 +75,9 @@
}
}
- D3D12_FEATURE_DATA_SHADER_MODEL knownShaderModels[] = {{D3D_SHADER_MODEL_6_2},
- {D3D_SHADER_MODEL_6_1},
- {D3D_SHADER_MODEL_6_0},
- {D3D_SHADER_MODEL_5_1}};
+ D3D12_FEATURE_DATA_SHADER_MODEL knownShaderModels[] = {
+ {D3D_SHADER_MODEL_6_4}, {D3D_SHADER_MODEL_6_3}, {D3D_SHADER_MODEL_6_2},
+ {D3D_SHADER_MODEL_6_1}, {D3D_SHADER_MODEL_6_0}, {D3D_SHADER_MODEL_5_1}};
uint32_t driverShaderModel = 0;
for (D3D12_FEATURE_DATA_SHADER_MODEL shaderModel : knownShaderModels) {
if (SUCCEEDED(adapter.GetDevice()->CheckFeatureSupport(
@@ -118,6 +117,8 @@
driverShaderModel >= D3D_SHADER_MODEL_6_2 && featureData4.Native16BitShaderOpsSupported;
}
+ info.supportsDP4a = driverShaderModel >= D3D_SHADER_MODEL_6_4;
+
return std::move(info);
}
diff --git a/src/dawn/native/d3d12/D3D12Info.h b/src/dawn/native/d3d12/D3D12Info.h
index c0ffc47..e0f2a66 100644
--- a/src/dawn/native/d3d12/D3D12Info.h
+++ b/src/dawn/native/d3d12/D3D12Info.h
@@ -33,6 +33,7 @@
uint32_t shaderModel;
PerStage<std::wstring> shaderProfiles;
bool supportsSharedResourceCapabilityTier1;
+ bool supportsDP4a;
};
ResultOrError<D3D12DeviceInfo> GatherDeviceInfo(const Adapter& adapter);
diff --git a/src/dawn/native/d3d12/DeviceD3D12.cpp b/src/dawn/native/d3d12/DeviceD3D12.cpp
index 96fbb3f..97f7128 100644
--- a/src/dawn/native/d3d12/DeviceD3D12.cpp
+++ b/src/dawn/native/d3d12/DeviceD3D12.cpp
@@ -217,9 +217,6 @@
MaybeError Device::ApplyUseDxcToggle() {
if (!ToBackend(GetAdapter())->GetBackend()->GetFunctions()->IsDXCAvailable()) {
ForceSetToggle(Toggle::UseDXC, false);
- } else if (IsFeatureEnabled(Feature::ShaderFloat16)) {
- // Currently we can only use DXC to compile HLSL shaders using float16.
- ForceSetToggle(Toggle::UseDXC, true);
}
if (IsToggleEnabled(Toggle::UseDXC)) {
diff --git a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
index 1a9592c..b9d4ef4 100644
--- a/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
+++ b/src/dawn/native/d3d12/ShaderModuleD3D12.cpp
@@ -31,6 +31,8 @@
#include "dawn/native/CacheKey.h"
#include "dawn/native/Pipeline.h"
#include "dawn/native/TintUtils.h"
+#include "dawn/native/d3d12/AdapterD3D12.h"
+#include "dawn/native/d3d12/BackendD3D12.h"
#include "dawn/native/d3d12/BindGroupLayoutD3D12.h"
#include "dawn/native/d3d12/D3D12Error.h"
#include "dawn/native/d3d12/DeviceD3D12.h"
@@ -45,19 +47,6 @@
namespace dawn::native::d3d12 {
namespace {
-ResultOrError<uint64_t> GetDXCompilerVersion(ComPtr<IDxcValidator> dxcValidator) {
- ComPtr<IDxcVersionInfo> versionInfo;
- DAWN_TRY(CheckHRESULT(dxcValidator.As(&versionInfo),
- "D3D12 QueryInterface IDxcValidator to IDxcVersionInfo"));
-
- uint32_t compilerMajor, compilerMinor;
- DAWN_TRY(CheckHRESULT(versionInfo->GetVersion(&compilerMajor, &compilerMinor),
- "IDxcVersionInfo::GetVersion"));
-
- // Pack both into a single version number.
- return (uint64_t(compilerMajor) << uint64_t(32)) + compilerMinor;
-}
-
uint64_t GetD3DCompilerVersion() {
return D3D_COMPILER_VERSION;
}
@@ -222,7 +211,8 @@
uint64_t dxcVersion = 0;
if (device->IsToggleEnabled(Toggle::UseDXC)) {
compiler = Compiler::DXC;
- DAWN_TRY_ASSIGN(dxcVersion, GetDXCompilerVersion(device->GetDxcValidator()));
+ DAWN_TRY_ASSIGN(dxcVersion,
+ ToBackend(device->GetAdapter())->GetBackend()->GetDXCompilerVersion());
} else {
compiler = Compiler::FXC;
}
diff --git a/src/dawn/native/d3d12/UtilsD3D12.cpp b/src/dawn/native/d3d12/UtilsD3D12.cpp
index 0833e01..0e761f8 100644
--- a/src/dawn/native/d3d12/UtilsD3D12.cpp
+++ b/src/dawn/native/d3d12/UtilsD3D12.cpp
@@ -386,4 +386,8 @@
object->SetPrivateData(WKPDID_D3DDebugObjectName, objectName.length(), objectName.c_str());
}
+uint64_t MakeDXCVersion(uint64_t majorVersion, uint64_t minorVersion) {
+ return (majorVersion << 32) + minorVersion;
+}
+
} // namespace dawn::native::d3d12
diff --git a/src/dawn/native/d3d12/UtilsD3D12.h b/src/dawn/native/d3d12/UtilsD3D12.h
index 0bc5afb..1418f54 100644
--- a/src/dawn/native/d3d12/UtilsD3D12.h
+++ b/src/dawn/native/d3d12/UtilsD3D12.h
@@ -68,6 +68,8 @@
void SetDebugName(Device* device, ID3D12Object* object, const char* prefix, std::string label = "");
+uint64_t MakeDXCVersion(uint64_t majorVersion, uint64_t minorVersion);
+
} // namespace dawn::native::d3d12
#endif // SRC_DAWN_NATIVE_D3D12_UTILSD3D12_H_
diff --git a/src/dawn/tests/BUILD.gn b/src/dawn/tests/BUILD.gn
index 9f86b59..e0dbdda 100644
--- a/src/dawn/tests/BUILD.gn
+++ b/src/dawn/tests/BUILD.gn
@@ -447,6 +447,7 @@
"end2end/DrawTests.cpp",
"end2end/DynamicBufferOffsetTests.cpp",
"end2end/EntryPointTests.cpp",
+ "end2end/ExperimentalDP4aTests.cpp",
"end2end/ExternalTextureTests.cpp",
"end2end/FirstIndexOffsetTests.cpp",
"end2end/GpuMemorySynchronizationTests.cpp",
diff --git a/src/dawn/tests/end2end/ExperimentalDP4aTests.cpp b/src/dawn/tests/end2end/ExperimentalDP4aTests.cpp
new file mode 100644
index 0000000..e789c07
--- /dev/null
+++ b/src/dawn/tests/end2end/ExperimentalDP4aTests.cpp
@@ -0,0 +1,106 @@
+// Copyright 2022 The Dawn Authors
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <vector>
+
+#include "dawn/tests/DawnTest.h"
+#include "dawn/utils/WGPUHelpers.h"
+
+namespace {
+using RequestDP4aExtension = bool;
+DAWN_TEST_PARAM_STRUCT(ExperimentalDP4aTestsParams, RequestDP4aExtension);
+
+} // anonymous namespace
+
+class ExperimentalDP4aTests : public DawnTestWithParams<ExperimentalDP4aTestsParams> {
+ protected:
+ std::vector<wgpu::FeatureName> GetRequiredFeatures() override {
+ mIsDP4aSupportedOnAdapter = SupportsFeatures({wgpu::FeatureName::ChromiumExperimentalDp4a});
+ if (!mIsDP4aSupportedOnAdapter) {
+ return {};
+ }
+
+ if (GetParam().mRequestDP4aExtension) {
+ return {wgpu::FeatureName::ChromiumExperimentalDp4a};
+ }
+
+ return {};
+ }
+
+ bool IsDP4aSupportedOnAdapter() const { return mIsDP4aSupportedOnAdapter; }
+
+ private:
+ bool mIsDP4aSupportedOnAdapter = false;
+};
+
+TEST_P(ExperimentalDP4aTests, BasicDP4aFeaturesTest) {
+ const char* computeShader = R"(
+ enable chromium_experimental_dp4a;
+
+ struct Buf {
+ data1 : i32,
+ data2 : u32,
+ data3 : i32,
+ data4 : u32,
+ }
+ @group(0) @binding(0) var<storage, read_write> buf : Buf;
+
+ @stage(compute) @workgroup_size(1)
+ fn main() {
+ var a = 0xFFFFFFFFu;
+ var b = 0xFFFFFFFEu;
+ var c = 0x01020304u;
+ buf.data1 = dot4I8Packed(a, b);
+ buf.data2 = dot4U8Packed(a, b);
+ buf.data3 = dot4I8Packed(a, c);
+ buf.data4 = dot4U8Packed(a, c);
+ }
+)";
+ if (!GetParam().mRequestDP4aExtension || !IsDP4aSupportedOnAdapter() ||
+ !HasToggleEnabled("use_dxc")) {
+ ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, computeShader));
+ return;
+ }
+
+ wgpu::BufferDescriptor bufferDesc;
+ bufferDesc.size = 4 * sizeof(uint32_t);
+ bufferDesc.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc;
+ wgpu::Buffer bufferOut = device.CreateBuffer(&bufferDesc);
+
+ wgpu::ComputePipelineDescriptor csDesc;
+ csDesc.compute.module = utils::CreateShaderModule(device, computeShader);
+ csDesc.compute.entryPoint = "main";
+ wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&csDesc);
+
+ wgpu::BindGroup bindGroup = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0),
+ {
+ {0, bufferOut},
+ });
+
+ wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+ wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+ pass.SetPipeline(pipeline);
+ pass.SetBindGroup(0, bindGroup);
+ pass.DispatchWorkgroups(1);
+ pass.End();
+ wgpu::CommandBuffer commands = encoder.Finish();
+ queue.Submit(1, &commands);
+
+ uint32_t expected[] = {5, 259845, static_cast<uint32_t>(-10), 2550};
+ EXPECT_BUFFER_U32_RANGE_EQ(expected, bufferOut, 0, 4);
+}
+
+DAWN_INSTANTIATE_TEST_P(ExperimentalDP4aTests,
+ {D3D12Backend(), D3D12Backend({"use_dxc"})},
+ {true, false});
diff --git a/src/dawn/wire/SupportedFeatures.cpp b/src/dawn/wire/SupportedFeatures.cpp
index 8dd6d58..0e5688a 100644
--- a/src/dawn/wire/SupportedFeatures.cpp
+++ b/src/dawn/wire/SupportedFeatures.cpp
@@ -36,6 +36,7 @@
case WGPUFeatureName_DawnShaderFloat16:
case WGPUFeatureName_DawnInternalUsages:
case WGPUFeatureName_DawnMultiPlanarFormats:
+ case WGPUFeatureName_ChromiumExperimentalDp4a:
return true;
}
