[llvm-branch-commits] [clang] 79147e4 - [clang][CUDA][Windows] Fix compilation error on Windows with `uint32_t __nvvm_get_smem_pointer`

[Tom Stellard via llvm-branch-commits]

Author: Evgeny Mankov
Date: 2022-06-02T22:27:35-07:00
New Revision: 79147e4722cc55b819b00007d0ac9a7d2f036956

URL: https://github.com/llvm/llvm-project/commit/79147e4722cc55b819b00007d0ac9a7d2f036956
DIFF: https://github.com/llvm/llvm-project/commit/79147e4722cc55b819b00007d0ac9a7d2f036956.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

(cherry picked from commit c23147106f7efc4b5e29c47a08951116b4d994ac)

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 e0875bbcbf4ae..3d2f5f40a435a 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"