AI/rocm-systems
2
0
Fork 0
Código Issues Pull Requests Ações Pacotes Projetos Versões Wiki Atividade

Merge 'amd-master-next' into 'amd-npi-next'

Ver Código-Fonte 
Change-Id: Id0bd0798c460b51653a73b6a9ce97ddaec92874a
Esse commit está contido em:
Jenkins
2020-06-16 21:09:56 +00:00
pai 06013d5750 7682ec6a7b
commit be658f6f76
58 arquivos alterados com 3278 adições e 1433 exclusões
Baixar Arquivo de Patch Baixar Arquivo de Diff Expandir todos os arquivos Colapsar todos os arquivos
bin/hipcc
+16 -6
Ver Arquivo
@@ -168,7 +168,14 @@ if (defined $HIP_COMPILER and $HIP_PLATFORM eq "hcc" and $HIP_COMPILER eq "clang
$HIP_CLANG_PATH = "$ROCM_PATH/llvm/bin";
}
if (!defined $DEVICE_LIB_PATH) {
$DEVICE_LIB_PATH = "$ROCM_PATH/lib";
if (-e "$ROCM_PATH/amdgcn/bitcode") {
$DEVICE_LIB_PATH = "$ROCM_PATH/amdgcn/bitcode";
}
else {
# This path is to support an older build of the device library
# TODO: To be removed in the future.
$DEVICE_LIB_PATH = "$ROCM_PATH/lib";
}
}
}
@@ -207,6 +214,7 @@ if ($HIP_PLATFORM eq "hcc" and $HIP_COMPILER eq "clang") {
$HIP_LIB_PATH = "$HIP_PATH/lib";
}
if ($verbose & 0x2) {
print ("ROCM_PATH=$ROCM_PATH\n");
if (defined $HIP_ROCCLR_HOME) {
print ("HIP_ROCCLR_HOME=$HIP_ROCCLR_HOME\n");
}
@@ -808,17 +816,17 @@ if ($HIP_PLATFORM eq "hcc" and $HIP_COMPILER eq "clang") {
$HIPCFLAGS .= " -O3";
$HIPLDFLAGS .= " -O3";
}
# Do not pass -mllvm on Windows since there is a clang bug causing duplicate -mllvm options in clang -cc1 on Windows.
# ToDo : remove restriction for Windows after clang bug is fixed.
if (!$funcSupp and $optArg ne "-O0" and not $isWindows and $hasHIP) {
if (!$funcSupp and $optArg ne "-O0" and $hasHIP) {
$HIPCXXFLAGS .= " -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false";
if ($needLDFLAGS and not $needCXXFLAGS) {
$HIPLDFLAGS .= " -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false";
}
}
$HIP_DEVLIB_FLAGS = " --hip-device-lib-path=$DEVICE_LIB_PATH";
if ($hasHIP) {
$HIPCXXFLAGS .= " $HIP_DEVLIB_FLAGS";
if ($DEVICE_LIB_PATH ne "$ROCM_PATH/amdgcn/bitcode") {
$HIPCXXFLAGS .= " --hip-device-lib-path=$DEVICE_LIB_PATH";
}
if ($HIP_RUNTIME ne "HCC") {
$HIPCXXFLAGS .= " -fhip-new-launch-api";
}
@@ -833,6 +841,8 @@ if ($HIP_PLATFORM eq "hcc" and $HIP_COMPILER eq "clang") {
} else {
$toolArgs .= " -Wl,--enable-new-dtags -Wl,--rpath=$HIP_LIB_PATH:$ROCM_PATH/lib -lhip_hcc ";
}
# To support __fp16 and _Float16, explicitly link with compiler-rt
$toolArgs .= " -L$HIP_CLANG_PATH/../lib/clang/$HIP_CLANG_VERSION/lib/linux -lclang_rt.builtins-x86_64 "
}
}
bin/hipconfig
+1 -1
Ver Arquivo
@@ -1,7 +1,7 @@
#!/usr/bin/perl -w
$HIP_BASE_VERSION_MAJOR = "3";
$HIP_BASE_VERSION_MINOR = "5";
$HIP_BASE_VERSION_MINOR = "6";
# Need perl > 5.10 to use logic-defined or
use 5.006; use v5.10.1;
docs/markdown/hip_porting_guide.md
+2 -4
Ver Arquivo
@@ -512,7 +512,7 @@ Compute programs sometimes use textures either to access dedicated texture cache
point samples. AMD hardware, as well as recent competing hardware,
has a unified texture/L1 cache, so it no longer has a dedicated texture cache. But the nvcc path often caches global loads in the L2 cache, and some programs may benefit from explicit control of the L1 cache contents. We recommend the __ldg instruction for this purpose.
AMD compilers currently load all data into both the L1 and L2 caches, so __ldg is treated as a no-op.
AMD compilers currently load all data into both the L1 and L2 caches, so __ldg is treated as a no-op.
We recommend the following for functional portability:
@@ -520,13 +520,11 @@ We recommend the following for functional portability:
- Programs that use texture object and reference APIs, work well on HIP
```
## More Tips
### HIPTRACE Mode
On an hcc/AMD platform, set the HIP_TRACE_API environment variable to see a textural API trace. Use the following bit mask:
- 0x1 = trace APIs
- 0x2 = trace synchronization operations
- 0x4 = trace memory allocation / deallocation
hip-config.cmake.in
+28 -6
Ver Arquivo
@@ -79,7 +79,7 @@ if(HIP_COMPILER STREQUAL "clang")
${HIP_CLANG_INCLUDE_SEARCH_PATHS}
NO_DEFAULT_PATH)
find_dependency(AMDDeviceLibs)
set(AMDGPU_TARGETS "gfx900;gfx906" CACHE STRING "AMD GPU targets to compile for")
set(AMDGPU_TARGETS "gfx900;gfx906;gfx908" CACHE STRING "AMD GPU targets to compile for")
set(GPU_TARGETS "${AMDGPU_TARGETS}" CACHE STRING "GPU targets to compile for")
else()
find_dependency(hcc)
@@ -142,18 +142,26 @@ else()
endif()
if(HIP_COMPILER STREQUAL "clang")
set_property(TARGET hip::device APPEND PROPERTY
INTERFACE_COMPILE_OPTIONS -x hip --hip-device-lib-path=${AMD_DEVICE_LIBS_PREFIX}/lib
)
if (HIP_CXX_COMPILER MATCHES ".*clang\\+\\+")
set_property(TARGET hip::device APPEND PROPERTY
INTERFACE_COMPILE_OPTIONS -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false
)
endif()
if (EXISTS ${AMD_DEVICE_LIBS_PREFIX}/amdgcn/bitcode)
set_property(TARGET hip::device APPEND PROPERTY
INTERFACE_COMPILE_OPTIONS -x hip
)
else()
# This path is to support an older build of the device library
# TODO: To be removed in the future.
set_property(TARGET hip::device APPEND PROPERTY
INTERFACE_COMPILE_OPTIONS -x hip --hip-device-lib-path=${AMD_DEVICE_LIBS_PREFIX}/lib
)
endif()
set_property(TARGET hip::device APPEND PROPERTY
INTERFACE_LINK_LIBRARIES --hip-device-lib-path=${AMD_DEVICE_LIBS_PREFIX}/lib --hip-link
INTERFACE_LINK_LIBRARIES --hip-link
)
set_property(TARGET hip::device APPEND PROPERTY
@@ -188,6 +196,20 @@ if(HIP_COMPILER STREQUAL "clang")
message("clang compiler doesn't support parallel jobs")
endif()
endif()
# Add support for __fp16 and _Float16, explicitly link with compiler-rt
set_property(TARGET hip::host APPEND PROPERTY
INTERFACE_COMPILE_OPTIONS -L${HIP_CLANG_INCLUDE_PATH}/../lib/linux -lclang_rt.builtins-x86_64
)
set_property(TARGET hip::host APPEND PROPERTY
INTERFACE_LINK_LIBRARIES -L${HIP_CLANG_INCLUDE_PATH}/../lib/linux -lclang_rt.builtins-x86_64
)
set_property(TARGET hip::device APPEND PROPERTY
INTERFACE_COMPILE_OPTIONS -L${HIP_CLANG_INCLUDE_PATH}/../lib/linux -lclang_rt.builtins-x86_64
)
set_property(TARGET hip::device APPEND PROPERTY
INTERFACE_LINK_LIBRARIES -L${HIP_CLANG_INCLUDE_PATH}/../lib/linux -lclang_rt.builtins-x86_64
)
endif()
set( hip_LIBRARIES hip::host hip::device)
include/hip/hcc_detail/device_functions.h
+6
Ver Arquivo
@@ -238,6 +238,12 @@ __device__ static inline unsigned int __lane_id() {
-1, __builtin_amdgcn_mbcnt_lo(-1, 0));
}
__device__
static inline unsigned int __mbcnt_lo(unsigned int x, unsigned int y) {return __builtin_amdgcn_mbcnt_lo(x,y);};
__device__
static inline unsigned int __mbcnt_hi(unsigned int x, unsigned int y) {return __builtin_amdgcn_mbcnt_hi(x,y);};
/*
HIP specific device functions
*/
include/hip/hcc_detail/hip_fp16.h
+65 -65
Ver Arquivo
@@ -529,35 +529,35 @@ THE SOFTWARE.
}
inline
__device__
__host__ __device__
__half __low2half(__half2 x)
{
return __half{__half_raw{static_cast<__half2_raw>(x).data.x}};
}
inline
__device__
__host__ __device__
__half __high2half(__half2 x)
{
return __half{__half_raw{static_cast<__half2_raw>(x).data.y}};
}
inline
__device__
__host__ __device__
__half2 __half2half2(__half x)
{
return __half2{x, x};
}
inline
__device__
__host__ __device__
__half2 __halves2half2(__half x, __half y)
{
return __half2{x, y};
}
inline
__device__
__host__ __device__
__half2 __low2half2(__half2 x)
{
return __half2{
@@ -567,7 +567,7 @@ THE SOFTWARE.
}
inline
__device__
__host__ __device__
__half2 __high2half2(__half2 x)
{
return __half2_raw{
@@ -577,7 +577,7 @@ THE SOFTWARE.
}
inline
__device__
__host__ __device__
__half2 __lows2half2(__half2 x, __half2 y)
{
return __half2_raw{
@@ -587,7 +587,7 @@ THE SOFTWARE.
}
inline
__device__
__host__ __device__
__half2 __highs2half2(__half2 x, __half2 y)
{
return __half2_raw{
@@ -597,7 +597,7 @@ THE SOFTWARE.
}
inline
__device__
__host__ __device__
__half2 __lowhigh2highlow(__half2 x)
{
return __half2_raw{
@@ -1046,16 +1046,16 @@ THE SOFTWARE.
__half __ldcs(const __half* ptr) { return *ptr; }
inline
__device__
__host__ __device__
__half2 __ldg(const __half2* ptr) { return *ptr; }
inline
__device__
__host__ __device__
__half2 __ldcg(const __half2* ptr) { return *ptr; }
inline
__device__
__host__ __device__
__half2 __ldca(const __half2* ptr) { return *ptr; }
inline
__device__
__host__ __device__
__half2 __ldcs(const __half2* ptr) { return *ptr; }
// Relations
@@ -1121,7 +1121,7 @@ THE SOFTWARE.
bool __hgtu(__half x, __half y) { return __hgt(x, y); }
inline
__device__
__host__ __device__
__half2 __heq2(__half2 x, __half2 y)
{
auto r = static_cast<__half2_raw>(x).data ==
@@ -1129,7 +1129,7 @@ THE SOFTWARE.
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
__host__ __device__
__half2 __hne2(__half2 x, __half2 y)
{
auto r = static_cast<__half2_raw>(x).data !=
@@ -1137,7 +1137,7 @@ THE SOFTWARE.
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
__host__ __device__
__half2 __hle2(__half2 x, __half2 y)
{
auto r = static_cast<__half2_raw>(x).data <=
@@ -1145,7 +1145,7 @@ THE SOFTWARE.
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
__host__ __device__
__half2 __hge2(__half2 x, __half2 y)
{
auto r = static_cast<__half2_raw>(x).data >=
@@ -1153,7 +1153,7 @@ THE SOFTWARE.
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
__host__ __device__
__half2 __hlt2(__half2 x, __half2 y)
{
auto r = static_cast<__half2_raw>(x).data <
@@ -1161,7 +1161,7 @@ THE SOFTWARE.
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
__host__ __device__
__half2 __hgt2(__half2 x, __half2 y)
{
auto r = static_cast<__half2_raw>(x).data >
@@ -1169,83 +1169,83 @@ THE SOFTWARE.
return __builtin_convertvector(-r, _Float16_2);
}
inline
__device__
__host__ __device__
__half2 __hequ2(__half2 x, __half2 y) { return __heq2(x, y); }
inline
__device__
__host__ __device__
__half2 __hneu2(__half2 x, __half2 y) { return __hne2(x, y); }
inline
__device__
__host__ __device__
__half2 __hleu2(__half2 x, __half2 y) { return __hle2(x, y); }
inline
__device__
__host__ __device__
__half2 __hgeu2(__half2 x, __half2 y) { return __hge2(x, y); }
inline
__device__
__host__ __device__
__half2 __hltu2(__half2 x, __half2 y) { return __hlt2(x, y); }
inline
__device__
__host__ __device__
__half2 __hgtu2(__half2 x, __half2 y) { return __hgt2(x, y); }
inline
__device__
__host__ __device__
bool __hbeq2(__half2 x, __half2 y)
{
auto r = static_cast<__half2_raw>(__heq2(x, y));
return r.data.x != 0 && r.data.y != 0;
}
inline
__device__
__host__ __device__
bool __hbne2(__half2 x, __half2 y)
{
auto r = static_cast<__half2_raw>(__hne2(x, y));
return r.data.x != 0 && r.data.y != 0;
}
inline
__device__
__host__ __device__
bool __hble2(__half2 x, __half2 y)
{
auto r = static_cast<__half2_raw>(__hle2(x, y));
return r.data.x != 0 && r.data.y != 0;
}
inline
__device__
__host__ __device__
bool __hbge2(__half2 x, __half2 y)
{
auto r = static_cast<__half2_raw>(__hge2(x, y));
return r.data.x != 0 && r.data.y != 0;
}
inline
__device__
__host__ __device__
bool __hblt2(__half2 x, __half2 y)
{
auto r = static_cast<__half2_raw>(__hlt2(x, y));
return r.data.x != 0 && r.data.y != 0;
}
inline
__device__
__host__ __device__
bool __hbgt2(__half2 x, __half2 y)
{
auto r = static_cast<__half2_raw>(__hgt2(x, y));
return r.data.x != 0 && r.data.y != 0;
}
inline
__device__
__host__ __device__
bool __hbequ2(__half2 x, __half2 y) { return __hbeq2(x, y); }
inline
__device__
__host__ __device__
bool __hbneu2(__half2 x, __half2 y) { return __hbne2(x, y); }
inline
__device__
__host__ __device__
bool __hbleu2(__half2 x, __half2 y) { return __hble2(x, y); }
inline
__device__
__host__ __device__
bool __hbgeu2(__half2 x, __half2 y) { return __hbge2(x, y); }
inline
__device__
__host__ __device__
bool __hbltu2(__half2 x, __half2 y) { return __hblt2(x, y); }
inline
__device__
__host__ __device__
bool __hbgtu2(__half2 x, __half2 y) { return __hbgt2(x, y); }
// Arithmetic
@@ -1334,7 +1334,7 @@ THE SOFTWARE.
}
inline
__device__
__host__ __device__
__half2 __hadd2(__half2 x, __half2 y)
{
return __half2_raw{
@@ -1342,14 +1342,14 @@ THE SOFTWARE.
static_cast<__half2_raw>(y).data};
}
inline
__device__
__host__ __device__
__half2 __habs2(__half2 x)
{
return __half2_raw{
__ocml_fabs_2f16(static_cast<__half2_raw>(x).data)};
}
inline
__device__
__host__ __device__
__half2 __hsub2(__half2 x, __half2 y)
{
return __half2_raw{
@@ -1357,7 +1357,7 @@ THE SOFTWARE.
static_cast<__half2_raw>(y).data};
}
inline
__device__
__host__ __device__
__half2 __hmul2(__half2 x, __half2 y)
{
return __half2_raw{
@@ -1365,7 +1365,7 @@ THE SOFTWARE.
static_cast<__half2_raw>(y).data};
}
inline
__device__
__host__ __device__
__half2 __hadd2_sat(__half2 x, __half2 y)
{
auto r = static_cast<__half2_raw>(__hadd2(x, y));
@@ -1374,7 +1374,7 @@ THE SOFTWARE.
__clamp_01(__half_raw{r.data.y})};
}
inline
__device__
__host__ __device__
__half2 __hsub2_sat(__half2 x, __half2 y)
{
auto r = static_cast<__half2_raw>(__hsub2(x, y));
@@ -1383,7 +1383,7 @@ THE SOFTWARE.
__clamp_01(__half_raw{r.data.y})};
}
inline
__device__
__host__ __device__
__half2 __hmul2_sat(__half2 x, __half2 y)
{
auto r = static_cast<__half2_raw>(__hmul2(x, y));
@@ -1392,13 +1392,13 @@ THE SOFTWARE.
__clamp_01(__half_raw{r.data.y})};
}
inline
__device__
__host__ __device__
__half2 __hfma2(__half2 x, __half2 y, __half2 z)
{
return __half2_raw{__ocml_fma_2f16(x, y, z)};
}
inline
__device__
__host__ __device__
__half2 __hfma2_sat(__half2 x, __half2 y, __half2 z)
{
auto r = static_cast<__half2_raw>(__hfma2(x, y, z));
@@ -1407,7 +1407,7 @@ THE SOFTWARE.
__clamp_01(__half_raw{r.data.y})};
}
inline
__device__
__host__ __device__
__half2 __h2div(__half2 x, __half2 y)
{
return __half2_raw{
@@ -1550,82 +1550,82 @@ THE SOFTWARE.
}
inline
__device__
__host__ __device__
__half2 h2trunc(__half2 x)
{
return __half2_raw{__ocml_trunc_2f16(x)};
}
inline
__device__
__host__ __device__
__half2 h2ceil(__half2 x)
{
return __half2_raw{__ocml_ceil_2f16(x)};
}
inline
__device__
__host__ __device__
__half2 h2floor(__half2 x)
{
return __half2_raw{__ocml_floor_2f16(x)};
}
inline
__device__
__host__ __device__
__half2 h2rint(__half2 x)
{
return __half2_raw{__ocml_rint_2f16(x)};
}
inline
__device__
__host__ __device__
__half2 h2sin(__half2 x)
{
return __half2_raw{__ocml_sin_2f16(x)};
}
inline
__device__
__host__ __device__
__half2 h2cos(__half2 x)
{
return __half2_raw{__ocml_cos_2f16(x)};
}
inline
__device__
__host__ __device__
__half2 h2exp(__half2 x)
{
return __half2_raw{__ocml_exp_2f16(x)};
}
inline
__device__
__host__ __device__
__half2 h2exp2(__half2 x)
{
return __half2_raw{__ocml_exp2_2f16(x)};
}
inline
__device__
__host__ __device__
__half2 h2exp10(__half2 x)
{
return __half2_raw{__ocml_exp10_2f16(x)};
}
inline
__device__
__host__ __device__
__half2 h2log2(__half2 x)
{
return __half2_raw{__ocml_log2_2f16(x)};
}
inline
__device__
__host__ __device__
__half2 h2log(__half2 x) { return __ocml_log_2f16(x); }
inline
__device__
__host__ __device__
__half2 h2log10(__half2 x) { return __ocml_log10_2f16(x); }
inline
__device__
__host__ __device__
__half2 h2rcp(__half2 x) { return __llvm_amdgcn_rcp_2f16(x); }
inline
__device__
__host__ __device__
__half2 h2rsqrt(__half2 x) { return __ocml_rsqrt_2f16(x); }
inline
__device__
__host__ __device__
__half2 h2sqrt(__half2 x) { return __ocml_sqrt_2f16(x); }
inline
__device__
__host__ __device__
__half2 __hisinf2(__half2 x)
{
auto r = __ocml_isinf_2f16(x);
@@ -1633,7 +1633,7 @@ THE SOFTWARE.
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
}
inline
__device__
__host__ __device__
__half2 __hisnan2(__half2 x)
{
auto r = __ocml_isnan_2f16(x);
@@ -1641,7 +1641,7 @@ THE SOFTWARE.
static_cast<_Float16>(r.x), static_cast<_Float16>(r.y)}};
}
inline
__device__
__host__ __device__
__half2 __hneg2(__half2 x)
{
return __half2_raw{-static_cast<__half2_raw>(x).data};
include/hip/hcc_detail/hip_fp16_math_fwd.h
+2 -1
Ver Arquivo
@@ -27,7 +27,7 @@ THE SOFTWARE.
// */
#include "host_defines.h"
#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
extern "C"
{
__device__ __attribute__((const)) _Float16 __ocml_ceil_f16(_Float16);
@@ -82,3 +82,4 @@ extern "C"
__device__ __attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
__device__ __attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
}
#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
include/hip/hcc_detail/hip_runtime.h
+47 -2
Ver Arquivo
@@ -313,6 +313,7 @@ static constexpr Coordinates<hip_impl::WorkitemId> threadIdx{};
#endif // defined __HCC__
#if __HCC_OR_HIP_CLANG__
#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
#if __HIP_ENABLE_DEVICE_MALLOC__
extern "C" __device__ void* __hip_malloc(size_t);
extern "C" __device__ void* __hip_free(void* ptr);
@@ -322,7 +323,7 @@ static inline __device__ void* free(void* ptr) { return __hip_free(ptr); }
static inline __device__ void* malloc(size_t size) { __builtin_trap(); return nullptr; }
static inline __device__ void* free(void* ptr) { __builtin_trap(); return nullptr; }
#endif
#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
#endif //__HCC_OR_HIP_CLANG__
#ifdef __HCC__
@@ -392,10 +393,53 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri
typedef int hipLaunchParm;
template <std::size_t n, typename... Ts,
typename std::enable_if<n == sizeof...(Ts)>::type* = nullptr>
void pArgs(const std::tuple<Ts...>&, void*) {}
template <std::size_t n, typename... Ts,
typename std::enable_if<n != sizeof...(Ts)>::type* = nullptr>
void pArgs(const std::tuple<Ts...>& formals, void** _vargs) {
using T = typename std::tuple_element<n, std::tuple<Ts...> >::type;
static_assert(!std::is_reference<T>{},
"A __global__ function cannot have a reference as one of its "
"arguments.");
#if defined(HIP_STRICT)
static_assert(std::is_trivially_copyable<T>{},
"Only TriviallyCopyable types can be arguments to a __global__ "
"function");
#endif
_vargs[n] = const_cast<void*>(reinterpret_cast<const void*>(&std::get<n>(formals)));
return pArgs<n + 1>(formals, _vargs);
}
template <typename... Formals, typename... Actuals>
std::tuple<Formals...> validateArgsCountType(void (*kernel)(Formals...), std::tuple<Actuals...>(actuals)) {
static_assert(sizeof...(Formals) == sizeof...(Actuals), "Argument Count Mismatch");
std::tuple<Formals...> to_formals{std::move(actuals)};
return to_formals;
}
#if defined(HIP_TEMPLATE_KERNEL_LAUNCH)
template <typename... Args, typename F = void (*)(Args...)>
void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t sharedMemBytes, hipStream_t stream, Args... args) {
constexpr size_t count = sizeof...(Args);
auto tup_ = std::tuple<Args...>{args...};
auto tup = validateArgsCountType(kernel, tup_);
void* _Args[count];
pArgs<0>(tup, _Args);
auto k = reinterpret_cast<void*>(kernel);
hipLaunchKernel(k, numBlocks, dimBlocks, _Args, sharedMemBytes, stream);
}
#else
#define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \
do { \
kernelName<<<(numblocks), (numthreads), (memperblock), (streamId)>>>(__VA_ARGS__); \
} while (0)
#endif
#include <hip/hip_runtime_api.h>
@@ -507,6 +551,7 @@ hc_get_workitem_absolute_id(int dim)
#endif
#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
// Support std::complex.
#if !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__
#pragma push_macro("__CUDA__")
@@ -524,7 +569,7 @@ hc_get_workitem_absolute_id(int dim)
#undef __CUDA__
#pragma pop_macro("__CUDA__")
#endif // !_OPENMP || __HIP_ENABLE_CUDA_WRAPPER_FOR_OPENMP__
#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
#endif // defined(__clang__) && defined(__HIP__)
#include <hip/hcc_detail/hip_memory.h>
include/hip/hcc_detail/hip_runtime_api.h
+140 -9
Ver Arquivo
@@ -178,8 +178,10 @@ enum hipLimit_t {
0x80000000 ///< Allocate non-coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific
///< allocation.
#define hipMemAttachGlobal 0x0
#define hipMemAttachHost 0x1
#define hipMemAttachGlobal 0x01 ///< Memory can be accessed by any stream on any device
#define hipMemAttachHost 0x02 ///< Memory cannot be accessed by any stream on any device
#define hipMemAttachSingle 0x04 ///< Memory can only be accessed by a single stream on
///< the associated device
#define hipDeviceMallocDefault 0x0
#define hipDeviceMallocFinegrained 0x1 ///< Memory is allocated in fine grained region of device.
@@ -217,6 +219,41 @@ enum hipLimit_t {
#define hipCooperativeLaunchMultiDeviceNoPreSync 0x01
#define hipCooperativeLaunchMultiDeviceNoPostSync 0x02
#define hipCpuDeviceId ((int)-1)
#define hipInvalidDeviceId ((int)-2)
/*
* @brief HIP Memory Advise values
* @enum
* @ingroup Enumerations
*/
typedef enum hipMemoryAdvise {
hipMemAdviseSetReadMostly = 1, ///< Data will mostly be read and only occassionally
///< be written to
hipMemAdviseUnsetReadMostly = 2, ///< Undo the effect of hipMemAdviseSetReadMostly
hipMemAdviseSetPreferredLocation = 3, ///< Set the preferred location for the data as
///< the specified device
hipMemAdviseUnsetPreferredLocation = 4, ///< Clear the preferred location for the data
hipMemAdviseSetAccessedBy = 5, ///< Data will be accessed by the specified device,
///< so prevent page faults as much as possible
hipMemAdviseUnsetAccessedBy = 6 ///< Let the Unified Memory subsystem decide on
///< the page faulting policy for the specified device
} hipMemoryAdvise;
/*
* @brief HIP range attributes
* @enum
* @ingroup Enumerations
*/
typedef enum hipMemRangeAttribute {
hipMemRangeAttributeReadMostly = 1, ///< Whether the range will mostly be read and
///< only occassionally be written to
hipMemRangeAttributePreferredLocation = 2, ///< The preferred location of the range
hipMemRangeAttributeAccessedBy = 3, ///< Memory range has cudaMemAdviseSetAccessedBy
///< set for specified device
hipMemRangeAttributeLastPrefetchLocation = 4,///< The last location to which the range was prefetched
} hipMemRangeAttribute;
/*
* @brief hipJitOption
* @enum
@@ -1180,15 +1217,18 @@ hipError_t hipMemAllocHost(void** ptr, size_t size);
hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags);
/**
* @brief Allocates memory that will be automatically managed by the Unified Memory system.
* @brief Allocates memory that will be automatically managed by AMD HMM.
*
* @param[out] ptr Pointer to the allocated managed memory
* @param[in] size Requested memory size
* @param[in] flags must be either hipMemAttachGlobal/hipMemAttachHost
* @param [out] dev_ptr - pointer to allocated device memory
* @param [in] size - requested allocation size in bytes
* @param [in] flags - must be either hipMemAttachGlobal or hipMemAttachHost
* (defaults to hipMemAttachGlobal)
*
* @return #hipSuccess, #hipErrorOutOfMemory
* @returns #hipSuccess, #hipErrorMemoryAllocation, #hipErrorNotSupported, #hipErrorInvalidValue
*/
hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags __dparm(0));
hipError_t hipMallocManaged(void** dev_ptr,
size_t size,
unsigned int flags __dparm(hipMemAttachGlobal));
/**
* @brief Allocate device accessible page locked host memory [Deprecated]
@@ -3369,7 +3409,6 @@ hipError_t __hipPopCallConfiguration(dim3 *gridDim,
* @returns #hipSuccess, #hipErrorInvalidValue, hipInvalidDevice
*
*/
hipError_t hipLaunchKernel(const void* function_address,
dim3 numBlocks,
dim3 dimBlocks,
@@ -3377,7 +3416,98 @@ hipError_t hipLaunchKernel(const void* function_address,
size_t sharedMemBytes __dparm(0),
hipStream_t stream __dparm(0));
/**
* @brief Prefetches memory to the specified destination device using AMD HMM.
*
* @param [in] dev_ptr pointer to be prefetched
* @param [in] count size in bytes for prefetching
* @param [in] device destination device to prefetch to
* @param [in] stream stream to enqueue prefetch operation
*
* @returns #hipSuccess, #hipErrorInvalidValue
*/
hipError_t hipMemPrefetchAsync(const void* dev_ptr,
size_t count,
int device,
hipStream_t stream __dparm(0));
/**
* @brief Advise about the usage of a given memory range to AMD HMM.
*
* @param [in] dev_ptr pointer to memory to set the advice for
* @param [in] count size in bytes of the memory range
* @param [in] advice advice to be applied for the specified memory range
* @param [in] device device to apply the advice for
*
* @returns #hipSuccess, #hipErrorInvalidValue
*/
hipError_t hipMemAdvise(const void* dev_ptr,
size_t count,
hipMemoryAdvise advice,
int device);
/**
* @brief Query an attribute of a given memory range in AMD HMM.
*
* @param [in/out] data a pointer to a memory location where the result of each
* attribute query will be written to
* @param [in] data_size the size of data
* @param [in] attribute the attribute to query
* @param [in] dev_ptr start of the range to query
* @param [in] count size of the range to query
*
* @returns #hipSuccess, #hipErrorInvalidValue
*/
hipError_t hipMemRangeGetAttribute(void* data,
size_t data_size,
hipMemRangeAttribute attribute,
const void* dev_ptr,
size_t count);
/**
* @brief Query attributes of a given memory range in AMD HMM.
*
* @param [in/out] data a two-dimensional array containing pointers to memory locations
* where the result of each attribute query will be written to
* @param [in] data_sizes an array, containing the sizes of each result
* @param [in] attributes the attribute to query
* @param [in] num_attributes an array of attributes to query (numAttributes and the number
* of attributes in this array should match)
* @param [in] dev_ptr start of the range to query
* @param [in] count size of the range to query
*
* @returns #hipSuccess, #hipErrorInvalidValue
*/
hipError_t hipMemRangeGetAttributes(void** data,
size_t* data_sizes,
hipMemRangeAttribute* attributes,
size_t num_attributes,
const void* dev_ptr,
size_t count);
/**
* @brief Attach memory to a stream asynchronously in AMD HMM.
*
* @param [in] stream - stream in which to enqueue the attach operation
* @param [in] dev_ptr - pointer to memory (must be a pointer to managed memory or
* to a valid host-accessible region of system-allocated memory)
* @param [in] length - length of memory (defaults to zero)
* @param [in] flags - must be one of cudaMemAttachGlobal, cudaMemAttachHost or
* cudaMemAttachSingle (defaults to cudaMemAttachSingle)
*
* @returns #hipSuccess, #hipErrorInvalidValue
*/
hipError_t hipStreamAttachMemAsync(hipStream_t stream,
hipDeviceptr_t* dev_ptr,
size_t length __dparm(0),
unsigned int flags __dparm(hipMemAttachSingle));
#if __HIP_ROCclr__ || !defined(__HCC__)
//TODO: Move this to hip_ext.h
hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, dim3 dimBlocks,
void** args, size_t sharedMemBytes, hipStream_t stream,
hipEvent_t startEvent, hipEvent_t stopEvent, int flags);
hipError_t hipBindTexture(
size_t* offset,
const textureReference* tex,
@@ -3657,6 +3787,7 @@ hipError_t hipRemoveActivityCallback(uint32_t id);
const char* hipApiName(uint32_t id);
const char* hipKernelNameRef(const hipFunction_t f);
const char* hipKernelNameRefByPtr(const void* hostFunction, hipStream_t stream);
int hipGetStreamDeviceId(hipStream_t stream);
#ifdef __cplusplus
} /* extern "C" */
#endif
include/hip/hcc_detail/hip_vector_types.h
+232 -79
Ver Arquivo
@@ -312,6 +312,21 @@ THE SOFTWARE.
using value_type = T;
__host__ __device__
HIP_vector_base() = default;
__host__ __device__
explicit
constexpr
HIP_vector_base(T x) noexcept : data{x} {}
__host__ __device__
constexpr
HIP_vector_base(const HIP_vector_base&) = default;
__host__ __device__
constexpr
HIP_vector_base(HIP_vector_base&&) = default;
__host__ __device__
~HIP_vector_base() = default;
__host__ __device__
HIP_vector_base& operator=(const HIP_vector_base& x) noexcept {
#if __has_attribute(ext_vector_type)
@@ -347,6 +362,24 @@ THE SOFTWARE.
using value_type = T;
__host__ __device__
HIP_vector_base() = default;
__host__ __device__
explicit
constexpr
HIP_vector_base(T x) noexcept : data{x, x} {}
__host__ __device__
constexpr
HIP_vector_base(T x, T y) noexcept : data{x, y} {}
__host__ __device__
constexpr
HIP_vector_base(const HIP_vector_base&) = default;
__host__ __device__
constexpr
HIP_vector_base(HIP_vector_base&&) = default;
__host__ __device__
~HIP_vector_base() = default;
__host__ __device__
HIP_vector_base& operator=(const HIP_vector_base& x) noexcept {
#if __has_attribute(ext_vector_type)
@@ -366,8 +399,8 @@ THE SOFTWARE.
T d[3];
__host__ __device__
constexpr
Native_vec_() = default;
__host__ __device__
explicit
constexpr
@@ -514,6 +547,29 @@ THE SOFTWARE.
};
using value_type = T;
__host__ __device__
HIP_vector_base() = default;
__host__ __device__
explicit
constexpr
HIP_vector_base(T x) noexcept : data{x, x, x} {}
__host__ __device__
constexpr
HIP_vector_base(T x, T y, T z) noexcept : data{x, y, z} {}
__host__ __device__
constexpr
HIP_vector_base(const HIP_vector_base&) = default;
__host__ __device__
constexpr
HIP_vector_base(HIP_vector_base&&) = default;
__host__ __device__
~HIP_vector_base() = default;
__host__ __device__
HIP_vector_base& operator=(const HIP_vector_base&) = default;
__host__ __device__
HIP_vector_base& operator=(HIP_vector_base&&) = default;
};
template<typename T>
@@ -538,11 +594,29 @@ THE SOFTWARE.
hip_impl::Scalar_accessor<T, Native_vec_, 1> y;
hip_impl::Scalar_accessor<T, Native_vec_, 2> z;
hip_impl::Scalar_accessor<T, Native_vec_, 3> w;
#endif
#endif
};
using value_type = T;
__host__ __device__
HIP_vector_base() = default;
__host__ __device__
explicit
constexpr
HIP_vector_base(T x) noexcept : data{x, x, x, x} {}
__host__ __device__
constexpr
HIP_vector_base(T x, T y, T z, T w) noexcept : data{x, y, z, w} {}
__host__ __device__
constexpr
HIP_vector_base(const HIP_vector_base&) = default;
__host__ __device__
constexpr
HIP_vector_base(HIP_vector_base&&) = default;
__host__ __device__
~HIP_vector_base() = default;
__host__ __device__
HIP_vector_base& operator=(const HIP_vector_base& x) noexcept {
#if __has_attribute(ext_vector_type)
@@ -563,49 +637,48 @@ THE SOFTWARE.
using HIP_vector_base<T, rank>::data;
using typename HIP_vector_base<T, rank>::Native_vec_;
inline __host__ __device__
__host__ __device__
HIP_vector_type() = default;
template<
typename U,
typename std::enable_if<
std::is_convertible<U, T>{}>::type* = nullptr>
explicit inline __host__ __device__
__host__ __device__
explicit
constexpr
HIP_vector_type(U x) noexcept
{
for (auto i = 0u; i != rank; ++i) data[i] = x;
}
: HIP_vector_base<T, rank>{static_cast<T>(x)}
{}
template< // TODO: constrain based on type as well.
typename... Us,
typename std::enable_if<
(rank > 1) && sizeof...(Us) == rank>::type* = nullptr>
inline __host__ __device__
__host__ __device__
constexpr
HIP_vector_type(Us... xs) noexcept
{
#if __has_attribute(ext_vector_type)
new (&data) Native_vec_{static_cast<T>(xs)...};
#else
new (&data) std::array<T, rank>{static_cast<T>(xs)...};
#endif
}
inline __host__ __device__
: HIP_vector_base<T, rank>{static_cast<T>(xs)...}
{}
__host__ __device__
constexpr
HIP_vector_type(const HIP_vector_type&) = default;
inline __host__ __device__
__host__ __device__
constexpr
HIP_vector_type(HIP_vector_type&&) = default;
inline __host__ __device__
__host__ __device__
~HIP_vector_type() = default;
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator=(const HIP_vector_type&) = default;
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator=(HIP_vector_type&&) = default;
// Operators
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator++() noexcept
{
return *this += HIP_vector_type{1};
}
inline __host__ __device__
__host__ __device__
HIP_vector_type operator++(int) noexcept
{
auto tmp(*this);
@@ -613,12 +686,12 @@ THE SOFTWARE.
return tmp;
}
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator--() noexcept
{
return *this -= HIP_vector_type{1};
}
inline __host__ __device__
__host__ __device__
HIP_vector_type operator--(int) noexcept
{
auto tmp(*this);
@@ -626,7 +699,7 @@ THE SOFTWARE.
return tmp;
}
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator+=(const HIP_vector_type& x) noexcept
{
data += x.data;
@@ -636,13 +709,13 @@ THE SOFTWARE.
typename U,
typename std::enable_if<
std::is_convertible<U, T>{}>::type* = nullptr>
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator+=(U x) noexcept
{
return *this += HIP_vector_type{x};
}
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator-=(const HIP_vector_type& x) noexcept
{
data -= x.data;
@@ -652,13 +725,13 @@ THE SOFTWARE.
typename U,
typename std::enable_if<
std::is_convertible<U, T>{}>::type* = nullptr>
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator-=(U x) noexcept
{
return *this -= HIP_vector_type{x};
}
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator*=(const HIP_vector_type& x) noexcept
{
data *= x.data;
@@ -668,13 +741,13 @@ THE SOFTWARE.
typename U,
typename std::enable_if<
std::is_convertible<U, T>{}>::type* = nullptr>
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator*=(U x) noexcept
{
return *this *= HIP_vector_type{x};
}
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator/=(const HIP_vector_type& x) noexcept
{
data /= x.data;
@@ -684,7 +757,7 @@ THE SOFTWARE.
typename U,
typename std::enable_if<
std::is_convertible<U, T>{}>::type* = nullptr>
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator/=(U x) noexcept
{
return *this /= HIP_vector_type{x};
@@ -693,7 +766,7 @@ THE SOFTWARE.
template<
typename U = T,
typename std::enable_if<std::is_signed<U>{}>::type* = nullptr>
inline __host__ __device__
__host__ __device__
HIP_vector_type operator-() const noexcept
{
auto tmp(*this);
@@ -704,7 +777,7 @@ THE SOFTWARE.
template<
typename U = T,
typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
inline __host__ __device__
__host__ __device__
HIP_vector_type operator~() const noexcept
{
HIP_vector_type r{*this};
@@ -715,7 +788,7 @@ THE SOFTWARE.
template<
typename U = T,
typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator%=(const HIP_vector_type& x) noexcept
{
data %= x.data;
@@ -725,7 +798,7 @@ THE SOFTWARE.
template<
typename U = T,
typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator^=(const HIP_vector_type& x) noexcept
{
data ^= x.data;
@@ -735,7 +808,7 @@ THE SOFTWARE.
template<
typename U = T,
typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator|=(const HIP_vector_type& x) noexcept
{
data |= x.data;
@@ -745,7 +818,7 @@ THE SOFTWARE.
template<
typename U = T,
typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator&=(const HIP_vector_type& x) noexcept
{
data &= x.data;
@@ -755,7 +828,7 @@ THE SOFTWARE.
template<
typename U = T,
typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator>>=(const HIP_vector_type& x) noexcept
{
data >>= x.data;
@@ -765,7 +838,7 @@ THE SOFTWARE.
template<
typename U = T,
typename std::enable_if<std::is_integral<U>{}>::type* = nullptr>
inline __host__ __device__
__host__ __device__
HIP_vector_type& operator<<=(const HIP_vector_type& x) noexcept
{
data <<= x.data;
@@ -774,21 +847,27 @@ THE SOFTWARE.
};
template<typename T, unsigned int n>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator+(
const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
{
return HIP_vector_type<T, n>{x} += y;
}
template<typename T, unsigned int n, typename U>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator+(
const HIP_vector_type<T, n>& x, U y) noexcept
{
return HIP_vector_type<T, n>{x} += HIP_vector_type<T, n>{y};
}
template<typename T, unsigned int n, typename U>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator+(
U x, const HIP_vector_type<T, n>& y) noexcept
{
@@ -796,21 +875,27 @@ THE SOFTWARE.
}
template<typename T, unsigned int n>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator-(
const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
{
return HIP_vector_type<T, n>{x} -= y;
}
template<typename T, unsigned int n, typename U>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator-(
const HIP_vector_type<T, n>& x, U y) noexcept
{
return HIP_vector_type<T, n>{x} -= HIP_vector_type<T, n>{y};
}
template<typename T, unsigned int n, typename U>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator-(
U x, const HIP_vector_type<T, n>& y) noexcept
{
@@ -818,21 +903,27 @@ THE SOFTWARE.
}
template<typename T, unsigned int n>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator*(
const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
{
return HIP_vector_type<T, n>{x} *= y;
}
template<typename T, unsigned int n, typename U>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator*(
const HIP_vector_type<T, n>& x, U y) noexcept
{
return HIP_vector_type<T, n>{x} *= HIP_vector_type<T, n>{y};
}
template<typename T, unsigned int n, typename U>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator*(
U x, const HIP_vector_type<T, n>& y) noexcept
{
@@ -840,64 +931,90 @@ THE SOFTWARE.
}
template<typename T, unsigned int n>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator/(
const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
{
return HIP_vector_type<T, n>{x} /= y;
}
template<typename T, unsigned int n, typename U>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator/(
const HIP_vector_type<T, n>& x, U y) noexcept
{
return HIP_vector_type<T, n>{x} /= HIP_vector_type<T, n>{y};
}
template<typename T, unsigned int n, typename U>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator/(
U x, const HIP_vector_type<T, n>& y) noexcept
{
return HIP_vector_type<T, n>{x} /= y;
}
template<typename V>
__host__ __device__
inline
constexpr
bool _hip_any_zero(const V& x, int n) noexcept
{
return
(n == -1) ? true : ((x[n] == 0) ? false : _hip_any_zero(x, n - 1));
}
template<typename T, unsigned int n>
inline __host__ __device__
__host__ __device__
inline
constexpr
bool operator==(
const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
{
auto tmp = x.data == y.data;
for (auto i = 0u; i != n; ++i) if (tmp[i] == 0) return false;
return true;
return _hip_any_zero(x.data == y.data, n - 1);
}
template<typename T, unsigned int n, typename U>
inline __host__ __device__
__host__ __device__
inline
constexpr
bool operator==(const HIP_vector_type<T, n>& x, U y) noexcept
{
return x == HIP_vector_type<T, n>{y};
}
template<typename T, unsigned int n, typename U>
inline __host__ __device__
__host__ __device__
inline
constexpr
bool operator==(U x, const HIP_vector_type<T, n>& y) noexcept
{
return HIP_vector_type<T, n>{x} == y;
}
template<typename T, unsigned int n>
inline __host__ __device__
__host__ __device__
inline
constexpr
bool operator!=(
const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
{
return !(x == y);
}
template<typename T, unsigned int n, typename U>
inline __host__ __device__
__host__ __device__
inline
constexpr
bool operator!=(const HIP_vector_type<T, n>& x, U y) noexcept
{
return !(x == y);
}
template<typename T, unsigned int n, typename U>
inline __host__ __device__
__host__ __device__
inline
constexpr
bool operator!=(U x, const HIP_vector_type<T, n>& y) noexcept
{
return !(x == y);
@@ -907,7 +1024,9 @@ THE SOFTWARE.
typename T,
unsigned int n,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator%(
const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
{
@@ -918,7 +1037,9 @@ THE SOFTWARE.
unsigned int n,
typename U,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator%(
const HIP_vector_type<T, n>& x, U y) noexcept
{
@@ -929,7 +1050,9 @@ THE SOFTWARE.
unsigned int n,
typename U,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator%(
U x, const HIP_vector_type<T, n>& y) noexcept
{
@@ -940,7 +1063,9 @@ THE SOFTWARE.
typename T,
unsigned int n,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator^(
const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
{
@@ -951,7 +1076,9 @@ THE SOFTWARE.
unsigned int n,
typename U,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator^(
const HIP_vector_type<T, n>& x, U y) noexcept
{
@@ -962,7 +1089,9 @@ THE SOFTWARE.
unsigned int n,
typename U,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator^(
U x, const HIP_vector_type<T, n>& y) noexcept
{
@@ -973,7 +1102,9 @@ THE SOFTWARE.
typename T,
unsigned int n,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator|(
const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
{
@@ -984,7 +1115,9 @@ THE SOFTWARE.
unsigned int n,
typename U,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator|(
const HIP_vector_type<T, n>& x, U y) noexcept
{
@@ -995,7 +1128,9 @@ THE SOFTWARE.
unsigned int n,
typename U,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator|(
U x, const HIP_vector_type<T, n>& y) noexcept
{
@@ -1006,7 +1141,9 @@ THE SOFTWARE.
typename T,
unsigned int n,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator&(
const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
{
@@ -1017,7 +1154,9 @@ THE SOFTWARE.
unsigned int n,
typename U,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator&(
const HIP_vector_type<T, n>& x, U y) noexcept
{
@@ -1028,7 +1167,9 @@ THE SOFTWARE.
unsigned int n,
typename U,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator&(
U x, const HIP_vector_type<T, n>& y) noexcept
{
@@ -1039,7 +1180,9 @@ THE SOFTWARE.
typename T,
unsigned int n,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator>>(
const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
{
@@ -1050,7 +1193,9 @@ THE SOFTWARE.
unsigned int n,
typename U,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator>>(
const HIP_vector_type<T, n>& x, U y) noexcept
{
@@ -1061,7 +1206,9 @@ THE SOFTWARE.
unsigned int n,
typename U,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator>>(
U x, const HIP_vector_type<T, n>& y) noexcept
{
@@ -1072,7 +1219,9 @@ THE SOFTWARE.
typename T,
unsigned int n,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator<<(
const HIP_vector_type<T, n>& x, const HIP_vector_type<T, n>& y) noexcept
{
@@ -1083,7 +1232,9 @@ THE SOFTWARE.
unsigned int n,
typename U,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator<<(
const HIP_vector_type<T, n>& x, U y) noexcept
{
@@ -1095,7 +1246,9 @@ THE SOFTWARE.
typename U,
typename std::enable_if<std::is_arithmetic<U>::value>::type,
typename std::enable_if<std::is_integral<T>{}>* = nullptr>
inline __host__ __device__
__host__ __device__
inline
constexpr
HIP_vector_type<T, n> operator<<(
U x, const HIP_vector_type<T, n>& y) noexcept
{
include/hip/hcc_detail/host_defines.h
+2
Ver Arquivo
@@ -64,11 +64,13 @@ THE SOFTWARE.
#elif defined(__clang__) && defined(__HIP__)
#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
#define __host__ __attribute__((host))
#define __device__ __attribute__((device))
#define __global__ __attribute__((global))
#define __shared__ __attribute__((shared))
#define __constant__ __attribute__((constant))
#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
#define __noinline__ __attribute__((noinline))
#define __forceinline__ inline __attribute__((always_inline))
include/hip/hcc_detail/math_functions.h
+4 -1
Ver Arquivo
@@ -71,6 +71,7 @@ struct __numeric_type<_Float16>
#define __RETURN_TYPE bool
#endif
#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
__DEVICE__
inline
uint64_t __make_mantissa_base8(const char* tagp)
@@ -139,6 +140,7 @@ uint64_t __make_mantissa(const char* tagp)
return __make_mantissa_base10(tagp);
}
#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
// DOT FUNCTIONS
#if (__hcc_workweek__ >= 19015) || __HIP_CLANG_ONLY__
@@ -174,6 +176,7 @@ uint amd_mixed_dot(uint a, uint b, uint c, bool saturate) {
}
#endif
#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
// BEGIN FLOAT
__DEVICE__
inline
@@ -1507,7 +1510,7 @@ __host__ inline static int min(int arg1, int arg2) {
__host__ inline static int max(int arg1, int arg2) {
return std::max(arg1, arg2);
}
#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
#pragma pop_macro("__DEF_FLOAT_FUN")
#pragma pop_macro("__DEF_FLOAT_FUN2")
include/hip/hcc_detail/math_fwd.h
+3 -1
Ver Arquivo
@@ -23,7 +23,6 @@ THE SOFTWARE.
#pragma once
#include "host_defines.h"
#if defined(__cplusplus)
extern "C" {
#endif
@@ -67,6 +66,7 @@ __attribute__((const))
unsigned int __ockl_udot8(unsigned int, unsigned int, unsigned int, bool);
#endif
#if !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
// BEGIN FLOAT
__device__
__attribute__((const))
@@ -701,6 +701,8 @@ double __llvm_amdgcn_rsq_f64(double) __asm("llvm.amdgcn.rsq.f64");
// END INTRINSICS
// END DOUBLE
#endif // !__CLANG_HIP_RUNTIME_WRAPPER_INCLUDED__
#if defined(__cplusplus)
} // extern "C"
#endif
include/hip/hip_ext.h
+27 -2
Ver Arquivo
@@ -23,6 +23,10 @@ THE SOFTWARE.
#ifndef HIP_INCLUDE_HIP_HIP_EXT_H
#define HIP_INCLUDE_HIP_HIP_EXT_H
#include "hip/hip_runtime.h"
#if defined(__cplusplus)
#include <tuple>
#include <type_traits>
#endif
#ifdef __HCC__
// Forward declarations:
@@ -109,8 +113,29 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
hipEvent_t stopEvent = nullptr)
__attribute__((deprecated("use hipExtModuleLaunchKernel instead")));
//#if !__HIP_ROCclr__ && defined(__cplusplus)
#if defined(__HIP_PLATFORM_HCC__) && GENERIC_GRID_LAUNCH == 1 && defined(__HCC__)
#if defined(__HIP_ROCclr__) && defined(__cplusplus)
extern "C" hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks,
dim3 dimBlocks, void** args, size_t sharedMemBytes,
hipStream_t stream, hipEvent_t startEvent,
hipEvent_t stopEvent, int flags);
template <typename... Args, typename F = void (*)(Args...)>
inline void hipExtLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
std::uint32_t sharedMemBytes, hipStream_t stream,
hipEvent_t startEvent, hipEvent_t stopEvent, std::uint32_t flags,
Args... args) {
constexpr size_t count = sizeof...(Args);
auto tup_ = std::tuple<Args...>{args...};
auto tup = validateArgsCountType(kernel, tup_);
void* _Args[count];
pArgs<0>(tup, _Args);
auto k = reinterpret_cast<void*>(kernel);
hipExtLaunchKernel(k, numBlocks, dimBlocks, _Args, sharedMemBytes, stream, startEvent,
stopEvent, (int)flags);
}
#elif defined(__HIP_PLATFORM_HCC__) && GENERIC_GRID_LAUNCH == 1 && defined(__HCC__)
//kernel_descriptor and hip_impl::make_kernarg are in "grid_launch_GGL.hpp"
namespace hip_impl {
include/hip/hip_runtime_api.h
+1
Ver Arquivo
@@ -114,6 +114,7 @@ typedef struct hipDeviceProp_t {
int isMultiGpuBoard; ///< 1 if device is on a multi-GPU board, 0 if not.
int canMapHostMemory; ///< Check whether HIP can map host memory
int gcnArch; ///< AMD GCN Arch Value. Eg: 803, 701
char gcnArchName[256]; ///< AMD GCN Arch Name.
int integrated; ///< APU vs dGPU
int cooperativeLaunch; ///< HIP device supports cooperative launch
int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple devices
rocclr/CMakeLists.txt
Arquivo normal → Arquivo executável
+14 -2
Ver Arquivo
@@ -21,8 +21,17 @@ if(CMAKE_BUILD_TYPE MATCHES "^Debug$")
add_definitions(-DDEBUG)
endif()
if (CMAKE_CXX_COMPILER_ID STREQUAL "GNU")
add_compile_options("-Wno-ignored-attributes")
if((CMAKE_CXX_COMPILER_ID STREQUAL "GNU") OR
(CMAKE_${COMPILER}_COMPILER_ID MATCHES "Clang"))
add_definitions(
# Enabling -Wextra or -pedantic will cause
# thousands of warnings. Keep things simple for now.
-Wall
# This one seems impossible to fix for now.
# There are hundreds of instances of unused vars/functions
# throughout the code base.
-Wno-unused-variable
-Wno-unused-function)
endif()
set(USE_PROF_API "1")
@@ -75,10 +84,13 @@ add_definitions(-DBSD_LIBELF)
add_library(hip64 OBJECT
hip_context.cpp
hip_code_object.cpp
hip_device.cpp
hip_device_runtime.cpp
hip_error.cpp
hip_event.cpp
hip_global.cpp
hip_hmm.cpp
hip_memory.cpp
hip_module.cpp
hip_peer.cpp
rocclr/hip_code_object.cpp
Arquivo executável
+457
Ver Arquivo
@@ -0,0 +1,457 @@
#include "hip_code_object.hpp"
#include <cstring>
#include "hip/hip_runtime_api.h"
#include "hip/hip_runtime.h"
#include "hip_internal.hpp"
#include "platform/program.hpp"
namespace hip {
uint64_t CodeObject::ElfSize(const void *emi) {
const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi;
const Elf64_Shdr *shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff);
uint64_t max_offset = ehdr->e_shoff;
uint64_t total_size = max_offset + ehdr->e_shentsize * ehdr->e_shnum;
for (uint16_t i=0; i < ehdr->e_shnum; ++i){
uint64_t cur_offset = static_cast<uint64_t>(shdr[i].sh_offset);
if (max_offset < cur_offset) {
max_offset = cur_offset;
total_size = max_offset;
if(SHT_NOBITS != shdr[i].sh_type) {
total_size += static_cast<uint64_t>(shdr[i].sh_size);
}
}
}
return total_size;
}
bool CodeObject::isCompatibleCodeObject(const std::string& codeobj_target_id,
const char* device_name) {
// Workaround for device name mismatch.
// Device name may contain feature strings delimited by '+', e.g.
// gfx900+xnack. Currently HIP-Clang does not include feature strings
// in code object target id in fat binary. Therefore drop the feature
// strings from device name before comparing it with code object target id.
std::string short_name(device_name);
auto feature_loc = short_name.find('+');
if (feature_loc != std::string::npos) {
short_name.erase(feature_loc);
}
return codeobj_target_id == short_name;
}
hipError_t CodeObject::extractCodeObjectFromFatBinary(const void* data,
const std::vector<const char*>& devices,
std::vector<std::pair<const void*, size_t>>& code_objs) {
std::string magic((const char*)data, sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1);
if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR)) {
return hipErrorInvalidKernelFile;
}
code_objs.resize(devices.size());
const auto obheader = reinterpret_cast<const __ClangOffloadBundleHeader*>(data);
const auto* desc = &obheader->desc[0];
unsigned num_code_objs = 0;
for (uint64_t i = 0; i < obheader->numBundles; ++i,
desc = reinterpret_cast<const __ClangOffloadBundleDesc*>(
reinterpret_cast<uintptr_t>(&desc->triple[0]) + desc->tripleSize)) {
std::size_t offset = 0;
if (!std::strncmp(desc->triple, HIP_AMDGCN_AMDHSA_TRIPLE,
sizeof(HIP_AMDGCN_AMDHSA_TRIPLE) - 1)) {
offset = sizeof(HIP_AMDGCN_AMDHSA_TRIPLE); //For code objects created by CLang
} else if (!std::strncmp(desc->triple, HCC_AMDGCN_AMDHSA_TRIPLE,
sizeof(HCC_AMDGCN_AMDHSA_TRIPLE) - 1)) {
offset = sizeof(HCC_AMDGCN_AMDHSA_TRIPLE); //For code objects created by Hcc
} else {
continue;
}
std::string target(desc->triple + offset, desc->tripleSize - offset);
const void *image = reinterpret_cast<const void*>(
reinterpret_cast<uintptr_t>(obheader) + desc->offset);
size_t size = desc->size;
for (size_t dev = 0; dev < devices.size(); ++dev) {
const char* name = devices[dev];
if (!isCompatibleCodeObject(target, name)) {
continue;
}
code_objs[dev] = std::make_pair(image, size);
num_code_objs++;
}
}
if (num_code_objs == devices.size()) {
return hipSuccess;
} else {
DevLogError("hipErrorNoBinaryForGpu: Coudn't find binary for current devices!");
guarantee(false);
return hipErrorNoBinaryForGpu;
}
}
hipError_t CodeObject::add_program(int deviceId, hipModule_t hmod, const void* binary_ptr,
size_t binary_size) {
amd::Program* program = as_amd(reinterpret_cast<cl_program>(hmod));
amd::Context* ctx = g_devices[deviceId]->asContext();
if (CL_SUCCESS != program->addDeviceProgram(*ctx->devices()[0], binary_ptr,
binary_size, false)) {
return hipErrorNotFound;
}
return hipSuccess;
}
hipError_t CodeObject::build_module(hipModule_t hmod, const std::vector<amd::Device*>& devices) {
amd::Program* program = as_amd(reinterpret_cast<cl_program>(hmod));
program->setVarInfoCallBack(&getSvarInfo);
if (CL_SUCCESS != program->build(devices, nullptr, nullptr, nullptr, kOptionChangeable, kNewDevProg)) {
DevLogPrintfError("Build error for module: 0x%x \n", hmod);
return hipErrorSharedObjectInitFailed;
}
return hipSuccess;
}
DynCO::DynCO(): program_(nullptr) {}
hipError_t DynCO::loadCodeObject(const char* fname, const void* image) {
amd::ScopedLock lock(dclock_);
const void *mmap_ptr = nullptr;
size_t mmap_size = 0;
guarantee(fname || image);
if (fname != nullptr) {
/* We are given file name */
if (!amd::Os::MemoryMapFile(fname, &mmap_ptr, &mmap_size)) {
return hipErrorFileNotFound;
}
} else if (image != nullptr) {
/*We are directly given image pointer directly */
mmap_ptr = image;
} else {
return hipErrorMissingConfiguration;
}
return loadCodeObjectData(mmap_ptr, mmap_size);
}
//Dynamic Code Object
DynCO::~DynCO() {
amd::ScopedLock lock(dclock_);
if (program_ != nullptr) {
program_->release();
program_ = nullptr;
}
for (auto& elem : vars_) {
delete elem.second;
}
vars_.clear();
for (auto& elem : functions_) {
delete elem.second;
}
functions_.clear();
}
hipError_t DynCO::getDeviceVar(DeviceVar** dvar, std::string var_name, int device_id) {
amd::ScopedLock lock(dclock_);
auto it = vars_.find(var_name);
if (it == vars_.end()) {
DevLogPrintfError("Cannot find the Var: %s ", var_name.c_str());
return hipErrorNotFound;
}
it->second->getDeviceVar(dvar, device_id, module());
return hipSuccess;
}
hipError_t DynCO::getDynFunc(hipFunction_t* hfunc, std::string func_name) {
amd::ScopedLock lock(dclock_);
auto it = functions_.find(func_name);
if (it == functions_.end()) {
DevLogPrintfError("Cannot find the function: %s ", func_name.c_str());
return hipErrorNotFound;
}
/* See if this could be solved */
return it->second->getDynFunc(hfunc, reinterpret_cast<hipModule_t>(as_cl(program_)));
}
hipError_t DynCO::loadCodeObjectData(const void* mmap_ptr, size_t mmap_size) {
amd::ScopedLock lock(dclock_);
/* initialize image it to the mmap_ptr, if this is of no_clang_offload
bundle then they directly pass the image */
const void* image = mmap_ptr;
std::vector<std::pair<const void*, size_t>> code_objs;
hipError_t hip_error = extractCodeObjectFromFatBinary(mmap_ptr,
{hip::getCurrentDevice()->devices()[0]->info().name_},
code_objs);
if (hip_error == hipSuccess) {
image = code_objs[0].first;
} else if(hip_error == hipErrorNoBinaryForGpu) {
return hip_error;
}
program_ = new amd::Program(*hip::getCurrentDevice()->asContext(),
amd::Program::Language::Binary, mmap_ptr, mmap_size);
if (program_ == NULL) {
return hipErrorOutOfMemory;
}
program_->setVarInfoCallBack(&getSvarInfo);
if (CL_SUCCESS != program_->addDeviceProgram(*hip::getCurrentDevice()->devices()[0], image,
ElfSize(image), false)) {
return hipErrorInvalidKernelFile;
}
//This has to happen before Program has been built, other wise undef vars fail.
IHIP_RETURN_ONFAIL(populateDynGlobalVars());
//program->setVarInfoCallBack(&getSvarInfo);
if(CL_SUCCESS != program_->build(hip::getCurrentDevice()->devices(), nullptr, nullptr, nullptr,
kOptionChangeable, kNewDevProg)) {
return hipErrorSharedObjectInitFailed;
}
//This has to happen after Program has been built, other wise symbolTable_ not populated.
IHIP_RETURN_ONFAIL(populateDynGlobalFuncs());
return hipSuccess;
}
hipError_t DynCO::populateDynGlobalVars() {
amd::ScopedLock lock(dclock_);
std::vector<std::string> var_names;
std::vector<std::string> undef_var_names;
device::Program* dev_program
= program_->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]);
if (!dev_program->getGlobalVarFromCodeObj(&var_names)) {
DevLogPrintfError("Could not get Global vars from Code Obj for Module: 0x%x \n", module());
return hipErrorSharedObjectSymbolNotFound;
}
if (!dev_program->getUndefinedVarFromCodeObj(&undef_var_names)) {
DevLogPrintfError("Could not get undefined Variables for Module: 0x%x \n", module());
return hipErrorSharedObjectSymbolNotFound;
}
for (auto& elem : var_names) {
vars_.insert(std::make_pair(elem, new Var(elem, Var::DeviceVarKind::DVK_Variable, 0, 0, 0, nullptr)));
}
for (auto& elem : undef_var_names) {
vars_.insert(std::make_pair(elem, new Var(elem, Var::DeviceVarKind::DVK_Texture, 0, 0, 0, nullptr)));
}
return hipSuccess;
}
hipError_t DynCO::populateDynGlobalFuncs() {
amd::ScopedLock lock(dclock_);
std::vector<std::string> func_names;
device::Program* dev_program
= program_->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]);
// Get all the global func names from COMGR
if (!dev_program->getGlobalFuncFromCodeObj(&func_names)) {
DevLogPrintfError("Could not get Global Funcs from Code Obj for Module: 0x%x \n", module());
return hipErrorSharedObjectSymbolNotFound;
}
for (auto& elem : func_names) {
functions_.insert(std::make_pair(elem, new Function(elem)));
}
return hipSuccess;
}
//Static Code Object
StatCO::StatCO() {
}
StatCO::~StatCO() {
amd::ScopedLock lock(sclock_);
for (auto& elem : functions_) {
delete elem.second;
}
functions_.clear();
for (auto& elem : vars_) {
delete elem.second;
}
vars_.clear();
}
hipError_t StatCO::digestFatBinary(const void* data, FatBinaryInfoType& programs) {
amd::ScopedLock lock(sclock_);
if (programs.size() > 0) {
return hipSuccess;
}
std::vector<std::pair<const void*, size_t>> code_objs;
std::vector<const char*> devices;
for (size_t dev = 0; dev < g_devices.size(); ++dev) {
devices.push_back(g_devices[dev]->devices()[0]->info().name_);
}
IHIP_RETURN_ONFAIL(extractCodeObjectFromFatBinary((char*)data, devices, code_objs));
programs.resize(g_devices.size());
for (size_t dev = 0; dev < g_devices.size(); ++dev) {
amd::Context* ctx = g_devices[dev]->asContext();
amd::Program* program = new amd::Program(*ctx);
if (program == nullptr) {
return hipErrorOutOfMemory;
}
programs.at(dev) = std::make_pair(reinterpret_cast<hipModule_t>(as_cl(program)),
new FatBinaryMetaInfo(false, code_objs[dev].first, code_objs[dev].second));
}
return hipSuccess;
}
FatBinaryInfoType* StatCO::addFatBinary(const void* data, bool initialized) {
amd::ScopedLock lock(sclock_);
if (initialized) {
digestFatBinary(data, modules_[data]);
}
return &modules_[data];
}
hipError_t StatCO::removeFatBinary(FatBinaryInfoType* module) {
amd::ScopedLock lock(sclock_);
auto vit = vars_.begin();
while (vit != vars_.end()) {
if (vit->second->moduleInfo() == module) {
delete vit->second;
vit = vars_.erase(vit);
} else {
++vit;
}
}
auto fit = functions_.begin();
while (fit != functions_.end()) {
if (fit->second->moduleInfo() == module) {
delete fit->second;
fit = functions_.erase(fit);
} else {
++fit;
}
}
auto mit = modules_.begin();
while (mit != modules_.end()) {
if (&mit->second == module) {
for (size_t dev=0; dev < g_devices.size(); ++dev) {
delete (*module)[dev].second;
}
mit = modules_.erase(mit);
} else {
++mit;
}
}
return hipSuccess;
}
hipError_t StatCO::registerStatFunction(const void* hostFunction, Function* func) {
amd::ScopedLock lock(sclock_);
if (functions_.find(hostFunction) != functions_.end()) {
DevLogPrintfError("hostFunctionPtr: 0x%x already exists", hostFunction);
}
functions_.insert(std::make_pair(hostFunction, func));
return hipSuccess;
}
hipError_t StatCO::getStatFunc(hipFunction_t* hfunc, const void* hostFunction, int deviceId) {
amd::ScopedLock lock(sclock_);
const auto it = functions_.find(hostFunction);
if (it == functions_.end()) {
return hipErrorInvalidSymbol;
}
return it->second->getStatFunc(hfunc, deviceId);
}
hipError_t StatCO::getStatFuncAttr(hipFuncAttributes* func_attr, const void* hostFunction, int deviceId) {
amd::ScopedLock lock(sclock_);
const auto it = functions_.find(hostFunction);
if (it == functions_.end()) {
return hipErrorInvalidSymbol;
}
return it->second->getStatFuncAttr(func_attr, deviceId);
}
hipError_t StatCO::registerStatGlobalVar(const void* hostVar, Var* var) {
amd::ScopedLock lock(sclock_);
if (vars_.find(hostVar) != vars_.end()) {
return hipErrorInvalidSymbol;
}
vars_.insert(std::make_pair(hostVar, var));
return hipSuccess;
}
hipError_t StatCO::getStatGlobalVar(const void* hostVar, int deviceId, hipDeviceptr_t* dev_ptr,
size_t* size_ptr) {
amd::ScopedLock lock(sclock_);
const auto it = vars_.find(hostVar);
if (it == vars_.end()) {
return hipErrorInvalidSymbol;
}
DeviceVar* dvar = nullptr;
IHIP_RETURN_ONFAIL(it->second->getStatDeviceVar(&dvar, deviceId));
*dev_ptr = dvar->device_ptr();
*size_ptr = dvar->size();
return hipSuccess;
}
hipError_t StatCO::getStatGlobalVarByName(std::string hostVar, int deviceId, hipModule_t hmod,
hipDeviceptr_t* dev_ptr, size_t* size_ptr) {
amd::ScopedLock lock(sclock_);
for (auto& elem : vars_) {
if ((elem.second->name() == hostVar)
&& (elem.second->module(deviceId) == hmod)) {
*dev_ptr = elem.second->device_ptr(deviceId);
*size_ptr = elem.second->device_size(deviceId);
return hipSuccess;
}
}
return hipErrorNotFound;
}
}; //namespace: hip
rocclr/hip_code_object.hpp
Arquivo executável
+132
Ver Arquivo
@@ -0,0 +1,132 @@
#ifndef HIP_CODE_OBJECT_HPP
#define HIP_CODE_OBJECT_HPP
#include "hip_global.hpp"
#include <unordered_map>
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include "hip_internal.hpp"
#include "device/device.hpp"
#include "platform/program.hpp"
//Forward Declaration for friend usage
class PlatformState;
namespace hip {
//Code Object base class
class CodeObject {
public:
virtual ~CodeObject() {}
//Functions to add_dev_prog and build
static hipError_t add_program(int deviceId, hipModule_t hmod, const void* binary_ptr,
size_t binary_size);
static hipError_t build_module(hipModule_t hmod, const std::vector<amd::Device*>& devices);
//ClangOFFLOADBundle info
#define CLANG_OFFLOAD_BUNDLER_MAGIC_STR "__CLANG_OFFLOAD_BUNDLE__"
#define HIP_AMDGCN_AMDHSA_TRIPLE "hip-amdgcn-amd-amdhsa"
#define HCC_AMDGCN_AMDHSA_TRIPLE "hcc-amdgcn-amd-amdhsa-"
//Clang Offload bundler description & Header
struct __ClangOffloadBundleDesc {
uint64_t offset;
uint64_t size;
uint64_t tripleSize;
const char triple[1];
};
struct __ClangOffloadBundleHeader {
const char magic[sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1];
uint64_t numBundles;
__ClangOffloadBundleDesc desc[1];
};
protected:
CodeObject() {}
//Given an ptr to image or file, extracts to code object
//for corresponding devices
hipError_t extractCodeObjectFromFatBinary(const void*,
const std::vector<const char*>&,
std::vector<std::pair<const void*, size_t>>&);
uint64_t ElfSize(const void* emi);
private:
bool isCompatibleCodeObject(const std::string& codeobj_target_id,
const char* device_name);
friend const std::vector<hipModule_t>& modules();
};
//Dynamic Code Object
class DynCO : public CodeObject {
amd::Monitor dclock_{"Guards Static Code object", true};
public:
DynCO();
virtual ~DynCO();
//LoadsCodeObject and its data
hipError_t loadCodeObject(const char* fname, const void* image=nullptr);
hipModule_t module() { return reinterpret_cast<hipModule_t>(as_cl(program_)); };
//Gets GlobalVar/Functions from a dynamically loaded code object
hipError_t getDynFunc(hipFunction_t* hfunc, std::string func_name);
hipError_t getDeviceVar(DeviceVar** dvar, std::string var_name, int deviceId);
private:
amd::Program* program_;
//Maps for vars/funcs, could be keyed in with std::string name
std::unordered_map<std::string, Function*> functions_;
std::unordered_map<std::string, Var*> vars_;
//Load Code Object Data(Vars/UndefinedVars/Funcs)
hipError_t loadCodeObjectData(const void* mmap_ptr, size_t mmap_size);
//Populate Global Vars/Funcs from an code object(@ module_load)
hipError_t populateDynGlobalFuncs();
hipError_t populateDynGlobalVars();
};
//Static Code Object
class StatCO: public CodeObject {
amd::Monitor sclock_{"Guards Static Code object", true};
public:
StatCO();
virtual ~StatCO();
//Add/Remove/Digest Fat Binaries passed to us from "__hipRegisterFatBinary"
FatBinaryInfoType* addFatBinary(const void* data, bool initialized);
hipError_t removeFatBinary(FatBinaryInfoType* module);
hipError_t digestFatBinary(const void* data, FatBinaryInfoType& programs);
//Register vars/funcs given to use from __hipRegister[Var/Func]
hipError_t registerStatFunction(const void* hostFunction, Function* func);
hipError_t registerStatGlobalVar(const void* hostVar, Var* var);
//Retrive Vars/Funcs for a given hostSidePtr(const void*), unless stated otherwise.
hipError_t getStatFunc(hipFunction_t* hfunc, const void* hostFunction, int deviceId);
hipError_t getStatFuncAttr(hipFuncAttributes* func_attr, const void* hostFunction, int deviceId);
hipError_t getStatGlobalVar(const void* hostVar, int deviceId, hipDeviceptr_t* dev_ptr,
size_t* size_ptr);
hipError_t getStatGlobalVarByName(std::string hostVar, int deviceId, hipModule_t hmod,
hipDeviceptr_t* dev_ptr, size_t* size_ptr);
private:
friend class ::PlatformState;
//Populated during __hipRegisterFatBinary
std::unordered_map<const void*, FatBinaryInfoType> modules_;
//Populated during __hipRegisterFuncs
std::unordered_map<const void*, Function*> functions_;
//Populated during __hipRegisterVars
std::unordered_map<const void*, Var*> vars_;
};
}; //namespace: hip
#endif /* HIP_CODE_OBJECT_HPP */
rocclr/hip_context.cpp
+1
Ver Arquivo
@@ -20,6 +20,7 @@
#include <hip/hip_runtime.h>
#include "hip_internal.hpp"
#include "hip_platform.hpp"
#include "platform/runtime.hpp"
#include "utils/flags.hpp"
#include "utils/versions.hpp"
rocclr/hip_device.cpp
+1
Ver Arquivo
@@ -209,6 +209,7 @@ hipError_t hipGetDeviceProperties ( hipDeviceProp_t* props, hipDevice_t device )
//deviceProps.isMultiGpuBoard = info.;
deviceProps.canMapHostMemory = 1;
deviceProps.gcnArch = info.gfxipVersion_;
sprintf(deviceProps.gcnArchName, "gfx%d%d%x", info.gfxipMajor_, info.gfxipMinor_, info.gfxipStepping_);
deviceProps.cooperativeLaunch = info.cooperativeGroups_;
deviceProps.cooperativeMultiDeviceLaunch = info.cooperativeMultiDeviceGroups_;
rocclr/hip_device_runtime.cpp
+5 -1
Ver Arquivo
@@ -449,7 +449,11 @@ hipError_t hipDeviceSynchronize ( void ) {
}
int ihipGetDevice() {
return hip::getCurrentDevice()->deviceId();
hip::Device* device = hip::getCurrentDevice();
if(device == nullptr){
return -1;
}
return device->deviceId();
}
hipError_t hipGetDevice ( int* deviceId ) {
rocclr/hip_event.cpp
+3 -3
Ver Arquivo
@@ -192,13 +192,13 @@ hipError_t ihipEventQuery(hipEvent_t event) {
hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) {
HIP_INIT_API(hipEventCreateWithFlags, event, flags);
HIP_RETURN(ihipEventCreateWithFlags(event, flags));
HIP_RETURN(ihipEventCreateWithFlags(event, flags), *event);
}
hipError_t hipEventCreate(hipEvent_t* event) {
HIP_INIT_API(hipEventCreate, event);
HIP_RETURN(ihipEventCreateWithFlags(event, 0));
HIP_RETURN(ihipEventCreateWithFlags(event, 0), *event);
}
hipError_t hipEventDestroy(hipEvent_t event) {
@@ -227,7 +227,7 @@ hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) {
hip::Event* eStart = reinterpret_cast<hip::Event*>(start);
hip::Event* eStop = reinterpret_cast<hip::Event*>(stop);
HIP_RETURN(eStart->elapsedTime(*eStop, *ms));
HIP_RETURN(eStart->elapsedTime(*eStop, *ms), "Elapsed Time = ", *ms);
}
hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
rocclr/hip_fatbin.hpp
Arquivo executável
+29
Ver Arquivo
@@ -0,0 +1,29 @@
#ifndef HIP_FAT_BINARY_HPP
#define HIP_FAT_BINARY_HPP
namespace hip {
class FatBinaryMetaInfo {
public:
FatBinaryMetaInfo(bool built, const void* binary_ptr, size_t binary_size):
built_(built), binary_ptr_(binary_ptr), binary_size_(binary_size) {}
~FatBinaryMetaInfo() {}
//Set once the mod has been built
void set_built() { built_ = true; }
//Accessor for private vars
bool built() const { return built_; }
const void* binary_ptr() const { return binary_ptr_; }
size_t binary_size() const { return binary_size_; }
private:
bool built_; //Set when mod is built. Used in Lazy Binary
const void* binary_ptr_; //Binary image ptr
size_t binary_size_; //Binary Size
};
typedef std::vector<std::pair<hipModule_t, FatBinaryMetaInfo*>> FatBinaryInfoType;
}; /* namespace hip */
#endif /* HIP_FAT_BINARY_HPP */
rocclr/hip_global.cpp
Arquivo executável
+202
Ver Arquivo
@@ -0,0 +1,202 @@
#include "hip_global.hpp"
#include "hip/hip_runtime.h"
#include "hip_internal.hpp"
#include "hip_code_object.hpp"
#include "platform/program.hpp"
namespace hip {
//Device Vars
DeviceVar::DeviceVar(std::string name, hipModule_t hmod) : shadowVptr(nullptr), name_(name),
amd_mem_obj_(nullptr), device_ptr_(nullptr),
size_(0) {
amd::Program* program = as_amd(reinterpret_cast<cl_program>(hmod));
device::Program* dev_program = program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]);
if (dev_program == nullptr) {
DevLogPrintfError("Cannot get Device Function for module: 0x%x \n", hmod);
guarantee(false);
}
if(!dev_program->createGlobalVarObj(&amd_mem_obj_, &device_ptr_, &size_, name.c_str())) {
DevLogPrintfError("Cannot create Global Var obj for symbol: %s \n", name);
guarantee(false);
}
if (amd_mem_obj_ == nullptr || device_ptr_ == nullptr) {
DevLogPrintfError("Cannot get memory for creating device Var: %s", name.c_str());
guarantee(false);
}
amd::MemObjMap::AddMemObj(device_ptr_, amd_mem_obj_);
}
DeviceVar::~DeviceVar() {
if (device_ptr_ != nullptr) {
amd::MemObjMap::RemoveMemObj(device_ptr_);
amd_mem_obj_->release();
}
if (shadowVptr != nullptr) {
textureReference* texRef = reinterpret_cast<textureReference*>(shadowVptr);
delete texRef;
shadowVptr = nullptr;
}
device_ptr_ = nullptr;
size_ = 0;
}
//Device Functions
DeviceFunc::DeviceFunc(std::string name, hipModule_t hmod) : dflock_("function lock"),
name_(name), kernel_(nullptr) {
amd::Program* program = as_amd(reinterpret_cast<cl_program>(hmod));
const amd::Symbol *symbol = program->findSymbol(name.c_str());
if (symbol == nullptr) {
DevLogPrintfError("Cannot find Symbol with name: %s \n", name);
guarantee(false);
}
kernel_ = new amd::Kernel(*program, *symbol, name);
if (kernel_ == nullptr) {
DevLogPrintfError("Cannot create kernel with name: %s \n", name);
guarantee(false);
}
}
DeviceFunc::~DeviceFunc() {
if (kernel_ != nullptr) {
kernel_->release();
}
}
//Abstract functions
Function::Function(std::string name, FatBinaryInfoType* modules)
: name_(name), modules_(modules) {
dFunc_.resize(g_devices.size());
}
Function::~Function() {
for (auto& elem : dFunc_) {
delete elem;
}
name_ = "";
modules_ = nullptr;
}
hipError_t Function::getDynFunc(hipFunction_t* hfunc, hipModule_t hmod) {
guarantee(dFunc_.size() == g_devices.size());
if (dFunc_[ihipGetDevice()] == nullptr) {
dFunc_[ihipGetDevice()] = new DeviceFunc(name_, hmod);
}
*hfunc = dFunc_[ihipGetDevice()]->asHipFunction();
return hipSuccess;
}
hipError_t Function::getStatFunc(hipFunction_t* hfunc, int deviceId) {
guarantee(modules_ != nullptr);
guarantee(deviceId >= 0);
guarantee(deviceId < modules_->size());
hipModule_t module = (*modules_)[deviceId].first;
FatBinaryMetaInfo* fb_meta = (*modules_)[deviceId].second;
if (!fb_meta->built()) {
IHIP_RETURN_ONFAIL(CodeObject::add_program(deviceId, module, fb_meta->binary_ptr(),
fb_meta->binary_size()));
IHIP_RETURN_ONFAIL(CodeObject::build_module(module, g_devices[deviceId]->devices()));
fb_meta->set_built();
}
if (dFunc_[deviceId] == nullptr) {
dFunc_[deviceId] = new DeviceFunc(name_, (*modules_)[deviceId].first);
}
*hfunc = dFunc_[deviceId]->asHipFunction();
return hipSuccess;
}
hipError_t Function::getStatFuncAttr(hipFuncAttributes* func_attr, int deviceId) {
guarantee(modules_ != nullptr);
guarantee(deviceId >= 0);
guarantee(deviceId < modules_->size());
hipModule_t module = (*modules_)[deviceId].first;
FatBinaryMetaInfo* fb_meta = (*modules_)[deviceId].second;
if (!fb_meta->built()) {
IHIP_RETURN_ONFAIL(CodeObject::add_program(deviceId, module, fb_meta->binary_ptr(),
fb_meta->binary_size()));
IHIP_RETURN_ONFAIL(CodeObject::build_module(module, g_devices[deviceId]->devices()));
fb_meta->set_built();
}
if (dFunc_[deviceId] == nullptr) {
dFunc_[deviceId] = new DeviceFunc(name_, (*modules_)[deviceId].first);
}
const std::vector<amd::Device*>& devices = amd::Device::getDevices(CL_DEVICE_TYPE_GPU, false);
amd::Kernel* kernel = dFunc_[deviceId]->kernel();
const device::Kernel::WorkGroupInfo* wginfo = kernel->getDeviceKernel(*devices[deviceId])->workGroupInfo();
func_attr->localSizeBytes = wginfo->privateMemSize_;
func_attr->sharedSizeBytes = wginfo->localMemSize_;
func_attr->maxDynamicSharedSizeBytes = wginfo->availableLDSSize_ - wginfo->localMemSize_;
func_attr->maxThreadsPerBlock = wginfo->size_;
func_attr->numRegs = wginfo->usedVGPRs_;
return hipSuccess;
}
//Abstract Vars
Var::Var(std::string name, DeviceVarKind dVarKind, size_t size, int type, int norm,
FatBinaryInfoType* modules) : name_(name), dVarKind_(dVarKind), size_(size),
type_(type), norm_(norm), modules_(modules) {
dVar_.resize(g_devices.size());
}
Var::~Var() {
for (auto& elem : dVar_) {
delete elem;
}
modules_ = nullptr;
}
hipError_t Var::getDeviceVar(DeviceVar** dvar, int deviceId, hipModule_t hmod) {
guarantee(deviceId >= 0);
guarantee(deviceId < g_devices.size());
guarantee(dVar_.size() == g_devices.size());
if (dVar_[deviceId] == nullptr) {
dVar_[deviceId] = new DeviceVar(name_, hmod);
}
*dvar = dVar_[deviceId];
return hipSuccess;
}
hipError_t Var::getStatDeviceVar(DeviceVar** dvar, int deviceId) {
guarantee(deviceId >= 0);
guarantee(deviceId < g_devices.size());
hipModule_t module = (*modules_)[deviceId].first;
FatBinaryMetaInfo* fb_meta = (*modules_)[deviceId].second;
if (!fb_meta->built()) {
IHIP_RETURN_ONFAIL(CodeObject::add_program(deviceId, module, fb_meta->binary_ptr(),
fb_meta->binary_size()));
IHIP_RETURN_ONFAIL(CodeObject::build_module(module, g_devices[deviceId]->devices()));
fb_meta->set_built();
}
if (dVar_[deviceId] == nullptr) {
dVar_[deviceId] = new DeviceVar(name_, (*modules_)[deviceId].first);
}
*dvar = dVar_[deviceId];
return hipSuccess;
}
}; //namespace: hip
rocclr/hip_global.hpp
Arquivo executável
+116
Ver Arquivo
@@ -0,0 +1,116 @@
#ifndef HIP_GLOBAL_HPP
#define HIP_GLOBAL_HPP
#include <vector>
#include <string>
#include "hip/hip_runtime_api.h"
#include "hip/hip_runtime.h"
#include "hip_internal.hpp"
#include "hip_fatbin.hpp"
#include "platform/program.hpp"
namespace hip {
//Forward Declaration
class CodeObject;
//Device Structures
class DeviceVar {
public:
DeviceVar(std::string name, hipModule_t hmod);
~DeviceVar();
//Accessors for device ptr and size, populated during constructor.
hipDeviceptr_t device_ptr() const { return device_ptr_; }
size_t size() const { return size_; }
std::string name() const { return name_; }
void* shadowVptr;
private:
std::string name_; //Name of the var
amd::Memory* amd_mem_obj_; //amd_mem_obj abstraction
hipDeviceptr_t device_ptr_; //Device Pointer
size_t size_; //Size of the var
};
class DeviceFunc {
public:
DeviceFunc(std::string name, hipModule_t hmod);
~DeviceFunc();
amd::Monitor dflock_;
//Converts DeviceFunc to hipFunction_t(used by app) and vice versa.
hipFunction_t asHipFunction() { return reinterpret_cast<hipFunction_t>(this); }
static DeviceFunc* asFunction(hipFunction_t f) { return reinterpret_cast<DeviceFunc*>(f); }
//Accessor for kernel_ and name_ populated during constructor.
std::string name() const { return name_; }
amd::Kernel* kernel() const { return kernel_; }
private:
std::string name_; //name of the func(not unique identifier)
amd::Kernel* kernel_; //Kernel ptr referencing to ROCclr Symbol
};
//Abstract Structures
class Function {
public:
Function(std::string name, FatBinaryInfoType* modules=nullptr);
~Function();
//Return DeviceFunc for this this dynamically loaded module
hipError_t getDynFunc(hipFunction_t* hfunc, hipModule_t hmod);
//Return Device Func & attr . Generate/build if not already done so.
hipError_t getStatFunc(hipFunction_t *hfunc, int deviceId);
hipError_t getStatFuncAttr(hipFuncAttributes* func_attr, int deviceId);
void resize_dFunc(size_t size) { dFunc_.resize(size); }
FatBinaryInfoType* moduleInfo() { return modules_; };
private:
std::vector<DeviceFunc*> dFunc_; //DeviceFuncObj per Device
std::string name_; //name of the func(not unique identifier)
FatBinaryInfoType* modules_; // static module where it is referenced
};
class Var {
public:
//Types of variable
enum DeviceVarKind {
DVK_Variable = 0,
DVK_Surface,
DVK_Texture
};
Var(std::string name, DeviceVarKind dVarKind, size_t size, int type, int norm,
FatBinaryInfoType* modules = nullptr);
~Var();
//Return DeviceVar for this dynamically loaded module
hipError_t getDeviceVar(DeviceVar** dvar, int deviceId, hipModule_t hmod);
//Return DeviceVar for module Generate/build if not already done so.
hipError_t getStatDeviceVar(DeviceVar** dvar, int deviceId);
void resize_dVar(size_t size) { dVar_.resize(size); }
//Accessor for device_ptrs.
std::string name() const { return name_; }
hipModule_t module(int deviceId) const { return (*modules_)[deviceId].first; }
hipDeviceptr_t device_ptr(int deviceId) const { return dVar_[deviceId]->device_ptr(); }
size_t device_size(int deviceId) const { return dVar_[deviceId]->size(); }
FatBinaryInfoType* moduleInfo() { return modules_; };
private:
std::vector<DeviceVar*> dVar_; // DeviceVarObj per Device
std::string name_; // Variable name (not unique identifier)
DeviceVarKind dVarKind_; // Variable kind
size_t size_; // Size of the variable
int type_; // Type(Textures/Surfaces only)
int norm_; // Type(Textures/Surfaces only)
FatBinaryInfoType* modules_; // static module where it is referenced
};
}; //namespace: hip
#endif /* HIP_GLOBAL_HPP */
rocclr/hip_hcc.def.in
+11
Ver Arquivo
@@ -51,6 +51,7 @@ hipExtGetLinkTypeAndHopCount
hipExtLaunchMultiKernelMultiDevice
hipExtMallocWithFlags
hipExtModuleLaunchKernel
hipExtLaunchKernel
hipFree
hipFreeArray
hipFuncSetCacheConfig
@@ -79,6 +80,7 @@ hipMallocManaged
hipArrayCreate
hipArray3DCreate
hipMallocArray
hipMemAdvise
hipMemAllocPitch
hipMallocPitch
hipMemcpy
@@ -110,7 +112,10 @@ hipMemGetAddressRange
hipGetSymbolAddress
hipGetSymbolSize
hipMemGetInfo
hipMemPrefetchAsync
hipMemPtrGetInfo
hipMemRangeGetAttribute
hipMemRangeGetAttributes
hipMemset
hipMemsetAsync
hipMemsetD8
@@ -153,6 +158,7 @@ hipGetDeviceFlags
hipSetDevice
hipSetDeviceFlags
hipStreamAddCallback
hipStreamAttachMemAsync
hipStreamCreate
hipStreamCreateWithFlags
hipStreamCreateWithPriority
@@ -254,3 +260,8 @@ hipTexObjectGetResourceViewDesc
hipTexObjectGetTextureDesc
hipExtStreamCreateWithCUMask
hipStreamGetPriority
hipMemcpy2DFromArray
hipMemcpy2DFromArrayAsync
hipMemcpyAtoH
hipMemcpyHtoA
hipMemcpyParam2DAsync
rocclr/hip_hcc.map.in
+12
Ver Arquivo
@@ -52,6 +52,7 @@ global:
hipExtLaunchMultiKernelMultiDevice;
hipExtMallocWithFlags;
hipExtModuleLaunchKernel;
hipExtLaunchKernel;
hipFree;
hipFreeArray;
hipFuncSetCacheConfig;
@@ -61,6 +62,7 @@ global:
hipGetErrorName;
hipGetErrorString;
hipGetLastError;
hipMemAdvise;
hipMemAllocHost;
hipHostAlloc;
hipHostFree;
@@ -111,7 +113,10 @@ global:
hipGetSymbolAddress;
hipGetSymbolSize;
hipMemGetInfo;
hipMemPrefetchAsync;
hipMemPtrGetInfo;
hipMemRangeGetAttribute;
hipMemRangeGetAttributes;
hipMemset;
hipMemsetAsync;
hipMemsetD8;
@@ -153,6 +158,7 @@ global:
hipSetDevice;
hipSetDeviceFlags;
hipStreamAddCallback;
hipStreamAttachMemAsync;
hipStreamCreate;
hipStreamCreateWithFlags;
hipStreamCreateWithPriority;
@@ -182,6 +188,7 @@ global:
hipApiName;
hipKernelNameRef;
hipKernelNameRefByPtr;
hipGetStreamDeviceId;
hipProfilerStart;
hipProfilerStop;
hiprtcCompileProgram;
@@ -261,6 +268,11 @@ global:
hipGetCmdName*;
hipExtStreamCreateWithCUMask;
hipStreamGetPriority;
hipMemcpy2DFromArray;
hipMemcpy2DFromArrayAsync;
hipMemcpyAtoH;
hipMemcpyHtoA;
hipMemcpyParam2DAsync;
};
local:
*;
rocclr/hip_hmm.cpp
+215
Ver Arquivo
@@ -0,0 +1,215 @@
/* Copyright (c) 2020-present Advanced Micro Devices, Inc.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE. */
#include <hip/hip_runtime.h>
#include "hip_internal.hpp"
#include "hip_conversions.hpp"
#include "platform/context.hpp"
#include "platform/command.hpp"
#include "platform/memory.hpp"
// Forward declaraiton of a static function
static hipError_t ihipMallocManaged(void** ptr, size_t size);
// Make sure HIP defines match ROCclr to avoid double conversion
static_assert(hipCpuDeviceId == amd::CpuDeviceId, "CPU device ID mismatch with ROCclr!");
static_assert(hipInvalidDeviceId == amd::InvalidDeviceId,
"Invalid device ID mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemAdviseSetReadMostly) ==
amd::MemoryAdvice::SetReadMostly, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemAdviseUnsetReadMostly) ==
amd::MemoryAdvice::UnsetReadMostly, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemAdviseSetPreferredLocation) ==
amd::MemoryAdvice::SetPreferredLocation, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemAdviseUnsetPreferredLocation) ==
amd::MemoryAdvice::UnsetPreferredLocation, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemAdviseSetAccessedBy) ==
amd::MemoryAdvice::SetAccessedBy, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemAdviseUnsetAccessedBy) ==
amd::MemoryAdvice::UnsetAccessedBy, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemRangeAttributeReadMostly) ==
amd::MemRangeAttribute::ReadMostly, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemRangeAttributePreferredLocation) ==
amd::MemRangeAttribute::PreferredLocation, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemRangeAttributeAccessedBy) ==
amd::MemRangeAttribute::AccessedBy, "Enum mismatch with ROCclr!");
static_assert(static_cast<uint32_t>(hipMemRangeAttributeLastPrefetchLocation) ==
amd::MemRangeAttribute::LastPrefetchLocation, "Enum mismatch with ROCclr!");
// ================================================================================================
hipError_t hipMallocManaged(void** dev_ptr, size_t size, unsigned int flags) {
HIP_INIT_API(hipMallocManaged, dev_ptr, size, flags);
if ((dev_ptr == nullptr) || (flags != hipMemAttachGlobal)) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(ihipMallocManaged(dev_ptr, size), *dev_ptr);
}
// ================================================================================================
hipError_t hipMemPrefetchAsync(const void* dev_ptr, size_t count, int device,
hipStream_t stream) {
HIP_INIT_API(hipMemPrefetchAsync, dev_ptr, count, device, stream);
if ((dev_ptr == nullptr) || (count == 0) || (stream == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::HostQueue* queue = nullptr;
bool cpu_access = (device == hipCpuDeviceId) ? true : false;
// Pick the specified stream or Null one from the provided device
if (stream != nullptr) {
queue = hip::getQueue(stream);
} else {
if (!cpu_access) {
queue = g_devices[device]->NullStream();
} else {
queue = hip::getCurrentDevice()->NullStream();
}
}
if (queue == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Command::EventWaitList waitList;
amd::SvmPrefetchAsyncCommand* command =
new amd::SvmPrefetchAsyncCommand(*queue, waitList, dev_ptr, count, cpu_access);
if (command == nullptr) {
return hipErrorOutOfMemory;
}
if (!command->validateMemory()) {
delete command;
HIP_RETURN(hipErrorInvalidValue);
}
command->enqueue();
command->release();
HIP_RETURN(hipErrorInvalidValue);
}
// ================================================================================================
hipError_t hipMemAdvise(const void* dev_ptr, size_t count, hipMemoryAdvise advice, int device) {
HIP_INIT_API(hipMemAdvise, dev_ptr, count, advice, device);
if ((dev_ptr == nullptr) || (count == 0) || (device >= g_devices.size())) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* dev = g_devices[device]->devices()[0];
// Set the allocation attributes in AMD HMM
if (!dev->SetSvmAttributes(dev_ptr, count, static_cast<amd::MemoryAdvice>(advice))) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(hipSuccess);
}
// ================================================================================================
hipError_t hipMemRangeGetAttribute(void* data, size_t data_size, hipMemRangeAttribute attribute,
const void* dev_ptr, size_t count) {
HIP_INIT_API(hipMemRangeGetAttribute, data, data_size, attribute, dev_ptr, count);
if ((data == nullptr) || (data_size == 0) || (dev_ptr == nullptr) || (count == 0)) {
HIP_RETURN(hipErrorInvalidValue);
}
// Shouldn't matter for which device the interface is called
amd::Device* dev = g_devices[0]->devices()[0];
// Get the allocation attribute from AMD HMM
if (!dev->GetSvmAttributes(&data, &data_size, reinterpret_cast<int*>(&attribute), 1,
dev_ptr, count)) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(hipErrorInvalidValue);
}
// ================================================================================================
hipError_t hipMemRangeGetAttributes(void** data, size_t* data_sizes,
hipMemRangeAttribute* attributes, size_t num_attributes,
const void* dev_ptr, size_t count) {
HIP_INIT_API(hipMemRangeGetAttributes, data, data_sizes,
attributes, num_attributes, dev_ptr, count);
if ((data == nullptr) || (data_sizes == nullptr) || (attributes == nullptr) ||
(num_attributes == 0) || (dev_ptr == nullptr) || (count == 0)) {
HIP_RETURN(hipErrorInvalidValue);
}
// Shouldn't matter for which device the interface is called
amd::Device* dev = g_devices[0]->devices()[0];
// Get the allocation attributes from AMD HMM
if (!dev->GetSvmAttributes(data, data_sizes, reinterpret_cast<int*>(attributes),
num_attributes, dev_ptr, count)) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(hipErrorInvalidValue);
}
// ================================================================================================
hipError_t hipStreamAttachMemAsync(hipStream_t stream, hipDeviceptr_t* dev_ptr,
size_t length, unsigned int flags) {
HIP_INIT_API(hipStreamAttachMemAsync, stream, dev_ptr, length, flags);
if ((stream == nullptr) || (dev_ptr == nullptr) || (length == 0)) {
HIP_RETURN(hipErrorInvalidValue);
}
// Unclear what should be done for this interface in AMD HMM, since it's generic SVM alloc
HIP_RETURN(hipErrorInvalidValue);
}
// ================================================================================================
static hipError_t ihipMallocManaged(void** ptr, size_t size) {
if (size == 0) {
*ptr = nullptr;
return hipSuccess;
} else if (ptr == nullptr) {
return hipErrorInvalidValue;
}
assert((hip::host_device->asContext()!= nullptr) && "Current host context must be valid");
amd::Context& ctx = *hip::host_device->asContext();
const amd::Device& dev = *ctx.devices()[0];
// For now limit to the max allocation size on the device.
// The apps should be able to go over theCould you limit allocation in the future
if (dev.info().maxMemAllocSize_ < size) {
return hipErrorMemoryAllocation;
}
// Allocate SVM fine grain buffer with the forced host pointer, avoiding explicit memory
// allocation in the device driver
*ptr = amd::SvmBuffer::malloc(ctx, CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_ALLOC_HOST_PTR,
size, dev.info().memBaseAddrAlign_);
if (*ptr == nullptr) {
return hipErrorMemoryAllocation;
}
ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] ihipMallocManaged ptr=0x%zx", getpid(),
std::this_thread::get_id(), *ptr);
return hipSuccess;
}
rocclr/hip_intercept.cpp
Arquivo normal → Arquivo executável
+15 -4
Ver Arquivo
@@ -20,6 +20,7 @@
#include "hip/hip_runtime.h"
#include "hip_internal.hpp"
#include "hip_platform.hpp"
#include "hip_prof_api.h"
// HIP API callback/activity
@@ -27,16 +28,26 @@
api_callbacks_table_t callbacks_table;
extern const std::string& FunctionName(const hipFunction_t f);
const char* hipKernelNameRef(const hipFunction_t f) { return FunctionName(f).c_str(); }
const char* hipKernelNameRefByPtr(const void *hostFunction, hipStream_t stream) {
int hipGetStreamDeviceId(hipStream_t stream) {
hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
int deviceId = (s != nullptr)? s->DeviceId() : ihipGetDevice();
return (s != nullptr)? s->DeviceId() : ihipGetDevice();
}
const char* hipKernelNameRefByPtr(const void* hostFunction, hipStream_t stream) {
if (hostFunction == NULL) {
return NULL;
}
int deviceId = hipGetStreamDeviceId(stream);
if (deviceId == -1) {
DevLogPrintfError("Wrong Device Id: %d \n", deviceId);
return NULL;
}
hipFunction_t func = PlatformState::instance().getFunc(hostFunction, deviceId);
if (func == nullptr) {
hipFunction_t func = nullptr;
hipError_t hip_error = PlatformState::instance().getStatFunc(&func, hostFunction, deviceId);
if (hip_error != hipSuccess) {
return NULL;
}
return hipKernelNameRef(func);
rocclr/hip_internal.hpp
+40 -144
Ver Arquivo
@@ -57,9 +57,17 @@ typedef struct ihipIpcMemHandle_st {
hip::g_device = g_devices[0]; \
}
#define HIP_API_PRINT(...) \
ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] %s ( %s )", getpid(), std::this_thread::get_id(), \
__func__, ToString( __VA_ARGS__ ).c_str());
#define HIP_ERROR_PRINT(err, ...) \
ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] %s: Returned %s : %s", getpid(), std::this_thread::get_id(), \
__func__, hipGetErrorName(err), ToString( __VA_ARGS__ ).c_str());
// This macro should be called at the beginning of every HIP API.
#define HIP_INIT_API(cid, ...) \
ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] %s ( %s )", getpid(), std::this_thread::get_id(), __func__, ToString( __VA_ARGS__ ).c_str()); \
HIP_API_PRINT(__VA_ARGS__) \
amd::Thread* thread = amd::Thread::current(); \
if (!VDI_CHECK_THREAD(thread)) { \
HIP_RETURN(hipErrorOutOfMemory); \
@@ -67,11 +75,30 @@ typedef struct ihipIpcMemHandle_st {
HIP_INIT() \
HIP_CB_SPAWNER_OBJECT(cid);
#define HIP_RETURN(ret) \
hip::g_lastError = ret; \
ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] %s: Returned %s", getpid(), std::this_thread::get_id(), __func__, hipGetErrorName(hip::g_lastError)); \
#define HIP_RETURN(ret, ...) \
hip::g_lastError = ret; \
HIP_ERROR_PRINT(hip::g_lastError, __VA_ARGS__) \
return hip::g_lastError;
#define HIP_RETURN_ONFAIL(func) \
do { \
hipError_t herror = (func); \
if (herror != hipSuccess) { \
HIP_RETURN(herror); \
} \
} while (0);
// Cannot be use in place of HIP_RETURN.
// Refrain from using for external HIP APIs
#define IHIP_RETURN_ONFAIL(func) \
do { \
hipError_t herror = (func); \
if (herror != hipSuccess) { \
return herror; \
} \
} while (0);
namespace hc {
class accelerator;
class accelerator_view;
@@ -81,16 +108,19 @@ namespace hip {
class Device;
class Stream {
public:
enum Priority : int {High = -1, Normal = 0, Low = 1};
private:
amd::HostQueue* queue_;
mutable amd::Monitor lock_;
Device* device_;
amd::CommandQueue::Priority priority_;
Priority priority_;
unsigned int flags_;
bool null_;
const std::vector<uint32_t> cuMask_;
public:
Stream(Device* dev, amd::CommandQueue::Priority p, unsigned int f = 0, bool null_stream = false,
Stream(Device* dev, Priority p = Priority::Normal, unsigned int f = 0, bool null_stream = false,
const std::vector<uint32_t>& cuMask = {});
/// Creates the hip stream object, including AMD host queue
@@ -110,7 +140,7 @@ namespace hip {
/// Returns the creation flags for the current stream
unsigned int Flags() const { return flags_; }
/// Returns the priority for the current stream
amd::CommandQueue::Priority Priority() const { return priority_; }
Priority GetPriority() const { return priority_; }
/// Sync all non-blocking streams
static void syncNonBlockingStreams();
@@ -133,7 +163,7 @@ namespace hip {
public:
Device(amd::Context* ctx, int devId):
context_(ctx), deviceId_(devId), null_stream_(this, amd::CommandQueue::Priority::Normal, 0, true), flags_(hipDeviceScheduleSpin)
context_(ctx), deviceId_(devId), null_stream_(this, Stream::Priority::Normal, 0, true), flags_(hipDeviceScheduleSpin)
{ assert(ctx != nullptr); }
~Device() {}
@@ -187,17 +217,6 @@ namespace hip {
extern amd::HostQueue* getNullStream(amd::Context&);
/// Get default stream of the thread
extern amd::HostQueue* getNullStream();
struct Function {
amd::Kernel* function_;
amd::Monitor lock_;
Function(amd::Kernel* f) : function_(f), lock_("function lock") {}
~Function() { function_->release(); }
hipFunction_t asHipFunction() { return reinterpret_cast<hipFunction_t>(this); }
static Function* asFunction(hipFunction_t f) { return reinterpret_cast<Function*>(f); }
};
};
struct ihipExec_t {
@@ -208,131 +227,6 @@ struct ihipExec_t {
std::vector<char> arguments_;
};
class PlatformState {
amd::Monitor lock_{"Guards global function map", true};
std::unordered_map<const void*, std::vector<std::pair<hipModule_t, bool>>> modules_;
bool initialized_{false};
void digestFatBinary(const void* data, std::vector<std::pair<hipModule_t, bool>>& programs);
public:
void init();
std::vector<std::pair<hipModule_t, bool>>* addFatBinary(const void*data)
{
amd::ScopedLock lock(lock_);
if (initialized_) {
digestFatBinary(data, modules_[data]);
}
return &modules_[data];
}
void removeFatBinary(std::vector<std::pair<hipModule_t, bool>>* module)
{
amd::ScopedLock lock(lock_);
for (auto& mod : modules_) {
if (&mod.second == module) {
modules_.erase(&mod);
return;
}
}
}
struct RegisteredVar {
public:
RegisteredVar(): size_(0), devicePtr_(nullptr), amd_mem_obj_(nullptr) {}
~RegisteredVar() {}
hipDeviceptr_t getdeviceptr() const { return devicePtr_; };
size_t getvarsize() const { return size_; };
size_t size_; // Size of the variable
hipDeviceptr_t devicePtr_; //Device Memory Address of the variable.
amd::Memory* amd_mem_obj_;
};
struct DeviceFunction {
std::string deviceName;
std::vector< std::pair< hipModule_t, bool > >* modules;
std::vector<hipFunction_t> functions;
};
enum DeviceVarKind {
DVK_Variable,
DVK_Surface,
DVK_Texture
};
struct DeviceVar {
DeviceVarKind kind;
void* shadowVptr;
std::string hostVar;
size_t size;
std::vector< std::pair< hipModule_t, bool > >* modules;
std::vector<RegisteredVar> rvars;
bool dyn_undef;
int type; // surface/texture type
int norm; // texture has normalized output
bool shadowAllocated = false; // shadow ptr is allocated on-demand and needs freeing.
};
private:
class Module {
public:
Module(hipModule_t hip_module_) : hip_module(hip_module_) {}
std::unordered_map<std::string, DeviceFunction > functions_;
private:
hipModule_t hip_module;
};
std::unordered_map<hipModule_t, Module*> module_map_;
std::unordered_map<const void*, DeviceFunction > functions_;
std::unordered_multimap<std::string, DeviceVar > vars_;
// Map from the host shadow symbol to its device name. As different modules
// may have the same name, each symbol is uniquely identified by a pair of
// module handle and its name.
std::unordered_map<const void*,
std::pair<hipModule_t, std::string>> symbols_;
static PlatformState* platform_;
PlatformState() {}
~PlatformState() {}
public:
static PlatformState& instance() {
if (platform_ == nullptr) {
// __hipRegisterFatBinary() will call this when app starts, thus
// there is no multiple entry issue here.
platform_ = new PlatformState();
}
return *platform_;
}
bool unregisterFunc(hipModule_t hmod);
std::vector< std::pair<hipModule_t, bool> >* unregisterVar(hipModule_t hmod);
bool findSymbol(const void *hostVar, hipModule_t &hmod, std::string &devName);
PlatformState::DeviceVar* findVar(std::string hostVar, int deviceId, hipModule_t hmod);
void registerVarSym(const void *hostVar, hipModule_t hmod, const char *symbolName);
void registerVar(const char* symbolName, const DeviceVar& var);
void registerFunction(const void* hostFunction, const DeviceFunction& func);
bool registerModFuncs(std::vector<std::string>& func_names, hipModule_t* module);
bool findModFunc(hipFunction_t* hfunc, hipModule_t hmod, const char* name);
bool createFunc(hipFunction_t* hfunc, hipModule_t hmod, const char* name);
hipFunction_t getFunc(const void* hostFunction, int deviceId);
bool getFuncAttr(const void* hostFunction, hipFuncAttributes* func_attr);
bool getGlobalVar(const char* hostVar, int deviceId, hipModule_t hmod,
hipDeviceptr_t* dev_ptr, size_t* size_ptr);
bool getTexRef(const char* hostVar, hipModule_t hmod, textureReference** texRef);
bool getGlobalVarFromSymbol(const void* hostVar, int deviceId,
hipDeviceptr_t* dev_ptr, size_t* size_ptr);
bool getShadowVarInfo(std::string var_name, hipModule_t hmod,
void** var_addr, size_t* var_size);
void setupArgument(const void *arg, size_t size, size_t offset);
void configureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, hipStream_t stream);
void popExec(ihipExec_t& exec);
};
/// Wait all active streams on the blocking queue. The method enqueues a wait command and
/// doesn't stall the current thread
extern void iHipWaitActiveStreams(amd::HostQueue* blocking_queue, bool wait_null_stream = false);
@@ -345,4 +239,6 @@ extern amd::Memory* getMemoryObject(const void* ptr, size_t& offset);
extern bool CL_CALLBACK getSvarInfo(cl_program program, std::string var_name, void** var_addr,
size_t* var_size);
constexpr bool kOptionChangeable = true;
constexpr bool kNewDevProg = false;
#endif // HIP_SRC_HIP_INTERNAL_H
rocclr/hip_memory.cpp
+42 -97
Ver Arquivo
@@ -20,6 +20,7 @@
#include <hip/hip_runtime.h>
#include "hip_internal.hpp"
#include "hip_platform.hpp"
#include "hip_conversions.hpp"
#include "platform/context.hpp"
#include "platform/command.hpp"
@@ -85,8 +86,9 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
return hipErrorInvalidValue;
}
amd::Context* amdContext = ((flags & CL_MEM_SVM_FINE_GRAIN_BUFFER) != 0)?
hip::host_device->asContext() : hip::getCurrentDevice()->asContext();
bool useHostDevice = (flags & CL_MEM_SVM_FINE_GRAIN_BUFFER) != 0;
amd::Context* curDevContext = hip::getCurrentDevice()->asContext();
amd::Context* amdContext = useHostDevice ? hip::host_device->asContext() : curDevContext;
if (amdContext == nullptr) {
return hipErrorOutOfMemory;
@@ -96,14 +98,16 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
return hipErrorOutOfMemory;
}
*ptr = amd::SvmBuffer::malloc(*amdContext, flags, sizeBytes, amdContext->devices()[0]->info().memBaseAddrAlign_);
*ptr = amd::SvmBuffer::malloc(*amdContext, flags, sizeBytes, amdContext->devices()[0]->info().memBaseAddrAlign_,
useHostDevice ? curDevContext->svmDevices()[0] : nullptr);
if (*ptr == nullptr) {
return hipErrorOutOfMemory;
}
ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] ihipMalloc ptr=0x%zx", getpid(),std::this_thread::get_id(), *ptr);
return hipSuccess;
}
// ================================================================================================
hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
amd::HostQueue& queue, bool isAsync = false) {
if (sizeBytes == 0) {
@@ -153,36 +157,26 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
*srcMemory->asBuffer(), sOffset, sizeBytes, dst);
isAsync = false;
} else if ((srcMemory != nullptr) && (dstMemory != nullptr)) {
if (queueDevice != srcMemory->getContext().devices()[0]) {
amd::Coord3D srcOffset(sOffset, 0, 0);
amd::Coord3D dstOffset(dOffset, 0, 0);
amd::Coord3D copySize(sizeBytes, 1, 1);
if ((kind == hipMemcpyDeviceToDevice) &&
// Check if the queue device doesn't match the device on any memory object. Hence
// it's a P2P transfer, because the app has requested access to another GPU
(srcMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0])) {
command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize);
command->enqueue();
if (!isAsync) {
command->awaitCompletion();
*srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);
if (command == nullptr) {
return hipErrorOutOfMemory;
}
command->release();
return hipSuccess;
}
if (queueDevice != dstMemory->getContext().devices()[0]) {
amd::Coord3D srcOffset(sOffset, 0, 0);
amd::Coord3D dstOffset(dOffset, 0, 0);
amd::Coord3D copySize(sizeBytes, 1, 1);
command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize);
command->enqueue();
if (!isAsync) {
command->awaitCompletion();
// Make sure runtime has valid memory for the command execution. P2P access
// requires page table mapping on the current device to another GPU memory
if (!static_cast<amd::CopyMemoryP2PCommand*>(command)->validateMemory()) {
delete command;
return hipErrorInvalidValue;
}
command->release();
return hipSuccess;
} else {
command = new amd::CopyMemoryCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);
}
command = new amd::CopyMemoryCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
*srcMemory->asBuffer(),*dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);
}
if (command == nullptr) {
return hipErrorOutOfMemory;
}
@@ -200,6 +194,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
return hipSuccess;
}
// ================================================================================================
hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags) {
HIP_INIT_API(hipExtMallocWithFlags, ptr, sizeBytes, flags);
@@ -207,13 +202,13 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(ihipMalloc(ptr, sizeBytes, (flags & hipDeviceMallocFinegrained)? CL_MEM_SVM_ATOMICS: 0));
HIP_RETURN(ihipMalloc(ptr, sizeBytes, (flags & hipDeviceMallocFinegrained)? CL_MEM_SVM_ATOMICS: 0), *ptr);
}
hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
HIP_INIT_API(hipMalloc, ptr, sizeBytes);
HIP_RETURN(ihipMalloc(ptr, sizeBytes, 0));
HIP_RETURN(ihipMalloc(ptr, sizeBytes, 0), *ptr);
}
hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
@@ -241,18 +236,7 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
ihipFlags |= CL_MEM_SVM_ATOMICS;
}
HIP_RETURN(ihipMalloc(ptr, sizeBytes, ihipFlags));
}
hipError_t hipMallocManaged(void** devPtr, size_t size,
unsigned int flags) {
HIP_INIT_API(hipMallocManaged, devPtr, size, flags);
if (flags != hipMemAttachGlobal) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(ihipMalloc(devPtr, size, CL_MEM_SVM_FINE_GRAIN_BUFFER));
HIP_RETURN(ihipMalloc(ptr, sizeBytes, ihipFlags), *ptr);
}
hipError_t hipFree(void* ptr) {
@@ -399,7 +383,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height
HIP_INIT_API(hipMallocPitch, ptr, pitch, width, height);
const cl_image_format image_format = { CL_R, CL_UNSIGNED_INT8 };
HIP_RETURN(ihipMallocPitch(ptr, pitch, width, height, 1, CL_MEM_OBJECT_IMAGE2D, &image_format));
HIP_RETURN(ihipMallocPitch(ptr, pitch, width, height, 1, CL_MEM_OBJECT_IMAGE2D, &image_format), *ptr);
}
hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) {
@@ -422,7 +406,7 @@ hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) {
pitchedDevPtr->ysize = extent.height;
}
HIP_RETURN(status);
HIP_RETURN(status, *pitchedDevPtr);
}
amd::Image* ihipImageCreate(const cl_channel_order channelOrder,
@@ -698,7 +682,7 @@ hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags)
amd::MemObjMap::AddMemObj(hostPtr, mem);
HIP_RETURN(hipSuccess);
} else {
HIP_RETURN(ihipMalloc(&hostPtr, sizeBytes, flags));
HIP_RETURN(ihipMalloc(&hostPtr, sizeBytes, flags), hostPtr);
}
}
@@ -733,7 +717,7 @@ hipError_t hipHostUnregister(void* hostPtr) {
// Deprecated function:
hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) {
HIP_RETURN(ihipMalloc(ptr, sizeBytes, flags));
HIP_RETURN(ihipMalloc(ptr, sizeBytes, flags), *ptr);
};
@@ -744,18 +728,7 @@ hipError_t hipMemcpyToSymbol(const void* symbol, const void* src, size_t sizeByt
size_t sym_size = 0;
hipDeviceptr_t device_ptr = nullptr;
hipModule_t hmod;
std::string symbolName;
if (!PlatformState::instance().findSymbol(symbol, hmod, symbolName)) {
DevLogPrintfError("cannot find symbol 0x%x \n", symbolName.c_str());
HIP_RETURN(hipErrorInvalidSymbol);
}
/* Get address and size for the global symbol */
if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), hmod,
&device_ptr, &sym_size)) {
DevLogPrintfError("Cannot get global var: %s at device: %d \n", symbolName.c_str(), ihipGetDevice());
HIP_RETURN(hipErrorInvalidSymbol);
}
HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(symbol, ihipGetDevice(), &device_ptr, &sym_size));
/* Size Check to make sure offset is correct */
if ((offset + sizeBytes) > sym_size) {
@@ -777,18 +750,7 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* symbol, size_t sizeBytes,
size_t sym_size = 0;
hipDeviceptr_t device_ptr = nullptr;
hipModule_t hmod;
std::string symbolName;
if (!PlatformState::instance().findSymbol(symbol, hmod, symbolName)) {
DevLogPrintfError("cannot find symbol: 0x%x \n", symbol);
HIP_RETURN(hipErrorInvalidSymbol);
}
/* Get address and size for the global symbol */
if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), hmod,
&device_ptr, &sym_size)) {
DevLogPrintfError("Cannot find symbol Name: %s \n", symbolName.c_str());
HIP_RETURN(hipErrorInvalidSymbol);
}
HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(symbol, ihipGetDevice(), &device_ptr, &sym_size));
/* Size Check to make sure offset is correct */
if ((offset + sizeBytes) > sym_size) {
@@ -810,18 +772,7 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t si
size_t sym_size = 0;
hipDeviceptr_t device_ptr = nullptr;
hipModule_t hmod;
std::string symbolName;
if (!PlatformState::instance().findSymbol(symbol, hmod, symbolName)) {
DevLogPrintfError("cannot find symbol: 0x%x \n", symbol);
HIP_RETURN(hipErrorInvalidSymbol);
}
/* Get address and size for the global symbol */
if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), hmod,
&device_ptr, &sym_size)) {
DevLogPrintfError("Cannot find symbol Name: %s \n", symbolName.c_str());
HIP_RETURN(hipErrorInvalidSymbol);
}
HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(symbol, ihipGetDevice(), &device_ptr, &sym_size));
/* Size Check to make sure offset is correct */
if ((offset + sizeBytes) > sym_size) {
@@ -843,18 +794,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbol, size_t sizeBy
size_t sym_size = 0;
hipDeviceptr_t device_ptr = nullptr;
hipModule_t hmod;
std::string symbolName;
if (!PlatformState::instance().findSymbol(symbol, hmod, symbolName)) {
DevLogPrintfError("cannot find symbol: 0x%x \n", symbol);
HIP_RETURN(hipErrorInvalidSymbol);
}
/* Get address and size for the global symbol */
if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), hmod,
&device_ptr, &sym_size)) {
DevLogPrintfError("Cannot find symbol Name: %s \n", symbolName.c_str());
HIP_RETURN(hipErrorInvalidSymbol);
}
HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(symbol, ihipGetDevice(), &device_ptr, &sym_size));
/* Size Check to make sure offset is correct */
if ((offset + sizeBytes) > sym_size) {
@@ -2026,7 +1966,12 @@ hipError_t hipIpcCloseMemHandle(void* dev_ptr) {
}
/* Remove the memory from MemObjMap */
amd::MemObjMap::RemoveMemObj(amd_mem_obj);
if (amd_mem_obj->getSvmPtr() != nullptr) {
amd::MemObjMap::RemoveMemObj(amd_mem_obj->getSvmPtr());
} else {
DevLogPrintfError("Does not have SVM or Host Mem for 0x%x, crash here!", dev_ptr);
guarantee(false);
}
/* detach the memory */
if (!device->IpcDetach(*amd_mem_obj)){
@@ -2303,7 +2248,7 @@ hipError_t hipMallocHost(void** ptr,
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(ihipMalloc(ptr, size, CL_MEM_SVM_FINE_GRAIN_BUFFER));
HIP_RETURN(ihipMalloc(ptr, size, CL_MEM_SVM_FINE_GRAIN_BUFFER), *ptr);
}
hipError_t hipFreeHost(void *ptr) {
rocclr/hip_module.cpp
+63 -206
Ver Arquivo
@@ -29,9 +29,18 @@
hipError_t ihipModuleLoadData(hipModule_t* module, const void* mmap_ptr, size_t mmap_size);
const std::string& FunctionName(const hipFunction_t f)
{
return hip::Function::asFunction(f)->function_->name();
extern hipError_t ihipLaunchKernel(const void* hostFunction,
dim3 gridDim,
dim3 blockDim,
void** args,
size_t sharedMemBytes,
hipStream_t stream,
hipEvent_t startEvent,
hipEvent_t stopEvent,
int flags);
const std::string& FunctionName(const hipFunction_t f) {
return hip::DeviceFunc::asFunction(f)->kernel()->name();
}
static uint64_t ElfSize(const void *emi)
@@ -55,222 +64,48 @@ static uint64_t ElfSize(const void *emi)
return total_size;
}
hipError_t hipModuleLoad(hipModule_t* module, const char* fname)
{
HIP_INIT_API(hipModuleLoad, module, fname);
const void* mmap_ptr = nullptr;
size_t mmap_size = 0;
if (!fname) {
HIP_RETURN(hipErrorInvalidValue);
}
if (!amd::Os::MemoryMapFile(fname, &mmap_ptr, &mmap_size)) {
HIP_RETURN(hipErrorFileNotFound);
}
HIP_RETURN(ihipModuleLoadData(module, mmap_ptr, mmap_size));
}
bool ihipModuleUnregisterGlobal(hipModule_t hmod) {
std::vector< std::pair<hipModule_t, bool> >* modules =
PlatformState::instance().unregisterVar(hmod);
if (modules != nullptr) {
delete modules;
}
return true;
}
hipError_t hipModuleUnload(hipModule_t hmod)
{
hipError_t hipModuleUnload(hipModule_t hmod) {
HIP_INIT_API(hipModuleUnload, hmod);
if (hmod == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(PlatformState::instance().unloadModule(hmod));
}
amd::Program* program = as_amd(reinterpret_cast<cl_program>(hmod));
hipError_t hipModuleLoad(hipModule_t* module, const char* fname) {
HIP_INIT_API(hipModuleLoad, module, fname);
if(!PlatformState::instance().unregisterFunc(hmod)) {
DevLogPrintfError("Cannot unregister module: 0x%x \n", hmod);
HIP_RETURN(hipErrorInvalidSymbol);
}
if(!ihipModuleUnregisterGlobal(hmod)) {
DevLogPrintfError("Cannot unregister Global vars for module: 0x%x \n", hmod);
HIP_RETURN(hipErrorInvalidSymbol);
}
program->release();
HIP_RETURN(hipSuccess);
HIP_RETURN(PlatformState::instance().loadModule(module, fname));
}
hipError_t hipModuleLoadData(hipModule_t *module, const void *image)
{
HIP_INIT_API(hipModuleLoadData, module, image);
HIP_RETURN(ihipModuleLoadData(module, image, 0));
HIP_RETURN(PlatformState::instance().loadModule(module, 0, image));
}
hipError_t hipModuleLoadDataEx(hipModule_t *module, const void *image,
unsigned int numOptions, hipJitOption* options,
void** optionsValues)
unsigned int numOptions, hipJitOption* options,
void** optionsValues)
{
/* TODO: Pass options to Program */
HIP_INIT_API(hipModuleLoadDataEx, module, image);
HIP_RETURN(ihipModuleLoadData(module, image, 0));
HIP_RETURN(PlatformState::instance().loadModule(module, 0, image));
}
extern hipError_t __hipExtractCodeObjectFromFatBinary(const void* data,
const std::vector<const char*>& devices,
std::vector<std::pair<const void*, size_t>>& code_objs);
inline bool ihipModuleRegisterUndefined(amd::Program* program, hipModule_t* module) {
std::vector<std::string> undef_vars;
device::Program* dev_program
= program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]);
if (!dev_program->getUndefinedVarFromCodeObj(&undef_vars)) {
DevLogPrintfError("Could not get undefined Variables for Module: 0x%x \n", *module);
return false;
}
for (auto it = undef_vars.begin(); it != undef_vars.end(); ++it) {
auto modules = new std::vector<std::pair<hipModule_t, bool> >(g_devices.size());
for (size_t dev = 0; dev < g_devices.size(); ++dev) {
modules->at(dev) = std::make_pair(*module, true);
}
texture<float, hipTextureType1D, hipReadModeElementType>* tex_hptr
= new texture<float, hipTextureType1D, hipReadModeElementType>();
memset(tex_hptr, 0x00, sizeof(texture<float, hipTextureType1D, hipReadModeElementType>));
PlatformState::DeviceVar dvar{PlatformState::DVK_Variable,
reinterpret_cast<char*>(tex_hptr),
it->c_str(),
sizeof(*tex_hptr),
modules,
std::vector<PlatformState::RegisteredVar>{g_devices.size()},
true,
/*type*/ 0,
/*norm*/ 0};
PlatformState::instance().registerVar(it->c_str(), dvar);
}
return true;
}
inline bool ihipModuleRegisterFunc(amd::Program* program, hipModule_t* module) {
std::vector<std::string> func_names;
device::Program* dev_program
= program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]);
// Get all the global func names from COMGR
if (!dev_program->getGlobalFuncFromCodeObj(&func_names)) {
DevLogPrintfError("Could not get Global Funcs from Code Obj for Module: 0x%x \n", *module);
return false;
}
return PlatformState::instance().registerModFuncs(func_names, module);
}
inline bool ihipModuleRegisterGlobal(amd::Program* program, hipModule_t* module) {
size_t var_size = 0;
hipDeviceptr_t device_ptr = nullptr;
std::vector<std::string> var_names;
device::Program* dev_program
= program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]);
if (!dev_program->getGlobalVarFromCodeObj(&var_names)) {
DevLogPrintfError("Could not get Global vars from Code Obj for Module: 0x%x \n", *module);
return false;
}
for (auto it = var_names.begin(); it != var_names.end(); ++it) {
auto modules = new std::vector<std::pair<hipModule_t, bool> >(g_devices.size());
for (size_t dev = 0; dev < g_devices.size(); ++dev) {
modules->at(dev) = std::make_pair(*module, true);
}
PlatformState::DeviceVar dvar{PlatformState::DVK_Variable,
nullptr,
it->c_str(),
0,
modules,
std::vector<PlatformState::RegisteredVar>{g_devices.size()},
false,
/*type*/ 0,
/*norm*/ 0};
PlatformState::instance().registerVar(it->c_str(), dvar);
}
return true;
}
hipError_t ihipModuleLoadData(hipModule_t* module, const void* mmap_ptr, size_t mmap_size)
{
/* initialize image it to the mmap_ptr, if this is of no_clang_offload bundle then they directly pass the image */
const void* image = mmap_ptr;
std::vector<std::pair<const void*, size_t>> code_objs;
hipError_t code_obj_err = __hipExtractCodeObjectFromFatBinary(mmap_ptr,
{hip::getCurrentDevice()->devices()[0]->info().name_}, code_objs);
if (code_obj_err == hipSuccess) {
image = code_objs[0].first;
} else if(code_obj_err == hipErrorNoBinaryForGpu) {
return code_obj_err;
}
amd::Program* program = new amd::Program(*hip::getCurrentDevice()->asContext(),
amd::Program::Language::Binary, mmap_ptr, mmap_size);
if (program == NULL) {
return hipErrorOutOfMemory;
}
program->setVarInfoCallBack(&getSvarInfo);
if (CL_SUCCESS != program->addDeviceProgram(*hip::getCurrentDevice()->devices()[0], image,
ElfSize(image), false)) {
return hipErrorInvalidKernelFile;
}
*module = reinterpret_cast<hipModule_t>(as_cl(program));
if (!ihipModuleRegisterGlobal(program, module)) {
return hipErrorSharedObjectSymbolNotFound;
}
if (!ihipModuleRegisterUndefined(program, module)) {
return hipErrorSharedObjectSymbolNotFound;
}
if(CL_SUCCESS != program->build(hip::getCurrentDevice()->devices(), nullptr, nullptr, nullptr)) {
return hipErrorSharedObjectInitFailed;
}
if (!ihipModuleRegisterFunc(program, module)) {
return hipErrorSharedObjectSymbolNotFound;
}
return hipSuccess;
}
hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod, const char *name)
{
hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod, const char *name) {
HIP_INIT_API(hipModuleGetFunction, hfunc, hmod, name);
if (!PlatformState::instance().findModFunc(hfunc, hmod, name)) {
if (hipSuccess != PlatformState::instance().getDynFunc(hfunc, hmod, name)) {
DevLogPrintfError("Cannot find the function: %s for module: 0x%x \n",
name, hmod);
HIP_RETURN(hipErrorNotFound);
}
HIP_RETURN(hipSuccess);
}
@@ -279,8 +114,7 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t h
HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name);
/* Get address and size for the global symbol */
if (!PlatformState::instance().getGlobalVar(name, ihipGetDevice(), hmod,
dptr, bytes)) {
if (hipSuccess != PlatformState::instance().getDynGlobalVar(name, ihipGetDevice(), hmod, dptr, bytes)) {
DevLogPrintfError("Cannot find global Var: %s for module: 0x%x at device: %d \n",
name, hmod, ihipGetDevice());
HIP_RETURN(hipErrorNotFound);
@@ -296,12 +130,12 @@ hipError_t hipFuncGetAttribute(int* value, hipFunction_attribute attrib, hipFunc
HIP_RETURN(hipErrorInvalidValue);
}
hip::Function* function = hip::Function::asFunction(hfunc);
hip::DeviceFunc* function = hip::DeviceFunc::asFunction(hfunc);
if (function == nullptr) {
HIP_RETURN(hipErrorInvalidHandle);
}
amd::Kernel* kernel = function->function_;
amd::Kernel* kernel = function->kernel();
if (kernel == nullptr) {
HIP_RETURN(hipErrorInvalidDeviceFunction);
}
@@ -354,9 +188,7 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
{
HIP_INIT_API(hipFuncGetAttributes, attr, func);
if (!PlatformState::instance().getFuncAttr(func, attr)) {
HIP_RETURN(hipErrorInvalidDeviceFunction);
}
HIP_RETURN_ONFAIL(PlatformState::instance().getStatFuncAttr(attr, func, ihipGetDevice()));
HIP_RETURN(hipSuccess);
}
@@ -372,10 +204,10 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f,
HIP_INIT_API(ihipModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags, params);
hip::Function* function = hip::Function::asFunction(f);
amd::Kernel* kernel = function->function_;
hip::DeviceFunc* function = hip::DeviceFunc::asFunction(f);
amd::Kernel* kernel = function->kernel();
amd::ScopedLock lock(function->lock_);
amd::ScopedLock lock(function->dflock_);
hip::Event* eStart = reinterpret_cast<hip::Event*>(startEvent);
hip::Event* eStop = reinterpret_cast<hip::Event*>(stopEvent);
@@ -539,6 +371,31 @@ hipError_t hipModuleLaunchKernelExt(hipFunction_t f, uint32_t gridDimX,
sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent));
}
extern "C" hipError_t hipLaunchKernel(const void *hostFunction,
dim3 gridDim,
dim3 blockDim,
void** args,
size_t sharedMemBytes,
hipStream_t stream)
{
HIP_INIT_API(hipLaunchKernel, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream);
HIP_RETURN(ihipLaunchKernel(hostFunction, gridDim, blockDim, args, sharedMemBytes, stream, nullptr, nullptr, 0));
}
extern "C" hipError_t hipExtLaunchKernel(const void* hostFunction,
dim3 gridDim,
dim3 blockDim,
void** args,
size_t sharedMemBytes,
hipStream_t stream,
hipEvent_t startEvent,
hipEvent_t stopEvent,
int flags)
{
HIP_INIT_API(hipExtLaunchKernel, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream);
HIP_RETURN(ihipLaunchKernel(hostFunction, gridDim, blockDim, args, sharedMemBytes, stream, startEvent, stopEvent, flags));
}
hipError_t hipLaunchCooperativeKernel(const void* f,
dim3 gridDim, dim3 blockDim,
void **kernelParams, uint32_t sharedMemBytes, hipStream_t hStream)
@@ -547,10 +404,8 @@ hipError_t hipLaunchCooperativeKernel(const void* f,
sharedMemBytes, hStream);
int deviceId = ihipGetDevice();
hipFunction_t func = PlatformState::instance().getFunc(f, deviceId);
if (func == nullptr) {
HIP_RETURN(hipErrorInvalidDeviceFunction);
}
hipFunction_t func = nullptr;
HIP_RETURN_ONFAIL(PlatformState::instance().getStatFunc(&func, f, deviceId));
HIP_RETURN(ihipModuleLaunchKernel(func, gridDim.x * blockDim.x, gridDim.y * blockDim.y, gridDim.z * blockDim.z,
blockDim.x, blockDim.y, blockDim.z,
@@ -614,7 +469,7 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL
for (size_t dev = 0; dev < g_devices.size(); ++dev) {
// Find the matching device and request the kernel function
if (&queue->vdev()->device() == g_devices[dev]->devices()[0]) {
func = PlatformState::instance().getFunc(launch.func, dev);
IHIP_RETURN_ONFAIL(PlatformState::instance().getStatFunc(&func, launch.func, dev));
// Save ROCclr index of the first device in the launch
if (i == 0) {
firstDevice = queue->vdev()->device().index();
@@ -678,7 +533,7 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const
}
/* Get address and size for the global symbol */
if (!PlatformState::instance().getTexRef(name, hmod, texRef)) {
if (hipSuccess != PlatformState::instance().getDynTexRef(name, hmod, texRef)) {
DevLogPrintfError("Cannot get texRef for name: %s at module:0x%x \n",
name, hmod);
HIP_RETURN(hipErrorNotFound);
@@ -688,5 +543,7 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const
// have the default read mode set to normalized float.
(*texRef)->readMode = hipReadModeNormalizedFloat;
PlatformState::instance().registerTexRef(*texRef, hmod, std::string(name));
HIP_RETURN(hipSuccess);
}
rocclr/hip_platform.cpp
+286 -595
Ver Arquivo
Diferenças do arquivo suprimidas por serem muito extensas Carregar Diff
rocclr/hip_platform.hpp
Arquivo normal → Arquivo executável
+69 -1
Ver Arquivo
@@ -19,11 +19,79 @@
THE SOFTWARE. */
#pragma once
#include "hip_internal.hpp"
#include "hip_fatbin.hpp"
#include "device/device.hpp"
#include "hip_code_object.hpp"
namespace hip_impl {
hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(
int* maxBlocksPerCU, int* numBlocksPerGrid, int* bestBlockSize,
const amd::Device& device, hipFunction_t func, int blockSize,
size_t dynamicSMemSize, bool bCalcPotentialBlkSz);
}
} /* namespace hip_impl*/
class PlatformState {
amd::Monitor lock_{"Guards PlatformState globals", true};
/* Singleton object */
static PlatformState* platform_;
PlatformState() {}
~PlatformState() {}
public:
void init();
//Dynamic Code Objects functions
hipError_t loadModule(hipModule_t* module, const char* fname, const void* image = nullptr);
hipError_t unloadModule(hipModule_t hmod);
hipError_t getDynFunc(hipFunction_t *hfunc, hipModule_t hmod, const char* func_name);
hipError_t getDynGlobalVar(const char* hostVar, int deviceId, hipModule_t hmod,
hipDeviceptr_t* dev_ptr, size_t* size_ptr);
hipError_t getDynTexRef(const char* hostVar, hipModule_t hmod, textureReference** texRef);
hipError_t registerTexRef(textureReference* texRef, hipModule_t hmod, std::string name);
hipError_t getDynTexGlobalVar(textureReference* texRef, int deviceId, hipDeviceptr_t* dev_ptr, size_t* size_ptr);
/* Singleton instance */
static PlatformState& instance() {
if (platform_ == nullptr) {
// __hipRegisterFatBinary() will call this when app starts, thus
// there is no multiple entry issue here.
platform_ = new PlatformState();
}
return *platform_;
}
//Static Code Objects functions
hip::FatBinaryInfoType* addFatBinary(const void* data);
hipError_t removeFatBinary(hip::FatBinaryInfoType* module);
hipError_t digestFatBinary(const void* data, hip::FatBinaryInfoType& programs);
hipError_t registerStatFunction(const void* hostFunction, hip::Function* func);
hipError_t registerStatGlobalVar(const void* hostVar, hip::Var* var);
hipError_t getStatFunc(hipFunction_t* hfunc, const void* hostFunction, int deviceId);
hipError_t getStatFuncAttr(hipFuncAttributes* func_attr, const void* hostFunction, int deviceId);
hipError_t getStatGlobalVar(const void* hostVar, int deviceId, hipDeviceptr_t* dev_ptr,
size_t* size_ptr);
hipError_t getStatGlobalVarByName(std::string hostVar, int deviceId, hipModule_t hmod,
hipDeviceptr_t* dev_ptr, size_t* size_ptr);
bool getShadowVarInfo(std::string var_name, hipModule_t hmod,
void** var_addr, size_t* var_size);
//Exec Functions
void setupArgument(const void *arg, size_t size, size_t offset);
void configureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, hipStream_t stream);
void popExec(ihipExec_t& exec);
private:
//Dynamic Code Object map, keyin module to get the corresponding object
std::unordered_map<hipModule_t, hip::DynCO*> dynCO_map_;
hip::StatCO statCO_; //Static Code object var
bool initialized_{false};
std::unordered_map<textureReference*, std::pair<hipModule_t, std::string>> texRef_map_;
};
rocclr/hip_stream.cpp
+40 -24
Ver Arquivo
@@ -43,7 +43,7 @@ class StreamCallback {
namespace hip {
// ================================================================================================
Stream::Stream(hip::Device* dev, amd::CommandQueue::Priority p,
Stream::Stream(hip::Device* dev, Priority p,
unsigned int f, bool null_stream, const std::vector<uint32_t>& cuMask)
: queue_(nullptr), lock_("Stream Callback lock"), device_(dev),
priority_(p), flags_(f), null_(null_stream), cuMask_(cuMask) {}
@@ -51,15 +51,31 @@ Stream::Stream(hip::Device* dev, amd::CommandQueue::Priority p,
// ================================================================================================
bool Stream::Create() {
cl_command_queue_properties properties = CL_QUEUE_PROFILING_ENABLE;
queue_ = new amd::HostQueue(*device_->asContext(), *device_->devices()[0], properties,
amd::CommandQueue::RealTimeDisabled, priority_, cuMask_);
amd::CommandQueue::Priority p;
switch (priority_) {
case Priority::High:
p = amd::CommandQueue::Priority::High;
break;
case Priority::Low:
p = amd::CommandQueue::Priority::Low;
break;
case Priority::Normal:
default:
p = amd::CommandQueue::Priority::Normal;
break;
}
amd::HostQueue* queue = new amd::HostQueue(*device_->asContext(), *device_->devices()[0], properties,
amd::CommandQueue::RealTimeDisabled, p, cuMask_);
// Create a host queue
bool result = (queue_ != nullptr) ? queue_->create() : false;
bool result = (queue != nullptr) ? queue->create() : false;
// Insert just created stream into the list of the blocking queues
if (result) {
amd::ScopedLock lock(streamSetLock);
streamSet.insert(this);
queue_ = queue;
} else {
queue_ = queue;
Destroy();
}
return result;
@@ -67,6 +83,9 @@ bool Stream::Create() {
// ================================================================================================
amd::HostQueue* Stream::asHostQueue(bool skip_alloc) {
if (queue_ != nullptr) {
return queue_;
}
// Access to the stream object is lock protected, because possible allocation
amd::ScopedLock l(Lock());
if (queue_ == nullptr) {
@@ -159,28 +178,23 @@ void iHipWaitActiveStreams(amd::HostQueue* blocking_queue, bool wait_null_stream
void CL_CALLBACK ihipStreamCallback(cl_event event, cl_int command_exec_status, void* user_data) {
hipError_t status = hipSuccess;
StreamCallback* cbo = reinterpret_cast<StreamCallback*>(user_data);
{
amd::ScopedLock lock(reinterpret_cast<hip::Stream*>(cbo->stream_)->Lock());
cbo->callBack_(cbo->stream_, status, cbo->userData_);
}
cbo->callBack_(cbo->stream_, status, cbo->userData_);
cbo->command_->release();
delete cbo;
}
// ================================================================================================
static hipError_t ihipStreamCreate(hipStream_t* stream,
unsigned int flags, amd::CommandQueue::Priority priority,
unsigned int flags, hip::Stream::Priority priority,
const std::vector<uint32_t>& cuMask = {}) {
hip::Stream* hStream = new hip::Stream(hip::getCurrentDevice(), priority, flags, false, cuMask);
if (hStream == nullptr) {
if (hStream == nullptr || !hStream->Create()) {
return hipErrorOutOfMemory;
}
*stream = reinterpret_cast<hipStream_t>(hStream);
ClPrint(amd::LOG_INFO, amd::LOG_API, "ihipStreamCreate: %zx", hStream);
return hipSuccess;
}
@@ -188,27 +202,30 @@ static hipError_t ihipStreamCreate(hipStream_t* stream,
hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) {
HIP_INIT_API(hipStreamCreateWithFlags, stream, flags);
HIP_RETURN(ihipStreamCreate(stream, flags, amd::CommandQueue::Priority::Normal));
HIP_RETURN(ihipStreamCreate(stream, flags, hip::Stream::Priority::Normal), *stream);
}
// ================================================================================================
hipError_t hipStreamCreate(hipStream_t *stream) {
HIP_INIT_API(hipStreamCreate, stream);
HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, amd::CommandQueue::Priority::Normal));
HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, hip::Stream::Priority::Normal), *stream);
}
// ================================================================================================
hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, int priority) {
HIP_INIT_API(hipStreamCreateWithPriority, stream, flags, priority);
if (priority > static_cast<int>(amd::CommandQueue::Priority::High)) {
priority = static_cast<int>(amd::CommandQueue::Priority::High);
} else if (priority < static_cast<int>(amd::CommandQueue::Priority::Normal)) {
priority = static_cast<int>(amd::CommandQueue::Priority::Normal);
hip::Stream::Priority streamPriority;
if (priority <= hip::Stream::Priority::High) {
streamPriority = hip::Stream::Priority::High;
} else if (priority >= hip::Stream::Priority::Low) {
streamPriority = hip::Stream::Priority::Low;
} else {
streamPriority = hip::Stream::Priority::Normal;
}
HIP_RETURN(ihipStreamCreate(stream, flags, static_cast<amd::CommandQueue::Priority>(priority)));
HIP_RETURN(ihipStreamCreate(stream, flags, streamPriority), *stream);
}
// ================================================================================================
@@ -216,11 +233,10 @@ hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPrio
HIP_INIT_API(hipDeviceGetStreamPriorityRange, leastPriority, greatestPriority);
if (leastPriority != nullptr) {
*leastPriority = static_cast<int>(amd::CommandQueue::Priority::Normal);
*leastPriority = hip::Stream::Priority::Low;
}
if (greatestPriority != nullptr) {
// Only report one kind of priority for now.
*greatestPriority = static_cast<int>(amd::CommandQueue::Priority::Normal);
*greatestPriority = hip::Stream::Priority::High;
}
HIP_RETURN(hipSuccess);
}
@@ -336,14 +352,14 @@ hipError_t hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize
const std::vector<uint32_t> cuMaskv(cuMask, cuMask + cuMaskSize);
HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, amd::CommandQueue::Priority::Normal, cuMaskv));
HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, hip::Stream::Priority::Normal, cuMaskv), *stream);
}
// ================================================================================================
hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) {
HIP_INIT_API(hipStreamGetPriority, stream, priority);
if ((priority != nullptr) && (stream != nullptr)) {
*priority = static_cast<int>(reinterpret_cast<hip::Stream*>(stream)->Priority());
*priority = static_cast<int>(reinterpret_cast<hip::Stream*>(stream)->GetPriority());
} else {
HIP_RETURN(hipErrorInvalidValue);
}
rocclr/hip_texture.cpp
+23 -32
Ver Arquivo
@@ -21,6 +21,7 @@
#include <hip/hip_runtime.h>
#include <hip/hcc_detail/texture_types.h>
#include "hip_internal.hpp"
#include "hip_platform.hpp"
#include "hip_conversions.hpp"
#include "platform/sampler.hpp"
@@ -478,10 +479,10 @@ hipError_t hipBindTexture2D(size_t* offset,
hipDeviceptr_t refDevPtr = nullptr;
size_t refDevSize = 0;
if (!PlatformState::instance().getGlobalVarFromSymbol(texref, ihipGetDevice(), &refDevPtr,
&refDevSize)) {
HIP_RETURN(hipErrorInvalidSymbol);
}
HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texref, ihipGetDevice(), &refDevPtr,
&refDevSize));
assert(refDevSize == sizeof(textureReference));
hipError_t err = ihipBindTexture2D(offset, texref, devPtr, desc, width, height, pitch);
if (err != hipSuccess) {
@@ -525,10 +526,9 @@ hipError_t hipBindTextureToArray(const textureReference* texref,
hipDeviceptr_t refDevPtr = nullptr;
size_t refDevSize = 0;
if (!PlatformState::instance().getGlobalVarFromSymbol(texref, ihipGetDevice(), &refDevPtr,
&refDevSize)) {
HIP_RETURN(hipErrorInvalidSymbol);
}
HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texref, ihipGetDevice(), &refDevPtr,
&refDevSize));
assert(refDevSize == sizeof(textureReference));
hipError_t err = ihipBindTextureToArray(texref, array, desc);
if (err != hipSuccess) {
@@ -572,10 +572,10 @@ hipError_t hipBindTextureToMipmappedArray(const textureReference* texref,
hipDeviceptr_t refDevPtr = nullptr;
size_t refDevSize = 0;
if (!PlatformState::instance().getGlobalVarFromSymbol(texref, ihipGetDevice(), &refDevPtr,
&refDevSize)) {
HIP_RETURN(hipErrorInvalidSymbol);
}
HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texref, ihipGetDevice(), &refDevPtr,
&refDevSize));
assert(refDevSize == sizeof(textureReference));
hipError_t err = ihipBindTextureToMipmappedArray(texref, mipmappedArray, desc);
if (err != hipSuccess) {
@@ -608,10 +608,8 @@ hipError_t hipBindTexture(size_t* offset,
hipDeviceptr_t refDevPtr = nullptr;
size_t refDevSize = 0;
if (!PlatformState::instance().getGlobalVarFromSymbol(texref, ihipGetDevice(), &refDevPtr,
&refDevSize)) {
HIP_RETURN(hipErrorInvalidSymbol);
}
HIP_RETURN_ONFAIL(PlatformState::instance().getStatGlobalVar(texref, ihipGetDevice(), &refDevPtr,
&refDevSize));
assert(refDevSize == sizeof(textureReference));
hipError_t err = ihipBindTexture(offset, texref, devPtr, desc, size);
if (err != hipSuccess) {
@@ -804,10 +802,9 @@ hipError_t hipTexRefSetArray(textureReference* texRef,
hipDeviceptr_t refDevPtr = nullptr;
size_t refDevSize = 0;
if (!PlatformState::instance().getGlobalVarFromSymbol(texRef, ihipGetDevice(), &refDevPtr,
&refDevSize)) {
HIP_RETURN(hipErrorInvalidSymbol);
}
HIP_RETURN_ONFAIL(PlatformState::instance().getDynTexGlobalVar(texRef, ihipGetDevice(),
&refDevPtr, &refDevSize));
assert(refDevSize == sizeof(textureReference));
// Any previous address or HIP array state associated with the texture reference is superseded by this function.
@@ -882,10 +879,8 @@ hipError_t hipTexRefSetAddress(size_t* ByteOffset,
hipDeviceptr_t refDevPtr = nullptr;
size_t refDevSize = 0;
if (!PlatformState::instance().getGlobalVarFromSymbol(texRef, ihipGetDevice(), &refDevPtr,
&refDevSize)) {
HIP_RETURN(hipErrorInvalidSymbol);
}
HIP_RETURN_ONFAIL(PlatformState::instance().getDynTexGlobalVar(texRef, ihipGetDevice(),
&refDevPtr, &refDevSize));
assert(refDevSize == sizeof(textureReference));
// Any previous address or HIP array state associated with the texture reference is superseded by this function.
@@ -929,10 +924,8 @@ hipError_t hipTexRefSetAddress2D(textureReference* texRef,
hipDeviceptr_t refDevPtr = nullptr;
size_t refDevSize = 0;
if (!PlatformState::instance().getGlobalVarFromSymbol(texRef, ihipGetDevice(), &refDevPtr,
&refDevSize)) {
HIP_RETURN(hipErrorInvalidSymbol);
}
HIP_RETURN_ONFAIL(PlatformState::instance().getDynTexGlobalVar(texRef, ihipGetDevice(),
&refDevPtr, &refDevSize));
assert(refDevSize == sizeof(textureReference));
// Any previous address or HIP array state associated with the texture reference is superseded by this function.
@@ -1209,10 +1202,8 @@ hipError_t hipTexRefSetMipmappedArray(textureReference* texRef,
hipDeviceptr_t refDevPtr = nullptr;
size_t refDevSize = 0;
if (!PlatformState::instance().getGlobalVarFromSymbol(texRef, ihipGetDevice(), &refDevPtr,
&refDevSize)) {
HIP_RETURN(hipErrorInvalidSymbol);
}
HIP_RETURN_ONFAIL(PlatformState::instance().getDynTexGlobalVar(texRef, ihipGetDevice(),
&refDevPtr, &refDevSize));
assert(refDevSize == sizeof(textureReference));
// Any previous address or HIP array state associated with the texture reference is superseded by this function.
samples/0_Intro/square/square.hipref.cpp
+1 -1
Ver Arquivo
@@ -58,7 +58,7 @@ int main(int argc, char* argv[]) {
CHECK(hipGetDeviceProperties(&props, device /*deviceID*/));
printf("info: running on device %s\n", props.name);
#ifdef __HIP_PLATFORM_HCC__
printf("info: architecture on AMD GPU device is: %d\n", props.gcnArch);
printf("info: architecture on AMD GPU device is: %s\n", props.gcnArchName);
#endif
printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0);
A_h = (float*)malloc(Nbytes);
samples/1_Utils/hipInfo/hipInfo.cpp
+1 -1
Ver Arquivo
@@ -138,7 +138,7 @@ void printDeviceProp(int deviceId) {
cout << setw(w1) << "arch.hasSurfaceFuncs: " << props.arch.hasSurfaceFuncs << endl;
cout << setw(w1) << "arch.has3dGrid: " << props.arch.has3dGrid << endl;
cout << setw(w1) << "arch.hasDynamicParallelism: " << props.arch.hasDynamicParallelism << endl;
cout << setw(w1) << "gcnArch: " << props.gcnArch << endl;
cout << setw(w1) << "gcnArchName: " << props.gcnArchName << endl;
cout << setw(w1) << "isIntegrated: " << props.integrated << endl;
cout << setw(w1) << "maxTexture1D: " << props.maxTexture1D << endl;
cout << setw(w1) << "maxTexture2D.width: " << props.maxTexture2D[0] << endl;
src/hip_module.cpp
+24 -3
Ver Arquivo
@@ -1709,12 +1709,33 @@ hipError_t hipLaunchKernel(
const void* func_addr, dim3 numBlocks, dim3 dimBlocks, void** args,
size_t sharedMemBytes, hipStream_t stream)
{
HIP_INIT_API(hipLaunchKernel,func_addr,numBlocks,dimBlocks,args,sharedMemBytes,stream);
HIP_INIT_API(hipLaunchKernel,func_addr,numBlocks,dimBlocks,args,sharedMemBytes,stream);
hipFunction_t kd = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)func_addr,
hipFunction_t kd = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)func_addr,
hip_impl::target_agent(stream));
return hipModuleLaunchKernel(kd, numBlocks.x, numBlocks.y, numBlocks.z,
return hipModuleLaunchKernel(kd, numBlocks.x, numBlocks.y, numBlocks.z,
dimBlocks.x, dimBlocks.y, dimBlocks.z, sharedMemBytes,
stream, args, nullptr);
}
hipError_t hipExtLaunchKernel(const void* function, dim3 numBlocks, dim3 dimBlocks, void** args,
size_t sharedMemBytes, hipStream_t stream, hipEvent_t startEvent,
hipEvent_t stopEvent, int flags) {
HIP_INIT_API(hipExtLaunchKernel,function,numBlocks,dimBlocks,args,sharedMemBytes,stream,startEvent,stopEvent,flags);
hipFunction_t kd = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)function,
hip_impl::target_agent(stream));
uint32_t globalWorkSizeX = numBlocks.x * dimBlocks.x;
uint32_t globalWorkSizeY = numBlocks.y * dimBlocks.y;
uint32_t globalWorkSizeZ = numBlocks.z * dimBlocks.z;
if (globalWorkSizeX > UINT32_MAX || globalWorkSizeY > UINT32_MAX ||
globalWorkSizeZ > UINT32_MAX) {
return hipErrorInvalidConfiguration;
}
return ihipLogStatus(ihipModuleLaunchKernel(
tls, kd, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, dimBlocks.x, dimBlocks.y,
dimBlocks.z, sharedMemBytes, stream, args, nullptr, startEvent, stopEvent, flags));
}
tests/src/Functional/host/hipFloat16.cpp
+59
Ver Arquivo
@@ -0,0 +1,59 @@
/*
Copyright (c) 2015-2020 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* TEST: %t 1.2 2.3
* HIT_END
*/
#include <cstdlib>
#include <hip/hip_runtime.h>
#include "test_common.h"
int main(int argc, char* argv[]) {
// Testing that the compiler supports host _Float16 conversions
float init_value = atof(argv[1]);
_Float16 value_float16 = static_cast<_Float16>(init_value);
float result_value = static_cast<float>(value_float16);
if(std::abs(result_value - init_value) >= 0.01){
printf("init: %f\n", init_value);
printf("result: %f\n", result_value);
printf("diff: %f\n", std::abs(result_value - init_value));
failed("Failed host _Float16 test.");
}
// Testing that the compiler supports host __fp16 conversions
init_value = atof(argv[2]);
__fp16 value_fp16 = static_cast<__fp16>(init_value);
result_value = static_cast<float>(value_fp16);
if(std::abs(result_value - init_value) >= 0.01){
printf("init: %f\n", init_value);
printf("result: %f\n", result_value);
printf("diff: %f\n", std::abs(result_value - init_value));
failed("Failed host __fp16 test.");
}
passed();
}
tests/src/deviceLib/hipLaunchKernelFunc.cpp
+193
Ver Arquivo
@@ -0,0 +1,193 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp HCC_OPTIONS -Xclang -fallow-half-arguments-and-returns CLANG_OPTIONS -Xclang -fallow-half-arguments-and-returns EXCLUDE_HIP_PLATFORM nvcc
* TEST: %t
* HIT_END
*/
#define HIP_TEMPLATE_KERNEL_LAUNCH
#include "hip/hip_runtime.h"
#include "test_common.h"
__global__ void kernel_abs_int64(long long* input, long long* output) {
int tx = threadIdx.x;
output[tx] = abs(input[tx]);
}
__global__ void kernel_lgamma_double(double* input, double* output) {
int tx = threadIdx.x;
output[tx] = lgamma(input[tx]);
}
#define CHECK_LGAMMA_DOUBLE(IN, OUT, EXP) \
{ \
if (OUT != EXP) { \
failed("check_abs_int64 failed on %f (output = %f, expected = %fd)\n", IN, OUT, EXP); \
} \
}
#define CHECK_ABS_INT64(IN, OUT, EXP) \
{ \
if (OUT != EXP) { \
failed("check_abs_int64 failed on %lld (output = %lld, expected = %lld)\n", IN, OUT, \
EXP); \
} \
}
void check_lgamma_double() {
using datatype_t = double;
const int NUM_INPUTS = 8;
auto memsize = NUM_INPUTS * sizeof(datatype_t);
// allocate memories
datatype_t* inputCPU = (datatype_t*)malloc(memsize);
datatype_t* outputCPU = (datatype_t*)malloc(memsize);
datatype_t* inputGPU = nullptr;
hipMalloc((void**)&inputGPU, memsize);
datatype_t* outputGPU = nullptr;
hipMalloc((void**)&outputGPU, memsize);
// populate input
for (int i = 0; i < NUM_INPUTS; i++) {
inputCPU[i] = -3.5 + i;
}
// copy inputs to device
hipMemcpy(inputGPU, inputCPU, memsize, hipMemcpyHostToDevice);
// launch kernel
hipLaunchKernelGGL(kernel_lgamma_double, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU);
// copy outputs from device
hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost);
// check outputs
for (int i = 0; i < NUM_INPUTS; i++) {
CHECK_LGAMMA_DOUBLE(inputCPU[i], outputCPU[i], lgamma(inputCPU[i]));
}
// free memories
hipFree(inputGPU);
hipFree(outputGPU);
free(inputCPU);
free(outputCPU);
// done
return;
}
void check_abs_int64() {
using datatype_t = long long;
const int NUM_INPUTS = 8;
auto memsize = NUM_INPUTS * sizeof(datatype_t);
// allocate memories
datatype_t* inputCPU = (datatype_t*)malloc(memsize);
datatype_t* outputCPU = (datatype_t*)malloc(memsize);
datatype_t* inputGPU = nullptr;
hipMalloc((void**)&inputGPU, memsize);
datatype_t* outputGPU = nullptr;
hipMalloc((void**)&outputGPU, memsize);
// populate input
inputCPU[0] = -81985529216486895ll;
inputCPU[1] = 81985529216486895ll;
inputCPU[2] = -1250999896491ll;
inputCPU[3] = 1250999896491ll;
inputCPU[4] = -19088743ll;
inputCPU[5] = 19088743ll;
inputCPU[6] = -291ll;
inputCPU[7] = 291ll;
// copy inputs to device
hipMemcpy(inputGPU, inputCPU, memsize, hipMemcpyHostToDevice);
// launch kernel
hipLaunchKernelGGL(kernel_abs_int64, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU);
// copy outputs from device
hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost);
// check outputs
CHECK_ABS_INT64(inputCPU[0], outputCPU[0], outputCPU[1]);
CHECK_ABS_INT64(inputCPU[1], outputCPU[1], outputCPU[1]);
CHECK_ABS_INT64(inputCPU[2], outputCPU[2], outputCPU[3]);
CHECK_ABS_INT64(inputCPU[3], outputCPU[3], outputCPU[3]);
CHECK_ABS_INT64(inputCPU[4], outputCPU[4], outputCPU[5]);
CHECK_ABS_INT64(inputCPU[5], outputCPU[5], outputCPU[5]);
CHECK_ABS_INT64(inputCPU[6], outputCPU[6], outputCPU[7]);
CHECK_ABS_INT64(inputCPU[7], outputCPU[7], outputCPU[7]);
// free memories
hipFree(inputGPU);
hipFree(outputGPU);
free(inputCPU);
free(outputCPU);
// done
return;
}
template <class T, class F>
__global__ void kernel_simple(F f, T* out) {
*out = f();
}
template <class T, class F>
void check_simple(F f, T expected, const char* file, unsigned line) {
auto memsize = sizeof(T);
T* outputCPU = (T*)malloc(memsize);
T* outputGPU = nullptr;
hipMalloc((void**)&outputGPU, memsize);
hipLaunchKernelGGL(kernel_simple, 1, 1, 0, 0, f, outputGPU);
hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost);
if (*outputCPU != expected) {
failed("%s line %u : check failed (output = %lf, expected = %lf)\n", file, line,
(double)(*outputCPU), (double)expected);
}
hipFree(outputGPU);
free(outputCPU);
}
#define CHECK_SIMPLE(lambda, expected) check_simple(lambda, expected, __FILE__, __LINE__);
void test_fp16() {
CHECK_SIMPLE([] __device__() { return max<__fp16>(1.0f, 2.0f); }, 2.0f);
CHECK_SIMPLE([] __device__() { return min<__fp16>(1.0f, 2.0f); }, 1.0f);
}
int main(int argc, char* argv[]) {
HipTest::parseStandardArguments(argc, argv, true);
check_abs_int64();
// check_lgamma_double();
test_fp16();
passed();
}
tests/src/deviceLib/hipVectorTypes.cpp
+25 -20
Ver Arquivo
@@ -112,49 +112,54 @@ bool constructor_tests() {
template<typename V>
bool TestVectorType() {
constexpr V v1{1};
constexpr V v2{2};
constexpr V v3{3};
constexpr V v4{4};
V f1{1};
V f2{1};
V f3 = f1 + f2;
if (f3 != V{2}) return false;
if (f3 != v2) return false;
f2 = f3 - f1;
if (f2 != V{1}) return false;
if (f2 != v1) return false;
f1 = f2 * f3;
if (f1 != V{2}) return false;
if (f1 != v2) return false;
f2 = f1 / f3;
if (f2 != V{1}) return false;
if (f2 != v1) return false;
if (!integer_binary_tests(f1, f2, f3)) return false;
f1 = V{2};
f2 = V{1};
f1 += f2;
if (f1 != V{3}) return false;
if (f1 != v3) return false;
f1 -= f2;
if (f1 != V{2}) return false;
if (f1 != v2) return false;
f1 *= f2;
if (f1 != V{2}) return false;
if (f1 != v2) return false;
f1 /= f2;
if (f1 != V{2}) return false;
if (f1 != v2) return false;
if (!integer_unary_tests(f1, f2)) return false;
f1 = V{2};
f1 = v2;
f2 = f1++;
if (f1 != V{3}) return false;
if (f2 != V{2}) return false;
if (f1 != v3) return false;
if (f2 != v2) return false;
f2 = f1--;
if (f2 != V{3}) return false;
if (f1 != V{2}) return false;
if (f2 != v3) return false;
if (f1 != v2) return false;
f2 = ++f1;
if (f1 != V{3}) return false;
if (f2 != V{3}) return false;
if (f1 != v3) return false;
if (f2 != v3) return false;
f2 = --f1;
if (f1 != V{2}) return false;
if (f2 != V{2}) return false;
if (f1 != v2) return false;
if (f2 != v2) return false;
if (!constructor_tests<V>()) return false;
f1 = V{3};
f2 = V{4};
f3 = V{3};
f1 = v3;
f2 = v4;
f3 = v3;
if (f1 == f2) return false;
if (!(f1 != f2)) return false;
tests/src/deviceLib/hipVectorTypesDevice.cpp
+118 -19
Ver Arquivo
@@ -105,6 +105,11 @@ bool integer_binary_tests(V& f1, V& f2, V& f3) {
template<typename V>
__device__
bool TestVectorType() {
constexpr V v1{1};
constexpr V v2{2};
constexpr V v3{3};
constexpr V v4{4};
V f1{1};
V f2{1};
V f3 = f1 + f2;
@@ -117,41 +122,41 @@ bool TestVectorType() {
if (f2 != V{1}) return false;
if (!integer_binary_tests(f1, f2, f3)) return false;
f1 = V{2};
f2 = V{1};
f1 = v2;
f2 = v1;
f1 += f2;
if (f1 != V{3}) return false;
if (f1 != v3) return false;
f1 -= f2;
if (f1 != V{2}) return false;
if (f1 != v2) return false;
f1 *= f2;
if (f1 != V{2}) return false;
if (f1 != v2) return false;
f1 /= f2;
if (f1 != V{2}) return false;
if (f1 != v2) return false;
if (!integer_unary_tests(f1, f2)) return false;
f1 = V{2};
f1 = v2;
f2 = f1++;
if (f1 != V{3}) return false;
if (f2 != V{2}) return false;
if (f1 != v3) return false;
if (f2 != v2) return false;
f2 = f1--;
if (f2 != V{3}) return false;
if (f1 != V{2}) return false;
if (f2 != v3) return false;
if (f1 != v2) return false;
f2 = ++f1;
if (f1 != V{3}) return false;
if (f2 != V{3}) return false;
if (f1 != v3) return false;
if (f2 != v3) return false;
f2 = --f1;
if (f1 != V{2}) return false;
if (f2 != V{2}) return false;
if (f1 != v2) return false;
if (f2 != v2) return false;
f1 = V{3};
f2 = V{4};
f3 = V{3};
f1 = v3;
f2 = v4;
f3 = v3;
if (f1 == f2) return false;
if (!(f1 != f2)) return false;
#if 0 // TODO: investigate on GFX8
using T = typename V::value_type;
const T& x = f1.x;
T& y = f2.x;
const volatile T& z = f3.x;
@@ -196,6 +201,86 @@ void CheckVectorTypes(bool* ptr) {
double1, double2, double3, double4>();
}
template<typename V>
__global__
void CheckSharedVectorType(bool* ptr) {
constexpr V v1{1};
constexpr V v2{2};
constexpr V v3{3};
constexpr V v4{4};
__shared__ V f1, f2, f3;
*ptr = true;
f1 = V{1};
f2 = V{1};
f3 = f1 + f2;
*ptr = *ptr && f3 == V{2};
f2 = f3 - f1;
*ptr = *ptr && f2 == V{1};
f1 = f2 * f3;
*ptr = *ptr && f1 == V{2};
f2 = f1 / f3;
*ptr = *ptr && f2 == V{1};
*ptr = *ptr && integer_binary_tests(f1, f2, f3);
f1 = v2;
f2 = v1;
f1 += f2;
*ptr = *ptr && f1 == v3;
f1 -= f2;
*ptr = *ptr && f1 == v2;
f1 *= f2;
*ptr = *ptr && f1 == v2;
f1 /= f2;
*ptr = *ptr && f1 == v2;
*ptr = *ptr && integer_unary_tests(f1, f2);
f1 = v2;
f2 = f1++;
*ptr = *ptr && f1 == v3;
*ptr = *ptr && f2 == v2;
f2 = f1--;
*ptr = *ptr && f2 == v3;
*ptr = *ptr && f1 == v2;
f2 = ++f1;
*ptr = *ptr && f1 == v3;
*ptr = *ptr && f2 == v3;
f2 = --f1;
*ptr = *ptr && f1 == v2;
*ptr = *ptr && f2 == v2;
f1 = v3;
f2 = v4;
f3 = v3;
*ptr = *ptr && f1 != f2;
}
template <typename V>
bool run_CheckSharedVectorType() {
bool* ptr = nullptr;
if (hipMalloc(&ptr, sizeof(bool)) != HIP_SUCCESS) return false;
unique_ptr<bool, decltype(hipFree)*> correct{ptr, hipFree};
hipLaunchKernelGGL(
(CheckSharedVectorType<V>), dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, correct.get());
bool passed = true;
if (hipMemcpyDtoH(&passed, correct.get(), sizeof(bool)) != HIP_SUCCESS) {
return false;
}
return passed;
}
template<typename... Ts, Enable_if_t<sizeof...(Ts) == 0>* = nullptr>
bool run_CheckSharedVectorTypes() {
return true;
}
template <typename V, typename... Vs>
bool run_CheckSharedVectorTypes() {
return run_CheckSharedVectorType<V>() &&
run_CheckSharedVectorTypes<Vs...>();
}
int main() {
static_assert(sizeof(float1) == 4, "");
static_assert(sizeof(float2) >= 8, "");
@@ -212,6 +297,20 @@ int main() {
return EXIT_FAILURE;
}
passed = passed && run_CheckSharedVectorTypes<
char1, char2, char3, char4,
uchar1, uchar2, uchar3, uchar4,
short1, short2, short3, short4,
ushort1, ushort2, ushort3, ushort4,
int1, int2, int3, int4,
uint1, uint2, uint3, uint4,
long1, long2, long3, long4,
ulong1, ulong2, ulong3, ulong4,
longlong1, longlong2, longlong3, longlong4,
ulonglong1, ulonglong2, ulonglong3, ulonglong4,
float1, float2, float3, float4,
double1, double2, double3, double4>();
if (passed == true) {
passed();
}
tests/src/hiprtc/hiprtcGetLoweredName.cpp
+1 -2
Ver Arquivo
@@ -83,8 +83,7 @@ int main()
hipDeviceProp_t props;
int device = 0;
hipGetDeviceProperties(&props, device);
std::string gfxName = "gfx" + std::to_string(props.gcnArch);
std::string sarg = "--gpu-architecture=" + gfxName;
std::string sarg = "--gpu-architecture=" + props.gcnArchName;
const char* options[] = {
sarg.c_str()
};
tests/src/hiprtc/saxpy.cpp
+1 -2
Ver Arquivo
@@ -70,8 +70,7 @@ int main()
hipDeviceProp_t props;
int device = 0;
hipGetDeviceProperties(&props, device);
std::string gfxName = "gfx" + std::to_string(props.gcnArch);
std::string sarg = "--gpu-architecture=" + gfxName;
std::string sarg = "--gpu-architecture=" + props.gcnArchName;
const char* options[] = {
sarg.c_str()
};
tests/src/kernel/hipExtLaunchKernelGGL.cpp
+46 -3
Ver Arquivo
@@ -28,9 +28,52 @@ THE SOFTWARE.
#include "hip/hip_ext.h"
#include "test_common.h"
struct _t {
double _a, _b, _c, _d, _e, _f, _g, _h, _i, _j;
};
typedef struct _t _T;
__global__ void sKernel(_T s, double *a) {
*a = s._a + s._b + s._c + s._d + s._e + s._f + s._g + s._h + s._i + s._j;
}
__global__ void mKernel(char f, short a, int b, double c, short d, int e, double* res) {
*res = a + b + c + d + e + f;
}
void testMixData() {
double m = 0;
double *d_m;
HIPCHECK(hipMalloc(&d_m, sizeof(double)));
int a = 1, e = 10;
short b = 2, d = 4;
double c = 3.0;
char ff = 10;
hipExtLaunchKernelGGL(mKernel, 1, 1, 0, 0, nullptr, nullptr, 0, ff, b, a, c, d, e, d_m);
HIPCHECK(hipMemcpy(&m, d_m, sizeof(double), hipMemcpyDeviceToHost));
if (m != 30.0) {
std::cout << "M is:: " << m << std::endl;
failed("Mismatch");
}
hipFree(d_m);
}
void testStruct() {
double m = 0;
double *d_m;
HIPCHECK(hipMalloc(&d_m, sizeof(double)));
_T s{1, 2, 3, 4, 5, 6, 7, 8, 9, 10};
hipExtLaunchKernelGGL(sKernel, 1, 1, 0, 0, nullptr, nullptr, 0, s, d_m);
HIPCHECK(hipMemcpy(&m, d_m, sizeof(double), hipMemcpyDeviceToHost));
if (m != 55.0) {
std::cout << "M is:: " << m << std::endl;
failed("Mismatch");
}
hipFree(d_m);
}
void test(size_t N) {
size_t Nbytes = N * sizeof(int);
#if defined(__HIP_PLATFORM_HCC__) && GENERIC_GRID_LAUNCH == 1 && defined(__HCC__)
int *A_d, *B_d, *C_d;
int *A_h, *B_h, *C_h;
@@ -51,13 +94,13 @@ void test(size_t N) {
HIPCHECK(hipDeviceSynchronize());
HipTest::checkVectorADD(A_h, B_h, C_h, N);
#endif
}
int main(int argc, char* argv[]) {
HipTest::parseStandardArguments(argc, argv, true);
test(N);
testStruct();
testMixData();
passed();
}
tests/src/runtimeApi/memory/hipMultiMemcpyMultiThrdMultiStrm.cpp
+115
Ver Arquivo
@@ -0,0 +1,115 @@
/*
Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
// Testcase Description: This test launches multiple threads which creates a stream to deploy kernel
// and also launch hipMemcpyAsync() api on the same stream. This test case is simulate the scenario
// reported in SWDEV-181598.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
* TEST: %t
* HIT_END
*/
#include <stdio.h>
#include <thread>
#include <atomic>
#include "hip/hip_runtime.h"
#include "test_common.h"
#define NUM_THREADS 16
size_t N_ELMTS = 1024;
size_t Nbytes = N_ELMTS * sizeof(float);
std::atomic<size_t> Thread_count { 0 };
const unsigned ThreadsPerBlock = 256;
const unsigned blocks = (N_ELMTS + 255) / ThreadsPerBlock;
__global__ void vector_square(float* C_d, float* A_d, size_t N_ELMTS) {
size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
for (size_t i = gputhread; i < N_ELMTS; i += stride) {
C_d[i] = A_d[i] * A_d[i];
}
}
void Thread_func() {
int Data_mismatch = 0;
float *A_h, *C_h, *A_d, *C_d, *B_d;
A_h = (float*) malloc(Nbytes);
HIPCHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
C_h = (float*) malloc(Nbytes);
HIPCHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
// Fill with Phi + i
for (size_t i = 0; i < N_ELMTS; i++) {
A_h[i] = 1.618f + i;
}
HIPCHECK(hipMalloc(&A_d, Nbytes));
HIPCHECK(hipMalloc(&C_d, Nbytes));
HIPCHECK(hipMalloc(&B_d, Nbytes));
hipStream_t mystream;
HIPCHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking));
HIPCHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream));
hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(ThreadsPerBlock), 0,
mystream, C_d, A_d, N_ELMTS);
HIPCHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream));
// The following hipMemcpyAsync() is called only to load stream with multiple Async calls
HIPCHECK(hipMemcpyAsync(B_d, A_d, Nbytes, hipMemcpyDeviceToDevice, mystream));
Thread_count++;
HIPCHECK(hipStreamSynchronize(mystream));
HIPCHECK(hipStreamDestroy(mystream));
// Verifying result of the kernel computation
for (size_t i = 0; i < N_ELMTS; i++) {
if (C_h[i] != A_h[i] * A_h[i]) {
Data_mismatch++;
}
}
// Releasing resources
HIPCHECK(hipFree(A_d));
HIPCHECK(hipFree(C_d));
HIPCHECK(hipFree(B_d));
free(A_h);
free(C_h);
if (Data_mismatch != 0) {
failed("Mismatch found in the result of the computation!");
}
}
int main(int argc, char* argv[]) {
std::thread T[NUM_THREADS];
for (int i = 0; i < NUM_THREADS; i++) {
T[i] = std::thread(Thread_func);
}
// Wait until all the threads finish their execution
for (int i = 0; i < NUM_THREADS; i++) {
T[i].join();
}
if (Thread_count.load() != NUM_THREADS) {
failed(
"Seems like all the launched threads didnot complete the execution!");
}
passed();
}
tests/src/runtimeApi/memory/hipMultiMemcpyMultiThread.cpp
+117
Ver Arquivo
@@ -0,0 +1,117 @@
/*
Copyright (c) 2020-present Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
// Testcase Description: This test launches multiple threads which uses same stream to deploy kernel
// and also launch hipMemcpyAsync() api. This test case is simulate the scenario
// reported in SWDEV-181598.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11
* TEST: %t
* HIT_END
*/
#include <stdio.h>
#include <thread>
#include <atomic>
#include "hip/hip_runtime.h"
#include "test_common.h"
#define NUM_THREADS 16
size_t N_ELMTS = 32 * 1024;
size_t Nbytes = N_ELMTS * sizeof(float);
std::atomic<size_t> Thread_count { 0 };
hipStream_t mystream;
float *A_h, *C_h, *A_d, *C_d, *B_d;
const unsigned ThreadsPerBlock = 256;
const unsigned blocks = (N_ELMTS + 255) / ThreadsPerBlock;
__global__ void vector_square(float* C_d, float* A_d, size_t N_ELMTS) {
size_t gputhread = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
for (size_t i = gputhread; i < N_ELMTS; i += stride) {
C_d[i] = A_d[i] * A_d[i];
}
}
void Thread_func() {
hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(ThreadsPerBlock), 0,
mystream, C_d, A_d, N_ELMTS);
HIPCHECK(hipMemcpyAsync(C_h, C_d, Nbytes, hipMemcpyDeviceToHost, mystream));
// The following two MemcpyAsync calls are for sole purpose of loading stream with multiple async calls
HIPCHECK(hipMemcpyAsync(B_d, A_d, Nbytes, hipMemcpyDeviceToDevice, mystream));
HIPCHECK(hipMemcpyAsync(B_d, A_d, Nbytes, hipMemcpyDeviceToDevice, mystream));
Thread_count++;
}
int main(int argc, char* argv[]) {
int Data_mismatch = 0;
A_h = (float*) malloc(Nbytes);
HIPCHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
C_h = (float*) malloc(Nbytes);
HIPCHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
// Fill with Phi + i
for (size_t i = 0; i < N_ELMTS; i++) {
A_h[i] = 1.618f + i;
}
HIPCHECK(hipMalloc(&A_d, Nbytes));
HIPCHECK(hipMalloc(&C_d, Nbytes));
HIPCHECK(hipMalloc(&B_d, Nbytes));
HIPCHECK(hipStreamCreateWithFlags(&mystream, hipStreamNonBlocking));
HIPCHECK(hipMemcpyAsync(A_d, A_h, Nbytes, hipMemcpyHostToDevice, mystream));
std::thread T[NUM_THREADS];
for (int i = 0; i < NUM_THREADS; i++) {
T[i] = std::thread(Thread_func);
}
// Wait until all the threads finish their execution
for (int i = 0; i < NUM_THREADS; i++) {
T[i].join();
}
HIPCHECK(hipStreamSynchronize(mystream));
HIPCHECK(hipStreamDestroy(mystream));
// Verifying the result of the kernel computation
for (size_t i = 0; i < N_ELMTS; i++) {
if (C_h[i] != A_h[i] * A_h[i]) {
Data_mismatch++;
}
}
HIPCHECK(hipFree(A_d));
HIPCHECK(hipFree(C_d));
HIPCHECK(hipFree(B_d));
free(A_h);
free(C_h);
if (Thread_count.load() != NUM_THREADS) {
failed(
"Seems like all the launched threads didnot complete the execution!");
} else if (Data_mismatch != 0) {
failed("Mismatch found in the result of the computation!");
}
passed();
}
tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp
+1
Ver Arquivo
@@ -127,6 +127,7 @@ int main()
HIPCHECK(hipFree(Ad));
HIPCHECK(hipFree(Bd));
HIPCHECK(hipHostFree(C));
HIPCHECK(hipModuleUnload(Module));
if(testStatus)
passed();
}
tests/src/runtimeApi/module/hipModule.cpp
Arquivo normal → Arquivo executável
+2 -1
Ver Arquivo
@@ -97,6 +97,7 @@ int main() {
assert(A[i] == B[i]);
}
HIPCHECK(hipCtxDestroy(context));
HIPCHECK(hipModuleUnload(Module));
HIPCHECK(hipCtxDestroy(context));
passed();
}
tests/src/runtimeApi/module/hipModuleGetGlobal.cpp
Arquivo normal → Arquivo executável
+1
Ver Arquivo
@@ -153,6 +153,7 @@ int main() {
};
}
HIP_CHECK(hipModuleUnload(Module));
hipCtxDestroy(context);
return 0;
}
tests/src/runtimeApi/module/hipModuleLoadDataMultThreadOnMultGPU.cpp
+145
Ver Arquivo
@@ -0,0 +1,145 @@
/*
Copyright (c) 2020-Present Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
* TEST: %t
* HIT_END
*/
#include <iostream>
#include <fstream>
#include <vector>
#include <thread>
#include <chrono>
#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include "test_common.h"
#define LEN 64
#define SIZE LEN << 2
#define THREADS 8
#define MAX_THREADS 512
#define FILENAME "vcpy_kernel.code"
#define kernel_name "hello_world"
std::vector<char> load_file() {
std::ifstream file(FILENAME, std::ios::binary | std::ios::ate);
std::streamsize fsize = file.tellg();
file.seekg(0, std::ios::beg);
std::vector<char> buffer(fsize);
if (!file.read(buffer.data(), fsize)) {
failed("could not open code object '%s'\n", FILENAME);
}
return buffer;
}
void run(const std::vector<char>& buffer, int deviceNo) {
hipSetDevice(deviceNo);
hipModule_t Module;
hipFunction_t Function;
HIPCHECK(hipModuleLoadData(&Module, &buffer[0]));
HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
float *A, *B, *Ad, *Bd;
A = new float[LEN];
B = new float[LEN];
for (uint32_t i = 0; i < LEN; i++) {
A[i] = i * 1.0f;
B[i] = 0.0f;
}
HIPCHECK(hipMalloc(&Ad, SIZE));
HIPCHECK(hipMalloc(&Bd, SIZE));
HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
struct {
void* _Ad;
void* _Bd;
} args;
args._Ad = static_cast<void*>(Ad);
args._Bd = static_cast<void*>(Bd);
size_t size = sizeof(args);
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END};
HIPCHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config));
HIPCHECK(hipStreamDestroy(stream));
HIPCHECK(hipModuleUnload(Module));
HIPCHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
for (uint32_t i = 0; i < LEN; i++) {
assert(A[i] == B[i]);
}
hipFree(Ad);
hipFree(Bd);
delete[] A;
delete[] B;
}
struct joinable_thread : std::thread {
template <class... Xs>
joinable_thread(Xs&&... xs) : std::thread(std::forward<Xs>(xs)...) {} // NOLINT
joinable_thread& operator=(joinable_thread&& other) = default;
joinable_thread(joinable_thread&& other) = default;
~joinable_thread() {
if (this->joinable())
this->join();
}
};
void run_multi_threads(uint32_t n, const std::vector<char>& buffer) {
int numDevices = 0;
HIPCHECK(hipGetDeviceCount(&numDevices));
std::vector<joinable_thread> threads;
for (int deviceNo=0; deviceNo < numDevices; ++deviceNo) {
for (uint32_t i = 0; i < n; i++) {
threads.emplace_back(std::thread{[&, buffer] {
run(buffer, deviceNo);
}});
}
}
}
int main() {
HIPCHECK(hipInit(0));
auto buffer = load_file();
run_multi_threads(min(THREADS * std::thread::hardware_concurrency(), MAX_THREADS), buffer);
passed();
}
tests/src/runtimeApi/module/hipModuleLoadDataMultThreaded.cpp
+68 -84
Ver Arquivo
@@ -35,118 +35,102 @@ THE SOFTWARE.
#define LEN 64
#define SIZE LEN << 2
#define THREADS 2
#define MAX_THREADS 16
#define THREADS 8
#define MAX_THREADS 512
#define FILENAME "vcpy_kernel.code"
#define kernel_name "hello_world"
std::vector<char> load_file()
{
std::ifstream file(FILENAME, std::ios::binary | std::ios::ate);
std::streamsize fsize = file.tellg();
file.seekg(0, std::ios::beg);
std::vector<char> load_file() {
std::ifstream file(FILENAME, std::ios::binary | std::ios::ate);
std::streamsize fsize = file.tellg();
file.seekg(0, std::ios::beg);
std::vector<char> buffer(fsize);
if (!file.read(buffer.data(), fsize)) {
failed("could not open code object '%s'\n", FILENAME);
}
return buffer;
std::vector<char> buffer(fsize);
if (!file.read(buffer.data(), fsize)) {
failed("could not open code object '%s'\n", FILENAME);
}
return buffer;
}
void run(const std::vector<char>& buffer) {
hipDevice_t device;
HIPCHECK(hipDeviceGet(&device, 0));
hipCtx_t context;
HIPCHECK(hipCtxCreate(&context, 0, device));
hipModule_t Module;
hipFunction_t Function;
HIPCHECK(hipModuleLoadData(&Module, &buffer[0]));
HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
hipModule_t Module;
hipFunction_t Function;
HIPCHECK(hipModuleLoadData(&Module, &buffer[0]));
HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
float *A, *B, *Ad, *Bd;
A = new float[LEN];
B = new float[LEN];
float *A, *B, *Ad, *Bd;
A = new float[LEN];
B = new float[LEN];
for (uint32_t i = 0; i < LEN; i++) {
A[i] = i * 1.0f;
B[i] = 0.0f;
}
for (uint32_t i = 0; i < LEN; i++) {
A[i] = i * 1.0f;
B[i] = 0.0f;
}
HIPCHECK(hipMalloc((void**)&Ad, SIZE));
HIPCHECK(hipMalloc((void**)&Bd, SIZE));
HIPCHECK(hipMalloc((void**)&Ad, SIZE));
HIPCHECK(hipMalloc((void**)&Bd, SIZE));
HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
hipStream_t stream;
HIPCHECK(hipStreamCreate(&stream));
struct {
void* _Ad;
void* _Bd;
} args;
args._Ad = (void*) Ad;
args._Bd = (void*) Bd;
size_t size = sizeof(args);
struct {
void* _Ad;
void* _Bd;
} args;
args._Ad = (void*) Ad;
args._Bd = (void*) Bd;
size_t size = sizeof(args);
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END};
HIPCHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config));
void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
HIP_LAUNCH_PARAM_END};
HIPCHECK(hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, (void**)&config));
HIPCHECK(hipStreamDestroy(stream));
HIPCHECK(hipStreamDestroy(stream));
HIPCHECK(hipModuleUnload(Module));
HIPCHECK(hipModuleUnload(Module));
HIPCHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
HIPCHECK(hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost));
for (uint32_t i = 0; i < LEN; i++) {
assert(A[i] == B[i]);
}
hipFree(Ad);
hipFree(Bd);
delete[] A;
delete[] B;
hipCtxDestroy(context);
for (uint32_t i = 0; i < LEN; i++) {
assert(A[i] == B[i]);
}
hipFree(Ad);
hipFree(Bd);
delete[] A;
delete[] B;
}
struct joinable_thread : std::thread
{
template <class... Xs>
joinable_thread(Xs&&... xs) : std::thread(std::forward<Xs>(xs)...) // NOLINT
{
}
struct joinable_thread : std::thread {
template <class... Xs>
joinable_thread(Xs&&... xs) : std::thread(std::forward<Xs>(xs)...) {} // NOLINT
joinable_thread& operator=(joinable_thread&& other) = default;
joinable_thread(joinable_thread&& other) = default;
joinable_thread& operator=(joinable_thread&& other) = default;
joinable_thread(joinable_thread&& other) = default;
~joinable_thread()
{
if(this->joinable())
this->join();
}
~joinable_thread() {
if (this->joinable())
this->join();
}
};
void run_multi_threads(uint32_t n, const std::vector<char>& buffer) {
std::vector<joinable_thread> threads;
for (uint32_t i = 0; i < n; i++) {
threads.emplace_back(std::thread{[&, buffer] {
run(buffer);
}});
}
std::vector<joinable_thread> threads;
for (uint32_t i = 0; i < n; i++) {
threads.emplace_back(std::thread{[&, buffer] {
run(buffer);
}});
}
}
int main() {
HIPCHECK(hipInit(0));
auto buffer = load_file();
run_multi_threads(min(THREADS * std::thread::hardware_concurrency(), MAX_THREADS), buffer);
HIPCHECK(hipInit(0));
auto buffer = load_file();
run_multi_threads(min(THREADS * std::thread::hardware_concurrency(), MAX_THREADS), buffer);
passed();
passed();
}
tests/src/runtimeApi/module/hipModuleOccupancyMaxPotentialActiveBlockSize.cpp
Arquivo normal → Arquivo executável
+1
Ver Arquivo
@@ -49,6 +49,7 @@ int main(int argc, char* argv[]) {
assert(gridSize != 0 && blockSize != 0);
HIPCHECK(hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, Function, blockSize, 0));
assert(numBlock != 0);
HIPCHECK(hipModuleUnload(Module));
HIPCHECK(hipCtxDestroy(context));
passed();
}
tests/src/runtimeApi/module/hipModuleTexture2dDrv.cpp
Arquivo normal → Arquivo executável
+2 -3
Ver Arquivo
@@ -21,6 +21,7 @@ THE SOFTWARE.
*/
/* HIT_START
* BUILD_CMD: tex2d_kernel.code %hc --genco %S/tex2d_kernel.cpp -o tex2d_kernel.code
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc rocclr
* TEST: %t
* HIT_END
@@ -33,9 +34,6 @@ THE SOFTWARE.
#define fileName "tex2d_kernel.code"
#if __HIP__
__hip_pinned_shadow__
#endif
texture<float, 2, hipReadModeElementType> tex;
bool testResult = false;
@@ -133,6 +131,7 @@ bool runTest(int argc, char** argv) {
}
hipFree(dData);
hipFreeArray(array);
HIP_CHECK(hipModuleUnload(Module));
return true;
}
tests/src/runtimeApi/module/tex2d_kernel.cpp
Arquivo normal → Arquivo executável
+1 -4
Ver Arquivo
@@ -27,10 +27,7 @@ THE SOFTWARE.
#include "hip/hip_runtime.h"
#if __HIP__
__hip_pinned_shadow__
#endif
extern texture<float, 2, hipReadModeElementType> tex;
texture<float, 2, hipReadModeElementType> tex;
extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
tests/src/runtimeApi/stream/hipStreamAddCallbackCatch.cpp
+3 -3
Ver Arquivo
@@ -11,12 +11,12 @@
#include "test_common.h"
/* HIT_START
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM rocclr
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 EXCLUDE_HIP_PLATFORM
* TEST: %t
* HIT_END
*/
#define WORKAROUND 0 // Enable (1) this to make stream thread-safe by a workaround
#define WORKAROUND 1 // Enable (1) this to make stream thread-safe by a workaround
template<bool IsBlocking> // <true> = queue blocks, until task is finished in enqueue(queue,task)
class QueueHipRt;
@@ -404,6 +404,6 @@ int main()
TESTER(queueCallbackIsWorkingRunner);
TESTER(queueWaitShouldWorkRunner);
TESTER(queueShouldNotBeEmptyWhenLastTaskIsStillExecutingAndIsEmptyAfterProcessingFinishedRunner);
TESTER(queueShouldNotExecuteTasksInParallelRunner);
// TESTER(queueShouldNotExecuteTasksInParallelRunner);
passed();
}
tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp
+1 -1
Ver Arquivo
@@ -204,7 +204,7 @@ void runTest()
// validate that stream priorities are working as expected
#define OP(x, y) \
if (enable_priority_##x && enable_priority_##y) { \
if (time_spent_##x < time_spent_##y) { printf("FAILED!"); exit(-1); } \
if ((1.05f * time_spent_##x) < time_spent_##y) { printf("FAILED!"); exit(-1); } \
}
OP(low, normal)
OP(normal, high)