From 3447a59895cbd459e83a088a95d563f9faa3bf10 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Mon, 14 Aug 2023 20:52:26 +0530 Subject: [PATCH] SWDEV-380340 - [catch2][dtest] DeviceLib tests migrated from direct to catch2 (#225) Change-Id: Ie2ec1c7dabdfedbe0bd36fd2525df7dc9d9ba2e5 --- catch/unit/deviceLib/CMakeLists.txt | 46 ++ catch/unit/deviceLib/hipBfloat16.cc | 130 ++++ catch/unit/deviceLib/hipComplex.cc | 438 ++++++++++++ catch/unit/deviceLib/hipDeviceMemcpy.cc | 60 ++ catch/unit/deviceLib/hipTestDevice.cc | 736 ++++++++++++++++++++ catch/unit/deviceLib/hipTestDeviceDouble.cc | 630 +++++++++++++++++ catch/unit/deviceLib/hipTestDeviceLimit.cc | 24 + catch/unit/deviceLib/hipTestDeviceSymbol.cc | 124 ++-- catch/unit/deviceLib/hipTestDotFunctions.cc | 54 ++ catch/unit/deviceLib/hipTestFMA.cc | 183 +++++ catch/unit/deviceLib/hipTestHalf.cc | 245 +++++++ catch/unit/deviceLib/hipTestHost.cc | 408 +++++++++++ catch/unit/deviceLib/hipTestIncludeMath.cc | 134 ++++ catch/unit/deviceLib/hipTestNativeHalf.cc | 169 +++++ catch/unit/deviceLib/hipTestNew.cc | 54 ++ catch/unit/deviceLib/hipThreadFence.cc | 76 ++ catch/unit/deviceLib/hipVectorTypes.cc | 204 ++++++ catch/unit/deviceLib/hip_test_make_type.cc | 334 +++++++++ catch/unit/deviceLib/hip_trig.cc | 86 +++ 19 files changed, 4074 insertions(+), 61 deletions(-) create mode 100644 catch/unit/deviceLib/hipBfloat16.cc create mode 100644 catch/unit/deviceLib/hipComplex.cc create mode 100644 catch/unit/deviceLib/hipDeviceMemcpy.cc create mode 100644 catch/unit/deviceLib/hipTestDevice.cc create mode 100644 catch/unit/deviceLib/hipTestDeviceDouble.cc create mode 100644 catch/unit/deviceLib/hipTestDeviceLimit.cc create mode 100644 catch/unit/deviceLib/hipTestDotFunctions.cc create mode 100644 catch/unit/deviceLib/hipTestFMA.cc create mode 100644 catch/unit/deviceLib/hipTestHalf.cc create mode 100644 catch/unit/deviceLib/hipTestHost.cc create mode 100644 catch/unit/deviceLib/hipTestIncludeMath.cc create mode 100644 catch/unit/deviceLib/hipTestNativeHalf.cc create mode 100644 catch/unit/deviceLib/hipTestNew.cc create mode 100644 catch/unit/deviceLib/hipThreadFence.cc create mode 100644 catch/unit/deviceLib/hipVectorTypes.cc create mode 100644 catch/unit/deviceLib/hip_test_make_type.cc create mode 100644 catch/unit/deviceLib/hip_trig.cc diff --git a/catch/unit/deviceLib/CMakeLists.txt b/catch/unit/deviceLib/CMakeLists.txt index b2718f39bb..5317d42ca5 100644 --- a/catch/unit/deviceLib/CMakeLists.txt +++ b/catch/unit/deviceLib/CMakeLists.txt @@ -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 diff --git a/catch/unit/deviceLib/hipBfloat16.cc b/catch/unit/deviceLib/hipBfloat16.cc new file mode 100644 index 0000000000..f87ab48ddf --- /dev/null +++ b/catch/unit/deviceLib/hipBfloat16.cc @@ -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 +#include +#include +#include +#include + +#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 gen(min, max); + return gen(rng); +} +__host__ __device__ bool testRelativeAccuracy(float a, hip_bfloat16 b) { + float c = static_cast(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(bf_a); + float fd = static_cast(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)); +} diff --git a/catch/unit/deviceLib/hipComplex.cc b/catch/unit/deviceLib/hipComplex.cc new file mode 100644 index 0000000000..04a4ecc737 --- /dev/null +++ b/catch/unit/deviceLib/hipComplex.cc @@ -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 +#include +#include +#include +#include + +#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 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(&Ad), size2)); + HIPCHECK(hipMalloc(reinterpret_cast(&Bd), size2)); + HIPCHECK(hipMalloc(reinterpret_cast(&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 +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(&Ad), size)); + HIPCHECK(hipMalloc(reinterpret_cast(&Bd), size)); + HIPCHECK(hipMalloc(reinterpret_cast(&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 +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(&Ad), size1)); + HIPCHECK(hipMalloc(reinterpret_cast(&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() && + test_makehipComplex_dev() && + test_makehipComplex_dev() && + test_makehipComplex_dev() && + test_complexMathFunc1_dev(COMPLEX_ADD) && + test_complexMathFunc1_dev(COMPLEX_ADD) + && test_complexMathFunc1_dev(COMPLEX_SUB) + && test_complexMathFunc1_dev + (COMPLEX_SUB) && test_complexMathFunc1_dev(COMPLEX_MUL) && test_complexMathFunc1_dev(COMPLEX_MUL) && test_complexMathFunc1_dev(COMPLEX_DIV) && test_complexMathFunc1_dev(COMPLEX_DIV) && test_complexMathFunc1_dev(COMPLEX_CONJ) && test_complexMathFunc1_dev< + hipDoubleComplex, double>(COMPLEX_CONJ) && test_complexMathFunc2_dev + (COMPLEX_REAL) && test_complexMathFunc2_dev + (COMPLEX_REAL) && test_complexMathFunc2_dev + (COMPLEX_IMAG) && test_complexMathFunc2_dev + (COMPLEX_IMAG) && test_complexMathFunc2_dev + (COMPLEX_SQABS) && test_complexMathFunc2_dev + (COMPLEX_SQABS) && test_complexMathFunc2_dev + (COMPLEX_ABS) && test_complexMathFunc2_dev + (COMPLEX_ABS) &&test_allcomplexMathFunc_host(); + REQUIRE(TestPassed == true); +} diff --git a/catch/unit/deviceLib/hipDeviceMemcpy.cc b/catch/unit/deviceLib/hipDeviceMemcpy.cc new file mode 100644 index 0000000000..df881aa454 --- /dev/null +++ b/catch/unit/deviceLib/hipDeviceMemcpy.cc @@ -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 + +#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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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; +} diff --git a/catch/unit/deviceLib/hipTestDevice.cc b/catch/unit/deviceLib/hipTestDevice.cc new file mode 100644 index 0000000000..606ad7dfb4 --- /dev/null +++ b/catch/unit/deviceLib/hipTestDevice.cc @@ -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 +#include + +#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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Cd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Cd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Dd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Cd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Cd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Dd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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); +} diff --git a/catch/unit/deviceLib/hipTestDeviceDouble.cc b/catch/unit/deviceLib/hipTestDeviceDouble.cc new file mode 100644 index 0000000000..865d54e9bd --- /dev/null +++ b/catch/unit/deviceLib/hipTestDeviceDouble.cc @@ -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 +#include + +#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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Cd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Cd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Dd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Cd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Bd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Cd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&Dd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&Ad), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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); +} diff --git a/catch/unit/deviceLib/hipTestDeviceLimit.cc b/catch/unit/deviceLib/hipTestDeviceLimit.cc new file mode 100644 index 0000000000..d085efde34 --- /dev/null +++ b/catch/unit/deviceLib/hipTestDeviceLimit.cc @@ -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 +TEST_CASE("Unit_hipTestDeviceLimit_Basic") { + size_t heap; + HIP_CHECK(hipDeviceGetLimit(&heap, hipLimitMallocHeapSize)); + REQUIRE(heap != NULL); +} diff --git a/catch/unit/deviceLib/hipTestDeviceSymbol.cc b/catch/unit/deviceLib/hipTestDeviceSymbol.cc index c7bfa59283..3a2a103906 100644 --- a/catch/unit/deviceLib/hipTestDeviceSymbol.cc +++ b/catch/unit/deviceLib/hipTestDeviceSymbol.cc @@ -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(&Ad), SIZE)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&Am), SIZE)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&Cm), SIZE)); for (size_t i = 0; i < NUM; i++) { A[i] = -1 * static_cast(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(&symbolAddress), + HIP_SYMBOL(globalConst))); + HIP_CHECK(hipMalloc(reinterpret_cast(&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(&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(&symbolVarAddress), HIP_SYMBOL(statConstVar))); + hipGetSymbolAddress(reinterpret_cast(&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(&addr), HIP_SYMBOL(notADeviceSymbol)), - hipErrorInvalidSymbol); + HIP_CHECK_ERROR(hipGetSymbolAddress(reinterpret_cast(&addr), + HIP_SYMBOL(notADeviceSymbol)), hipErrorInvalidSymbol); } SECTION("Nullptr symbol") { int* addr{nullptr}; - HIP_CHECK_ERROR(hipGetSymbolAddress(reinterpret_cast(&addr), nullptr), - hipErrorInvalidSymbol); + HIP_CHECK_ERROR(hipGetSymbolAddress(reinterpret_cast(&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") { diff --git a/catch/unit/deviceLib/hipTestDotFunctions.cc b/catch/unit/deviceLib/hipTestDotFunctions.cc new file mode 100644 index 0000000000..43f0492da5 --- /dev/null +++ b/catch/unit/deviceLib/hipTestDotFunctions.cc @@ -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 +#include +#include + +__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); +} diff --git a/catch/unit/deviceLib/hipTestFMA.cc b/catch/unit/deviceLib/hipTestFMA.cc new file mode 100644 index 0000000000..3e496f4d17 --- /dev/null +++ b/catch/unit/deviceLib/hipTestFMA.cc @@ -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 +#include + +#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(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(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(f), + static_cast(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(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(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(f), 3) == doubleResult); + Check(fma(d, static_cast(f), static_cast(3)) == doubleResult); + Check(fma(d, static_cast(f), (unsigned char)3) == doubleResult); + Check(fma(d, static_cast(f), (int32_t)3) == doubleResult); + Check(fma(d, static_cast(f), (uint32_t)3) == doubleResult); + Check(fma(d, static_cast(f), static_cast(3)) == doubleResult); + Check(fma(d, static_cast(f), (unsigned int)3) == doubleResult); + Check(fma(d, static_cast(f), (int64_t)3) == doubleResult); + Check(fma(d, static_cast(f), (int64_t)3) == doubleResult); + Check(fma(d, static_cast(f), true) == + fma(static_cast(f), static_cast(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(&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 . + if (sizeof(decltype(fma(h, h, 3))) == 8) { + Check(fma(h, h, 3) == doubleResult); + Check(fma(h, h, static_cast(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(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(h), + static_cast(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::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(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(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(&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(); + } +} diff --git a/catch/unit/deviceLib/hipTestHalf.cc b/catch/unit/deviceLib/hipTestHalf.cc new file mode 100644 index 0000000000..936625ef56 --- /dev/null +++ b/catch/unit/deviceLib/hipTestHalf.cc @@ -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 +#include + +__device__ void test_convert() { + __half x; + float y = static_cast(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(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(&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 (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(&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 (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(&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)); +} diff --git a/catch/unit/deviceLib/hipTestHost.cc b/catch/unit/deviceLib/hipTestHost.cc new file mode 100644 index 0000000000..e1e77ff9e3 --- /dev/null +++ b/catch/unit/deviceLib/hipTestHost.cc @@ -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 + +#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(&devicePtr), sizeof(bool))); + HIP_CHECK(hipMalloc(reinterpret_cast(&Afd), sizeof(float)*N)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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; +} diff --git a/catch/unit/deviceLib/hipTestIncludeMath.cc b/catch/unit/deviceLib/hipTestIncludeMath.cc new file mode 100644 index 0000000000..2bb78b03b2 --- /dev/null +++ b/catch/unit/deviceLib/hipTestIncludeMath.cc @@ -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 + #include + +// 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(-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(FloatMathPreciseKernel), + dim3(1, 1, 1), + dim3(1, 1, 1), 0, 0, 0); + REQUIRE(err == hipSuccess); +} diff --git a/catch/unit/deviceLib/hipTestNativeHalf.cc b/catch/unit/deviceLib/hipTestNativeHalf.cc new file mode 100644 index 0000000000..521b09bca0 --- /dev/null +++ b/catch/unit/deviceLib/hipTestNativeHalf.cc @@ -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 +#include +#include + +__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(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)); +} diff --git a/catch/unit/deviceLib/hipTestNew.cc b/catch/unit/deviceLib/hipTestNew.cc new file mode 100644 index 0000000000..35f0efb195 --- /dev/null +++ b/catch/unit/deviceLib/hipTestNew.cc @@ -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 + +#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(&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; +} diff --git a/catch/unit/deviceLib/hipThreadFence.cc b/catch/unit/deviceLib/hipThreadFence.cc new file mode 100644 index 0000000000..4272b6aa2d --- /dev/null +++ b/catch/unit/deviceLib/hipThreadFence.cc @@ -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 + +#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(&In1d), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&In2d), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&In3d), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&In4d), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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)); +} diff --git a/catch/unit/deviceLib/hipVectorTypes.cc b/catch/unit/deviceLib/hipVectorTypes.cc new file mode 100644 index 0000000000..09d19b0397 --- /dev/null +++ b/catch/unit/deviceLib/hipVectorTypes.cc @@ -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 +#include + +#include +#include +#include + +using namespace std; // NOLINT + +template().x)>{}>* = nullptr> +bool integer_unary_tests(V&, V&) { + return true; +} + +template().x)>{}>* = nullptr> +bool integer_binary_tests(V&, V&, V&...) { + return true; +} + +template().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().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 +bool constructor_tests() { + if (is_constructible{} && + is_constructible{} && + is_constructible{} && + is_constructible{} && + is_constructible{} && + is_constructible{} && + is_constructible{} && + is_constructible{} && + is_constructible{} && + is_constructible{} && + is_constructible{} && + is_constructible{}) { + return true; + } +} + +template +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()) 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* = nullptr> +bool TestVectorTypes() { + return true; +} + +template +bool TestVectorTypes() { + if (!TestVectorType()) return false; + return TestVectorTypes(); +} + +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); +} diff --git a/catch/unit/deviceLib/hip_test_make_type.cc b/catch/unit/deviceLib/hip_test_make_type.cc new file mode 100644 index 0000000000..32e889daee --- /dev/null +++ b/catch/unit/deviceLib/hip_test_make_type.cc @@ -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 +#include + +#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 +bool dataTypesRunChar1() { + T* hostA; + T* hostB; + T* hostC; + + T* deviceA; + T* deviceB; + T* deviceC; + + int i; + int errors; + + hostA = reinterpret_cast(malloc(NUM * sizeof(T))); + hostB = reinterpret_cast(malloc(NUM * sizeof(T))); + hostC = reinterpret_cast(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(&deviceA), NUM * sizeof(T))); + HIP_CHECK(hipMalloc(reinterpret_cast(&deviceB), NUM * sizeof(T))); + HIP_CHECK(hipMalloc(reinterpret_cast(&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 +bool dataTypesRunChar2() { + T* hostA; + T* hostB; + T* hostC; + + T* deviceA; + T* deviceB; + T* deviceC; + + int i; + int errors; + + hostA = reinterpret_cast(malloc(NUM * sizeof(T))); + hostB = reinterpret_cast(malloc(NUM * sizeof(T))); + hostC = reinterpret_cast(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(&deviceA), NUM * sizeof(T))); + HIP_CHECK(hipMalloc(reinterpret_cast(&deviceB), NUM * sizeof(T))); + HIP_CHECK(hipMalloc(reinterpret_cast(&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 +bool dataTypesRunChar3() { + T* hostA; + T* hostB; + T* hostC; + + T* deviceA; + T* deviceB; + T* deviceC; + + int i; + int errors; + + hostA = reinterpret_cast(malloc(NUM * sizeof(T))); + hostB = reinterpret_cast(malloc(NUM * sizeof(T))); + hostC = reinterpret_cast(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(&deviceA), NUM * sizeof(T))); + HIP_CHECK(hipMalloc(reinterpret_cast(&deviceB), NUM * sizeof(T))); + HIP_CHECK(hipMalloc(reinterpret_cast(&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 +bool dataTypesRunChar4() { + char4* hostA; + char4* hostB; + char4* hostC; + + char4* deviceA; + char4* deviceB; + char4* deviceC; + + int i; + int errors; + + hostA = reinterpret_cast(malloc(NUM * sizeof(T))); + hostB = reinterpret_cast(malloc(NUM * sizeof(T))); + hostC = reinterpret_cast(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(&deviceA), NUM * sizeof(T))); + HIP_CHECK(hipMalloc(reinterpret_cast(&deviceB), NUM * sizeof(T))); + HIP_CHECK(hipMalloc(reinterpret_cast(&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() && dataTypesRunChar2() && + dataTypesRunChar3() && dataTypesRunChar4(); + + REQUIRE(errors == true); +} diff --git a/catch/unit/deviceLib/hip_trig.cc b/catch/unit/deviceLib/hip_trig.cc new file mode 100644 index 0000000000..a264d69783 --- /dev/null +++ b/catch/unit/deviceLib/hip_trig.cc @@ -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 +#include +#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(&In_d), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&sin_d), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&cos_d), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&tan_d), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&sin_pd), SIZE)); + HIP_CHECK(hipMalloc(reinterpret_cast(&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)); +}