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

[Evgeny Mankov via cfe-commits]

Author: Evgeny Mankov
Date: 2022-04-21T00:41:20+03:00
New Revision: c23147106f7efc4b5e29c47a08951116b4d994ac

URL: https://github.com/llvm/llvm-project/commit/c23147106f7efc4b5e29c47a08951116b4d994ac
DIFF: https://github.com/llvm/llvm-project/commit/c23147106f7efc4b5e29c47a08951116b4d994ac.diff

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

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;

Reviewed By: tra

Differential Revision: https://reviews.llvm.org/D122897

Added: 
    

Modified: 
    clang/lib/Headers/__clang_cuda_intrinsics.h

Removed: 
    


################################################################################
diff  --git a/clang/lib/Headers/__clang_cuda_intrinsics.h b/clang/lib/Headers/__clang_cuda_intrinsics.h
index eee2930ece85a..cfd5eb869e340 100644
--- a/clang/lib/Headers/__clang_cuda_intrinsics.h
+++ b/clang/lib/Headers/__clang_cuda_intrinsics.h
@@ -509,7 +509,7 @@ __device__ inline void *__nv_cvta_constant_to_generic_impl(size_t __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"