Emit array<u32> Storage Buffer Sizes in HLSL and MSL Transforms The HLSL and MSL immediate transforms stored the storage buffer sizes as array<vec4<u32>>, which only forced a 16-byte alignment without benefit since immediate data is read at a 4-byte stride. Emit a tightly packed array<u32> from both transforms and drop the per-backend element-type branch in the shared ArrayLengthFromImmediates read. On Metal, upload the sizes at the matching 4-byte-aligned offset. Use Align for the Metal setBytes upload size and backing array capacity. Make IsPowerOfTwo constexpr so Align can be evaluated at compile time. Bug: 366291600, 536438952 Change-Id: Iaac084eec79ab1843edc3be7781a6ed61ceae2fa Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/324277 Commit-Queue: Shaobo Yan <shaoboyan@microsoft.com> Reviewed-by: Corentin Wallez <cwallez@chromium.org>
diff --git a/src/dawn/native/metal/CommandBufferMTL.mm b/src/dawn/native/metal/CommandBufferMTL.mm index 546adfc..79bc5b4 100644 --- a/src/dawn/native/metal/CommandBufferMTL.mm +++ b/src/dawn/native/metal/CommandBufferMTL.mm
@@ -450,13 +450,8 @@ } } -// Metal uses a physical addressing mode which means buffers in the shading language are -// just pointers to the virtual address of their start. This means there is no way to know -// the length of a buffer to compute the arrayLength() of unsized arrays at the end of storage -// buffers. Tint implements the arrayLength() of unsized arrays by requiring immediates -// that stores the length of other buffers. This structure that keeps track of the -// length of storage buffers and apply them to the reserved "immediate blocks" when -// needed for a draw or a dispatch. +// Metal buffer pointers do not expose their length, so Tint receives storage buffer lengths +// through immediate data for arrayLength(). struct StorageBufferLengthTracker { StorageBufferLengthTracker() = delete; explicit StorageBufferLengthTracker(DeviceBase* device) { @@ -466,15 +461,10 @@ wgpu::ShaderStage dirtyStages = wgpu::ShaderStage::None; - // The lengths of buffers are stored as 32bit integers because that is the width the - // MSL code generated by Tint expects. - // UBOs require we align the max buffer count to 4 elements (16 bytes). - static constexpr size_t MaxBufferCount = ((kGenericMetalBufferSlots + 3) / 4) * 4; + // Tint expects one uint32_t length per possible buffer binding. + static constexpr size_t MaxBufferCount = kGenericMetalBufferSlots; PerStage<std::array<uint32_t, MaxBufferCount>> data; - // The actual size in bytes of the buffer length data to upload for each shader stage. - // This is calculated as sizeof(uint32_t) * aligned_buffer_count and represents the - // number of bytes that need to be uploaded to the GPU buffer containing buffer lengths. - // This size accounts for the 4-element (16-byte) alignment requirement for UBOs. + // Number of bytes to upload for each shader stage. PerStage<uint32_t> dataSize; // TODO(crbug.com/366291600): Remove this logic when merging @@ -506,15 +496,13 @@ bufferCount += uint32_t{pipeline->GetVertexBufferCount()}; } - bufferCount = Align(bufferCount, 4); DAWN_ASSERT(bufferCount <= data[SingleShaderStage::Vertex].size()); dataSize[SingleShaderStage::Vertex] = sizeof(uint32_t) * bufferCount; } if (dirtyStages & wgpu::ShaderStage::Fragment) { - uint32_t bufferCount = Align(ToBackend(pipeline->GetLayout()) - ->GetBufferBindingCount(SingleShaderStage::Fragment), - 4); + uint32_t bufferCount = ToBackend(pipeline->GetLayout()) + ->GetBufferBindingCount(SingleShaderStage::Fragment); DAWN_ASSERT(bufferCount <= data[SingleShaderStage::Fragment].size()); dataSize[SingleShaderStage::Fragment] = sizeof(uint32_t) * bufferCount; } @@ -527,8 +515,8 @@ return; } - uint32_t bufferCount = Align( - ToBackend(pipeline->GetLayout())->GetBufferBindingCount(SingleShaderStage::Compute), 4); + uint32_t bufferCount = + ToBackend(pipeline->GetLayout())->GetBufferBindingCount(SingleShaderStage::Compute); DAWN_ASSERT(bufferCount <= data[SingleShaderStage::Compute].size()); dataSize[SingleShaderStage::Compute] = sizeof(uint32_t) * bufferCount; } @@ -552,11 +540,8 @@ using ComputeImmediatesTracker = UserImmediatesTrackerBase<ComputeImmediates, ComputePipelineBase>; -// Template class that manages immediates for Metal backend. -// This tracker combines immediate data with buffer length information, -// uploading both to a single buffer that can be accessed by shaders. -// Template parameter T should be either RenderImmediatesTrackerBase -// or ComputeImmediatesTrackerBase. +// Uploads immediates, both external and internal (including storage buffer lengths) through one +// Metal setBytes call per stage. // TODO(crbug.com/366291600): Merge StorageBufferLength in ImmediateTracker template <typename T, typename EncoderType> class ImmediateTracker : public T { @@ -587,12 +572,8 @@ size * kImmediateElementByteSize); } - // Calculate buffer sizes start offset based on ImmediateBlock layouts - // describes in PipelineLayoutMTL.h - // - must be 16-byte aligned for UBO requirements - static_assert(16 % kImmediateElementByteSize == 0); - size_t bufferSizeOffsetElements = - RoundUp(pipelineMask.count(), 16 / kImmediateElementByteSize); + uint32_t bufferSizeByteOffset = GetImmediateBufferSizesByteOffset(pipelineMask); + size_t bufferSizeOffsetElements = bufferSizeByteOffset / kImmediateElementByteSize; // Update storage buffer length data that are needed and changed. for (auto stage : IterateStages(lengthTracker->dirtyStages)) { @@ -603,11 +584,9 @@ lengthTracker->data[stage].data(), lengthTracker->dataSize[stage]); } - // Update per stage dirty size. lengthTracker always keeps last valid length of buffer - // sizes. for (auto stage : IterateStages(stages)) { - dirtySize[stage] = bufferSizeOffsetElements * kImmediateElementByteSize + - lengthTracker->dataSize[stage]; + dirtySize[stage] = + Align(bufferSizeByteOffset + lengthTracker->dataSize[stage], kSetBytesAlignment); } // Reset StorageBufferLengthTracker dirty stages @@ -623,13 +602,17 @@ static constexpr bool kIsRenderImmediates = std::is_same_v<T, RenderImmediatesTracker>; static constexpr bool kIsComputeImmediates = std::is_same_v<T, ComputeImmediatesTracker>; - // The lengths of buffers are stored as 32bit integers because that is the width the - // MSL code generated by Tint expects. - // UBOs require we align the max buffer count to 4 elements (16 bytes). + // Round up to 16 bytes so the setBytes length covers the size of the constant-buffer struct + // argument, whose size is rounded up to its largest member alignment (at most 16 bytes, for + // vec4-class members). + static constexpr size_t kSetBytesAlignment = 16; + static constexpr size_t MaxBufferCount = StorageBufferLengthTracker::MaxBufferCount; + static constexpr size_t kImmediateStructU32 = kIsRenderImmediates + ? sizeof(RenderImmediates) / sizeof(uint32_t) + : sizeof(UserImmediates) / sizeof(uint32_t); static constexpr size_t kMaxImmediateBlockSize = - kIsRenderImmediates ? sizeof(RenderImmediates) / sizeof(uint32_t) + MaxBufferCount - : sizeof(UserImmediates) / sizeof(uint32_t) + MaxBufferCount; + Align(kImmediateStructU32 + MaxBufferCount, kSetBytesAlignment / kImmediateElementByteSize); // Writes data to the immediate block content for the specified shader stages. // This is used for both immediates and storage buffer length data.
diff --git a/src/dawn/native/metal/PipelineLayoutMTL.h b/src/dawn/native/metal/PipelineLayoutMTL.h index 5923099..a1c9618 100644 --- a/src/dawn/native/metal/PipelineLayoutMTL.h +++ b/src/dawn/native/metal/PipelineLayoutMTL.h
@@ -30,8 +30,10 @@ #import <Metal/Metal.h> +#include "src/dawn/common/Constants.h" #include "src/dawn/common/ityp_stack_vec.h" #include "src/dawn/native/BindingInfo.h" +#include "src/dawn/native/IntegerTypes.h" #include "src/dawn/native/PerStage.h" #include "src/dawn/native/PipelineLayout.h" @@ -41,13 +43,7 @@ // The number of Metal buffers usable by applications in general inline constexpr size_t kMetalBufferTableSize = 31; -// The Metal buffer slot that Dawn reserves for immediate block. -// The layout of ImmediateBlock: -// struct ImmediateBlock { -// - Normal render/compute immediates, ref to ImmediateLayout.h -// - Optional Paddings to align the following vec4 to 16 bytes -// - Storage Buffer sizes - vec4<u32> arrays -// }; +// Immediate blocks contain pipeline immediates followed by tightly packed storage buffer sizes. inline constexpr size_t kImmediateBlockBufferSlot = kMetalBufferTableSize - 1; // The number of Metal buffers Dawn can use in a generic way (i.e. that aren't reserved) inline constexpr size_t kGenericMetalBufferSlots = kMetalBufferTableSize - 1; @@ -55,6 +51,10 @@ // The Last buffer slot to be used by argument buffers inline constexpr size_t kArgumentBufferSlotMax = kImmediateBlockBufferSlot - 1; +inline uint32_t GetImmediateBufferSizesByteOffset(ImmediateMask pipelineImmediateMask) { + return static_cast<uint32_t>(pipelineImmediateMask.count()) * kImmediateElementByteSize; +} + inline constexpr BindGroupIndex kPullingBufferBindingSet = BindGroupIndex(kMaxBindGroups); class PipelineLayout final : public PipelineLayoutBase {
diff --git a/src/dawn/native/metal/ShaderModuleMTL.mm b/src/dawn/native/metal/ShaderModuleMTL.mm index b1235cf..89885fa 100644 --- a/src/dawn/native/metal/ShaderModuleMTL.mm +++ b/src/dawn/native/metal/ShaderModuleMTL.mm
@@ -33,7 +33,6 @@ #include "dawn/platform/DawnPlatform.h" #include "src/dawn/common/MatchVariant.h" -#include "src/dawn/common/Math.h" #include "src/dawn/common/Range.h" #include "src/dawn/native/Adapter.h" #include "src/dawn/native/BindGroupLayout.h" @@ -281,10 +280,8 @@ } if (!arrayLengthFromConstants.bindpoint_to_size_index.empty()) { - // Based on Immediate block layouts describes in PipelineLayoutMTL.h, it requires - // vec4<u32> array aligns to 16 bytes. arrayLengthFromConstants.buffer_sizes_offset = - RoundUp(pipelineImmediateMask.count() * kImmediateElementByteSize, 16); + GetImmediateBufferSizesByteOffset(pipelineImmediateMask); } // Type should match src/tint/lang/msl/writer/common/options.h
diff --git a/src/tint/lang/core/ir/transform/array_length_from.cc b/src/tint/lang/core/ir/transform/array_length_from.cc index 082895a..354fbb0 100644 --- a/src/tint/lang/core/ir/transform/array_length_from.cc +++ b/src/tint/lang/core/ir/transform/array_length_from.cc
@@ -34,6 +34,9 @@ #include "src/tint/lang/core/ir/module.h" #include "src/tint/lang/core/ir/transform/prepare_immediate_data.h" #include "src/tint/lang/core/ir/validator.h" +#include "src/tint/lang/core/type/array.h" +#include "src/tint/lang/core/type/pointer.h" +#include "src/tint/lang/core/type/struct.h" using namespace tint::core::fluent_types; // NOLINT using namespace tint::core::number_suffixes; // NOLINT @@ -458,25 +461,22 @@ TINT_IR_ASSERT(ir, bindpoint_to_size_index.contains(info.binding_point)); TINT_IR_ASSERT(ir, bindpoint_to_length_member_index.Contains(info.binding_point)); - // Load the total storage buffer size from the uniform buffer. - // The sizes are packed into vec4s to satisfy the 16-byte alignment requirement for - // array elements in uniform buffers, so we have to find the vector and element that - // correspond to the index that we want. + // Uniform data packs sizes into vec4s; immediate data uses tightly packed u32s. const uint32_t size_index = bindpoint_to_size_index.at(info.binding_point); - const uint32_t array_index = size_index / 4; - const uint32_t vec_index = size_index % 4; Value* total_buffer_size = nullptr; if (from_uniform) { + const uint32_t array_index = size_index / 4; + const uint32_t vec_index = size_index % 4; auto* vec_ptr = b.Access<ptr<uniform, vec4u>>(BufferSizes(), u32(array_index)); total_buffer_size = b.LoadVectorElement(vec_ptr, u32(vec_index))->Result(); } else { auto* buffer_sizes = b.Access( - ty.ptr(immediate, ty.array(ty.vec4u(), buffer_sizes_array_elements_num)), + ty.ptr(immediate, ty.array(ty.u32(), buffer_sizes_array_elements_num)), immediate_data_layout.var, u32(immediate_data_layout.IndexOf(buffer_sizes_offset))); - auto* vec_ptr = b.Access(ty.ptr(immediate, ty.vec4u()), buffer_sizes->Result(), - u32(array_index)); - total_buffer_size = b.LoadVectorElement(vec_ptr, u32(vec_index))->Result(); + auto* size_ptr = b.Access(ty.ptr(immediate, ty.u32()), buffer_sizes->Result(), + u32(size_index)); + total_buffer_size = b.Load(size_ptr)->Result(); } // Calculate actual array length:
diff --git a/src/tint/lang/core/ir/transform/array_length_from.h b/src/tint/lang/core/ir/transform/array_length_from.h index bcb27df..d43948f 100644 --- a/src/tint/lang/core/ir/transform/array_length_from.h +++ b/src/tint/lang/core/ir/transform/array_length_from.h
@@ -78,7 +78,7 @@ /// @group(0) @binding(30) /// struct tint_immediate_data_struct { /// ... -/// buffer_sizes: array<vec4<u32>, 8>; // offset is provided via config +/// buffer_sizes: array<u32, 8>; // offset is provided via config // }; /// var<immediate> tint_immediate_data : tint_immediate_data_struct; /// ``` @@ -91,8 +91,7 @@ /// @param bindpoint_to_size_index The map from binding point to an index which holds the size /// of that buffer. /// @param buffer_sizes_offset The offset in immediate block where buffer sizes start. -/// @param buffer_sizes_array_elements_num the number of vec4s used to store buffer sizes that will -/// be set into the immediate block. +/// @param buffer_sizes_array_elements_num number of u32 buffer-size elements in the immediate block /// @returns the transform result or failure /// TODO(crbug.com/366291600): Replace ArrayLengthFromUniform. Result<ArrayLengthResult> ArrayLengthFromImmediates(
diff --git a/src/tint/lang/core/ir/transform/array_length_from_immediate_test.cc b/src/tint/lang/core/ir/transform/array_length_from_immediate_test.cc index 50bae8c..c11a5e4 100644 --- a/src/tint/lang/core/ir/transform/array_length_from_immediate_test.cc +++ b/src/tint/lang/core/ir/transform/array_length_from_immediate_test.cc
@@ -48,7 +48,7 @@ for (auto& entry : bindpoint_to_size_index) { max_index = std::max(max_index, entry.second); } - return (max_index / 4) + 1; + return max_index + 1; } TEST_F(IR_ArrayLengthFromImmediatesTest, NoModify_UserFunction) { @@ -152,8 +152,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } tint_array_lengths_struct = struct @align(4) { @@ -167,9 +167,9 @@ %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B2: { - %4:ptr<immediate, array<vec4<u32>, 1>, read> = access %tint_immediate_data, 0u - %5:ptr<immediate, vec4<u32>, read> = access %4, 0u - %6:u32 = load_vector_element %5, 0u + %4:ptr<immediate, array<u32, 1>, read> = access %tint_immediate_data, 0u + %5:ptr<immediate, u32, read> = access %4, 0u + %6:u32 = load %5 %7:u32 = div %6, 4u %8:tint_array_lengths_struct = construct %7 %9:u32 = access %8, 0u @@ -187,7 +187,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -229,8 +229,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 2> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 8> @offset(16) } tint_array_lengths_struct = struct @align(4) { @@ -244,9 +244,9 @@ %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B2: { - %4:ptr<immediate, array<vec4<u32>, 2>, read> = access %tint_immediate_data, 0u - %5:ptr<immediate, vec4<u32>, read> = access %4, 1u - %6:u32 = load_vector_element %5, 3u + %4:ptr<immediate, array<u32, 8>, read> = access %tint_immediate_data, 0u + %5:ptr<immediate, u32, read> = access %4, 7u + %6:u32 = load %5 %7:u32 = div %6, 4u %8:tint_array_lengths_struct = construct %7 %9:u32 = access %8, 0u @@ -264,7 +264,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -427,8 +427,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } tint_array_lengths_struct = struct @align(4) { @@ -448,9 +448,9 @@ } %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B3: { - %7:ptr<immediate, array<vec4<u32>, 1>, read> = access %tint_immediate_data, 0u - %8:ptr<immediate, vec4<u32>, read> = access %7, 0u - %9:u32 = load_vector_element %8, 0u + %7:ptr<immediate, array<u32, 1>, read> = access %tint_immediate_data, 0u + %8:ptr<immediate, u32, read> = access %7, 0u + %9:u32 = load %8 %10:u32 = div %9, 4u %11:tint_array_lengths_struct = construct %10 %12:u32 = call %bar, %11 @@ -468,7 +468,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -523,8 +523,8 @@ a:array<i32> @offset(0) } -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } tint_array_lengths_struct = struct @align(4) { @@ -538,9 +538,9 @@ %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B2: { - %4:ptr<immediate, array<vec4<u32>, 1>, read> = access %tint_immediate_data, 0u - %5:ptr<immediate, vec4<u32>, read> = access %4, 0u - %6:u32 = load_vector_element %5, 0u + %4:ptr<immediate, array<u32, 1>, read> = access %tint_immediate_data, 0u + %5:ptr<immediate, u32, read> = access %4, 0u + %6:u32 = load %5 %7:u32 = sub %6, 0u %8:u32 = div %7, 4u %9:tint_array_lengths_struct = construct %8 @@ -560,7 +560,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -630,8 +630,8 @@ a:array<i32> @offset(20) } -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } tint_array_lengths_struct = struct @align(4) { @@ -645,9 +645,9 @@ %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B2: { - %4:ptr<immediate, array<vec4<u32>, 1>, read> = access %tint_immediate_data, 0u - %5:ptr<immediate, vec4<u32>, read> = access %4, 0u - %6:u32 = load_vector_element %5, 0u + %4:ptr<immediate, array<u32, 1>, read> = access %tint_immediate_data, 0u + %5:ptr<immediate, u32, read> = access %4, 0u + %6:u32 = load %5 %7:u32 = sub %6, 20u %8:u32 = div %7, 4u %9:tint_array_lengths_struct = construct %8 @@ -667,7 +667,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -710,8 +710,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } tint_array_lengths_struct = struct @align(4) { @@ -725,9 +725,9 @@ %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B2: { - %4:ptr<immediate, array<vec4<u32>, 1>, read> = access %tint_immediate_data, 0u - %5:ptr<immediate, vec4<u32>, read> = access %4, 0u - %6:u32 = load_vector_element %5, 0u + %4:ptr<immediate, array<u32, 1>, read> = access %tint_immediate_data, 0u + %5:ptr<immediate, u32, read> = access %4, 0u + %6:u32 = load %5 %7:u32 = div %6, 4u %8:tint_array_lengths_struct = construct %7 %let:ptr<storage, array<i32>, read_write> = let %buffer @@ -746,7 +746,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -802,8 +802,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } tint_array_lengths_struct = struct @align(4) { @@ -822,9 +822,9 @@ } %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B3: { - %7:ptr<immediate, array<vec4<u32>, 1>, read> = access %tint_immediate_data, 0u - %8:ptr<immediate, vec4<u32>, read> = access %7, 0u - %9:u32 = load_vector_element %8, 0u + %7:ptr<immediate, array<u32, 1>, read> = access %tint_immediate_data, 0u + %8:ptr<immediate, u32, read> = access %7, 0u + %9:u32 = load %8 %10:u32 = div %9, 4u %11:tint_array_lengths_struct = construct %10 %12:u32 = access %11, 0u @@ -843,7 +843,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -913,8 +913,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } tint_array_lengths_struct = struct @align(4) { @@ -939,9 +939,9 @@ } %foo_2 = @compute @workgroup_size(1u, 1u, 1u) func():void { # %foo_2: 'foo' $B4: { - %11:ptr<immediate, array<vec4<u32>, 1>, read> = access %tint_immediate_data, 0u - %12:ptr<immediate, vec4<u32>, read> = access %11, 0u - %13:u32 = load_vector_element %12, 0u + %11:ptr<immediate, array<u32, 1>, read> = access %tint_immediate_data, 0u + %12:ptr<immediate, u32, read> = access %11, 0u + %13:u32 = load %12 %14:u32 = div %13, 4u %15:tint_array_lengths_struct = construct %14 %16:u32 = access %15, 0u @@ -960,7 +960,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -1099,8 +1099,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } tint_array_lengths_struct = struct @align(4) { @@ -1121,9 +1121,9 @@ } %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B3: { - %9:ptr<immediate, array<vec4<u32>, 1>, read> = access %tint_immediate_data, 0u - %10:ptr<immediate, vec4<u32>, read> = access %9, 0u - %11:u32 = load_vector_element %10, 0u + %9:ptr<immediate, array<u32, 1>, read> = access %tint_immediate_data, 0u + %10:ptr<immediate, u32, read> = access %9, 0u + %11:u32 = load %10 %12:u32 = div %11, 4u %13:tint_array_lengths_struct = construct %12 %14:u32 = access %13, 0u @@ -1142,7 +1142,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -1206,8 +1206,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } tint_array_lengths_struct = struct @align(4) { @@ -1228,9 +1228,9 @@ } %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B3: { - %13:ptr<immediate, array<vec4<u32>, 1>, read> = access %tint_immediate_data, 0u - %14:ptr<immediate, vec4<u32>, read> = access %13, 0u - %15:u32 = load_vector_element %14, 0u + %13:ptr<immediate, array<u32, 1>, read> = access %tint_immediate_data, 0u + %14:ptr<immediate, u32, read> = access %13, 0u + %15:u32 = load %14 %16:u32 = div %15, 4u %17:tint_array_lengths_struct = construct %16 %18:u32 = access %17, 0u @@ -1251,7 +1251,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -1329,8 +1329,8 @@ a:array<i32> @offset(4) } -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } tint_array_lengths_struct = struct @align(4) { @@ -1351,9 +1351,9 @@ } %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B3: { - %9:ptr<immediate, array<vec4<u32>, 1>, read> = access %tint_immediate_data, 0u - %10:ptr<immediate, vec4<u32>, read> = access %9, 0u - %11:u32 = load_vector_element %10, 0u + %9:ptr<immediate, array<u32, 1>, read> = access %tint_immediate_data, 0u + %10:ptr<immediate, u32, read> = access %9, 0u + %11:u32 = load %10 %12:u32 = sub %11, 4u %13:u32 = div %12, 4u %14:tint_array_lengths_struct = construct %13 @@ -1375,7 +1375,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -1416,8 +1416,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } tint_array_lengths_struct = struct @align(4) { @@ -1431,9 +1431,9 @@ %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B2: { - %4:ptr<immediate, array<vec4<u32>, 1>, read> = access %tint_immediate_data, 0u - %5:ptr<immediate, vec4<u32>, read> = access %4, 0u - %6:u32 = load_vector_element %5, 0u + %4:ptr<immediate, array<u32, 1>, read> = access %tint_immediate_data, 0u + %5:ptr<immediate, u32, read> = access %4, 0u + %6:u32 = load %5 %7:u32 = div %6, 16u %8:tint_array_lengths_struct = construct %7 %9:u32 = access %8, 0u @@ -1451,7 +1451,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -1514,8 +1514,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 2> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 6> @offset(16) } tint_array_lengths_struct = struct @align(4) { @@ -1537,25 +1537,25 @@ %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B2: { - %8:ptr<immediate, array<vec4<u32>, 2>, read> = access %tint_immediate_data, 0u - %9:ptr<immediate, vec4<u32>, read> = access %8, 0u - %10:u32 = load_vector_element %9, 0u + %8:ptr<immediate, array<u32, 6>, read> = access %tint_immediate_data, 0u + %9:ptr<immediate, u32, read> = access %8, 0u + %10:u32 = load %9 %11:u32 = div %10, 4u - %12:ptr<immediate, array<vec4<u32>, 2>, read> = access %tint_immediate_data, 0u - %13:ptr<immediate, vec4<u32>, read> = access %12, 1u - %14:u32 = load_vector_element %13, 1u + %12:ptr<immediate, array<u32, 6>, read> = access %tint_immediate_data, 0u + %13:ptr<immediate, u32, read> = access %12, 5u + %14:u32 = load %13 %15:u32 = div %14, 4u - %16:ptr<immediate, array<vec4<u32>, 2>, read> = access %tint_immediate_data, 0u - %17:ptr<immediate, vec4<u32>, read> = access %16, 0u - %18:u32 = load_vector_element %17, 3u + %16:ptr<immediate, array<u32, 6>, read> = access %tint_immediate_data, 0u + %17:ptr<immediate, u32, read> = access %16, 3u + %18:u32 = load %17 %19:u32 = div %18, 4u - %20:ptr<immediate, array<vec4<u32>, 2>, read> = access %tint_immediate_data, 0u - %21:ptr<immediate, vec4<u32>, read> = access %20, 0u - %22:u32 = load_vector_element %21, 2u + %20:ptr<immediate, array<u32, 6>, read> = access %tint_immediate_data, 0u + %21:ptr<immediate, u32, read> = access %20, 2u + %22:u32 = load %21 %23:u32 = div %22, 4u - %24:ptr<immediate, array<vec4<u32>, 2>, read> = access %tint_immediate_data, 0u - %25:ptr<immediate, vec4<u32>, read> = access %24, 1u - %26:u32 = load_vector_element %25, 0u + %24:ptr<immediate, array<u32, 6>, read> = access %tint_immediate_data, 0u + %25:ptr<immediate, u32, read> = access %24, 4u + %26:u32 = load %25 %27:u32 = div %26, 4u %28:tint_array_lengths_struct = construct %11, %15, %19, %23, %27 %29:u32 = access %28, 0u @@ -1580,7 +1580,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -1619,8 +1619,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } $B1: { # root @@ -1644,7 +1644,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -1687,8 +1687,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } $B1: { # root @@ -1713,7 +1713,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -1770,8 +1770,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } $B1: { # root @@ -1801,7 +1801,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -1842,8 +1842,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } $B1: { # root @@ -1867,7 +1867,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -1935,8 +1935,8 @@ b:array<u32> @offset(4) } -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } tint_array_lengths_struct = struct @align(4) { @@ -1950,9 +1950,9 @@ %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B2: { - %4:ptr<immediate, array<vec4<u32>, 1>, read> = access %tint_immediate_data, 0u - %5:ptr<immediate, vec4<u32>, read> = access %4, 0u - %6:u32 = load_vector_element %5, 0u + %4:ptr<immediate, array<u32, 1>, read> = access %tint_immediate_data, 0u + %5:ptr<immediate, u32, read> = access %4, 0u + %6:u32 = load %5 %7:tint_array_lengths_struct = construct %6 %offset:u32 = let 16u %9:u32 = access %7, 0u @@ -1974,7 +1974,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -2020,8 +2020,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } $B1: { # root @@ -2046,7 +2046,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -2147,8 +2147,8 @@ length:u32 @offset(12) } -tint_immediate_data_struct = struct @align(16), @block { - tint_storage_buffer_sizes:array<vec4<u32>, 1> @offset(16) +tint_immediate_data_struct = struct @align(4), @block { + tint_storage_buffer_sizes:array<u32, 1> @offset(16) } tint_array_lengths_struct = struct @align(4) { @@ -2173,9 +2173,9 @@ } %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B3: { - %13:ptr<immediate, array<vec4<u32>, 1>, read> = access %tint_immediate_data, 0u - %14:ptr<immediate, vec4<u32>, read> = access %13, 0u - %15:u32 = load_vector_element %14, 0u + %13:ptr<immediate, array<u32, 1>, read> = access %tint_immediate_data, 0u + %14:ptr<immediate, u32, read> = access %13, 0u + %15:u32 = load %14 %16:tint_array_lengths_struct = construct %15 %17:u32 = access %16, 0u %18:u32 = sub %17, 0u @@ -2197,7 +2197,7 @@ uint32_t num_elements = GetBufferSizesNumElements(bindpoint_to_index); ASSERT_EQ(immediate_data_config.AddInternalImmediateData( buffer_size_start_offset, mod.symbols.New("tint_storage_buffer_sizes"), - ty.array(ty.vec4u(), num_elements)), + ty.array(ty.u32(), num_elements)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success);
diff --git a/src/tint/lang/hlsl/writer/raise/array_offset_from_immediate.cc b/src/tint/lang/hlsl/writer/raise/array_offset_from_immediate.cc index af7ce4b..12ff915 100644 --- a/src/tint/lang/hlsl/writer/raise/array_offset_from_immediate.cc +++ b/src/tint/lang/hlsl/writer/raise/array_offset_from_immediate.cc
@@ -57,7 +57,7 @@ /// The offset in immediate block for buffer offsets array. uint32_t buffer_offsets_offset = 0; - /// The total number of vec4s used to store buffer offsets provided in the immediate block. + /// The total number of u32 elements used to store buffer offsets in the immediate block. uint32_t buffer_offsets_array_elements_num = 0; /// The map from binding point to the element index which holds the offset into that buffer. @@ -73,10 +73,8 @@ void Process() { // Validate that buffer_offsets_array_elements_num is large enough for (const auto& [binding_point, offset_index] : bindpoint_to_offset_index) { - uint32_t vec4_index = offset_index / 4; - if (vec4_index >= buffer_offsets_array_elements_num) { + if (offset_index >= buffer_offsets_array_elements_num) { TINT_ICE() << "ArrayOffsetFromImmediates: offset_index " << offset_index - << " requires vec4 element " << vec4_index << " but buffer_offsets_array_elements_num is " << buffer_offsets_array_elements_num; } @@ -218,18 +216,12 @@ /// Loads the storage buffer dynamic offset from the immediate block. /// @returns the loaded dynamic offset value Value* LoadDynamicOffset(uint32_t offset_index) { - // Load the dynamic offset from the immediate block. - // The offsets are packed into vec4s to satisfy the 16-byte alignment requirement for - // array elements in immediate block, so we have to find the vector and element that - // correspond to the index that we want. - const uint32_t array_index = offset_index / 4; - const uint32_t vec_index = offset_index % 4; auto* buffer_offsets = b.Access( - ty.ptr(immediate, ty.array(ty.vec4u(), buffer_offsets_array_elements_num)), + ty.ptr(immediate, ty.array(ty.u32(), buffer_offsets_array_elements_num)), immediate_data_layout.var, u32(immediate_data_layout.IndexOf(buffer_offsets_offset))); - auto* vec_ptr = - b.Access(ty.ptr(immediate, ty.vec4u()), buffer_offsets->Result(), u32(array_index)); - return b.LoadVectorElement(vec_ptr, u32(vec_index))->Result(); + auto* offset_ptr = + b.Access(ty.ptr(immediate, ty.u32()), buffer_offsets->Result(), u32(offset_index)); + return b.Load(offset_ptr)->Result(); } };
diff --git a/src/tint/lang/hlsl/writer/raise/array_offset_from_immediate.h b/src/tint/lang/hlsl/writer/raise/array_offset_from_immediate.h index 0490624..88f39fc 100644 --- a/src/tint/lang/hlsl/writer/raise/array_offset_from_immediate.h +++ b/src/tint/lang/hlsl/writer/raise/array_offset_from_immediate.h
@@ -50,7 +50,7 @@ /// ``` /// struct tint_immediate_data_struct { /// ... -/// buffer_offsets: array<vec4<u32>, 8>; // offset is provided via config +/// buffer_offsets: array<u32, 8>; // offset is provided via config /// }; /// var<immediate> tint_immediate_data : tint_immediate_data_struct; /// ``` @@ -61,8 +61,8 @@ /// @param module the module to transform /// @param immediate_data_layout The immediate data layout information. /// @param buffer_offsets_offset The offset in immediate block where buffer offsets start. -/// @param buffer_offsets_array_elements_num the number of vec4s used to store buffer offsets that -/// will be set into the immediate block. +/// @param buffer_offsets_array_elements_num number of u32 buffer-offset elements in the immediate +/// block /// @param bindpoint_to_offset_index The map from binding point to an index which holds the offset /// of that buffer. /// @returns the transform result or failure
diff --git a/src/tint/lang/hlsl/writer/raise/array_offset_from_immediate_test.cc b/src/tint/lang/hlsl/writer/raise/array_offset_from_immediate_test.cc index fb062c3..35fc010 100644 --- a/src/tint/lang/hlsl/writer/raise/array_offset_from_immediate_test.cc +++ b/src/tint/lang/hlsl/writer/raise/array_offset_from_immediate_test.cc
@@ -67,7 +67,7 @@ core::ir::transform::PrepareImmediateDataConfig immediate_data_config; ASSERT_EQ(immediate_data_config.AddInternalImmediateData(0, mod.symbols.New("buffer_offsets"), - ty.array(ty.vec4u(), 6)), + ty.array(ty.u32(), 3)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); ASSERT_EQ(immediate_data, Success); @@ -79,13 +79,13 @@ bindpoint_to_offset_index[{0, 0}] = 2; auto result = - ArrayOffsetFromImmediates(mod, immediate_data.Get(), 0, 6, bindpoint_to_offset_index); + ArrayOffsetFromImmediates(mod, immediate_data.Get(), 0, 3, bindpoint_to_offset_index); ASSERT_EQ(result, Success); EXPECT_EQ(str(), R"( -tint_immediate_data_struct = struct @align(16), @block { - buffer_offsets:array<vec4<u32>, 6> @offset(0) +tint_immediate_data_struct = struct @align(4), @block { + buffer_offsets:array<u32, 3> @offset(0) } $B1: { # root @@ -95,9 +95,9 @@ %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B2: { - %4:ptr<immediate, array<vec4<u32>, 6>, read> = access %tint_immediate_data, 0u - %5:ptr<immediate, vec4<u32>, read> = access %4, 0u - %6:u32 = load_vector_element %5, 2u + %4:ptr<immediate, array<u32, 3>, read> = access %tint_immediate_data, 0u + %5:ptr<immediate, u32, read> = access %4, 2u + %6:u32 = load %5 %7:u32 = add 42u, %6 %8:u32 = %buffer.Load %7 ret @@ -137,8 +137,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - buffer_offsets:array<vec4<u32>, 1> @offset(0) +tint_immediate_data_struct = struct @align(4), @block { + buffer_offsets:array<u32, 1> @offset(0) } $B1: { # root @@ -156,7 +156,7 @@ core::ir::transform::PrepareImmediateDataConfig immediate_data_config; ASSERT_EQ(immediate_data_config.AddInternalImmediateData(0, mod.symbols.New("buffer_offsets"), - ty.array(ty.vec4u(), 1)), + ty.array(ty.u32(), 1)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -198,8 +198,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - buffer_offsets:array<vec4<u32>, 6> @offset(0) +tint_immediate_data_struct = struct @align(4), @block { + buffer_offsets:array<u32, 21> @offset(0) } $B1: { # root @@ -217,14 +217,14 @@ core::ir::transform::PrepareImmediateDataConfig immediate_data_config; ASSERT_EQ(immediate_data_config.AddInternalImmediateData(0, mod.symbols.New("buffer_offsets"), - ty.array(ty.vec4u(), 6)), + ty.array(ty.u32(), 21)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); std::unordered_map<BindingPoint, uint32_t> bindpoint_to_offset_index; bindpoint_to_offset_index[{0, 1}] = 20; // Doesn't match binding point - Run(ArrayOffsetFromImmediates, immediate_data.Get(), 0u, 6u, bindpoint_to_offset_index); + Run(ArrayOffsetFromImmediates, immediate_data.Get(), 0u, 21u, bindpoint_to_offset_index); EXPECT_EQ(expect, str()); } @@ -261,8 +261,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - buffer_offsets:array<vec4<u32>, 6> @offset(0) +tint_immediate_data_struct = struct @align(4), @block { + buffer_offsets:array<u32, 21> @offset(0) } $B1: { # root @@ -281,14 +281,14 @@ core::ir::transform::PrepareImmediateDataConfig immediate_data_config; ASSERT_EQ(immediate_data_config.AddInternalImmediateData(0, mod.symbols.New("buffer_offsets"), - ty.array(ty.vec4u(), 6)), + ty.array(ty.u32(), 21)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); std::unordered_map<BindingPoint, uint32_t> bindpoint_to_offset_index; bindpoint_to_offset_index[{0, 0}] = 20; - Run(ArrayOffsetFromImmediates, immediate_data.Get(), 0u, 6u, bindpoint_to_offset_index); + Run(ArrayOffsetFromImmediates, immediate_data.Get(), 0u, 21u, bindpoint_to_offset_index); EXPECT_EQ(expect, str()); } @@ -352,8 +352,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - buffer_offsets:array<vec4<u32>, 15> @offset(0) +tint_immediate_data_struct = struct @align(4), @block { + buffer_offsets:array<u32, 58> @offset(0) } $B1: { # root @@ -364,44 +364,44 @@ %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B2: { - %5:ptr<immediate, array<vec4<u32>, 15>, read> = access %tint_immediate_data, 0u - %6:ptr<immediate, vec4<u32>, read> = access %5, 6u - %7:u32 = load_vector_element %6, 2u + %5:ptr<immediate, array<u32, 58>, read> = access %tint_immediate_data, 0u + %6:ptr<immediate, u32, read> = access %5, 26u + %7:u32 = load %6 %8:u32 = add 42u, %7 %9:u32 = %buffer.Load %8 - %10:ptr<immediate, array<vec4<u32>, 15>, read> = access %tint_immediate_data, 0u - %11:ptr<immediate, vec4<u32>, read> = access %10, 6u - %12:u32 = load_vector_element %11, 2u + %10:ptr<immediate, array<u32, 58>, read> = access %tint_immediate_data, 0u + %11:ptr<immediate, u32, read> = access %10, 26u + %12:u32 = load %11 %13:u32 = add 43u, %12 %14:vec2<u32> = %buffer.Load2 %13 - %15:ptr<immediate, array<vec4<u32>, 15>, read> = access %tint_immediate_data, 0u - %16:ptr<immediate, vec4<u32>, read> = access %15, 6u - %17:u32 = load_vector_element %16, 2u + %15:ptr<immediate, array<u32, 58>, read> = access %tint_immediate_data, 0u + %16:ptr<immediate, u32, read> = access %15, 26u + %17:u32 = load %16 %18:u32 = add 44u, %17 %19:vec3<u32> = %buffer.Load3 %18 - %20:ptr<immediate, array<vec4<u32>, 15>, read> = access %tint_immediate_data, 0u - %21:ptr<immediate, vec4<u32>, read> = access %20, 6u - %22:u32 = load_vector_element %21, 2u + %20:ptr<immediate, array<u32, 58>, read> = access %tint_immediate_data, 0u + %21:ptr<immediate, u32, read> = access %20, 26u + %22:u32 = load %21 %23:u32 = add 45u, %22 %24:vec4<u32> = %buffer.Load4 %23 - %25:ptr<immediate, array<vec4<u32>, 15>, read> = access %tint_immediate_data, 0u - %26:ptr<immediate, vec4<u32>, read> = access %25, 14u - %27:u32 = load_vector_element %26, 1u + %25:ptr<immediate, array<u32, 58>, read> = access %tint_immediate_data, 0u + %26:ptr<immediate, u32, read> = access %25, 57u + %27:u32 = load %26 %28:u32 = add 46u, %27 %29:void = %buffer_1.Store %28, 123u - %30:ptr<immediate, array<vec4<u32>, 15>, read> = access %tint_immediate_data, 0u - %31:ptr<immediate, vec4<u32>, read> = access %30, 14u - %32:u32 = load_vector_element %31, 1u + %30:ptr<immediate, array<u32, 58>, read> = access %tint_immediate_data, 0u + %31:ptr<immediate, u32, read> = access %30, 57u + %32:u32 = load %31 %33:u32 = add 47u, %32 %34:void = %buffer_1.Store2 %33, vec2<u32>(123u, 124u) - %35:ptr<immediate, array<vec4<u32>, 15>, read> = access %tint_immediate_data, 0u - %36:ptr<immediate, vec4<u32>, read> = access %35, 14u - %37:u32 = load_vector_element %36, 1u + %35:ptr<immediate, array<u32, 58>, read> = access %tint_immediate_data, 0u + %36:ptr<immediate, u32, read> = access %35, 57u + %37:u32 = load %36 %38:u32 = add 48u, %37 %39:void = %buffer_1.Store3 %38, vec3<u32>(123u, 124u, 125u) - %40:ptr<immediate, array<vec4<u32>, 15>, read> = access %tint_immediate_data, 0u - %41:ptr<immediate, vec4<u32>, read> = access %40, 14u - %42:u32 = load_vector_element %41, 1u + %40:ptr<immediate, array<u32, 58>, read> = access %tint_immediate_data, 0u + %41:ptr<immediate, u32, read> = access %40, 57u + %42:u32 = load %41 %43:u32 = add 49u, %42 %44:void = %buffer_1.Store4 %43, vec4<u32>(123u, 124u, 125u, 126u) ret @@ -411,7 +411,7 @@ core::ir::transform::PrepareImmediateDataConfig immediate_data_config; ASSERT_EQ(immediate_data_config.AddInternalImmediateData(0, mod.symbols.New("buffer_offsets"), - ty.array(ty.vec4u(), 15)), + ty.array(ty.u32(), 58)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); @@ -419,7 +419,7 @@ std::unordered_map<BindingPoint, uint32_t> bindpoint_to_offset_index; bindpoint_to_offset_index[{5, 6}] = 26; bindpoint_to_offset_index[{7, 8}] = 57; - Run(ArrayOffsetFromImmediates, immediate_data.Get(), 0u, 15u, bindpoint_to_offset_index); + Run(ArrayOffsetFromImmediates, immediate_data.Get(), 0u, 58u, bindpoint_to_offset_index); EXPECT_EQ(expect, str()); } @@ -460,14 +460,14 @@ core::ir::transform::PrepareImmediateDataConfig immediate_data_config; ASSERT_EQ(immediate_data_config.AddInternalImmediateData(0, mod.symbols.New("buffer_offsets"), - ty.array(ty.vec4u(), 3)), + ty.array(ty.u32(), 8)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); std::unordered_map<BindingPoint, uint32_t> bindpoint_to_offset_index; bindpoint_to_offset_index[{0, 0}] = 7; - Run(ArrayOffsetFromImmediates, immediate_data.Get(), 0u, 3u, bindpoint_to_offset_index); + Run(ArrayOffsetFromImmediates, immediate_data.Get(), 0u, 8u, bindpoint_to_offset_index); // Just verify it doesn't crash - detailed output checking would be very long EXPECT_NE(str(), ""); @@ -518,8 +518,8 @@ EXPECT_EQ(src, str()); auto* expect = R"( -tint_immediate_data_struct = struct @align(16), @block { - buffer_offsets:array<vec4<u32>, 3> @offset(0) +tint_immediate_data_struct = struct @align(4), @block { + buffer_offsets:array<u32, 10> @offset(0) } $B1: { # root @@ -531,19 +531,19 @@ %foo = @compute @workgroup_size(1u, 1u, 1u) func():void { $B2: { - %6:ptr<immediate, array<vec4<u32>, 3>, read> = access %tint_immediate_data, 0u - %7:ptr<immediate, vec4<u32>, read> = access %6, 0u - %8:u32 = load_vector_element %7, 1u + %6:ptr<immediate, array<u32, 10>, read> = access %tint_immediate_data, 0u + %7:ptr<immediate, u32, read> = access %6, 1u + %8:u32 = load %7 %9:u32 = add 0u, %8 %10:u32 = %buffer0.Load %9 - %11:ptr<immediate, array<vec4<u32>, 3>, read> = access %tint_immediate_data, 0u - %12:ptr<immediate, vec4<u32>, read> = access %11, 1u - %13:u32 = load_vector_element %12, 1u + %11:ptr<immediate, array<u32, 10>, read> = access %tint_immediate_data, 0u + %12:ptr<immediate, u32, read> = access %11, 5u + %13:u32 = load %12 %14:u32 = add 0u, %13 %15:u32 = %buffer1.Load %14 - %16:ptr<immediate, array<vec4<u32>, 3>, read> = access %tint_immediate_data, 0u - %17:ptr<immediate, vec4<u32>, read> = access %16, 2u - %18:u32 = load_vector_element %17, 1u + %16:ptr<immediate, array<u32, 10>, read> = access %tint_immediate_data, 0u + %17:ptr<immediate, u32, read> = access %16, 9u + %18:u32 = load %17 %19:u32 = add 0u, %18 %20:u32 = %buffer2.Load %19 ret @@ -553,16 +553,16 @@ core::ir::transform::PrepareImmediateDataConfig immediate_data_config; ASSERT_EQ(immediate_data_config.AddInternalImmediateData(0, mod.symbols.New("buffer_offsets"), - ty.array(ty.vec4u(), 3)), + ty.array(ty.u32(), 10)), Success); auto immediate_data = PrepareImmediateData(mod, immediate_data_config); EXPECT_EQ(immediate_data, Success); std::unordered_map<BindingPoint, uint32_t> bindpoint_to_offset_index; - bindpoint_to_offset_index[{0, 0}] = 1; // vec4[0].y - bindpoint_to_offset_index[{0, 1}] = 5; // vec4[1].y - bindpoint_to_offset_index[{0, 2}] = 9; // vec4[2].y - Run(ArrayOffsetFromImmediates, immediate_data.Get(), 0u, 3u, bindpoint_to_offset_index); + bindpoint_to_offset_index[{0, 0}] = 1; + bindpoint_to_offset_index[{0, 1}] = 5; + bindpoint_to_offset_index[{0, 2}] = 9; + Run(ArrayOffsetFromImmediates, immediate_data.Get(), 0u, 10u, bindpoint_to_offset_index); EXPECT_EQ(expect, str()); }
diff --git a/src/tint/lang/hlsl/writer/raise/raise.cc b/src/tint/lang/hlsl/writer/raise/raise.cc index f6f27ea..fd98408 100644 --- a/src/tint/lang/hlsl/writer/raise/raise.cc +++ b/src/tint/lang/hlsl/writer/raise/raise.cc
@@ -136,39 +136,29 @@ } if (array_length_from_uniform_options.buffer_sizes_offset) { - // Find the largest index declared in the map, in order to determine the number of - // elements needed in the array of buffer sizes. The buffer sizes will be packed into - // vec4s to satisfy the 16-byte alignment requirement for array elements in constant - // buffers. uint32_t max_index = 0; for (const auto& entry : array_length_from_uniform_options.bindpoint_to_size_index) { max_index = std::max(max_index, entry.second); } - buffer_sizes_array_elements_num = (max_index / 4) + 1; + buffer_sizes_array_elements_num = max_index + 1; TINT_CHECK_RESULT(immediate_data_config.AddInternalImmediateData( array_length_from_uniform_options.buffer_sizes_offset.value(), module.symbols.New("buffer_sizes"), - module.Types().array(module.Types().vec4<core::u32>(), - buffer_sizes_array_elements_num))); + module.Types().array(module.Types().u32(), buffer_sizes_array_elements_num))); } if (array_offset_from_uniform_options.buffer_offsets_offset) { - // Find the largest index declared in the map, in order to determine the number of - // elements needed in the array of buffer offsets. The buffer offsets will be packed into - // vec4s to satisfy the 16-byte alignment requirement for array elements in constant - // buffers. uint32_t max_index = 0; for (const auto& entry : array_offset_from_uniform_options.bindpoint_to_offset_index) { max_index = std::max(max_index, entry.second); } - buffer_offsets_array_elements_num = (max_index / 4) + 1; + buffer_offsets_array_elements_num = max_index + 1; TINT_CHECK_RESULT(immediate_data_config.AddInternalImmediateData( array_offset_from_uniform_options.buffer_offsets_offset.value(), module.symbols.New("buffer_offsets"), - module.Types().array(module.Types().vec4<core::u32>(), - buffer_offsets_array_elements_num))); + module.Types().array(module.Types().u32(), buffer_offsets_array_elements_num))); } TINT_CHECK_RESULT_UNWRAP(immediate_data_layout, core::ir::transform::PrepareImmediateData(
diff --git a/src/tint/lang/msl/writer/raise/raise.cc b/src/tint/lang/msl/writer/raise/raise.cc index 71400bc..56b679a 100644 --- a/src/tint/lang/msl/writer/raise/raise.cc +++ b/src/tint/lang/msl/writer/raise/raise.cc
@@ -105,27 +105,21 @@ PopulateBindingRelatedOptions(options, remapper_data, multiplanar_map, array_length_from_constants); - // The number of vec4s used to store buffer sizes that will be set into the immediate block. uint32_t buffer_sizes_array_elements_num = 0; // PrepareImmediateData must come before any transform that needs internal immediates. core::ir::transform::PrepareImmediateDataConfig immediate_data_config; if (array_length_from_constants.buffer_sizes_offset) { - // Find the largest index declared in the map, in order to determine the number of - // elements needed in the array of buffer sizes. The buffer sizes will be packed into - // vec4s to satisfy the 16-byte alignment requirement for array elements in uniform - // buffers. uint32_t max_index = 0; for (auto& entry : array_length_from_constants.bindpoint_to_size_index) { max_index = std::max(max_index, entry.second); } - buffer_sizes_array_elements_num = (max_index / 4) + 1; + buffer_sizes_array_elements_num = max_index + 1; TINT_CHECK_RESULT(immediate_data_config.AddInternalImmediateData( array_length_from_constants.buffer_sizes_offset.value(), module.symbols.New("tint_storage_buffer_sizes"), - module.Types().array(module.Types().vec4<core::u32>(), - buffer_sizes_array_elements_num))); + module.Types().array(module.Types().u32(), buffer_sizes_array_elements_num))); } if (options.depth_range_offsets) { TINT_CHECK_RESULT(immediate_data_config.AddInternalImmediateData(
diff --git a/src/tint/lang/msl/writer/writer_test.cc b/src/tint/lang/msl/writer/writer_test.cc index 03626fd..e253505 100644 --- a/src/tint/lang/msl/writer/writer_test.cc +++ b/src/tint/lang/msl/writer/writer_test.cc
@@ -204,7 +204,7 @@ }; struct tint_immediate_data_struct { - tint_array<uint4, 1> tint_storage_buffer_sizes; + tint_array<uint, 1> tint_storage_buffer_sizes; }; struct tint_module_vars_struct { @@ -257,7 +257,7 @@ struct tint_immediate_data_struct { /* 0x0000 */ tint_array<int8_t, 64> tint_pad; - /* 0x0040 */ tint_array<uint4, 1> tint_storage_buffer_sizes; + /* 0x0040 */ tint_array<uint, 1> tint_storage_buffer_sizes; }; struct tint_module_vars_struct { @@ -272,7 +272,7 @@ [[max_total_threads_per_threadgroup(1)]] kernel void entry(device tint_array<uint, 1>* a [[buffer(0)]], const constant tint_immediate_data_struct* tint_immediate_data [[buffer(30)]]) { tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.a=a, .tint_immediate_data=tint_immediate_data}; - (*tint_module_vars.a)[0u] = tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_immediate_data).tint_storage_buffer_sizes[0u].x / 4u)}.tint_array_length_0_0; + (*tint_module_vars.a)[0u] = tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_immediate_data).tint_storage_buffer_sizes[0u] / 4u)}.tint_array_length_0_0; } )"); EXPECT_TRUE(output_.needs_storage_buffer_sizes); @@ -482,7 +482,7 @@ struct tint_immediate_data_struct { /* 0x0000 */ tint_array<int8_t, 64> tint_pad; - /* 0x0040 */ tint_array<uint4, 1> tint_storage_buffer_sizes; + /* 0x0040 */ tint_array<uint, 1> tint_storage_buffer_sizes; }; struct tint_module_vars_struct { @@ -499,7 +499,7 @@ }; float4 entry_inner(uint tint_vertex_index, tint_module_vars_struct tint_module_vars) { - return float4(as_type<float>((*tint_module_vars.tint_vertex_buffer_0)[min(tint_vertex_index, (tint_array_lengths_struct{.tint_array_length_0_1=((*tint_module_vars.tint_immediate_data).tint_storage_buffer_sizes[0u].x / 4u)}.tint_array_length_0_1 - 1u))]), 0.0f, 0.0f, 1.0f); + return float4(as_type<float>((*tint_module_vars.tint_vertex_buffer_0)[min(tint_vertex_index, (tint_array_lengths_struct{.tint_array_length_0_1=((*tint_module_vars.tint_immediate_data).tint_storage_buffer_sizes[0u] / 4u)}.tint_array_length_0_1 - 1u))]), 0.0f, 0.0f, 1.0f); } vertex entry_outputs entry(uint tint_vertex_index [[vertex_id]], const device tint_array<uint, 1>* tint_vertex_buffer_0 [[buffer(1)]], const constant tint_immediate_data_struct* tint_immediate_data [[buffer(30)]]) {