composable_kernel/include/ck/tensor_operation/gpu/device/impl at 0658353babf6722060acde0c0effd32b024cbf36 - composable_kernel - Public git mirror

ROCm/composable_kernel

mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-16 02:54:21 +00:00

Files

History

Bartłomiej Kocot 1ade932aed Add dynamic elementwise op (#1426 )

* Add dynamic elementwise op

Co-authored-by: ThruptiRajLakshmanaGowda <thruptiraj.lakshmanagowda@amd.com>

* CI issues fix

* Custom parameter value for dynamic functions - Comments addressed

---------

Co-authored-by: ThruptiRajLakshmanaGowda <thruptiraj.lakshmanagowda@amd.com>
Co-authored-by: ThruptiRajLakshmanaGowda <tlakshma@amd.com>

[ROCm/composable_kernel commit: 31bf253aeb]

2024-10-26 15:22:37 +02:00

..

codegen_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432 )

2024-08-06 10:06:10 +02:00

device_avgpool2d_bwd_nhwc_nhwc.hpp

Pool2d max/avg kernel in the BWD version (#1494 )

2024-09-12 11:47:52 +02:00

device_avgpool3d_bwd_ndhwc_ndhwc.hpp

…

device_batched_contraction_multiple_d_wmma_cshuffle.hpp

Merging the gfx12 code into public repo. (#1362 )

2024-06-27 00:33:34 -07:00

device_batched_contraction_multiple_d_xdl_cshuffle.hpp

…

device_batched_gemm_e_permute_xdl.hpp

…

device_batched_gemm_gemm_xdl_cshuffle.hpp

…

device_batched_gemm_multi_d_xdl.hpp

…

device_batched_gemm_multiple_d_dl.hpp

Merging the gfx12 code into public repo. (#1362 )

2024-06-27 00:33:34 -07:00

device_batched_gemm_multiple_d_gemm_multiple_d_xdl_cshuffle.hpp

…

device_batched_gemm_reduce_xdl_cshuffle.hpp

…

device_batched_gemm_softmax_gemm_permute_wmma_cshuffle.hpp

Merging the gfx12 code into public repo. (#1362 )

2024-06-27 00:33:34 -07:00

device_batched_gemm_softmax_gemm_permute_xdl_cshuffle.hpp

…

device_batched_gemm_softmax_gemm_xdl_cshuffle.hpp

…

device_batched_gemm_xdl.hpp

…

device_batchnorm_backward_impl.hpp

…

device_batchnorm_forward_impl_obsolete.hpp

…

device_batchnorm_forward_impl.hpp

…

device_cgemm_4gemm_xdl_cshuffle.hpp

Implement GetWorkSpaceSize from BaseOperator. (#1564 )

2024-10-12 14:05:11 +08:00

device_column_to_image_impl.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432 )

2024-08-06 10:06:10 +02:00

device_contraction_multiple_abd_xdl_cshuffle.hpp

…

device_contraction_multiple_d_xdl_cshuffle.hpp

Merging the gfx12 code into public repo. (#1362 )

2024-06-27 00:33:34 -07:00

device_contraction_utils.hpp

…

device_conv2d_backward_weight_xdl_c_shuffle_nhwc_kyxc_nhwk.hpp

…

device_conv2d_bwd_data_xdl_nhwc_kyxc_nhwk.hpp

…

device_conv2d_fwd_xdl_c_shuffle_bias_activation_add_nhwc_kyxc_nhwk.hpp

…

device_conv2d_fwd_xdl_c_shuffle_bias_activation_nhwc_kyxc_nhwk.hpp

…

device_conv2d_fwd_xdl_c_shuffle_nhwc_kyxc_nhwk.hpp

…

device_conv2d_fwd_xdl_nhwc_kyxc_nhwk.hpp

…

device_conv3d_fwd_naive_ndhwc_kzyxc_ndhwk.hpp

…

device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk.hpp

…

device_convnd_bwd_data_nwc_kxc_nwk_dl.hpp

Merging the gfx12 code into public repo. (#1362 )

2024-06-27 00:33:34 -07:00

device_convnd_bwd_data_nwc_kxc_nwk_xdl.hpp

…

device_elementwise_dynamic_vector_dims_impl.hpp

…

device_elementwise_normalization_impl.hpp

…

device_elementwise_scale_impl.hpp

…

device_fpAintB_gemm_wmma.hpp

Merging the gfx12 code into public repo. (#1362 )

2024-06-27 00:33:34 -07:00

device_gemm_bias_add_reduce_xdl_cshuffle.hpp

…

device_gemm_dl.hpp

Merging the gfx12 code into public repo. (#1362 )

2024-06-27 00:33:34 -07:00

device_gemm_dpp.hpp

…

device_gemm_multiple_abd_xdl_cshuffle.hpp

…

device_gemm_multiple_d_dl.hpp

Merging the gfx12 code into public repo. (#1362 )

2024-06-27 00:33:34 -07:00

device_gemm_multiple_d_layernorm_xdl_cshuffle.hpp

…

device_gemm_multiple_d_multiple_r_xdl_cshuffle.hpp

…

device_gemm_multiple_d_wmma_cshuffle.hpp

Merging the gfx12 code into public repo. (#1362 )

2024-06-27 00:33:34 -07:00

device_gemm_multiple_d_xdl_cshuffle_lds_direct_load.hpp

…

device_gemm_multiple_d_xdl_cshuffle_v3_ab_scale.hpp

[GEMM] F8 GEMM, performance optimized. (#1384 )

2024-07-19 22:06:52 +08:00

device_gemm_multiple_d_xdl_cshuffle_v3.hpp

Add gemm universal bf16 instances (#1484 )

2024-09-04 20:58:54 -07:00

device_gemm_multiple_d_xdl_cshuffle.hpp

…

device_gemm_reduce_xdl_cshuffle.hpp

…

device_gemm_wmma.hpp

Merging the gfx12 code into public repo. (#1362 )

2024-06-27 00:33:34 -07:00

device_gemm_xdl_cshuffle_lds_direct_load.hpp

…

device_gemm_xdl_cshuffle_streamk_v3.hpp

Universal streamk with atomics (#1360 )

2024-07-05 21:40:30 -07:00

device_gemm_xdl_cshuffle_v2.hpp

…

device_gemm_xdl_cshuffle_v3.hpp

Add gemm universal bf16 instances (#1484 )

2024-09-04 20:58:54 -07:00

device_gemm_xdl_cshuffle_v3r1.hpp

Universal gemm splitk using reduce (with multi-d) (#1341 )

2024-07-19 22:01:22 +08:00

device_gemm_xdl_cshuffle.hpp

…

device_gemm_xdl_layernorm_cshuffle.hpp

…

device_gemm_xdl_skip_b_lds.hpp

…

device_gemm_xdl_splitk_c_shuffle_lds_direct_load.hpp

…

device_gemm_xdl_splitk_c_shuffle.hpp

…

device_gemm_xdl_streamk.hpp

…

device_gemm_xdl_waveletmodel_cshuffle.hpp

…

device_gemm_xdl.hpp

…

device_grouped_contraction_multiple_d_xdl_cshuffle.hpp

…

device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle.hpp

Merging the gfx12 code into public repo. (#1362 )

2024-06-27 00:33:34 -07:00

device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp

Add read_first_lane function for int64 (#1347 )

2024-06-18 15:05:30 -05:00

device_grouped_conv_bwd_weight_dl.hpp

Add support for NGCHW in grouped conv bwd wei (#1491 )

2024-09-03 10:52:03 +02:00

device_grouped_conv_bwd_weight_multiple_d_xdl_cshuffle.hpp

Add support for NGCHW in grouped conv bwd wei (#1491 )

2024-09-03 10:52:03 +02:00

device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp

Add support for NGCHW in grouped conv fwd (#1499 )

2024-09-20 10:45:46 +02:00

device_grouped_conv_bwd_weight_wmma_cshuffle.hpp

Add support for NGCHW in grouped conv bwd wei (#1491 )

2024-09-03 10:52:03 +02:00

device_grouped_conv_bwd_weight_xdl_cshuffle.hpp

Add support for NGCHW in grouped conv bwd wei (#1491 )

2024-09-03 10:52:03 +02:00

device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432 )

2024-08-06 10:06:10 +02:00

device_grouped_conv_fwd_dl_nhwc_kyxc_nhwk.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432 )

2024-08-06 10:06:10 +02:00

device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp

Add support for NGCHW in grouped conv fwd (#1499 )

2024-09-20 10:45:46 +02:00

device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp

Add dynamic elementwise op (#1426 )

2024-10-26 15:22:37 +02:00

device_grouped_conv_fwd_multiple_d_multiple_r_xdl_cshuffle.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432 )

2024-08-06 10:06:10 +02:00

device_grouped_conv_fwd_multiple_d_multiple_r.hpp

…

device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432 )

2024-08-06 10:06:10 +02:00

device_grouped_conv_fwd_multiple_d_xdl_cshuffle.hpp

…

device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp

Add support for NGCHW in grouped conv bwd wei (#1491 )

2024-09-03 10:52:03 +02:00

device_grouped_conv_utils.hpp

Add support for NGCHW in grouped conv fwd (#1499 )

2024-09-20 10:45:46 +02:00

device_grouped_gemm_multi_abd_xdl_fixed_nk.hpp

Fix grouped gemm check to avoid overflow (#1545 )

2024-10-04 17:32:43 +02:00

device_grouped_gemm_multiple_d_dl.hpp

Merging the gfx12 code into public repo. (#1362 )

2024-06-27 00:33:34 -07:00

device_grouped_gemm_multiple_d_splitk_xdl_cshuffle_two_stage.hpp

Fix grouped gemm check to avoid overflow (#1545 )

2024-10-04 17:32:43 +02:00

device_grouped_gemm_multiple_d_xdl_cshuffle_tile_loop.hpp

Fix grouped gemm check to avoid overflow (#1545 )

2024-10-04 17:32:43 +02:00

device_grouped_gemm_softmax_gemm_permute_xdl_cshuffle.hpp

…

device_grouped_gemm_xdl_fixed_nk.hpp

Fix grouped gemm check to avoid overflow (#1545 )

2024-10-04 17:32:43 +02:00

device_grouped_gemm_xdl_splitk_cshuffle.hpp

…

device_grouped_gemm_xdl.hpp

…

device_grouped_query_attention_forward_wmma.hpp

Merging the gfx12 code into public repo. (#1362 )

2024-06-27 00:33:34 -07:00

device_image_to_column_impl.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432 )

2024-08-06 10:06:10 +02:00

device_max_pool_bwd_impl.hpp

…

device_multi_query_attention_forward_wmma.hpp

Merging the gfx12 code into public repo. (#1362 )

2024-06-27 00:33:34 -07:00

device_multiple_reduce_multiblock.hpp

…

device_multiple_reduce_threadwise.hpp

…

device_normalization_bwd_data_impl.hpp

…

device_normalization_bwd_gamma_beta_impl.hpp

…

device_normalization_fwd_impl.hpp

…

device_normalization_fwd_splitk_impl.hpp

…

device_permute_impl.hpp

…

device_pool2d_fwd_nhwc_nhwc.hpp

Rewrite pool2d fwd (#1462 )

2024-09-11 15:21:00 +02:00

device_pool3d_fwd_ndhwc_ndhwc.hpp

…

device_put_element_impl.hpp

…

device_reduce_common.hpp

Support large: 12d tensor size for reduction kenrel (#1465 )

2024-08-13 16:15:47 +02:00

device_reduce_multiblock.hpp

Support large: 12d tensor size for reduction kenrel (#1465 )

2024-08-13 16:15:47 +02:00

device_reduce_threadwise_multi_d.hpp

Support large: 12d tensor size for reduction kenrel (#1465 )

2024-08-13 16:15:47 +02:00

device_reduce_threadwise.hpp

Support large: 12d tensor size for reduction kenrel (#1465 )

2024-08-13 16:15:47 +02:00

device_softmax_impl.hpp

…

device_sparse_embeddings_forward_layernorm.hpp

…

device_splitk_contraction_multiple_d_xdl_cshuffle.hpp

…