From cf8ec29fe217b4b82bf9bc601600133142e07823 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 5 Aug 2019 16:05:39 +0300 Subject: [PATCH 1/9] [HIP][doc] Populate README.md with the supported CUDA APIs links --- README.md | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/README.md b/README.md index 200b2f7474..145f950fef 100644 --- a/README.md +++ b/README.md @@ -37,6 +37,15 @@ HIP releases are typically of two types. The tag naming convention is different - [HIP Debugging](docs/markdown/hip_debugging.md) - [HIP Terminology](docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenCL) - [hipify-clang](hipify-clang/README.md) +- Supported CUDA APIs: + * [Runtime API](docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md) + * [Driver API](docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md) + * [cuComplex API](docs/markdown/cuComplex_API_supported_by_HIP.md) + * [cuBLAS](docs/markdown/CUBLAS_API_supported_by_HIP.md) + * [cuRAND](docs/markdown/CURAND_API_supported_by_HIP.md) + * [cuDNN](docs/markdown/CUDNN_API_supported_by_HIP.md) + * [cuFFT](docs/markdown/CUFFT_API_supported_by_HIP.md) + * [cuSPARSE](docs/markdown/CUSPARSE_API_supported_by_HIP.md) - [Developer/CONTRIBUTING Info](CONTRIBUTING.md) - [Release Notes](RELEASE.md) From 59bda1497996731ce3336b3f422128f50b1a4796 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 7 Aug 2019 04:00:32 +0530 Subject: [PATCH 2/9] Enable temporarily disabled device properties on HIP/VDI --- include/hip/hip_runtime_api.h | 2 -- 1 file changed, 2 deletions(-) diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 1c8e971545..1f7183acf0 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -117,13 +117,11 @@ typedef struct hipDeviceProp_t { int integrated; ///< APU vs dGPU int cooperativeLaunch; ///< HIP device supports cooperative launch int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple devices -#if !__HIP_VDI__ // Temporarily disable the following five new fields for HIP/VDI runtime int maxTexture1D; ///< Maximum number of elements in 1D images int maxTexture2D[2]; ///< Maximum dimensions (width, height) of 2D images, in image elements int maxTexture3D[3]; ///< Maximum dimensions (width, height, depth) of 3D images, in image elements unsigned int* hdpMemFlushCntl; ///< Addres of HDP_MEM_COHERENCY_FLUSH_CNTL register unsigned int* hdpRegFlushCntl; ///< Addres of HDP_REG_COHERENCY_FLUSH_CNTL register -#endif } hipDeviceProp_t; From 435badcb0edf3f1e02323a5718e31739a8d093e5 Mon Sep 17 00:00:00 2001 From: lthakur Date: Wed, 7 Aug 2019 15:54:46 +0530 Subject: [PATCH 3/9] Fixed compilation errors being observed on windows --- .../deviceLib/hipDoublePrecisionMathHost.cpp | 300 +++++++++--------- .../deviceLib/hipSinglePrecisionMathHost.cpp | 296 ++++++++--------- 2 files changed, 310 insertions(+), 286 deletions(-) mode change 100644 => 100755 tests/src/deviceLib/hipDoublePrecisionMathHost.cpp mode change 100644 => 100755 tests/src/deviceLib/hipSinglePrecisionMathHost.cpp diff --git a/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp b/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp old mode 100644 new mode 100755 index a3283622ff..643b300f2f --- a/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp +++ b/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp @@ -1,142 +1,158 @@ -/* -Copyright (c) 2015-2016 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. -*/ - -/* HIT_START - * BUILD: %t %s ../test_common.cpp - * TEST: %t - * HIT_END - */ - -#include -//#include -#include "test_common.h" -#include - -#pragma GCC diagnostic ignored "-Wall" -#pragma clang diagnostic ignored "-Wunused-variable" - -__host__ void double_precision_math_functions() { - int iX; - double fX, fY; - - acos(1.0); - acosh(1.0); - asin(0.0); - asinh(0.0); - atan(0.0); - atan2(0.0, 1.0); - atanh(0.0); - cbrt(0.0); - ceil(0.0); - copysign(1.0, -2.0); - cos(0.0); - cosh(0.0); - // cospi(0.0); - // cyl_bessel_i0(0.0); - // cyl_bessel_i1(0.0); - erf(0.0); - erfc(0.0); - // erfcinv(2.0); - // erfcx(0.0); - // erfinv(1.0); - exp(0.0); - exp10(0.0); - exp2(0.0); - expm1(0.0); - fabs(1.0); - fdim(1.0, 0.0); - floor(0.0); - fma(1.0, 2.0, 3.0); - fmax(0.0, 0.0); - fmin(0.0, 0.0); - fmod(0.0, 1.0); - frexp(0.0, &iX); - hypot(1.0, 0.0); - ilogb(1.0); - std::isfinite(0.0); - std::isinf(0.0); - std::isnan(0.0); - j0(0.0); - j1(0.0); - jn(-1.0, 1.0); - ldexp(0.0, 0); - // lgamma(1.0); - llrint(0.0); - llround(0.0); - log(1.0); - log10(1.0); - log1p(-1.0); - log2(1.0); - logb(1.0); - lrint(0.0); - lround(0.0); - modf(0.0, &fX); - nan("1"); - nearbyint(0.0); - // nextafter(0.0); - fX = 1.0; // norm(1, &fX); -#if defined(__HIP_PLATFORM_HCC__) - // norm3d(1.0, 0.0, 0.0); - // norm4d(1.0, 0.0, 0.0, 0.0); -#endif - // normcdf(0.0); - // normcdfinv(1.0); - pow(1.0, 0.0); - // rcbrt(1.0); - - remainder(2.0, 1.0); - remquo(1.0, 2.0, &iX); -#if defined(__HIP_PLATFORM_HCC__) - // rhypot(0.0, 1.0); -#endif - rint(1.0); -#if defined(__HIP_PLATFORM_HCC__) - fX = 1.0; // rnorm(1, &fX); - // rnorm3d(0.0, 0.0, 1.0); - // rnorm4d(0.0, 0.0, 0.0, 1.0); -#endif - round(0.0); - // rsqrt(1.0); - scalbln(0.0, 1); - scalbn(0.0, 1); - std::signbit(1.0); - sin(0.0); - sincos(0.0, &fX, &fY); - // sincospi(0.0, &fX, &fY); - sinh(0.0); - // sinpi(0.0); - sqrt(0.0); - tan(0.0); - tanh(0.0); - tgamma(2.0); - trunc(0.0); - y0(1.0); - y1(1.0); - yn(1, 1.0); -} - -static void compileOnHost() { double_precision_math_functions(); } - -int main() { - compileOnHost(); - passed(); -} +/* +Copyright (c) 2015-2016 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. +*/ + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include +//#include +#include "test_common.h" +#include + +#pragma GCC diagnostic ignored "-Wall" +#pragma clang diagnostic ignored "-Wunused-variable" + +__host__ void double_precision_math_functions() { + int iX; + double fX, fY; + + acos(1.0); + acosh(1.0); + asin(0.0); + asinh(0.0); + atan(0.0); + atan2(0.0, 1.0); + atanh(0.0); + cbrt(0.0); + ceil(0.0); + copysign(1.0, -2.0); + cos(0.0); + cosh(0.0); + // cospi(0.0); + // cyl_bessel_i0(0.0); + // cyl_bessel_i1(0.0); + erf(0.0); + erfc(0.0); + // erfcinv(2.0); + // erfcx(0.0); + // erfinv(1.0); + exp(0.0); + #ifdef __unix__ + exp10(0.0); + #endif + exp2(0.0); + expm1(0.0); + fabs(1.0); + fdim(1.0, 0.0); + floor(0.0); + fma(1.0, 2.0, 3.0); + fmax(0.0, 0.0); + fmin(0.0, 0.0); + fmod(0.0, 1.0); + frexp(0.0, &iX); + hypot(1.0, 0.0); + ilogb(1.0); + std::isfinite(0.0); + std::isinf(0.0); + std::isnan(0.0); + #ifdef __unix__ + j0(1.0); + j1(1.0); + jn(-1.0, 1.0); + #elif _WIN64 + _j0(0.0); + _j1(0.0); + _jn(-1.0, 1.0); + #endif + ldexp(0.0, 0); + // lgamma(1.0); + llrint(0.0); + llround(0.0); + log(1.0); + log10(1.0); + log1p(-1.0); + log2(1.0); + logb(1.0); + lrint(0.0); + lround(0.0); + modf(0.0, &fX); + nan("1"); + nearbyint(0.0); + // nextafter(0.0); + fX = 1.0; // norm(1, &fX); +#if defined(__HIP_PLATFORM_HCC__) + // norm3d(1.0, 0.0, 0.0); + // norm4d(1.0, 0.0, 0.0, 0.0); +#endif + // normcdf(0.0); + // normcdfinv(1.0); + pow(1.0, 0.0); + // rcbrt(1.0); + + remainder(2.0, 1.0); + remquo(1.0, 2.0, &iX); +#if defined(__HIP_PLATFORM_HCC__) + // rhypot(0.0, 1.0); +#endif + rint(1.0); +#if defined(__HIP_PLATFORM_HCC__) + fX = 1.0; // rnorm(1, &fX); + // rnorm3d(0.0, 0.0, 1.0); + // rnorm4d(0.0, 0.0, 0.0, 1.0); +#endif + round(0.0); + // rsqrt(1.0); + scalbln(0.0, 1); + scalbn(0.0, 1); + std::signbit(1.0); + sin(0.0); + #ifdef __unix__ + sincos(0.0, &fX, &fY); + #endif + // sincospi(0.0, &fX, &fY); + sinh(0.0); + // sinpi(0.0); + sqrt(0.0); + tan(0.0); + tanh(0.0); + tgamma(2.0); + trunc(0.0); + #ifdef __unix__ + y0(1.0); + y1(1.0); + yn(1, 1.0); + #elif _WIN64 + _y0(1.0); + _y1(1.0); + _yn(1, 1.0); + #endif +} + +static void compileOnHost() { double_precision_math_functions(); } + +int main() { + compileOnHost(); + passed(); +} diff --git a/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp b/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp old mode 100644 new mode 100755 index 1dd5737f86..cd827bc318 --- a/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp +++ b/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp @@ -1,144 +1,152 @@ -/* -Copyright (c) 2015-2016 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. -*/ - -/* HIT_START - * BUILD: %t %s ../test_common.cpp - * TEST: %t - * HIT_END - */ - -#include -//#include -#include "test_common.h" -#include - -#pragma GCC diagnostic ignored "-Wall" -#pragma clang diagnostic ignored "-Wunused-variable" - -__host__ void single_precision_math_functions() { - 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); - ceilf(0.0f); - 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); - fabsf(1.0f); - fdimf(1.0f, 0.0f); -#if defined(__HIP_PLATFORM_HCC__) - // fdividef(0.0f, 1.0f); -#endif - floorf(0.0f); - fmaf(1.0f, 2.0f, 3.0f); - fmaxf(0.0f, 0.0f); - fminf(0.0f, 0.0f); - fmodf(0.0f, 1.0f); - frexpf(0.0f, &iX); - hypotf(1.0f, 0.0f); - ilogbf(1.0f); - std::isfinite(0.0f); - std::isinf(0.0f); - std::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); - nanf("1"); - nearbyintf(0.0f); - // nextafterf(0.0f); -#if defined(__HIP_PLATFORM_HCC__) - // norm3df(1.0f, 0.0f, 0.0f); - // norm4df(1.0f, 0.0f, 0.0f, 0.0f); -#endif - // 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); -#if defined(__HIP_PLATFORM_HCC__) - // rhypotf(0.0f, 1.0f); -#endif - rintf(1.0f); -#if defined(__HIP_PLATFORM_HCC__) - // rnorm3df(0.0f, 0.0f, 1.0f); - // rnorm4df(0.0f, 0.0f, 0.0f, 1.0f); - fX = 1.0f; // rnormf(1, &fX); -#endif - roundf(0.0f); - /// rsqrtf(1.0f); - scalblnf(0.0f, 1); - scalbnf(0.0f, 1); - std::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); - truncf(0.0f); - y0f(1.0f); - y1f(1.0f); - ynf(1, 1.0f); -} - -static void compileOnHost() { single_precision_math_functions(); } - -int main() { - compileOnHost(); - passed(); -} +/* +Copyright (c) 2015-2016 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. +*/ + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include +//#include +#include "test_common.h" +#include + +#pragma GCC diagnostic ignored "-Wall" +#pragma clang diagnostic ignored "-Wunused-variable" + +__host__ void single_precision_math_functions() { + 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); + ceilf(0.0f); + 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); + #ifdef __unix__ + exp10f(0.0f); + #endif + exp2f(0.0f); + expf(0.0f); + expm1f(0.0f); + fabsf(1.0f); + fdimf(1.0f, 0.0f); +#if defined(__HIP_PLATFORM_HCC__) + // fdividef(0.0f, 1.0f); +#endif + floorf(0.0f); + fmaf(1.0f, 2.0f, 3.0f); + fmaxf(0.0f, 0.0f); + fminf(0.0f, 0.0f); + fmodf(0.0f, 1.0f); + frexpf(0.0f, &iX); + hypotf(1.0f, 0.0f); + ilogbf(1.0f); + std::isfinite(0.0f); + std::isinf(0.0f); + std::isnan(0.0f); + #ifdef __unix__ + j0f(0.0f); + j1f(0.0f); + jnf(-1.0f, 1.0f); + #endif + 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); + nanf("1"); + nearbyintf(0.0f); + // nextafterf(0.0f); +#if defined(__HIP_PLATFORM_HCC__) + // norm3df(1.0f, 0.0f, 0.0f); + // norm4df(1.0f, 0.0f, 0.0f, 0.0f); +#endif + // 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); +#if defined(__HIP_PLATFORM_HCC__) + // rhypotf(0.0f, 1.0f); +#endif + rintf(1.0f); +#if defined(__HIP_PLATFORM_HCC__) + // rnorm3df(0.0f, 0.0f, 1.0f); + // rnorm4df(0.0f, 0.0f, 0.0f, 1.0f); + fX = 1.0f; // rnormf(1, &fX); +#endif + roundf(0.0f); + /// rsqrtf(1.0f); + scalblnf(0.0f, 1); + scalbnf(0.0f, 1); + std::signbit(1.0f); + #ifdef __unix__ + sincosf(0.0f, &fX, &fY); + #endif + // 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); + truncf(0.0f); + #ifdef __unix__ + y0f(1.0f); + y1f(1.0f); + ynf(1, 1.0f); + #endif +} + +static void compileOnHost() { single_precision_math_functions(); } + +int main() { + compileOnHost(); + passed(); +} From a90310162d06167f6a0fb817be6c01e1e0d1794f Mon Sep 17 00:00:00 2001 From: David Rohr Date: Wed, 7 Aug 2019 19:11:12 +0200 Subject: [PATCH 4/9] Treat .cxx files like .cpp files --- bin/hipcc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/bin/hipcc b/bin/hipcc index 1eda99b936..6c8dc20597 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -624,7 +624,7 @@ foreach $arg (@ARGV) #print "O: <$arg>\n"; } else { # input files and libraries - if (($arg =~ /\.cpp$/) or ($arg =~ /\.c$/) or ($arg =~ /\.cc$/) ) { + if (($arg =~ /\.cpp$/) or ($arg =~ /\.cxx$/) or ($arg =~ /\.c$/) or ($arg =~ /\.cc$/) ) { $hasC = 1; $needCXXFLAGS = 1; if ($HIP_PLATFORM eq 'clang') { From 80a3897f9ee9d907a128561b7e4c76c7417c3f12 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 7 Aug 2019 16:59:17 -0400 Subject: [PATCH 5/9] Fix HIP-VDI using incorrect include files When HIP-VDI depends on an llvm-amdgpu package, the include clang directory is not properly assessed. --- bin/hipcc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/bin/hipcc b/bin/hipcc index 1eda99b936..e51d9dde09 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -118,10 +118,10 @@ if ($HIP_RUNTIME eq "VDI" and !defined $HIP_VDI_HOME) { if (defined $HIP_VDI_HOME) { if (!defined $HIP_CLANG_PATH and (-e "$HIP_VDI_HOME/bin/clang" or -e "$HIP_VDI_HOME/bin/clang.exe")) { $HIP_CLANG_PATH = "$HIP_VDI_HOME/bin"; - } - # With HIP_VDI_HOME defined, assume the installation of clang components, including headers. - if (!defined $HIP_CLANG_INCLUDE_PATH) { - $HIP_CLANG_INCLUDE_PATH = "$HIP_VDI_HOME/include/clang"; + # With HIP_VDI_HOME defined, assume the installation of clang components, including headers. + if (!defined $HIP_CLANG_INCLUDE_PATH) { + $HIP_CLANG_INCLUDE_PATH = "$HIP_VDI_HOME/include/clang"; + } } if (!defined $DEVICE_LIB_PATH and -e "$HIP_VDI_HOME/lib/bitcode") { $DEVICE_LIB_PATH = "$HIP_VDI_HOME/lib/bitcode"; From d18160920e51d5f589f4a0b4700525724e78030c Mon Sep 17 00:00:00 2001 From: lthakur Date: Thu, 8 Aug 2019 11:26:26 +0530 Subject: [PATCH 6/9] Resubmitting the fix so as to address the changed format in the last submission --- .../deviceLib/hipDoublePrecisionMathHost.cpp | 316 +++++++++--------- .../deviceLib/hipSinglePrecisionMathHost.cpp | 304 ++++++++--------- 2 files changed, 310 insertions(+), 310 deletions(-) mode change 100755 => 100644 tests/src/deviceLib/hipDoublePrecisionMathHost.cpp mode change 100755 => 100644 tests/src/deviceLib/hipSinglePrecisionMathHost.cpp diff --git a/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp b/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp old mode 100755 new mode 100644 index 643b300f2f..295d2f5c95 --- a/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp +++ b/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp @@ -1,158 +1,158 @@ -/* -Copyright (c) 2015-2016 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. -*/ - -/* HIT_START - * BUILD: %t %s ../test_common.cpp - * TEST: %t - * HIT_END - */ - -#include -//#include -#include "test_common.h" -#include - -#pragma GCC diagnostic ignored "-Wall" -#pragma clang diagnostic ignored "-Wunused-variable" - -__host__ void double_precision_math_functions() { - int iX; - double fX, fY; - - acos(1.0); - acosh(1.0); - asin(0.0); - asinh(0.0); - atan(0.0); - atan2(0.0, 1.0); - atanh(0.0); - cbrt(0.0); - ceil(0.0); - copysign(1.0, -2.0); - cos(0.0); - cosh(0.0); - // cospi(0.0); - // cyl_bessel_i0(0.0); - // cyl_bessel_i1(0.0); - erf(0.0); - erfc(0.0); - // erfcinv(2.0); - // erfcx(0.0); - // erfinv(1.0); - exp(0.0); - #ifdef __unix__ - exp10(0.0); - #endif - exp2(0.0); - expm1(0.0); - fabs(1.0); - fdim(1.0, 0.0); - floor(0.0); - fma(1.0, 2.0, 3.0); - fmax(0.0, 0.0); - fmin(0.0, 0.0); - fmod(0.0, 1.0); - frexp(0.0, &iX); - hypot(1.0, 0.0); - ilogb(1.0); - std::isfinite(0.0); - std::isinf(0.0); - std::isnan(0.0); - #ifdef __unix__ - j0(1.0); - j1(1.0); - jn(-1.0, 1.0); - #elif _WIN64 - _j0(0.0); - _j1(0.0); - _jn(-1.0, 1.0); - #endif - ldexp(0.0, 0); - // lgamma(1.0); - llrint(0.0); - llround(0.0); - log(1.0); - log10(1.0); - log1p(-1.0); - log2(1.0); - logb(1.0); - lrint(0.0); - lround(0.0); - modf(0.0, &fX); - nan("1"); - nearbyint(0.0); - // nextafter(0.0); - fX = 1.0; // norm(1, &fX); -#if defined(__HIP_PLATFORM_HCC__) - // norm3d(1.0, 0.0, 0.0); - // norm4d(1.0, 0.0, 0.0, 0.0); -#endif - // normcdf(0.0); - // normcdfinv(1.0); - pow(1.0, 0.0); - // rcbrt(1.0); - - remainder(2.0, 1.0); - remquo(1.0, 2.0, &iX); -#if defined(__HIP_PLATFORM_HCC__) - // rhypot(0.0, 1.0); -#endif - rint(1.0); -#if defined(__HIP_PLATFORM_HCC__) - fX = 1.0; // rnorm(1, &fX); - // rnorm3d(0.0, 0.0, 1.0); - // rnorm4d(0.0, 0.0, 0.0, 1.0); -#endif - round(0.0); - // rsqrt(1.0); - scalbln(0.0, 1); - scalbn(0.0, 1); - std::signbit(1.0); - sin(0.0); - #ifdef __unix__ - sincos(0.0, &fX, &fY); - #endif - // sincospi(0.0, &fX, &fY); - sinh(0.0); - // sinpi(0.0); - sqrt(0.0); - tan(0.0); - tanh(0.0); - tgamma(2.0); - trunc(0.0); - #ifdef __unix__ - y0(1.0); - y1(1.0); - yn(1, 1.0); - #elif _WIN64 - _y0(1.0); - _y1(1.0); - _yn(1, 1.0); - #endif -} - -static void compileOnHost() { double_precision_math_functions(); } - -int main() { - compileOnHost(); - passed(); -} +/* +Copyright (c) 2015-2016 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. +*/ + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include +//#include +#include "test_common.h" +#include + +#pragma GCC diagnostic ignored "-Wall" +#pragma clang diagnostic ignored "-Wunused-variable" + +__host__ void double_precision_math_functions() { + int iX; + double fX, fY; + + acos(1.0); + acosh(1.0); + asin(0.0); + asinh(0.0); + atan(0.0); + atan2(0.0, 1.0); + atanh(0.0); + cbrt(0.0); + ceil(0.0); + copysign(1.0, -2.0); + cos(0.0); + cosh(0.0); + // cospi(0.0); + // cyl_bessel_i0(0.0); + // cyl_bessel_i1(0.0); + erf(0.0); + erfc(0.0); + // erfcinv(2.0); + // erfcx(0.0); + // erfinv(1.0); + exp(0.0); + #ifdef __unix__ + exp10(0.0); + #endif + exp2(0.0); + expm1(0.0); + fabs(1.0); + fdim(1.0, 0.0); + floor(0.0); + fma(1.0, 2.0, 3.0); + fmax(0.0, 0.0); + fmin(0.0, 0.0); + fmod(0.0, 1.0); + frexp(0.0, &iX); + hypot(1.0, 0.0); + ilogb(1.0); + std::isfinite(0.0); + std::isinf(0.0); + std::isnan(0.0); + #ifdef __unix__ + j0(0.0); + j1(0.0); + jn(-1.0, 1.0); + #elif _WIN64 + _j0(0.0); + _j1(0.0); + _jn(-1.0, 1.0); + #endif + ldexp(0.0, 0); + // lgamma(1.0); + llrint(0.0); + llround(0.0); + log(1.0); + log10(1.0); + log1p(-1.0); + log2(1.0); + logb(1.0); + lrint(0.0); + lround(0.0); + modf(0.0, &fX); + nan("1"); + nearbyint(0.0); + // nextafter(0.0); + fX = 1.0; // norm(1, &fX); +#if defined(__HIP_PLATFORM_HCC__) + // norm3d(1.0, 0.0, 0.0); + // norm4d(1.0, 0.0, 0.0, 0.0); +#endif + // normcdf(0.0); + // normcdfinv(1.0); + pow(1.0, 0.0); + // rcbrt(1.0); + + remainder(2.0, 1.0); + remquo(1.0, 2.0, &iX); +#if defined(__HIP_PLATFORM_HCC__) + // rhypot(0.0, 1.0); +#endif + rint(1.0); +#if defined(__HIP_PLATFORM_HCC__) + fX = 1.0; // rnorm(1, &fX); + // rnorm3d(0.0, 0.0, 1.0); + // rnorm4d(0.0, 0.0, 0.0, 1.0); +#endif + round(0.0); + // rsqrt(1.0); + scalbln(0.0, 1); + scalbn(0.0, 1); + std::signbit(1.0); + sin(0.0); + #ifdef _unix__ + sincos(0.0, &fX, &fY); + #endif + // sincospi(0.0, &fX, &fY); + sinh(0.0); + // sinpi(0.0); + sqrt(0.0); + tan(0.0); + tanh(0.0); + tgamma(2.0); + trunc(0.0); + #ifdef __unix__ + y0(1.0); + y1(1.0); + yn(1, 1.0); + #elif _WIN64 + _y0(1.0); + _y1(1.0); + _yn(1, 1.0); + #endif +} + +static void compileOnHost() { double_precision_math_functions(); } + +int main() { + compileOnHost(); + passed(); +} diff --git a/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp b/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp old mode 100755 new mode 100644 index cd827bc318..78c2099e8f --- a/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp +++ b/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp @@ -1,152 +1,152 @@ -/* -Copyright (c) 2015-2016 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. -*/ - -/* HIT_START - * BUILD: %t %s ../test_common.cpp - * TEST: %t - * HIT_END - */ - -#include -//#include -#include "test_common.h" -#include - -#pragma GCC diagnostic ignored "-Wall" -#pragma clang diagnostic ignored "-Wunused-variable" - -__host__ void single_precision_math_functions() { - 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); - ceilf(0.0f); - 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); - #ifdef __unix__ - exp10f(0.0f); - #endif - exp2f(0.0f); - expf(0.0f); - expm1f(0.0f); - fabsf(1.0f); - fdimf(1.0f, 0.0f); -#if defined(__HIP_PLATFORM_HCC__) - // fdividef(0.0f, 1.0f); -#endif - floorf(0.0f); - fmaf(1.0f, 2.0f, 3.0f); - fmaxf(0.0f, 0.0f); - fminf(0.0f, 0.0f); - fmodf(0.0f, 1.0f); - frexpf(0.0f, &iX); - hypotf(1.0f, 0.0f); - ilogbf(1.0f); - std::isfinite(0.0f); - std::isinf(0.0f); - std::isnan(0.0f); - #ifdef __unix__ - j0f(0.0f); - j1f(0.0f); - jnf(-1.0f, 1.0f); - #endif - 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); - nanf("1"); - nearbyintf(0.0f); - // nextafterf(0.0f); -#if defined(__HIP_PLATFORM_HCC__) - // norm3df(1.0f, 0.0f, 0.0f); - // norm4df(1.0f, 0.0f, 0.0f, 0.0f); -#endif - // 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); -#if defined(__HIP_PLATFORM_HCC__) - // rhypotf(0.0f, 1.0f); -#endif - rintf(1.0f); -#if defined(__HIP_PLATFORM_HCC__) - // rnorm3df(0.0f, 0.0f, 1.0f); - // rnorm4df(0.0f, 0.0f, 0.0f, 1.0f); - fX = 1.0f; // rnormf(1, &fX); -#endif - roundf(0.0f); - /// rsqrtf(1.0f); - scalblnf(0.0f, 1); - scalbnf(0.0f, 1); - std::signbit(1.0f); - #ifdef __unix__ - sincosf(0.0f, &fX, &fY); - #endif - // 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); - truncf(0.0f); - #ifdef __unix__ - y0f(1.0f); - y1f(1.0f); - ynf(1, 1.0f); - #endif -} - -static void compileOnHost() { single_precision_math_functions(); } - -int main() { - compileOnHost(); - passed(); -} +/* +Copyright (c) 2015-2016 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. +*/ + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * TEST: %t + * HIT_END + */ + +#include +//#include +#include "test_common.h" +#include + +#pragma GCC diagnostic ignored "-Wall" +#pragma clang diagnostic ignored "-Wunused-variable" + +__host__ void single_precision_math_functions() { + 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); + ceilf(0.0f); + 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); + #ifdef __unix__ + exp10f(0.0f); + #endif + exp2f(0.0f); + expf(0.0f); + expm1f(0.0f); + fabsf(1.0f); + fdimf(1.0f, 0.0f); +#if defined(__HIP_PLATFORM_HCC__) + // fdividef(0.0f, 1.0f); +#endif + floorf(0.0f); + fmaf(1.0f, 2.0f, 3.0f); + fmaxf(0.0f, 0.0f); + fminf(0.0f, 0.0f); + fmodf(0.0f, 1.0f); + frexpf(0.0f, &iX); + hypotf(1.0f, 0.0f); + ilogbf(1.0f); + std::isfinite(0.0f); + std::isinf(0.0f); + std::isnan(0.0f); + #ifdef __unix__ + j0f(0.0f); + j1f(0.0f); + jnf(-1.0f, 1.0f); + #endif + 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); + nanf("1"); + nearbyintf(0.0f); + // nextafterf(0.0f); +#if defined(__HIP_PLATFORM_HCC__) + // norm3df(1.0f, 0.0f, 0.0f); + // norm4df(1.0f, 0.0f, 0.0f, 0.0f); +#endif + // 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); +#if defined(__HIP_PLATFORM_HCC__) + // rhypotf(0.0f, 1.0f); +#endif + rintf(1.0f); +#if defined(__HIP_PLATFORM_HCC__) + // rnorm3df(0.0f, 0.0f, 1.0f); + // rnorm4df(0.0f, 0.0f, 0.0f, 1.0f); + fX = 1.0f; // rnormf(1, &fX); +#endif + roundf(0.0f); + /// rsqrtf(1.0f); + scalblnf(0.0f, 1); + scalbnf(0.0f, 1); + std::signbit(1.0f); + #ifdef __unix__ + sincosf(0.0f, &fX, &fY); + #endif + // 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); + truncf(0.0f); + #ifdef __unix__ + y0f(1.0f); + y1f(1.0f); + ynf(1, 1.0f); + #endif +} + +static void compileOnHost() { single_precision_math_functions(); } + +int main() { + compileOnHost(); + passed(); +} From 6ce86f409d3635b2a5808f74f6809e3192fdee14 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 8 Aug 2019 01:27:41 -0700 Subject: [PATCH 7/9] Add support for hipFuncGetAttribute (#1279) * Add support for hipFunGetAttribute * Support NVCC path * Test using sample module_api_global * Try fixing CI build failure due to hip_prof_gen scan * Fix for CI build issue * Resolve conflict * Rebase and resolve conflicts with master * Fix build error * Fix NVCC path build error --- include/hip/hcc_detail/driver_types.h | 14 ++++++ include/hip/hcc_detail/hip_runtime_api.h | 15 +++++- include/hip/nvcc_detail/hip_runtime_api.h | 18 ++++++++ .../0_Intro/module_api_global/runKernel.cpp | 5 ++ src/hip_module.cpp | 46 +++++++++++++++++++ 5 files changed, 96 insertions(+), 2 deletions(-) diff --git a/include/hip/hcc_detail/driver_types.h b/include/hip/hcc_detail/driver_types.h index 18f344d642..510d3d058e 100644 --- a/include/hip/hcc_detail/driver_types.h +++ b/include/hip/hcc_detail/driver_types.h @@ -319,4 +319,18 @@ static inline struct hipExtent make_hipExtent(size_t w, size_t h, size_t d) { return e; } +typedef enum hipFunction_attribute { + HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK, + HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES, + HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES, + HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES, + HIP_FUNC_ATTRIBUTE_NUM_REGS, + HIP_FUNC_ATTRIBUTE_PTX_VERSION, + HIP_FUNC_ATTRIBUTE_BINARY_VERSION, + HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA, + HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES, + HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT, + HIP_FUNC_ATTRIBUTE_MAX +}hipFunction_attribute; + #endif diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index e8871912fa..e71ac8d8fb 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2603,16 +2603,27 @@ hipError_t hipModuleUnload(hipModule_t module); hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, const char* kname); /** - * @bried Find out attributes for a given function. + * @brief Find out attributes for a given function. * * @param [out] attr * @param [in] func * - * @returns hipSuccess, hipErrorInvalidDeviceFunction + * @returns hipSuccess, hipErrorInvalidValue, hipErrorInvalidDeviceFunction */ hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func); +/** + * @brief Find out a specific attribute for a given function. + * + * @param [out] value + * @param [in] attrib + * @param [in] hfunc + * + * @returns hipSuccess, hipErrorInvalidValue, hipErrorInvalidDeviceFunction + */ +hipError_t hipFuncGetAttribute(int* value, hipFunction_attribute attrib, hipFunction_t hfunc); + #if !__HIP_VDI__ #if defined(__cplusplus) } // extern "C" diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index fb4987feee..a77d5560d8 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -162,6 +162,7 @@ typedef CUdeviceptr hipDeviceptr_t; typedef struct cudaArray hipArray; typedef struct cudaArray* hipArray_const_t; typedef cudaFuncAttributes hipFuncAttributes; +typedef CUfunction_attribute hipFunction_attribute; #define hip_Memcpy2D CUDA_MEMCPY2D #define hipMemcpy3DParms cudaMemcpy3DParms #define hipArrayDefault cudaArrayDefault @@ -197,6 +198,19 @@ typedef cudaSurfaceObject_t hipSurfaceObject_t; #define hipSharedMemBankSizeFourByte cudaSharedMemBankSizeFourByte #define hipSharedMemBankSizeEightByte cudaSharedMemBankSizeEightByte +//Function Attributes +#define HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK CU_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK +#define HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES CU_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES +#define HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES CU_FUNC_ATTRIBUTE_CONST_SIZE_BYTES +#define HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES +#define HIP_FUNC_ATTRIBUTE_NUM_REGS CU_FUNC_ATTRIBUTE_NUM_REGS +#define HIP_FUNC_ATTRIBUTE_PTX_VERSION CU_FUNC_ATTRIBUTE_PTX_VERSION +#define HIP_FUNC_ATTRIBUTE_BINARY_VERSION CU_FUNC_ATTRIBUTE_BINARY_VERSION +#define HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA CU_FUNC_ATTRIBUTE_CACHE_MODE_CA +#define HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES CU_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES +#define HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT CU_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT +#define HIP_FUNC_ATTRIBUTE_MAX CU_FUNC_ATTRIBUTE_MAX + inline static hipError_t hipCUDAErrorTohipError(cudaError_t cuError) { switch (cuError) { case cudaSuccess: @@ -1197,6 +1211,10 @@ inline static hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const voi return hipCUDAErrorTohipError(cudaFuncGetAttributes(attr, func)); } +inline static hipError_t hipFuncGetAttribute (int* value, hipFunction_attribute attrib, hipFunction_t hfunc) { + return hipCUResultTohipError(cuFuncGetAttribute(value, attrib, hfunc)); +} + inline static hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod, const char* name) { return hipCUResultTohipError(cuModuleGetGlobal(dptr, bytes, hmod, name)); diff --git a/samples/0_Intro/module_api_global/runKernel.cpp b/samples/0_Intro/module_api_global/runKernel.cpp index 7031b9f898..7aa7f7ba2d 100644 --- a/samples/0_Intro/module_api_global/runKernel.cpp +++ b/samples/0_Intro/module_api_global/runKernel.cpp @@ -124,6 +124,11 @@ int main() { { hipFunction_t Function; HIP_CHECK(hipModuleGetFunction(&Function, Module, "test_globals")); + int val =-1; + HIP_CHECK(hipFuncGetAttribute(&val, HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES,Function)); + printf("Shared Size Bytes = %d\n",val); + HIP_CHECK(hipFuncGetAttribute(&val, HIP_FUNC_ATTRIBUTE_NUM_REGS, Function)); + printf("Num Regs = %d\n",val); HIP_CHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config)); hipMemcpyDtoH(B, Bd, SIZE); diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 04cdedbb75..93751167ee 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -803,6 +803,52 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func) return ihipLogStatus(hipSuccess); } +hipError_t hipFuncGetAttribute(int* value, hipFunction_attribute attrib, hipFunction_t hfunc) +{ + HIP_INIT_API(hipFuncGetAttribute, value, attrib, hfunc); + using namespace hip_impl; + + hipError_t retVal = hipSuccess; + if (!value) return ihipLogStatus(hipErrorInvalidValue); + hipFuncAttributes attr{}; + attr = make_function_attributes(tls, *hfunc); + switch(attrib) { + case HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES: + *value = (int) attr.sharedSizeBytes; + break; + case HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK: + *value = attr.maxThreadsPerBlock; + break; + case HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES: + *value = (int) attr.constSizeBytes; + break; + case HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES: + *value = (int) attr.localSizeBytes; + break; + case HIP_FUNC_ATTRIBUTE_NUM_REGS: + *value = attr.numRegs; + break; + case HIP_FUNC_ATTRIBUTE_PTX_VERSION: + *value = attr.ptxVersion; + break; + case HIP_FUNC_ATTRIBUTE_BINARY_VERSION: + *value = attr.binaryVersion; + break; + case HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA: + *value = attr.cacheModeCA; + break; + case HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES: + *value = attr.maxDynamicSharedSizeBytes; + break; + case HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT: + *value = attr.preferredShmemCarveout; + break; + default: + retVal = hipErrorInvalidValue; + } + return ihipLogStatus(retVal); +} + hipError_t ihipModuleLoadData(TlsData *tls, hipModule_t* module, const void* image) { using namespace hip_impl; From 83af327ef20b6e82086c5004898b76ee4bd5a854 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Thu, 8 Aug 2019 04:33:04 -0400 Subject: [PATCH 8/9] Compile HIP runtime with hidden visibility by default (#1303) * add default visibility to most APIs in program_state * remove unwanted C++ headers * Add symbol visibility pragmas and compiler flags * Add visibility attribute to APIs in channel_descriptor and hip_hcc * remove unused headers * simplify build flags with hcc * add pragma visibility hidden to functional_grid_launch * [CMake] add gfx908 back --- CMakeLists.txt | 8 +++----- include/hip/hcc_detail/channel_descriptor.h | 2 ++ .../hip/hcc_detail/functional_grid_launch.hpp | 19 ++++++------------- include/hip/hcc_detail/hip_runtime_api.h | 14 ++++++++------ include/hip/hcc_detail/program_state.hpp | 12 ++++++++++++ include/hip/hip_common.h | 8 ++++++++ include/hip/hip_hcc.h | 4 ++++ src/functional_grid_launch.inl | 2 +- src/hip_hcc.cpp | 1 + 9 files changed, 45 insertions(+), 25 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index 187b32045b..8b9a603ad8 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -221,7 +221,7 @@ if(HIP_PLATFORM STREQUAL "hcc") set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -DHIP_VERSION_MAJOR=${HIP_VERSION_MAJOR} -DHIP_VERSION_MINOR=${HIP_VERSION_MINOR} -DHIP_VERSION_PATCH=${HIP_VERSION_PATCH}") # Add remaining flags - execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --cxxflags OUTPUT_VARIABLE HCC_CXX_FLAGS) + set(HCC_CXX_FLAGS "-hc -fno-gpu-rdc --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908 ") set(HIP_HCC_BUILD_FLAGS "${HIP_HCC_BUILD_FLAGS} -fPIC ${HCC_CXX_FLAGS} -I${HSA_PATH}/include") # Set compiler and compiler flags @@ -252,9 +252,6 @@ if(HIP_PLATFORM STREQUAL "hcc") src/env.cpp src/h2f.cpp) - execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS) - set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS}") - set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906 --amdgpu-target=gfx908") if(COMPILE_HIP_ATP_MARKER) set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L/opt/rocm/profiler/CXLActivityLogger/bin/x86_64 -lCXLActivityLogger") endif() @@ -270,7 +267,8 @@ if(HIP_PLATFORM STREQUAL "hcc") PRIVATE ${PROJECT_SOURCE_DIR}/include ${HSA_PATH}/include) target_link_libraries(hiprtc PUBLIC stdc++fs) endif() - + set_target_properties(hip_hcc PROPERTIES CXX_VISIBILITY_PRESET hidden) + set_target_properties(hip_hcc PROPERTIES VISIBILITY_INLINES_HIDDEN 1) if(HIP_PLATFORM STREQUAL "hcc") find_package(amd_comgr REQUIRED CONFIG diff --git a/include/hip/hcc_detail/channel_descriptor.h b/include/hip/hcc_detail/channel_descriptor.h index de290fafc1..38acff9951 100644 --- a/include/hip/hcc_detail/channel_descriptor.h +++ b/include/hip/hcc_detail/channel_descriptor.h @@ -23,11 +23,13 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H #define HIP_INCLUDE_HIP_HCC_DETAIL_CHANNEL_DESCRIPTOR_H +#include #include #include #ifdef __cplusplus +HIP_PUBLIC_API hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f); static inline hipChannelFormatDesc hipCreateChannelDescHalf() { diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp index f502fddf42..883470b1fb 100644 --- a/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/include/hip/hcc_detail/functional_grid_launch.hpp @@ -25,25 +25,16 @@ THE SOFTWARE. #include "concepts.hpp" #include "helpers.hpp" #include "program_state.hpp" +#include "hip_runtime_api.h" -#include "hc.hpp" -#include "hip/hip_hcc.h" -#include "hip_runtime.h" - -#include -#include #include #include -#include -#include -#include #include -#include #include #include -#include #include -#include + +#pragma GCC visibility push(hidden) namespace hip_impl { template {}>::type* = nullptr> @@ -114,7 +105,7 @@ inline hip_impl::kernarg make_kernarg( } -hsa_agent_t target_agent(hipStream_t stream); +HIP_INTERNAL_EXPORTED_API hsa_agent_t target_agent(hipStream_t stream); inline __attribute__((visibility("hidden"))) @@ -194,3 +185,5 @@ inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBloc hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{}, std::move(args)...); } + +#pragma GCC visibility pop diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index e71ac8d8fb..cec9eb9c94 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -67,19 +67,18 @@ THE SOFTWARE. #define HIP_LAUNCH_PARAM_END ((void*)0x03) #ifdef __cplusplus - #include - #include - #include - #include - #include - #define __dparm(x) \ = x #else #define __dparm(x) #endif +#ifdef __GNUC__ +#pragma GCC visibility push (default) +#endif + #ifdef __cplusplus + namespace hip_impl { hipError_t hip_init(); } // namespace hip_impl @@ -3233,6 +3232,9 @@ hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject); #endif +#ifdef __GNUC__ +#pragma GCC visibility pop +#endif /** *------------------------------------------------------------------------------------------------- diff --git a/include/hip/hcc_detail/program_state.hpp b/include/hip/hcc_detail/program_state.hpp index 4e0fbb76a4..fca88f8094 100644 --- a/include/hip/hcc_detail/program_state.hpp +++ b/include/hip/hcc_detail/program_state.hpp @@ -31,11 +31,19 @@ THE SOFTWARE. #include #include +#include + struct ihipModuleSymbol_t; using hipFunction_t = ihipModuleSymbol_t*; namespace hip_impl { +// This section contains internal APIs that +// needs to be exported +#ifdef __GNUC__ +#pragma GCC visibility push (default) +#endif + struct kernarg_impl; class kernarg { public: @@ -83,6 +91,10 @@ private: friend kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t); }; +#ifdef __GNUC__ +#pragma GCC visibility pop +#endif + inline __attribute__((visibility("hidden"))) program_state& get_program_state() { diff --git a/include/hip/hip_common.h b/include/hip/hip_common.h index 71285fc8dc..79c787be4e 100644 --- a/include/hip/hip_common.h +++ b/include/hip/hip_common.h @@ -45,6 +45,14 @@ THE SOFTWARE. #define __HIP_DEVICE_COMPILE__ 1 #endif +#ifdef __GNUC__ +#define HIP_PUBLIC_API __attribute__ ((visibility ("default"))) +#define HIP_INTERNAL_EXPORTED_API __attribute__ ((visibility ("default"))) +#else +#define HIP_PUBLIC_API +#define HIP_INTERNAL_EXPORTED_API +#endif + #if __HIP_DEVICE_COMPILE__ == 0 // 32-bit Atomics #define __HIP_ARCH_HAS_GLOBAL_INT32_ATOMICS__ (0) diff --git a/include/hip/hip_hcc.h b/include/hip/hip_hcc.h index 3935f4903b..a82a8c622b 100644 --- a/include/hip/hip_hcc.h +++ b/include/hip/hip_hcc.h @@ -48,6 +48,7 @@ class accelerator_view; * @brief Return hc::accelerator associated with the specified deviceId * @return #hipSuccess, #hipErrorInvalidDevice */ +HIP_PUBLIC_API hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc); /** @@ -56,6 +57,7 @@ hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc); * If stream is 0, the accelerator_view for the default stream is returned. * @return #hipSuccess */ +HIP_PUBLIC_API hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** av); @@ -89,6 +91,7 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** a * HIP/ROCm actually updates the start event when the associated kernel completes. */ +HIP_PUBLIC_API hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, uint32_t localWorkSizeX, uint32_t localWorkSizeY, @@ -98,6 +101,7 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, hipEvent_t stopEvent = nullptr, uint32_t flags = 0); +HIP_PUBLIC_API hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, uint32_t localWorkSizeX, uint32_t localWorkSizeY, diff --git a/src/functional_grid_launch.inl b/src/functional_grid_launch.inl index 3c611ed468..a99a15db28 100644 --- a/src/functional_grid_launch.inl +++ b/src/functional_grid_launch.inl @@ -41,7 +41,7 @@ using namespace std; namespace hip_impl { - hsa_agent_t target_agent(hipStream_t stream) + HIP_INTERNAL_EXPORTED_API hsa_agent_t target_agent(hipStream_t stream) { if (stream) { return *static_cast( diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index dcfdfb6137..da0b38ed4f 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -46,6 +46,7 @@ THE SOFTWARE. #include "hsa/hsa_ext_image.h" #include "hip/hip_runtime.h" #include "hip_hcc_internal.h" +#include "hip/hip_hcc.h" #include "trace_helper.h" #include "env.h" From 05b685cc6b15e100450e2a34d54aaa88b7970c87 Mon Sep 17 00:00:00 2001 From: Rene van Oostrum Date: Thu, 8 Aug 2019 10:46:45 -0500 Subject: [PATCH 9/9] Add missing hipHostRegister flags to hipify-perl for flags that are defined and supported in hip_runtime_api.h --- bin/hipify-perl | 2 ++ 1 file changed, 2 insertions(+) diff --git a/bin/hipify-perl b/bin/hipify-perl index e375d22928..625d2b6fd2 100755 --- a/bin/hipify-perl +++ b/bin/hipify-perl @@ -252,6 +252,8 @@ while (@ARGV) { $ft{'mem'} += s/\bcudaHostAllocPortable\b/hipHostMallocPortable/g; $ft{'mem'} += s/\bcudaHostAllocMapped\b/hipHostMallocMapped/g; $ft{'mem'} += s/\bcudaHostAllocWriteCombined\b/hipHostMallocWriteCombined/g; + $ft{'mem'} += s/\bcudaHostRegisterDefault\b/hipHostRegisterDefault/g; + $ft{'mem'} += s/\bcudaHostRegisterPortable\b/hipHostRegisterPortable/g; $ft{'mem'} += s/\bcudaHostRegisterMapped\b/hipHostRegisterMapped/g; $ft{'mem'} += s/\bcudaHostRegister\b/hipHostRegister/g; $ft{'mem'} += s/\bcudaHostUnregister\b/hipHostUnregister/g;