From 041aa70cc64e661d387a0e4e289430b3af16d501 Mon Sep 17 00:00:00 2001 From: Ammar ELWazir Date: Sat, 15 Jul 2023 20:27:11 +0000 Subject: [PATCH] SWDEV-411283: Adding ROCTx Test Change-Id: Icf99ef584ff8ff38e16bc731951a1f47b7de90b0 --- tests-v2/featuretests/tracer/CMakeLists.txt | 21 +++ .../tracer/apps/MatrixTranspose.cpp | 143 ++++++++++++++++++ .../matrix_transpose_golden_trace.txt | 80 ++++++++++ tests-v2/featuretests/tracer/tracer_gtest.cpp | 90 ++++++++++- 4 files changed, 333 insertions(+), 1 deletion(-) create mode 100755 tests-v2/featuretests/tracer/apps/MatrixTranspose.cpp create mode 100644 tests-v2/featuretests/tracer/apps/goldentraces/matrix_transpose_golden_trace.txt diff --git a/tests-v2/featuretests/tracer/CMakeLists.txt b/tests-v2/featuretests/tracer/CMakeLists.txt index 266712e68e..12b106406a 100644 --- a/tests-v2/featuretests/tracer/CMakeLists.txt +++ b/tests-v2/featuretests/tracer/CMakeLists.txt @@ -59,6 +59,27 @@ install( target_link_libraries(copy_on_engine hsa-runtime64::hsa-runtime64 Threads::Threads dl stdc++fs) +# Compile MatrixTranspose App with ROCTX +find_library(ROCTX_LIBRARY NAMES roctx64 HINTS ${ROCM_PATH}/lib) +if(ROCTX_LIBRARY) + set_source_files_properties(apps/MatrixTranspose.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) + hip_add_executable(tracer_matrix_transpose apps/MatrixTranspose.cpp) + set_target_properties( + tracer_matrix_transpose + PROPERTIES RUNTIME_OUTPUT_DIRECTORY + "${PROJECT_BINARY_DIR}/tests-v2/featuretests/tracer/apps") + target_link_options(tracer_matrix_transpose PRIVATE "-Wl,--build-id=md5") + target_include_directories( + tracer_matrix_transpose PRIVATE ${ROCM_PATH}) + target_link_libraries(tracer_matrix_transpose ${ROCTX_LIBRARY}) + install( + TARGETS tracer_matrix_transpose + RUNTIME + DESTINATION + ${CMAKE_INSTALL_DATAROOTDIR}/${PROJECT_NAME}/tests/featuretests/tracer/apps + COMPONENT tests) +endif() + # Add test cpp file add_executable(runTracerFeatureTests tracer_gtest.cpp ${GTEST_MAIN_SRC_FILE} ${TEST_UTILS_SRC_FILES}) diff --git a/tests-v2/featuretests/tracer/apps/MatrixTranspose.cpp b/tests-v2/featuretests/tracer/apps/MatrixTranspose.cpp new file mode 100755 index 0000000000..d6452395d6 --- /dev/null +++ b/tests-v2/featuretests/tracer/apps/MatrixTranspose.cpp @@ -0,0 +1,143 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include + +// hip header file +#include +// roctx header file +#include + + +#define WIDTH 1024 + + +#define NUM (WIDTH * WIDTH) + +#define THREADS_PER_BLOCK_X 4 +#define THREADS_PER_BLOCK_Y 4 +#define THREADS_PER_BLOCK_Z 1 + +// Device (Kernel) function, it must be void +__global__ void matrixTranspose(float* out, float* in, const int width) { + int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + out[y * width + x] = in[x * width + y]; +} + +// CPU implementation of matrix transpose +void matrixTransposeCPUReference(float* output, float* input, const unsigned int width) { + for (unsigned int j = 0; j < width; j++) { + for (unsigned int i = 0; i < width; i++) { + output[i * width + j] = input[j * width + i]; + } + } +} + +int main() { + float* Matrix; + float* TransposeMatrix; + float* cpuTransposeMatrix; + + float* gpuMatrix; + float* gpuTransposeMatrix; + + hipDeviceProp_t devProp; + hipGetDeviceProperties(&devProp, 0); + + std::cout << "Device name " << devProp.name << std::endl; + + int i; + int errors; + + Matrix = (float*)malloc(NUM * sizeof(float)); + TransposeMatrix = (float*)malloc(NUM * sizeof(float)); + cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float)); + + // initialize the input data + for (i = 0; i < NUM; i++) { + Matrix[i] = (float)i * 10.0f; + } + + // allocate the memory on the device side + hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); + hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + + uint32_t iterations = 10; + while (iterations-- > 0) { + std::cout << "## Iteration (" << iterations << ") #################" << std::endl; + + // Memory transfer from host to device + hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); + + roctxMark("ROCTX-MARK: before hipLaunchKernel"); + roctxRangePush("ROCTX-RANGE: hipLaunchKernel"); + + roctx_range_id_t roctx_id = roctxRangeStartA("roctx_range with id"); + + // Lauching kernel from host + hipLaunchKernelGGL(matrixTranspose, dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, gpuTransposeMatrix, + gpuMatrix, WIDTH); + + roctxRangeStop(roctx_id); + roctxMark("ROCTX-MARK: after hipLaunchKernel"); + + // Memory transfer from device to host + roctxRangePush("ROCTX-RANGE: hipMemcpy"); + + hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + + roctxRangePop(); // for "hipMemcpy" + roctxRangePop(); // for "hipLaunchKernel" + + // CPU MatrixTranspose computation + matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); + + // verify the results + errors = 0; + double eps = 1.0E-6; + for (i = 0; i < NUM; i++) { + if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) { + errors++; + } + } + if (errors != 0) { + printf("FAILED: %d errors\n", errors); + } else { + printf("PASSED!\n"); + } + + } + + // free the resources on device side + hipFree(gpuMatrix); + hipFree(gpuTransposeMatrix); + + // free the resources on host side + free(Matrix); + free(TransposeMatrix); + free(cpuTransposeMatrix); + + return errors; +} diff --git a/tests-v2/featuretests/tracer/apps/goldentraces/matrix_transpose_golden_trace.txt b/tests-v2/featuretests/tracer/apps/goldentraces/matrix_transpose_golden_trace.txt new file mode 100644 index 0000000000..320775e069 --- /dev/null +++ b/tests-v2/featuretests/tracer/apps/goldentraces/matrix_transpose_golden_trace.txt @@ -0,0 +1,80 @@ +Domain(ROCTX_DOMAIN), Timestamp(2143674432317505), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: before hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674432324258), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674432326692), ROCTX_ID(1), ROCTX_Message(roctx_range with id) +Domain(ROCTX_DOMAIN), Timestamp(2143674432737213), ROCTX_ID(1), +Domain(ROCTX_DOMAIN), Timestamp(2143674432737534), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: after hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674432737774), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipMemcpy) +Domain(ROCTX_DOMAIN), Timestamp(2143674434723226), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674434723687), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674444387378), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: before hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674444387789), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674444387999), ROCTX_ID(2), ROCTX_Message(roctx_range with id) +Domain(ROCTX_DOMAIN), Timestamp(2143674444394461), ROCTX_ID(2), +Domain(ROCTX_DOMAIN), Timestamp(2143674444394622), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: after hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674444394792), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipMemcpy) +Domain(ROCTX_DOMAIN), Timestamp(2143674445256099), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674445256269), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674453836781), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: before hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674453837042), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674453837302), ROCTX_ID(3), ROCTX_Message(roctx_range with id) +Domain(ROCTX_DOMAIN), Timestamp(2143674453840137), ROCTX_ID(3), +Domain(ROCTX_DOMAIN), Timestamp(2143674453840288), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: after hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674453840488), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipMemcpy) +Domain(ROCTX_DOMAIN), Timestamp(2143674454689492), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674454689663), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674463340274), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: before hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674463340495), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674463340745), ROCTX_ID(4), ROCTX_Message(roctx_range with id) +Domain(ROCTX_DOMAIN), Timestamp(2143674463343360), ROCTX_ID(4), +Domain(ROCTX_DOMAIN), Timestamp(2143674463343510), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: after hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674463343771), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipMemcpy) +Domain(ROCTX_DOMAIN), Timestamp(2143674464197424), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674464197614), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674472801650), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: before hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674472801950), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674472802141), ROCTX_ID(5), ROCTX_Message(roctx_range with id) +Domain(ROCTX_DOMAIN), Timestamp(2143674472804525), ROCTX_ID(5), +Domain(ROCTX_DOMAIN), Timestamp(2143674472804676), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: after hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674472804916), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipMemcpy) +Domain(ROCTX_DOMAIN), Timestamp(2143674473667145), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674473667315), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674482305554), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: before hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674482305874), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674482306075), ROCTX_ID(6), ROCTX_Message(roctx_range with id) +Domain(ROCTX_DOMAIN), Timestamp(2143674482312076), ROCTX_ID(6), +Domain(ROCTX_DOMAIN), Timestamp(2143674482312236), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: after hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674482312397), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipMemcpy) +Domain(ROCTX_DOMAIN), Timestamp(2143674483171179), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674483171390), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674491782799), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: before hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674491783029), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674491783219), ROCTX_ID(7), ROCTX_Message(roctx_range with id) +Domain(ROCTX_DOMAIN), Timestamp(2143674491785684), ROCTX_ID(7), +Domain(ROCTX_DOMAIN), Timestamp(2143674491785834), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: after hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674491786004), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipMemcpy) +Domain(ROCTX_DOMAIN), Timestamp(2143674492644777), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674492644977), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674501291111), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: before hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674501291362), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674501291562), ROCTX_ID(8), ROCTX_Message(roctx_range with id) +Domain(ROCTX_DOMAIN), Timestamp(2143674501293906), ROCTX_ID(8), +Domain(ROCTX_DOMAIN), Timestamp(2143674501294067), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: after hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674501294237), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipMemcpy) +Domain(ROCTX_DOMAIN), Timestamp(2143674502139965), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674502140176), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674510759599), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: before hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674510759840), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674510760040), ROCTX_ID(9), ROCTX_Message(roctx_range with id) +Domain(ROCTX_DOMAIN), Timestamp(2143674510762735), ROCTX_ID(9), +Domain(ROCTX_DOMAIN), Timestamp(2143674510762885), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: after hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674510763056), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipMemcpy) +Domain(ROCTX_DOMAIN), Timestamp(2143674511618703), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674511618863), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674520249468), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: before hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674520249688), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674520249879), ROCTX_ID(10), ROCTX_Message(roctx_range with id) +Domain(ROCTX_DOMAIN), Timestamp(2143674520252674), ROCTX_ID(10), +Domain(ROCTX_DOMAIN), Timestamp(2143674520252824), ROCTX_ID(0), ROCTX_Message(ROCTX-MARK: after hipLaunchKernel) +Domain(ROCTX_DOMAIN), Timestamp(2143674520252984), ROCTX_ID(0), ROCTX_Message(ROCTX-RANGE: hipMemcpy) +Domain(ROCTX_DOMAIN), Timestamp(2143674521109282), ROCTX_ID(0), +Domain(ROCTX_DOMAIN), Timestamp(2143674521109493), ROCTX_ID(0), \ No newline at end of file diff --git a/tests-v2/featuretests/tracer/tracer_gtest.cpp b/tests-v2/featuretests/tracer/tracer_gtest.cpp index d8c4855f22..9d5d7a649d 100644 --- a/tests-v2/featuretests/tracer/tracer_gtest.cpp +++ b/tests-v2/featuretests/tracer/tracer_gtest.cpp @@ -76,15 +76,28 @@ void ApplicationParser::SetApplicationEnv(const char* app_name, const char* trac if (trace_type.find("hip") != std::string::npos) { // set --hip-api option + unsetenv("ROCPROFILER_HSA_API_TRACE"); + unsetenv("ROCPROFILER_HSA_ACTIVITY_TRACE"); + unsetenv("ROCPROFILER_ROCTX_TRACE"); setenv("ROCPROFILER_HIP_API_TRACE", "1", true); } if (trace_type.find("hsa") != std::string::npos) { // set --hsa-api and --hsa-activity + unsetenv("ROCPROFILER_HIP_API_TRACE"); + unsetenv("ROCPROFILER_ROCTX_TRACE"); setenv("ROCPROFILER_HSA_API_TRACE", "1", true); setenv("ROCPROFILER_HSA_ACTIVITY_TRACE", "1", true); } + if (trace_type.find("roctx") != std::string::npos) { + // set --roctx-trace option + unsetenv("ROCPROFILER_HIP_API_TRACE"); + unsetenv("ROCPROFILER_HSA_API_TRACE"); + unsetenv("ROCPROFILER_HSA_ACTIVITY_TRACE"); + setenv("ROCPROFILER_ROCTX_TRACE", "1", true); + } + std::stringstream os; os << app_path << test_app_path << app_name; @@ -248,7 +261,7 @@ TEST_F(HelloWorldTest, WhenRunningTracerWithAppThenKernelDurationShouldBePositiv /* * ################################################### - * ############ Async COopy HSA Tests ################ + * ############ Async Copy HSA Tests ################ * ################################################### */ @@ -299,4 +312,79 @@ TEST_F(AsyncCopyTest, DISABLED_WhenRunningTracerWithAppThenAsyncCorelationCountI corelation_count--; EXPECT_EQ(corelation_count, corelation_pair.size()); +} + +/* + * ################################################### + * ######### Matrix Transpose ROCTX Tests ############ + * ################################################### + */ + +constexpr auto GoldenOutputMatrixTranspose = "matrix_transpose_golden_trace.txt"; + +class ROCTXTest : public Tracertest { + protected: + void SetUp() { Tracertest::SetUp("tracer_matrix_transpose", "roctx"); } + void TearDown() {} +}; + +TEST_F(ROCTXTest, WhenRunningTracerWithAppThenROCTxOutputIsgenerated) { + std::string golden_file_path = + GetRunningPath(running_path).append(golden_trace_path).append(GoldenOutputMatrixTranspose); + std::ifstream golden_file(golden_file_path); + std::string golden_output_line; + std::vector roctx_output; + for (auto& itr : output_lines) { + if (itr.find("ROCTX") != std::string::npos) roctx_output.push_back(itr); + } + uint32_t i = 0; + while (!golden_file.eof()) { + ASSERT_TRUE(i < roctx_output.size()) + << "Current Output number of records is less than golden output number of records" + << std::endl; + std::string roctx_output_line = roctx_output[i]; + getline(golden_file, golden_output_line); + std::vector golden_output_split; + std::vector roctx_output_split; + std::string temp = ""; + for (int j = 0; j < (int)roctx_output_line.size(); j++) { + if (roctx_output_line[j] != ',') { + if (roctx_output_line[j] == ' ') continue; + temp += roctx_output_line[j]; + } else { + roctx_output_split.push_back(temp); + temp = ""; + } + } + roctx_output_split.push_back(temp); + temp = ""; + for (int j = 0; j < (int)golden_output_line.size(); j++) { + if (golden_output_line[j] != ',') { + if (golden_output_line[j] == ' ') continue; + temp += golden_output_line[j]; + } else { + golden_output_split.push_back(temp); + temp = ""; + } + } + golden_output_split.push_back(temp); + ASSERT_TRUE(roctx_output_split.size() == golden_output_split.size()) + << "Current Output number of records(" << roctx_output_split.size() + << ") is not equal to golden output number of records (" << golden_output_split.size() + << ")" << std::endl; + uint32_t j = 0; + for (auto& roctx_record : roctx_output_split) { + ASSERT_TRUE(j < golden_output_split.size()) + << "ROCTX record not found: " << roctx_record << std::endl; + std::string golden_output_record = golden_output_split[j++]; + if (roctx_record.find("Timestamp") == std::string::npos) + EXPECT_EQ(roctx_record, golden_output_record) << "ROCTX record mismatch" << std::endl; + else + continue; + } + i++; + } + EXPECT_EQ(roctx_output.size(), i) + << "Current Output number of records is greater than golden output number of records" + << std::endl; } \ No newline at end of file