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

[emankov]

Summary

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.

[Solution]
Replace all sptr occurrences with _sptr to eliminate the below clang compilation error on Windows.

In file included from C:\GIT\LLVM\trunk-for-submits\llvm-64-release-vs2017-15.9.5\dist\lib\clang\9.0.0\include\__clang_cuda_runtime_wrapper.h:162:
C:\GIT\LLVM\trunk-for-submits\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);
                                ^

[How to repro]

clang++.exe -x cuda "c:\ProgramData\NVIDIA Corporation\CUDA Samples\v9.0\0_Simple\simplePrintf\simplePrintf.cu" -I"c:\ProgramData\NVIDIA Corporation\CUDA Samples\v9.0\common\inc" --cuda-gpu-arch=sm_50 --cuda-path="C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0" -L"c:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0\lib\x64" -lcudart.lib  -v

[tra]

___ stands out as a sore thumb and raises unnecessary questions -- "why does it have three underscores, while __cptr is fine with two?".
Perhaps for consistency sake it would be better to replace __sptr -> __s and __cptr -> __c.
Given that we're just passing the args through, we don't really need to have ptr here.

[emankov]

Perhaps for consistency sake it would be better to replace __sptr -> __s and __cptr -> __c.

Well, it came from NVIDIA code, you know, I mean all those double underscores. As long as _sptr is also a keyword reserved by Microsoft, I think __s and __c are a good choice. Moreover, there are also double *__b and int *__c presented as device functions' arguments.

[emankov]

As long as _sptr is also a keyword reserved by Microsoft, I agree, that float *__s and float *__c are a good choice. Moreover, there are also double *__b and int *__c presented as device functions' arguments.

[tra]

In D59423#1431151, @emankov wrote:

Perhaps for consistency sake it would be better to replace __sptr -> __s and __cptr -> __c.

Well, it came from NVIDIA code, you know, I mean all those double underscores.

Not really. Function *signatures* are determined by CUDA, but argument names are not part of them and can be anything. Use of __ prefix is reserved for internal compiler use. Alas, in this case internal use of __sptr by clang happened to clash with the use of __sptr on windows. Neither is wrong.

[emankov]

Use of __ prefix is reserved for internal compiler use.

It is generally known. What is unknown, at least for me, is why in those clang's headers __ prefix is used so heavily?

[tra]

The intent is to avoid unintentional clashes with the preprocessor macros the user may have defined.
https://reviews.llvm.org/rL260647

[emankov]

The intent is to avoid unintentional clashes with the preprocessor macros the user may have defined.
https://reviews.llvm.org/rL260647

Ok, thanks!

lib/Headers/__clang_cuda_device_functions.h

 #define __prof_trigger(__a) asm __volatile__("pmevent \t%0;" ::"i"(__a))
 __DEVICE__ int __rhadd(int __a, int __b) { return __nv_rhadd(__a, __b); }
 __DEVICE__ unsigned int __sad(int __a, int __b, unsigned int __c) {
   return __nv_sad(__a, __b, __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); }
 __DEVICE__ int __syncthreads_count(int __a) { return __nvvm_bar0_popc(__a); }
 __DEVICE__ int __syncthreads_or(int __a) { return __nvvm_bar0_or(__a); }
 __DEVICE__ float __tanf(float __a) { return __nv_fast_tanf(__a); }
 __DEVICE__ void __threadfence(void) { __nvvm_membar_gl(); }
 __DEVICE__ void __threadfence_block(void) { __nvvm_membar_cta(); };

 __DEVICE__ float scalblnf(float __a, long __b) {
   if (__b > INT_MAX)
     return __a > 0 ? HUGE_VALF : -HUGE_VALF;
   if (__b < INT_MIN)
     return __a > 0 ? 0.f : -0.f;
   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);
 }
 __DEVICE__ double sinh(double __a) { return __nv_sinh(__a); }
 __DEVICE__ float sinhf(float __a) { return __nv_sinhf(__a); }
 __DEVICE__ double sinpi(double __a) { return __nv_sinpi(__a); }
 __DEVICE__ float sinpif(float __a) { return __nv_sinpif(__a); }

lib/Headers/__clang_cuda_libdevice_declares.h

 __device__ float __nv_fast_cosf(float __a);
 __device__ float __nv_fast_exp10f(float __a);
 __device__ float __nv_fast_expf(float __a);
 __device__ float __nv_fast_fdividef(float __a, float __b);
 __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);
 __device__ float __nv_fdimf(float __a, float __b);
 __device__ float __nv_fdiv_rd(float __a, float __b);
 __device__ float __nv_fdiv_rn(float __a, float __b);
 __device__ float __nv_fdiv_ru(float __a, float __b);
 __device__ float __nv_fdiv_rz(float __a, float __b);