diff --git a/backends/vulkan/runtime/graph/ops/DispatchNode.cpp b/backends/vulkan/runtime/graph/ops/DispatchNode.cpp index 6730d851483..51ff0c122b0 100644 --- a/backends/vulkan/runtime/graph/ops/DispatchNode.cpp +++ b/backends/vulkan/runtime/graph/ops/DispatchNode.cpp @@ -21,10 +21,10 @@ DispatchNode::DispatchNode( const utils::uvec3& local_workgroup_size, const std::vector& args, const vkapi::ParamsBindList& params, + const std::vector& push_constants, const vkapi::SpecVarList& spec_vars, - const ResizeFunction& resize_fn, const std::vector& resize_args, - const std::vector& push_constants) + const ResizeFunction& resize_fn) : ExecuteNode(resize_fn, resize_args, args, shader.kernel_name), shader_(shader), global_workgroup_size_(global_workgroup_size), diff --git a/backends/vulkan/runtime/graph/ops/DispatchNode.h b/backends/vulkan/runtime/graph/ops/DispatchNode.h index e3794e9a9e4..172ab49a98a 100644 --- a/backends/vulkan/runtime/graph/ops/DispatchNode.h +++ b/backends/vulkan/runtime/graph/ops/DispatchNode.h @@ -33,10 +33,10 @@ class DispatchNode final : public ExecuteNode { const utils::uvec3& local_workgroup_size, const std::vector& args, const vkapi::ParamsBindList& params, + const std::vector& push_constants = {}, const vkapi::SpecVarList& spec_vars = {}, - const ResizeFunction& resize_fn = nullptr, const std::vector& resize_args = {}, - const std::vector& push_constants = {}); + const ResizeFunction& resize_fn = nullptr); ~DispatchNode() override = default; diff --git a/backends/vulkan/runtime/graph/ops/impl/Arange.cpp b/backends/vulkan/runtime/graph/ops/impl/Arange.cpp index db519563614..490def4860a 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Arange.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Arange.cpp @@ -94,16 +94,19 @@ void add_arange_node( graph.create_global_wg_size(out), graph.create_local_wg_size(out), // Inputs and Outputs - {{out, vkapi::MemoryAccessType::WRITE}}, + {{out, vkapi::kWrite}}, // Shader params buffers {t_out->sizes_ubo(), graph.create_params_buffer(start_val), graph.create_params_buffer(step_val)}, + // Push Constants + {}, // Specialization Constants {}, + // Resize Args + {start, end, step}, // Resizing Logic - resize_arange_node, - {start, end, step})); + resize_arange_node)); } void arange(ComputeGraph& graph, const std::vector& args) { diff --git a/backends/vulkan/runtime/graph/ops/impl/BatchNorm.cpp b/backends/vulkan/runtime/graph/ops/impl/BatchNorm.cpp index ec7b6c2fc12..81cbd62d90c 100644 --- a/backends/vulkan/runtime/graph/ops/impl/BatchNorm.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/BatchNorm.cpp @@ -90,12 +90,19 @@ void add_native_batch_norm_node( VK_KERNEL_FROM_STR(kernel_name), graph.create_global_wg_size(out_ref), graph.create_local_wg_size(out_ref), - {{out_ref, vkapi::MemoryAccessType::WRITE}, - {{in_ref, arg_weight, arg_bias, arg_mean, arg_var}, - vkapi::MemoryAccessType::READ}}, + {{out_ref, vkapi::kWrite}, + {{in_ref, arg_weight, arg_bias, arg_mean, arg_var}, vkapi::kRead}}, {t_out->logical_limits_ubo(), graph.create_params_buffer(epsilon), - graph.create_params_buffer(num_texel_per_batch)})); + graph.create_params_buffer(num_texel_per_batch)}, + // Push Constants + {}, + // Specialization Constants + {}, + // Resize Args + {}, + // Resizing Logic + nullptr)); } void native_batch_norm(ComputeGraph& graph, const std::vector& args) { diff --git a/backends/vulkan/runtime/graph/ops/impl/BinaryOp.cpp b/backends/vulkan/runtime/graph/ops/impl/BinaryOp.cpp index c3c686772e1..ff6b54c5289 100644 --- a/backends/vulkan/runtime/graph/ops/impl/BinaryOp.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/BinaryOp.cpp @@ -84,19 +84,20 @@ void add_binary_op_texture_node( graph.create_global_wg_size(out), graph.create_local_wg_size(out), // Inputs and Outputs - {{out, vkapi::MemoryAccessType::WRITE}, - {{arg1, arg2}, vkapi::MemoryAccessType::READ}}, + {{out, vkapi::kWrite}, {{arg1, arg2}, vkapi::kRead}}, // Shader params buffers {}, - // Specialization Constants - {t_out->hashed_layout(), t_in1->hashed_layout(), t_in2->hashed_layout()}, - // Resizing Logic - resize_binary_op_node, - {}, + // Push Constants {{graph.sizes_pc_of(out), graph.sizes_pc_of(arg1), graph.sizes_pc_of(arg2), - PushConstantDataInfo(&binary_ops_params, sizeof(binary_ops_params))}})); + PushConstantDataInfo(&binary_ops_params, sizeof(binary_ops_params))}}, + // Specialization Constants + {t_out->hashed_layout(), t_in1->hashed_layout(), t_in2->hashed_layout()}, + // Resize Args + {}, + // Resizing Logic + resize_binary_op_node)); } void add_binary_op_buffer_node( @@ -127,17 +128,10 @@ void add_binary_op_buffer_node( graph.create_global_wg_size(out), graph.create_local_wg_size(out), // Inputs and Outputs - {{out, vkapi::MemoryAccessType::WRITE}, - {{in1, in2}, vkapi::MemoryAccessType::READ}}, + {{out, vkapi::kWrite}, {{in1, in2}, vkapi::kRead}}, // Shader params buffers {}, - // Specialization Constants - {graph.packed_dim_of(out), - graph.packed_dim_of(in1), - graph.packed_dim_of(in2)}, - // Resizing Logic - resize_binary_op_node, - {}, + // Push Constants {{ graph.sizes_pc_of(in1), graph.sizes_pc_of(in2), @@ -146,7 +140,15 @@ void add_binary_op_buffer_node( graph.strides_pc_of(in2), graph.numel_pc_of(out), PushConstantDataInfo(&alpha_val, sizeof(float)), - }})); + }}, + // Specialization Constants + {graph.packed_dim_of(out), + graph.packed_dim_of(in1), + graph.packed_dim_of(in2)}, + // Resize Args + {}, + // Resizing Logic + resize_binary_op_node)); } void add_binary_op_node( diff --git a/backends/vulkan/runtime/graph/ops/impl/Clone.cpp b/backends/vulkan/runtime/graph/ops/impl/Clone.cpp index c763588043f..b547bc3572d 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Clone.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Clone.cpp @@ -50,8 +50,12 @@ void add_clone_node( {{out, vkapi::kWrite}, {in, vkapi::kRead}}, // Parameter Buffers {t_out->logical_limits_ubo()}, + // Push Constants + {}, // Specialization Constants {}, + // Resize Args + {}, // Resizing Logic resize_clone_node)); } @@ -74,8 +78,12 @@ void add_image_to_buffer_node( {{buffer, vkapi::kWrite}, {image, vkapi::kRead}}, // Parameter Buffers {graph.sizes_ubo(image), graph.strides_ubo(buffer)}, + // Push Constants + {}, // Specialization Constants {graph.hashed_layout_of(image)}, + // Resize Args + {}, // Resizing Logic resize_clone_node)); } @@ -98,8 +106,12 @@ void add_buffer_to_image_node( {{image, vkapi::kWrite}, {buffer, vkapi::kRead}}, // Parameter Buffers {graph.sizes_ubo(image), graph.strides_ubo(buffer)}, + // Push Constants + {}, // Specialization Constants {graph.hashed_layout_of(image)}, + // Resize Args + {}, // Resizing Logic resize_clone_node)); } diff --git a/backends/vulkan/runtime/graph/ops/impl/Convolution.cpp b/backends/vulkan/runtime/graph/ops/impl/Convolution.cpp index 060f5028c02..5d4537f9938 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Convolution.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Convolution.cpp @@ -444,16 +444,17 @@ void add_conv2d_node( wg_size, graph.create_local_wg_size(wg_size), // Inputs and Outputs - {{out, vkapi::MemoryAccessType::WRITE}, - {{in, arg_weight, arg_bias}, vkapi::MemoryAccessType::READ}}, + {{out, vkapi::kWrite}, {{in, arg_weight, arg_bias}, vkapi::kRead}}, // Shader params buffers param_buffers, + // Push Constants + push_constants, // Specialization Constants {}, - // Resizing Logic - resize_conv2d_node, + // Resize Args {weight_data, stride, padding, dilation, transposed, output_padding}, - push_constants)); + // Resizing Logic + resize_conv2d_node)); } void add_conv1d_node( @@ -541,8 +542,7 @@ void add_conv1d_node( global_size, local_size, // Inputs and Outputs - {{out, vkapi::MemoryAccessType::WRITE}, - {{in, arg_weight, arg_bias}, vkapi::MemoryAccessType::READ}}, + {{out, vkapi::kWrite}, {{in, arg_weight, arg_bias}, vkapi::kRead}}, // Shader params buffers { t_out->logical_limits_ubo(), @@ -550,14 +550,17 @@ void add_conv1d_node( graph.create_params_buffer(kernel_params), graph.create_params_buffer(out_params), }, + // Push Constants + {}, // Specialization Constants {t_out->hashed_layout(), t_in->hashed_layout(), t_weight->hashed_layout(), t_bias->hashed_layout()}, + // Resize Args + {weight, stride, padding, dilation}, // Resizing Logic - resize_conv1d_node, - {weight, stride, padding, dilation})); + resize_conv1d_node)); } void conv(ComputeGraph& graph, const std::vector& args) { diff --git a/backends/vulkan/runtime/graph/ops/impl/Copy.cpp b/backends/vulkan/runtime/graph/ops/impl/Copy.cpp index 80379880b0f..c4f37bd9386 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Copy.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Copy.cpp @@ -50,19 +50,22 @@ void add_copy_offset_node( }, // Parameter buffers {}, + // Push Constants + { + PushConstantDataInfo(&range, sizeof(range), sizeof(ivec4)), + PushConstantDataInfo(&src_offset, sizeof(src_offset), sizeof(ivec4)), + PushConstantDataInfo(&dst_offset, sizeof(dst_offset), sizeof(ivec4)), + }, // Specialization Constants {graph.hashed_layout_of(out), graph.hashed_layout_of(in), (calc_out_pos_using_src_chnl ? 1 : calc_in_pos_using_dst_chnl ? 2 : 0)}, - nullptr, + // Resize Args {}, - { - PushConstantDataInfo(&range, sizeof(range), sizeof(ivec4)), - PushConstantDataInfo(&src_offset, sizeof(src_offset), sizeof(ivec4)), - PushConstantDataInfo(&dst_offset, sizeof(dst_offset), sizeof(ivec4)), - })); + // Resizing Logic + nullptr)); } void add_copy_packed_dim_offset_node( @@ -138,22 +141,25 @@ void add_copy_packed_dim_offset_node( graph.create_local_wg_size(global_wg_size), // Inputs and Outputs { - {out, vkapi::MemoryAccessType::WRITE}, - {out, vkapi::MemoryAccessType::READ}, - {in, vkapi::MemoryAccessType::READ}, + {out, vkapi::kWrite}, + {out, vkapi::kRead}, + {in, vkapi::kRead}, }, // Parameter buffers {}, - // Specialization Constants - {graph.hashed_layout_of(out), graph.hashed_layout_of(in)}, - nullptr, - {}, + // Push Constants { PushConstantDataInfo( &final_range, sizeof(final_range), sizeof(ivec4)), PushConstantDataInfo(&src_offset, sizeof(src_offset), sizeof(ivec4)), PushConstantDataInfo(&dst_offset, sizeof(dst_offset), sizeof(ivec4)), - })); + }, + // Specialization Constants + {graph.hashed_layout_of(out), graph.hashed_layout_of(in)}, + // Resize Args + {}, + // Resizing Logic + nullptr)); } void add_copy_channel_offset_node( @@ -248,22 +254,24 @@ void add_copy_channel_offset_node( local_size, // Inputs and Outputs { - {out, vkapi::MemoryAccessType::WRITE}, - {out, vkapi::MemoryAccessType::READ}, - {in, vkapi::MemoryAccessType::READ}, + {out, vkapi::kWrite}, + {out, vkapi::kRead}, + {in, vkapi::kRead}, }, // Parameter buffers {}, - // Specialization Constants - {graph.hashed_layout_of(out), graph.hashed_layout_of(in)}, - nullptr, - {}, + // Push Constants {graph.sizes_pc_of(out), graph.sizes_pc_of(in), PushConstantDataInfo(&range_params, sizeof(range_params)), PushConstantDataInfo(&offset_params, sizeof(offset_params)), - PushConstantDataInfo( - &src_channel_offset, sizeof(src_channel_offset))})); + PushConstantDataInfo(&src_channel_offset, sizeof(src_channel_offset))}, + // Specialization Constants + {graph.hashed_layout_of(out), graph.hashed_layout_of(in)}, + // Resize Args + {}, + // Resizing Logic + nullptr)); } } diff --git a/backends/vulkan/runtime/graph/ops/impl/Embedding.cpp b/backends/vulkan/runtime/graph/ops/impl/Embedding.cpp index 8160908cc59..85c80e01c27 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Embedding.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Embedding.cpp @@ -58,9 +58,16 @@ void add_embedding_node( { t_out->sizes_ubo(), }, + // Push Constants + {}, + // Specialization Constants {t_out->hashed_layout(), t_in->hashed_layout(), - t_weight->hashed_layout()})); + t_weight->hashed_layout()}, + // Resize Args + {}, + // Resizing Logic + nullptr)); } void embedding(ComputeGraph& graph, const std::vector& args) { diff --git a/backends/vulkan/runtime/graph/ops/impl/Flip.cpp b/backends/vulkan/runtime/graph/ops/impl/Flip.cpp index 3d0ceedd0a3..04aac2484ac 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Flip.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Flip.cpp @@ -74,8 +74,12 @@ void add_flip_node( graph.sizes_ubo(out), graph.create_params_buffer(dim_bitmap), }, + // Push Constants + {}, // Specialization Constants {}, + // Resize Args + {}, // Resizing Logic resize_flip_node)); } diff --git a/backends/vulkan/runtime/graph/ops/impl/Full.cpp b/backends/vulkan/runtime/graph/ops/impl/Full.cpp index 377fdba53a2..3ed18445463 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Full.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Full.cpp @@ -50,14 +50,17 @@ void add_full_node( graph.create_global_wg_size(out), graph.create_local_wg_size(out), // Inputs and Outputs - {{out, vkapi::MemoryAccessType::WRITE}}, + {{out, vkapi::kWrite}}, // Shader params buffers {t_out->sizes_ubo(), graph.create_params_buffer(fill_value_val)}, + // Push Constants + {}, // Specialization Constants {SV(t_out->packed_dim())}, + // Resize Args + {size_or_in}, // Resizing Logic - resize_full_node, - {size_or_in})); + resize_full_node)); } void full(ComputeGraph& graph, const std::vector& args) { diff --git a/backends/vulkan/runtime/graph/ops/impl/GridPriors.cpp b/backends/vulkan/runtime/graph/ops/impl/GridPriors.cpp index d2c758de918..0624020c872 100644 --- a/backends/vulkan/runtime/graph/ops/impl/GridPriors.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/GridPriors.cpp @@ -55,7 +55,7 @@ void add_grid_priors_node( graph.create_local_wg_size(out), // Inputs and Outputs { - {out, vkapi::MemoryAccessType::WRITE}, + {out, vkapi::kWrite}, }, // Shader params buffers { @@ -63,10 +63,14 @@ void add_grid_priors_node( t_out->sizes_ubo(), graph.create_params_buffer(param), }, + // Push Constants + {}, // Specialization Constants {}, - resize_grid_priors_node, - {in})); + // Resize Args + {in}, + // Resizing Logic + resize_grid_priors_node)); } void grid_priors(ComputeGraph& graph, const std::vector& args) { diff --git a/backends/vulkan/runtime/graph/ops/impl/IndexSelect.cpp b/backends/vulkan/runtime/graph/ops/impl/IndexSelect.cpp index 1f56d3c45d3..8203829c50f 100644 --- a/backends/vulkan/runtime/graph/ops/impl/IndexSelect.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/IndexSelect.cpp @@ -46,9 +46,16 @@ void add_index_select_channel_node( VK_KERNEL_FROM_STR(kernel_name), graph.create_global_wg_size(out), graph.create_local_wg_size(out), - {{out, vkapi::MemoryAccessType::WRITE}, - {{in, idx}, vkapi::MemoryAccessType::READ}}, - {t_out->sizes_ubo(), t_in->sizes_ubo()})); + {{out, vkapi::kWrite}, {{in, idx}, vkapi::kRead}}, + {t_out->sizes_ubo(), t_in->sizes_ubo()}, + // Push Constants + {}, + // Specialization Constants + {}, + // Resize Args + {}, + // Resizing Logic + nullptr)); } struct IndexSelectParams final { @@ -95,9 +102,16 @@ void add_index_select_node( VK_KERNEL_FROM_STR(kernel_name), graph.create_global_wg_size(out), graph.create_local_wg_size(out), - {{out, vkapi::MemoryAccessType::WRITE}, - {{in, idx}, vkapi::MemoryAccessType::READ}}, - {t_out->sizes_ubo(), graph.create_params_buffer(params)})); + {{out, vkapi::kWrite}, {{in, idx}, vkapi::kRead}}, + {t_out->sizes_ubo(), graph.create_params_buffer(params)}, + // Push Constants + {}, + // Specialization Constants + {}, + // Resize Args + {}, + // Resizing Logic + nullptr)); } int64_t get_dim_idx(ComputeGraph& graph, ValueRef in, ValueRef dim_ref) { diff --git a/backends/vulkan/runtime/graph/ops/impl/Linear.cpp b/backends/vulkan/runtime/graph/ops/impl/Linear.cpp index ddcdb41ece8..86df735acbe 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Linear.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Linear.cpp @@ -113,8 +113,7 @@ void add_addmm_naive_texture_node( global_wg_size, graph.create_local_wg_size(global_wg_size), // Inputs and Outputs - {{out, vkapi::MemoryAccessType::WRITE}, - {{mat1, mat2, self}, vkapi::MemoryAccessType::READ}}, + {{out, vkapi::kWrite}, {{mat1, mat2, self}, vkapi::kRead}}, // Shader params buffers { graph.sizes_ubo(out), @@ -124,14 +123,17 @@ void add_addmm_naive_texture_node( graph.sizes_ubo(self), graph.create_params_buffer(params), }, + // Push Constants + {}, // Specialization Constants {graph.hashed_layout_of(out), graph.hashed_layout_of(mat1), graph.hashed_layout_of(mat2), graph.hashed_layout_of(self)}, + // Resize Args + {mat2_is_transposed}, // Resizing Logic - resize_addmm_node, - {mat2_is_transposed})); + resize_addmm_node)); } void add_addmm_naive_buffer_node( @@ -190,11 +192,14 @@ void add_addmm_naive_buffer_node( graph.numel_ubo(out), graph.create_params_buffer(params), }, + // Push Constants + {}, // Specialization Constants {mat2_is_transposed_val}, + // Resize Args + {mat2_is_transposed}, // Resizing Logic - resize_addmm_node, - {mat2_is_transposed})); + resize_addmm_node)); } void add_addmm_optimized_node( @@ -269,8 +274,8 @@ void add_addmm_optimized_node( global_size, local_size, // Inputs and Outputs - {{out, vkapi::MemoryAccessType::WRITE}, - {{mat1_W_packed, mat2_packed, self}, vkapi::MemoryAccessType::READ}}, + {{out, vkapi::kWrite}, + {{mat1_W_packed, mat2_packed, self}, vkapi::kRead}}, // Shader params buffers { graph.sizes_ubo(out), @@ -279,14 +284,17 @@ void add_addmm_optimized_node( graph.sizes_ubo(self), graph.create_params_buffer(params), }, + // Push Constants + {}, // Specialization Constants {graph.hashed_layout_of(out), graph.hashed_layout_of(mat1_W_packed), graph.hashed_layout_of(mat2_packed), graph.hashed_layout_of(self)}, + // Resize Args + {mat2_is_transposed}, // Resizing Logic - resize_addmm_node, - {mat2_is_transposed})); + resize_addmm_node)); } void add_addmm_node( diff --git a/backends/vulkan/runtime/graph/ops/impl/MatMul.cpp b/backends/vulkan/runtime/graph/ops/impl/MatMul.cpp index a515f454266..724f4630264 100644 --- a/backends/vulkan/runtime/graph/ops/impl/MatMul.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/MatMul.cpp @@ -88,8 +88,7 @@ void add_matmul_naive_buffer_node( global_size, graph.create_local_wg_size(global_size), // Inputs and Outputs - {{out, vkapi::MemoryAccessType::WRITE}, - {{mat1, mat2}, vkapi::MemoryAccessType::READ}}, + {{out, vkapi::kWrite}, {{mat1, mat2}, vkapi::kRead}}, // Shader params buffers { graph.sizes_ubo(out), @@ -100,11 +99,14 @@ void add_matmul_naive_buffer_node( graph.strides_ubo(mat2), graph.numel_ubo(out), }, + // Push Constants + {}, // Specialization Constants {mat2_is_transposed_val}, + // Resize Args + {mat2_is_transposed}, // Resizing Logic - resize_matmul_node, - {mat2_is_transposed})); + resize_matmul_node)); } void add_matmul_naive_texture3d_node( @@ -134,8 +136,7 @@ void add_matmul_naive_texture3d_node( global_wg_size, graph.create_local_wg_size(global_wg_size), // Inputs and Outputs - {{out, vkapi::MemoryAccessType::WRITE}, - {{mat1, mat2}, vkapi::MemoryAccessType::READ}}, + {{out, vkapi::kWrite}, {{mat1, mat2}, vkapi::kRead}}, // Shader params buffers { graph.sizes_ubo(out), @@ -143,13 +144,16 @@ void add_matmul_naive_texture3d_node( graph.sizes_ubo(mat1), graph.sizes_ubo(mat2), }, + // Push Constants + {}, // Specialization Constants {graph.hashed_layout_of(out), graph.hashed_layout_of(mat1), graph.hashed_layout_of(mat2)}, + // Resize Args + {mat2_is_transposed}, // Resizing Logic - resize_matmul_node, - {mat2_is_transposed})); + resize_matmul_node)); } void add_matmul_optimized_node( @@ -228,21 +232,23 @@ void add_matmul_optimized_node( global_size, local_size, // Inputs and Outputs - {{out, vkapi::MemoryAccessType::WRITE}, - {{mat1_W_packed, mat2_packed}, vkapi::MemoryAccessType::READ}}, + {{out, vkapi::kWrite}, {{mat1_W_packed, mat2_packed}, vkapi::kRead}}, // Shader params buffers { graph.sizes_ubo(out), graph.sizes_ubo(mat1_W_packed), graph.sizes_ubo(mat2_packed), }, + // Push Constants + {}, // Specialization Constants {graph.hashed_layout_of(out), graph.hashed_layout_of(mat1_W_packed), graph.hashed_layout_of(mat2_packed)}, + // Resize Args + {mat2_is_transposed}, // Resizing Logic - resize_matmul_node, - {mat2_is_transposed})); + resize_matmul_node)); } void add_matmul_node( diff --git a/backends/vulkan/runtime/graph/ops/impl/NativeLayerNorm.cpp b/backends/vulkan/runtime/graph/ops/impl/NativeLayerNorm.cpp index f2e8eff763a..100d6e33931 100644 --- a/backends/vulkan/runtime/graph/ops/impl/NativeLayerNorm.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/NativeLayerNorm.cpp @@ -108,24 +108,25 @@ void add_native_layer_norm_node( global_size, local_size, // Inputs and Outputs - {{{out_val->at(0), out_val->at(1), out_val->at(2)}, - vkapi::MemoryAccessType::WRITE}, - {{in, arg_weight, arg_bias}, vkapi::MemoryAccessType::READ}}, + {{{out_val->at(0), out_val->at(1), out_val->at(2)}, vkapi::kWrite}, + {{in, arg_weight, arg_bias}, vkapi::kRead}}, // Shader params buffers {}, + // Push Constants + { + graph.logical_limits_pc_of(out_val->at(0)), + graph.sizes_pc_of(out_val->at(0)), + PushConstantDataInfo(&epsilon, sizeof(epsilon)), + }, // Specialization Constants { t_input->hashed_layout(), t_out->hashed_layout(), }, - // Resizing Logic - resize_native_layer_norm_node, + // Resize Args {normalized_shape}, - { - graph.logical_limits_pc_of(out_val->at(0)), - graph.sizes_pc_of(out_val->at(0)), - PushConstantDataInfo(&epsilon, sizeof(epsilon)), - })); + // Resizing Logic + resize_native_layer_norm_node)); } void native_layer_norm(ComputeGraph& graph, const std::vector& args) { diff --git a/backends/vulkan/runtime/graph/ops/impl/Pad.cpp b/backends/vulkan/runtime/graph/ops/impl/Pad.cpp index 4e4bff189ba..8f3ba7532a9 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Pad.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Pad.cpp @@ -84,17 +84,20 @@ void add_constant_pad_nd_node( graph.create_global_wg_size(out), graph.create_local_wg_size(out), // Inputs and Outputs - {{out, vkapi::MemoryAccessType::WRITE}, - {in, vkapi::MemoryAccessType::READ}}, + {{out, vkapi::kWrite}, {in, vkapi::kRead}}, // Shader params buffers {t_out->sizes_ubo(), t_in->sizes_ubo(), graph.create_params_buffer(pad_param), graph.create_params_buffer(fill_value_val)}, + // Push Constants + {}, // Specialization Constants {}, - resize_constant_pad_node, - {pad})); + // Resize Args + {pad}, + // Resizing Logic + resize_constant_pad_node)); } void constant_pad_nd(ComputeGraph& graph, const std::vector& args) { diff --git a/backends/vulkan/runtime/graph/ops/impl/Permute.cpp b/backends/vulkan/runtime/graph/ops/impl/Permute.cpp index 4352e98de0b..8e2c72d7627 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Permute.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Permute.cpp @@ -85,18 +85,19 @@ void add_permute_node( VK_KERNEL_FROM_STR(kernel_name), graph.create_global_wg_size(out), graph.create_local_wg_size(out), - {{out, vkapi::MemoryAccessType::WRITE}, - {in, vkapi::MemoryAccessType::READ}}, - {}, - // Specialization Constants - spec_vars, - // Resizing Logic - nullptr, + {{out, vkapi::kWrite}, {in, vkapi::kRead}}, {}, + // Push Constants {{graph.logical_limits_pc_of(out), graph.sizes_pc_of(in), PushConstantDataInfo(&out_dims, sizeof(out_dims)), - PushConstantDataInfo(&channel_info, sizeof(channel_info))}})); + PushConstantDataInfo(&channel_info, sizeof(channel_info))}}, + // Specialization Constants + spec_vars, + // Resize Args + {}, + // Resizing Logic + nullptr)); } void add_permute_node( diff --git a/backends/vulkan/runtime/graph/ops/impl/Pool.cpp b/backends/vulkan/runtime/graph/ops/impl/Pool.cpp index b7015d2b1a0..e8afafa9a45 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Pool.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Pool.cpp @@ -98,19 +98,21 @@ void add_max_pool2d_node( global_size, local_size, // Inputs and Outputs - {{{out_val->at(0), out_val->at(1)}, vkapi::MemoryAccessType::WRITE}, - {in, vkapi::MemoryAccessType::READ}}, + {{{out_val->at(0), out_val->at(1)}, vkapi::kWrite}, {in, vkapi::kRead}}, // Shader params buffers { t_out->logical_limits_ubo(), t_in->sizes_ubo(), graph.create_params_buffer(kernel_params), }, + // Push Constants + {}, // Specialization Constants {}, + // Resize Args + {kernel_size, stride, padding, dilation, ceil_mode}, // Resizing Logic - resize_pool2d_node, - {kernel_size, stride, padding, dilation, ceil_mode})); + resize_pool2d_node)); } void max_pool2d(ComputeGraph& graph, const std::vector& args) { @@ -171,22 +173,24 @@ void add_avg_pool2d_node( global_size, local_size, // Inputs and Outputs - {{out, vkapi::MemoryAccessType::WRITE}, - {in, vkapi::MemoryAccessType::READ}}, + {{out, vkapi::kWrite}, {in, vkapi::kRead}}, // Shader params buffers {t_out->logical_limits_ubo(), t_in->sizes_ubo(), graph.create_params_buffer(kernel_params), graph.create_params_buffer(divisor_params)}, + // Push Constants + {}, // Specialization Constants {}, - // Resizing Logic - resize_pool2d_node, + // Resize Args {kernel_size, stride, padding, /*dilation= */ kDummyValueRef, - ceil_mode})); + ceil_mode}, + // Resizing Logic + resize_pool2d_node)); } void avg_pool2d(ComputeGraph& graph, const std::vector& args) { diff --git a/backends/vulkan/runtime/graph/ops/impl/QuantizedLinearQCSNW.cpp b/backends/vulkan/runtime/graph/ops/impl/QuantizedLinearQCSNW.cpp index 7269b75ae6e..85695488dfc 100644 --- a/backends/vulkan/runtime/graph/ops/impl/QuantizedLinearQCSNW.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/QuantizedLinearQCSNW.cpp @@ -128,12 +128,14 @@ void add_linear_qcs8w_node( {{mat1_W_packed, q_mat2, scales}, vkapi::MemoryAccessType::READ}}, // Shader params buffers {}, + // Push Constants + pcs, // Specialization Constants {}, - // Resizing Logic - resize_linear_qcs8w_node, + // Resize Args {}, - pcs)); + // Resizing Logic + resize_linear_qcs8w_node)); if (!graph.is_buffer_storage(out) && graph.packed_dim_of(out) != WHCN::kWidthDim) { viewFn(graph, {out_W_packed, graph.add_none(), out}); @@ -215,13 +217,14 @@ void add_linear_qcs8w_tiled_node( {{out, vkapi::kWrite}, {{mat1, q_mat2, scales}, vkapi::kRead}}, // Shader params buffers {}, + // Push Constants + {{graph.sizes_pc_of(out), graph.sizes_pc_of(mat1)}}, // Specialization Constants {}, - // Resizing Logic - resize_linear_qcs8w_node, + // Resize Args {}, - // Push Constants - {{graph.sizes_pc_of(out), graph.sizes_pc_of(mat1)}})); + // Resizing Logic + resize_linear_qcs8w_node)); } bool can_use_tiled_impl( diff --git a/backends/vulkan/runtime/graph/ops/impl/QuantizedLinearQGANW.cpp b/backends/vulkan/runtime/graph/ops/impl/QuantizedLinearQGANW.cpp index ec718bea7da..b3ead94d8ff 100644 --- a/backends/vulkan/runtime/graph/ops/impl/QuantizedLinearQGANW.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/QuantizedLinearQGANW.cpp @@ -173,15 +173,16 @@ void add_linear_qga4w_node( {{out, vkapi::kWrite}, {{mat1, mat2, scales_and_zeros}, vkapi::kRead}}, // Shader params buffers {}, - // Specialization Constants - {SV(group_size_val)}, - // Resizing Logic - resize_linear_qga4w_node, - {}, // Push Constants {graph.sizes_pc_of(out), graph.sizes_pc_of(mat1), - graph.sizes_pc_of(mat2)})); + graph.sizes_pc_of(mat2)}, + // Specialization Constants + {SV(group_size_val)}, + // Resize Args + {}, + // Resizing Logic + resize_linear_qga4w_node)); } void linear_weight_int4( diff --git a/backends/vulkan/runtime/graph/ops/impl/Reduce.cpp b/backends/vulkan/runtime/graph/ops/impl/Reduce.cpp index 9b1cdf824d2..8fcd4a0609c 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Reduce.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Reduce.cpp @@ -93,11 +93,14 @@ void add_reduce_node( {{out, vkapi::kWrite}, {in, vkapi::kRead}}, // Shader params buffers {graph.logical_limits_ubo(in), graph.sizes_ubo(in)}, + // Push Constants + {}, // Specialization Constants {graph.packed_dim_of(out), reduce_dim, group_dim}, + // Resize Args + {dim}, // Resizing Logic - resize_reduce_node, - {dim})); + resize_reduce_node)); } #define DEFINE_REDUCE_FN(op_name, out_arg_idx) \ diff --git a/backends/vulkan/runtime/graph/ops/impl/Repeat.cpp b/backends/vulkan/runtime/graph/ops/impl/Repeat.cpp index 24e51e99c73..f472e4dad0d 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Repeat.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Repeat.cpp @@ -101,17 +101,20 @@ void add_repeat_node( }, // Parameter buffers {}, - // Specialization Constants - {graph.hashed_layout_of(out), graph.hashed_layout_of(in)}, - nullptr, - {}, + // Push Constants { PushConstantDataInfo(&wg_size, sizeof(wg_size), sizeof(utils::ivec4)), PushConstantDataInfo( &src_dims, sizeof(src_dims), sizeof(utils::ivec4)), PushConstantDataInfo( &dst_repeats, sizeof(dst_repeats), sizeof(utils::ivec4)), - })); + }, + // Specialization Constants + {graph.hashed_layout_of(out), graph.hashed_layout_of(in)}, + // Resize Args + {}, + // Resizing Logic + nullptr)); } void repeat(ComputeGraph& graph, const std::vector& args) { diff --git a/backends/vulkan/runtime/graph/ops/impl/RepeatInterleave.cpp b/backends/vulkan/runtime/graph/ops/impl/RepeatInterleave.cpp index 5e4608a65bb..5bfadf43160 100644 --- a/backends/vulkan/runtime/graph/ops/impl/RepeatInterleave.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/RepeatInterleave.cpp @@ -64,14 +64,17 @@ void add_repeat_interleave_node( {in, vkapi::MemoryAccessType::READ}}, // Parameter buffers {graph.logical_limits_ubo(in)}, + // Push Constants + {}, // Specialization Constants {graph.hashed_layout_of(out), graph.hashed_layout_of(in), nrepeats, repeat_dim}, + // Resize Args + {num_repeats, dim}, // Resizing Logic - resize_repeat_interleave_node, - {num_repeats, dim})); + resize_repeat_interleave_node)); } void repeat_interleave(ComputeGraph& graph, const std::vector& args) { diff --git a/backends/vulkan/runtime/graph/ops/impl/RotaryEmbedding.cpp b/backends/vulkan/runtime/graph/ops/impl/RotaryEmbedding.cpp index 859a3d98aac..ee40a043ee5 100644 --- a/backends/vulkan/runtime/graph/ops/impl/RotaryEmbedding.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/RotaryEmbedding.cpp @@ -67,8 +67,12 @@ void add_rotary_embedding_node( {{xq, xk, freqs_cos, freqs_sin}, vkapi::kRead}}, // Parameter buffers {graph.logical_limits_ubo(xq_out), graph.logical_limits_ubo(xk_out)}, + // Push Constants + {}, // Specialization Constants {}, + // Resize Args + {}, // Resizing Logic resize_rotary_embedding_node)); } diff --git a/backends/vulkan/runtime/graph/ops/impl/SDPA.cpp b/backends/vulkan/runtime/graph/ops/impl/SDPA.cpp index 6dcf2fc4f45..5ef84347181 100644 --- a/backends/vulkan/runtime/graph/ops/impl/SDPA.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/SDPA.cpp @@ -59,11 +59,14 @@ void add_kv_cache_update_node( {{cache, vkapi::kWrite}, {projected, vkapi::kRead}}, // Shader param buffers param_ubos, + // Push Constants + {}, // Specialization Constants {}, + // Resize Args + {}, // Resizing Logic - nullptr, - {})); + nullptr)); } void add_attn_weight_scale_and_mask_node( @@ -113,11 +116,14 @@ void add_attn_weight_scale_and_mask_node( {{attn_weight, vkapi::kReadWrite}}, // Shader param buffers param_ubos, + // Push Constants + {}, // Specialization Constants {}, + // Resize Args + {}, // Resizing Logic - nullptr, - {})); + nullptr)); } std::vector get_cache_slice_sizes( diff --git a/backends/vulkan/runtime/graph/ops/impl/Select.cpp b/backends/vulkan/runtime/graph/ops/impl/Select.cpp index 2e1ae534acd..a83e986e414 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Select.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Select.cpp @@ -119,8 +119,14 @@ void add_select_int_node( // created. graph.create_params_buffer( utils::make_ivec4({index, num_batches, num_texel_per_batch, 0}))}, + // Push Constants + {}, // Specialization Constants - {})); + {}, + // Resize Args + {}, + // Resizing Logic + nullptr)); } void select_int(ComputeGraph& graph, const std::vector& args) { diff --git a/backends/vulkan/runtime/graph/ops/impl/Slice.cpp b/backends/vulkan/runtime/graph/ops/impl/Slice.cpp index efda6e04992..c40e16f7c0a 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Slice.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Slice.cpp @@ -105,7 +105,10 @@ void add_slice_tensor_copy_node( {t_out->sizes_ubo(), t_in->sizes_ubo(), graph.create_params_buffer(params)}, - spec_vars)); + {}, + spec_vars, + {}, + nullptr)); } else { // GPU's coordinate is in x = 0, y = 1, z = 2, w = 3 @@ -146,7 +149,10 @@ void add_slice_tensor_copy_node( {{out, vkapi::MemoryAccessType::WRITE}, {in, vkapi::MemoryAccessType::READ}}, {t_out->sizes_ubo(), graph.create_params_buffer(params)}, - spec_vars)); + {}, + spec_vars, + {}, + nullptr)); } } diff --git a/backends/vulkan/runtime/graph/ops/impl/Softmax.cpp b/backends/vulkan/runtime/graph/ops/impl/Softmax.cpp index b4fede79b12..7469cbb0eb2 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Softmax.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Softmax.cpp @@ -98,8 +98,12 @@ void add_softmax_node( {in, vkapi::MemoryAccessType::READ}}, // Shader params buffers {graph.logical_limits_ubo(out), graph.sizes_ubo(in)}, + // Push Constants + {}, // Specialization Constants {graph.packed_dim_of(out), reduce_dim, group_dim}, + // Resize Args + {}, // Resizing Logic resize_softmax_node)); } diff --git a/backends/vulkan/runtime/graph/ops/impl/Staging.cpp b/backends/vulkan/runtime/graph/ops/impl/Staging.cpp index f59d1cd65d9..32e63baeafc 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Staging.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Staging.cpp @@ -47,11 +47,14 @@ void add_staging_to_tensor_node( {{out_tensor, vkapi::kWrite}, {in_staging, vkapi::kRead}}, // Parameter Buffers ubos, + // Push Constants + {}, // Specialization Constants {graph.hashed_layout_of(out_tensor)}, + // Resize Args + {}, // Resizing Logic - nullptr, - {})); + nullptr)); } const std::string kBitw8PrefixStr = "bitw8_image_to_nchw_nobitw8buffer"; @@ -106,8 +109,14 @@ void add_tensor_to_staging_node( {{out_staging, vkapi::kWrite}, {in_tensor, vkapi::kRead}}, // Parameter Buffers ubos, + // Push Constants + {}, // Specialization Constants - {graph.hashed_layout_of(in_tensor)})); + {graph.hashed_layout_of(in_tensor)}, + // Resize Args + {}, + // Resizing Logic + nullptr)); } void add_prepack_standard_node( diff --git a/backends/vulkan/runtime/graph/ops/impl/UnaryOp.cpp b/backends/vulkan/runtime/graph/ops/impl/UnaryOp.cpp index 9a3ab002403..bffa8e2a181 100644 --- a/backends/vulkan/runtime/graph/ops/impl/UnaryOp.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/UnaryOp.cpp @@ -61,8 +61,12 @@ void add_unary_op_node( {in, vkapi::MemoryAccessType::READ}}, // Shader params buffers ubos, + // Push Constants + {}, // Specialization Constants {}, + // Resize Args + {}, // Resizing Logic resize_unary_op_node)); } diff --git a/backends/vulkan/runtime/graph/ops/impl/Upsample.cpp b/backends/vulkan/runtime/graph/ops/impl/Upsample.cpp index 79777b3f9ac..d098ed94c7f 100644 --- a/backends/vulkan/runtime/graph/ops/impl/Upsample.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/Upsample.cpp @@ -126,10 +126,14 @@ void add_upsample_nearest2d_node( {graph.logical_limits_ubo(out), graph.logical_limits_ubo(in), graph.create_params_buffer(recip_scales)}, + // Push Constants + {}, // Specialization Constants {align_corners_val}, - resize_upsample_nearest2d_node, - {output_sizes, scale_factors})); + // Resize Args + {output_sizes, scale_factors}, + // Resizing Logic + resize_upsample_nearest2d_node)); } void upsample_nearest2d( diff --git a/backends/vulkan/runtime/graph/ops/impl/View.cpp b/backends/vulkan/runtime/graph/ops/impl/View.cpp index fc5c7075222..ef71f8d6d29 100644 --- a/backends/vulkan/runtime/graph/ops/impl/View.cpp +++ b/backends/vulkan/runtime/graph/ops/impl/View.cpp @@ -77,13 +77,14 @@ void add_view_node( {in, vkapi::MemoryAccessType::READ}}, // Parameter Buffers {}, + // Push Constants + {{graph.sizes_pc_of(out), graph.sizes_pc_of(in)}}, // Specialization Constants {SV(t_in->packed_dim()), SV(t_out->packed_dim())}, - // Resizing Logic - resize_view_node, + // Resize Args {sizes}, - // Push Constants - {{graph.sizes_pc_of(out), graph.sizes_pc_of(in)}})); + // Resizing Logic + resize_view_node)); } void view(ComputeGraph& graph, const std::vector& args) { diff --git a/backends/vulkan/test/vulkan_compute_api_test.cpp b/backends/vulkan/test/vulkan_compute_api_test.cpp index 84e9826ec87..143e6704889 100644 --- a/backends/vulkan/test/vulkan_compute_api_test.cpp +++ b/backends/vulkan/test/vulkan_compute_api_test.cpp @@ -1471,11 +1471,13 @@ TEST(VulkanComputeGraphTest, test_simple_graph_with_symint) { // Shader params buffers {graph.logical_limits_ubo(a.value), graph.get_or_create_int_param_buffer(scalar)}, + // Push constants + {}, // Specialization Constants {}, // Resizing Logic - nullptr, - {})); + {}, + nullptr)); out.staging = graph.set_output_tensor(out.value);