| // Copyright 2023 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 <algorithm> |
| #include <array> |
| #include <string> |
| #include <vector> |
| |
| #include "dawn/common/Assert.h" |
| #include "dawn/common/Constants.h" |
| #include "dawn/common/Math.h" |
| #include "dawn/tests/DawnTest.h" |
| #include "dawn/utils/ComboRenderPipelineDescriptor.h" |
| #include "dawn/utils/WGPUHelpers.h" |
| |
| namespace dawn { |
| namespace { |
| |
| constexpr wgpu::TextureFormat kDefaultFormat = wgpu::TextureFormat::RGBA8Unorm; |
| |
| wgpu::Texture Create2DTexture(wgpu::Device device, |
| uint32_t width, |
| uint32_t height, |
| uint32_t arrayLayerCount, |
| uint32_t mipLevelCount, |
| uint32_t sampleCount, |
| wgpu::TextureUsage usage) { |
| wgpu::TextureDescriptor descriptor; |
| descriptor.dimension = wgpu::TextureDimension::e2D; |
| descriptor.size.width = width; |
| descriptor.size.height = height; |
| descriptor.size.depthOrArrayLayers = arrayLayerCount; |
| descriptor.sampleCount = sampleCount; |
| descriptor.format = kDefaultFormat; |
| descriptor.mipLevelCount = mipLevelCount; |
| descriptor.usage = usage; |
| return device.CreateTexture(&descriptor); |
| } |
| |
| class TextureShaderBuiltinTests : public DawnTest { |
| protected: |
| wgpu::Texture CreateTexture(uint32_t arrayLayerCount, |
| uint32_t mipLevelCount, |
| uint32_t sampleCount) { |
| DAWN_ASSERT(arrayLayerCount > 0 && mipLevelCount > 0); |
| DAWN_ASSERT(sampleCount == 1 || sampleCount == 4); |
| |
| const uint32_t textureWidthLevel0 = 1 << mipLevelCount; |
| const uint32_t textureHeightLevel0 = 1 << mipLevelCount; |
| constexpr wgpu::TextureUsage kUsage = wgpu::TextureUsage::CopyDst | |
| wgpu::TextureUsage::TextureBinding | |
| wgpu::TextureUsage::RenderAttachment; |
| return Create2DTexture(device, textureWidthLevel0, textureHeightLevel0, arrayLayerCount, |
| mipLevelCount, sampleCount, kUsage); |
| } |
| |
| wgpu::TextureView CreateTextureView(const wgpu::Texture& tex, |
| wgpu::TextureViewDimension dimension, |
| uint32_t baseMipLevel = 0, |
| uint32_t mipLevelCount = wgpu::kMipLevelCountUndefined) { |
| wgpu::TextureViewDescriptor descriptor; |
| descriptor.dimension = dimension; |
| // textureNumLevels return texture view levels |
| descriptor.baseMipLevel = baseMipLevel; |
| descriptor.mipLevelCount = mipLevelCount; |
| return tex.CreateView(&descriptor); |
| } |
| }; |
| |
| // Note: the following tests testing textureNumLevels and textureNumSamples behavior is mainly |
| // targeted at OpenGL/OpenGLES backend without native GLSL support for these builtins. |
| // These tests should be trivial for otherbackend, and thus can be used as control case. |
| |
| // Test calling textureNumLevels & textureNumSamples in one shader. |
| TEST_P(TextureShaderBuiltinTests, Basic) { |
| constexpr uint32_t kLayers = 3; |
| constexpr uint32_t kMipLevels = 2; |
| wgpu::Texture tex1 = CreateTexture(kLayers, kMipLevels, 1); |
| wgpu::TextureView texView1 = CreateTextureView(tex1, wgpu::TextureViewDimension::e2DArray); |
| |
| constexpr uint32_t kSampleCount = 4; |
| wgpu::Texture tex2 = CreateTexture(1, 1, kSampleCount); |
| wgpu::TextureView texView2 = tex2.CreateView(); |
| |
| constexpr uint32_t kMipLevelsView = 1; |
| wgpu::Texture tex3 = CreateTexture(kLayers, kMipLevels, 1); |
| wgpu::TextureView texView3 = |
| CreateTextureView(tex3, wgpu::TextureViewDimension::e2D, 1, kMipLevelsView); |
| |
| const uint32_t expected[] = { |
| kLayers, |
| kMipLevels, |
| kSampleCount, |
| kMipLevelsView, |
| }; |
| |
| wgpu::BufferDescriptor bufferDesc; |
| bufferDesc.size = sizeof(expected); |
| bufferDesc.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc; |
| wgpu::Buffer buffer = device.CreateBuffer(&bufferDesc); |
| |
| wgpu::ComputePipelineDescriptor pipelineDescriptor; |
| pipelineDescriptor.compute.module = utils::CreateShaderModule(device, R"( |
| @group(0) @binding(0) var<storage, read_write> dstBuf : array<u32>; |
| @group(0) @binding(1) var tex1 : texture_2d_array<f32>; |
| // Use sparse binding to test impact of binding remapping |
| @group(0) @binding(4) var tex2 : texture_multisampled_2d<f32>; |
| @group(1) @binding(3) var tex3 : texture_2d<f32>; |
| |
| @compute @workgroup_size(1, 1, 1) fn main() { |
| dstBuf[0] = textureNumLayers(tex1); // control case |
| dstBuf[1] = textureNumLevels(tex1); |
| dstBuf[2] = textureNumSamples(tex2); |
| dstBuf[3] = textureNumLevels(tex3); |
| } |
| )"); |
| pipelineDescriptor.compute.entryPoint = "main"; |
| wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&pipelineDescriptor); |
| |
| wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); |
| { |
| wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); |
| pass.SetPipeline(pipeline); |
| pass.SetBindGroup(0, utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), |
| { |
| {0, buffer}, |
| {1, texView1}, |
| {4, texView2}, |
| })); |
| pass.SetBindGroup(1, utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(1), |
| { |
| {3, texView3}, |
| })); |
| pass.DispatchWorkgroups(1); |
| pass.End(); |
| } |
| |
| wgpu::CommandBuffer commands = encoder.Finish(); |
| queue.Submit(1, &commands); |
| |
| EXPECT_BUFFER_U32_RANGE_EQ(expected, buffer, 0, sizeof(expected) / sizeof(uint32_t)); |
| } |
| |
| // Test calling textureNumLevels & textureNumSamples inside function and taking a function param as |
| // the argument. |
| TEST_P(TextureShaderBuiltinTests, BuiltinCallInFunction) { |
| constexpr uint32_t kLayers = 3; |
| constexpr uint32_t kMipLevels1 = 2; |
| wgpu::Texture tex1 = CreateTexture(kLayers, kMipLevels1, 1); |
| wgpu::TextureView texView1 = CreateTextureView(tex1, wgpu::TextureViewDimension::e2DArray); |
| constexpr uint32_t kMipLevels2 = 5; |
| wgpu::Texture tex2 = CreateTexture(1, kMipLevels2, 1); |
| wgpu::TextureView texView2 = CreateTextureView(tex2, wgpu::TextureViewDimension::e2DArray); |
| |
| const uint32_t expected[] = { |
| kLayers, kMipLevels1, kMipLevels1, kMipLevels2, kMipLevels1 + 100u, |
| }; |
| |
| wgpu::BufferDescriptor bufferDesc; |
| bufferDesc.size = sizeof(expected); |
| bufferDesc.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc; |
| wgpu::Buffer buffer = device.CreateBuffer(&bufferDesc); |
| |
| wgpu::ComputePipelineDescriptor pipelineDescriptor; |
| pipelineDescriptor.compute.module = utils::CreateShaderModule(device, R"( |
| @group(0) @binding(0) var<storage, read_write> dstBuf : array<u32>; |
| @group(0) @binding(1) var tex1 : texture_2d_array<f32>; |
| @group(0) @binding(2) var tex2 : texture_2d_array<f32>; |
| |
| fn f(tex: texture_2d_array<f32>) -> u32 { |
| // TODO(tint:2006) Workaround to preserve usage of tex param. Remove when bug is fixed. |
| var result = textureNumLayers(tex); |
| result = textureNumLevels(tex); |
| return result; |
| } |
| |
| fn f_nested(tex: texture_2d_array<f32>, d: u32) -> u32 { |
| return f(tex) + d; |
| } |
| |
| @compute @workgroup_size(1, 1, 1) fn main() { |
| dstBuf[0] = textureNumLayers(tex1); // control case |
| dstBuf[1] = textureNumLevels(tex1); |
| dstBuf[2] = f(tex1); |
| dstBuf[3] = f(tex2); |
| dstBuf[4] = f_nested(tex1, 100u); |
| } |
| )"); |
| pipelineDescriptor.compute.entryPoint = "main"; |
| wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&pipelineDescriptor); |
| |
| wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); |
| { |
| wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); |
| pass.SetPipeline(pipeline); |
| pass.SetBindGroup(0, utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), |
| { |
| {0, buffer}, |
| {1, texView1}, |
| {2, texView2}, |
| })); |
| pass.DispatchWorkgroups(1); |
| pass.End(); |
| } |
| |
| wgpu::CommandBuffer commands = encoder.Finish(); |
| queue.Submit(1, &commands); |
| |
| EXPECT_BUFFER_U32_RANGE_EQ(expected, buffer, 0, sizeof(expected) / sizeof(uint32_t)); |
| } |
| |
| // Test the internal uniform buffer data is properly updated between dispatches |
| // When the same pipeline is set only once. |
| TEST_P(TextureShaderBuiltinTests, OnePipelineMultipleDispatches) { |
| const char* shader = R"( |
| @group(0) @binding(0) var<storage, read_write> dstBuf : array<u32>; |
| @group(0) @binding(1) var tex1 : texture_2d_array<f32>; |
| // Use sparse binding to test impact of binding remapping |
| @group(0) @binding(4) var tex2 : texture_multisampled_2d<f32>; |
| @group(1) @binding(3) var tex3 : texture_2d_array<f32>; |
| |
| @compute @workgroup_size(1, 1, 1) fn main() { |
| dstBuf[0] = textureNumLayers(tex1); // control case |
| dstBuf[1] = textureNumLevels(tex1); |
| dstBuf[2] = textureNumSamples(tex2); |
| dstBuf[3] = textureNumLevels(tex3); |
| } |
| )"; |
| |
| constexpr uint32_t kLayers_1 = 3; |
| constexpr uint32_t kMipLevels_1 = 2; |
| wgpu::Texture tex1_1 = CreateTexture(kLayers_1, kMipLevels_1, 1); |
| wgpu::TextureView texView1_1 = CreateTextureView(tex1_1, wgpu::TextureViewDimension::e2DArray); |
| constexpr uint32_t kLayers_2 = 5; |
| constexpr uint32_t kMipLevels_2 = 4; |
| wgpu::Texture tex1_2 = CreateTexture(kLayers_2, kMipLevels_2, 1); |
| wgpu::TextureView texView1_2 = CreateTextureView(tex1_2, wgpu::TextureViewDimension::e2DArray); |
| |
| constexpr uint32_t kSampleCount_1 = 4; |
| wgpu::Texture tex2_1 = CreateTexture(1, 1, kSampleCount_1); |
| wgpu::TextureView texView2_1 = tex2_1.CreateView(); |
| constexpr uint32_t kSampleCount_2 = 4; |
| wgpu::Texture tex2_2 = CreateTexture(1, 1, kSampleCount_2); |
| wgpu::TextureView texView2_2 = tex2_2.CreateView(); |
| |
| constexpr uint32_t kMipLevelsView_1 = 1; |
| wgpu::Texture tex3_1 = CreateTexture(kLayers_1, kMipLevels_1, 1); |
| wgpu::TextureView texView3_1 = |
| CreateTextureView(tex3_1, wgpu::TextureViewDimension::e2DArray, 0, kMipLevelsView_1); |
| constexpr uint32_t kMipLevelsView_2 = 2; |
| wgpu::Texture tex3_2 = CreateTexture(kLayers_2, kMipLevels_2, 1); |
| wgpu::TextureView texView3_2 = |
| CreateTextureView(tex3_2, wgpu::TextureViewDimension::e2DArray, 0, kMipLevelsView_2); |
| |
| constexpr uint32_t expected_1[] = { |
| // Output from first dispatch |
| kLayers_1, |
| kMipLevels_1, |
| kSampleCount_1, |
| kMipLevelsView_1, |
| }; |
| constexpr uint32_t expected_2[] = { |
| // Output from second dispatch with different bind group |
| kLayers_2, |
| kMipLevels_2, |
| kSampleCount_2, |
| kMipLevelsView_2, |
| }; |
| constexpr uint32_t expected_3[] = { |
| // Output from third dispatch with bind group partially reset |
| kLayers_1, |
| kMipLevels_1, |
| kSampleCount_1, |
| kMipLevelsView_2, |
| }; |
| DAWN_ASSERT(sizeof(expected_1) == sizeof(expected_2)); |
| DAWN_ASSERT(sizeof(expected_1) == sizeof(expected_3)); |
| |
| wgpu::BufferDescriptor bufferDesc; |
| bufferDesc.size = sizeof(expected_1); |
| bufferDesc.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc; |
| wgpu::Buffer buffer_1 = device.CreateBuffer(&bufferDesc); |
| wgpu::Buffer buffer_2 = device.CreateBuffer(&bufferDesc); |
| wgpu::Buffer buffer_3 = device.CreateBuffer(&bufferDesc); |
| |
| wgpu::ComputePipeline pipeline; |
| { |
| wgpu::ComputePipelineDescriptor pipelineDescriptor; |
| pipelineDescriptor.compute.module = utils::CreateShaderModule(device, shader); |
| pipelineDescriptor.compute.entryPoint = "main"; |
| pipeline = device.CreateComputePipeline(&pipelineDescriptor); |
| } |
| |
| wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); |
| { |
| wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); |
| pass.SetPipeline(pipeline); |
| |
| pass.SetBindGroup(0, utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), |
| { |
| {0, buffer_1}, |
| {1, texView1_1}, |
| {4, texView2_1}, |
| })); |
| pass.SetBindGroup(1, utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(1), |
| { |
| {3, texView3_1}, |
| })); |
| pass.DispatchWorkgroups(1); |
| |
| pass.SetBindGroup(0, utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), |
| { |
| {0, buffer_2}, |
| {1, texView1_2}, |
| {4, texView2_2}, |
| })); |
| pass.SetBindGroup(1, utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(1), |
| { |
| {3, texView3_2}, |
| })); |
| pass.DispatchWorkgroups(1); |
| |
| pass.SetBindGroup(0, utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), |
| { |
| {0, buffer_3}, |
| {1, texView1_1}, |
| {4, texView2_1}, |
| })); |
| // Note: bind group 1 is not set |
| pass.DispatchWorkgroups(1); |
| |
| pass.End(); |
| } |
| |
| wgpu::CommandBuffer commands = encoder.Finish(); |
| queue.Submit(1, &commands); |
| |
| EXPECT_BUFFER_U32_RANGE_EQ(expected_1, buffer_1, 0, sizeof(expected_1) / sizeof(uint32_t)); |
| EXPECT_BUFFER_U32_RANGE_EQ(expected_2, buffer_2, 0, sizeof(expected_2) / sizeof(uint32_t)); |
| EXPECT_BUFFER_U32_RANGE_EQ(expected_3, buffer_3, 0, sizeof(expected_3) / sizeof(uint32_t)); |
| } |
| |
| // Test textureNumLevels & textureNumSamples results correctness used in multiple pipelines sharing |
| // same shader module. |
| TEST_P(TextureShaderBuiltinTests, OneShaderModuleMultipleEntryPoints) { |
| const char* shader = R"( |
| @group(0) @binding(0) var<storage, read_write> dstBuf : array<u32>; |
| @group(0) @binding(1) var tex1 : texture_2d_array<f32>; |
| // Use sparse binding to test impact of binding remapping |
| @group(0) @binding(4) var tex2 : texture_multisampled_2d<f32>; |
| @group(1) @binding(3) var tex3 : texture_2d<f32>; |
| |
| @compute @workgroup_size(1, 1, 1) fn main1() { |
| dstBuf[0] = textureNumLayers(tex1); // control case |
| dstBuf[1] = textureNumLevels(tex1); |
| dstBuf[2] = textureNumSamples(tex2); |
| dstBuf[3] = textureNumLevels(tex3); |
| } |
| |
| @compute @workgroup_size(1, 1, 1) fn main2() { |
| dstBuf[0] = textureNumLayers(tex1); // control case |
| dstBuf[1] = textureNumLevels(tex1); |
| dstBuf[2] = textureNumSamples(tex2); |
| // _ = textureNumLevels(tex3); |
| dstBuf[3] = 99; |
| } |
| )"; |
| |
| constexpr uint32_t kLayers_1 = 3; |
| constexpr uint32_t kMipLevels_1 = 2; |
| wgpu::Texture tex1_1 = CreateTexture(kLayers_1, kMipLevels_1, 1); |
| wgpu::TextureView texView1_1 = CreateTextureView(tex1_1, wgpu::TextureViewDimension::e2DArray); |
| constexpr uint32_t kLayers_2 = 5; |
| constexpr uint32_t kMipLevels_2 = 4; |
| wgpu::Texture tex1_2 = CreateTexture(kLayers_2, kMipLevels_2, 1); |
| wgpu::TextureView texView1_2 = CreateTextureView(tex1_2, wgpu::TextureViewDimension::e2DArray); |
| |
| constexpr uint32_t kSampleCount_1 = 4; |
| wgpu::Texture tex2_1 = CreateTexture(1, 1, kSampleCount_1); |
| wgpu::TextureView texView2_1 = tex2_1.CreateView(); |
| // constexpr uint32_t kSampleCount_2 = 1; |
| constexpr uint32_t kSampleCount_2 = 4; |
| wgpu::Texture tex2_2 = CreateTexture(1, 1, kSampleCount_2); |
| wgpu::TextureView texView2_2 = tex2_2.CreateView(); |
| |
| constexpr uint32_t kMipLevelsView_1 = 1; |
| wgpu::Texture tex3_1 = CreateTexture(kLayers_1, kMipLevels_1, 1); |
| wgpu::TextureView texView3_1 = |
| CreateTextureView(tex3_1, wgpu::TextureViewDimension::e2D, 0, kMipLevelsView_1); |
| constexpr uint32_t kMipLevelsView_2 = 1; |
| wgpu::Texture tex3_2 = CreateTexture(kLayers_2, kMipLevels_2, 1); |
| wgpu::TextureView texView3_2 = |
| CreateTextureView(tex3_2, wgpu::TextureViewDimension::e2D, 0, kMipLevelsView_2); |
| |
| constexpr uint32_t expected_1[] = { |
| // Output from first dispatch |
| kLayers_1, |
| kMipLevels_1, |
| kSampleCount_1, |
| kMipLevelsView_1, |
| }; |
| constexpr uint32_t expected_2[] = { |
| // Output from second dispatch with different bind group |
| kLayers_2, |
| kMipLevels_2, |
| kSampleCount_2, |
| 99, |
| }; |
| DAWN_ASSERT(sizeof(expected_1) == sizeof(expected_2)); |
| |
| wgpu::BufferDescriptor bufferDesc; |
| bufferDesc.size = sizeof(expected_1); |
| bufferDesc.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc; |
| wgpu::Buffer buffer_1 = device.CreateBuffer(&bufferDesc); |
| wgpu::Buffer buffer_2 = device.CreateBuffer(&bufferDesc); |
| |
| wgpu::ShaderModule module = utils::CreateShaderModule(device, shader); |
| wgpu::ComputePipeline pipeline_1; |
| { |
| wgpu::ComputePipelineDescriptor pipelineDescriptor; |
| pipelineDescriptor.compute.module = module; |
| pipelineDescriptor.compute.entryPoint = "main1"; |
| pipeline_1 = device.CreateComputePipeline(&pipelineDescriptor); |
| } |
| wgpu::ComputePipeline pipeline_2; |
| { |
| wgpu::ComputePipelineDescriptor pipelineDescriptor; |
| pipelineDescriptor.compute.module = module; |
| pipelineDescriptor.compute.entryPoint = "main2"; |
| pipeline_2 = device.CreateComputePipeline(&pipelineDescriptor); |
| } |
| |
| wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); |
| { |
| wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); |
| |
| pass.SetPipeline(pipeline_1); |
| pass.SetBindGroup(0, utils::MakeBindGroup(device, pipeline_1.GetBindGroupLayout(0), |
| { |
| {0, buffer_1}, |
| {1, texView1_1}, |
| {4, texView2_1}, |
| })); |
| pass.SetBindGroup(1, utils::MakeBindGroup(device, pipeline_1.GetBindGroupLayout(1), |
| { |
| {3, texView3_1}, |
| })); |
| pass.DispatchWorkgroups(1); |
| |
| pass.SetPipeline(pipeline_2); |
| pass.SetBindGroup(0, utils::MakeBindGroup(device, pipeline_2.GetBindGroupLayout(0), |
| { |
| {0, buffer_2}, |
| {1, texView1_2}, |
| {4, texView2_2}, |
| })); |
| pass.DispatchWorkgroups(1); |
| |
| pass.End(); |
| } |
| |
| wgpu::CommandBuffer commands = encoder.Finish(); |
| queue.Submit(1, &commands); |
| |
| EXPECT_BUFFER_U32_RANGE_EQ(expected_1, buffer_1, 0, sizeof(expected_1) / sizeof(uint32_t)); |
| EXPECT_BUFFER_U32_RANGE_EQ(expected_2, buffer_2, 0, sizeof(expected_2) / sizeof(uint32_t)); |
| } |
| |
| DAWN_INSTANTIATE_TEST(TextureShaderBuiltinTests, |
| D3D11Backend(), |
| D3D12Backend(), |
| MetalBackend(), |
| OpenGLBackend(), |
| OpenGLESBackend(), |
| VulkanBackend()); |
| |
| } // anonymous namespace |
| } // namespace dawn |
