From 2a5a2c66c47072f035de0873b1155418b69a46ba Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 19 Sep 2018 15:06:22 +0530 Subject: [PATCH 1/7] Update Jenkinsfile [ci] Update list of disable tests for automation --- Jenkinsfile | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index 2432cea38e..3eb24eebdd 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -167,7 +167,8 @@ def docker_build_inside_image( def build_image, String inside_args, String platf } // Cap the maximum amount of testing, in case of hangs - // Excluding hipPrintfKernel test from automation; variable fails on CI test machines + // Excluding hipVectorTypes test from automation; due to regression from HCC commit 2367133 + // Excluding hipFloatMath test from automation; due to regression from ROCDL commit 2fc04e1 timeout(time: 1, unit: 'HOURS') { stage("${platform} unit testing") @@ -177,7 +178,7 @@ def docker_build_inside_image( def build_image, String inside_args, String platf cd ${build_dir_rel} make install -j\$(nproc) make build_tests -i -j\$(nproc) - ctest -E hipVectorTypes + ctest -E "(hipVectorTypes.tst|hipVectorTypesDevice.tst|hipFloatMath.tst)" """ // If unit tests output a junit or xunit file in the future, jenkins can parse that file // to display test results on the dashboard From bd622a4b4ab4679c8d902f797b3cd60b95f13d9c Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Fri, 10 Aug 2018 11:12:53 -0400 Subject: [PATCH 2/7] Add fma function with float and _Float16 arguments --- include/hip/hcc_detail/math_functions.h | 10 ++ tests/src/deviceLib/hipTestFMA.cpp | 142 ++++++++++++++++++++++++ 2 files changed, 152 insertions(+) create mode 100644 tests/src/deviceLib/hipTestFMA.cpp diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index 702c120b86..c1adef68fd 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/include/hip/hcc_detail/math_functions.h @@ -1166,6 +1166,16 @@ long long llabs(long long x) #endif // END INTEGER +__DEVICE__ +inline _Float16 fma(_Float16 x, _Float16 y, _Float16 z) { + return __ocml_fma_f16(x, y, z); +} + +__DEVICE__ +inline float fma(float x, float y, float z) { + return fmaf(x, y, z); +} + #pragma push_macro("__DEF_FLOAT_FUN") #pragma push_macro("__DEF_FLOAT_FUN2") #pragma push_macro("__DEF_FLOAT_FUN2I") diff --git a/tests/src/deviceLib/hipTestFMA.cpp b/tests/src/deviceLib/hipTestFMA.cpp new file mode 100644 index 0000000000..5e1913a5c7 --- /dev/null +++ b/tests/src/deviceLib/hipTestFMA.cpp @@ -0,0 +1,142 @@ +/* +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 + * RUN: %t + * HIT_END + */ + +#include "test_common.h" +#include +#include +#include + +#define HIP_ASSERT(status) assert(status == hipSuccess) + +#define LEN 50 +#define SIZE (LEN * sizeof(bool)) + +struct TestFMA { + static __global__ void kernel(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); + + // check promote to float. + Check(fma(f, f, 3) == floatResult); + Check(fma(f, f, (char)3) == floatResult); + Check(fma(f, f, (unsigned char)3) == floatResult); + Check(fma(f, f, (short)3) == floatResult); + Check(fma(f, f, (unsigned short)3) == floatResult); + Check(fma(f, f, (int)3) == floatResult); + Check(fma(f, f, (unsigned int)3) == floatResult); + Check(fma(f, f, (long)3) == floatResult); + Check(fma(f, f, (unsigned long)3) == floatResult); + Check(fma(f, f, true) == fma(f, f, 1.0f)); + + // check promote to double. + Check(fma(d, (double)f, 3) == doubleResult); + Check(fma(d, (double)f, (char)3) == doubleResult); + Check(fma(d, (double)f, (unsigned char)3) == doubleResult); + Check(fma(d, (double)f, (short)3) == doubleResult); + Check(fma(d, (double)f, (unsigned short)3) == doubleResult); + Check(fma(d, (double)f, (int)3) == doubleResult); + Check(fma(d, (double)f, (unsigned int)3) == doubleResult); + Check(fma(d, (double)f, (long)3) == doubleResult); + Check(fma(d, (double)f, (unsigned long)3) == doubleResult); + Check(fma(d, (double)f, true) == fma((double)f, (double)f, 1.0)); + + while (i < LEN) + Check(true); + } + void run() { + bool *Ad; + bool A[LEN]; + for (unsigned i = 0; i < LEN; i++) { + A[i] = 0; + } + + HIP_ASSERT(hipMalloc((void **)&Ad, SIZE)); + hipLaunchKernelGGL(kernel, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, Ad); + HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost)); + + for (unsigned i = 0; i < LEN; i++) { + assert(A[i]); + } + } +}; + +struct TestHalfFMA { + static __global__ void kernel(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. + Check(fma(h, h, 3) == halfResult); + Check(fma(h, h, (char)3) == halfResult); + Check(fma(h, h, (unsigned char)3) == halfResult); + Check(fma(h, h, (short)3) == halfResult); + Check(fma(h, h, (unsigned short)3) == halfResult); + Check(fma(h, h, (int)3) == halfResult); + Check(fma(h, h, (unsigned int)3) == halfResult); + Check(fma(h, h, (long)3) == halfResult); + Check(fma(h, h, (unsigned long)3) == halfResult); + Check(fma(h, h, true) == fma(h, h, (_Float16)1)); + + while (i < LEN) + Check(true); + } + + void run() { + bool *Ad; + bool A[LEN]; + for (unsigned i = 0; i < LEN; i++) { + A[i] = 0; + } + + HIP_ASSERT(hipMalloc((void **)&Ad, SIZE)); + hipLaunchKernelGGL(kernel, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, Ad); + HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost)); + + for (unsigned i = 0; i < LEN; i++) { + assert(A[i]); + } + } +}; + +int main() { + TestFMA().run(); + TestHalfFMA().run(); + passed(); +} From a5c961e26c0f6220cf1b2634d86f75a5e8426a5c Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Wed, 19 Sep 2018 10:38:48 -0400 Subject: [PATCH 3/7] Silent warnings about duplicate static keyword static is already in __DEVICE__, so should be removed. --- include/hip/hcc_detail/math_functions.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index 702c120b86..b618790133 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/include/hip/hcc_detail/math_functions.h @@ -1359,10 +1359,10 @@ __DEVICE__ inline static unsigned long long max(long long arg1, unsigned long lo return max((unsigned long long) arg1, arg2); }*/ #else -__DEVICE__ inline static int min(int arg1, int arg2) { +__DEVICE__ inline int min(int arg1, int arg2) { return (arg1 < arg2) ? arg1 : arg2; } -__DEVICE__ inline static int max(int arg1, int arg2) { +__DEVICE__ inline int max(int arg1, int arg2) { return (arg1 > arg2) ? arg1 : arg2; } From e01c53b185c3e530530ebeebeffb307749f47790 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 19 Sep 2018 18:29:56 +0000 Subject: [PATCH 4/7] Create separate config files for hcc and clang The state of HIP_COMPILER is not passing into config files, therefore config files cannot use if statements to determine dependency. For HIP-Clang, we should remove find_dependency(hcc), so we create separate config files depending on compiler path. This fixes issue in apps that uses hip-config.cmake. --- CMakeLists.txt | 12 +++- ...nfig.cmake.in => hip-config-clang.cmake.in | 3 - hip-config-hcc.cmake.in | 65 +++++++++++++++++++ 3 files changed, 76 insertions(+), 4 deletions(-) rename hip-config.cmake.in => hip-config-clang.cmake.in (97%) create mode 100644 hip-config-hcc.cmake.in diff --git a/CMakeLists.txt b/CMakeLists.txt index edb8c2d238..3c62ea4365 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -288,12 +288,22 @@ if(HIP_PLATFORM STREQUAL "hcc") install(EXPORT hip-targets DESTINATION ${CONFIG_PACKAGE_INSTALL_DIR} NAMESPACE hip::) include(CMakePackageConfigHelpers) + if(HIP_COMPILER STREQUAL "hcc") configure_package_config_file( - hip-config.cmake.in + hip-config-hcc.cmake.in ${CMAKE_CURRENT_BINARY_DIR}/hip-config.cmake INSTALL_DESTINATION ${CONFIG_PACKAGE_INSTALL_DIR} PATH_VARS LIB_INSTALL_DIR INCLUDE_INSTALL_DIR BIN_INSTALL_DIR ) + elseif(HIP_COMPILER STREQUAL "clang") + configure_package_config_file( + hip-config-clang.cmake.in + ${CMAKE_CURRENT_BINARY_DIR}/hip-config.cmake + INSTALL_DESTINATION ${CONFIG_PACKAGE_INSTALL_DIR} + PATH_VARS LIB_INSTALL_DIR INCLUDE_INSTALL_DIR BIN_INSTALL_DIR + ) + endif() + write_basic_package_version_file( ${CMAKE_CURRENT_BINARY_DIR}/hip-config-version.cmake VERSION "${HIP_VERSION}" diff --git a/hip-config.cmake.in b/hip-config-clang.cmake.in similarity index 97% rename from hip-config.cmake.in rename to hip-config-clang.cmake.in index d5dc6803fc..240f01f60e 100644 --- a/hip-config.cmake.in +++ b/hip-config-clang.cmake.in @@ -48,9 +48,6 @@ set_and_check( hip_BIN_INSTALL_DIR "@PACKAGE_BIN_INSTALL_DIR@" ) set_and_check(hip_HIPCC_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipcc") set_and_check(hip_HIPCONFIG_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipconfig") -if(HIP_COMPILER STREQUAL "hcc") - find_dependency(hcc) -endif() include( "${CMAKE_CURRENT_LIST_DIR}/hip-targets.cmake" ) set( hip_LIBRARIES hip::host hip::device) diff --git a/hip-config-hcc.cmake.in b/hip-config-hcc.cmake.in new file mode 100644 index 0000000000..efcdf708bb --- /dev/null +++ b/hip-config-hcc.cmake.in @@ -0,0 +1,65 @@ +@PACKAGE_INIT@ + +include(CMakeFindDependencyMacro OPTIONAL RESULT_VARIABLE _CMakeFindDependencyMacro_FOUND) +if (NOT _CMakeFindDependencyMacro_FOUND) + macro(find_dependency dep) + if (NOT ${dep}_FOUND) + set(cmake_fd_version) + if (${ARGC} GREATER 1) + set(cmake_fd_version ${ARGV1}) + endif() + set(cmake_fd_exact_arg) + if(${CMAKE_FIND_PACKAGE_NAME}_FIND_VERSION_EXACT) + set(cmake_fd_exact_arg EXACT) + endif() + set(cmake_fd_quiet_arg) + if(${CMAKE_FIND_PACKAGE_NAME}_FIND_QUIETLY) + set(cmake_fd_quiet_arg QUIET) + endif() + set(cmake_fd_required_arg) + if(${CMAKE_FIND_PACKAGE_NAME}_FIND_REQUIRED) + set(cmake_fd_required_arg REQUIRED) + endif() + find_package(${dep} ${cmake_fd_version} + ${cmake_fd_exact_arg} + ${cmake_fd_quiet_arg} + ${cmake_fd_required_arg} + ) + string(TOUPPER ${dep} cmake_dep_upper) + if (NOT ${dep}_FOUND AND NOT ${cmake_dep_upper}_FOUND) + set(${CMAKE_FIND_PACKAGE_NAME}_NOT_FOUND_MESSAGE "${CMAKE_FIND_PACKAGE_NAME} could not be found because dependency ${dep} could not be found.") + set(${CMAKE_FIND_PACKAGE_NAME}_FOUND False) + return() + endif() + set(cmake_fd_version) + set(cmake_fd_required_arg) + set(cmake_fd_quiet_arg) + set(cmake_fd_exact_arg) + endif() + endmacro() +endif() + + +set_and_check( hip_INCLUDE_DIR "@PACKAGE_INCLUDE_INSTALL_DIR@" ) +set_and_check( hip_INCLUDE_DIRS "${hip_INCLUDE_DIR}" ) +set_and_check( hip_LIB_INSTALL_DIR "@PACKAGE_LIB_INSTALL_DIR@" ) +set_and_check( hip_BIN_INSTALL_DIR "@PACKAGE_BIN_INSTALL_DIR@" ) + +set_and_check(hip_HIPCC_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipcc") +set_and_check(hip_HIPCONFIG_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipconfig") + +find_dependency(hcc) +include( "${CMAKE_CURRENT_LIST_DIR}/hip-targets.cmake" ) + +set( hip_LIBRARIES hip::host hip::device) +set( hip_LIBRARY ${hip_LIBRARIES}) + +set(HIP_INCLUDE_DIR ${hip_INCLUDE_DIR}) +set(HIP_INCLUDE_DIRS ${hip_INCLUDE_DIRS}) +set(HIP_LIB_INSTALL_DIR ${hip_LIB_INSTALL_DIR}) +set(HIP_BIN_INSTALL_DIR ${hip_BIN_INSTALL_DIR}) +set(HIP_LIBRARIES ${hip_LIBRARIES}) +set(HIP_LIBRARY ${hip_LIBRARY}) +set(HIP_HIPCC_EXECUTABLE ${hip_HIPCC_EXECUTABLE}) +set(HIP_HIPCONFIG_EXECUTABLE ${hip_HIPCONFIG_EXECUTABLE}) + From ecd6a212c7e1cbac802b415adf3bd9ecb72648f8 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 20 Sep 2018 11:23:51 +0530 Subject: [PATCH 5/7] Update hipTestFMA.cpp Fix the test so that it works on nvcc path as well. --- tests/src/deviceLib/hipTestFMA.cpp | 21 +++++++++------------ 1 file changed, 9 insertions(+), 12 deletions(-) diff --git a/tests/src/deviceLib/hipTestFMA.cpp b/tests/src/deviceLib/hipTestFMA.cpp index 5e1913a5c7..1f1a5bc921 100644 --- a/tests/src/deviceLib/hipTestFMA.cpp +++ b/tests/src/deviceLib/hipTestFMA.cpp @@ -33,8 +33,7 @@ THE SOFTWARE. #define LEN 50 #define SIZE (LEN * sizeof(bool)) -struct TestFMA { - static __global__ void kernel(bool *Ad) { + __global__ void kernelTestFMA(bool *Ad) { float f = 1.0f / 3.0f; double d = f; int i = 0; @@ -71,7 +70,8 @@ struct TestFMA { while (i < LEN) Check(true); } - void run() { + + void runTestFMA() { bool *Ad; bool A[LEN]; for (unsigned i = 0; i < LEN; i++) { @@ -79,17 +79,15 @@ struct TestFMA { } HIP_ASSERT(hipMalloc((void **)&Ad, SIZE)); - hipLaunchKernelGGL(kernel, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, Ad); + hipLaunchKernelGGL(kernelTestFMA, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, Ad); HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost)); for (unsigned i = 0; i < LEN; i++) { assert(A[i]); } } -}; -struct TestHalfFMA { - static __global__ void kernel(bool *Ad) { + __global__ void kernelTestHalfFMA(bool *Ad) { _Float16 h = (_Float16)(1.0f/3.0f); float f = h; double d = f; @@ -118,7 +116,7 @@ struct TestHalfFMA { Check(true); } - void run() { + void runTestHalfFMA() { bool *Ad; bool A[LEN]; for (unsigned i = 0; i < LEN; i++) { @@ -126,17 +124,16 @@ struct TestHalfFMA { } HIP_ASSERT(hipMalloc((void **)&Ad, SIZE)); - hipLaunchKernelGGL(kernel, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, Ad); + hipLaunchKernelGGL(kernelTestHalfFMA, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, Ad); HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost)); for (unsigned i = 0; i < LEN; i++) { assert(A[i]); } } -}; int main() { - TestFMA().run(); - TestHalfFMA().run(); + runTestFMA(); + runTestHalfFMA(); passed(); } From 6d794cd5057508c5f33d9ad966c32bfde3b60f64 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 21 Sep 2018 10:02:26 +0530 Subject: [PATCH 6/7] [ci] Renable nvcc testing Change-Id: I7d720b41a3ddc99453fee8b9be30494bfec3a808 --- Jenkinsfile | 36 ++++++++++++++++++++++++++++++++++++ 1 file changed, 36 insertions(+) diff --git a/Jenkinsfile b/Jenkinsfile index 3eb24eebdd..ab6dd6d67f 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -386,4 +386,40 @@ rocm_head: docker_clean_images( job_name, hip_image_name ) */ } +}, +cuda_9_x: +{ + node('hip-cuda') + { + //////////////////////////////////////////////////////////////////////// + // Block of string constants customizing behavior for cuda + String nvcc_ver = 'nvcc-9.x' + String from_image = 'ci_test_nodes/cuda-9.x/ubuntu-16.04:latest' + String inside_args = '--runtime=nvidia'; + + // Checkout source code, dependencies and version files + String source_hip_rel = checkout_and_version( nvcc_ver ) + + // Create/reuse a docker image that represents the hip build environment + def hip_build_image = docker_build_image( nvcc_ver, 'hip', '', source_hip_rel, from_image ) + + // Print system information for the log + hip_build_image.inside( inside_args ) + { + sh """#!/usr/bin/env bash + set -x + nvidia-smi + nvcc --version + """ + } + + // Conctruct a binary directory path based on build config + String build_hip_rel = build_directory_rel( build_config ); + + // Build hip inside of the build environment + docker_build_inside_image( hip_build_image, inside_args, nvcc_ver, "-DHIP_NVCC_FLAGS=--Wno-deprecated-gpu-targets", build_config, source_hip_rel, build_hip_rel ) + + // Clean docker image + docker_clean_images( 'hip', docker_build_image_name( ) ) + } } From 255589ae15467b5e9d1be454588f1fc7ca72cd58 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 21 Sep 2018 11:00:08 +0530 Subject: [PATCH 7/7] Update hipTestFMA.cpp --- tests/src/deviceLib/hipTestFMA.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/src/deviceLib/hipTestFMA.cpp b/tests/src/deviceLib/hipTestFMA.cpp index 1f1a5bc921..2771fac585 100644 --- a/tests/src/deviceLib/hipTestFMA.cpp +++ b/tests/src/deviceLib/hipTestFMA.cpp @@ -18,7 +18,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp + * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc * RUN: %t * HIT_END */