Skip to content

Commit

Permalink
[HIP] Move std headers after device malloc/free
Browse files Browse the repository at this point in the history
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
  • Loading branch information
aaronenyeshi committed Jul 9, 2021
1 parent 47aeeff commit ccb1026
Show file tree
Hide file tree
Showing 2 changed files with 56 additions and 28 deletions.
68 changes: 40 additions & 28 deletions clang/lib/Headers/__clang_hip_runtime_wrapper.h
Original file line number Diff line number Diff line change
Expand Up @@ -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))
Expand Down Expand Up @@ -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>

Expand Down
16 changes: 16 additions & 0 deletions clang/test/Headers/hip-header.hip
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}

0 comments on commit ccb1026

Please sign in to comment.