Merge pull request #586 from ThePortlandGroup/nv_stage

Pull 2018-09-12T15-41 Recent NVIDIA Changes

Author: sscalpone

runtime/libpgmath/lib/common/arm64intrin.h

@@ -2608,7 +2608,7 @@ _mm_min_sd(const __m128d& a, const __m128d& b)
 {
   double aa = vec_extract(a, 0);
   double bb = vec_extract(b, 0);
-  aa < bb ? aa : bb;
+  aa = aa < bb ? aa : bb;
   return vec_insert(aa, a, 0);
 }
 
@@ -2618,7 +2618,7 @@ _mm_max_sd(const __m128d& a, const __m128d& b)
 {
   double aa = vec_extract(a, 0);
   double bb = vec_extract(b, 0);
-  aa > bb ? aa : bb;
+  aa = aa > bb ? aa : bb;
   return vec_insert(aa, a, 0);
 }
 

runtime/libpgmath/lib/common/log/fma3/fdlog1.cpp

@@ -36,7 +36,7 @@ __m128d __internal_fast_int2dbl(__m128i a)
     __m128d const INT2DBL    = (__m128d)_mm_set1_epi64x(INT2DBL_D);
 
     __m128i t = _mm_xor_si128(INT2DBL_LO, a);
-    t = _mm_blend_epi32(INT2DBL_HI, t, H55_D); 
+    t = _mm_blend_epi32(INT2DBL_HI, t, 0x5);
     return _mm_sub_sd((__m128d)t, INT2DBL);
 }
 

runtime/libpgmath/lib/common/log/fma3/fdlog2.cpp

@@ -36,7 +36,7 @@ __m128d __internal_fast_int2dbl(__m128i a)
     __m128d const INT2DBL    = (__m128d)_mm_set1_epi64x(INT2DBL_D);
 
     __m128i t = _mm_xor_si128(INT2DBL_LO, a);
-    t = _mm_blend_epi32(INT2DBL_HI, t, H55_D); 
+    t = _mm_blend_epi32(INT2DBL_HI, t, 0x5); 
     return _mm_sub_pd((__m128d)t, INT2DBL);
 }
 

runtime/libpgmath/lib/common/log/fma3/fdlog4.cpp

@@ -35,7 +35,7 @@ __m256d __internal_fast_int2dbl(__m256i a)
     __m256d const INT2DBL    = (__m256d)_mm256_set1_epi64x(INT2DBL_D);
 
     __m256i t = _mm256_xor_si256(INT2DBL_LO, a);
-    t = _mm256_blend_epi32(INT2DBL_HI, t, H55_D); 
+    t = _mm256_blend_epi32(INT2DBL_HI, t, 0x55);
     return _mm256_sub_pd((__m256d)t, INT2DBL);
 }
 

runtime/libpgmath/lib/common/log/fma3/fdlog8.cpp

@@ -37,7 +37,7 @@ __m512d __internal_fast_int2dbl(__m512i a)
     __m512d const INT2DBL    = (__m512d)_mm512_set1_epi64(INT2DBL_D);
 
     __m512i t = _mm512_xor_si512(INT2DBL_LO, a);
-    t = _mm512_mask_blend_epi32((H55_D<<8)|H55_D, INT2DBL_HI, t);
+    t = _mm512_mask_blend_epi32(0x5555, INT2DBL_HI, t);
     return _mm512_sub_pd((__m512d)t, INT2DBL);
 }
 

runtime/libpgmath/lib/common/log/fma3/fdlog_defs.h

@@ -1,5 +1,5 @@
 /* 
- * Copyright (c) 2017, NVIDIA CORPORATION.  All rights reserved.
+ * Copyright (c) 2017-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.
@@ -20,8 +20,6 @@
 #define INT2DBL_LO_D              0x0000000080000000
 #define INT2DBL_D                 0x4330000080000000
 
-#define H55_D                     0x55
-
 #define ZERO_D                    0.0
 #define ALL_ONES_EXPONENT_D       0x7FF0000000000000
 #define NAN_VAL_D                 0xFFF8000000000000

tools/flang1/flang1exe/symacc.h

@@ -181,7 +181,7 @@ typedef struct {
   const char *ocnames[OC_MAX + 1];
   const char *scnames[SC_MAX + 1];
   const char *tynames[TY_MAX + 1];
-  int i0, i1;
+  SPTR i0, i1;
   int k0, k1;
   SPTR flt0, dbl0, quad0;
   SPTR fltm0, dblm0, quadm0; /* floating point minus 0 */