encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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>
This commit is contained in:
Jiawei Shao 2022-05-20 08:21:00 +00:00 committed by Dawn LUCI CQ
parent 0c5aaf4833
commit 9ebba367e8
17 changed files with 243 additions and 82 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
dawn.json
View File

@ -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": {

17
src/dawn/native/Device.cpp
View File

@ -182,6 +182,8 @@ DeviceBase::DeviceBase(AdapterBase* adapter, const DeviceDescriptor* descriptor)
if (togglesDesc != nullptr) {
ApplyToggleOverrides(togglesDesc);
}
SetDefaultToggles();
ApplyFeatures(descriptor);
DawnCacheDeviceDescriptor defaultCacheDesc = {};
@ -198,7 +200,6 @@ DeviceBase::DeviceBase(AdapterBase* adapter, const DeviceDescriptor* descriptor)
}
mFormatTable = BuildFormatTable(this);
SetDefaultToggles();
SetWGSLExtensionAllowList();
@ -1234,13 +1235,25 @@ void DeviceBase::ApplyFeatures(const DeviceDescriptor* deviceDescriptor) {
}
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 {

4
src/dawn/native/Device.h
View File

@ -265,7 +265,11 @@ class DeviceBase : public RefCounted {
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();

125
src/dawn/native/Features.cpp
View File

@ -32,64 +32,69 @@ struct FeatureEnumAndInfo {
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 @@ Feature FromAPIFeature(wgpu::FeatureName feature) {
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 @@ wgpu::FeatureName ToAPIFeature(Feature feature) {
return wgpu::FeatureName::DawnMultiPlanarFormats;
case Feature::DawnNative:
return wgpu::FeatureName::DawnNative;
case Feature::ChromiumExperimentalDp4a:
return wgpu::FeatureName::ChromiumExperimentalDp4a;
case Feature::EnumCount:
break;

1
src/dawn/native/Features.h
View File

@ -36,6 +36,7 @@ enum class Feature {
DepthClamping,
Depth24UnormStencil8,
Depth32FloatStencil8,
ChromiumExperimentalDp4a,
// Dawn-specific
DawnInternalUsages,

12
src/dawn/native/d3d12/AdapterD3D12.cpp
View File

@ -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 @@ MaybeError Adapter::InitializeSupportedFeaturesImpl() {
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 {};
}

16
src/dawn/native/d3d12/BackendD3D12.cpp
View File

@ -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 @@ ComPtr<IDxcValidator> Backend::GetDxcValidator() const {
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();
}

1
src/dawn/native/d3d12/BackendD3D12.h
View File

@ -40,6 +40,7 @@ class Backend : public BackendConnection {
ComPtr<IDxcLibrary> GetDxcLibrary() const;
ComPtr<IDxcCompiler> GetDxcCompiler() const;
ComPtr<IDxcValidator> GetDxcValidator() const;
ResultOrError<uint64_t> GetDXCompilerVersion();
const PlatformFunctions* GetFunctions() const;

9
src/dawn/native/d3d12/D3D12Info.cpp
View File

@ -75,10 +75,9 @@ ResultOrError<D3D12DeviceInfo> GatherDeviceInfo(const Adapter& adapter) {
}
}
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 @@ ResultOrError<D3D12DeviceInfo> GatherDeviceInfo(const Adapter& adapter) {
driverShaderModel >= D3D_SHADER_MODEL_6_2 && featureData4.Native16BitShaderOpsSupported;
}
info.supportsDP4a = driverShaderModel >= D3D_SHADER_MODEL_6_4;
return std::move(info);
}

1
src/dawn/native/d3d12/D3D12Info.h
View File

@ -33,6 +33,7 @@ struct D3D12DeviceInfo {
uint32_t shaderModel;
PerStage<std::wstring> shaderProfiles;
bool supportsSharedResourceCapabilityTier1;
bool supportsDP4a;
};
ResultOrError<D3D12DeviceInfo> GatherDeviceInfo(const Adapter& adapter);

3
src/dawn/native/d3d12/DeviceD3D12.cpp
View File

@ -217,9 +217,6 @@ ComPtr<IDXGIFactory4> Device::GetFactory() const {
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)) {

18
src/dawn/native/d3d12/ShaderModuleD3D12.cpp
View File

@ -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 @@ struct ShaderCompilationRequest {
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;
}

4
src/dawn/native/d3d12/UtilsD3D12.cpp
View File

@ -386,4 +386,8 @@ void SetDebugName(Device* device, ID3D12Object* object, const char* prefix, std:
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

2
src/dawn/native/d3d12/UtilsD3D12.h
View File

@ -68,6 +68,8 @@ void RecordBufferTextureCopy(BufferTextureCopyDirection direction,
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_

1
src/dawn/tests/BUILD.gn
View File

@ -447,6 +447,7 @@ source_set("end2end_tests_sources") {
"end2end/DrawTests.cpp",
"end2end/DynamicBufferOffsetTests.cpp",
"end2end/EntryPointTests.cpp",
"end2end/ExperimentalDP4aTests.cpp",
"end2end/ExternalTextureTests.cpp",
"end2end/FirstIndexOffsetTests.cpp",
"end2end/GpuMemorySynchronizationTests.cpp",

106
src/dawn/tests/end2end/ExperimentalDP4aTests.cpp Normal file
View File

@ -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});

1
src/dawn/wire/SupportedFeatures.cpp
View File

@ -36,6 +36,7 @@ bool IsFeatureSupported(WGPUFeatureName feature) {
case WGPUFeatureName_DawnShaderFloat16:
case WGPUFeatureName_DawnInternalUsages:
case WGPUFeatureName_DawnMultiPlanarFormats:
case WGPUFeatureName_ChromiumExperimentalDp4a:
return true;
}