AI/rocm-systems
2
0
Fork 0
Código Issues Pull Requests Ações Pacotes Projetos Versões Wiki Atividade

SWDEV-380340 - [catch2][dtest] DeviceLib tests migrated from direct to catch2 (#225)

Ver Código-Fonte 
Change-Id: Ie2ec1c7dabdfedbe0bd36fd2525df7dc9d9ba2e5
Esse commit está contido em:
ROCm CI Service Account
2023-08-14 20:52:26 +05:30
commit de GitHub
pai cdf434b357
commit 3447a59895
19 arquivos alterados com 4074 adições e 61 exclusões
Baixar Arquivo de Patch Baixar Arquivo de Diff Expandir todos os arquivos Colapsar todos os arquivos
catch/unit/deviceLib/CMakeLists.txt
+46
Ver Arquivo
@@ -1,3 +1,23 @@
# 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.
# Common Tests - Test independent of all platforms
set(TEST_SRC
floatMath.cc
@@ -27,7 +47,26 @@ set(TEST_SRC
hipTestAtomicAdd.cc
hipStdComplex.cc
hipTestClock.cc
hip_trig.cc
hipDeviceMemcpy.cc
hipTestIncludeMath.cc
hipTestDotFunctions.cc
hipTestDeviceSymbol.cc
hipTestNew.cc
hipThreadFence.cc
hipTestDevice.cc
hipTestDeviceLimit.cc
hipTestDeviceDouble.cc
hipTestHost.cc
)
if(HIP_PLATFORM MATCHES "nvidia")
set_source_files_properties(hipTestHost.cc PROPERTIES COMPILE_OPTIONS "--expt-relaxed-constexpr")
endif()
if(UNIX)
set(TEST_SRC ${TEST_SRC}
deviceAllocation.cc)
endif()
# AMD only tests
set(AMD_TEST_SRC
@@ -38,6 +77,13 @@ set(AMD_TEST_SRC
floatTM.cc
hipMathFunctions.cc
hmax_hmin.cc
hipBfloat16.cc
hipVectorTypes.cc
hipTestHalf.cc
hipComplex.cc
hipTestFMA.cc
hipTestNativeHalf.cc
hip_test_make_type.cc
bfloat16.cc
)
set(AMD_ARCH_SPEC_TEST_SRC
catch/unit/deviceLib/hipBfloat16.cc
+130
Ver Arquivo
@@ -0,0 +1,130 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip/hip_bfloat16.h>
#include <type_traits>
#include <random>
#include <climits>
#define SIZE 100
static std::random_device dev;
static std::mt19937 rng(dev());
inline float getRandomFloat(int16_t min = 10, int64_t max = LONG_MAX) {
std::uniform_real_distribution<float> gen(min, max);
return gen(rng);
}
__host__ __device__ bool testRelativeAccuracy(float a, hip_bfloat16 b) {
float c = static_cast<float>(b);
// float relative error should be less than 1/(2^7) since bfloat16
// has 7 bits mantissa.
if (fabs(c - a) / a <= 1.0 / 128) {
return true;
}
return false;
}
__host__ __device__ bool testOperations(const float &fa, const float &fb) {
bool testPass = true;
hip_bfloat16 bf_a(fa);
hip_bfloat16 bf_b(fb);
float fc = static_cast<float>(bf_a);
float fd = static_cast<float>(bf_b);
testPass &= testRelativeAccuracy(fa, bf_a);
testPass &= testRelativeAccuracy(fb, bf_b);
testPass &= testRelativeAccuracy(fc + fd, bf_a + bf_b);
// when checked as above for add, operation sub fails on GPU
if (hip_bfloat16(fc - fd) == (bf_a - bf_b)) {
testPass &= true;
}
testPass &= testRelativeAccuracy(fc * fd, bf_a * bf_b);
testPass &= testRelativeAccuracy(fc / fd, bf_a / bf_b);
hip_bfloat16 bf_x;
bf_x = bf_a;
bf_x++;
bf_x--;
++bf_x;
--bf_x;
// hip_bfloat16 is converted to float and then inc/decremented,
// hence check with reduced precision
testPass &= testRelativeAccuracy(bf_x, bf_a);
bf_x = bf_a;
bf_x += bf_b;
bf_x = bf_a;
bf_x -= bf_b;
bf_x = bf_a;
bf_x *= bf_b;
bf_x = bf_a;
bf_x /= bf_b;
hip_bfloat16 bf_rounded = hip_bfloat16::round_to_bfloat16(fa);
if (std::isnan(bf_rounded)) {
if (std::isnan(bf_rounded) || std::isinf(bf_rounded)) {
testPass &= true;
}
}
return testPass;
}
__global__ void testOperationsGPU(float* d_a, float* d_b, bool *testPass) {
int id = threadIdx.x;
if (id > SIZE) return;
float &a = d_a[id];
float &b = d_b[id];
*testPass = testOperations(a, b);
}
TEST_CASE("Unit_hipBfloat16") {
float *h_fa, *h_fb;
float *d_fa, *d_fb;
bool *d_fc, h_fc = false;
h_fa = new float[SIZE];
h_fb = new float[SIZE];
bool result = false;
for (int i = 0; i < SIZE; i++) {
h_fa[i] = getRandomFloat();
h_fb[i] = getRandomFloat();
result = testOperations(h_fa[i], h_fb[i]);
REQUIRE(result == true);
}
HIP_CHECK(hipMalloc(&d_fa, sizeof(float) * SIZE));
HIP_CHECK(hipMalloc(&d_fb, sizeof(float) * SIZE));
HIP_CHECK(hipMalloc(&d_fc, sizeof(bool)));
HIP_CHECK(hipMemcpy(d_fa, h_fa, sizeof(float) * SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(d_fb, h_fb, sizeof(float) * SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(d_fc, &h_fc, sizeof(bool), hipMemcpyHostToDevice));
hipLaunchKernelGGL(testOperationsGPU, 1, SIZE, 0, 0, d_fa, d_fb, d_fc);
HIP_CHECK(hipDeviceSynchronize());
HIP_CHECK(hipMemcpy(&h_fc, d_fc, sizeof(bool), hipMemcpyDeviceToHost));
REQUIRE(h_fc == true);
delete[] h_fa;
delete[] h_fb;
HIP_CHECK(hipFree(d_fa));
HIP_CHECK(hipFree(d_fb));
HIP_CHECK(hipFree(d_fc));
}
catch/unit/deviceLib/hipComplex.cc
+438
Ver Arquivo
@@ -0,0 +1,438 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip/hip_complex.h>
#include <math.h>
#include <iostream>
#include <type_traits>
#define LEN 64
/* Comparing 2 floating point/double variables using floating point
precision. The precision is set at compile time using EPSILON. */
#define COMPARE_REALNUM(A, B, EPSILON) (fabs(A-B) < EPSILON)
enum ComplexFuncType {
COMPLEX_ADD,
COMPLEX_SUB,
COMPLEX_MUL,
COMPLEX_DIV,
COMPLEX_CONJ,
COMPLEX_REAL,
COMPLEX_IMAG,
COMPLEX_SQABS,
COMPLEX_ABS
};
__global__ static void testMakeComplexFunc(float* A, float* B,
hipFloatComplex* C) {
int tx = threadIdx.x + blockIdx.x * blockDim.x;
C[tx] = make_hipFloatComplex(A[tx], B[tx]);
}
__global__ static void testMakeComplexFunc(double* A, double* B,
hipDoubleComplex* C) {
int tx = threadIdx.x + blockIdx.x * blockDim.x;
C[tx] = make_hipDoubleComplex(A[tx], B[tx]);
}
__global__ static void testComplexMathFunc1(hipFloatComplex* A,
hipFloatComplex* B,
hipFloatComplex* C,
enum ComplexFuncType type) {
int tx = threadIdx.x + blockIdx.x * blockDim.x;
switch (type) {
case COMPLEX_ADD:
C[tx] = hipCaddf(A[tx], B[tx]);
break;
case COMPLEX_SUB:
C[tx] = hipCsubf(A[tx], B[tx]);
break;
case COMPLEX_MUL:
C[tx] = hipCmulf(A[tx], B[tx]);
break;
case COMPLEX_DIV:
C[tx] = hipCdivf(A[tx], B[tx]);
break;
case COMPLEX_CONJ:
C[tx] = hipConjf(A[tx]);
break;
default:
break;
}
}
__global__ static void testComplexMathFunc1(hipDoubleComplex* A,
hipDoubleComplex* B,
hipDoubleComplex* C,
enum ComplexFuncType type) {
int tx = threadIdx.x + blockIdx.x * blockDim.x;
switch (type) {
case COMPLEX_ADD:
C[tx] = hipCadd(A[tx], B[tx]);
break;
case COMPLEX_SUB:
C[tx] = hipCsub(A[tx], B[tx]);
break;
case COMPLEX_MUL:
C[tx] = hipCmul(A[tx], B[tx]);
break;
case COMPLEX_DIV:
C[tx] = hipCdiv(A[tx], B[tx]);
break;
case COMPLEX_CONJ:
C[tx] = hipConj(A[tx]);
break;
default:
break;
}
}
__global__ static void testComplexMathFunc2(hipFloatComplex* A,
float* B,
enum ComplexFuncType type) {
int tx = threadIdx.x + blockIdx.x * blockDim.x;
switch (type) {
case COMPLEX_REAL:
B[tx] = hipCrealf(A[tx]);
break;
case COMPLEX_IMAG:
B[tx] = hipCimagf(A[tx]);
break;
case COMPLEX_SQABS:
B[tx] = hipCsqabsf(A[tx]);
break;
case COMPLEX_ABS:
B[tx] = hipCabsf(A[tx]);
break;
default:
break;
}
}
__global__ static void testComplexMathFunc2(hipDoubleComplex* A,
double* B,
enum ComplexFuncType type) {
int tx = threadIdx.x + blockIdx.x * blockDim.x;
switch (type) {
case COMPLEX_REAL:
B[tx] = hipCreal(A[tx]);
break;
case COMPLEX_IMAG:
B[tx] = hipCimag(A[tx]);
break;
case COMPLEX_SQABS:
B[tx] = hipCsqabs(A[tx]);
break;
case COMPLEX_ABS:
B[tx] = hipCabs(A[tx]);
break;
default:
break;
}
}
/**
* Validates all hipComplex inline functions on device
* Functions validated are: make_hipDoubleComplex, make_hipFloatComplex
*/
template<typename T1, typename T2> bool test_makehipComplex_dev() {
T2 *A, *Ad, *B, *Bd;
T1 *C, *Cd;
bool TestPassed = true;
A = new T2[LEN];
B = new T2[LEN];
C = new T1[LEN];
for (uint32_t i = 0; i < LEN; i++) {
A[i] = 2*i*1.0;
B[i] = (2*i + 1)*1.0;
}
unsigned int size2 = LEN * sizeof(T2);
unsigned int size1 = LEN * sizeof(T1);
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&Ad), size2));
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&Bd), size2));
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&Cd), size1));
HIPCHECK(hipMemcpy(Ad, A, size2, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(Bd, B, size2, hipMemcpyHostToDevice));
hipLaunchKernelGGL(testMakeComplexFunc, dim3(1), dim3(LEN),
0, 0, Ad, Bd, Cd);
HIPCHECK(hipMemcpy(C, Cd, size1, hipMemcpyDeviceToHost));
// Validate the output of the kernel functions.
for (uint32_t i = 0; i < LEN; i++) {
if ((A[i] != C[i].x) || (B[i] != C[i].y)) {
TestPassed = false;
break;
}
}
HIPCHECK(hipFree(Cd));
HIPCHECK(hipFree(Bd));
HIPCHECK(hipFree(Ad));
delete[] C;
delete[] B;
delete[] A;
return TestPassed;
}
/**
* Validates all hipComplex inline functions on device
* Functions validated are: hipCaddf, hipCsubf, hipCmulf and hipCdivf
* hipCadd, hipCsub, hipCmul, hipCdiv
*/
template<typename T1, typename T2>
bool test_complexMathFunc1_dev(enum ComplexFuncType mathFuncType) {
T1 *A, *Ad, *B, *Bd;
T1 *C, *Cd;
bool TestPassed = true;
A = new T1[LEN];
B = new T1[LEN];
C = new T1[LEN];
for (uint32_t i = 0; i < LEN; i++) {
A[i].x = 2*i*1.0;
A[i].y = (2*i + 1)*1.0;
B[i].x = 2*i*1.0 + 0.5;
B[i].y = (2*i + 1)*1.0 + 0.5;
}
unsigned int size = LEN * sizeof(T1);
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&Ad), size));
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&Bd), size));
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&Cd), size));
HIPCHECK(hipMemcpy(Ad, A, size, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(Bd, B, size, hipMemcpyHostToDevice));
hipLaunchKernelGGL(testComplexMathFunc1, dim3(1), dim3(LEN),
0, 0, Ad, Bd, Cd, mathFuncType);
HIPCHECK(hipMemcpy(C, Cd, size, hipMemcpyDeviceToHost));
// Validate the output of the kernel functions.
T2 epsilon = 0.0001f;
T2 real, imag;
for (uint32_t i = 0; i < LEN; i++) {
if (mathFuncType == COMPLEX_ADD) {
real = (A[i].x + B[i].x);
imag = (A[i].y + B[i].y);
} else if (mathFuncType == COMPLEX_SUB) {
real = (A[i].x - B[i].x);
imag = (A[i].y - B[i].y);
} else if (mathFuncType == COMPLEX_MUL) {
real = (A[i].x*B[i].x - A[i].y*B[i].y);
imag = (A[i].y*B[i].x + A[i].x*B[i].y);
} else if (mathFuncType == COMPLEX_DIV) {
T2 sqabs = (B[i].x*B[i].x + B[i].y*B[i].y);
real = (A[i].x * B[i].x + A[i].y * B[i].y)/sqabs;
imag = (A[i].y * B[i].x - A[i].x * B[i].y)/sqabs;
} else if (mathFuncType == COMPLEX_CONJ) {
real = A[i].x;
imag = -A[i].y;
}
if (!COMPARE_REALNUM(real, C[i].x, epsilon) ||
!COMPARE_REALNUM(imag, C[i].y, epsilon)) {
TestPassed = false;
break;
}
}
HIPCHECK(hipFree(Cd));
HIPCHECK(hipFree(Bd));
HIPCHECK(hipFree(Ad));
delete[] C;
delete[] B;
delete[] A;
return TestPassed;
}
/**
* Validates all hipComplex inline functions on device
* Functions validated are: hipCrealf, hipCimagf, hipCsqabsf and hipCabsf
* hipCreal, hipCimag, hipCsqabs, hipCabs
*/
template<typename T1, typename T2>
bool test_complexMathFunc2_dev(enum ComplexFuncType mathFuncType) {
T1 *A, *Ad;
T2 *B, *Bd;
bool TestPassed = true;
A = new T1[LEN];
B = new T2[LEN];
for (uint32_t i = 0; i < LEN; i++) {
A[i].x = 2*i*1.0;
A[i].y = (2*i + 1)*1.0;
}
unsigned int size1 = LEN * sizeof(T1);
unsigned int size2 = LEN * sizeof(T2);
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&Ad), size1));
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&Bd), size2));
HIPCHECK(hipMemcpy(Ad, A, size1, hipMemcpyHostToDevice));
hipLaunchKernelGGL(testComplexMathFunc2, dim3(1), dim3(LEN),
0, 0, Ad, Bd, mathFuncType);
HIPCHECK(hipMemcpy(B, Bd, size2, hipMemcpyDeviceToHost));
// Validate the output of the kernel functions.
T2 epsilon = 0.0001f;
if (mathFuncType == COMPLEX_REAL) {
for (uint32_t i = 0; i < LEN; i++) {
if (!COMPARE_REALNUM(A[i].x, B[i], epsilon)) {
TestPassed = false;
break;
}
}
} else if (mathFuncType == COMPLEX_IMAG) {
for (uint32_t i = 0; i < LEN; i++) {
if (!COMPARE_REALNUM(A[i].y, B[i], epsilon)) {
TestPassed = false;
break;
}
}
} else if (mathFuncType == COMPLEX_SQABS) {
for (uint32_t i = 0; i < LEN; i++) {
T2 sqabs = A[i].x * A[i].x + A[i].y * A[i].y;
#ifdef __HIP_PLATFORM_NVCC__
/* Setting the Floating Point precision to 0.01 as this scenario
is failing on NVIDIA targets. */
epsilon = 0.01f;
#endif
if (!COMPARE_REALNUM(sqabs, B[i], epsilon)) {
TestPassed = false;
break;
}
}
} else if (mathFuncType == COMPLEX_ABS) {
for (uint32_t i = 0; i < LEN; i++) {
T2 sqabs = A[i].x * A[i].x + A[i].y * A[i].y;
if (!COMPARE_REALNUM(sqrtf(sqabs), B[i], epsilon)) {
TestPassed = false;
break;
}
}
}
HIPCHECK(hipFree(Bd));
HIPCHECK(hipFree(Ad));
delete[] B;
delete[] A;
return TestPassed;
}
// Validates all hipComplex inline functions on host
static bool test_allcomplexMathFunc_host() {
bool TestPassed = true;
float fa = 2.0, fb = 3.0;
hipFloatComplex fc = make_hipFloatComplex(fa, fb);
if ((fc.x != fa) || (fc.y != fb)) {
TestPassed &= false;
}
double da = 2.0, db = 3.0;
hipDoubleComplex dc = make_hipDoubleComplex(da, db);
if ((dc.x != da) || (dc.y != db)) {
TestPassed &= false;
}
hipFloatComplex fp, fq, fx;
fp.x = 2.0;
fp.y = 3.0;
fq.x = 4.0;
fq.y = 5.0;
fx = hipCaddf(fp, fq);
if ((fx.x != (fp.x + fq.x)) || (fx.y != (fp.y + fq.y))) {
TestPassed &= false;
}
fx = hipCsubf(fp, fq);
if ((fx.x != (fp.x - fq.x)) || (fx.y != (fp.y - fq.y))) {
TestPassed &= false;
}
fx = hipCmulf(fp, fq);
if ((fx.x != (fp.x*fq.x - fp.y*fq.y)) ||
(fx.y != (fp.y*fq.x + fp.x*fq.y))) {
TestPassed &= false;
}
fx = hipCdivf(fp, fq);
float fsqabs = fq.x*fq.x + fq.y*fq.y;
float epsilon = 0.0001f;
if ((!COMPARE_REALNUM(fx.x, (fp.x*fq.x + fp.y*fq.y)/fsqabs, epsilon)) ||
(!COMPARE_REALNUM(fx.y, (fp.y*fq.x - fp.x*fq.y)/fsqabs, epsilon))) {
TestPassed &= false;
}
if ((fp.x != hipCrealf(fp)) || (fp.y != hipCimagf(fp))) {
TestPassed &= false;
}
fx = hipConjf(fp);
if ((fx.x != fp.x) || (fx.y != -fp.y)) {
TestPassed &= false;
}
if (!COMPARE_REALNUM((fp.x*fp.x + fp.y*fp.y), hipCsqabsf(fp), epsilon)) {
TestPassed &= false;
}
if (!COMPARE_REALNUM(sqrtf(fp.x*fp.x + fp.y*fp.y), hipCabsf(fp), epsilon)) {
TestPassed &= false;
}
hipDoubleComplex dp, dq, dx;
dp.x = 2.0;
dp.y = 3.0;
dq.x = 4.0;
dq.y = 5.0;
dx = hipCadd(dp, dq);
if ((dx.x != (dp.x + dq.x)) || (dx.y != (dp.y + dq.y))) {
TestPassed &= false;
}
dx = hipCsub(dp, dq);
if ((dx.x != (dp.x - dq.x)) || (dx.y != (dp.y - dq.y))) {
TestPassed &= false;
}
dx = hipCmul(dp, dq);
if ((dx.x != (dp.x*dq.x - dp.y*dq.y)) ||
(dx.y != (dp.y*dq.x + dp.x*dq.y))) {
TestPassed &= false;
}
dx = hipCdiv(dp, dq);
float dsqabs = dq.x*dq.x + dq.y*dq.y;
if ((!COMPARE_REALNUM(dx.x, (dp.x*dq.x + dp.y*dq.y)/dsqabs, epsilon)) ||
(!COMPARE_REALNUM(dx.y, (dp.y*dq.x - dp.x*dq.y)/dsqabs, epsilon))) {
TestPassed &= false;
}
if ((dp.x != hipCreal(dp)) || (dp.y != hipCimag(dp))) {
TestPassed &= false;
}
dx = hipConj(dp);
if ((dx.x != dp.x) || (dx.y != -dp.y)) {
TestPassed &= false;
}
if (!COMPARE_REALNUM((dp.x*dp.x + dp.y*dp.y), hipCsqabs(dp), epsilon)) {
TestPassed &= false;
}
if (!COMPARE_REALNUM(sqrtf(dp.x*dp.x + dp.y*dp.y), hipCabs(dp), epsilon)) {
TestPassed &= false;
}
return TestPassed;
}
TEST_CASE("Unit_TestMathFuncComplex") {
bool TestPassed = false;
TestPassed = test_makehipComplex_dev<hipFloatComplex, float>() &&
test_makehipComplex_dev<float2, float>() &&
test_makehipComplex_dev<hipDoubleComplex, double>() &&
test_makehipComplex_dev<double2, double>() &&
test_complexMathFunc1_dev<hipFloatComplex, float>(COMPLEX_ADD) &&
test_complexMathFunc1_dev<hipDoubleComplex, double>(COMPLEX_ADD)
&& test_complexMathFunc1_dev<hipFloatComplex, float>(COMPLEX_SUB)
&& test_complexMathFunc1_dev<hipDoubleComplex, double>
(COMPLEX_SUB) && test_complexMathFunc1_dev<hipFloatComplex,
float>(COMPLEX_MUL) && test_complexMathFunc1_dev<hipDoubleComplex,
double>(COMPLEX_MUL) && test_complexMathFunc1_dev<hipFloatComplex,
float>(COMPLEX_DIV) && test_complexMathFunc1_dev<hipDoubleComplex,
double>(COMPLEX_DIV) && test_complexMathFunc1_dev<hipFloatComplex,
float>(COMPLEX_CONJ) && test_complexMathFunc1_dev<
hipDoubleComplex, double>(COMPLEX_CONJ) && test_complexMathFunc2_dev
<hipFloatComplex, float>(COMPLEX_REAL) && test_complexMathFunc2_dev
<hipDoubleComplex, double>(COMPLEX_REAL) && test_complexMathFunc2_dev
<hipFloatComplex, float>(COMPLEX_IMAG) && test_complexMathFunc2_dev
<hipDoubleComplex, double>(COMPLEX_IMAG) && test_complexMathFunc2_dev
<hipFloatComplex, float>(COMPLEX_SQABS) && test_complexMathFunc2_dev
<hipDoubleComplex, double>(COMPLEX_SQABS) && test_complexMathFunc2_dev
<hipFloatComplex, float>(COMPLEX_ABS) && test_complexMathFunc2_dev
<hipDoubleComplex, double>(COMPLEX_ABS) &&test_allcomplexMathFunc_host();
REQUIRE(TestPassed == true);
}
catch/unit/deviceLib/hipDeviceMemcpy.cc
+60
Ver Arquivo
@@ -0,0 +1,60 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#define LEN 1024
#define SIZE (LEN << 2)
__global__ static void cpy(uint32_t* Out, uint32_t* In) {
int tx = threadIdx.x;
memcpy(Out + tx, In + tx, sizeof(uint32_t));
}
__global__ static void set(uint32_t* ptr, uint8_t val) {
int tx = threadIdx.x;
memset(ptr + tx, val, sizeof(uint32_t));
}
TEST_CASE("Unit_ToAndFroMemCpyToDevice") {
uint32_t *A, *Ad, *B, *Bd;
A = new uint32_t[LEN];
B = new uint32_t[LEN];
for (int i = 0; i < LEN; i++) {
A[i] = i;
B[i] = 0;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(cpy, dim3(1), dim3(LEN), 0, 0, Bd, Ad);
HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
for (int i = LEN - 16; i < LEN; i++) {
REQUIRE(A[i] == B[i]);
}
hipLaunchKernelGGL(set, dim3(1), dim3(LEN), 0, 0, Bd, 0x1);
HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
for (int i = LEN - 16; i < LEN; i++) {
REQUIRE(0x01010101 == B[i]);
}
delete[] A;
delete[] B;
}
catch/unit/deviceLib/hipTestDevice.cc
+736
Ver Arquivo
@@ -0,0 +1,736 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip/math_functions.h>
#define N 512
#define SIZE (N * sizeof(float))
__global__ void test_sincosf(float* a, float* b, float* c) {
int tid = threadIdx.x;
sincosf(a[tid], b + tid, c + tid);
}
__global__ void test_sincospif(float* a, float* b, float* c) {
int tid = threadIdx.x;
sincospif(a[tid], b + tid, c + tid);
}
__global__ void test_fdividef(float* a, float* b, float* c) {
int tid = threadIdx.x;
c[tid] = fdividef(a[tid], b[tid]);
}
__global__ void test_llrintf(float* a, int64_t* b) {
int tid = threadIdx.x;
b[tid] = llrintf(a[tid]);
}
__global__ void test_lrintf(float* a, int64_t* b) {
int tid = threadIdx.x;
b[tid] = lrintf(a[tid]);
}
__global__ void test_rintf(float* a, float* b) {
int tid = threadIdx.x;
b[tid] = rintf(a[tid]);
}
__global__ void test_llroundf(float* a, int64_t* b) {
int tid = threadIdx.x;
b[tid] = llroundf(a[tid]);
}
__global__ void test_lroundf(float* a, int64_t* b) {
int tid = threadIdx.x;
b[tid] = lroundf(a[tid]);
}
__global__ void test_rhypotf(float* a, float* b, float* c) {
int tid = threadIdx.x;
c[tid] = rhypotf(a[tid], b[tid]);
}
__global__ void test_norm3df(float* a, float* b, float* c, float* d) {
int tid = threadIdx.x;
d[tid] = norm3df(a[tid], b[tid], c[tid]);
}
__global__ void test_norm4df(float* a, float* b, float* c, float* d, float* e) {
int tid = threadIdx.x;
e[tid] = norm4df(a[tid], b[tid], c[tid], d[tid]);
}
__global__ void test_normf(float* a, float* b) {
int tid = threadIdx.x;
b[tid] = normf(N, a);
}
__global__ void test_rnorm3df(float* a, float* b, float* c, float* d) {
int tid = threadIdx.x;
d[tid] = rnorm3df(a[tid], b[tid], c[tid]);
}
__global__ void test_rnorm4df(float* a, float* b, float* c, float* d,
float* e) {
int tid = threadIdx.x;
e[tid] = rnorm4df(a[tid], b[tid], c[tid], d[tid]);
}
__global__ void test_rnormf(float* a, float* b) {
int tid = threadIdx.x;
b[tid] = rnormf(N, a);
}
__global__ void test_erfinvf(float* a, float* b) {
int tid = threadIdx.x;
b[tid] = erff(erfinvf(a[tid]));
}
bool run_sincosf() {
float *A, *Ad, *B, *C, *Bd, *Cd;
A = new float[N];
B = new float[N];
C = new float[N];
for (int i = 0; i < N; i++) {
A[i] = 1.0f;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Cd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_sincosf, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (B[i] == sinf(1.0f)) {
passed = 1;
}
}
passed = 0;
for (int i = 0; i < 512; i++) {
if (C[i] == cosf(1.0f)) {
passed = 1;
}
}
delete[] A;
delete[] B;
delete[] C;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
if (passed == 1) {
return true;
} else {
return false;
}
}
bool run_sincospif() {
float *A, *Ad, *B, *C, *Bd, *Cd;
A = new float[N];
B = new float[N];
C = new float[N];
for (int i = 0; i < N; i++) {
A[i] = 1.0f;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Cd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_sincospif, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (B[i] - sinf(3.14 * 1.0f) < 0.1) {
passed = 1;
}
}
passed = 0;
for (int i = 0; i < 512; i++) {
if (C[i] - cosf(3.14 * 1.0f) < 0.1) {
passed = 1;
}
}
delete[] A;
delete[] B;
delete[] C;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
if (passed == 1) {
return true;
} else {
return false;
}
}
bool run_fdividef() {
float *A, *Ad, *B, *C, *Bd, *Cd;
A = new float[N];
B = new float[N];
C = new float[N];
for (int i = 0; i < N; i++) {
A[i] = 1.0f;
B[i] = 2.0f;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Cd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_fdividef, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
HIP_CHECK(hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (C[i] == A[i] / B[i]) {
passed = 1;
}
}
delete[] A;
delete[] B;
delete[] C;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
if (passed == 1) {
return true;
} else {
return false;
}
}
bool run_llrintf() {
float *A, *Ad;
int64_t *B, *Bd;
A = new float[N];
B = new int64_t[N];
for (int i = 0; i < N; i++) {
A[i] = 1.345f;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), N * sizeof(int64_t)));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_llrintf, dim3(1), dim3(N), 0, 0, Ad, Bd);
HIP_CHECK(hipMemcpy(B, Bd, N * sizeof(int64_t), hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
int x = roundf(A[i]);
if (B[i] == x) {
passed = 1;
}
}
delete[] A;
delete[] B;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
if (passed == 1) {
return true;
} else {
return false;
}
}
bool run_lrintf() {
float *A, *Ad;
int64_t *B, *Bd;
A = new float[N];
B = new int64_t[N];
for (int i = 0; i < N; i++) {
A[i] = 1.345f;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), N * sizeof(int64_t)));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_lrintf, dim3(1), dim3(N), 0, 0, Ad, Bd);
HIP_CHECK(hipMemcpy(B, Bd, N * sizeof(int64_t), hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
int x = roundf(A[i]);
if (B[i] == x) {
passed = 1;
}
}
delete[] A;
delete[] B;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
if (passed == 1) {
return true;
} else {
return false;
}
}
bool run_rintf() {
float *A, *Ad;
float *B, *Bd;
A = new float[N];
B = new float[N];
for (int i = 0; i < N; i++) {
A[i] = 1.345f;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_rintf, dim3(1), dim3(N), 0, 0, Ad, Bd);
HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
float x = roundf(A[i]);
if (B[i] == x) {
passed = 1;
}
}
delete[] A;
delete[] B;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
if (passed == 1) {
return true;
} else {
return false;
}
}
bool run_llroundf() {
float *A, *Ad;
int64_t *B, *Bd;
A = new float[N];
B = new int64_t[N];
for (int i = 0; i < N; i++) {
A[i] = 1.345f;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), N * sizeof(int64_t)));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_llroundf, dim3(1), dim3(N), 0, 0, Ad, Bd);
HIP_CHECK(hipMemcpy(B, Bd, N * sizeof(int64_t), hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
int x = roundf(A[i]);
if (B[i] == x) {
passed = 1;
}
}
delete[] A;
delete[] B;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
if (passed == 1) {
return true;
} else {
return false;
}
}
bool run_lroundf() {
float *A, *Ad;
int64_t *B, *Bd;
A = new float[N];
B = new int64_t[N];
for (int i = 0; i < N; i++) {
A[i] = 1.345f;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), N * sizeof(int64_t)));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_lroundf, dim3(1), dim3(N), 0, 0, Ad, Bd);
HIP_CHECK(hipMemcpy(B, Bd, N * sizeof(int64_t), hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
int x = roundf(A[i]);
if (B[i] == x) {
passed = 1;
}
}
delete[] A;
delete[] B;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
if (passed == 1) {
return true;
} else {
return false;
}
}
bool run_norm3df() {
float *A, *Ad, *B, *Bd, *C, *Cd, *D, *Dd;
A = new float[N];
B = new float[N];
C = new float[N];
D = new float[N];
float val = 0.0f;
for (int i = 0; i < N; i++) {
A[i] = 1.0f;
B[i] = 2.0f;
C[i] = 3.0f;
}
val = sqrtf(1.0f + 4.0f + 9.0f);
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Cd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Dd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_norm3df, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd);
HIP_CHECK(hipMemcpy(D, Dd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (D[i] - val < 0.000001) {
passed = 1;
}
}
delete[] A;
delete[] B;
delete[] C;
delete[] D;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
HIP_CHECK(hipFree(Dd));
if (passed == 1) {
return true;
} else {
return false;
}
}
bool run_norm4df() {
float *A, *Ad, *B, *Bd, *C, *Cd, *D, *Dd, *E, *Ed;
A = new float[N];
B = new float[N];
C = new float[N];
D = new float[N];
E = new float[N];
float val = 0.0f;
for (int i = 0; i < N; i++) {
A[i] = 1.0f;
B[i] = 2.0f;
C[i] = 3.0f;
D[i] = 4.0f;
}
val = sqrtf(1.0f + 4.0f + 9.0f + 16.0f);
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Cd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Dd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ed), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Dd, D, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_norm4df, dim3(1), dim3(N), 0, 0, Ad,
Bd, Cd, Dd, Ed);
HIP_CHECK(hipMemcpy(E, Ed, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (E[i] - val < 0.000001) {
passed = 1;
}
}
delete[] A;
delete[] B;
delete[] C;
delete[] D;
delete[] E;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
HIP_CHECK(hipFree(Dd));
HIP_CHECK(hipFree(Ed));
if (passed == 1) {
return true;
} else {
return false;
}
}
bool run_normf() {
float *A, *Ad, *B, *Bd;
A = new float[N];
B = new float[N];
float val = 0.0f;
for (int i = 0; i < N; i++) {
A[i] = 1.0f;
B[i] = 0.0f;
val += 1.0f;
}
val = sqrtf(val);
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_normf, dim3(1), dim3(N), 0, 0, Ad, Bd);
HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (B[0] - val < 0.000001) {
passed = 1;
}
}
delete[] A;
delete[] B;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
if (passed == 1) {
return true;
} else {
return false;
}
}
bool run_rhypotf() {
float *A, *Ad, *B, *Bd, *C, *Cd;
A = new float[N];
B = new float[N];
C = new float[N];
float val = 0.0f;
for (int i = 0; i < N; i++) {
A[i] = 1.0f;
B[i] = 2.0f;
}
val = 1 / sqrtf(1.0f + 4.0f);
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Cd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_rhypotf, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
HIP_CHECK(hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (C[i] - val < 0.000001) {
passed = 1;
}
}
delete[] A;
delete[] B;
delete[] C;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
if (passed == 1) {
return true;
} else {
return false;
}
}
bool run_rnorm3df() {
float *A, *Ad, *B, *Bd, *C, *Cd, *D, *Dd;
A = new float[N];
B = new float[N];
C = new float[N];
D = new float[N];
float val = 0.0f;
for (int i = 0; i < N; i++) {
A[i] = 1.0f;
B[i] = 2.0f;
C[i] = 3.0f;
}
val = 1 / sqrtf(1.0f + 4.0f + 9.0f);
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Cd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Dd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_rnorm3df, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd);
HIP_CHECK(hipMemcpy(D, Dd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (D[i] - val < 0.000001) {
passed = 1;
}
}
delete[] A;
delete[] B;
delete[] C;
delete[] D;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
HIP_CHECK(hipFree(Dd));
if (passed == 1) {
return true;
} else {
return false;
}
}
bool run_rnorm4df() {
float *A, *Ad, *B, *Bd, *C, *Cd, *D, *Dd, *E, *Ed;
A = new float[N];
B = new float[N];
C = new float[N];
D = new float[N];
E = new float[N];
float val = 0.0f;
for (int i = 0; i < N; i++) {
A[i] = 1.0f;
B[i] = 2.0f;
C[i] = 3.0f;
D[i] = 4.0f;
}
val = 1 / sqrtf(1.0f + 4.0f + 9.0f + 16.0f);
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Cd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Dd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ed), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Dd, D, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_rnorm4df, dim3(1), dim3(N), 0, 0, Ad,
Bd, Cd, Dd, Ed);
HIP_CHECK(hipMemcpy(E, Ed, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (E[i] - val < 0.000001) {
passed = 1;
}
}
delete[] A;
delete[] B;
delete[] C;
delete[] D;
delete[] E;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
HIP_CHECK(hipFree(Dd));
HIP_CHECK(hipFree(Ed));
if (passed == 1) {
return true;
} else {
return false;
}
}
bool run_rnormf() {
float *A, *Ad, *B, *Bd;
A = new float[N];
B = new float[N];
float val = 0.0f;
for (int i = 0; i < N; i++) {
A[i] = 1.0f;
B[i] = 0.0f;
val += 1.0f;
}
val = 1 / sqrtf(val);
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_rnormf, dim3(1), dim3(N), 0, 0, Ad, Bd);
HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (B[0] - val < 0.000001) {
passed = 1;
}
}
delete[] A;
delete[] B;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
if (passed == 1) {
return true;
} else {
return false;
}
}
bool run_erfinvf() {
float *A, *Ad, *B, *Bd;
A = new float[N];
B = new float[N];
for (int i = 0; i < N; i++) {
A[i] = -0.6f;
B[i] = 0.0f;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_erfinvf, dim3(1), dim3(N), 0, 0, Ad, Bd);
HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (B[i] - A[i] < 0.000001) {
passed = 1;
}
}
delete[] A;
delete[] B;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
if (passed == 1) {
return true;
} else {
return false;
}
}
TEST_CASE("Unit_hipDeviceTrigFunc_Float") {
bool result = false;
result = run_sincosf() && run_sincospif() && run_fdividef() &&
run_llrintf() && run_norm3df() && run_norm4df() &&
run_normf() && run_rnorm3df() && run_rnorm4df() &&
run_rnormf() && run_lroundf() && run_llroundf() &&
run_rintf() && run_rhypotf() && run_erfinvf();
REQUIRE(result == true);
}
catch/unit/deviceLib/hipTestDeviceDouble.cc
+630
Ver Arquivo
@@ -0,0 +1,630 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip/math_functions.h>
#define N 512
#define SIZE (N * sizeof(double))
__global__ void test_sincos(double* a, double* b, double* c) {
int tid = threadIdx.x;
sincos(a[tid], b + tid, c + tid);
}
__global__ void test_sincospi(double* a, double* b, double* c) {
int tid = threadIdx.x;
sincospi(a[tid], b + tid, c + tid);
}
__global__ void test_llrint(double* a, int64_t* b) {
int tid = threadIdx.x;
b[tid] = llrint(a[tid]);
}
__global__ void test_lrint(double* a, int64_t* b) {
int tid = threadIdx.x;
b[tid] = lrint(a[tid]);
}
__global__ void test_rint(double* a, double* b) {
int tid = threadIdx.x;
b[tid] = rint(a[tid]);
}
__global__ void test_llround(double* a, int64_t* b) {
int tid = threadIdx.x;
b[tid] = llround(a[tid]);
}
__global__ void test_lround(double* a, int64_t* b) {
int tid = threadIdx.x;
b[tid] = lround(a[tid]);
}
__global__ void test_rhypot(double* a, double* b, double* c) {
int tid = threadIdx.x;
c[tid] = rhypot(a[tid], b[tid]);
}
__global__ void test_norm3d(double* a, double* b, double* c, double* d) {
int tid = threadIdx.x;
d[tid] = norm3d(a[tid], b[tid], c[tid]);
}
__global__ void test_norm4d(double* a, double* b, double* c, double* d,
double* e) {
int tid = threadIdx.x;
e[tid] = norm4d(a[tid], b[tid], c[tid], d[tid]);
}
__global__ void test_rnorm3d(double* a, double* b, double* c, double* d) {
int tid = threadIdx.x;
d[tid] = rnorm3d(a[tid], b[tid], c[tid]);
}
__global__ void test_rnorm4d(double* a, double* b, double* c, double* d,
double* e) {
int tid = threadIdx.x;
e[tid] = rnorm4d(a[tid], b[tid], c[tid], d[tid]);
}
__global__ void test_rnorm(double* a, double* b) {
int tid = threadIdx.x;
b[tid] = rnorm(N, a);
}
__global__ void test_erfinv(double* a, double* b) {
int tid = threadIdx.x;
b[tid] = erf(erfinv(a[tid]));
}
bool run_sincos() {
double *A, *Ad, *B, *C, *Bd, *Cd;
A = new double[N];
B = new double[N];
C = new double[N];
for (int i = 0; i < N; i++) {
A[i] = 1.0;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Cd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_sincos, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if ((B[i] == sin(1.0)) && (C[i] == cos(1.0))) {
passed = 1;
}
}
delete[] A;
delete[] B;
delete[] C;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
if (passed == 1) {
return true;
}
return false;
}
bool run_sincospi() {
double *A, *Ad, *B, *C, *Bd, *Cd;
A = new double[N];
B = new double[N];
C = new double[N];
for (int i = 0; i < N; i++) {
A[i] = 1.0;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Cd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_sincospi, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if ((B[i] - sin(3.14 * 1.0) < 0.1) && (C[i] - cos(3.14 * 1.0) < 0.1)) {
passed = 1;
}
}
delete[] A;
delete[] B;
delete[] C;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
if (passed == 1) {
return true;
}
return false;
}
bool run_llrint() {
double *A, *Ad;
int64_t *B, *Bd;
A = new double[N];
B = new int64_t[N];
for (int i = 0; i < N; i++) {
A[i] = 1.345;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd),
N * sizeof(int64_t)));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_llrint, dim3(1), dim3(N), 0, 0, Ad, Bd);
HIP_CHECK(hipMemcpy(B, Bd, N * sizeof(int64_t),
hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
int x = round(A[i]);
int64_t y = x;
if (B[i] == x) {
passed = 1;
}
}
delete[] A;
delete[] B;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
if (passed == 1) {
return true;
}
return false;
}
bool run_lrint() {
double *A, *Ad;
int64_t *B, *Bd;
A = new double[N];
B = new int64_t[N];
for (int i = 0; i < N; i++) {
A[i] = 1.345;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), N * sizeof(int64_t)));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_lrint, dim3(1), dim3(N), 0, 0, Ad, Bd);
HIP_CHECK(hipMemcpy(B, Bd, N * sizeof(int64_t), hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
int64_t x = round(A[i]);
if (B[i] == x) {
passed = 1;
}
}
delete[] A;
delete[] B;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
if (passed == 1) {
return true;
}
return false;
}
bool run_rint() {
double *A, *Ad;
double *B, *Bd;
A = new double[N];
B = new double[N];
for (int i = 0; i < N; i++) {
A[i] = 1.345;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_rint, dim3(1), dim3(N), 0, 0, Ad, Bd);
HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
double x = round(A[i]);
if (B[i] == x) {
passed = 1;
}
}
delete[] A;
delete[] B;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
if (passed == 1) {
return true;
}
return false;
}
bool run_llround() {
double *A, *Ad;
int64_t *B, *Bd;
A = new double[N];
B = new int64_t[N];
for (int i = 0; i < N; i++) {
A[i] = 1.345;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd),
N * sizeof(int64_t)));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_llround, dim3(1), dim3(N), 0, 0, Ad, Bd);
HIP_CHECK(hipMemcpy(B, Bd, N * sizeof(int64_t),
hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
int64_t x = round(A[i]);
if (B[i] == x) {
passed = 1;
}
}
delete[] A;
delete[] B;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
if (passed == 1) {
return true;
}
return false;
}
bool run_lround() {
double *A, *Ad;
int64_t *B, *Bd;
A = new double[N];
B = new int64_t[N];
for (int i = 0; i < N; i++) {
A[i] = 1.345;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), N * sizeof(int64_t)));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_lround, dim3(1), dim3(N), 0, 0, Ad, Bd);
HIP_CHECK(hipMemcpy(B, Bd, N * sizeof(int64_t), hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
int64_t x = round(A[i]);
if (B[i] == x) {
passed = 1;
}
}
delete[] A;
delete[] B;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
if (passed == 1) {
return true;
}
return false;
}
bool run_norm3d() {
double *A, *Ad, *B, *Bd, *C, *Cd, *D, *Dd;
A = new double[N];
B = new double[N];
C = new double[N];
D = new double[N];
double val = 0.0;
for (int i = 0; i < N; i++) {
A[i] = 1.0;
B[i] = 2.0;
C[i] = 3.0;
}
val = sqrt(1.0 + 4.0 + 9.0);
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Cd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Dd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_norm3d, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd);
HIP_CHECK(hipMemcpy(D, Dd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (D[i] - val < 0.000001) {
passed = 1;
}
}
delete[] A;
delete[] B;
delete[] C;
delete[] D;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
HIP_CHECK(hipFree(Dd));
if (passed == 1) {
return true;
}
return false;
}
bool run_norm4d() {
double *A, *Ad, *B, *Bd, *C, *Cd, *D, *Dd, *E, *Ed;
A = new double[N];
B = new double[N];
C = new double[N];
D = new double[N];
E = new double[N];
double val = 0.0;
for (int i = 0; i < N; i++) {
A[i] = 1.0;
B[i] = 2.0;
C[i] = 3.0;
D[i] = 4.0;
}
val = sqrt(1.0 + 4.0 + 9.0 + 16.0);
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Cd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Dd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ed), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Dd, D, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_norm4d, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd, Ed);
HIP_CHECK(hipMemcpy(E, Ed, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (E[i] - val < 0.000001) {
passed = 1;
}
}
delete[] A;
delete[] B;
delete[] C;
delete[] D;
delete[] E;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
HIP_CHECK(hipFree(Dd));
HIP_CHECK(hipFree(Ed));
if (passed == 1) {
return true;
}
return false;
}
bool run_rhypot() {
double *A, *Ad, *B, *Bd, *C, *Cd;
A = new double[N];
B = new double[N];
C = new double[N];
double val = 0.0;
for (int i = 0; i < N; i++) {
A[i] = 1.0;
B[i] = 2.0;
}
val = 1 / sqrt(1.0 + 4.0);
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Cd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_rhypot, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd);
HIP_CHECK(hipMemcpy(C, Cd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (C[i] - val < 0.000001) {
passed = 1;
}
}
delete[] A;
delete[] B;
delete[] C;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
if (passed == 1) {
return true;
}
return false;
}
bool run_rnorm3d() {
double *A, *Ad, *B, *Bd, *C, *Cd, *D, *Dd;
A = new double[N];
B = new double[N];
C = new double[N];
D = new double[N];
double val = 0.0;
for (int i = 0; i < N; i++) {
A[i] = 1.0;
B[i] = 2.0;
C[i] = 3.0;
}
val = 1 / sqrt(1.0 + 4.0 + 9.0);
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Cd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Dd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_rnorm3d, dim3(1), dim3(N), 0, 0, Ad, Bd, Cd, Dd);
HIP_CHECK(hipMemcpy(D, Dd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (D[i] - val < 0.000001) {
passed = 1;
}
}
delete[] A;
delete[] B;
delete[] C;
delete[] D;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
HIP_CHECK(hipFree(Dd));
if (passed == 1) {
return true;
}
return false;
}
bool run_rnorm4d() {
double *A, *Ad, *B, *Bd, *C, *Cd, *D, *Dd, *E, *Ed;
A = new double[N];
B = new double[N];
C = new double[N];
D = new double[N];
E = new double[N];
double val = 0.0;
for (int i = 0; i < N; i++) {
A[i] = 1.0;
B[i] = 2.0;
C[i] = 3.0;
D[i] = 4.0;
}
val = 1 / sqrt(1.0 + 4.0 + 9.0 + 16.0);
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Cd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Dd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ed), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Cd, C, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Dd, D, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_rnorm4d, dim3(1), dim3(N), 0, 0, Ad,
Bd, Cd, Dd, Ed);
HIP_CHECK(hipMemcpy(E, Ed, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (E[i] - val < 0.000001) {
passed = 1;
}
}
delete[] A;
delete[] B;
delete[] C;
delete[] D;
delete[] E;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
HIP_CHECK(hipFree(Cd));
HIP_CHECK(hipFree(Dd));
HIP_CHECK(hipFree(Ed));
if (passed == 1) {
return true;
}
return false;
}
bool run_rnorm() {
double *A, *Ad, *B, *Bd;
A = new double[N];
B = new double[N];
double val = 0.0;
for (int i = 0; i < N; i++) {
A[i] = 1.0;
B[i] = 0.0;
val += 1.0;
}
val = 1 / sqrt(val);
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_rnorm, dim3(1), dim3(N), 0, 0, Ad, Bd);
HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (B[0] - val < 0.000001) {
passed = 1;
}
}
delete[] A;
delete[] B;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
if (passed == 1) {
return true;
}
return false;
}
bool run_erfinv() {
double *A, *Ad, *B, *Bd;
A = new double[N];
B = new double[N];
for (int i = 0; i < N; i++) {
A[i] = -0.6;
B[i] = 0.0;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Bd), SIZE));
HIP_CHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(test_erfinv, dim3(1), dim3(N), 0, 0, Ad, Bd);
HIP_CHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
int passed = 0;
for (int i = 0; i < 512; i++) {
if (B[i] - A[i] < 0.000001) {
passed = 1;
}
}
delete[] A;
delete[] B;
HIP_CHECK(hipFree(Ad));
HIP_CHECK(hipFree(Bd));
if (passed == 1) {
return true;
}
return false;
}
TEST_CASE("Unit_hipTrigDeviceFunc_Double") {
bool result = false;
result = run_sincos() && run_sincospi() && run_llrint() &&
run_norm3d() && run_norm4d() && run_rnorm3d() &&
run_rnorm4d() && run_rnorm() && run_lround() && run_llround()
&& run_rint() && run_rhypot() && run_erfinv();
REQUIRE(result == true);
}
catch/unit/deviceLib/hipTestDeviceLimit.cc
+24
Ver Arquivo
@@ -0,0 +1,24 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
TEST_CASE("Unit_hipTestDeviceLimit_Basic") {
size_t heap;
HIP_CHECK(hipDeviceGetLimit(&heap, hipLimitMallocHeapSize));
REQUIRE(heap != NULL);
}
catch/unit/deviceLib/hipTestDeviceSymbol.cc
+63 -61
Ver Arquivo
@@ -1,5 +1,5 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
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
@@ -33,25 +33,32 @@ constexpr size_t SIZE = 1024 * 4;
__device__ int globalIn[NUM];
__device__ int globalOut[NUM];
__global__ void Assign(int* Out) {
__global__ static void Assign(int* Out) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
Out[tid] = globalIn[tid];
globalOut[tid] = globalIn[tid];
}
__device__ __constant__ int globalConst[NUM];
__device__ static __constant__ float statConstVar[NUM];
__global__ void checkAddress(int* addr, bool* out) { *out = (globalConst == addr); }
__global__ void checkAddress(int* addr, bool* out) {
*out = (globalConst == addr);
}
__global__ void checkStaticConstVarAddress(float* addr, bool* out) {
*out = (statConstVar == addr);
}
TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
int *A{nullptr}, *Am{nullptr}, *B{nullptr}, *Ad{nullptr}, *C{nullptr}, *Cm{nullptr};
int *A{nullptr}, *Am{nullptr}, *B{nullptr}, *Ad{nullptr},
*C{nullptr}, *Cm{nullptr};
A = new int[NUM];
B = new int[NUM];
C = new int[NUM];
HIP_CHECK(hipMalloc((void**)&Ad, SIZE));
HIP_CHECK(hipHostMalloc((void**)&Am, SIZE));
HIP_CHECK(hipHostMalloc((void**)&Cm, SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&Am), SIZE));
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&Cm), SIZE));
for (size_t i = 0; i < NUM; i++) {
A[i] = -1 * static_cast<int>(i);
@@ -66,13 +73,14 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
hipStream_t stream{};
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(
hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), Am, SIZE, 0, hipMemcpyHostToDevice, stream));
hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), Am, SIZE, 0,
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipStreamSynchronize(stream));
hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpyFromSymbolAsync(Cm, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost,
stream));
HIP_CHECK(hipMemcpyFromSymbolAsync(Cm, HIP_SYMBOL(globalOut), SIZE, 0,
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamDestroy(stream));
for (size_t i = 0; i < NUM; i++) {
@@ -82,11 +90,13 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
}
SECTION("Calling hipMemcpyTo/FromSymbol - validate value in host memory") {
HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(globalIn), A, SIZE, 0,
hipMemcpyHostToDevice));
hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpyFromSymbol(C, HIP_SYMBOL(globalOut), SIZE, 0,
hipMemcpyDeviceToHost));
for (size_t i = 0; i < NUM; i++) {
REQUIRE(A[i] == B[i]);
@@ -98,13 +108,15 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
hipStream_t stream{};
HIP_CHECK(hipStreamCreate(&stream));
HIP_CHECK(
hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice, stream));
hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0,
hipMemcpyHostToDevice, stream));
HIP_CHECK(hipStreamSynchronize(stream));
hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(
hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost, stream));
hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0,
hipMemcpyDeviceToHost, stream));
HIP_CHECK(hipStreamSynchronize(stream));
HIP_CHECK(hipStreamDestroy(stream));
@@ -115,14 +127,14 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
}
SECTION("Calling hipMemcpyTo/FromSymbol using hipStreamPerThread") {
HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0, hipMemcpyHostToDevice,
hipStreamPerThread));
HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(globalIn), A, SIZE, 0,
hipMemcpyHostToDevice, hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
hipLaunchKernelGGL(Assign, dim3(1, 1, 1), dim3(NUM, 1, 1), 0, 0, Ad);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(B, Ad, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0, hipMemcpyDeviceToHost,
hipStreamPerThread));
HIP_CHECK(hipMemcpyFromSymbolAsync(C, HIP_SYMBOL(globalOut), SIZE, 0,
hipMemcpyDeviceToHost, hipStreamPerThread));
HIP_CHECK(hipStreamSynchronize(hipStreamPerThread));
for (size_t i = 0; i < NUM; i++) {
@@ -140,14 +152,18 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
size_t symbolSize = 0;
int* symbolAddress{nullptr};
HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConst)));
HIP_CHECK(hipGetSymbolAddress((void**)&symbolAddress, HIP_SYMBOL(globalConst)));
HIP_CHECK(hipMalloc((void**)&checkOkD, sizeof(bool)));
hipLaunchKernelGGL(checkAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, symbolAddress, checkOkD);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost));
HIP_CHECK(hipGetSymbolAddress(reinterpret_cast<void**>(&symbolAddress),
HIP_SYMBOL(globalConst)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&checkOkD),
sizeof(bool)));
hipLaunchKernelGGL(checkAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
symbolAddress, checkOkD);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool),
hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(checkOkD));
HIP_ASSERT(checkOk);
HIP_ASSERT((symbolSize == SIZE));
REQUIRE(checkOk);
REQUIRE((symbolSize == SIZE));
}
HIP_CHECK(hipHostFree(Am));
@@ -157,11 +173,9 @@ TEST_CASE("Unit_hipMemcpyToSymbolAsync_ToNFrom") {
delete[] B;
delete[] C;
}
/**
1) Validate get symbol address/size for global const array.
2) Validate get symbol address/size for static const variable.
*/
/*
1) Validate get symbol address/size for static const variable.
*/
TEST_CASE("Unit_hipGetSymbolAddressAndSize_Validation") {
bool* checkOkD{nullptr};
bool checkOk = false;
@@ -169,32 +183,20 @@ TEST_CASE("Unit_hipGetSymbolAddressAndSize_Validation") {
int* symbolArrAddress{};
float* symbolVarAddress{};
SECTION("Validate symbol size/address of global const array") {
HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConstArr)));
HIP_CHECK(hipGetSymbolAddress(reinterpret_cast<void**>(&symbolArrAddress),
HIP_SYMBOL(globalConstArr)));
HIP_CHECK(hipMalloc(&checkOkD, sizeof(bool)));
hipLaunchKernelGGL(checkGlobalConstAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
symbolArrAddress, checkOkD);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(checkOkD));
HIP_ASSERT(checkOk);
HIP_ASSERT(symbolSize == SIZE);
}
SECTION("Validate symbol size/address of static const variable") {
HIP_CHECK(hipGetSymbolSize(&symbolSize, HIP_SYMBOL(statConstVar)));
HIP_CHECK(
hipGetSymbolAddress(reinterpret_cast<void**>(&symbolVarAddress), HIP_SYMBOL(statConstVar)));
hipGetSymbolAddress(reinterpret_cast<void**>(&symbolVarAddress),
HIP_SYMBOL(statConstVar)));
HIP_CHECK(hipMalloc(&checkOkD, sizeof(bool)));
hipLaunchKernelGGL(checkStaticConstVarAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0,
symbolVarAddress, checkOkD);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost));
hipLaunchKernelGGL(checkStaticConstVarAddress, dim3(1, 1, 1),
dim3(1, 1, 1), 0, 0, symbolVarAddress, checkOkD);
HIP_CHECK(hipGetLastError());
HIP_CHECK(hipMemcpy(&checkOk, checkOkD, sizeof(bool),
hipMemcpyDeviceToHost));
HIP_CHECK(hipFree(checkOkD));
HIP_ASSERT(checkOk);
HIP_ASSERT(symbolSize == sizeof(float));
REQUIRE(checkOk);
REQUIRE(symbolSize == SIZE);
}
}
@@ -202,15 +204,14 @@ TEST_CASE("Unit_hipGetSymbolAddress_Negative") {
SECTION("Invalid symbol") {
int notADeviceSymbol{0};
int* addr{nullptr};
HIP_CHECK_ERROR(
hipGetSymbolAddress(reinterpret_cast<void**>(&addr), HIP_SYMBOL(notADeviceSymbol)),
hipErrorInvalidSymbol);
HIP_CHECK_ERROR(hipGetSymbolAddress(reinterpret_cast<void**>(&addr),
HIP_SYMBOL(notADeviceSymbol)), hipErrorInvalidSymbol);
}
SECTION("Nullptr symbol") {
int* addr{nullptr};
HIP_CHECK_ERROR(hipGetSymbolAddress(reinterpret_cast<void**>(&addr), nullptr),
hipErrorInvalidSymbol);
HIP_CHECK_ERROR(hipGetSymbolAddress(reinterpret_cast<void**>(&addr),
nullptr), hipErrorInvalidSymbol);
}
}
@@ -218,7 +219,8 @@ TEST_CASE("Unit_hipGetSymbolSize_Negative") {
SECTION("Invalid symbol") {
int notADeviceSymbol{0};
size_t dsize{0};
HIP_CHECK_ERROR(hipGetSymbolSize(&dsize, HIP_SYMBOL(notADeviceSymbol)), hipErrorInvalidSymbol);
HIP_CHECK_ERROR(hipGetSymbolSize(&dsize, HIP_SYMBOL(notADeviceSymbol)),
hipErrorInvalidSymbol);
}
SECTION("Nullptr symbol") {
catch/unit/deviceLib/hipTestDotFunctions.cc
+54
Ver Arquivo
@@ -0,0 +1,54 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip/hip_runtime.h>
#include <hip/math_functions.h>
#include <hip_test_common.hh>
__global__ static void DotFunctions(bool* result) {
// Dot Functions
#if HT_AMD
short2 sa{1}, sb{1};
result[0] = amd_mixed_dot(sa, sb, 1, result[0]) && result[0];
ushort2 usa{1}, usb{1};
result[0] = amd_mixed_dot(usa, usb, (uint) 1, result[0]) && result[0];
char4 ca{1}, cb{1};
result[0] = amd_mixed_dot(ca, cb, 1, result[0]) && result[0];
uchar4 uca{1}, ucb{1};
result[0] = amd_mixed_dot(uca, ucb, (uint) 1, result[0]) && result[0];
int ia{1}, ib{1};
result[0] = amd_mixed_dot(ia, ib, 1, result[0]) && result[0];
uint ua{1}, ub{1};
result[0] = amd_mixed_dot(ua, ub, (uint) 1, result[0]) && result[0];
#endif
}
TEST_CASE("Unit_hipTestDotFunctions") {
bool* result{nullptr};
hipHostMalloc(&result, 1);
result[0] = true;
hipLaunchKernelGGL(DotFunctions, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result);
hipDeviceSynchronize();
REQUIRE(result[0] == true);
hipHostFree(result);
}
catch/unit/deviceLib/hipTestFMA.cc
+183
Ver Arquivo
@@ -0,0 +1,183 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <iostream>
#define LEN 50
#define SIZE (LEN * sizeof(bool))
__global__ void kernelTestFMA(bool *Ad) {
float f = 1.0f / 3.0f;
double d = f;
int i = 0;
auto Check = [&](bool Cond) { Ad[i++] = Cond; };
// f * f + 3.0f will be different if promoted to double.
float floatResult = fma(f, f, 3.0f);
double doubleResult = fma(d, d, 3.0);
Check(floatResult != doubleResult);
if (sizeof(decltype(fma(f, f, 3))) == 8) {
// To align with libcxx, if any argument has integral type,
// it is cast to double.
// Check type promotes to double.
Check(fma(f, f, 3) == doubleResult);
Check(fma(f, f, static_cast<char>(3)) == doubleResult);
Check(fma(f, f, (unsigned char)3) == doubleResult);
Check(fma(f, f, (int32_t)3) == doubleResult);
Check(fma(f, f, (uint32_t)3) == doubleResult);
Check(fma(f, f, static_cast<int>(3)) == doubleResult);
Check(fma(f, f, (unsigned int)3) == doubleResult);
Check(fma(f, f, (int64_t)3) == doubleResult);
Check(fma(f, f, (uint64_t)3) == doubleResult);
Check(fma(f, f, true) == fma(static_cast<double>(f),
static_cast<double>(f), 1.0));
} else if (sizeof(decltype(fma(f, f, 3))) == 4) {
// Previous HIP headers returns float type.
// Delete this to support backwards compatibility.
// check promote to float.
Check(fma(f, f, 3) == floatResult);
Check(fma(f, f, static_cast<char>(3)) == floatResult);
Check(fma(f, f, (unsigned char)3) == floatResult);
Check(fma(f, f, (int32_t)3) == floatResult);
Check(fma(f, f, (uint32_t)3) == floatResult);
Check(fma(f, f, static_cast<int>(3)) == floatResult);
Check(fma(f, f, (unsigned int)3) == floatResult);
Check(fma(f, f, (int64_t)3) == floatResult);
Check(fma(f, f, (uint64_t)3) == floatResult);
Check(fma(f, f, true) == fma(f, f, 1.0f));
} else {
Check(false);
}
Check(fma(d, static_cast<double>(f), 3) == doubleResult);
Check(fma(d, static_cast<double>(f), static_cast<char>(3)) == doubleResult);
Check(fma(d, static_cast<double>(f), (unsigned char)3) == doubleResult);
Check(fma(d, static_cast<double>(f), (int32_t)3) == doubleResult);
Check(fma(d, static_cast<double>(f), (uint32_t)3) == doubleResult);
Check(fma(d, static_cast<double>(f), static_cast<int>(3)) == doubleResult);
Check(fma(d, static_cast<double>(f), (unsigned int)3) == doubleResult);
Check(fma(d, static_cast<double>(f), (int64_t)3) == doubleResult);
Check(fma(d, static_cast<double>(f), (int64_t)3) == doubleResult);
Check(fma(d, static_cast<double>(f), true) ==
fma(static_cast<double>(f), static_cast<double>(f), 1.0));
while (i < LEN)
Check(true);
}
void runTestFMA() {
bool *Ad;
bool A[LEN];
for (unsigned i = 0; i < LEN; i++) {
A[i] = 0;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void **>(&Ad), SIZE));
hipLaunchKernelGGL(kernelTestFMA, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, Ad);
HIP_CHECK(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost));
for (unsigned i = 0; i < LEN; i++) {
REQUIRE(A[i] == true);
}
}
__global__ void kernelTestHalfFMA(bool *Ad) {
_Float16 h = (_Float16)(1.0f/3.0f);
float f = h;
double d = f;
int i = 0;
auto Check = [&](bool Cond) { Ad[i++] = Cond; };
// h * h + 3 will be different if promoted to float.
_Float16 halfResult = fma(h, h, (_Float16)3);
float floatResult = fma(f, f, 3.0f);
double doubleResult = fma(d, d, 3.0);
Check(halfResult != floatResult);
Check(halfResult != doubleResult);
// check promote to half.
// fma(_Float16, _Float16, int) should resolve to
// fma(double, double, double). This is similar to
// fma(float, float, int) resolving to fma(double, double, double)
// as required Standard C++ header <cmath>.
if (sizeof(decltype(fma(h, h, 3))) == 8) {
Check(fma(h, h, 3) == doubleResult);
Check(fma(h, h, static_cast<char>(3)) == doubleResult);
Check(fma(h, h, (unsigned char)3) == doubleResult);
Check(fma(h, h, (int32_t)3) == doubleResult);
Check(fma(h, h, (uint32_t)3) == doubleResult);
Check(fma(h, h, static_cast<int>(3)) == doubleResult);
Check(fma(h, h, (unsigned int)3) == doubleResult);
Check(fma(h, h, (int64_t)3) == doubleResult);
Check(fma(h, h, (uint64_t)3) == doubleResult);
Check(fma(h, h, true) == fma(static_cast<double>(h),
static_cast<double>(h), 1.0));
} else if (sizeof(decltype(fma(h, h, 3))) == 2) {
// ToDo: Currently there is a bug in clang header
// __clang_hip_cmath.h due to using
// std::numeric_limits<T>::is_specified to define
// overloaded math functions. Since numeric_limits is
// not specicialized for _Float16, overloaded template
// functions with argument promotion are not defined
// for _Float16. As a result, fma(_Float16, _Float16, int)
// is resolved to fma(_Float16, _Float16, _Float16).
// This part should be removed after __clang_hip_cmath.h
// is fixed.
Check(fma(h, h, 3) == halfResult);
Check(fma(h, h, static_cast<char>(3)) == halfResult);
Check(fma(h, h, (unsigned char)3) == halfResult);
Check(fma(h, h, (int32_t)3) == halfResult);
Check(fma(h, h, (uint32_t)3) == halfResult);
Check(fma(h, h, static_cast<int>(3)) == halfResult);
Check(fma(h, h, (unsigned int)3) == halfResult);
Check(fma(h, h, (int64_t)3) == halfResult);
Check(fma(h, h, (int64_t)3) == halfResult);
Check(fma(h, h, true) == fma(h, h, (_Float16)1));
} else {
Check(false);
}
while (i < LEN)
Check(true);
}
void runTestHalfFMA() {
bool *Ad;
bool A[LEN];
for (unsigned i = 0; i < LEN; i++) {
A[i] = 0;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void **>(&Ad), SIZE));
hipLaunchKernelGGL(kernelTestHalfFMA, dim3(1, 1, 1), dim3(1, 1, 1),
0, 0, Ad);
HIP_CHECK(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost));
for (unsigned i = 0; i < LEN; i++) {
REQUIRE(A[i] == true);
}
}
TEST_CASE("Unit_hipTestFMA") {
SECTION("test FMA") {
runTestFMA();
}
SECTION("test HalfFMA") {
runTestHalfFMA();
}
}
catch/unit/deviceLib/hipTestHalf.cc
+245
Ver Arquivo
@@ -0,0 +1,245 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip/hip_fp16.h>
#include <hip_test_common.hh>
__device__ void test_convert() {
__half x;
float y = static_cast<float>(x);
}
__global__
void __halfMath(bool* result, __half a) {
result[0] = __heq(__hadd(a, __half{1}), __half{2});
result[0] = __heq(__hadd_sat(a, __half{1}), __half{1}) && result[0];
result[0] = __heq(__hfma(a, __half{2}, __half{3}), __half{5}) && result[0];
result[0] =
__heq(__hfma_sat(a, __half{2}, __half{3}), __half{1}) && result[0];
result[0] = __heq(__hsub(a, __half{1}), __half{0}) && result[0];
result[0] = __heq(__hsub_sat(a, __half{2}), __half{0}) && result[0];
result[0] = __heq(__hmul(a, __half{2}), __half{2}) && result[0];
result[0] = __heq(__hmul_sat(a, __half{2}), __half{1}) && result[0];
result[0] = __heq(__hdiv(a, __half{2}), __half{0.5}) && result[0];
}
__device__
bool to_bool(const __half2& x) {
auto r = static_cast<const __half2_raw&>(x);
return r.data.x != 0 && r.data.y != 0;
}
__global__
void __half2Math(bool* result, __half2 a) {
result[0] =
to_bool(__heq2(__hadd2(a, __half2{1, 1}), __half2{2, 2}));
result[0] = to_bool(__heq2(__hadd2_sat(a, __half2{1, 1}), __half2{1, 1})) &&
result[0];
result[0] = to_bool(__heq2(
__hfma2(a, __half2{2, 2}, __half2{3, 3}), __half2{5, 5})) && result[0];
result[0] = to_bool(__heq2(
__hfma2_sat(a, __half2{2, 2}, __half2{3, 3}), __half2{1, 1})) && result[0];
result[0] = to_bool(__heq2(__hsub2(a, __half2{1, 1}), __half2{0, 0})) &&
result[0];
result[0] = to_bool(__heq2(__hsub2_sat(a, __half2{2, 2}), __half2{0, 0})) &&
result[0];
result[0] = to_bool(__heq2(__hmul2(a, __half2{2, 2}), __half2{2, 2})) &&
result[0];
result[0] = to_bool(__heq2(__hmul2_sat(a, __half2{2, 2}), __half2{1, 1})) &&
result[0];
result[0] = to_bool(__heq2(__h2div(a, __half2{2, 2}), __half2{0.5, 0.5})) &&
result[0];
}
__global__
void kernel_hisnan(__half* input, int* output) {
int tx = threadIdx.x;
output[tx] = __hisnan(input[tx]);
}
__global__
void kernel_hisinf(__half* input, int* output) {
int tx = threadIdx.x;
output[tx] = __hisinf(input[tx]);
}
__global__ void testHalfAbs(float* p) {
auto a = __float2half(*p);
a = __habs(a);
*p = __half2float(a);
}
__global__ void testHalf2Abs(float2* p) {
auto a = __float22half2_rn(*p);
a = __habs2(a);
*p = __half22float2(a);
}
__half host_ushort_as_half(uint32_t s) {
union {__half h; uint32_t s; } converter;
converter.s = s;
return converter.h;
}
void check_hisnan(int NUM_INPUTS, __half* inputCPU, __half* inputGPU) {
// allocate memory
auto memsize = NUM_INPUTS * sizeof(int);
int* outputGPU = nullptr;
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&outputGPU), memsize));
// launch the kernel
hipLaunchKernelGGL(
kernel_hisnan, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU);
// copy output from device
int* outputCPU = reinterpret_cast<int*> (malloc(memsize));
HIP_CHECK(hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost));
// check output
for (int i=0; i < NUM_INPUTS; i++) {
if ((2 <= i) && (i <= 5)) { // inputs are nan, output should be true
REQUIRE(outputCPU[i] == true);
} else { // inputs are NOT nan, output should be false
REQUIRE(outputCPU[i] == false);
}
}
// free memory
free(outputCPU);
HIP_CHECK(hipFree(outputGPU));
}
void check_hisinf(int NUM_INPUTS, __half* inputCPU, __half* inputGPU) {
// allocate memory
auto memsize = NUM_INPUTS * sizeof(int);
int* outputGPU = nullptr;
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&outputGPU), memsize));
// launch the kernel
hipLaunchKernelGGL(
kernel_hisinf, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU);
// copy output from device
int* outputCPU = reinterpret_cast<int*> (malloc(memsize));
HIP_CHECK(hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost));
// check output
for (int i=0; i < NUM_INPUTS; i++) {
if ((0 <= i) && (i <= 1)) { // inputs are inf, output should be true
REQUIRE(outputCPU[i] == true);
} else { // inputs are NOT inf, output should be false
REQUIRE(outputCPU[i] == false);
}
}
// free memory
free(outputCPU);
HIP_CHECK(hipFree(outputGPU));
}
void checkFunctional() {
// allocate memory
const int NUM_INPUTS = 16;
auto memsize = NUM_INPUTS * sizeof(__half);
__half* inputCPU = reinterpret_cast<__half*> (malloc(memsize));
// populate inputs
inputCPU[0] = host_ushort_as_half(0x7c00); // inf
inputCPU[1] = host_ushort_as_half(0xfc00); // -inf
inputCPU[2] = host_ushort_as_half(0x7c01); // nan
inputCPU[3] = host_ushort_as_half(0x7e00); // nan
inputCPU[4] = host_ushort_as_half(0xfc01); // nan
inputCPU[5] = host_ushort_as_half(0xfe00); // nan
inputCPU[6] = host_ushort_as_half(0x0000); // 0
inputCPU[7] = host_ushort_as_half(0x8000); // -0
inputCPU[8] = host_ushort_as_half(0x7bff); // max +ve normal
inputCPU[9] = host_ushort_as_half(0xfbff); // max -ve normal
inputCPU[10] = host_ushort_as_half(0x0400); // min +ve normal
inputCPU[11] = host_ushort_as_half(0x8400); // min -ve normal
inputCPU[12] = host_ushort_as_half(0x03ff); // max +ve sub-normal
inputCPU[13] = host_ushort_as_half(0x83ff); // max -ve sub-normal
inputCPU[14] = host_ushort_as_half(0x0001); // min +ve sub-normal
inputCPU[15] = host_ushort_as_half(0x8001); // min -ve sub-normal
// copy inputs to the GPU
__half* inputGPU = nullptr;
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&inputGPU), memsize));
HIP_CHECK(hipMemcpy(inputGPU, inputCPU, memsize, hipMemcpyHostToDevice));
// run checks
check_hisnan(NUM_INPUTS, inputCPU, inputGPU);
check_hisinf(NUM_INPUTS, inputCPU, inputGPU);
// free memory
HIP_CHECK(hipFree(inputGPU));
free(inputCPU);
}
void checkHalfAbs() {
SECTION("Half Abs") {
float *p;
HIP_CHECK(hipMalloc(&p, sizeof(float)));
float pp = -2.1f;
HIP_CHECK(hipMemcpy(p, &pp, sizeof(float), hipMemcpyDefault));
hipLaunchKernelGGL(testHalfAbs, 1, 1, 0, 0, p);
HIP_CHECK(hipMemcpy(&pp, p, sizeof(float), hipMemcpyDefault));
HIP_CHECK(hipFree(p));
REQUIRE(pp >= 0.0f);
}
SECTION("Half2 Abs") {
float2 *p;
HIP_CHECK(hipMalloc(&p, sizeof(float2)));
float2 pp;
pp.x = -2.1f;
pp.y = -1.1f;
HIP_CHECK(hipMemcpy(p, &pp, sizeof(float2), hipMemcpyDefault));
hipLaunchKernelGGL(testHalf2Abs, 1, 1, 0, 0, p);
HIP_CHECK(hipMemcpy(&pp, p, sizeof(float2), hipMemcpyDefault));
HIP_CHECK(hipFree(p));
bool result = true;
if (pp.x < 0.0f || pp.y < 0.0f) { result = false; }
REQUIRE(result == true);
}
}
TEST_CASE("Unit_hipTestHalf") {
bool* result{nullptr};
HIP_CHECK(hipHostMalloc(&result, sizeof(result)));
SECTION("Test half math") {
result[0] = false;
hipLaunchKernelGGL(
__halfMath, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result, __half{1});
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(result[0] == true);
}
SECTION("Test half math") {
result[0] = false;
hipLaunchKernelGGL(
__half2Math, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result, __half2{1, 1});
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(result[0] == true);
}
SECTION("Functional checks") {
checkFunctional();
checkHalfAbs();
}
HIP_CHECK(hipHostFree(result));
}
catch/unit/deviceLib/hipTestHost.cc
+408
Ver Arquivo
@@ -0,0 +1,408 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#define N 512
__device__ bool check_erfcinvf() {
uint32_t len = 4;
float Val[] = {0.1, 1.2, 1, 0.9};
float Out[] = {1.16309, -0.179144, 0, 0.0889};
for (int i = 0; i < len; i++) {
if ((Out[i] - erfcinvf(Val[i])) > 0.0001) {
return false;
}
}
return true;
}
__device__ bool check_erfcxf() {
uint32_t len = 4;
float Val[] = {-0.5, 15, 3.2, 1};
float Out[] = {1.9524, 0.0375, 0.1687, 0.4276};
for (int i = 0; i < len; i++) {
if (Out[i] - erfcxf(Val[i]) > 0.0001) {
return false;
}
}
return true;
}
__device__ bool check_erfinvf() {
uint32_t len = 4;
float Val[] = {0, -0.5, 0.9, -0.2};
float Out[] = {0, -0.4769, 1.1631, -0.1791};
for (int i = 0; i < len; i++) {
if (Out[i] - erfinvf(Val[i]) > 0.0001) {
return false;
}
}
return true;
}
__device__ bool check_fdividef() {
uint32_t len = 4;
float Val[] = {0, -0.5, 0.9, -0.2};
float Out[] = {1, -0.4769, 1.1631, -0.1791};
for (int i = 0; i < len; i++) {
if (Val[i] / Out[i] - fdividef(Val[i], Out[i]) > 0.0001) {
return false;
}
}
return true;
}
__device__ bool check_erfcinv() {
uint32_t len = 4;
double Val[] = {0.1, 1.2, 1, 0.9};
double Out[] = {1.16309, -0.179144, 0, 0.0889};
for (int i = 0; i < len; i++) {
if (Out[i] - erfcinv(Val[i]) > 0.0001) {
return false;
}
}
return true;
}
__device__ bool check_erfcx() {
uint32_t len = 4;
double Val[] = {-0.5, 15, 3.2, 1};
double Out[] = {1.9524, 0.0375, 0.1687, 0.4276};
for (int i = 0; i < len; i++) {
if (Out[i] - erfcx(Val[i]) > 0.0001) {
return false;
}
}
return true;
}
__device__ bool check_erfinv() {
uint32_t len = 4;
double Val[] = {0, -0.5, 0.9, -0.2};
double Out[] = {0, -0.4769, 1.1631, -0.1791};
for (int i = 0; i < len; i++) {
if (Out[i] - erfinv(Val[i]) > 0.0001) {
return false;
}
}
return true;
}
__device__ bool check_fdivide() {
uint32_t len = 4;
double Val[] = {0, -0.5, 0.9, -0.2};
double Out[] = {1, -0.4769, 1.1631, -0.1791};
for (int i = 0; i < len; i++) {
if (Val[i] / Out[i] - fdividef(Val[i], Out[i]) > 0.0001) {
return false;
}
}
return true;
}
__device__ bool check_modff() {
uint32_t len = 4;
float Val[] = {0, -0.5, 0.9, -0.2};
float iPtr[] = {0, 0, 0, 0};
float frac[] = {0, -0.5, 0.9, -0.2};
float Out[] = {1, 1, 1, 1};
for (int i = 0; i < len; i++) {
if (frac[i] - modff(Val[i], Out + i) > 0.0001 && iPtr[i] == Out[i]) {
return false;
}
}
return true;
}
__device__ bool check_modf() {
uint32_t len = 4;
double Val[] = {0, -0.5, 0.9, -0.2};
double iPtr[] = {0, 0, 0, 0};
double frac[] = {0, -0.5, 0.9, -0.2};
double Out[] = {1, 1, 1, 1};
for (int i = 0; i < len; i++) {
if (frac[i] - modf(Val[i], Out + i) > 0.0001 && iPtr[i] == Out[i]) {
return false;
}
}
return true;
}
__device__ bool check_nextafterf() {
uint32_t len = 4;
float Val[] = {0, -0.5, 0.9, -0.2};
float iPtr[] = {0, 0, 0, 0};
float frac[] = {0, -0.5, 0.9, -0.2};
float Out[] = {1, 1, 1, 1};
for (int i = 0; i < len; i++) {
if (nextafterf(Val[i], 1) - Val[i] > 0.0001) {
return false;
}
}
return true;
}
__device__ bool check_nextafter() {
uint32_t len = 4;
double Val[] = {0, -0.5, 0.9, -0.2};
double iPtr[] = {0, 0, 0, 0};
double frac[] = {0, -0.5, 0.9, -0.2};
double Out[] = {1, 1, 1, 1};
for (int i = 0; i < len; i++) {
if (nextafter(Val[i], 1) - Val[i] > 0.0001) {
return false;
}
}
return true;
}
__device__ bool check_norm3df(float* A) {
float f = norm3df(A[0], A[1], A[2]);
float out = sqrt(A[0] * A[0] + A[1] * A[1] + A[2] * A[2]);
if (f - out > 0.0001) {
return false;
}
return true;
}
__device__ bool check_norm3d(double* A) {
double f = norm3d(A[0], A[1], A[2]);
double out = sqrt(A[0] * A[0] + A[1] * A[1] + A[2] * A[2]);
if (f - out > 0.0001) {
return false;
}
return true;
}
__device__ bool check_norm4df(float* A) {
float f = norm4df(A[0], A[1], A[2], A[3]);
float out = sqrt(A[0] * A[0] + A[1] * A[1] + A[2] * A[2] + A[3] * A[3]);
if (f - out > 0.0001) {
return false;
}
return true;
}
__device__ bool check_norm4d(double* A) {
double f = norm4d(A[0], A[1], A[2], A[3]);
double out = sqrt(A[0] * A[0] + A[1] * A[1] + A[2] * A[2] + A[3] * A[3]);
if (f - out > 0.0001) {
return false;
}
return true;
}
__device__ bool check_normcdff() {
uint32_t len = 2;
float Val[] = {0, 1};
float Out[] = {0.5, 0.8413};
for (int i = 0; i < len; i++) {
if (Out[i] - normcdff(Val[i]) > 0.0001) {
return false;
}
}
return true;
}
__device__ bool check_normcdf() {
uint32_t len = 2;
float Val[] = {0, 1};
float Out[] = {0.5, 0.8413};
for (int i = 0; i < len; i++) {
if (Out[i] - normcdf(Val[i]) > 0.0001) {
return false;
}
}
return true;
}
__device__ bool check_normcdfinvf() {
uint32_t len = 2;
double Val[] = {0.5, 0.8413};
for (int i = 0; i < len; i++) {
if (Val[i] - normcdfinvf(normcdff(Val[i])) > 0.0001) {
return false;
}
}
return true;
}
__device__ bool check_normcdfinv() {
uint32_t len = 2;
double Val[] = {0.5, 0.8413};
for (int i = 0; i < len; i++) {
if (Val[i] - normcdfinv(normcdf(Val[i])) > 0.0001) {
return false;
}
}
return true;
}
__device__ bool check_rcbrtf() {
float f = 1.0f;
if (rcbrtf(f) != 1.0f) {
return false;
}
return true;
}
__device__ bool check_rcbrt() {
double f = 1.0;
if (rcbrt(f) != 1.0) {
return false;
}
return true;
}
__device__ bool check_rhypotf() {
float f = 1.0f;
float g = 2.0f;
float val = rhypotf(f, g);
float sq = f * f + g * g;
if (1 / (val * val) - sq > 0.0001) {
return false;
}
return true;
}
__device__ bool check_rhypot() {
double f = 1.0f;
double g = 2.0f;
double val = rhypot(f, g);
double sq = f * f + g * g;
if (1 / (val * val) - sq > 0.0001) {
return false;
}
return true;
}
__device__ bool check_rnorm3df(float* A) {
float f = rnorm3df(A[0], A[1], A[2]);
float out = sqrt(A[0] * A[0] + A[1] * A[1] + A[2] * A[2]);
if (f - 1 / out > 0.0001) {
return false;
}
return true;
}
__device__ bool check_rnorm3d(double* A) {
double f = rnorm3d(A[0], A[1], A[2]);
double out = sqrt(A[0] * A[0] + A[1] * A[1] + A[2] * A[2]);
if (f - 1 / out > 0.0001) {
return false;
}
return true;
}
__device__ bool check_rnorm4df(float* A) {
float f = rnorm4df(A[0], A[1], A[2], A[3]);
float out = sqrt(A[0] * A[0] + A[1] * A[1] + A[2] * A[2] + A[3] * A[3]);
if (f - 1 / out > 0.0001) {
return false;
}
return true;
}
__device__ bool check_rnorm4d(double* A) {
double f = rnorm4d(A[0], A[1], A[2], A[3]);
double out = sqrt(A[0] * A[0] + A[1] * A[1] + A[2] * A[2] + A[3] * A[3]);
if (f - 1 / out > 0.0001) {
return false;
}
return true;
}
__device__ bool check_rnormf(float* A) {
return (rnorm3df(A[0], A[1], A[2]) - rnormf(3, A) < 0.0001) &&
(rnorm4df(A[0], A[1], A[2], A[3]) - rnormf(4, A) < 0.0001);
}
__device__ bool check_rnorm(double* A) {
return (rnorm3d(A[0], A[1], A[2]) - rnorm(3, A) < 0.0001) &&
(rnorm4d(A[0], A[1], A[2], A[3]) - rnorm(4, A) < 0.0001);
}
__device__ bool check_sincospif() {
float s1, c1, s2, c2;
float in1 = 1, in2 = 0.5;
sincospif(in1, &s1, &c1);
sincospif(in2, &s2, &c2);
if ((s1 - 0 < 0.00001) && (s2 - 1 < 0.00001) &&
(c1 + 1 < 0.00001) && (c2 - 0 < 0.00001)) {
return true;
}
return false;
}
__device__ bool check_sincospi() {
double s1, c1, s2, c2;
double in1 = 1, in2 = 0.5;
sincospi(in1, &s1, &c1);
sincospi(in2, &s2, &c2);
if ((s1 - 0 < 0.00001) && (s2 - 1 < 0.00001) &&
(c1 + 1 < 0.00001) && (c2 - 0 < 0.00001)) {
return true;
}
return false;
}
__global__ void testFunctions(bool *result, float *Af, double *A) {
result[0] &= check_erfcinvf() && check_erfcxf() && check_erfcinvf()
&& check_erfcinv() && check_erfcx() && check_erfcinv()
&& check_fdividef() && check_fdivide() && check_modff()
&& check_modf() && check_nextafterf() && check_norm3df(Af)
&& check_norm3d(A) && check_norm4df(Af) && check_norm4d(A)
&& check_normcdff() && check_normcdf() && check_normcdfinvf()
&& check_normcdfinv() && check_rcbrtf() && check_rcbrt() &&
check_rhypotf() && check_rhypot() && check_rnorm3df(Af) &&
check_rnorm3d(A) && check_rnorm4df(Af) && check_rnorm4d(A) &&
check_rnormf(Af) && check_rnorm(A) && check_sincospif() &&
check_sincospi() && check_nextafter();
}
TEST_CASE("Unit_TestDevice_DoublePrecisionMathFunc") {
float* Af = new float[N];
double* A = new double[N];
for (int i = 0; i < N; i++) {
Af[i] = i * 1.0f;
A[i] = i * 1.0;
}
float *Afd;
double *Ad;
bool *srcPtr, *devicePtr;
srcPtr = new bool;
srcPtr[0] = true;
// Device pointers
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&devicePtr), sizeof(bool)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Afd), sizeof(float)*N));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), sizeof(double)*N));
// MemCpy
HIP_CHECK(hipMemcpy(devicePtr, srcPtr, sizeof(bool), hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Afd, Af, sizeof(float)*N, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(Ad, A, sizeof(double)*N, hipMemcpyHostToDevice));
// Kernel Launch
hipLaunchKernelGGL(testFunctions, dim3(1), dim3(1), 0, 0, devicePtr, Afd, Ad);
HIP_CHECK(hipMemcpy(srcPtr, devicePtr, sizeof(bool), hipMemcpyDeviceToHost));
// Validation
REQUIRE(srcPtr[0] == true);
HIP_CHECK(hipFree(devicePtr));
delete srcPtr;
delete [] Af;
delete [] A;
}
catch/unit/deviceLib/hipTestIncludeMath.cc
+134
Ver Arquivo
@@ -0,0 +1,134 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <new>
// Test __HIP_DEVICE_COMPILE__ is defined after math_functions.h
// is included.
__device__ __host__ inline void throw_std_bad_alloc() {
#ifndef __HIP_DEVICE_COMPILE__
throw std::bad_alloc();
#else
std::size_t kHuge = static_cast<std::size_t>(-1);
new int[kHuge];
#endif
}
__global__ void FloatMathPreciseKernel() {
int iX;
float fX, fY;
acosf(1.0f);
acoshf(1.0f);
asinf(0.0f);
asinhf(0.0f);
atan2f(0.0f, 1.0f);
atanf(0.0f);
atanhf(0.0f);
cbrtf(0.0f);
fX = ceilf(0.0f);
fX = copysignf(1.0f, -2.0f);
cosf(0.0f);
coshf(0.0f);
cospif(0.0f);
cyl_bessel_i0f(0.0f);
cyl_bessel_i1f(0.0f);
erfcf(0.0f);
erfcinvf(2.0f);
erfcxf(0.0f);
erff(0.0f);
erfinvf(1.0f);
exp10f(0.0f);
exp2f(0.0f);
expf(0.0f);
expm1f(0.0f);
fX = fabsf(1.0f);
fdimf(1.0f, 0.0f);
fdividef(0.0f, 1.0f);
fX = floorf(0.0f);
fmaf(1.0f, 2.0f, 3.0f);
fX = fmaxf(0.0f, 0.0f);
fX = fminf(0.0f, 0.0f);
fmodf(0.0f, 1.0f);
frexpf(0.0f, &iX);
hypotf(1.0f, 0.0f);
ilogbf(1.0f);
isfinite(0.0f);
fX = isinf(0.0f);
fX = isnan(0.0f);
j0f(0.0f);
j1f(0.0f);
jnf(-1.0f, 1.0f);
ldexpf(0.0f, 0);
lgammaf(1.0f);
llrintf(0.0f);
llroundf(0.0f);
log10f(1.0f);
log1pf(-1.0f);
log2f(1.0f);
logbf(1.0f);
logf(1.0f);
lrintf(0.0f);
lroundf(0.0f);
modff(0.0f, &fX);
fX = nanf("1");
fX = nearbyintf(0.0f);
nextafterf(0.0f, 0.0f);
norm3df(1.0f, 0.0f, 0.0f);
norm4df(1.0f, 0.0f, 0.0f, 0.0f);
normcdff(0.0f);
normcdfinvf(1.0f);
fX = 1.0f;
normf(1, &fX);
powf(1.0f, 0.0f);
rcbrtf(1.0f);
remainderf(2.0f, 1.0f);
remquof(1.0f, 2.0f, &iX);
rhypotf(0.0f, 1.0f);
fY = rintf(1.0f);
rnorm3df(0.0f, 0.0f, 1.0f);
rnorm4df(0.0f, 0.0f, 0.0f, 1.0f);
fX = 1.0f;
rnormf(1, &fX);
fY = roundf(0.0f);
rsqrtf(1.0f);
scalblnf(0.0f, 1);
scalbnf(0.0f, 1);
signbit(1.0f);
sincosf(0.0f, &fX, &fY);
sincospif(0.0f, &fX, &fY);
sinf(0.0f);
sinhf(0.0f);
sinpif(0.0f);
sqrtf(0.0f);
tanf(0.0f);
tanhf(0.0f);
tgammaf(2.0f);
fY = truncf(0.0f);
y0f(1.0f);
y1f(1.0f);
ynf(1, 1.0f);
}
TEST_CASE("Unit_TestIncludeMathPreciseFloat") {
hipError_t err;
err = hipLaunchKernel(reinterpret_cast<void *>(FloatMathPreciseKernel),
dim3(1, 1, 1),
dim3(1, 1, 1), 0, 0, 0);
REQUIRE(err == hipSuccess);
}
catch/unit/deviceLib/hipTestNativeHalf.cc
+169
Ver Arquivo
@@ -0,0 +1,169 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip/hip_fp16.h>
#include <hip_test_common.hh>
#include <type_traits>
__global__
void __halfTest(bool* result, __half a) {
// Construction
result[0] &= std::is_default_constructible<__half>{};
result[0] &= std::is_copy_constructible<__half>{};
result[0] &= std::is_move_constructible<__half>{};
result[0] &= std::is_constructible<__half, float>{};
result[0] &= std::is_constructible<__half, double>{};
result[0] &= std::is_constructible<__half, uint32_t>{};
result[0] &= std::is_constructible<__half, int32_t>{};
result[0] &= std::is_constructible<__half, uint32_t>{};
result[0] &= std::is_constructible<__half, int>{};
result[0] &= std::is_constructible<__half, uint64_t>{};
result[0] &= std::is_constructible<__half, int64_t>{};
result[0] &= std::is_constructible<__half, int64_t>{};
result[0] &= std::is_constructible<__half, uint64_t>{};
result[0] &= std::is_constructible<__half, __half_raw>{};
// Assignment
result[0] &= std::is_copy_assignable<__half>{};
result[0] &= std::is_move_assignable<__half>{};
result[0] &= std::is_assignable<__half, float>{};
result[0] &= std::is_assignable<__half, double>{};
result[0] &= std::is_assignable<__half, uint32_t>{};
result[0] &= std::is_assignable<__half, int32_t>{};
result[0] &= std::is_assignable<__half, uint32_t>{};
result[0] &= std::is_assignable<__half, int>{};
result[0] &= std::is_assignable<__half, uint64_t>{};
result[0] &= std::is_assignable<__half, int64_t>{};
result[0] &= std::is_assignable<__half, int64_t>{};
result[0] &= std::is_assignable<__half, uint64_t>{};
result[0] &= std::is_assignable<__half, __half_raw>{};
result[0] &= std::is_assignable<__half, volatile __half_raw&>{};
result[0] &= std::is_assignable<__half, volatile __half_raw&&>{};
// Conversion
result[0] &= std::is_convertible<__half, float>{};
result[0] &= std::is_convertible<__half, uint32_t>{};
result[0] &= std::is_convertible<__half, int32_t>{};
result[0] &= std::is_convertible<__half, uint32_t>{};
result[0] &= std::is_convertible<__half, int>{};
result[0] &= std::is_convertible<__half, uint64_t>{};
result[0] &= std::is_convertible<__half, int64_t>{};
result[0] &= std::is_convertible<__half, int64_t>{};
result[0] &= std::is_convertible<__half, bool>{};
result[0] &= std::is_convertible<__half, uint64_t>{};
result[0] &= std::is_convertible<__half, __half_raw>{};
result[0] &= std::is_convertible<__half, volatile __half_raw>{};
// Nullary
result[0] &= __heq(a, +a) && result[0];
result[0] &= __heq(__hneg(a), -a) && result[0];
// Unary arithmetic
result[0] &= __heq(a += 0, a) && result[0];
result[0] &= __heq(a -= 0, a) && result[0];
result[0] &= __heq(a *= 1, a) && result[0];
result[0] &= __heq(a /= 1, a) && result[0];
// Binary arithmetic
result[0] &= __heq((a + a), __hadd(a, a)) && result[0];
result[0] &= __heq((a - a), __hsub(a, a)) && result[0];
result[0] &= __heq((a * a), __hmul(a, a)) && result[0];
result[0] &= __heq((a / a), __hdiv(a, a)) && result[0];
// Relations
result[0] &= (a == a) && result[0];
result[0] &= !(a != a) && result[0];
result[0] &= (a <= a) && result[0];
result[0] &= (a >= a) && result[0];
result[0] &= !(a < a) && result[0];
result[0] &= !(a > a) && result[0];
}
__device__
static bool to_bool(const __half2& x) {
auto r = static_cast<const __half2_raw&>(x);
return r.data.x != 0 && r.data.y != 0;
}
__global__
void __half2Test(bool* result, __half2 a) {
// Construction
result[0] &= std::is_default_constructible<__half2>{};
result[0] &= std::is_copy_constructible<__half2>{};
result[0] &= std::is_move_constructible<__half2>{};
result[0] &= std::is_constructible<__half2, __half, __half>{};
result[0] &= std::is_constructible<__half2, __half2_raw>{};
// Assignment
result[0] &= std::is_copy_assignable<__half2>{};
result[0] &= std::is_move_assignable<__half2>{};
result[0] &= std::is_assignable<__half2, __half2_raw>{};
// Conversion
result[0] &= std::is_convertible<__half2, __half2_raw>{};
// Nullary
result[0] &= to_bool(__heq2(a, +a)) && result[0];
result[0] &= to_bool(__heq2(__hneg2(a), -a)) && result[0];
// Unary arithmetic
result[0] &= to_bool(__heq2(a += 0, a)) && result[0];
result[0] &= to_bool(__heq2(a -= 0, a)) && result[0];
result[0] &= to_bool(__heq2(a *= 1, a)) && result[0];
result[0] &= to_bool(__heq2(a /= 1, a)) && result[0];
// Binary arithmetic
result[0] &= to_bool(__heq2((a + a), __hadd2(a, a))) && result[0];
result[0] &= to_bool(__heq2((a - a), __hsub2(a, a))) && result[0];
result[0] &= to_bool(__heq2((a * a), __hmul2(a, a))) && result[0];
result[0] &= to_bool(__heq2((a / a), __h2div(a, a))) && result[0];
// Relations
result[0] &= (a == a) && result[0];
result[0] &= !(a != a) && result[0];
result[0] &= (a <= a) && result[0];
result[0] &= (a >= a) && result[0];
result[0] &= !(a < a) && result[0];
result[0] &= !(a > a) && result[0];
// Dot Functions
result[0] &= amd_mixed_dot(a, a, 1, 1) && result[0];
half X = a.x;
half Y = a.y;
}
TEST_CASE("Unit_hipTestNativeHalf") {
bool* result{nullptr};
HIP_CHECK(hipHostMalloc(&result, 1));
SECTION("Half Test") {
result[0] = true;
hipLaunchKernelGGL(
__halfTest, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result, __half{1});
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(result[0] == true);
}
SECTION("Half2 Test") {
result[0] = true;
hipLaunchKernelGGL(
__half2Test, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result, __half2{1, 1});
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(result[0] == true);
}
HIP_CHECK(hipHostFree(result));
}
catch/unit/deviceLib/hipTestNew.cc
+54
Ver Arquivo
@@ -0,0 +1,54 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#define LEN 512
#define SIZE 2048
class A {
public:
__device__ A() {
a = threadIdx.x + blockIdx.x * blockDim.x;
}
private:
int a;
};
static __global__ void kernel(int* Ad) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
new(Ad+tid) A();
}
TEST_CASE("Unit_hipTest_DeviceNewOperator") {
int *A, *Ad;
A = new int[LEN];
for (unsigned i = 0; i < LEN; i++) {
A[i] = 0;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad), SIZE));
hipLaunchKernelGGL(kernel, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, Ad);
HIP_CHECK(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost));
// Validation
for (unsigned i = 0; i < LEN; i++) {
REQUIRE(i == A[i]);
}
delete[] A;
}
catch/unit/deviceLib/hipThreadFence.cc
+76
Ver Arquivo
@@ -0,0 +1,76 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#define NUM 1024
#define SIZE (NUM * sizeof(float))
__global__ static void vAdd(float* In1, float* In2, float* In3,
float* In4, float* Out) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
In4[tid] = In1[tid] + In2[tid];
__threadfence();
In3[tid] = In3[tid] + In4[tid];
__threadfence_block();
Out[tid] = In4[tid] + In3[tid];
}
TEST_CASE("Unit_hipThreadFence") {
float* In1 = new float[NUM];
float* In2 = new float[NUM];
float* In3 = new float[NUM];
float* In4 = new float[NUM];
float* Out = new float[NUM];
// Initialization
for (uint32_t i = 0; i < NUM; i++) {
In1[i] = 1.0f;
In2[i] = 1.0f;
In3[i] = 1.0f;
In4[i] = 1.0f;
}
float *In1d, *In2d, *In3d, *In4d, *Outd;
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&In1d), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&In2d), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&In3d), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&In4d), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Outd), SIZE));
HIP_CHECK(hipMemcpy(In1d, In1, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(In2d, In2, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(In3d, In3, SIZE, hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(In4d, In4, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(vAdd, dim3(32, 1, 1), dim3(32, 1, 1), 0, 0,
In1d, In2d, In3d, In4d, Outd);
HIP_CHECK(hipMemcpy(Out, Outd, SIZE, hipMemcpyDeviceToHost));
for (uint32_t i = 0; i < NUM; i++) {
REQUIRE(Out[i] == 2 * In1[i] + 2 * In2[i] + In3[i]);
}
delete[] In1;
delete[] In2;
delete[] In3;
delete[] In4;
delete[] Out;
HIP_CHECK(hipFree(In1d));
HIP_CHECK(hipFree(In2d));
HIP_CHECK(hipFree(In3d));
HIP_CHECK(hipFree(In4d));
HIP_CHECK(hipFree(Outd));
}
catch/unit/deviceLib/hipVectorTypes.cc
+204
Ver Arquivo
@@ -0,0 +1,204 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip/hip_vector_types.h>
#include <string>
#include <sstream>
#include <type_traits>
using namespace std; // NOLINT
template<typename V,
enable_if_t<!is_integral<decltype(declval<V>().x)>{}>* = nullptr>
bool integer_unary_tests(V&, V&) {
return true;
}
template<typename V,
enable_if_t<!is_integral<decltype(declval<V>().x)>{}>* = nullptr>
bool integer_binary_tests(V&, V&, V&...) {
return true;
}
template<typename V,
enable_if_t<is_integral<decltype(declval<V>().x)>{}>* = nullptr>
bool integer_unary_tests(V f1, V f2) {
f1 %= f2;
if (f1 != V{0}) return false;
f1 &= f2;
if (f1 != V{0}) return false;
f1 |= f2;
if (f1 != V{1}) return false;
f1 ^= f2;
if (f1 != V{0}) return false;
f1 = V{1};
f1 <<= f2;
if (f1 != V{2}) return false;
f1 >>= f2;
if (f1 != V{1}) return false;
f2 = ~f1;
return f2 == V{~1};
}
template<typename V,
enable_if_t<is_integral<decltype(declval<V>().x)>{}>* = nullptr>
bool integer_binary_tests(V f1, V f2, V f3) {
f3 = f1 % f2;
if (f3 != V{0}) return false;
f1 = f3 & f2;
if (f1 != V{0}) return false;
f2 = f1 ^ f3;
if (f2 != V{0}) return false;
f1 = V{1};
f2 = V{2};
f3 = f1 << f2;
if (f3 != V{4}) return false;
f2 = f3 >> f1;
return f2 == V{2};
}
template<typename V>
bool constructor_tests() {
if (is_constructible<V, unsigned char>{} &&
is_constructible<V, signed char>{} &&
is_constructible<V, uint32_t>{} &&
is_constructible<V, int32_t>{} &&
is_constructible<V, unsigned int>{} &&
is_constructible<V, signed int>{} &&
is_constructible<V, uint64_t>{} &&
is_constructible<V, int64_t>{} &&
is_constructible<V, uint64_t>{} &&
is_constructible<V, int64_t>{} &&
is_constructible<V, float>{} &&
is_constructible<V, double>{}) {
return true;
}
}
template<typename V>
bool TestVectorType() {
constexpr V v1{1};
constexpr V v2{2};
constexpr V v3{3};
constexpr V v4{4};
V f1{1};
V f2{1};
V f3 = f1 + f2;
if (f3 != v2) return false;
f2 = f3 - f1;
if (f2 != v1) return false;
f1 = f2 * f3;
if (f1 != v2) return false;
f2 = f1 / f3;
if (f2 != v1) return false;
if (!integer_binary_tests(f1, f2, f3)) return false;
f1 = V{2};
f2 = V{1};
f1 += f2;
if (f1 != v3) return false;
f1 -= f2;
if (f1 != v2) return false;
f1 *= f2;
if (f1 != v2) return false;
f1 /= f2;
if (f1 != v2) return false;
if (!integer_unary_tests(f1, f2)) return false;
f1 = v2;
f2 = f1++;
if (f1 != v3) return false;
if (f2 != v2) return false;
f2 = f1--;
if (f2 != v3) return false;
if (f1 != v2) return false;
f2 = ++f1;
if (f1 != v3) return false;
if (f2 != v3) return false;
f2 = --f1;
if (f1 != v2) return false;
if (f2 != v2) return false;
if (!constructor_tests<V>()) return false;
f1 = v3;
f2 = v4;
f3 = v3;
if (f1 == f2) return false;
if (!(f1 != f2)) return false;
using T = typename V::value_type;
const T& x = f1.x;
T& y = f2.x;
const volatile T& z = f3.x;
volatile T& w = f2.x;
if (x != T{3}) return false;
if (y != T{4}) return false;
if (z != T{3}) return false;
if (w != T{4}) return false;
stringstream str;
str << f1.x;
str >> f2.x;
if (f1.x != f2.x) return false;
return true;
}
template<typename... Ts, enable_if_t<sizeof...(Ts) == 0>* = nullptr>
bool TestVectorTypes() {
return true;
}
template<typename T, typename... Ts>
bool TestVectorTypes() {
if (!TestVectorType<T>()) return false;
return TestVectorTypes<Ts...>();
}
bool CheckVectorTypes() {
return TestVectorTypes<
char1, char2, char3, char4,
uchar1, uchar2, uchar3, uchar4,
short1, short2, short3, short4,
ushort1, ushort2, ushort3, ushort4,
int1, int2, int3, int4,
uint1, uint2, uint3, uint4,
long1, long2, long3, long4,
ulong1, ulong2, ulong3, ulong4,
longlong1, longlong2, longlong3, longlong4,
ulonglong1, ulonglong2, ulonglong3, ulonglong4,
float1, float2, float3, float4,
double1, double2, double3, double4>();
}
TEST_CASE("Unit_TestVectorTypes") {
REQUIRE(sizeof(float1) == 4);
REQUIRE(sizeof(float2) >= 8);
REQUIRE(sizeof(float3) == 12);
REQUIRE(sizeof(float4) >= 16);
bool result = false;
result = CheckVectorTypes();
REQUIRE(result == true);
}
catch/unit/deviceLib/hip_test_make_type.cc
+334
Ver Arquivo
@@ -0,0 +1,334 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip/hip_runtime.h>
#include <hip_test_common.hh>
#define WIDTH 8
#define HEIGHT 8
#define NUM (WIDTH * HEIGHT)
#define THREADS_PER_BLOCK_X 8
#define THREADS_PER_BLOCK_Y 8
#define THREADS_PER_BLOCK_Z 1
__global__ void vectoradd_char1(char1* a, const char1* bm, const char1* cm,
int width, int height) {
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
int i = y * width + x;
if (i < (width * height)) {
a[i] = make_char1(bm[i].x) + make_char1(cm[i].x);
}
}
__global__ void vectoradd_char2(char2* a, const char2* bm, const char2* cm,
int width, int height) {
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
int i = y * width + x;
if (i < (width * height)) {
a[i] = make_char2(bm[i].x, bm[i].y) + make_char2(cm[i].x, cm[i].y);
}
}
__global__ void vectoradd_char3(char3* a, const char3* bm, const char3* cm,
int width, int height) {
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
int i = y * width + x;
if (i < (width * height)) {
a[i] = make_char3(bm[i].x, bm[i].y, bm[i].z) + make_char3(cm[i].x,
cm[i].y, cm[i].z);
}
}
__global__ void vectoradd_char4(char4* a, const char4* bm, const char4* cm,
int width, int height) {
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
int i = y * width + x;
if (i < (width * height)) {
a[i] = make_char4(bm[i].x, bm[i].y, bm[i].z, bm[i].w) +
make_char4(cm[i].x, cm[i].y, cm[i].z, cm[i].w);
}
}
template <typename T>
bool dataTypesRunChar1() {
T* hostA;
T* hostB;
T* hostC;
T* deviceA;
T* deviceB;
T* deviceC;
int i;
int errors;
hostA = reinterpret_cast<T*>(malloc(NUM * sizeof(T)));
hostB = reinterpret_cast<T*>(malloc(NUM * sizeof(T)));
hostC = reinterpret_cast<T*>(malloc(NUM * sizeof(T)));
// initialize the input data
for (i = 0; i < NUM; i++) {
hostB[i] = (T)i;
hostC[i] = (T)i;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&deviceA), NUM * sizeof(T)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&deviceB), NUM * sizeof(T)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&deviceC), NUM * sizeof(T)));
HIP_CHECK(hipMemcpy(deviceB, hostB, NUM * sizeof(T), hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(deviceC, hostC, NUM * sizeof(T), hipMemcpyHostToDevice));
hipLaunchKernelGGL(HIP_KERNEL_NAME(vectoradd_char1),
dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT /
THREADS_PER_BLOCK_Y), dim3(THREADS_PER_BLOCK_X,
THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB, deviceC,
WIDTH, HEIGHT);
HIP_CHECK(hipMemcpy(hostA, deviceA, NUM * sizeof(T), hipMemcpyDeviceToHost));
bool ret = false;
// verify the results
errors = 0;
for (i = 0; i < NUM; i++) {
if (hostA[i] != (hostB[i] + hostC[i])) {
errors++;
}
}
if (errors != 0) {
ret = false;
} else {
ret = true;
}
HIP_CHECK(hipFree(deviceA));
HIP_CHECK(hipFree(deviceB));
HIP_CHECK(hipFree(deviceC));
free(hostA);
free(hostB);
free(hostC);
return ret;
}
template <typename T>
bool dataTypesRunChar2() {
T* hostA;
T* hostB;
T* hostC;
T* deviceA;
T* deviceB;
T* deviceC;
int i;
int errors;
hostA = reinterpret_cast<T*>(malloc(NUM * sizeof(T)));
hostB = reinterpret_cast<T*>(malloc(NUM * sizeof(T)));
hostC = reinterpret_cast<T*>(malloc(NUM * sizeof(T)));
// initialize the input data
for (i = 0; i < NUM; i++) {
hostB[i] = (T)i;
hostC[i] = (T)i;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&deviceA), NUM * sizeof(T)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&deviceB), NUM * sizeof(T)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&deviceC), NUM * sizeof(T)));
HIP_CHECK(hipMemcpy(deviceB, hostB, NUM * sizeof(T), hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(deviceC, hostC, NUM * sizeof(T), hipMemcpyHostToDevice));
hipLaunchKernelGGL(HIP_KERNEL_NAME(vectoradd_char2),
dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT /
THREADS_PER_BLOCK_Y), dim3(THREADS_PER_BLOCK_X,
THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB, deviceC,
WIDTH, HEIGHT);
HIP_CHECK(hipMemcpy(hostA, deviceA, NUM * sizeof(T), hipMemcpyDeviceToHost));
bool ret = false;
// verify the results
errors = 0;
for (i = 0; i < NUM; i++) {
if (hostA[i] != (hostB[i] + hostC[i])) {
errors++;
}
}
if (errors != 0) {
ret = false;
} else {
ret = true;
}
HIP_CHECK(hipFree(deviceA));
HIP_CHECK(hipFree(deviceB));
HIP_CHECK(hipFree(deviceC));
free(hostA);
free(hostB);
free(hostC);
return ret;
}
template <typename T>
bool dataTypesRunChar3() {
T* hostA;
T* hostB;
T* hostC;
T* deviceA;
T* deviceB;
T* deviceC;
int i;
int errors;
hostA = reinterpret_cast<T*>(malloc(NUM * sizeof(T)));
hostB = reinterpret_cast<T*>(malloc(NUM * sizeof(T)));
hostC = reinterpret_cast<T*>(malloc(NUM * sizeof(T)));
// initialize the input data
for (i = 0; i < NUM; i++) {
hostB[i] = (T)i;
hostC[i] = (T)i;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&deviceA), NUM * sizeof(T)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&deviceB), NUM * sizeof(T)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&deviceC), NUM * sizeof(T)));
HIP_CHECK(hipMemcpy(deviceB, hostB, NUM * sizeof(T), hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(deviceC, hostC, NUM * sizeof(T), hipMemcpyHostToDevice));
hipLaunchKernelGGL(HIP_KERNEL_NAME(vectoradd_char3),
dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT /
THREADS_PER_BLOCK_Y), dim3(THREADS_PER_BLOCK_X,
THREADS_PER_BLOCK_Y), 0, 0, deviceA, deviceB, deviceC,
WIDTH, HEIGHT);
HIP_CHECK(hipMemcpy(hostA, deviceA, NUM * sizeof(T), hipMemcpyDeviceToHost));
bool ret = false;
// verify the results
errors = 0;
for (i = 0; i < NUM; i++) {
if (hostA[i] != (hostB[i] + hostC[i])) {
errors++;
}
}
if (errors != 0) {
ret = false;
} else {
ret = true;
}
HIP_CHECK(hipFree(deviceA));
HIP_CHECK(hipFree(deviceB));
HIP_CHECK(hipFree(deviceC));
free(hostA);
free(hostB);
free(hostC);
return ret;
}
template <typename T>
bool dataTypesRunChar4() {
char4* hostA;
char4* hostB;
char4* hostC;
char4* deviceA;
char4* deviceB;
char4* deviceC;
int i;
int errors;
hostA = reinterpret_cast<T*>(malloc(NUM * sizeof(T)));
hostB = reinterpret_cast<T*>(malloc(NUM * sizeof(T)));
hostC = reinterpret_cast<T*>(malloc(NUM * sizeof(T)));
// initialize the input data
for (i = 0; i < NUM; i++) {
hostB[i] = (T)i;
hostC[i] = (T)i;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&deviceA), NUM * sizeof(T)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&deviceB), NUM * sizeof(T)));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&deviceC), NUM * sizeof(T)));
HIP_CHECK(hipMemcpy(deviceB, hostB, NUM * sizeof(T), hipMemcpyHostToDevice));
HIP_CHECK(hipMemcpy(deviceC, hostC, NUM * sizeof(T), hipMemcpyHostToDevice));
hipLaunchKernelGGL(HIP_KERNEL_NAME(vectoradd_char4),
dim3(WIDTH / THREADS_PER_BLOCK_X, HEIGHT /
THREADS_PER_BLOCK_Y), dim3(THREADS_PER_BLOCK_X,
THREADS_PER_BLOCK_Y), 0, 0, deviceA,
deviceB, deviceC, WIDTH, HEIGHT);
HIP_CHECK(hipMemcpy(hostA, deviceA, NUM * sizeof(T), hipMemcpyDeviceToHost));
bool ret = false;
// verify the results
errors = 0;
for (i = 0; i < NUM; i++) {
if (hostA[i] != (hostB[i] + hostC[i])) {
errors++;
}
}
if (errors != 0) {
ret = false;
} else {
ret = true;
}
HIP_CHECK(hipFree(deviceA));
HIP_CHECK(hipFree(deviceB));
HIP_CHECK(hipFree(deviceC));
free(hostA);
free(hostB);
free(hostC);
return ret;
}
TEST_CASE("Unit_Test_makechar_functionality") {
bool errors;
errors = dataTypesRunChar1<char1>() && dataTypesRunChar2<char2>() &&
dataTypesRunChar3<char3>() && dataTypesRunChar4<char4>();
REQUIRE(errors == true);
}
catch/unit/deviceLib/hip_trig.cc
+86
Ver Arquivo
@@ -0,0 +1,86 @@
/*
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 WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip/device_functions.h>
#define LEN 512
#define SIZE (LEN << 2)
__global__ static void kernel_trig(float* In, float* sin_d,
float* cos_d, float* tan_d,
float* sin_pd, float* cos_pd) {
int tid = threadIdx.x + blockIdx.x * blockDim.x;
sin_d[tid] = sinf(In[tid]);
cos_d[tid] = cosf(In[tid]);
tan_d[tid] = tanf(In[tid]);
sincosf(In[tid], &sin_pd[tid], &cos_pd[tid]);
}
TEST_CASE("Unit_kernel_trigger") {
float *In, *sin_h, *cos_h, *tan_h, *sin_ph, *cos_ph;
float *In_d, *sin_d, *cos_d, *tan_d, *sin_pd, *cos_pd;
int errors = 0;
In = new float[LEN];
sin_h = new float[LEN];
cos_h = new float[LEN];
tan_h = new float[LEN];
sin_ph = new float[LEN];
cos_ph = new float[LEN];
for (int i = 0; i < LEN; i++) {
In[i] = 1.0f;
sin_h[i] = 0.0f;
cos_h[i] = 0.0f;
tan_h[i] = 0.0f;
sin_ph[i] = 0.0f;
cos_ph[i] = 0.0f;
}
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&In_d), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&sin_d), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&cos_d), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&tan_d), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&sin_pd), SIZE));
HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&cos_pd), SIZE));
HIP_CHECK(hipMemcpy(In_d, In, SIZE, hipMemcpyHostToDevice));
hipLaunchKernelGGL(kernel_trig, dim3(LEN, 1, 1), dim3(1, 1, 1), 0, 0,
In_d, sin_d, cos_d, tan_d,
sin_pd, cos_pd);
HIP_CHECK(hipMemcpy(sin_h, sin_d, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(cos_h, cos_d, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(tan_h, tan_d, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(sin_ph, sin_pd, SIZE, hipMemcpyDeviceToHost));
HIP_CHECK(hipMemcpy(cos_ph, cos_pd, SIZE, hipMemcpyDeviceToHost));
// Validation
for (int i = 0; i < LEN; i++) {
if (sin_h[i] != sin_ph[i] || cos_h[i] != cos_ph[i] || tan_h[i] *
cos_h[i] != sin_h[i]) {
errors++;
}
}
REQUIRE(errors == 0);
delete[] In;
delete[] sin_h;
delete[] cos_h;
delete[] tan_h;
delete[] sin_ph;
delete[] cos_ph;
HIP_CHECK(hipFree(In_d));
HIP_CHECK(hipFree(sin_d));
HIP_CHECK(hipFree(cos_d));
HIP_CHECK(hipFree(tan_d));
HIP_CHECK(hipFree(sin_pd));
HIP_CHECK(hipFree(cos_pd));
}