diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 9b10f551f4..fef4175e66 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'}; @@ -56,7 +55,7 @@ if ($verbose & 0x2) { print ("CUDA_PATH=$CUDA_PATH\n"); } - +$enablestdcpplib = 0; if ($HIP_PLATFORM eq "hcc") { $HSA_PATH=$ENV{'HSA_PATH'}; @@ -67,15 +66,17 @@ 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; $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: @@ -142,6 +143,11 @@ foreach $arg (@ARGV) $needLDFLAGS = 1; } + if($arg eq '-stdlib=libstdc++' and $enablestdcpplib eq 0) + { + $HIPCXXFLAGS .= " -stdlib=libstdc++"; + $enablestdcpplib = 1; + } if ($arg =~ m/^-/) { # options start with - @@ -181,13 +187,18 @@ if ($hasCU and $HIP_PLATFORM eq 'hcc') { $HIPCXXFLAGS .= " -x c++"; } +if ($enablestdcpplib eq 0 and $HIP_PLATFORM eq 'hcc') +{ + $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; #$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"; } 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/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 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; } 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 ) + #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} $ ) 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) \