From: Evgeny Mankov <evgeny.mankov@gmail.com>
Date: Fri, 15 Mar 2019 19:04:46 +0000 (+0000)
Subject: [CUDA][Windows] Partial fix for bug 38811 (Step 2 of 3)
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=fc62308da45f483cd97bcf823f5a01b025f6a8ba;p=clang

[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

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@356291 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/lib/Headers/__clang_cuda_device_functions.h b/lib/Headers/__clang_cuda_device_functions.h
index 28b214c147..2389605f2e 100644
--- a/lib/Headers/__clang_cuda_device_functions.h
+++ b/lib/Headers/__clang_cuda_device_functions.h
@@ -520,8 +520,8 @@ __DEVICE__ unsigned int __sad(int __a, int __b, unsigned int __c) {
 __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, long __b) {
   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);
diff --git a/lib/Headers/__clang_cuda_libdevice_declares.h b/lib/Headers/__clang_cuda_libdevice_declares.h
index 71df7f849d..9c428a2162 100644
--- a/lib/Headers/__clang_cuda_libdevice_declares.h
+++ b/lib/Headers/__clang_cuda_libdevice_declares.h
@@ -141,7 +141,7 @@ __device__ float __nv_fast_log10f(float __a);
 __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);