From a091652370cb2a2c29d60100253fd6fba2882307 Mon Sep 17 00:00:00 2001 From: Dylan Lim Date: Mon, 27 Jan 2025 20:57:10 -0800 Subject: [PATCH] branch merge and test fixes --- lib/kernels/include/kernels/accessor.h | 76 ++++-------- lib/kernels/include/kernels/flat_kernels.h | 2 +- .../include/kernels/loss_function_kernels.h | 2 +- .../include/kernels/managed_ff_stream.h | 1 + .../kernels/managed_per_device_ff_handle.h | 1 + lib/kernels/include/kernels/metrics_kernels.h | 6 +- lib/kernels/include/kernels/pool_2d_kernels.h | 2 +- lib/kernels/src/cuda/metrics_functions.cu | 10 +- lib/kernels/src/hip/embedding_kernels.cpp | 30 ++--- .../test/src/test_batch_norm_kernel.cc | 2 +- lib/kernels/test/src/test_concat_kernel.cc | 2 +- lib/kernels/test/src/test_flat_kernel.cc | 2 +- .../test/src/test_layer_norm_kernels.cc | 2 +- .../src/test_managed_per_device_ff_handle.cc | 7 +- lib/kernels/test/src/test_partition_kernel.cc | 2 +- lib/kernels/test/src/test_pool_2d_kernels.cc | 2 +- lib/kernels/test/src/test_reduction_kernel.cc | 2 +- lib/kernels/test/src/test_reverse_kernels.cc | 2 +- lib/kernels/test/src/test_split_kernel.cc | 2 +- lib/kernels/test/src/test_transpose_kernel.cc | 3 +- lib/kernels/test/src/test_utils.cc | 18 ++- lib/kernels/test/src/test_utils.h | 2 +- .../test/src/test_local_cost_estimator.cc | 115 +++++++++--------- .../include/op-attrs/aggregate_op.enum.toml | 2 - .../include/op-attrs/make_datatype_value.h | 16 --- ...ke_datatype_value.cc => datatype_value.cc} | 2 +- .../test/src/op-attrs/datatype_value.cc | 68 +++++++++++ lib/pcg/include/pcg/metric.enum.toml | 26 ++++ lib/pcg/include/pcg/metric.h | 72 ----------- lib/pcg/include/pcg/metric_attrs.h | 28 +++++ lib/pcg/include/pcg/strided_rectangle.h | 17 --- lib/pcg/src/pcg/computation_graph_builder.cc | 2 +- lib/pcg/src/pcg/metric.cc | 8 +- .../parallel_computation_graph_builder.cc | 2 +- lib/pcg/src/pcg/strided_rectangle_side.cc | 17 --- lib/pcg/src/strided_rectangle.cc | 35 ------ lib/pcg/test/src/test_machine_view.cc | 74 ----------- lib/pcg/test/src/test_strided_rectangle.cc | 37 ------ 38 files changed, 263 insertions(+), 436 deletions(-) delete mode 100644 lib/op-attrs/include/op-attrs/make_datatype_value.h rename lib/op-attrs/src/op-attrs/{make_datatype_value.cc => datatype_value.cc} (92%) create mode 100644 lib/op-attrs/test/src/op-attrs/datatype_value.cc create mode 100644 lib/pcg/include/pcg/metric.enum.toml delete mode 100644 lib/pcg/include/pcg/metric.h create mode 100644 lib/pcg/include/pcg/metric_attrs.h delete mode 100644 lib/pcg/include/pcg/strided_rectangle.h delete mode 100644 lib/pcg/src/pcg/strided_rectangle_side.cc delete mode 100644 lib/pcg/src/strided_rectangle.cc delete mode 100644 lib/pcg/test/src/test_machine_view.cc delete mode 100644 lib/pcg/test/src/test_strided_rectangle.cc diff --git a/lib/kernels/include/kernels/accessor.h b/lib/kernels/include/kernels/accessor.h index 487bc1f8f0..a6fc4129e0 100644 --- a/lib/kernels/include/kernels/accessor.h +++ b/lib/kernels/include/kernels/accessor.h @@ -11,6 +11,28 @@ namespace FlexFlow { +inline int calculate_accessor_offset(std::vector const &indices, + ArrayShape const &shape) { + int offset = 0; + int multiplier = 1; + + for (int i = 0; i < shape.num_dims(); i++) { + if (indices.at(i) >= shape.at(legion_dim_t{i})) { + throw mk_runtime_error( + fmt::format("In {} dimension, attempting to access index {} " + "when only {} indexes exist", + i, + indices.at(i), + shape.at(legion_dim_t{i}))); + } + + offset += indices.at(i) * multiplier; + multiplier *= shape.at(legion_dim_t{i}); + } + + return offset; +} + class GenericTensorAccessorR { public: template @@ -57,23 +79,7 @@ class GenericTensorAccessorR { using T = real_type_t
; T const *data_ptr = static_cast(this->ptr); - - int offset = 0; - int multiplier = 1; - for (int i = 0; i < this->shape.num_dims(); i++) { - if (indices.at(i) >= this->shape.at(legion_dim_t{i})) { - throw mk_runtime_error( - fmt::format("In {} dimension, attempting to access index {} " - "when only {} indexes exist", - i, - indices.at(i), - this->shape.at(legion_dim_t{i}))); - } - - offset += indices.at(i) * multiplier; - multiplier *= this->shape.at(legion_dim_t{i}); - } - + int offset = calculate_accessor_offset(indices, this->shape); return data_ptr[offset]; } @@ -141,24 +147,8 @@ class GenericTensorAccessorW { } using T = real_type_t
; - T *data_ptr = static_cast(this->ptr); - int offset = 0; - int multiplier = 1; - for (int i = 0; i < this->shape.num_dims(); i++) { - if (indices.at(i) >= this->shape.at(legion_dim_t{i})) { - throw mk_runtime_error( - fmt::format("In {} dimension, attempting to access index {} " - "when only {} indexes exist", - i, - indices.at(i), - this->shape.at(legion_dim_t{i}))); - } - - offset += indices.at(i) * multiplier; - multiplier *= this->shape.at(legion_dim_t{i}); - } - + int offset = calculate_accessor_offset(indices, this->shape); return data_ptr[offset]; } @@ -179,24 +169,8 @@ class GenericTensorAccessorW { } using T = real_type_t
; - T const *data_ptr = static_cast(this->ptr); - int offset = 0; - int multiplier = 1; - for (int i = 0; i < this->shape.num_dims(); i++) { - if (indices.at(i) >= this->shape.at(legion_dim_t{i})) { - throw mk_runtime_error( - fmt::format("In {} dimension, attempting to access index {} " - "when only {} indexes exist", - i, - indices.at(i), - this->shape.at(legion_dim_t{i}))); - } - - offset += indices.at(i) * multiplier; - multiplier *= this->shape.at(legion_dim_t{i}); - } - + int offset = calculate_accessor_offset(indices, this->shape); return data_ptr[offset]; } diff --git a/lib/kernels/include/kernels/flat_kernels.h b/lib/kernels/include/kernels/flat_kernels.h index d60a1a5157..54839bd7fa 100644 --- a/lib/kernels/include/kernels/flat_kernels.h +++ b/lib/kernels/include/kernels/flat_kernels.h @@ -10,7 +10,7 @@ void forward_kernel(ffStream_t stream, GenericTensorAccessorR input, float *output_ptr); -void backward_kernel(cudaStream_t stream, +void backward_kernel(ffStream_t stream, GenericTensorAccessorR input, float const *output_grad_ptr, float *input_grad_ptr); diff --git a/lib/kernels/include/kernels/loss_function_kernels.h b/lib/kernels/include/kernels/loss_function_kernels.h index 9e0dbd4ba1..bab404f884 100644 --- a/lib/kernels/include/kernels/loss_function_kernels.h +++ b/lib/kernels/include/kernels/loss_function_kernels.h @@ -1,7 +1,7 @@ #ifndef _FLEXFLOW_KERNELS_INCLUDE_KERNELS_LOSS_FUNCTION_KERNELS_H #define _FLEXFLOW_KERNELS_INCLUDE_KERNELS_LOSS_FUNCTION_KERNELS_H -#include "device.h" +#include "kernels/device.h" namespace FlexFlow { diff --git a/lib/kernels/include/kernels/managed_ff_stream.h b/lib/kernels/include/kernels/managed_ff_stream.h index 26d5fb4911..7f103ea560 100644 --- a/lib/kernels/include/kernels/managed_ff_stream.h +++ b/lib/kernels/include/kernels/managed_ff_stream.h @@ -19,6 +19,7 @@ struct ManagedFFStream { ffStream_t const &raw_stream() const; +private: void cleanup(); private: diff --git a/lib/kernels/include/kernels/managed_per_device_ff_handle.h b/lib/kernels/include/kernels/managed_per_device_ff_handle.h index 035ea574de..9bd9370685 100644 --- a/lib/kernels/include/kernels/managed_per_device_ff_handle.h +++ b/lib/kernels/include/kernels/managed_per_device_ff_handle.h @@ -24,6 +24,7 @@ struct ManagedPerDeviceFFHandle { PerDeviceFFHandle const &raw_handle() const; +private: void cleanup(); private: diff --git a/lib/kernels/include/kernels/metrics_kernels.h b/lib/kernels/include/kernels/metrics_kernels.h index d961ee7503..430608db55 100644 --- a/lib/kernels/include/kernels/metrics_kernels.h +++ b/lib/kernels/include/kernels/metrics_kernels.h @@ -2,20 +2,20 @@ #define _FLEXFLOW_KERNELS_INCLUDE_KERNELS_METRICS_KERNELS_H #include "kernels/perf_metrics.h" -#include "pcg/metric.h" +#include "pcg/metric_attrs.h" namespace FlexFlow { void update_metrics_sparse_label_kernel_wrapper(float const *logit_ptr, int const *label_ptr, - MetricsAttrs const *me, + MetricsAttrs const &me, int num_effective_samples, int num_classes, PerfMetrics &perf_zc); void update_metrics_label_kernel_wrapper(float const *logit_ptr, float const *label_ptr, - MetricsAttrs const *me, + MetricsAttrs const &me, int num_samples, int num_classes, PerfMetrics &perf_zc); diff --git a/lib/kernels/include/kernels/pool_2d_kernels.h b/lib/kernels/include/kernels/pool_2d_kernels.h index ad0a52efb9..9650859a18 100644 --- a/lib/kernels/include/kernels/pool_2d_kernels.h +++ b/lib/kernels/include/kernels/pool_2d_kernels.h @@ -67,7 +67,7 @@ void forward_kernel(ffStream_t stream, void const *input_ptr, void *output_ptr); -void backward_kernel(cudaStream_t stream, +void backward_kernel(ffStream_t stream, Pool2DPerDeviceState const &m, void const *output_ptr, void const *output_grad_ptr, diff --git a/lib/kernels/src/cuda/metrics_functions.cu b/lib/kernels/src/cuda/metrics_functions.cu index 0250f829ec..112f84c90c 100644 --- a/lib/kernels/src/cuda/metrics_functions.cu +++ b/lib/kernels/src/cuda/metrics_functions.cu @@ -16,7 +16,7 @@ #include "device.h" #include "kernels/metrics_kernels.h" #include "kernels/perf_metrics.h" -#include "pcg/metric.h" +#include "pcg/metric_attrs.h" namespace FlexFlow { @@ -163,7 +163,7 @@ __global__ void update_metrics_label_kernel(float const *logits, void update_metrics_sparse_label_kernel_wrapper(float const *logit_ptr, int const *label_ptr, - MetricsAttrs const *me, + MetricsAttrs const &me, int num_effective_samples, int num_classes, PerfMetrics &perf_zc) { @@ -179,7 +179,7 @@ void update_metrics_sparse_label_kernel_wrapper(float const *logit_ptr, CUDA_NUM_THREADS, 0, stream>>>( - logit_ptr, label_ptr, perf_cuda, *me, num_effective_samples, num_classes); + logit_ptr, label_ptr, perf_cuda, me, num_effective_samples, num_classes); checkCUDA(cudaStreamSynchronize(stream)); checkCUDA(cudaMemcpy( &perf, perf_cuda, sizeof(CUDAPerfMetrics), cudaMemcpyDeviceToHost)); @@ -188,7 +188,7 @@ void update_metrics_sparse_label_kernel_wrapper(float const *logit_ptr, void update_metrics_label_kernel_wrapper(float const *logit_ptr, float const *label_ptr, - MetricsAttrs const *me, + MetricsAttrs const &me, int num_samples, int num_classes, PerfMetrics &perf_zc) { @@ -201,7 +201,7 @@ void update_metrics_label_kernel_wrapper(float const *logit_ptr, cudaStream_t stream; checkCUDA(get_legion_stream(&stream)); update_metrics_label_kernel<<>>( - logit_ptr, label_ptr, perf_cuda, *me, num_samples, num_classes); + logit_ptr, label_ptr, perf_cuda, me, num_samples, num_classes); checkCUDA(cudaStreamSynchronize(stream)); checkCUDA(cudaMemcpy( &perf, perf_cuda, sizeof(CUDAPerfMetrics), cudaMemcpyDeviceToHost)); diff --git a/lib/kernels/src/hip/embedding_kernels.cpp b/lib/kernels/src/hip/embedding_kernels.cpp index 7ca3149f2f..06b42d420a 100644 --- a/lib/kernels/src/hip/embedding_kernels.cpp +++ b/lib/kernels/src/hip/embedding_kernels.cpp @@ -364,8 +364,8 @@ struct ForwardKernel { weight.data_type == DataType::FLOAT || weight.data_type == DataType::DOUBLE); - if (aggr == AggregateOp::NONE) { - hipLaunchKernelGGL(HIP_KERNEL_NAME(embed_forward_no_aggr), + if (aggr == AggregateOp::AVG || aggr == AggregateOp::SUM) { + hipLaunchKernelGGL(HIP_KERNEL_NAME(embed_forward_with_aggr), GET_BLOCKS(output.shape.get_volume()), CUDA_NUM_THREADS, 0, @@ -374,10 +374,11 @@ struct ForwardKernel { output.get(), weight.get(), out_dim, - batch_size); + in_dim, + batch_size, + aggr); } else { - assert(aggr == AggregateOp::AVG || aggr == AggregateOp::SUM); - hipLaunchKernelGGL(HIP_KERNEL_NAME(embed_forward_with_aggr), + hipLaunchKernelGGL(HIP_KERNEL_NAME(embed_forward_no_aggr), GET_BLOCKS(output.shape.get_volume()), CUDA_NUM_THREADS, 0, @@ -386,9 +387,7 @@ struct ForwardKernel { output.get(), weight.get(), out_dim, - in_dim, - batch_size, - aggr); + batch_size); } } } @@ -408,8 +407,9 @@ struct BackwardKernel { assert(output.data_type == DataType::HALF || output.data_type == DataType::FLOAT || output.data_type == DataType::DOUBLE); - if (aggr == AggregateOp::NONE) { - hipLaunchKernelGGL(HIP_KERNEL_NAME(embed_backward_no_aggr), + + if (aggr == AggregateOp::AVG || aggr == AggregateOp::SUM) { + hipLaunchKernelGGL(HIP_KERNEL_NAME(embed_backward_with_aggr), GET_BLOCKS(output.shape.get_volume()), CUDA_NUM_THREADS, 0, @@ -418,9 +418,11 @@ struct BackwardKernel { output.get(), weight_grad.get(), out_dim, - batch_size); + in_dim, + batch_size, + aggr); } else { - hipLaunchKernelGGL(HIP_KERNEL_NAME(embed_backward_with_aggr), + hipLaunchKernelGGL(HIP_KERNEL_NAME(embed_backward_no_aggr), GET_BLOCKS(output.shape.get_volume()), CUDA_NUM_THREADS, 0, @@ -429,9 +431,7 @@ struct BackwardKernel { output.get(), weight_grad.get(), out_dim, - in_dim, - batch_size, - aggr); + batch_size); } } } diff --git a/lib/kernels/test/src/test_batch_norm_kernel.cc b/lib/kernels/test/src/test_batch_norm_kernel.cc index 03a3a1ad40..270fad7bb6 100644 --- a/lib/kernels/test/src/test_batch_norm_kernel.cc +++ b/lib/kernels/test/src/test_batch_norm_kernel.cc @@ -1,6 +1,6 @@ #include "doctest/doctest.h" #include "kernels/batch_norm_kernels.h" -#include "op-attrs/make_datatype_value.h" +#include "op-attrs/datatype_value.h" #include "test_utils.h" using namespace ::FlexFlow; diff --git a/lib/kernels/test/src/test_concat_kernel.cc b/lib/kernels/test/src/test_concat_kernel.cc index 4607171a54..5447b12fc5 100644 --- a/lib/kernels/test/src/test_concat_kernel.cc +++ b/lib/kernels/test/src/test_concat_kernel.cc @@ -8,7 +8,7 @@ TEST_SUITE(FF_TEST_SUITE) { TEST_CASE("Test concat kernel forward and backward") { size_t num_inputs = 2; size_t size_per_input = 10; - ff_dim_t concat_axis = ff_dim_t{1}; + ff_dim_t concat_axis = ff_dim_t{nonnegative_int{1}}; ManagedPerDeviceFFHandle managed_handle{ /*workSpaceSize=*/1024 * 1024, diff --git a/lib/kernels/test/src/test_flat_kernel.cc b/lib/kernels/test/src/test_flat_kernel.cc index 0bb69aa1dc..bbeb349ced 100644 --- a/lib/kernels/test/src/test_flat_kernel.cc +++ b/lib/kernels/test/src/test_flat_kernel.cc @@ -1,6 +1,6 @@ #include "doctest/doctest.h" #include "kernels/flat_kernels.h" -#include "op-attrs/make_datatype_value.h" +#include "op-attrs/datatype_value.h" #include "test_utils.h" using namespace ::FlexFlow; diff --git a/lib/kernels/test/src/test_layer_norm_kernels.cc b/lib/kernels/test/src/test_layer_norm_kernels.cc index 7d7298f83d..80a046fe37 100644 --- a/lib/kernels/test/src/test_layer_norm_kernels.cc +++ b/lib/kernels/test/src/test_layer_norm_kernels.cc @@ -1,6 +1,6 @@ #include "doctest/doctest.h" #include "kernels/layer_norm_kernels.h" -#include "op-attrs/make_datatype_value.h" +#include "op-attrs/datatype_value.h" #include "test_utils.h" using namespace ::FlexFlow; diff --git a/lib/kernels/test/src/test_managed_per_device_ff_handle.cc b/lib/kernels/test/src/test_managed_per_device_ff_handle.cc index de3e5b72b1..d081a0b07c 100644 --- a/lib/kernels/test/src/test_managed_per_device_ff_handle.cc +++ b/lib/kernels/test/src/test_managed_per_device_ff_handle.cc @@ -5,7 +5,8 @@ using namespace ::FlexFlow; TEST_SUITE(FF_TEST_SUITE) { TEST_CASE("ManagedPerDeviceFFHandle") { - ManagedPerDeviceFFHandle base_handle{1024 * 1024, true}; + ManagedPerDeviceFFHandle base_handle{/*workSpaceSize=*/1024 * 1024, + /*allowTensorOpMathConversion=*/true}; PerDeviceFFHandle const *base_handle_ptr = &base_handle.raw_handle(); SUBCASE("constructor") { @@ -22,7 +23,9 @@ TEST_SUITE(FF_TEST_SUITE) { SUBCASE("move assignment operator") { SUBCASE("move assign to other") { - ManagedPerDeviceFFHandle new_handle{1024 * 1024, true}; + ManagedPerDeviceFFHandle new_handle{ + /*workSpaceSize=*/1024 * 1024, + /*allowTensorOpMathConversion=*/true}; new_handle = std::move(base_handle); CHECK(&base_handle.raw_handle() == nullptr); diff --git a/lib/kernels/test/src/test_partition_kernel.cc b/lib/kernels/test/src/test_partition_kernel.cc index e88c811803..25264b7a58 100644 --- a/lib/kernels/test/src/test_partition_kernel.cc +++ b/lib/kernels/test/src/test_partition_kernel.cc @@ -1,6 +1,6 @@ #include "doctest/doctest.h" #include "kernels/partition_kernels.h" -#include "op-attrs/make_datatype_value.h" +#include "op-attrs/datatype_value.h" #include "test_utils.h" using namespace ::FlexFlow; diff --git a/lib/kernels/test/src/test_pool_2d_kernels.cc b/lib/kernels/test/src/test_pool_2d_kernels.cc index 00fa968235..eb0702a970 100644 --- a/lib/kernels/test/src/test_pool_2d_kernels.cc +++ b/lib/kernels/test/src/test_pool_2d_kernels.cc @@ -1,6 +1,6 @@ #include "doctest/doctest.h" #include "kernels/pool_2d_kernels.h" -#include "op-attrs/make_datatype_value.h" +#include "op-attrs/datatype_value.h" #include "test_utils.h" using namespace ::FlexFlow; diff --git a/lib/kernels/test/src/test_reduction_kernel.cc b/lib/kernels/test/src/test_reduction_kernel.cc index 1c389cb20d..a33748c0de 100644 --- a/lib/kernels/test/src/test_reduction_kernel.cc +++ b/lib/kernels/test/src/test_reduction_kernel.cc @@ -1,6 +1,6 @@ #include "doctest/doctest.h" #include "kernels/reduction_kernels.h" -#include "op-attrs/make_datatype_value.h" +#include "op-attrs/datatype_value.h" #include "test_utils.h" using namespace ::FlexFlow; diff --git a/lib/kernels/test/src/test_reverse_kernels.cc b/lib/kernels/test/src/test_reverse_kernels.cc index 4adf79847a..c06919d603 100644 --- a/lib/kernels/test/src/test_reverse_kernels.cc +++ b/lib/kernels/test/src/test_reverse_kernels.cc @@ -1,7 +1,7 @@ #include "doctest/doctest.h" #include "kernels/reverse_kernels.h" #include "kernels/reverse_kernels_cpu.h" -#include "op-attrs/make_datatype_value.h" +#include "op-attrs/datatype_value.h" #include "test_utils.h" using namespace ::FlexFlow; diff --git a/lib/kernels/test/src/test_split_kernel.cc b/lib/kernels/test/src/test_split_kernel.cc index 34993fa151..e94d102b71 100644 --- a/lib/kernels/test/src/test_split_kernel.cc +++ b/lib/kernels/test/src/test_split_kernel.cc @@ -1,6 +1,6 @@ #include "doctest/doctest.h" #include "kernels/split_kernels.h" -#include "op-attrs/make_datatype_value.h" +#include "op-attrs/datatype_value.h" #include "test_utils.h" #include "utils/containers/repeat.h" diff --git a/lib/kernels/test/src/test_transpose_kernel.cc b/lib/kernels/test/src/test_transpose_kernel.cc index 0bc85cb8e0..f87fb67921 100644 --- a/lib/kernels/test/src/test_transpose_kernel.cc +++ b/lib/kernels/test/src/test_transpose_kernel.cc @@ -7,7 +7,8 @@ TEST_SUITE(FF_TEST_SUITE) { TEST_CASE("Test Transpose Kernel Operations") { std::size_t num_dims = 2; - std::vector perm = {ff_dim_t{0}, ff_dim_t{1}}; + std::vector perm = {ff_dim_t{nonnegative_int{0}}, + ff_dim_t{nonnegative_int{1}}}; ManagedPerDeviceFFHandle managed_handle{ /*workSpaceSize=*/1024 * 1024, diff --git a/lib/kernels/test/src/test_utils.cc b/lib/kernels/test/src/test_utils.cc index bfed1241ba..c75abd50ff 100644 --- a/lib/kernels/test/src/test_utils.cc +++ b/lib/kernels/test/src/test_utils.cc @@ -1,5 +1,6 @@ #include "test_utils.h" #include "op-attrs/tensor_shape.h" +#include "utils/join_strings.h" #include namespace FlexFlow { @@ -140,21 +141,16 @@ template struct Print2DCPUAccessorR { void operator()(GenericTensorAccessorR const &accessor, std::ostream &stream) { - using T = real_type_t
; - - T const *data_ptr = accessor.get
(); int rows = accessor.shape.at(legion_dim_t{0}); int cols = accessor.shape.at(legion_dim_t{1}); - for (int i = 0; i < rows; i++) { - for (int j = 0; j < cols; j++) { - stream << data_ptr[i * cols + j]; + std::vector indices(cols); + std::iota(indices.begin(), indices.end(), 0); - if (j < cols - 1) { - stream << " "; - } - } - stream << std::endl; + for (int i = 0; i < rows; i++) { + stream << join_strings(indices, " ", [&](int k) { + return accessor.at
({i, k}); + }) << std::endl; } } }; diff --git a/lib/kernels/test/src/test_utils.h b/lib/kernels/test/src/test_utils.h index 19599d2900..a41bfc3aff 100644 --- a/lib/kernels/test/src/test_utils.h +++ b/lib/kernels/test/src/test_utils.h @@ -15,7 +15,7 @@ #include #include -using namespace FlexFlow; +using namespace ::FlexFlow; GenericTensorAccessorW create_random_filled_accessor_w(TensorShape const &shape, Allocator &allocator); diff --git a/lib/local-execution/test/src/test_local_cost_estimator.cc b/lib/local-execution/test/src/test_local_cost_estimator.cc index 512c1ef33b..9f8b4092c1 100644 --- a/lib/local-execution/test/src/test_local_cost_estimator.cc +++ b/lib/local-execution/test/src/test_local_cost_estimator.cc @@ -13,71 +13,70 @@ // TEST_CASE("Local Cost Estimator") { // // local backing initialization // ManagedPerDeviceFFHandle managed_handle{ -/*workSpaceSize=*/1024 * 1024, - /*allowTensorOpMathConversion=*/true -} -; +// /*workSpaceSize=*/1024 * 1024, +// /*allowTensorOpMathConversion=*/true}; -// RuntimeArgConfig runtime_arg_config = RuntimeArgConfig{ -// DeviceSpecific::create(managed_handle.raw_handle()), -// EnableProfiling::YES, -// ProfilingSettings{/*warmup_iters=*/0, -// /*measure_iters=*/1}}; +// RuntimeArgConfig runtime_arg_config = RuntimeArgConfig{ +// DeviceSpecific::create(managed_handle.raw_handle()), +// EnableProfiling::YES, +// ProfilingSettings{/*warmup_iters=*/0, +// /*measure_iters=*/1}}; -// LocalCostEstimator cost_estimator = -// LocalCostEstimator{runtime_arg_config}; +// LocalCostEstimator cost_estimator = +// LocalCostEstimator{runtime_arg_config}; -// SUBCASE("Estimate cost -- Attention Op") { -// int embed_dim = 32; -// int num_heads = 10; -// MultiHeadAttentionAttrs attrs = MultiHeadAttentionAttrs{ -// /*embed_dim=*/embed_dim, -// /*num_heads=*/num_heads, -// /*kdim=*/embed_dim, -// /*vdim=*/embed_dim, -// /*dropout=*/0.0, -// /*bias=*/true, -// /*add_bias_kv=*/false, -// /*add_zero_attn=*/false, -// }; +// SUBCASE("Estimate cost -- Attention Op") { +// int embed_dim = 32; +// int num_heads = 10; +// MultiHeadAttentionAttrs attrs = MultiHeadAttentionAttrs{ +// /*embed_dim=*/embed_dim, +// /*num_heads=*/num_heads, +// /*kdim=*/embed_dim, +// /*vdim=*/embed_dim, +// /*dropout=*/0.0, +// /*bias=*/true, +// /*add_bias_kv=*/false, +// /*add_zero_attn=*/false, +// }; -// size_t batch_size = 40; -// size_t seq_len = 48; -// size_t feature_size = 36; +// size_t batch_size = 40; +// size_t seq_len = 48; +// size_t feature_size = 36; -// DataType dtype = DataType::FLOAT; -// ParallelTensorShape inputs_shape = lift_to_parallel(TensorShape{ -// TensorDims{FFOrdered{batch_size, seq_len, feature_size}}, -// DataType::FLOAT, -// }); +// DataType dtype = DataType::FLOAT; +// ParallelTensorShape inputs_shape = lift_to_parallel(TensorShape{ +// TensorDims{FFOrdered{batch_size, seq_len, +// feature_size}}, DataType::FLOAT, +// }); -// ParallelTensorShape weights_shape = throw_if_unexpected( -// get_weights_shape(attrs, inputs_shape, inputs_shape, -// inputs_shape)); -// ParallelTensorAttrs weight_attrs = -// ParallelTensorAttrs{weights_shape, -// /*sync_type=*/std::nullopt, -// /*initializer=*/std::nullopt, -// CreateGrad::YES}; +// ParallelTensorShape weights_shape = throw_if_unexpected( +// get_weights_shape(attrs, inputs_shape, inputs_shape, +// inputs_shape)); +// ParallelTensorAttrs weight_attrs = +// ParallelTensorAttrs{weights_shape, +// /*sync_type=*/std::nullopt, +// /*initializer=*/std::nullopt, +// CreateGrad::YES}; -// ParallelTensorShape output_shape = throw_if_unexpected( -// get_output_shape(attrs, inputs_shape, inputs_shape, inputs_shape)); -// ParallelTensorAttrs output_attrs = -// ParallelTensorAttrs{output_shape, -// /*sync_type=*/std::nullopt, -// /*initializer=*/std::nullopt, -// CreateGrad::YES}; +// ParallelTensorShape output_shape = throw_if_unexpected( +// get_output_shape(attrs, inputs_shape, inputs_shape, +// inputs_shape)); +// ParallelTensorAttrs output_attrs = +// ParallelTensorAttrs{output_shape, +// /*sync_type=*/std::nullopt, +// /*initializer=*/std::nullopt, +// CreateGrad::YES}; -// CostDetails result = cost_estimator.estimate_cost( -// PCGOperatorAttrs{attrs}, -// std::vector{ -// inputs_shape, inputs_shape, inputs_shape}, -// std::vector{weight_attrs}, -// std::vector{output_attrs}, -// make_1d_machine_view(gpu_id_t{0}, gpu_id_t{1})); +// CostDetails result = cost_estimator.estimate_cost( +// PCGOperatorAttrs{attrs}, +// std::vector{ +// inputs_shape, inputs_shape, inputs_shape}, +// std::vector{weight_attrs}, +// std::vector{output_attrs}, +// make_1d_machine_view(gpu_id_t{0}, gpu_id_t{1})); -// CHECK(result.total_elapsed_time > 0); -// CHECK(result.total_mem_usage > 0); +// CHECK(result.total_elapsed_time > 0); +// CHECK(result.total_mem_usage > 0); +// } +// } // } -// } -// } diff --git a/lib/op-attrs/include/op-attrs/aggregate_op.enum.toml b/lib/op-attrs/include/op-attrs/aggregate_op.enum.toml index 2c524c120a..09ee99915d 100644 --- a/lib/op-attrs/include/op-attrs/aggregate_op.enum.toml +++ b/lib/op-attrs/include/op-attrs/aggregate_op.enum.toml @@ -13,5 +13,3 @@ name = "SUM" [[values]] name = "AVG" -[[values]] -name = "NONE" diff --git a/lib/op-attrs/include/op-attrs/make_datatype_value.h b/lib/op-attrs/include/op-attrs/make_datatype_value.h deleted file mode 100644 index af4792dd9e..0000000000 --- a/lib/op-attrs/include/op-attrs/make_datatype_value.h +++ /dev/null @@ -1,16 +0,0 @@ -#ifndef _FLEXFLOW_LIB_OP_ATTRS_INCLUDE_OP_ATTRS_MAKE_DATATYPE_VALUE_H -#define _FLEXFLOW_LIB_OP_ATTRS_INCLUDE_OP_ATTRS_MAKE_DATATYPE_VALUE_H - -#include "op-attrs/datatype_value.dtg.h" - -namespace FlexFlow { - -DataTypeValue make_float_data_type_value(float value); -DataTypeValue make_double_data_type_value(double value); -DataTypeValue make_int32_data_type_value(int32_t value); -DataTypeValue make_int64_data_type_value(int64_t value); -DataTypeValue make_bool_data_type_value(bool value); - -} // namespace FlexFlow - -#endif // _FLEXFLOW_LIB_OP_ATTRS_INCLUDE_OP_ATTRS_MAKE_DATATYPE_VALUE_H diff --git a/lib/op-attrs/src/op-attrs/make_datatype_value.cc b/lib/op-attrs/src/op-attrs/datatype_value.cc similarity index 92% rename from lib/op-attrs/src/op-attrs/make_datatype_value.cc rename to lib/op-attrs/src/op-attrs/datatype_value.cc index 76d712949a..4604ef0b4e 100644 --- a/lib/op-attrs/src/op-attrs/make_datatype_value.cc +++ b/lib/op-attrs/src/op-attrs/datatype_value.cc @@ -1,4 +1,4 @@ -#include "op-attrs/make_datatype_value.h" +#include "op-attrs/datatype_value.h" namespace FlexFlow { diff --git a/lib/op-attrs/test/src/op-attrs/datatype_value.cc b/lib/op-attrs/test/src/op-attrs/datatype_value.cc new file mode 100644 index 0000000000..9b0e90b601 --- /dev/null +++ b/lib/op-attrs/test/src/op-attrs/datatype_value.cc @@ -0,0 +1,68 @@ +#include "op-attrs/datatype_value.h" +#include + +using namespace ::FlexFlow; + +TEST_SUITE(FF_TEST_SUITE) { + TEST_CASE("test make_data_type_value") { + SUBCASE("make_float_data_type_value") { + float value = 1.0f; + DataTypeValue data_type_value = make_float_data_type_value(value); + + CHECK(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK(data_type_value.get() == value); + } + + SUBCASE("make_double_data_type_value") { + double value = 2.71828; + DataTypeValue data_type_value = make_double_data_type_value(value); + + CHECK(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK(data_type_value.get() == value); + } + + SUBCASE("make_int32_data_type_value") { + int32_t value = -42; + DataTypeValue data_type_value = make_int32_data_type_value(value); + + CHECK(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK(data_type_value.get() == value); + } + + SUBCASE("make_int64_data_type_value") { + int64_t value = 1LL << 40; + DataTypeValue data_type_value = make_int64_data_type_value(value); + + CHECK(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK(data_type_value.get() == value); + } + + SUBCASE("make_bool_data_type_value") { + bool value = true; + DataTypeValue data_type_value = make_bool_data_type_value(value); + + CHECK(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK_FALSE(data_type_value.has()); + CHECK(data_type_value.get() == value); + } + } +} diff --git a/lib/pcg/include/pcg/metric.enum.toml b/lib/pcg/include/pcg/metric.enum.toml new file mode 100644 index 0000000000..ebb2323203 --- /dev/null +++ b/lib/pcg/include/pcg/metric.enum.toml @@ -0,0 +1,26 @@ +namespace = "FlexFlow" +name = "Metric" +features = [ + "hash", + "json", + "rapidcheck", + "fmt", +] + +[[values]] +name = "ACCURACY" + +[[values]] +name = "CATEGORICAL_CROSSENTROPY" + +[[values]] +name = "SPARSE_CATEGORICAL_CROSSENTROPY" + +[[values]] +name = "MEAN_SQUARED_ERROR" + +[[values]] +name = "ROOT_MEAN_SQUARED_ERROR" + +[[values]] +name = "MEAN_ABSOLUTE_ERROR" diff --git a/lib/pcg/include/pcg/metric.h b/lib/pcg/include/pcg/metric.h deleted file mode 100644 index 718919112f..0000000000 --- a/lib/pcg/include/pcg/metric.h +++ /dev/null @@ -1,72 +0,0 @@ -#ifndef _FF_METRICS_H_ -#define _FF_METRICS_H_ - -#include "op-attrs/ops/loss_functions/loss_functions.h" -#include "utils/fmt.h" -#include - -namespace FlexFlow { - -enum class Metric { - ACCURACY, - CATEGORICAL_CROSSENTROPY, - SPARSE_CATEGORICAL_CROSSENTROPY, - MEAN_SQUARED_ERROR, - ROOT_MEAN_SQUARED_ERROR, - MEAN_ABSOLUTE_ERROR, -}; - -class MetricsAttrs { -public: - MetricsAttrs() = delete; - MetricsAttrs(LossFunction, std::vector const &); - -public: - LossFunction loss_type; - bool measure_accuracy; - bool measure_categorical_crossentropy; - bool measure_sparse_categorical_crossentropy; - bool measure_mean_squared_error; - bool measure_root_mean_squared_error; - bool measure_mean_absolute_error; -}; - -} // namespace FlexFlow - -namespace fmt { - -template <> -struct formatter<::FlexFlow::Metric> : formatter { - template - auto format(::FlexFlow::Metric m, FormatContext &ctx) const - -> decltype(ctx.out()) { - using namespace FlexFlow; - - string_view name = "unknown"; - switch (m) { - case Metric::ACCURACY: - name = "Accuracy"; - break; - case Metric::CATEGORICAL_CROSSENTROPY: - name = "CategoricalCrossEntropy"; - break; - case Metric::SPARSE_CATEGORICAL_CROSSENTROPY: - name = "SparseCategoricalCrossEntropy"; - break; - case Metric::MEAN_SQUARED_ERROR: - name = "MeanSquaredError"; - break; - case Metric::ROOT_MEAN_SQUARED_ERROR: - name = "RootMeanSquaredError"; - break; - case Metric::MEAN_ABSOLUTE_ERROR: - name = "MeanAbsoluteError"; - break; - } - return formatter::format(name, ctx); - } -}; - -} // namespace fmt - -#endif diff --git a/lib/pcg/include/pcg/metric_attrs.h b/lib/pcg/include/pcg/metric_attrs.h new file mode 100644 index 0000000000..343c2154dd --- /dev/null +++ b/lib/pcg/include/pcg/metric_attrs.h @@ -0,0 +1,28 @@ +#ifndef _FF_METRICS_H_ +#define _FF_METRICS_H_ + +#include "op-attrs/ops/loss_functions/loss_functions.h" +#include "pcg/metric.dtg.h" +#include "utils/fmt.h" +#include + +namespace FlexFlow { + +class MetricsAttrs { +public: + MetricsAttrs() = delete; + MetricsAttrs(LossFunction, std::unordered_set const &); + +public: + LossFunction loss_type; + bool measure_accuracy; + bool measure_categorical_crossentropy; + bool measure_sparse_categorical_crossentropy; + bool measure_mean_squared_error; + bool measure_root_mean_squared_error; + bool measure_mean_absolute_error; +}; + +} // namespace FlexFlow + +#endif diff --git a/lib/pcg/include/pcg/strided_rectangle.h b/lib/pcg/include/pcg/strided_rectangle.h deleted file mode 100644 index 9c3b8eeda9..0000000000 --- a/lib/pcg/include/pcg/strided_rectangle.h +++ /dev/null @@ -1,17 +0,0 @@ -#ifndef _FLEXFLOW_PCG_INCLUDE_PCG_STRIDED_RECTANGLE_H -#define _FLEXFLOW_PCG_INCLUDE_PCG_STRIDED_RECTANGLE_H - -#include "op-attrs/ff_dim.dtg.h" -#include "pcg/side_size_t.dtg.h" -#include "pcg/strided_rectangle.dtg.h" - -namespace FlexFlow { - -size_t get_num_dims(StridedRectangle const &); -StridedRectangleSide get_side_at_idx(StridedRectangle const &rect, - ff_dim_t const &idx); -num_points_t get_num_points(StridedRectangle const &rect); - -} // namespace FlexFlow - -#endif diff --git a/lib/pcg/src/pcg/computation_graph_builder.cc b/lib/pcg/src/pcg/computation_graph_builder.cc index 7ff5bec2f7..09772fa9d9 100644 --- a/lib/pcg/src/pcg/computation_graph_builder.cc +++ b/lib/pcg/src/pcg/computation_graph_builder.cc @@ -1,9 +1,9 @@ #include "pcg/computation_graph_builder.h" #include "op-attrs/computation_graph_op_attrs.h" +#include "op-attrs/datatype_value.h" #include "op-attrs/get_incoming_tensor_roles.h" #include "op-attrs/get_op_type.h" #include "op-attrs/get_output_shapes.h" -#include "op-attrs/make_datatype_value.h" #include "op-attrs/ops/attention.h" #include "op-attrs/ops/batch_norm.h" #include "op-attrs/ops/broadcast.h" diff --git a/lib/pcg/src/pcg/metric.cc b/lib/pcg/src/pcg/metric.cc index 69aba90d12..9a93e75350 100644 --- a/lib/pcg/src/pcg/metric.cc +++ b/lib/pcg/src/pcg/metric.cc @@ -1,8 +1,8 @@ -#include "pcg/metric.h" +#include "pcg/metric_attrs.h" namespace FlexFlow { MetricsAttrs::MetricsAttrs(LossFunction _loss_type, - std::vector const &metrics) + std::unordered_set const &metrics) : loss_type(_loss_type), measure_accuracy(false), measure_categorical_crossentropy(false), measure_sparse_categorical_crossentropy(false), @@ -29,8 +29,8 @@ MetricsAttrs::MetricsAttrs(LossFunction _loss_type, measure_mean_absolute_error = true; continue; default: - throw mk_runtime_error( - "Initializing MetricsAttrs with unrecogonized metrics type"); + throw mk_runtime_error(fmt::format( + "Initializing MetricsAttrs with unrecogonized metrics type {}", m)); } } } diff --git a/lib/pcg/src/pcg/parallel_computation_graph/parallel_computation_graph_builder.cc b/lib/pcg/src/pcg/parallel_computation_graph/parallel_computation_graph_builder.cc index 79ac43ae66..e2f4555328 100644 --- a/lib/pcg/src/pcg/parallel_computation_graph/parallel_computation_graph_builder.cc +++ b/lib/pcg/src/pcg/parallel_computation_graph/parallel_computation_graph_builder.cc @@ -1,6 +1,6 @@ #include "pcg/parallel_computation_graph/parallel_computation_graph_builder.h" +#include "op-attrs/datatype_value.h" #include "op-attrs/get_incoming_tensor_roles.h" -#include "op-attrs/make_datatype_value.h" #include "op-attrs/ops/attention.h" #include "op-attrs/ops/batch_matmul.h" #include "op-attrs/ops/batch_norm.h" diff --git a/lib/pcg/src/pcg/strided_rectangle_side.cc b/lib/pcg/src/pcg/strided_rectangle_side.cc deleted file mode 100644 index e6caf4cb86..0000000000 --- a/lib/pcg/src/pcg/strided_rectangle_side.cc +++ /dev/null @@ -1,17 +0,0 @@ -#include "pcg/strided_rectangle_side.h" -#include "utils/exception.h" - -namespace FlexFlow { - -StridedRectangleSide strided_side_from_size_and_stride(side_size_t side_size, - int stride) { - assert((side_size.unwrapped % stride) == 0); - return StridedRectangleSide{num_points_t{side_size.unwrapped / stride}, - stride}; -} - -side_size_t get_side_size(StridedRectangleSide const &s) { - return side_size_t{s.num_points.unwrapped * s.stride}; -} - -} // namespace FlexFlow diff --git a/lib/pcg/src/strided_rectangle.cc b/lib/pcg/src/strided_rectangle.cc deleted file mode 100644 index 1c61424ab9..0000000000 --- a/lib/pcg/src/strided_rectangle.cc +++ /dev/null @@ -1,35 +0,0 @@ -#include "pcg/strided_rectangle.h" -#include "op-attrs/dim_ordered/transform.h" -#include "utils/containers.h" - -namespace FlexFlow { - -/* size_t StridedRectangle::at(FFOrdered const &coord) const { */ -/* assert(coord.size() == this->num_dims()); */ - -/* size_t _1d_stride = 1; */ -/* size_t idx = 0; */ -/* for (auto dim : inner_to_outer_idxs(this->sides)) { */ -/* idx += this->sides.at(dim).at(coord.at(dim)).value() * _1d_stride; */ -/* _1d_stride *= this->sides.at(dim).get_size().value(); */ -/* } */ -/* return idx; */ -/* } */ - -size_t get_num_dims(StridedRectangle const &rect) { - return rect.sides.size(); -} - -num_points_t get_num_points(StridedRectangle const &rect) { - return num_points_t{ - product(transform(rect.sides, [](StridedRectangleSide const &side) { - return side.num_points.unwrapped; - }))}; -} - -StridedRectangleSide get_side_at_idx(StridedRectangle const &rect, - ff_dim_t const &idx) { - return rect.sides.at(idx); -} - -} // namespace FlexFlow diff --git a/lib/pcg/test/src/test_machine_view.cc b/lib/pcg/test/src/test_machine_view.cc deleted file mode 100644 index 92a96d5e9a..0000000000 --- a/lib/pcg/test/src/test_machine_view.cc +++ /dev/null @@ -1,74 +0,0 @@ -#include "doctest/doctest.h" -#include "pcg/machine_view.h" -#include "pcg/strided_rectangle.h" -#include "pcg/strided_rectangle_side.h" - -TEST_SUITE(FF_TEST_SUITE) { - TEST_CASE("MachineView general util functions") { - StridedRectangle rect{{StridedRectangleSide{num_points_t{7}, 5}, - StridedRectangleSide{num_points_t{10}, 2}}}; - gpu_id_t start(1); - MachineView mv{device_id_t{start}, rect}; - SUBCASE("num_dims") { - CHECK(num_dims(mv) == 2); - } - SUBCASE("num_devices") { - CHECK(num_devices(mv) == 7 * 10); - } - SUBCASE("get_device_type") { - CHECK(get_device_type(mv) == DeviceType::GPU); - } - } - - TEST_CASE("MachineView make_1d_machine_view - GPU") { - StridedRectangle rect{{StridedRectangleSide{num_points_t{7}, 5}}}; - device_id_t start_gpu{gpu_id_t{1}}; - MachineView gpu_mv{start_gpu, rect}; - - SUBCASE("make_1d_machine_view(gpu_id_t start, gpu_id_t stop, int stride)") { - MachineView result = - make_1d_machine_view(start_gpu, device_id_t{gpu_id_t(1 + 7 * 5)}, 5); - MachineView correct = gpu_mv; - CHECK(result == correct); - } - SUBCASE("make_1d_machine_view(gpu_id_t start, num_points_t num_points, int " - "stride)") { - MachineView result = make_1d_machine_view(start_gpu, num_points_t{7}, 5); - MachineView correct = gpu_mv; - CHECK(result == correct); - } - SUBCASE("make_1d_machine_view(gpu_id_t start, side_size_t interval_size, " - "int stride)") { - MachineView result = make_1d_machine_view( - start_gpu, get_side_size(rect.sides.at(ff_dim_t{0})), 5); - MachineView correct = gpu_mv; - CHECK(result == correct); - } - } - - TEST_CASE("MachineView make_1d_machine_view - CPU") { - StridedRectangle rect{{StridedRectangleSide{num_points_t{11}, 4}}}; - device_id_t start_cpu{cpu_id_t{2}}; - MachineView cpu_mv{start_cpu, rect}; - - SUBCASE("make_1d_machine_view(cpu_id_t start, cpu_id_t stop, int stride)") { - MachineView result = - make_1d_machine_view(start_cpu, device_id_t{cpu_id_t(2 + 11 * 4)}, 4); - MachineView correct = cpu_mv; - CHECK(result == correct); - } - SUBCASE("make_1d_machine_view(cpu_id_t start, num_points_t num_points, int " - "stride)") { - MachineView result = make_1d_machine_view(start_cpu, num_points_t{11}, 4); - MachineView correct = cpu_mv; - CHECK(result == correct); - } - SUBCASE("make_1d_machine_view(cpu_id_t start, side_size_t interval_size, " - "int stride)") { - MachineView result = make_1d_machine_view( - start_cpu, get_side_size(rect.sides.at(ff_dim_t{0})), 4); - MachineView correct = cpu_mv; - CHECK(result == correct); - } - } -} diff --git a/lib/pcg/test/src/test_strided_rectangle.cc b/lib/pcg/test/src/test_strided_rectangle.cc deleted file mode 100644 index ef342944de..0000000000 --- a/lib/pcg/test/src/test_strided_rectangle.cc +++ /dev/null @@ -1,37 +0,0 @@ -#include "doctest/doctest.h" -#include "pcg/strided_rectangle.h" -#include "pcg/strided_rectangle_side.h" - -TEST_SUITE(FF_TEST_SUITE) { - TEST_CASE("get_side_size(StridedRectangleSide)") { - StridedRectangleSide side{num_points_t{7}, 5}; - - CHECK(get_side_size(side) == side_size_t{7 * 5}); - } - TEST_CASE("strided_side_from_size_and_stride") { - StridedRectangleSide correct{num_points_t{10}, 3}; - StridedRectangleSide result = - strided_side_from_size_and_stride(side_size_t{10 * 3}, 3); - CHECK(result == correct); - } - - TEST_CASE("StridedRectangle - helper functions") { - - StridedRectangleSide s0{num_points_t{7}, 5}; - StridedRectangleSide s1{num_points_t{10}, 2}; - StridedRectangleSide s2{num_points_t{8}, 1}; - StridedRectangle rect{{s0, s1, s2}}; - - SUBCASE("get_num_dims") { - CHECK(get_num_dims(rect) == 3); - } - SUBCASE("get_num_points") { - CHECK(get_num_points(rect) == num_points_t{7 * 8 * 10}); - } - SUBCASE("get_side_at_idx") { - CHECK(get_side_at_idx(rect, ff_dim_t{0}) == s0); - CHECK(get_side_at_idx(rect, ff_dim_t{1}) == s1); - CHECK(get_side_at_idx(rect, ff_dim_t{2}) == s2); - } - } -}