diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index 7f80325779..d99907d695 100644 --- a/projects/clr/hipamd/CMakeLists.txt +++ b/projects/clr/hipamd/CMakeLists.txt @@ -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) diff --git a/projects/clr/hipamd/bin/hipcc b/projects/clr/hipamd/bin/hipcc index 2c95a21a99..33041d240f 100755 --- a/projects/clr/hipamd/bin/hipcc +++ b/projects/clr/hipamd/bin/hipcc @@ -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 (); diff --git a/projects/clr/hipamd/bin/hipconfig b/projects/clr/hipamd/bin/hipconfig index 49c385753d..4f4559cb8e 100755 --- a/projects/clr/hipamd/bin/hipconfig +++ b/projects/clr/hipamd/bin/hipconfig @@ -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"; diff --git a/projects/clr/hipamd/clang-hipify/src/Cuda2Hip.cpp b/projects/clr/hipamd/clang-hipify/src/Cuda2Hip.cpp index a258825091..02edeef95b 100644 --- a/projects/clr/hipamd/clang-hipify/src/Cuda2Hip.cpp +++ b/projects/clr/hipamd/clang-hipify/src/Cuda2Hip.cpp @@ -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 diff --git a/projects/clr/hipamd/packaging/hip_hcc.txt b/projects/clr/hipamd/packaging/hip_hcc.txt index 63637b1e0a..556bade072 100644 --- a/projects/clr/hipamd/packaging/hip_hcc.txt +++ b/projects/clr/hipamd/packaging/hip_hcc.txt @@ -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) #############################