Skip to content

Commit c0f8d11

Browse files
committed
Update on "[ET-VK] Replace Uniform buffers with push constants for view op"
This diff replaces uniform buffers with push constants for view op in the Vulkan backend of Executorch. The changes include updating the GLSL code to use push constants instead of uniform buffers and updating the C++ code to pass the sizes as push constants to the shader. Differential Revision: [D66733658](https://our.internmc.facebook.com/intern/diff/D66733658/) [ghstack-poisoned]
2 parents 2dab102 + 125354d commit c0f8d11

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)