ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-19 02:01:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

integrate tensor initialization

Browse Source
This commit is contained in:
Kevin Abraham
2025-12-10 13:52:55 +00:00
parent ee08e55e48
commit 60a476420d
4 changed files with 20 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

13
experimental/builder/include/ck_tile/builder/testing/conv_fwd.hpp
View File

@@ -239,7 +239,20 @@ UniqueInputs<SIGNATURE> alloc_inputs(const Args<SIGNATURE>& args)
};
}
/// @brief `init_inputs()` specialization for forward convolution.
///
/// @tparam SIGNATURE Forward convolution signature.
///
/// @see alloc_inputs()
template <auto SIGNATURE>
requires ValidConvSignature<SIGNATURE> && ConvDirectionIsForward<SIGNATURE> &&
ValidUniqueInputs<SIGNATURE>
void init_inputs(const Args<SIGNATURE>& args, UniqueInputs<SIGNATURE>& inputs)
{
init_tensor_buffer_uniform_fp(inputs.input_buf, args.make_input_descriptor(), -2.0f, 2.0f);
init_tensor_buffer_uniform_fp(inputs.weight_buf, args.make_weight_descriptor(), -2.0f, 2.0f);
}
/// @brief `alloc_outputs()` specialization for forward convolution.
///

6
experimental/builder/include/ck_tile/builder/testing/tensor_initialization.hpp
View File

@@ -20,7 +20,7 @@
namespace ck_tile::builder::test {
template <DataType DT>
void init_tensor_buffer_uniform(const DeviceBuffer& buf, const TensorDescriptor<DT>& descriptor, int min_val, int max_val)
void init_tensor_buffer_uniform_int(const DeviceBuffer& buf, const TensorDescriptor<DT>& descriptor, int min_val, int max_val)
{
size_t size = descriptor.get_element_space_size_in_bytes();
@@ -50,7 +50,7 @@ void init_tensor_buffer_uniform(const DeviceBuffer& buf, const TensorDescriptor<
template <DataType DT>
void init_tensor_buffer_uniform(const DeviceBuffer& buf, const TensorDescriptor<DT>& descriptor, float min_value, float max_value)
void init_tensor_buffer_uniform_fp(const DeviceBuffer& buf, const TensorDescriptor<DT>& descriptor, float min_value, float max_value)
{
size_t size = descriptor.get_element_space_size_in_bytes();
@@ -69,7 +69,7 @@ void init_tensor_buffer_uniform(const DeviceBuffer& buf, const TensorDescriptor<
size_t packed_size = ck::packed_size_v<ck_type>;
fill_tensor_uniform_rand_fp_values<<<256, 256>>>(
static_cast<ck_type*>(buf.get()), min_value, max_value, (size * packed_size) / sizeof(ck_type));
reinterpret_cast<ck_type*>(buf.get()), min_value, max_value, (size * packed_size) / sizeof(ck_type));
}
template <DataType DT>

3
experimental/builder/include/ck_tile/builder/testing/testing.hpp
View File

@@ -214,9 +214,10 @@ UniqueInputs<SIGNATURE> alloc_inputs(const Args<SIGNATURE>& args);
///
/// @see Inputs
/// @see UniqueInputs
/// @see tensor_initialization
template <auto SIGNATURE>
requires ValidUniqueInputs<SIGNATURE>
UniqueInputs<SIGNATURE> init_inputs(const Args<SIGNATURE>& args);
void init_inputs(const Args<SIGNATURE>& args, UniqueInputs<SIGNATURE>& inputs);
/// @brief Allocate outputs corresponding to a signature.
///

2
experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp16.cpp
View File

@@ -81,6 +81,8 @@ TEST(Fwd2DFp16_CShufV3_GNHWC, EndToEnd)
auto inputs = alloc_inputs(args);
auto outputs = alloc_outputs(args);
init_inputs(args, inputs);
auto conv = Instance{};
ckt::run(conv, args, inputs.get(), outputs.get());
}