Author: Aaron En Ye Shi Date: 2021-07-09T21:20:16Z New Revision: ccb10266f56bc34123eb7741c6ebcd5ba8ed3dcc
URL: https://github.com/llvm/llvm-project/commit/ccb10266f56bc34123eb7741c6ebcd5ba8ed3dcc DIFF: https://github.com/llvm/llvm-project/commit/ccb10266f56bc34123eb7741c6ebcd5ba8ed3dcc.diff LOG: [HIP] Move std headers after device malloc/free Set the device malloc and free functions as weak, and move the std headers after device malloc/free to avoid issues with std malloc/free. Fixes: SWDEV-293590 Reviewed By: yaxunl Differential Revision: https://reviews.llvm.org/D105707 Added: Modified: clang/lib/Headers/__clang_hip_runtime_wrapper.h clang/test/Headers/hip-header.hip Removed: ################################################################################ diff --git a/clang/lib/Headers/__clang_hip_runtime_wrapper.h b/clang/lib/Headers/__clang_hip_runtime_wrapper.h index c557796c8fa0a..6aefbeaec566f 100644 --- a/clang/lib/Headers/__clang_hip_runtime_wrapper.h +++ b/clang/lib/Headers/__clang_hip_runtime_wrapper.h @@ -18,28 +18,6 @@ #if __HIP__ -#if !defined(__HIPCC_RTC__) -#include <cmath> -#include <cstdlib> -#include <stdlib.h> -#else -typedef __SIZE_TYPE__ size_t; -// Define macros which are needed to declare HIP device API's without standard -// C/C++ headers. This is for readability so that these API's can be written -// the same way as non-hipRTC use case. These macros need to be popped so that -// they do not pollute users' name space. -#pragma push_macro("NULL") -#pragma push_macro("uint32_t") -#pragma push_macro("uint64_t") -#pragma push_macro("CHAR_BIT") -#pragma push_macro("INT_MAX") -#define NULL (void *)0 -#define uint32_t __UINT32_TYPE__ -#define uint64_t __UINT64_TYPE__ -#define CHAR_BIT __CHAR_BIT__ -#define INT_MAX __INTMAX_MAX__ -#endif // __HIPCC_RTC__ - #define __host__ __attribute__((host)) #define __device__ __attribute__((device)) #define __global__ __attribute__((global)) @@ -68,24 +46,58 @@ extern "C" { } #endif //__cplusplus +typedef __SIZE_TYPE__ __hip_size_t; + +#ifdef __cplusplus +extern "C" { +#endif //__cplusplus + #if __HIP_ENABLE_DEVICE_MALLOC__ -extern "C" __device__ void *__hip_malloc(size_t __size); -extern "C" __device__ void *__hip_free(void *__ptr); -static inline __device__ void *malloc(size_t __size) { +__device__ void *__hip_malloc(__hip_size_t __size); +__device__ void *__hip_free(void *__ptr); +__attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) { return __hip_malloc(__size); } -static inline __device__ void *free(void *__ptr) { return __hip_free(__ptr); } +__attribute__((weak)) inline __device__ void *free(void *__ptr) { + return __hip_free(__ptr); +} #else -static inline __device__ void *malloc(size_t __size) { +__attribute__((weak)) inline __device__ void *malloc(__hip_size_t __size) { __builtin_trap(); return nullptr; } -static inline __device__ void *free(void *__ptr) { +__attribute__((weak)) inline __device__ void *free(void *__ptr) { __builtin_trap(); return nullptr; } #endif +#ifdef __cplusplus +} // extern "C" +#endif //__cplusplus + +#if !defined(__HIPCC_RTC__) +#include <cmath> +#include <cstdlib> +#include <stdlib.h> +#else +typedef __SIZE_TYPE__ size_t; +// Define macros which are needed to declare HIP device API's without standard +// C/C++ headers. This is for readability so that these API's can be written +// the same way as non-hipRTC use case. These macros need to be popped so that +// they do not pollute users' name space. +#pragma push_macro("NULL") +#pragma push_macro("uint32_t") +#pragma push_macro("uint64_t") +#pragma push_macro("CHAR_BIT") +#pragma push_macro("INT_MAX") +#define NULL (void *)0 +#define uint32_t __UINT32_TYPE__ +#define uint64_t __UINT64_TYPE__ +#define CHAR_BIT __CHAR_BIT__ +#define INT_MAX __INTMAX_MAX__ +#endif // __HIPCC_RTC__ + #include <__clang_hip_libdevice_declares.h> #include <__clang_hip_math.h> diff --git a/clang/test/Headers/hip-header.hip b/clang/test/Headers/hip-header.hip index 0e95d58d55700..aa7abcedb7ae4 100644 --- a/clang/test/Headers/hip-header.hip +++ b/clang/test/Headers/hip-header.hip @@ -115,3 +115,19 @@ __device__ double test_isnan() { return r ; } + +// Check that device malloc and free do not conflict with std headers. +#include <cstdlib> +// CHECK-LABEL: define{{.*}}@_Z11test_malloc +// CHECK: call {{.*}}i8* @malloc(i64 +// CHECK: define weak {{.*}}i8* @malloc(i64 +__device__ void test_malloc(void *a) { + a = malloc(42); +} + +// CHECK-LABEL: define{{.*}}@_Z9test_free +// CHECK: call {{.*}}i8* @free(i8* +// CHECK: define weak {{.*}}i8* @free(i8* +__device__ void test_free(void *a) { + free(a); +} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
