diff --git a/catch/hipTestMain/config/config_amd_linux_common.json b/catch/hipTestMain/config/config_amd_linux_common.json index 2e7d092244..39c403ef89 100644 --- a/catch/hipTestMain/config/config_amd_linux_common.json +++ b/catch/hipTestMain/config/config_amd_linux_common.json @@ -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" ] } diff --git a/catch/unit/rtc/CMakeLists.txt b/catch/unit/rtc/CMakeLists.txt index 57b17b7807..62fd20bf7e 100644 --- a/catch/unit/rtc/CMakeLists.txt +++ b/catch/unit/rtc/CMakeLists.txt @@ -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} diff --git a/catch/unit/rtc/RtcConfig.json b/catch/unit/rtc/RtcConfig.json new file mode 100644 index 0000000000..f0d00f704f --- /dev/null +++ b/catch/unit/rtc/RtcConfig.json @@ -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" + } +] diff --git a/catch/unit/rtc/RtcFunctions.cpp b/catch/unit/rtc/RtcFunctions.cpp new file mode 100644 index 0000000000..e73e195c07 --- /dev/null +++ b/catch/unit/rtc/RtcFunctions.cpp @@ -0,0 +1,3301 @@ +/* +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 contains functions for idividual HIPRTC supported compiler options +validation. For PASS senario the function returns 1 or 0 otherwise. +*/ + +#include +#include +#include +#include +#include +#include +#include +#include +#include "headers/RtcUtility.h" +#include "headers/RtcFunctions.h" +#include "headers/RtcKernels.h" +#include +#include "headers/printf_common.h" + +bool check_architecture(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "architecture"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string retrieved_CO = get_string_parameters("compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME " << block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + std::string actual_architecture = prop.gcnArchName; + std::string complete_CO = retrieved_CO + actual_architecture; + const char* compiler_option = complete_CO.c_str(); + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, max_thread_string, + kername, 0, NULL, NULL)); + if (Combination_CO_size != -1) { + hiprtcResult compileResult{hiprtcCompileProgram(prog, + Combination_CO_size, + Combination_CO)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } else { + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, + &compiler_option)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } + return 1; +} + +bool check_rdc(const char** Combination_CO, int Combination_CO_size, + int max_thread_pos, int fast_math_present) { + std::string block_name = "rdc"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string CO = get_string_parameters("compiler_option", + block_name); + if (CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + const char* compiler_opt = CO.c_str(); + float *A_d, *B_d, *C_d; + float *A_h, *B_h, *C_h, *result; + float Nbytes = sizeof(float); + A_h = new float[1]; + B_h = new float[1]; + C_h = new float[1]; + result = new float[1]; + for (int i = 0; i < 1; i++) { + A_h[i] = 4; + B_h[i] = 4; + result[i] = 16; + } + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&B_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, rdc_string, kername, 0, NULL, NULL)); + if (Combination_CO_size != -1) { + hiprtcResult compileResult{hiprtcCompileProgram(prog, Combination_CO_size, + Combination_CO)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_opt); + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } else { + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, &compiler_opt)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_opt); + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } + void* kernelParam[] = {A_d, B_d, C_d}; + auto size = sizeof(kernelParam); + void* kernel_parameter[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &kernelParam, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + size_t codeSize; + HIPRTC_CHECK(hiprtcGetBitcodeSize(prog, &codeSize)); + std::vector codec(codeSize); + HIPRTC_CHECK(hiprtcGetBitcode(prog, codec.data())); + float wall_time; + int reg_count = 2; + int max_thread = 1; + unsigned int log_size = 5120; + char error_log[5120]; + char info_log[5120]; + std::vector jit_options = {HIPRTC_JIT_MAX_REGISTERS, + HIPRTC_JIT_THREADS_PER_BLOCK, + HIPRTC_JIT_WALL_TIME, + HIPRTC_JIT_INFO_LOG_BUFFER, + HIPRTC_JIT_INFO_LOG_BUFFER_SIZE_BYTES, + HIPRTC_JIT_ERROR_LOG_BUFFER, + HIPRTC_JIT_ERROR_LOG_BUFFER_SIZE_BYTES, + HIPRTC_JIT_LOG_VERBOSE}; + const void* lopts[] = {reinterpret_cast(®_count), + reinterpret_cast(&max_thread), + reinterpret_cast(&wall_time), + info_log, + reinterpret_cast(log_size), + error_log, + reinterpret_cast(log_size), + reinterpret_cast(1)}; + hiprtcLinkState rtc_link_state; + void* binary; + size_t binarySize; + int pass_count = 0; + hipModule_t module; + hipFunction_t function; + for (int i = 0; i < 2; i++) { + switch (i) { + case 0 : + HIPRTC_CHECK(hiprtcLinkCreate(0, nullptr, nullptr, &rtc_link_state)); + HIPRTC_CHECK(hiprtcLinkAddData(rtc_link_state, + HIPRTC_JIT_INPUT_LLVM_BITCODE, + codec.data(), codeSize, 0, 0, 0, 0)); + HIPRTC_CHECK(hiprtcLinkComplete(rtc_link_state, &binary, &binarySize)); + HIP_CHECK(hipModuleLoadData(&module, binary)); + HIP_CHECK(hipModuleGetFunction(&function, module, kername)); + HIP_CHECK(hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, 0, + nullptr, kernel_parameter)); + pass_count++; + break; + case 1 : + HIPRTC_CHECK(hiprtcLinkCreate(8, jit_options.data(), + reinterpret_cast(&lopts), + &rtc_link_state)); + HIPRTC_CHECK(hiprtcLinkAddData(rtc_link_state, + HIPRTC_JIT_INPUT_LLVM_BITCODE, + codec.data(), codeSize, 0, 0, 0, 0)); + HIPRTC_CHECK(hiprtcLinkComplete(rtc_link_state, &binary, &binarySize)); + HIP_CHECK(hipModuleLoadData(&module, binary)); + HIP_CHECK(hipModuleGetFunction(&function, module, kername)); + HIP_CHECK(hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, 0, + nullptr, kernel_parameter)); + pass_count++; + break; + default: + WARN(" NOT VALID INPUT "); + break; + } + } + HIP_CHECK(hipMemcpy(result, C_d, Nbytes, hipMemcpyDeviceToHost)); + for (int i = 0 ; i< 1; i++) { + if (result[i] != ((A_h[i] * B_h[i]))) { + WARN("Compiler Option : " << compiler_opt); + WARN("EXPECTED RESULT DOES NOT MATCH "); + WARN("INPUT A & B : " << A_h[i] <<" , "<< B_h[i]); + WARN("EXPECTED RES : " << (A_h[i] * B_h[i])); + WARN("OBTAINED RES : " << result[i]); + return 0; + } + } + if (pass_count == 2) { + return 1; + } else { + WARN(" pass_count IS NOT MATCHING "); + return 0; + } +} + +bool check_denormals_enabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "denormals"; + std::string retrieved_CO = get_string_parameters("compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::string kernel_name = get_string_parameters("kernel_name", block_name); + picojson::array Input_Vals = get_array_parameters("Input_Vals", block_name); + picojson::array Expected_Results = get_array_parameters("Expected_Results", + block_name); + const char* kername = kernel_name.c_str(); + const char* compiler_option = retrieved_CO.c_str(); + std::vector double_vec_input; + for (auto& indx : Input_Vals) { + double_vec_input.push_back(indx.get()); + } + std::vector Input_Vals_int; + for (auto& indx : double_vec_input) { + Input_Vals_int.push_back(static_cast(indx)); + } + std::vector double_vec_expected; + for (auto& indx : Expected_Results) { + double_vec_expected.push_back(indx.get()); + } + std::vector Expected_Results_int; + for (auto& indx : double_vec_expected) { + Expected_Results_int.push_back(static_cast(indx)); + } + int pass_count = 0; + int test_case, res_inc; + for (test_case = 0, res_inc = 0; test_case < Input_Vals_int.size() && + res_inc < Expected_Results_int.size(); test_case+=2, res_inc++) { + double *base_h, *power_h, *result_h; + double *base_d, *power_d, *result_d; + double Nbytes = sizeof(double); + base_h = new double[1]; + power_h = new double[1]; + result_h = new double[1]; + *base_h = Input_Vals_int[test_case]; + *power_h = Input_Vals_int[test_case+1]; + *result_h = 1; + HIP_CHECK(hipMalloc(&base_d, Nbytes)); + HIP_CHECK(hipMalloc(&power_d, Nbytes)); + HIP_CHECK(hipMalloc(&result_d, Nbytes)); + HIP_CHECK(hipMemcpy(base_d, base_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(power_d, power_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(result_d, result_h, Nbytes, hipMemcpyHostToDevice)); + hiprtcProgram program; + HIPRTC_CHECK(hiprtcCreateProgram(&program, denormals_string, + "denormals", 0, NULL, NULL)); + if (Combination_CO_size != -1) { + hiprtcResult compileResult{hiprtcCompileProgram(program, + Combination_CO_size, + Combination_CO)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(program, &log[0])); + WARN(log); + } + return 0; + } + } else { + hiprtcResult compileResult{hiprtcCompileProgram(program, 1, + &compiler_option)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(program, &log[0])); + WARN(log); + } + return 0; + } + } + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(program, &codeSize)); + std::vector codec(codeSize); + HIPRTC_CHECK(hiprtcGetCode(program, codec.data())); + void* kernelParam[] = {base_d, power_d, result_d}; + auto size = sizeof(kernelParam); + void* kernel_parameter[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &kernelParam, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + hipModule_t module; + hipFunction_t function; + HIP_CHECK(hipModuleLoadData(&module, codec.data())); + HIP_CHECK(hipModuleGetFunction(&function, module, kername)); + hipError_t status = hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, 0, + nullptr, kernel_parameter); + HIP_CHECK(hipMemcpy(result_h, result_d, sizeof(double), + hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipModuleUnload(module)); + HIPRTC_CHECK(hiprtcDestroyProgram(&program)); + if (*result_h != Expected_Results_int[res_inc]) { + WARN("Compiler Option : " << compiler_option); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("EXPECTED RESULT DOES NOT MATCH FOR " << res_inc); + WARN("th ITERATION (start iteration is 0 ) "); + WARN("INPUT : pow(2, " << *power_h << ") "); + WARN("EXPECTED OP: " << Expected_Results_int[res_inc]); + WARN("OBTAINED OP: " << *result_h); + return 0; + } + } + return 1; +} + +bool check_denormals_disabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "denormals"; + std::string retrieved_CO = get_string_parameters("reverse_compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::string kernel_name = get_string_parameters("kernel_name", block_name); + picojson::array Input_Vals = get_array_parameters("Input_Vals", block_name); + picojson::array Expected_Results_for_no = get_array_parameters( + "Expected_Results_for_no", block_name); + const char* kername = kernel_name.c_str(); + const char* compiler_option = retrieved_CO.c_str(); + int CO_size = 1; + std::vector double_vec_input; + for (auto& indx : Input_Vals) { + double_vec_input.push_back(indx.get()); + } + std::vector Input_Vals_int; + for (auto& indx : double_vec_input) { + Input_Vals_int.push_back(static_cast(indx)); + } + std::vector double_vec_expected_for_no; + for (auto& indx : Expected_Results_for_no) { + double_vec_expected_for_no.push_back(indx.get()); + } + std::vector Expected_Results_for_no_int; + for (auto& indx : double_vec_expected_for_no) { + Expected_Results_for_no_int.push_back(static_cast(indx)); + } + int pass_count = 0; + int test_case, res_inc; + for (test_case = 0, res_inc = 0; test_case < Input_Vals_int.size() && + res_inc < Expected_Results_for_no_int.size(); test_case+=2, res_inc++) { + double *base_h, *power_h, *result_h; + double *base_d, *power_d, *result_d; + double Nbytes = sizeof(double); + base_h = new double[1]; + power_h = new double[1]; + result_h = new double[1]; + *base_h = Input_Vals_int[test_case]; + *power_h = Input_Vals_int[test_case+1]; + *result_h = 0; + HIP_CHECK(hipMalloc(&base_d, Nbytes)); + HIP_CHECK(hipMalloc(&power_d, Nbytes)); + HIP_CHECK(hipMalloc(&result_d, Nbytes)); + HIP_CHECK(hipMemcpy(base_d, base_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(power_d, power_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(result_d, result_h, Nbytes, hipMemcpyHostToDevice)); + hiprtcProgram program; + HIPRTC_CHECK(hiprtcCreateProgram(&program, denormals_string, + "denormals", 0, NULL, NULL)); + if (Combination_CO_size != -1) { + hiprtcResult compileResult{hiprtcCompileProgram(program, + Combination_CO_size, + Combination_CO)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(program, &log[0])); + WARN(log); + } + return 0; + } + } else { + hiprtcResult compileResult{hiprtcCompileProgram(program, 1, + &compiler_option)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(program, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(program, &log[0])); + WARN(log); + } + return 0; + } + } + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(program, &codeSize)); + std::vector codec(codeSize); + HIPRTC_CHECK(hiprtcGetCode(program, codec.data())); + void* kernelParam[] = {base_d, power_d, result_d}; + auto size = sizeof(kernelParam); + void* kernel_parameter[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &kernelParam, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + hipModule_t module; + hipFunction_t function; + HIP_CHECK(hipModuleLoadData(&module, codec.data())); + HIP_CHECK(hipModuleGetFunction(&function, module, kername)); + hipError_t status = hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, 0, + nullptr, kernel_parameter); + HIP_CHECK(hipMemcpy(result_h, result_d, sizeof(double), + hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipModuleUnload(module)); + HIPRTC_CHECK(hiprtcDestroyProgram(&program)); + if (*result_h != Expected_Results_for_no_int[res_inc]) { + WARN("Compiler Option : " << compiler_option); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("EXPECTED RESULT DOES NOT MATCH FOR " << res_inc); + WARN("th ITERATION (start iteration is 0 ) "); + WARN("INPUT : pow(2, " << *power_h << ") "); + WARN("EXPECTED OP: "<< Expected_Results_for_no_int[res_inc]); + WARN("OBTAINED OP: "<< *result_h); + return 0; + } + } + return 1; +} + +bool check_ffp_contract_off(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "ffp_contract"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + picojson::array retrieved_CO = get_array_parameters("compiler_option", + block_name); + if (retrieved_CO.size() < 3) { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::vector CO_vec; + for (auto& indx : retrieved_CO) { + CO_vec.push_back(indx.get()); + } + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + std::string hold = CO_vec[0]; + CO_IRadded[0] = hold.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + std::string data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO, Combination_CO_size); + if (data == "") { + WARN("Compiler option : " << retrieved_CO[0]); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (data.find("fmul contract") != -1 && + data.find("@llvm.fmuladd.f32") != -1) { + WARN("Compiler option : " << retrieved_CO[0]); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR CONTAIN EITHER"); + WARN("'fmul contract' or '@llvm.fmuladd.f32' or both "); + WARN("WHICH IS NOT EXPECTED"); + return 0; + } else { + return 1; + } +} + +bool check_ffp_contract_on(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "ffp_contract"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + picojson::array retrieved_CO = get_array_parameters("compiler_option", + block_name); + if (retrieved_CO.size() < 3) { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::vector CO_vec; + for (auto& indx : retrieved_CO) { + CO_vec.push_back(indx.get()); + } + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + std::string hold = CO_vec[1]; + CO_IRadded[0] = hold.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + std::string data = checking_IR(kername, CO_IRadded, + CO_IRadded_size, Combination_CO, + Combination_CO_size); + if (data == "") { + WARN("Compiler option : " << retrieved_CO[1]); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (fast_math_present!= -1) { + if (fast_math_present == 0 && data.find("@llvm.fmuladd.f32")!= -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO[1]); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN '@llvm.fmuladd.f32' "); + return 0; + } + } else { + if (data.find("@llvm.fmuladd.f32") != -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO[1]); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN '@llvm.fmuladd.f32' "); + return 0; + } + } +} + +bool check_ffp_contract_fast(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "ffp_contract"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + picojson::array retrieved_CO = get_array_parameters("compiler_option", + block_name); + if (retrieved_CO.size() < 3) { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::vector CO_vec; + for (auto& indx : retrieved_CO) { + CO_vec.push_back(indx.get()); + } + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + std::string hold = CO_vec[2]; + CO_IRadded[0] = hold.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + std::string data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO, Combination_CO_size); + if (data == "") { + WARN("Compiler option : " << retrieved_CO[2]); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (fast_math_present!= -1) { + if (fast_math_present == 1 && data.find("contract")!= -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO[2]); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'fmul contract' "); + return 0; + } + } else { + if (data.find("fmul contract") != -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO[2]); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'fmul contract' "); + return 0; + } + } +} + +bool check_fast_math_enabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "fast_math"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string retrieved_CO = get_string_parameters("compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + std::string data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO, Combination_CO_size); + if (data == "") { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (data.find("fmul fast")!= -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'fmul fast' "); + return 0; + } +} + +bool check_fast_math_disabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "fast_math"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string retrieved_CO = get_string_parameters("reverse_compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + std::string data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO, Combination_CO_size); + if (data == "") { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (data.find("fmul fast")!= -1) { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'fmul fast' "); + return 0; + } else { + return 1; + } +} + +bool check_slp_vectorize_enabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "slp_vectorize"; + std::string retrieved_CO = get_string_parameters("compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + __half2 *a_d, *x_d, *y_d; + __half2 a_h, x_h, *y_h; + a_h.data.x = 1.5; + x_h.data.y = 3.0; + CaptureStream capture(stderr); + HIP_CHECK(hipMalloc(&a_d, sizeof(__half2))); + HIP_CHECK(hipMalloc(&x_d, sizeof(__half2))); + HIP_CHECK(hipMalloc(&y_d, sizeof(__half2))); + HIP_CHECK(hipMemcpy(a_d, &a_h, sizeof(__half2), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(x_d, &x_h, sizeof(__half2), hipMemcpyHostToDevice)); + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, slp_vectorize_string, + kername, 0, NULL, NULL)); + if (Combination_CO_size != -1) { + int Combination_CO_IRadded_size = Combination_CO_size+3; + int b = 0; + std::string add_ir_forcombi[Combination_CO_size+3]; + const char** Combination_CO_IRadded = + new const char*[Combination_CO_size+3]; + for (int i = 0; i < Combination_CO_size+3; ++i) { + if (i == Combination_CO_size) { + Combination_CO_IRadded[i] = "-fno-signed-zeros"; + Combination_CO_IRadded[i+1] = "-mllvm"; + Combination_CO_IRadded[i+2] = "-print-after=constmerge"; + break; + } + add_ir_forcombi[i] = Combination_CO[b]; + Combination_CO_IRadded[i] = add_ir_forcombi[i].c_str(); + b++; + } + capture.Begin(); + hiprtcResult compileResult{hiprtcCompileProgram(prog, + Combination_CO_IRadded_size, + Combination_CO_IRadded)}; + capture.End(); + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler option : " << retrieved_CO); + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size+3; i++) { + WARN(Combination_CO_IRadded[i]); + } + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } else { + capture.Begin(); + hiprtcResult compileResult{hiprtcCompileProgram(prog, CO_IRadded_size, + CO_IRadded)}; + capture.End(); + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler option : " << retrieved_CO); + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } + std::string data = capture.getData(); + std::stringstream dataStream; + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + std::vector codec(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, codec.data())); + void* kernelParam[] = {reinterpret_cast(a_d), + reinterpret_cast(x_d), + reinterpret_cast(y_d)}; + auto size = sizeof(kernelParam); + void* kernel_parameter[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &kernelParam, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + hipModule_t module; + hipFunction_t function; + HIP_CHECK(hipModuleLoadData(&module, codec.data())); + HIP_CHECK(hipModuleGetFunction(&function, module, kername)); + HIP_CHECK(hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, 0, nullptr, + kernel_parameter)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipModuleUnload(module)); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + if (data == "") { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + int times = 0; + if (data.find("contract <2 x half>", 0) != -1) { + times++; + } + int start = data.find("contract <2 x half>", 0) + 1; + while (data.find("contract <2 x half>", start) != -1) { + times++; + start = data.find("contract <2 x half>", start)+1; + } + if (times == 1) { + return 1; + } else if (times == 0) { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'fadd contract <2 x half>' "); + return 0; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR CONTAIN 'fadd contract <2 x half>' " << times << "times"); + WARN(" WHICH IS NOT EXPECTED (IT SHOULD BE PRESENT ONCE)"); + return 0; + } +} + +bool check_slp_vectorize_disabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "slp_vectorize"; + std::string retrieved_CO = get_string_parameters("reverse_compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME " << block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + __half2 *a_d, *x_d, *y_d; + __half2 a_h, x_h, *y_h; + a_h.data.x = 1.5; + x_h.data.y = 3.0; + CaptureStream capture(stderr); + HIP_CHECK(hipMalloc(&a_d, sizeof(__half2))); + HIP_CHECK(hipMalloc(&x_d, sizeof(__half2))); + HIP_CHECK(hipMalloc(&y_d, sizeof(__half2))); + HIP_CHECK(hipMemcpy(a_d, &a_h, sizeof(__half2), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(x_d, &x_h, sizeof(__half2), hipMemcpyHostToDevice)); + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, slp_vectorize_string, + kername, 0, NULL, NULL)); + if (Combination_CO_size != -1) { + int Combination_CO_IRadded_size = Combination_CO_size+3; + int b = 0; + std::string add_ir_forcombi[Combination_CO_size+3]; + const char** Combination_CO_IRadded = + new const char*[Combination_CO_size+3]; + for (int i = 0; i < Combination_CO_size+3; ++i) { + if (i == Combination_CO_size) { + Combination_CO_IRadded[i] = "-fno-signed-zeros"; + Combination_CO_IRadded[i+1] = "-mllvm"; + Combination_CO_IRadded[i+2] = "-print-after=constmerge"; + break; + } + add_ir_forcombi[i] = Combination_CO[b]; + Combination_CO_IRadded[i] = add_ir_forcombi[i].c_str(); + b++; + } + capture.Begin(); + hiprtcResult compileResult{hiprtcCompileProgram(prog, + Combination_CO_IRadded_size, + Combination_CO_IRadded)}; + capture.End(); + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler option : " << retrieved_CO); + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size+3; i++) { + WARN(Combination_CO_IRadded[i]); + } + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } else { + capture.Begin(); + hiprtcResult compileResult{hiprtcCompileProgram(prog, CO_IRadded_size, + CO_IRadded)}; + capture.End(); + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler option : " << retrieved_CO); + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } + std::string data = capture.getData(); + std::stringstream dataStream; + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + std::vector codec(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, codec.data())); + void* kernelParam[] = {reinterpret_cast(a_d), + reinterpret_cast(x_d), + reinterpret_cast(y_d)}; + auto size = sizeof(kernelParam); + void* kernel_parameter[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &kernelParam, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + hipModule_t module; + hipFunction_t function; + HIP_CHECK(hipModuleLoadData(&module, codec.data())); + HIP_CHECK(hipModuleGetFunction(&function, module, kername)); + HIP_CHECK(hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, 0, nullptr, + kernel_parameter)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipModuleUnload(module)); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + int times = 0; + if (data.find("contract <2 x half>", 0) != -1) { + times++; + } + int start = data.find("contract <2 x half>", 0) + 1; + while (data.find("contract <2 x half>", start) != -1) { + times++; + start = data.find("contract <2 x half>", start)+1; + } + if (times == 2) { + return 1; + } else if (times < 2) { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR CONTAIN 'fadd contract <2 x half>' " << times << "times"); + WARN(" WHICH IS NOT EXPECTED(IT SHOULD BE PRESENT TWICE)"); + return 0; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR CONTAIN 'fadd contract <2 x half>' " << times << "times"); + WARN(" WHICH IS NOT EXPECTED(IT SHOULD BE PRESENT TWICE)"); + return 0; + } +} + +bool check_macro(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "macro"; + std::string retrieved_CO = get_string_parameters("compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME " << block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::string kernel_name = get_string_parameters("kernel_name", block_name); + picojson::array Expected_Results = get_array_parameters("Expected_Results", + block_name); + const char* kername = kernel_name.c_str(); + std::vector double_vec_expected; + for (auto& indx : Expected_Results) { + double_vec_expected.push_back(indx.get()); + } + std::vector Expected_Results_int; + for (auto& indx : double_vec_expected) { + Expected_Results_int.push_back(static_cast(indx)); + } + const char* compiler_option = retrieved_CO.c_str(); + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, macro_string, + kername, 0, NULL, NULL)); + if (Combination_CO_size != -1) { + hiprtcResult compileResult{hiprtcCompileProgram(prog, Combination_CO_size, + Combination_CO)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + WARN("hiprtcCompileProgram() api failed!! with error code: "); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } else { + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, + &compiler_option)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("hiprtcCompileProgram() api failed!! with error code: "); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } + int *macro_value_h; + int *macro_value_d; + macro_value_h = new int[1]; + HIP_CHECK(hipMalloc(¯o_value_d, sizeof(int))); + *macro_value_h = 0; + HIP_CHECK(hipMemcpy(macro_value_d, macro_value_h, sizeof(int), + hipMemcpyHostToDevice)); + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + std::vector codec(codeSize); + hiprtcGetCode(prog, codec.data()); + void* kernelParam[] = {macro_value_d}; + auto size = sizeof(kernelParam); + void* kernel_parameter[]={HIP_LAUNCH_PARAM_BUFFER_POINTER, &kernelParam, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + hipModule_t module; + hipFunction_t function; + HIP_CHECK(hipModuleLoadData(&module, codec.data())); + HIP_CHECK(hipModuleGetFunction(&function, module, kername)); + HIP_CHECK(hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, 0, nullptr, + kernel_parameter)); + HIP_CHECK(hipMemcpy(macro_value_h, macro_value_d, sizeof(int), + hipMemcpyDeviceToHost)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipModuleUnload(module)); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + if (*macro_value_h != Expected_Results_int[0]) { + WARN("Compiler Option : " << compiler_option); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("EXPECTED RESULT DOES NOT MATCH"); + WARN("INPUT: " << compiler_option); + WARN("EXPECTED OP : "<< Expected_Results_int[0]); + WARN("OBTAINED OP: "<< *macro_value_h); + return 0; + } else { + return 1; + } +} + +bool check_undef_macro(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "undef_macro"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + picojson::array comp_opt = get_array_parameters("compiler_option", + block_name); + if (comp_opt.size() < 2) { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME " << block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::vector compiler_option; + for (auto& indx : comp_opt) { + compiler_option.push_back(indx.get()); + } + std::string variable[compiler_option.size()]; + const char** appended_compiler_options = + new const char*[compiler_option.size()]; + for (int i = 0; i < compiler_option.size(); ++i) { + variable[i] = compiler_option[i]; + appended_compiler_options[i] = variable[i].c_str(); + } + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, undef_macro_string, + kername, 0, NULL, NULL)); + if (Combination_CO_size != -1) { + hiprtcResult compileResult{hiprtcCompileProgram(prog, Combination_CO_size, + Combination_CO)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + if (log.find("undeclared identifier")) { + return 1; + } + } else { + WARN("Compiler Option : " << appended_compiler_options[1]); + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + WARN("Expected error : 'undeclared identifier' NOT GENERATED"); + return 0; + } + } + } else { + hiprtcResult compileResult{hiprtcCompileProgram(prog, + compiler_option.size(), + appended_compiler_options)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + if (log.find("undeclared identifier")) { + return 1; + } + } else { + WARN("Compiler Option : " << appended_compiler_options[0]); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("Expected error : 'undeclared identifier' NOT GENERATED"); + return 0; + } + } + } + WARN("Compiler Option : " << appended_compiler_options[0]); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("EXPECTED ERROR WAS NOT GENERATED"); + return 0; +} + +bool check_header_dir(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "header_dir"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string compiler_option = get_string_parameters("compiler_option", + block_name); + if (compiler_option == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + picojson::array Headers = get_array_parameters("Headers", block_name); + picojson::array depending_comp_optn = + get_array_parameters("depending_comp_optn", block_name); + picojson::array Src_headers = + get_array_parameters("Src_headers", block_name); + picojson::array Input_Thrd_Vals = + get_array_parameters("Input_Vals", block_name); + picojson::array Expected_Results = + get_array_parameters("Expected_Results", block_name); + 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 = compiler_option + " " + break_dir + append_str; + const char* appended_CO = CO.c_str(); + std::vector Headers_list; + for (auto& indx : Headers) { + Headers_list.push_back(indx.get()); + } + std::vector Src_headers_list; + for (auto& indx : Src_headers) { + Src_headers_list.push_back(indx.get()); + } + std::vector depending_co_list; + for (auto& indx : depending_comp_optn) { + depending_co_list.push_back(indx.get()); + } + std::vector double_vec_target; + for (auto& indx : Input_Thrd_Vals) { + double_vec_target.push_back(indx.get()); + } + std::vector Input_Thrd_Vals_int; + for (auto& indx : double_vec_target) { + Input_Thrd_Vals_int.push_back(static_cast(indx)); + } + std::vector double_vec_expected; + for (auto& indx : Expected_Results) { + double_vec_expected.push_back(indx.get()); + } + std::vector Expected_Results_int; + for (auto& indx : double_vec_expected) { + Expected_Results_int.push_back(static_cast(indx)); + } + std::string src_var_hdr_lst[Src_headers_list.size()]; + const char** src_hder_lst = new const char*[Src_headers_list.size()]; + for (int i = 0; i < Src_headers_list.size(); ++i) { + src_var_hdr_lst[i] = Src_headers_list[i]; + src_hder_lst[i] = src_var_hdr_lst[i].c_str(); + } + std::string var_hdr_lst[Headers_list.size()]; + const char** hder_lst = new const char*[Headers_list.size()]; + for (int i = 0; i < Headers_list.size(); ++i) { + var_hdr_lst[i] = Headers_list[i]; + hder_lst[i] = var_hdr_lst[i].c_str(); + } + int pass_count = 0; + for (int senario = 0; senario< Input_Thrd_Vals_int.size(); senario++) { + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, header_dir_string, + kername, Headers_list.size(), + src_hder_lst, hder_lst)); + if (Combination_CO_size != -1) { + hiprtcResult compileResult{hiprtcCompileProgram(prog, + Combination_CO_size, + Combination_CO)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << appended_CO); + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } else { + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, + &appended_CO)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << appended_CO); + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + std::vector codec(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, codec.data())); + int value_h = 0; + int* ptr_value_h = &value_h; + int input_h = Input_Thrd_Vals_int[senario]; + int* ptr_input_h = &input_h; + int* value_d; + int* input_d; + HIP_CHECK(hipMalloc(&value_d, sizeof(int))); + HIP_CHECK(hipMalloc(&input_d, sizeof(int))); + HIP_CHECK(hipMemcpy(value_d, ptr_value_h, sizeof(int), + hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(input_d, ptr_input_h, sizeof(int), + hipMemcpyHostToDevice)); + void* kernelParam[] = {value_d, input_d}; + auto size = sizeof(kernelParam); + void* kernel_parameter[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &kernelParam, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + hipModule_t module; + hipFunction_t function; + HIP_CHECK(hipModuleLoadData(&module, codec.data())); + HIP_CHECK(hipModuleGetFunction(&function, module, kername)); + HIP_CHECK(hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, 0, nullptr, + kernel_parameter)); + HIP_CHECK(hipMemcpy(ptr_value_h, value_d, sizeof(int), + hipMemcpyDeviceToHost)); + if (*ptr_value_h != Expected_Results_int[senario]) { + WARN("Compiler Option : " << appended_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN(" EXPECTED RESULT DOES NOT MATCH FOR " << senario); + WARN("th ITERATION (start iteration is 0 ) "); + WARN(" INPUT: " << Input_Thrd_Vals_int[senario]); + WARN(" EXPECTED OP: "<< Expected_Results_int[senario]); + WARN(" OBTAINED OP: "<< *ptr_value_h); + return 0; + } + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipModuleUnload(module)); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + } + return 1; +} + +bool check_warning(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "warning"; + std::string retrieved_CO = + get_string_parameters("compiler_option", block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME " << block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + const char* compiler_option = retrieved_CO.c_str(); + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, warning_string, kername, + 0, NULL, NULL)); + if (Combination_CO_size != -1) { + hiprtcResult compileResult{hiprtcCompileProgram(prog, Combination_CO_size, + Combination_CO)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } else { + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, + &compiler_option)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + if (-1 != log.find("#warning")) { + WARN("Compiler Option : " << compiler_option); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN(" WARNING MESSAGE IS PRINTING WHICH IS NOT SUPRESSED "); + return 0; + } else { + return 1; + } + } else { + return 1; + } +} + +bool check_Rpass_inline(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "Rpass_inline"; + std::string retrieved_CO = + get_string_parameters("compiler_option", block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME " << block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + const char* compiler_option = retrieved_CO.c_str(); + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, max_thread_string, + kername, 0, NULL, NULL)); + if (Combination_CO_size != -1) { + hiprtcResult compileResult{hiprtcCompileProgram(prog, Combination_CO_size, + Combination_CO)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } else { + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, + &compiler_option)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + if (log.find("inlined into")) { + return 1; + } else { + WARN("Compiler Option : " << compiler_option); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("EXPECTED STRING 'inlined into' IS NOT PRESENT IN LOG "); + return 0; + } + } else { + WARN("Compiler Option : " << compiler_option); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN(" LOG WITH EXPECTED STRING 'inlined into' IS NOT PRESENT "); + return 0; + } +} + +bool check_conversionerror_enabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "error"; + picojson::array retrieved_CO = get_array_parameters("compiler_option", + block_name); + if (retrieved_CO.size() < 4) { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::vector CO_vec; + for (auto& indx : retrieved_CO) { + CO_vec.push_back(indx.get()); + } + std::string variable = CO_vec[0]; + const char* compiler_option = variable.c_str(); + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, error_string, + kername, 0, NULL, NULL)); + if (Combination_CO_size != -1) { + hiprtcResult compileResult{hiprtcCompileProgram(prog, Combination_CO_size, + Combination_CO)}; + } else { + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, + &compiler_option)}; + } + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + std::string variable = "error"; + if (-1 != log.find(variable)) { + return 1; + } else { + WARN("Compiler Option : " << compiler_option); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("ERROR MSG : '" << variable <<"' NOT FOUND"); + return 0; + } + } else { + WARN("Compiler Option : " << compiler_option); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("LOG IS NOT GENERATED"); + WARN("maybe due to presence of '-w' compiler option"); + return 0; + } +} + +bool check_conversionerror_disabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "error"; + picojson::array retrieved_CO = get_array_parameters("compiler_option", + block_name); + if (retrieved_CO.size() < 4) { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::vector CO_vec; + for (auto& indx : retrieved_CO) { + CO_vec.push_back(indx.get()); + } + std::string variable = CO_vec[1]; + const char* compiler_option = variable.c_str(); + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, error_string, + kername, 0, NULL, NULL)); + if (Combination_CO_size != -1) { + hiprtcResult compileResult{hiprtcCompileProgram(prog, Combination_CO_size, + Combination_CO)}; + } else { + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, + &compiler_option)}; + }size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + if (-1 != log.find("error")) { + WARN("Compiler Option : " << compiler_option); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("LOG IS PRESENT WITH ERROR WHICH IS NOT EXPECTED : "); + WARN("maybe due to presence of '-w' compiler option"); + return 0; + } else { + return 1; + } + } else { + return 1; + } +} + +bool check_conversionwarning_enabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "error"; + picojson::array retrieved_CO = get_array_parameters("compiler_option", + block_name); + if (retrieved_CO.size() < 4) { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::vector CO_vec; + for (auto& indx : retrieved_CO) { + CO_vec.push_back(indx.get()); + } + std::string variable = CO_vec[2]; + const char* compiler_option = variable.c_str(); + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, error_string, + kername, 0, NULL, NULL)); + if (Combination_CO_size != -1) { + hiprtcResult compileResult{hiprtcCompileProgram(prog, Combination_CO_size, + Combination_CO)}; + } else { + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, + &compiler_option)}; + }size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + std::string variable = "warning"; + if (-1 != log.find(variable)) { + return 1; + } else { + WARN("Compiler Option : " << compiler_option); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("LOG DOESN'T CONTAIN WARNING AS EXP : " << compiler_option); + return 0; + } + } else { + WARN("Compiler Option : " << compiler_option); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("LOG IS NOT GENERATED"); + return 0; + } +} + +bool check_conversionwarning_disabled(const char** Combination_CO, + int Combination_CO_size, + int max_thread_pos, + int fast_math_present) { + std::string block_name = "error"; + picojson::array retrieved_CO = get_array_parameters("compiler_option", + block_name); + if (retrieved_CO.size() < 4) { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::vector CO_vec; + for (auto& indx : retrieved_CO) { + CO_vec.push_back(indx.get()); + } + std::string variable = CO_vec[3]; + const char* compiler_option = variable.c_str(); + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, error_string, + kername, 0, NULL, NULL)); + if (Combination_CO_size != -1) { + hiprtcResult compileResult{hiprtcCompileProgram(prog, Combination_CO_size, + Combination_CO)}; + } else { + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, + &compiler_option)}; + }size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + if (-1 != log.find("warning")) { + WARN("Compiler Option : " << compiler_option); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("WARNING IS GENERATED WHICH IS NOT EXPECTED"); + WARN(compiler_option); + return 0; + } else { + return 1; + } + } else { + return 1; + } +} + +bool check_max_thread(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "max_thread"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + std::string default_CO = get_string_parameters("kernel_name", block_name); + picojson::array Target_Thrd_Vals = get_array_parameters("Target_Vals", + block_name); + picojson::array Input_Thrd_Vals = get_array_parameters("Input_Vals", + block_name); + picojson::array Expected_Results = get_array_parameters("Expected_Results", + block_name); + const char* kername = kernel_name.c_str(); + std::string compiler_option = get_string_parameters("compiler_option", + block_name); + if (compiler_option == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::vector double_vec_target; + for (auto& indx : Target_Thrd_Vals) { + double_vec_target.push_back(indx.get()); + } + std::vector Target_Thrd_Vals_int; + for (auto& indx : double_vec_target) { + Target_Thrd_Vals_int.push_back(static_cast(indx)); + } + int a = 0; + std::string variable[Target_Thrd_Vals_int.size()]; + const char** appended_compiler_options = + new const char*[Target_Thrd_Vals_int.size()]; + for (int i = 0; i < Target_Thrd_Vals_int.size() ; i++) { + variable[i] = compiler_option + std::to_string(Target_Thrd_Vals_int[i]); + appended_compiler_options[i] = variable[i].c_str(); + } + std::vector double_vec_input; + for (auto& indx : Input_Thrd_Vals) { + double_vec_input.push_back(indx.get()); + } + std::vector Input_Thrd_Vals_int; + for (auto& indx : double_vec_input) { + Input_Thrd_Vals_int.push_back(static_cast(indx)); + } + std::vector double_vec_expected; + for (auto& indx : Expected_Results) { + double_vec_expected.push_back(indx.get()); + } + std::vector Expected_Results_int; + for (auto& indx : double_vec_expected) { + Expected_Results_int.push_back(static_cast(indx)); + } + int pass_count = 0; + int inc = (Input_Thrd_Vals_int.size()/Target_Thrd_Vals_int.size()); + int start = 0; + int check, test_case; + for (int senario = 0; senario < Target_Thrd_Vals_int.size(); senario++) { + if (Target_Thrd_Vals_int[senario] == 0) { + check = 0; + for (test_case = start; test_case< (start+inc); test_case++) { + if (check == Expected_Results_int[test_case]) { + pass_count++; + } + } + start+= inc; + continue; + } + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, max_thread_string, + kername, 0, NULL, NULL)); + if (Combination_CO_size != -1) { + std::string max_thread_string = variable[senario]; + Combination_CO[max_thread_pos] = max_thread_string.c_str(); + hiprtcResult compileResult{hiprtcCompileProgram(prog, + Combination_CO_size, + Combination_CO)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << appended_compiler_options[senario]); + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } else { + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, + &appended_compiler_options[senario])}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << appended_compiler_options[senario]); + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + std::vector codec(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, codec.data())); + for (test_case = start; test_case< (start+inc); test_case++) { + int num_threads_h = 0; + int* ptr_num_threads_h = &num_threads_h; + int* Thread_count_d; + HIP_CHECK(hipMalloc(&Thread_count_d, sizeof(int))); + HIP_CHECK(hipMemcpy(Thread_count_d, ptr_num_threads_h, sizeof(int), + hipMemcpyHostToDevice)); + void* kernelParam[] = {Thread_count_d}; + auto size = sizeof(kernelParam); + void* kernel_parameter[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &kernelParam, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + hipModule_t module; + hipFunction_t function; + HIP_CHECK(hipModuleLoadData(&module, codec.data())); + HIP_CHECK(hipModuleGetFunction(&function, module, kername)); + hipError_t status = hipModuleLaunchKernel(function, 1, 1, 1, + Input_Thrd_Vals_int[test_case], + 1, 1, 0, 0, nullptr, + kernel_parameter); + HIP_CHECK(hipMemcpy(ptr_num_threads_h, Thread_count_d, sizeof(int), + hipMemcpyDeviceToHost)); + if ((status == hipSuccess) && + (num_threads_h <= Target_Thrd_Vals_int[senario])) { + check = 1; + } else { + check = 0; + } + if (check != Expected_Results_int[test_case]) { + WARN("Compiler Option : " << appended_compiler_options[senario]); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + std::string max_thread_string = variable[senario]; + Combination_CO[max_thread_pos] = max_thread_string.c_str(); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("EXPECTED RESULT DOES NOT MATCH FOR " << test_case); + WARN("th ITERATION (start iteration is 0 ) "); + WARN("IP THREAD VAL: " << Input_Thrd_Vals_int[test_case]); + WARN("EXPECTED OP: "<< Expected_Results_int[test_case]); + WARN("OBTAINED OP: "<< check); + return 0; + } + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipModuleUnload(module)); + } + start+=inc; + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + } + return 1; +} + +bool check_unsafe_atomic_enabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "unsafe_atomic"; + std::string compiler_option = get_string_parameters("compiler_option", + block_name); + if (compiler_option == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + const char *compiler_option_cstr = compiler_option.c_str(); + float *A_d; + int N = 1000; + float A_h[N]; + float Nbytes = N * sizeof(float); + double sum_w = 0, sum_wo = 0, sum_tocheck = 0; + for (int i = 0; i < N; i++) { + A_h[i] = 0.1f; + sum_tocheck += A_h[i] + 0.2f; + } + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + for (int senario = 0; senario < 2; senario ++) { + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, unsafe_atomic_string, + kername, 0, NULL, NULL)); + if (Combination_CO_size != -1) { + hiprtcResult compileResult{hiprtcCompileProgram(prog, + Combination_CO_size, + Combination_CO)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } else { + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, + &compiler_option_cstr)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + std::vector codec(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, codec.data())); + void* kernelParam[] = {A_d}; + auto size = sizeof(kernelParam); + void* kernel_parameter[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &kernelParam, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + hipModule_t module; + hipFunction_t function; + HIP_CHECK(hipModuleLoadData(&module, codec.data())); + HIP_CHECK(hipModuleGetFunction(&function, module, kername)); + HIP_CHECK(hipModuleLaunchKernel(function, N, 1, 1, N, 1, 1, 0, 0, + nullptr, kernel_parameter)); + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + for (int i = 0; i < N; i++) { + if (senario == 0) { + sum_wo += A_h[i]; + } else { + sum_w += A_h[i]; + } + } + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipModuleUnload(module)); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + } + if (sum_w != sum_tocheck) { + return 1; + } else { + WARN("Compiler Option : " << compiler_option); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("EXPECTED : " << sum_w << " != " << sum_tocheck); + return 0; + } +} + +bool check_unsafe_atomic_disabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "unsafe_atomic"; + std::string retrieved_CO = get_string_parameters("reverse_compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + const char* compiler_option = retrieved_CO.c_str(); + float *A_d; + int N = 1000; + float A_h[N]; + float Nbytes = N * sizeof(float); + double sum = 0, sum_tocheck = 0; + for (int i = 0; i < N; i++) { + A_h[i] = 0.1f; + sum_tocheck += A_h[i] + 0.2f; + } + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, unsafe_atomic_string, + kername, 0, NULL, NULL)); + if (Combination_CO_size != -1) { + hiprtcResult compileResult{hiprtcCompileProgram(prog, + Combination_CO_size, + Combination_CO)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } else { + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, &compiler_option)}; + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler Option : " << compiler_option); + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return 0; + } + } + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + std::vector codec(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, codec.data())); + void* kernelParam[] = {A_d}; + auto size = sizeof(kernelParam); + void* kernel_parameter[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &kernelParam, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + hipModule_t module; + hipFunction_t function; + HIP_CHECK(hipModuleLoadData(&module, codec.data())); + HIP_CHECK(hipModuleGetFunction(&function, module, kername)); + HIP_CHECK(hipModuleLaunchKernel(function, N, 1, 1, N, 1, 1, 0, 0, + nullptr, kernel_parameter)); + HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost)); + for (int i = 0; i < N; i++) { + sum += A_h[i]; + } + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipModuleUnload(module)); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + if (sum == sum_tocheck) { + return 1; + } else { + WARN("Compiler Option : " << compiler_option); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("EXPECTED RESULT IS NOT OBTAINED "); + WARN("EXPECTED RESULT: "<< sum_tocheck); + WARN("OBTAINED RESULT: "<< sum); + return 0; + } +} + +bool check_infinite_num_enabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "infinite_num"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string retrieved_CO = get_string_parameters("compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + std::string data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO, Combination_CO_size); + if (data == "") { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (fast_math_present != -1) { + if (fast_math_present == 0 && data.find("contract") != -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'contract' "); + return 0; + } + } else { + if (data.find("ninf")!= -1) { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'ninf' "); + return 0; + } else { + return 1; + } + } +} + +bool check_infinite_num_disabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "infinite_num"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string retrieved_CO = get_string_parameters("reverse_compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + std::string data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO, Combination_CO_size); + if (data == "") { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (fast_math_present != -1) { + if (fast_math_present == 1 && data.find("fmul fast")!= -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'fmul fast' "); + return 0; + } + } else { + if (data.find("ninf")!= -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'ninf' "); + return 0; + } + } +} + +bool check_NAN_num_enabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "NAN_num"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string retrieved_CO = get_string_parameters("compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME "); + WARN(block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + std::string data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO, Combination_CO_size); + if (data == "") { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (fast_math_present!= -1) { + if (fast_math_present == 0 && data.find("contract")!= -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'contract' "); + return 0; + } + } else { + if (data.find("nnan")!= -1) { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'nnan' "); + return 0; + } else { + return 1; + } + } +} + +bool check_NAN_num_disabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "NAN_num"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string retrieved_CO = get_string_parameters("reverse_compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME " << block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + std::string data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO, Combination_CO_size); + if (data == "") { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (fast_math_present!= -1) { + if (fast_math_present == 1 && data.find("fmul fast")!= -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'fmul fast' "); + return 0; + } + } else { + if (data.find("nnan")!= -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'nnan' "); + return 0; + } + } +} + +bool check_finite_math_enabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "finite_math"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string retrieved_CO = get_string_parameters("compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME " << block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + std::string data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO, Combination_CO_size); + if (data == "") { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (fast_math_present!= -1) { + if (fast_math_present == 1 && data.find("fmul fast")!= -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'fmul fast'"); + return 0; + } + } else { + if (data.find("nnan")!= -1 && (data.find("ninf") != -1)) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'nnan' or 'ninf' or both "); + return 0; + } + } +} + +bool check_finite_math_disabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "finite_math"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string retrieved_CO = get_string_parameters("reverse_compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME " << block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + std::string data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO, Combination_CO_size); + if (data == "") { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (fast_math_present!= -1) { + if (fast_math_present == 0 && data.find("contract")!= -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'contract'"); + return 0; + } + } else { + if (data.find("nnan")!= -1 && (data.find("ninf") != -1)) { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR CONTAIN 'nnan' or 'ninf' or both WHICH IS NOT EXPECTED "); + return 0; + } else { + return 1; + } + } +} + +bool check_associative_math_enabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "associative_math"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string retrieved_CO = get_string_parameters("compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME " << block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + int CO_IRadded_size = 4, a = 0; + const char** CO_IRadded = new const char*[4]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-fno-signed-zeros"; + CO_IRadded[2] = "-mllvm"; + CO_IRadded[3] = "-print-after=constmerge"; + std::string data; + if (Combination_CO_size != -1) { + int Combination_CO_IRadded_size = Combination_CO_size+1; + int b = 0; + std::string add_ir_forcombi[Combination_CO_size+1]; + const char** Combination_CO_IRadded = + new const char*[Combination_CO_size+1]; + for (int i = 0; i < Combination_CO_size+1; ++i) { + if (i == Combination_CO_size) { + Combination_CO_IRadded[i] = "-fno-signed-zeros"; + break; + } + add_ir_forcombi[i] = Combination_CO[b]; + Combination_CO_IRadded[i] = add_ir_forcombi[i].c_str(); + b++; + } + data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO_IRadded, + Combination_CO_IRadded_size); + } else { + data = checking_IR(kername, CO_IRadded, CO_IRadded_size, Combination_CO, + Combination_CO_size); + } + if (data == "") { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (fast_math_present!= -1) { + if (fast_math_present == 1 && data.find("fmul fast")!= -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'fmul fast' "); + return 0; + } + } else { + if (data.find("reassoc") != -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'reassoc' "); + WARN(data); + return 0; + } + } +} + +bool check_associative_math_disabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "associative_math"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string retrieved_CO = get_string_parameters("reverse_compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME " << block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + int CO_IRadded_size = 4, a = 0; + const char** CO_IRadded = new const char*[4]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-fno-signed-zeros"; + CO_IRadded[2] = "-mllvm"; + CO_IRadded[3] = "-print-after=constmerge"; + std::string data; + if (Combination_CO_size != -1) { + int Combination_CO_IRadded_size = Combination_CO_size+1; + int b = 0; + std::string add_ir_forcombi[Combination_CO_size+1]; + const char** Combination_CO_IRadded = + new const char*[Combination_CO_size+1]; + for (int i = 0; i < Combination_CO_size+1; ++i) { + if (i == Combination_CO_size) { + Combination_CO_IRadded[i] = "-fno-signed-zeros"; + break; + } + add_ir_forcombi[i] = Combination_CO[b]; + Combination_CO_IRadded[i] = add_ir_forcombi[i].c_str(); + b++; + } + data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO_IRadded, + Combination_CO_IRadded_size); + } else { + data = checking_IR(kername, CO_IRadded, CO_IRadded_size, Combination_CO, + Combination_CO_size); + } + if (data == "") { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (fast_math_present!= -1) { + if (fast_math_present == 0 && data.find("contract")!= -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'contract' "); + return 0; + } + } else { + if (data.find("reassoc")!= -1) { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR CONTAIN 'reassoc' WHICH IS NOT EXPECTED "); + return 0; + } else { + return 1; + } + } +} + +bool check_signed_zeros_enabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "signed_zeros"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string retrieved_CO = get_string_parameters("compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME " << block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + std::string data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO, Combination_CO_size); + if (data == "") { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (fast_math_present!= -1) { + if (fast_math_present == 0 && data.find("contract")!= -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'contract' "); + return 0; + } + } else { + if (data.find("nsz") != -1) { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR CONTAIN 'nsz' WHICH IS NOT EXPECTED "); + return 0; + } else { + return 1; + } + } +} + +bool check_signed_zeros_disabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "signed_zeros"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string retrieved_CO = get_string_parameters("reverse_compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME " << block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + std::string data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO, Combination_CO_size); + if (data == "") { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (fast_math_present!= -1) { + if (fast_math_present == 1 && data.find("fmul fast")!= -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'fmul fast' "); + return 0; + } + } else { + if (data.find("nsz") != -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN 'nsz' "); + return 0; + } + } +} + +bool check_trapping_math_enabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "trapping_math"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string retrieved_CO = get_string_parameters("compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME " << block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + std::string data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO, Combination_CO_size); + if (data == "") { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (data.find("\"no-trapping-math\"=\"true\"") != -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN '\"no-trapping-math\"=\"true\"'"); + return 0; + } +} + +bool check_trapping_math_disabled(const char** Combination_CO, + int Combination_CO_size, int max_thread_pos, + int fast_math_present) { + std::string block_name = "trapping_math"; + std::string kernel_name = get_string_parameters("kernel_name", block_name); + const char* kername = kernel_name.c_str(); + std::string retrieved_CO = get_string_parameters("reverse_compiler_option", + block_name); + if (retrieved_CO == "") { + WARN("COMPILER OPTION NOT PROVIDED FOR BLOCK NAME " << block_name); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + return 0; + } + int CO_IRadded_size = 3, a = 0; + const char** CO_IRadded = new const char*[3]; + CO_IRadded[0] = retrieved_CO.c_str(); + CO_IRadded[1] = "-mllvm"; + CO_IRadded[2] = "-print-after=constmerge"; + std::string data = checking_IR(kername, CO_IRadded, CO_IRadded_size, + Combination_CO, Combination_CO_size); + if (data == "") { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR NOT GENERATED"); + return 0; + } + if (data.find("\"no-trapping-math\"=\"true\"") != -1) { + return 1; + } else { + WARN("Compiler option : " << retrieved_CO); + if (Combination_CO_size != -1) { + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + } + WARN("IR DOESN'T CONTAIN '\"no-trapping-math\"=\"true\"'"); + return 0; + } +} + +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) { + float *A_d, *B_d, *C_d; + float *A_h, *B_h, *C_h, *result; + float Nbytes = sizeof(float); + A_h = new float[1]; + B_h = new float[1]; + C_h = new float[1]; + result = new float[1]; + for (int i = 0; i < 1; i++) { + A_h[i] = 0.1f; + B_h[i] = 0.1f; + C_h[i] = 0.1f; + result[i] = 0.2f; + } + HIP_CHECK(hipMalloc(&A_d, Nbytes)); + HIP_CHECK(hipMalloc(&B_d, Nbytes)); + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(C_d, C_h, Nbytes, hipMemcpyHostToDevice)); + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, ffp_contract_string, + kername, 0, NULL, NULL)); + int Combination_CO_IRadded_size; + CaptureStream capture(stderr); + if (Combination_CO_size != -1) { + Combination_CO_IRadded_size = Combination_CO_size+2; + int b = 0; + std::string add_ir_forcombi[Combination_CO_size+2]; + const char** Combination_CO_IRadded = + new const char*[Combination_CO_size+2]; + for (int i = 0; i < Combination_CO_size+2; ++i) { + if (i == Combination_CO_size) { + Combination_CO_IRadded[i] = "-mllvm"; + Combination_CO_IRadded[i+1] = "-print-after=constmerge"; + break; + } + add_ir_forcombi[i] = Combination_CO[b]; + Combination_CO_IRadded[i] = add_ir_forcombi[i].c_str(); + b++; + } + capture.Begin(); + hiprtcResult compileResult{hiprtcCompileProgram(prog, + Combination_CO_IRadded_size, + Combination_CO_IRadded)}; + capture.End(); + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("Compiler option : " << extra_CO_IRadded[0]); + WARN("FAILED IN COMBINATION :"); + for (int i = 0; i < Combination_CO_size; i++) { + WARN(Combination_CO[i]); + } + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return ""; + } + } else { + capture.Begin(); + hiprtcResult compileResult{hiprtcCompileProgram(prog, + extra_CO_IRadded_size, + extra_CO_IRadded)}; + capture.End(); + if (!(compileResult == HIPRTC_SUCCESS)) { + WARN("hiprtcCompileProgram() api failed!! with error code: "); + WARN(compileResult); + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + return""; + } + } + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + std::vector codec(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, codec.data())); + void* kernelParam[] = {A_d, B_d, C_d}; + auto size = sizeof(kernelParam); + void* kernel_parameter[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &kernelParam, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END}; + hipModule_t module; + hipFunction_t function; + HIP_CHECK(hipModuleLoadData(&module, codec.data())); + HIP_CHECK(hipModuleGetFunction(&function, module, kername)); + HIP_CHECK(hipModuleLaunchKernel(function, 1, 1, 1, 1, 1, 1, 0, 0, nullptr, + kernel_parameter)); + HIP_CHECK(hipMemcpy(result, C_d, Nbytes, hipMemcpyDeviceToHost)); + for (int i = 0; i< 1; i++) { + if (result[i] != ((A_h[i] * B_h[i]) + C_h[i])) { + return ""; + } + } + std::string data = capture.getData(); + std::stringstream dataStream; + HIP_CHECK(hipModuleUnload(module)); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + return data; +} diff --git a/catch/unit/rtc/RtcUtility.cpp b/catch/unit/rtc/RtcUtility.cpp new file mode 100644 index 0000000000..2463a73fd1 --- /dev/null +++ b/catch/unit/rtc/RtcUtility.cpp @@ -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 +#include +#include +#include +#include +#include +#include +#include "headers/RtcUtility.h" +#include "headers/RtcFunctions.h" +#include "headers/RtcKernels.h" +#include +#include "headers/printf_common.h" + +std::vector get_combi_string_vec() { + picojson::array combi_string = get_array_parameters("Combi_CO", + "all_compier_options"); + std::vector combi_string_list; + for (auto& indx : combi_string) { + combi_string_list.push_back(indx.get()); + } + return combi_string_list; +} + +int split_comb_string(std::string option) { + int start_collon_index = option.find(':'); + int start_index = 0; + std::vector 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 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 undef_CO_vec; + for (auto& indx : undef_compiler_option) { + undef_CO_vec.push_back(indx.get()); + } + 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 CO_vec; + for (auto& indx : compiler_option) { + CO_vec.push_back(indx.get()); + } + 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 CO_vec; + for (auto& indx : compiler_option) { + CO_vec.push_back(indx.get()); + } + 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 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 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(json_file)), + std::istreambuf_iterator()); + 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(); + 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(); + std::string blk_name = block_obj.at("block_name").get(); + if (blk_name == block_name) { + if (para_name_to_retrieve == "compiler_option") { + std::string compiler_opt = + block_obj.at("compiler_option").get(); + returnValue += compiler_opt; + } else if (para_name_to_retrieve == "Target_Vals") { + std::string Target_Vals = + block_obj.at("Target_Vals").get(); + returnValue += Target_Vals; + } else if (para_name_to_retrieve == "kernel_name") { + std::string ker_name = block_obj.at("kernel_name").get(); + returnValue += ker_name; + } else if (para_name_to_retrieve == "reverse_compiler_option") { + std::string reverse = + block_obj.at("reverse_compiler_option").get(); + returnValue += reverse; + } else if (para_name_to_retrieve == "ready_compiler_option") { + std::string ready_CO = + block_obj.at("ready_compiler_option").get(); + 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(); + std::string blk_name = block_obj.at("block_name").get(); + if (blk_name == block_name) { + if (para_name_to_retrieve == "Target_Vals") { + picojson::array& Target_Vals = + block_obj.at("Target_Vals").get(); + return Target_Vals; + } else if (para_name_to_retrieve == "single_CO") { + picojson::array& single_CO = + block_obj.at("single_CO").get(); + return single_CO; + } else if (para_name_to_retrieve == "Combi_CO") { + picojson::array& Combi_CO = + block_obj.at("Combi_CO").get(); + return Combi_CO; + } else if (para_name_to_retrieve == "Input_Vals") { + picojson::array& Input_Vals = + block_obj.at("Input_Vals").get(); + return Input_Vals; + } else if (para_name_to_retrieve == "Expected_Results") { + picojson::array& Expected = + block_obj.at("Expected_Results").get(); + 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(); + return Expected_for_no; + } else if (para_name_to_retrieve == "compiler_option") { + picojson::array& compiler_option = + block_obj.at("compiler_option").get(); + return compiler_option; + } else if (para_name_to_retrieve == "reverse_compiler_option") { + picojson::array& reverse_compiler_option = + block_obj.at("reverse_compiler_option").get(); + return reverse_compiler_option; + } else if (para_name_to_retrieve == "Headers") { + picojson::array& Headers = + block_obj.at("Headers").get(); + return Headers; + } else if (para_name_to_retrieve == "Src_headers") { + picojson::array& Src_headers = + block_obj.at("Src_headers").get(); + return Src_headers; + } else if (para_name_to_retrieve == "depending_comp_optn") { + picojson::array& depending_comp_optn = + block_obj.at("depending_comp_optn").get(); + 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(); +} diff --git a/catch/unit/rtc/headers/RtcFact.h b/catch/unit/rtc/headers/RtcFact.h new file mode 100644 index 0000000000..43e19bff41 --- /dev/null +++ b/catch/unit/rtc/headers/RtcFact.h @@ -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_ diff --git a/catch/unit/rtc/headers/RtcFunctions.h b/catch/unit/rtc/headers/RtcFunctions.h new file mode 100644 index 0000000000..46f0d27810 --- /dev/null +++ b/catch/unit/rtc/headers/RtcFunctions.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 + +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_ diff --git a/catch/unit/rtc/headers/RtcKernels.h b/catch/unit/rtc/headers/RtcKernels.h new file mode 100644 index 0000000000..f3ff6aac83 --- /dev/null +++ b/catch/unit/rtc/headers/RtcKernels.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 +#include +#include + +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_ diff --git a/catch/unit/rtc/headers/RtcUtility.h b/catch/unit/rtc/headers/RtcUtility.h new file mode 100644 index 0000000000..c7fdd71372 --- /dev/null +++ b/catch/unit/rtc/headers/RtcUtility.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 +#include +#include + +std::vector get_combi_string_vec(); + +int split_comb_string(std::string option); + +int calling_combination_function(std::vector 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_ diff --git a/catch/unit/rtc/headers/printf_common.h b/catch/unit/rtc/headers/printf_common.h new file mode 100644 index 0000000000..bd17d24f33 --- /dev/null +++ b/catch/unit/rtc/headers/printf_common.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 +#include +#include +#include +#include +#include +#include +#include + +#if defined(_WIN32) +#include +#else +#include +#include +#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(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(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_ diff --git a/catch/unit/rtc/hiprtcComplrOptnTesting.cc b/catch/unit/rtc/hiprtcComplrOptnTesting.cc new file mode 100644 index 0000000000..53809539cd --- /dev/null +++ b/catch/unit/rtc/hiprtcComplrOptnTesting.cc @@ -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 +#include +#include +#include +#include +#include +#include +#include +#include +#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 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); +}