Map device built-ins to compiler built-ins #763
Conversation
linehill
force-pushed
the
use-builtins
branch
from
6aa1000 to
a8aff74
Compare
pvelesko
force-pushed
the
use-builtins
branch
from
65808d5 to
b6becef
Compare
|
|
||
| if(CHIP_ENABLE_NON_COMPLIANT_DEVICELIB_CODE) | ||
| # add support for calling certain LLVM builtins that are not supported by HIP/CUDA | ||
| add_custom_command( | ||
| OUTPUT "${CMAKE_CURRENT_BINARY_DIR}/BC/c_to_opencl.bc" | ||
| DEPENDS "${CMAKE_CURRENT_SOURCE_DIR}/c_to_opencl.c" | ||
| COMMAND "${CMAKE_CXX_COMPILER}" ${BITCODE_C_COMPILE_FLAGS} | ||
| -o "${CMAKE_CURRENT_BINARY_DIR}/BC/c_to_opencl.bc" | ||
| -c "${CMAKE_CURRENT_SOURCE_DIR}/c_to_opencl.c" | ||
| COMMENT "Building c_to_opencl.bc" | ||
| VERBATIM) | ||
| list(APPEND DEPEND_LIST "${CMAKE_CURRENT_BINARY_DIR}/BC/c_to_opencl.bc") | ||
| endif() | ||
| add_custom_command( |
There was a problem hiding this comment.
Why are we getting rid of the option? These are still non-compliant, right?
There was a problem hiding this comment.
Why are we getting rid of the option?
These are required for using compiler built-ins (see c_to_opencl.def for explanation).
These are still non-compliant, right?
Non-compliant respect to what?
There was a problem hiding this comment.
So why not keep the option? In atomics.hh we still rely on this option:
#ifdef CHIP_ENABLE_NON_COMPLIANT_DEVICELIB_CODE
// At least rocPRIM tests call the unsigned long variant although it's not
// listed in the user manual. Annoyingly, size_t is typically defined as
// unsigned long.
// FIXME: We should check that unsigned long is 64bits for the host.
extern "C++" inline __device__ unsigned long atomicAdd(unsigned long *address,
unsigned long val) {
return __chip_atomic_add_l(address, val);
}
#endif
Non-compliant respect to what?
"when CHIP_ENABLE_NON_COMPLIANT_DEVICELIB_CODE` was defined, which brought in the atomicAdd() for longs (see 71b5c82) which is not defined in the HIP spec."
There was a problem hiding this comment.
If the concern here is that removing the option in the bitcode library would expose non-compliant functions to the users, that shouldn’t be happening because they are not declared in the HIP headers. Therefore, removing the CHIP_ENABLE_NON_COMPLIANT_DEVICELIB_CODE option check shouldn't have negative consequences, AFAICS.
bitcode/devicelib.cl
Outdated
|
|
||
| #ifdef CHIP_ENABLE_NON_COMPLIANT_DEVICELIB_CODE | ||
|
|
||
| // See c_to_opencl.def for details. | ||
| #define DEF_UNARY_FN_MAP(NAME_, TYPE_) \ | ||
| TYPE_ MAP_PREFIX##NAME_(TYPE_ x) { return NAME_(x); } | ||
| #define DEF_BINARY_FN_MAP(NAME_, TYPE_) \ | ||
| TYPE_ MAP_PREFIX##NAME_(TYPE_ x, TYPE_ y) { return NAME_(x, y); } | ||
| #define MAP_PREFIX __chip_c2ocl_ | ||
| #define DEF_UNARY_FN_MAP(FROM_FN_, TO_FN_, TYPE_) \ | ||
| TYPE_ __chip_c2ocl_##FROM_FN_(TYPE_ x) { return TO_FN_(x); } | ||
| #define DEF_BINARY_FN_MAP(FROM_FN_, TO_FN_, TYPE_) \ | ||
| TYPE_ __chip_c2ocl_##FROM_FN_(TYPE_ x, TYPE_ y) { return TO_FN_(x, y); } | ||
| #include "c_to_opencl.def" | ||
| #undef UNARY_FN | ||
| #undef BINARY_FN |
There was a problem hiding this comment.
Can you remind me why we have this here? If we already have this in c_to_opencl.c which we compile and then link into the final .bc?
There was a problem hiding this comment.
The reason is explained in the c_to_opencl.def.
And prepare for follow-up patch where HIP built-ins are mapped to compiler built-ins. Some of the compiler built-ins may be lowered to C math calls by the compiler which the cmath->OpenCL mapping takes care of.
Map a set of HIP built-ins referenced by HeCBench benchmarks to compiler built-ins. The myocyte-hip benchmark seems to benefit from this the most - most likely due to: * HIP built-in calls with compile time values which are constant-folded. * HIP built-in calls which are simplified such as `pow(x, 2) --> x*x`.
Also, just in case mark all HIP->compiler built-in functions 'static inline'.
pvelesko
force-pushed
the
use-builtins
branch
from
0dba21d to
bef642c
Compare
|
rebased |
Map a set of HIP built-ins referenced by HeCBench benchmarks to compiler built-ins. The myocyte-hip benchmark seems to benefit from this the most - most likely due to:
HIP built-in calls with compile time values which are constant-folded.
HIP built-in calls which are simplified such as
pow(x, 2) --> x*x.