AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge 'master' into 'amd-master'

Browse Source 
Change-Id: I3de7b9ac03c05e50e7fbdb855e120b2dff8d23b4
This commit is contained in:
Jenkins
2018-09-23 20:41:55 -05:00
parent 175d39cb11 0e3ba5a47d
commit 91d77b2351
6 changed files with 266 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files
CMakeLists.txt
+11 -1
View File
@@ -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}"
Jenkinsfile
Vendored
+39 -2
View File
@@ -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
@@ -385,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( ) )
}
}
hip-config.cmake.in → hip-config-clang.cmake.in
-3
View File
@@ -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)
hip-config-hcc.cmake.in
+65
View File
@@ -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})
include/hip/hcc_detail/math_functions.h
+12 -2
View File
@@ -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")
@@ -1359,10 +1369,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;
}
tests/src/deviceLib/hipTestFMA.cpp
+139
View File
@@ -0,0 +1,139 @@
/*
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 EXCLUDE_HIP_PLATFORM nvcc
* RUN: %t
* HIT_END
*/
#include "test_common.h"
#include <hip/hip_runtime.h>
#include <hip/hip_runtime_api.h>
#include <iostream>
#define HIP_ASSERT(status) assert(status == hipSuccess)
#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);
// 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 runTestFMA() {
bool *Ad;
bool A[LEN];
for (unsigned i = 0; i < LEN; i++) {
A[i] = 0;
}
HIP_ASSERT(hipMalloc((void **)&Ad, SIZE));
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]);
}
}
__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.
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 runTestHalfFMA() {
bool *Ad;
bool A[LEN];
for (unsigned i = 0; i < LEN; i++) {
A[i] = 0;
}
HIP_ASSERT(hipMalloc((void **)&Ad, SIZE));
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() {
runTestFMA();
runTestHalfFMA();
passed();
}