[clang][CUDA][Windows] Fix compilation error on Windows with `uint32_t __nvvm_get_smem_pointer`

[emankov]

Summary

The change fixes https://github.com/llvm/llvm-project/issues/54609 (the second reported issue) by eliminating a compilation error occurring only on Windows while trying to compile any CUDA source file by clang (-x cuda).

[Repro]
clang -x cuda <any_cu_source>

[Error]

__clang_cuda_runtime_wrapper.h:473:
__clang_cuda_intrinsics.h(517,19): error GC871EEFB: unknown type name 'uint32_t'; did you mean 'cuuint32_t'?
__device__ inline uint32_t __nvvm_get_smem_pointer(void *__ptr) {
                          ^
C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v11.6/include\cuda.h:57:26: note: 'cuuint32_t' declared here
typedef unsigned __int32 cuuint32_t;

[tra]

clang/lib/Headers/__clang_cuda_intrinsics.h

489

We should include stdint.h unconditionally, and move inclusion to the top of the file.
Special-casing windows does not buy us anything here.

490

Include should be done outside of extern "C"

[emankov]

What about using cuuint32_t instead of uint32_t in __device__ inline uint32_t __nvvm_get_smem_pointer(void *__ptr)? Actually, I like it more. Let's CUDA do all the needed typedefs.

[tra]

In D122897#3423574, @emankov wrote:

What about using cuuint32_t instead of uint32_t in __device__ inline uint32_t __nvvm_get_smem_pointer(void *__ptr)? Actually, I like it more. Let's CUDA do all the needed typedefs.

That should work. cuuint32_t is defined in cuda.h which we always include during CUDA compilation.

I wonder why NVIDIA decided not to include stdint.h on windows:

#ifdef _MSC_VER
typedef unsigned __int32 cuuint32_t;
typedef unsigned __int64 cuuint64_t;
#else
#include <stdint.h>
typedef uint32_t cuuint32_t;
typedef uint64_t cuuint64_t;
#endif

[emankov]

I wonder why NVIDIA decided not to include stdint.h on windows:

#ifdef _MSC_VER
typedef unsigned __int32 cuuint32_t;
typedef unsigned __int64 cuuint64_t;
#else
#include <stdint.h>
typedef uint32_t cuuint32_t;
typedef uint64_t cuuint64_t;
#endif

Because CUDA uses Microsoft specific types unsigned __int32 and unsigned __int64 instead, and it looks that they are all what is needed, so there is no need in including <stdint.h> on Win path.
That is why I prefer using cuuint32_t instead of uint32_t, as it is the type proposed to be used in CUDA-related stuff.

[emankov]

Use cuuint32_t instead of uint32_t, as it is the type proposed to be used in CUDA-related stuff. The decision is based on cuda.h header file that is always included when compiling CUDA by clang -x cuda.

cuda.h:

#ifdef _MSC_VER
typedef unsigned __int32 cuuint32_t;
typedef unsigned __int64 cuuint64_t;
#else
#include <stdint.h>
typedef uint32_t cuuint32_t;
typedef uint64_t cuuint64_t;
#endif
``

clang/lib/Headers/__clang_cuda_intrinsics.h

inline __device__ unsigned __funnelshift_rc(unsigned low32, unsigned high32,
   return ret;
 }
 
 #endif // !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 320
 
 #if CUDA_VERSION >= 11000
 extern "C" {
 __device__ inline size_t __nv_cvta_generic_to_global_impl(const void *__ptr) {
   return (size_t)(void __attribute__((address_space(1))) *)__ptr;
 }
 __device__ inline size_t __nv_cvta_generic_to_shared_impl(const void *__ptr) {
   return (size_t)(void __attribute__((address_space(3))) *)__ptr;
 }
 __device__ inline size_t __nv_cvta_generic_to_constant_impl(const void *__ptr) {
   return (size_t)(void __attribute__((address_space(4))) *)__ptr;
 }
 __device__ inline size_t __nv_cvta_generic_to_local_impl(const void *__ptr) {
   return (size_t)(void __attribute__((address_space(5))) *)__ptr;
 }
 __device__ inline void *__nv_cvta_global_to_generic_impl(size_t __ptr) {
   return (void *)(void __attribute__((address_space(1))) *)__ptr;
 }
 __device__ inline void *__nv_cvta_shared_to_generic_impl(size_t __ptr) {
   return (void *)(void __attribute__((address_space(3))) *)__ptr;
 }
 __device__ inline void *__nv_cvta_constant_to_generic_impl(size_t __ptr) {
   return (void *)(void __attribute__((address_space(4))) *)__ptr;
 }
 __device__ inline void *__nv_cvta_local_to_generic_impl(size_t __ptr) {
   return (void *)(void __attribute__((address_space(5))) *)__ptr;
 }
-__device__ inline uint32_t __nvvm_get_smem_pointer(void *__ptr) {
+__device__ inline cuuint32_t __nvvm_get_smem_pointer(void *__ptr) {
   return __nv_cvta_generic_to_shared_impl(__ptr);
 }
 } // extern "C"
 #endif // CUDA_VERSION >= 11000
 
 #endif // defined(__CLANG_CUDA_INTRINSICS_H__)