From 812f5b3a9730605841a4a4cad258f58dfbbaef18 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Sat, 2 Apr 2016 22:32:05 +0530 Subject: [PATCH 1/8] Add c++abi to linker options for HCC [ROCm/hip commit: 8e3465f6aed7932089de54c97ac793b2acfda6b5] --- projects/hip/bin/hipcc | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 9b10f551f4..289574d8e5 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -74,8 +74,10 @@ 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"; - # 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"; + # 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"; + # Satisfy HCC dependencies + $HIPLDFLAGS .= " -lc++abi"; $HIPLDFLAGS .= " -L$HSA_PATH/lib -lhsa-runtime64 -lhc_am"; # Add trace marker library: @@ -187,7 +189,7 @@ if ($needHipHcc) { #$HIPLDFLAGS .= " -L/opt/hip/lib -lhip_hcc" ; if ($HIP_USE_SHARED_LIBRARY) { - $HIPLDFLAGS .= " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc" + $HIPLDFLAGS .= " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc"; } else { $HIPLDFLAGS .= " $HIP_PATH/lib/device_util.cpp.o $HIP_PATH/lib/hip_device.cpp.o $HIP_PATH/lib/hip_error.cpp.o $HIP_PATH/lib/hip_event.cpp.o $HIP_PATH/lib/hip_hcc.cpp.o $HIP_PATH/lib/hip_memory.cpp.o $HIP_PATH/lib/hip_peer.cpp.o $HIP_PATH/lib/hip_stream.cpp.o $HIP_PATH/lib/staging_buffer.cpp.o"; } From efdb946fba6244b8d3b99af1d5c4fda7b8d3d164 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 4 Apr 2016 13:53:24 +0530 Subject: [PATCH 2/8] Enable directed tests to build HIP locally via cmake option [ROCm/hip commit: 033269c5135159379abd30a5858936f3fb2038bd] --- projects/hip/tests/src/CMakeLists.txt | 15 ++++++++++++--- 1 file changed, 12 insertions(+), 3 deletions(-) diff --git a/projects/hip/tests/src/CMakeLists.txt b/projects/hip/tests/src/CMakeLists.txt index 8d5dae6328..a0413eaad8 100644 --- a/projects/hip/tests/src/CMakeLists.txt +++ b/projects/hip/tests/src/CMakeLists.txt @@ -8,7 +8,13 @@ include_directories( ${PROJECT_SOURCE_DIR}/include ) set (HIP_Unit_Test_VERSION_MAJOR 1) set (HIP_Unit_Test_VERSION_MINOR 0) -set (HIP_BUILD_LOCAL 0) +if(NOT DEFINED HIP_BUILD_LOCAL) + if(NOT DEFINED ENV{HIP_BUILD_LOCAL}) + set(HIP_BUILD_LOCAL 0 CACHE BOOL "Build HIP in local folder") + else() + set(HIP_BUILD_LOCAL $ENV{HIP_BUILD_LOCAL} CACHE BOOL "Build HIP in local folder") + endif() +endif() set(HIP_PATH $ENV{HIP_PATH}) if (NOT DEFINED HIP_PATH) @@ -75,9 +81,12 @@ add_library(test_common OBJECT test_common.cpp ) # usage : make_hip_executable (exe_name CPP_FILES) macro (make_hip_executable exe cpp) if (${HIP_PLATFORM} STREQUAL "hcc") - add_executable (${exe} ${cpp} ${ARGN} $ ) + #add_executable (${exe} ${cpp} ${ARGN} $ ) if (${HIP_BUILD_LOCAL}) - target_link_libraries(${exe} hip_hcc) + #target_link_libraries(${exe} hip_hcc) + add_executable (${exe} ${cpp} ${ARGN} $ $ ) + else() + add_executable (${exe} ${cpp} ${ARGN} $ ) endif() else() add_executable (${exe} ${cpp} ${ARGN} $ ) From 5f1f0b34d06fbe3bb55218e5257d7f262a5df129 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 4 Apr 2016 14:47:02 +0530 Subject: [PATCH 3/8] Remove deprecated KERNELBEGIN and KERNELEND from bit_extract sample [ROCm/hip commit: 70f8236ac5a487525ac9820ba37eaaaee39cc2f7] --- projects/hip/samples/0_Intro/bit_extract/bit_extract.cpp | 4 ---- 1 file changed, 4 deletions(-) diff --git a/projects/hip/samples/0_Intro/bit_extract/bit_extract.cpp b/projects/hip/samples/0_Intro/bit_extract/bit_extract.cpp index 14b5be66d4..5545a99c0f 100644 --- a/projects/hip/samples/0_Intro/bit_extract/bit_extract.cpp +++ b/projects/hip/samples/0_Intro/bit_extract/bit_extract.cpp @@ -33,8 +33,6 @@ THE SOFTWARE. void __global__ bit_extract_kernel(hipLaunchParm lp, uint32_t *C_d, const uint32_t *A_d, size_t N) { - KERNELBEGIN; - size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); size_t stride = hipBlockDim_x * hipGridDim_x ; @@ -45,8 +43,6 @@ bit_extract_kernel(hipLaunchParm lp, uint32_t *C_d, const uint32_t *A_d, size_t C_d[i] = ((A_d[i] & 0xf00) >> 8); #endif } - - KERNELEND; } From cf9cefcb6ae3157962fba25a4861d6129553708f Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Wed, 6 Apr 2016 14:18:51 -0500 Subject: [PATCH 4/8] Added feature for cross-compilation and seperated hcc api from hip api [ROCm/hip commit: 1dcdfc93f135a6fe9570f80cd909f76881bed0f5] --- projects/hip/bin/hipcc | 10 ++++++++-- projects/hip/include/hcc.h | 8 ++++++++ projects/hip/include/hcc_detail/hcc_acc.h | 20 +++++++++++++++++++ .../hip/include/hcc_detail/hip_runtime_api.h | 14 ------------- projects/hip/tests/src/hipHcc.cpp | 2 +- 5 files changed, 37 insertions(+), 17 deletions(-) create mode 100644 projects/hip/include/hcc.h create mode 100644 projects/hip/include/hcc_detail/hcc_acc.h diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 289574d8e5..e4d203f685 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -40,7 +40,6 @@ $CODEXL_PATH = $ENV{'CODEXL_PATH'}; $CODEXL_PATH = "/opt/AMD/CodeXL" unless defined $CODEXL_PATH; $marker_path = "$CODEXL_PATH/SDK/AMDTActivityLogger"; - #--- #HIP_PLATFORM controls whether to use NVCC or HCC for compilation: $HIP_PLATFORM=$ENV{'HIP_PLATFORM'}; @@ -67,7 +66,7 @@ if ($HIP_PLATFORM eq "hcc") { # HCC* may be used to compile src/hip_hcc.o (and also feed the HIPCXXFLAGS below) $HCC = "$HCC_HOME/bin/hcc"; - $HCCFLAGS = " -hc -I$HCC_HOME/include -stdlib=libc++"; + $HCCFLAGS = " -hc -I$HCC_HOME/include "; $HIPCC=$HCC; $HIPCXXFLAGS = $HCCFLAGS; @@ -144,6 +143,13 @@ foreach $arg (@ARGV) $needLDFLAGS = 1; } + if($arg eq '-stdlib=libstdc++') + { + $HIPCXXFLAGS .= " -stdlib=libstdc++"; + }else + { + $HIPCXXFLAGS .= " -stdlib=libc++"; + } if ($arg =~ m/^-/) { # options start with - diff --git a/projects/hip/include/hcc.h b/projects/hip/include/hcc.h new file mode 100644 index 0000000000..9766e95935 --- /dev/null +++ b/projects/hip/include/hcc.h @@ -0,0 +1,8 @@ +#ifndef HCC_H +#define HCC_H + +#if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__) +#include +#endif + +#endif diff --git a/projects/hip/include/hcc_detail/hcc_acc.h b/projects/hip/include/hcc_detail/hcc_acc.h new file mode 100644 index 0000000000..c66b340579 --- /dev/null +++ b/projects/hip/include/hcc_detail/hcc_acc.h @@ -0,0 +1,20 @@ +#ifndef HCC_ACC_H +#define HCC_ACC_H +#include "hip_runtime_api.h" + +#if __cplusplus +#ifdef __HCC__ +#include +/** + * @brief Return hc::accelerator associated with the specified deviceId + */ +hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc); + +/** + * @brief Return hc::accelerator_view associated with the specified stream + */ +hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av); +#endif +#endif + +#endif diff --git a/projects/hip/include/hcc_detail/hip_runtime_api.h b/projects/hip/include/hcc_detail/hip_runtime_api.h index eecab9abc3..b5fdb312a4 100644 --- a/projects/hip/include/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hcc_detail/hip_runtime_api.h @@ -1041,20 +1041,6 @@ hipError_t hipDriverGetVersion(int *driverVersion) ; * @endcode * */ -#if __cplusplus -#ifdef __HCC__ -#include -/** - * @brief Return hc::accelerator associated with the specified deviceId - */ -hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc); - -/** - * @brief Return hc::accelerator_view associated with the specified stream - */ -hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **av); -#endif -#endif // end-group HCC_Specific /** diff --git a/projects/hip/tests/src/hipHcc.cpp b/projects/hip/tests/src/hipHcc.cpp index b3580db660..e3740fa27b 100644 --- a/projects/hip/tests/src/hipHcc.cpp +++ b/projects/hip/tests/src/hipHcc.cpp @@ -24,7 +24,7 @@ THE SOFTWARE. #include #include #include - +#include #include "test_common.h" #define CHECK(error) \ From 71ef0b9802c5218b28364311c8ea3934105176dd Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Wed, 6 Apr 2016 14:50:27 -0500 Subject: [PATCH 5/8] Corrected corner cases [ROCm/hip commit: 3304092b4958000d915d96f4cb61bf381028cd23] --- projects/hip/bin/hipcc | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index e4d203f685..14a6cfe2f2 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -55,7 +55,7 @@ if ($verbose & 0x2) { print ("CUDA_PATH=$CUDA_PATH\n"); } - +$enablestdcpplib = 0; if ($HIP_PLATFORM eq "hcc") { $HSA_PATH=$ENV{'HSA_PATH'}; @@ -143,12 +143,10 @@ foreach $arg (@ARGV) $needLDFLAGS = 1; } - if($arg eq '-stdlib=libstdc++') + if($arg eq '-stdlib=libstdc++' and $enablestdcpplib eq 0) { $HIPCXXFLAGS .= " -stdlib=libstdc++"; - }else - { - $HIPCXXFLAGS .= " -stdlib=libc++"; + $enablestdcpplib = 1; } if ($arg =~ m/^-/) { @@ -189,6 +187,11 @@ if ($hasCU and $HIP_PLATFORM eq 'hcc') { $HIPCXXFLAGS .= " -x c++"; } +if ($enablestdcpplib eq 0) +{ + $HIPCXXFLAGS .= " -stdlib=libc++"; +} + if ($needHipHcc) { $HIP_USE_SHARED_LIBRARY = $ENV{'HIP_USE_SHARED_LIBRARY'}; $HIP_USE_SHARED_LIBRARY = 0 unless defined $HIP_USE_SHARED_LIBRARY; From d06cd35623a8d097ba58ae4338fcb7cbe96d8c63 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 7 Apr 2016 10:48:16 +0530 Subject: [PATCH 6/8] Fix broken CUDA path in hipcc [ROCm/hip commit: 1e6b79ea6482910fb0f0140d8f73f25dfe9ff799] --- projects/hip/bin/hipcc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 14a6cfe2f2..fef4175e66 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -187,7 +187,7 @@ if ($hasCU and $HIP_PLATFORM eq 'hcc') { $HIPCXXFLAGS .= " -x c++"; } -if ($enablestdcpplib eq 0) +if ($enablestdcpplib eq 0 and $HIP_PLATFORM eq 'hcc') { $HIPCXXFLAGS .= " -stdlib=libc++"; } From fc289181005e6280e7c5a030124e540a48bb5b60 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 7 Apr 2016 04:57:31 -0500 Subject: [PATCH 7/8] C proofing hip [ROCm/hip commit: 91d1974204ee55e9df9baa3d6886424a9cdd8eee] --- projects/hip/include/hcc_detail/hip_vector_types.h | 7 ++++++- projects/hip/include/hip_vector_types.h | 2 ++ 2 files changed, 8 insertions(+), 1 deletion(-) diff --git a/projects/hip/include/hcc_detail/hip_vector_types.h b/projects/hip/include/hcc_detail/hip_vector_types.h index 29d2d53a8d..7d3ed98431 100644 --- a/projects/hip/include/hcc_detail/hip_vector_types.h +++ b/projects/hip/include/hcc_detail/hip_vector_types.h @@ -29,8 +29,12 @@ THE SOFTWARE. #error("This version of HIP requires a newer version of HCC."); #endif +#if __cplusplus #include +using namespace hc::short_vector; +#endif + //-- Signed // Define char vector types typedef hc::short_vector::char1 char1; @@ -108,7 +112,7 @@ typedef hc::short_vector::double2 double2; typedef hc::short_vector::double3 double3; typedef hc::short_vector::double4 double4; - +/* ///--- // Inline functions for creating vector types from basic types #define ONE_COMPONENT_ACCESS(T, VT) inline VT make_ ##VT (T x) { VT t; t.x = x; return t; }; @@ -191,3 +195,4 @@ ONE_COMPONENT_ACCESS (double, double1); TWO_COMPONENT_ACCESS (double, double2); THREE_COMPONENT_ACCESS(double, double3); FOUR_COMPONENT_ACCESS (double, double4); +*/ diff --git a/projects/hip/include/hip_vector_types.h b/projects/hip/include/hip_vector_types.h index 61d7f6c7d9..64702b8655 100644 --- a/projects/hip/include/hip_vector_types.h +++ b/projects/hip/include/hip_vector_types.h @@ -27,7 +27,9 @@ THE SOFTWARE. #if defined(__HIP_PLATFORM_HCC__) && !defined (__HIP_PLATFORM_NVCC__) +#if __cplusplus #include +#endif #elif defined(__HIP_PLATFORM_NVCC__) && !defined (__HIP_PLATFORM_HCC__) #include #else From bdef1bb4e8edba9e7527688693360f641714a66c Mon Sep 17 00:00:00 2001 From: pensun Date: Thu, 7 Apr 2016 09:46:00 -0500 Subject: [PATCH 8/8] clean up unused comments [ROCm/hip commit: 4b2c5976cea7bfcde4d4803da70beff66f0579af] --- projects/hip/src/hip_hcc.cpp | 10 +--------- 1 file changed, 1 insertion(+), 9 deletions(-) diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 6c80b84e0e..3c087467d7 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -137,7 +137,7 @@ ihipStream_t::ihipStream_t(unsigned device_index, hc::accelerator_view av, unsig //--- ihipStream_t::~ihipStream_t() -{ +{ } @@ -899,7 +899,6 @@ void ihipInit() for (int i=0; i