AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Merge 'master' into 'amd-master'

Просмотр исходного кода 
Change-Id: I4f744ea43f46f148c348d41d19252c191fc7019a
Этот коммит содержится в:
Jenkins
2019-02-15 04:09:54 -06:00
родитель 59c54da4f0 438cf792b7
Коммит e2243d3d7f
14 изменённых файлов: 81 добавлений и 37 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
Jenkinsfile
поставляемый
+4 -4
Просмотреть файл
@@ -295,7 +295,7 @@ def docker_upload_dockerhub( String local_org, String image_name, String remote_
String build_config = 'Release'
String job_name = env.JOB_NAME.toLowerCase( )
// The following launches 3 builds in parallel: rocm-head, rocm-2.1.x and cuda-9.x
// The following launches 3 builds in parallel: rocm-head, rocm-2.1.x and cuda-10.x
parallel rocm_2_1:
{
node('hip-rocm')
@@ -386,14 +386,14 @@ rocm_head:
*/
}
},
cuda_9_x:
cuda_10_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 nvcc_ver = 'cuda-10.x'
String from_image = 'ci_test_nodes/cuda-10.x/ubuntu-16.04:latest'
String inside_args = '--runtime=nvidia';
// Checkout source code, dependencies and version files
bin/hipcc
+35 -10
Просмотреть файл
@@ -53,11 +53,14 @@ sub parse_config_file {
$verbose = $ENV{'HIPCC_VERBOSE'} // 0;
# Verbose: 0x1=commands, 0x2=paths, 0x4=hipcc args
$isWindows = $^O eq 'MSWin32';
$HIPCC_COMPILE_FLAGS_APPEND=$ENV{'HIPCC_COMPILE_FLAGS_APPEND'};
$HIPCC_LINK_FLAGS_APPEND=$ENV{'HIPCC_LINK_FLAGS_APPEND'};
$HIP_PATH=$ENV{'HIP_PATH'} // dirname (dirname $0); # use parent directory of hipcc
$HIP_VDI_HOME=$ENV{'HIP_VDI_HOME'};
$HIP_LIB_PATH=$ENV{'HIP_LIB_PATH'};
$HIP_CLANG_PATH=$ENV{'HIP_CLANG_PATH'};
$DEVICE_LIB_PATH=$ENV{'DEVICE_LIB_PATH'};
$HIP_CLANG_HCC_COMPAT_MODE=$ENV{'HIP_CLANG_HCC_COMPAT_MODE'}; # HCC compatibility mode
@@ -120,7 +123,9 @@ if (defined $HIP_VDI_HOME) {
}
$HIP_CLANG_INCLUDE_PATH = "$HIP_VDI_HOME/include/clang";
$HIP_INCLUDE_PATH = "$HIP_VDI_HOME/include";
$HIP_LIB_PATH = "$HIP_VDI_HOME/lib" . $bits;
if (!defined $HIP_LIB_PATH) {
$HIP_LIB_PATH = "$HIP_VDI_HOME/lib" . $bits;
}
}
if ($HIP_COMPILER eq "clang") {
@@ -188,13 +193,20 @@ if ($HIP_PLATFORM eq "clang") {
# is required.
$HIP_DEVLIB_FLAGS = " --hip-device-lib-path=$DEVICE_LIB_PATH";
$HIPCXXFLAGS .= " $HIP_DEVLIB_FLAGS -std=c++11 -isystem $HIP_CLANG_INCLUDE_PATH";
$HIPLDFLAGS .= " $HIP_DEVLIB_FLAGS -L$HIP_LIB_PATH -Wl,--rpath=$HIP_LIB_PATH -lhip_hcc";
$HIPLDFLAGS .= " $HIP_DEVLIB_FLAGS -L$HIP_LIB_PATH";
if (not $isWindows) {
$HIPLDFLAGS .= " -Wl,--rpath=$HIP_LIB_PATH";
}
$HIPLDFLAGS .= " -lhip_hcc";
if ($HIP_CLANG_HCC_COMPAT_MODE) {
## Allow __fp16 as function parameter and return type.
$HIPCXXFLAGS .= " -Xclang -fallow-half-arguments-and-returns -D__HIP_HCC_COMPAT_MODE__=1";
}
} elsif ($HIP_PLATFORM eq "hcc") {
$HIP_INCLUDE_PATH = "$HIP_PATH/include";
if (! defined $HIP_LIB_PATH) {
$HIP_LIB_PATH = "$HIP_PATH/lib";
}
$HSA_PATH=$ENV{'HSA_PATH'} // "/opt/rocm/hsa";
$HCC_HOME=$ENV{'HCC_HOME'} // $hipConfig{'HCC_HOME'} // "/opt/rocm/hcc";
@@ -268,7 +280,9 @@ if ($HIP_PLATFORM eq "clang") {
$HIPLDFLAGS .= " -L$marker_lib_path -lCXLActivityLogger -Wl,--rpath=$marker_lib_path";
}
$HIPLDFLAGS .= " -lm";
if (not $isWindows) {
$HIPLDFLAGS .= " -lm";
}
if ($verbose & 0x2) {
print ("HSA_PATH=$HSA_PATH\n");
@@ -306,6 +320,7 @@ my $runCmd = 1;
my $buildDeps = 0;
my $linkType = 1;
my $setLinkType = 0;
my $useCodeObjectV3 = 0;
my @options = ();
my @inputs = ();
@@ -409,6 +424,10 @@ foreach $arg (@ARGV)
$swallowArg = 1;
}
if (($trimarg eq '-mcode-object-v3') and ($HIP_PLATFORM eq 'clang') ) {
$useCodeObjectV3 = 1;
}
if (($arg =~ /--genco/) and $HIP_PLATFORM eq 'clang' ) {
$arg = "--cuda-device-only";
}
@@ -742,6 +761,12 @@ if ($buildDeps and $HIP_PLATFORM eq 'clang') {
$HIPCXXFLAGS .= " --cuda-host-only";
}
if ($useCodeObjectV3 and $HIP_PLATFORM eq 'clang') {
$HIPCXXFLAGS .= " -mcode-object-v3";
} elsif ($HIP_PLATFORM eq 'clang') {
$HIPCXXFLAGS .= " -mno-code-object-v3";
}
# Add --hip-link only if there are no source files.
if (!$needCXXFLAGS and $HIP_PLATFORM eq 'clang') {
$HIPLDFLAGS .= " --hip-link";
@@ -754,9 +779,9 @@ if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc')
if ($needHipHcc) {
if ($linkType eq 0) {
substr($HIPLDFLAGS,0,0) = " $HIP_PATH/lib/libhip_hcc_static.a " ;
substr($HIPLDFLAGS,0,0) = " $HIP_LIB_PATH/libhip_hcc_static.a " ;
} else {
substr($HIPLDFLAGS,0,0) = " -Wl,--rpath=$HIP_PATH/lib $HIP_PATH/lib/libhip_hcc.so ";
substr($HIPLDFLAGS,0,0) = " -Wl,--rpath=$HIP_LIB_PATH $HIP_LIB_PATH/libhip_hcc.so ";
}
}
@@ -766,13 +791,13 @@ if ($needHipHcc) {
# pass-through CPP mode.
if ($HIP_PLATFORM eq "clang") {
$HIPLDFLAGS .= " -lgcc_s -lgcc -lpthread -lm";
if ($isWindows) {
$HIPCXXFLAGS .= " -std=c++14 -fms-extensions -fms-compatibility";
} else {
$HIPLDFLAGS .= " -lgcc_s -lgcc -lpthread -lm";
}
}
# Windows specific options
if ($^O eq 'MSWin32' and $HIP_PLATFORM eq "clang") {
$HIPCXXFLAGS .= " -std=c++14 -fms-extensions -fms-compatibility"
}
if ($HIPCC_COMPILE_FLAGS_APPEND) {
$HIPCXXFLAGS .= " $HIPCC_COMPILE_FLAGS_APPEND";
docs/markdown/CURAND_API_supported_by_HIP.md
+2
Просмотреть файл
@@ -134,6 +134,8 @@
| `curandSetPseudoRandomGeneratorSeed` | `hiprandSetPseudoRandomGeneratorSeed` |
| `curandSetQuasiRandomGeneratorDimensions` | `hiprandSetQuasiRandomGeneratorDimensions` |
| `curandSetStream` | `hiprandSetStream` |
| `curandMakeMTGP32Constants` | `hiprandMakeMTGP32Constants` |
| `curandMakeMTGP32KernelState` | `hiprandMakeMTGP32KernelState` |
## **3. Device API Functions**
hipify-clang/src/CUDA2HIP.cpp
+2 -2
Просмотреть файл
@@ -48,9 +48,9 @@ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP{
{"curand_lognormal.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_mrg32k3a.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_mtgp32.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_mtgp32_host.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_mtgp32_host.h", {"hiprand_mtgp32_host.h", "", CONV_INCLUDE, API_RAND}},
{"curand_mtgp32_kernel.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_mtgp32dc_p_11213.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_mtgp32dc_p_11213.h", {"rocrand_mtgp32_11213.h", "", CONV_INCLUDE, API_RAND}},
{"curand_normal.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_normal_static.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
{"curand_philox4x32_x.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}},
hipify-clang/src/CUDA2HIP_RAND_API_functions.cpp
+2
Просмотреть файл
@@ -51,6 +51,8 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RAND_FUNCTION_MAP{
{"curandSetPseudoRandomGeneratorSeed", {"hiprandSetPseudoRandomGeneratorSeed", "", CONV_LIB_FUNC, API_RAND}},
{"curandSetQuasiRandomGeneratorDimensions", {"hiprandSetQuasiRandomGeneratorDimensions", "", CONV_LIB_FUNC, API_RAND}},
{"curandSetStream", {"hiprandSetStream", "", CONV_LIB_FUNC, API_RAND}},
{"curandMakeMTGP32Constants", {"hiprandMakeMTGP32Constants", "", CONV_LIB_FUNC, API_RAND}},
{"curandMakeMTGP32KernelState", {"hiprandMakeMTGP32KernelState", "", CONV_LIB_FUNC, API_RAND}},
// RAND Device functions
{"curand", {"hiprand", "", CONV_LIB_DEVICE_FUNC, API_RAND}},
hipify-clang/src/HipifyAction.cpp
+4 -2
Просмотреть файл
@@ -191,8 +191,10 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) {
case CONV_INCLUDE:
switch (hipToken.apiType) {
case API_RAND:
if (insertedRAND_kernelHeader) { return true; }
insertedRAND_kernelHeader = true;
if (hipToken.hipName == "hiprand_kernel.h") {
if (insertedRAND_kernelHeader) { return true; }
insertedRAND_kernelHeader = true;
}
return false;
default:
return false;
include/hip/hcc_detail/hip_atomic.h
+1 -1
Просмотреть файл
@@ -112,7 +112,7 @@ unsigned int atomicExch(unsigned int* address, unsigned int val)
}
__device__
inline
unsigned int atomicExch(unsigned long long* address, unsigned long long val)
unsigned long long atomicExch(unsigned long long* address, unsigned long long val)
{
return __atomic_exchange_n(address, val, __ATOMIC_RELAXED);
}
include/hip/hcc_detail/hip_vector_types.h
+1 -1
Просмотреть файл
@@ -34,7 +34,7 @@ THE SOFTWARE.
#include "hip/hcc_detail/host_defines.h"
#if !defined(_MSC_VER) || __HIP_DEVICE_COMPILE__
#if !defined(_MSC_VER) || __clang__
#if defined(__clang__)
#define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n)))
#elif defined(__GNUC__) // N.B.: GCC does not support .xyzw syntax.
include/hip/hcc_detail/math_functions.h
+7 -2
Просмотреть файл
@@ -899,6 +899,7 @@ __DEVICE__
inline
double nan(const char* tagp)
{
#if !_WIN32
union {
double val;
struct ieee_double {
@@ -906,8 +907,7 @@ double nan(const char* tagp)
uint32_t quiet : 1;
uint32_t exponent : 11;
uint32_t sign : 1;
} bits;
} bits;
static_assert(sizeof(double) == sizeof(ieee_double), "");
} tmp;
@@ -917,6 +917,11 @@ double nan(const char* tagp)
tmp.bits.mantissa = __make_mantissa(tagp);
return tmp.val;
#else
uint64_t val = __make_mantissa(tagp);
val |= 0xFFF << 51;
return reinterpret_cast<double>(val);
#endif
}
__DEVICE__
inline
include/hip/hcc_detail/texture_functions.h
+1 -1
Просмотреть файл
@@ -50,7 +50,7 @@ union TData {
#define __TEXTURE_FUNCTIONS_DECL__ static inline __device__
#if (__hcc_workweek__ >= 18114)
#if (__hcc_workweek__ >= 18114) || __clang__
#define ADDRESS_SPACE_CONSTANT __attribute__((address_space(4)))
#else
#define ADDRESS_SPACE_CONSTANT __attribute__((address_space(2)))
include/hip/hip_runtime_api.h
+5 -1
Просмотреть файл
@@ -318,9 +318,13 @@ enum hipComputeMode {
*
* Perform automatic type conversion to eliminate need for excessive typecasting (ie void**)
*
* __HIP_DISABLE_CPP_FUNCTIONS__ macro can be defined to suppress these
* wrappers. It is useful for applications which need to obtain decltypes of
* HIP runtime APIs.
*
* @see hipMalloc
*/
#ifdef __cplusplus
#if defined(__cplusplus) && !defined(__HIP_DISABLE_CPP_FUNCTIONS__)
template <class T>
static inline hipError_t hipMalloc(T** devPtr, size_t size) {
return hipMalloc((void**)devPtr, size);
packaging/hip_base.prerm
+8 -4
Просмотреть файл
@@ -10,14 +10,17 @@ popd () {
ROCMDIR=/opt/rocm
HIPDIR=$ROCMDIR/hip
([ ! -d $ROCMDIR ] || [ ! -d $HIPDIR ]) && exit 0
# Remove soft-link to bin files
HIPBINFILES=$HIPDIR/bin/*
HIPBINDIR=$HIPDIR/bin
ROCMBINDIR=$ROCMDIR/bin
([ ! -d $ROCMBINDIR ] || [ ! -d $HIPBINDIR ]) && exit 0
HIPBINFILES=$HIPBINDIR/*
pushd $ROCMBINDIR
for f in $HIPBINFILES
do
rm $(basename $f)
for f in $HIPBINFILES; do
[ -e $f ] || continue
rm $(basename $f)
done
rm .hipVersion
popd
@@ -26,6 +29,7 @@ rmdir --ignore-fail-on-non-empty $ROCMBINDIR
# Remove soft-link to headers
HIPINCDIR=$HIPDIR/include
ROCMINCDIR=$ROCMDIR/include
([ ! -d $ROCMINCDIR ] || [ ! -d $HIPINCDIR ]) && exit 0
pushd $ROCMINCDIR
rm hip
popd
packaging/hip_hcc.prerm
+5 -3
Просмотреть файл
@@ -12,13 +12,15 @@ ROCMDIR=/opt/rocm
ROCMLIBDIR=$ROCMDIR/lib
HIPDIR=$ROCMDIR/hip
HIPLIBDIR=$ROCMDIR/hip/lib
([ ! -d $ROCMDIR ] || [ ! -d $HIPDIR ]) && exit 0
([ ! -d $ROCMLIBDIR ] || [ ! -d $HIPLIBDIR ]) && exit 0
# Remove soft-links to libraries
HIPLIBFILES=$(ls -aF $HIPLIBDIR | grep -v [-/$])
pushd $ROCMLIBDIR
for f in $HIPLIBFILES
do
rm $(basename $f)
for f in $HIPLIBFILES; do
[ -e $f ] || continue
rm $(basename $f)
done
pushd cmake
unlink hip
tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp
+4 -6
Просмотреть файл
@@ -37,9 +37,9 @@
#include <curand.h>
// CHECK: #include <hiprand_kernel.h>
#include <curand_kernel.h>
// CHECK-NOT: #include <curand_mtgp32_host.h>
// CHECK-NOT: #include <curand_mtgp32dc_p_11213.h>
// CHECK: #include <hiprand_mtgp32_host.h>
#include <curand_mtgp32_host.h>
// CHECK: #include <rocrand_mtgp32_11213.h>
#include <curand_mtgp32dc_p_11213.h>
// CHECK: if ((x) != hipSuccess) {
@@ -201,11 +201,9 @@ struct runner<curandStateMtgp32_t>
CUDA_CALL(cudaMalloc((void **)&states, states_size * sizeof(curandStateMtgp32_t)));
// CHECK: CUDA_CALL(hipMalloc((void **)&d_param, sizeof(mtgp32_kernel_params)));
CUDA_CALL(cudaMalloc((void **)&d_param, sizeof(mtgp32_kernel_params)));
// curandMakeMTGP32Constants is yet unsupported by HIP
// CHECK-NOT: CURAND_CALL(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, d_param));
// CHECK: CURAND_CALL(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, d_param));
CURAND_CALL(curandMakeMTGP32Constants(mtgp32dc_params_fast_11213, d_param));
// curandMakeMTGP32KernelState is yet unsupported by HIP
// CHECK-NOT: CURAND_CALL(hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, d_param, states_size, seed));
// CHECK: CURAND_CALL(hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, d_param, states_size, seed));
CURAND_CALL(curandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, d_param, states_size, seed));
}