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

SWDEV-351966 - Dispatch table for hip runtime

Browse Source 
Change-Id: Ie4a44fa8cf1ff9c152146070bbbf6b0636d4e325


[ROCm/clr commit: 5e21f0c6bd]
This commit is contained in:
Anusha GodavarthySurya
2023-04-21 10:46:05 +00:00
committed by Anusha Godavarthy Surya
parent cea00f53fb
commit 3bdedf0cc7
37 changed files with 4496 additions and 192 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp
+1318
View File
File diff suppressed because it is too large Load Diff
projects/clr/hipamd/src/CMakeLists.txt
+3 -1
View File
@@ -122,7 +122,9 @@ target_sources(amdhip64 PRIVATE
hip_surface.cpp
hip_texture.cpp
hip_gl.cpp
hip_vm.cpp)
hip_vm.cpp
hip_api_trace.cpp
hip_table_interface.cpp)
if(WIN32)
target_sources(amdhip64 PRIVATE
projects/clr/hipamd/src/amdhip.def
+1 -1
View File
@@ -240,7 +240,7 @@ hipTexRefGetMaxAnisotropy
hipTexRefGetMipmapFilterMode
hipTexRefGetMipmapLevelBias
hipTexRefGetMipmapLevelClamp
hipTexRefGetMipmappedArray
hipTexRefGetMipMappedArray
hipTexRefSetAddress
hipTexRefSetAddress2D
hipTexRefSetAddressMode
projects/clr/hipamd/src/hip_api_trace.cpp
+1260
View File
File diff suppressed because it is too large Load Diff
projects/clr/hipamd/src/hip_code_object.cpp
+2 -5
View File
@@ -30,7 +30,7 @@ THE SOFTWARE.
#include "hip_internal.hpp"
#include "platform/program.hpp"
#include <elf/elf.hpp>
namespace hip {
hipError_t ihipFree(void* ptr);
// forward declaration of methods required for managed variables
hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align = 0);
@@ -59,8 +59,6 @@ struct __ClangOffloadBundleHeader {
};
} // namespace
namespace hip {
bool CodeObject::IsClangOffloadMagicBundle(const void* data) {
std::string magic(reinterpret_cast<const char*>(data), kOffloadBundleMagicStrSize - 1);
return magic.compare(kOffloadBundleMagicStr) ? false : true;
@@ -532,7 +530,6 @@ hipError_t CodeObject::extractCodeObjectFromFatBinary(
LogPrintfError(" %s - [Unsupported]", bundleEntryId.c_str());
}
}
return hipErrorNoBinaryForGpu;
}
}
@@ -906,4 +903,4 @@ hipError_t StatCO::initStatManagedVarDevicePtr(int deviceId) {
}
return err;
}
}; // namespace hip
} // namespace hip
projects/clr/hipamd/src/hip_code_object.hpp
+2 -3
View File
@@ -34,11 +34,10 @@ THE SOFTWARE.
#include "device/device.hpp"
#include "platform/program.hpp"
namespace hip {
//Forward Declaration for friend usage
class PlatformState;
namespace hip {
//Code Object base class
class CodeObject {
public:
@@ -150,7 +149,7 @@ public:
//pointer to the alocated managed memory has to be copied to the address of symbol
hipError_t initStatManagedVarDevicePtr(int deviceId);
private:
friend class ::PlatformState;
friend class hip::PlatformState;
//Populated during __hipRegisterFatBinary
std::unordered_map<const void*, FatBinaryInfo*> modules_;
//Populated during __hipRegisterFuncs
projects/clr/hipamd/src/hip_context.cpp
+10 -14
View File
@@ -25,13 +25,14 @@
#include "utils/flags.hpp"
#include "utils/versions.hpp"
std::vector<hip::Device*> g_devices;
std::once_flag g_ihipInitialized;
namespace hip {
std::vector<hip::Device*> g_devices;
thread_local TlsAggregator tls;
amd::Context* host_context = nullptr;
//init() is only to be called from the HIP_INIT macro only once
// init() is only to be called from the HIP_INIT macro only once
void init(bool* status) {
amd::IS_HIP = true;
GPU_NUM_MEM_DEPENDENCY = 0;
@@ -78,12 +79,10 @@ void init(bool* status) {
return;
}
Device* getCurrentDevice() {
return tls.device_;
}
Device* getCurrentDevice() { return tls.device_; }
void setCurrentDevice(unsigned int index) {
assert(index<g_devices.size());
assert(index < g_devices.size());
tls.device_ = g_devices[index];
uint32_t preferredNumaNode = (tls.device_)->devices()[0]->getPreferredNumaNode();
amd::Os::setPreferredNumaNode(preferredNumaNode);
@@ -134,10 +133,6 @@ hip::Stream* getNullStream() {
return device ? device->NullStream() : nullptr;
}
};
using namespace hip;
hipError_t hipInit(unsigned int flags) {
HIP_INIT_API(hipInit, flags);
@@ -148,7 +143,7 @@ hipError_t hipInit(unsigned int flags) {
HIP_RETURN(hipSuccess);
}
hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device) {
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()) {
@@ -169,12 +164,12 @@ hipError_t hipCtxSetCurrent(hipCtx_t ctx) {
HIP_INIT_API(hipCtxSetCurrent, ctx);
if (ctx == nullptr) {
if(!tls.ctxt_stack_.empty()) {
if (!tls.ctxt_stack_.empty()) {
tls.ctxt_stack_.pop();
}
} else {
hip::tls.device_ = reinterpret_cast<hip::Device*>(ctx);
if(!tls.ctxt_stack_.empty()) {
if (!tls.ctxt_stack_.empty()) {
tls.ctxt_stack_.pop();
}
tls.ctxt_stack_.push(hip::getCurrentDevice());
@@ -199,7 +194,7 @@ hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) {
HIP_RETURN(hipSuccess);
}
hipError_t hipRuntimeGetVersion(int *runtimeVersion) {
hipError_t hipRuntimeGetVersion(int* runtimeVersion) {
HIP_INIT_API_NO_RETURN(hipRuntimeGetVersion, runtimeVersion);
if (!runtimeVersion) {
@@ -399,3 +394,4 @@ hipError_t hipDevicePrimaryCtxSetFlags(hipDevice_t dev, unsigned int flags) {
HIP_RETURN(hipErrorContextAlreadyInUse);
}
}
} // namespace hip
projects/clr/hipamd/src/hip_device.cpp
+2 -3
View File
@@ -147,8 +147,6 @@ Device::~Device() {
}
}
} // namespace hip
void ihipDestroyDevice() {
for (auto deviceHandle : g_devices) {
delete deviceHandle;
@@ -648,4 +646,5 @@ extern "C" hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, hipDevice_t
*props = deviceProps;
return hipSuccess;
}
}
} // namespace hip
projects/clr/hipamd/src/hip_device_runtime.cpp
+3
View File
@@ -22,6 +22,8 @@
#include "hip_internal.hpp"
namespace hip {
hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* properties) {
HIP_INIT_API(hipChooseDevice, device, properties);
@@ -689,3 +691,4 @@ hipError_t hipSetValidDevices(int* device_arr, int len) {
HIP_RETURN(hipErrorNotSupported);
}
} //namespace hip
projects/clr/hipamd/src/hip_error.cpp
+3
View File
@@ -22,6 +22,8 @@
#include "hip_internal.hpp"
namespace hip {
hipError_t hipGetLastError()
{
HIP_INIT_API(hipGetLastError);
@@ -380,3 +382,4 @@ hipError_t hipDrvGetErrorString(hipError_t hip_error, const char** errStr)
return hipErrorInvalidValue;
}
}
} //namespace hip
projects/clr/hipamd/src/hip_event.cpp
+1 -1
View File
@@ -277,7 +277,6 @@ bool isValid(hipEvent_t event) {
return true;
}
} // namespace hip
// ================================================================================================
hipError_t ihipEventCreateWithFlags(hipEvent_t* event, unsigned flags) {
unsigned supportedFlags = hipEventDefault | hipEventBlockingSync | hipEventDisableTiming |
@@ -458,3 +457,4 @@ hipError_t hipEventQuery(hipEvent_t event) {
HIP_INIT_API(hipEventQuery, event);
HIP_RETURN(ihipEventQuery(event));
}
} // namespace hip
projects/clr/hipamd/src/hip_event.hpp
+2 -2
View File
@@ -25,6 +25,7 @@
#include "thread/monitor.hpp"
// Internal structure for stream callback handler
namespace hip {
class StreamCallback {
protected:
void* userData_;
@@ -64,7 +65,6 @@ class LaunchHostFuncCallback : public StreamCallback {
void CL_CALLBACK ihipStreamCallback(cl_event event, cl_int command_exec_status, void* user_data);
namespace hip {
#define IPC_SIGNALS_PER_EVENT 32
typedef struct ihipIpcEventShmem_s {
@@ -226,11 +226,11 @@ class IPCEvent : public Event {
hipError_t enqueueRecordCommand(hipStream_t stream, amd::Command* command, bool record);
};
}; // namespace hip
struct CallbackData {
int previous_read_index;
hip::ihipIpcEventShmem_t* shmem;
};
} // namespace hip
#endif // HIP_EVEMT_H
projects/clr/hipamd/src/hip_event_ipc.cpp
+2 -4
View File
@@ -28,11 +28,10 @@
#endif
// ================================================================================================
namespace hip {
hipError_t ihipEventCreateWithFlags(hipEvent_t* event, unsigned flags);
namespace hip {
bool IPCEvent::createIpcEventShmemIfNeeded() {
if (ipc_evt_.ipc_shmem_) {
// ipc_shmem_ already created, no need to create it again
@@ -219,8 +218,6 @@ hipError_t IPCEvent::OpenHandle(ihipIpcEventHandle_t* handle) {
return status;
}
} // namespace hip
// ================================================================================================
hipError_t hipIpcGetEventHandle(hipIpcEventHandle_t* handle, hipEvent_t event) {
@@ -248,3 +245,4 @@ hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle)
ihipIpcEventHandle_t* iHandle = reinterpret_cast<ihipIpcEventHandle_t*>(&handle);
HIP_RETURN(e->OpenHandle(iHandle));
}
} // namespace hip
projects/clr/hipamd/src/hip_gl.cpp
+3
View File
@@ -30,6 +30,8 @@
namespace amd {
static std::once_flag interopOnce;
}
namespace hip {
// Sets up GL context association with amd context.
// NOTE: Refer to Context setup code in OCLTestImp.cpp
void setupGLInteropOnce() {
@@ -767,3 +769,4 @@ hipError_t hipGraphicsUnregisterResource(hipGraphicsResource_t resource) {
HIP_RETURN(hipSuccess);
}
} // namespace hip
projects/clr/hipamd/src/hip_graph.cpp
+5 -2
View File
@@ -26,6 +26,8 @@
#include "hip_event.hpp"
#include "hip_mempool_impl.hpp"
namespace hip {
std::vector<hip::Stream*> g_captureStreams;
amd::Monitor g_captureStreamsLock{"StreamCaptureGlobalList"};
amd::Monitor g_streamSetLock{"StreamCaptureset"};
@@ -2477,7 +2479,7 @@ hipError_t hipGraphRetainUserObject(hipGraph_t graph, hipUserObject_t object, un
HIP_RETURN(hipSuccess);
}
if (flags != hipGraphUserObjectMove) {
status = hipUserObjectRetain(object, count);
status = hip::hipUserObjectRetain(object, count);
if (status != hipSuccess) {
HIP_RETURN(status);
}
@@ -2505,7 +2507,7 @@ hipError_t hipGraphReleaseUserObject(hipGraph_t graph, hipUserObject_t object, u
if (userObject->referenceCount() == releaseCount) {
g->RemoveUserObjGraph(userObject);
}
hipError_t status = hipUserObjectRelease(object, count);
hipError_t status = hip::hipUserObjectRelease(object, count);
HIP_RETURN(status);
}
@@ -2595,3 +2597,4 @@ hipError_t hipGraphUpload(hipGraphExec_t graphExec, hipStream_t stream) {
// memory for memAlloc nodes if any when support is added with mempool feature
HIP_RETURN(hipSuccess);
}
} // namespace hip
projects/clr/hipamd/src/hip_graph_capture.hpp
+2
View File
@@ -19,6 +19,7 @@
THE SOFTWARE. */
#pragma once
namespace hip {
// forward declaration of capture methods
hipError_t capturehipLaunchKernel(hipStream_t& stream, const void*& hostFunction, dim3& gridDim,
dim3& blockDim, void**& args, size_t& sharedMemBytes);
@@ -97,3 +98,4 @@ hipError_t capturehipLaunchHostFunc(hipStream_t& stream, hipHostFn_t& fn, void*&
hipError_t capturehipMallocAsync(hipStream_t stream, hipMemPool_t mem_pool, size_t size, void** dev_ptr);
hipError_t capturehipFreeAsync(hipStream_t stream, void* dev_ptr);
}
projects/clr/hipamd/src/hip_graph_helper.hpp
+9 -6
View File
@@ -22,6 +22,7 @@ THE SOFTWARE.
#include "hip_conversions.hpp"
namespace hip {
hipError_t ihipMemcpy3D_validate(const hipMemcpy3DParms* p);
hipError_t ihipDrvMemcpy3D_validate(const HIP_MEMCPY3D* pCopy);
@@ -51,11 +52,11 @@ hipError_t ihipMemset3D_validate(hipPitchedPtr pitchedDevPtr, int value, hipExte
hipError_t ihipLaunchKernelCommand(amd::Command*& command, hipFunction_t f,
uint32_t globalWorkSizeX, uint32_t globalWorkSizeY,
uint32_t globalWorkSizeZ, uint32_t blockDimX, uint32_t blockDimY,
uint32_t blockDimZ, uint32_t sharedMemBytes,
hip::Stream* stream, void** kernelParams, void** extra,
hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags,
uint32_t params, uint32_t gridId, uint32_t numGrids,
uint64_t prevGridSum, uint64_t allGridSum, uint32_t firstDevice);
uint32_t blockDimZ, uint32_t sharedMemBytes, hip::Stream* stream,
void** kernelParams, void** extra, hipEvent_t startEvent,
hipEvent_t stopEvent, uint32_t flags, uint32_t params,
uint32_t gridId, uint32_t numGrids, uint64_t prevGridSum,
uint64_t allGridSum, uint32_t firstDevice);
hipError_t ihipMemcpy3DCommand(amd::Command*& command, const hipMemcpy3DParms* p,
hip::Stream* stream);
@@ -67,7 +68,8 @@ hipError_t ihipMemsetCommand(std::vector<amd::Command*>& commands, void* dst, in
size_t valueSize, size_t sizeBytes, hip::Stream* stream);
hipError_t ihipMemset3DCommand(std::vector<amd::Command*>& commands, hipPitchedPtr pitchedDevPtr,
int value, hipExtent extent, hip::Stream* stream, size_t elementSize = 1);
int value, hipExtent extent, hip::Stream* stream,
size_t elementSize = 1);
hipError_t ihipMemcpySymbol_validate(const void* symbol, size_t sizeBytes, size_t offset,
size_t& sym_size, hipDeviceptr_t& device_ptr);
@@ -121,3 +123,4 @@ hipError_t ihipMemcpyAtoHValidate(hipArray_t srcArray, void* dstHost, amd::Coord
size_t &start);
hipError_t ihipGraphMemsetParams_validate(const hipMemsetParams* pNodeParams);
} // namespace hip
projects/clr/hipamd/src/hip_hcc.map.in
+1 -1
View File
@@ -235,7 +235,7 @@ global:
hipTexRefGetMipmapFilterMode;
hipTexRefGetMipmapLevelBias;
hipTexRefGetMipmapLevelClamp;
hipTexRefGetMipmappedArray;
hipTexRefGetMipMappedArray;
hipTexRefSetAddress;
hipTexRefSetAddress2D;
hipTexRefSetAddressMode;
projects/clr/hipamd/src/hip_hmm.cpp
+3
View File
@@ -25,6 +25,8 @@
#include "platform/command.hpp"
#include "platform/memory.hpp"
namespace hip {
// Forward declaraiton of a function
hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align = 0);
@@ -304,3 +306,4 @@ hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align) {
ClPrint(amd::LOG_INFO, amd::LOG_API, "ihipMallocManaged ptr=0x%zx", *ptr);
return hipSuccess;
}
} //namespace hip
projects/clr/hipamd/src/hip_intercept.cpp
+6 -6
View File
@@ -24,11 +24,10 @@
#include "hip_prof_api.h"
// HIP API callback/activity
namespace hip {
extern const std::string& FunctionName(const hipFunction_t f);
extern "C" {
int hipGetStreamDeviceId(hipStream_t stream) {
if (!hip::isValid(stream)) {
return -1;
@@ -47,11 +46,12 @@ const char* hipKernelNameRefByPtr(const void* host_function, hipStream_t stream)
: nullptr;
}
void hipRegisterTracerCallback(int (*function)(activity_domain_t domain, uint32_t operation_id,
void* data)) {
activity_prof::report_activity.store(function, std::memory_order_relaxed);
void hipRegisterTracerCallback(const void* function) {
typedef int (*fptr)(activity_domain_t domain, uint32_t operation_id, void* data);
fptr my_fptr = reinterpret_cast<fptr>(function);
activity_prof::report_activity.store(my_fptr, std::memory_order_relaxed);
}
const char* hipApiName(uint32_t id) { return hip_api_name(id); }
} // extern "C"
} // namespace hip
projects/clr/hipamd/src/hip_internal.hpp
+74 -74
View File
@@ -52,20 +52,7 @@
#define IHIP_IPC_MEM_HANDLE_SIZE 32
#define IHIP_IPC_MEM_RESERVED_SIZE LP64_SWITCH(20,12)
namespace hip {
struct Graph;
struct GraphNode;
struct GraphExec;
struct UserObject;
}
typedef struct ihipIpcMemHandle_st {
char ipc_handle[IHIP_IPC_MEM_HANDLE_SIZE]; ///< ipc memory handle on ROCr
size_t psize;
size_t poffset;
int owners_process_id;
char reserved[IHIP_IPC_MEM_RESERVED_SIZE];
} ihipIpcMemHandle_t;
extern std::once_flag g_ihipInitialized;
typedef struct hipArray {
void* data; // FIXME: generalize this
@@ -81,6 +68,19 @@ typedef struct hipArray {
unsigned int flags;
}hipArray;
namespace hip {
struct Graph;
struct GraphNode;
struct GraphExec;
struct UserObject;
typedef struct ihipIpcMemHandle_st {
char ipc_handle[IHIP_IPC_MEM_HANDLE_SIZE]; ///< ipc memory handle on ROCr
size_t psize;
size_t poffset;
int owners_process_id;
char reserved[IHIP_IPC_MEM_RESERVED_SIZE];
} ihipIpcMemHandle_t;
#define IHIP_IPC_EVENT_HANDLE_SIZE 32
#define IHIP_IPC_EVENT_RESERVED_SIZE LP64_SWITCH(28,24)
typedef struct ihipIpcEventHandle_st {
@@ -91,18 +91,18 @@ typedef struct ihipIpcEventHandle_st {
}ihipIpcEventHandle_t;
const char* ihipGetErrorName(hipError_t hip_error);
}
extern std::once_flag g_ihipInitialized;
#define HIP_INIT(noReturn) \
{ \
bool status = true; \
std::call_once(g_ihipInitialized, hip::init, &status); \
if (!status && !noReturn) { \
if (!status && !noReturn) { \
HIP_RETURN(hipErrorInvalidDevice); \
} \
if (hip::tls.device_ == nullptr && g_devices.size() > 0) { \
hip::tls.device_ = g_devices[0]; \
amd::Os::setPreferredNumaNode(g_devices[0]->devices()[0]->getPreferredNumaNode()); \
if (hip::tls.device_ == nullptr && hip::g_devices.size() > 0) { \
hip::tls.device_ = hip::g_devices[0]; \
amd::Os::setPreferredNumaNode(hip::g_devices[0]->devices()[0]->getPreferredNumaNode()); \
} \
}
@@ -110,9 +110,9 @@ extern std::once_flag g_ihipInitialized;
{ \
bool status = true; \
std::call_once(g_ihipInitialized, hip::init, &status); \
if (hip::tls.device_ == nullptr && g_devices.size() > 0) { \
hip::tls.device_ = g_devices[0]; \
amd::Os::setPreferredNumaNode(g_devices[0]->devices()[0]->getPreferredNumaNode()); \
if (hip::tls.device_ == nullptr && hip::g_devices.size() > 0) { \
hip::tls.device_ = hip::g_devices[0]; \
amd::Os::setPreferredNumaNode(hip::g_devices[0]->devices()[0]->getPreferredNumaNode()); \
} \
}
@@ -125,37 +125,36 @@ extern std::once_flag g_ihipInitialized;
#define HIP_ERROR_PRINT(err, ...) \
ClPrint(amd::LOG_INFO, amd::LOG_API, "%s: Returned %s : %s", \
__func__, ihipGetErrorName(err), ToString( __VA_ARGS__ ).c_str());
__func__, hip::ihipGetErrorName(err), ToString( __VA_ARGS__ ).c_str());
#define HIP_INIT_API_INTERNAL(noReturn, cid, ...) \
amd::Thread* thread = amd::Thread::current(); \
if (!VDI_CHECK_THREAD(thread)) { \
ClPrint(amd::LOG_NONE, amd::LOG_ALWAYS, "An internal error has occurred." \
" This may be due to insufficient memory."); \
if (!noReturn) { \
return hipErrorOutOfMemory; \
} \
} \
HIP_INIT(noReturn) \
HIP_API_PRINT(__VA_ARGS__) \
#define HIP_INIT_API_INTERNAL(noReturn, cid, ...) \
amd::Thread* thread = amd::Thread::current(); \
if (!VDI_CHECK_THREAD(thread)) { \
ClPrint(amd::LOG_NONE, amd::LOG_ALWAYS, \
"An internal error has occurred." \
" This may be due to insufficient memory."); \
if (!noReturn) { \
return hipErrorOutOfMemory; \
} \
} \
HIP_INIT(noReturn) \
HIP_API_PRINT(__VA_ARGS__) \
HIP_CB_SPAWNER_OBJECT(cid);
// This macro should be called at the beginning of every HIP API.
#define HIP_INIT_API(cid, ...) \
HIP_INIT_API_INTERNAL(0, cid, __VA_ARGS__) \
if (g_devices.size() == 0) { \
if (hip::g_devices.size() == 0) { \
HIP_RETURN(hipErrorNoDevice); \
}
#define HIP_INIT_API_NO_RETURN(cid, ...) \
HIP_INIT_API_INTERNAL(1, cid, __VA_ARGS__)
#define HIP_RETURN_DURATION(ret, ...) \
hip::tls.last_error_ = ret; \
HIPPrintDuration(amd::LOG_INFO, amd::LOG_API, &startTimeUs, \
"%s: Returned %s : %s", \
__func__, ihipGetErrorName(hip::tls.last_error_), \
ToString( __VA_ARGS__ ).c_str()); \
#define HIP_RETURN_DURATION(ret, ...) \
hip::tls.last_error_ = ret; \
HIPPrintDuration(amd::LOG_INFO, amd::LOG_API, &startTimeUs, "%s: Returned %s : %s", __func__, \
hip::ihipGetErrorName(hip::tls.last_error_), ToString(__VA_ARGS__).c_str()); \
return hip::tls.last_error_;
#define HIP_RETURN(ret, ...) \
@@ -202,11 +201,11 @@ extern std::once_flag g_ihipInitialized;
}
#define STREAM_CAPTURE(name, stream, ...) \
getStreamPerThread(stream); \
hip::getStreamPerThread(stream); \
if (stream != nullptr && \
reinterpret_cast<hip::Stream*>(stream)->GetCaptureStatus() == \
hipStreamCaptureStatusActive) { \
hipError_t status = capture##name(stream, ##__VA_ARGS__); \
hipError_t status = hip::capture##name(stream, ##__VA_ARGS__); \
return status; \
}
@@ -228,6 +227,7 @@ struct ihipExec_t {
std::vector<char> arguments_;
};
namespace hip {
class stream_per_thread {
private:
std::vector<hipStream_t> m_streams;
@@ -239,7 +239,6 @@ public:
hipStream_t get();
};
namespace hip {
class Device;
class MemoryPool;
class Stream : public amd::HostQueue {
@@ -581,40 +580,41 @@ namespace hip {
extern bool isValid(hipEvent_t event);
extern amd::Monitor hipArraySetLock;
extern std::unordered_set<hipArray*> hipArraySet;
}; // namespace hip
extern void WaitThenDecrementSignal(hipStream_t stream, hipError_t status, void* user_data);
extern void WaitThenDecrementSignal(hipStream_t stream, hipError_t status, void* user_data);
/// Wait all active streams on the blocking queue. The method enqueues a wait command and
/// doesn't stall the current thread
extern void iHipWaitActiveStreams(hip::Stream* blocking_stream, bool wait_null_stream = false);
/// Wait all active streams on the blocking queue. The method enqueues a wait command and
/// doesn't stall the current thread
extern void iHipWaitActiveStreams(hip::Stream* blocking_stream, bool wait_null_stream = false);
extern std::vector<hip::Device*> g_devices;
extern hipError_t ihipDeviceGetCount(int* count);
extern int ihipGetDevice();
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, size_t size = 0);
extern amd::Memory* getMemoryObjectWithOffset(const void* ptr, const size_t size = 0);
extern void getStreamPerThread(hipStream_t& stream);
extern hipStream_t getPerThreadDefaultStream();
extern hipError_t ihipUnbindTexture(textureReference* texRef);
extern hipError_t ihipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags);
extern hipError_t ihipHostUnregister(void* hostPtr);
extern hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, hipDevice_t device);
extern hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags);
extern amd::Memory* getMemoryObject(const void* ptr, size_t& offset, size_t size = 0);
extern amd::Memory* getMemoryObjectWithOffset(const void* ptr, const size_t size = 0);
extern void getStreamPerThread(hipStream_t& stream);
extern hipStream_t getPerThreadDefaultStream();
extern hipError_t ihipUnbindTexture(textureReference* texRef);
extern hipError_t ihipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags);
extern hipError_t ihipHostUnregister(void* hostPtr);
extern hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, hipDevice_t device);
extern hipError_t ihipDeviceGet(hipDevice_t* device, int deviceId);
extern hipError_t ihipStreamOperation(hipStream_t stream, cl_command_type cmdType, void* ptr,
uint64_t value, uint64_t mask, unsigned int flags, size_t sizeBytes);
hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
hip::Stream& stream, bool isHostAsync = false, bool isGPUAsync = true);
constexpr bool kOptionChangeable = true;
constexpr bool kNewDevProg = false;
extern hipError_t ihipDeviceGet(hipDevice_t* device, int deviceId);
extern hipError_t ihipStreamOperation(hipStream_t stream, cl_command_type cmdType, void* ptr,
uint64_t value, uint64_t mask, unsigned int flags,
size_t sizeBytes);
hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
hip::Stream& stream, bool isHostAsync = false, bool isGPUAsync = true);
constexpr bool kOptionChangeable = true;
constexpr bool kNewDevProg = false;
constexpr bool kMarkerDisableFlush = true; //!< Avoids command batch flush in ROCclr
constexpr bool kMarkerDisableFlush = true; //!< Avoids command batch flush in ROCclr
extern std::vector<hip::Stream*> g_captureStreams;
extern amd::Monitor g_captureStreamsLock;
extern amd::Monitor g_streamSetLock;
extern std::unordered_set<hip::Stream*> g_allCapturingStreams;
#endif // HIP_SRC_HIP_INTERNAL_H
extern std::vector<hip::Stream*> g_captureStreams;
extern amd::Monitor g_captureStreamsLock;
extern amd::Monitor g_streamSetLock;
extern std::unordered_set<hip::Stream*> g_allCapturingStreams;
} // namespace hip
#endif // HIP_SRC_HIP_INTERNAL_H
projects/clr/hipamd/src/hip_memory.cpp
+8 -7
View File
@@ -26,9 +26,10 @@
#include "platform/command.hpp"
#include "platform/memory.hpp"
#include "platform/external_memory.hpp"
namespace hip {
amd::Monitor hip::hipArraySetLock{"Guards global hipArray set"};
std::unordered_set<hipArray*> hip::hipArraySet;
amd::Monitor hipArraySetLock{"Guards global hipArray set"};
std::unordered_set<hipArray*> hipArraySet;
// ================================================================================================
amd::Memory* getMemoryObject(const void* ptr, size_t& offset, size_t size) {
@@ -727,7 +728,7 @@ hipError_t ihipArrayDestroy(hipArray_t array) {
return hipErrorInvalidValue;
}
{
amd::ScopedLock lock(hip::hipArraySetLock);
amd::ScopedLock lock(hipArraySetLock);
if (hip::hipArraySet.find(array) == hip::hipArraySet.end()) {
return hipErrorContextIsDestroyed;
} else {
@@ -1101,7 +1102,7 @@ hipError_t ihipArrayCreate(hipArray_t* array,
(*array)->NumChannels = pAllocateArray->NumChannels;
(*array)->flags = pAllocateArray->Flags;
{
amd::ScopedLock lock(hip::hipArraySetLock);
amd::ScopedLock lock(hipArraySetLock);
hip::hipArraySet.insert(*array);
}
return hipSuccess;
@@ -3859,7 +3860,7 @@ hipError_t hipArrayDestroy(hipArray_t array) {
hipError_t ihipArray3DGetDescriptor(HIP_ARRAY3D_DESCRIPTOR* desc,
hipArray_t array) {
{
amd::ScopedLock lock(hip::hipArraySetLock);
amd::ScopedLock lock(hipArraySetLock);
if (hip::hipArraySet.find(array) == hip::hipArraySet.end()) {
return hipErrorInvalidHandle;
}
@@ -4267,7 +4268,7 @@ hipError_t ihipMipmappedArrayGetLevel(hipArray_t* level_array_pptr,
(*level_array_pptr)->textureType = 0;
(*level_array_pptr)->flags = mipmapped_array_ptr->flags;
amd::ScopedLock lock(hip::hipArraySetLock);
amd::ScopedLock lock(hipArraySetLock);
hip::hipArraySet.insert(*level_array_pptr);
return hipSuccess;
@@ -4360,4 +4361,4 @@ hipError_t hipExternalMemoryGetMappedMipmappedArray(
HIP_RETURN(ihipMipmapArrayCreate(mipmap, &allocateArray, mipmapDesc->numLevels,
(size_t)mipmapDesc->offset, buf));
}
} // namespace hip
projects/clr/hipamd/src/hip_mempool.cpp
+2
View File
@@ -20,6 +20,7 @@
#include "hip_mempool_impl.hpp"
namespace hip {
/**
* API interfaces
*/
@@ -386,3 +387,4 @@ hipError_t hipMemPoolImportPointer(
mpool->retain();
HIP_RETURN(hipSuccess);
}
} // namespace hip
projects/clr/hipamd/src/hip_module.cpp
+8 -6
View File
@@ -27,6 +27,7 @@
#include "hip_event.hpp"
#include "hip_platform.hpp"
namespace hip {
hipError_t ihipModuleLoadData(hipModule_t* module, const void* mmap_ptr, size_t mmap_size);
extern hipError_t ihipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDim,
@@ -460,7 +461,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX, uint32_t gr
hStream, kernelParams, extra, nullptr, nullptr));
}
hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
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,
@@ -484,7 +485,7 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
}
hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
hipError_t __hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ,
uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ,
size_t sharedMemBytes, hipStream_t hStream, void** kernelParams,
@@ -655,7 +656,7 @@ hipError_t hipModuleLaunchCooperativeKernelMultiDevice(hipFunctionLaunchParams*
}
extern "C" hipError_t hipLaunchKernel_common(const void* hostFunction, dim3 gridDim, dim3 blockDim,
hipError_t hipLaunchKernel_common(const void* hostFunction, dim3 gridDim, dim3 blockDim,
void** args, size_t sharedMemBytes,
hipStream_t stream) {
STREAM_CAPTURE(hipLaunchKernel, stream, hostFunction, gridDim, blockDim, args, sharedMemBytes);
@@ -663,20 +664,20 @@ extern "C" hipError_t hipLaunchKernel_common(const void* hostFunction, dim3 grid
nullptr, 0);
}
extern "C" hipError_t hipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDim,
hipError_t hipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDim,
void** args, size_t sharedMemBytes, hipStream_t stream) {
HIP_INIT_API(hipLaunchKernel, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream);
HIP_RETURN(hipLaunchKernel_common(hostFunction, gridDim, blockDim, args, sharedMemBytes, stream));
}
extern "C" hipError_t hipLaunchKernel_spt(const void* hostFunction, dim3 gridDim, dim3 blockDim,
hipError_t hipLaunchKernel_spt(const void* hostFunction, dim3 gridDim, dim3 blockDim,
void** args, size_t sharedMemBytes, hipStream_t stream) {
HIP_INIT_API(hipLaunchKernel, hostFunction, gridDim, blockDim, args, sharedMemBytes, stream);
PER_THREAD_DEFAULT_STREAM(stream);
HIP_RETURN(hipLaunchKernel_common(hostFunction, gridDim, blockDim, args, sharedMemBytes, stream));
}
extern "C" hipError_t hipExtLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDim,
hipError_t hipExtLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDim,
void** args, size_t sharedMemBytes, hipStream_t stream,
hipEvent_t startEvent, hipEvent_t stopEvent, int flags) {
HIP_INIT_API(hipExtLaunchKernel, hostFunction, gridDim, blockDim, args, sharedMemBytes,
@@ -840,3 +841,4 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const
HIP_RETURN(err);
}
} // namespace hip
projects/clr/hipamd/src/hip_peer.cpp
+2 -25
View File
@@ -22,31 +22,7 @@
#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);
}
namespace hip {
hipError_t canAccessPeer(int* canAccessPeer, int deviceId, int peerDeviceId){
amd::Device* device = nullptr;
@@ -254,3 +230,4 @@ hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx) {
HIP_RETURN(hipSuccess);
}
} // namespace hip
projects/clr/hipamd/src/hip_platform.cpp
+43 -17
View File
@@ -26,7 +26,7 @@
#include "platform/runtime.hpp"
#include <unordered_map>
namespace hip {
constexpr unsigned __hipFatMAGIC2 = 0x48495046; // "HIPF"
PlatformState* PlatformState::platform_; // Initiaized as nullptr by default
@@ -67,17 +67,17 @@ static bool isCompatibleCodeObject(const std::string& codeobj_target_id, const c
return codeobj_target_id == short_name;
}
extern "C" hip::FatBinaryInfo** __hipRegisterFatBinary(const void* data) {
void** __hipRegisterFatBinary(const void* data) {
const __CudaFatBinaryWrapper* fbwrapper = reinterpret_cast<const __CudaFatBinaryWrapper*>(data);
if (fbwrapper->magic != __hipFatMAGIC2 || fbwrapper->version != 1) {
LogPrintfError("Cannot Register fat binary. FatMagic: %u version: %u ", fbwrapper->magic,
fbwrapper->version);
return nullptr;
}
return PlatformState::instance().addFatBinary(fbwrapper->binary);
return reinterpret_cast<void**>(PlatformState::instance().addFatBinary(fbwrapper->binary));
}
extern "C" void __hipRegisterFunction(hip::FatBinaryInfo** modules, const void* hostFunction,
void __hipRegisterFunction(hip::FatBinaryInfo** modules, const void* hostFunction,
char* deviceFunction, const char* deviceName,
unsigned int threadLimit, uint3* tid, uint3* bid,
dim3* blockDim, dim3* gridDim, int* wSize) {
@@ -107,7 +107,7 @@ extern "C" void __hipRegisterFunction(hip::FatBinaryInfo** modules, const void*
// global variable in host code. The shadow host variable is used to keep
// track of the value of the device side global variable between kernel
// executions.
extern "C" void __hipRegisterVar(
void __hipRegisterVar(
hip::FatBinaryInfo** modules, // The device modules containing code object
void* var, // The shadow variable in host code
char* hostVar, // Variable name in host code
@@ -123,7 +123,7 @@ extern "C" void __hipRegisterVar(
guarantee((err == hipSuccess), "Cannot register Static Global Var, error:%d \n", err);
}
extern "C" void __hipRegisterSurface(
void __hipRegisterSurface(
hip::FatBinaryInfo** modules, // The device modules containing code object
void* var, // The shadow variable in host code
char* hostVar, // Variable name in host code
@@ -135,7 +135,7 @@ extern "C" void __hipRegisterSurface(
guarantee((err == hipSuccess), "Cannot register Static Glbal Var, err:%d \n", err);
}
extern "C" void __hipRegisterManagedVar(
void __hipRegisterManagedVar(
void* hipModule, // Pointer to hip module returned from __hipRegisterFatbinary
void** pointer, // Pointer to a chunk of managed memory with size \p size and alignment \p
// align HIP runtime allocates such managed memory and assign it to \p pointer
@@ -162,7 +162,7 @@ extern "C" void __hipRegisterManagedVar(
guarantee((status == hipSuccess), "Cannot register Static Managed Var, error: %d \n", status);
}
extern "C" void __hipRegisterTexture(
void __hipRegisterTexture(
hip::FatBinaryInfo** modules, // The device modules containing code object
void* var, // The shadow variable in host code
char* hostVar, // Variable name in host code
@@ -174,12 +174,38 @@ extern "C" void __hipRegisterTexture(
guarantee((err == hipSuccess), "Cannot register Static Global Var, status: %d \n", err);
}
extern "C" void __hipUnregisterFatBinary(hip::FatBinaryInfo** modules) {
void __hipUnregisterFatBinary(hip::FatBinaryInfo** modules) {
hipError_t err = PlatformState::instance().removeFatBinary(modules);
guarantee((err == hipSuccess), "Cannot Unregister Fat Binary, error:%d \n", err);
}
extern "C" hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem,
void __hipRegisterFunction(void** modules, const void* hostFunction, char* deviceFunction,
const char* deviceName, unsigned int threadLimit, uint3* tid, uint3* bid,
dim3* blockDim, dim3* gridDim, int* wSize) {
return __hipRegisterFunction(reinterpret_cast<hip::FatBinaryInfo**>(modules), hostFunction,
deviceFunction, deviceName, threadLimit, tid, bid, blockDim, gridDim,
wSize);
}
void __hipRegisterSurface(void** modules, void* var, char* hostVar, char* deviceVar, int type,
int ext) {
return __hipRegisterSurface(reinterpret_cast<hip::FatBinaryInfo**>(modules), var, hostVar,
deviceVar, type, ext);
}
void __hipRegisterTexture(void** modules, void* var, char* hostVar, char* deviceVar, int type,
int norm, int ext) {
return __hipRegisterTexture(reinterpret_cast<hip::FatBinaryInfo**>(modules), var, hostVar,
deviceVar, type, norm, ext);
}
void __hipRegisterVar(void** modules, void* var, char* hostVar, char* deviceVar, int ext,
size_t size, int constant, int global) {
return __hipRegisterVar(reinterpret_cast<hip::FatBinaryInfo**>(modules), var, hostVar,
deviceVar, ext, size, constant, global);
}
void __hipUnregisterFatBinary(void** modules) {
return __hipUnregisterFatBinary(reinterpret_cast<hip::FatBinaryInfo**>(modules));
}
hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem,
hipStream_t stream) {
HIP_INIT_API(hipConfigureCall, gridDim, blockDim, sharedMem, stream);
@@ -188,7 +214,7 @@ extern "C" hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t share
HIP_RETURN(hipSuccess);
}
extern "C" hipError_t __hipPushCallConfiguration(dim3 gridDim, dim3 blockDim, size_t sharedMem,
hipError_t __hipPushCallConfiguration(dim3 gridDim, dim3 blockDim, size_t sharedMem,
hipStream_t stream) {
HIP_INIT_API(__hipPushCallConfiguration, gridDim, blockDim, sharedMem, stream);
@@ -197,7 +223,7 @@ extern "C" hipError_t __hipPushCallConfiguration(dim3 gridDim, dim3 blockDim, si
HIP_RETURN(hipSuccess);
}
extern "C" hipError_t __hipPopCallConfiguration(dim3* gridDim, dim3* blockDim, size_t* sharedMem,
hipError_t __hipPopCallConfiguration(dim3* gridDim, dim3* blockDim, size_t* sharedMem,
hipStream_t* stream) {
HIP_INIT_API(__hipPopCallConfiguration, gridDim, blockDim, sharedMem, stream);
@@ -211,7 +237,7 @@ extern "C" hipError_t __hipPopCallConfiguration(dim3* gridDim, dim3* blockDim, s
HIP_RETURN(hipSuccess);
}
extern "C" hipError_t hipSetupArgument(const void* arg, size_t size, size_t offset) {
hipError_t hipSetupArgument(const void* arg, size_t size, size_t offset) {
HIP_INIT_API(hipSetupArgument, arg, size, offset);
PlatformState::instance().setupArgument(arg, size, offset);
@@ -219,7 +245,7 @@ extern "C" hipError_t hipSetupArgument(const void* arg, size_t size, size_t offs
HIP_RETURN(hipSuccess);
}
extern "C" hipError_t hipLaunchByPtr(const void* hostFunction) {
hipError_t hipLaunchByPtr(const void* hostFunction) {
HIP_INIT_API(hipLaunchByPtr, hostFunction);
ihipExec_t exec;
@@ -294,7 +320,7 @@ hipError_t ihipCreateGlobalVarObj(const char* name, hipModule_t hmod, amd::Memor
HIP_RETURN(hipSuccess);
}
} // namespace hip
namespace hip_impl {
hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(
@@ -389,7 +415,7 @@ hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(
}
} // namespace hip_impl
extern "C" {
namespace hip {
hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize, const void* f,
size_t dynSharedMemPerBlk, int blockSizeLimit) {
HIP_INIT_API(hipOccupancyMaxPotentialBlockSize, f, dynSharedMemPerBlk, blockSizeLimit);
@@ -558,7 +584,6 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks,
*numBlocks = num_blocks;
HIP_RETURN(ret);
}
}
hipError_t ihipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDim, void** args,
size_t sharedMemBytes, hipStream_t stream, hipEvent_t startEvent,
@@ -927,3 +952,4 @@ bool PlatformState::CloseUniqueFileHandle(const std::shared_ptr<UniqueFD>& ufd)
}
return true;
}
} //namespace hip
projects/clr/hipamd/src/hip_platform.hpp
+3 -1
View File
@@ -29,7 +29,7 @@ namespace hip_impl {
hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(
int* maxBlocksPerCU, int* numBlocksPerGrid, int* bestBlockSize, const amd::Device& device,
hipFunction_t func, int inputBlockSize, size_t dynamicSMemSize, bool bCalcPotentialBlkSz);
} /* namespace hip_impl*/
} // namespace hip_impl
// Unique file descriptor class
struct UniqueFD {
@@ -41,6 +41,7 @@ struct UniqueFD {
const size_t fsize_; //!< File Size
};
namespace hip {
class PlatformState {
amd::Monitor lock_{"Guards PlatformState globals", true};
@@ -114,3 +115,4 @@ class PlatformState {
std::unordered_map<std::string, std::shared_ptr<UniqueFD>> ufd_map_; //!< Unique File Desc Map
};
} // namespace hip
projects/clr/hipamd/src/hip_prof_gen.py
+6 -1
View File
@@ -370,7 +370,12 @@ def parse_src(api_map, src_path, src_patt, out):
if pattern.search(fnm):
file = root + '/' + fnm
message(file)
content = parse_content(file, api_map, out);
content = ''
filename = os.path.basename(file)
if re.match("hip_table_interface.cpp", filename):
message("SKIP FILE:" + filename)
else:
content = parse_content(file, api_map, out);
if (hip_patch_mode != 0) and (content != ''):
f = open(file, 'w')
f.write(content)
projects/clr/hipamd/src/hip_profile.cpp
+2
View File
@@ -22,6 +22,7 @@
#include "hip_internal.hpp"
namespace hip {
hipError_t hipProfilerStart() {
HIP_INIT_API(hipProfilerStart);
@@ -38,3 +39,4 @@ hipError_t hipProfilerStop() {
HIP_RETURN(hipErrorNotSupported);
}
} //namespace hip
projects/clr/hipamd/src/hip_runtime.cpp
+3 -2
View File
@@ -22,8 +22,9 @@
#include <windows.h>
#include <iostream>
namespace hip {
void ihipDestroyDevice();
}
#ifdef DEBUG
static int reportHook(int reportType, char* message, int* returnValue) {
@@ -52,7 +53,7 @@ extern "C" BOOL WINAPI DllMain(HINSTANCE hinst, DWORD reason, LPVOID reserved) {
((thread = new amd::HostThread()) != nullptr && thread == amd::Thread::current()))) {
return true;
}
ihipDestroyDevice();
hip::ihipDestroyDevice();
} break;
case DLL_THREAD_DETACH: {
amd::Thread* thread = amd::Thread::current();
projects/clr/hipamd/src/hip_stream.cpp
+2 -3
View File
@@ -24,9 +24,9 @@
#include "thread/monitor.hpp"
#include "hip_prof_api.h"
namespace hip {
static amd::Monitor streamSetLock{"Guards global stream set"};
static std::unordered_set<hip::Stream*> streamSet;
namespace hip {
// ================================================================================================
Stream::Stream(hip::Device* dev, Priority p, unsigned int f, bool null_stream,
@@ -194,8 +194,6 @@ bool Stream::existsActiveStreamForDevice(hip::Device* device) {
return false;
}
};// hip namespace
// ================================================================================================
void iHipWaitActiveStreams(hip::Stream* blocking_stream, bool wait_null_stream) {
amd::Command::EventWaitList eventWaitList(0);
@@ -864,3 +862,4 @@ hipError_t hipStreamGetDevice(hipStream_t stream, hipDevice_t* device) {
HIP_RETURN(hipSuccess);
}
} // hip namespace
projects/clr/hipamd/src/hip_stream_ops.cpp
+2
View File
@@ -22,6 +22,7 @@
#include "hip_internal.hpp"
#include "platform/command_utils.hpp"
namespace hip {
hipError_t ihipStreamOperation(hipStream_t stream, cl_command_type cmdType, void* ptr,
uint64_t value, uint64_t mask, unsigned int flags, size_t sizeBytes) {
size_t offset = 0;
@@ -135,3 +136,4 @@ hipError_t hipStreamWriteValue64(hipStream_t stream, void* ptr, uint64_t value,
0, // flags un-used for now set it to 0
sizeof(uint64_t)));
}
} // namespace hip
projects/clr/hipamd/src/hip_surface.cpp
+4 -2
View File
@@ -23,8 +23,6 @@
#include "hip_internal.hpp"
#include <hip/surface_types.h>
hipError_t ihipFree(void* ptr);
struct __hip_surface {
uint32_t imageSRD[HIP_IMAGE_OBJECT_SIZE_DWORD];
amd::Image* image;
@@ -40,6 +38,9 @@ struct __hip_surface {
}
};
namespace hip {
hipError_t ihipFree(void* ptr);
hipError_t ihipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject,
const hipResourceDesc* pResDesc) {
amd::Device* device = hip::getCurrentDevice()->devices()[0];
@@ -103,3 +104,4 @@ hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject) {
HIP_RETURN(ihipDestroySurfaceObject(surfaceObject));
}
} //namespace hip
projects/clr/hipamd/src/hip_table_interface.cpp
+1689
View File
File diff suppressed because it is too large Load Diff
projects/clr/hipamd/src/hip_texture.cpp
+5 -3
View File
@@ -25,8 +25,6 @@
#include "hip_conversions.hpp"
#include "platform/sampler.hpp"
hipError_t ihipFree(void* ptr);
struct __hip_texture {
uint32_t imageSRD[HIP_IMAGE_OBJECT_SIZE_DWORD];
uint32_t samplerSRD[HIP_SAMPLER_OBJECT_SIZE_DWORD];
@@ -57,6 +55,9 @@ struct __hip_texture {
}
};
namespace hip {
hipError_t ihipFree(void* ptr);
amd::Image* ihipImageCreate(const cl_channel_order channelOrder,
const cl_channel_type channelType,
const cl_mem_object_type imageType,
@@ -1317,7 +1318,7 @@ hipError_t hipTexRefGetMipmapLevelClamp(float* pminMipmapLevelClamp,
HIP_RETURN(hipErrorInvalidValue);
}
hipError_t hipTexRefGetMipmappedArray(hipMipmappedArray_t* pArray,
hipError_t hipTexRefGetMipMappedArray(hipMipmappedArray_t* pArray,
const textureReference* texRef) {
// TODO overload operator<<(ostream&, textureReference&).
HIP_INIT_API(hipTexRefGetMipmappedArray, pArray, &texRef);
@@ -1579,3 +1580,4 @@ hipError_t hipTexObjectGetTextureDesc(HIP_TEXTURE_DESC* pTexDesc,
HIP_RETURN(hipSuccess);
}
} // namespace hip
projects/clr/hipamd/src/hip_vm.cpp
+3
View File
@@ -21,6 +21,7 @@
#include <hip/hip_runtime.h>
#include "hip_internal.hpp"
#include "hip_vm.hpp"
namespace hip {
static_assert(static_cast<uint32_t>(hipMemAccessFlagsProtNone)
== static_cast<uint32_t>(amd::Device::VmmAccess::kNone),
@@ -324,3 +325,5 @@ hipError_t hipMemUnmap(void* ptr, size_t size) {
HIP_RETURN(hipSuccess);
}
} //namespace hip
projects/clr/hipamd/src/hip_vm.hpp
+2 -2
View File
@@ -26,10 +26,10 @@
#include "platform/object.hpp"
hipError_t ihipFree(void* ptr);
namespace hip {
hipError_t ihipFree(void* ptr);
struct MemMapAllocUserData {
void* ptr_; // Original pointer of the allocation
size_t size_; // Aligned size of the allocation