37 namespace experimental
39 namespace dynamic_fusion
60 const auto dst_h =
static_cast<int32_t
>(_dst->
dimension(2));
62 auto const_0_i32 = writer->declare_constant_tile(ckw::ConstantData({{0}}, ckw::DataType::Int32));
63 auto const_pos_1_i32 = writer->declare_constant_tile(ckw::ConstantData({{1}}, ckw::DataType::Int32));
64 auto const_dst_h_i32 = writer->declare_constant_tile(ckw::ConstantData({{dst_h}}, ckw::DataType::Int32));
71 const auto &tile_src =
src->tile();
72 auto &sampler_src =
src->tensor_sampler();
74 const auto dst_n0 =
static_cast<int32_t
>(tile_src.tile_info().width());
75 const auto dst_m0 =
static_cast<int32_t
>(tile_src.tile_info().height());
76 const int32_t dst_n0_partial = _dst->
dimension(0) % dst_n0;
77 const int32_t dst_shift_back = (dst_n0 - dst_n0_partial) % dst_n0;
82 auto const_n0_i32 = writer->declare_constant_tile(ckw::ConstantData({{dst_n0}}, ckw::DataType::Int32));
83 auto const_m0_i32 = writer->declare_constant_tile(ckw::ConstantData({{dst_m0}}, ckw::DataType::Int32));
84 auto const_shift_back_n0_i32 =
85 writer->declare_constant_tile(ckw::ConstantData({{dst_shift_back}}, ckw::DataType::Int32));
100 auto tile_gid_0 = writer->declare_tile(
"gid_0", ckw::TileInfo(ckw::DataType::Int32));
101 auto tile_gid_1 = writer->declare_tile(
"gid_1", ckw::TileInfo(ckw::DataType::Int32));
102 auto tile_gid_2 = writer->declare_tile(
"gid_2", ckw::TileInfo(ckw::DataType::Int32));
104 writer->op_get_global_id(tile_gid_0, 0);
105 writer->op_get_global_id(tile_gid_1, 1);
106 writer->op_get_global_id(tile_gid_2, 2);
108 auto tile_nout0 = writer->declare_tile(
"cout0", ckw::TileInfo(ckw::DataType::Int32));
109 auto tile_mout0 = writer->declare_tile(
"mout0", ckw::TileInfo(ckw::DataType::Int32));
110 auto tile_mout1 = writer->declare_tile(
"mout1", ckw::TileInfo(ckw::DataType::Int32));
111 auto tile_bout0 = writer->declare_tile(
"bout0", ckw::TileInfo(ckw::DataType::Int32));
119 if (sampler_src.format() == ckw::TensorSamplerFormat::Dim0_Dim1xDim2_1)
121 writer->op_assign(tile_mout1, const_0_i32);
124 else if (sampler_src.format() == ckw::TensorSamplerFormat::Dim0_Dim1_Dim2)
127 writer->op_binary(tile_mout1, ckw::BinaryOp::Mod, tile_gid_2, const_dst_h_i32);
128 writer->op_binary(tile_bout0, ckw::BinaryOp::Div, tile_gid_2, const_dst_h_i32);
134 writer->op_store(
dst->tensor(), tile_src, sampler_src, tile_nout0, tile_mout0, tile_mout1, tile_bout0);