SWDEV-472723 - Correct file format and remove trailing spaces
Change-Id: Ie40c763e9391fa36d6c890cd0a171659a1502a83
This commit is contained in:
Julia Jiang
@@ -1,178 +1,178 @@
|
||||
/*
|
||||
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
The Functions defined in RtcFunctions.cpp are declared here in RtcFunctions.h.
|
||||
*/
|
||||
|
||||
#ifndef CATCH_UNIT_RTC_HEADERS_RTCFUNCTIONS_H_
|
||||
#define CATCH_UNIT_RTC_HEADERS_RTCFUNCTIONS_H_
|
||||
#include <string>
|
||||
|
||||
bool check_architecture(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_rdc(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_denormals_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_denormals_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_ffp_contract_off(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_ffp_contract_on(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_ffp_contract_fast(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_fast_math_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_fast_math_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_slp_vectorize_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_slp_vectorize_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_macro(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_undef_macro(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_header_dir(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_warning(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_Rpass_inline(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_conversionerror_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_conversionerror_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size,
|
||||
int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_conversionwarning_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size,
|
||||
int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_conversionwarning_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size,
|
||||
int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_max_thread(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_unsafe_atomic_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_unsafe_atomic_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_infinite_num_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_infinite_num_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_NAN_num_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_NAN_num_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_finite_math_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_finite_math_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_associative_math_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size,
|
||||
int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_associative_math_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size,
|
||||
int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_signed_zeros_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size,
|
||||
int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_signed_zeros_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_trapping_math_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_trapping_math_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
std::string checking_IR(const char* kername, const char** extra_CO_IRadded,
|
||||
int extra_CO_IRadded_size, const char** Combination_CO,
|
||||
int Combination_CO_size);
|
||||
|
||||
#endif // CATCH_UNIT_RTC_HEADERS_RTCFUNCTIONS_H_
|
||||
/*
|
||||
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
The Functions defined in RtcFunctions.cpp are declared here in RtcFunctions.h.
|
||||
*/
|
||||
|
||||
#ifndef CATCH_UNIT_RTC_HEADERS_RTCFUNCTIONS_H_
|
||||
#define CATCH_UNIT_RTC_HEADERS_RTCFUNCTIONS_H_
|
||||
#include <string>
|
||||
|
||||
bool check_architecture(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_rdc(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_denormals_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_denormals_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_ffp_contract_off(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_ffp_contract_on(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_ffp_contract_fast(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_fast_math_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_fast_math_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_slp_vectorize_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_slp_vectorize_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_macro(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_undef_macro(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_header_dir(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_warning(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_Rpass_inline(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_conversionerror_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_conversionerror_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size,
|
||||
int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_conversionwarning_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size,
|
||||
int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_conversionwarning_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size,
|
||||
int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_max_thread(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_unsafe_atomic_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_unsafe_atomic_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_infinite_num_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_infinite_num_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_NAN_num_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_NAN_num_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_finite_math_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_finite_math_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_associative_math_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size,
|
||||
int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_associative_math_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size,
|
||||
int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_signed_zeros_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size,
|
||||
int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_signed_zeros_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_trapping_math_enabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
bool check_trapping_math_disabled(const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_pos,
|
||||
int fast_math_present);
|
||||
|
||||
std::string checking_IR(const char* kername, const char** extra_CO_IRadded,
|
||||
int extra_CO_IRadded_size, const char** Combination_CO,
|
||||
int Combination_CO_size);
|
||||
|
||||
#endif // CATCH_UNIT_RTC_HEADERS_RTCFUNCTIONS_H_
|
||||
|
||||
+163
-163
@@ -1,163 +1,163 @@
|
||||
/*
|
||||
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
RtcKernels.h contains the string's with the which includes the kernel code.
|
||||
They are utilized by the compiler option functions, defined in RtcFunctions.cpp
|
||||
*/
|
||||
|
||||
#ifndef CATCH_UNIT_RTC_HEADERS_RTCKERNELS_H_
|
||||
#define CATCH_UNIT_RTC_HEADERS_RTCKERNELS_H_
|
||||
#include <hip/hiprtc.h>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <math.h>
|
||||
|
||||
static constexpr auto max_thread_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void max_thread(int* a) {
|
||||
int BD = blockDim.x;
|
||||
*a = BD;
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto denormals_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void denormals(double* base, double* power, double* result) {
|
||||
float denorm = powf(*base, *power);
|
||||
if (*result == 0 || *result ==1 )
|
||||
*result = (denorm==0) ? 0 : 1;
|
||||
else
|
||||
*result = powf(*base, *power);
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto warning_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void warning() {
|
||||
#warning "Just printing a WARNING message onto the terminal";
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto fp32_div_sqrt_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void fp32_div_sqrt(float* result) {
|
||||
float input = 109.6209;
|
||||
*result = sqrt(input);
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto error_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void error() {
|
||||
unsigned int a = -1;
|
||||
unsigned int b = +1;
|
||||
signed int c = -1;
|
||||
signed int d = +1;
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto macro_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void macro(int *result) {
|
||||
*result = PI;
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto undef_macro_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void undef_macro() {
|
||||
int a = Z;
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto header_dir_string {
|
||||
R"(
|
||||
#include "RtcFact.h"
|
||||
extern "C"
|
||||
__global__ void header_dir(int* a, int* val) {
|
||||
*a = fact(*val);
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto rdc_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void rdc(float* a, float* b, float* c) {
|
||||
*c = *a * *b;
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto ffp_contract_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void ffp_contract(float* a, float* b, float* c) {
|
||||
*c = *a * *b + *c;
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto slp_vectorize_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void slp_vectorize(__half2 a, __half2 x, __half2 *y) {
|
||||
(*y).data.x = x.data.x + a.data.x;
|
||||
(*y).data.y = x.data.y + a.data.y;
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto unsafe_atomic_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void unsafe_atomic(float* a) {
|
||||
int id = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
if (id < 1000) {
|
||||
unsafeAtomicAdd(&a[id], 0.2f);
|
||||
}
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto amdgpu_ieee_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void amdgpu_ieee(float* a, float* b, float* c) {
|
||||
*c = sqrt(*a / *b);
|
||||
printf("sqrt(a * b) = %f\n", *c);
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto associative_math_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void associative_math(int* check) {
|
||||
double x = 0.1f;
|
||||
double y = 0.2f;
|
||||
double z = 0.3f;
|
||||
if((x*y)*z != x*(y*z))
|
||||
*check = 1;
|
||||
else *check = 0;
|
||||
}
|
||||
)"};
|
||||
|
||||
#endif // CATCH_UNIT_RTC_HEADERS_RTCKERNELS_H_
|
||||
/*
|
||||
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
RtcKernels.h contains the string's with the which includes the kernel code.
|
||||
They are utilized by the compiler option functions, defined in RtcFunctions.cpp
|
||||
*/
|
||||
|
||||
#ifndef CATCH_UNIT_RTC_HEADERS_RTCKERNELS_H_
|
||||
#define CATCH_UNIT_RTC_HEADERS_RTCKERNELS_H_
|
||||
#include <hip/hiprtc.h>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <math.h>
|
||||
|
||||
static constexpr auto max_thread_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void max_thread(int* a) {
|
||||
int BD = blockDim.x;
|
||||
*a = BD;
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto denormals_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void denormals(double* base, double* power, double* result) {
|
||||
float denorm = powf(*base, *power);
|
||||
if (*result == 0 || *result ==1 )
|
||||
*result = (denorm==0) ? 0 : 1;
|
||||
else
|
||||
*result = powf(*base, *power);
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto warning_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void warning() {
|
||||
#warning "Just printing a WARNING message onto the terminal";
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto fp32_div_sqrt_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void fp32_div_sqrt(float* result) {
|
||||
float input = 109.6209;
|
||||
*result = sqrt(input);
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto error_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void error() {
|
||||
unsigned int a = -1;
|
||||
unsigned int b = +1;
|
||||
signed int c = -1;
|
||||
signed int d = +1;
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto macro_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void macro(int *result) {
|
||||
*result = PI;
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto undef_macro_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void undef_macro() {
|
||||
int a = Z;
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto header_dir_string {
|
||||
R"(
|
||||
#include "RtcFact.h"
|
||||
extern "C"
|
||||
__global__ void header_dir(int* a, int* val) {
|
||||
*a = fact(*val);
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto rdc_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void rdc(float* a, float* b, float* c) {
|
||||
*c = *a * *b;
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto ffp_contract_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void ffp_contract(float* a, float* b, float* c) {
|
||||
*c = *a * *b + *c;
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto slp_vectorize_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void slp_vectorize(__half2 a, __half2 x, __half2 *y) {
|
||||
(*y).data.x = x.data.x + a.data.x;
|
||||
(*y).data.y = x.data.y + a.data.y;
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto unsafe_atomic_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void unsafe_atomic(float* a) {
|
||||
int id = threadIdx.x + blockIdx.x * blockDim.x;
|
||||
if (id < 1000) {
|
||||
unsafeAtomicAdd(&a[id], 0.2f);
|
||||
}
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto amdgpu_ieee_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void amdgpu_ieee(float* a, float* b, float* c) {
|
||||
*c = sqrt(*a / *b);
|
||||
printf("sqrt(a * b) = %f\n", *c);
|
||||
}
|
||||
)"};
|
||||
|
||||
static constexpr auto associative_math_string {
|
||||
R"(
|
||||
extern "C"
|
||||
__global__ void associative_math(int* check) {
|
||||
double x = 0.1f;
|
||||
double y = 0.2f;
|
||||
double z = 0.3f;
|
||||
if((x*y)*z != x*(y*z))
|
||||
*check = 1;
|
||||
else *check = 0;
|
||||
}
|
||||
)"};
|
||||
|
||||
#endif // CATCH_UNIT_RTC_HEADERS_RTCKERNELS_H_
|
||||
|
||||
@@ -1,53 +1,53 @@
|
||||
/*
|
||||
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
The Functions defined in RtcUtility.cpp are declared here in RtcUtility.h.
|
||||
*/
|
||||
|
||||
#ifndef CATCH_UNIT_RTC_HEADERS_RTCUTILITY_H_
|
||||
#define CATCH_UNIT_RTC_HEADERS_RTCUTILITY_H_
|
||||
#include <picojson.h>
|
||||
#include <vector>
|
||||
#include <string>
|
||||
|
||||
std::vector<std::string> get_combi_string_vec();
|
||||
|
||||
int split_comb_string(std::string option);
|
||||
|
||||
int calling_combination_function(std::vector<std::string> combi_vec_list);
|
||||
|
||||
int check_positive_CO_present(std::string find_string);
|
||||
|
||||
int check_negative_CO_present(std::string find_string);
|
||||
|
||||
bool calling_resp_function(const std::string block_name,
|
||||
const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_position,
|
||||
int fast_math_present);
|
||||
|
||||
picojson::array getblock_fromconfig();
|
||||
|
||||
std::string get_string_parameters(std::string para_name_to_retrieve,
|
||||
std::string block_name);
|
||||
|
||||
picojson::array get_array_parameters(std::string para_name_to_retrieve,
|
||||
std::string block_name);
|
||||
|
||||
#endif // CATCH_UNIT_RTC_HEADERS_RTCUTILITY_H_
|
||||
/*
|
||||
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/*
|
||||
The Functions defined in RtcUtility.cpp are declared here in RtcUtility.h.
|
||||
*/
|
||||
|
||||
#ifndef CATCH_UNIT_RTC_HEADERS_RTCUTILITY_H_
|
||||
#define CATCH_UNIT_RTC_HEADERS_RTCUTILITY_H_
|
||||
#include <picojson.h>
|
||||
#include <vector>
|
||||
#include <string>
|
||||
|
||||
std::vector<std::string> get_combi_string_vec();
|
||||
|
||||
int split_comb_string(std::string option);
|
||||
|
||||
int calling_combination_function(std::vector<std::string> combi_vec_list);
|
||||
|
||||
int check_positive_CO_present(std::string find_string);
|
||||
|
||||
int check_negative_CO_present(std::string find_string);
|
||||
|
||||
bool calling_resp_function(const std::string block_name,
|
||||
const char** Combination_CO,
|
||||
int Combination_CO_size, int max_thread_position,
|
||||
int fast_math_present);
|
||||
|
||||
picojson::array getblock_fromconfig();
|
||||
|
||||
std::string get_string_parameters(std::string para_name_to_retrieve,
|
||||
std::string block_name);
|
||||
|
||||
picojson::array get_array_parameters(std::string para_name_to_retrieve,
|
||||
std::string block_name);
|
||||
|
||||
#endif // CATCH_UNIT_RTC_HEADERS_RTCUTILITY_H_
|
||||
|
||||
Reference in New Issue
Block a user