EXSWHTEC-288 - Implement tests for special-purpose device math function #235
Change-Id: Ifa601f462a5291ef6454240191d10e245eef8d7c
[ROCm/hip-tests commit: ce167afa87]
Este commit está contenido en:
Nives Vukovic
Rakesh Roy
@@ -28,6 +28,7 @@ set(TEST_SRC
|
||||
root_funcs.cc
|
||||
pow_funcs.cc
|
||||
log_funcs.cc
|
||||
special_funcs.cc
|
||||
)
|
||||
|
||||
if(HIP_PLATFORM MATCHES "nvidia")
|
||||
@@ -96,3 +97,7 @@ add_test(NAME Unit_Device_log_Negative
|
||||
COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py
|
||||
${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH}
|
||||
math_log_negative_kernels.cc 24)
|
||||
add_test(NAME Unit_Device_special_funcs_Negative
|
||||
COMMAND python3 ${CMAKE_CURRENT_SOURCE_DIR}/../compileAndCaptureOutput.py
|
||||
${CMAKE_CURRENT_SOURCE_DIR} ${HIP_PLATFORM} ${HIP_PATH}
|
||||
math_special_func_kernels.cc 76)
|
||||
|
||||
@@ -0,0 +1,60 @@
|
||||
/*
|
||||
Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
|
||||
#define NEGATIVE_KERNELS_SHELL_ONE_ARG(func_name) \
|
||||
__global__ void func_name##_kernel_v1(double* x) { double result = func_name(x); } \
|
||||
__global__ void func_name##_kernel_v2(Dummy x) { double result = func_name(x); } \
|
||||
__global__ void func_name##f_kernel_v1(float* x) { float result = func_name##f(x); } \
|
||||
__global__ void func_name##f_kernel_v2(Dummy x) { float result = func_name##f(x); }
|
||||
|
||||
#define NEGATIVE_KERNELS_SHELL_TWO_ARGS(func_name) \
|
||||
__global__ void func_name##_kernel_v1(int* x, double y) { double result = func_name(x, y); } \
|
||||
__global__ void func_name##_kernel_v2(int x, double* y) { double result = func_name(x, y); } \
|
||||
__global__ void func_name##_kernel_v3(Dummy x, double y) { double result = func_name(x, y); } \
|
||||
__global__ void func_name##_kernel_v4(int x, Dummy y) { double result = func_name(x, y); } \
|
||||
__global__ void func_name##f_kernel_v1(int* x, float y) { float result = func_name##f(x, y); } \
|
||||
__global__ void func_name##f_kernel_v2(int x, float* y) { float result = func_name##f(x, y); } \
|
||||
__global__ void func_name##f_kernel_v3(Dummy x, float y) { float result = func_name##f(x, y); } \
|
||||
__global__ void func_name##f_kernel_v4(int x, Dummy y) { float result = func_name##f(x, y); }
|
||||
|
||||
NEGATIVE_KERNELS_SHELL_ONE_ARG(erf)
|
||||
NEGATIVE_KERNELS_SHELL_ONE_ARG(erfc)
|
||||
NEGATIVE_KERNELS_SHELL_ONE_ARG(erfinv)
|
||||
NEGATIVE_KERNELS_SHELL_ONE_ARG(erfcinv)
|
||||
NEGATIVE_KERNELS_SHELL_ONE_ARG(erfcx)
|
||||
NEGATIVE_KERNELS_SHELL_ONE_ARG(normcdf)
|
||||
NEGATIVE_KERNELS_SHELL_ONE_ARG(normcdfinv)
|
||||
NEGATIVE_KERNELS_SHELL_ONE_ARG(lgamma)
|
||||
NEGATIVE_KERNELS_SHELL_ONE_ARG(tgamma)
|
||||
NEGATIVE_KERNELS_SHELL_ONE_ARG(j0)
|
||||
NEGATIVE_KERNELS_SHELL_ONE_ARG(j1)
|
||||
NEGATIVE_KERNELS_SHELL_TWO_ARGS(jn)
|
||||
NEGATIVE_KERNELS_SHELL_ONE_ARG(y0)
|
||||
NEGATIVE_KERNELS_SHELL_ONE_ARG(y1)
|
||||
NEGATIVE_KERNELS_SHELL_TWO_ARGS(yn)
|
||||
NEGATIVE_KERNELS_SHELL_ONE_ARG(cyl_bessel_i0)
|
||||
NEGATIVE_KERNELS_SHELL_ONE_ARG(cyl_bessel_i1)
|
||||
@@ -0,0 +1,236 @@
|
||||
/*
|
||||
Copyright (c) 2021 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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
/*
|
||||
Negative kernels used for the math special function negative Test Cases that are using RTC.
|
||||
*/
|
||||
|
||||
static constexpr auto kErf{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void erf_kernel_v1(double* x) { double result = erf(x); }
|
||||
__global__ void erf_kernel_v2(Dummy x) { double result = erf(x); }
|
||||
__global__ void erff_kernel_v1(float* x) { float result = erff(x); }
|
||||
__global__ void erff_kernel_v2(Dummy x) { float result = erff(x); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kErfc{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void erfc_kernel_v1(double* x) { double result = erfc(x); }
|
||||
__global__ void erfc_kernel_v2(Dummy x) { double result = erfc(x); }
|
||||
__global__ void erfcf_kernel_v1(float* x) { float result = erfcf(x); }
|
||||
__global__ void erfcf_kernel_v2(Dummy x) { float result = erfcf(x); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kErfinv{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void erfinv_kernel_v1(double* x) { double result = erfinv(x); }
|
||||
__global__ void erfinv_kernel_v2(Dummy x) { double result = erfinv(x); }
|
||||
__global__ void erfinvf_kernel_v1(float* x) { float result = erfinvf(x); }
|
||||
__global__ void erfinvf_kernel_v2(Dummy x) { float result = erfinvf(x); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kErfcinv{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void erfcinv_kernel_v1(double* x) { double result = erfcinv(x); }
|
||||
__global__ void erfcinv_kernel_v2(Dummy x) { double result = erfcinv(x); }
|
||||
__global__ void erfcinvf_kernel_v1(float* x) { float result = erfcinvf(x); }
|
||||
__global__ void erfcinvf_kernel_v2(Dummy x) { float result = erfcinvf(x); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kErfcx{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void erfcx_kernel_v1(double* x) { double result = erfcx(x); }
|
||||
__global__ void erfcx_kernel_v2(Dummy x) { double result = erfcx(x); }
|
||||
__global__ void erfcxf_kernel_v1(float* x) { float result = erfcxf(x); }
|
||||
__global__ void erfcxf_kernel_v2(Dummy x) { float result = erfcxf(x); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kNormcdf{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void normcdf_kernel_v1(double* x) { double result = normcdf(x); }
|
||||
__global__ void normcdf_kernel_v2(Dummy x) { double result = normcdf(x); }
|
||||
__global__ void normcdff_kernel_v1(float* x) { float result = normcdff(x); }
|
||||
__global__ void normcdff_kernel_v2(Dummy x) { float result = normcdff(x); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kNormcdfinv{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void normcdfinv_kernel_v1(double* x) { double result = normcdfinv(x); }
|
||||
__global__ void normcdfinv_kernel_v2(Dummy x) { double result = normcdfinv(x); }
|
||||
__global__ void normcdfinvf_kernel_v1(float* x) { float result = normcdfinvf(x); }
|
||||
__global__ void normcdfinvf_kernel_v2(Dummy x) { float result = normcdfinvf(x); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kLgamma{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void lgamma_kernel_v1(double* x) { double result = lgamma(x); }
|
||||
__global__ void lgamma_kernel_v2(Dummy x) { double result = lgamma(x); }
|
||||
__global__ void lgammaf_kernel_v1(float* x) { float result = lgammaf(x); }
|
||||
__global__ void lgammaf_kernel_v2(Dummy x) { float result = lgammaf(x); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kTgamma{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void tgamma_kernel_v1(double* x) { double result = tgamma(x); }
|
||||
__global__ void tgamma_kernel_v2(Dummy x) { double result = tgamma(x); }
|
||||
__global__ void tgammaf_kernel_v1(float* x) { float result = tgammaf(x); }
|
||||
__global__ void tgammaf_kernel_v2(Dummy x) { float result = tgammaf(x); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kJ0{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void j0_kernel_v1(double* x) { double result = j0(x); }
|
||||
__global__ void j0_kernel_v2(Dummy x) { double result = j0(x); }
|
||||
__global__ void j0f_kernel_v1(float* x) { float result = j0f(x); }
|
||||
__global__ void j0f_kernel_v2(Dummy x) { float result = j0f(x); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kJ1{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void j1_kernel_v1(double* x) { double result = j1(x); }
|
||||
__global__ void j1_kernel_v2(Dummy x) { double result = j1(x); }
|
||||
__global__ void j1f_kernel_v1(float* x) { float result = j1f(x); }
|
||||
__global__ void j1f_kernel_v2(Dummy x) { float result = j1f(x); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kJn{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void jn_kernel_v1(int* x, double y) { double result = jn(x, y); }
|
||||
__global__ void jn_kernel_v2(int x, double* y) { double result = jn(x, y); }
|
||||
__global__ void jn_kernel_v3(Dummy x, double y) { double result = jn(x, y); }
|
||||
__global__ void jn_kernel_v4(int x, Dummy y) { double result = jn(x, y); }
|
||||
__global__ void jnf_kernel_v1(int* x, float y) { float result = jnf(x, y); }
|
||||
__global__ void jnf_kernel_v2(int x, float* y) { float result = jnf(x, y); }
|
||||
__global__ void jnf_kernel_v3(Dummy x, float y) { float result = jnf(x, y); }
|
||||
__global__ void jnf_kernel_v4(int x, Dummy y) { float result = jnf(x, y); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kY0{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void y0_kernel_v1(double* x) { double result = y0(x); }
|
||||
__global__ void y0_kernel_v2(Dummy x) { double result = y0(x); }
|
||||
__global__ void y0f_kernel_v1(float* x) { float result = y0f(x); }
|
||||
__global__ void y0f_kernel_v2(Dummy x) { float result = y0f(x); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kY1{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void y1_kernel_v1(double* x) { double result = y1(x); }
|
||||
__global__ void y1_kernel_v2(Dummy x) { double result = y1(x); }
|
||||
__global__ void y1f_kernel_v1(float* x) { float result = y1f(x); }
|
||||
__global__ void y1f_kernel_v2(Dummy x) { float result = y1f(x); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kYn{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void yn_kernel_v1(int* x, double y) { double result = yn(x, y); }
|
||||
__global__ void yn_kernel_v2(int x, double* y) { double result = yn(x, y); }
|
||||
__global__ void yn_kernel_v3(Dummy x, double y) { double result = yn(x, y); }
|
||||
__global__ void yn_kernel_v4(int x, Dummy y) { double result = yn(x, y); }
|
||||
__global__ void ynf_kernel_v1(int* x, float y) { float result = ynf(x, y); }
|
||||
__global__ void ynf_kernel_v2(int x, float* y) { float result = ynf(x, y); }
|
||||
__global__ void ynf_kernel_v3(Dummy x, float y) { float result = ynf(x, y); }
|
||||
__global__ void ynf_kernel_v4(int x, Dummy y) { float result = ynf(x, y); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kCylBesselI0{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void cyl_bessel_i0_kernel_v1(double* x) { double result = cyl_bessel_i0(x); }
|
||||
__global__ void cyl_bessel_i0_kernel_v2(Dummy x) { double result = cyl_bessel_i0(x); }
|
||||
__global__ void cyl_bessel_i0f_kernel_v1(float* x) { float result = cyl_bessel_i0f(x); }
|
||||
__global__ void cyl_bessel_i0f_kernel_v2(Dummy x) { float result = cyl_bessel_i0f(x); }
|
||||
)"};
|
||||
|
||||
static constexpr auto kCylBesselI1{R"(
|
||||
class Dummy {
|
||||
public:
|
||||
__device__ Dummy() {}
|
||||
__device__ ~Dummy() {}
|
||||
};
|
||||
__global__ void cyl_bessel_i1_kernel_v1(double* x) { double result = cyl_bessel_i1(x); }
|
||||
__global__ void cyl_bessel_i1_kernel_v2(Dummy x) { double result = cyl_bessel_i1(x); }
|
||||
__global__ void cyl_bessel_i1f_kernel_v1(float* x) { float result = cyl_bessel_i1f(x); }
|
||||
__global__ void cyl_bessel_i1f_kernel_v2(Dummy x) { float result = cyl_bessel_i1f(x); }
|
||||
)"};
|
||||
@@ -0,0 +1,145 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
#include "math_common.hh"
|
||||
#include "math_special_values.hh"
|
||||
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
|
||||
namespace cg = cooperative_groups;
|
||||
|
||||
#define MATH_BESSEL_N_KERNEL_DEF(func_name) \
|
||||
template <typename T> \
|
||||
__global__ void func_name##_kernel(T* const ys, const size_t num_xs, int* n, T* const xs) { \
|
||||
const auto tid = cg::this_grid().thread_rank(); \
|
||||
const auto stride = cg::this_grid().size(); \
|
||||
\
|
||||
for (auto i = tid; i < num_xs; i += stride) { \
|
||||
if constexpr (std::is_same_v<float, T>) { \
|
||||
ys[i] = func_name##f(n[i], xs[i]); \
|
||||
} else if constexpr (std::is_same_v<double, T>) { \
|
||||
ys[i] = func_name(n[i], xs[i]); \
|
||||
} \
|
||||
} \
|
||||
}
|
||||
|
||||
template <typename T> using kernel_bessel_n_sig = void (*)(T*, const size_t, int*, T*);
|
||||
|
||||
template <typename T> using ref_bessel_n_sig = T (*)(int, T);
|
||||
|
||||
template <typename ValidatorBuilder>
|
||||
void BesselDoublePrecisionBruteForceTest(kernel_bessel_n_sig<double> kernel,
|
||||
ref_bessel_n_sig<long double> ref_func,
|
||||
const ValidatorBuilder& validator_builder, int n_input = 0,
|
||||
const double a = std::numeric_limits<double>::lowest(),
|
||||
const double b = std::numeric_limits<double>::max()) {
|
||||
const auto [grid_size, block_size] = GetOccupancyMaxPotentialBlockSize(kernel);
|
||||
const uint64_t num_iterations = GetTestIterationCount();
|
||||
const auto max_batch_size = std::min(
|
||||
GetMaxAllowedDeviceMemoryUsage() / (sizeof(double) * 2 + sizeof(int)), num_iterations);
|
||||
LinearAllocGuard<int> x1s{LinearAllocs::hipHostMalloc, max_batch_size * sizeof(int)};
|
||||
LinearAllocGuard<double> x2s{LinearAllocs::hipHostMalloc, max_batch_size * sizeof(double)};
|
||||
|
||||
MathTest math_test(kernel, max_batch_size);
|
||||
std::fill_n(x1s.ptr(), max_batch_size, n_input);
|
||||
|
||||
auto batch_size = max_batch_size;
|
||||
const auto num_threads = thread_pool.thread_count();
|
||||
for (uint64_t i = 0ul; i < num_iterations; i += batch_size) {
|
||||
batch_size = std::min<uint64_t>(max_batch_size, num_iterations - i);
|
||||
|
||||
const auto min_sub_batch_size = batch_size / num_threads;
|
||||
const auto tail = batch_size % num_threads;
|
||||
|
||||
auto base_idx = 0u;
|
||||
for (auto i = 0u; i < num_threads; ++i) {
|
||||
const auto sub_batch_size = min_sub_batch_size + (i < tail);
|
||||
thread_pool.Post([=, &x2s] {
|
||||
const auto generator = [=] {
|
||||
static thread_local std::mt19937 rng(std::random_device{}());
|
||||
std::uniform_real_distribution<RefType_t<double>> unif_dist(a, b);
|
||||
return static_cast<double>(unif_dist(rng));
|
||||
};
|
||||
std::generate(x2s.ptr() + base_idx, x2s.ptr() + base_idx + sub_batch_size, generator);
|
||||
});
|
||||
base_idx += sub_batch_size;
|
||||
}
|
||||
|
||||
thread_pool.Wait();
|
||||
|
||||
math_test.Run(validator_builder, grid_size, block_size, ref_func, batch_size, x1s.ptr(),
|
||||
x2s.ptr());
|
||||
}
|
||||
}
|
||||
|
||||
template <typename ValidatorBuilder>
|
||||
void BesselSinglePrecisionRangeTest(kernel_bessel_n_sig<float> kernel,
|
||||
ref_bessel_n_sig<double> ref_func,
|
||||
const ValidatorBuilder& validator_builder, int n_input,
|
||||
const float a, const float b) {
|
||||
const auto [grid_size, block_size] = GetOccupancyMaxPotentialBlockSize(kernel);
|
||||
const auto max_batch_size = GetMaxAllowedDeviceMemoryUsage() / (sizeof(float) * 2 + sizeof(int));
|
||||
LinearAllocGuard<int> x1s{LinearAllocs::hipHostMalloc, max_batch_size * sizeof(int)};
|
||||
LinearAllocGuard<float> x2s{LinearAllocs::hipHostMalloc, max_batch_size * sizeof(float)};
|
||||
|
||||
MathTest math_test(kernel, max_batch_size);
|
||||
std::fill_n(x1s.ptr(), max_batch_size, n_input);
|
||||
|
||||
size_t inserted = 0u;
|
||||
for (float v = a; v != b; v = std::nextafter(v, b)) {
|
||||
x2s.ptr()[inserted++] = v;
|
||||
if (inserted < max_batch_size) continue;
|
||||
|
||||
math_test.Run(validator_builder, grid_size, block_size, ref_func, inserted, x1s.ptr(),
|
||||
x2s.ptr());
|
||||
inserted = 0u;
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T, typename F, typename ValidatorBuilder>
|
||||
void SpecialSimpleTest(F kernel, const ValidatorBuilder& validator_builder, const T* x,
|
||||
const T* ref, size_t num_args) {
|
||||
LinearAllocGuard<T> x_dev{LinearAllocs::hipMalloc, num_args * sizeof(T)};
|
||||
LinearAllocGuard<T> y{LinearAllocs::hipHostMalloc, num_args * sizeof(T)};
|
||||
LinearAllocGuard<T> y_dev{LinearAllocs::hipMalloc, num_args * sizeof(T)};
|
||||
|
||||
HIP_CHECK(hipMemcpy(x_dev.ptr(), x, num_args * sizeof(T), hipMemcpyHostToDevice));
|
||||
|
||||
kernel<<<1, num_args>>>(y_dev.ptr(), num_args, x_dev.ptr());
|
||||
HIP_CHECK(hipGetLastError());
|
||||
|
||||
HIP_CHECK(hipMemcpy(y.ptr(), y_dev.ptr(), num_args * sizeof(T), hipMemcpyDeviceToHost));
|
||||
|
||||
for (auto i = 0u; i < num_args; ++i) {
|
||||
const auto actual_val = y.ptr()[i];
|
||||
const auto ref_val = ref[i];
|
||||
const auto validator = validator_builder(ref_val);
|
||||
|
||||
if (!validator->match(actual_val)) {
|
||||
std::stringstream ss;
|
||||
ss << "Input value(s): " << std::scientific
|
||||
<< std::setprecision(std::numeric_limits<T>::max_digits10 - 1);
|
||||
ss << x[i] << " " << actual_val << " " << ref_val << "\n";
|
||||
INFO(ss.str());
|
||||
REQUIRE(false);
|
||||
}
|
||||
}
|
||||
}
|
||||
La diferencia del archivo ha sido suprimido porque es demasiado grande
Cargar Diff
Referencia en una nueva incidencia
Block a user