| /aosp_15_r20/external/executorch/kernels/portable/ |
| H A D | functions.yaml | 23 kernel_name: torch::executor::_cdist_forward_out
28 kernel_name: torch::executor::log_softmax_out
33 kernel_name: torch::executor::_native_batch_norm_legit_out
38 kernel_name: torch::executor::_native_batch_norm_legit_no_stats_out
43 kernel_name: torch::executor::_native_batch_norm_legit_no_training_out
48 kernel_name: torch::executor::_pdist_forward_out
53 kernel_name: torch::executor::softmax_out
58 kernel_name: torch::executor::to_copy_out
63 kernel_name: torch::executor::abs_out
68 kernel_name: torch::executor::acos_out
[all …]
|
| /aosp_15_r20/external/executorch/backends/cadence/aot/ |
| H A D | functions.yaml | 18 kernel_name: torch::executor::to_copy_out
23 kernel_name: torch::executor::softmax_out
28 kernel_name: torch::executor::add_out
33 kernel_name: torch::executor::bmm_out
38 kernel_name: torch::executor::cat_out
43 kernel_name: torch::executor::clone_out
48 kernel_name: torch::executor::div_out
53 kernel_name: torch::executor::div_out_mode
58 kernel_name: torch::executor::embedding_out
63 kernel_name: torch::executor::empty_out
[all …]
|
| H A D | functions_hifi.yaml | 18 kernel_name: torch::executor::to_copy_out
23 kernel_name: torch::executor::softmax_out
28 kernel_name: cadence::impl::HiFi::add_out
33 kernel_name: torch::executor::bmm_out
38 kernel_name: torch::executor::cat_out
43 kernel_name: torch::executor::clone_out
48 kernel_name: cadence::impl::HiFi::div_out
53 kernel_name: cadence::impl::HiFi::div_out_mode
58 kernel_name: torch::executor::embedding_out
63 kernel_name: torch::executor::full_out
[all …]
|
| /aosp_15_r20/external/executorch/backends/vulkan/runtime/graph/ops/utils/ |
| H A D | StagingUtils.cpp | 26 std::string kernel_name; in get_nchw_to_tensor_shader() local
27 kernel_name.reserve(kShaderNameReserve); in get_nchw_to_tensor_shader()
31 kernel_name = "nchw_to_bitw8_image_nobitw8buffer"; in get_nchw_to_tensor_shader()
32 add_storage_type_suffix(kernel_name, v_dst); in get_nchw_to_tensor_shader()
33 add_dtype_suffix(kernel_name, v_dst); in get_nchw_to_tensor_shader()
34 return VK_KERNEL_FROM_STR(kernel_name); in get_nchw_to_tensor_shader()
38 kernel_name = "nchw_to_buffer"; in get_nchw_to_tensor_shader()
39 add_dtype_suffix(kernel_name, v_dst); in get_nchw_to_tensor_shader()
40 return VK_KERNEL_FROM_STR(kernel_name); in get_nchw_to_tensor_shader()
43 kernel_name = "nchw_to_image"; in get_nchw_to_tensor_shader()
[all …]
|
| H A D | ShaderNameUtils.cpp | 14 std::string& kernel_name, in add_storage_type_suffix() argument
18 kernel_name += "_buffer"; in add_storage_type_suffix()
21 kernel_name += "_texture3d"; in add_storage_type_suffix()
24 kernel_name += "_texture2d"; in add_storage_type_suffix()
30 std::string& kernel_name, in add_storage_type_suffix() argument
32 return add_storage_type_suffix(kernel_name, tensor.storage_type()); in add_storage_type_suffix()
35 void add_dtype_suffix(std::string& kernel_name, const vkapi::ScalarType dtype) { in add_dtype_suffix() argument
38 kernel_name += "_float"; in add_dtype_suffix()
41 kernel_name += "_half"; in add_dtype_suffix()
44 kernel_name += "_int"; in add_dtype_suffix()
[all …]
|
| /aosp_15_r20/external/pytorch/torch/_inductor/codegen/ |
| H A D | cpp_wrapper_cuda.py | 36 kernel_name: str,
42 self.kernel_name = kernel_name
47 params = CudaKernelParamCache.get(self.kernel_name)
50 ), f"{self.kernel_name} not found in CudaKernelParamCache"
54 ), f"{key} not found in CudaKernelParamCache[{self.kernel_name}]"
61 return DeferredCudaKernelLine(self.kernel_name, line, self.keys)
71 kernel_name: str,
76 self.kernel_name = kernel_name
97 params = CudaKernelParamCache.get(self.kernel_name)
100 ), f"{self.kernel_name} not found in CudaKernelParamCache"
[all …]
|
| H A D | debug_utils.py | 34 kernel_name: str = "",
42 self.kernel_name = kernel_name
50 self.kernel_name,
55 def __exit__(self, args_to_print_or_save, kernel_name, arg_signatures): argument
58 kernel_name,
66 kernel_name, argument
76 self.kernel_name,
84 self.kernel_name,
101 kernel_name: str,
112 self.kernel_name = kernel_name
[all …]
|
| H A D | multi_kernel.py | 103 kernel_names = tuple(k.kernel_name for k in kernels)
150 self.kernel_name = V.graph.wrapper_code.multi_kernel_state.define_kernel(
158 def call_kernel(self, kernel_name): argument
163 assert kernel_name == self.kernel_name
176 picked_kernel = MultiKernelCall.lookup_choice(kernel_name)
177 kernel_name = self.kernels[picked_kernel].kernel_name
181 kernel_name, call_args, arg_types, grid
184 grid = V.graph.wrapper_code.generate_default_grid(kernel_name, grid)
186 kernel_name,
225 def warn_mix_layout(self, kernel_name: str):
[all …]
|
| /aosp_15_r20/external/executorch/kernels/optimized/ |
| H A D | optimized.yaml | 8 kernel_name: torch::executor::opt_log_softmax_out
13 kernel_name: torch::executor::opt_add_out
18 kernel_name: torch::executor::opt_add_scalar_out
23 kernel_name: torch::executor::opt_bmm_out
28 kernel_name: torch::executor::opt_div_out
33 kernel_name: torch::executor::opt_div_scalar_out
38 kernel_name: torch::executor::opt_exp_out
43 kernel_name: torch::executor::opt_sigmoid_out
48 kernel_name: torch::executor::opt_gelu_out
53 kernel_name: torch::executor::opt_le_scalar_out
[all …]
|
| H A D | optimized-oss.yaml | 11 kernel_name: torch::executor::opt_add_out
16 kernel_name: torch::executor::opt_add_scalar_out
21 kernel_name: torch::executor::opt_bmm_out
26 kernel_name: torch::executor::opt_div_out
31 kernel_name: torch::executor::opt_div_scalar_out
36 kernel_name: torch::executor::opt_exp_out
41 kernel_name: torch::executor::opt_sigmoid_out
46 kernel_name: torch::executor::opt_le_scalar_out
51 kernel_name: torch::executor::opt_le_tensor_out
56 kernel_name: torch::executor::opt_linear_out
[all …]
|
| /aosp_15_r20/external/executorch/kernels/quantized/ |
| H A D | quantized.yaml | 5 kernel_name: torch::executor::quantized_add_out
11 kernel_name: torch::executor::choose_qparams_tensor_out
17 kernel_name: torch::executor::dequantize_per_tensor_out
23 kernel_name: torch::executor::dequantize_per_tensor_tensor_args_out
29 kernel_name: torch::executor::quantize_per_channel_out
35 kernel_name: torch::executor::dequantize_per_channel_out
41 kernel_name: torch::executor::quantized_embedding_byte_out
47 kernel_name: torch::executor::quantized_embedding_byte_dtype_out
53 kernel_name: torch::executor::quantized_embedding_2bit_out
59 kernel_name: torch::executor::quantized_embedding_2bit_dtype_out
[all …]
|
| /aosp_15_r20/external/executorch/backends/vulkan/runtime/graph/ops/impl/ |
| H A D | QuantizedLinear.cpp | 89 std::string kernel_name = "q_8w_linear"; in add_q_8w_linear_node() local
90 kernel_name.reserve(kShaderNameReserve); in add_q_8w_linear_node()
91 add_packed_dim_suffix(kernel_name, graph.packed_dim_of(mat1_W_packed)); in add_q_8w_linear_node()
92 add_packed_dim_suffix(kernel_name, graph.packed_dim_of(q_mat2)); in add_q_8w_linear_node()
93 add_dtype_suffix(kernel_name, graph.dtype_of(out_W_packed)); in add_q_8w_linear_node()
94 add_storage_type_suffix(kernel_name, graph.storage_type_of(out_W_packed)); in add_q_8w_linear_node()
118 VK_KERNEL_FROM_STR(kernel_name), in add_q_8w_linear_node()
160 std::string kernel_name = "q_8w_linear_optimized"; in add_q_8w_linear_optimized_node() local
161 kernel_name.reserve(kShaderNameReserve); in add_q_8w_linear_optimized_node()
162 add_packed_dim_suffix(kernel_name, graph.packed_dim_of(mat1_W_packed)); in add_q_8w_linear_optimized_node()
[all …]
|
| H A D | MatMul.cpp | 72 std::string kernel_name = "matmul_naive_buffer"; in add_matmul_naive_buffer_node() local
73 add_dtype_suffix(kernel_name, graph.dtype_of(out)); in add_matmul_naive_buffer_node()
87 VK_KERNEL_FROM_STR(kernel_name), in add_matmul_naive_buffer_node()
123 std::string kernel_name = graph.get_bool(mat2_is_transposed) in add_matmul_naive_texture3d_node() local
126 kernel_name.reserve(kShaderNameReserve); in add_matmul_naive_texture3d_node()
127 add_storage_type_suffix(kernel_name, graph.storage_type_of(out)); in add_matmul_naive_texture3d_node()
128 add_dtype_suffix(kernel_name, graph.dtype_of(out)); in add_matmul_naive_texture3d_node()
133 VK_KERNEL_FROM_STR(kernel_name), in add_matmul_naive_texture3d_node()
184 std::string kernel_name = mat2_is_transposed_val in add_matmul_optimized_node() local
191 kernel_name = "batch_" + kernel_name; in add_matmul_optimized_node()
[all …]
|
| H A D | Convolution.cpp | 130 std::string kernel_name; in get_conv2d_shader() local
131 kernel_name.reserve(kShaderNameReserve); in get_conv2d_shader()
134 kernel_name = "conv2d_dw"; in get_conv2d_shader()
138 kernel_name += "_output_tile_3x3"; in get_conv2d_shader()
141 kernel_name += "_output_tile_5x5"; in get_conv2d_shader()
147 kernel_name = "conv2d"; in get_conv2d_shader()
149 kernel_name = "conv2d_pw"; in get_conv2d_shader()
153 kernel_name = "conv2d"; in get_conv2d_shader()
156 kernel_name = "conv_transpose2d"; in get_conv2d_shader()
160 kernel_name += "_prepack_weights"; in get_conv2d_shader()
[all …]
|
| H A D | Select.cpp | 66 std::string kernel_name; in add_select_int_node() local
75 kernel_name = "select_channel_3d"; in add_select_int_node()
77 kernel_name = "select_height_3d"; in add_select_int_node()
79 kernel_name = "select_width_3d"; in add_select_int_node()
89 kernel_name = "select_batch_4d"; in add_select_int_node()
91 kernel_name = "select_channel_4d"; in add_select_int_node()
93 kernel_name = "select_height_4d"; in add_select_int_node()
95 kernel_name = "select_width_4d"; in add_select_int_node()
102 kernel_name.reserve(kShaderNameReserve); in add_select_int_node()
103 add_dtype_suffix(kernel_name, *t_out); in add_select_int_node()
[all …]
|
| H A D | Copy.cpp | 31 std::string kernel_name = "copy_offset"; in add_copy_offset_node() local
32 kernel_name.reserve(kShaderNameReserve); in add_copy_offset_node()
33 add_dtype_suffix(kernel_name, *t_out); in add_copy_offset_node()
34 add_storage_type_suffix(kernel_name, *t_out); in add_copy_offset_node()
46 auto shader = VK_KERNEL_FROM_STR(kernel_name); in add_copy_offset_node()
50 VK_KERNEL_FROM_STR(kernel_name), in add_copy_offset_node()
114 std::string kernel_name = "copy_channel_offset"; in add_copy_channel_offset_node() local
115 kernel_name.reserve(kShaderNameReserve); in add_copy_channel_offset_node()
116 add_dtype_suffix(kernel_name, *t_out); in add_copy_channel_offset_node()
154 auto shader = VK_KERNEL_FROM_STR(kernel_name); in add_copy_channel_offset_node()
[all …]
|
| H A D | Linear.cpp | 103 std::string kernel_name = in add_addmm_naive_node() local
105 kernel_name.reserve(kShaderNameReserve); in add_addmm_naive_node()
106 add_storage_type_suffix(kernel_name, graph.storage_type_of(out)); in add_addmm_naive_node()
107 add_dtype_suffix(kernel_name, graph.dtype_of(out)); in add_addmm_naive_node()
112 VK_KERNEL_FROM_STR(kernel_name), in add_addmm_naive_node()
169 std::string kernel_name = graph.get_bool(mat2_is_transposed) in add_addmm_optimized_node() local
176 kernel_name = "batch_" + kernel_name; in add_addmm_optimized_node()
179 kernel_name += "_tile_row_2"; in add_addmm_optimized_node()
181 kernel_name += "_tile_row_4"; in add_addmm_optimized_node()
184 add_dtype_suffix(kernel_name, graph.dtype_of(out)); in add_addmm_optimized_node()
[all …]
|
| /aosp_15_r20/external/executorch/backends/vulkan/test/utils/ |
| H A D | test_utils.cpp | 120 std::string kernel_name = "bitw8_image_to_nchw_nobitw8buffer"; in record_bitw8_image_to_nchw_nobitw8buffer_op() local
121 add_storage_type_suffix(kernel_name, v_src); in record_bitw8_image_to_nchw_nobitw8buffer_op()
122 add_dtype_suffix(kernel_name, v_src); in record_bitw8_image_to_nchw_nobitw8buffer_op()
125 VK_KERNEL_FROM_STR(kernel_name), in record_bitw8_image_to_nchw_nobitw8buffer_op()
146 std::string kernel_name; in record_conv2d_prepack_weights_op() local
148 kernel_name = "conv_transpose2d"; in record_conv2d_prepack_weights_op()
150 kernel_name = "conv2d"; in record_conv2d_prepack_weights_op()
152 kernel_name += "_prepack_weights"; in record_conv2d_prepack_weights_op()
153 add_dtype_suffix(kernel_name, v_dst); in record_conv2d_prepack_weights_op()
154 vkapi::ShaderInfo shader = VK_KERNEL_FROM_STR(kernel_name); in record_conv2d_prepack_weights_op()
[all …]
|
| /aosp_15_r20/external/pytorch/torch/_inductor/ |
| H A D | autotune_process.py | 481 kernel_name: str,
487 self.kernel_name = kernel_name
609 kernel_name: str,
620 super().__init__(kernel_name, input_tensor_meta, output_tensor_meta, extra_args)
638 run_method = getattr(mod, self.kernel_name).run
674 getattr(mod, self.kernel_name).precompile()
677 return f"{self.kernel_name=}, {self.module_path=}, {self.module_cache_key=}"
686 kernel_name: str,
692 super().__init__(kernel_name, input_tensor_meta, output_tensor_meta, extra_args)
719 …"make_run_fn: self.kernel_name=%s, self.source_file=%s, self.hash_key=%s, self.DLL=%s, args=%s, se…
[all …]
|
| /aosp_15_r20/external/executorch/shim/xplat/executorch/kernels/test/ |
| H A D | util.bzl | 3 def op_test(name, deps = [], kernel_name = "portable", use_kernel_prefix = False):
13 kernel_name: The name string as in //executorch/kernels/<kernel_name>.
21 if kernel_name == "aten":
31 "//executorch/kernels/{}/cpu:{}".format(kernel_name, op_root),
32 "//executorch/kernels/{}:generated_lib_headers".format(kernel_name),
33 "//executorch/kernels/{}/test:supported_features".format(kernel_name),
38 if kernel_name == "aten":
43 name_prefix = kernel_name + "_"
116 def codegen_function_header_wrapper(kernel_path, kernel_name):
121 Use target "function_header_wrapper_<kernel_name>" in tests.
[all …]
|
| /aosp_15_r20/external/gemmlowp/meta/generators/ |
| H A D | common.py | 90 def __init__(self, emitter, kernel_name, output_stream_name): argument
91 self.kernel_name = kernel_name
99 in_type, out_type, self.kernel_name, self.output_stream_name, kernel_m,
105 'const FusedKernelParams<%s, %s>&' % (self.kernel_name,
110 _TemplateName(self.kernel_name + self.output_stream_name,
119 def __init__(self, emitter, kernel_name): argument
120 self.kernel_name = kernel_name
127 in_type, out_type, self.kernel_name, kernel_size, leftovers
132 ['const %s&' % self.kernel_name, 'params'],
135 _TemplateName(self.kernel_name, template_params))
|
| /aosp_15_r20/external/OpenCL-CTS/test_conformance/math_brute_force/ |
| H A D | common.cpp | 123 std::string GetUnaryKernel(const std::string &kernel_name, const char *builtin, in GetUnaryKernel() argument
136 __kernel void )", kernel_name.c_str(), R"((__global RETTYPE* out, in GetUnaryKernel()
145 __kernel void )", kernel_name.c_str(), R"((__global RETTYPE_SCALAR* out, in GetUnaryKernel()
197 std::string GetUnaryKernel(const std::string &kernel_name, const char *builtin, in GetUnaryKernel() argument
212 __kernel void )", kernel_name.c_str(), R"((__global RETTYPE1* out1, in GetUnaryKernel()
222 __kernel void )", kernel_name.c_str(), R"((__global RETTYPE1_SCALAR* out1, in GetUnaryKernel()
280 std::string GetBinaryKernel(const std::string &kernel_name, const char *builtin, in GetBinaryKernel() argument
312 __kernel void )", kernel_name.c_str(), R"((__global RETTYPE* out, in GetBinaryKernel()
322 __kernel void )", kernel_name.c_str(), R"((__global RETTYPE_SCALAR* out, in GetBinaryKernel()
379 std::string GetBinaryKernel(const std::string &kernel_name, const char *builtin, in GetBinaryKernel() argument
[all …]
|
| /aosp_15_r20/external/pytorch/torch/_inductor/codegen/cuda/ |
| H A D | cuda_cpp_scheduling.py | 52 kernel_name = wrapper.src_to_kernel[src_code]
59 kernel_name = "_".join(["cuda", fused_name, wrapper.next_kernel_suffix()])
61 wrapper.src_to_kernel[src_code] = kernel_name
62 src_code = src_code.replace("KERNEL_NAME", kernel_name)
75 kernel_name, compile_wrapper.getvalue(), metadata_comment
77 return kernel_name
102 kernel_name = self.define_kernel(src_code, node_schedule)
108 call_args, kernel_name, arg_signatures, kernel
111 kernel.call_kernel(kernel_name, ctb)
|
| /aosp_15_r20/external/executorch/kernels/quantized/test/ |
| H A D | targets.bzl | 6 op_test("op_quantize_test", kernel_name = "quantized")
7 op_test("op_dequantize_test", kernel_name = "quantized")
8 op_test("op_choose_qparams_test", kernel_name = "quantized")
9 op_test("op_add_test", kernel_name = "quantized", deps = [
18 op_test("op_embedding_test", kernel_name = "quantized", deps = [
28 op_test("op_embedding2b_test", kernel_name = "quantized")
29 op_test("op_embedding4b_test", kernel_name = "quantized")
30 op_test("op_mixed_mm_test", kernel_name = "quantized", deps = [
36 op_test("op_mixed_linear_test", kernel_name = "quantized", deps = [
|
| /aosp_15_r20/external/ComputeLibrary/src/gpu/cl/kernels/ |
| H A D | ClIm2ColKernel.cpp | 56 std::string kernel_name{}; member
162 std::string kernel_name = "im2col_generic_"; in configure_opencl_kernel() local
197 kernel_name = "im2col3x3_"; in configure_opencl_kernel()
202 kernel_name = "im2col9x9_"; in configure_opencl_kernel()
233 kernel_name = "im2col1x1_stridex1_"; in configure_opencl_kernel()
239 kernel_name = "im2col3x3_"; in configure_opencl_kernel()
244 kernel_name = "im2col5x5_"; in configure_opencl_kernel()
252 kernel_name = "im2col11x11_padx0_pady0_"; in configure_opencl_kernel()
258 kernel_name = "im2col_generic_"; in configure_opencl_kernel()
266 kernel_name = "im2col_generic_padx0_pady0_"; in configure_opencl_kernel()
[all …]
|