hipcc uses -stdlib=stdlibc++ by default.
hipcc switches to -stdlib=libc++ if user requests it. hack to disable LDG.
This commit is contained in:
+15
-5
@@ -55,7 +55,8 @@ if ($verbose & 0x2) {
|
||||
print ("CUDA_PATH=$CUDA_PATH\n");
|
||||
}
|
||||
|
||||
$enablestdcpplib = 0;
|
||||
# set if user explicitly requests -stdlib=libc++. (else we default to libstdc++ for better interop with g++):
|
||||
$enablelibcpp = 0; # TODO - set to 0
|
||||
|
||||
if ($HIP_PLATFORM eq "hcc") {
|
||||
$HSA_PATH=$ENV{'HSA_PATH'};
|
||||
@@ -73,6 +74,10 @@ if ($HIP_PLATFORM eq "hcc") {
|
||||
|
||||
$HIPCC=$HCC;
|
||||
$HIPCXXFLAGS = $HCCFLAGS;
|
||||
|
||||
# To interoperate cleanly with G++, we default to using libc++ library:
|
||||
$HIPCXXFLAGS .= " -stdlib=libstdc++";
|
||||
|
||||
$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";
|
||||
@@ -146,10 +151,10 @@ foreach $arg (@ARGV)
|
||||
$needLDFLAGS = 1;
|
||||
}
|
||||
|
||||
if($arg eq '-stdlib=libstdc++' and $enablestdcpplib eq 0)
|
||||
if($arg eq '-stdlib=libc++' and $enablelibcpp eq 0)
|
||||
{
|
||||
$HIPCXXFLAGS .= " -stdlib=libstdc++";
|
||||
$enablestdcpplib = 1;
|
||||
$HIPCXXFLAGS .= " -stdlib=libc++";
|
||||
$enablelibcpp = 1;
|
||||
}
|
||||
|
||||
if ($arg =~ m/^-/) {
|
||||
@@ -190,8 +195,13 @@ if ($hasCU and $HIP_PLATFORM eq 'hcc') {
|
||||
$HIPCXXFLAGS .= " -x c++";
|
||||
}
|
||||
|
||||
if ($enablestdcpplib eq 0 and $HIP_PLATFORM eq 'hcc')
|
||||
if ($enablelibcpp eq 0 and $HIP_PLATFORM eq 'hcc')
|
||||
{
|
||||
$HIPCXXFLAGS .= " -stdlib=libstdc++";
|
||||
}
|
||||
# TODO -remove me.
|
||||
if ($enablelibcpp eq 2 and $HIP_PLATFORM eq 'hcc')
|
||||
{
|
||||
$HIPCXXFLAGS .= " -stdlib=libc++";
|
||||
}
|
||||
|
||||
|
||||
@@ -21,9 +21,10 @@ THE SOFTWARE.
|
||||
#define HIP_LDG_H
|
||||
|
||||
#if __HCC__
|
||||
#if defined (USE_LDG)
|
||||
#include"hip_vector_types.h"
|
||||
#include"host_defines.h"
|
||||
#if __hcc_workweek__ >= 16164
|
||||
#if __hcc_workweek__ >= 16164
|
||||
#include"hip/hip_vector_types.h"
|
||||
#include"hip/hcc_detail/host_defines.h"
|
||||
|
||||
@@ -101,5 +102,7 @@ __device__ double4 __ldg(const double4* );
|
||||
|
||||
#endif // __HCC__
|
||||
|
||||
#endif
|
||||
|
||||
#endif // HIP_LDG_H
|
||||
|
||||
|
||||
@@ -45,11 +45,9 @@ THE SOFTWARE.
|
||||
//---
|
||||
// Remainder of this file only compiles with HCC
|
||||
#ifdef __HCC__
|
||||
//#if __cplusplus
|
||||
//#include <hc.hpp>
|
||||
//#endif
|
||||
#include <grid_launch.h>
|
||||
#include <kalmar_defines.h>
|
||||
|
||||
extern int HIP_TRACE_API;
|
||||
|
||||
//TODO-HCC-GL - change this to typedef.
|
||||
@@ -112,6 +110,7 @@ extern int HIP_TRACE_API;
|
||||
#define __HCC_C__
|
||||
#endif
|
||||
|
||||
#ifdef DISABLED_MATH_FUNC // TODO - remove this flag.
|
||||
__device__ float acosf(float x);
|
||||
__device__ float acoshf(float x);
|
||||
__device__ float asinf(float x);
|
||||
@@ -445,7 +444,7 @@ __device__ float __dsqrt_rd(double x);
|
||||
__device__ float __dsqrt_rn(double x);
|
||||
__device__ float __dsqrt_ru(double x);
|
||||
__device__ float __dsqrt_rz(double x);
|
||||
#endif
|
||||
#endif /* DISABLED_MATH_FUNC*/
|
||||
/**
|
||||
* Kernel launching
|
||||
*/
|
||||
@@ -575,3 +574,4 @@ do {\
|
||||
|
||||
|
||||
#endif
|
||||
#endif
|
||||
|
||||
@@ -17,8 +17,10 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hc.hpp>
|
||||
#include"hcc_detail/hip_ldg.h"
|
||||
#if __hcc_workweek__ >= 16164
|
||||
|
||||
#if __hcc_workweek__ >= 16164 and defined (USE_LDG)
|
||||
__device__ char __ldg(const char* ptr)
|
||||
{
|
||||
return ptr[0];
|
||||
|
||||
Reference in New Issue
Block a user