SWDEV-394940 - [catch2][dtest] Adding basic testcases with HIPRTC apis for supported compiler options & also with combination of compiler options (#327)

Change-Id: Ie393e91afe195387152b588780e3109f21d1ad94
2023-08-14 21:17:40 +05:30
@@ -111,22 +111,26 @@
 	   "Unit_deviceAllocation_Malloc_ComplexDataType",
 	   "Unit_deviceAllocation_New_ComplexDataType",
 	   "Unit_hipStreamValue_Wait32_Blocking_Mask_Eq_1",
+	   "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/327 ===",
+	   "Unit_hiprtcGpuRdcComplrOptnTst",
+	   "Unit_hiprtcDisabledSlpVectorizeComplrOptnTst",
+	   "Unit_hiprtcRpassInlineComplrOptnTst",
+	   "Unit_hiprtcCombiComplrOptnTst",
 	   "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/92 ===",
 	   "Unit_hipGetChannelDesc_Negative_Parameters",
 	   "Unit_hipGraphAddChildGraphNode_CmplxNstGrph_UpdKerFun_Clone",
 	   "=== Below tests fail in stress test on 24/07/23 ===",
 	   "Unit_hipStreamCreateWithPriority_ValidateWithEvents",
 	   "Unit_hipEventIpc",
-           "=== Below hiprtc tests are disabled temporarily,  will be renabled once patches for SWDEV-395996 are merged ===",
-           "Unit_hiprtc_saxpy.Unit_hiprtc_saxpy",
-           "Unit_hiprtc_warpsize.Unit_hiprtc_warpsize",
-           "Unit_hiprtc_functional.Unit_hiprtc_functional",
-           "Unit_hipStreamCaptureRtc.Unit_hipStreamCaptureRtc",
-           "Unit_hiprtc_cpp17.Unit_hiprtc_cpp17",
-           "Unit_hiprtc_namehandling.Unit_hiprtc_namehandling",
-           "Unit_hiprtc_getloweredname.Unit_hiprtc_getloweredname",
-           "Unit_hiprtc_test_hip_bfloat16.Unit_hiprtc_test_hip_bfloat16",
-           "Unit_RTC_LinkerAPI.Unit_RTC_LinkerAPI"
-
+	   "=== Below hiprtc tests are disabled temporarily,  will be renabled once patches for SWDEV-395996 are merged ===",
+	   "Unit_hiprtc_saxpy.Unit_hiprtc_saxpy",
+	   "Unit_hiprtc_warpsize.Unit_hiprtc_warpsize",
+	   "Unit_hiprtc_functional.Unit_hiprtc_functional",
+	   "Unit_hipStreamCaptureRtc.Unit_hipStreamCaptureRtc",
+	   "Unit_hiprtc_cpp17.Unit_hiprtc_cpp17",
+	   "Unit_hiprtc_namehandling.Unit_hiprtc_namehandling",
+	   "Unit_hiprtc_getloweredname.Unit_hiprtc_getloweredname",
+	   "Unit_hiprtc_test_hip_bfloat16.Unit_hiprtc_test_hip_bfloat16",
+	   "Unit_RTC_LinkerAPI.Unit_RTC_LinkerAPI"
 	]
 }
@@ -14,6 +14,13 @@ set(AMD_TEST_SRC
    shfl.cc
 )

+if(UNIX)
+   set(AMD_TEST_SRC ${TEST_SRC}
+       RtcFunctions.cpp
+       RtcUtility.cpp
+       hiprtcComplrOptnTesting.cc)
+endif()
+
 if(HIP_PLATFORM MATCHES "nvidia")
  hip_add_exe_to_target(NAME RTC
                      TEST_SRC ${TEST_SRC}
@@ -0,0 +1,175 @@
+[  {
+      "block_name" : "all_compier_options",
+      "single_CO" : ["architecture", "rdc", "no_denormals",
+                     "denormals", "off_ffp_contract",
+                     "on_ffp_contract", "fast_ffp_contract",
+                     "no_fast_math", "fast_math",
+                     "no_slp_vectorize", "slp_vectorize",
+                     "macro", "undef_macro",
+                     "header_dir", "warning", "Rpass_inline",
+                     "conversion_error", "conversion_no_error",
+                     "conversion_no_warning", "conversion_warning",
+                     "max_thread", "no_unsafe_atomic", "unsafe_atomic",
+                     "no_infinite_num", "infinite_num",
+                     "no_NAN_num", "NAN_num",
+                     "no_finite_math", "finite_math",
+                     "no_associative_math",  "associative_math",
+                     "no_signed_zeros","signed_zeros",
+                     "no_trapping_math", "trapping_math"
+                     ],
+      "Combi_CO" : ["header_dir:no_denormals:undef_macro:slp_vectorize:warning:conversion_no_warning:no_associative_math",
+                    "max_thread:warning:architecture:macro:conversion_no_error:associative_math",
+                    "denormals:macro:warning:header_dir:conversion_no_warning:unsafe_atomic:no_slp_vectorize",
+                    "macro:warning:header_dir:undef_macro:architecture:on_ffp_contract:max_thread",
+                    "no_NAN_num:max_thread:undef_macro:header_dir:warning",
+                    "no_unsafe_atomic:fast_ffp_contract:undef_macro:fast_math",
+                    "warning:max_thread:denormals:header_dir:no_infinite_num",
+                    "no_denormals:header_dir:conversion_warning:off_ffp_contract",
+                    "max_thread:fast_math:macro:no_denormals:no_unsafe_atomic",
+                    "Rpass_inline:no_finite_math:NAN_num:warning:no_fast_math",
+                    "no_infinite_num:no_trapping_math:conversion_no_error",
+                    "infinite_num:no_trapping_math:undef_macro:header_dir",
+                    "infinite_num:NAN_num:no_finite_math:max_thread:NAN_num",
+                    "undef_macro:warning:header_dir:no_denormals:finite_math:associative_math",
+                    "header_dir:Rpass_inline:no_fast_math:infinite_num:NAN_num:no_finite_math",
+                    "denormals:no_unsafe_atomic:fast_ffp_contract",
+                    "fast_math:no_infinite_num:no_NAN_num:finite_math:associative_math:no_signed_zeros:no_trapping_math:fast_ffp_contract",
+                    "no_fast_math:infinite_num:NAN_num:no_finite_math:no_associative_math:signed_zeros:trapping_math:on_ffp_contract"
+                    ]
+   },
+  {
+      "block_name" : "architecture",
+      "compiler_option" : "--gpu-architecture=",
+      "kernel_name" : "max_thread"
+    },
+  {
+      "block_name" : "rdc",
+      "compiler_option" : "-fgpu-rdc",
+      "kernel_name" : "rdc"
+    },
+  {
+      "block_name" : "denormals",
+      "compiler_option" : "-fgpu-flush-denormals-to-zero",
+      "reverse_compiler_option" : "-fno-gpu-flush-denormals-to-zero",
+      "Input_Vals" : [2, -125, 2, -126, 2, -149, 2, -150, 2, -1000, 2, -2000, 2, -128],
+      "Expected_Results" : [1, 0, 0, 0, 0, 0, 0 ],
+      "Expected_Results_for_no" : [1, 1, 1, 0, 0, 0, 1 ],
+      "kernel_name" : "denormals"
+    },
+  {
+      "block_name" : "max_thread",
+      "compiler_option" : "--gpu-max-threads-per-block=",
+      "Target_Vals" : [10, 100],
+      "Input_Vals" : [ 1, 2, 3, 50, 30, 50, 100, 110],
+      "Expected_Results" : [1, 1, 1, 0, 1, 1, 1, 0 ],
+      "ready_compiler_option" : "--gpu-max-threads-per-block=1000",
+      "kernel_name" : "max_thread"
+    },
+  {
+      "block_name" : "fp32_div_sqrt",
+      "compiler_option" : "-fhip-fp32-correctly-rounded-divide-sqrt",
+      "reverse_compiler_option" : "-fno-hip-fp32-correctly-rounded-divide-sqrt",
+      "kernel_name" : "fp32_div_sqrt"
+    },
+  {
+      "block_name" : "ffp_contract",
+      "compiler_option" : ["-ffp-contract=off", "-ffp-contract=on", "-ffp-contract=fast", "-ffp-contract=fast-honor-pragmas"],
+      "kernel_name" : "ffp_contract"
+    },
+  {
+      "block_name" : "fast_math",
+      "compiler_option" : "-ffast-math",
+      "reverse_compiler_option" : "-fno-fast-math",
+      "kernel_name" : "ffp_contract"
+    },
+  {
+      "block_name" : "warning",
+      "compiler_option" : "-w",
+      "kernel_name" : "warning"
+    },
+  {
+      "block_name" : "error",
+      "compiler_option" : ["-Werror=conversion", "-Wno-error=conversion", "-Wconversion", "-Wno-conversion"],
+      "kernel_name" : "error"
+    },
+  {
+      "block_name" : "Rpass_inline",
+      "compiler_option" : "-Rpass=inline",
+      "kernel_name" : "max_thread"
+    },
+  {
+      "block_name" : "macro",
+      "compiler_option" : "-DPI=50",
+      "Expected_Results" : [50],
+      "kernel_name" : "macro"
+    },
+  {
+      "block_name" : "undef_macro",
+      "compiler_option" : ["-DZ=10", "-UZ"],
+      "kernel_name" : "undef_macro"
+    },
+  {
+      "block_name" : "header_dir",
+      "compiler_option" : "-I",
+      "depending_comp_optn" : ["-includefact.h"],
+  	  "Headers" : ["RtcFact.h"],
+      "Src_headers" : ["__device__ int fact(int num) {int fact =1; for (int i=1 ; i<= num ; i++){fact*=i;}return fact;}"],
+    	"Input_Vals" : [5, 10],
+     	"Expected_Results" : [120, 3628800],
+      "kernel_name" : "header_dir"
+    },
+  {
+      "block_name" : "amdgpu_ieee",
+      "compiler_option" : ["-mamdgpu-ieee"],
+      "reverse_compiler_option" : ["-mno-amdgpu-ieee", "-fno-honor-nans"],
+      "kernel_name" : "amdgpu_ieee"
+    },
+  {
+      "block_name" : "slp_vectorize",
+      "compiler_option" : "-fslp-vectorize",
+      "reverse_compiler_option" : "-fno-slp-vectorize",
+      "kernel_name" : "slp_vectorize"
+    },
+  {
+      "block_name" : "unsafe_atomic",
+      "compiler_option" : "-munsafe-fp-atomics",
+      "reverse_compiler_option" : "-mno-unsafe-fp-atomics",
+      "kernel_name" : "unsafe_atomic"
+    },
+  {
+      "block_name" : "infinite_num",
+      "compiler_option" : "-fhonor-infinities",
+      "reverse_compiler_option" : "-fno-honor-infinities",
+      "kernel_name" : "ffp_contract"
+    },
+  {
+      "block_name" : "NAN_num",
+      "compiler_option" : "-fhonor-nans",
+      "reverse_compiler_option" : "-fno-honor-nans",
+      "kernel_name" : "ffp_contract"
+    },
+  {
+      "block_name" : "finite_math",
+      "compiler_option" : "-ffinite-math-only",
+      "reverse_compiler_option" : "-fno-finite-math-only",
+      "kernel_name" : "ffp_contract"
+    },
+  {
+      "block_name" : "associative_math",
+      "compiler_option" : "-fassociative-math",
+      "reverse_compiler_option" : "-fno-associative-math",
+      "kernel_name" : "ffp_contract"
+    },
+  {
+      "block_name" : "signed_zeros",
+      "compiler_option" : "-fsigned-zeros",
+      "reverse_compiler_option" : "-fno-signed-zeros",
+      "kernel_name" : "ffp_contract"
+    },
+  {
+      "block_name" : "trapping_math",
+      "compiler_option" : "-ftrapping-math",
+      "reverse_compiler_option" : "-fno-trapping-math",
+      "kernel_name" : "ffp_contract"
+  }
+]
@@ -0,0 +1,507 @@
+/*
+Copyright (c) 2023 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 sindxl
+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.
+*/
+
+/*
+This file has definition of functions for the following functinality:
+
+1) get_combi_string_vec() : Retrieve the combination string which contains
+contains the combination of block name which indicate the respective compiler
+option seperated by ':' from RtcConfig.jason file and returns them in the
+form of vectors.
+
+2) split_comb_string() : The combination of blockname which are seperated by
+':' has to split so that their respective compiler option can be retrieved
+from the json file. This functn internally calls calling_combination_function()
+for each of the combination of compiler options. This function returns a
+int value i.e the total failed cases in that combination which is obtained
+by calling_combination_function() function.
+
+3) calling_combination_function() : This function takes the combination of
+blockname as the input. The respective compiler option for that block name is
+retrieved from the json file and store the compiler options in a array.
+calling_resp_function() is called which mapps the compiler option function
+which has to be called with a set of required parameters
+(combination of compiler options is one among them). this function returns
+the status of execution ie 1 or 0 (bool).
+
+4) getblock_fromconfig() : This function is used to open the RtcConfig.json
+file and return the blocks.
+
+5) get_string_parameters() and get_array_parameters() : retrieved the
+parameters of the respective block name.
+
+*/
+
+#include <hip/hiprtc.h>
+#include <hip/hip_runtime.h>
+#include <picojson.h>
+#include <vector>
+#include <string>
+#include <fstream>
+#include <iostream>
+#include "headers/RtcUtility.h"
+#include "headers/RtcFunctions.h"
+#include "headers/RtcKernels.h"
+#include <hip_test_common.hh>
+#include "headers/printf_common.h"
+
+std::vector<std::string> get_combi_string_vec() {
+  picojson::array combi_string = get_array_parameters("Combi_CO",
+                                                      "all_compier_options");
+  std::vector<std::string> combi_string_list;
+  for (auto& indx : combi_string) {
+    combi_string_list.push_back(indx.get<std::string>());
+  }
+  return combi_string_list;
+}
+
+int split_comb_string(std::string option) {
+  int start_collon_index = option.find(':');
+  int start_index = 0;
+  std::vector<std::string> combi_block_name;
+  while (start_collon_index != std::string::npos) {
+    std::string singleoption = option.substr(start_index,
+                               start_collon_index - start_index);
+    combi_block_name.push_back(singleoption);
+    start_index = start_collon_index + 1;
+    start_collon_index = option.find(':', start_index);
+  }
+  std::string last_option = option.substr(start_index,
+                            option.length() - start_index);
+  combi_block_name.push_back(last_option);
+  return calling_combination_function(combi_block_name);
+}
+
+int calling_combination_function(std::vector<std::string> combi_vec_list) {
+  int combi_size = combi_vec_list.size();
+  int fast_math_present = -1, undef_present = 0;
+  int a = 0;
+  int max_thread_position;
+  std::string hold_CO[combi_size];
+  const char** Combination_CO = new const char*[combi_size];
+  picojson::array undef_compiler_option = get_array_parameters(
+                                          "compiler_option", "undef_macro");
+  std::vector<std::string> undef_CO_vec;
+  for (auto& indx : undef_compiler_option) {
+    undef_CO_vec.push_back(indx.get<std::string>());
+  }
+  for (int i=0; i< combi_size; i++) {
+    if (combi_vec_list[i] == "max_thread") {
+      std::string ready_CO = get_string_parameters("ready_compiler_option",
+                                                    combi_vec_list[i]);
+      hold_CO[i] = ready_CO;
+      if (combi_vec_list[i] == "max_thread") {
+        max_thread_position = i;
+      }
+    } else if (combi_vec_list[i] == "header_dir") {
+      std::string retrived_CO = get_string_parameters("compiler_option",
+                                                      "header_dir");
+      std::string str = "pwd";
+      const char *cmd = str.c_str();
+      CaptureStream capture(stdout);
+      capture.Begin();
+      system(cmd);
+      capture.End();
+      std::string wor_dir = capture.getData();
+      std::string break_dir = wor_dir.substr(0, wor_dir.find("build"));
+      std::string append_str = "catch/unit/rtc/headers";
+      std::string CO = retrived_CO + " " + break_dir + append_str;
+      hold_CO[i] = CO;
+    } else if (combi_vec_list[i] == "architecture") {
+      std::string retrived_CO = get_string_parameters("compiler_option",
+                                                      "architecture");
+      hipDeviceProp_t prop;
+      HIP_CHECK(hipGetDeviceProperties(&prop, 0));
+      std::string actual_architecture = prop.gcnArchName;
+      std::string complete_CO = retrived_CO + actual_architecture;
+      hold_CO[i] = complete_CO;
+    } else if (check_positive_CO_present(combi_vec_list[i]) == 1) {
+      std::string positive_CO = get_string_parameters("compiler_option",
+                                                      combi_vec_list[i]);
+      hold_CO[i] = positive_CO;
+      if (combi_vec_list[i] == "fast_math")
+        fast_math_present = 1;
+    } else if (check_negative_CO_present(combi_vec_list[i]) == 1) {
+      std::string split_block_name = combi_vec_list[i].substr(3,
+                                               combi_vec_list[i].length() - 3);
+      std::string negative_CO = get_string_parameters(
+                                "reverse_compiler_option", split_block_name);
+      hold_CO[i] = negative_CO;
+      if (split_block_name == "fast_math")
+        fast_math_present = 0;
+    } else if ( combi_vec_list[i] == "conversion_error"
+             || combi_vec_list[i] == "conversion_no_error"
+             || combi_vec_list[i] == "conversion_no_warning"
+             || combi_vec_list[i] == "conversion_warning") {
+      picojson::array compiler_option = get_array_parameters("compiler_option",
+                                                             "error");
+      std::vector<std::string> CO_vec;
+      for (auto& indx : compiler_option) {
+        CO_vec.push_back(indx.get<std::string>());
+      }
+      if (combi_vec_list[i] == "conversion_error") {
+        hold_CO[i] = CO_vec[0];
+      } else if (combi_vec_list[i] == "conversion_no_error") {
+        hold_CO[i] = CO_vec[1];
+      } else if (combi_vec_list[i] == "conversion_warning") {
+        hold_CO[i] = CO_vec[2];
+      } else if (combi_vec_list[i] == "conversion_no_warning") {
+        hold_CO[i] = CO_vec[3];
+      }
+    } else if (combi_vec_list[i] == "off_ffp_contract"
+             || combi_vec_list[i] == "on_ffp_contract"
+             || combi_vec_list[i] == "fast_ffp_contract"
+             || combi_vec_list[i] == "pragmas_ffp_contract") {
+      picojson::array compiler_option = get_array_parameters("compiler_option",
+                                                             "ffp_contract");
+      std::vector<std::string> CO_vec;
+      for (auto& indx : compiler_option) {
+        CO_vec.push_back(indx.get<std::string>());
+      }
+      if (combi_vec_list[i] == "off_ffp_contract") {
+        hold_CO[i] = CO_vec[0];
+      } else if (combi_vec_list[i] == "on_ffp_contract") {
+        hold_CO[i] = CO_vec[1];
+      } else if (combi_vec_list[i] == "fast_ffp_contract") {
+        hold_CO[i] = CO_vec[2];
+      } else if (combi_vec_list[i] == "pragmas_ffp_contract") {
+        hold_CO[i] = CO_vec[3];
+      }
+    } else if (combi_vec_list[i] =="undef_macro") {
+      hold_CO[i] = undef_CO_vec[1].c_str();
+      undef_present = 1;
+    } else {
+      WARN("BLOCK NAME " << combi_vec_list[i] << " NOT PRESENT");
+    }
+    Combination_CO[i] = hold_CO[i].c_str();
+  }
+  int errors = 0;
+  for (int j = 0; j< combi_size; j++) {
+    std::string block_name = combi_vec_list[j].c_str();
+    if (!calling_resp_function(block_name, Combination_CO, combi_size,
+                               max_thread_position, fast_math_present)) {
+      errors++;
+    }
+    Combination_CO[j] = hold_CO[j].c_str();
+  }
+  return errors;
+}
+
+int check_positive_CO_present(std::string find_string) {
+  static std::vector<std::string> positive_CO = {"macro", "warning", "rdc",
+                                                 "denormals", "fp32_div_sqrt",
+                                                 "Rpass_inline", "fast_math",
+                                                 "slp_vectorize",
+                                                 "amdgpu_ieee",
+                                                 "unsafe_atomic",
+                                                 "infinite_num", "NAN_num",
+                                                 "slp_vectorize", "math_errno",
+                                                 "associative_math",
+                                                 "signed_zeros", "finite_math",
+                                                 "trapping_math"};
+  if (std::find(positive_CO.begin(), positive_CO.end(),
+      find_string) != positive_CO.end())
+    return 1;
+  else
+    return 0;
+}
+
+int check_negative_CO_present(std::string find_string) {
+  static std::vector<std::string> negative_CO = {"no_fast_math",
+                                                 "no_fp32_div_sqrt",
+                                                 "no_denormals",
+                                                 "no_slp_vectorize",
+                                                 "no_amdgpu_ieee",
+                                                 "no_unsafe_atomic",
+                                                 "no_infinite_num",
+                                                 "no_slp_vectorize",
+                                                 "no_NAN_num",
+                                                 "no_math_errno",
+                                                 "no_associative_math",
+                                                 "no_signed_zeros",
+                                                 "no_finite_math",
+                                                 "no_trapping_math"};
+  if (std::find(negative_CO.begin(), negative_CO.end(),
+      find_string) != negative_CO.end())
+    return 1;
+  else
+    return 0;
+}
+
+bool  calling_resp_function(const std::string block_name,
+                           const char** Combination_CO,
+                           int Combination_CO_size, int max_thread_position,
+                           int fast_math_present) {
+  if (block_name == "max_thread") {
+    return check_max_thread(Combination_CO, Combination_CO_size,
+                            max_thread_position, fast_math_present);
+  } else if (block_name == "architecture") {
+    return check_architecture(Combination_CO, Combination_CO_size,
+                              max_thread_position, fast_math_present);
+  } else if (block_name == "rdc") {
+    return check_rdc(Combination_CO, Combination_CO_size,
+                     max_thread_position, fast_math_present);
+  } else if (block_name == "denormals") {
+    return check_denormals_enabled(Combination_CO, Combination_CO_size,
+                                   max_thread_position, fast_math_present);
+  } else if (block_name == "no_denormals") {
+    return check_denormals_disabled(Combination_CO, Combination_CO_size,
+                                    max_thread_position, fast_math_present);
+  } else if (block_name == "warning") {
+    return check_warning(Combination_CO, Combination_CO_size,
+                         max_thread_position, fast_math_present);
+  } else if (block_name == "conversion_error") {
+    return check_conversionerror_enabled(Combination_CO, Combination_CO_size,
+                                         max_thread_position,
+                                         fast_math_present);
+  } else if (block_name == "conversion_no_error") {
+    return check_conversionerror_disabled(Combination_CO, Combination_CO_size,
+                                          max_thread_position,
+                                          fast_math_present);
+  } else if (block_name == "conversion_warning") {
+    return check_conversionwarning_enabled(Combination_CO, Combination_CO_size,
+                                           max_thread_position,
+                                           fast_math_present);
+  } else if (block_name == "conversion_no_warning") {
+    return check_conversionwarning_disabled(Combination_CO,
+                                            Combination_CO_size,
+                                            max_thread_position,
+                                            fast_math_present);
+  } else if (block_name == "Rpass_inline") {
+    return check_Rpass_inline(Combination_CO, Combination_CO_size,
+                              max_thread_position, fast_math_present);
+  } else if (block_name == "macro") {
+    return check_macro(Combination_CO, Combination_CO_size,
+                       max_thread_position, fast_math_present);
+  } else if (block_name == "undef_macro") {
+    return check_undef_macro(Combination_CO, Combination_CO_size,
+                             max_thread_position, fast_math_present);
+  } else if (block_name == "header_dir") {
+    return check_header_dir(Combination_CO, Combination_CO_size,
+                            max_thread_position, fast_math_present);
+  } else if (block_name == "no_fast_math") {
+    return check_fast_math_disabled(Combination_CO, Combination_CO_size,
+                                    max_thread_position, fast_math_present);
+  } else if (block_name == "fast_math") {
+    return check_fast_math_enabled(Combination_CO, Combination_CO_size,
+                                   max_thread_position, fast_math_present);
+  } else if (block_name == "off_ffp_contract") {
+    return check_ffp_contract_off(Combination_CO, Combination_CO_size,
+                                  max_thread_position, fast_math_present);
+  } else if (block_name == "on_ffp_contract") {
+    return check_ffp_contract_on(Combination_CO, Combination_CO_size,
+                                 max_thread_position, fast_math_present);
+  } else if (block_name == "fast_ffp_contract") {
+    return check_ffp_contract_fast(Combination_CO, Combination_CO_size,
+                                   max_thread_position, fast_math_present);
+  } else if (block_name == "no_unsafe_atomic") {
+    return check_unsafe_atomic_disabled(Combination_CO, Combination_CO_size,
+                                        max_thread_position,
+                                        fast_math_present);
+  } else if (block_name == "unsafe_atomic") {
+    return check_unsafe_atomic_enabled(Combination_CO, Combination_CO_size,
+                                       max_thread_position,
+                                       fast_math_present);
+  } else if (block_name == "no_slp_vectorize") {
+    return check_slp_vectorize_disabled(Combination_CO, Combination_CO_size,
+                                        max_thread_position,
+                                        fast_math_present);
+  } else if (block_name == "slp_vectorize") {
+    return check_slp_vectorize_enabled(Combination_CO, Combination_CO_size,
+                                       max_thread_position,
+                                       fast_math_present);
+  } else if (block_name == "infinite_num") {
+    return check_infinite_num_enabled(Combination_CO, Combination_CO_size,
+                                      max_thread_position,
+                                      fast_math_present);
+  } else if (block_name == "no_infinite_num") {
+    return check_infinite_num_disabled(Combination_CO, Combination_CO_size,
+                                       max_thread_position,
+                                       fast_math_present);
+  } else if (block_name == "NAN_num") {
+    return check_NAN_num_enabled(Combination_CO, Combination_CO_size,
+                                 max_thread_position, fast_math_present);
+  } else if (block_name == "no_NAN_num") {
+    return check_NAN_num_disabled(Combination_CO, Combination_CO_size,
+                                  max_thread_position, fast_math_present);
+  } else if (block_name == "finite_math") {
+    return check_finite_math_enabled(Combination_CO, Combination_CO_size,
+                                     max_thread_position, fast_math_present);
+  } else if (block_name == "no_finite_math") {
+    return check_finite_math_disabled(Combination_CO, Combination_CO_size,
+                                      max_thread_position, fast_math_present);
+  } else if (block_name == "associative_math") {
+    return check_associative_math_enabled(Combination_CO, Combination_CO_size,
+                                          max_thread_position,
+                                          fast_math_present);
+  } else if (block_name == "no_associative_math") {
+    return check_associative_math_disabled(Combination_CO, Combination_CO_size,
+                                           max_thread_position,
+                                           fast_math_present);
+  } else if (block_name == "signed_zeros") {
+    return check_signed_zeros_enabled(Combination_CO, Combination_CO_size,
+                                      max_thread_position,
+                                      fast_math_present);
+  } else if (block_name == "no_signed_zeros") {
+    return check_signed_zeros_disabled(Combination_CO, Combination_CO_size,
+                                       max_thread_position,
+                                       fast_math_present);
+  } else if (block_name == "trapping_math") {
+    return check_trapping_math_enabled(Combination_CO, Combination_CO_size,
+                                       max_thread_position,
+                                       fast_math_present);
+  } else if (block_name == "no_trapping_math") {
+    return check_trapping_math_disabled(Combination_CO, Combination_CO_size,
+                                        max_thread_position,
+                                        fast_math_present);
+  } else {
+    WARN("BLOCK NAME '" << block_name << "' not found");
+    return 0;
+  }
+}
+
+picojson::array getblock_fromconfig() {
+  std::string str = "pwd";
+  const char *cmd = str.c_str();
+  CaptureStream capture(stdout);
+  capture.Begin();
+  system(cmd);
+  capture.End();
+  std::string wor_dir = capture.getData();
+  std::string break_dir = wor_dir.substr(0, wor_dir.find("build"));
+  std::string append_str = "catch/unit/rtc/RtcConfig.json";
+  std::string config_path = break_dir + append_str;
+  std::string returnValue = "";
+  std::ifstream json_file(config_path.c_str());
+  if (!json_file.is_open()) {
+    WARN("Error loading config.jason");
+    exit(0);
+  }
+  std::string json_str((std::istreambuf_iterator<char>(json_file)),
+                        std::istreambuf_iterator<char>());
+  picojson::value v;
+  std::string err = picojson::parse(v, json_str);
+  if (!err.empty()) {
+    WARN("empty config.jason");
+    exit(0);
+  }
+  picojson::array& blocks = v.get<picojson::array>();
+  return blocks;
+}
+
+std::string get_string_parameters(std::string para_name_to_retrieve,
+                                  std::string block_name) {
+  std::string returnValue = "";
+  picojson::array blocks = getblock_fromconfig();
+  for (picojson::value& block : blocks) {
+    picojson::object& block_obj = block.get<picojson::object>();
+    std::string blk_name = block_obj.at("block_name").get<std::string>();
+    if (blk_name == block_name) {
+      if (para_name_to_retrieve == "compiler_option") {
+        std::string compiler_opt =
+                            block_obj.at("compiler_option").get<std::string>();
+        returnValue += compiler_opt;
+      } else if (para_name_to_retrieve == "Target_Vals") {
+        std::string Target_Vals =
+                                block_obj.at("Target_Vals").get<std::string>();
+        returnValue += Target_Vals;
+      } else if (para_name_to_retrieve == "kernel_name") {
+        std::string ker_name = block_obj.at("kernel_name").get<std::string>();
+        returnValue += ker_name;
+      } else if (para_name_to_retrieve == "reverse_compiler_option") {
+        std::string reverse =
+                    block_obj.at("reverse_compiler_option").get<std::string>();
+        returnValue += reverse;
+      } else if (para_name_to_retrieve == "ready_compiler_option") {
+        std::string ready_CO =
+                      block_obj.at("ready_compiler_option").get<std::string>();
+        returnValue += ready_CO;
+      } else {
+        WARN("REQUESTED FIELD not present : " << para_name_to_retrieve);
+      }
+    } else {
+      continue;
+    }
+  }
+  return returnValue;
+}
+
+picojson::array get_array_parameters(std::string para_name_to_retrieve,
+                                     std::string block_name) {
+  std::string returnValue = "";
+  picojson::array blocks = getblock_fromconfig();
+  for (picojson::value& block : blocks) {
+    picojson::object& block_obj = block.get<picojson::object>();
+    std::string blk_name = block_obj.at("block_name").get<std::string>();
+    if (blk_name == block_name) {
+      if (para_name_to_retrieve == "Target_Vals") {
+        picojson::array& Target_Vals =
+                    block_obj.at("Target_Vals").get<picojson::array>();
+        return Target_Vals;
+      } else if (para_name_to_retrieve == "single_CO") {
+        picojson::array& single_CO =
+                         block_obj.at("single_CO").get<picojson::array>();
+        return single_CO;
+      } else if (para_name_to_retrieve == "Combi_CO") {
+        picojson::array& Combi_CO =
+                         block_obj.at("Combi_CO").get<picojson::array>();
+        return Combi_CO;
+      } else if (para_name_to_retrieve == "Input_Vals") {
+        picojson::array& Input_Vals =
+                         block_obj.at("Input_Vals").get<picojson::array>();
+        return Input_Vals;
+      } else if (para_name_to_retrieve == "Expected_Results") {
+        picojson::array& Expected =
+                     block_obj.at("Expected_Results").get<picojson::array>();
+        return Expected;
+      } else if (para_name_to_retrieve == "Expected_Results_for_no") {
+        picojson::array& Expected_for_no =
+                block_obj.at("Expected_Results_for_no").get<picojson::array>();
+        return Expected_for_no;
+      } else if (para_name_to_retrieve == "compiler_option") {
+        picojson::array& compiler_option =
+                       block_obj.at("compiler_option").get<picojson::array>();
+        return compiler_option;
+      } else if (para_name_to_retrieve == "reverse_compiler_option") {
+        picojson::array& reverse_compiler_option =
+                block_obj.at("reverse_compiler_option").get<picojson::array>();
+        return reverse_compiler_option;
+      } else if (para_name_to_retrieve == "Headers") {
+        picojson::array& Headers =
+                         block_obj.at("Headers").get<picojson::array>();
+        return Headers;
+      } else if (para_name_to_retrieve == "Src_headers") {
+        picojson::array& Src_headers =
+                         block_obj.at("Src_headers").get<picojson::array>();
+        return Src_headers;
+      } else if (para_name_to_retrieve == "depending_comp_optn") {
+        picojson::array& depending_comp_optn =
+                 block_obj.at("depending_comp_optn").get<picojson::array>();
+        return depending_comp_optn;
+      } else {
+        WARN("REQUESTED FIELD not present : " << para_name_to_retrieve);
+        return picojson::array();
+      }
+    } else {
+      continue;
+    }
+  }
+  WARN("REQUESTED BLOCK " << block_name << " is not present ");
+  return picojson::array();
+}
@@ -0,0 +1,38 @@
+/*
+Copyright (c) 2023 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.
+*/
+
+/*
+This file is being read by a function defined check_headers which is present in
+RtcFunctions.cpp file, it requires a function named 'fact' to be present
+in a separate file. The -I compiler option uses this function fact's path as an
+input to find this file and access it.
+*/
+
+#ifndef CATCH_UNIT_RTC_HEADERS_RTCFACT_H_
+#define CATCH_UNIT_RTC_HEADERS_RTCFACT_H_
+
+__device__ int fact(int num) {
+  int fact = 1;
+  for (int i = 1; i <= num; i++) {
+    fact *= i;
+  }
+  return fact;
+}
+
+#endif  // CATCH_UNIT_RTC_HEADERS_RTCFACT_H_
@@ -0,0 +1,178 @@
+/*
+Copyright (c) 2023 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.
+*/
+
+/*
+The Functions defined in RtcFunctions.cpp are declared here in RtcFunctions.h.
+*/
+
+#ifndef CATCH_UNIT_RTC_HEADERS_RTCFUNCTIONS_H_
+#define CATCH_UNIT_RTC_HEADERS_RTCFUNCTIONS_H_
+#include <string>
+
+bool check_architecture(const char** Combination_CO,
+                       int Combination_CO_size, int max_thread_pos,
+                       int fast_math_present);
+
+bool check_rdc(const char** Combination_CO,
+               int Combination_CO_size, int max_thread_pos,
+               int fast_math_present);
+
+bool check_denormals_enabled(const char** Combination_CO,
+                             int Combination_CO_size, int max_thread_pos,
+                             int fast_math_present);
+
+bool check_denormals_disabled(const char** Combination_CO,
+                              int Combination_CO_size, int max_thread_pos,
+                              int fast_math_present);
+
+bool check_ffp_contract_off(const char** Combination_CO,
+                            int Combination_CO_size, int max_thread_pos,
+                            int fast_math_present);
+
+bool check_ffp_contract_on(const char** Combination_CO,
+                           int Combination_CO_size, int max_thread_pos,
+                           int fast_math_present);
+
+bool check_ffp_contract_fast(const char** Combination_CO,
+                             int Combination_CO_size, int max_thread_pos,
+                             int fast_math_present);
+
+bool check_fast_math_enabled(const char** Combination_CO,
+                             int Combination_CO_size, int max_thread_pos,
+                             int fast_math_present);
+
+bool check_fast_math_disabled(const char** Combination_CO,
+                              int Combination_CO_size, int max_thread_pos,
+                              int fast_math_present);
+
+bool check_slp_vectorize_enabled(const char** Combination_CO,
+                                 int Combination_CO_size, int max_thread_pos,
+                                 int fast_math_present);
+
+bool check_slp_vectorize_disabled(const char** Combination_CO,
+                                  int Combination_CO_size, int max_thread_pos,
+                                  int fast_math_present);
+
+bool check_macro(const char** Combination_CO,
+                 int Combination_CO_size, int max_thread_pos,
+                 int fast_math_present);
+
+bool check_undef_macro(const char** Combination_CO,
+                       int Combination_CO_size, int max_thread_pos,
+                       int fast_math_present);
+
+bool check_header_dir(const char** Combination_CO,
+                      int Combination_CO_size, int max_thread_pos,
+                      int fast_math_present);
+
+bool check_warning(const char** Combination_CO,
+                   int Combination_CO_size, int max_thread_pos,
+                   int fast_math_present);
+
+bool check_Rpass_inline(const char** Combination_CO,
+                        int Combination_CO_size, int max_thread_pos,
+                        int fast_math_present);
+
+bool check_conversionerror_enabled(const char** Combination_CO,
+                                   int Combination_CO_size, int max_thread_pos,
+                                   int fast_math_present);
+
+bool check_conversionerror_disabled(const char** Combination_CO,
+                                    int Combination_CO_size,
+                                    int max_thread_pos,
+                                    int fast_math_present);
+
+bool check_conversionwarning_enabled(const char** Combination_CO,
+                                     int Combination_CO_size,
+                                     int max_thread_pos,
+                                     int fast_math_present);
+
+bool check_conversionwarning_disabled(const char** Combination_CO,
+                                      int Combination_CO_size,
+                                      int max_thread_pos,
+                                      int fast_math_present);
+
+bool check_max_thread(const char** Combination_CO,
+                      int Combination_CO_size, int max_thread_pos,
+                      int fast_math_present);
+
+bool check_unsafe_atomic_enabled(const char** Combination_CO,
+                                 int Combination_CO_size, int max_thread_pos,
+                                 int fast_math_present);
+
+bool check_unsafe_atomic_disabled(const char** Combination_CO,
+                                  int Combination_CO_size, int max_thread_pos,
+                                  int fast_math_present);
+
+bool check_infinite_num_enabled(const char** Combination_CO,
+                                int Combination_CO_size, int max_thread_pos,
+                                int fast_math_present);
+
+bool check_infinite_num_disabled(const char** Combination_CO,
+                                 int Combination_CO_size, int max_thread_pos,
+                                 int fast_math_present);
+
+bool check_NAN_num_enabled(const char** Combination_CO,
+                           int Combination_CO_size, int max_thread_pos,
+                           int fast_math_present);
+
+bool check_NAN_num_disabled(const char** Combination_CO,
+                            int Combination_CO_size, int max_thread_pos,
+                            int fast_math_present);
+
+bool check_finite_math_enabled(const char** Combination_CO,
+                               int Combination_CO_size, int max_thread_pos,
+                               int fast_math_present);
+
+bool check_finite_math_disabled(const char** Combination_CO,
+                                int Combination_CO_size, int max_thread_pos,
+                                int fast_math_present);
+
+bool check_associative_math_enabled(const char** Combination_CO,
+                                    int Combination_CO_size,
+                                    int max_thread_pos,
+                                    int fast_math_present);
+
+bool check_associative_math_disabled(const char** Combination_CO,
+                                     int Combination_CO_size,
+                                     int max_thread_pos,
+                                     int fast_math_present);
+
+bool check_signed_zeros_enabled(const char** Combination_CO,
+                                int Combination_CO_size,
+                                int max_thread_pos,
+                                int fast_math_present);
+
+bool check_signed_zeros_disabled(const char** Combination_CO,
+                                 int Combination_CO_size, int max_thread_pos,
+                                 int fast_math_present);
+
+bool check_trapping_math_enabled(const char** Combination_CO,
+                                 int Combination_CO_size, int max_thread_pos,
+                                 int fast_math_present);
+
+bool check_trapping_math_disabled(const char** Combination_CO,
+                                  int Combination_CO_size, int max_thread_pos,
+                                  int fast_math_present);
+
+std::string checking_IR(const char* kername, const char** extra_CO_IRadded,
+                        int extra_CO_IRadded_size, const char** Combination_CO,
+                        int Combination_CO_size);
+
+#endif  // CATCH_UNIT_RTC_HEADERS_RTCFUNCTIONS_H_
@@ -0,0 +1,163 @@
+/*
+Copyright (c) 2023 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.
+*/
+
+/*
+RtcKernels.h contains the string's with the which includes the kernel code.
+They are utilized by the compiler option functions, defined in RtcFunctions.cpp
+*/
+
+#ifndef CATCH_UNIT_RTC_HEADERS_RTCKERNELS_H_
+#define CATCH_UNIT_RTC_HEADERS_RTCKERNELS_H_
+#include <hip/hiprtc.h>
+#include <hip/hip_runtime.h>
+#include <math.h>
+
+static constexpr auto max_thread_string {
+R"(
+extern "C"
+__global__ void max_thread(int* a) {
+  int BD = blockDim.x;
+  *a = BD;
+}
+)"};
+
+static constexpr auto denormals_string {
+R"(
+extern "C"
+__global__ void denormals(double* base, double* power, double* result) {
+  float denorm = powf(*base, *power);
+  if (*result == 0 || *result ==1 )
+    *result = (denorm==0) ? 0 : 1;
+  else
+    *result = powf(*base, *power);
+}
+)"};
+
+static constexpr auto warning_string {
+R"(
+extern "C"
+__global__ void warning() {
+  #warning "Just printing a WARNING message onto the terminal";
+}
+)"};
+
+static constexpr auto fp32_div_sqrt_string {
+R"(
+extern "C"
+__global__ void fp32_div_sqrt(float* result) {
+  float input = 109.6209;
+  *result = sqrt(input);
+}
+)"};
+
+static constexpr auto error_string {
+R"(
+extern "C"
+__global__ void error() {
+  unsigned int a = -1;
+  unsigned int b = +1;
+  signed int c = -1;
+  signed int d = +1;
+}
+)"};
+
+static constexpr auto macro_string {
+R"(
+extern "C"
+__global__ void macro(int *result) {
+  *result = PI;
+}
+)"};
+
+static constexpr auto undef_macro_string {
+R"(
+extern "C"
+__global__ void undef_macro() {
+  int a = Z;
+}
+)"};
+
+static constexpr auto header_dir_string {
+R"(
+#include "RtcFact.h"
+extern "C"
+__global__ void header_dir(int* a, int* val) {
+  *a = fact(*val);
+}
+)"};
+
+static constexpr auto rdc_string {
+R"(
+extern "C"
+__global__ void rdc(float* a, float* b, float* c) {
+  *c = *a * *b;
+}
+)"};
+
+static constexpr auto ffp_contract_string {
+R"(
+extern "C"
+__global__ void ffp_contract(float* a, float* b, float* c) {
+  *c = *a * *b + *c;
+}
+)"};
+
+static constexpr auto slp_vectorize_string {
+R"(
+extern "C"
+__global__ void slp_vectorize(__half2 a, __half2 x, __half2 *y) {
+  (*y).data.x = x.data.x + a.data.x;
+  (*y).data.y = x.data.y + a.data.y;
+}
+)"};
+
+static constexpr auto unsafe_atomic_string {
+R"(
+extern "C"
+__global__ void unsafe_atomic(float* a) {
+  int id = threadIdx.x + blockIdx.x * blockDim.x;
+  if (id < 1000) {
+    unsafeAtomicAdd(&a[id], 0.2f);
+  }
+}
+)"};
+
+static constexpr auto amdgpu_ieee_string {
+R"(
+extern "C"
+__global__ void amdgpu_ieee(float* a, float* b, float* c) {
+  *c = sqrt(*a / *b);
+  printf("sqrt(a * b) = %f\n", *c);
+}
+)"};
+
+static constexpr auto associative_math_string {
+R"(
+extern "C"
+__global__ void associative_math(int* check) {
+  double x = 0.1f;
+  double y = 0.2f;
+  double z = 0.3f;
+  if((x*y)*z != x*(y*z))
+    *check = 1;
+  else *check = 0;
+}
+)"};
+
+#endif  // CATCH_UNIT_RTC_HEADERS_RTCKERNELS_H_
@@ -0,0 +1,53 @@
+/*
+Copyright (c) 2023 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.
+*/
+
+/*
+The Functions defined in RtcUtility.cpp are declared here in RtcUtility.h.
+*/
+
+#ifndef CATCH_UNIT_RTC_HEADERS_RTCUTILITY_H_
+#define CATCH_UNIT_RTC_HEADERS_RTCUTILITY_H_
+#include <picojson.h>
+#include <vector>
+#include <string>
+
+std::vector<std::string> get_combi_string_vec();
+
+int split_comb_string(std::string option);
+
+int calling_combination_function(std::vector<std::string> combi_vec_list);
+
+int check_positive_CO_present(std::string find_string);
+
+int check_negative_CO_present(std::string find_string);
+
+bool calling_resp_function(const std::string block_name,
+                           const char** Combination_CO,
+                           int Combination_CO_size, int max_thread_position,
+                           int fast_math_present);
+
+picojson::array getblock_fromconfig();
+
+std::string get_string_parameters(std::string para_name_to_retrieve,
+                                  std::string block_name);
+
+picojson::array get_array_parameters(std::string para_name_to_retrieve,
+                                     std::string block_name);
+
+#endif  // CATCH_UNIT_RTC_HEADERS_RTCUTILITY_H_
@@ -0,0 +1,181 @@
+/*
+Copyright (c) 2023 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.
+*/
+#ifndef CATCH_UNIT_RTC_HEADERS_PRINTF_COMMON_H_
+#define CATCH_UNIT_RTC_HEADERS_PRINTF_COMMON_H_
+#include <sys/stat.h>
+#include <stdlib.h>
+#include <fcntl.h>
+#include <errno.h>
+#include <fstream>
+#include <iostream>
+#include <map>
+#include <string>
+
+#if defined(_WIN32)
+#include <io.h>
+#else
+#include <error.h>
+#include <unistd.h>
+#endif
+
+#if defined(_WIN32)
+class CaptureStream {
+ private:
+  FILE* stream;
+  int fdPipe[2];
+  int fd;
+
+  static constexpr size_t bufferSize = 25 * 1024 * 1024;
+
+ public:
+  explicit CaptureStream(FILE *original) {
+    stream = original;
+
+    if (pipe(fdPipe, bufferSize, O_TEXT) != 0) {
+      fprintf(stderr, "pipe(3) failed with error %d\n", errno);
+      assert(false);
+    }
+
+    if ((fd = dup(fileno(stream))) == -1) {
+      fprintf(stderr, "dup(1) failed with error %d\n", errno);
+      assert(false);
+    }
+  }
+
+  ~CaptureStream() {
+    close(fd);
+    close(fdPipe[1]);
+    close(fdPipe[0]);
+  }
+
+  void Begin() {
+    fflush(stream);
+
+    if (dup2(fdPipe[1], fileno(stream)) == -1) {
+      fprintf(stderr, "dup2(2) failed with error %d\n", errno);
+      assert(false);
+    }
+
+    setvbuf(stream, NULL, _IONBF, 0);
+  }
+
+  void End() {
+    if (dup2(fd, fileno(stream)) == -1) {
+      fprintf(stderr, "dup2(2) failed with error %d\n", errno);
+      assert(false);
+    }
+  }
+
+  std::string getData() {
+    std::string data;
+    data.resize(bufferSize);
+
+    int numRead = read(fdPipe[0], const_cast<char*>(data.c_str()), bufferSize);
+    data[numRead] = '\0';
+
+    data.resize(strlen(data.c_str()));
+    data.shrink_to_fit();
+
+    return data;
+  }
+};
+#else
+struct CaptureStream {
+  int saved_fd;
+  int orig_fd;
+  int temp_fd;
+
+  char tempname[13] = "mytestXXXXXX";
+
+  explicit CaptureStream(FILE *original) {
+    orig_fd = fileno(original);
+    saved_fd = dup(orig_fd);
+
+    if ((temp_fd = mkstemp(tempname)) == -1) {
+      error(0, errno, "Error");
+      assert(false);
+    }
+  }
+
+  void Begin() {
+    fflush(nullptr);
+    if (dup2(temp_fd, orig_fd) == -1) {
+      error(0, errno, "Error");
+      assert(false);
+    }
+    if (close(temp_fd) != 0) {
+      error(0, errno, "Error");
+      assert(false);
+    }
+  }
+
+  void End() {
+    fflush(nullptr);
+    if (dup2(saved_fd, orig_fd) == -1) {
+      error(0, errno, "Error");
+      assert(false);
+    }
+    if (close(saved_fd) != 0) {
+      error(0, errno, "Error");
+      assert(false);
+    }
+  }
+
+  std::string getData() {
+    std::ifstream tmpFileStream(tempname);
+    std::stringstream strStream;
+    strStream << tmpFileStream.rdbuf();
+    return strStream.str();
+  }
+
+  ~CaptureStream() {
+    if (remove(tempname) != 0) {
+      error(0, errno, "Error");
+      assert(false);
+    }
+  }
+
+  // Truncate the file up to size if we don't want too long log
+  void Truncate(size_t size) {
+    struct stat sb = { 0 };
+    if (::stat(tempname, &sb) == -1) {
+      std::cout << "failed lstat " << tempname;
+      std::cout << "with error: " << ::strerror(errno) << std::endl;
+      return;
+    }
+    if (sb.st_size > size) {
+      if (::truncate(tempname, static_cast<off_t>(size)) == -1) {
+        std::cout << "failed truncate " << tempname;
+        std::cout << "with error: " << ::strerror(errno) << std::endl;
+        return;
+      }
+    }
+  }
+};
+#endif
+
+#define DECLARE_DATA()                                                         \
+  const char *msg_short = "Carpe diem.";                                       \
+  const char *msg_long1 = "Lorem ipsum dolor sit amet, consectetur nullam. "   \
+                          "In mollis imperdiet nibh nec ullamcorper.";         \
+  const char *msg_long2 = "Curabitur nec metus sit amet augue vehicula "       \
+                          "ultrices ut id leo. Lorem ipsum dolor sit amet, "   \
+                          "consectetur adipiscing elit amet.";
+
+#endif  // CATCH_UNIT_RTC_HEADERS_PRINTF_COMMON_H_
@@ -0,0 +1,244 @@
+/*
+Copyright (c) 2023 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/hiprtc.h>
+#include <hip/hip_runtime.h>
+#include <hip/hip_fp16.h>
+#include <picojson.h>
+#include <vector>
+#include <string>
+#include <iostream>
+#include <fstream>
+#include <hip_test_common.hh>
+#include "headers/RtcUtility.h"
+#include "headers/RtcFunctions.h"
+#include "headers/RtcKernels.h"
+#include "headers/printf_common.h"
+
+/*
+Unit_hiprtcSingleComplrOptnTst is a test scenario which validates each
+HIPRTC supported compiler option idividually.
+*/
+// SINGLE COMPILER OPTION TESTING
+const char** null = {};
+TEST_CASE("Unit_hiprtcGpuArchComplrOptnTst") {
+  INFO("Testing '--gpu-architecture=gfx906:sramecc+:xnack-' compiler opt")
+  REQUIRE(check_architecture(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcGpuRdcComplrOptnTst") {
+  INFO("Testing '-fgpu-rdc' compiler option")
+  REQUIRE(check_rdc(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcEnabledDenormalsComplrOptnTst") {
+  INFO("Testing '-fgpu-flush-denormals-to-zero' compiler option")
+  REQUIRE(check_denormals_enabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcDisabledDenormalsComplrOptnTst") {
+  INFO("Testing '-fno-gpu-flush-denormals-to-zero' compiler option")
+  REQUIRE(check_denormals_disabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcOff_ffpContractComplrOptnTst") {
+  INFO("Testing '-ffp-contract=off' compiler option")
+  REQUIRE(check_ffp_contract_off(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcOnffpContractComplrOptnTst") {
+  INFO("Testing '-ffp-contract=on' compiler option")
+  REQUIRE(check_ffp_contract_on(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcFastffpContractComplrOptnTst") {
+  INFO("Testing '-ffp-contract=fast' compiler option")
+  REQUIRE(check_ffp_contract_fast(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcEnabledFastMathComplrOptnTst") {
+  INFO("Testing '-ffast-math' compiler option")
+  REQUIRE(check_fast_math_enabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcDisabledFastMathComplrOptnTst") {
+  INFO("Testing '-fno-fast-math' compiler option")
+  REQUIRE(check_fast_math_disabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcEnabledSlpVectorizeComplrOptnTst") {
+  INFO("Testing '-fslp-vectorize' compiler option")
+  REQUIRE(check_slp_vectorize_enabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcDisabledSlpVectorizeComplrOptnTst") {
+  INFO("Testing '-fno-slp-vectorize' compiler option")
+  REQUIRE(check_slp_vectorize_disabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcDefineMacroComplrOptnTst") {
+  INFO("Testing '-D' compiler option")
+  REQUIRE(check_macro(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcUndefMacroComplrOptnTst") {
+  INFO("Testing '-U' compiler option")
+  REQUIRE(check_undef_macro(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcHeaderDirectoryComplrOptnTst") {
+  INFO("Testing '-I' compiler option")
+  REQUIRE(check_header_dir(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcWarningComplrOptnTst") {
+  INFO("Testing '-w' compiler option")
+  REQUIRE(check_warning(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcRpassInlineComplrOptnTst") {
+  INFO("Testing '-Rpass=inline' compiler option")
+  REQUIRE(check_Rpass_inline(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcEnabledConversionErrComplrOptnTst") {
+  INFO("Testing '-Werror=conversion' compiler option")
+  REQUIRE(check_conversionerror_enabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcDisabledConversionErrComplrOptnTst") {
+  INFO("Testing '-Wno-error=conversion' compiler option")
+  REQUIRE(check_conversionerror_disabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcEnabledConversionWarningComplrOptnTst") {
+  INFO("Testing '-Wconversion' compiler option")
+  REQUIRE(check_conversionwarning_enabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcDisabledConversionWarningComplrOptnTst") {
+  INFO("Testing '-Wno-conversion' compiler option")
+  REQUIRE(check_conversionwarning_disabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcGpuMaxThreadPerBlockComplrOptnTst") {
+  INFO("Testing '--gpu-max-threads-per-block=n' compiler option")
+  REQUIRE(check_max_thread(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcEnabledUnsafeAtomicComplrOptnTst") {
+  INFO("Testing '-munsafe-fp-atomics' compiler option")
+  REQUIRE(check_unsafe_atomic_enabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcDisabledUnsafeAtomicComplrOptnTst") {
+  INFO("Testing '-mno-unsafe-fp-atomics' compiler option")
+  REQUIRE(check_unsafe_atomic_disabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcEnabledInfiniteNumComplrOptnTst") {
+  INFO("Testing '-fhonor-infinities' compiler option")
+  REQUIRE(check_infinite_num_enabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcDisabledInfiniteNumComplrOptnTst") {
+  INFO("Testing '-fno-honor-infinities' compiler option")
+  REQUIRE(check_infinite_num_disabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcEnabledNANComplrOptnTst") {
+  INFO("Testing '-fhonor-nans' compiler option")
+  REQUIRE(check_NAN_num_enabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcDisabledNANComplrOptnTst") {
+  INFO("Testing '-fno-honor-nans' compiler option")
+  REQUIRE(check_NAN_num_disabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcEnabledFiniteMathComplrOptnTst") {
+  INFO("Testing '-ffinite-math-only' compiler option")
+  REQUIRE(check_finite_math_enabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcDisabledFiniteMathComplrOptnTst") {
+  INFO("Testing '-fno-finite-math-only' compiler option")
+  REQUIRE(check_finite_math_disabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcEnabledAssociativeMathComplrOptnTst") {
+  INFO("Testing '-fassociative-math' compiler option")
+  REQUIRE(check_associative_math_enabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcDisabledAssociativeMathComplrOptnTst") {
+  INFO("Testing '-fno-associative-math' compiler option")
+  REQUIRE(check_associative_math_disabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcEnabledSignedZerosComplrOptnTst") {
+  INFO("Testing '-fsigned-zeros' compiler option")
+  REQUIRE(check_signed_zeros_enabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcDisabledSignedZerosComplrOptnTst") {
+  INFO("Testing '-fno-signed-zeros' compiler option")
+  REQUIRE(check_signed_zeros_disabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcEnabledTrappingMathComplrOptnTst") {
+  INFO("Testing '-ftrapping-math' compiler option")
+  REQUIRE(check_trapping_math_enabled(null, -1, -1, -1));
+}
+
+TEST_CASE("Unit_hiprtcDisabledTrappingMathComplrOptnTst") {
+  INFO("Testing '-fno-trapping-math' compiler option")
+  REQUIRE(check_trapping_math_disabled(null, -1, -1, -1));
+}
+
+/*
+Unit_hiprtcCombiComplrOptnTst is a test scenario which validates
+a combination of HIPRTC supported compiler options which a retrieved from
+RtcConfig.jason file.
+*/
+
+TEST_CASE("Unit_hiprtcCombiComplrOptnTst") {
+  // COMBINATION COMPILER OPTIONS
+  std::vector<std::string> CombiCompOptions = get_combi_string_vec();
+  int TotalCombos = CombiCompOptions.size();
+  REQUIRE(TotalCombos != -1);
+  /*
+  use '-Werror=conversion' and '-Wconversion' compiler option individually as
+  the generate ERROR and WARNING message which might effect when used in
+  combination.
+
+  These can be used only if the ERROR and WARNING messages are required.
+  '-fgpu-rdc' has to be tested in ISOLATION, cannot be validated with
+  combi compiler options.
+  */
+  int TotalErrors = 0;
+  for (int i = 0; i < TotalCombos; i++) {
+    std::string one_combi = CombiCompOptions[i];
+    TotalErrors += split_comb_string(one_combi);
+  }
+  if (TotalErrors) {
+    WARN("TOTAL FAILED CASES : " << TotalErrors);
+  }
+  REQUIRE(!TotalErrors);
+}