AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

name change vdi to rocclr

Browse Source 
Change-Id: I06d198bbb4a499e153b290b73a92afed3553b252
This commit is contained in:
Payam
2020-04-29 23:28:46 -04:00
parent b1e13d8af2
commit c5f76c3de3
71 changed files with 155 additions and 149 deletions
Download Patch File Download Diff File Expand all files Collapse all files
rocclr/CMakeLists.txt
+232
View File
@@ -0,0 +1,232 @@
#project("hip")
cmake_minimum_required(VERSION 3.5.1)
set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -Wl,--no-keep-memory -Wl,-Bsymbolic -Wl,--unresolved-symbols=report-all -Wl,--version-script=${CMAKE_CURRENT_LIST_DIR}/hip_hcc.map.in")
if(CMAKE_CXX_FLAGS MATCHES "fsanitize=address")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -shared-libasan")
endif()
set (CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib)
set (CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/lib)
set(LIB_INSTALL_DIR ${CMAKE_INSTALL_PREFIX}/lib)
set(CONFIG_PACKAGE_INSTALL_DIR ${LIB_INSTALL_DIR}/cmake/hip)
find_package(PythonInterp REQUIRED)
add_definitions(-D__HIP_ROCclr__ -D__HIP_PLATFORM_HCC__ -DLINUX -D__x86_64__ -D__AMD64__ -DUNIX_OS -DqLittleEndian -DOPENCL_MAJOR=2 -DOPENCL_MINOR=0 -DCL_TARGET_OPENCL_VERSION=220 -DWITH_AQL -DWITH_ONLINE_COMPILER -DATI_OS_LINUX -DATI_ARCH_X86 -DLITTLEENDIAN_CPU -DATI_BITS_64 -DATI_COMP_GCC -DWITH_HSA_DEVICE -DWITH_TARGET_AMDGCN -DOPENCL_EXPORTS -DCL_USE_DEPRECATED_OPENCL_1_0_APIS -DCL_USE_DEPRECATED_OPENCL_1_1_APIS -DCL_USE_DEPRECATED_OPENCL_1_2_APIS -DCL_USE_DEPRECATED_OPENCL_2_0_APIS -DVEGA10_ONLY=false -DWITH_LIGHTNING_COMPILER -DUSE_PROF_API)
if(CMAKE_BUILD_TYPE MATCHES "^Debug$")
add_definitions(-DDEBUG)
endif()
if (CMAKE_CXX_COMPILER_ID STREQUAL "GNU")
add_compile_options("-Wno-ignored-attributes")
endif()
set(USE_PROF_API "1")
if(NOT DEFINED LIBROCclr_STATIC_DIR)
find_path(LIBROCclr_STATIC_DIR
NAMES libamdrocclr_static.a
PATHS /opt/rocm/rocclr
PATH_SUFFIXES lib
)
endif()
if(NOT DEFINED ROCclr_DIR)
find_path(ROCclr_DIR
NAMES top.hpp
PATH_SUFFIXES include
PATHS /opt/rocm/rocclr
)
endif()
message("Found Static rocclr lib:${LIBROCclr_STATIC_DIR} and rocclr includes: ${ROCclr_DIR}")
set(PROF_API_HEADER_PATH ${ROCclr_DIR}/platform)
#############################
# Profiling API support
#############################
# Generate profiling API macros/structures header
# FIXME: This should not be writing to the source directory
set(PROF_API_STR "${PROJECT_BINARY_DIR}/include/hip/hcc_detail/hip_prof_str.h")
set(PROF_API_HDR "${CMAKE_CURRENT_SOURCE_DIR}/../include/hip/hcc_detail/hip_runtime_api.h")
set(PROF_API_SRC "${CMAKE_CURRENT_SOURCE_DIR}")
set(PROF_API_GEN "${CMAKE_CURRENT_SOURCE_DIR}/hip_prof_gen.py")
set(PROF_API_LOG "${PROJECT_BINARY_DIR}/hip_prof_gen.log.txt")
add_custom_command(OUTPUT ${PROF_API_STR}
COMMAND ${PYTHON_EXECUTABLE} ${PROF_API_GEN} -v -t --priv ${OPT_PROF_API} ${PROF_API_HDR} ${PROF_API_SRC} ${PROF_API_STR}
OUTPUT_FILE ${PROF_API_LOG}
DEPENDS ${PROF_API_HDR} ${PROF_API_GEN}
COMMENT "Generating profiling primitives: ${PROF_API_STR}")
add_custom_target(gen-prof-api-str-header ALL
DEPENDS ${PROF_API_STR}
SOURCES ${PROF_API_HDR})
# Enable profiling API
if(USE_PROF_API EQUAL 1)
find_path(PROF_API_HEADER_DIR prof_protocol.h
HINTS
${PROF_API_HEADER_PATH}
PATHS
/opt/rocm/roctracer
PATH_SUFFIXES
include/ext
)
if(NOT PROF_API_HEADER_DIR)
MESSAGE(WARNING "Profiling API header not found. Disabling roctracer integration. Use -DPROF_API_HEADER_PATH=<path to prof_protocol.h header>")
else()
add_definitions(-DUSE_PROF_API=1)
include_directories(${PROF_API_HEADER_DIR})
MESSAGE(STATUS "Profiling API: ${PROF_API_HEADER_DIR}")
endif()
endif()
if(NOT DEFINED ROCclr_DIR OR NOT DEFINED LIBOCL_STATIC_DIR OR NOT DEFINED LIBROCclr_STATIC_DIR )
# message(FATAL_ERROR "define ROCclr_DIR, LIBOCL_STATIC_DIR\n")
endif()
list ( APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules" )
set(CMAKE_MODULE_PATH${CMAKE_MODULE_PATH} "${CMAKE_CURRENT_SOURCE_DIR}/cmake" "${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules")
include_directories(${ROCR_INCLUDES})
if (DEFINED LLVM_INCLUDES AND NOT ${LLVM_INCLUDES} STREQUAL "")
message(STATUS "LLVM includes found ${LLVM_INCLUDES}")
include_directories(${LLVM_INCLUDES})
endif() # if (DEFINED LLVM_INCLUDES AND NOT ${LLVM_INCLUDES} STREQUAL "")
include_directories(${CMAKE_SOURCE_DIR})
include_directories(${CMAKE_SOURCE_DIR}/include)
include_directories(${PROJECT_BINARY_DIR}/include)
include_directories(${CMAKE_SOURCE_DIR}/elfio)
include_directories(${CMAKE_SOURCE_DIR}/amdocl)
include_directories(${CMAKE_SOURCE_DIR}/include/hip/hcc_detail/elfio)
include_directories(${ROCclr_DIR})
include_directories(${ROCclr_DIR}/include)
include_directories(${ROCclr_DIR}/compiler/lib)
include_directories(${ROCclr_DIR}/compiler/lib/include)
include_directories(${ROCclr_DIR}/elf/utils/common)
include_directories(${ROCclr_DIR}/elf/utils/libelf)
add_definitions(-DUSE_COMGR_LIBRARY -DCOMGR_DYN_DLL)
find_package(amd_comgr REQUIRED CONFIG
PATHS
/opt/rocm/
PATH_SUFFIXES
cmake/amd_comgr
lib/cmake/amd_comgr
)
MESSAGE(STATUS "Code Object Manager found at ${amd_comgr_DIR}.")
include_directories("$<TARGET_PROPERTY:amd_comgr,INTERFACE_INCLUDE_DIRECTORIES>")
add_definitions(-DBSD_LIBELF)
add_library(hip64 OBJECT
hip_context.cpp
hip_device.cpp
hip_device_runtime.cpp
hip_error.cpp
hip_event.cpp
hip_memory.cpp
hip_module.cpp
hip_peer.cpp
hip_platform.cpp
hip_profile.cpp
hip_stream.cpp
hip_surface.cpp
hip_texture.cpp
hip_activity.cpp
hip_intercept.cpp
hip_rtc.cpp
cl_gl.cpp
cl_lqdflash_amd.cpp
fixme.cpp
)
set_target_properties(hip64 PROPERTIES POSITION_INDEPENDENT_CODE ON)
set_target_properties(
hip64 PROPERTIES
CXX_STANDARD 14
CXX_STANDARD_REQUIRED ON
CXX_EXTENSIONS OFF
)
add_dependencies(hip64 gen-prof-api-str-header)
set(THREADS_PREFER_PTHREAD_FLAG ON)
find_package(Threads REQUIRED)
include(${LIBROCclr_STATIC_DIR}/amdrocclr_staticTargets.cmake)
add_library(amdhip64 SHARED
$<TARGET_OBJECTS:hip64>
)
add_library(amdhip64_static STATIC
$<TARGET_OBJECTS:hip64>
)
set_target_properties(
amdhip64 PROPERTIES
VERSION ${HIP_LIB_VERSION_STRING}
SOVERSION ${HIP_LIB_VERSION_MAJOR}
)
set_target_properties(hip64 PROPERTIES PUBLIC_HEADER ${PROF_API_STR})
set_target_properties(amdhip64 PROPERTIES PUBLIC_HEADER ${PROF_API_STR})
set_target_properties(amdhip64_static PROPERTIES PUBLIC_HEADER ${PROF_API_STR})
# We expect amdhip64_static to contain objects of rocclr and hip. But linker
# let amdhip64_static contain objects of hip only. So we will use a
# a custom amdhip64_static_combiner to combine objects of vid and hip into
# amdhip64_static. To avoid amdhip64_static contains itself,
# amdhip64_static_temp is created internally.
add_library(amdhip64_static_temp STATIC
$<TARGET_OBJECTS:hip64>
)
add_library(host INTERFACE)
target_link_libraries(host INTERFACE amdhip64)
add_library(device INTERFACE)
target_link_libraries(device INTERFACE host)
# TODO: we may create host_static and device_static to let app
# link amdhip64_static
<<<<<<< HEAD:vdi/CMakeLists.txt
# FIXME: Linux convention is to create static library with same base
# filename.
target_link_libraries(amdhip64 PRIVATE amdvdi_static Threads::Threads dl)
target_link_libraries(amdhip64_static PRIVATE Threads::Threads dl)
target_link_libraries(amdhip64_static_temp PRIVATE Threads::Threads dl)
=======
target_link_libraries(amdhip64 PRIVATE amdrocclr_static pthread dl)
target_link_libraries(amdhip64_static PRIVATE pthread dl)
target_link_libraries(amdhip64_static_temp PRIVATE pthread dl)
>>>>>>> bf7c645f... name change vdi to rocclr:rocclr/CMakeLists.txt
# combine objects of vid and hip into amdhip64_static
add_custom_target(
amdhip64_static_combiner
ALL
COMMAND rm -f $<TARGET_FILE:amdhip64_static> # Must remove old one, otherwise the new one will contain obsolete stuff
COMMAND ${CMAKE_AR} -rcsT $<TARGET_FILE:amdhip64_static> $<TARGET_FILE:amdhip64_static_temp> $<TARGET_FILE:amdrocclr_static>
DEPENDS amdhip64_static amdhip64_static_temp amdrocclr_static # To make sure this is the last step
COMMENT "Combining static libs into amdhip64_static"
)
INSTALL(PROGRAMS $<TARGET_FILE:amdhip64_static> DESTINATION lib COMPONENT MAIN)
INSTALL(PROGRAMS $<TARGET_FILE:amdhip64> DESTINATION lib COMPONENT MAIN)
INSTALL(CODE "execute_process( COMMAND ${CMAKE_COMMAND} -E create_symlink libamdhip64.so lib/libhip_hcc.so )" DESTINATION lib COMPONENT MAIN)
INSTALL(CODE "execute_process( COMMAND ${CMAKE_COMMAND} -E create_symlink libamdhip64.so lib/libhiprtc.so )" DESTINATION lib COMPONENT MAIN)
INSTALL(FILES ${CMAKE_BINARY_DIR}/lib/libhip_hcc.so DESTINATION lib COMPONENT MAIN)
INSTALL(FILES ${CMAKE_BINARY_DIR}/lib/libhiprtc.so DESTINATION lib COMPONENT MAIN)
INSTALL(TARGETS amdhip64 amdhip64_static host device EXPORT hip-targets DESTINATION ${LIB_INSTALL_DIR})
INSTALL(EXPORT hip-targets DESTINATION ${CONFIG_PACKAGE_INSTALL_DIR} NAMESPACE hip::)
rocclr/cl_gl.cpp
+2432
View File
File diff suppressed because it is too large Load Diff
rocclr/cl_gl_amd.hpp
+379
View File
@@ -0,0 +1,379 @@
/* Copyright (c) 2010-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. */
#ifndef CL_GL_AMD_HPP_
#define CL_GL_AMD_HPP_
#ifdef _WIN32
#include <windows.h>
#else //!_WIN32
#include <dlfcn.h>
#endif //!_WIN32
#include <GL/gl.h>
#include <GL/glext.h>
#include "CL/cl_gl.h"
#ifndef _WIN32
#include <GL/glx.h>
#endif //!_WIN32
#include <EGL/egl.h>
#include <EGL/eglext.h>
#include <EGL/eglplatform.h>
#include "platform/context.hpp"
#include "platform/command.hpp"
namespace amd
{
//! Class GLObject keeps all the info about the GL object
//! from which the CL object is created
class GLObject : public InteropObject
{
protected:
cl_gl_object_type clGLType_; //!< CL GL object type
GLenum glTarget_;
GLuint gluiName_;
GLint gliMipLevel_;
GLenum glInternalFormat_;
GLint gliWidth_;
GLint gliHeight_;
GLint gliDepth_;
GLenum glCubemapFace_;
GLsizei glNumSamples_;
public:
//! GLObject constructor initializes member variables
GLObject(
GLenum glTarget,
GLuint gluiName,
GLint gliMipLevel,
GLenum glInternalFormat,
GLint gliWidth,
GLint gliHeight,
GLint gliDepth,
cl_gl_object_type clGLType,
GLenum glCubemapFace,
GLsizei glNumSamples
): // Initialization of member variables
clGLType_(clGLType),
glTarget_(glTarget),
gluiName_(gluiName),
gliMipLevel_(gliMipLevel),
glInternalFormat_(glInternalFormat),
gliWidth_(gliWidth),
gliHeight_(gliHeight),
gliDepth_(gliDepth),
glCubemapFace_(glCubemapFace),
glNumSamples_(glNumSamples)
{
}
virtual ~GLObject() {}
virtual GLObject* asGLObject() {return this;}
//! GLObject query functions to get GL info from member variables
GLenum getGLTarget() const {return glTarget_;}
GLuint getGLName() const {return gluiName_;}
GLint getGLMipLevel() const {return gliMipLevel_;}
GLenum getGLInternalFormat() const {return glInternalFormat_;}
GLint getGLSize() const {return gliWidth_;}
GLint getGLWidth() const {return gliWidth_;}
GLint getGLHeight() const {return gliHeight_;}
GLint getGLDepth() const {return gliDepth_;}
cl_gl_object_type getCLGLObjectType() const { return clGLType_; }
GLenum getCubemapFace() const {return glCubemapFace_;}
GLsizei getNumSamples() const { return glNumSamples_;}
};
//! Class BufferGL is drived from classes Buffer and GLObject
//! where the former keeps all data for CL object and
//! the latter keeps all data for GL object
class BufferGL : public Buffer, public GLObject
{
protected:
//! Initializes the device memory array which is nested
// after'BufferGL' object in memory layout.
virtual void initDeviceMemory();
public:
//! BufferGL constructor just calls constructors of base classes
//! to pass down the parameters
BufferGL(
Context& amdContext,
cl_mem_flags clFlags,
size_t uiSizeInBytes,
GLenum glTarget,
GLuint gluiName)
: // Call base classes constructors
Buffer(
amdContext,
clFlags,
uiSizeInBytes
),
GLObject(
glTarget,
gluiName,
0, // Mipmap level default
GL_ARRAY_BUFFER, // Just init to some value
(GLint) uiSizeInBytes,
1,
1,
CL_GL_OBJECT_BUFFER,
0,
0
)
{
setInteropObj(this);
}
virtual ~BufferGL() {}
virtual BufferGL* asBufferGL() { return this; }
};
//! Class ImageGL is derived from classes Image and GLObject
//! where the former keeps all data for CL object and
//! the latter keeps all data for GL object
class ImageGL : public Image, public GLObject
{
public:
//! ImageGL constructor just calls constructors of base classes
//! to pass down the parameters
ImageGL(
Context& amdContext,
cl_mem_object_type clType,
cl_mem_flags clFlags,
const Format& format,
size_t width,
size_t height,
size_t depth,
GLenum glTarget,
GLuint gluiName,
GLint gliMipLevel,
GLenum glInternalFormat,
cl_gl_object_type clGLType,
GLsizei numSamples,
GLenum glCubemapFace = 0)
: Image(amdContext, clType, clFlags, format, width, height, depth,
Format(format).getElementSize() * width,
Format(format).getElementSize() * width * depth)
, GLObject(glTarget, gluiName, gliMipLevel, glInternalFormat,
static_cast<GLint>(width), static_cast<GLint>(height),
static_cast<GLint>(depth), clGLType, glCubemapFace,numSamples)
{
setInteropObj(this);
}
virtual ~ImageGL() {}
protected:
//! Initializes the device memory array which is nested
// after'BufferGL' object in memory layout.
virtual void initDeviceMemory();
};
#ifdef _WIN32
#define APICALL WINAPI
#define GETPROCADDRESS GetProcAddress
#define API_GETPROCADDR "wglGetProcAddress"
#define FCN_STR_TYPE LPCSTR
typedef PROC (WINAPI* PFN_xxxGetProcAddress) (LPCSTR fcnName);
typedef HGLRC (APICALL* PFN_wglCreateContext) (HDC hdc);
typedef HGLRC (APICALL* PFN_wglGetCurrentContext) (void);
typedef HDC (APICALL* PFN_wglGetCurrentDC) (void);
typedef BOOL (APICALL* PFN_wglDeleteContext) (HGLRC hglrc);
typedef BOOL (APICALL* PFN_wglMakeCurrent) (HDC hdc, HGLRC hglrc);
typedef BOOL (APICALL* PFN_wglShareLists) (HGLRC hglrc1, HGLRC hglrc2);
#else //!_WIN32
#define APICALL // __stdcall //??? todo odintsov
#define API_GETPROCADDR "glXGetProcAddress"
#define GETPROCADDRESS dlsym
#define FCN_STR_TYPE const GLubyte*
#define WINAPI
#define PROC void*
typedef void* (*PFN_xxxGetProcAddress) (const GLubyte* procName);
// X11 typedef
typedef Display* (*PFNXOpenDisplay)(_Xconst char* display_name );
typedef int (*PFNXCloseDisplay)(Display* display );
//glx typedefs
typedef GLXDrawable (*PFNglXGetCurrentDrawable)();
typedef Display* (*PFNglXGetCurrentDisplay)();
typedef GLXContext (*PFNglXGetCurrentContext)( void );
typedef XVisualInfo* (*PFNglXChooseVisual)(Display *dpy, int screen, int *attribList);
typedef GLXContext(*PFNglXCreateContext)(Display* dpy,XVisualInfo* vis,GLXContext shareList,Bool direct);
typedef void(*PFNglXDestroyContext)(Display* dpy, GLXContext ctx);
typedef Bool(*PFNglXMakeCurrent)( Display* dpy, GLXDrawable drawable, GLXContext ctx);
typedef void* HMODULE;
#endif //!_WIN32
#define GLPREFIX(rtype, fcn, dclargs) \
typedef rtype (APICALL* PFN_##fcn) dclargs;
// Declare prototypes for GL functions
#include "gl_functions.hpp"
class GLFunctions
{
public:
//! Locks any access to the virtual GPUs
class SetIntEnv : public amd::StackObject {
public:
//! Default constructor
SetIntEnv(GLFunctions* env);
//! Destructor
~SetIntEnv();
//! Checks if the environment setup was successful
bool isValid() const { return isValid_; }
private:
GLFunctions* env_; //!< GL environment
bool isValid_; //!< If TRUE, then it's a valid setup
};
private:
HMODULE libHandle_;
int missed_; // Indicates how many GL functions not init'ed, if any
amd::Monitor lock_;
EGLDisplay eglDisplay_;
EGLContext eglOriginalContext_;
EGLContext eglInternalContext_;
EGLContext eglTempContext_;
bool isEGL_;
#ifdef _WIN32
HGLRC hOrigGLRC_;
HDC hDC_;
HGLRC hIntGLRC_; // handle for internal GLRC to access shared context
HDC tempDC_;
HGLRC tempGLRC_;
PFN_wglCreateContext wglCreateContext_;
PFN_wglGetCurrentContext wglGetCurrentContext_;
PFN_wglGetCurrentDC wglGetCurrentDC_;
PFN_wglDeleteContext wglDeleteContext_;
PFN_wglMakeCurrent wglMakeCurrent_;
PFN_wglShareLists wglShareLists_;
#else
public:
Display* Dpy_;
GLXDrawable Drawable_;
GLXContext origCtx_;
Display* intDpy_;
Window intDrawable_;
GLXContext intCtx_;
Display* tempDpy_;
GLXDrawable tempDrawable_;
GLXContext tempCtx_;
//pointers to X11 functions
PFNXOpenDisplay XOpenDisplay_;
PFNXCloseDisplay XCloseDisplay_;
//pointers to GLX functions
PFNglXGetCurrentDrawable glXGetCurrentDrawable_;
PFNglXGetCurrentDisplay glXGetCurrentDisplay_;
PFNglXGetCurrentContext glXGetCurrentContext_;
PFNglXChooseVisual glXChooseVisual_;
PFNglXCreateContext glXCreateContext_;
PFNglXDestroyContext glXDestroyContext_;
PFNglXMakeCurrent glXMakeCurrent_;
#endif
public:
GLFunctions(HMODULE h, bool isEGL);
~GLFunctions();
// Query CL-GL context association
bool isAssociated() const
{
if (isEGL_ && eglDisplay_ && eglOriginalContext_) return true;
#ifdef _WIN32
if(hDC_ && hOrigGLRC_) return true;
#else //!_WIN32
if(Dpy_ && origCtx_) return true;
#endif //!_WIN32
return false;
}
bool isEGL() const
{
return isEGL_;
}
// Accessor methods
#ifdef _WIN32
HGLRC getOrigGLRC() const {return hOrigGLRC_;}
HDC getDC() const {return hDC_;}
HGLRC getIntGLRC() const {return hIntGLRC_;}
#else //!_WIN32
Display* getDpy() const {return Dpy_;}
GLXDrawable getDrawable() const {return Drawable_;}
GLXContext getOrigCtx() const {return origCtx_;}
Display* getIntDpy() const {return intDpy_;}
GLXDrawable getIntDrawable() const {return intDrawable_;}
GLXContext getIntCtx() const {return intCtx_;}
EGLDisplay getEglDpy() const { return eglDisplay_; }
EGLContext getEglOrigCtx() const { return eglOriginalContext_; }
#endif //!_WIN32
// Initialize GL dynamic library and function pointers
bool init(intptr_t hdc, intptr_t hglrc);
// Return true if successful, false - if error occurred
bool setIntEnv();
bool restoreEnv();
amd::Monitor& getLock() { return lock_; }
PFN_xxxGetProcAddress GetProcAddress_;
#define GLPREFIX(rtype, fcn, dclargs) \
PFN_##fcn fcn##_;
// Declare pointers to GL functions
#include "gl_functions.hpp"
};
//! Functions for executing the GL related stuff
cl_mem clCreateFromGLBufferAMD(Context& amdContext, cl_mem_flags flags,
GLuint bufobj, cl_int* errcode_ret);
cl_mem clCreateFromGLTextureAMD(Context& amdContext, cl_mem_flags flags,
GLenum target, GLint miplevel, GLuint texture, int* errcode_ret);
cl_mem clCreateFromGLRenderbufferAMD(Context& amdContext, cl_mem_flags flags,
GLuint renderbuffer, int* errcode_ret);
bool
getCLFormatFromGL(
const Context& amdContext,
GLint gliInternalFormat,
cl_image_format* pclImageFormat,
int* piBytesPerPixel,
cl_mem_flags flags
);
} //namespace amd
#endif //CL_GL_AMD_HPP_
rocclr/cl_lqdflash_amd.cpp
+310
View File
@@ -0,0 +1,310 @@
/* Copyright (c) 2015-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 "cl_common.hpp"
#include <CL/cl_ext.h>
#include "platform/object.hpp"
#include "cl_lqdflash_amd.h"
#if (!defined(BUILD_HSA_TARGET) && defined(WITH_HSA_DEVICE) && \
defined(WITH_AMDGPU_PRO)) || defined(_WIN32) || defined(WITH_PAL_DEVICE)
#define WITH_LIQUID_FLASH 1
#endif // _WIN32
#if defined(WITH_LIQUID_FLASH)
#include "lf.h"
#include <locale>
#include <codecvt>
#endif // WITH_LIQUID_FLASH
namespace amd {
LiquidFlashFile::~LiquidFlashFile() { close(); }
bool LiquidFlashFile::open() {
#if defined WITH_LIQUID_FLASH
lf_status err;
lf_file_flags flags = 0;
switch (flags_) {
case CL_FILE_READ_ONLY_AMD:
flags = LF_READ;
break;
case CL_FILE_WRITE_ONLY_AMD:
flags = LF_WRITE;
break;
case CL_FILE_READ_WRITE_AMD:
flags = LF_READ | LF_WRITE;
break;
}
#ifdef ATI_OS_LINUX
assert(sizeof(wchar_t) != sizeof(lf_char));
std::string name_char;
std::wstring_convert<std::codecvt_utf8<wchar_t>, wchar_t> cv;
name_char = cv.to_bytes(name_);
handle_ = lfOpenFile(name_char.c_str(), flags, &err);
#else
handle_ = lfOpenFile(name_.c_str(), flags, &err);
#endif
if (err != lf_success) {
return false;
}
if (lfGetFileBlockSize((lf_file)handle_, &blockSize_) != lf_success) {
return false;
}
if (lfGetFileSize((lf_file)handle_, &fileSize_) != lf_success) {
return false;
}
return true;
#else
return false;
#endif // WITH_LIQUID_FLASH
}
void LiquidFlashFile::close() {
#if defined WITH_LIQUID_FLASH
if (handle_ != NULL) {
lfReleaseFile((lf_file)handle_);
handle_ = NULL;
}
#endif // WITH_LIQUID_FLASH
}
bool LiquidFlashFile::transferBlock(bool writeBuffer, void* srcDst, uint64_t bufferSize,
uint64_t fileOffset, uint64_t bufferOffset,
uint64_t size) const {
#if defined WITH_LIQUID_FLASH
lf_status status;
lf_region_descriptor region = {fileOffset / blockSize(), bufferOffset / blockSize(),
size / blockSize()};
if (writeBuffer) {
status = lfReadFile(srcDst, bufferSize, (lf_file)handle_, 1, &region, NULL);
} else {
status = lfWriteFile(srcDst, bufferSize, (lf_file)handle_, 1, &region, NULL);
}
if (lf_success == status) {
return true;
} else {
return false;
}
#else
return false;
#endif // WITH_LIQUID_FLASH
}
} // namespace amd
/*! \addtogroup API
* @{
*
* \addtogroup AMD_Extensions
* @{
*
*/
RUNTIME_ENTRY_RET(cl_file_amd, clCreateSsgFileObjectAMD,
(cl_context context, cl_file_flags_amd flags, const wchar_t* file_name,
cl_int* errcode_ret)) {
#if defined WITH_LIQUID_FLASH && defined ATI_OS_LINUX
if (!is_valid(context)) {
*not_null(errcode_ret) = CL_INVALID_CONTEXT;
LogWarning("invalid parameter \"context\"");
return (cl_file_amd)0;
}
const std::vector<amd::Device*>& devices = as_amd(context)->devices();
bool supportPass = false;
for (auto& dev : devices) {
if (lf_success == lfCheckExtensionSupportForDevice(dev->info().pcieDeviceId_,
dev->info().pcieRevisionId_)) {
supportPass = true;
break;
}
}
if (!supportPass) {
*not_null(errcode_ret) = CL_INVALID_DEVICE;
LogWarning("SSG isn't supported");
return (cl_file_amd)0;
}
#endif
amd::LiquidFlashFile* file = new amd::LiquidFlashFile(file_name, flags);
if (file == NULL) {
*not_null(errcode_ret) = CL_OUT_OF_HOST_MEMORY;
return (cl_file_amd)0;
}
if (!file->open()) {
*not_null(errcode_ret) = CL_INVALID_VALUE;
delete file;
return (cl_file_amd)0;
}
*not_null(errcode_ret) = CL_SUCCESS;
return as_cl(file);
}
RUNTIME_EXIT
RUNTIME_ENTRY(cl_int, clGetSsgFileObjectInfoAMD,
(cl_file_amd file, cl_file_info_amd param_name, size_t param_value_size,
void* param_value, size_t* param_value_size_ret)) {
if (!is_valid(file)) {
return CL_INVALID_FILE_OBJECT_AMD;
}
switch (param_name) {
case CL_FILE_BLOCK_SIZE_AMD: {
cl_uint blockSize = as_amd(file)->blockSize();
return amd::clGetInfo(blockSize, param_value_size, param_value, param_value_size_ret);
}
case CL_FILE_SIZE_AMD: {
cl_ulong fileSize = as_amd(file)->fileSize();
return amd::clGetInfo(fileSize, param_value_size, param_value, param_value_size_ret);
}
default:
break;
}
return CL_INVALID_VALUE;
}
RUNTIME_EXIT
RUNTIME_ENTRY(cl_int, clRetainSsgFileObjectAMD, (cl_file_amd file)) {
if (!is_valid(file)) {
return CL_INVALID_FILE_OBJECT_AMD;
}
as_amd(file)->retain();
return CL_SUCCESS;
}
RUNTIME_EXIT
RUNTIME_ENTRY(cl_int, clReleaseSsgFileObjectAMD, (cl_file_amd file)) {
if (!is_valid(file)) {
return CL_INVALID_FILE_OBJECT_AMD;
}
as_amd(file)->release();
return CL_SUCCESS;
}
RUNTIME_EXIT
static cl_int EnqueueTransferBufferFromSsgFileAMD(
cl_bool isWrite, cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write,
size_t buffer_offset, size_t cb, cl_file_amd file, size_t file_offset,
cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) {
if (!is_valid(command_queue)) {
return CL_INVALID_COMMAND_QUEUE;
}
if (!is_valid(buffer)) {
return CL_INVALID_MEM_OBJECT;
}
amd::Buffer* pBuffer = as_amd(buffer)->asBuffer();
if (pBuffer == NULL) {
return CL_INVALID_MEM_OBJECT;
}
if (pBuffer->getMemFlags() & (CL_MEM_HOST_READ_ONLY | CL_MEM_HOST_NO_ACCESS)) {
return CL_INVALID_OPERATION;
}
amd::HostQueue* queue = as_amd(command_queue)->asHostQueue();
if (NULL == queue) {
return CL_INVALID_COMMAND_QUEUE;
}
amd::HostQueue& hostQueue = *queue;
if (hostQueue.context() != pBuffer->getContext()) {
return CL_INVALID_CONTEXT;
}
if (!is_valid(file)) {
return CL_INVALID_FILE_OBJECT_AMD;
}
amd::LiquidFlashFile* amdFile = as_amd(file);
amd::Coord3D bufferOffset(buffer_offset, 0, 0);
amd::Coord3D bufferSize(cb, 1, 1);
if ((!pBuffer->validateRegion(bufferOffset, bufferSize)) ||
// LF library supports aligned sizes only
((buffer_offset % amdFile->blockSize()) != 0) || ((cb % amdFile->blockSize()) != 0) ||
((file_offset % amdFile->blockSize()) != 0)) {
return CL_INVALID_VALUE;
}
amd::Command::EventWaitList eventWaitList;
cl_int err = amd::clSetEventWaitList(eventWaitList, hostQueue, num_events_in_wait_list,
event_wait_list);
if (err != CL_SUCCESS) {
return err;
}
amd::TransferBufferFileCommand* command;
command = new amd::TransferBufferFileCommand(
isWrite ? CL_COMMAND_READ_SSG_FILE_AMD : CL_COMMAND_WRITE_SSG_FILE_AMD, hostQueue,
eventWaitList, *pBuffer, bufferOffset, bufferSize, amdFile, file_offset);
if (command == NULL) {
return CL_OUT_OF_HOST_MEMORY;
}
// Make sure we have memory for the command execution
if (!command->validateMemory()) {
delete command;
return CL_MEM_OBJECT_ALLOCATION_FAILURE;
}
command->enqueue();
if (blocking_write) {
command->awaitCompletion();
}
*not_null(event) = as_cl(&command->event());
if (event == NULL) {
command->release();
}
return CL_SUCCESS;
}
RUNTIME_ENTRY(cl_int, clEnqueueReadSsgFileAMD,
(cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write,
size_t buffer_offset, size_t cb, cl_file_amd file, size_t file_offset,
cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event)) {
return EnqueueTransferBufferFromSsgFileAMD(CL_TRUE, command_queue, buffer, blocking_write,
buffer_offset, cb, file, file_offset,
num_events_in_wait_list, event_wait_list, event);
}
RUNTIME_EXIT
RUNTIME_ENTRY(cl_int, clEnqueueWriteSsgFileAMD,
(cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write,
size_t buffer_offset, size_t cb, cl_file_amd file, size_t file_offset,
cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event)) {
return EnqueueTransferBufferFromSsgFileAMD(CL_FALSE, command_queue, buffer, blocking_write,
buffer_offset, cb, file, file_offset,
num_events_in_wait_list, event_wait_list, event);
}
RUNTIME_EXIT
rocclr/cl_lqdflash_amd.h
+58
View File
@@ -0,0 +1,58 @@
/* Copyright (c) 2015-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. */
#ifndef __CL_LQDFLASH_AMD_H
#define __CL_LQDFLASH_AMD_H
#include "CL/cl_ext.h"
#ifdef __cplusplus
extern "C" {
#endif /*__cplusplus*/
extern CL_API_ENTRY cl_file_amd CL_API_CALL
clCreateSsgFileObjectAMD(cl_context context, cl_file_flags_amd flags, const wchar_t* file_name,
cl_int* errcode_ret) CL_EXT_SUFFIX__VERSION_1_2;
extern CL_API_ENTRY cl_int CL_API_CALL clGetSsgFileObjectInfoAMD(
cl_file_amd file, cl_file_info_amd param_name, size_t param_value_size, void* param_value,
size_t* param_value_size_ret) CL_EXT_SUFFIX__VERSION_1_2;
extern CL_API_ENTRY cl_int CL_API_CALL clRetainSsgFileObjectAMD(cl_file_amd file)
CL_EXT_SUFFIX__VERSION_1_2;
extern CL_API_ENTRY cl_int CL_API_CALL clReleaseSsgFileObjectAMD(cl_file_amd file)
CL_EXT_SUFFIX__VERSION_1_2;
extern CL_API_ENTRY cl_int CL_API_CALL clEnqueueReadSsgFileAMD(
cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write, size_t buffer_offset,
size_t cb, cl_file_amd file, size_t file_offset, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) CL_EXT_SUFFIX__VERSION_1_2;
extern CL_API_ENTRY cl_int CL_API_CALL clEnqueueWriteSsgFileAMD(
cl_command_queue command_queue, cl_mem buffer, cl_bool blocking_write, size_t buffer_offset,
size_t cb, cl_file_amd file, size_t file_offset, cl_uint num_events_in_wait_list,
const cl_event* event_wait_list, cl_event* event) CL_EXT_SUFFIX__VERSION_1_2;
#ifdef __cplusplus
} /*extern "C"*/
#endif /*__cplusplus*/
#endif
rocclr/fixme.cpp
+32
View File
@@ -0,0 +1,32 @@
/* Copyright (c) 2015-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 "vdi_common.hpp"
#include <icd/loader/icd_dispatch.h>
cl_icd_dispatch amd::ICDDispatchedObject::icdVendorDispatch_[] = {0};
amd::PlatformIDS amd::PlatformID::Platform = {amd::ICDDispatchedObject::icdVendorDispatch_};
RUNTIME_ENTRY(cl_int, clGetDeviceIDs,
(cl_platform_id platform, cl_device_type device_type, cl_uint num_entries,
cl_device_id* devices, cl_uint* num_devices)) {
return CL_SUCCESS;
}
RUNTIME_EXIT
rocclr/hip_activity.cpp
+35
View File
@@ -0,0 +1,35 @@
/* Copyright (c) 2015-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 "platform/activity.hpp"
extern "C" void hipInitActivityCallback(void* id_callback, void* op_callback, void* arg) {
activity_prof::CallbacksTable::init(reinterpret_cast<activity_prof::id_callback_fun_t>(id_callback),
reinterpret_cast<activity_prof::callback_fun_t>(op_callback),
arg);
}
extern "C" bool hipEnableActivityCallback(unsigned op, bool enable) {
return activity_prof::CallbacksTable::SetEnabled(op, enable);
}
extern "C" const char* hipGetCmdName(unsigned op) {
return getOclCommandKindString(static_cast<uint32_t>(op));
}
rocclr/hip_context.cpp
Executable
+372
View File
@@ -0,0 +1,372 @@
/* Copyright (c) 2015-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 "platform/runtime.hpp"
#include "utils/flags.hpp"
#include "utils/versions.hpp"
std::vector<hip::Device*> g_devices;
namespace hip {
thread_local Device* g_device = nullptr;
thread_local std::stack<Device*> g_ctxtStack;
thread_local hipError_t g_lastError = hipSuccess;
std::once_flag g_ihipInitialized;
Device* host_device = nullptr;
void init() {
if (!amd::Runtime::initialized()) {
amd::IS_HIP = true;
GPU_NUM_MEM_DEPENDENCY = 0;
amd::Runtime::init();
}
const std::vector<amd::Device*>& devices = amd::Device::getDevices(CL_DEVICE_TYPE_GPU, false);
for (unsigned int i=0; i<devices.size(); i++) {
const std::vector<amd::Device*> device(1, devices[i]);
amd::Context* context = new amd::Context(device, amd::Context::Info());
if (!context) return;
// Enable active wait on the device by default
devices[i]->SetActiveWait(true);
if (context && CL_SUCCESS != context->create(nullptr)) {
context->release();
} else {
g_devices.push_back(new Device(context, i));
}
}
amd::Context* hContext = new amd::Context(devices, amd::Context::Info());
if (!hContext) return;
if (CL_SUCCESS != hContext->create(nullptr)) {
hContext->release();
}
host_device = new Device(hContext, -1);
PlatformState::instance().init();
}
Device* getCurrentDevice() {
return g_device;
}
void setCurrentDevice(unsigned int index) {
assert(index<g_devices.size());
g_device = g_devices[index];
}
amd::HostQueue* getQueue(hipStream_t stream) {
if (stream == nullptr) {
return getNullStream();
} else {
constexpr bool WaitNullStreamOnly = true;
amd::HostQueue* queue = reinterpret_cast<hip::Stream*>(stream)->asHostQueue();
iHipWaitActiveStreams(queue, WaitNullStreamOnly);
return queue;
}
}
amd::HostQueue* getNullStream(amd::Context& ctx) {
for (auto& it : g_devices) {
if (it->asContext() == &ctx) {
return it->NullStream();
}
}
return nullptr;
}
amd::HostQueue* getNullStream() {
Device* device = getCurrentDevice();
return device ? device->NullStream() : nullptr;
}
};
using namespace hip;
hipError_t hipInit(unsigned int flags) {
HIP_INIT_API(hipInit, flags);
HIP_RETURN(hipSuccess);
}
hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device) {
HIP_INIT_API(hipCtxCreate, ctx, flags, device);
if (static_cast<size_t>(device) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidValue);
}
*ctx = reinterpret_cast<hipCtx_t>(g_devices[device]);
// Increment ref count for device primary context
g_devices[device]->retain();
g_ctxtStack.push(g_devices[device]);
HIP_RETURN(hipSuccess);
}
hipError_t hipCtxSetCurrent(hipCtx_t ctx) {
HIP_INIT_API(hipCtxSetCurrent, ctx);
if (ctx == nullptr) {
if(!g_ctxtStack.empty()) {
g_ctxtStack.pop();
}
} else {
hip::g_device = reinterpret_cast<hip::Device*>(ctx);
if(!g_ctxtStack.empty()) {
g_ctxtStack.pop();
}
g_ctxtStack.push(hip::getCurrentDevice());
}
HIP_RETURN(hipSuccess);
}
hipError_t hipCtxGetCurrent(hipCtx_t* ctx) {
HIP_INIT_API(hipCtxGetCurrent, ctx);
*ctx = reinterpret_cast<hipCtx_t>(hip::getCurrentDevice());
HIP_RETURN(hipSuccess);
}
hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) {
HIP_INIT_API(hipCtxGetSharedMemConfig, pConfig);
*pConfig = hipSharedMemBankSizeFourByte;
HIP_RETURN(hipSuccess);
}
hipError_t hipRuntimeGetVersion(int *runtimeVersion) {
HIP_INIT_API(hipRuntimeGetVersion, runtimeVersion);
if (!runtimeVersion) {
HIP_RETURN(hipErrorInvalidValue);
}
*runtimeVersion = AMD_PLATFORM_BUILD_NUMBER;
HIP_RETURN(hipSuccess);
}
hipError_t hipCtxDestroy(hipCtx_t ctx) {
HIP_INIT_API(hipCtxDestroy, ctx);
hip::Device* dev = reinterpret_cast<hip::Device*>(ctx);
if (dev == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
// Release last tracked command
hip::getNullStream()->setLastQueuedCommand(nullptr);
// Need to remove the ctx of calling thread if its the top one
if (!g_ctxtStack.empty() && g_ctxtStack.top() == dev) {
g_ctxtStack.pop();
}
// Remove context from global context list
for (unsigned int i = 0; i < g_devices.size(); i++) {
if (g_devices[i] == dev) {
// Decrement ref count for device primary context
dev->release();
}
}
HIP_RETURN(hipSuccess);
}
hipError_t hipCtxPopCurrent(hipCtx_t* ctx) {
HIP_INIT_API(hipCtxPopCurrent, ctx);
hip::Device** dev = reinterpret_cast<hip::Device**>(ctx);
if (dev == nullptr) {
HIP_RETURN(hipErrorInvalidContext);
}
if (!g_ctxtStack.empty()) {
*dev = g_ctxtStack.top();
g_ctxtStack.pop();
} else {
DevLogError("Context Stack empty \n");
HIP_RETURN(hipErrorInvalidContext);
}
HIP_RETURN(hipSuccess);
}
hipError_t hipCtxPushCurrent(hipCtx_t ctx) {
HIP_INIT_API(hipCtxPushCurrent, ctx);
hip::Device* dev = reinterpret_cast<hip::Device*>(ctx);
if (dev == nullptr) {
HIP_RETURN(hipErrorInvalidContext);
}
hip::g_device = dev;
g_ctxtStack.push(hip::getCurrentDevice());
HIP_RETURN(hipSuccess);
}
hipError_t hipDriverGetVersion(int* driverVersion) {
HIP_INIT_API(hipDriverGetVersion, driverVersion);
auto* deviceHandle = g_devices[0]->devices()[0];
const auto& info = deviceHandle->info();
if (driverVersion) {
*driverVersion = AMD_PLATFORM_BUILD_NUMBER * 100 +
AMD_PLATFORM_REVISION_NUMBER;
} else {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(hipSuccess);
}
hipError_t hipCtxGetDevice(hipDevice_t* device) {
HIP_INIT_API(hipCtxGetDevice, device);
if (device != nullptr) {
*device = hip::getCurrentDevice()->deviceId();
HIP_RETURN(hipSuccess);
} else {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(hipErrorInvalidContext);
}
hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) {
HIP_INIT_API(hipCtxGetApiVersion, apiVersion);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) {
HIP_INIT_API(hipCtxGetCacheConfig, cacheConfig);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) {
HIP_INIT_API(hipCtxSetCacheConfig, cacheConfig);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) {
HIP_INIT_API(hipCtxSetSharedMemConfig, config);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipCtxSynchronize(void) {
HIP_INIT_API(hipCtxSynchronize, 1);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipCtxGetFlags(unsigned int* flags) {
HIP_INIT_API(hipCtxGetFlags, flags);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int* flags, int* active) {
HIP_INIT_API(hipDevicePrimaryCtxGetState, dev, flags, active);
if (static_cast<unsigned int>(dev) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
}
if (flags != nullptr) {
*flags = 0;
}
if (active != nullptr) {
*active = (g_devices[dev] == hip::getCurrentDevice())? 1 : 0;
}
HIP_RETURN(hipSuccess);
}
hipError_t hipDevicePrimaryCtxRelease(hipDevice_t dev) {
HIP_INIT_API(hipDevicePrimaryCtxRelease, dev);
if (static_cast<unsigned int>(dev) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
}
HIP_RETURN(hipSuccess);
}
hipError_t hipDevicePrimaryCtxRetain(hipCtx_t* pctx, hipDevice_t dev) {
HIP_INIT_API(hipDevicePrimaryCtxRetain, pctx, dev);
if (static_cast<unsigned int>(dev) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
}
if (pctx == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
*pctx = reinterpret_cast<hipCtx_t>(g_devices[dev]);
HIP_RETURN(hipSuccess);
}
hipError_t hipDevicePrimaryCtxReset(hipDevice_t dev) {
HIP_INIT_API(hipDevicePrimaryCtxReset, dev);
HIP_RETURN(hipSuccess);
}
hipError_t hipDevicePrimaryCtxSetFlags(hipDevice_t dev, unsigned int flags) {
HIP_INIT_API(hipDevicePrimaryCtxSetFlags, dev, flags);
if (static_cast<unsigned int>(dev) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
} else {
HIP_RETURN(hipErrorContextAlreadyInUse);
}
}
rocclr/hip_conversions.hpp
+903
View File
@@ -0,0 +1,903 @@
/*
Copyright (c) 2015 - 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 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.
*/
#pragma once
#include <hip/hcc_detail/driver_types.h>
#include <hip/hcc_detail/texture_types.h>
namespace hip
{
inline
cl_channel_type getCLChannelType(const hipArray_Format hipFormat,
const hipTextureReadMode hipReadMode) {
if (hipReadMode == hipReadModeElementType) {
switch (hipFormat) {
case HIP_AD_FORMAT_UNSIGNED_INT8:
return CL_UNSIGNED_INT8;
case HIP_AD_FORMAT_SIGNED_INT8:
return CL_SIGNED_INT8;
case HIP_AD_FORMAT_UNSIGNED_INT16:
return CL_UNSIGNED_INT16;
case HIP_AD_FORMAT_SIGNED_INT16:
return CL_SIGNED_INT16;
case HIP_AD_FORMAT_UNSIGNED_INT32:
return CL_UNSIGNED_INT32;
case HIP_AD_FORMAT_SIGNED_INT32:
return CL_SIGNED_INT32;
case HIP_AD_FORMAT_HALF:
return CL_HALF_FLOAT;
case HIP_AD_FORMAT_FLOAT:
return CL_FLOAT;
}
} else if (hipReadMode == hipReadModeNormalizedFloat) {
switch (hipFormat) {
case HIP_AD_FORMAT_UNSIGNED_INT8:
return CL_UNORM_INT8;
case HIP_AD_FORMAT_SIGNED_INT8:
return CL_SNORM_INT8;
case HIP_AD_FORMAT_UNSIGNED_INT16:
return CL_UNORM_INT16;
case HIP_AD_FORMAT_SIGNED_INT16:
return CL_SNORM_INT16;
case HIP_AD_FORMAT_UNSIGNED_INT32:
return CL_UNSIGNED_INT32;
case HIP_AD_FORMAT_SIGNED_INT32:
return CL_SIGNED_INT32;
case HIP_AD_FORMAT_HALF:
return CL_HALF_FLOAT;
case HIP_AD_FORMAT_FLOAT:
return CL_FLOAT;
}
}
ShouldNotReachHere();
return {};
}
inline
cl_channel_order getCLChannelOrder(const unsigned int hipNumChannels,
const int sRGB) {
switch (hipNumChannels) {
case 1:
return CL_R;
case 2:
return CL_RG;
case 4:
return (sRGB == 1) ? CL_sRGBA : CL_RGBA;
default:
break;
}
ShouldNotReachHere();
return {};
}
inline
cl_mem_object_type getCLMemObjectType(const unsigned int hipWidth,
const unsigned int hipHeight,
const unsigned int hipDepth,
const unsigned int flags) {
if (flags == hipArrayDefault) {
if ((hipWidth != 0) && (hipHeight == 0) && (hipDepth == 0)) {
return CL_MEM_OBJECT_IMAGE1D;
} else if ((hipWidth != 0) && (hipHeight != 0) && (hipDepth == 0)) {
return CL_MEM_OBJECT_IMAGE2D;
} else if ((hipWidth != 0) && (hipHeight != 0) && (hipDepth != 0)) {
return CL_MEM_OBJECT_IMAGE3D;
}
} else if (flags == hipArrayLayered) {
if ((hipWidth != 0) && (hipHeight == 0) && (hipDepth != 0)) {
return CL_MEM_OBJECT_IMAGE1D_ARRAY;
} else if ((hipWidth != 0) && (hipHeight != 0) && (hipDepth != 0)) {
return CL_MEM_OBJECT_IMAGE2D_ARRAY;
}
}
ShouldNotReachHere();
return {};
}
inline
cl_addressing_mode getCLAddressingMode(const hipTextureAddressMode hipAddressMode) {
switch (hipAddressMode) {
case hipAddressModeWrap:
return CL_ADDRESS_REPEAT;
case hipAddressModeClamp:
return CL_ADDRESS_CLAMP;
case hipAddressModeMirror:
return CL_ADDRESS_MIRRORED_REPEAT;
case hipAddressModeBorder:
return CL_ADDRESS_CLAMP_TO_EDGE;
}
ShouldNotReachHere();
return {};
}
inline
cl_filter_mode getCLFilterMode(const hipTextureFilterMode hipFilterMode) {
switch (hipFilterMode) {
case hipFilterModePoint:
return CL_FILTER_NEAREST;
case hipFilterModeLinear:
return CL_FILTER_LINEAR;
}
ShouldNotReachHere();
return {};
}
inline
cl_mem_object_type getCLMemObjectType(const hipResourceType hipResType) {
switch (hipResType) {
case hipResourceTypeLinear:
return CL_MEM_OBJECT_IMAGE1D_BUFFER;
case hipResourceTypePitch2D:
return CL_MEM_OBJECT_IMAGE2D;
default:
break;
}
ShouldNotReachHere();
return {};
}
inline
size_t getElementSize(const hipArray_const_t array) {
switch (array->Format) {
case HIP_AD_FORMAT_UNSIGNED_INT8:
case HIP_AD_FORMAT_SIGNED_INT8:
return 1 * array->NumChannels;
case HIP_AD_FORMAT_UNSIGNED_INT16:
case HIP_AD_FORMAT_SIGNED_INT16:
case HIP_AD_FORMAT_HALF:
return 2 * array->NumChannels;
case HIP_AD_FORMAT_UNSIGNED_INT32:
case HIP_AD_FORMAT_SIGNED_INT32:
case HIP_AD_FORMAT_FLOAT:
return 4 * array->NumChannels;
}
ShouldNotReachHere();
return {};
}
inline
hipChannelFormatDesc getChannelFormatDesc(int numChannels,
hipArray_Format arrayFormat) {
switch (arrayFormat) {
case HIP_AD_FORMAT_UNSIGNED_INT8:
switch (numChannels) {
case 1:
return {8, 0, 0, 0, hipChannelFormatKindUnsigned};
case 2:
return {8, 8, 0, 0, hipChannelFormatKindUnsigned};
case 4:
return {8, 8, 8, 8, hipChannelFormatKindUnsigned};
}
case HIP_AD_FORMAT_SIGNED_INT8:
switch (numChannels) {
case 1:
return {8, 0, 0, 0, hipChannelFormatKindSigned};
case 2:
return {8, 8, 0, 0, hipChannelFormatKindSigned};
case 4:
return {8, 8, 8, 8, hipChannelFormatKindSigned};
}
case HIP_AD_FORMAT_UNSIGNED_INT16:
switch (numChannels) {
case 1:
return {16, 0, 0, 0, hipChannelFormatKindUnsigned};
case 2:
return {16, 16, 0, 0, hipChannelFormatKindUnsigned};
case 4:
return {16, 16, 16, 16, hipChannelFormatKindUnsigned};
}
case HIP_AD_FORMAT_SIGNED_INT16:
switch (numChannels) {
case 1:
return {16, 0, 0, 0, hipChannelFormatKindSigned};
case 2:
return {16, 16, 0, 0, hipChannelFormatKindSigned};
case 4:
return {16, 16, 16, 16, hipChannelFormatKindSigned};
}
case HIP_AD_FORMAT_UNSIGNED_INT32:
switch (numChannels) {
case 1:
return {32, 0, 0, 0, hipChannelFormatKindUnsigned};
case 2:
return {32, 32, 0, 0, hipChannelFormatKindUnsigned};
case 4:
return {32, 32, 32, 32, hipChannelFormatKindUnsigned};
}
case HIP_AD_FORMAT_SIGNED_INT32:
switch (numChannels) {
case 1:
return {32, 0, 0, 0, hipChannelFormatKindSigned};
case 2:
return {32, 32, 0, 0, hipChannelFormatKindSigned};
case 4:
return {32, 32, 32, 32, hipChannelFormatKindSigned};
}
case HIP_AD_FORMAT_HALF:
switch (numChannels) {
case 1:
return {16, 0, 0, 0, hipChannelFormatKindFloat};
case 2:
return {16, 16, 0, 0, hipChannelFormatKindFloat};
case 4:
return {16, 16, 16, 16, hipChannelFormatKindFloat};
}
case HIP_AD_FORMAT_FLOAT:
switch (numChannels) {
case 1:
return {32, 0, 0, 0, hipChannelFormatKindFloat};
case 2:
return {32, 32, 0, 0, hipChannelFormatKindFloat};
case 4:
return {32, 32, 32, 32, hipChannelFormatKindFloat};
}
}
ShouldNotReachHere();
return {};
}
inline
unsigned int getNumChannels(const hipChannelFormatDesc& desc) {
return ((desc.x != 0) + (desc.y != 0) + (desc.z != 0) + (desc.w != 0));
}
inline
hipArray_Format getArrayFormat(const hipChannelFormatDesc& desc) {
switch (desc.f) {
case hipChannelFormatKindUnsigned:
switch (desc.x) {
case 8:
return HIP_AD_FORMAT_UNSIGNED_INT8;
case 16:
return HIP_AD_FORMAT_UNSIGNED_INT16;
case 32:
return HIP_AD_FORMAT_UNSIGNED_INT32;
}
case hipChannelFormatKindSigned:
switch (desc.x) {
case 8:
return HIP_AD_FORMAT_SIGNED_INT8;
case 16:
return HIP_AD_FORMAT_SIGNED_INT16;
case 32:
return HIP_AD_FORMAT_SIGNED_INT32;
}
case hipChannelFormatKindFloat:
switch (desc.x) {
case 16:
return HIP_AD_FORMAT_HALF;
case 32:
return HIP_AD_FORMAT_FLOAT;
}
default:
break;
}
ShouldNotReachHere();
return {};
}
inline
int getNumChannels(const hipResourceViewFormat hipFormat) {
switch (hipFormat) {
case hipResViewFormatUnsignedChar1:
case hipResViewFormatSignedChar1:
case hipResViewFormatUnsignedShort1:
case hipResViewFormatSignedShort1:
case hipResViewFormatUnsignedInt1:
case hipResViewFormatSignedInt1:
case hipResViewFormatHalf1:
case hipResViewFormatFloat1:
return 1;
case hipResViewFormatUnsignedChar2:
case hipResViewFormatSignedChar2:
case hipResViewFormatUnsignedShort2:
case hipResViewFormatSignedShort2:
case hipResViewFormatUnsignedInt2:
case hipResViewFormatSignedInt2:
case hipResViewFormatHalf2:
case hipResViewFormatFloat2:
return 2;
case hipResViewFormatUnsignedChar4:
case hipResViewFormatSignedChar4:
case hipResViewFormatUnsignedShort4:
case hipResViewFormatSignedShort4:
case hipResViewFormatUnsignedInt4:
case hipResViewFormatSignedInt4:
case hipResViewFormatHalf4:
case hipResViewFormatFloat4:
return 4;
default:
break;
}
ShouldNotReachHere();
return {};
}
inline
hipArray_Format getArrayFormat(const hipResourceViewFormat hipFormat) {
switch (hipFormat) {
case hipResViewFormatUnsignedChar1:
case hipResViewFormatUnsignedChar2:
case hipResViewFormatUnsignedChar4:
return HIP_AD_FORMAT_UNSIGNED_INT8;
case hipResViewFormatSignedChar1:
case hipResViewFormatSignedChar2:
case hipResViewFormatSignedChar4:
return HIP_AD_FORMAT_SIGNED_INT8;
case hipResViewFormatUnsignedShort1:
case hipResViewFormatUnsignedShort2:
case hipResViewFormatUnsignedShort4:
return HIP_AD_FORMAT_UNSIGNED_INT16;
case hipResViewFormatSignedShort1:
case hipResViewFormatSignedShort2:
case hipResViewFormatSignedShort4:
return HIP_AD_FORMAT_SIGNED_INT16;
case hipResViewFormatUnsignedInt1:
case hipResViewFormatUnsignedInt2:
case hipResViewFormatUnsignedInt4:
return HIP_AD_FORMAT_UNSIGNED_INT32;
case hipResViewFormatSignedInt1:
case hipResViewFormatSignedInt2:
case hipResViewFormatSignedInt4:
return HIP_AD_FORMAT_SIGNED_INT32;
case hipResViewFormatHalf1:
case hipResViewFormatHalf2:
case hipResViewFormatHalf4:
return HIP_AD_FORMAT_HALF;
case hipResViewFormatFloat1:
case hipResViewFormatFloat2:
case hipResViewFormatFloat4:
return HIP_AD_FORMAT_FLOAT;
default:
break;
}
ShouldNotReachHere();
return {};
}
inline
hipResourceViewFormat getResourceViewFormat(const hipChannelFormatDesc& desc) {
switch (desc.f) {
case hipChannelFormatKindUnsigned:
switch (getNumChannels(desc)) {
case 1:
switch (desc.x) {
case 8:
return hipResViewFormatUnsignedChar1;
case 16:
return hipResViewFormatUnsignedShort1;
case 32:
return hipResViewFormatUnsignedInt1;
}
case 2:
switch (desc.x) {
case 8:
return hipResViewFormatUnsignedChar2;
case 16:
return hipResViewFormatUnsignedShort2;
case 32:
return hipResViewFormatUnsignedInt2;
}
case 4:
switch (desc.x) {
case 8:
return hipResViewFormatUnsignedChar4;
case 16:
return hipResViewFormatUnsignedShort4;
case 32:
return hipResViewFormatUnsignedInt4;
}
}
case hipChannelFormatKindSigned:
switch (getNumChannels(desc)) {
case 1:
switch (desc.x) {
case 8:
return hipResViewFormatSignedChar1;
case 16:
return hipResViewFormatSignedShort1;
case 32:
return hipResViewFormatSignedInt1;
}
case 2:
switch (desc.x) {
case 8:
return hipResViewFormatSignedChar2;
case 16:
return hipResViewFormatSignedShort2;
case 32:
return hipResViewFormatSignedInt2;
}
case 4:
switch (desc.x) {
case 8:
return hipResViewFormatSignedChar4;
case 16:
return hipResViewFormatSignedShort4;
case 32:
return hipResViewFormatSignedInt4;
}
}
case hipChannelFormatKindFloat:
switch (getNumChannels(desc)) {
case 1:
switch (desc.x) {
case 16:
return hipResViewFormatHalf1;
case 32:
return hipResViewFormatFloat1;
}
case 2:
switch (desc.x) {
case 16:
return hipResViewFormatHalf2;
case 32:
return hipResViewFormatFloat2;
}
case 4:
switch (desc.x) {
case 16:
return hipResViewFormatHalf4;
case 32:
return hipResViewFormatFloat4;
}
}
default:
break;
}
ShouldNotReachHere();
return {};
}
inline
hipTextureDesc getTextureDesc(const textureReference* texRef) {
hipTextureDesc texDesc = {};
std::memcpy(texDesc.addressMode, texRef->addressMode, sizeof(texDesc.addressMode));
texDesc.filterMode = texRef->filterMode;
texDesc.readMode = texRef->readMode;
texDesc.sRGB = texRef->sRGB;
texDesc.normalizedCoords = texRef->normalized;
texDesc.maxAnisotropy = texRef->maxAnisotropy;
texDesc.mipmapFilterMode = texRef->mipmapFilterMode;
texDesc.mipmapLevelBias = texRef->mipmapLevelBias;
texDesc.minMipmapLevelClamp = texRef->minMipmapLevelClamp;
texDesc.maxMipmapLevelClamp = texRef->maxMipmapLevelClamp;
return texDesc;
}
inline
hipResourceViewDesc getResourceViewDesc(hipArray_const_t array,
const hipResourceViewFormat format) {
hipResourceViewDesc resViewDesc = {};
resViewDesc.format = format;
resViewDesc.width = array->width;
resViewDesc.height = array->height;
resViewDesc.depth = array->depth;
resViewDesc.firstMipmapLevel = 0;
resViewDesc.lastMipmapLevel = 0;
resViewDesc.firstLayer = 0;
resViewDesc.lastLayer = 0; /* TODO add hipArray::numLayers */
return resViewDesc;
}
inline
hipResourceViewDesc getResourceViewDesc(hipMipmappedArray_const_t array,
const hipResourceViewFormat format) {
hipResourceViewDesc resViewDesc = {};
resViewDesc.format = format;
resViewDesc.width = array->width;
resViewDesc.height = array->height;
resViewDesc.depth = array->depth;
resViewDesc.firstMipmapLevel = 0;
resViewDesc.lastMipmapLevel = 0; /* TODO add hipMipmappedArray::numMipLevels */
resViewDesc.firstLayer = 0;
resViewDesc.lastLayer = 0; /* TODO add hipArray::numLayers */
return resViewDesc;
}
inline
std::pair<hipMemoryType, hipMemoryType> getMemoryType(const hipMemcpyKind kind) {
switch (kind) {
case hipMemcpyHostToHost:
return {hipMemoryTypeHost, hipMemoryTypeHost};
case hipMemcpyHostToDevice:
return {hipMemoryTypeHost, hipMemoryTypeDevice};
case hipMemcpyDeviceToHost:
return {hipMemoryTypeDevice, hipMemoryTypeHost};
case hipMemcpyDeviceToDevice:
return {hipMemoryTypeDevice, hipMemoryTypeDevice};
case hipMemcpyDefault:
return {hipMemoryTypeUnified, hipMemoryTypeUnified};
}
ShouldNotReachHere();
return {};
}
inline
HIP_MEMCPY3D getDrvMemcpy3DDesc(const hip_Memcpy2D& desc2D) {
HIP_MEMCPY3D desc3D = {};
desc3D.srcXInBytes = desc2D.srcXInBytes;
desc3D.srcY = desc2D.srcY;
desc3D.srcZ = 0;
desc3D.srcLOD = 0;
desc3D.srcMemoryType = desc2D.srcMemoryType;
desc3D.srcHost = desc2D.srcHost;
desc3D.srcDevice = desc2D.srcDevice;
desc3D.srcArray = desc2D.srcArray;
desc3D.srcPitch = desc2D.srcPitch;
desc3D.srcHeight = 0;
desc3D.dstXInBytes = desc2D.dstXInBytes;
desc3D.dstY = desc2D.dstY;
desc3D.dstZ = 0;
desc3D.dstLOD = 0;
desc3D.dstMemoryType = desc2D.dstMemoryType;
desc3D.dstHost = desc2D.dstHost;
desc3D.dstDevice = desc2D.dstDevice;
desc3D.dstArray = desc2D.dstArray;
desc3D.dstPitch = desc2D.dstPitch;
desc3D.dstHeight = 0;
desc3D.WidthInBytes = desc2D.WidthInBytes;
desc3D.Height = desc2D.Height;
desc3D.Depth = 0;
return desc3D;
}
inline
HIP_MEMCPY3D getDrvMemcpy3DDesc(const hipMemcpy3DParms& desc) {
HIP_MEMCPY3D descDrv = {};
descDrv.WidthInBytes = desc.extent.width;
descDrv.Height = desc.extent.height;
descDrv.Depth = desc.extent.depth;
descDrv.srcXInBytes = desc.srcPos.x;
descDrv.srcY = desc.srcPos.y;
descDrv.srcZ = desc.srcPos.z;
descDrv.srcLOD = 0;
descDrv.dstXInBytes = desc.dstPos.x;
descDrv.dstY = desc.dstPos.y;
descDrv.dstZ = desc.dstPos.z;
descDrv.dstLOD = 0;
if (desc.srcArray != nullptr) {
descDrv.srcMemoryType = hipMemoryTypeArray;
descDrv.srcArray = desc.srcArray;
// When reffering to array memory, hipPos::x is in elements.
descDrv.srcXInBytes *= getElementSize(desc.srcArray);
}
if (desc.srcPtr.ptr != nullptr) {
descDrv.srcMemoryType = std::get<0>(hip::getMemoryType(desc.kind));
descDrv.srcHost = desc.srcPtr.ptr;
descDrv.srcDevice = desc.srcPtr.ptr;
descDrv.srcPitch = desc.srcPtr.pitch;
descDrv.srcHeight = desc.srcPtr.ysize;
}
if (desc.dstArray != nullptr) {
descDrv.dstMemoryType = hipMemoryTypeArray;
descDrv.dstArray = desc.dstArray;
// When reffering to array memory, hipPos::x is in elements.
descDrv.dstXInBytes *= getElementSize(desc.dstArray);
}
if (desc.dstPtr.ptr != nullptr) {
descDrv.dstMemoryType = std::get<1>(getMemoryType(desc.kind));
descDrv.dstHost = desc.dstPtr.ptr;
descDrv.dstDevice = desc.dstPtr.ptr;
descDrv.dstPitch = desc.dstPtr.pitch;
descDrv.dstHeight = desc.dstPtr.ysize;
}
// If a HIP array is participating in the copy, the extent is defined in terms of that array's elements.
if ((desc.srcArray != nullptr) && (desc.dstArray == nullptr)) {
descDrv.WidthInBytes *= getElementSize(desc.srcArray);
} else if ((desc.srcArray == nullptr) && (desc.dstArray != nullptr)) {
descDrv.WidthInBytes *= getElementSize(desc.dstArray);
} else if ((desc.srcArray != nullptr) && (desc.dstArray != nullptr)) {
descDrv.WidthInBytes *= getElementSize(desc.dstArray);
}
return descDrv;
}
inline
hipResourceType getResourceType(const HIPresourcetype resType) {
// These two enums should be isomorphic.
return static_cast<hipResourceType>(resType);
}
inline
HIPresourcetype getResourceType(const hipResourceType resType) {
// These two enums should be isomorphic.
return static_cast<HIPresourcetype>(resType);
}
inline
hipResourceDesc getResourceDesc(const HIP_RESOURCE_DESC& resDesc) {
hipResourceDesc desc;
desc.resType = getResourceType(resDesc.resType);
switch (desc.resType) {
case hipResourceTypeArray:
desc.res.array.array = resDesc.res.array.hArray;
break;
case hipResourceTypeMipmappedArray:
desc.res.mipmap.mipmap = resDesc.res.mipmap.hMipmappedArray;
break;
case hipResourceTypeLinear:
desc.res.linear.devPtr = resDesc.res.linear.devPtr;
desc.res.linear.desc = getChannelFormatDesc(resDesc.res.linear.numChannels, resDesc.res.linear.format);
desc.res.linear.sizeInBytes = resDesc.res.linear.sizeInBytes;
break;
case hipResourceTypePitch2D:
desc.res.pitch2D.devPtr = resDesc.res.pitch2D.devPtr;
desc.res.pitch2D.desc = getChannelFormatDesc(resDesc.res.pitch2D.numChannels, resDesc.res.pitch2D.format);
desc.res.pitch2D.width = resDesc.res.pitch2D.width;
desc.res.pitch2D.height = resDesc.res.pitch2D.height;
desc.res.pitch2D.pitchInBytes = resDesc.res.pitch2D.pitchInBytes;
break;
default:
break;
}
return desc;
}
inline
HIP_RESOURCE_DESC getResourceDesc(const hipResourceDesc& resDesc) {
HIP_RESOURCE_DESC desc;
desc.resType = getResourceType(resDesc.resType);
switch (desc.resType) {
case HIP_RESOURCE_TYPE_ARRAY:
desc.res.array.hArray = resDesc.res.array.array;
break;
case HIP_RESOURCE_TYPE_MIPMAPPED_ARRAY:
desc.res.mipmap.hMipmappedArray = resDesc.res.mipmap.mipmap;
break;
case HIP_RESOURCE_TYPE_LINEAR:
desc.res.linear.devPtr = resDesc.res.linear.devPtr;
desc.res.linear.numChannels = getNumChannels(resDesc.res.linear.desc);
desc.res.linear.format = getArrayFormat(resDesc.res.linear.desc);
desc.res.linear.sizeInBytes = resDesc.res.linear.sizeInBytes;
break;
case HIP_RESOURCE_TYPE_PITCH2D:
desc.res.pitch2D.devPtr = resDesc.res.pitch2D.devPtr;
desc.res.pitch2D.numChannels = getNumChannels(resDesc.res.pitch2D.desc);
desc.res.pitch2D.format = getArrayFormat(resDesc.res.pitch2D.desc);
desc.res.pitch2D.width = resDesc.res.pitch2D.width;
desc.res.pitch2D.height = resDesc.res.pitch2D.height;
desc.res.pitch2D.pitchInBytes = resDesc.res.pitch2D.pitchInBytes;
break;
default:
break;
}
return desc;
}
inline
hipTextureAddressMode getAddressMode(const HIPaddress_mode mode) {
// These two enums should be isomorphic.
return static_cast<hipTextureAddressMode>(mode);
}
inline
HIPaddress_mode getAddressMode(const hipTextureAddressMode mode) {
// These two enums should be isomorphic.
return static_cast<HIPaddress_mode>(mode);
}
inline
hipTextureFilterMode getFilterMode(const HIPfilter_mode mode) {
// These two enums should be isomorphic.
return static_cast<hipTextureFilterMode>(mode);
}
inline
HIPfilter_mode getFilterMode(const hipTextureFilterMode mode) {
// These two enums should be isomorphic.
return static_cast<HIPfilter_mode>(mode);
}
inline
hipTextureReadMode getReadMode(const unsigned int flags) {
if (flags & HIP_TRSF_READ_AS_INTEGER) {
return hipReadModeElementType;
} else {
return hipReadModeNormalizedFloat;
}
}
inline
unsigned int getReadMode(const hipTextureReadMode mode) {
if (mode == hipReadModeElementType) {
return HIP_TRSF_READ_AS_INTEGER;
} else {
return 0;
}
}
inline
int getsRGB(const unsigned int flags) {
if (flags & HIP_TRSF_SRGB) {
return 1;
} else {
return 0;
}
}
inline
unsigned int getsRGB(const int sRGB) {
if (sRGB == 1) {
return HIP_TRSF_SRGB;
} else {
return 0;
}
}
inline
int getNormalizedCoords(const unsigned int flags) {
if (flags & HIP_TRSF_NORMALIZED_COORDINATES) {
return 1;
} else {
return 0;
}
}
inline
unsigned int getNormalizedCoords(const int normalizedCoords) {
if (normalizedCoords == 1) {
return HIP_TRSF_NORMALIZED_COORDINATES;
} else {
return 0;
}
}
inline
hipTextureDesc getTextureDesc(const HIP_TEXTURE_DESC& texDesc) {
hipTextureDesc desc;
desc.addressMode[0] = getAddressMode(texDesc.addressMode[0]);
desc.addressMode[1] = getAddressMode(texDesc.addressMode[1]);
desc.addressMode[2] = getAddressMode(texDesc.addressMode[2]);
desc.filterMode = getFilterMode(texDesc.filterMode);
desc.readMode = getReadMode(texDesc.flags);
desc.sRGB = getsRGB(texDesc.flags);
std::memcpy(desc.borderColor, texDesc.borderColor, sizeof(desc.borderColor));
desc.normalizedCoords = getNormalizedCoords(texDesc.flags);
desc.maxAnisotropy = texDesc.maxAnisotropy;
desc.mipmapFilterMode = getFilterMode(texDesc.mipmapFilterMode);
desc.mipmapLevelBias = texDesc.mipmapLevelBias;
desc.minMipmapLevelClamp = texDesc.minMipmapLevelClamp;
desc.maxMipmapLevelClamp = texDesc.maxMipmapLevelClamp;
return desc;
}
inline
HIP_TEXTURE_DESC getTextureDesc(const hipTextureDesc& texDesc) {
HIP_TEXTURE_DESC desc;
desc.addressMode[0] = getAddressMode(texDesc.addressMode[0]);
desc.addressMode[1] = getAddressMode(texDesc.addressMode[1]);
desc.addressMode[2] = getAddressMode(texDesc.addressMode[2]);
desc.filterMode = getFilterMode(texDesc.filterMode);
desc.flags = 0;
desc.flags |= getReadMode(texDesc.readMode);
desc.flags |= getsRGB(texDesc.sRGB);
desc.flags |= getNormalizedCoords(texDesc.normalizedCoords);
desc.maxAnisotropy = texDesc.maxAnisotropy;
desc.mipmapFilterMode = getFilterMode(texDesc.mipmapFilterMode);
desc.mipmapLevelBias = texDesc.mipmapLevelBias;
desc.minMipmapLevelClamp = texDesc.minMipmapLevelClamp;
desc.maxMipmapLevelClamp = texDesc.maxMipmapLevelClamp;
std::memcpy(desc.borderColor, texDesc.borderColor, sizeof(desc.borderColor));
return desc;
}
inline
hipResourceViewFormat getResourceViewFormat(const HIPresourceViewFormat format) {
// These two enums should be isomorphic.
return static_cast<hipResourceViewFormat>(format);
}
inline
HIPresourceViewFormat getResourceViewFormat(const hipResourceViewFormat format) {
// These two enums should be isomorphic.
return static_cast<HIPresourceViewFormat>(format);
}
inline
hipResourceViewDesc getResourceViewDesc(const HIP_RESOURCE_VIEW_DESC& resViewDesc) {
hipResourceViewDesc desc;
desc.format = getResourceViewFormat(resViewDesc.format);
desc.width = resViewDesc.width;
desc.height = resViewDesc.height;
desc.depth = resViewDesc.depth;
desc.firstMipmapLevel = resViewDesc.firstMipmapLevel;
desc.lastMipmapLevel = resViewDesc.lastMipmapLevel;
desc.firstLayer = resViewDesc.firstLayer;
desc.lastLayer = resViewDesc.lastLayer;
return desc;
}
inline
HIP_RESOURCE_VIEW_DESC getResourceViewDesc(const hipResourceViewDesc& resViewDesc) {
HIP_RESOURCE_VIEW_DESC desc;
desc.format = getResourceViewFormat(resViewDesc.format);
desc.width = resViewDesc.width;
desc.height = resViewDesc.height;
desc.depth = resViewDesc.depth;
desc.firstMipmapLevel = resViewDesc.firstMipmapLevel;
desc.lastMipmapLevel = resViewDesc.lastMipmapLevel;
desc.firstLayer = resViewDesc.firstLayer;
desc.lastLayer = resViewDesc.lastLayer;
return desc;
}
inline
size_t getElementSize(const hipChannelFormatDesc &desc) {
return (desc.x / 4) * getNumChannels(desc);
}
};
rocclr/hip_device.cpp
+254
View File
@@ -0,0 +1,254 @@
/* Copyright (c) 2018-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"
namespace hip {
// ================================================================================================
amd::HostQueue* Device::NullStream(bool skip_alloc) {
amd::HostQueue* null_queue = null_stream_.asHostQueue(skip_alloc);
if (null_queue == nullptr) {
return nullptr;
}
// Wait for all active streams before executing commands on the default
iHipWaitActiveStreams(null_queue);
return null_queue;
}
}
hipError_t hipDeviceGet(hipDevice_t *device, int deviceId) {
HIP_INIT_API(hipDeviceGet, device, deviceId);
if (device != nullptr) {
*device = deviceId;
} else {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(hipSuccess);
};
hipError_t hipFuncSetCacheConfig (const void* func, hipFuncCache_t cacheConfig) {
HIP_INIT_API(hipFuncSetCacheConfig, cacheConfig);
// No way to set cache config yet.
HIP_RETURN(hipSuccess);
}
hipError_t hipDeviceTotalMem (size_t *bytes, hipDevice_t device) {
HIP_INIT_API(hipDeviceTotalMem, bytes, device);
if (device < 0 || static_cast<size_t>(device) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
}
if (bytes == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
auto* deviceHandle = g_devices[device]->devices()[0];
const auto& info = deviceHandle->info();
*bytes = info.globalMemSize_;
HIP_RETURN(hipSuccess);
}
hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device) {
HIP_INIT_API(hipDeviceComputeCapability, major, minor, device);
if (device < 0 || static_cast<size_t>(device) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
}
if (major == nullptr || minor == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
auto* deviceHandle = g_devices[device]->devices()[0];
const auto& info = deviceHandle->info();
*major = info.gfxipVersion_ / 100;
*minor = info.gfxipVersion_ % 100;
HIP_RETURN(hipSuccess);
}
hipError_t hipDeviceGetCount(int* count) {
HIP_INIT_API(NONE, count);
HIP_RETURN(ihipDeviceGetCount(count));
}
hipError_t ihipDeviceGetCount(int* count) {
if (count == nullptr) {
return hipErrorInvalidValue;
}
// Get all available devices
*count = g_devices.size();
if (*count < 1) {
return hipErrorNoDevice;
}
return hipSuccess;
}
hipError_t hipDeviceGetName(char *name, int len, hipDevice_t device) {
HIP_INIT_API(hipDeviceGetName, (void*)name, len, device);
if (device < 0 || static_cast<size_t>(device) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
}
if (name == nullptr || len <= 0) {
HIP_RETURN(hipErrorInvalidValue);
}
auto* deviceHandle = g_devices[device]->devices()[0];
const auto& info = deviceHandle->info();
const auto nameLen = ::strlen(info.boardName_);
// Make sure that the size of `dest` is big enough to hold `src` including
// trailing zero byte
if (nameLen > (cl_uint)(len - 1)) {
HIP_RETURN(hipErrorInvalidValue);
}
::strncpy(name, info.boardName_, (nameLen + 1));
HIP_RETURN(hipSuccess);
}
hipError_t hipGetDeviceProperties ( hipDeviceProp_t* props, hipDevice_t device ) {
HIP_INIT_API(hipGetDeviceProperties, props, device);
if (props == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
if (unsigned(device) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
}
auto* deviceHandle = g_devices[device]->devices()[0];
hipDeviceProp_t deviceProps = {0};
const auto& info = deviceHandle->info();
::strncpy(deviceProps.name, info.boardName_, 128);
deviceProps.totalGlobalMem = info.globalMemSize_;
deviceProps.sharedMemPerBlock = info.localMemSizePerCU_;
deviceProps.regsPerBlock = info.availableSGPRs_;
deviceProps.warpSize = info.wavefrontWidth_;
deviceProps.maxThreadsPerBlock = info.maxWorkGroupSize_;
deviceProps.maxThreadsDim[0] = info.maxWorkItemSizes_[0];
deviceProps.maxThreadsDim[1] = info.maxWorkItemSizes_[1];
deviceProps.maxThreadsDim[2] = info.maxWorkItemSizes_[2];
deviceProps.maxGridSize[0] = INT32_MAX;
deviceProps.maxGridSize[1] = INT32_MAX;
deviceProps.maxGridSize[2] = INT32_MAX;
deviceProps.clockRate = info.maxEngineClockFrequency_ * 1000;
deviceProps.memoryClockRate = info.maxMemoryClockFrequency_ * 1000;
deviceProps.memoryBusWidth = info.globalMemChannels_ * 32;
deviceProps.totalConstMem = info.maxConstantBufferSize_;
deviceProps.major = info.gfxipVersion_ / 100;
deviceProps.minor = info.gfxipVersion_ % 100;
deviceProps.multiProcessorCount = info.maxComputeUnits_;
deviceProps.l2CacheSize = info.l2CacheSize_;
deviceProps.maxThreadsPerMultiProcessor = info.maxThreadsPerCU_;
deviceProps.computeMode = 0;
deviceProps.clockInstructionRate = info.timeStampFrequency_;
deviceProps.arch.hasGlobalInt32Atomics = 1;
deviceProps.arch.hasGlobalFloatAtomicExch = 1;
deviceProps.arch.hasSharedInt32Atomics = 1;
deviceProps.arch.hasSharedFloatAtomicExch = 1;
deviceProps.arch.hasFloatAtomicAdd = 0;
deviceProps.arch.hasGlobalInt64Atomics = 1;
deviceProps.arch.hasSharedInt64Atomics = 1;
deviceProps.arch.hasDoubles = 1;
deviceProps.arch.hasWarpVote = 0;
deviceProps.arch.hasWarpBallot = 0;
deviceProps.arch.hasWarpShuffle = 0;
deviceProps.arch.hasFunnelShift = 0;
deviceProps.arch.hasThreadFenceSystem = 1;
deviceProps.arch.hasSyncThreadsExt = 0;
deviceProps.arch.hasSurfaceFuncs = 0;
deviceProps.arch.has3dGrid = 1;
deviceProps.arch.hasDynamicParallelism = 0;
deviceProps.concurrentKernels = 1;
deviceProps.pciDomainID = info.deviceTopology_.pcie.function;
deviceProps.pciBusID = info.deviceTopology_.pcie.bus;
deviceProps.pciDeviceID = info.deviceTopology_.pcie.device;
deviceProps.maxSharedMemoryPerMultiProcessor = info.localMemSizePerCU_;
//deviceProps.isMultiGpuBoard = info.;
deviceProps.canMapHostMemory = 1;
deviceProps.gcnArch = info.gfxipVersion_;
deviceProps.cooperativeLaunch = info.cooperativeGroups_;
deviceProps.cooperativeMultiDeviceLaunch = info.cooperativeMultiDeviceGroups_;
deviceProps.cooperativeMultiDeviceUnmatchedFunc = info.cooperativeMultiDeviceGroups_;
deviceProps.cooperativeMultiDeviceUnmatchedGridDim = info.cooperativeMultiDeviceGroups_;
deviceProps.cooperativeMultiDeviceUnmatchedBlockDim = info.cooperativeMultiDeviceGroups_;
deviceProps.cooperativeMultiDeviceUnmatchedSharedMem = info.cooperativeMultiDeviceGroups_;
deviceProps.maxTexture1D = info.imageMaxBufferSize_;
deviceProps.maxTexture2D[0] = info.image2DMaxWidth_;
deviceProps.maxTexture2D[1] = info.image2DMaxHeight_;
deviceProps.maxTexture3D[0] = info.image3DMaxWidth_;
deviceProps.maxTexture3D[1] = info.image3DMaxHeight_;
deviceProps.maxTexture3D[2] = info.image3DMaxDepth_;
deviceProps.hdpMemFlushCntl = nullptr;
deviceProps.hdpRegFlushCntl = nullptr;
deviceProps.memPitch = info.maxMemAllocSize_;
deviceProps.textureAlignment = info.imageBaseAddressAlignment_;
deviceProps.texturePitchAlignment = info.imagePitchAlignment_;
deviceProps.kernelExecTimeoutEnabled = 0;
deviceProps.ECCEnabled = info.errorCorrectionSupport_? 1:0;
deviceProps.isLargeBar = info.largeBar_ ? 1 : 0;
*props = deviceProps;
HIP_RETURN(hipSuccess);
}
hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator* acc) {
HIP_INIT_API(NONE, deviceId, acc);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view** av) {
HIP_INIT_API(NONE, stream, av);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
rocclr/hip_device_runtime.cpp
+570
View File
@@ -0,0 +1,570 @@
/* Copyright (c) 2018-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"
hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* properties) {
HIP_INIT_API(hipChooseDevice, device, properties);
if (device == nullptr || properties == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
*device = 0;
cl_uint maxMatchedCount = 0;
int count = 0;
ihipDeviceGetCount(&count);
for (cl_int i = 0; i< count; ++i) {
hipDeviceProp_t currentProp = {0};
cl_uint validPropCount = 0;
cl_uint matchedCount = 0;
hipError_t err = hipGetDeviceProperties(&currentProp, i);
if (properties->major != 0) {
validPropCount++;
if(currentProp.major >= properties->major) {
matchedCount++;
}
}
if (properties->minor != 0) {
validPropCount++;
if(currentProp.minor >= properties->minor) {
matchedCount++;
}
}
if(properties->totalGlobalMem != 0) {
validPropCount++;
if(currentProp.totalGlobalMem >= properties->totalGlobalMem) {
matchedCount++;
}
}
if(properties->sharedMemPerBlock != 0) {
validPropCount++;
if(currentProp.sharedMemPerBlock >= properties->sharedMemPerBlock) {
matchedCount++;
}
}
if(properties->maxThreadsPerBlock != 0) {
validPropCount++;
if(currentProp.maxThreadsPerBlock >= properties->maxThreadsPerBlock ) {
matchedCount++;
}
}
if(properties->totalConstMem != 0) {
validPropCount++;
if(currentProp.totalConstMem >= properties->totalConstMem ) {
matchedCount++;
}
}
if(properties->multiProcessorCount != 0) {
validPropCount++;
if(currentProp.multiProcessorCount >=
properties->multiProcessorCount ) {
matchedCount++;
}
}
if(properties->maxThreadsPerMultiProcessor != 0) {
validPropCount++;
if(currentProp.maxThreadsPerMultiProcessor >=
properties->maxThreadsPerMultiProcessor ) {
matchedCount++;
}
}
if(properties->memoryClockRate != 0) {
validPropCount++;
if(currentProp.memoryClockRate >= properties->memoryClockRate ) {
matchedCount++;
}
}
if(properties->memoryBusWidth != 0) {
validPropCount++;
if(currentProp.memoryBusWidth >= properties->memoryBusWidth ) {
matchedCount++;
}
}
if(properties->l2CacheSize != 0) {
validPropCount++;
if(currentProp.l2CacheSize >= properties->l2CacheSize ) {
matchedCount++;
}
}
if(properties->regsPerBlock != 0) {
validPropCount++;
if(currentProp.regsPerBlock >= properties->regsPerBlock ) {
matchedCount++;
}
}
if(properties->maxSharedMemoryPerMultiProcessor != 0) {
validPropCount++;
if(currentProp.maxSharedMemoryPerMultiProcessor >=
properties->maxSharedMemoryPerMultiProcessor ) {
matchedCount++;
}
}
if(properties->warpSize != 0) {
validPropCount++;
if(currentProp.warpSize >= properties->warpSize ) {
matchedCount++;
}
}
if(validPropCount == matchedCount) {
*device = matchedCount > maxMatchedCount ? i : *device;
maxMatchedCount = std::max(matchedCount, maxMatchedCount);
}
}
HIP_RETURN(hipSuccess);
}
hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) {
HIP_INIT_API(hipDeviceGetAttribute, pi, attr, device);
if (pi == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
int count = 0;
ihipDeviceGetCount(&count);
if (device < 0 || device >= count) {
HIP_RETURN(hipErrorInvalidDevice);
}
//FIXME: should we cache the props, or just select from deviceHandle->info_?
hipDeviceProp_t prop = {0};
hipError_t err = hipGetDeviceProperties(&prop, device);
if (err != hipSuccess) {
HIP_RETURN(err);
}
switch (attr) {
case hipDeviceAttributeMaxThreadsPerBlock:
*pi = prop.maxThreadsPerBlock;
break;
case hipDeviceAttributeMaxBlockDimX:
*pi = prop.maxThreadsDim[0];
break;
case hipDeviceAttributeMaxBlockDimY:
*pi = prop.maxThreadsDim[1];
break;
case hipDeviceAttributeMaxBlockDimZ:
*pi = prop.maxThreadsDim[2];
break;
case hipDeviceAttributeMaxGridDimX:
*pi = prop.maxGridSize[0];
break;
case hipDeviceAttributeMaxGridDimY:
*pi = prop.maxGridSize[1];
break;
case hipDeviceAttributeMaxGridDimZ:
*pi = prop.maxGridSize[2];
break;
case hipDeviceAttributeMaxSharedMemoryPerBlock:
*pi = prop.sharedMemPerBlock;
break;
case hipDeviceAttributeTotalConstantMemory:
*pi = prop.totalConstMem;
break;
case hipDeviceAttributeWarpSize:
*pi = prop.warpSize;
break;
case hipDeviceAttributeMaxRegistersPerBlock:
*pi = prop.regsPerBlock;
break;
case hipDeviceAttributeClockRate:
*pi = prop.clockRate;
break;
case hipDeviceAttributeMemoryClockRate:
*pi = prop.memoryClockRate;
break;
case hipDeviceAttributeMemoryBusWidth:
*pi = prop.memoryBusWidth;
break;
case hipDeviceAttributeMultiprocessorCount:
*pi = prop.multiProcessorCount;
break;
case hipDeviceAttributeComputeMode:
*pi = prop.computeMode;
break;
case hipDeviceAttributeL2CacheSize:
*pi = prop.l2CacheSize;
break;
case hipDeviceAttributeMaxThreadsPerMultiProcessor:
*pi = prop.maxThreadsPerMultiProcessor;
break;
case hipDeviceAttributeComputeCapabilityMajor:
*pi = prop.major;
break;
case hipDeviceAttributeComputeCapabilityMinor:
*pi = prop.minor;
break;
case hipDeviceAttributePciBusId:
*pi = prop.pciBusID;
break;
case hipDeviceAttributeConcurrentKernels:
*pi = prop.concurrentKernels;
break;
case hipDeviceAttributePciDeviceId:
*pi = prop.pciDeviceID;
break;
case hipDeviceAttributeMaxSharedMemoryPerMultiprocessor:
*pi = prop.maxSharedMemoryPerMultiProcessor;
break;
case hipDeviceAttributeIsMultiGpuBoard:
*pi = prop.isMultiGpuBoard;
break;
case hipDeviceAttributeCooperativeLaunch:
*pi = prop.cooperativeLaunch;
break;
case hipDeviceAttributeCooperativeMultiDeviceLaunch:
*pi = prop.cooperativeMultiDeviceLaunch;
break;
case hipDeviceAttributeIntegrated:
*pi = prop.integrated;
break;
case hipDeviceAttributeMaxTexture1DWidth:
*pi = prop.maxTexture1D;
break;
case hipDeviceAttributeMaxTexture2DWidth:
*pi = prop.maxTexture2D[0];
break;
case hipDeviceAttributeMaxTexture2DHeight:
*pi = prop.maxTexture2D[1];
break;
case hipDeviceAttributeMaxTexture3DWidth:
*pi = prop.maxTexture3D[0];
break;
case hipDeviceAttributeMaxTexture3DHeight:
*pi = prop.maxTexture3D[1];
break;
case hipDeviceAttributeMaxTexture3DDepth:
*pi = prop.maxTexture3D[2];
break;
case hipDeviceAttributeHdpMemFlushCntl:
*reinterpret_cast<unsigned int**>(pi) = prop.hdpMemFlushCntl;
break;
case hipDeviceAttributeHdpRegFlushCntl:
*reinterpret_cast<unsigned int**>(pi) = prop.hdpRegFlushCntl;
break;
case hipDeviceAttributeMaxPitch:
*pi = prop.memPitch;
break;
case hipDeviceAttributeTextureAlignment:
*pi = prop.textureAlignment;
break;
case hipDeviceAttributeTexturePitchAlignment:
*pi = prop.texturePitchAlignment;
break;
case hipDeviceAttributeKernelExecTimeout:
*pi = prop.kernelExecTimeoutEnabled;
break;
case hipDeviceAttributeCanMapHostMemory:
*pi = prop.canMapHostMemory;
break;
case hipDeviceAttributeEccEnabled:
*pi = prop.ECCEnabled;
break;
case hipDeviceAttributeCooperativeMultiDeviceUnmatchedFunc:
*pi = prop.cooperativeMultiDeviceUnmatchedFunc;
break;
case hipDeviceAttributeCooperativeMultiDeviceUnmatchedGridDim:
*pi = prop.cooperativeMultiDeviceUnmatchedGridDim;
break;
case hipDeviceAttributeCooperativeMultiDeviceUnmatchedBlockDim:
*pi = prop.cooperativeMultiDeviceUnmatchedBlockDim;
break;
case hipDeviceAttributeCooperativeMultiDeviceUnmatchedSharedMem:
*pi = prop.cooperativeMultiDeviceUnmatchedSharedMem;
break;
default:
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(hipSuccess);
}
hipError_t hipDeviceGetByPCIBusId(int* device, const char*pciBusIdstr) {
HIP_INIT_API(hipDeviceGetByPCIBusId, device, pciBusIdstr);
if (device == nullptr || pciBusIdstr == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
int pciBusID = -1;
int pciDeviceID = -1;
int pciDomainID = -1;
if (sscanf (pciBusIdstr, "%04x:%02x:%02x", &pciDomainID, &pciBusID, &pciDeviceID) == 0x3) {
int count = 0;
ihipDeviceGetCount(&count);
for (cl_int i = 0; i < count; i++) {
int pi = 0;
hipDevice_t dev;
hipDeviceGet(&dev, i);
hipDeviceGetAttribute(&pi, hipDeviceAttributePciBusId, dev);
if (pciBusID == pi) {
*device = i;
break;
}
}
}
HIP_RETURN(hipSuccess);
}
hipError_t hipDeviceGetCacheConfig ( hipFuncCache_t * cacheConfig ) {
HIP_INIT_API(hipDeviceGetCacheConfig, cacheConfig);
if(cacheConfig == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
*cacheConfig = hipFuncCache_t();
HIP_RETURN(hipSuccess);
}
hipError_t hipDeviceGetLimit ( size_t* pValue, hipLimit_t limit ) {
HIP_INIT_API(hipDeviceGetLimit, pValue, limit);
if(pValue == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
if(limit == hipLimitMallocHeapSize) {
hipDeviceProp_t prop;
hipGetDeviceProperties(&prop, ihipGetDevice());
*pValue = prop.totalGlobalMem;
HIP_RETURN(hipSuccess);
} else {
HIP_RETURN(hipErrorUnsupportedLimit);
}
}
/**
hipError_t hipDeviceGetP2PAttribute ( int* value, hipDeviceP2PAttr attr, int srcDevice, int dstDevice ) {
assert(0);
HIP_RETURN(hipSuccess);
}
**/
hipError_t hipDeviceGetPCIBusId ( char* pciBusId, int len, int device ) {
HIP_INIT_API(hipDeviceGetPCIBusId, (void*)pciBusId, len, device);
int count;
ihipDeviceGetCount(&count);
if (device < 0 || device > count) {
HIP_RETURN(hipErrorInvalidDevice);
}
if (pciBusId == nullptr || len < 0) {
HIP_RETURN(hipErrorInvalidValue);
}
hipDeviceProp_t prop;
hipGetDeviceProperties(&prop, device);
snprintf (pciBusId, len, "%04x:%02x:%02x.0",
prop.pciDomainID,
prop.pciBusID,
prop.pciDeviceID);
HIP_RETURN(hipSuccess);
}
hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig ) {
HIP_INIT_API(hipDeviceGetSharedMemConfig, pConfig);
*pConfig = hipSharedMemBankSizeFourByte;
HIP_RETURN(hipSuccess);
}
hipError_t hipDeviceReset ( void ) {
HIP_INIT_API(hipDeviceReset);
/* FIXME */
HIP_RETURN(hipSuccess);
}
hipError_t hipDeviceSetCacheConfig ( hipFuncCache_t cacheConfig ) {
HIP_INIT_API(hipDeviceSetCacheConfig, cacheConfig);
// No way to set cache config yet.
HIP_RETURN(hipSuccess);
}
hipError_t hipDeviceSetLimit ( hipLimit_t limit, size_t value ) {
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ) {
HIP_INIT_API(hipDeviceSetSharedMemConfig, config);
// No way to set cache config yet.
HIP_RETURN(hipSuccess);
}
hipError_t hipDeviceSynchronize ( void ) {
HIP_INIT_API(hipDeviceSynchronize);
amd::HostQueue* queue = hip::getNullStream();
if (!queue) {
HIP_RETURN(hipErrorOutOfMemory);
}
queue->finish();
HIP_RETURN(hipSuccess);
}
int ihipGetDevice() {
return hip::getCurrentDevice()->deviceId();
}
hipError_t hipGetDevice ( int* deviceId ) {
HIP_INIT_API(hipGetDevice, deviceId);
if (deviceId != nullptr) {
int dev = ihipGetDevice();
if (dev == -1) {
HIP_RETURN(hipErrorNoDevice);
}
*deviceId = dev;
HIP_RETURN(hipSuccess);
} else {
HIP_RETURN(hipErrorInvalidValue);
}
}
hipError_t hipGetDeviceCount ( int* count ) {
HIP_INIT_API(hipGetDeviceCount, count);
HIP_RETURN(ihipDeviceGetCount(count));
}
hipError_t hipGetDeviceFlags ( unsigned int* flags ) {
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipIpcGetEventHandle ( hipIpcEventHandle_t* handle, hipEvent_t event ) {
HIP_INIT_API(NONE, handle, event);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipIpcOpenEventHandle ( hipEvent_t* event, hipIpcEventHandle_t handle ) {
HIP_INIT_API(NONE, event, handle);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipSetDevice ( int device ) {
HIP_INIT_API(hipSetDevice, device);
if (static_cast<unsigned int>(device) < g_devices.size()) {
hip::setCurrentDevice(device);
HIP_RETURN(hipSuccess);
}
HIP_RETURN(hipErrorInvalidDevice);
}
hipError_t hipSetDeviceFlags ( unsigned int flags ) {
HIP_INIT_API(hipSetDeviceFlags, flags);
constexpr uint32_t supportedFlags =
hipDeviceScheduleMask | hipDeviceMapHost | hipDeviceLmemResizeToMax;
if (flags & ~supportedFlags) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* device = hip::getCurrentDevice()->devices()[0];
switch (flags & hipDeviceScheduleMask) {
case hipDeviceScheduleAuto:
// Current behavior is different from the spec, due to MT usage in runtime
if (hip::host_device->devices().size() >= std::thread::hardware_concurrency()) {
device->SetActiveWait(false);
break;
}
// Fall through for active wait...
case hipDeviceScheduleSpin:
case hipDeviceScheduleYield:
// The both options falls into yield, because MT usage in runtime
device->SetActiveWait(true);
break;
case hipDeviceScheduleBlockingSync:
device->SetActiveWait(false);
break;
default:
break;
}
HIP_RETURN(hipSuccess);
}
hipError_t hipSetValidDevices ( int* device_arr, int len ) {
HIP_INIT_API(NONE, device_arr, len);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipExtGetLinkTypeAndHopCount(int device1, int device2, uint32_t* linktype, uint32_t* hopcount) {
HIP_INIT_API(hipExtGetLinkTypeAndHopCount, device1, device2, linktype, hopcount);
amd::Device* amd_dev_obj1 = nullptr;
amd::Device* amd_dev_obj2 = nullptr;
const int numDevices = static_cast<int>(g_devices.size());
if ((device1 < 0) || (device1 >= numDevices) || (device2 < 0) || (device2 >= numDevices)) {
HIP_RETURN(hipErrorInvalidDevice);
}
if ((linktype == nullptr) || (hopcount == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd_dev_obj1 = g_devices[device1]->devices()[0];
amd_dev_obj2 = g_devices[device2]->devices()[0];
if (!amd_dev_obj1->findLinkTypeAndHopCount(amd_dev_obj2, linktype, hopcount)) {
HIP_RETURN(hipErrorInvalidHandle);
}
HIP_RETURN(hipSuccess);
}
rocclr/hip_error.cpp
+172
View File
@@ -0,0 +1,172 @@
/* Copyright (c) 2015-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"
hipError_t hipGetLastError()
{
HIP_INIT_API(hipGetLastError);
hipError_t err = hip::g_lastError;
hip::g_lastError = hipSuccess;
return err;
}
hipError_t hipPeekAtLastError()
{
HIP_INIT_API(hipPeekAtLastError);
hipError_t err = hip::g_lastError;
HIP_RETURN(err);
}
const char *hipGetErrorName(hipError_t hip_error)
{
switch (hip_error) {
case hipSuccess:
return "hipSuccess";
case hipErrorInvalidValue:
return "hipErrorInvalidValue";
case hipErrorOutOfMemory:
return "hipErrorOutOfMemory";
case hipErrorNotInitialized:
return "hipErrorNotInitialized";
case hipErrorDeinitialized:
return "hipErrorDeinitialized";
case hipErrorProfilerDisabled:
return "hipErrorProfilerDisabled";
case hipErrorProfilerNotInitialized:
return "hipErrorProfilerNotInitialized";
case hipErrorProfilerAlreadyStarted:
return "hipErrorProfilerAlreadyStarted";
case hipErrorProfilerAlreadyStopped:
return "hipErrorProfilerAlreadyStopped";
case hipErrorInvalidConfiguration:
return "hipErrorInvalidConfiguration";
case hipErrorInvalidSymbol:
return "hipErrorInvalidSymbol";
case hipErrorInvalidDevicePointer:
return "hipErrorInvalidDevicePointer";
case hipErrorInvalidMemcpyDirection:
return "hipErrorInvalidMemcpyDirection";
case hipErrorInsufficientDriver:
return "hipErrorInsufficientDriver";
case hipErrorMissingConfiguration:
return "hipErrorMissingConfiguration";
case hipErrorPriorLaunchFailure:
return "hipErrorPriorLaunchFailure";
case hipErrorInvalidDeviceFunction:
return "hipErrorInvalidDeviceFunction";
case hipErrorNoDevice:
return "hipErrorNoDevice";
case hipErrorInvalidDevice:
return "hipErrorInvalidDevice";
case hipErrorInvalidImage:
return "hipErrorInvalidImage";
case hipErrorInvalidContext:
return "hipErrorInvalidContext";
case hipErrorContextAlreadyCurrent:
return "hipErrorContextAlreadyCurrent";
case hipErrorMapFailed:
return "hipErrorMapFailed";
case hipErrorUnmapFailed:
return "hipErrorUnmapFailed";
case hipErrorArrayIsMapped:
return "hipErrorArrayIsMapped";
case hipErrorAlreadyMapped:
return "hipErrorAlreadyMapped";
case hipErrorNoBinaryForGpu:
return "hipErrorNoBinaryForGpu";
case hipErrorAlreadyAcquired:
return "hipErrorAlreadyAcquired";
case hipErrorNotMapped:
return "hipErrorNotMapped";
case hipErrorNotMappedAsArray:
return "hipErrorNotMappedAsArray";
case hipErrorNotMappedAsPointer:
return "hipErrorNotMappedAsPointer";
case hipErrorECCNotCorrectable:
return "hipErrorECCNotCorrectable";
case hipErrorUnsupportedLimit:
return "hipErrorUnsupportedLimit";
case hipErrorContextAlreadyInUse:
return "hipErrorContextAlreadyInUse";
case hipErrorPeerAccessUnsupported:
return "hipErrorPeerAccessUnsupported";
case hipErrorInvalidKernelFile:
return "hipErrorInvalidKernelFile";
case hipErrorInvalidGraphicsContext:
return "hipErrorInvalidGraphicsContext";
case hipErrorInvalidSource:
return "hipErrorInvalidSource";
case hipErrorFileNotFound:
return "hipErrorFileNotFound";
case hipErrorSharedObjectSymbolNotFound:
return "hipErrorSharedObjectSymbolNotFound";
case hipErrorSharedObjectInitFailed:
return "hipErrorSharedObjectInitFailed";
case hipErrorOperatingSystem:
return "hipErrorOperatingSystem";
case hipErrorInvalidHandle:
return "hipErrorInvalidHandle";
case hipErrorNotFound:
return "hipErrorNotFound";
case hipErrorNotReady:
return "hipErrorNotReady";
case hipErrorIllegalAddress:
return "hipErrorIllegalAddress";
case hipErrorLaunchOutOfResources:
return "hipErrorLaunchOutOfResources";
case hipErrorLaunchTimeOut:
return "hipErrorLaunchTimeOut";
case hipErrorPeerAccessAlreadyEnabled:
return "hipErrorPeerAccessAlreadyEnabled";
case hipErrorPeerAccessNotEnabled:
return "hipErrorPeerAccessNotEnabled";
case hipErrorSetOnActiveProcess:
return "hipErrorSetOnActiveProcess";
case hipErrorAssert:
return "hipErrorAssert";
case hipErrorHostMemoryAlreadyRegistered:
return "hipErrorHostMemoryAlreadyRegistered";
case hipErrorHostMemoryNotRegistered:
return "hipErrorHostMemoryNotRegistered";
case hipErrorLaunchFailure:
return "hipErrorLaunchFailure";
case hipErrorNotSupported:
return "hipErrorNotSupported";
case hipErrorUnknown:
return "hipErrorUnknown";
case hipErrorRuntimeMemory:
return "hipErrorRuntimeMemory";
case hipErrorRuntimeOther:
return "hipErrorRuntimeOther";
case hipErrorTbd:
return "hipErrorTbd";
default:
return "hipErrorUnknown";
};
}
const char *hipGetErrorString(hipError_t hip_error)
{
return hipGetErrorName(hip_error);
}
rocclr/hip_event.cpp
+251
View File
@@ -0,0 +1,251 @@
/* Copyright (c) 2015-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_event.hpp"
namespace hip {
bool Event::ready() {
if (event_->status() != CL_COMPLETE) {
event_->notifyCmdQueue();
}
return (event_->status() == CL_COMPLETE);
}
hipError_t Event::query() {
amd::ScopedLock lock(lock_);
// If event is not recorded, event_ is null, hence return hipSuccess
if (event_ == nullptr) {
return hipSuccess;
}
return ready() ? hipSuccess : hipErrorNotReady;
}
hipError_t Event::synchronize() {
amd::ScopedLock lock(lock_);
// If event is not recorded, event_ is null, hence return hipSuccess
if (event_ == nullptr) {
return hipSuccess;
}
event_->awaitCompletion();
return hipSuccess;
}
hipError_t Event::elapsedTime(Event& eStop, float& ms) {
amd::ScopedLock startLock(lock_);
if (this == &eStop) {
if (event_ == nullptr) {
return hipErrorInvalidHandle;
}
if (flags & hipEventDisableTiming) {
return hipErrorInvalidHandle;
}
if (!ready()) {
return hipErrorNotReady;
}
ms = 0.f;
return hipSuccess;
}
amd::ScopedLock stopLock(eStop.lock_);
if (event_ == nullptr ||
eStop.event_ == nullptr) {
return hipErrorInvalidHandle;
}
if ((flags | eStop.flags) & hipEventDisableTiming) {
return hipErrorInvalidHandle;
}
if (!ready() || !eStop.ready()) {
return hipErrorNotReady;
}
ms = static_cast<float>(static_cast<int64_t>(eStop.event_->profilingInfo().end_ -
event_->profilingInfo().start_))/1000000.f;
return hipSuccess;
}
hipError_t Event::streamWait(amd::HostQueue* hostQueue, uint flags) {
if ((event_ == nullptr) || (event_->command().queue() == hostQueue)) {
return hipSuccess;
}
amd::ScopedLock lock(lock_);
bool retain = false;
if (!event_->notifyCmdQueue()) {
return hipErrorLaunchOutOfResources;
}
amd::Command::EventWaitList eventWaitList;
eventWaitList.push_back(event_);
amd::Command* command = new amd::Marker(*hostQueue, false, eventWaitList);
if (command == NULL) {
return hipErrorOutOfMemory;
}
command->enqueue();
command->release();
return hipSuccess;
}
void Event::addMarker(amd::HostQueue* queue, amd::Command* command) {
amd::ScopedLock lock(lock_);
if (event_ == &command->event()) return;
if (event_ != nullptr) {
event_->release();
}
event_ = &command->event();
}
}
hipError_t ihipEventCreateWithFlags(hipEvent_t* event, unsigned flags) {
if (event == nullptr) {
return hipErrorInvalidValue;
}
unsigned supportedFlags = hipEventDefault | hipEventBlockingSync | hipEventDisableTiming |
hipEventReleaseToDevice | hipEventReleaseToSystem;
const unsigned releaseFlags = (hipEventReleaseToDevice | hipEventReleaseToSystem);
const bool illegalFlags =
(flags & ~supportedFlags) || // can't set any unsupported flags.
(flags & releaseFlags) == releaseFlags; // can't set both release flags
if (!illegalFlags) {
hip::Event* e = new hip::Event(flags);
if (e == nullptr) {
return hipErrorOutOfMemory;
}
*event = reinterpret_cast<hipEvent_t>(e);
} else {
return hipErrorInvalidValue;
}
return hipSuccess;
}
hipError_t ihipEventQuery(hipEvent_t event) {
if (event == nullptr) {
return hipErrorInvalidHandle;
}
hip::Event* e = reinterpret_cast<hip::Event*>(event);
return e->query();
}
hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags) {
HIP_INIT_API(hipEventCreateWithFlags, event, flags);
HIP_RETURN(ihipEventCreateWithFlags(event, flags));
}
hipError_t hipEventCreate(hipEvent_t* event) {
HIP_INIT_API(hipEventCreate, event);
HIP_RETURN(ihipEventCreateWithFlags(event, 0));
}
hipError_t hipEventDestroy(hipEvent_t event) {
HIP_INIT_API(hipEventDestroy, event);
if (event == nullptr) {
HIP_RETURN(hipErrorInvalidHandle);
}
delete reinterpret_cast<hip::Event*>(event);
HIP_RETURN(hipSuccess);
}
hipError_t hipEventElapsedTime(float *ms, hipEvent_t start, hipEvent_t stop) {
HIP_INIT_API(hipEventElapsedTime, ms, start, stop);
if (start == nullptr || stop == nullptr) {
HIP_RETURN(hipErrorInvalidHandle);
}
if (ms == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
hip::Event* eStart = reinterpret_cast<hip::Event*>(start);
hip::Event* eStop = reinterpret_cast<hip::Event*>(stop);
HIP_RETURN(eStart->elapsedTime(*eStop, *ms));
}
hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) {
HIP_INIT_API(hipEventRecord, event, stream);
if (event == nullptr) {
HIP_RETURN(hipErrorInvalidHandle);
}
amd::HostQueue* queue = hip::getQueue(stream);
amd::Command* command = queue->getLastQueuedCommand(true);
if (command == nullptr) {
command = new amd::Marker(*queue, false);
command->enqueue();
}
hip::Event* e = reinterpret_cast<hip::Event*>(event);
e->addMarker(queue, command);
HIP_RETURN(hipSuccess);
}
hipError_t hipEventSynchronize(hipEvent_t event) {
HIP_INIT_API(hipEventSynchronize, event);
if (event == nullptr) {
HIP_RETURN(hipErrorInvalidHandle);
}
hip::Event* e = reinterpret_cast<hip::Event*>(event);
HIP_RETURN(e->synchronize());
}
hipError_t hipEventQuery(hipEvent_t event) {
HIP_INIT_API(hipEventQuery, event);
HIP_RETURN(ihipEventQuery(event));
}
rocclr/hip_event.hpp
+68
View File
@@ -0,0 +1,68 @@
/* Copyright (c) 2015-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. */
#ifndef HIP_EVENT_H
#define HIP_EVENT_H
#include "hip_internal.hpp"
#include "thread/monitor.hpp"
namespace hip {
class TimerMarker: public amd::Marker {
public:
TimerMarker(amd::HostQueue& queue) : amd::Marker(queue, false) {
profilingInfo_.enabled_ = true;
profilingInfo_.callback_ = nullptr;
profilingInfo_.start_ = profilingInfo_.end_ = 0;
}
};
class Event {
public:
Event(unsigned int flags) : flags(flags), lock_("hipEvent_t"), event_(nullptr) {
// No need to init event_ here as addMarker does that
}
~Event() {
if (event_ != nullptr) {
event_->release();
}
}
unsigned int flags;
hipError_t query();
hipError_t synchronize();
hipError_t elapsedTime(Event& stop, float& ms);
hipError_t streamWait(amd::HostQueue* queue, uint flags);
void addMarker(amd::HostQueue* queue, amd::Command* command);
private:
amd::Monitor lock_;
amd::HostQueue* stream_;
amd::Event* event_;
bool ready();
};
};
#endif // HIP_EVEMT_H
rocclr/hip_formatting.hpp
+843
View File
@@ -0,0 +1,843 @@
/* Copyright (c) 2015-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/hiprtc.h>
#include <hip/hip_runtime_api.h>
inline std::ostream& operator<<(std::ostream& os, const hipTextureFilterMode& s) {
switch (s) {
case hipFilterModePoint:
os << "hipFilterModePoint";
break;
case hipFilterModeLinear:
os << "hipFilterModeLinear";
break;
default:
os << "hipFilterModePoint";
};
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipTextureReadMode& s) {
switch (s) {
case hipReadModeElementType:
os << "hipReadModeElementType";
break;
case hipReadModeNormalizedFloat:
os << "hipReadModeNormalizedFloat";
break;
default:
os << "hipReadModeElementType";
};
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipTextureAddressMode& s) {
switch (s) {
case hipAddressModeWrap:
os << "hipAddressModeWrap";
break;
case hipAddressModeClamp:
os << "hipAddressModeClamp";
break;
case hipAddressModeMirror:
os << "hipAddressModeMirror";
break;
case hipAddressModeBorder:
os << "hipAddressModeBorder";
break;
default:
os << "hipAddressModeWrap";
};
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipMemcpyKind& s) {
switch (s) {
case hipMemcpyHostToHost:
os << "hipMemcpyHostToHost";
break;
case hipMemcpyHostToDevice:
os << "hipMemcpyHostToDevice";
break;
case hipMemcpyDeviceToHost:
os << "hipMemcpyDeviceToHost";
break;
case hipMemcpyDeviceToDevice:
os << "hipMemcpyDeviceToDevice";
break;
case hipMemcpyDefault:
os << "hipMemcpyDefault";
break;
default:
os << "hipMemcpyDefault";
};
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipChannelFormatKind& s) {
switch (s) {
case hipChannelFormatKindSigned:
os << "hipChannelFormatKindSigned";
break;
case hipChannelFormatKindUnsigned:
os << "hipMemcpyHostToDevice";
break;
case hipChannelFormatKindFloat:
os << "hipChannelFormatKindFloat";
break;
case hipChannelFormatKindNone:
os << "hipChannelFormatKindNone";
break;
default:
os << "hipChannelFormatKindNone";
};
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipArray_Format& s) {
switch (s) {
case HIP_AD_FORMAT_UNSIGNED_INT8:
os << "HIP_AD_FORMAT_UNSIGNED_INT8";
break;
case HIP_AD_FORMAT_UNSIGNED_INT16:
os << "HIP_AD_FORMAT_UNSIGNED_INT16";
break;
case HIP_AD_FORMAT_UNSIGNED_INT32:
os << "HIP_AD_FORMAT_UNSIGNED_INT32";
break;
case HIP_AD_FORMAT_SIGNED_INT8:
os << "HIP_AD_FORMAT_SIGNED_INT8";
break;
case HIP_AD_FORMAT_SIGNED_INT16:
os << "HIP_AD_FORMAT_SIGNED_INT16";
break;
case HIP_AD_FORMAT_SIGNED_INT32:
os << "HIP_AD_FORMAT_SIGNED_INT32";
break;
case HIP_AD_FORMAT_HALF:
os << "HIP_AD_FORMAT_HALF";
break;
case HIP_AD_FORMAT_FLOAT:
os << "HIP_AD_FORMAT_FLOAT";
break;
default:
os << "HIP_AD_FORMAT_FLOAT";
};
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipResourceViewFormat& s) {
switch (s) {
case hipResViewFormatNone:
os << "hipResViewFormatNone";
break;
case hipResViewFormatUnsignedChar1:
os << "hipResViewFormatUnsignedChar1";
break;
case hipResViewFormatUnsignedChar2:
os << "hipResViewFormatUnsignedChar2";
break;
case hipResViewFormatUnsignedChar4:
os << "hipResViewFormatUnsignedChar4";
break;
case hipResViewFormatSignedChar1:
os << "hipResViewFormatSignedChar1";
break;
case hipResViewFormatSignedChar2:
os << "hipResViewFormatSignedChar2";
break;
case hipResViewFormatSignedChar4:
os << "hipResViewFormatSignedChar4";
break;
case hipResViewFormatUnsignedShort1:
os << "hipResViewFormatUnsignedShort1";
break;
case hipResViewFormatUnsignedShort2:
os << "hipResViewFormatUnsignedShort2";
break;
case hipResViewFormatUnsignedShort4:
os << "hipResViewFormatUnsignedShort4";
break;
case hipResViewFormatSignedShort1:
os << "hipResViewFormatSignedShort1";
break;
case hipResViewFormatSignedShort2:
os << "hipResViewFormatSignedShort2";
break;
case hipResViewFormatSignedShort4:
os << "hipResViewFormatSignedShort4";
break;
case hipResViewFormatUnsignedInt1:
os << "hipResViewFormatUnsignedInt1";
break;
case hipResViewFormatUnsignedInt2:
os << "hipResViewFormatUnsignedInt2";
break;
case hipResViewFormatUnsignedInt4:
os << "hipResViewFormatUnsignedInt4";
break;
case hipResViewFormatSignedInt1:
os << "hipResViewFormatSignedInt1";
break;
case hipResViewFormatSignedInt2:
os << "hipResViewFormatSignedInt2";
break;
case hipResViewFormatSignedInt4:
os << "hipResViewFormatSignedInt4";
break;
case hipResViewFormatHalf1:
os << "hipResViewFormatHalf1";
break;
case hipResViewFormatHalf2:
os << "hipResViewFormatHalf2";
break;
case hipResViewFormatHalf4:
os << "hipResViewFormatHalf4";
break;
case hipResViewFormatFloat1:
os << "hipResViewFormatFloat1";
break;
case hipResViewFormatFloat2:
os << "hipResViewFormatFloat2";
break;
case hipResViewFormatFloat4:
os << "hipResViewFormatFloat4";
break;
case hipResViewFormatUnsignedBlockCompressed1:
os << "hipResViewFormatUnsignedBlockCompressed1";
break;
case hipResViewFormatUnsignedBlockCompressed2:
os << "hipResViewFormatUnsignedBlockCompressed2";
break;
case hipResViewFormatUnsignedBlockCompressed3:
os << "hipResViewFormatUnsignedBlockCompressed3";
break;
case hipResViewFormatUnsignedBlockCompressed4:
os << "hipResViewFormatUnsignedBlockCompressed4";
break;
case hipResViewFormatSignedBlockCompressed4:
os << "hipResViewFormatSignedBlockCompressed4";
break;
case hipResViewFormatUnsignedBlockCompressed5:
os << "hipResViewFormatUnsignedBlockCompressed5";
break;
case hipResViewFormatSignedBlockCompressed5:
os << "hipResViewFormatSignedBlockCompressed5";
break;
case hipResViewFormatUnsignedBlockCompressed6H:
os << "hipResViewFormatUnsignedBlockCompressed6H";
break;
case hipResViewFormatSignedBlockCompressed6H:
os << "hipResViewFormatSignedBlockCompressed6H";
break;
case hipResViewFormatUnsignedBlockCompressed7:
os << "hipResViewFormatUnsignedBlockCompressed7";
break;
default:
os << "hipResViewFormatNone";
};
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipFunction_attribute& s) {
switch (s) {
case HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK:
os << "HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK";
break;
case HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES:
os << "HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES";
break;
case HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES:
os << "HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES";
break;
case HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES:
os << "HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES";
break;
case HIP_FUNC_ATTRIBUTE_NUM_REGS:
os << "HIP_FUNC_ATTRIBUTE_NUM_REGS";
break;
case HIP_FUNC_ATTRIBUTE_PTX_VERSION:
os << "HIP_FUNC_ATTRIBUTE_PTX_VERSION";
break;
case HIP_FUNC_ATTRIBUTE_BINARY_VERSION:
os << "HIP_FUNC_ATTRIBUTE_BINARY_VERSION";
break;
case HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA:
os << "HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA";
break;
case HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES:
os << "HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES";
break;
case HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT:
os << "HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT";
break;
case HIP_FUNC_ATTRIBUTE_MAX:
os << "HIP_FUNC_ATTRIBUTE_MAX";
break;
default:
os << "HIP_FUNC_ATTRIBUTE_MAX";
};
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hiprtcResult& s) {
switch (s) {
case HIPRTC_SUCCESS:
os << "HIPRTC_SUCCESS";
break;
case HIPRTC_ERROR_OUT_OF_MEMORY:
os << "HIPRTC_ERROR_OUT_OF_MEMORY";
break;
case HIPRTC_ERROR_PROGRAM_CREATION_FAILURE:
os << "HIPRTC_ERROR_PROGRAM_CREATION_FAILURE";
break;
case HIPRTC_ERROR_INVALID_INPUT:
os << "HIPRTC_ERROR_INVALID_INPUT";
break;
case HIPRTC_ERROR_INVALID_PROGRAM:
os << "HIPRTC_ERROR_INVALID_PROGRAM";
break;
case HIPRTC_ERROR_INVALID_OPTION:
os << "HIPRTC_ERROR_INVALID_OPTION";
break;
case HIPRTC_ERROR_COMPILATION:
os << "HIPRTC_ERROR_COMPILATION";
break;
case HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE:
os << "HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE";
break;
case HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION:
os << "HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION";
break;
case HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION:
os << "IPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION";
break;
case HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID:
os << "HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID";
break;
case HIPRTC_ERROR_INTERNAL_ERROR:
os << "HIPRTC_ERROR_INTERNAL_ERROR";
break;
default:
os << "HIPRTC_ERROR_INTERNAL_ERROR";
};
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipJitOption& s) {
switch (s) {
case hipJitOptionMaxRegisters:
os << "hipJitOptionMaxRegisters";
break;
case hipJitOptionThreadsPerBlock:
os << "hipJitOptionThreadsPerBlock";
break;
case hipJitOptionWallTime:
os << "hipJitOptionWallTime";
break;
case hipJitOptionInfoLogBuffer:
os << "hipJitOptionInfoLogBuffer";
break;
case hipJitOptionInfoLogBufferSizeBytes:
os << "hipJitOptionInfoLogBufferSizeBytes";
break;
case hipJitOptionErrorLogBuffer:
os << "hipJitOptionErrorLogBuffer";
break;
case hipJitOptionErrorLogBufferSizeBytes:
os << "hipJitOptionErrorLogBufferSizeBytes";
break;
case hipJitOptionOptimizationLevel:
os << "hipJitOptionOptimizationLevel";
break;
case hipJitOptionTargetFromContext:
os << "hipJitOptionTargetFromContext";
break;
case hipJitOptionTarget:
os << "hipJitOptionTarget";
break;
case hipJitOptionFallbackStrategy:
os << "hipJitOptionFallbackStrategy";
break;
case hipJitOptionGenerateDebugInfo:
os << "hipJitOptionGenerateDebugInfo";
break;
case hipJitOptionCacheMode:
os << "hipJitOptionCacheMode";
break;
case hipJitOptionSm3xOpt:
os << "hipJitOptionSm3xOpt";
break;
case hipJitOptionFastCompile:
os << "hipJitOptionFastCompile";
break;
case hipJitOptionNumOptions:
os << "hipJitOptionNumOptions";
break;
default:
os << "hipJitOptionMaxRegisters";
};
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipFuncCache_t& s) {
switch (s) {
case hipFuncCachePreferNone:
os << "hipFuncCachePreferNone";
break;
case hipFuncCachePreferShared:
os << "hipFuncCachePreferShared";
break;
case hipFuncCachePreferL1:
os << "hipFuncCachePreferL1";
break;
case hipFuncCachePreferEqual:
os << "hipFuncCachePreferEqual";
break;
default:
os << "hipFuncCachePreferNone";
};
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipSharedMemConfig& s) {
switch (s) {
case hipSharedMemBankSizeDefault:
os << "hipSharedMemBankSizeDefault";
break;
case hipSharedMemBankSizeFourByte:
os << "hipSharedMemBankSizeFourByte";
break;
case hipSharedMemBankSizeEightByte:
os << "hipSharedMemBankSizeEightByte";
break;
default:
os << "hipSharedMemBankSizeDefault";
};
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipDataType& s) {
switch (s) {
case HIP_R_16F:
os << "HIP_R_16F";
break;
case HIP_R_32F:
os << "HIP_R_32F";
break;
case HIP_R_64F:
os << "HIP_R_64F";
break;
case HIP_C_16F:
os << "HIP_C_16F";
break;
case HIP_C_32F:
os << "HIP_C_32F";
break;
case HIP_C_64F:
os << "HIP_C_64F";
break;
default:
os << "HIP_R_16F";
};
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipLibraryPropertyType& s) {
switch (s) {
case HIP_LIBRARY_MAJOR_VERSION:
os << "HIP_LIBRARY_MAJOR_VERSION";
break;
case HIP_LIBRARY_MINOR_VERSION:
os << "HIP_LIBRARY_MINOR_VERSION";
break;
case HIP_LIBRARY_PATCH_LEVEL:
os << "HIP_LIBRARY_PATCH_LEVEL";
break;
default:
os << "HIP_LIBRARY_MAJOR_VERSION";
};
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hip_api_id_t& s) {
os << hip_api_name(s);
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hip_api_id_t* s) {
if (s) {
os << *s;
} else {
os << "nullptr";
}
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipTextureDesc& s) {
os << '{'
<< '{'
<< s.addressMode[0]
<< ','
<< s.addressMode[1]
<< ','
<< s.addressMode[2]
<< '}'
<< ','
<< s.filterMode
<< ','
<< s.readMode
<< ','
<< s.sRGB
<< ','
<< '{'
<< s.borderColor[0]
<< ','
<< s.borderColor[1]
<< ','
<< s.borderColor[2]
<< ','
<< s.borderColor[3]
<< '}'
<< ','
<< s.normalizedCoords
<< ','
<< s.mipmapFilterMode
<< ','
<< s.mipmapLevelBias
<< ','
<< s.minMipmapLevelClamp
<< ','
<< s.maxMipmapLevelClamp
<< '}';
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipTextureDesc* s) {
if (s) {
os << *s;
} else {
os << "nullptr";
}
return os;
}
inline std::ostream& operator<<(std::ostream& os, const dim3& s) {
os << '{'
<< s.x
<< ','
<< s.y
<< ','
<< s.z
<< '}';
return os;
}
inline std::ostream& operator<<(std::ostream& os, const dim3* s) {
if (s) {
os << *s;
} else {
os << "nullptr";
}
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipChannelFormatDesc& s) {
os << '{'
<< s.x
<< ','
<< s.y
<< ','
<< s.z
<< ','
<< s.w
<< ','
<< s.f
<< '}';
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipChannelFormatDesc* s) {
if (s) {
os << *s;
} else {
os << "nullptr";
}
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipMipmappedArray& s) {
os << '{'
<< s.data
<< ','
<< s.desc
<< ','
<< s.width
<< ','
<< s.height
<< ','
<< s.depth
<< '}';
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipMipmappedArray* s) {
if (s) {
os << *s;
} else {
os << "nullptr";
}
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipResourceDesc& s) {
os << '{'
<< s.resType
<< ','
<< '{';
switch (s.resType) {
case hipResourceTypeLinear:
os << s.res.linear.devPtr
<< ','
<< s.res.linear.desc
<< ','
<< s.res.linear.sizeInBytes;
break;
case hipResourceTypePitch2D:
os << s.res.pitch2D.devPtr
<< ','
<< s.res.pitch2D.desc
<< ','
<< s.res.pitch2D.width
<< ','
<< s.res.pitch2D.height
<< ','
<< s.res.pitch2D.pitchInBytes;
break;
case hipResourceTypeArray:
os << s.res.array.array;
break;
case hipResourceTypeMipmappedArray:
os <<s.res.mipmap.mipmap;
break;
default:
break;
}
os << '}';
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipResourceDesc* s) {
if (s) {
os << *s;
} else {
os << "nullptr";
}
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipArray& s) {
os << '{'
<< s.data
<< ','
<< s.desc
<< ','
<< s.type
<< ','
<< s.width
<< ','
<< s.height
<< ','
<< s.depth
<< ','
<< s.Format
<< ','
<< s.NumChannels
<< ','
<< s.isDrv
<< ','
<< s.textureType
<< '}';
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipArray* s) {
if (s) {
os << *s;
} else {
os << "nullptr";
}
return os;
}
inline std::ostream& operator<<(std::ostream& os, const textureReference& s) {
os << '{'
<< s.normalized
<< ','
<< s.readMode
<< ','
<< s.filterMode
<< ','
<< '{'
<< s.addressMode[0]
<< ','
<< s.addressMode[1]
<< ','
<< s.addressMode[2]
<< '}'
<< ','
<< s.channelDesc
<< ','
<< s.sRGB
<< ','
<< s.maxAnisotropy
<< ','
<< s.mipmapFilterMode
<< ','
<< s.mipmapLevelBias
<< ','
<< s.minMipmapLevelClamp
<< ','
<< s.maxMipmapLevelClamp
<< ','
<< s.textureObject
<< '}';
return os;
}
inline std::ostream& operator<<(std::ostream& os, const textureReference* s) {
if (s) {
os << *s;
} else {
os << "nullptr";
}
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipError_t& s) {
os << hipGetErrorName(s);
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipError_t* s) {
if (s) {
os << *s;
} else {
os << "nullptr";
}
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipResourceViewDesc& s) {
os << '{'
<< s.format
<< ','
<< s.width
<< ','
<< s.height
<< ','
<< s.depth
<< ','
<< s.firstMipmapLevel
<< ','
<< s.lastMipmapLevel
<< ','
<< s.firstLayer
<< ','
<< s.lastLayer
<< '}';
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipResourceViewDesc* s) {
if (s) {
os << *s;
} else {
os << "nullptr";
}
return os;
}
inline std::ostream& operator<<(std::ostream& os, const HIP_ARRAY_DESCRIPTOR& s) {
os << '{'
<< s.Width
<< ','
<< s.Height
<< ','
<< s.Format
<< ','
<< s.NumChannels
<< '}';
return os;
}
inline std::ostream& operator<<(std::ostream& os, const HIP_ARRAY_DESCRIPTOR* s) {
if (s) {
os << *s;
} else {
os << "nullptr";
}
return os;
}
inline std::ostream& operator<<(std::ostream& os, const HIP_ARRAY3D_DESCRIPTOR& s) {
os << '{'
<< s.Width
<< ','
<< s.Height
<< ','
<< s.Depth
<< ','
<< s.Format
<< ','
<< s.NumChannels
<< ','
<< s.Flags
<< '}';
return os;
}
inline std::ostream& operator<<(std::ostream& os, const HIP_ARRAY3D_DESCRIPTOR* s) {
if (s) {
os << *s;
} else {
os << "nullptr";
}
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipExtent& s) {
os << '{'
<< s.width
<< ','
<< s.height
<< ','
<< s.depth
<< '}';
return os;
}
rocclr/hip_hcc.def.in
Executable
+253
View File
@@ -0,0 +1,253 @@
EXPORTS
hipChooseDevice
hipCtxCreate
hipCtxDestroy
hipCtxDisablePeerAccess
hipCtxEnablePeerAccess
hipCtxGetApiVersion
hipCtxGetCacheConfig
hipCtxGetCurrent
hipCtxGetDevice
hipCtxGetFlags
hipCtxGetSharedMemConfig
hipCtxPopCurrent
hipCtxPushCurrent
hipCtxSetCacheConfig
hipCtxSetCurrent
hipCtxSetSharedMemConfig
hipCtxSynchronize
hipDeviceCanAccessPeer
hipDeviceComputeCapability
hipDeviceDisablePeerAccess
hipDeviceEnablePeerAccess
hipDeviceGet
hipDeviceGetAttribute
hipDeviceGetByPCIBusId
hipDeviceGetCacheConfig
hipDeviceGetStreamPriorityRange
hipDeviceGetLimit
hipDeviceGetName
hipDeviceGetPCIBusId
hipDeviceGetSharedMemConfig
hipDevicePrimaryCtxGetState
hipDevicePrimaryCtxRelease
hipDevicePrimaryCtxReset
hipDevicePrimaryCtxRetain
hipDevicePrimaryCtxSetFlags
hipDeviceReset
hipDeviceSetCacheConfig
hipDeviceSetSharedMemConfig
hipDeviceSynchronize
hipDeviceTotalMem
hipDriverGetVersion
hipEventCreate
hipEventCreateWithFlags
hipEventDestroy
hipEventElapsedTime
hipEventQuery
hipEventRecord
hipEventSynchronize
hipExtGetLinkTypeAndHopCount
hipExtLaunchMultiKernelMultiDevice
hipExtMallocWithFlags
hipExtModuleLaunchKernel
hipFree
hipFreeArray
hipFuncSetCacheConfig
hipGetDevice
hipGetDeviceCount
hipGetDeviceProperties
hipGetErrorName
hipGetErrorString
hipGetLastError
hipMemAllocHost
hipHostAlloc
hipHostFree
hipHostGetDevicePointer
hipHostGetFlags
hipHostMalloc
hipHostRegister
hipHostUnregister
hipInit
hipIpcCloseMemHandle
hipIpcGetMemHandle
hipIpcOpenMemHandle
hipMalloc
hipMalloc3D
hipMalloc3DArray
hipMallocManaged
hipArrayCreate
hipArray3DCreate
hipMallocArray
hipMemAllocPitch
hipMallocPitch
hipMemcpy
hipMemcpyWithStream
hipMemcpyParam2D
hipMemcpy2D
hipMemcpy2DAsync
hipMemcpy2DToArray
hipMemcpy3D
hipMemcpy3DAsync
hipDrvMemcpy3D
hipDrvMemcpy3DAsync
hipMemcpyAsync
hipMemcpyDtoD
hipMemcpyDtoDAsync
hipMemcpyDtoH
hipMemcpyDtoHAsync
hipMemcpyFromSymbol
hipMemcpyFromSymbolAsync
hipMemcpyHtoD
hipMemcpyHtoDAsync
hipMemcpyPeer
hipMemcpyPeerAsync
hipMemcpyToArray
hipMemcpyFromArray
hipMemcpyToSymbol
hipMemcpyToSymbolAsync
hipMemGetAddressRange
hipGetSymbolAddress
hipGetSymbolSize
hipMemGetInfo
hipMemPtrGetInfo
hipMemset
hipMemsetAsync
hipMemsetD8
hipMemsetD8Async
hipMemsetD16
hipMemsetD16Async
hipMemsetD32
hipMemsetD32Async
hipMemset2D
hipMemset2DAsync
hipMemset3D
hipMemset3DAsync
hipModuleGetFunction
hipModuleGetGlobal
hipModuleGetTexRef
hipModuleLaunchKernel
hipModuleLaunchKernelExt
hipLaunchCooperativeKernel
hipLaunchCooperativeKernelMultiDevice
hipHccModuleLaunchKernel
hipModuleLoad
hipModuleLoadData
hipModuleLoadDataEx
hipModuleUnload
hipModuleOccupancyMaxPotentialBlockSize
hipModuleOccupancyMaxPotentialBlockSizeWithFlags
hipModuleOccupancyMaxActiveBlocksPerMultiprocessor
hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags
hipOccupancyMaxPotentialBlockSize
hipOccupancyMaxActiveBlocksPerMultiprocessor
hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags
hipFuncGetAttribute
hipFuncGetAttributes
hipPeekAtLastError
hipPointerGetAttributes
hipProfilerStart
hipProfilerStop
hipRuntimeGetVersion
hipSetDevice
hipSetDeviceFlags
hipStreamAddCallback
hipStreamCreate
hipStreamCreateWithFlags
hipStreamCreateWithPriority
hipStreamDestroy
hipStreamGetFlags
hipStreamQuery
hipStreamSynchronize
hipStreamWaitEvent
__hipPopCallConfiguration
__hipPushCallConfiguration
__hipRegisterFatBinary
__hipRegisterFunction
__hipRegisterVar
__hipRegisterSurface
__hipRegisterTexture
__hipUnregisterFatBinary
__gnu_h2f_ieee
__gnu_f2h_ieee
hipConfigureCall
hipSetupArgument
hipLaunchByPtr
hipLaunchKernel
hipRegisterApiCallback
hipRemoveApiCallback
hipRegisterActivityCallback
hipRemoveActivityCallback
hipApiName
hipKernelNameRef
hipBindTexture
hipBindTexture2D
hipBindTextureToArray
hipBindTextureToMipmappedArray
hipGetTextureAlignmentOffset
hipGetTextureReference
hipUnbindTexture
hipCreateChannelDesc
hipCreateTextureObject
hipDestroyTextureObject
hipGetChannelDesc
hipGetTextureObjectResourceDesc
hipGetTextureObjectResourceViewDesc
hipGetTextureObjectTextureDesc
hipTexRefGetAddress
hipTexRefGetAddressMode
hipTexRefGetArray
hipTexRefGetBorderColor
hipTexRefGetFilterMode
hipTexRefGetFlags
hipTexRefGetFormat
hipTexRefGetMaxAnisotropy
hipTexRefGetMipmapFilterMode
hipTexRefGetMipmapLevelBias
hipTexRefGetMipmapLevelClamp
hipTexRefGetMipmappedArray
hipTexRefSetAddress
hipTexRefSetAddress2D
hipTexRefSetAddressMode
hipTexRefSetArray
hipTexRefSetBorderColor
hipTexRefSetFilterMode
hipTexRefSetFlags
hipTexRefSetFormat
hipTexRefSetMaxAnisotropy
hipTexRefSetMipmapFilterMode
hipTexRefSetMipmapLevelBias
hipTexRefSetMipmapLevelClamp
hipTexRefSetMipmappedArray
hipProfilerStart
hipProfilerStop
hipHccGetAccelerator
hipHccGetAcceleratorView
hipCreateSurfaceObject
hipDestroySurfaceObject
hipInitActivityCallback
hipEnableActivityCallback
hipGetCmdName
hiprtcAddNameExpression
hiprtcCompileProgram
hiprtcCreateProgram
hiprtcDestroyProgram
hiprtcGetLoweredName
hiprtcGetProgramLog
hiprtcGetProgramLogSize
hiprtcGetCode
hiprtcGetCodeSize
hiprtcGetErrorString
hipMipmappedArrayCreate
hipMallocMipmappedArray
hipMipmappedArrayDestroy
hipFreeMipmappedArray
hipMipmappedArrayGetLevel
hipGetMipmappedArrayLevel
hipMallocHost
hipFreeHost
hipTexObjectCreate
hipTexObjectDestroy
hipTexObjectGetResourceDesc
hipTexObjectGetResourceViewDesc
hipTexObjectGetTextureDesc
rocclr/hip_hcc.map.in
Executable
+263
View File
@@ -0,0 +1,263 @@
{
global:
hipChooseDevice;
hipCtxCreate;
hipCtxDestroy;
hipCtxDisablePeerAccess;
hipCtxEnablePeerAccess;
hipCtxGetApiVersion;
hipCtxGetCacheConfig;
hipCtxGetCurrent;
hipCtxGetDevice;
hipCtxGetFlags;
hipCtxGetSharedMemConfig;
hipCtxPopCurrent;
hipCtxPushCurrent;
hipCtxSetCacheConfig;
hipCtxSetCurrent;
hipCtxSetSharedMemConfig;
hipCtxSynchronize;
hipDeviceCanAccessPeer;
hipDeviceComputeCapability;
hipDeviceDisablePeerAccess;
hipDeviceEnablePeerAccess;
hipDeviceGet;
hipDeviceGetAttribute;
hipDeviceGetByPCIBusId;
hipDeviceGetCacheConfig;
hipDeviceGetStreamPriorityRange;
hipDeviceGetLimit;
hipDeviceGetName;
hipDeviceGetPCIBusId;
hipDeviceGetSharedMemConfig;
hipDevicePrimaryCtxGetState;
hipDevicePrimaryCtxRelease;
hipDevicePrimaryCtxReset;
hipDevicePrimaryCtxRetain;
hipDevicePrimaryCtxSetFlags;
hipDeviceReset;
hipDeviceSetCacheConfig;
hipDeviceSetSharedMemConfig;
hipDeviceSynchronize;
hipDeviceTotalMem;
hipDriverGetVersion;
hipEventCreate;
hipEventCreateWithFlags;
hipEventDestroy;
hipEventElapsedTime;
hipEventQuery;
hipEventRecord;
hipEventSynchronize;
hipExtGetLinkTypeAndHopCount;
hipExtLaunchMultiKernelMultiDevice;
hipExtMallocWithFlags;
hipExtModuleLaunchKernel;
hipFree;
hipFreeArray;
hipFuncSetCacheConfig;
hipGetDevice;
hipGetDeviceCount;
hipGetDeviceProperties;
hipGetErrorName;
hipGetErrorString;
hipGetLastError;
hipMemAllocHost;
hipHostAlloc;
hipHostFree;
hipHostGetDevicePointer;
hipHostGetFlags;
hipHostMalloc;
hipHostRegister;
hipHostUnregister;
hipInit;
hipIpcCloseMemHandle;
hipIpcGetMemHandle;
hipIpcOpenMemHandle;
hipMalloc;
hipMalloc3D;
hipMalloc3DArray;
hipMallocManaged;
hipArrayCreate;
hipArray3DCreate;
hipMallocArray;
hipMallocPitch;
hipMemAllocPitch;
hipMemcpy;
hipMemcpyWithStream;
hipMemcpyParam2D;
hipMemcpy2D;
hipMemcpy2DAsync;
hipMemcpy2DToArray;
hipMemcpy3D;
hipMemcpy3DAsync;
hipDrvMemcpy3D;
hipDrvMemcpy3DAsync;
hipMemcpyAsync;
hipMemcpyDtoD;
hipMemcpyDtoDAsync;
hipMemcpyDtoH;
hipMemcpyDtoHAsync;
hipMemcpyFromSymbol;
hipMemcpyFromSymbolAsync;
hipMemcpyHtoD;
hipMemcpyHtoDAsync;
hipMemcpyPeer;
hipMemcpyPeerAsync;
hipMemcpyToArray;
hipMemcpyFromArray;
hipMemcpyToSymbol;
hipMemcpyToSymbolAsync;
hipMemGetAddressRange;
hipGetSymbolAddress;
hipGetSymbolSize;
hipMemGetInfo;
hipMemPtrGetInfo;
hipMemset;
hipMemsetAsync;
hipMemsetD8;
hipMemsetD8Async;
hipMemsetD16;
hipMemsetD16Async;
hipMemsetD32;
hipMemsetD32Async;
hipMemset2D;
hipMemset2DAsync;
hipMemset3D;
hipMemset3DAsync;
hipModuleGetFunction;
hipModuleGetGlobal;
hipModuleGetTexRef;
hipModuleLaunchKernel;
hipModuleLaunchKernelExt;
hipLaunchCooperativeKernel;
hipLaunchCooperativeKernelMultiDevice;
hipModuleLoad;
hipModuleLoadData;
hipModuleLoadDataEx;
hipModuleUnload;
hipModuleOccupancyMaxPotentialBlockSize;
hipModuleOccupancyMaxPotentialBlockSizeWithFlags;
hipModuleOccupancyMaxActiveBlocksPerMultiprocessor;
hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags;
hipOccupancyMaxPotentialBlockSize;
hipOccupancyMaxActiveBlocksPerMultiprocessor;
hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags;
hipFuncGetAttribute;
hipFuncGetAttributes;
hipPeekAtLastError;
hipPointerGetAttributes;
hipProfilerStart;
hipProfilerStop;
hipRuntimeGetVersion;
hipSetDevice;
hipSetDeviceFlags;
hipStreamAddCallback;
hipStreamCreate;
hipStreamCreateWithFlags;
hipStreamCreateWithPriority;
hipStreamDestroy;
hipStreamGetFlags;
hipStreamQuery;
hipStreamSynchronize;
hipStreamWaitEvent;
__hipPopCallConfiguration;
__hipPushCallConfiguration;
__hipRegisterFatBinary;
__hipRegisterFunction;
__hipRegisterVar;
__hipRegisterSurface;
__hipRegisterTexture;
__hipUnregisterFatBinary;
__gnu_h2f_ieee;
__gnu_f2h_ieee;
hipConfigureCall;
hipSetupArgument;
hipLaunchByPtr;
hipLaunchKernel;
hipRegisterApiCallback;
hipRemoveApiCallback;
hipRegisterActivityCallback;
hipRemoveActivityCallback;
hipApiName;
hipKernelNameRef;
hipProfilerStart;
hipProfilerStop;
hiprtcCompileProgram;
hiprtcCreateProgram;
hiprtcDestroyProgram;
hiprtcGetLoweredName;
hiprtcGetProgramLog;
hiprtcGetProgramLogSize;
hiprtcGetCode;
hiprtcGetCodeSize;
hiprtcGetErrorString;
hiprtcAddNameExpression;
hiprtcVersion;
hipBindTexture;
hipBindTexture2D;
hipBindTextureToArray;
hipBindTextureToMipmappedArray;
hipGetTextureAlignmentOffset;
hipGetTextureReference;
hipUnbindTexture;
hipCreateChannelDesc;
hipCreateTextureObject;
hipDestroyTextureObject;
hipGetChannelDesc;
hipGetTextureObjectResourceDesc;
hipGetTextureObjectResourceViewDesc;
hipGetTextureObjectTextureDesc;
hipTexRefGetAddress;
hipTexRefGetAddressMode;
hipTexRefGetArray;
hipTexRefGetBorderColor;
hipTexRefGetFilterMode;
hipTexRefGetFlags;
hipTexRefGetFormat;
hipTexRefGetMaxAnisotropy;
hipTexRefGetMipmapFilterMode;
hipTexRefGetMipmapLevelBias;
hipTexRefGetMipmapLevelClamp;
hipTexRefGetMipmappedArray;
hipTexRefSetAddress;
hipTexRefSetAddress2D;
hipTexRefSetAddressMode;
hipTexRefSetArray;
hipTexRefSetBorderColor;
hipTexRefSetFilterMode;
hipTexRefSetFlags;
hipTexRefSetFormat;
hipTexRefSetMaxAnisotropy;
hipTexRefSetMipmapFilterMode;
hipTexRefSetMipmapLevelBias;
hipTexRefSetMipmapLevelClamp;
hipTexRefSetMipmappedArray;
hipMipmappedArrayCreate;
hipMallocMipmappedArray;
hipMipmappedArrayDestroy;
hipFreeMipmappedArray;
hipMipmappedArrayGetLevel;
hipGetMipmappedArrayLevel;
hipMallocHost;
hipFreeHost;
hipTexObjectCreate;
hipTexObjectDestroy;
hipTexObjectGetResourceDesc;
hipTexObjectGetResourceViewDesc;
hipTexObjectGetTextureDesc;
extern "C++" {
hip_impl::hipLaunchKernelGGLImpl*;
hip_impl::demangle*;
hipHccGetAccelerator*;
hipHccGetAcceleratorView*;
hipCreateSurfaceObject*;
hipDestroySurfaceObject*;
hipHccModuleLaunchKernel*;
hipExtModuleLaunchKernel*;
hipInitActivityCallback*;
hipEnableActivityCallback*;
hipGetCmdName*;
};
local:
*;
};
rocclr/hip_hcc.rc
+75
View File
@@ -0,0 +1,75 @@
#define STR(__macro__) #__macro__
#define XSTR(__macro__) STR(__macro__)
#if defined(_DEBUG)
#define DEBUG_ONLY(x) x
#else
#define DEBUG_ONLY(x)
#endif
#define VERSION_PREFIX_MAJOR 2
#define VERSION_PREFIX_MINOR 0
#define APSTUDIO_READONLY_SYMBOLS
/////////////////////////////////////////////////////////////////////////////
//
// Generated from the TEXTINCLUDE 2 resource.
//
#include "winresrc.h"
#include "utils/versions.hpp"
/////////////////////////////////////////////////////////////////////////////
#undef APSTUDIO_READONLY_SYMBOLS
/////////////////////////////////////////////////////////////////////////////
// English (U.S.) resources
#if !defined(AFX_RESOURCE_DLL) || defined(AFX_TARG_ENU)
#ifdef _WIN32
LANGUAGE LANG_ENGLISH, SUBLANG_ENGLISH_US
#pragma code_page(1252)
#endif //_WIN32
/////////////////////////////////////////////////////////////////////////////
//
// Version
//
VS_VERSION_INFO VERSIONINFO
FILEVERSION 10,0,AMD_PLATFORM_BUILD_NUMBER,AMD_PLATFORM_REVISION_NUMBER
PRODUCTVERSION 10,0,AMD_PLATFORM_BUILD_NUMBER,AMD_PLATFORM_REVISION_NUMBER
FILEFLAGSMASK 0x3fL
#ifdef _DEBUG
FILEFLAGS 0x1L
#else
FILEFLAGS 0x0L
#endif
FILEOS 0x40004L
FILETYPE 0x2L
FILESUBTYPE 0x0L
BEGIN
BLOCK "StringFileInfo"
BEGIN
BLOCK "040904b0"
BEGIN
VALUE "Comments", " \0"
VALUE "CompanyName", "Advanced Micro Devices Inc.\0"
VALUE "FileDescription", AMD_PLATFORM_NAME " OpenCL " XSTR(VERSION_PREFIX_MAJOR) "." XSTR(VERSION_PREFIX_MINOR) " Runtime\0"
VALUE "FileVersion", "10.0." XSTR(AMD_PLATFORM_BUILD_NUMBER) "." XSTR(AMD_PLATFORM_REVISION_NUMBER)
VALUE "InternalName", "OpenCL"
VALUE "LegalCopyright", "Copyright (C) 2011 Advanced Micro Devices Inc.\0"
VALUE "OriginalFilename", "OpenCL.dll"
VALUE "ProductName", "OpenCL " XSTR(VERSION_PREFIX_MAJOR) "." XSTR(VERSION_PREFIX_MINOR) " " AMD_PLATFORM_INFO "\0"
VALUE "ProductVersion", "10.0." XSTR(AMD_PLATFORM_BUILD_NUMBER) "." XSTR(AMD_PLATFORM_REVISION_NUMBER)
END
END
BLOCK "VarFileInfo"
BEGIN
VALUE "Translation", 0x409, 1200
END
END
#endif // English (U.S.) resources
/////////////////////////////////////////////////////////////////////////////
rocclr/hip_intercept.cpp
+56
View File
@@ -0,0 +1,56 @@
/* Copyright (c) 2019-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_prof_api.h"
// HIP API callback/activity
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(); }
hipError_t hipRegisterApiCallback(uint32_t id, void* fun, void* arg) {
return callbacks_table.set_callback(id, reinterpret_cast<api_callbacks_table_t::fun_t>(fun), arg) ?
hipSuccess : hipErrorInvalidValue;
}
hipError_t hipRemoveApiCallback(uint32_t id) {
return callbacks_table.set_callback(id, NULL, NULL) ? hipSuccess : hipErrorInvalidValue;
}
hipError_t hipRegisterActivityCallback(uint32_t id, void* fun, void* arg) {
return callbacks_table.set_activity(id, reinterpret_cast<api_callbacks_table_t::act_t>(fun), arg) ?
hipSuccess : hipErrorInvalidValue;
}
hipError_t hipRemoveActivityCallback(uint32_t id) {
return callbacks_table.set_activity(id, NULL, NULL) ? hipSuccess : hipErrorInvalidValue;
}
hipError_t hipEnableTracing(bool enabled) {
callbacks_table.set_enabled(enabled);
return hipSuccess;
}
const char* hipApiName(uint32_t id) {
return hip_api_name(id);
}
rocclr/hip_internal.hpp
Executable
+335
View File
@@ -0,0 +1,335 @@
/* Copyright (c) 2015-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. */
#ifndef HIP_SRC_HIP_INTERNAL_H
#define HIP_SRC_HIP_INTERNAL_H
#include "vdi_common.hpp"
#include "hip_prof_api.h"
#include "trace_helper.h"
#include "utils/debug.hpp"
#include "hip_formatting.hpp"
#include <unordered_set>
#include <thread>
#include <stack>
#include <mutex>
#include <iterator>
#ifdef _WIN32
#include <process.h>
#else
#include <unistd.h>
#endif
/*! IHIP IPC MEMORY Structure */
#define IHIP_IPC_MEM_HANDLE_SIZE 32
#define IHIP_IPC_MEM_RESERVED_SIZE LP64_SWITCH(28,24)
typedef struct ihipIpcMemHandle_st {
char ipc_handle[IHIP_IPC_MEM_HANDLE_SIZE]; ///< ipc memory handle on ROCr
size_t psize;
char reserved[IHIP_IPC_MEM_RESERVED_SIZE];
} ihipIpcMemHandle_t;
#ifdef _WIN32
inline int getpid() { return _getpid(); }
#endif
#define HIP_INIT() \
std::call_once(hip::g_ihipInitialized, hip::init); \
if (hip::g_device == nullptr && g_devices.size() > 0) { \
hip::g_device = g_devices[0]; \
}
// 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()); \
amd::Thread* thread = amd::Thread::current(); \
if (!VDI_CHECK_THREAD(thread)) { \
HIP_RETURN(hipErrorOutOfMemory); \
} \
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)); \
return hip::g_lastError;
namespace hc {
class accelerator;
class accelerator_view;
};
namespace hip {
class Device;
class Stream {
amd::HostQueue* queue_;
mutable amd::Monitor lock_;
Device* device_;
amd::CommandQueue::Priority priority_;
unsigned int flags_;
bool null_;
public:
Stream(Device* dev, amd::CommandQueue::Priority p, unsigned int f = 0, bool null_stream = false);
/// Creates the hip stream object, including AMD host queue
bool Create();
/// Get device AMD host queue object. The method can allocate the queue
amd::HostQueue* asHostQueue(bool skip_alloc = false);
void Destroy();
void Finish() const;
/// Get device ID associated with the current stream;
int DeviceId() const;
/// Returns if stream is null stream
bool Null() const { return null_; }
/// Returns the lock object for the current stream
amd::Monitor& Lock() const { return lock_; }
/// Returns the creation flags for the current stream
unsigned int Flags() const { return flags_; }
};
/// HIP Device class
class Device {
amd::Monitor lock_{"Device lock"};
/// VDI context
amd::Context* context_;
/// Device's ID
/// Store it here so we don't have to loop through the device list every time
int deviceId_;
/// VDI host queue for default streams
Stream null_stream_;
//Maintain list of user enabled peers
std::list<int> userEnabledPeers;
public:
Device(amd::Context* ctx, int devId):
context_(ctx), deviceId_(devId), null_stream_(this, amd::CommandQueue::Priority::Normal, 0, true)
{ assert(ctx != nullptr); }
~Device() {}
amd::Context* asContext() const { return context_; }
int deviceId() const { return deviceId_; }
void retain() const { context_->retain(); }
void release() const { context_->release(); }
const std::vector<amd::Device*>& devices() const { return context_->devices(); }
hipError_t EnablePeerAccess(int peerDeviceId){
amd::ScopedLock lock(lock_);
bool found = (std::find(userEnabledPeers.begin(), userEnabledPeers.end(), peerDeviceId) != userEnabledPeers.end());
if (found) {
return hipErrorPeerAccessAlreadyEnabled;
}
userEnabledPeers.push_back(peerDeviceId);
return hipSuccess;
}
hipError_t DisablePeerAccess(int peerDeviceId) {
amd::ScopedLock lock(lock_);
bool found = (std::find(userEnabledPeers.begin(), userEnabledPeers.end(), peerDeviceId) != userEnabledPeers.end());
if (found) {
userEnabledPeers.remove(peerDeviceId);
return hipSuccess;
} else {
return hipErrorPeerAccessNotEnabled;
}
}
amd::HostQueue* NullStream(bool skip_alloc = false);
};
extern std::once_flag g_ihipInitialized;
/// Current thread's device
extern thread_local Device* g_device;
extern thread_local hipError_t g_lastError;
/// Device representing the host - for pinned memory
extern Device* host_device;
extern void init();
extern Device* getCurrentDevice();
extern void setCurrentDevice(unsigned int index);
/// Get VDI queue associated with hipStream
/// Note: This follows the CUDA spec to sync with default streams
/// and Blocking streams
extern amd::HostQueue* getQueue(hipStream_t s);
/// Get default stream associated with the VDI context
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 {
dim3 gridDim_;
dim3 blockDim_;
size_t sharedMem_;
hipStream_t hStream_;
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)
{
if (initialized_) {
digestFatBinary(data, modules_[data]);
}
return &modules_[data];
}
void removeFatBinary(std::vector<std::pair<hipModule_t, bool>>* module)
{
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);
extern std::vector<hip::Device*> g_devices;
extern hipError_t ihipDeviceGetCount(int* count);
extern int ihipGetDevice();
extern hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags);
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);
#endif // HIP_SRC_HIP_INTERNAL_H
rocclr/hip_memory.cpp
Executable
+2286
View File
File diff suppressed because it is too large Load Diff
rocclr/hip_module.cpp
Executable
+692
View File
@@ -0,0 +1,692 @@
/* Copyright (c) 2015-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 <libelf.h>
#include <fstream>
#include "hip_internal.hpp"
#include "platform/program.hpp"
#include "hip_event.hpp"
#include "hip_platform.hpp"
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();
}
static uint64_t 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;
}
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)
{
HIP_INIT_API(hipModuleUnload, hmod);
if (hmod == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Program* program = as_amd(reinterpret_cast<cl_program>(hmod));
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);
}
hipError_t hipModuleLoadData(hipModule_t *module, const void *image)
{
HIP_INIT_API(hipModuleLoadData, module, image);
HIP_RETURN(ihipModuleLoadData(module, image, 0));
}
hipError_t hipModuleLoadDataEx(hipModule_t *module, const void *image,
unsigned int numOptions, hipJitOption* options,
void** optionsValues)
{
/* TODO: Pass options to Program */
HIP_INIT_API(hipModuleLoadData, module, image);
HIP_RETURN(ihipModuleLoadData(module, image, 0));
}
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)
{
const void* image = nullptr;
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)
{
HIP_INIT_API(hipModuleGetFunction, hfunc, hmod, name);
if (!PlatformState::instance().findModFunc(hfunc, hmod, name)) {
DevLogPrintfError("Cannot find the function: %s for module: 0x%x \n",
name, hmod);
HIP_RETURN(hipErrorNotFound);
}
HIP_RETURN(hipSuccess);
}
hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod, const char* name)
{
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)) {
DevLogPrintfError("Cannot find global Var: %s for module: 0x%x at device: %d \n",
name, hmod, ihipGetDevice());
HIP_RETURN(hipErrorNotFound);
}
HIP_RETURN(hipSuccess);
}
hipError_t hipFuncGetAttribute(int* value, hipFunction_attribute attrib, hipFunction_t hfunc) {
HIP_INIT_API(hipFuncGetAttribute, value, attrib, hfunc);
if ((value == nullptr) || (hfunc == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
hip::Function* function = hip::Function::asFunction(hfunc);
if (function == nullptr) {
HIP_RETURN(hipErrorInvalidHandle);
}
amd::Kernel* kernel = function->function_;
if (kernel == nullptr) {
HIP_RETURN(hipErrorInvalidDeviceFunction);
}
const device::Kernel::WorkGroupInfo* wrkGrpInfo
= kernel->getDeviceKernel(*(hip::getCurrentDevice()->devices()[0]))->workGroupInfo();
if (wrkGrpInfo == nullptr) {
HIP_RETURN(hipErrorMissingConfiguration);
}
switch(attrib) {
case HIP_FUNC_ATTRIBUTE_SHARED_SIZE_BYTES:
*value = static_cast<int>(wrkGrpInfo->localMemSize_
- wrkGrpInfo->privateMemSize_);
break;
case HIP_FUNC_ATTRIBUTE_MAX_THREADS_PER_BLOCK:
*value = static_cast<int>(wrkGrpInfo->wavefrontPerSIMD_
* wrkGrpInfo->wavefrontSize_);
break;
case HIP_FUNC_ATTRIBUTE_CONST_SIZE_BYTES:
*value = 0;
break;
case HIP_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES:
*value = static_cast<int>(wrkGrpInfo->localMemSize_);
break;
case HIP_FUNC_ATTRIBUTE_NUM_REGS:
*value = static_cast<int>(wrkGrpInfo->availableGPRs_);
break;
case HIP_FUNC_ATTRIBUTE_PTX_VERSION:
*value = 30; // Defaults to 3.0 as HCC
break;
case HIP_FUNC_ATTRIBUTE_BINARY_VERSION:
*value = static_cast<int>(kernel->signature().version());
break;
case HIP_FUNC_ATTRIBUTE_CACHE_MODE_CA:
*value = 0;
break;
case HIP_FUNC_ATTRIBUTE_MAX_DYNAMIC_SHARED_SIZE_BYTES:
*value = static_cast<int>(wrkGrpInfo->availableLDSSize_);
break;
case HIP_FUNC_ATTRIBUTE_PREFERRED_SHARED_MEMORY_CARVEOUT:
*value = 0;
break;
default:
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(hipSuccess);
}
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(hipSuccess);
}
hipError_t ihipModuleLaunchKernel(hipFunction_t f,
uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ,
uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ,
uint32_t sharedMemBytes, hipStream_t hStream,
void **kernelParams, void **extra,
hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags = 0,
uint32_t params = 0, uint32_t gridId = 0, uint32_t numGrids = 0,
uint64_t prevGridSum = 0, uint64_t allGridSum = 0, uint32_t firstDevice = 0) {
HIP_INIT_API(NONE, 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_;
amd::ScopedLock lock(function->lock_);
hip::Event* eStart = reinterpret_cast<hip::Event*>(startEvent);
hip::Event* eStop = reinterpret_cast<hip::Event*>(stopEvent);
amd::HostQueue* queue = hip::getQueue(hStream);
const amd::Device& device = queue->vdev()->device();
// Make sure dispatch doesn't exceed max workgroup size limit
if (blockDimX * blockDimY * blockDimZ > device.info().maxWorkGroupSize_) {
return hipErrorInvalidConfiguration;
}
if (params & amd::NDRangeKernelCommand::CooperativeGroups) {
if (!device.info().cooperativeGroups_) {
return hipErrorLaunchFailure;
}
int num_blocks = 0;
int num_grids = 0;
int block_size = blockDimX * blockDimY * blockDimZ;
hip_impl::ihipOccupancyMaxActiveBlocksPerMultiprocessor(
&num_blocks, &num_grids, device, f, block_size, sharedMemBytes, true);
if (((gridDimX * gridDimY * gridDimZ) / block_size) > unsigned(num_grids)) {
return hipErrorCooperativeLaunchTooLarge;
}
}
if (params & amd::NDRangeKernelCommand::CooperativeMultiDeviceGroups) {
if (!device.info().cooperativeMultiDeviceGroups_) {
return hipErrorLaunchFailure;
}
}
if (!queue) {
return hipErrorOutOfMemory;
}
size_t globalWorkOffset[3] = {0};
size_t globalWorkSize[3] = { gridDimX, gridDimY, gridDimZ };
size_t localWorkSize[3] = { blockDimX, blockDimY, blockDimZ };
amd::NDRangeContainer ndrange(3, globalWorkOffset, globalWorkSize, localWorkSize);
amd::Command::EventWaitList waitList;
address kernargs = nullptr;
// 'extra' is a struct that contains the following info: {
// HIP_LAUNCH_PARAM_BUFFER_POINTER, kernargs,
// HIP_LAUNCH_PARAM_BUFFER_SIZE, &kernargs_size,
// HIP_LAUNCH_PARAM_END }
if (extra != nullptr) {
if (extra[0] != HIP_LAUNCH_PARAM_BUFFER_POINTER ||
extra[2] != HIP_LAUNCH_PARAM_BUFFER_SIZE || extra[4] != HIP_LAUNCH_PARAM_END) {
return hipErrorNotInitialized;
}
kernargs = reinterpret_cast<address>(extra[1]);
}
const amd::KernelSignature& signature = kernel->signature();
for (size_t i = 0; i < signature.numParameters(); ++i) {
const amd::KernelParameterDescriptor& desc = signature.at(i);
if (kernelParams == nullptr) {
assert(kernargs != nullptr);
kernel->parameters().set(i, desc.size_, kernargs + desc.offset_,
desc.type_ == T_POINTER/*svmBound*/);
} else {
assert(extra == nullptr);
kernel->parameters().set(i, desc.size_, kernelParams[i], desc.type_ == T_POINTER/*svmBound*/);
}
}
amd::NDRangeKernelCommand* command = new amd::NDRangeKernelCommand(
*queue, waitList, *kernel, ndrange, sharedMemBytes,
params, gridId, numGrids, prevGridSum, allGridSum, firstDevice);
if (!command) {
return hipErrorOutOfMemory;
}
// Capture the kernel arguments
if (CL_SUCCESS != command->captureAndValidate()) {
delete command;
return hipErrorOutOfMemory;
}
command->enqueue();
if(startEvent != nullptr) {
eStart->addMarker(queue, command);
command->retain();
}
if(stopEvent != nullptr) {
eStop->addMarker(queue, command);
command->retain();
}
command->release();
return hipSuccess;
}
hipError_t hipModuleLaunchKernel(hipFunction_t f,
uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ,
uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ,
uint32_t sharedMemBytes, hipStream_t hStream,
void **kernelParams, void **extra)
{
HIP_INIT_API(hipModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ,
blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream,
kernelParams, extra);
HIP_RETURN(ihipModuleLaunchKernel(f, gridDimX * blockDimX, gridDimY * blockDimY, gridDimZ * blockDimZ,
blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream, kernelParams, extra, nullptr, nullptr));
}
hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t localWorkSizeX, uint32_t localWorkSizeY,
uint32_t localWorkSizeZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags)
{
HIP_INIT_API(NONE, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ,
localWorkSizeX, localWorkSizeY, localWorkSizeZ,
sharedMemBytes, hStream,
kernelParams, extra, startEvent, stopEvent, flags);
HIP_RETURN(ihipModuleLaunchKernel(f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY,
localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags));
}
hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX,
uint32_t gridDimY, uint32_t gridDimZ,
uint32_t blockDimX, uint32_t blockDimY,
uint32_t blockDimZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent,
hipEvent_t stopEvent)
{
HIP_INIT_API(NONE, f, gridDimX, gridDimY, gridDimZ,
blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream,
kernelParams, extra, startEvent, stopEvent);
HIP_RETURN(ihipModuleLaunchKernel(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent));
}
hipError_t hipModuleLaunchKernelExt(hipFunction_t f, uint32_t gridDimX,
uint32_t gridDimY, uint32_t gridDimZ,
uint32_t blockDimX, uint32_t blockDimY,
uint32_t blockDimZ, size_t sharedMemBytes,
hipStream_t hStream, void** kernelParams, void** extra,
hipEvent_t startEvent,
hipEvent_t stopEvent)
{
HIP_INIT_API(NONE, f, gridDimX, gridDimY, gridDimZ,
blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream,
kernelParams, extra, startEvent, stopEvent);
HIP_RETURN(ihipModuleLaunchKernel(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ,
sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent));
}
hipError_t hipLaunchCooperativeKernel(const void* f,
dim3 gridDim, dim3 blockDim,
void **kernelParams, uint32_t sharedMemBytes, hipStream_t hStream)
{
HIP_INIT_API(hipLaunchCooperativeKernel, f, gridDim, blockDim,
sharedMemBytes, hStream);
int deviceId = ihipGetDevice();
hipFunction_t func = PlatformState::instance().getFunc(f, deviceId);
if (func == nullptr) {
HIP_RETURN(hipErrorInvalidDeviceFunction);
}
HIP_RETURN(ihipModuleLaunchKernel(func, gridDim.x * blockDim.x, gridDim.y * blockDim.y, gridDim.z * blockDim.z,
blockDim.x, blockDim.y, blockDim.z,
sharedMemBytes, hStream, kernelParams, nullptr, nullptr, nullptr, 0,
amd::NDRangeKernelCommand::CooperativeGroups));
}
hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices, unsigned int flags, uint32_t extFlags)
{
int numActiveGPUs = 0;
ihipDeviceGetCount(&numActiveGPUs);
if ((numDevices > numActiveGPUs) || (launchParamsList == nullptr)) {
return hipErrorInvalidValue;
}
hipError_t result = hipErrorUnknown;
uint64_t allGridSize = 0;
std::vector<const amd::Device*> mgpu_list(numDevices);
for (int i = 0; i < numDevices; ++i) {
const hipLaunchParams& launch = launchParamsList[i];
allGridSize += launch.gridDim.x * launch.gridDim.y * launch.gridDim.z;
// Make sure block dimensions are valid
if (0 == launch.blockDim.x * launch.blockDim.y * launch.blockDim.z) {
return hipErrorInvalidConfiguration;
}
if (launch.stream != nullptr) {
// Validate devices to make sure it dosn't have duplicates
amd::HostQueue* queue = reinterpret_cast<hip::Stream*>(launch.stream)->asHostQueue();
auto device = &queue->vdev()->device();
for (int j = 0; j < numDevices; ++j) {
if (mgpu_list[j] == device) {
return hipErrorInvalidDevice;
}
}
mgpu_list[i] = device;
} else {
return hipErrorInvalidResourceHandle;
}
}
uint64_t prevGridSize = 0;
uint32_t firstDevice = 0;
// Sync the execution streams on all devices
if ((flags & hipCooperativeLaunchMultiDeviceNoPreSync) == 0) {
for (int i = 0; i < numDevices; ++i) {
amd::HostQueue* queue =
reinterpret_cast<hip::Stream*>(launchParamsList[i].stream)->asHostQueue();
queue->finish();
}
}
for (int i = 0; i < numDevices; ++i) {
const hipLaunchParams& launch = launchParamsList[i];
amd::HostQueue* queue = reinterpret_cast<hip::Stream*>(launch.stream)->asHostQueue();
hipFunction_t func = nullptr;
// The order of devices in the launch may not match the order in the global array
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);
// Save ROCclr index of the first device in the launch
if (i == 0) {
firstDevice = queue->vdev()->device().index();
}
break;
}
}
if (func == nullptr) {
result = hipErrorInvalidDeviceFunction;
HIP_RETURN(result);
}
result = ihipModuleLaunchKernel(func,
launch.gridDim.x * launch.blockDim.x,
launch.gridDim.y * launch.blockDim.y,
launch.gridDim.z * launch.blockDim.z,
launch.blockDim.x, launch.blockDim.y, launch.blockDim.z,
launch.sharedMem, launch.stream, launch.args, nullptr, nullptr, nullptr,
flags, extFlags, i, numDevices, prevGridSize, allGridSize, firstDevice);
if (result != hipSuccess) {
break;
}
prevGridSize += launch.gridDim.x * launch.gridDim.y * launch.gridDim.z;
}
// Sync the execution streams on all devices
if ((flags & hipCooperativeLaunchMultiDeviceNoPostSync) == 0) {
for (int i = 0; i < numDevices; ++i) {
amd::HostQueue* queue =
reinterpret_cast<hip::Stream*>(launchParamsList[i].stream)->asHostQueue();
queue->finish();
}
}
return result;
}
hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices, unsigned int flags)
{
HIP_INIT_API(hipLaunchCooperativeKernelMultiDevice, launchParamsList, numDevices, flags);
return ihipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags,
(amd::NDRangeKernelCommand::CooperativeGroups |
amd::NDRangeKernelCommand::CooperativeMultiDeviceGroups));
}
hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
int numDevices, unsigned int flags) {
HIP_INIT_API(hipExtLaunchMultiKernelMultiDevice, launchParamsList, numDevices, flags);
return ihipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags, 0);
}
hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name) {
HIP_INIT_API(hipModuleGetTexRef, texRef, hmod, name);
/* input args check */
if ((texRef == nullptr) || (name == nullptr)) {
HIP_RETURN(hipErrorInvalidValue);
}
/* Get address and size for the global symbol */
if (!PlatformState::instance().getTexRef(name, hmod, texRef)) {
DevLogPrintfError("Cannot get texRef for name: %s at module:0x%x \n",
name, hmod);
HIP_RETURN(hipErrorNotFound);
}
// Texture references created by HIP driver API
// have the default read mode set to normalized float.
(*texRef)->readMode = hipReadModeNormalizedFloat;
HIP_RETURN(hipSuccess);
}
rocclr/hip_peer.cpp
Executable
+127
View File
@@ -0,0 +1,127 @@
/* Copyright (c) 2015-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"
hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx) {
HIP_INIT_API(NONE, canAccessPeer, thisCtx, peerCtx);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipMemcpyPeer(void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t srcCtx,
size_t sizeBytes) {
HIP_INIT_API(NONE, dst, dstCtx, src, srcCtx, sizeBytes);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipMemcpyPeerAsync(void* dst, hipCtx_t dstDevice, const void* src, hipCtx_t srcDevice,
size_t sizeBytes, hipStream_t stream) {
HIP_INIT_API(NONE, dst, dstDevice, src, srcDevice, sizeBytes, stream);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
hipError_t canAccessPeer(int* canAccessPeer, int deviceId, int peerDeviceId){
amd::Device* device = nullptr;
amd::Device* peer_device = nullptr;
if (canAccessPeer == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
/* Peer cannot be self */
if (deviceId == peerDeviceId) {
*canAccessPeer = 0;
HIP_RETURN(hipSuccess);
}
/* Cannot exceed the max number of devices */
if (static_cast<size_t>(deviceId) >= g_devices.size()
|| static_cast<size_t>(peerDeviceId) >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
}
device = g_devices[deviceId]->devices()[0];
peer_device = g_devices[peerDeviceId]->devices()[0];
*canAccessPeer = static_cast<int>(std::find(device->p2pDevices_.begin(),
device->p2pDevices_.end(), as_cl(peer_device))
!= device->p2pDevices_.end());
HIP_RETURN(hipSuccess);
}
hipError_t hipDeviceCanAccessPeer(int* canAccess, int deviceId, int peerDeviceId) {
HIP_INIT_API(hipDeviceCanAccessPeer, canAccess, deviceId, peerDeviceId);
HIP_RETURN(canAccessPeer(canAccess, deviceId, peerDeviceId));
}
hipError_t hipDeviceDisablePeerAccess(int peerDeviceId) {
HIP_INIT_API(hipDeviceDisablePeerAccess, peerDeviceId);
int deviceId = hip::getCurrentDevice()->deviceId();
int canAccess = 0;
if ((hipSuccess != canAccessPeer(&canAccess, deviceId, peerDeviceId)) || (canAccess == 0)) {
HIP_RETURN(hipErrorInvalidDevice);
}
HIP_RETURN(hip::getCurrentDevice()->DisablePeerAccess(peerDeviceId));
}
hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags) {
HIP_INIT_API(hipDeviceEnablePeerAccess, peerDeviceId, flags);
int deviceId = hip::getCurrentDevice()->deviceId();
int canAccess = 0;
if (flags != 0) {
HIP_RETURN(hipErrorInvalidValue);
}
if ((hipSuccess != canAccessPeer(&canAccess, deviceId, peerDeviceId)) || (canAccess == 0)) {
HIP_RETURN(hipErrorInvalidDevice);
}
HIP_RETURN(hip::getCurrentDevice()->EnablePeerAccess(peerDeviceId));
}
hipError_t hipMemcpyPeer(void* dst, int dstDevice, const void* src, int srcDevice,
size_t sizeBytes) {
HIP_INIT_API(hipMemcpyPeer, dst, dstDevice, src, srcDevice, sizeBytes);
HIP_RETURN(hipMemcpy(dst, src, sizeBytes, hipMemcpyDeviceToDevice));
}
hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int srcDevice,
size_t sizeBytes, hipStream_t stream) {
HIP_INIT_API(hipMemcpyPeerAsync, dst, dstDevice, src, srcDevice, sizeBytes, stream);
HIP_RETURN(hipMemcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToDevice, stream));
}
hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) {
HIP_INIT_API(hipCtxEnablePeerAccess, peerCtx, flags);
HIP_RETURN(hipSuccess);
}
hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx) {
HIP_INIT_API(hipCtxDisablePeerAccess, peerCtx);
HIP_RETURN(hipSuccess);
}
rocclr/hip_platform.cpp
Executable
+1347
View File
File diff suppressed because it is too large Load Diff
rocclr/hip_platform.hpp
+29
View File
@@ -0,0 +1,29 @@
/* Copyright (c) 2015-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. */
#pragma once
#include "device/device.hpp"
namespace hip_impl {
hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(
int* numBlocks, int* numGrids,
const amd::Device& device, hipFunction_t func, int blockSize,
size_t dynamicSMemSize, bool bCalcPotentialBlkSz);
}
rocclr/hip_prof_api.h
+250
View File
@@ -0,0 +1,250 @@
/* Copyright (c) 2019-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. */
#ifndef HIP_SRC_HIP_PROF_API_H
#define HIP_SRC_HIP_PROF_API_H
#include <atomic>
#include <iostream>
#include <mutex>
#if USE_PROF_API
#include "hip/hcc_detail/hip_prof_str.h"
#include "platform/prof_protocol.h"
// HIP API callbacks spawner object macro
#define HIP_CB_SPAWNER_OBJECT(CB_ID) \
api_callbacks_spawner_t<HIP_API_ID_##CB_ID> __api_tracer; \
{ \
hip_api_data_t* api_data = __api_tracer.get_api_data_ptr(); \
if (api_data != NULL) { \
hip_api_data_t& api_data_ref = *api_data; \
INIT_CB_ARGS_DATA(CB_ID, api_data_ref); \
__api_tracer.call(); \
} \
}
static const uint32_t HIP_DOMAIN_ID = ACTIVITY_DOMAIN_HIP_API;
typedef activity_record_t hip_api_record_t;
typedef activity_rtapi_callback_t hip_api_callback_t;
typedef activity_sync_callback_t hip_act_callback_t;
class api_callbacks_table_t {
public:
typedef std::mutex mutex_t;
typedef hip_api_record_t record_t;
typedef hip_api_callback_t fun_t;
typedef hip_act_callback_t act_t;
// HIP API callbacks table
struct hip_cb_table_entry_t {
volatile std::atomic<bool> sync;
volatile std::atomic<uint32_t> sem;
act_t act;
void* a_arg;
fun_t fun;
void* arg;
};
struct hip_cb_table_t {
hip_cb_table_entry_t arr[HIP_API_ID_NUMBER];
};
api_callbacks_table_t() {
memset(&callbacks_table_, 0, sizeof(callbacks_table_));
}
bool set_activity(uint32_t id, act_t fun, void* arg) {
std::lock_guard<mutex_t> lock(mutex_);
bool ret = true;
if (id < HIP_API_ID_NUMBER) {
cb_sync(id);
callbacks_table_.arr[id].act = fun;
callbacks_table_.arr[id].a_arg = arg;
enabled_ = true;
cb_release(id);
} else {
ret = false;
}
return ret;
}
bool set_callback(uint32_t id, fun_t fun, void* arg) {
std::lock_guard<mutex_t> lock(mutex_);
bool ret = true;
if (id < HIP_API_ID_NUMBER) {
cb_sync(id);
callbacks_table_.arr[id].fun = fun;
callbacks_table_.arr[id].arg = arg;
enabled_ = true;
cb_release(id);
} else {
ret = false;
}
return ret;
}
void set_enabled(const bool& enabled) {
enabled_ = enabled;
}
inline hip_cb_table_entry_t& entry(const uint32_t& id) {
return callbacks_table_.arr[id];
}
inline void sem_sync(const uint32_t& id) {
sem_increment(id);
if (entry(id).sync.load() == true) sync_wait(id);
}
inline void sem_release(const uint32_t& id) {
sem_decrement(id);
}
inline bool is_enabled() const {
return enabled_;
}
private:
inline void cb_sync(const uint32_t& id) {
entry(id).sync.store(true);
while (entry(id).sem.load() != 0) {}
}
inline void cb_release(const uint32_t& id) {
entry(id).sync.store(false);
}
inline void sem_increment(const uint32_t& id) {
const uint32_t prev = entry(id).sem.fetch_add(1);
if (prev == UINT32_MAX) {
std::cerr << "sem overflow id = " << id << std::endl << std::flush;
abort();
}
}
inline void sem_decrement(const uint32_t& id) {
const uint32_t prev = entry(id).sem.fetch_sub(1);
if (prev == 0) {
std::cerr << "sem corrupted id = " << id << std::endl << std::flush;
abort();
}
}
void sync_wait(const uint32_t& id) {
sem_decrement(id);
while (entry(id).sync.load() == true) {}
sem_increment(id);
}
mutex_t mutex_;
hip_cb_table_t callbacks_table_;
bool enabled_;
};
extern api_callbacks_table_t callbacks_table;
template <int cid_>
class api_callbacks_spawner_t {
public:
api_callbacks_spawner_t() :
api_data_(NULL)
{
if (!is_enabled()) return;
if (cid_ >= HIP_API_ID_NUMBER) {
fprintf(stderr, "HIP %s bad id %d\n", __FUNCTION__, cid_);
abort();
}
callbacks_table.sem_sync(cid_);
hip_act_callback_t act = entry(cid_).act;
if (act != NULL) api_data_ = (hip_api_data_t*) act(cid_, NULL, NULL, NULL);
}
void call() {
hip_api_callback_t fun = entry(cid_).fun;
void* arg = entry(cid_).arg;
if (fun != NULL) {
fun(HIP_DOMAIN_ID, cid_, api_data_, arg);
api_data_->phase = ACTIVITY_API_PHASE_EXIT;
}
}
~api_callbacks_spawner_t() {
if (!is_enabled()) return;
if (api_data_ != NULL) {
hip_api_callback_t fun = entry(cid_).fun;
void* arg = entry(cid_).arg;
hip_act_callback_t act = entry(cid_).act;
void* a_arg = entry(cid_).a_arg;
if (fun != NULL) fun(HIP_DOMAIN_ID, cid_, api_data_, arg);
if (act != NULL) act(cid_, NULL, NULL, a_arg);
}
callbacks_table.sem_release(cid_);
}
hip_api_data_t* get_api_data_ptr() {
return api_data_;
}
bool is_enabled() const {
return callbacks_table.is_enabled();
}
private:
inline api_callbacks_table_t::hip_cb_table_entry_t& entry(const uint32_t& id) {
return callbacks_table.entry(id);
}
hip_api_data_t* api_data_;
};
template <>
class api_callbacks_spawner_t<HIP_API_ID_NUMBER> {
public:
api_callbacks_spawner_t() {}
void call() {}
hip_api_data_t* get_api_data_ptr() { return NULL; }
bool is_enabled() const { return false; }
};
#else
#define HIP_CB_SPAWNER_OBJECT(x) do {} while(0)
class api_callbacks_table_t {
public:
typedef void* act_t;
typedef void* fun_t;
bool set_activity(uint32_t id, act_t fun, void* arg) { return false; }
bool set_callback(uint32_t id, fun_t fun, void* arg) { return false; }
};
#endif
#endif // HIP_SRC_HIP_PROF_API_H
rocclr/hip_prof_gen.py
Executable
+617
View File
@@ -0,0 +1,617 @@
#!/usr/bin/python
# Copyright (c) 2019-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 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.
import os, sys, re
PROF_HEADER = "hip_prof_str.h"
OUTPUT = PROF_HEADER
REC_MAX_LEN = 1024
# Recursive sources processing
recursive_mode = 0
# HIP_INIT_API macro patching
hip_patch_mode = 0
# API matching types check
types_check_mode = 0
# Private API check
private_check_mode = 0
# Messages and errors controll
verbose = 0
errexit = 0
inp_file = 'none'
line_num = -1
# Verbose message
def message(msg):
if verbose: sys.stdout.write(msg + '\n')
# Fatal error termination
def error(msg):
if line_num != -1:
msg += ", file '" + inp_file + "', line (" + str(line_num) + ")"
if errexit:
msg = " Error: " + msg
else:
msg = " Warning: " + msg
sys.stdout.write(msg + '\n')
sys.stderr.write(sys.argv[0] + msg +'\n')
def fatal(msg):
error(msg)
sys.exit(1)
#############################################################
# Normalizing API name
def filtr_api_name(name):
name = re.sub(r'\s*$', r'', name);
return name
def filtr_api_decl(record):
record = re.sub("\s__dparm\([^\)]*\)", '', record);
record = re.sub("\(void\*\)", '', record);
return record
# Normalizing API arguments
def filtr_api_args(args_str):
args_str = re.sub(r'^\s*', r'', args_str);
args_str = re.sub(r'\s*$', r'', args_str);
args_str = re.sub(r'\s*,\s*', r',', args_str);
args_str = re.sub(r'\s+', r' ', args_str);
args_str = re.sub(r'\s*(\*+)\s*', r'\1 ', args_str);
args_str = re.sub(r'(enum|struct) ', '', args_str);
return args_str
# Normalizing types
def norm_api_types(type_str):
type_str = re.sub(r'uint32_t', r'unsigned int', type_str)
type_str = re.sub(r'^unsigned$', r'unsigned int', type_str)
return type_str
# Creating a list of arguments [(type, name), ...]
def list_api_args(args_str):
args_str = filtr_api_args(args_str)
args_list = []
if args_str != '':
for arg_pair in args_str.split(','):
if arg_pair == 'void': continue
arg_pair = re.sub(r'\s*=\s*\S+$','', arg_pair);
m = re.match("^(.*)\s(\S+)$", arg_pair);
if m:
arg_type = norm_api_types(m.group(1))
arg_name = m.group(2)
args_list.append((arg_type, arg_name))
else:
fatal("bad args: args_str: '" + args_str + "' arg_pair: '" + arg_pair + "'")
return args_list;
# Creating arguments string "type0, type1, ..."
def filtr_api_types(args_str):
args_list = list_api_args(args_str)
types_str = ''
for arg_tuple in args_list:
types_str += arg_tuple[0] + ', '
return types_str
# Creating options list [opt0, opt1, ...]
def filtr_api_opts(args_str):
args_list = list_api_args(args_str)
opts_list = []
for arg_tuple in args_list:
opts_list.append(arg_tuple[1])
return opts_list
#############################################################
# Parsing API header
# hipError_t hipSetupArgument(const void* arg, size_t size, size_t offset);
def parse_api(inp_file_p, out):
global inp_file
global line_num
inp_file = inp_file_p
beg_pattern = re.compile("^(hipError_t|const char\s*\*)\s+([^\(]+)\(");
api_pattern = re.compile("^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)");
end_pattern = re.compile("Texture");
hidden_pattern = re.compile(r'__attribute__\(\(visibility\("hidden"\)\)\)')
nms_open_pattern = re.compile(r'namespace hip_impl {')
nms_close_pattern = re.compile(r'}')
inp = open(inp_file, 'r')
found = 0
hidden = 0
nms_level = 0;
record = ""
line_num = -1
for line in inp.readlines():
record += re.sub(r'^\s+', r' ', line[:-1])
line_num += 1
if len(record) > REC_MAX_LEN:
fatal("bad record \"" + record + "\"")
m = beg_pattern.match(line)
if m:
name = m.group(2)
if hidden != 0:
message("api: " + name + " - hidden")
elif nms_level != 0:
message("api: " + name + " - hip_impl")
else:
message("api: " + name)
found = 1
if found != 0:
record = re.sub("\s__dparm\([^\)]*\)", '', record);
m = api_pattern.match(record)
if m:
found = 0
if end_pattern.search(record): break
api_name = filtr_api_name(m.group(2))
api_args = m.group(3)
if not api_name in out:
out[api_name] = api_args
else: continue
hidden = 0
if hidden_pattern.match(line): hidden = 1
if nms_open_pattern.match(line): nms_level += 1
if (nms_level > 0) and nms_close_pattern.match(line): nms_level -= 1
if nms_level < 0:
fatal("nms level < 0")
record = ""
inp.close()
line_num = -1
#############################################################
# Parsing API implementation
# hipError_t hipSetupArgument(const void* arg, size_t size, size_t offset) {
# HIP_INIT_API(hipSetupArgument, arg, size, offset);
# inp_file - input implementation source file
# api_map - input public API map [<api name>] => <api args>
# out - output map [<api name>] => [opt0, opt1, ...]
def parse_content(inp_file_p, api_map, out):
global hip_patch_mode
global types_check_mode
global private_check_mode
global inp_file
global line_num
inp_file = inp_file_p
# API method begin pattern
beg_pattern = re.compile("^(hipError_t|const char\s*\*)\s+[^\(]+\(");
# API declaration pattern
decl_pattern = re.compile("^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)\s*;");
# API definition pattern
api_pattern = re.compile("^(hipError_t|const char\s*\*)\s+([^\(]+)\(([^\)]*)\)\s*{");
# API init macro pattern
init_pattern = re.compile("(^\s*HIP_INIT_API\s*)\((([^,]+)(,.*|)|)(\);|,)\s*$");
# Open input file
inp = open(inp_file, 'r')
# API name
api_name = ""
# Valid public API found flag
api_valid = 0
# Input file patched content
content = ''
# Sub content for found API defiition
sub_content = ''
# Current record, accumulating several API definition related lines
record = ''
# Current input file line number
line_num = -1
# API beginning found flag
found = 0
# Reading input file
for line in inp.readlines():
# Accumulating record
record += re.sub(r'^\s+', r' ', line[:-1])
line_num += 1
if len(record) > REC_MAX_LEN:
fatal("bad record \"" + record + "\"")
break;
# Looking for API begin
if found == 0:
if beg_pattern.match(record):
found = 1
record = filtr_api_decl(record)
# Matching API declaration
if found == 1:
if decl_pattern.match(record):
found = 0
# Matching API definition
if found == 1:
m = api_pattern.match(record)
# Checking if complete API matched
if m:
found = 2
api_name = filtr_api_name(m.group(2))
# Checking if API name is in the API map
if (private_check_mode == 0) or (api_name in api_map):
if not api_name in api_map: api_map[api_name] = ''
# Getting API arguments
api_args = m.group(3)
# Getting etalon arguments from the API map
eta_args = api_map[api_name]
if eta_args == '':
eta_args = api_args
api_map[api_name] = eta_args
# Normalizing API arguments
api_types = filtr_api_types(api_args)
# Normalizing etalon arguments
eta_types = filtr_api_types(eta_args)
if (api_types == eta_types) or ((types_check_mode == 0) and (not api_name in out)):
# API is already found and not is mismatched
if (api_name in out):
fatal("API redefined \"" + api_name + "\", record \"" + record + "\"")
# Set valid public API found flag
api_valid = 1
# Set output API map with API arguments list
out[api_name] = filtr_api_opts(api_args)
# Register missmatched API methods
else:
# Warning about mismatched API, possible non public overloaded version
api_diff = '\t\t' + inp_file + " line(" + str(line_num) + ")\n\t\tapi: " + api_types + "\n\t\teta: " + eta_types
message("\t" + api_name + ' args mismatch:\n' + api_diff + '\n')
if hip_patch_mode != 0:
# Looking for INIT macro
m = init_pattern.match(line)
if m:
if api_valid == 0: api_name = 'NONE'
if api_name == m.group(3):
if hip_patch_mode == 1: hip_patch_mode = 0
else: fatal("patching failed")
else:
hip_patch_mode = 2
init_args = m.group(2)
if init_args != '': init_args = ', ' + init_args
line = m.group(1) + '(' + api_name + init_args + m.group(5) + '\n'
non_public_api = 0
# API found action
if found == 2:
# Looking for INIT macro
m = init_pattern.match(line)
if m:
found = 0
non_public_api = 0
if api_valid == 1:
api_valid = 0
message("\t" + api_name)
else:
non_public_api = 1
if non_public_api == 1:
# Registering dummy API for non public API if the name in INIT is not NONE
init_name = m.group(3)
# Ignore if it is initialized as NONE
if init_name != 'NONE':
# Check if init name matching API name
if init_name != api_name:
fatal("init name mismatch: '" + init_name + "' <> '" + api_name + "'")
# If init name is not in public API map then it is private API
# else it was not identified and will be checked on finish
if not init_name in api_map:
if init_name in out:
fatal("API reinit \"" + api_name + "\", record \"" + record + "\"")
out[init_name] = []
elif re.search('}', line):
found = 0
# Expect INIT macro for valid public API
# Removing and registering non-conformant APIs with missing HIP_INIT macro
if api_valid == 1:
api_valid = 0
if api_name in out:
del out[api_name]
del api_map[api_name]
# Registering non-conformant APIs
out['.' + api_name] = 1
else:
fatal("API is not in out \"" + api_name + "\", record \"" + record + "\"")
if found != 1: record = ""
content += line
inp.close()
line_num = -1
if len(out) != 0:
return content
else:
return ''
# src path walk
def parse_src(api_map, src_path, src_patt, out):
global recursive_mode
pattern = re.compile(src_patt)
src_path = re.sub(r'\s', '', src_path)
for src_dir in src_path.split(':'):
message("Parsing " + src_dir + " for '" + src_patt + "'")
for root, dirs, files in os.walk(src_dir):
for fnm in files:
if pattern.search(fnm):
file = root + '/' + fnm
message(file)
content = parse_content(file, api_map, out);
if (hip_patch_mode != 0) and (content != ''):
f = open(file, 'w')
f.write(content)
f.close()
if recursive_mode == 0: break
#############################################################
# Generating profiling primitives header
# api_map - public API map [<api name>] => [(type, name), ...]
# opts_map - opts map [<api name>] => [opt0, opt1, ...]
def generate_prof_header(f, api_map, opts_map):
# Private API list
priv_lst = []
f.write('// automatically generated sources\n')
f.write('#ifndef _HIP_PROF_STR_H\n');
f.write('#define _HIP_PROF_STR_H\n');
f.write('#define HIP_PROF_VER 1\n')
# Generating dummy macro for non-public API
f.write('\n// Dummy API primitives\n')
f.write('#define INIT_NONE_CB_ARGS_DATA(cb_data) {};\n')
for name in opts_map:
if not name in api_map:
opts_lst = opts_map[name]
if len(opts_lst) != 0:
fatal("bad dummy API \"" + name + "\", args: " + str(opts_lst))
f.write('#define INIT_'+ name + '_CB_ARGS_DATA(cb_data) {};\n')
priv_lst.append(name)
for name in priv_lst:
message("Private: " + name)
# Generating the callbacks ID enumaration
f.write('\n// HIP API callbacks ID enumaration\n')
f.write('enum hip_api_id_t {\n')
cb_id = 0
for name in api_map.keys():
f.write(' HIP_API_ID_' + name + ' = ' + str(cb_id) + ',\n')
cb_id += 1
f.write(' HIP_API_ID_NUMBER = ' + str(cb_id) + ',\n')
f.write('\n')
f.write(' HIP_API_ID_NONE = HIP_API_ID_NUMBER,\n')
for name in priv_lst:
f.write(' HIP_API_ID_' + name + ' = HIP_API_ID_NUMBER,\n')
f.write('};\n')
# Generating the callbacks ID enumaration
f.write('\n// Return HIP API string\n')
f.write('inline const char* hip_api_name(const uint32_t id) {\n')
f.write(' switch(id) {\n')
for name in api_map.keys():
f.write(' case HIP_API_ID_' + name + ': return "' + name + '";\n')
f.write(' };\n')
f.write(' return "unknown";\n')
f.write('};\n')
# Generating the callbacks data structure
f.write('\n// HIP API callbacks data structure\n')
f.write(
'typedef struct hip_api_data_s {\n' +
' uint64_t correlation_id;\n' +
' uint32_t phase;\n' +
' union {\n'
)
for name, args in api_map.items():
if len(args) != 0:
f.write(' struct {\n')
for arg_tuple in args:
if arg_tuple[0] == "hipLimit_t":
f.write(' enum ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n')
else:
f.write(' ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n')
f.write(' } ' + name + ';\n')
f.write(
' } args;\n' +
'} hip_api_data_t;\n'
)
# Generating the callbacks args data filling macros
f.write('\n// HIP API callbacks args data filling macros\n')
for name, args in api_map.items():
f.write('// ' + name + str(args) + '\n')
f.write('#define INIT_' + name + '_CB_ARGS_DATA(cb_data) { \\\n')
if name in opts_map:
opts_list = opts_map[name]
if len(args) != len(opts_list):
fatal("\"" + name + "\" API args and opts mismatch, args: " + str(args) + ", opts: " + str(opts_list))
# API args iterating:
# type is args[<ind>][0]
# name is args[<ind>][1]
for ind in range(0, len(args)):
arg_tuple = args[ind]
arg_type = arg_tuple[0]
fld_name = arg_tuple[1]
arg_name = opts_list[ind]
f.write(' cb_data.args.' + name + '.' + fld_name + ' = (' + arg_type + ')' + arg_name + '; \\\n')
f.write('};\n')
f.write('#define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data)\n')
# Generating the method for the API string, name and parameters
f.write('\n')
f.write('#if HIP_PROF_HIP_API_STRING\n')
f.write('#include <sstream>\n');
f.write('#include <string>\n');
f.write('// HIP API string method, method name and parameters\n')
f.write('const char* hipApiString(hip_api_id_t id, const hip_api_data_t* data) {\n')
f.write(' std::ostringstream oss;\n')
f.write(' switch (id) {\n')
for name, args in api_map.items():
f.write(' case HIP_API_ID_' + name + ':\n')
f.write(' oss << "' + name + '("')
for ind in range(0, len(args)):
arg_tuple = args[ind]
arg_name = arg_tuple[1]
if ind != 0: f.write(' << ","')
f.write('\n << " ' + arg_name + '=" << data->args.' + name + '.' + arg_name)
f.write('\n << ")";\n')
f.write(' break;\n')
f.write(' default: oss << "unknown";\n')
f.write(' };\n')
f.write(' return strdup(oss.str().c_str());\n')
f.write('};\n')
f.write('#endif // HIP_PROF_HIP_API_STRING\n')
f.write('#endif // _HIP_PROF_STR_H\n');
#############################################################
# main
while len(sys.argv) > 1:
if not re.match(r'-', sys.argv[1]): break
if (sys.argv[1] == '-v'):
verbose = 1
sys.argv.pop(1)
if (sys.argv[1] == '-r'):
recursive_mode = 1
sys.argv.pop(1)
if (sys.argv[1] == '-t'):
types_check_mode = 1
sys.argv.pop(1)
if (sys.argv[1] == '--priv'):
private_check_mode = 1
sys.argv.pop(1)
if (sys.argv[1] == '-e'):
errexit = 1
sys.argv.pop(1)
if (sys.argv[1] == '-p'):
hip_patch_mode = 1
sys.argv.pop(1)
# Usage
if (len(sys.argv) < 3):
fatal ("Usage: " + sys.argv[0] + " [-v] <input HIP API .h file> <patched srcs path> [<output>]\n" +
" -v - verbose messages\n" +
" -r - process source directory recursively\n" +
" -t - API types matching check\n" +
" --priv - private API check\n" +
" -e - on error exit mode\n" +
" -p - HIP_INIT_API macro patching mode\n" +
"\n" +
" Example:\n" +
" $ " + sys.argv[0] + " -v -p -t --priv ./api/hip/include/hip/hcc_detail/hip_runtime_api.h ./api/hip ./api/hip/include/hip/hcc_detail/hip_prof_str.h");
# API header file given as an argument
src_pat = "\.cpp$"
api_hfile = sys.argv[1]
if not os.path.isfile(api_hfile):
fatal("input file '" + api_hfile + "' not found")
# Srcs directory given as an argument
src_dir = sys.argv[2]
if not os.path.isdir(src_dir):
fatal("src directory " + src_dir + "' not found")
if len(sys.argv) > 3: OUTPUT = sys.argv[3]
# API declaration map
api_map = {
'hipSetupArgument': '',
'hipMalloc3DArray': '',
'hipFuncGetAttribute': '',
'hipMemset3DAsync': '',
'hipKernelNameRef': '',
'hipStreamGetPriority': '',
'hipLaunchByPtr': '',
'hipFreeHost': '',
'hipGetErrorName': '',
'hipMemcpy3DAsync': '',
'hipMemcpyParam2DAsync': '',
'hipArray3DCreate': '',
'hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags': '',
'hipOccupancyMaxPotentialBlockSize': '',
'hipMallocManaged': '',
'hipOccupancyMaxActiveBlocksPerMultiprocessor': '',
'hipGetErrorString': '',
'hipMallocHost': '',
'hipModuleLoadDataEx': '',
'hipGetDeviceProperties': '',
'hipConfigureCall': '',
'hipHccModuleLaunchKernel': '',
'hipExtModuleLaunchKernel': '',
}
# API options map
opts_map = {}
# Parsing API header
parse_api(api_hfile, api_map)
# Parsing sources
parse_src(api_map, src_dir, src_pat, opts_map)
# Checking for non-conformant APIs with missing HIP_INIT macro
for name in list(opts_map.keys()):
m = re.match(r'\.(\S*)', name)
if m:
message("Init missing: " + m.group(1))
del opts_map[name]
# Converting api map to map of lists
# Checking for not found APIs
not_found = 0
if len(opts_map) != 0:
for name in api_map.keys():
args_str = api_map[name];
api_map[name] = list_api_args(args_str)
if not name in opts_map:
error("implementation not found: " + name)
not_found += 1
if not_found != 0:
error(str(not_found) + " API calls missing in interception layer")
# The output subdirectory seems to exist or not depending on the
# version of cmake.
output_dir = os.path.dirname(OUTPUT)
if not os.path.exists(output_dir):
os.makedirs(output_dir)
# Generating output header file
with open(OUTPUT, 'w') as f:
generate_prof_header(f, api_map, opts_map)
# Successfull exit
sys.exit(0)
rocclr/hip_profile.cpp
+40
View File
@@ -0,0 +1,40 @@
/* Copyright (c) 2015-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"
hipError_t hipProfilerStart() {
HIP_INIT_API(hipProfilerStart);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipProfilerStop() {
HIP_INIT_API(hipProfilerStop);
assert(0 && "Unimplemented");
HIP_RETURN(hipErrorNotSupported);
}
rocclr/hip_rtc.cpp
Executable
+397
View File
@@ -0,0 +1,397 @@
/* Copyright (c) 2015-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 "hiprtc_internal.hpp"
#include <hip/hiprtc.h>
#include "platform/program.hpp"
namespace hiprtc {
thread_local hiprtcResult g_lastRtcError = HIPRTC_SUCCESS;
}
class ProgramState {
amd::Monitor lock_;
private:
static ProgramState* programState_;
ProgramState() : lock_("Guards program state") {}
~ProgramState() {}
public:
std::unordered_map<amd::Program*,
std::pair<std::vector<std::string>, std::vector<std::string>>> progHeaders_;
std::map<std::string, std::pair<std::string, std::string>> nameExpresssion_;
static ProgramState& instance();
void createProgramHeaders(amd::Program* program, int numHeaders,
const char** headers, const char** headerNames);
void getProgramHeaders(amd::Program* program, int* numHeaders, char** headers, char ** headerNames);
uint32_t addNameExpression(const char* name_expression);
char* getLoweredName(const char* name_expression);
};
ProgramState* ProgramState::programState_ = nullptr;
ProgramState& ProgramState::instance() {
if (programState_ == nullptr) {
programState_ = new ProgramState;
}
return *programState_;
}
void ProgramState::createProgramHeaders(amd::Program* program, int numHeaders,
const char** headers, const char** headerNames) {
amd::ScopedLock lock(lock_);
std::vector<std::string> vHeaderNames;
std::vector<std::string> vHeaders;
for (auto i = 0; i != numHeaders; ++i) {
vHeaders.emplace_back(headers[i]);
vHeaderNames.emplace_back(headerNames[i]);
progHeaders_[program] = std::make_pair(std::move(vHeaders), std::move(vHeaderNames));
}
}
void ProgramState::getProgramHeaders(amd::Program* program, int* numHeaders,
char** headers, char ** headerNames) {
amd::ScopedLock lock(lock_);
const auto it = progHeaders_.find(program);
if (it != progHeaders_.cend()) {
*numHeaders = it->second.first.size();
*headers = reinterpret_cast<char*>(it->second.first.data());
*headerNames = reinterpret_cast<char*>(it->second.second.data());
}
}
uint32_t ProgramState::addNameExpression(const char* name_expression) {
amd::ScopedLock lock(lock_);
// Strip clean of any '(' or ')' or '&'
std::string strippedName(name_expression);
if (strippedName.back() == ')') {
strippedName.pop_back();
strippedName.erase(0, strippedName.find('('));
}
if (strippedName.front() == '&') {
strippedName.erase(0, 1);
}
auto it = nameExpresssion_.find(name_expression);
if (it == nameExpresssion_.end()) {
nameExpresssion_.insert(std::pair<std::string, std::pair<std::string, std::string>>
(name_expression, std::make_pair(strippedName,"")));
}
return nameExpresssion_.size();
}
char* demangle(const char* loweredName) {
if (!loweredName) {
return nullptr;
}
#if __linux__
int status = 0;
char* demangledName = DEMANGLE(loweredName, nullptr, nullptr, &status);
if (status != 0) {
DevLogPrintfError("Cannot demangle loweredName: %s \n", loweredName);
return nullptr;
}
#elif defined(_WIN32)
char* demangledName = (char*)malloc(UNDECORATED_SIZE);
if (!UnDecorateSymbolName(loweredName, demangledName,
UNDECORATED_SIZE/ sizeof(*demangledName), UNDNAME_COMPLETE))
{
free(demangledName);
DevLogPrintfError("Cannot undecorate loweredName: %s demangledName: %s \n",
loweredName, demangledName);
return nullptr;
}
#else
#error "Only Linux and Windows are supported"
#endif // __linux__
return demangledName;
}
static std::string handleMangledName(std::string name) {
std::string loweredName;
char* demangled = demangle(name.c_str());
loweredName.assign(demangled == nullptr ? std::string() : demangled);
free(demangled);
if (loweredName.empty()) {
return name;
}
if (loweredName.find(".kd") != std::string::npos) {
return {};
}
if (loweredName.find("void ") == 0) {
loweredName.erase(0, strlen("void "));
}
auto dx{loweredName.find_first_of("(<")};
if (dx == std::string::npos) {
return loweredName;
}
if (loweredName[dx] == '<') {
uint32_t count = 1;
do {
++dx;
count += (loweredName[dx] == '<') ? 1 : ((loweredName[dx] == '>') ? -1 : 0);
} while (count);
loweredName.erase(++dx);
} else {
loweredName.erase(dx);
}
return loweredName;
}
const char* hiprtcGetErrorString(hiprtcResult x) {
switch (x) {
case HIPRTC_SUCCESS:
return "HIPRTC_SUCCESS";
case HIPRTC_ERROR_OUT_OF_MEMORY:
return "HIPRTC_ERROR_OUT_OF_MEMORY";
case HIPRTC_ERROR_PROGRAM_CREATION_FAILURE:
return "HIPRTC_ERROR_PROGRAM_CREATION_FAILURE";
case HIPRTC_ERROR_INVALID_INPUT:
return "HIPRTC_ERROR_INVALID_INPUT";
case HIPRTC_ERROR_INVALID_PROGRAM:
return "HIPRTC_ERROR_INVALID_PROGRAM";
case HIPRTC_ERROR_INVALID_OPTION:
return "HIPRTC_ERROR_INVALID_OPTION";
case HIPRTC_ERROR_COMPILATION:
return "HIPRTC_ERROR_COMPILATION";
case HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE:
return "HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE";
case HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION:
return "HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION";
case HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION:
return "HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION";
case HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID:
return "HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID";
case HIPRTC_ERROR_INTERNAL_ERROR:
return "HIPRTC_ERROR_INTERNAL_ERROR";
default:
DevLogPrintfError("Invalid HIPRTC error code: %d \n", x);
return nullptr;
};
ShouldNotReachHere();
return nullptr;
}
hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog, const char* src, const char* name,
int numHeaders, const char** headers, const char** headerNames) {
HIPRTC_INIT_API(prog, src, name, numHeaders, headers, headerNames);
if (prog == nullptr) {
HIPRTC_RETURN(HIPRTC_ERROR_INVALID_PROGRAM);
}
if (numHeaders < 0) {
HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT);
}
if (numHeaders && (headers == nullptr || headerNames == nullptr)) {
HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT);
}
amd::Program* program = new amd::Program(*hip::getCurrentDevice()->asContext(), src, amd::Program::HIP);
if (program == NULL) {
HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT);
}
if (CL_SUCCESS != program->addDeviceProgram(*hip::getCurrentDevice()->devices()[0])) {
program->release();
HIPRTC_RETURN(HIPRTC_ERROR_PROGRAM_CREATION_FAILURE);
}
ProgramState::instance().createProgramHeaders(program, numHeaders, headers, headerNames);
*prog = reinterpret_cast<hiprtcProgram>(as_cl(program));
HIPRTC_RETURN(HIPRTC_SUCCESS);
}
hiprtcResult hiprtcCompileProgram(hiprtcProgram prog, int numOptions, const char** options) {
// FIXME[skudchad] Add headers to amd::Program::build and device::Program::build,
// pass the saved from ProgramState to amd::Program::build
HIPRTC_INIT_API(prog, numOptions, options);
amd::Program* program = as_amd(reinterpret_cast<cl_program>(prog));
std::ostringstream ostrstr;
std::vector<const char*> oarr(&options[0], &options[numOptions]);
std::copy(oarr.begin(), oarr.end(), std::ostream_iterator<std::string>(ostrstr, " "));
ostrstr.str().append(" -DHIP_VERSION_MAJOR=9");
ostrstr.str().append(" -DHIP_VERSION_MINOR=0");
std::vector<amd::Device*> devices{hip::getCurrentDevice()->devices()[0]};
if (CL_SUCCESS != program->build(devices, ostrstr.str().c_str(), nullptr, nullptr)) {
HIPRTC_RETURN(HIPRTC_ERROR_COMPILATION);
}
HIPRTC_RETURN(HIPRTC_SUCCESS);
}
hiprtcResult hiprtcAddNameExpression(hiprtcProgram prog, const char* name_expression) {
HIPRTC_INIT_API(prog, name_expression);
if (name_expression == nullptr) {
HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT);
}
amd::Program* program = as_amd(reinterpret_cast<cl_program>(prog));
uint32_t id = ProgramState::instance().addNameExpression(name_expression);
const auto var{"__hiprtc_" + std::to_string(id)};
const auto code{"\nextern \"C\" constexpr auto " + var + " = " + name_expression + ';'};
program->appendToSource(code.c_str());
HIPRTC_RETURN(HIPRTC_SUCCESS);
}
hiprtcResult hiprtcGetLoweredName(hiprtcProgram prog, const char* name_expression,
const char** loweredName) {
HIPRTC_INIT_API(prog, name_expression, loweredName);
if (name_expression == nullptr || loweredName == nullptr) {
HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT);
}
amd::Program* program = as_amd(reinterpret_cast<cl_program>(prog));
device::Program* dev_program
= program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]);
auto it = ProgramState::instance().nameExpresssion_.find(name_expression);
if (it == ProgramState::instance().nameExpresssion_.end()) {
return HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID;
}
std::string strippedName = it->second.first;
std::vector<std::string> mangledNames;
if (!dev_program->getLoweredNames(&mangledNames)) {
HIPRTC_RETURN(HIPRTC_ERROR_COMPILATION);
}
for (auto &name : mangledNames) {
std::string demangledName = handleMangledName(name);
if (demangledName == strippedName) {
it->second.second.assign(name);
}
}
*loweredName = it->second.second.c_str();
HIPRTC_RETURN(HIPRTC_SUCCESS);
}
hiprtcResult hiprtcDestroyProgram(hiprtcProgram* prog) {
HIPRTC_INIT_API(prog);
if (prog == NULL) {
HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT);
}
// Release program. hiprtcProgram is a double pointer so free *prog
amd::Program* program = as_amd(reinterpret_cast<cl_program>(*prog));
program->release();
HIPRTC_RETURN(HIPRTC_SUCCESS);
}
hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* binaryMem) {
HIPRTC_INIT_API(prog, binaryMem);
amd::Program* program = as_amd(reinterpret_cast<cl_program>(prog));
const device::Program::binary_t& binary =
program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0])->binary();
::memcpy(binaryMem, binary.first, binary.second);
HIPRTC_RETURN(HIPRTC_SUCCESS);
}
hiprtcResult hiprtcGetCodeSize(hiprtcProgram prog, size_t* binarySizeRet) {
HIPRTC_INIT_API(prog, binarySizeRet);
amd::Program* program = as_amd(reinterpret_cast<cl_program>(prog));
*binarySizeRet =
program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0])->binary().second;
HIPRTC_RETURN(HIPRTC_SUCCESS);
}
hiprtcResult hiprtcGetProgramLog(hiprtcProgram prog, char* dst) {
HIPRTC_INIT_API(prog, dst);
amd::Program* program = as_amd(reinterpret_cast<cl_program>(prog));
const device::Program* devProgram =
program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]);
auto log = program->programLog() + devProgram->buildLog().c_str();
log.copy(dst, log.size());
dst[log.size()] = '\0';
HIPRTC_RETURN(HIPRTC_SUCCESS);
}
hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram prog, size_t* logSizeRet) {
HIPRTC_INIT_API(prog, logSizeRet);
amd::Program* program = as_amd(reinterpret_cast<cl_program>(prog));
const device::Program* devProgram =
program->getDeviceProgram(*hip::getCurrentDevice()->devices()[0]);
auto log = program->programLog() + devProgram->buildLog().c_str();
*logSizeRet = log.size() + 1;
HIPRTC_RETURN(HIPRTC_SUCCESS);
}
hiprtcResult hiprtcVersion(int* major, int* minor) {
HIPRTC_INIT_API(major, minor);
if (major == nullptr || minor == nullptr) {
HIPRTC_RETURN(HIPRTC_ERROR_INVALID_INPUT);
}
*major = 9;
*minor = 0;
HIPRTC_RETURN(HIPRTC_SUCCESS);
}
rocclr/hip_stream.cpp
+313
View File
@@ -0,0 +1,313 @@
/* Copyright (c) 2015-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_event.hpp"
#include "thread/monitor.hpp"
static amd::Monitor streamSetLock("Guards global stream set");
static std::unordered_set<hip::Stream*> streamSet;
// Internal structure for stream callback handler
class StreamCallback {
public:
StreamCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData,
amd::Command* command)
: stream_(stream), callBack_(callback),
userData_(userData), command_(command) {
};
hipStream_t stream_;
hipStreamCallback_t callBack_;
void* userData_;
amd::Command* command_;
};
namespace hip {
// ================================================================================================
Stream::Stream(hip::Device* dev, amd::CommandQueue::Priority p,
unsigned int f, bool null_stream)
: queue_(nullptr), lock_("Stream Callback lock"), device_(dev),
priority_(p), flags_(f), null_(null_stream) {}
// ================================================================================================
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_);
// Create a host queue
bool result = (queue_ != nullptr) ? queue_->create() : false;
// Insert just created stream into the list of the blocking queues
if (result) {
if (!(flags_ & hipStreamNonBlocking)) {
amd::ScopedLock lock(streamSetLock);
streamSet.insert(this);
}
} else {
Destroy();
}
return result;
}
// ================================================================================================
amd::HostQueue* Stream::asHostQueue(bool skip_alloc) {
// Access to the stream object is lock protected, because possible allocation
amd::ScopedLock l(Lock());
if (queue_ == nullptr) {
// Create the host queue for the first time
if (!skip_alloc) {
Create();
}
}
return queue_;
}
// ================================================================================================
void Stream::Destroy() {
if (queue_ != nullptr) {
queue_->release();
queue_ = nullptr;
amd::ScopedLock lock(streamSetLock);
streamSet.erase(this);
}
delete this;
}
// ================================================================================================
void Stream::Finish() const {
if (queue_ != nullptr) {
queue_->finish();
}
}
// ================================================================================================
int Stream::DeviceId() const {
return device_->deviceId();
}
};
// ================================================================================================
void iHipWaitActiveStreams(amd::HostQueue* blocking_queue, bool wait_null_stream) {
amd::Command::EventWaitList eventWaitList;
{
amd::ScopedLock lock(streamSetLock);
for (const auto& stream : streamSet) {
amd::HostQueue* active_queue = stream->asHostQueue();
// If it's the current device
if ((&active_queue->device() == &blocking_queue->device()) &&
// and it's not the current stream
(active_queue != blocking_queue) &&
// check for a wait on the null stream
(stream->Null() == wait_null_stream)) {
// Get the last valid command
amd::Command* command = active_queue->getLastQueuedCommand(true);
if ((command != nullptr) &&
// Check the current active status
(command->status() != CL_COMPLETE)) {
eventWaitList.push_back(command);
}
}
}
}
// Check if we have to wait anything
if (eventWaitList.size() > 0) {
amd::Command* command = new amd::Marker(*blocking_queue, false, eventWaitList);
if (command != nullptr) {
command->enqueue();
command->release();
}
}
// Release all active commands. It's safe after the marker was enqueued
for (const auto& it : eventWaitList) {
it->release();
}
}
// ================================================================================================
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->command_->release();
delete cbo;
}
// ================================================================================================
static hipError_t ihipStreamCreate(hipStream_t* stream,
unsigned int flags, amd::CommandQueue::Priority priority) {
hip::Stream* hStream = new hip::Stream(hip::getCurrentDevice(), priority, flags);
if (hStream == nullptr) {
return hipErrorOutOfMemory;
}
*stream = reinterpret_cast<hipStream_t>(hStream);
ClPrint(amd::LOG_INFO, amd::LOG_API, "ihipStreamCreate: %zx", hStream);
return hipSuccess;
}
// ================================================================================================
hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags) {
HIP_INIT_API(hipStreamCreateWithFlags, stream, flags);
HIP_RETURN(ihipStreamCreate(stream, flags, amd::CommandQueue::Priority::Normal));
}
// ================================================================================================
hipError_t hipStreamCreate(hipStream_t *stream) {
HIP_INIT_API(hipStreamCreate, stream);
HIP_RETURN(ihipStreamCreate(stream, hipStreamDefault, amd::CommandQueue::Priority::Normal));
}
// ================================================================================================
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);
}
return HIP_RETURN(ihipStreamCreate(stream, flags, static_cast<amd::CommandQueue::Priority>(priority)));
}
// ================================================================================================
hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority) {
HIP_INIT_API(hipDeviceGetStreamPriorityRange, leastPriority, greatestPriority);
if (leastPriority != nullptr) {
*leastPriority = static_cast<int>(amd::CommandQueue::Priority::Normal);
}
if (greatestPriority != nullptr) {
// Only report one kind of priority for now.
*greatestPriority = static_cast<int>(amd::CommandQueue::Priority::Normal);
}
return HIP_RETURN(hipSuccess);
}
// ================================================================================================
hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags) {
HIP_INIT_API(hipStreamGetFlags, stream, flags);
if ((flags != nullptr) && (stream != nullptr)) {
*flags = reinterpret_cast<hip::Stream*>(stream)->Flags();
} else {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(hipSuccess);
}
// ================================================================================================
hipError_t hipStreamSynchronize(hipStream_t stream) {
HIP_INIT_API(hipStreamSynchronize, stream);
// Wait for the current host queue
hip::getQueue(stream)->finish();
HIP_RETURN(hipSuccess);
}
// ================================================================================================
hipError_t hipStreamDestroy(hipStream_t stream) {
HIP_INIT_API(hipStreamDestroy, stream);
if (stream == nullptr) {
HIP_RETURN(hipErrorInvalidHandle);
}
reinterpret_cast<hip::Stream*>(stream)->Destroy();
HIP_RETURN(hipSuccess);
}
// ================================================================================================
hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) {
HIP_INIT_API(hipStreamWaitEvent, stream, event, flags);
if (event == nullptr) {
HIP_RETURN(hipErrorInvalidHandle);
}
amd::HostQueue* queue = hip::getQueue(stream);
hip::Event* e = reinterpret_cast<hip::Event*>(event);
HIP_RETURN(e->streamWait(queue, flags));
}
// ================================================================================================
hipError_t hipStreamQuery(hipStream_t stream) {
HIP_INIT_API(hipStreamQuery, stream);
amd::HostQueue* hostQueue = hip::getQueue(stream);
amd::Command* command = hostQueue->getLastQueuedCommand(true);
if (command == nullptr) {
// Nothing was submitted to the queue
HIP_RETURN(hipSuccess);
}
amd::Event& event = command->event();
if (command->type() != 0) {
event.notifyCmdQueue();
}
hipError_t status = (command->status() == CL_COMPLETE) ? hipSuccess : hipErrorNotReady;
command->release();
HIP_RETURN(status);
}
// ================================================================================================
hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData,
unsigned int flags) {
HIP_INIT_API(hipStreamAddCallback, stream, callback, userData, flags);
amd::HostQueue* hostQueue = hip::getQueue(stream);
amd::Command* command = hostQueue->getLastQueuedCommand(true);
if (command == nullptr) {
amd::Command::EventWaitList eventWaitList;
command = new amd::Marker(*hostQueue, false, eventWaitList);
command->enqueue();
}
amd::Event& event = command->event();
StreamCallback* cbo = new StreamCallback(stream, callback, userData, command);
if(!event.setCallback(CL_COMPLETE, ihipStreamCallback, reinterpret_cast<void*>(cbo))) {
command->release();
return hipErrorInvalidHandle;
}
event.notifyCmdQueue();
HIP_RETURN(hipSuccess);
}
rocclr/hip_surface.cpp
+37
View File
@@ -0,0 +1,37 @@
/* Copyright (c) 2015-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/hcc_detail/hip_surface_types.h>
hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject,
const hipResourceDesc* pResDesc) {
HIP_INIT_API(hipCreateSurfaceObject, pSurfObject, pResDesc);
HIP_RETURN(hipErrorNotSupported);
}
hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject) {
HIP_INIT_API(hipDestroySurfaceObject, surfaceObject);
HIP_RETURN(hipErrorNotSupported);
}
rocclr/hip_texture.cpp
Executable
+1305
View File
File diff suppressed because it is too large Load Diff
rocclr/hiprtc_internal.hpp
+65
View File
@@ -0,0 +1,65 @@
/* Copyright (c) 2015-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. */
#ifndef HIPRTC_SRC_HIP_INTERNAL_H
#define HIPRTC_SRC_HIP_INTERNAL_H
#include "hip_internal.hpp"
#if __linux__
#include <cstdlib>
#if HIPRTC_USE_CXXABI
#include <cxxabi.h>
#define DEMANGLE abi::__cxa_demangle
#else
extern "C" char * __cxa_demangle(const char *mangled_name, char *output_buffer,
size_t *length, int *status);
#define DEMANGLE __cxa_demangle
#endif //HIPRTC_USE_CXXABI
#elif defined(_WIN32)
#include <Windows.h>
#include <DbgHelp.h>
#define UNDECORATED_SIZE 4096
#endif // __linux__
// This macro should be called at the beginning of every HIP RTC API.
#define HIPRTC_INIT_API(...) \
ClPrint(amd::LOG_INFO, amd::LOG_API, "[%zx] %s ( %s )", std::this_thread::get_id(), __func__, ToString( __VA_ARGS__ ).c_str()); \
amd::Thread* thread = amd::Thread::current(); \
if (!VDI_CHECK_THREAD(thread)) { \
HIPRTC_RETURN(HIPRTC_ERROR_INTERNAL_ERROR); \
} \
HIP_INIT();
#define HIPRTC_RETURN(ret) \
hiprtc::g_lastRtcError = ret; \
ClPrint(amd::LOG_INFO, amd::LOG_API, "[%zx] %s: Returned %s", std::this_thread::get_id(), __func__, \
hiprtcGetErrorString(hiprtc::g_lastRtcError)); \
return hiprtc::g_lastRtcError;
#endif // HIPRTC_SRC_HIP_INTERNAL_H
rocclr/trace_helper.h
+245
View File
@@ -0,0 +1,245 @@
/* Copyright (c) 2015-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. */
#pragma once
#include <iostream>
#include <iomanip>
#include <sstream>
#include <string>
//---
// Helper functions to convert HIP function arguments into strings.
// Handles POD data types as well as enumerations (ie hipMemcpyKind).
// The implementation uses C++11 variadic templates and template specialization.
// The hipMemcpyKind example below is a good example that shows how to implement conversion for a
// new HSA type.
// Handy macro to convert an enumeration to a stringified version of same:
#define CASE_STR(x) \
case x: \
return #x;
inline const char* ihipErrorString(hipError_t hip_error) {
switch (hip_error) {
CASE_STR(hipSuccess);
CASE_STR(hipErrorOutOfMemory);
CASE_STR(hipErrorNotInitialized);
CASE_STR(hipErrorDeinitialized);
CASE_STR(hipErrorProfilerDisabled);
CASE_STR(hipErrorProfilerNotInitialized);
CASE_STR(hipErrorProfilerAlreadyStarted);
CASE_STR(hipErrorProfilerAlreadyStopped);
CASE_STR(hipErrorInvalidImage);
CASE_STR(hipErrorInvalidContext);
CASE_STR(hipErrorContextAlreadyCurrent);
CASE_STR(hipErrorMapFailed);
CASE_STR(hipErrorUnmapFailed);
CASE_STR(hipErrorArrayIsMapped);
CASE_STR(hipErrorAlreadyMapped);
CASE_STR(hipErrorNoBinaryForGpu);
CASE_STR(hipErrorAlreadyAcquired);
CASE_STR(hipErrorNotMapped);
CASE_STR(hipErrorNotMappedAsArray);
CASE_STR(hipErrorNotMappedAsPointer);
CASE_STR(hipErrorECCNotCorrectable);
CASE_STR(hipErrorUnsupportedLimit);
CASE_STR(hipErrorContextAlreadyInUse);
CASE_STR(hipErrorPeerAccessUnsupported);
CASE_STR(hipErrorInvalidKernelFile);
CASE_STR(hipErrorInvalidGraphicsContext);
CASE_STR(hipErrorInvalidSource);
CASE_STR(hipErrorFileNotFound);
CASE_STR(hipErrorSharedObjectSymbolNotFound);
CASE_STR(hipErrorSharedObjectInitFailed);
CASE_STR(hipErrorOperatingSystem);
CASE_STR(hipErrorSetOnActiveProcess);
CASE_STR(hipErrorInvalidHandle);
CASE_STR(hipErrorNotFound);
CASE_STR(hipErrorIllegalAddress);
CASE_STR(hipErrorMissingConfiguration);
CASE_STR(hipErrorLaunchFailure);
CASE_STR(hipErrorPriorLaunchFailure);
CASE_STR(hipErrorLaunchTimeOut);
CASE_STR(hipErrorLaunchOutOfResources);
CASE_STR(hipErrorInvalidDeviceFunction);
CASE_STR(hipErrorInvalidConfiguration);
CASE_STR(hipErrorInvalidDevice);
CASE_STR(hipErrorInvalidValue);
CASE_STR(hipErrorInvalidDevicePointer);
CASE_STR(hipErrorInvalidMemcpyDirection);
CASE_STR(hipErrorUnknown);
CASE_STR(hipErrorNotReady);
CASE_STR(hipErrorNoDevice);
CASE_STR(hipErrorPeerAccessAlreadyEnabled);
CASE_STR(hipErrorPeerAccessNotEnabled);
CASE_STR(hipErrorRuntimeMemory);
CASE_STR(hipErrorRuntimeOther);
CASE_STR(hipErrorHostMemoryAlreadyRegistered);
CASE_STR(hipErrorHostMemoryNotRegistered);
CASE_STR(hipErrorTbd);
default:
return "hipErrorUnknown";
};
};
// Building block functions:
template <typename T>
inline std::string ToHexString(T v) {
std::ostringstream ss;
ss << "0x" << std::hex << v;
return ss.str();
};
template <typename T>
inline std::string ToString(T* v) {
std::ostringstream ss;
if (v == NULL) {
ss << "char array:<null>";
} else {
ss << v;
}
return ss.str();
};
template <typename T>
inline std::string ToString(T** v) {
std::ostringstream ss;
if (v == NULL) {
ss << "char array:<null>";
} else {
ss << v;
}
return ss.str();
};
//---
// Template overloads for ToString to handle specific types
// This is the default which works for most types:
template <typename T>
inline std::string ToString(T v) {
std::ostringstream ss;
ss << v;
return ss.str();
};
template <>
inline std::string ToString(hipFunction_t v) {
std::ostringstream ss;
ss << "0x" << std::hex << static_cast<void*>(v);
return ss.str();
};
// hipEvent_t specialization. TODO - maybe add an event ID for debug?
template <>
inline std::string ToString(hipEvent_t v) {
std::ostringstream ss;
ss << "event:" << std::hex << static_cast<void*>(v);
return ss.str();
};
// hipStream_t
template <>
inline std::string ToString(hipStream_t v) {
std::ostringstream ss;
if (v == NULL) {
ss << "stream:<null>";
} else {
ss << "stream:" << std::hex << static_cast<void*>(v);
}
return ss.str();
};
// hipCtx_t
template <>
inline std::string ToString(hipCtx_t v) {
std::ostringstream ss;
if (v == NULL) {
ss << "context:<null>";
} else {
ss << "context:" << std::hex << static_cast<void*>(v);
}
return ss.str();
};
// hipPitchedPtr
template <>
inline std::string ToString(hipPitchedPtr v) {
std::ostringstream ss;
ss << "pitchPtr:" << std::hex << static_cast<void*>(v.ptr);
return ss.str();
};
// hipMemcpyKind specialization
template <>
inline std::string ToString(hipMemcpyKind v) {
switch (v) {
CASE_STR(hipMemcpyHostToHost);
CASE_STR(hipMemcpyHostToDevice);
CASE_STR(hipMemcpyDeviceToHost);
CASE_STR(hipMemcpyDeviceToDevice);
CASE_STR(hipMemcpyDefault);
default:
return ToHexString(v);
};
};
template <>
inline std::string ToString(hipFuncCache_t v) {
switch (v) {
CASE_STR(hipFuncCachePreferNone);
CASE_STR(hipFuncCachePreferShared);
CASE_STR(hipFuncCachePreferL1);
CASE_STR(hipFuncCachePreferEqual);
default:
return ToHexString(v);
};
};
template <>
inline std::string ToString(hipSharedMemConfig v) {
switch (v) {
CASE_STR(hipSharedMemBankSizeDefault);
CASE_STR(hipSharedMemBankSizeFourByte);
CASE_STR(hipSharedMemBankSizeEightByte);
default:
return ToHexString(v);
};
};
template <>
inline std::string ToString(hipError_t v) {
return ihipErrorString(v);
};
// Catch empty arguments case
inline std::string ToString() { return (""); }
//---
// C++11 variadic template - peels off first argument, converts to string, and calls itself again to
// peel the next arg. Strings are automatically separated by comma+space.
template <typename T, typename... Args>
inline std::string ToString(T first, Args... args) {
return ToString(first) + ", " + ToString(args...);
}