r356291 - [CUDA][Windows] Partial fix for bug 38811 (Step 2 of 3)

[Evgeny Mankov via cfe-commits]

Author: emankov
Date: Fri Mar 15 12:04:46 2019
New Revision: 356291

URL: http://llvm.org/viewvc/llvm-project?rev=356291&view=rev
Log:
[CUDA][Windows] Partial fix for bug 38811 (Step 2 of 3)

Partial fix for the clang Bug 38811 "Clang fails to compile with CUDA-9.x on Windows".

[Synopsis]
__sptr is a new Microsoft specific modifier (https://docs.microsoft.com/en-us/cpp/cpp/sptr-uptr?view=vs-2017).

[Solution]
Replace all `__sptr` occurrences with `__s` (and all `__cptr` with `__c` as well) to eliminate the below clang compilation error on Windows.

In file included from C:\GIT\LLVM\trunk\llvm-64-release-vs2017-15.9.5\dist\lib\clang\9.0.0\include\__clang_cuda_runtime_wrapper.h:162:
C:\GIT\LLVM\trunk\llvm-64-release-vs2017-15.9.5\dist\lib\clang\9.0.0\include\__clang_cuda_device_functions.h:524:33: error: expected expression
  return __nv_fast_sincosf(__a, __sptr, __cptr);
                                ^
Reviewed by: Artem Belevich

Differential Revision: http://reviews.llvm.org/D59423

Modified:
    cfe/trunk/lib/Headers/__clang_cuda_device_functions.h
    cfe/trunk/lib/Headers/__clang_cuda_libdevice_declares.h

Modified: cfe/trunk/lib/Headers/__clang_cuda_device_functions.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/__clang_cuda_device_functions.h?rev=356291&r1=356290&r2=356291&view=diff
==============================================================================
--- cfe/trunk/lib/Headers/__clang_cuda_device_functions.h (original)
+++ cfe/trunk/lib/Headers/__clang_cuda_device_functions.h Fri Mar 15 12:04:46 2019
@@ -520,8 +520,8 @@ __DEVICE__ unsigned int __sad(int __a, i
 __DEVICE__ float __saturatef(float __a) { return __nv_saturatef(__a); }
 __DEVICE__ int __signbitd(double __a) { return __nv_signbitd(__a); }
 __DEVICE__ int __signbitf(float __a) { return __nv_signbitf(__a); }
-__DEVICE__ void __sincosf(float __a, float *__sptr, float *__cptr) {
-  return __nv_fast_sincosf(__a, __sptr, __cptr);
+__DEVICE__ void __sincosf(float __a, float *__s, float *__c) {
+  return __nv_fast_sincosf(__a, __s, __c);
 }
 __DEVICE__ float __sinf(float __a) { return __nv_fast_sinf(__a); }
 __DEVICE__ int __syncthreads_and(int __a) { return __nvvm_bar0_and(__a); }
@@ -1713,17 +1713,17 @@ __DEVICE__ float scalblnf(float __a, lon
   return scalbnf(__a, (int)__b);
 }
 __DEVICE__ double sin(double __a) { return __nv_sin(__a); }
-__DEVICE__ void sincos(double __a, double *__sptr, double *__cptr) {
-  return __nv_sincos(__a, __sptr, __cptr);
+__DEVICE__ void sincos(double __a, double *__s, double *__c) {
+  return __nv_sincos(__a, __s, __c);
 }
-__DEVICE__ void sincosf(float __a, float *__sptr, float *__cptr) {
-  return __FAST_OR_SLOW(__nv_fast_sincosf, __nv_sincosf)(__a, __sptr, __cptr);
+__DEVICE__ void sincosf(float __a, float *__s, float *__c) {
+  return __FAST_OR_SLOW(__nv_fast_sincosf, __nv_sincosf)(__a, __s, __c);
 }
-__DEVICE__ void sincospi(double __a, double *__sptr, double *__cptr) {
-  return __nv_sincospi(__a, __sptr, __cptr);
+__DEVICE__ void sincospi(double __a, double *__s, double *__c) {
+  return __nv_sincospi(__a, __s, __c);
 }
-__DEVICE__ void sincospif(float __a, float *__sptr, float *__cptr) {
-  return __nv_sincospif(__a, __sptr, __cptr);
+__DEVICE__ void sincospif(float __a, float *__s, float *__c) {
+  return __nv_sincospif(__a, __s, __c);
 }
 __DEVICE__ float sinf(float __a) {
   return __FAST_OR_SLOW(__nv_fast_sinf, __nv_sinf)(__a);

Modified: cfe/trunk/lib/Headers/__clang_cuda_libdevice_declares.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Headers/__clang_cuda_libdevice_declares.h?rev=356291&r1=356290&r2=356291&view=diff
==============================================================================
--- cfe/trunk/lib/Headers/__clang_cuda_libdevice_declares.h (original)
+++ cfe/trunk/lib/Headers/__clang_cuda_libdevice_declares.h Fri Mar 15 12:04:46 2019
@@ -141,7 +141,7 @@ __device__ float __nv_fast_log10f(float
 __device__ float __nv_fast_log2f(float __a);
 __device__ float __nv_fast_logf(float __a);
 __device__ float __nv_fast_powf(float __a, float __b);
-__device__ void __nv_fast_sincosf(float __a, float *__sptr, float *__cptr);
+__device__ void __nv_fast_sincosf(float __a, float *__s, float *__c);
 __device__ float __nv_fast_sinf(float __a);
 __device__ float __nv_fast_tanf(float __a);
 __device__ double __nv_fdim(double __a, double __b);