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

Merge branch 'amd-develop' into amd-master

Browse Source 
Change-Id: I07d5840e1213f28865a3130d45da061f7dee8349
This commit is contained in:
Maneesh Gupta
2016-07-13 11:27:23 +05:30
parent 41e80065ca b0dc9d208a
commit f0e4eae7f8
5 changed files with 104 additions and 50 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/CMakeLists.txt
+3 -3
View File
@@ -5,7 +5,7 @@ project(hip)
# Configure variables
#############################
# Determine HIP_VERSION
execute_process(COMMAND ${CMAKE_CURRENT_SOURCE_DIR}/bin/hipcc --short-version
execute_process(COMMAND ${CMAKE_CURRENT_SOURCE_DIR}/bin/hipconfig --version
OUTPUT_VARIABLE HIP_VERSION
OUTPUT_STRIP_TRAILING_WHITESPACE)
string(REPLACE "." ";" VERSION_LIST ${HIP_VERSION})
@@ -118,7 +118,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
set(CMAKE_C_FLAGS " -hc -I${HCC_HOME}/include -I${HSA_PATH}/include -I/opt/rocm/libhsakmt/include/libhsakmt -stdlib=libc++ -DHIP_HCC")
set(SOURCE_FILES src/device_util.cpp
src/hip_hcc.cpp
src/hip_hcc.cpp
src/hip_device.cpp
src/hip_error.cpp
src/hip_event.cpp
@@ -126,7 +126,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
src/hip_memory.cpp
src/hip_peer.cpp
src/hip_stream.cpp
src/hip_fp16.cpp
src/hip_fp16.cpp
src/staging_buffer.cpp)
if(${HIP_USE_SHARED_LIBRARY} EQUAL 1)
hipamd/bin/hipcc
+26 -31
View File
@@ -2,7 +2,7 @@
use File::Basename;
# HIP compiler driver
# HIP compiler driver
# Will call NVCC or HCC (depending on target) and pass the appropriate include and library options for
# the target compiler and HIP infrastructure.
@@ -14,7 +14,7 @@ use File::Basename;
# If HIP_PLATFORM is not set hipcc will attempt auto-detect based on if nvcc is found.
#
# Other environment variable controls:
# HIP_PATH : Path to HIP directory, default is one dir level above location of this script
# HIP_PATH : Path to HIP directory, default is one dir level above location of this script
# CUDA_PATH : Path to CUDA SDK (default /usr/local/cuda). Used on NVIDIA platforms only.
# HCC_HOME : Path to HCC SDK (default /opt/rocm/hcc). Used on AMD platforms only.
# HSA_PATH : Path to HSA dir (default /opt/rocm/hsa). Used on AMD platforms only.
@@ -24,10 +24,6 @@ print "No Arguments passed, exiting ...\n";
exit(-1);
}
$HIP_VERSION_MAJOR = "0";
$HIP_VERSION_MINOR = "90";
$HIP_VERSION_PATCH = "0";
$verbose = $ENV{'HIPCC_VERBOSE'};
$verbose = 0 unless defined $verbose;
# Verbose: 0x1=commands, 0x2=paths, 0x4=hippc args
@@ -35,9 +31,6 @@ $verbose = 0 unless defined $verbose;
$HIP_PATH=$ENV{'HIP_PATH'};
$HIP_PATH=dirname (dirname $0) unless defined $HIP_PATH; # use parent directory of hipcc
#print "HIP_PATH=$HIP_PATH\n";
$CODEXL_PATH = $ENV{'CODEXL_PATH'};
$CODEXL_PATH = "/opt/AMD/CodeXL" unless defined $CODEXL_PATH;
$marker_path = "$CODEXL_PATH/SDK/AMDTActivityLogger";
@@ -45,21 +38,21 @@ $marker_path = "$CODEXL_PATH/SDK/AMDTActivityLogger";
#---
#HIP_PLATFORM controls whether to use NVCC or HCC for compilation:
$HIP_PLATFORM= `$HIP_PATH/bin/hipconfig --platform`;
$HIP_VERSION= `$HIP_PATH/bin/hipconfig --version`;
$HIP_PLATFORM="hcc" unless defined $HIP_PLATFORM;
if ($verbose & 0x2) {
print ("HIP_PATH=$HIP_PATH\n");
print ("HIP_PLATFORM=$HIP_PLATFORM\n");
print ("HIP_PATH=$HIP_PATH\n");
print ("HIP_PLATFORM=$HIP_PLATFORM\n");
}
# set if user explicitly requests -stdlib=libc++. (else we default to libstdc++ for better interop with g++):
$setStdLib = 0; # TODO - set to 0
if ($HIP_PLATFORM eq "hcc") {
$HSA_PATH=$ENV{'HSA_PATH'};
$HSA_PATH="/opt/rocm/hsa" unless defined $HSA_PATH;
$HSA_PATH=$ENV{'HSA_PATH'};
$HSA_PATH="/opt/rocm/hsa" unless defined $HSA_PATH;
$HCC_HOME=$ENV{'HCC_HOME'};
$HCC_HOME="/opt/rocm/hcc" unless defined $HCC_HOME;
@@ -76,12 +69,14 @@ if ($HIP_PLATFORM eq "hcc") {
$HIPCXXFLAGS .= " -I$HIP_PATH/include/hcc_detail/cuda";
$HIPCXXFLAGS .= " -I$HSA_PATH/include";
$HIPLDFLAGS = "-hc -L$HCC_HOME/lib -Wl,--rpath=$HCC_HOME/lib -lc++ -ldl -lpthread -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive";
$HIPLDFLAGS = "-hc -L$HCC_HOME/lib -Wl,--rpath=$HCC_HOME/lib -lc++ -ldl -lpthread -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive";
# Suppress linker warnings in case HCC distribution contains OpenCL/SPIR symbols
$HIPLDFLAGS .= " -Wl,--defsym=_binary_kernel_spir_end=0 -Wl,--defsym=_binary_kernel_spir_start=0 -Wl,--defsym=_binary_kernel_cl_start=0 -Wl,--defsym=_binary_kernel_cl_end=0";
#$HIPLDFLAGS .= " -Wl,--defsym=_binary_kernel_spir_end=0 -Wl,--defsym=_binary_kernel_spir_start=0 -Wl,--defsym=_binary_kernel_cl_start=0 -Wl,--defsym=_binary_kernel_cl_end=0";
# Satisfy HCC dependencies
$HIPLDFLAGS .= " -lc++abi";
$HIPLDFLAGS .= " -L$HSA_PATH/lib -L$ROCM_PATH/lib -lhsa-runtime64 -lhc_am -lhsakmt";
$HIPLDFLAGS .= " -L$HSA_PATH/lib -L$ROCM_PATH/lib -lhsa-runtime64 -lhc_am -lhsakmt";
# Add trace marker library:
# TODO - once we cleanly separate the HIP API headers from HIP library headers this logic should move to CMakebuild option - apps do not need to see the marker library.
@@ -95,9 +90,9 @@ if ($HIP_PLATFORM eq "hcc") {
$HIPLDFLAGS .= " -L$marker_lib_path -lAMDTActivityLogger";
}
# Add C++ libs for GCC.
$HIPLDFLAGS .= " -lstdc++";
$HIPLDFLAGS .= " -lm";
# Add C++ libs for GCC.
$HIPLDFLAGS .= " -lstdc++";
$HIPLDFLAGS .= " -lm";
if ($verbose & 0x2) {
print ("HSA_PATH=$HSA_PATH\n");
@@ -114,10 +109,10 @@ if ($HIP_PLATFORM eq "hcc") {
$HIPCC="$CUDA_PATH/bin/nvcc";
$HIPCXXFLAGS .= " -I$CUDA_PATH/include";
$HIPLDFLAGS = "";
$HIPLDFLAGS = "";
} else {
printf ("error: unknown HIP_PLATFORM = '$HIP_PLATFORM'");
exit (-1);
printf ("error: unknown HIP_PLATFORM = '$HIP_PLATFORM'");
exit (-1);
}
@@ -137,11 +132,11 @@ my @options = ();
my @inputs = ();
if ($verbose & 0x4) {
print "hipcc-args: ", join (" ", @ARGV), "\n";
print "hipcc-args: ", join (" ", @ARGV), "\n";
}
my $toolArgs = ""; # arguments to pass to the hcc or nvcc tool
foreach $arg (@ARGV)
foreach $arg (@ARGV)
{
my $swallowArg = 0;
if ($arg eq '-c') {
@@ -173,7 +168,7 @@ foreach $arg (@ARGV)
if ($arg =~ m/^-/) {
# options start with -
# Process HIPCC options here:
if ($arg =~ m/^--hipcc/) {
$swallowArg = 1;
@@ -186,11 +181,11 @@ foreach $arg (@ARGV)
}
#print "O: <$arg>\n";
} else {
# input files and libraries
# input files and libraries
if (($arg =~ /\.cpp$/) or ($arg =~ /\.c$/) or ($arg =~ /\.cc$/) ) {
$hasC = 1;
$needCXXFLAGS = 1;
}
}
if (($arg =~ /\.cu$/) or ($arg =~ /\.cuh$/)) {
$hasCU = 1;
$needCXXFLAGS = 1;
@@ -228,7 +223,7 @@ if ($needHipHcc) {
# hipcc currrently requires separate compilation of source files, ie it is not possible to pass
# CPP files combined with .O files
# Reason is that NVCC uses the file extension to determine whether to compile in CUDA mode or
# Reason is that NVCC uses the file extension to determine whether to compile in CUDA mode or
# pass-through CPP mode.
@@ -236,7 +231,7 @@ if ($needHipHcc) {
my $CMD="$HIPCC";
if ($needCXXFLAGS) {
$CMD .= " $HIPCXXFLAGS";
}
}
if ($needLDFLAGS and not $compileOnly) {
$CMD .= " $HIPLDFLAGS";
}
@@ -250,7 +245,7 @@ if ($printHipVersion) {
if ($runCmd) {
print "HIP version: "
}
print $HIP_VERSION_MAJOR, ".", $HIP_VERSION_MINOR, ".", $HIP_VERSION_PATCH, "\n";
print $HIP_VERSION, "\n";
}
if ($runCmd) {
system ("$CMD") and die ();
hipamd/bin/hipconfig
+13 -1
View File
@@ -1,4 +1,9 @@
#!/usr/bin/perl -w
$HIP_VERSION_MAJOR = "0";
$HIP_VERSION_MINOR = "90";
$HIP_VERSION_PATCH = "0";
use Getopt::Long;
use Cwd;
@@ -15,6 +20,7 @@ GetOptions(
,"platform|P" => \$p_platform
,"cpp_config|cxx_config|C" => \$p_cpp_config
,"full|f|info" => \$p_full,
,"version|v" => \$p_version,
,"check" => \$p_check,
,"newline|n" => \$p_newline
);
@@ -26,12 +32,14 @@ if ($p_help) {
print " --compiler, -c : print compiler (hcc or nvcc)\n";
print " --platform, -P : print platform (hcc or nvcc)\n";
print " --full, -f : print full config\n";
print " --version, -v : print hip version\n";
print " --check : check configuration\n";
print " --newline, -n : print newline\n";
print " --help, -h : print help message\n";
exit();
}
$HIP_VERSION="$HIP_VERSION_MAJOR.$HIP_VERSION_MINOR.$HIP_VERSION_PATCH";
$CUDA_PATH=$ENV{'CUDA_PATH'};
$CUDA_PATH='/usr/local/cuda' unless defined $CUDA_PATH;
@@ -83,9 +91,13 @@ if ($p_compiler or $p_platform) {
$printed = 1;
}
if ($p_version) {
print $HIP_VERSION;
$printed = 1;
}
if (!$printed or $p_full) {
print "HIP version : ", $HIP_VERSION, "\n\n";
print "== hipconfig\n";
print "HIP_PATH : ", $HIP_PATH, "\n";
print "HIP_PLATFORM : ", $HIP_PLATFORM, "\n";
hipamd/clang-hipify/src/Cuda2Hip.cpp
+62 -11
View File
@@ -95,7 +95,8 @@ struct cuda2hipMap {
cuda2hipRename["cuda_runtime_api.h"] = {"hip_runtime_api.h", CONV_INCLUDE};
// HIP includes
cuda2hipRename["cudacommon.h.prehip"] = {"cudacommon.h", CONV_INCLUDE};
// TODO: uncomment this when hip/cudacommon.h will be renamed to hip/hipcommon.h
//cuda2hipRename["cudacommon.h"] = {"hipcommon.h", CONV_INCLUDE};
// CUBLAS includes
cuda2hipRename["cublas.h"] = {"hipblas.h", CONV_INCLUDE};
@@ -134,23 +135,44 @@ struct cuda2hipMap {
// Memcpy
cuda2hipRename["cudaMemcpy"] = {"hipMemcpy", CONV_MEM};
cuda2hipRename["cudaMemcpyHostToHost"] = {"hipMemcpyHostToHost", CONV_MEM};
cuda2hipRename["cudaMemcpyHostToDevice"] = {"hipMemcpyHostToDevice", CONV_MEM};
cuda2hipRename["cudaMemcpyDeviceToHost"] = {"hipMemcpyDeviceToHost", CONV_MEM};
cuda2hipRename["cudaMemcpyDeviceToDevice"] = {"hipMemcpyDeviceToDevice", CONV_MEM};
cuda2hipRename["cudaMemcpyDefault"] = {"hipMemcpyDefault", CONV_MEM};
cuda2hipRename["cudaMemcpyToSymbol"] = {"hipMemcpyToSymbol", CONV_MEM};
cuda2hipRename["cudaMemset"] = {"hipMemset", CONV_MEM};
cuda2hipRename["cudaMemsetAsync"] = {"hipMemsetAsync", CONV_MEM};
cuda2hipRename["cudaMemcpyAsync"] = {"hipMemcpyAsync", CONV_MEM};
cuda2hipRename["cudaMemGetInfo"] = {"hipMemGetInfo", CONV_MEM};
// Memcpy kind
cuda2hipRename["cudaMemcpyKind"] = {"hipMemcpyKind", CONV_MEM};
cuda2hipRename["cudaMemcpyHostToHost"] = {"hipMemcpyHostToHost", CONV_MEM};
cuda2hipRename["cudaMemcpyHostToDevice"] = {"hipMemcpyHostToDevice", CONV_MEM};
cuda2hipRename["cudaMemcpyDeviceToHost"] = {"hipMemcpyDeviceToHost", CONV_MEM};
cuda2hipRename["cudaMemcpyDeviceToDevice"] = {"hipMemcpyDeviceToDevice", CONV_MEM};
cuda2hipRename["cudaMemcpyDefault"] = {"hipMemcpyDefault", CONV_MEM};
// Memory management
cuda2hipRename["cudaMalloc"] = {"hipMalloc", CONV_MEM};
cuda2hipRename["cudaMallocHost"] = {"hipHostAlloc", CONV_MEM};
cuda2hipRename["cudaFree"] = {"hipFree", CONV_MEM};
cuda2hipRename["cudaFreeHost"] = {"hipHostFree", CONV_MEM};
cuda2hipRename["cudaMalloc"] = {"hipMalloc", CONV_MEM};
cuda2hipRename["cudaMallocHost"] = {"hipHostMalloc", CONV_MEM};
cuda2hipRename["cudaFree"] = {"hipFree", CONV_MEM};
cuda2hipRename["cudaFreeHost"] = {"hipHostFree", CONV_MEM};
cuda2hipRename["cudaHostRegister"] = {"hipHostRegister", CONV_MEM};
cuda2hipRename["cudaHostUnregister"] = {"hipHostUnregister", CONV_MEM};
// Memory types
cuda2hipRename["cudaMemoryType"] = {"hipMemoryType", CONV_MEM};
cuda2hipRename["cudaMemoryTypeHost"] = {"hipMemoryTypeHost", CONV_MEM};
cuda2hipRename["cudaMemoryTypeDevice"] = {"hipMemoryTypeDevice", CONV_MEM};
// Host Malloc Flags
cuda2hipRename["cudaHostAllocDefault"] = {"hipHostMallocDefault", CONV_MEM};
cuda2hipRename["cudaHostAllocPortable"] = {"hipHostMallocPortable", CONV_MEM};
cuda2hipRename["cudaHostAllocMapped"] = {"hipHostMallocMapped", CONV_MEM};
cuda2hipRename["cudaHostAllocWriteCombined"] = {"hipHostMallocWriteCombined", CONV_MEM};
// Host Register Flags
cuda2hipRename["cudaHostGetFlags"] = {"hipHostGetFlags", CONV_MEM};
cuda2hipRename["cudaHostRegisterDefault"] = {"hipHostRegisterDefault", CONV_MEM};
cuda2hipRename["cudaHostRegisterPortable"] = {"hipHostRegisterPortable", CONV_MEM};
cuda2hipRename["cudaHostRegisterMapped"] = {"hipHostRegisterMapped", CONV_MEM};
cuda2hipRename["cudaHostRegisterIoMemory"] = {"hipHostRegisterIoMemory", CONV_MEM};
// Coordinate Indexing and Dimensions
cuda2hipRename["threadIdx.x"] = {"hipThreadIdx_x", CONV_COORD_FUNC};
@@ -191,14 +213,22 @@ struct cuda2hipMap {
cuda2hipRename["cudaEventRecord"] = {"hipEventRecord", CONV_EVENT};
cuda2hipRename["cudaEventElapsedTime"] = {"hipEventElapsedTime", CONV_EVENT};
cuda2hipRename["cudaEventSynchronize"] = {"hipEventSynchronize", CONV_EVENT};
cuda2hipRename["cudaEventQuery"] = {"hipEventQuery", CONV_EVENT};
// Event Flags
cuda2hipRename["cudaEventDefault"] = {"hipEventDefault", CONV_EVENT};
cuda2hipRename["cudaEventBlockingSync"] = {"hipEventBlockingSync", CONV_EVENT};
cuda2hipRename["cudaEventDisableTiming"] = {"hipEventDisableTiming", CONV_EVENT};
cuda2hipRename["cudaEventInterprocess"] = {"hipEventInterprocess", CONV_EVENT};
// Streams
cuda2hipRename["cudaStream_t"] = {"hipStream_t", CONV_STREAM};
cuda2hipRename["cudaStreamCreate"] = {"hipStreamCreate", CONV_STREAM};
cuda2hipRename["cudaStreamCreateWithFlags"] = {"hipStreamCreateWithFlags", CONV_STREAM};
cuda2hipRename["cudaStreamDestroy"] = {"hipStreamDestroy", CONV_STREAM};
cuda2hipRename["cudaStreamWaitEvent"] = {"hipStreamWaitEven", CONV_STREAM};
cuda2hipRename["cudaStreamWaitEvent"] = {"hipStreamWaitEvent", CONV_STREAM};
cuda2hipRename["cudaStreamSynchronize"] = {"hipStreamSynchronize", CONV_STREAM};
// Stream Flags
cuda2hipRename["cudaStreamGetFlags"] = {"hipStreamGetFlags", CONV_STREAM};
cuda2hipRename["cudaStreamDefault"] = {"hipStreamDefault", CONV_STREAM};
cuda2hipRename["cudaStreamNonBlocking"] = {"hipStreamNonBlocking", CONV_STREAM};
@@ -241,10 +271,31 @@ struct cuda2hipMap {
cuda2hipRename["cudaDevAttrMaxSharedMemoryPerMultiprocessor"] = {"hipDeviceAttributeMaxSharedMemoryPerMultiprocessor", CONV_DEV};
cuda2hipRename["cudaDevAttrIsMultiGpuBoard"] = {"hipDeviceAttributeIsMultiGpuBoard", CONV_DEV};
// Pointer Attributes
cuda2hipRename["cudaPointerAttributes"] = {"hipPointerAttribute_t", CONV_MEM};
cuda2hipRename["cudaPointerGetAttributes"] = {"hipPointerGetAttributes", CONV_MEM};
cuda2hipRename["cudaHostGetDevicePointer"] = {"hipHostGetDevicePointer", CONV_MEM};
// Device
cuda2hipRename["cudaDeviceProp"] = {"hipDeviceProp_t", CONV_DEV};
cuda2hipRename["cudaGetDeviceProperties"] = {"hipGetDeviceProperties", CONV_DEV};
// Device Flags
cuda2hipRename["cudaSetDeviceFlags"] = {"hipSetDeviceFlags", CONV_DEV};
cuda2hipRename["cudaDeviceScheduleAuto"] = {"hipDeviceScheduleAuto", CONV_DEV};
cuda2hipRename["cudaDeviceScheduleSpin"] = {"hipDeviceScheduleSpin", CONV_DEV};
cuda2hipRename["cudaDeviceScheduleYield"] = {"hipDeviceScheduleYield", CONV_DEV};
// deprecated as of CUDA 4.0 and replaced with cudaDeviceScheduleBlockingSync
cuda2hipRename["cudaDeviceBlockingSync"] = {"hipDeviceBlockingSync", CONV_DEV};
// unsupported yet
//cuda2hipRename["cudaDeviceScheduleBlockingSync"] = {"hipDeviceScheduleBlockingSync", CONV_DEV};
//cuda2hipRename["cudaDeviceScheduleMask"] = {"hipDeviceScheduleMask", CONV_DEV};
cuda2hipRename["cudaDeviceMapHost"] = {"hipDeviceMapHost", CONV_DEV};
// unsupported yet
//cuda2hipRename["cudaDeviceLmemResizeToMax"] = {"hipDeviceLmemResizeToMax", CONV_DEV};
//cuda2hipRename["cudaDeviceMask"] = {"hipDeviceMask", CONV_DEV};
// Cache config
cuda2hipRename["cudaDeviceSetCacheConfig"] = {"hipDeviceSetCacheConfig", CONV_DEV};
// translate deprecated
hipamd/packaging/hip_hcc.txt
-4
View File
@@ -1,10 +1,6 @@
cmake_minimum_required(VERSION 2.8.3)
project(hip_hcc)
#execute_process(COMMAND @CMAKE_CURRENT_SOURCE_DIR@/bin/hipconfig --full 2>/dev/null | grep clang | cut -d" " -f9
# OUTPUT_VARIABLE HCC_VERSION
# OUTPUT_STRIP_TRAILING_WHITESPACE)
install(DIRECTORY @PROJECT_BINARY_DIR@/CMakeFiles/hip_hcc.dir/src/ DESTINATION lib)
#############################