Merge pull request #639 from ThePortlandGroup/nv_stage

Pull 2018-12-17T09-39 Recent NVIDIA Changes

Author: sscalpone

runtime/flang/ieee_arithmetic.s

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017, NVIDIA CORPORATION.  All rights reserved.
+ * Copyright (c) 1995-2018, NVIDIA CORPORATION.  All rights reserved.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -404,7 +404,11 @@ __fort_getenv(const char *nm)
   int n;
 
   n = strlen(nm);
+#if defined(TARGET_OSX)
   e = env;
+#else             
+  e = environ;
+#endif
   while (*e != NULL) {
     if ((strncmp(*e, nm, n) == 0) && ((*((*e) + n)) == '=')) {
       return ((*e) + n + 1);

runtime/flang/ieee_exceptions.s

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017, NVIDIA CORPORATION.  All rights reserved.
+ * Copyright (c) 1998-2018, NVIDIA CORPORATION.  All rights reserved.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -25,7 +25,8 @@
 
 /* get environ */
 
-#if defined(WIN32) || defined(WIN64)
+#if defined(WIN64)
+char * * * __cdecl __p__environ(void);
 /*
  * enclose _fileno within parens to ensure calling the function rather than
  * the _fileno function macro (if/when it exists).
@@ -45,7 +46,9 @@ extern char **environ;
 char **
 __io_environ()
 {
-#if   !defined(TARGET_OSX)
+#if defined(TARGET_WIN)
+  return(*__p__environ());
+#elif !defined(TARGET_OSX)
   return (environ);
 #else
   return (*_NSGetEnviron());
@@ -287,10 +290,11 @@ __io_timezone(void *tm)
 }
 
 #if  (defined(WIN32) || defined(WIN64))
-/* OT 10 */
+/* wrappers for stderr, stdin, stdout : include
+  pgc/port/pgi_iobuf.h after stdio.h 
+ */
 void * 
 _pgi_get_iob(int xx) {
-	 return & __iob_func()[xx];
+	 return __acrt_iob_func (xx);
 }
-
 #endif

runtime/flang/initpar.c

@@ -39,7 +39,7 @@ if(${CMAKE_SYSTEM_PROCESSOR} MATCHES "x86_64" AND NOT ${LIBPGMATH_WITH_GENERIC})
     add_subdirectory("tanf")
 
     # Specific definition, will probably need to remove when we create a single file for each def
-    set_property(SOURCE dispatch.c APPEND_STRING PROPERTY COMPILE_FLAGS "-fno-builtin-c{sqrt,pow,log,exp,acos,asin,atan,cos,sin,tan,cosh,sinh,tanh}{,f}")
+    set_property(SOURCE dispatch.c APPEND_STRING PROPERTY COMPILE_FLAGS "-fno-builtin-c{sqrt,pow,log,exp,acos,asin,atan,cos,sin,tan,cosh,sinh,tanh}{,f} -fno-builtin-{ceil,floor}{,f}")
     set_property(SOURCE dispatch.c APPEND PROPERTY COMPILE_DEFINITIONS FOR_LIBPGC PGFLANG)
     set_property(SOURCE mth_xintrinsics.c APPEND_STRING PROPERTY COMPILE_FLAGS "-march=core2 ")
     set_property(SOURCE mth_yintrinsics.c APPEND_STRING PROPERTY COMPILE_FLAGS "-march=sandybridge ")
@@ -79,7 +79,7 @@ if(${CMAKE_SYSTEM_PROCESSOR} MATCHES "x86_64" AND NOT ${LIBPGMATH_WITH_GENERIC})
     add_subdirectory("tanf")
 
     # Specific definition, will probably need to remove when we create a single file for each def
-    set_property(SOURCE dispatch.c APPEND_STRING PROPERTY COMPILE_FLAGS "-fno-builtin-c{sqrt,pow,log,exp,acos,asin,atan,cos,sin,tan,cosh,sinh,tanh}{,f}")
+    set_property(SOURCE dispatch.c APPEND_STRING PROPERTY COMPILE_FLAGS "-fno-builtin-c{sqrt,pow,log,exp,acos,asin,atan,cos,sin,tan,cosh,sinh,tanh}{,f} -fno-builtin-{ceil,floor}{,f}")
     set_property(SOURCE dispatch.c APPEND PROPERTY COMPILE_DEFINITIONS FOR_LIBPGC PGFLANG)
     set_property(SOURCE mth_xintrinsics.c APPEND_STRING PROPERTY COMPILE_FLAGS "-msse2 ")
     set_property(SOURCE mth_yintrinsics.c APPEND_STRING PROPERTY COMPILE_FLAGS "-mavx ")
@@ -218,6 +218,7 @@ set(SRCS
   bessel_y0.c
   bessel_y1.c
   bessel_yn.c
+  ceil.c
   ${MTH_CMPLX_SRCS}
   cosd.c
   cvt.c
@@ -231,7 +232,9 @@ set(SRCS
   dbessel_y0.c
   dbessel_y1.c
   dbessel_yn.c
+  dceil.c
   dcosd.c
+  dfloor.c
   dmod.c
   dpowi.c
   dpowk.c
@@ -244,6 +247,7 @@ set(SRCS
   erfc_scaledf.c
   erfcf.c
   erff.c
+  floor.c
   fpcvt.c
   gamma.c
   gammaf.c

runtime/flangrti/iostdinit.c

@@ -0,0 +1,27 @@
+/*
+ * Copyright (c) 2018, NVIDIA CORPORATION.  All rights reserved.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *
+ */
+
+#include "mthdecls.h"
+#if     defined(__SSE4_1__) || defined(__AVX__)
+#include    <immintrin.h>
+#endif
+
+float
+__mth_i_ceil(float x)
+{
+  return ceilf(x);
+}

runtime/libpgmath/lib/common/CMakeLists.txt

@@ -0,0 +1,41 @@
+/*
+ * Copyright (c) 2018, NVIDIA CORPORATION.  All rights reserved.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *
+ */
+
+#include "mthdecls.h"
+#if     defined(__SSE4_1__) || defined(__AVX__)
+#include    <immintrin.h>
+#endif
+
+#if     defined(__AVX__)
+double
+__mth_i_dceil_avx(double x)
+{
+  return _mm_cvtsd_f64(_mm_ceil_sd(_mm_set1_pd(x), _mm_set1_pd(x)));
+}
+#elif   defined(__SSE4_1__)
+double
+__mth_i_dceil_sse(double x)
+{
+  return _mm_cvtsd_f64(_mm_ceil_sd(_mm_set1_pd(x), _mm_set1_pd(x)));
+}
+#else
+double
+__mth_i_dceil(double x)
+{
+  return ceil(x);
+}
+#endif

runtime/libpgmath/lib/common/ceil.c

@@ -0,0 +1,41 @@
+/*
+ * Copyright (c) 2018, NVIDIA CORPORATION.  All rights reserved.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *
+ */
+
+#include "mthdecls.h"
+#if     defined(__SSE4_1__) || defined(__AVX__)
+#include    <immintrin.h>
+#endif
+
+#if     defined(__AVX__)
+double
+__mth_i_dfloor_avx(double x)
+{
+  return _mm_cvtsd_f64(_mm_floor_sd(_mm_set1_pd(x), _mm_set1_pd(x)));
+}
+#elif   defined(__SSE4_1__)
+double
+__mth_i_dfloor_sse(double x)
+{
+  return _mm_cvtsd_f64(_mm_floor_sd(_mm_set1_pd(x), _mm_set1_pd(x)));
+}
+#else
+double
+__mth_i_dfloor(double x)
+{
+  return floor(x);
+}
+#endif

runtime/libpgmath/lib/common/dceil.c

@@ -0,0 +1,41 @@
+/*
+ * Copyright (c) 2018, NVIDIA CORPORATION.  All rights reserved.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ *
+ */
+
+#include "mthdecls.h"
+#if     defined(__SSE4_1__) || defined(__AVX__)
+#include    <immintrin.h>
+#endif
+
+#if     defined(__AVX__)
+float
+__mth_i_floor_avx(float x)
+{
+  return _mm_cvtss_f32(_mm_floor_ss(_mm_set1_ps(x), _mm_set1_ps(x)));
+}
+#elif   defined(__SSE4_1__)
+float
+__mth_i_floor_sse(float x)
+{
+  return _mm_cvtss_f32(_mm_floor_ss(_mm_set1_ps(x), _mm_set1_ps(x)));
+}
+#else
+float
+__mth_i_floor(float x)
+{
+  return floorf(x);
+}
+#endif

runtime/libpgmath/lib/common/dfloor.c

@@ -522,6 +522,8 @@ float __mth_i_bessel_y0(float arg);
 float __mth_i_bessel_y1(float arg);
 float __mth_i_bessel_yn(int n, float arg);
 float __f90_bessel_yn(int n1, int n2, float f);
+float __mth_i_ceil(float);
+float __mth_i_floor(float);
 
 int __mth_i_idnint(double d);
 int __mth_i_mod(int i, int j);
@@ -575,6 +577,8 @@ double __mth_i_dbessel_y0(double arg);
 double __mth_i_dbessel_y1(double arg);
 double __mth_i_dbessel_yn(int n, double arg);
 double __f90_dbessel_yn(int n1, int n, double d);
+double __mth_i_dceil(double);
+double __mth_i_dfloor(double);
 
 #if	! defined (TARGET_X8664) && ! defined(LINUX8664)
 /*