Merge branch 'origin/pghafari/hip-vdi' into lmoriche/amd-master-next
Change-Id: I22c145d39f430ca571a981687bcb034ea6e3b8a2
[ROCm/hip commit: a8b81e96b4]
Este cometimento está contido em:
Laurent Morichetti
@@ -0,0 +1,198 @@
|
||||
#project("hip")
|
||||
cmake_minimum_required(VERSION 3.5.1)
|
||||
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11")
|
||||
|
||||
set(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)
|
||||
|
||||
add_definitions(-D__HIP_VDI__ -D__HIP_PLATFORM_HCC__ -DLINUX -D__x86_64__ -D__AMD64__ -DUNIX_OS -DqLittleEndian -DOPENCL_MAJOR=2 -DOPENCL_MINOR=0 -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()
|
||||
|
||||
#command example:
|
||||
#cmake -DVDI_DIR=/home/pghafari/vdi/vdi -DLIBOCL_STATIC_DIR=/home/pghafari/git-ocl/build/amdocl -DLIBVDI_STATIC_DIR=/home/pghafari/vdi/vdi/build -DUSE_COMGR_LIBRARY=yes -DCMAKE_BUILD_TYPE=Debug -DCMAKE_PREFIX_PATH="/home/pghafari/rocmgitvdi/hsa-runtime/opensrc/hsa-runtime/build;/home/pghafari/rocmgitvdi/hsa-runtime/opensrc;/home/pghafari/rocmgitvdi/support/lib/comgr;/home/pghafari/rocmgitvdi/support/lib/comgr/build;" -DCMAKE_MODULE_PATH=/home/pghafari/rocmgit/opencl/cmake/modules -DLLVM_INCLUDES=/home/pghafari/rocmgit/llvm-project/llvm/include ..
|
||||
|
||||
set(USE_PROF_API "1")
|
||||
|
||||
find_package(ROCT REQUIRED)
|
||||
find_package(ROCR REQUIRED)
|
||||
|
||||
if(NOT DEFINED LIBVDI_STATIC_DIR)
|
||||
find_path(LIBVDI_STATIC_DIR
|
||||
NAMES libamdvdi_static.a
|
||||
PATHS /opt/rocm/vdi
|
||||
PATH_SUFFIXES lib
|
||||
)
|
||||
endif()
|
||||
|
||||
if(NOT DEFINED VDI_DIR)
|
||||
find_path(VDI_DIR
|
||||
NAMES top.hpp
|
||||
PATH_SUFFIXES include
|
||||
PATHS /opt/rocm/vdi
|
||||
)
|
||||
endif()
|
||||
message("vdilib:${LIBVDI_STATIC_DIR} ,found vdi includes ${VDI_DIR}")
|
||||
set(PROF_API_HEADER_PATH ${VDI_DIR}/platform)
|
||||
#############################
|
||||
# Profiling API support
|
||||
#############################
|
||||
# Generate profiling API macros/structures header
|
||||
set(PROF_API_STR "${CMAKE_CURRENT_SOURCE_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")
|
||||
set(PROF_API_CMD "${PROF_API_GEN} -v -t --priv ${OPT_PROF_API} ${PROF_API_HDR} ${PROF_API_SRC} ${PROF_API_STR} >${PROF_API_LOG}")
|
||||
MESSAGE(STATUS "Generating profiling promitives: ${PROF_API_STR}")
|
||||
execute_process(COMMAND sh -c "rm -f ${PROF_API_STR}; ${PROF_API_CMD}")
|
||||
MESSAGE(COMMAND sh -c "rm -f ${PROF_API_STR}; ${PROF_API_CMD}")
|
||||
set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS ${PROF_API_GEN} ${PROF_API_HDR} ${PROF_API_STR})
|
||||
|
||||
# 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 VDI_DIR OR NOT DEFINED LIBOCL_STATIC_DIR OR NOT DEFINED LIBVDI_STATIC_DIR )
|
||||
# message(FATAL_ERROR "define VDI_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 "")
|
||||
|
||||
#find_package(hsa-runtime REQUIRED CONFIG PATHS ${HSA_RUNTIME} NO_DEFAULT_PATH)
|
||||
include_directories(${CMAKE_SOURCE_DIR})
|
||||
include_directories(${CMAKE_SOURCE_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(${VDI_DIR})
|
||||
include_directories(${VDI_DIR}/include)
|
||||
include_directories(${VDI_DIR}/compiler/lib)
|
||||
include_directories(${VDI_DIR}/compiler/lib/include)
|
||||
include_directories(${VDI_DIR}/compiler/lib/loaders)
|
||||
include_directories(${VDI_DIR}/compiler/lib/loaders/elf/utils/common)
|
||||
include_directories(${VDI_DIR}/compiler/lib/loaders/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}.")
|
||||
|
||||
#find_package(amd_comgr REQUIRED CONFIG)
|
||||
include_directories("$<TARGET_PROPERTY:amd_comgr,INTERFACE_INCLUDE_DIRECTORIES>")
|
||||
|
||||
add_definitions(-DBSD_LIBELF)
|
||||
#find_package( Threads REQUIRED CONFIG)
|
||||
#find_package( LibElf REQUIRED CONFIG)
|
||||
|
||||
add_library(amdhip64 SHARED
|
||||
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_icd.cpp
|
||||
cl_lqdflash_amd.cpp
|
||||
)
|
||||
add_library(amdvdi_static STATIC IMPORTED)
|
||||
set_target_properties(amdvdi_static PROPERTIES IMPORTED_LOCATION "${LIBVDI_STATIC_DIR}/libamdvdi_static.a")
|
||||
#add_library(amdocl64 SHARED IMPORTED)
|
||||
#set_target_properties(amdocl64 PROPERTIES IMPORTED_LOCATION "/opt/rocm/opencl/lib/x86_64/libamdocl64.so")
|
||||
#add_library(oclelf STATIC IMPORTED)
|
||||
#set_target_properties(oclelf PROPERTIES IMPORTED_LOCATION "${LIBOCLELF_STATIC_DIR}/liboclelf.a")
|
||||
|
||||
add_library(host INTERFACE)
|
||||
target_link_libraries(host INTERFACE amdhip64)
|
||||
add_library(device INTERFACE)
|
||||
target_link_libraries(device INTERFACE host)
|
||||
#target_link_libraries(amdhip64 PRIVATE amd_comgr)
|
||||
|
||||
target_link_libraries(amdhip64 amdvdi_static pthread dl ${ROCT_LIBRARIES} ${ROCR_LIBRARIES})
|
||||
#target_link_libraries(amdhip64 amdvdi_static amdocl64_static pthread dl ${ROCT_LIBRARIES} ${ROCR_LIBRARIES})
|
||||
|
||||
#install(CODE "execute_process( COMMAND ${CMAKE_COMMAND} -E create_symlink libamdhip64.so lib/libhip_hcc.so)" )
|
||||
install(TARGETS amdhip64 LIBRARY DESTINATION lib/x86_64 COMPONENT applications)
|
||||
|
||||
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(FILES ${CMAKE_BINARY_DIR}/lib/libhip_hcc.so DESTINATION lib COMPONENT MAIN)
|
||||
|
||||
|
||||
if(0)
|
||||
set(CPACK_GENERATOR "DEB;RPM" CACHE STRING "Default packaging generators")
|
||||
|
||||
set(CPACK_PACKAGE_CONTACT "Advanced Micro Devices Inc.")
|
||||
set(CPACK_PACKAGE_VENDOR "AMD")
|
||||
set(CPACK_PACKAGING_INSTALL_PREFIX "/opt/rocm/hip-on-vdi")
|
||||
|
||||
set(CPACK_PACKAGE_VERSION_MAJOR "1")
|
||||
set(CPACK_PACKAGE_VERSION_MINOR "0")
|
||||
set(CPACK_PACKAGE_VERSION_PATCH "0")
|
||||
|
||||
# Debian CPACK variables
|
||||
set(CPACK_DEB_COMPONENT_INSTALL ON)
|
||||
|
||||
set(CPACK_DEBIAN_MAIN_FILE_NAME "hip-on-vdi-1.0.0.deb")
|
||||
set(CPACK_DEBIAN_MAIN_PACKAGE_NAME "hip-on-vdi")
|
||||
#set(CPACK_DEBIAN_MAIN_PACKAGE_CONTROL_EXTRA "${CMAKE_CURRENT_SOURCE_DIR}/DEB/postinst;${CMAKE_CURRENT_SOURCE_DIR}/DEB/prerm")
|
||||
|
||||
# RPM CPACK variables
|
||||
set(CPACK_RPM_COMPONENT_INSTALL ON)
|
||||
|
||||
set(CPACK_RPM_MAIN_FILE_NAME "hip-on-vdi-1.0.0.rpm")
|
||||
set(CPACK_RPM_MAIN_PACKAGE_NAME "hip-on-vdi")
|
||||
#set(CPACK_RPM_MAIN_POST_INSTALL_SCRIPT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/RPM/rpm_post")
|
||||
#set(CPACK_RPM_MAIN_POST_UNINSTALL_SCRIPT_FILE "${CMAKE_CURRENT_SOURCE_DIR}/RPM/rpm_postun")
|
||||
|
||||
set(CPACK_COMPONENTS_ALL MAIN)
|
||||
include(CPack)
|
||||
endif()
|
||||
A apresentação das diferenças no ficheiro foi suprimida por ser demasiado grande
Carregar diff
@@ -0,0 +1,363 @@
|
||||
//
|
||||
// Copyright 2010 Advanced Micro Devices, Inc. All rights reserved.
|
||||
//
|
||||
|
||||
#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_
|
||||
@@ -0,0 +1,279 @@
|
||||
//
|
||||
// Copyright (c) 2008 Advanced Micro Devices, Inc. All rights reserved.
|
||||
//
|
||||
|
||||
#include "cl_common.hpp"
|
||||
#ifdef _WIN32
|
||||
#include <d3d10_1.h>
|
||||
#include "cl_d3d9_amd.hpp"
|
||||
#include "cl_d3d10_amd.hpp"
|
||||
#include "cl_d3d11_amd.hpp"
|
||||
#endif //_WIN32
|
||||
|
||||
#include <icd/loader/icd_dispatch.h>
|
||||
|
||||
#include <mutex>
|
||||
|
||||
amd::PlatformIDS amd::PlatformID::Platform = //{ NULL };
|
||||
{amd::ICDDispatchedObject::icdVendorDispatch_};
|
||||
|
||||
static cl_int CL_API_CALL icdGetPlatformInfo(cl_platform_id platform, cl_platform_info param_name,
|
||||
size_t param_value_size, void* param_value,
|
||||
size_t* param_value_size_ret) {
|
||||
return NULL ; //clGetPlatformInfo(NULL, param_name, param_value_size, param_value, param_value_size_ret);
|
||||
}
|
||||
|
||||
static cl_int CL_API_CALL icdGetDeviceIDs(cl_platform_id platform, cl_device_type device_type,
|
||||
cl_uint num_entries, cl_device_id* devices,
|
||||
cl_uint* num_devices) {
|
||||
return NULL ;//clGetDeviceIDs(NULL, device_type, num_entries, devices, num_devices);
|
||||
}
|
||||
|
||||
static cl_int CL_API_CALL icdGetDeviceInfo(cl_device_id device, cl_device_info param_name,
|
||||
size_t param_value_size, void* param_value,
|
||||
size_t* param_value_size_ret) {
|
||||
if (param_name == CL_DEVICE_PLATFORM) {
|
||||
// Return the ICD platform instead of the default NULL platform.
|
||||
cl_platform_id platform = reinterpret_cast<cl_platform_id>(&amd::PlatformID::Platform);
|
||||
return amd::clGetInfo(platform, param_value_size, param_value, param_value_size_ret);
|
||||
}
|
||||
|
||||
return NULL ; //clGetDeviceInfo(device, param_name, param_value_size, param_value, param_value_size_ret);
|
||||
}
|
||||
|
||||
KHRicdVendorDispatch amd::ICDDispatchedObject::icdVendorDispatch_[] = {
|
||||
{NULL /* should not get called */, icdGetPlatformInfo, icdGetDeviceIDs, icdGetDeviceInfo}
|
||||
#if 0
|
||||
clCreateContext, clCreateContextFromType, clRetainContext, clReleaseContext, clGetContextInfo,
|
||||
clCreateCommandQueue, clRetainCommandQueue, clReleaseCommandQueue, clGetCommandQueueInfo,
|
||||
clSetCommandQueueProperty, clCreateBuffer, clCreateImage2D, clCreateImage3D, clRetainMemObject,
|
||||
clReleaseMemObject, clGetSupportedImageFormats, clGetMemObjectInfo, clGetImageInfo,
|
||||
clCreateSampler, clRetainSampler, clReleaseSampler, clGetSamplerInfo,
|
||||
clCreateProgramWithSource, clCreateProgramWithBinary, clRetainProgram, clReleaseProgram,
|
||||
clBuildProgram, clUnloadCompiler, clGetProgramInfo, clGetProgramBuildInfo, clCreateKernel,
|
||||
clCreateKernelsInProgram, clRetainKernel, clReleaseKernel, clSetKernelArg, clGetKernelInfo,
|
||||
clGetKernelWorkGroupInfo, clWaitForEvents, clGetEventInfo, clRetainEvent, clReleaseEvent,
|
||||
clGetEventProfilingInfo, clFlush, clFinish, clEnqueueReadBuffer, clEnqueueWriteBuffer,
|
||||
clEnqueueCopyBuffer, clEnqueueReadImage, clEnqueueWriteImage, clEnqueueCopyImage,
|
||||
clEnqueueCopyImageToBuffer, clEnqueueCopyBufferToImage, clEnqueueMapBuffer, clEnqueueMapImage,
|
||||
clEnqueueUnmapMemObject, clEnqueueNDRangeKernel, clEnqueueTask, clEnqueueNativeKernel,
|
||||
clEnqueueMarker, clEnqueueWaitForEvents, clEnqueueBarrier, clGetExtensionFunctionAddress,
|
||||
clCreateFromGLBuffer, clCreateFromGLTexture2D, clCreateFromGLTexture3D,
|
||||
clCreateFromGLRenderbuffer, clGetGLObjectInfo, clGetGLTextureInfo, clEnqueueAcquireGLObjects,
|
||||
clEnqueueReleaseGLObjects, clGetGLContextInfoKHR,
|
||||
WINDOWS_SWITCH(clGetDeviceIDsFromD3D10KHR, NULL),
|
||||
WINDOWS_SWITCH(clCreateFromD3D10BufferKHR, NULL),
|
||||
WINDOWS_SWITCH(clCreateFromD3D10Texture2DKHR, NULL),
|
||||
WINDOWS_SWITCH(clCreateFromD3D10Texture3DKHR, NULL),
|
||||
WINDOWS_SWITCH(clEnqueueAcquireD3D10ObjectsKHR, NULL),
|
||||
WINDOWS_SWITCH(clEnqueueReleaseD3D10ObjectsKHR, NULL), clSetEventCallback, clCreateSubBuffer,
|
||||
clSetMemObjectDestructorCallback, clCreateUserEvent, clSetUserEventStatus,
|
||||
clEnqueueReadBufferRect, clEnqueueWriteBufferRect, clEnqueueCopyBufferRect,
|
||||
NULL, NULL, NULL, clCreateEventFromGLsyncKHR,
|
||||
|
||||
/* OpenCL 1.2*/
|
||||
clCreateSubDevices, clRetainDevice, clReleaseDevice, clCreateImage,
|
||||
clCreateProgramWithBuiltInKernels, clCompileProgram, clLinkProgram, clUnloadPlatformCompiler,
|
||||
clGetKernelArgInfo, clEnqueueFillBuffer, clEnqueueFillImage, clEnqueueMigrateMemObjects,
|
||||
clEnqueueMarkerWithWaitList, clEnqueueBarrierWithWaitList,
|
||||
clGetExtensionFunctionAddressForPlatform, clCreateFromGLTexture,
|
||||
|
||||
WINDOWS_SWITCH(clGetDeviceIDsFromD3D11KHR, NULL),
|
||||
WINDOWS_SWITCH(clCreateFromD3D11BufferKHR, NULL),
|
||||
WINDOWS_SWITCH(clCreateFromD3D11Texture2DKHR, NULL),
|
||||
WINDOWS_SWITCH(clCreateFromD3D11Texture3DKHR, NULL),
|
||||
WINDOWS_SWITCH(clCreateFromDX9MediaSurfaceKHR, NULL),
|
||||
WINDOWS_SWITCH(clEnqueueAcquireD3D11ObjectsKHR, NULL),
|
||||
WINDOWS_SWITCH(clEnqueueReleaseD3D11ObjectsKHR, NULL),
|
||||
|
||||
WINDOWS_SWITCH(clGetDeviceIDsFromDX9MediaAdapterKHR,
|
||||
NULL), // KHRpfn_clGetDeviceIDsFromDX9MediaAdapterKHR
|
||||
// clGetDeviceIDsFromDX9MediaAdapterKHR;
|
||||
WINDOWS_SWITCH(
|
||||
clEnqueueAcquireDX9MediaSurfacesKHR,
|
||||
NULL), // KHRpfn_clEnqueueAcquireDX9MediaSurfacesKHR clEnqueueAcquireDX9MediaSurfacesKHR;
|
||||
WINDOWS_SWITCH(
|
||||
clEnqueueReleaseDX9MediaSurfacesKHR,
|
||||
NULL), // KHRpfn_clEnqueueReleaseDX9MediaSurfacesKHR clEnqueueReleaseDX9MediaSurfacesKHR;
|
||||
|
||||
NULL,
|
||||
NULL, NULL, NULL,
|
||||
|
||||
clCreateCommandQueueWithProperties, clCreatePipe, clGetPipeInfo, clSVMAlloc, clSVMFree,
|
||||
clEnqueueSVMFree, clEnqueueSVMMemcpy, clEnqueueSVMMemFill, clEnqueueSVMMap, clEnqueueSVMUnmap,
|
||||
clCreateSamplerWithProperties, clSetKernelArgSVMPointer, clSetKernelExecInfo,
|
||||
clGetKernelSubGroupInfo,
|
||||
clCloneKernel,
|
||||
clCreateProgramWithIL,
|
||||
clEnqueueSVMMigrateMem,
|
||||
clGetDeviceAndHostTimer,
|
||||
clGetHostTimer,
|
||||
clGetKernelSubGroupInfo,
|
||||
clSetDefaultDeviceCommandQueue,
|
||||
|
||||
clSetProgramReleaseCallback,
|
||||
clSetProgramSpecializationConstant }
|
||||
#endif
|
||||
};
|
||||
|
||||
#if defined(ATI_OS_WIN)
|
||||
#include <Shlwapi.h>
|
||||
|
||||
#pragma comment(lib, "shlwapi.lib")
|
||||
|
||||
static bool ShouldLoadPlatform() {
|
||||
// Get the OpenCL ICD registry values
|
||||
HKEY platformsKey = NULL;
|
||||
if (RegOpenKeyExA(HKEY_LOCAL_MACHINE, "SOFTWARE\\Khronos\\OpenCL\\Vendors", 0, KEY_READ,
|
||||
&platformsKey) != ERROR_SUCCESS)
|
||||
return true;
|
||||
|
||||
std::vector<std::string> registryValues;
|
||||
DWORD dwIndex = 0;
|
||||
while (true) {
|
||||
char cszLibraryName[1024] = {0};
|
||||
DWORD dwLibraryNameSize = sizeof(cszLibraryName);
|
||||
DWORD dwLibraryNameType = 0;
|
||||
DWORD dwValue = 0;
|
||||
DWORD dwValueSize = sizeof(dwValue);
|
||||
|
||||
if (RegEnumValueA(platformsKey, dwIndex++, cszLibraryName, &dwLibraryNameSize, NULL,
|
||||
&dwLibraryNameType, (LPBYTE)&dwValue, &dwValueSize) != ERROR_SUCCESS)
|
||||
break;
|
||||
// Require that the value be a DWORD and equal zero
|
||||
if (dwLibraryNameType != REG_DWORD || dwValue != 0) {
|
||||
continue;
|
||||
}
|
||||
registryValues.push_back(cszLibraryName);
|
||||
}
|
||||
RegCloseKey(platformsKey);
|
||||
|
||||
HMODULE hm = NULL;
|
||||
if (!GetModuleHandleExA(
|
||||
GET_MODULE_HANDLE_EX_FLAG_FROM_ADDRESS | GET_MODULE_HANDLE_EX_FLAG_UNCHANGED_REFCOUNT,
|
||||
(LPCSTR)&ShouldLoadPlatform, &hm))
|
||||
return true;
|
||||
|
||||
char cszDllPath[1024] = {0};
|
||||
if (!GetModuleFileNameA(hm, cszDllPath, sizeof(cszDllPath))) return true;
|
||||
|
||||
// If we are loaded from the DriverStore, then there should be a registry
|
||||
// value matching our current module absolute path.
|
||||
if (std::find(registryValues.begin(), registryValues.end(), cszDllPath) == registryValues.end())
|
||||
return true;
|
||||
|
||||
LPSTR cszFileName;
|
||||
char buffer[1024] = {0};
|
||||
if (!GetFullPathNameA(cszDllPath, sizeof(buffer), buffer, &cszFileName)) return true;
|
||||
|
||||
// We found an absolute path in the registry that matched this DLL, now
|
||||
// check if there is also an entry with the same filename.
|
||||
if (std::find(registryValues.begin(), registryValues.end(), cszFileName) == registryValues.end())
|
||||
return true;
|
||||
|
||||
// Lastly, check if there is a DLL with the same name in the System folder.
|
||||
char cszSystemPath[1024] = {0};
|
||||
#if defined(ATI_BITS_32)
|
||||
if (!GetSystemWow64DirectoryA(cszSystemPath, sizeof(cszSystemPath)))
|
||||
#endif // defined(ATI_BITS_32)
|
||||
if (!GetSystemDirectoryA(cszSystemPath, sizeof(cszSystemPath))) return true;
|
||||
|
||||
std::string systemDllPath;
|
||||
systemDllPath.append(cszSystemPath).append("\\").append(cszFileName);
|
||||
if (!PathFileExistsA(systemDllPath.c_str())) {
|
||||
return true;
|
||||
}
|
||||
|
||||
// If we get here, then all 3 conditions are true:
|
||||
// - An entry in the registry with an absolute path matches the current DLL
|
||||
// - An entry in the registry with a relative path matches the current DLL
|
||||
// - A DLL with the same name was found in the system directory
|
||||
//
|
||||
// We should not load this platform!
|
||||
|
||||
return false;
|
||||
}
|
||||
|
||||
#else
|
||||
|
||||
#include <dlfcn.h>
|
||||
|
||||
// If there is only one platform, load it.
|
||||
// If there is more than one platform, only load platforms that have visible devices
|
||||
// If all platforms have no devices available, only load the PAL platform
|
||||
static bool ShouldLoadPlatform() {
|
||||
bool shouldLoad = true;
|
||||
|
||||
if (!amd::Runtime::initialized()) {
|
||||
amd::Runtime::init();
|
||||
}
|
||||
const int numDevices = amd::Device::numDevices(CL_DEVICE_TYPE_GPU, false);
|
||||
|
||||
void *otherPlatform = nullptr;
|
||||
if (amd::IS_LEGACY) {
|
||||
otherPlatform = dlopen("libamdocl64.so", RTLD_LAZY);
|
||||
if (otherPlatform != nullptr) { // Present platform exists
|
||||
shouldLoad = numDevices > 0;
|
||||
}
|
||||
} else {
|
||||
otherPlatform = dlopen("libamdocl-orca64.so", RTLD_LAZY);
|
||||
if (otherPlatform != nullptr) { // Legacy platform exists
|
||||
// gcc4.8 doesn't support casting void* to a function pointer
|
||||
// Work around this by creating a typedef untill we upgrade the compiler
|
||||
typedef void*(*clGetFunctionAddress_t)(const char *);
|
||||
typedef cl_int(*clIcdGetPlatformIDs_t)(cl_uint, cl_platform_id *, cl_uint *);
|
||||
|
||||
clGetFunctionAddress_t legacyGetFunctionAddress =
|
||||
reinterpret_cast<clGetFunctionAddress_t>(dlsym(otherPlatform, "clGetExtensionFunctionAddress"));
|
||||
clIcdGetPlatformIDs_t legacyGetPlatformIDs =
|
||||
reinterpret_cast<clIcdGetPlatformIDs_t>(legacyGetFunctionAddress("clIcdGetPlatformIDsKHR"));
|
||||
|
||||
cl_uint numLegacyPlatforms = 0;
|
||||
legacyGetPlatformIDs(0, nullptr, &numLegacyPlatforms);
|
||||
|
||||
shouldLoad = (numDevices > 0) || (numLegacyPlatforms == 0);
|
||||
}
|
||||
}
|
||||
|
||||
if (otherPlatform != nullptr) {
|
||||
dlclose(otherPlatform);
|
||||
}
|
||||
|
||||
return shouldLoad;
|
||||
}
|
||||
|
||||
#endif // defined(ATI_OS_WIN)
|
||||
|
||||
CL_API_ENTRY cl_int CL_API_CALL clIcdGetPlatformIDsKHR(cl_uint num_entries,
|
||||
cl_platform_id* platforms,
|
||||
cl_uint* num_platforms) {
|
||||
if (((num_entries > 0 || num_platforms == NULL) && platforms == NULL) ||
|
||||
(num_entries == 0 && platforms != NULL)) {
|
||||
return CL_INVALID_VALUE;
|
||||
}
|
||||
|
||||
static bool shouldLoad = true;
|
||||
|
||||
static std::once_flag initOnce;
|
||||
std::call_once(initOnce, [](){ shouldLoad = ShouldLoadPlatform(); });
|
||||
|
||||
if (!shouldLoad) {
|
||||
*not_null(num_platforms) = 0;
|
||||
return CL_SUCCESS;
|
||||
}
|
||||
|
||||
if (!amd::Runtime::initialized()) {
|
||||
amd::Runtime::init();
|
||||
}
|
||||
|
||||
if (num_platforms != NULL && platforms == NULL) {
|
||||
*num_platforms = 1;
|
||||
return CL_SUCCESS;
|
||||
}
|
||||
|
||||
assert(platforms != NULL && "check the code above");
|
||||
*platforms = reinterpret_cast<cl_platform_id>(&amd::PlatformID::Platform);
|
||||
|
||||
*not_null(num_platforms) = 1;
|
||||
return CL_SUCCESS;
|
||||
}
|
||||
@@ -0,0 +1,294 @@
|
||||
//
|
||||
// Copyright (c) 2015 Advanced Micro Devices, Inc. All rights reserved.
|
||||
//
|
||||
|
||||
#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, ®ion, NULL);
|
||||
} else {
|
||||
status = lfWriteFile(srcDst, bufferSize, (lf_file)handle_, 1, ®ion, 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
|
||||
@@ -0,0 +1,42 @@
|
||||
//
|
||||
// Copyright (c) 2015 Advanced Micro Devices, Inc. All rights reserved.
|
||||
//
|
||||
|
||||
#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
|
||||
Ficheiro executável
+610
@@ -0,0 +1,610 @@
|
||||
#!/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'void \*', r'void* ', 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');
|
||||
|
||||
# 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(
|
||||
'struct hip_api_data_t {\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' +
|
||||
'};\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
|
||||
if False:
|
||||
f.write('\n')
|
||||
f.write('#if 0\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\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': '',
|
||||
}
|
||||
# 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")
|
||||
|
||||
# Generating output header file
|
||||
with open(OUTPUT, 'w') as f:
|
||||
generate_prof_header(f, api_map, opts_map)
|
||||
|
||||
# Successfull exit
|
||||
sys.exit(0)
|
||||
Criar uma nova questão referindo esta
Bloquear um utilizador