1 // Copyright 2021 The Dawn Authors
2 //
3 // Licensed under the Apache License, Version 2.0 (the "License");
4 // you may not use this file except in compliance with the License.
5 // You may obtain a copy of the License at
6 //
7 // http://www.apache.org/licenses/LICENSE-2.0
8 //
9 // Unless required by applicable law or agreed to in writing, software
10 // distributed under the License is distributed on an "AS IS" BASIS,
11 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12 // See the License for the specific language governing permissions and
13 // limitations under the License.
14
15 #include "tests/DawnTest.h"
16
17 #include "common/Math.h"
18 #include "common/Platform.h"
19 #include "utils/WGPUHelpers.h"
20
21 class MaxLimitTests : public DawnTest {
22 public:
GetRequiredLimits(const wgpu::SupportedLimits & supported)23 wgpu::RequiredLimits GetRequiredLimits(const wgpu::SupportedLimits& supported) override {
24 wgpu::RequiredLimits required = {};
25 required.limits = supported.limits;
26 return required;
27 }
28 };
29
30 // Test using the maximum amount of workgroup memory works
TEST_P(MaxLimitTests,MaxComputeWorkgroupStorageSize)31 TEST_P(MaxLimitTests, MaxComputeWorkgroupStorageSize) {
32 uint32_t maxComputeWorkgroupStorageSize =
33 GetSupportedLimits().limits.maxComputeWorkgroupStorageSize;
34
35 std::string shader = R"(
36 [[block]] struct Dst {
37 value0 : u32;
38 value1 : u32;
39 };
40
41 [[group(0), binding(0)]] var<storage, write> dst : Dst;
42
43 struct WGData {
44 value0 : u32;
45 // padding such that value0 and value1 are the first and last bytes of the memory.
46 [[size()" + std::to_string(maxComputeWorkgroupStorageSize / 4 - 2) +
47 R"()]] padding : u32;
48 value1 : u32;
49 };
50 var<workgroup> wg_data : WGData;
51
52 [[stage(compute), workgroup_size(2,1,1)]]
53 fn main([[builtin(local_invocation_index)]] LocalInvocationIndex : u32) {
54 if (LocalInvocationIndex == 0u) {
55 // Put data into the first and last byte of workgroup memory.
56 wg_data.value0 = 79u;
57 wg_data.value1 = 42u;
58 }
59
60 workgroupBarrier();
61
62 if (LocalInvocationIndex == 1u) {
63 // Read data out of workgroup memory into a storage buffer.
64 dst.value0 = wg_data.value0;
65 dst.value1 = wg_data.value1;
66 }
67 }
68 )";
69 wgpu::ComputePipelineDescriptor csDesc;
70 csDesc.compute.module = utils::CreateShaderModule(device, shader.c_str());
71 csDesc.compute.entryPoint = "main";
72 wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&csDesc);
73
74 // Set up dst storage buffer
75 wgpu::BufferDescriptor dstDesc;
76 dstDesc.size = 8;
77 dstDesc.usage =
78 wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst;
79 wgpu::Buffer dst = device.CreateBuffer(&dstDesc);
80
81 // Set up bind group and issue dispatch
82 wgpu::BindGroup bindGroup = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0),
83 {
84 {0, dst},
85 });
86
87 wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
88 wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
89 pass.SetPipeline(pipeline);
90 pass.SetBindGroup(0, bindGroup);
91 pass.Dispatch(1);
92 pass.EndPass();
93 wgpu::CommandBuffer commands = encoder.Finish();
94 queue.Submit(1, &commands);
95
96 EXPECT_BUFFER_U32_EQ(79, dst, 0);
97 EXPECT_BUFFER_U32_EQ(42, dst, 4);
98 }
99
100 // Test using the maximum uniform/storage buffer binding size works
TEST_P(MaxLimitTests,MaxBufferBindingSize)101 TEST_P(MaxLimitTests, MaxBufferBindingSize) {
102 // The uniform buffer layout used in this test is not supported on ES.
103 DAWN_TEST_UNSUPPORTED_IF(IsOpenGLES());
104
105 // TODO(crbug.com/dawn/1172)
106 DAWN_SUPPRESS_TEST_IF(IsWindows() && IsVulkan() && IsIntel());
107
108 // TODO(crbug.com/dawn/1217): Remove this suppression.
109 DAWN_SUPPRESS_TEST_IF(IsWindows() && IsVulkan() && IsNvidia());
110
111 for (wgpu::BufferUsage usage : {wgpu::BufferUsage::Storage, wgpu::BufferUsage::Uniform}) {
112 uint64_t maxBufferBindingSize;
113 std::string shader;
114 switch (usage) {
115 case wgpu::BufferUsage::Storage:
116 maxBufferBindingSize = GetSupportedLimits().limits.maxStorageBufferBindingSize;
117 // TODO(crbug.com/dawn/1160): Usually can't actually allocate a buffer this large
118 // because allocating the buffer for zero-initialization fails.
119 maxBufferBindingSize =
120 std::min(maxBufferBindingSize, uint64_t(2) * 1024 * 1024 * 1024);
121 // With WARP or on 32-bit platforms, such large buffer allocations often fail.
122 #ifdef DAWN_PLATFORM_32_BIT
123 if (IsWindows()) {
124 continue;
125 }
126 #endif
127 if (IsWARP()) {
128 maxBufferBindingSize =
129 std::min(maxBufferBindingSize, uint64_t(512) * 1024 * 1024);
130 }
131 shader = R"(
132 [[block]] struct Buf {
133 values : array<u32>;
134 };
135
136 [[block]] struct Result {
137 value0 : u32;
138 value1 : u32;
139 };
140
141 [[group(0), binding(0)]] var<storage, read> buf : Buf;
142 [[group(0), binding(1)]] var<storage, write> result : Result;
143
144 [[stage(compute), workgroup_size(1,1,1)]]
145 fn main() {
146 result.value0 = buf.values[0];
147 result.value1 = buf.values[arrayLength(&buf.values) - 1u];
148 }
149 )";
150 break;
151 case wgpu::BufferUsage::Uniform:
152 maxBufferBindingSize = GetSupportedLimits().limits.maxUniformBufferBindingSize;
153
154 // Clamp to not exceed the maximum i32 value for the WGSL [[size(x)]] annotation.
155 maxBufferBindingSize = std::min(maxBufferBindingSize,
156 uint64_t(std::numeric_limits<int32_t>::max()) + 8);
157
158 shader = R"(
159 [[block]] struct Buf {
160 value0 : u32;
161 // padding such that value0 and value1 are the first and last bytes of the memory.
162 [[size()" +
163 std::to_string(maxBufferBindingSize - 8) + R"()]] padding : u32;
164 value1 : u32;
165 };
166
167 [[block]] struct Result {
168 value0 : u32;
169 value1 : u32;
170 };
171
172 [[group(0), binding(0)]] var<uniform> buf : Buf;
173 [[group(0), binding(1)]] var<storage, write> result : Result;
174
175 [[stage(compute), workgroup_size(1,1,1)]]
176 fn main() {
177 result.value0 = buf.value0;
178 result.value1 = buf.value1;
179 }
180 )";
181 break;
182 default:
183 UNREACHABLE();
184 }
185
186 device.PushErrorScope(wgpu::ErrorFilter::OutOfMemory);
187
188 wgpu::BufferDescriptor bufDesc;
189 bufDesc.size = Align(maxBufferBindingSize, 4);
190 bufDesc.usage = usage | wgpu::BufferUsage::CopyDst;
191 wgpu::Buffer buffer = device.CreateBuffer(&bufDesc);
192
193 WGPUErrorType oomResult;
194 device.PopErrorScope([](WGPUErrorType type, const char*,
195 void* userdata) { *static_cast<WGPUErrorType*>(userdata) = type; },
196 &oomResult);
197 FlushWire();
198 // Max buffer size is smaller than the max buffer binding size.
199 DAWN_TEST_UNSUPPORTED_IF(oomResult == WGPUErrorType_OutOfMemory);
200
201 wgpu::BufferDescriptor resultBufDesc;
202 resultBufDesc.size = 8;
203 resultBufDesc.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc;
204 wgpu::Buffer resultBuffer = device.CreateBuffer(&resultBufDesc);
205
206 uint32_t value0 = 89234;
207 queue.WriteBuffer(buffer, 0, &value0, sizeof(value0));
208
209 uint32_t value1 = 234;
210 uint64_t value1Offset = Align(maxBufferBindingSize - sizeof(value1), 4);
211 queue.WriteBuffer(buffer, value1Offset, &value1, sizeof(value1));
212
213 wgpu::ComputePipelineDescriptor csDesc;
214 csDesc.compute.module = utils::CreateShaderModule(device, shader.c_str());
215 csDesc.compute.entryPoint = "main";
216 wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&csDesc);
217
218 wgpu::BindGroup bindGroup = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0),
219 {{0, buffer}, {1, resultBuffer}});
220
221 wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
222 wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
223 pass.SetPipeline(pipeline);
224 pass.SetBindGroup(0, bindGroup);
225 pass.Dispatch(1);
226 pass.EndPass();
227 wgpu::CommandBuffer commands = encoder.Finish();
228 queue.Submit(1, &commands);
229
230 EXPECT_BUFFER_U32_EQ(value0, resultBuffer, 0)
231 << "maxBufferBindingSize=" << maxBufferBindingSize << "; offset=" << 0
232 << "; usage=" << usage;
233 EXPECT_BUFFER_U32_EQ(value1, resultBuffer, 4)
234 << "maxBufferBindingSize=" << maxBufferBindingSize << "; offset=" << value1Offset
235 << "; usage=" << usage;
236 }
237 }
238
239 DAWN_INSTANTIATE_TEST(MaxLimitTests,
240 D3D12Backend(),
241 MetalBackend(),
242 OpenGLBackend(),
243 OpenGLESBackend(),
244 VulkanBackend());
245