Skip to content

Commit c9d7d5d

Browse files
committed
Update on "[ET-VK] Add PushConstantDataInfo and vector to hold push constants data in DispatchNode."
This diff adds a new class called `PushConstantDataInfo` to the `DispatchNode` class in the Vulkan backend for Executorch. This class represents a push constant data entry, which can either be a shared pointer to a tensor's uniform data with an attribute or data with a maximum size of 16 bytes. The `write` method is also added to this class, which writes the data to a destination buffer. Differential Revision: [D66796049](https://our.internmc.facebook.com/intern/diff/D66796049/) [ghstack-poisoned]
2 parents 8a0cb03 + 4b7dd07 commit c9d7d5d

File tree

2 files changed

+7
-4
lines changed

2 files changed

+7
-4
lines changed

backends/vulkan/runtime/graph/ops/DispatchNode.cpp

Lines changed: 6 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -75,19 +75,21 @@ void DispatchNode::encode(ComputeGraph* graph) {
7575

7676
bind_params_to_descriptor_set(params_, descriptor_set, idx);
7777

78-
uint8_t push_constants_data[128];
78+
std::array<uint8_t, kMaxPushConstantSize> push_constants_data;
7979
uint32_t push_constants_offset = 0;
8080

8181
for (const auto& push_constant : push_constants_) {
82-
push_constants_offset +=
83-
push_constant.write(push_constants_data, push_constants_offset, 128);
82+
push_constants_offset += push_constant.write(
83+
push_constants_data.data(),
84+
push_constants_offset,
85+
kMaxPushConstantSize);
8486
}
8587
context->register_shader_dispatch(
8688
descriptor_set,
8789
pipeline_barrier,
8890
shader_,
8991
global_workgroup_size_,
90-
push_constants_data,
92+
push_constants_data.data(),
9193
push_constants_offset);
9294

9395
context->report_shader_dispatch_end();

backends/vulkan/runtime/graph/ops/DispatchNode.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -18,6 +18,7 @@ namespace vkcompute {
1818

1919
class ComputeGraph;
2020

21+
constexpr uint32_t kMaxPushConstantSize = 128;
2122
/*
2223
* Represents a push constant data entry
2324
* Which is either shared pointer to a tensor's uniform data with an attribute

0 commit comments

Comments
 (0)