Add more function declarations to hip-clang
Cette révision appartient à :
@@ -90,7 +90,7 @@ if ($HIP_PLATFORM eq "clang") {
|
||||
}
|
||||
$ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm";
|
||||
$HIPCC="$HIP_CLANG_PATH/clang++";
|
||||
$HIPCXXFLAGS .= " -I$HIP_PATH/include";
|
||||
$HIPCXXFLAGS .= "-std=c++11 -I$HIP_PATH/include";
|
||||
|
||||
$HIPLDFLAGS = "--hip-link --hip-device-lib-path=$DEVICE_LIB_PATH -L$HIP_PATH/lib -lhip_hcc";
|
||||
} elsif ($HIP_PLATFORM eq "hcc") {
|
||||
|
||||
@@ -29,6 +29,20 @@ THE SOFTWARE.
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_RUNTIME_H
|
||||
|
||||
#if defined(__HCC__)
|
||||
#define __HCC_OR_HIP_CLANG__ 1
|
||||
#define __HCC_ONLY__ 1
|
||||
#define __HIP_CLANG_ONLY__ 0
|
||||
#elif defined(__clang__) && defined(__HIP__)
|
||||
#define __HCC_OR_HIP_CLANG__ 1
|
||||
#define __HCC_ONLY__ 0
|
||||
#define __HIP_CLANG_ONLY__ 1
|
||||
#else
|
||||
#define __HCC_OR_HIP_CLANG__ 0
|
||||
#define __HCC_ONLY__ 0
|
||||
#define __HIP_CLANG_ONLY__ 0
|
||||
#endif
|
||||
|
||||
//---
|
||||
// Top part of file can be compiled with any compiler
|
||||
|
||||
@@ -41,15 +55,16 @@ THE SOFTWARE.
|
||||
#include <stddef.h>
|
||||
#endif //__cplusplus
|
||||
|
||||
#if __HCC__
|
||||
#if __HCC_OR_HIP_CLANG__
|
||||
|
||||
// Define NVCC_COMPAT for CUDA compatibility
|
||||
#define NVCC_COMPAT
|
||||
#define CUDA_SUCCESS hipSuccess
|
||||
|
||||
#include <hip/hip_runtime_api.h>
|
||||
#endif // __HCC_OR_HIP_CLANG__
|
||||
|
||||
|
||||
#if __HCC__
|
||||
// define HIP_ENABLE_PRINTF to enable printf
|
||||
#ifdef HIP_ENABLE_PRINTF
|
||||
#define HCC_ENABLE_ACCELERATOR_PRINTF 1
|
||||
@@ -164,6 +179,10 @@ extern int HIP_TRACE_API;
|
||||
#define __HCC_C__
|
||||
#endif
|
||||
|
||||
#endif // defined __HCC__
|
||||
|
||||
#if __HCC_OR_HIP_CLANG__
|
||||
|
||||
// TODO - hipify-clang - change to use the function call.
|
||||
//#define warpSize hc::__wavesize()
|
||||
static constexpr int warpSize = 64;
|
||||
@@ -371,6 +390,10 @@ __device__ void __threadfence_system(void);
|
||||
* @}
|
||||
*/
|
||||
|
||||
#endif // __HCC_OR_CLANG__
|
||||
|
||||
#if defined __HCC__
|
||||
|
||||
template <
|
||||
typename std::common_type<decltype(hc_get_group_id), decltype(hc_get_group_size),
|
||||
decltype(hc_get_num_groups), decltype(hc_get_workitem_id)>::type f>
|
||||
@@ -628,6 +651,9 @@ __DEVICE__ void inline __assert_fail(const char * __assertion,
|
||||
// Ignore all the args for now.
|
||||
__device_trap();
|
||||
}
|
||||
|
||||
__DEVICE__ void __syncthreads();
|
||||
|
||||
#pragma push_macro("__DEVICE__")
|
||||
|
||||
#include <hip/hcc_detail/math_functions.h>
|
||||
|
||||
Référencer dans un nouveau ticket
Bloquer un utilisateur