add tensor initialization to builder

experimental/builder/include/ck_tile/builder/testing/conv_fwd.hpp

@@ -9,6 +9,7 @@
 #include "ck_tile/builder/testing/testing.hpp"
 #include "ck_tile/builder/testing/extent.hpp"
 #include "ck_tile/builder/testing/tensor_buffer.hpp"
+#include "ck_tile/builder/testing/tensor_initialization.hpp"
 #include "ck/library/utility/convolution_parameter.hpp"
 #include "ck/library/utility/convolution_host_tensor_descriptor_helper.hpp"
 /// This file implements common functionality for invoking/testing grouped

experimental/builder/include/ck_tile/builder/testing/tensor_initialization.hpp

@@ -0,0 +1,91 @@
+// Copyright (c) advanced Micro Devices, Inc. All rights reserved.
+// SPDX-License-Identifier: MIT
+
+#pragma once
+
+#include <stdexcept>
+#include <memory>
+#include <numeric>
+#include <span>
+#include <concepts>
+#include <hip/hip_runtime.h>
+#include "ck_tile/builder/conv_signature_concepts.hpp"
+#include "ck_tile/builder/factory/helpers/ck/conv_tensor_type.hpp"
+#include "ck_tile/builder/testing/type_traits.hpp"
+#include "ck_tile/host/host_tensor.hpp"
+#include "ck/utility/data_type.hpp"
+
+#include "ck/library/utility/device_tensor_generator.hpp"
+
+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)
+{
+    size_t size = descriptor.get_element_space_size_in_bytes();
+
+   if(size % sizeof(DT) != 0)
+    {
+        throw std::runtime_error("size must be a multiple of datatype size");
+    }
+    if(max_val - min_val <= 1)
+    {
+        throw std::runtime_error("Error while filling device tensor with random integer data: max "
+                                 "value must be at least 2 greater than min value, otherwise "
+                                 "tensor will be filled by a constant value (end is exclusive)");
+    }
+
+    using ck_type = factory::internal::DataTypeToCK<DT>::type;
+
+    // we might be asked to generate int values on fp data types that don't have the required precision
+    if(static_cast<ck_type>(max_val - 1) == static_cast<ck_type>(min_val))
+    {
+        throw std::runtime_error("Error while filling device tensor with random integer data: "
+                                 "insufficient precision in specified range");
+    }
+    size_t packed_size = ck::packed_size_v<ck_type>;
+    fill_tensor_uniform_rand_int_values<<<256, 256>>>(
+        static_cast<ck_type>(buf.get()), min_val, max_val, (size * packed_size) / sizeof(ck_type));
+}
+
+
+template <DataType DT>
+void init_tensor_buffer_uniform(const DeviceBuffer& buf, const TensorDescriptor<DT>& descriptor, float min_value, float max_value)
+{
+    size_t size = descriptor.get_element_space_size_in_bytes();
+
+   if(size % sizeof(DT) != 0)
+    {
+        throw std::runtime_error("size must be a multiple of datatype size");
+    }
+    using ck_type = factory::internal::DataTypeToCK<DT>::type;
+
+    if(std::nextafter(static_cast<ck_type>(min_value), static_cast<ck_type>(max_value)) >= max_value)
+    {
+        // nextafter is not defined for non-standard data types, skip test for now
+        // throw std::runtime_error("Error while filling device tensor with random integer data:
+        // insufficient precision in specified range");
+    }
+
+    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));
+}
+
+template <DataType DT>
+void init_tensor_buffer_normal_fp(const DeviceBuffer& buf, const TensorDescriptor<DT>& descriptor, float sigma, float mean)
+{
+    size_t size = descriptor.get_element_space_size_in_bytes();
+
+    if(size % sizeof(DT) != 0)
+    {
+        throw std::runtime_error("size must be a multiple of datatype size");
+    }
+    using ck_type = factory::internal::DataTypeToCK<DT>::type;
+    size_t packed_size = ck::packed_size_v<ck_type>;
+    fill_tensor_norm_rand_fp_values<<<256, 256>>>(
+        static_cast<ck_type*>(buf.get()), sigma, mean, (size * packed_size) / sizeof(ck_type));
+
+}
+
+} // namespace ck_tile::builder::test