From 06e54694c590457a544eb7fed8a8b88113bd9e7f Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Thu, 28 Apr 2016 13:17:49 -0500
Subject: [PATCH] hipcc uses -stdlib=stdlibc++ by default.

hipcc switches to -stdlib=libc++ if user requests it.
hack to disable LDG.
---
 hipamd/bin/hipcc                        | 20 +++++++++++++++-----
 hipamd/include/hcc_detail/hip_ldg.h     |  5 ++++-
 hipamd/include/hcc_detail/hip_runtime.h |  8 ++++----
 hipamd/src/hip_ldg.cpp                  |  4 +++-
 4 files changed, 26 insertions(+), 11 deletions(-)

diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc
index f09e673e7c..1cf43ecd78 100755
--- a/hipamd/bin/hipcc
+++ b/hipamd/bin/hipcc
@@ -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++";
 }
 
diff --git a/hipamd/include/hcc_detail/hip_ldg.h b/hipamd/include/hcc_detail/hip_ldg.h
index 4dab90b4e8..2a041cf9cb 100644
--- a/hipamd/include/hcc_detail/hip_ldg.h
+++ b/hipamd/include/hcc_detail/hip_ldg.h
@@ -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
 
diff --git a/hipamd/include/hcc_detail/hip_runtime.h b/hipamd/include/hcc_detail/hip_runtime.h
index 07cb9cc73e..1509f412cc 100644
--- a/hipamd/include/hcc_detail/hip_runtime.h
+++ b/hipamd/include/hcc_detail/hip_runtime.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
diff --git a/hipamd/src/hip_ldg.cpp b/hipamd/src/hip_ldg.cpp
index 434c2efdc3..e3af9860e3 100644
--- a/hipamd/src/hip_ldg.cpp
+++ b/hipamd/src/hip_ldg.cpp
@@ -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];