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

Port ComputeSharedMemoryTests to WGSL 

Tint now supports barriers

Bug: dawn:572
Change-Id: I1b5cdd9026e0ead7ecd743295c81ef147bdc5080
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/45602
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Commit-Queue: Corentin Wallez <cwallez@chromium.org>
This commit is contained in:
Ben Clayton 2021-03-24 09:48:11 +00:00 committed by Commit Bot service account
parent fcafc6e347
commit 218d48890a
1 changed files with 21 additions and 19 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

40
src/tests/end2end/ComputeSharedMemoryTests.cpp
View File

@ -27,7 +27,7 @@ class ComputeSharedMemoryTests : public DawnTest {
void ComputeSharedMemoryTests::BasicTest(const char* shader) {
// Set up shader and pipeline
auto module = utils::CreateShaderModule(device, utils::SingleShaderStage::Compute, shader);
auto module = utils::CreateShaderModuleFromWGSL(device, shader);
wgpu::ComputePipelineDescriptor csDesc;
csDesc.computeStage.module = module;
@ -70,31 +70,33 @@ void ComputeSharedMemoryTests::BasicTest(const char* shader) {
// Basic shared memory test
TEST_P(ComputeSharedMemoryTests, Basic) {
// TODO(crbug.com/tint/375): Implement barriers in Tint.
DAWN_SKIP_TEST_IF(HasToggleEnabled("use_tint_generator"));
BasicTest(R"(
#version 450
const uint kTileSize = 4;
const uint kInstances = 11;
const kTileSize : u32 = 4;
const kInstances : u32 = 11;
layout(local_size_x = kTileSize, local_size_y = kTileSize, local_size_z = 1) in;
layout(std140, set = 0, binding = 0) buffer Dst { uint x; } dst;
shared uint tmp;
[[builtin(local_invocation_id)]] var<in> LocalInvocationID : vec3<u32>;
void main() {
uint index = gl_LocalInvocationID.y * kTileSize + gl_LocalInvocationID.x;
if (index == 0) {
tmp = 0;
[[block]] struct Dst {
x : u32;
};
[[group(0), binding(0)]] var<storage> dst : Dst;
var<workgroup> tmp : u32;
[[stage(compute), workgroup_size(4,4,1)]]
fn main() -> void {
var index : u32 = LocalInvocationID.y * kTileSize + LocalInvocationID.x;
if (index == 0u) {
tmp = 0u;
}
barrier();
for (uint i = 0; i < kInstances; ++i) {
workgroupBarrier();
for (var i : u32 = 0u; i < kInstances; i = i + 1u) {
if (i == index) {
tmp++;
tmp = tmp + 1u;
}
barrier();
workgroupBarrier();
}
if (index == 0) {
if (index == 0u) {
dst.x = tmp;
}
})");