AI/rocm-systems
2
0
انشعاب 0
کد مسائل تقاضاهای واکشی Actions Packages پروژه‌ها انتشارها دانشنامه فعالیت

Merge commit '0ac913e64c988e6b9ea8e86d74292cdf44c5d255' into develop

فهرست منبع
This commit is contained in:
systems-assistant[bot]
2025-08-18 18:08:03 +00:00
والد 24332bb65b 0ac913e64c
کامیت 23d4f23d7f
64فایلهای تغییر یافته به همراه2064 افزوده شده و 458 حذف شده
دانلود پرونده وصله دانلود فایل تغییرات diff Expand all files Collapse all files
projects/clr/CHANGELOG.md
+39 -7
مشاهده پرونده
@@ -2,6 +2,37 @@
Full documentation for HIP is available at [rocm.docs.amd.com](https://rocm.docs.amd.com/projects/HIP/en/latest/index.html)
## HIP 7.1 for ROCm 7.1
### Added
* New HIP APIs
- `hipModuleGetFunctionCount` returns the number of functions within a module
- `hipMemsetD2D8` Used for setting 2D memory range with specified 8-bit values
- `hipMemsetD2D8Async` Used for setting 2D memory range with specified 8-bit values asynchronously
- `hipMemsetD2D16` Used for setting 2D memory range with specified 16-bit values
- `hipMemsetD2D16Async` Used for setting 2D memory range with specified 16-bit values asynchronously
- `hipMemsetD2D32` Used for setting 2D memory range with specified 32-bit values
- `hipMemsetD2D32Async` Used for setting 2D memory range with specified 32-bit values asynchronously
- `hipStreamSetAttribute` sets attributes such as synchronization policy for a given stream
- `hipStreamGetAttribute` returns attributes such as priority for a given stream
- `hipModuleLoadFatBinary` loads fatbin binary to a module
- `hipMemcpyBatchAsync` Performs a batch of 1D or 2D memory copied asynchronously
- `hipMemcpy3DBatchAsync` Performs a batch of 3D memory copied asynchronously
- `hipMemcpy3DPeer` Copies memory between devices
- `hipMemcpy3DPeerAsync`Copied memory between devices asynchronously
- `hipMemsetD2D32Async` Used for setting 2D memory range with specified 32-bit values
asynchronously
- `hipMemPrefetchAsync_v2` prefetches memory to the specified location
- `hipMemAdvise_v2` advise about the usage of a given memory range
* Changed HIP APIs
- `hipMemCreate` now can take hipDeviceMallocUncached as a flag to allocate uncached memory
### Optimized
* Improved hip module loading latency
* Optimized kernel metadata retrieval during module post load
## HIP 7.0 for ROCm 7.0
### Added
@@ -12,15 +43,16 @@ Full documentation for HIP is available at [rocm.docs.amd.com](https://rocm.docs
- `hipDrvLaunchKernelEx` dispatches the device kernel represented by a HIP function object.
- `hipMemGetHandleForAddressRange` gets a handle for the address range requested.
- `num_threads` Total number of threads in the group. The legacy API size is alias.
- `hipGetDriverEntryPoint ` gets function pointer of a HIP API.
* New support for Open Compute Project (OCP) floating-point `FP4`/`FP6`/`FP8` as the following. For details, see [Low precision floating point document](https://rocm.docs.amd.com/projects/HIP/en/latest/reference/low_fp_types.html).
- Data types for `FP4`/`FP6`/`FP8`.
- HIP APIs for `FP4`/`FP6`/`FP8`, which are compatible with corresponding CUDA APIs.
- HIP Extensions APIs for microscaling formats, which are supported on AMD GPUs.
* New `wptr` and `rptr` values in `ClPrint`, for better logging in dispatch barrier methods.
* New debug mask, to print precise code object information for logging.
* The `_sync()` version of crosslane builtins such as `shfl_sync()` and `__reduce_add_sync` are enabled by default. These can be disabled by setting the preprocessor macro `HIP_DISABLE_WARP_SYNC_BUILTINS`.
* The `_sync()` version of crosslane builtins such as `shfl_sync()` are enabled by default. These can be disabled by setting the preprocessor macro `HIP_DISABLE_WARP_SYNC_BUILTINS`.
* Added `constexpr` operators for `fp16`/`bf16`.
* Added `__syncwarp` operation.
* Added warp level primitives: `__syncwarp` and reduce intrinsics (e.g. `__reduce_add_sync()`)
* Extended fine grained system memory pool.
* `num_threads` total number of threads in the group. The legacy API size is alias.
* Added PCI CHIP ID information as the device attribute.
@@ -145,9 +177,9 @@ HIP runtime has the following functional improvements which greatly improve runt
* Refactored memory validation, creates a unique function to validate a variety of memory copy operations.
* Improved kernel logging using demangling shader names.
* Advanced support for SPIRV, now kernel compilation caching is enabled by default. This feature is controlled by the environment variable `AMD_COMGR_CACHE`, for details, see [hip_rtc document](https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_rtc.html).
* Programmatic support for scratch limits on MI300 and MI350 series up GPU devices. More enumeration values were added in `hipLimit_t` as following,
- `hipExtLimitScratchMin`, minimum allowed value in bytes for scratch limit on the device.
- `hipExtLimitScratchMax`, maximum allowed value in bytes for scratch limit on the device.
* Programmatic support for scratch limits on MI300 and MI350 series up GPU devices. More enumeration values were added in `hipLimit_t` as following,
- `hipExtLimitScratchMin`, minimum allowed value in bytes for scratch limit on the device.
- `hipExtLimitScratchMax`, maximum allowed value in bytes for scratch limit on the device.
- `hipExtLimitScratchCurrent`, current scratch limit threshold in bytes on the device. Must be between the value `hipExtLimitScratchMin` and `hipExtLimitScratchMax`.
Developers can now use the environment variable `HSA_SCRATCH_SINGLE_LIMIT_ASYNC` to change the default allocation size with expected scratch limit in ROCR runtime. On top of it, this value can also be overwritten programmatically in the application using the HIP API `hipDeviceSetLimit(hipExtLimitScratchCurrent, value)` to reset the scratch limit value.
* HIP runtime now enables peer-to-peer (P2P) memory copies to utilize all available SDMA engines, rather than being limited to a single engine. It also selects the best engine first to give optimal bandwidth.
@@ -189,7 +221,7 @@ HIP runtime has the following functional improvements which greatly improve runt
* The memory leak in virtual memory management (VMM). HIP runtime now uses the size of handle for allocated memory range instead of actual size for physical memory, which fixed the issue of address clash with VMM.
* Large memory allocation issue. HIP runtime now checks GPU video RAM and system RAM properly and sets size limits during memory allocation either on the host or the GPU device.
* Support of `hipDeviceMallocContiguous` flags in `hipExtMallocWithFlags()`. It now enables `HSA_AMD_MEMORY_POOL_CONTIGUOUS_FLAG` in the memory pool allocation on GPU device.
* Radom memory segmentation fault in handling `GraphExec` object release and `hipDeviceSyncronization`. HIP runtime now uses internal device synchronize function in `__hipUnregisterFatBinary`.
* Radom memory segmentation fault in handling `GraphExec` object release and `hipDeviceSyncronization`. HIP runtime now uses internal device synchronize function in `__hipUnregisterFatBinary`.
## HIP 6.4.1 for ROCm 6.4.1
@@ -279,7 +311,7 @@ The following are the list of backwards incompatible changes planned for the upc
- `hipModuleLoad`
- `hipLaunchCooperativeKernelMultiDevice`
- `hipExtLaunchCooperativeKernelMultiDevice`
* HIPRTC implementation, the compilation of hiprtc now uses namespace ` __hip_internal`, instead of the standard headers `std`.
* Stream capture mode update in the following hip APIs. Stream can only be captured in relax mode, to match the behavior of the corresponding CUDA APIs,
- `hipMallocManaged`
projects/clr/hipamd/include/hip/amd_detail/amd_hip_runtime_pt_api.h
+5
مشاهده پرونده
@@ -82,6 +82,9 @@ THE SOFTWARE.
#define hipStreamIsCapturing __HIP_API_SPT(hipStreamIsCapturing)
#define hipStreamGetCaptureInfo __HIP_API_SPT(hipStreamGetCaptureInfo)
#define hipStreamGetCaptureInfo_v2 __HIP_API_SPT(hipStreamGetCaptureInfo_v2)
// Driver Entry Point API
#define hipGetDriverEntryPoint __HIP_API_SPT(hipGetDriverEntryPoint)
#endif
#ifdef __cplusplus
@@ -192,6 +195,8 @@ hipError_t hipStreamGetCaptureInfo_v2_spt(hipStream_t stream, hipStreamCaptureSt
const hipGraphNode_t** dependencies_out,
size_t* numDependencies_out);
hipError_t hipLaunchHostFunc_spt(hipStream_t stream, hipHostFn_t fn, void* userData);
hipError_t hipGetDriverEntryPoint_spt(const char* symbol, void** funcPtr, unsigned long long flags,
hipDriverEntryPointQueryResult* status);
#ifdef __cplusplus
projects/clr/hipamd/include/hip/amd_detail/amd_surface_functions.h
+9 -8
مشاهده پرونده
@@ -114,10 +114,11 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ void surf1Dread(T* data, hipSurfaceObject_t surfObj, int x,
int boundaryMode = hipBoundaryModeZero) {
__HIP_SURFACE_OBJECT_PARAMETERS_INIT
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
auto tmp = __ockl_image_load_1D(i, x);
*data = __hipMapFrom<T>(tmp);
__HIP_SURFACE_OBJECT_PARAMETERS_INIT;
(void)boundaryMode;
x = __hipGetPixelAddr(x, __ockl_image_channel_data_type_1D(i), __ockl_image_channel_order_1D(i));
auto tmp = __ockl_image_load_1D(i, x);
*data = __hipMapFrom<T>(tmp);
}
/** \brief Writes the value data to the one-dimensional surface at coordinate x.
@@ -157,7 +158,7 @@ static __device__ __hip_img_chk__ void surf2Dread(T* data, hipSurfaceObject_t su
*data = __hipMapFrom<T>(tmp);
}
/** \brief Writes the value data to the two-dimensional surface at coordinate
/** \brief Writes the value data to the two-dimensional surface at coordinate
* x, y.
*
* \tparam T The data type of the surface.
@@ -177,7 +178,7 @@ static __device__ __hip_img_chk__ void surf2Dwrite(T data, hipSurfaceObject_t su
__ockl_image_store_2D(i, get_native_vector(coords), tmp);
}
/** \brief Reads the value from the three-dimensional surface at coordinate
/** \brief Reads the value from the three-dimensional surface at coordinate
* x, y, z.
*
* \tparam T The data type of the surface.
@@ -238,7 +239,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredread(T* data, hipSurfaceObje
*data = __hipMapFrom<T>(tmp);
}
/** \brief Writes the value data to the one-dimensional layered surface at
/** \brief Writes the value data to the one-dimensional layered surface at
* coordinate x and layer index.
*
* \tparam T The data type of the surface.
@@ -257,7 +258,7 @@ static __device__ __hip_img_chk__ void surf1DLayeredwrite(T data, hipSurfaceObje
__ockl_image_store_lod_1D(i, x, layer, tmp);
}
/** \brief Reads the value from the two-dimensional layered surface at
/** \brief Reads the value from the two-dimensional layered surface at
* coordinate x, y and layer index.
*
* \tparam T The data type of the surface.
projects/clr/hipamd/include/hip/amd_detail/hip_api_trace.hpp
+61 -3
مشاهده پرونده
@@ -63,7 +63,7 @@
#define HIP_API_TABLE_STEP_VERSION 0
#define HIP_COMPILER_API_TABLE_STEP_VERSION 0
#define HIP_TOOLS_API_TABLE_STEP_VERSION 0
#define HIP_RUNTIME_API_TABLE_STEP_VERSION 13
#define HIP_RUNTIME_API_TABLE_STEP_VERSION 14
// HIP API interface
// HIP compiler dispatch functions
@@ -484,6 +484,8 @@ typedef hipError_t (*t_hipMemAddressReserve)(void** ptr, size_t size, size_t ali
unsigned long long flags);
typedef hipError_t (*t_hipMemAdvise)(const void* dev_ptr, size_t count, hipMemoryAdvise advice,
int device);
typedef hipError_t (*t_hipMemAdvise_v2)(const void* dev_ptr, size_t count, hipMemoryAdvise advice,
hipMemLocation device);
typedef hipError_t (*t_hipMemAllocHost)(void** ptr, size_t size);
typedef hipError_t (*t_hipMemAllocPitch)(hipDeviceptr_t* dptr, size_t* pitch, size_t widthInBytes,
size_t height, unsigned int elementSizeBytes);
@@ -535,6 +537,9 @@ typedef hipError_t (*t_hipMemPoolSetAttribute)(hipMemPool_t mem_pool, hipMemPool
typedef hipError_t (*t_hipMemPoolTrimTo)(hipMemPool_t mem_pool, size_t min_bytes_to_hold);
typedef hipError_t (*t_hipMemPrefetchAsync)(const void* dev_ptr, size_t count, int device,
hipStream_t stream);
typedef hipError_t (*t_hipMemPrefetchAsync_v2)(const void* dev_ptr, size_t count,
hipMemLocation location, unsigned int flags,
hipStream_t stream);
typedef hipError_t (*t_hipMemPtrGetInfo)(void* ptr, size_t* size);
typedef hipError_t (*t_hipMemRangeGetAttribute)(void* data, size_t data_size,
hipMemRangeAttribute attribute, const void* dev_ptr,
@@ -635,6 +640,7 @@ typedef hipError_t (*t_hipMipmappedArrayGetLevel)(hipArray_t* pLevelArray,
unsigned int level);
typedef hipError_t (*t_hipModuleGetFunction)(hipFunction_t* function, hipModule_t module,
const char* kname);
typedef hipError_t (*t_hipModuleGetFunctionCount)(unsigned int* count, hipModule_t module);
typedef hipError_t (*t_hipModuleGetGlobal)(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod,
const char* name);
typedef hipError_t (*t_hipModuleGetTexRef)(textureReference** texRef, hipModule_t hmod,
@@ -1045,6 +1051,39 @@ typedef hipError_t (*t_hipMemGetHandleForAddressRange)(void* handle, hipDevicept
size_t size,
hipMemRangeHandleType handleType,
unsigned long long flags);
typedef hipError_t (*t_hipMemsetD2D8)(hipDeviceptr_t dst, size_t dstPitch, unsigned char value,
size_t width, size_t height);
typedef hipError_t (*t_hipMemsetD2D8Async)(hipDeviceptr_t dst, size_t dstPitch, unsigned char value,
size_t width, size_t height, hipStream_t stream);
typedef hipError_t (*t_hipMemsetD2D16)(hipDeviceptr_t dst, size_t dstPitch, unsigned short value,
size_t width, size_t height);
typedef hipError_t (*t_hipMemsetD2D16Async)(hipDeviceptr_t dst, size_t dstPitch,
unsigned short value, size_t width, size_t height,
hipStream_t stream);
typedef hipError_t (*t_hipMemsetD2D32)(hipDeviceptr_t dst, size_t dstPitch, unsigned int value,
size_t width, size_t height);
typedef hipError_t (*t_hipMemsetD2D32Async)(hipDeviceptr_t dst, size_t dstPitch, unsigned int value,
size_t width, size_t height, hipStream_t stream);
typedef hipError_t (*t_hipStreamSetAttribute)(hipStream_t stream, hipStreamAttrID attr,
const hipStreamAttrValue *value);
typedef hipError_t (*t_hipStreamGetAttribute)(hipStream_t stream, hipStreamAttrID attr,
hipStreamAttrValue *value_out);
typedef hipError_t (*t_hipModuleLoadFatBinary)(hipModule_t* module, const void* fatbin);
typedef hipError_t (*t_hipMemcpyBatchAsync) (void **dsts, void **srcs, size_t *sizes, size_t count,
hipMemcpyAttributes *attrs, size_t *attrsIdxs,
size_t numAttrs, size_t *failIdx, hipStream_t stream);
typedef hipError_t (*t_hipMemcpy3DBatchAsync) (size_t numOps, struct hipMemcpy3DBatchOp *opList,
size_t *failIdx, unsigned long long flags,
hipStream_t stream);
typedef hipError_t (*t_hipMemcpy3DPeer) (hipMemcpy3DPeerParms *p);
typedef hipError_t (*t_hipMemcpy3DPeerAsync) (hipMemcpy3DPeerParms *p, hipStream_t stream);
typedef hipError_t (*t_hipGetDriverEntryPoint)(const char* symbol, void** funcPtr,
unsigned long long flags,
hipDriverEntryPointQueryResult* status);
typedef hipError_t (*t_hipGetDriverEntryPoint_spt)(const char* symbol, void** funcPtr,
unsigned long long flags,
hipDriverEntryPointQueryResult* status);
// HIP Compiler dispatch table
struct HipCompilerDispatchTable {
@@ -1588,10 +1627,29 @@ struct HipDispatchTable {
t_hipMemGetHandleForAddressRange hipMemGetHandleForAddressRange_fn;
// HIP_RUNTIME_API_TABLE_STEP_VERSION = 13
t_hipModuleGetFunctionCount hipModuleGetFunctionCount_fn;
t_hipMemsetD2D8 hipMemsetD2D8_fn;
t_hipMemsetD2D8Async hipMemsetD2D8Async_fn;
t_hipMemsetD2D16 hipMemsetD2D16_fn;
t_hipMemsetD2D16Async hipMemsetD2D16Async_fn;
t_hipMemsetD2D32 hipMemsetD2D32_fn;
t_hipMemsetD2D32Async hipMemsetD2D32Async_fn;
t_hipStreamGetAttribute hipStreamGetAttribute_fn;
t_hipStreamSetAttribute hipStreamSetAttribute_fn;
t_hipModuleLoadFatBinary hipModuleLoadFatBinary_fn;
t_hipMemcpyBatchAsync hipMemcpyBatchAsync_fn;
t_hipMemcpy3DBatchAsync hipMemcpy3DBatchAsync_fn;
t_hipMemcpy3DPeer hipMemcpy3DPeer_fn;
t_hipMemcpy3DPeerAsync hipMemcpy3DPeerAsync_fn;
t_hipGetDriverEntryPoint hipGetDriverEntryPoint_fn;
t_hipGetDriverEntryPoint_spt hipGetDriverEntryPoint_spt_fn;
t_hipMemPrefetchAsync_v2 hipMemPrefetchAsync_v2_fn;
t_hipMemAdvise_v2 hipMemAdvise_v2_fn;
// HIP_RUNTIME_API_TABLE_STEP_VERSION = 14
// removed HIP_MEMSET_NODE_PARAMS replaced by hipMemsetParams
// DO NOT EDIT ABOVE!
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 13
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 14
// ******************************************************************************************* //
//
@@ -1628,4 +1686,4 @@ struct HipToolsDispatchTable {
// 4) GENERATE COMMENT FOR NEXT STEP VERSION
// 5) ADD "DO NOT EDIT ABOVE!" COMMENT
// ******************************************************************************************* //
};
};
projects/clr/hipamd/include/hip/amd_detail/hip_prof_str.h
+531 -8
مشاهده پرونده
@@ -438,7 +438,24 @@ enum hip_api_id_t {
HIP_API_ID_hipLinkDestroy = 418,
HIP_API_ID_hipLaunchKernelExC = 419,
HIP_API_ID_hipDrvLaunchKernelEx = 420,
HIP_API_ID_LAST = 420,
HIP_API_ID_hipModuleGetFunctionCount = 421,
HIP_API_ID_hipMemsetD2D16 = 422,
HIP_API_ID_hipMemsetD2D16Async = 423,
HIP_API_ID_hipMemsetD2D32 = 424,
HIP_API_ID_hipMemsetD2D32Async = 425,
HIP_API_ID_hipMemsetD2D8 = 426,
HIP_API_ID_hipMemsetD2D8Async = 427,
HIP_API_ID_hipStreamGetAttribute = 428,
HIP_API_ID_hipStreamSetAttribute = 429,
HIP_API_ID_hipModuleLoadFatBinary = 430,
HIP_API_ID_hipMemcpy3DBatchAsync = 431,
HIP_API_ID_hipMemcpy3DPeer = 432,
HIP_API_ID_hipMemcpy3DPeerAsync = 433,
HIP_API_ID_hipMemcpyBatchAsync = 434,
HIP_API_ID_hipGetDriverEntryPoint = 435,
HIP_API_ID_hipMemPrefetchAsync_v2 = 436,
HIP_API_ID_hipMemAdvise_v2 = 437,
HIP_API_ID_LAST = 437,
HIP_API_ID_hipChooseDevice = HIP_API_ID_CONCAT(HIP_API_ID_,hipChooseDevice),
HIP_API_ID_hipGetDeviceProperties = HIP_API_ID_CONCAT(HIP_API_ID_,hipGetDeviceProperties),
@@ -451,6 +468,7 @@ enum hip_api_id_t {
HIP_API_ID_hipDestroyTextureObject = HIP_API_ID_NONE,
HIP_API_ID_hipDeviceGetCount = HIP_API_ID_NONE,
HIP_API_ID_hipDeviceGetTexture1DLinearMaxWidth = HIP_API_ID_NONE,
HIP_API_ID_hipGetDriverEntryPoint_spt = HIP_API_ID_NONE,
HIP_API_ID_hipGetTextureAlignmentOffset = HIP_API_ID_NONE,
HIP_API_ID_hipGetTextureObjectResourceDesc = HIP_API_ID_NONE,
HIP_API_ID_hipGetTextureObjectResourceViewDesc = HIP_API_ID_NONE,
@@ -588,6 +606,7 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipGetDeviceFlags: return "hipGetDeviceFlags";
case HIP_API_ID_hipGetDevicePropertiesR0000: return "hipGetDevicePropertiesR0000";
case HIP_API_ID_hipGetDevicePropertiesR0600: return "hipGetDevicePropertiesR0600";
case HIP_API_ID_hipGetDriverEntryPoint: return "hipGetDriverEntryPoint";
case HIP_API_ID_hipGetErrorString: return "hipGetErrorString";
case HIP_API_ID_hipGetFuncBySymbol: return "hipGetFuncBySymbol";
case HIP_API_ID_hipGetLastError: return "hipGetLastError";
@@ -727,6 +746,7 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipMemAddressFree: return "hipMemAddressFree";
case HIP_API_ID_hipMemAddressReserve: return "hipMemAddressReserve";
case HIP_API_ID_hipMemAdvise: return "hipMemAdvise";
case HIP_API_ID_hipMemAdvise_v2: return "hipMemAdvise_v2";
case HIP_API_ID_hipMemAllocHost: return "hipMemAllocHost";
case HIP_API_ID_hipMemAllocPitch: return "hipMemAllocPitch";
case HIP_API_ID_hipMemCreate: return "hipMemCreate";
@@ -751,6 +771,7 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipMemPoolSetAttribute: return "hipMemPoolSetAttribute";
case HIP_API_ID_hipMemPoolTrimTo: return "hipMemPoolTrimTo";
case HIP_API_ID_hipMemPrefetchAsync: return "hipMemPrefetchAsync";
case HIP_API_ID_hipMemPrefetchAsync_v2: return "hipMemPrefetchAsync_v2";
case HIP_API_ID_hipMemPtrGetInfo: return "hipMemPtrGetInfo";
case HIP_API_ID_hipMemRangeGetAttribute: return "hipMemRangeGetAttribute";
case HIP_API_ID_hipMemRangeGetAttributes: return "hipMemRangeGetAttributes";
@@ -768,11 +789,15 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipMemcpy2DToArrayAsync: return "hipMemcpy2DToArrayAsync";
case HIP_API_ID_hipMemcpy3D: return "hipMemcpy3D";
case HIP_API_ID_hipMemcpy3DAsync: return "hipMemcpy3DAsync";
case HIP_API_ID_hipMemcpy3DBatchAsync: return "hipMemcpy3DBatchAsync";
case HIP_API_ID_hipMemcpy3DPeer: return "hipMemcpy3DPeer";
case HIP_API_ID_hipMemcpy3DPeerAsync: return "hipMemcpy3DPeerAsync";
case HIP_API_ID_hipMemcpyAsync: return "hipMemcpyAsync";
case HIP_API_ID_hipMemcpyAtoA: return "hipMemcpyAtoA";
case HIP_API_ID_hipMemcpyAtoD: return "hipMemcpyAtoD";
case HIP_API_ID_hipMemcpyAtoH: return "hipMemcpyAtoH";
case HIP_API_ID_hipMemcpyAtoHAsync: return "hipMemcpyAtoHAsync";
case HIP_API_ID_hipMemcpyBatchAsync: return "hipMemcpyBatchAsync";
case HIP_API_ID_hipMemcpyDtoA: return "hipMemcpyDtoA";
case HIP_API_ID_hipMemcpyDtoD: return "hipMemcpyDtoD";
case HIP_API_ID_hipMemcpyDtoDAsync: return "hipMemcpyDtoDAsync";
@@ -801,6 +826,12 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipMemsetAsync: return "hipMemsetAsync";
case HIP_API_ID_hipMemsetD16: return "hipMemsetD16";
case HIP_API_ID_hipMemsetD16Async: return "hipMemsetD16Async";
case HIP_API_ID_hipMemsetD2D16: return "hipMemsetD2D16";
case HIP_API_ID_hipMemsetD2D16Async: return "hipMemsetD2D16Async";
case HIP_API_ID_hipMemsetD2D32: return "hipMemsetD2D32";
case HIP_API_ID_hipMemsetD2D32Async: return "hipMemsetD2D32Async";
case HIP_API_ID_hipMemsetD2D8: return "hipMemsetD2D8";
case HIP_API_ID_hipMemsetD2D8Async: return "hipMemsetD2D8Async";
case HIP_API_ID_hipMemsetD32: return "hipMemsetD32";
case HIP_API_ID_hipMemsetD32Async: return "hipMemsetD32Async";
case HIP_API_ID_hipMemsetD8: return "hipMemsetD8";
@@ -814,6 +845,7 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipModuleLaunchCooperativeKernel: return "hipModuleLaunchCooperativeKernel";
case HIP_API_ID_hipModuleLaunchCooperativeKernelMultiDevice: return "hipModuleLaunchCooperativeKernelMultiDevice";
case HIP_API_ID_hipModuleLaunchKernel: return "hipModuleLaunchKernel";
case HIP_API_ID_hipModuleLoadFatBinary: return "hipModuleLoadFatBinary";
case HIP_API_ID_hipModuleLoad: return "hipModuleLoad";
case HIP_API_ID_hipModuleLoadData: return "hipModuleLoadData";
case HIP_API_ID_hipModuleLoadDataEx: return "hipModuleLoadDataEx";
@@ -847,6 +879,7 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipStreamCreateWithPriority: return "hipStreamCreateWithPriority";
case HIP_API_ID_hipStreamDestroy: return "hipStreamDestroy";
case HIP_API_ID_hipStreamEndCapture: return "hipStreamEndCapture";
case HIP_API_ID_hipStreamGetAttribute: return "hipStreamGetAttribute";
case HIP_API_ID_hipStreamGetCaptureInfo: return "hipStreamGetCaptureInfo";
case HIP_API_ID_hipStreamGetCaptureInfo_v2: return "hipStreamGetCaptureInfo_v2";
case HIP_API_ID_hipStreamGetDevice: return "hipStreamGetDevice";
@@ -854,6 +887,7 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipStreamGetPriority: return "hipStreamGetPriority";
case HIP_API_ID_hipStreamIsCapturing: return "hipStreamIsCapturing";
case HIP_API_ID_hipStreamQuery: return "hipStreamQuery";
case HIP_API_ID_hipStreamSetAttribute: return "hipStreamSetAttribute";
case HIP_API_ID_hipStreamSynchronize: return "hipStreamSynchronize";
case HIP_API_ID_hipStreamUpdateCaptureDependencies: return "hipStreamUpdateCaptureDependencies";
case HIP_API_ID_hipStreamWaitEvent: return "hipStreamWaitEvent";
@@ -885,6 +919,7 @@ static inline const char* hip_api_name(const uint32_t id) {
case HIP_API_ID_hipUserObjectRelease: return "hipUserObjectRelease";
case HIP_API_ID_hipUserObjectRetain: return "hipUserObjectRetain";
case HIP_API_ID_hipWaitExternalSemaphoresAsync: return "hipWaitExternalSemaphoresAsync";
case HIP_API_ID_hipModuleGetFunctionCount: return "hipModuleGetFunctionCount";
};
return "unknown";
};
@@ -1003,6 +1038,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipGetDeviceFlags", name) == 0) return HIP_API_ID_hipGetDeviceFlags;
if (strcmp("hipGetDevicePropertiesR0000", name) == 0) return HIP_API_ID_hipGetDevicePropertiesR0000;
if (strcmp("hipGetDevicePropertiesR0600", name) == 0) return HIP_API_ID_hipGetDevicePropertiesR0600;
if (strcmp("hipGetDriverEntryPoint", name) == 0) return HIP_API_ID_hipGetDriverEntryPoint;
if (strcmp("hipGetErrorString", name) == 0) return HIP_API_ID_hipGetErrorString;
if (strcmp("hipGetFuncBySymbol", name) == 0) return HIP_API_ID_hipGetFuncBySymbol;
if (strcmp("hipGetLastError", name) == 0) return HIP_API_ID_hipGetLastError;
@@ -1142,6 +1178,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipMemAddressFree", name) == 0) return HIP_API_ID_hipMemAddressFree;
if (strcmp("hipMemAddressReserve", name) == 0) return HIP_API_ID_hipMemAddressReserve;
if (strcmp("hipMemAdvise", name) == 0) return HIP_API_ID_hipMemAdvise;
if (strcmp("hipMemAdvise_v2", name) == 0) return HIP_API_ID_hipMemAdvise_v2;
if (strcmp("hipMemAllocHost", name) == 0) return HIP_API_ID_hipMemAllocHost;
if (strcmp("hipMemAllocPitch", name) == 0) return HIP_API_ID_hipMemAllocPitch;
if (strcmp("hipMemCreate", name) == 0) return HIP_API_ID_hipMemCreate;
@@ -1166,6 +1203,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipMemPoolSetAttribute", name) == 0) return HIP_API_ID_hipMemPoolSetAttribute;
if (strcmp("hipMemPoolTrimTo", name) == 0) return HIP_API_ID_hipMemPoolTrimTo;
if (strcmp("hipMemPrefetchAsync", name) == 0) return HIP_API_ID_hipMemPrefetchAsync;
if (strcmp("hipMemPrefetchAsync_v2", name) == 0) return HIP_API_ID_hipMemPrefetchAsync_v2;
if (strcmp("hipMemPtrGetInfo", name) == 0) return HIP_API_ID_hipMemPtrGetInfo;
if (strcmp("hipMemRangeGetAttribute", name) == 0) return HIP_API_ID_hipMemRangeGetAttribute;
if (strcmp("hipMemRangeGetAttributes", name) == 0) return HIP_API_ID_hipMemRangeGetAttributes;
@@ -1183,11 +1221,15 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipMemcpy2DToArrayAsync", name) == 0) return HIP_API_ID_hipMemcpy2DToArrayAsync;
if (strcmp("hipMemcpy3D", name) == 0) return HIP_API_ID_hipMemcpy3D;
if (strcmp("hipMemcpy3DAsync", name) == 0) return HIP_API_ID_hipMemcpy3DAsync;
if (strcmp("hipMemcpy3DBatchAsync", name) == 0) return HIP_API_ID_hipMemcpy3DBatchAsync;
if (strcmp("hipMemcpy3DPeer", name) == 0) return HIP_API_ID_hipMemcpy3DPeer;
if (strcmp("hipMemcpy3DPeerAsync", name) == 0) return HIP_API_ID_hipMemcpy3DPeerAsync;
if (strcmp("hipMemcpyAsync", name) == 0) return HIP_API_ID_hipMemcpyAsync;
if (strcmp("hipMemcpyAtoA", name) == 0) return HIP_API_ID_hipMemcpyAtoA;
if (strcmp("hipMemcpyAtoD", name) == 0) return HIP_API_ID_hipMemcpyAtoD;
if (strcmp("hipMemcpyAtoH", name) == 0) return HIP_API_ID_hipMemcpyAtoH;
if (strcmp("hipMemcpyAtoHAsync", name) == 0) return HIP_API_ID_hipMemcpyAtoHAsync;
if (strcmp("hipMemcpyBatchAsync", name) == 0) return HIP_API_ID_hipMemcpyBatchAsync;
if (strcmp("hipMemcpyDtoA", name) == 0) return HIP_API_ID_hipMemcpyDtoA;
if (strcmp("hipMemcpyDtoD", name) == 0) return HIP_API_ID_hipMemcpyDtoD;
if (strcmp("hipMemcpyDtoDAsync", name) == 0) return HIP_API_ID_hipMemcpyDtoDAsync;
@@ -1216,6 +1258,12 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipMemsetAsync", name) == 0) return HIP_API_ID_hipMemsetAsync;
if (strcmp("hipMemsetD16", name) == 0) return HIP_API_ID_hipMemsetD16;
if (strcmp("hipMemsetD16Async", name) == 0) return HIP_API_ID_hipMemsetD16Async;
if (strcmp("hipMemsetD2D16", name) == 0) return HIP_API_ID_hipMemsetD2D16;
if (strcmp("hipMemsetD2D16Async", name) == 0) return HIP_API_ID_hipMemsetD2D16Async;
if (strcmp("hipMemsetD2D32", name) == 0) return HIP_API_ID_hipMemsetD2D32;
if (strcmp("hipMemsetD2D32Async", name) == 0) return HIP_API_ID_hipMemsetD2D32Async;
if (strcmp("hipMemsetD2D8", name) == 0) return HIP_API_ID_hipMemsetD2D8;
if (strcmp("hipMemsetD2D8Async", name) == 0) return HIP_API_ID_hipMemsetD2D8Async;
if (strcmp("hipMemsetD32", name) == 0) return HIP_API_ID_hipMemsetD32;
if (strcmp("hipMemsetD32Async", name) == 0) return HIP_API_ID_hipMemsetD32Async;
if (strcmp("hipMemsetD8", name) == 0) return HIP_API_ID_hipMemsetD8;
@@ -1229,6 +1277,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipModuleLaunchCooperativeKernel", name) == 0) return HIP_API_ID_hipModuleLaunchCooperativeKernel;
if (strcmp("hipModuleLaunchCooperativeKernelMultiDevice", name) == 0) return HIP_API_ID_hipModuleLaunchCooperativeKernelMultiDevice;
if (strcmp("hipModuleLaunchKernel", name) == 0) return HIP_API_ID_hipModuleLaunchKernel;
if (strcmp("hipModuleLoadFatBinary", name) == 0) return HIP_API_ID_hipModuleLoadFatBinary;
if (strcmp("hipModuleLoad", name) == 0) return HIP_API_ID_hipModuleLoad;
if (strcmp("hipModuleLoadData", name) == 0) return HIP_API_ID_hipModuleLoadData;
if (strcmp("hipModuleLoadDataEx", name) == 0) return HIP_API_ID_hipModuleLoadDataEx;
@@ -1262,6 +1311,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipStreamCreateWithPriority", name) == 0) return HIP_API_ID_hipStreamCreateWithPriority;
if (strcmp("hipStreamDestroy", name) == 0) return HIP_API_ID_hipStreamDestroy;
if (strcmp("hipStreamEndCapture", name) == 0) return HIP_API_ID_hipStreamEndCapture;
if (strcmp("hipStreamGetAttribute", name) == 0) return HIP_API_ID_hipStreamGetAttribute;
if (strcmp("hipStreamGetCaptureInfo", name) == 0) return HIP_API_ID_hipStreamGetCaptureInfo;
if (strcmp("hipStreamGetCaptureInfo_v2", name) == 0) return HIP_API_ID_hipStreamGetCaptureInfo_v2;
if (strcmp("hipStreamGetDevice", name) == 0) return HIP_API_ID_hipStreamGetDevice;
@@ -1269,6 +1319,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipStreamGetPriority", name) == 0) return HIP_API_ID_hipStreamGetPriority;
if (strcmp("hipStreamIsCapturing", name) == 0) return HIP_API_ID_hipStreamIsCapturing;
if (strcmp("hipStreamQuery", name) == 0) return HIP_API_ID_hipStreamQuery;
if (strcmp("hipStreamSetAttribute", name) == 0) return HIP_API_ID_hipStreamSetAttribute;
if (strcmp("hipStreamSynchronize", name) == 0) return HIP_API_ID_hipStreamSynchronize;
if (strcmp("hipStreamUpdateCaptureDependencies", name) == 0) return HIP_API_ID_hipStreamUpdateCaptureDependencies;
if (strcmp("hipStreamWaitEvent", name) == 0) return HIP_API_ID_hipStreamWaitEvent;
@@ -1300,6 +1351,7 @@ static inline uint32_t hipApiIdByName(const char* name) {
if (strcmp("hipUserObjectRelease", name) == 0) return HIP_API_ID_hipUserObjectRelease;
if (strcmp("hipUserObjectRetain", name) == 0) return HIP_API_ID_hipUserObjectRetain;
if (strcmp("hipWaitExternalSemaphoresAsync", name) == 0) return HIP_API_ID_hipWaitExternalSemaphoresAsync;
if (strcmp("hipModuleGetFunctionCount", name) == 0) return HIP_API_ID_hipModuleGetFunctionCount;
return HIP_API_ID_NONE;
}
@@ -1866,6 +1918,15 @@ typedef struct hip_api_data_s {
hipDeviceProp_tR0600 prop__val;
int deviceId;
} hipGetDevicePropertiesR0600;
struct {
const char* symbol;
char symbol__val;
void** funcPtr;
void* funcPtr__val;
unsigned long long flags;
hipDriverEntryPointQueryResult* driverStatus;
hipDriverEntryPointQueryResult driverStatus__val;
} hipGetDriverEntryPoint;
struct {
hipFunction_t* functionPtr;
hipFunction_t functionPtr__val;
@@ -2759,6 +2820,12 @@ typedef struct hip_api_data_s {
hipMemoryAdvise advice;
int device;
} hipMemAdvise;
struct {
const void* dev_ptr;
size_t count;
hipMemoryAdvise advice;
hipMemLocation location;
} hipMemAdvise_v2;
struct {
void** ptr;
void* ptr__val;
@@ -2905,6 +2972,13 @@ typedef struct hip_api_data_s {
int device;
hipStream_t stream;
} hipMemPrefetchAsync;
struct {
const void* dev_ptr;
size_t count;
hipMemLocation location;
unsigned int flags;
hipStream_t stream;
} hipMemPrefetchAsync_v2;
struct {
void* ptr;
size_t* size;
@@ -3034,6 +3108,24 @@ typedef struct hip_api_data_s {
hipMemcpy3DParms p__val;
hipStream_t stream;
} hipMemcpy3DAsync;
struct {
size_t numOps;
hipMemcpy3DBatchOp* opList;
hipMemcpy3DBatchOp opList__val;
size_t* failIdx;
size_t failIdx__val;
unsigned long long flags;
hipStream_t stream;
} hipMemcpy3DBatchAsync;
struct {
hipMemcpy3DPeerParms* p;
hipMemcpy3DPeerParms p__val;
} hipMemcpy3DPeer;
struct {
hipMemcpy3DPeerParms* p;
hipMemcpy3DPeerParms p__val;
hipStream_t stream;
} hipMemcpy3DPeerAsync;
struct {
void* dst;
const void* src;
@@ -3067,6 +3159,23 @@ typedef struct hip_api_data_s {
size_t ByteCount;
hipStream_t stream;
} hipMemcpyAtoHAsync;
struct {
void** dsts;
void* dsts__val;
void** srcs;
void* srcs__val;
size_t* sizes;
size_t sizes__val;
size_t count;
hipMemcpyAttributes* attrs;
hipMemcpyAttributes attrs__val;
size_t* attrsIdxs;
size_t attrsIdxs__val;
size_t numAttrs;
size_t* failIdx;
size_t failIdx__val;
hipStream_t stream;
} hipMemcpyBatchAsync;
struct {
hipArray_t dstArray;
size_t dstOffset;
@@ -3244,6 +3353,51 @@ typedef struct hip_api_data_s {
size_t count;
hipStream_t stream;
} hipMemsetD16Async;
struct {
hipDeviceptr_t dst;
size_t dstPitch;
unsigned short value;
size_t width;
size_t height;
} hipMemsetD2D16;
struct {
hipDeviceptr_t dst;
size_t dstPitch;
unsigned short value;
size_t width;
size_t height;
hipStream_t stream;
} hipMemsetD2D16Async;
struct {
hipDeviceptr_t dst;
size_t dstPitch;
unsigned int value;
size_t width;
size_t height;
} hipMemsetD2D32;
struct {
hipDeviceptr_t dst;
size_t dstPitch;
unsigned int value;
size_t width;
size_t height;
hipStream_t stream;
} hipMemsetD2D32Async;
struct {
hipDeviceptr_t dst;
size_t dstPitch;
unsigned char value;
size_t width;
size_t height;
} hipMemsetD2D8;
struct {
hipDeviceptr_t dst;
size_t dstPitch;
unsigned char value;
size_t width;
size_t height;
hipStream_t stream;
} hipMemsetD2D8Async;
struct {
hipDeviceptr_t dest;
int value;
@@ -3289,6 +3443,11 @@ typedef struct hip_api_data_s {
const char* kname;
char kname__val;
} hipModuleGetFunction;
struct {
unsigned int* count;
unsigned int count__val;
hipModule_t mod;
} hipModuleGetFunctionCount;
struct {
hipDeviceptr_t* dptr;
hipDeviceptr_t dptr__val;
@@ -3339,6 +3498,11 @@ typedef struct hip_api_data_s {
void** extra;
void* extra__val;
} hipModuleLaunchKernel;
struct {
hipModule_t* module;
hipModule_t module__val;
const void* fatbin;
} hipModuleLoadFatBinary;
struct {
hipModule_t* module;
hipModule_t module__val;
@@ -3520,6 +3684,12 @@ typedef struct hip_api_data_s {
hipGraph_t* pGraph;
hipGraph_t pGraph__val;
} hipStreamEndCapture;
struct {
hipStream_t stream;
hipLaunchAttributeID attr;
const hipLaunchAttributeValue* value_out;
hipLaunchAttributeValue value_out__val;
} hipStreamGetAttribute;
struct {
hipStream_t stream;
hipStreamCaptureStatus* pCaptureStatus;
@@ -3563,6 +3733,12 @@ typedef struct hip_api_data_s {
struct {
hipStream_t stream;
} hipStreamQuery;
struct {
hipStream_t stream;
hipLaunchAttributeID attr;
const hipLaunchAttributeValue* value;
hipLaunchAttributeValue value__val;
} hipStreamSetAttribute;
struct {
hipStream_t stream;
} hipStreamSynchronize;
@@ -4337,6 +4513,13 @@ typedef struct hip_api_data_s {
cb_data.args.hipGetDevicePropertiesR0600.prop = (hipDeviceProp_tR0600*)prop; \
cb_data.args.hipGetDevicePropertiesR0600.deviceId = (int)device; \
};
// hipGetDriverEntryPoint[('const char*', 'symbol'), ('void**', 'funcPtr'), ('unsigned long long', 'flags'), ('hipDriverEntryPointQueryResult*', 'driverStatus')]
#define INIT_hipGetDriverEntryPoint_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipGetDriverEntryPoint.symbol = (symbol) ? strdup(symbol) : NULL; \
cb_data.args.hipGetDriverEntryPoint.funcPtr = (void**)funcPtr; \
cb_data.args.hipGetDriverEntryPoint.flags = (unsigned long long)flags; \
cb_data.args.hipGetDriverEntryPoint.driverStatus = (hipDriverEntryPointQueryResult*)status; \
};
// hipGetErrorString[]
#define INIT_hipGetErrorString_CB_ARGS_DATA(cb_data) { \
};
@@ -5206,6 +5389,13 @@ typedef struct hip_api_data_s {
cb_data.args.hipMemAdvise.advice = (hipMemoryAdvise)advice; \
cb_data.args.hipMemAdvise.device = (int)device; \
};
// hipMemAdvise_v2[('const void*', 'dev_ptr'), ('size_t', 'count'), ('hipMemoryAdvise', 'advice'), ('hipMemLocation', 'location')]
#define INIT_hipMemAdvise_v2_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemAdvise_v2.dev_ptr = (const void*)dev_ptr; \
cb_data.args.hipMemAdvise_v2.count = (size_t)count; \
cb_data.args.hipMemAdvise_v2.advice = (hipMemoryAdvise)advice; \
cb_data.args.hipMemAdvise_v2.location = (hipMemLocation)location; \
};
// hipMemAllocHost[('void**', 'ptr'), ('size_t', 'size')]
#define INIT_hipMemAllocHost_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemAllocHost.ptr = (void**)ptr; \
@@ -5351,6 +5541,14 @@ typedef struct hip_api_data_s {
cb_data.args.hipMemPrefetchAsync.device = (int)device; \
cb_data.args.hipMemPrefetchAsync.stream = (hipStream_t)stream; \
};
// hipMemPrefetchAsync_v2[('const void*', 'dev_ptr'), ('size_t', 'count'), ('hipMemLocation', 'location'), ('unsigned int', 'flags'), ('hipStream_t', 'stream')]
#define INIT_hipMemPrefetchAsync_v2_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemPrefetchAsync_v2.dev_ptr = (const void*)dev_ptr; \
cb_data.args.hipMemPrefetchAsync_v2.count = (size_t)count; \
cb_data.args.hipMemPrefetchAsync_v2.location = (hipMemLocation)location; \
cb_data.args.hipMemPrefetchAsync_v2.flags = (unsigned int)flags; \
cb_data.args.hipMemPrefetchAsync_v2.stream = (hipStream_t)stream; \
};
// hipMemPtrGetInfo[('void*', 'ptr'), ('size_t*', 'size')]
#define INIT_hipMemPtrGetInfo_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemPtrGetInfo.ptr = (void*)ptr; \
@@ -5489,6 +5687,23 @@ typedef struct hip_api_data_s {
cb_data.args.hipMemcpy3DAsync.p = (const hipMemcpy3DParms*)p; \
cb_data.args.hipMemcpy3DAsync.stream = (hipStream_t)stream; \
};
// hipMemcpy3DBatchAsync[('size_t', 'numOps'), ('hipMemcpy3DBatchOp*', 'opList'), ('size_t*', 'failIdx'), ('unsigned long long', 'flags'), ('hipStream_t', 'stream')]
#define INIT_hipMemcpy3DBatchAsync_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemcpy3DBatchAsync.numOps = (size_t)numOps; \
cb_data.args.hipMemcpy3DBatchAsync.opList = (hipMemcpy3DBatchOp*)opList; \
cb_data.args.hipMemcpy3DBatchAsync.failIdx = (size_t*)failIdx; \
cb_data.args.hipMemcpy3DBatchAsync.flags = (unsigned long long)flags; \
cb_data.args.hipMemcpy3DBatchAsync.stream = (hipStream_t)stream; \
};
// hipMemcpy3DPeer[('hipMemcpy3DPeerParms*', 'p')]
#define INIT_hipMemcpy3DPeer_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemcpy3DPeer.p = (hipMemcpy3DPeerParms*)p; \
};
// hipMemcpy3DPeerAsync[('hipMemcpy3DPeerParms*', 'p'), ('hipStream_t', 'stream')]
#define INIT_hipMemcpy3DPeerAsync_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemcpy3DPeerAsync.p = (hipMemcpy3DPeerParms*)p; \
cb_data.args.hipMemcpy3DPeerAsync.stream = (hipStream_t)stream; \
};
// hipMemcpyAsync[('void*', 'dst'), ('const void*', 'src'), ('size_t', 'sizeBytes'), ('hipMemcpyKind', 'kind'), ('hipStream_t', 'stream')]
#define INIT_hipMemcpyAsync_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemcpyAsync.dst = (void*)dst; \
@@ -5527,6 +5742,18 @@ typedef struct hip_api_data_s {
cb_data.args.hipMemcpyAtoHAsync.ByteCount = (size_t)ByteCount; \
cb_data.args.hipMemcpyAtoHAsync.stream = (hipStream_t)stream; \
};
// hipMemcpyBatchAsync[('void**', 'dsts'), ('void**', 'srcs'), ('size_t*', 'sizes'), ('size_t', 'count'), ('hipMemcpyAttributes*', 'attrs'), ('size_t*', 'attrsIdxs'), ('size_t', 'numAttrs'), ('size_t*', 'failIdx'), ('hipStream_t', 'stream')]
#define INIT_hipMemcpyBatchAsync_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemcpyBatchAsync.dsts = (void**)dsts; \
cb_data.args.hipMemcpyBatchAsync.srcs = (void**)srcs; \
cb_data.args.hipMemcpyBatchAsync.sizes = (size_t*)sizes; \
cb_data.args.hipMemcpyBatchAsync.count = (size_t)count; \
cb_data.args.hipMemcpyBatchAsync.attrs = (hipMemcpyAttributes*)attrs; \
cb_data.args.hipMemcpyBatchAsync.attrsIdxs = (size_t*)attrsIdxs; \
cb_data.args.hipMemcpyBatchAsync.numAttrs = (size_t)numAttrs; \
cb_data.args.hipMemcpyBatchAsync.failIdx = (size_t*)failIdx; \
cb_data.args.hipMemcpyBatchAsync.stream = (hipStream_t)stream; \
};
// hipMemcpyDtoA[('hipArray_t', 'dstArray'), ('size_t', 'dstOffset'), ('hipDeviceptr_t', 'srcDevice'), ('size_t', 'ByteCount')]
#define INIT_hipMemcpyDtoA_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemcpyDtoA.dstArray = (hipArray_t)dstArray; \
@@ -5601,16 +5828,16 @@ typedef struct hip_api_data_s {
cb_data.args.hipMemcpyHtoAAsync.ByteCount = (size_t)ByteCount; \
cb_data.args.hipMemcpyHtoAAsync.stream = (hipStream_t)stream; \
};
// hipMemcpyHtoD[('hipDeviceptr_t', 'dst'), ('void*', 'src'), ('size_t', 'sizeBytes')]
// hipMemcpyHtoD[('hipDeviceptr_t', 'dst'), ('const void*', 'src'), ('size_t', 'sizeBytes')]
#define INIT_hipMemcpyHtoD_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemcpyHtoD.dst = (hipDeviceptr_t)dstDevice; \
cb_data.args.hipMemcpyHtoD.src = (void*)srcHost; \
cb_data.args.hipMemcpyHtoD.src = (const void*)srcHost; \
cb_data.args.hipMemcpyHtoD.sizeBytes = (size_t)ByteCount; \
};
// hipMemcpyHtoDAsync[('hipDeviceptr_t', 'dst'), ('void*', 'src'), ('size_t', 'sizeBytes'), ('hipStream_t', 'stream')]
// hipMemcpyHtoDAsync[('hipDeviceptr_t', 'dst'), ('const void*', 'src'), ('size_t', 'sizeBytes'), ('hipStream_t', 'stream')]
#define INIT_hipMemcpyHtoDAsync_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemcpyHtoDAsync.dst = (hipDeviceptr_t)dstDevice; \
cb_data.args.hipMemcpyHtoDAsync.src = (void*)srcHost; \
cb_data.args.hipMemcpyHtoDAsync.src = (const void*)srcHost; \
cb_data.args.hipMemcpyHtoDAsync.sizeBytes = (size_t)ByteCount; \
cb_data.args.hipMemcpyHtoDAsync.stream = (hipStream_t)stream; \
};
@@ -5730,6 +5957,57 @@ typedef struct hip_api_data_s {
cb_data.args.hipMemsetD16Async.count = (size_t)count; \
cb_data.args.hipMemsetD16Async.stream = (hipStream_t)stream; \
};
// hipMemsetD2D16[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned short', 'value'), ('size_t', 'width'), ('size_t', 'height')]
#define INIT_hipMemsetD2D16_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemsetD2D16.dst = (hipDeviceptr_t)dst; \
cb_data.args.hipMemsetD2D16.dstPitch = (size_t)dstPitch; \
cb_data.args.hipMemsetD2D16.value = (unsigned short)value; \
cb_data.args.hipMemsetD2D16.width = (size_t)width; \
cb_data.args.hipMemsetD2D16.height = (size_t)height; \
};
// hipMemsetD2D16Async[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned short', 'value'), ('size_t', 'width'), ('size_t', 'height'), ('hipStream_t', 'stream')]
#define INIT_hipMemsetD2D16Async_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemsetD2D16Async.dst = (hipDeviceptr_t)dst; \
cb_data.args.hipMemsetD2D16Async.dstPitch = (size_t)dstPitch; \
cb_data.args.hipMemsetD2D16Async.value = (unsigned short)value; \
cb_data.args.hipMemsetD2D16Async.width = (size_t)width; \
cb_data.args.hipMemsetD2D16Async.height = (size_t)height; \
cb_data.args.hipMemsetD2D16Async.stream = (hipStream_t)stream; \
};
// hipMemsetD2D32[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned int', 'value'), ('size_t', 'width'), ('size_t', 'height')]
#define INIT_hipMemsetD2D32_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemsetD2D32.dst = (hipDeviceptr_t)dst; \
cb_data.args.hipMemsetD2D32.dstPitch = (size_t)dstPitch; \
cb_data.args.hipMemsetD2D32.value = (unsigned int)value; \
cb_data.args.hipMemsetD2D32.width = (size_t)width; \
cb_data.args.hipMemsetD2D32.height = (size_t)height; \
};
// hipMemsetD2D32Async[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned int', 'value'), ('size_t', 'width'), ('size_t', 'height'), ('hipStream_t', 'stream')]
#define INIT_hipMemsetD2D32Async_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemsetD2D32Async.dst = (hipDeviceptr_t)dst; \
cb_data.args.hipMemsetD2D32Async.dstPitch = (size_t)dstPitch; \
cb_data.args.hipMemsetD2D32Async.value = (unsigned int)value; \
cb_data.args.hipMemsetD2D32Async.width = (size_t)width; \
cb_data.args.hipMemsetD2D32Async.height = (size_t)height; \
cb_data.args.hipMemsetD2D32Async.stream = (hipStream_t)stream; \
};
// hipMemsetD2D8[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned char', 'value'), ('size_t', 'width'), ('size_t', 'height')]
#define INIT_hipMemsetD2D8_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemsetD2D8.dst = (hipDeviceptr_t)dst; \
cb_data.args.hipMemsetD2D8.dstPitch = (size_t)dstPitch; \
cb_data.args.hipMemsetD2D8.value = (unsigned char)value; \
cb_data.args.hipMemsetD2D8.width = (size_t)width; \
cb_data.args.hipMemsetD2D8.height = (size_t)height; \
};
// hipMemsetD2D8Async[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned char', 'value'), ('size_t', 'width'), ('size_t', 'height'), ('hipStream_t', 'stream')]
#define INIT_hipMemsetD2D8Async_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemsetD2D8Async.dst = (hipDeviceptr_t)dst; \
cb_data.args.hipMemsetD2D8Async.dstPitch = (size_t)dstPitch; \
cb_data.args.hipMemsetD2D8Async.value = (unsigned char)value; \
cb_data.args.hipMemsetD2D8Async.width = (size_t)width; \
cb_data.args.hipMemsetD2D8Async.height = (size_t)height; \
cb_data.args.hipMemsetD2D8Async.stream = (hipStream_t)stream; \
};
// hipMemsetD32[('hipDeviceptr_t', 'dest'), ('int', 'value'), ('size_t', 'count')]
#define INIT_hipMemsetD32_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipMemsetD32.dest = (hipDeviceptr_t)dst; \
@@ -5824,6 +6102,11 @@ typedef struct hip_api_data_s {
cb_data.args.hipModuleLaunchKernel.kernelParams = (void**)kernelParams; \
cb_data.args.hipModuleLaunchKernel.extra = (void**)extra; \
};
// hipModuleLoadFatBinary[('hipModule_t*', 'module'), ('const void*', 'fatbin')]
#define INIT_hipModuleLoadFatBinary_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipModuleLoadFatBinary.module = (hipModule_t*)module; \
cb_data.args.hipModuleLoadFatBinary.fatbin = (const void*)fatbin; \
};
// hipModuleLoad[('hipModule_t*', 'module'), ('const char*', 'fname')]
#define INIT_hipModuleLoad_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipModuleLoad.module = (hipModule_t*)module; \
@@ -6016,6 +6299,9 @@ typedef struct hip_api_data_s {
cb_data.args.hipStreamEndCapture.stream = (hipStream_t)stream; \
cb_data.args.hipStreamEndCapture.pGraph = (hipGraph_t*)pGraph; \
};
// hipStreamGetAttribute[('hipStream_t', 'stream'), ('hipLaunchAttributeID', 'attr'), ('const hipLaunchAttributeValue*', 'value_out')]
#define INIT_hipStreamGetAttribute_CB_ARGS_DATA(cb_data) { \
};
// hipStreamGetCaptureInfo[('hipStream_t', 'stream'), ('hipStreamCaptureStatus*', 'pCaptureStatus'), ('unsigned long long*', 'pId')]
#define INIT_hipStreamGetCaptureInfo_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipStreamGetCaptureInfo.stream = (hipStream_t)stream; \
@@ -6055,6 +6341,9 @@ typedef struct hip_api_data_s {
#define INIT_hipStreamQuery_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipStreamQuery.stream = (hipStream_t)stream; \
};
// hipStreamSetAttribute[('hipStream_t', 'stream'), ('hipLaunchAttributeID', 'attr'), ('const hipLaunchAttributeValue*', 'value')]
#define INIT_hipStreamSetAttribute_CB_ARGS_DATA(cb_data) { \
};
// hipStreamSynchronize[('hipStream_t', 'stream')]
#define INIT_hipStreamSynchronize_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipStreamSynchronize.stream = (hipStream_t)stream; \
@@ -6236,6 +6525,12 @@ typedef struct hip_api_data_s {
cb_data.args.hipWaitExternalSemaphoresAsync.numExtSems = (unsigned int)numExtSems; \
cb_data.args.hipWaitExternalSemaphoresAsync.stream = (hipStream_t)stream; \
};
// hipModuleGetFunctionCount[('unsigned int*', 'count'), ('hipModule_t', 'mod')]
#define INIT_hipModuleGetFunctionCount_CB_ARGS_DATA(cb_data) { \
cb_data.args.hipModuleGetFunctionCount.count = (unsigned int*)count; \
cb_data.args.hipModuleGetFunctionCount.mod = (hipModule_t)mod; \
};
#define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data)
// Macros for non-public API primitives
@@ -6255,6 +6550,8 @@ typedef struct hip_api_data_s {
#define INIT_hipDeviceGetCount_CB_ARGS_DATA(cb_data) {};
// hipDeviceGetTexture1DLinearMaxWidth()
#define INIT_hipDeviceGetTexture1DLinearMaxWidth_CB_ARGS_DATA(cb_data) {};
// hipGetDriverEntryPoint_spt()
#define INIT_hipGetDriverEntryPoint_spt_CB_ARGS_DATA(cb_data) {};
// hipGetTextureAlignmentOffset()
#define INIT_hipGetTextureAlignmentOffset_CB_ARGS_DATA(cb_data) {};
// hipGetTextureObjectResourceDesc()
@@ -6366,7 +6663,7 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
// hipCtxEnablePeerAccess[('hipCtx_t', 'peerCtx'), ('unsigned int', 'flags')]
case HIP_API_ID_hipCtxEnablePeerAccess:
break;
// hipCtxGetApiVersion[('hipCtx_t', 'ctx'), ('int*', 'apiVersion')]
// hipCtxGetApiVersion[('hipCtx_t', 'ctx'), ('unsigned int*', 'apiVersion')]
case HIP_API_ID_hipCtxGetApiVersion:
if (data->args.hipCtxGetApiVersion.apiVersion) data->args.hipCtxGetApiVersion.apiVersion__val = *(data->args.hipCtxGetApiVersion.apiVersion);
break;
@@ -6728,6 +7025,12 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
case HIP_API_ID_hipGetDevicePropertiesR0600:
if (data->args.hipGetDevicePropertiesR0600.prop) data->args.hipGetDevicePropertiesR0600.prop__val = *(data->args.hipGetDevicePropertiesR0600.prop);
break;
// hipGetDriverEntryPoint[('const char*', 'symbol'), ('void**', 'funcPtr'), ('unsigned long long', 'flags'), ('hipDriverEntryPointQueryResult*', 'driverStatus')]
case HIP_API_ID_hipGetDriverEntryPoint:
if (data->args.hipGetDriverEntryPoint.symbol) data->args.hipGetDriverEntryPoint.symbol__val = *(data->args.hipGetDriverEntryPoint.symbol);
if (data->args.hipGetDriverEntryPoint.funcPtr) data->args.hipGetDriverEntryPoint.funcPtr__val = *(data->args.hipGetDriverEntryPoint.funcPtr);
if (data->args.hipGetDriverEntryPoint.driverStatus) data->args.hipGetDriverEntryPoint.driverStatus__val = *(data->args.hipGetDriverEntryPoint.driverStatus);
break;
// hipGetErrorString[]
case HIP_API_ID_hipGetErrorString:
break;
@@ -7306,6 +7609,9 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
// hipMemAdvise[('const void*', 'dev_ptr'), ('size_t', 'count'), ('hipMemoryAdvise', 'advice'), ('int', 'device')]
case HIP_API_ID_hipMemAdvise:
break;
// hipMemAdvise_v2[('const void*', 'dev_ptr'), ('size_t', 'count'), ('hipMemoryAdvise', 'advice'), ('hipMemLocation', 'location')]
case HIP_API_ID_hipMemAdvise_v2:
break;
// hipMemAllocHost[('void**', 'ptr'), ('size_t', 'size')]
case HIP_API_ID_hipMemAllocHost:
if (data->args.hipMemAllocHost.ptr) data->args.hipMemAllocHost.ptr__val = *(data->args.hipMemAllocHost.ptr);
@@ -7403,6 +7709,9 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
// hipMemPrefetchAsync[('const void*', 'dev_ptr'), ('size_t', 'count'), ('int', 'device'), ('hipStream_t', 'stream')]
case HIP_API_ID_hipMemPrefetchAsync:
break;
// hipMemPrefetchAsync_v2[('const void*', 'dev_ptr'), ('size_t', 'count'), ('hipMemLocation', 'location'), ('unsigned int', 'flags'), ('hipStream_t', 'stream')]
case HIP_API_ID_hipMemPrefetchAsync_v2:
break;
// hipMemPtrGetInfo[('void*', 'ptr'), ('size_t*', 'size')]
case HIP_API_ID_hipMemPtrGetInfo:
if (data->args.hipMemPtrGetInfo.size) data->args.hipMemPtrGetInfo.size__val = *(data->args.hipMemPtrGetInfo.size);
@@ -7462,6 +7771,19 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
case HIP_API_ID_hipMemcpy3DAsync:
if (data->args.hipMemcpy3DAsync.p) data->args.hipMemcpy3DAsync.p__val = *(data->args.hipMemcpy3DAsync.p);
break;
// hipMemcpy3DBatchAsync[('size_t', 'numOps'), ('hipMemcpy3DBatchOp*', 'opList'), ('size_t*', 'failIdx'), ('unsigned long long', 'flags'), ('hipStream_t', 'stream')]
case HIP_API_ID_hipMemcpy3DBatchAsync:
if (data->args.hipMemcpy3DBatchAsync.opList) data->args.hipMemcpy3DBatchAsync.opList__val = *(data->args.hipMemcpy3DBatchAsync.opList);
if (data->args.hipMemcpy3DBatchAsync.failIdx) data->args.hipMemcpy3DBatchAsync.failIdx__val = *(data->args.hipMemcpy3DBatchAsync.failIdx);
break;
// hipMemcpy3DPeer[('hipMemcpy3DPeerParms*', 'p')]
case HIP_API_ID_hipMemcpy3DPeer:
if (data->args.hipMemcpy3DPeer.p) data->args.hipMemcpy3DPeer.p__val = *(data->args.hipMemcpy3DPeer.p);
break;
// hipMemcpy3DPeerAsync[('hipMemcpy3DPeerParms*', 'p'), ('hipStream_t', 'stream')]
case HIP_API_ID_hipMemcpy3DPeerAsync:
if (data->args.hipMemcpy3DPeerAsync.p) data->args.hipMemcpy3DPeerAsync.p__val = *(data->args.hipMemcpy3DPeerAsync.p);
break;
// hipMemcpyAsync[('void*', 'dst'), ('const void*', 'src'), ('size_t', 'sizeBytes'), ('hipMemcpyKind', 'kind'), ('hipStream_t', 'stream')]
case HIP_API_ID_hipMemcpyAsync:
break;
@@ -7477,6 +7799,15 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
// hipMemcpyAtoHAsync[('void*', 'dstHost'), ('hipArray_t', 'srcArray'), ('size_t', 'srcOffset'), ('size_t', 'ByteCount'), ('hipStream_t', 'stream')]
case HIP_API_ID_hipMemcpyAtoHAsync:
break;
// hipMemcpyBatchAsync[('void**', 'dsts'), ('void**', 'srcs'), ('size_t*', 'sizes'), ('size_t', 'count'), ('hipMemcpyAttributes*', 'attrs'), ('size_t*', 'attrsIdxs'), ('size_t', 'numAttrs'), ('size_t*', 'failIdx'), ('hipStream_t', 'stream')]
case HIP_API_ID_hipMemcpyBatchAsync:
if (data->args.hipMemcpyBatchAsync.dsts) data->args.hipMemcpyBatchAsync.dsts__val = *(data->args.hipMemcpyBatchAsync.dsts);
if (data->args.hipMemcpyBatchAsync.srcs) data->args.hipMemcpyBatchAsync.srcs__val = *(data->args.hipMemcpyBatchAsync.srcs);
if (data->args.hipMemcpyBatchAsync.sizes) data->args.hipMemcpyBatchAsync.sizes__val = *(data->args.hipMemcpyBatchAsync.sizes);
if (data->args.hipMemcpyBatchAsync.attrs) data->args.hipMemcpyBatchAsync.attrs__val = *(data->args.hipMemcpyBatchAsync.attrs);
if (data->args.hipMemcpyBatchAsync.attrsIdxs) data->args.hipMemcpyBatchAsync.attrsIdxs__val = *(data->args.hipMemcpyBatchAsync.attrsIdxs);
if (data->args.hipMemcpyBatchAsync.failIdx) data->args.hipMemcpyBatchAsync.failIdx__val = *(data->args.hipMemcpyBatchAsync.failIdx);
break;
// hipMemcpyDtoA[('hipArray_t', 'dstArray'), ('size_t', 'dstOffset'), ('hipDeviceptr_t', 'srcDevice'), ('size_t', 'ByteCount')]
case HIP_API_ID_hipMemcpyDtoA:
break;
@@ -7507,10 +7838,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
// hipMemcpyHtoAAsync[('hipArray_t', 'dstArray'), ('size_t', 'dstOffset'), ('const void*', 'srcHost'), ('size_t', 'ByteCount'), ('hipStream_t', 'stream')]
case HIP_API_ID_hipMemcpyHtoAAsync:
break;
// hipMemcpyHtoD[('hipDeviceptr_t', 'dst'), ('void*', 'src'), ('size_t', 'sizeBytes')]
// hipMemcpyHtoD[('hipDeviceptr_t', 'dst'), ('const void*', 'src'), ('size_t', 'sizeBytes')]
case HIP_API_ID_hipMemcpyHtoD:
break;
// hipMemcpyHtoDAsync[('hipDeviceptr_t', 'dst'), ('void*', 'src'), ('size_t', 'sizeBytes'), ('hipStream_t', 'stream')]
// hipMemcpyHtoDAsync[('hipDeviceptr_t', 'dst'), ('const void*', 'src'), ('size_t', 'sizeBytes'), ('hipStream_t', 'stream')]
case HIP_API_ID_hipMemcpyHtoDAsync:
break;
// hipMemcpyParam2D[('const hip_Memcpy2D*', 'pCopy')]
@@ -7563,6 +7894,24 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
// hipMemsetD16Async[('hipDeviceptr_t', 'dest'), ('unsigned short', 'value'), ('size_t', 'count'), ('hipStream_t', 'stream')]
case HIP_API_ID_hipMemsetD16Async:
break;
// hipMemsetD2D16[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned short', 'value'), ('size_t', 'width'), ('size_t', 'height')]
case HIP_API_ID_hipMemsetD2D16:
break;
// hipMemsetD2D16Async[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned short', 'value'), ('size_t', 'width'), ('size_t', 'height'), ('hipStream_t', 'stream')]
case HIP_API_ID_hipMemsetD2D16Async:
break;
// hipMemsetD2D32[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned int', 'value'), ('size_t', 'width'), ('size_t', 'height')]
case HIP_API_ID_hipMemsetD2D32:
break;
// hipMemsetD2D32Async[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned int', 'value'), ('size_t', 'width'), ('size_t', 'height'), ('hipStream_t', 'stream')]
case HIP_API_ID_hipMemsetD2D32Async:
break;
// hipMemsetD2D8[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned char', 'value'), ('size_t', 'width'), ('size_t', 'height')]
case HIP_API_ID_hipMemsetD2D8:
break;
// hipMemsetD2D8Async[('hipDeviceptr_t', 'dst'), ('size_t', 'dstPitch'), ('unsigned char', 'value'), ('size_t', 'width'), ('size_t', 'height'), ('hipStream_t', 'stream')]
case HIP_API_ID_hipMemsetD2D8Async:
break;
// hipMemsetD32[('hipDeviceptr_t', 'dest'), ('int', 'value'), ('size_t', 'count')]
case HIP_API_ID_hipMemsetD32:
break;
@@ -7616,6 +7965,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
if (data->args.hipModuleLaunchKernel.kernelParams) data->args.hipModuleLaunchKernel.kernelParams__val = *(data->args.hipModuleLaunchKernel.kernelParams);
if (data->args.hipModuleLaunchKernel.extra) data->args.hipModuleLaunchKernel.extra__val = *(data->args.hipModuleLaunchKernel.extra);
break;
// hipModuleLoadFatBinary[('hipModule_t*', 'module'), ('const void*', 'fatbin')]
case HIP_API_ID_hipModuleLoadFatBinary:
if (data->args.hipModuleLoadFatBinary.module) data->args.hipModuleLoadFatBinary.module__val = *(data->args.hipModuleLoadFatBinary.module);
break;
// hipModuleLoad[('hipModule_t*', 'module'), ('const char*', 'fname')]
case HIP_API_ID_hipModuleLoad:
if (data->args.hipModuleLoad.module) data->args.hipModuleLoad.module__val = *(data->args.hipModuleLoad.module);
@@ -7743,6 +8096,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
case HIP_API_ID_hipStreamEndCapture:
if (data->args.hipStreamEndCapture.pGraph) data->args.hipStreamEndCapture.pGraph__val = *(data->args.hipStreamEndCapture.pGraph);
break;
// hipStreamGetAttribute[('hipStream_t', 'stream'), ('hipLaunchAttributeID', 'attr'), ('const hipLaunchAttributeValue*', 'value_out')]
case HIP_API_ID_hipStreamGetAttribute:
if (data->args.hipStreamGetAttribute.value_out) data->args.hipStreamGetAttribute.value_out__val = *(data->args.hipStreamGetAttribute.value_out);
break;
// hipStreamGetCaptureInfo[('hipStream_t', 'stream'), ('hipStreamCaptureStatus*', 'pCaptureStatus'), ('unsigned long long*', 'pId')]
case HIP_API_ID_hipStreamGetCaptureInfo:
if (data->args.hipStreamGetCaptureInfo.pCaptureStatus) data->args.hipStreamGetCaptureInfo.pCaptureStatus__val = *(data->args.hipStreamGetCaptureInfo.pCaptureStatus);
@@ -7775,6 +8132,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
// hipStreamQuery[('hipStream_t', 'stream')]
case HIP_API_ID_hipStreamQuery:
break;
// hipStreamSetAttribute[('hipStream_t', 'stream'), ('hipLaunchAttributeID', 'attr'), ('const hipLaunchAttributeValue*', 'value')]
case HIP_API_ID_hipStreamSetAttribute:
if (data->args.hipStreamSetAttribute.value) data->args.hipStreamSetAttribute.value__val = *(data->args.hipStreamSetAttribute.value);
break;
// hipStreamSynchronize[('hipStream_t', 'stream')]
case HIP_API_ID_hipStreamSynchronize:
break;
@@ -7907,6 +8268,10 @@ static inline void hipApiArgsInit(hip_api_id_t id, hip_api_data_t* data) {
if (data->args.hipWaitExternalSemaphoresAsync.extSemArray) data->args.hipWaitExternalSemaphoresAsync.extSemArray__val = *(data->args.hipWaitExternalSemaphoresAsync.extSemArray);
if (data->args.hipWaitExternalSemaphoresAsync.paramsArray) data->args.hipWaitExternalSemaphoresAsync.paramsArray__val = *(data->args.hipWaitExternalSemaphoresAsync.paramsArray);
break;
// hipModuleGetFunctionCount[('unsigned int*', 'count'), ('hipModule_t', 'mod')]
case HIP_API_ID_hipModuleGetFunctionCount:
if (data->args.hipModuleGetFunctionCount.count) data->args.hipModuleGetFunctionCount.count__val = *(data->args.hipModuleGetFunctionCount.count);
break;
default: break;
};
}
@@ -8705,6 +9070,17 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
oss << ", deviceId="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDevicePropertiesR0600.deviceId);
oss << ")";
break;
case HIP_API_ID_hipGetDriverEntryPoint:
oss << "hipGetDriverEntryPoint(";
if (data->args.hipGetDriverEntryPoint.symbol == NULL) oss << "symbol=NULL";
else { oss << "symbol="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDriverEntryPoint.symbol__val); }
if (data->args.hipGetDriverEntryPoint.funcPtr == NULL) oss << ", funcPtr=NULL";
else { oss << ", funcPtr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDriverEntryPoint.funcPtr__val); }
oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDriverEntryPoint.flags);
if (data->args.hipGetDriverEntryPoint.driverStatus == NULL) oss << ", driverStatus=NULL";
else { oss << ", driverStatus="; roctracer::hip_support::detail::operator<<(oss, data->args.hipGetDriverEntryPoint.driverStatus__val); }
oss << ")";
break;
case HIP_API_ID_hipGetErrorString:
oss << "hipGetErrorString(";
oss << ")";
@@ -9880,6 +10256,14 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
oss << ", device="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemAdvise.device);
oss << ")";
break;
case HIP_API_ID_hipMemAdvise_v2:
oss << "hipMemAdvise_v2(";
oss << "dev_ptr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemAdvise_v2.dev_ptr);
oss << ", count="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemAdvise_v2.count);
oss << ", advice="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemAdvise_v2.advice);
oss << ", location="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemAdvise_v2.location);
oss << ")";
break;
case HIP_API_ID_hipMemAllocHost:
oss << "hipMemAllocHost(";
if (data->args.hipMemAllocHost.ptr == NULL) oss << "ptr=NULL";
@@ -10074,6 +10458,15 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemPrefetchAsync.stream);
oss << ")";
break;
case HIP_API_ID_hipMemPrefetchAsync_v2:
oss << "hipMemPrefetchAsync_v2(";
oss << "dev_ptr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemPrefetchAsync_v2.dev_ptr);
oss << ", count="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemPrefetchAsync_v2.count);
oss << ", location="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemPrefetchAsync_v2.location);
oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemPrefetchAsync_v2.flags);
oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemPrefetchAsync_v2.stream);
oss << ")";
break;
case HIP_API_ID_hipMemPtrGetInfo:
oss << "hipMemPtrGetInfo(";
oss << "ptr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemPtrGetInfo.ptr);
@@ -10237,6 +10630,30 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy3DAsync.stream);
oss << ")";
break;
case HIP_API_ID_hipMemcpy3DBatchAsync:
oss << "hipMemcpy3DBatchAsync(";
oss << "numOps="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy3DBatchAsync.numOps);
if (data->args.hipMemcpy3DBatchAsync.opList == NULL) oss << ", opList=NULL";
else { oss << ", opList="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy3DBatchAsync.opList__val); }
if (data->args.hipMemcpy3DBatchAsync.failIdx == NULL) oss << ", failIdx=NULL";
else { oss << ", failIdx="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy3DBatchAsync.failIdx__val); }
oss << ", flags="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy3DBatchAsync.flags);
oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy3DBatchAsync.stream);
oss << ")";
break;
case HIP_API_ID_hipMemcpy3DPeer:
oss << "hipMemcpy3DPeer(";
if (data->args.hipMemcpy3DPeer.p == NULL) oss << "p=NULL";
else { oss << "p="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy3DPeer.p__val); }
oss << ")";
break;
case HIP_API_ID_hipMemcpy3DPeerAsync:
oss << "hipMemcpy3DPeerAsync(";
if (data->args.hipMemcpy3DPeerAsync.p == NULL) oss << "p=NULL";
else { oss << "p="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy3DPeerAsync.p__val); }
oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpy3DPeerAsync.stream);
oss << ")";
break;
case HIP_API_ID_hipMemcpyAsync:
oss << "hipMemcpyAsync(";
oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAsync.dst);
@@ -10280,6 +10697,25 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyAtoHAsync.stream);
oss << ")";
break;
case HIP_API_ID_hipMemcpyBatchAsync:
oss << "hipMemcpyBatchAsync(";
if (data->args.hipMemcpyBatchAsync.dsts == NULL) oss << "dsts=NULL";
else { oss << "dsts="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyBatchAsync.dsts__val); }
if (data->args.hipMemcpyBatchAsync.srcs == NULL) oss << ", srcs=NULL";
else { oss << ", srcs="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyBatchAsync.srcs__val); }
if (data->args.hipMemcpyBatchAsync.sizes == NULL) oss << ", sizes=NULL";
else { oss << ", sizes="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyBatchAsync.sizes__val); }
oss << ", count="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyBatchAsync.count);
if (data->args.hipMemcpyBatchAsync.attrs == NULL) oss << ", attrs=NULL";
else { oss << ", attrs="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyBatchAsync.attrs__val); }
if (data->args.hipMemcpyBatchAsync.attrsIdxs == NULL) oss << ", attrsIdxs=NULL";
else { oss << ", attrsIdxs="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyBatchAsync.attrsIdxs__val); }
oss << ", numAttrs="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyBatchAsync.numAttrs);
if (data->args.hipMemcpyBatchAsync.failIdx == NULL) oss << ", failIdx=NULL";
else { oss << ", failIdx="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyBatchAsync.failIdx__val); }
oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyBatchAsync.stream);
oss << ")";
break;
case HIP_API_ID_hipMemcpyDtoA:
oss << "hipMemcpyDtoA(";
oss << "dstArray="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemcpyDtoA.dstArray);
@@ -10513,6 +10949,63 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD16Async.stream);
oss << ")";
break;
case HIP_API_ID_hipMemsetD2D16:
oss << "hipMemsetD2D16(";
oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16.dst);
oss << ", dstPitch="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16.dstPitch);
oss << ", value="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16.value);
oss << ", width="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16.width);
oss << ", height="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16.height);
oss << ")";
break;
case HIP_API_ID_hipMemsetD2D16Async:
oss << "hipMemsetD2D16Async(";
oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16Async.dst);
oss << ", dstPitch="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16Async.dstPitch);
oss << ", value="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16Async.value);
oss << ", width="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16Async.width);
oss << ", height="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16Async.height);
oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D16Async.stream);
oss << ")";
break;
case HIP_API_ID_hipMemsetD2D32:
oss << "hipMemsetD2D32(";
oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32.dst);
oss << ", dstPitch="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32.dstPitch);
oss << ", value="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32.value);
oss << ", width="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32.width);
oss << ", height="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32.height);
oss << ")";
break;
case HIP_API_ID_hipMemsetD2D32Async:
oss << "hipMemsetD2D32Async(";
oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32Async.dst);
oss << ", dstPitch="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32Async.dstPitch);
oss << ", value="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32Async.value);
oss << ", width="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32Async.width);
oss << ", height="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32Async.height);
oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D32Async.stream);
oss << ")";
break;
case HIP_API_ID_hipMemsetD2D8:
oss << "hipMemsetD2D8(";
oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8.dst);
oss << ", dstPitch="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8.dstPitch);
oss << ", value="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8.value);
oss << ", width="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8.width);
oss << ", height="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8.height);
oss << ")";
break;
case HIP_API_ID_hipMemsetD2D8Async:
oss << "hipMemsetD2D8Async(";
oss << "dst="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8Async.dst);
oss << ", dstPitch="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8Async.dstPitch);
oss << ", value="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8Async.value);
oss << ", width="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8Async.width);
oss << ", height="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8Async.height);
oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD2D8Async.stream);
oss << ")";
break;
case HIP_API_ID_hipMemsetD32:
oss << "hipMemsetD32(";
oss << "dest="; roctracer::hip_support::detail::operator<<(oss, data->args.hipMemsetD32.dest);
@@ -10634,6 +11127,13 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
else { oss << ", extra="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLaunchKernel.extra__val); }
oss << ")";
break;
case HIP_API_ID_hipModuleLoadFatBinary:
oss << "hipModuleLoadFatBinary(";
if (data->args.hipModuleLoadFatBinary.module == NULL) oss << "module=NULL";
else { oss << "module="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLoadFatBinary.module__val); }
oss << ", fatbin="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleLoadFatBinary.fatbin);
oss << ")";
break;
case HIP_API_ID_hipModuleLoad:
oss << "hipModuleLoad(";
if (data->args.hipModuleLoad.module == NULL) oss << "module=NULL";
@@ -10887,6 +11387,14 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
else { oss << ", pGraph="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamEndCapture.pGraph__val); }
oss << ")";
break;
case HIP_API_ID_hipStreamGetAttribute:
oss << "hipStreamGetAttribute(";
oss << "stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamGetAttribute.stream);
oss << ", attr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamGetAttribute.attr);
if (data->args.hipStreamGetAttribute.value_out == NULL) oss << ", value_out=NULL";
else { oss << ", value_out="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamGetAttribute.value_out__val); }
oss << ")";
break;
case HIP_API_ID_hipStreamGetCaptureInfo:
oss << "hipStreamGetCaptureInfo(";
oss << "stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamGetCaptureInfo.stream);
@@ -10944,6 +11452,14 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
oss << "stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamQuery.stream);
oss << ")";
break;
case HIP_API_ID_hipStreamSetAttribute:
oss << "hipStreamSetAttribute(";
oss << "stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamSetAttribute.stream);
oss << ", attr="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamSetAttribute.attr);
if (data->args.hipStreamSetAttribute.value == NULL) oss << ", value=NULL";
else { oss << ", value="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamSetAttribute.value__val); }
oss << ")";
break;
case HIP_API_ID_hipStreamSynchronize:
oss << "hipStreamSynchronize(";
oss << "stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipStreamSynchronize.stream);
@@ -11195,6 +11711,13 @@ static inline const char* hipApiString(hip_api_id_t id, const hip_api_data_t* da
oss << ", stream="; roctracer::hip_support::detail::operator<<(oss, data->args.hipWaitExternalSemaphoresAsync.stream);
oss << ")";
break;
case HIP_API_ID_hipModuleGetFunctionCount:
oss << "hipModuleGetFunctionCount(";
if (data->args.hipModuleGetFunctionCount.count == NULL) oss << "count=NULL";
else { oss << "count="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleGetFunctionCount.count__val); }
oss << ", mod="; roctracer::hip_support::detail::operator<<(oss, data->args.hipModuleGetFunctionCount.mod);
oss << ")";
break;
default: oss << "unknown";
};
return strdup(oss.str().c_str());
projects/clr/hipamd/include/hip/amd_detail/texture_fetch_functions.h
+15 -3
مشاهده پرونده
@@ -31,9 +31,10 @@ THE SOFTWARE.
#include <type_traits>
#endif // !defined(__HIPCC_RTC__)
#define TEXTURE_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)t.textureObject; \
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
#define TEXTURE_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)t.textureObject; \
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \
(void)s;
template<typename T>
struct __hip_is_tex_surf_scalar_channel_type
@@ -376,6 +377,11 @@ template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemapGrad(texture<T, hipTextureTypeCubemap, readMode> t, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_PARAMETERS_INIT;
(void)x;
(void)y;
(void)z;
(void)dPdx;
(void)dPdy;
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CM(i, s, get_native_vector(float4(x, y, z, 0.0f)),
// get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x,
@@ -387,6 +393,12 @@ template <typename T, hipTextureReadMode readMode>
static __forceinline__ __device__ __hip_img_chk__ __hip_tex_ret_t<T, readMode> texCubemapLayeredGrad(texture<T, hipTextureTypeCubemapLayered, readMode> t, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
{
TEXTURE_PARAMETERS_INIT;
(void)x;
(void)y;
(void)z;
(void)layer;
(void)dPdx;
(void)dPdy;
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CMa(i, s, get_native_vector(float4(x, y, z, layer)),
// get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x,
projects/clr/hipamd/include/hip/amd_detail/texture_indirect_functions.h
+45 -30
مشاهده پرونده
@@ -32,9 +32,10 @@ THE SOFTWARE.
#include <type_traits>
#endif // !defined(__HIPCC_RTC__)
#define TEXTURE_OBJECT_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD;
#define TEXTURE_OBJECT_PARAMETERS_INIT \
unsigned int ADDRESS_SPACE_CONSTANT* i = (unsigned int ADDRESS_SPACE_CONSTANT*)textureObject; \
unsigned int ADDRESS_SPACE_CONSTANT* s = i + HIP_SAMPLER_OBJECT_OFFSET_DWORD; \
(void)s;
template <
typename T,
@@ -287,10 +288,11 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex1DLayeredLod(hipTextureObject_t textureObject, float x, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
float2 coords{x, layer};
auto tmp = __ockl_image_sample_1Da(i, s, get_native_vector(coords));
return __hipMapFrom<T>(tmp);
TEXTURE_OBJECT_PARAMETERS_INIT;
(void)level;
float2 coords{x, layer};
auto tmp = __ockl_image_sample_1Da(i, s, get_native_vector(coords));
return __hipMapFrom<T>(tmp);
}
template <
@@ -306,10 +308,11 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex2DLayeredLod(hipTextureObject_t textureObject, float x, float y, int layer, float level)
{
TEXTURE_OBJECT_PARAMETERS_INIT
float4 coords{x, y, layer, 0.0f};
auto tmp = __ockl_image_sample_2Da(i, s, get_native_vector(coords));
return __hipMapFrom<T>(tmp);
TEXTURE_OBJECT_PARAMETERS_INIT;
(void)level;
float4 coords{x, y, layer, 0.0f};
auto tmp = __ockl_image_sample_2Da(i, s, get_native_vector(coords));
return __hipMapFrom<T>(tmp);
}
template <
@@ -344,12 +347,17 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T texCubemapGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CM(i, s, get_native_vector(float4(x, y, z, 0.0f)),
// get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x,
// dPdy.y, dPdy.z, 0.0f))); return __hipMapFrom<T>(tmp);
return {};
TEXTURE_OBJECT_PARAMETERS_INIT;
(void)x;
(void)y;
(void)z;
(void)dPdx;
(void)dPdy;
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CM(i, s, get_native_vector(float4(x, y, z, 0.0f)),
// get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x,
// dPdy.y, dPdy.z, 0.0f))); return __hipMapFrom<T>(tmp);
return {};
}
template <
@@ -422,13 +430,14 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T tex3DGrad(hipTextureObject_t textureObject, float x, float y, float z, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
float4 coords{x, y, z, 0.0f};
float4 gradx{dPdy.x, dPdy.y, dPdy.z, 0.0f};
float4 grady{dPdy.x, dPdy.y, dPdy.z, 0.0f};
auto tmp = __ockl_image_sample_grad_3D(i, s, get_native_vector(coords),
get_native_vector(gradx), get_native_vector(grady));
return __hipMapFrom<T>(tmp);
TEXTURE_OBJECT_PARAMETERS_INIT;
(void)dPdx;
float4 coords{x, y, z, 0.0f};
float4 gradx{dPdy.x, dPdy.y, dPdy.z, 0.0f};
float4 grady{dPdy.x, dPdy.y, dPdy.z, 0.0f};
auto tmp = __ockl_image_sample_grad_3D(i, s, get_native_vector(coords), get_native_vector(gradx),
get_native_vector(grady));
return __hipMapFrom<T>(tmp);
}
template <
@@ -483,12 +492,18 @@ template <
typename __hip_internal::enable_if<__hip_is_tex_surf_channel_type<T>::value>::type* = nullptr>
static __device__ __hip_img_chk__ T texCubemapLayeredGrad(hipTextureObject_t textureObject, float x, float y, float z, int layer, float4 dPdx, float4 dPdy)
{
TEXTURE_OBJECT_PARAMETERS_INIT
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CMa(i, s, get_native_vector(float4(x, y, z, layer)),
// get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x,
// dPdy.y, dPdy.z, 0.0f))); return __hipMapFrom<T>(tmp);
return {};
TEXTURE_OBJECT_PARAMETERS_INIT;
(void)x;
(void)y;
(void)z;
(void)layer;
(void)dPdx;
(void)dPdy;
// TODO missing in device libs.
// auto tmp = __ockl_image_sample_grad_CMa(i, s, get_native_vector(float4(x, y, z, layer)),
// get_native_vector(float4(dPdx.x, dPdx.y, dPdx.z, 0.0f)), get_native_vector(float4(dPdy.x,
// dPdy.y, dPdy.z, 0.0f))); return __hipMapFrom<T>(tmp);
return {};
}
template <
projects/clr/hipamd/src/CMakeLists.txt
+4 -7
مشاهده پرونده
@@ -178,20 +178,17 @@ if(DISABLE_DIRECT_DISPATCH)
target_compile_definitions(amdhip64 PRIVATE DISABLE_DIRECT_DISPATCH)
endif()
# hipamd will reference llvm symbols, so we need install rocm-llvm-dev package
find_package(LLVM REQUIRED CONFIG
HINTS
${ROCM_PATH}/llvm)
message(STATUS "Found LLVM_INCLUDE_DIRS=" ${LLVM_INCLUDE_DIRS} )
target_include_directories(amdhip64 PRIVATE ${LLVM_INCLUDE_DIRS})
# Short-Term solution for pre-compiled headers for online compilation
# Enable pre compiled header
if(__HIP_ENABLE_PCH)
find_package(LLVM REQUIRED CONFIG
HINTS
${ROCM_PATH}/llvm)
# find_package(LLVM) returns the lib/cmake/llvm location. We require the root.
if(NOT DEFINED HIP_LLVM_ROOT)
set(HIP_LLVM_ROOT "${LLVM_DIR}/../../..")
endif()
execute_process(COMMAND sh -c "${CMAKE_CURRENT_SOURCE_DIR}/hip_embed_pch.sh ${HIP_COMMON_INCLUDE_DIR} ${PROJECT_BINARY_DIR}/include ${PROJECT_SOURCE_DIR}/include ${HIP_LLVM_ROOT}" COMMAND_ECHO STDERR RESULT_VARIABLE EMBED_PCH_RC WORKING_DIRECTORY ${CMAKE_BINARY_DIR})
if (EMBED_PCH_RC AND NOT EMBED_PCH_RC EQUAL 0)
message(FATAL_ERROR "Failed to embed PCH")
projects/clr/hipamd/src/amd_hsa_elf.hpp
+162
مشاهده پرونده
@@ -0,0 +1,162 @@
/*
Copyright (c) 2015 - 2021 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
// This header file is partially copied from
// https://github.com/llvm/llvm-project/blob/main/llvm/include/llvm/BinaryFormat/ELF.h
// AMDGPU OS for HSA compatible compute kernels.
enum { ELFOSABI_AMDGPU_HSA = 64, ELFOSABI_AMDGPU_PAL = 65, ELFOSABI_AMDGPU_MESA3D = 66 };
enum {
ELFABIVERSION_AMDGPU_HSA_V2 = 0,
ELFABIVERSION_AMDGPU_HSA_V3 = 1,
ELFABIVERSION_AMDGPU_HSA_V4 = 2,
ELFABIVERSION_AMDGPU_HSA_V5 = 3,
ELFABIVERSION_AMDGPU_HSA_V6 = 4,
};
// AMDGPU specific e_flags
enum : unsigned {
EF_AMDGPU_MACH = 0x0ff,
// AMDGPU processors
EF_AMDGPU_MACH_NONE = 0x000,
EF_AMDGPU_MACH_R600_R600 = 0x001,
EF_AMDGPU_MACH_R600_R630 = 0x002,
EF_AMDGPU_MACH_R600_RS880 = 0x003,
EF_AMDGPU_MACH_R600_RV670 = 0x004,
EF_AMDGPU_MACH_R600_RV710 = 0x005,
EF_AMDGPU_MACH_R600_RV730 = 0x006,
EF_AMDGPU_MACH_R600_RV770 = 0x007,
EF_AMDGPU_MACH_R600_CEDAR = 0x008,
EF_AMDGPU_MACH_R600_CYPRESS = 0x009,
EF_AMDGPU_MACH_R600_JUNIPER = 0x00a,
EF_AMDGPU_MACH_R600_REDWOOD = 0x00b,
EF_AMDGPU_MACH_R600_SUMO = 0x00c,
EF_AMDGPU_MACH_R600_BARTS = 0x00d,
EF_AMDGPU_MACH_R600_CAICOS = 0x00e,
EF_AMDGPU_MACH_R600_CAYMAN = 0x00f,
EF_AMDGPU_MACH_R600_TURKS = 0x010,
EF_AMDGPU_MACH_R600_RESERVED_FIRST = 0x011,
EF_AMDGPU_MACH_R600_RESERVED_LAST = 0x01f,
EF_AMDGPU_MACH_R600_FIRST = EF_AMDGPU_MACH_R600_R600,
EF_AMDGPU_MACH_R600_LAST = EF_AMDGPU_MACH_R600_TURKS,
// AMDGCN-based processors.
EF_AMDGPU_MACH_AMDGCN_GFX600 = 0x020,
EF_AMDGPU_MACH_AMDGCN_GFX601 = 0x021,
EF_AMDGPU_MACH_AMDGCN_GFX700 = 0x022,
EF_AMDGPU_MACH_AMDGCN_GFX701 = 0x023,
EF_AMDGPU_MACH_AMDGCN_GFX702 = 0x024,
EF_AMDGPU_MACH_AMDGCN_GFX703 = 0x025,
EF_AMDGPU_MACH_AMDGCN_GFX704 = 0x026,
EF_AMDGPU_MACH_AMDGCN_RESERVED_0X27 = 0x027,
EF_AMDGPU_MACH_AMDGCN_GFX801 = 0x028,
EF_AMDGPU_MACH_AMDGCN_GFX802 = 0x029,
EF_AMDGPU_MACH_AMDGCN_GFX803 = 0x02a,
EF_AMDGPU_MACH_AMDGCN_GFX810 = 0x02b,
EF_AMDGPU_MACH_AMDGCN_GFX900 = 0x02c,
EF_AMDGPU_MACH_AMDGCN_GFX902 = 0x02d,
EF_AMDGPU_MACH_AMDGCN_GFX904 = 0x02e,
EF_AMDGPU_MACH_AMDGCN_GFX906 = 0x02f,
EF_AMDGPU_MACH_AMDGCN_GFX908 = 0x030,
EF_AMDGPU_MACH_AMDGCN_GFX909 = 0x031,
EF_AMDGPU_MACH_AMDGCN_GFX90C = 0x032,
EF_AMDGPU_MACH_AMDGCN_GFX1010 = 0x033,
EF_AMDGPU_MACH_AMDGCN_GFX1011 = 0x034,
EF_AMDGPU_MACH_AMDGCN_GFX1012 = 0x035,
EF_AMDGPU_MACH_AMDGCN_GFX1030 = 0x036,
EF_AMDGPU_MACH_AMDGCN_GFX1031 = 0x037,
EF_AMDGPU_MACH_AMDGCN_GFX1032 = 0x038,
EF_AMDGPU_MACH_AMDGCN_GFX1033 = 0x039,
EF_AMDGPU_MACH_AMDGCN_GFX602 = 0x03a,
EF_AMDGPU_MACH_AMDGCN_GFX705 = 0x03b,
EF_AMDGPU_MACH_AMDGCN_GFX805 = 0x03c,
EF_AMDGPU_MACH_AMDGCN_GFX1035 = 0x03d,
EF_AMDGPU_MACH_AMDGCN_GFX1034 = 0x03e,
EF_AMDGPU_MACH_AMDGCN_GFX90A = 0x03f,
EF_AMDGPU_MACH_AMDGCN_GFX940 = 0x040,
EF_AMDGPU_MACH_AMDGCN_GFX1100 = 0x041,
EF_AMDGPU_MACH_AMDGCN_GFX1013 = 0x042,
EF_AMDGPU_MACH_AMDGCN_GFX1150 = 0x043,
EF_AMDGPU_MACH_AMDGCN_GFX1103 = 0x044,
EF_AMDGPU_MACH_AMDGCN_GFX1036 = 0x045,
EF_AMDGPU_MACH_AMDGCN_GFX1101 = 0x046,
EF_AMDGPU_MACH_AMDGCN_GFX1102 = 0x047,
EF_AMDGPU_MACH_AMDGCN_GFX1200 = 0x048,
EF_AMDGPU_MACH_AMDGCN_RESERVED_0X49 = 0x049,
EF_AMDGPU_MACH_AMDGCN_GFX1151 = 0x04a,
EF_AMDGPU_MACH_AMDGCN_GFX941 = 0x04b,
EF_AMDGPU_MACH_AMDGCN_GFX942 = 0x04c,
EF_AMDGPU_MACH_AMDGCN_RESERVED_0X4D = 0x04d,
EF_AMDGPU_MACH_AMDGCN_GFX1201 = 0x04e,
EF_AMDGPU_MACH_AMDGCN_GFX950 = 0x04f,
EF_AMDGPU_MACH_AMDGCN_RESERVED_0X50 = 0x050,
EF_AMDGPU_MACH_AMDGCN_GFX9_GENERIC = 0x051,
EF_AMDGPU_MACH_AMDGCN_GFX10_1_GENERIC = 0x052,
EF_AMDGPU_MACH_AMDGCN_GFX10_3_GENERIC = 0x053,
EF_AMDGPU_MACH_AMDGCN_GFX11_GENERIC = 0x054,
EF_AMDGPU_MACH_AMDGCN_GFX1152 = 0x055,
EF_AMDGPU_MACH_AMDGCN_RESERVED_0X56 = 0x056,
EF_AMDGPU_MACH_AMDGCN_RESERVED_0X57 = 0x057,
EF_AMDGPU_MACH_AMDGCN_RESERVED_0X58 = 0x058,
EF_AMDGPU_MACH_AMDGCN_GFX12_GENERIC = 0x059,
EF_AMDGPU_MACH_AMDGCN_GFX9_4_GENERIC = 0x05f,
EF_AMDGPU_MACH_AMDGCN_FIRST = EF_AMDGPU_MACH_AMDGCN_GFX600,
EF_AMDGPU_MACH_AMDGCN_LAST = EF_AMDGPU_MACH_AMDGCN_GFX9_4_GENERIC,
// Indicates if the "xnack" target feature is enabled for all code contained
// in the object.
//
// Only valid for ELFOSABI_AMDGPU_HSA and ELFABIVERSION_AMDGPU_HSA_V3.
EF_AMDGPU_FEATURE_XNACK_V3 = 0x100,
// Indicates if the "sramecc" target feature is enabled for all code
// contained in the object.
//
// Only valid for ELFOSABI_AMDGPU_HSA and ELFABIVERSION_AMDGPU_HSA_V3.
EF_AMDGPU_FEATURE_SRAMECC_V3 = 0x200,
// Only valid for ELFOSABI_AMDGPU_HSA and ELFABIVERSION_AMDGPU_HSA_V4,
// ELFABIVERSION_AMDGPU_HSA_V5 and ELFABIVERSION_AMDGPU_HSA_V6.
EF_AMDGPU_FEATURE_XNACK_V4 = 0x300,
EF_AMDGPU_FEATURE_XNACK_UNSUPPORTED_V4 = 0x000,
EF_AMDGPU_FEATURE_XNACK_ANY_V4 = 0x100,
EF_AMDGPU_FEATURE_XNACK_OFF_V4 = 0x200,
EF_AMDGPU_FEATURE_XNACK_ON_V4 = 0x300,
// SRAMECC selection mask for EF_AMDGPU_FEATURE_SRAMECC_* values.
// Only valid for ELFOSABI_AMDGPU_HSA and ELFABIVERSION_AMDGPU_HSA_V4,
// ELFABIVERSION_AMDGPU_HSA_V5 and ELFABIVERSION_AMDGPU_HSA_V6.
EF_AMDGPU_FEATURE_SRAMECC_V4 = 0xc00,
EF_AMDGPU_FEATURE_SRAMECC_UNSUPPORTED_V4 = 0x000,
EF_AMDGPU_FEATURE_SRAMECC_ANY_V4 = 0x400,
EF_AMDGPU_FEATURE_SRAMECC_OFF_V4 = 0x800,
EF_AMDGPU_FEATURE_SRAMECC_ON_V4 = 0xc00,
// Generic target versioning. This is contained in the list byte of EFLAGS.
EF_AMDGPU_GENERIC_VERSION = 0xff000000,
EF_AMDGPU_GENERIC_VERSION_OFFSET = 24,
EF_AMDGPU_GENERIC_VERSION_MIN = 1,
EF_AMDGPU_GENERIC_VERSION_MAX = 0xff,
};
projects/clr/hipamd/src/amdhip.def
+18
مشاهده پرونده
@@ -492,3 +492,21 @@ hipLinkCreate
hipLinkDestroy
hipLaunchKernelExC
hipDrvLaunchKernelEx
hipModuleGetFunctionCount
hipMemsetD2D8
hipMemsetD2D8Async
hipMemsetD2D16
hipMemsetD2D16Async
hipMemsetD2D32
hipMemsetD2D32Async
hipStreamGetAttribute
hipStreamSetAttribute
hipModuleLoadFatBinary
hipMemcpyBatchAsync
hipMemcpy3DBatchAsync
hipMemcpy3DPeer
hipMemcpy3DPeerAsync
hipGetDriverEntryPoint
hipGetDriverEntryPoint_spt
hipMemPrefetchAsync_v2
hipMemAdvise_v2
projects/clr/hipamd/src/hip_api_trace.cpp
+73 -4
مشاهده پرونده
@@ -186,6 +186,10 @@ hipError_t hipGetDeviceCount(int* count);
hipError_t hipGetDeviceFlags(unsigned int* flags);
hipError_t hipGetDevicePropertiesR0600(hipDeviceProp_tR0600* prop, int deviceId);
hipError_t hipGetDevicePropertiesR0000(hipDeviceProp_tR0000* prop, int device);
hipError_t hipGetDriverEntryPoint(const char* symbol, void** funcPtr, unsigned long long flags,
hipDriverEntryPointQueryResult* status);
hipError_t hipGetDriverEntryPoint_spt(const char* symbol, void** funcPtr, unsigned long long flags,
hipDriverEntryPointQueryResult* status);
const char* hipGetErrorName(hipError_t hip_error);
const char* hipGetErrorString(hipError_t hipError);
hipError_t hipGetLastError(void);
@@ -396,6 +400,8 @@ hipError_t hipMemAddressFree(void* devPtr, size_t size);
hipError_t hipMemAddressReserve(void** ptr, size_t size, size_t alignment, void* addr,
unsigned long long flags);
hipError_t hipMemAdvise(const void* dev_ptr, size_t count, hipMemoryAdvise advice, int device);
hipError_t hipMemAdvise_v2(const void* dev_ptr, size_t count, hipMemoryAdvise advice,
hipMemLocation location);
hipError_t hipMemAllocHost(void** ptr, size_t size);
hipError_t hipMemAllocPitch(hipDeviceptr_t* dptr, size_t* pitch, size_t widthInBytes, size_t height,
unsigned int elementSizeBytes);
@@ -437,6 +443,8 @@ hipError_t hipMemPoolSetAccess(hipMemPool_t mem_pool, const hipMemAccessDesc* de
hipError_t hipMemPoolSetAttribute(hipMemPool_t mem_pool, hipMemPoolAttr attr, void* value);
hipError_t hipMemPoolTrimTo(hipMemPool_t mem_pool, size_t min_bytes_to_hold);
hipError_t hipMemPrefetchAsync(const void* dev_ptr, size_t count, int device, hipStream_t stream);
hipError_t hipMemPrefetchAsync_v2(const void* dev_ptr, size_t count, hipMemLocation location,
unsigned int flags, hipStream_t stream);
hipError_t hipMemPtrGetInfo(void* ptr, size_t* size);
hipError_t hipMemRangeGetAttribute(void* data, size_t data_size, hipMemRangeAttribute attribute,
const void* dev_ptr, size_t count);
@@ -519,6 +527,7 @@ hipError_t hipMipmappedArrayDestroy(hipMipmappedArray_t hMipmappedArray);
hipError_t hipMipmappedArrayGetLevel(hipArray_t* pLevelArray, hipMipmappedArray_t hMipMappedArray,
unsigned int level);
hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, const char* kname);
hipError_t hipModuleGetFunctionCount(unsigned int* count, hipModule_t mod);
hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod,
const char* name);
hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name);
@@ -534,6 +543,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigne
unsigned int blockDimY, unsigned int blockDimZ,
unsigned int sharedMemBytes, hipStream_t stream,
void** kernelParams, void** extra);
hipError_t hipModuleLoadFatBinary(hipModule_t* module, const void* fatbin);
hipError_t hipModuleLoad(hipModule_t* module, const char* fname);
hipError_t hipModuleLoadData(hipModule_t* module, const void* image);
hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned int numOptions,
@@ -827,10 +837,32 @@ hipError_t hipEventRecordWithFlags(hipEvent_t event, hipStream_t stream, unsigne
hipError_t hipLaunchKernelExC(const hipLaunchConfig_t* config, const void* fPtr, void** args);
hipError_t hipDrvLaunchKernelEx(const HIP_LAUNCH_CONFIG* config, hipFunction_t f, void** params,
void** extra);
hipError_t hipMemGetHandleForAddressRange(void* handle, hipDeviceptr_t dptr, size_t size,
hipError_t hipMemGetHandleForAddressRange(void* handle, hipDeviceptr_t dptr, size_t size,
hipMemRangeHandleType handleType,
unsigned long long flags);
hipError_t hipMemsetD2D8(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, size_t width,
size_t height);
hipError_t hipMemsetD2D8Async(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, size_t width,
size_t height, hipStream_t stream);
hipError_t hipMemsetD2D16(hipDeviceptr_t dst, size_t dstPitch, unsigned short value, size_t width,
size_t height);
hipError_t hipMemsetD2D16Async(hipDeviceptr_t dst, size_t dstPitch, unsigned short value,
size_t width, size_t height, hipStream_t stream);
hipError_t hipMemsetD2D32(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, size_t width,
size_t height);
hipError_t hipMemsetD2D32Async(hipDeviceptr_t dst, size_t dstPitch, unsigned int value,
size_t width, size_t height, hipStream_t stream);
hipError_t hipStreamGetAttribute(hipStream_t stream, hipStreamAttrID attr,
hipStreamAttrValue *value);
hipError_t hipStreamSetAttribute(hipStream_t stream, hipStreamAttrID attr,
const hipStreamAttrValue *value);
hipError_t hipMemcpyBatchAsync(void **dsts, void **srcs, size_t *sizes, size_t count,
hipMemcpyAttributes *attrs, size_t *attrsIdxs, size_t numAttrs,
size_t *failIdx, hipStream_t stream);
hipError_t hipMemcpy3DBatchAsync(size_t numOps, struct hipMemcpy3DBatchOp *opList, size_t *failIdx,
unsigned long long flags, hipStream_t stream);
hipError_t hipMemcpy3DPeer(hipMemcpy3DPeerParms *p);
hipError_t hipMemcpy3DPeerAsync(hipMemcpy3DPeerParms *p, hipStream_t stream);
} // namespace hip
namespace hip {
@@ -1096,6 +1128,7 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) {
ptrDispatchTable->hipMemAddressFree_fn = hip::hipMemAddressFree;
ptrDispatchTable->hipMemAddressReserve_fn = hip::hipMemAddressReserve;
ptrDispatchTable->hipMemAdvise_fn = hip::hipMemAdvise;
ptrDispatchTable->hipMemAdvise_v2_fn = hip::hipMemAdvise_v2;
ptrDispatchTable->hipMemAllocHost_fn = hip::hipMemAllocHost;
ptrDispatchTable->hipMemAllocPitch_fn = hip::hipMemAllocPitch;
ptrDispatchTable->hipMemCreate_fn = hip::hipMemCreate;
@@ -1122,6 +1155,7 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) {
ptrDispatchTable->hipMemPoolSetAttribute_fn = hip::hipMemPoolSetAttribute;
ptrDispatchTable->hipMemPoolTrimTo_fn = hip::hipMemPoolTrimTo;
ptrDispatchTable->hipMemPrefetchAsync_fn = hip::hipMemPrefetchAsync;
ptrDispatchTable->hipMemPrefetchAsync_v2_fn = hip::hipMemPrefetchAsync_v2;
ptrDispatchTable->hipMemPtrGetInfo_fn = hip::hipMemPtrGetInfo;
ptrDispatchTable->hipMemRangeGetAttribute_fn = hip::hipMemRangeGetAttribute;
ptrDispatchTable->hipMemRangeGetAttributes_fn = hip::hipMemRangeGetAttributes;
@@ -1174,12 +1208,14 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) {
ptrDispatchTable->hipMipmappedArrayDestroy_fn = hip::hipMipmappedArrayDestroy;
ptrDispatchTable->hipMipmappedArrayGetLevel_fn = hip::hipMipmappedArrayGetLevel;
ptrDispatchTable->hipModuleGetFunction_fn = hip::hipModuleGetFunction;
ptrDispatchTable->hipModuleGetFunctionCount_fn = hip::hipModuleGetFunctionCount;
ptrDispatchTable->hipModuleGetGlobal_fn = hip::hipModuleGetGlobal;
ptrDispatchTable->hipModuleGetTexRef_fn = hip::hipModuleGetTexRef;
ptrDispatchTable->hipModuleLaunchCooperativeKernel_fn = hip::hipModuleLaunchCooperativeKernel;
ptrDispatchTable->hipModuleLaunchCooperativeKernelMultiDevice_fn =
hip::hipModuleLaunchCooperativeKernelMultiDevice;
ptrDispatchTable->hipModuleLaunchKernel_fn = hip::hipModuleLaunchKernel;
ptrDispatchTable->hipModuleLoadFatBinary_fn = hip::hipModuleLoadFatBinary;
ptrDispatchTable->hipModuleLoad_fn = hip::hipModuleLoad;
ptrDispatchTable->hipModuleLoadData_fn = hip::hipModuleLoadData;
ptrDispatchTable->hipModuleLoadDataEx_fn = hip::hipModuleLoadDataEx;
@@ -1314,6 +1350,8 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) {
ptrDispatchTable->hipGetStreamDeviceId_fn = hip::hipGetStreamDeviceId;
ptrDispatchTable->hipDrvGraphAddMemsetNode_fn = hip::hipDrvGraphAddMemsetNode;
ptrDispatchTable->hipGetDevicePropertiesR0000_fn = hip::hipGetDevicePropertiesR0000;
ptrDispatchTable->hipGetDriverEntryPoint_fn = hip::hipGetDriverEntryPoint;
ptrDispatchTable->hipGetDriverEntryPoint_spt_fn = hip::hipGetDriverEntryPoint_spt;
ptrDispatchTable->hipExtGetLastError_fn = hip::hipExtGetLastError;
ptrDispatchTable->hipTexRefGetBorderColor_fn = hip::hipTexRefGetBorderColor;
ptrDispatchTable->hipTexRefGetArray_fn = hip::hipTexRefGetArray;
@@ -1346,6 +1384,18 @@ void UpdateDispatchTable(HipDispatchTable* ptrDispatchTable) {
ptrDispatchTable->hipLaunchKernelExC_fn = hip::hipLaunchKernelExC;
ptrDispatchTable->hipDrvLaunchKernelEx_fn = hip::hipDrvLaunchKernelEx;
ptrDispatchTable->hipMemGetHandleForAddressRange_fn = hip::hipMemGetHandleForAddressRange;
ptrDispatchTable->hipMemsetD2D8_fn = hip::hipMemsetD2D8;
ptrDispatchTable->hipMemsetD2D8Async_fn = hip::hipMemsetD2D8Async;
ptrDispatchTable->hipMemsetD2D16_fn = hip::hipMemsetD2D16;
ptrDispatchTable->hipMemsetD2D16Async_fn = hip::hipMemsetD2D16Async;
ptrDispatchTable->hipMemsetD2D32_fn = hip::hipMemsetD2D32;
ptrDispatchTable->hipMemsetD2D32Async_fn = hip::hipMemsetD2D32Async;
ptrDispatchTable->hipStreamGetAttribute_fn = hip::hipStreamGetAttribute;
ptrDispatchTable->hipStreamSetAttribute_fn = hip::hipStreamSetAttribute;
ptrDispatchTable->hipMemcpyBatchAsync_fn = hip::hipMemcpyBatchAsync;
ptrDispatchTable->hipMemcpy3DBatchAsync_fn = hip::hipMemcpy3DBatchAsync;
ptrDispatchTable->hipMemcpy3DPeer_fn = hip::hipMemcpy3DPeer;
ptrDispatchTable->hipMemcpy3DPeerAsync_fn = hip::hipMemcpy3DPeerAsync;
}
#if HIP_ROCPROFILER_REGISTER > 0
@@ -1989,15 +2039,34 @@ HIP_ENFORCE_ABI(HipDispatchTable, hipLaunchKernelExC_fn, 474);
HIP_ENFORCE_ABI(HipDispatchTable, hipDrvLaunchKernelEx_fn, 475);
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 12
HIP_ENFORCE_ABI(HipDispatchTable, hipMemGetHandleForAddressRange_fn, 476);
// HIP_RUNTIME_API_TABLE_STEP_VERSION == 13
HIP_ENFORCE_ABI(HipDispatchTable, hipModuleGetFunctionCount_fn, 477);
HIP_ENFORCE_ABI(HipDispatchTable, hipMemsetD2D8_fn, 478);
HIP_ENFORCE_ABI(HipDispatchTable, hipMemsetD2D8Async_fn, 479);
HIP_ENFORCE_ABI(HipDispatchTable, hipMemsetD2D16_fn, 480);
HIP_ENFORCE_ABI(HipDispatchTable, hipMemsetD2D16Async_fn, 481);
HIP_ENFORCE_ABI(HipDispatchTable, hipMemsetD2D32_fn, 482);
HIP_ENFORCE_ABI(HipDispatchTable, hipMemsetD2D32Async_fn, 483);
HIP_ENFORCE_ABI(HipDispatchTable, hipStreamGetAttribute_fn, 484);
HIP_ENFORCE_ABI(HipDispatchTable, hipStreamSetAttribute_fn, 485);
HIP_ENFORCE_ABI(HipDispatchTable, hipModuleLoadFatBinary_fn, 486);
HIP_ENFORCE_ABI(HipDispatchTable, hipMemcpyBatchAsync_fn, 487);
HIP_ENFORCE_ABI(HipDispatchTable, hipMemcpy3DBatchAsync_fn, 488);
HIP_ENFORCE_ABI(HipDispatchTable, hipMemcpy3DPeer_fn, 489);
HIP_ENFORCE_ABI(HipDispatchTable, hipMemcpy3DPeerAsync_fn, 490);
HIP_ENFORCE_ABI(HipDispatchTable, hipGetDriverEntryPoint_fn, 491);
HIP_ENFORCE_ABI(HipDispatchTable, hipGetDriverEntryPoint_spt_fn, 492);
HIP_ENFORCE_ABI(HipDispatchTable, hipMemPrefetchAsync_v2_fn, 493);
HIP_ENFORCE_ABI(HipDispatchTable, hipMemAdvise_v2_fn, 494);
// if HIP_ENFORCE_ABI entries are added for each new function pointer in the table, the number below
// will be +1 of the number in the last HIP_ENFORCE_ABI line. E.g.:
//
// HIP_ENFORCE_ABI(<table>, <functor>, 8)
//
// HIP_ENFORCE_ABI_VERSIONING(<table>, 9) <- 8 + 1 = 9
HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 477)
HIP_ENFORCE_ABI_VERSIONING(HipDispatchTable, 495)
static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 13,
static_assert(HIP_RUNTIME_API_TABLE_MAJOR_VERSION == 0 && HIP_RUNTIME_API_TABLE_STEP_VERSION == 14,
"If you get this error, add new HIP_ENFORCE_ABI(...) code for the new function "
"pointers and then update this check so it is true");
#endif
projects/clr/hipamd/src/hip_code_object.cpp
+20 -9
مشاهده پرونده
@@ -19,17 +19,19 @@ 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.
*/
#define LLVM_DISABLE_ABI_BREAKING_CHECKS_ENFORCING 1
#include "llvm/BinaryFormat/ELF.h"
#if defined(_WIN32)
#if defined(__has_attribute)
// MS compiler doesn't support __has_attribute
#undef __has_attribute
#endif
#endif
#include "hip_code_object.hpp"
#include "amd_hsa_elf.hpp"
#include <cstring>
#include <hip/driver_types.h>
#include "hip/hip_runtime_api.h"
#include "hip/hip_runtime.h"
#include "hip_internal.hpp"
#include "platform/program.hpp"
#include <elf/elf.hpp>
#include "comgrctx.hpp"
#include "hip_comgr_helper.hpp"
using namespace llvm::ELF;
namespace hip {
hipError_t ihipFree(void* ptr);
@@ -122,6 +124,15 @@ hipError_t DynCO::getDynFunc(hipFunction_t* hfunc, std::string func_name) {
return it->second->getDynFunc(hfunc, module_);
}
hipError_t DynCO::getFuncCount(unsigned int* count) {
amd::ScopedLock lock(dclock_);
if (count == nullptr) {
return hipErrorInvalidValue;
}
*count = functions_.size();
return hipSuccess;
}
bool DynCO::isValidDynFunc(const void* hfunc) {
amd::ScopedLock lock(dclock_);
return std::any_of(functions_.begin(), functions_.end(),
projects/clr/hipamd/src/hip_code_object.hpp
+1
مشاهده پرونده
@@ -111,6 +111,7 @@ public:
//Gets GlobalVar/Functions from a dynamically loaded code object
hipError_t getDynFunc(hipFunction_t* hfunc, std::string func_name);
hipError_t getFuncCount(unsigned int* count);
bool isValidDynFunc(const void* hfunc);
hipError_t getDeviceVar(DeviceVar** dvar, std::string var_name);
projects/clr/hipamd/src/hip_comgr_helper.cpp
+2 -9
مشاهده پرونده
@@ -19,18 +19,11 @@ 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.
*/
#define LLVM_DISABLE_ABI_BREAKING_CHECKS_ENFORCING 1
#include "llvm/BinaryFormat/ELF.h"
#include "hip_comgr_helper.hpp"
#if defined(_WIN32)
#include <io.h>
#if defined(__has_attribute)
// MS compiler doesn't support __has_attribute
#undef __has_attribute
#endif
#endif
#include "hip_comgr_helper.hpp"
using namespace llvm::ELF;
#include "../src/amd_hsa_elf.hpp"
namespace hip {
std::unordered_set<LinkProgram*> LinkProgram::linker_set_;
projects/clr/hipamd/src/hip_context.cpp
+5 -1
مشاهده پرونده
@@ -45,7 +45,11 @@ void init(bool* status) {
#if DISABLE_DIRECT_DISPATCH
constexpr bool kDirectDispatch = false;
#else
constexpr bool kDirectDispatch = IS_LINUX;
#ifndef WITHOUT_HSA_BACKEND
constexpr bool kDirectDispatch = true;
#else
constexpr bool kDirectDispatch = false;
#endif
#endif
AMD_DIRECT_DISPATCH = flagIsDefault(AMD_DIRECT_DISPATCH) ? kDirectDispatch : AMD_DIRECT_DISPATCH;
if (!amd::Runtime::init()) {
projects/clr/hipamd/src/hip_conversions.hpp
+74 -1
مشاهده پرونده
@@ -978,4 +978,77 @@ inline
size_t getElementSize(const hipChannelFormatDesc &desc) {
return (desc.x / 8) * getNumChannels(desc);
}
};
inline
hipMemcpy3DParms getMemcpy3DParms(const hipMemcpy3DBatchOp& desc) {
hipMemcpy3DParms params;
params.extent = desc.extent;
params.kind = hipMemcpyDefault;
// infer elementSize
size_t elementSize = 1;
if (desc.src.type == hipMemcpyOperandTypeArray) {
elementSize = getElementSize(desc.src.op.array.array);
} else if (desc.dst.type == hipMemcpyOperandTypeArray) {
elementSize = getElementSize(desc.dst.op.array.array);
}
// source
if (desc.src.type == hipMemcpyOperandTypePointer) {
size_t row = desc.src.op.ptr.rowLength;
size_t height = desc.src.op.ptr.layerHeight;
size_t spitch = (row ? row : desc.extent.width) * elementSize;
size_t swidth = (row ? row : desc.extent.width);
size_t sheight = (height ? height : desc.extent.height);
params.srcPtr = make_hipPitchedPtr(
desc.src.op.ptr.ptr,
spitch,
swidth,
sheight
);
} else if (desc.src.type == hipMemcpyOperandTypeArray) {
params.srcArray = desc.src.op.array.array;
params.srcPos = make_hipPos(
desc.src.op.array.offset.x,
desc.src.op.array.offset.y,
desc.src.op.array.offset.z
);
}
// dest
if (desc.dst.type == hipMemcpyOperandTypePointer) {
size_t row = desc.dst.op.ptr.rowLength;
size_t height = desc.dst.op.ptr.layerHeight;
size_t spitch = (row ? row : desc.extent.width) * elementSize;
size_t swidth = (row ? row : desc.extent.width);
size_t sheight = (height ? height : desc.extent.height);
params.dstPtr = make_hipPitchedPtr(
desc.dst.op.ptr.ptr,
spitch,
swidth,
sheight
);
} else if (desc.dst.type == hipMemcpyOperandTypeArray) {
params.dstArray = desc.dst.op.array.array;
params.dstPos = make_hipPos(
desc.dst.op.array.offset.x,
desc.dst.op.array.offset.y,
desc.dst.op.array.offset.z
);
}
return params;
}
inline
hipMemcpy3DParms getMemcpy3DParms(const hipMemcpy3DPeerParms& desc) {
hipMemcpy3DParms params;
params.srcArray = desc.srcArray;
params.srcPos = desc.srcPos;
params.srcPtr = desc.srcPtr;
params.dstArray = desc.dstArray;
params.dstPos = desc.dstPos;
params.dstPtr = desc.dstPtr;
params.extent = desc.extent;
params.kind = hipMemcpyDeviceToDevice;
return params;
}
};
projects/clr/hipamd/src/hip_device_runtime.cpp
+55
مشاهده پرونده
@@ -21,6 +21,7 @@
#include <hip/hip_runtime.h>
#include "hip_internal.hpp"
#include "hip_platform.hpp"
#undef hipChooseDevice
#undef hipDeviceProp_t
@@ -716,6 +717,60 @@ hipError_t hipGetDeviceFlags(unsigned int* flags) {
HIP_RETURN(hipSuccess);
}
hipError_t hipGetDriverEntryPoint_common(const char* symbol, void** funcPtr, unsigned long long flags,
hipDriverEntryPointQueryResult* status) {
std::string symbolString = symbol;
if (symbol == nullptr || symbolString == "" || funcPtr == nullptr) {
return hipErrorInvalidValue;
}
if (flags != hipEnableDefault && flags != hipEnableLegacyStream
&& flags != hipEnablePerThreadDefaultStream) {
return hipErrorInvalidValue;
}
void* handle = hip::PlatformState::instance().getDynamicLibraryHandle();
if (handle == nullptr) {
return hipErrorInvalidValue;
}
if (flags == hipEnablePerThreadDefaultStream) {
symbolString += "_spt";
}
*funcPtr = amd::Os::getSymbol(handle, symbolString.c_str());
if (funcPtr == nullptr) {
if (flags == hipEnablePerThreadDefaultStream) {
*funcPtr = amd::Os::getSymbol(handle, symbol);
}
if (funcPtr == nullptr) {
if (status != nullptr) {
*status = hipDriverEntryPointSymbolNotFound;
}
return hipErrorInvalidValue;
}
}
if (status != nullptr) {
*status = hipDriverEntryPointSuccess;
}
return hipSuccess;
}
hipError_t hipGetDriverEntryPoint(const char* symbol, void** funcPtr, unsigned long long flags,
hipDriverEntryPointQueryResult* status) {
HIP_INIT_API(hipGetDriverEntryPoint, symbol, funcPtr, flags, status);
HIP_RETURN(hipGetDriverEntryPoint_common(symbol, funcPtr, flags, status));
}
hipError_t hipGetDriverEntryPoint_spt(const char* symbol, void** funcPtr, unsigned long long flags,
hipDriverEntryPointQueryResult* status) {
HIP_INIT_API(hipGetDriverEntryPoint, symbol, funcPtr, flags, status);
flags = (flags == hipEnableDefault) ? hipEnablePerThreadDefaultStream : flags;
HIP_RETURN(hipGetDriverEntryPoint_common(symbol, funcPtr, flags, status));
}
hipError_t hipSetDevice(int device) {
HIP_INIT_API_NO_RETURN(hipSetDevice, device);
projects/clr/hipamd/src/hip_event.cpp
+14 -9
مشاهده پرونده
@@ -19,8 +19,9 @@
THE SOFTWARE. */
#include <hip/hip_runtime.h>
#include "hip_event.hpp"
#include "hip_graph_internal.hpp"
#if !defined(_MSC_VER)
#include <unistd.h>
#endif
@@ -76,7 +77,9 @@ hipError_t Event::synchronize() {
auto hip_device = g_devices[deviceId()];
// Check HW status of the ROCcrl event. Note: not all ROCclr modes support HW status
static constexpr bool kWaitCompletion = true;
if (!hip_device->devices()[0]->IsHwEventReady(*event_, kWaitCompletion, flags_)) {
amd::SyncPolicy policy = (flags_ == hipEventBlockingSync) ? amd::SyncPolicy::Blocking :
amd::SyncPolicy::Auto;
if (!hip_device->devices()[0]->IsHwEventReady(*event_, kWaitCompletion, policy)) {
event_->awaitCompletion();
}
return hipSuccess;
@@ -88,7 +91,9 @@ bool Event::awaitEventCompletion() {
}
bool EventDD::awaitEventCompletion() {
return g_devices[deviceId()]->devices()[0]->IsHwEventReady(*event_, true, flags_);
amd::SyncPolicy policy = (flags_ == hipEventBlockingSync) ? amd::SyncPolicy::Blocking :
amd::SyncPolicy::Auto;
return g_devices[deviceId()]->devices()[0]->IsHwEventReady(*event_, true, policy);
}
hipError_t Event::elapsedTime(Event& eStop, float& ms) {
@@ -402,16 +407,16 @@ hipError_t hipEventRecord_common(hipEvent_t event, hipStream_t stream, unsigned
std::vector<hip::GraphNode*> lastCapturedNodes = s->GetLastCapturedNodes();
e->SetNodesPrevToRecorded(lastCapturedNodes);
if (flags == hipEventRecordExternal) {
hip::GraphNode* pGraphNode;
status =
hipGraphAddEventRecordNode((hipGraphNode_t*)&pGraphNode, (hipGraph_t)s->GetCaptureGraph(),
(hipGraphNode_t*)s->GetLastCapturedNodes().data(),
s->GetLastCapturedNodes().size(), (hipEvent_t)e);
hip::GraphNode* node = new hip::GraphEventRecordNode(reinterpret_cast<hipEvent_t>(e));
hipError_t status = hip::ihipGraphAddNode(
node, reinterpret_cast<hip::Graph*>(s->GetCaptureGraph()),
reinterpret_cast<hip::GraphNode* const*>(s->GetLastCapturedNodes().data()),
s->GetLastCapturedNodes().size(), false);
if (status != hipSuccess) {
ClPrint(amd::LOG_ERROR, amd::LOG_API, "hipEventRecord add external event node failed");
return status;
}
s->SetLastCapturedNode(pGraphNode);
s->SetLastCapturedNode(node);
}
} else {
if (e->deviceId() != hip_stream->DeviceId()) {
projects/clr/hipamd/src/hip_event.hpp
+13 -1
مشاهده پرونده
@@ -24,6 +24,10 @@
#include "hip_internal.hpp"
#include "thread/monitor.hpp"
#if !defined(_MSC_VER)
#include <sys/mman.h>
#endif
// Internal structure for stream callback handler
namespace hip {
class StreamCallback {
@@ -98,7 +102,9 @@ class Event {
std::vector<hip::GraphNode*> nodesPrevToRecorded_;
protected:
bool CheckHwEvent() {
return g_devices[deviceId()]->devices()[0]->IsHwEventReady(*event_, false, flags_);
amd::SyncPolicy policy = (flags_ == hipEventBlockingSync) ? amd::SyncPolicy::Blocking :
amd::SyncPolicy::Auto;
return g_devices[deviceId()]->devices()[0]->IsHwEventReady(*event_, false, policy);
}
public:
@@ -204,6 +210,12 @@ class IPCEvent : public Event {
amd::Os::shm_unlink(ipc_evt_.ipc_name_);
}
}
#if !defined(_MSC_VER)
// Clean up the POSIX shared memory object
if (!ipc_evt_.ipc_name_.empty() && ipc_evt_.ipc_name_ != "dummy") {
shm_unlink(ipc_evt_.ipc_name_.c_str());
}
#endif
}
IPCEvent() : Event(hipEventInterprocess) {}
bool createIpcEventShmemIfNeeded();
projects/clr/hipamd/src/hip_event_ipc.cpp
+22 -17
مشاهده پرونده
@@ -38,15 +38,16 @@ bool IPCEvent::createIpcEventShmemIfNeeded() {
return true;
}
char name_template[] = "/tmp/eventXXXXXX";
#if !defined(_MSC_VER)
int temp_fd = mkstemp(name_template);
static std::atomic<int> counter{0};
ipc_evt_.ipc_name_ = "/hip_" + std::to_string(getpid()) + "_" + std::to_string(counter++);
#else
char name_template[] = "/hip_XXXXXX";
_mktemp_s(name_template, sizeof(name_template));
#endif
ipc_evt_.ipc_name_ = name_template;
ipc_evt_.ipc_name_.replace(0, 5, "/hip_");
#endif
if (!amd::Os::MemoryMapFileTruncated(
ipc_evt_.ipc_name_.c_str(),
const_cast<const void**>(reinterpret_cast<void**>(&(ipc_evt_.ipc_shmem_))),
@@ -54,10 +55,6 @@ bool IPCEvent::createIpcEventShmemIfNeeded() {
return false;
}
#if !defined(_MSC_VER)
close(temp_fd);
#endif
ipc_evt_.ipc_shmem_->owners = 1;
ipc_evt_.ipc_shmem_->read_index = -1;
ipc_evt_.ipc_shmem_->write_index = 0;
@@ -134,21 +131,28 @@ hipError_t IPCEvent::enqueueRecordCommand(hip::Stream* stream, amd::Command* com
ipc_evt_.ipc_shmem_->owners_device_id = deviceId();
command->enqueue();
// device writes 0 to signal after the hipEventRecord command is completed
// the signal value is checked by WaitThenDecrementSignal cb
hipError_t status = ihipStreamOperation(
reinterpret_cast<hipStream_t>(stream), ROCCLR_COMMAND_STREAM_WRITE_VALUE,
&(ipc_evt_.ipc_shmem_->signal[offset]), 0, 0, 0, sizeof(uint32_t));
if (status != hipSuccess) {
return status;
}
// Set event_ in order to release marked command when event is destroyed
if (event_ != nullptr) {
event_->release();
}
event_ = &command->event();
// device writes 0 to signal after the hipEventRecord command is completed
// the signal value is checked by WaitThenDecrementSignal cb
hipError_t status =
ihipStreamOperation(reinterpret_cast<hipStream_t>(stream), ROCCLR_COMMAND_STREAM_WRITE_VALUE,
&(ipc_evt_.ipc_shmem_->signal[offset]), 0, 0, 0, sizeof(uint32_t));
if (status != hipSuccess) {
return status;
}
// Update read index to indicate new signal.
int expected = write_index - 1;
while (!ipc_evt_.ipc_shmem_->read_index.compare_exchange_weak(expected, write_index)) {
amd::Os::sleep(1);
}
return hipSuccess;
}
@@ -213,6 +217,7 @@ hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle)
hip::Event* e = reinterpret_cast<hip::Event*>(*event);
ihipIpcEventHandle_t* iHandle = reinterpret_cast<ihipIpcEventHandle_t*>(&handle);
status = e->OpenHandle(iHandle);
// Free the event in case of failure
if (status != hipSuccess) {
projects/clr/hipamd/src/hip_fatbin.cpp
+2 -5
مشاهده پرونده
@@ -21,16 +21,13 @@ THE SOFTWARE.
*/
#include "hip/hip_runtime_api.h"
#include "llvm/BinaryFormat/ELF.h"
#include "hip_fatbin.hpp"
#include "hip_global.hpp"
#include <unordered_map>
#include "hip_code_object.hpp"
#include "hip_platform.hpp"
#include "comgrctx.hpp"
#include "amd_hsa_elf.hpp"
namespace hip {
namespace comgr_helper {
@@ -251,7 +248,7 @@ static bool IsCodeObjectCompressed(const void* image) {
static bool IsCodeObjectElf(const void* image) {
const amd::Elf64_Ehdr* ehdr = reinterpret_cast<const amd::Elf64_Ehdr*>(image);
return ehdr->e_machine == EM_AMDGPU && ehdr->e_ident[EI_OSABI] == llvm::ELF::ELFOSABI_AMDGPU_HSA;
return ehdr->e_machine == EM_AMDGPU && ehdr->e_ident[EI_OSABI] == ELFOSABI_AMDGPU_HSA;
}
static bool UncompressAndPopulateCodeObject(
projects/clr/hipamd/src/hip_formatting.hpp
+26
مشاهده پرونده
@@ -937,3 +937,29 @@ inline std::ostream& operator<<(std::ostream& os, const hipPitchedPtr* p) {
}
return os;
}
inline std::ostream& operator<<(std::ostream& os, const hipMemLocation& s) {
os << "{type=";
switch (s.type) {
case hipMemLocationTypeInvalid:
os << "hipMemLocationTypeInvalid";
break;
case hipMemLocationTypeDevice:
os << "hipMemLocationTypeDevice";
break;
case hipMemLocationTypeHost:
os << "hipMemLocationTypeHost";
break;
case hipMemLocationTypeHostNuma:
os << "hipMemLocationTypeHostNuma";
break;
case hipMemLocationTypeHostNumaCurrent:
os << "hipMemLocationTypeHostNumaCurrent";
break;
default:
os << static_cast<int>(s.type);
break;
}
os << ", id=" << s.id << "}";
return os;
}
projects/clr/hipamd/src/hip_graph_internal.hpp
+3
مشاهده پرونده
@@ -47,6 +47,9 @@ class GraphExec;
class UserObject;
class GraphKernelNode;
typedef GraphNode* Node;
hipError_t ihipGraphAddNode(hip::GraphNode* graphNode, hip::Graph* graph,
hip::GraphNode* const* pDependencies, size_t numDependencies,
bool capture);
class UserObject : public amd::ReferenceCountedObject {
typedef void (*UserCallbackDestructor)(void* data);
projects/clr/hipamd/src/hip_hcc.map.in
+24
مشاهده پرونده
@@ -605,3 +605,27 @@ global:
local:
*;
} hip_6.4;
hip_7.1 {
global:
hipModuleGetFunctionCount;
hipMemsetD2D8;
hipMemsetD2D8Async;
hipMemsetD2D16;
hipMemsetD2D16Async;
hipMemsetD2D32;
hipMemsetD2D32Async;
hipStreamGetAttribute;
hipStreamSetAttribute;
hipModuleLoadFatBinary;
hipMemcpyBatchAsync;
hipMemcpy3DBatchAsync;
hipMemcpy3DPeer;
hipMemcpy3DPeerAsync;
hipGetDriverEntryPoint;
hipGetDriverEntryPoint_spt;
hipMemPrefetchAsync_v2;
hipMemAdvise_v2;
local:
*;
} hip_6.5;
projects/clr/hipamd/src/hip_hmm.cpp
+152 -77
مشاهده پرونده
@@ -29,6 +29,10 @@ namespace hip {
// Forward declaraiton of a function
hipError_t ihipMallocManaged(void** ptr, size_t size, size_t align = 0, bool use_host_ptr = 0);
hipError_t ihipMemPrefetchAsync(const void* dev_ptr, size_t count, hipMemLocation location,
hipStream_t stream);
hipError_t ihipMemAdvise(const void* dev_ptr, size_t count, hipMemoryAdvise advice,
hipMemLocation location);
// Make sure HIP defines match ROCclr to avoid double conversion
static_assert(hipCpuDeviceId == amd::CpuDeviceId, "CPU device ID mismatch with ROCclr!");
@@ -83,98 +87,51 @@ hipError_t hipMallocManaged(void** dev_ptr, size_t size, unsigned int flags) {
hipError_t hipMemPrefetchAsync(const void* dev_ptr, size_t count, int device,
hipStream_t stream) {
HIP_INIT_API(hipMemPrefetchAsync, dev_ptr, count, device, stream);
if ((dev_ptr == nullptr) || (count == 0)) {
HIP_RETURN(hipErrorInvalidValue);
}
getStreamPerThread(stream);
size_t offset = 0;
amd::Memory* memObj = getMemoryObject(dev_ptr, offset);
if ((memObj != nullptr) && (count > (memObj->getSize() - offset))) {
HIP_RETURN(hipErrorInvalidValue);
}
if (device != hipCpuDeviceId && (static_cast<size_t>(device) >= g_devices.size())) {
HIP_RETURN(hipErrorInvalidDevice);
}
hip::Stream* hip_stream = nullptr;
amd::Device* dev = nullptr;
bool cpu_access = false;
if ((memObj == nullptr) && (device != hipCpuDeviceId) &&
(!g_devices[device]->devices()[0]->info().hmmCpuMemoryAccessible_)) {
HIP_RETURN(hipErrorNotSupported);
}
// Pick the specified stream or Null one from the provided device
CHECK_STREAM_CAPTURE_SUPPORTED();
hipMemLocation location;
if (device == hipCpuDeviceId) {
cpu_access = true;
hip_stream = (stream == nullptr || stream == hipStreamLegacy) ?
hip::getCurrentDevice()->NullStream() : hip::getStream(stream);
location.type = hipMemLocationTypeHost;
location.id = hipCpuDeviceId;
} else {
dev = g_devices[device]->devices()[0];
hip_stream = (stream == nullptr || stream == hipStreamLegacy) ?
g_devices[device]->NullStream() : hip::getStream(stream);
location.type = hipMemLocationTypeDevice;
location.id = device;
}
HIP_RETURN(ihipMemPrefetchAsync(dev_ptr, count, location, stream));
}
if (hip_stream == nullptr) {
// ================================================================================================
hipError_t hipMemPrefetchAsync_v2(const void* dev_ptr, size_t count, hipMemLocation location,
unsigned int flags, hipStream_t stream) {
HIP_INIT_API(hipMemPrefetchAsync_v2, dev_ptr, count, location, flags, stream);
CHECK_STREAM_CAPTURE_SUPPORTED();
if (flags != 0) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Command::EventWaitList waitList;
amd::SvmPrefetchAsyncCommand* command =
new amd::SvmPrefetchAsyncCommand(*hip_stream, waitList, dev_ptr, count, dev, cpu_access);
if (command == nullptr) {
return hipErrorOutOfMemory;
}
command->enqueue();
command->release();
HIP_RETURN(hipSuccess);
HIP_RETURN(ihipMemPrefetchAsync(dev_ptr, count, location, stream));
}
// ================================================================================================
hipError_t hipMemAdvise(const void* dev_ptr, size_t count, hipMemoryAdvise advice, int device) {
HIP_INIT_API(hipMemAdvise, dev_ptr, count, advice, device);
CHECK_STREAM_CAPTURE_SUPPORTED();
bool isAdviseReadMostly = (advice == hipMemAdviseSetReadMostly) ||
(advice == hipMemAdviseUnsetReadMostly);
if (!isAdviseReadMostly && ((device != hipCpuDeviceId) &&
(static_cast<size_t>(device) >= g_devices.size()))) {
HIP_RETURN(hipErrorInvalidDevice);
hipMemLocation location;
if (device == hipCpuDeviceId) {
location.type = hipMemLocationTypeHost;
location.id = hipCpuDeviceId;
} else {
location.type = hipMemLocationTypeDevice;
location.id = device;
}
if ((dev_ptr == nullptr) || (count == 0)) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(ihipMemAdvise(dev_ptr, count, advice, location));
}
if (!hip::tls.capture_streams_.empty() || !g_captureStreams.empty()) {
HIP_RETURN(hipErrorStreamCaptureUnsupported);
}
size_t offset = 0;
amd::Memory* memObj = getMemoryObject(dev_ptr, offset);
if (memObj && count > (memObj->getSize() - offset)) {
HIP_RETURN(hipErrorInvalidValue);
}
amd::Device* dev = (device == hipCpuDeviceId || isAdviseReadMostly) ?
g_devices[0]->devices()[0] : g_devices[device]->devices()[0];
bool use_cpu = (device == hipCpuDeviceId) ? true : false;
// Set the allocation attributes in AMD HMM
if (!dev->SetSvmAttributes(dev_ptr, count, static_cast<amd::MemoryAdvice>(advice), use_cpu)) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(hipSuccess);
// ================================================================================================
hipError_t hipMemAdvise_v2(const void* dev_ptr, size_t count, hipMemoryAdvise advice,
hipMemLocation location) {
HIP_INIT_API(hipMemAdvise_v2, dev_ptr, count, advice, location);
CHECK_STREAM_CAPTURE_SUPPORTED();
HIP_RETURN(ihipMemAdvise(dev_ptr, count, advice, location));
}
// ================================================================================================
@@ -322,4 +279,122 @@ hipError_t ihipMallocManaged(void** ptr, size_t size, size_t align, bool use_hos
ClPrint(amd::LOG_INFO, amd::LOG_API, "ihipMallocManaged ptr=0x%zx", *ptr);
return hipSuccess;
}
// ================================================================================================
hipError_t ihipMemPrefetchAsync(const void* dev_ptr, size_t count, hipMemLocation location,
hipStream_t stream) {
if ((dev_ptr == nullptr) || (count == 0)) {
return hipErrorInvalidValue;
}
getStreamPerThread(stream);
size_t offset = 0;
amd::Memory* memObj = getMemoryObject(dev_ptr, offset);
if ((memObj != nullptr) && (count > (memObj->getSize() - offset))) {
return hipErrorInvalidValue;
}
// Compute the type of prefetch
const bool isHost = (location.type == hipMemLocationTypeHost);
const bool isHostNuma = (location.type == hipMemLocationTypeHostNuma);
const bool isHostCurrent = (location.type == hipMemLocationTypeHostNumaCurrent);
const bool cpuAccess = isHost || isHostNuma || isHostCurrent;
// Determine the target device index:
// - for host-prefetch and host-current, always use device 0
// - for host-NUMA or device-prefetch, use the provided id
int targetDevice = (isHost || isHostCurrent) ? hipCpuDeviceId : location.id;
amd::Device* dev = nullptr;
if (cpuAccess == false) {
if (static_cast<size_t>(targetDevice) >= g_devices.size()) {
return hipErrorInvalidDevice;
}
dev = g_devices[targetDevice]->devices()[0];
if (memObj == nullptr && !dev->info().hmmCpuMemoryAccessible_) {
return hipErrorNotSupported;
}
}
hip::Stream* hip_stream = nullptr;
// Pick the specified stream or Null one from the provided target device
if (cpuAccess == true) {
hip_stream = (stream == nullptr || stream == hipStreamLegacy)
? hip::getCurrentDevice()->NullStream()
: hip::getStream(stream);
} else {
dev = g_devices[targetDevice]->devices()[0];
hip_stream = (stream == nullptr || stream == hipStreamLegacy)
? g_devices[targetDevice]->NullStream()
: hip::getStream(stream);
}
if (hip_stream == nullptr) {
return hipErrorInvalidValue;
}
amd::Command::EventWaitList waitList;
amd::SvmPrefetchAsyncCommand* command = new amd::SvmPrefetchAsyncCommand(
*hip_stream, waitList, dev_ptr, count, dev, cpuAccess, targetDevice);
if (command == nullptr) {
return hipErrorOutOfMemory;
}
command->enqueue();
command->release();
return hipSuccess;
}
// ================================================================================================
hipError_t ihipMemAdvise(const void* dev_ptr, size_t count, hipMemoryAdvise advice,
hipMemLocation location) {
if ((dev_ptr == nullptr) || (count == 0)) {
return hipErrorInvalidValue;
}
if (!hip::tls.capture_streams_.empty() || !g_captureStreams.empty()) {
return hipErrorStreamCaptureUnsupported;
}
// Determine device and CPU access from location
int targetDevice = hipCpuDeviceId;
bool use_cpu = true;
bool isAdviseReadMostly =
(advice == hipMemAdviseSetReadMostly) || (advice == hipMemAdviseUnsetReadMostly);
switch (location.type) {
case hipMemLocationTypeDevice:
targetDevice = location.id;
use_cpu = false;
break;
case hipMemLocationTypeHostNuma:
targetDevice = location.id; // NUMA node ID
use_cpu = true;
break;
case hipMemLocationTypeHost:
case hipMemLocationTypeHostNumaCurrent:
targetDevice = hipCpuDeviceId;
use_cpu = true;
break;
default:
return hipErrorInvalidValue;
}
if (!isAdviseReadMostly && !use_cpu && (static_cast<size_t>(targetDevice) >= g_devices.size())) {
return hipErrorInvalidDevice;
}
size_t offset = 0;
amd::Memory* memObj = getMemoryObject(dev_ptr, offset);
if (memObj && count > (memObj->getSize() - offset)) {
return hipErrorInvalidValue;
}
amd::Device* dev = (use_cpu || isAdviseReadMostly) ? g_devices[0]->devices()[0]
: g_devices[targetDevice]->devices()[0];
// Set the allocation attributes in AMD HMM
if (!dev->SetSvmAttributes(dev_ptr, count, static_cast<amd::MemoryAdvice>(advice), use_cpu,
targetDevice)) {
return hipErrorInvalidValue;
}
return hipSuccess;
}
} //namespace hip
projects/clr/hipamd/src/hip_internal.hpp
+2
مشاهده پرونده
@@ -680,6 +680,8 @@ public:
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);
hipError_t ihipMemcpy3D(const hipMemcpy3DParms* p, hipStream_t stream = nullptr,
bool isAsync = false);
constexpr bool kOptionChangeable = true;
constexpr bool kNewDevProg = false;
projects/clr/hipamd/src/hip_memory.cpp
+125 -34
مشاهده پرونده
@@ -1087,18 +1087,11 @@ amd::Image* ihipImageCreate(const cl_channel_order channelOrder,
case CL_MEM_OBJECT_IMAGE1D:
case CL_MEM_OBJECT_IMAGE2D:
case CL_MEM_OBJECT_IMAGE3D:
image = new (context) amd::Image(*buffer->asBuffer(),
imageType,
CL_MEM_READ_WRITE,
imageFormat,
imageWidth,
(imageHeight == 0) ? 1 : imageHeight,
(imageDepth == 0) ? 1 : imageDepth,
imageRowPitch,
imageSlicePitch,
numMipLevels,
offset);
break;
image = new (buffer->getContext())
amd::Image(*buffer->asBuffer(), imageType, CL_MEM_READ_WRITE, imageFormat, imageWidth,
(imageHeight == 0) ? 1 : imageHeight, (imageDepth == 0) ? 1 : imageDepth,
imageRowPitch, imageSlicePitch, numMipLevels, offset);
break;
default:
LogPrintfError("Cannot create image of imageType: 0x%x for external buffer", imageType);
}
@@ -1106,17 +1099,10 @@ amd::Image* ihipImageCreate(const cl_channel_order channelOrder,
switch (imageType) {
case CL_MEM_OBJECT_IMAGE1D_BUFFER:
case CL_MEM_OBJECT_IMAGE2D:
image = new (context) amd::Image(*buffer->asBuffer(),
imageType,
CL_MEM_READ_WRITE,
imageFormat,
(imageWidth == 0) ? 1 : imageWidth,
(imageHeight == 0) ? 1 : imageHeight,
(imageDepth == 0) ? 1 : imageDepth,
imageRowPitch,
imageSlicePitch,
numMipLevels,
offset);
image = new (buffer->getContext()) amd::Image(
*buffer->asBuffer(), imageType, CL_MEM_READ_WRITE, imageFormat,
(imageWidth == 0) ? 1 : imageWidth, (imageHeight == 0) ? 1 : imageHeight,
(imageDepth == 0) ? 1 : imageDepth, imageRowPitch, imageSlicePitch, numMipLevels, offset);
break;
default:
LogPrintfError("Cannot create image of imageType: 0x%x", imageType);
@@ -2773,7 +2759,7 @@ hipError_t ihipMemcpy3DCommand(amd::Command*& command, const hipMemcpy3DParms* p
return ihipGetMemcpyParam3DCommand(command, &desc, stream);
}
hipError_t ihipMemcpy3D(const hipMemcpy3DParms* p, hipStream_t stream, bool isAsync = false) {
hipError_t ihipMemcpy3D(const hipMemcpy3DParms* p, hipStream_t stream, bool isAsync) {
hipError_t status = ihipMemcpy3D_validate(p);
if (status != hipSuccess) {
return status;
@@ -2822,10 +2808,72 @@ hipError_t hipDrvMemcpy3D(const HIP_MEMCPY3D* pCopy) {
hipError_t hipDrvMemcpy3DAsync(const HIP_MEMCPY3D* pCopy, hipStream_t stream) {
HIP_INIT_API(hipDrvMemcpy3DAsync, pCopy, stream);
HIP_RETURN_DURATION(ihipMemcpyParam3D(pCopy, stream, true));
}
hipError_t hipMemcpyBatchAsync(void **dsts, void **srcs, size_t *sizes, size_t count,
hipMemcpyAttributes *attrs, size_t *attrsIdxs, size_t numAttrs,
size_t *failIdx, hipStream_t stream) {
HIP_INIT_API(hipMemcpyBatchAsync, dsts, srcs, sizes, count, attrs, attrsIdxs, numAttrs, failIdx,
stream);
// validate stream
if(!hip::isValid(stream)) {
HIP_RETURN(hipErrorInvalidResourceHandle);
}
// validate inputs
if (dsts == nullptr || srcs == nullptr || sizes == nullptr || failIdx == nullptr ||
count == 0) {
HIP_RETURN(hipErrorInvalidValue);
}
// no support for memcpy attributes
if (numAttrs > 0) {
HIP_RETURN(hipErrorNotSupported);
}
hipError_t status = hipSuccess;
*failIdx = SIZE_MAX;
for (int i = 0; i < count; ++i) {
if (sizes[i] == 0) {
*failIdx = i;
status = hipErrorInvalidValue;
break;
}
status = ihipMemcpy(dsts[i], srcs[i], sizes[i], hipMemcpyDefault, *hip::getStream(stream),
true, true);
if (status != hipSuccess) {
*failIdx = i;
break;
}
}
HIP_RETURN(status);
}
hipError_t hipMemcpy3DBatchAsync(size_t numOps, struct hipMemcpy3DBatchOp *opList, size_t *failIdx,
unsigned long long flags, hipStream_t stream) {
HIP_INIT_API(hipMemcpy3DBatchAsync, numOps, opList, failIdx, flags, stream);
if (flags != 0 || opList == nullptr || numOps == 0) {
HIP_RETURN(hipErrorInvalidValue);
}
if (!hip::isValid(stream)) {
HIP_RETURN(hipErrorInvalidResourceHandle);
}
hipError_t status = hipSuccess;
*failIdx = SIZE_MAX;
for (int i = 0; i < numOps; ++i) {
hipMemcpy3DParms parms = getMemcpy3DParms(opList[i]);
status = ihipMemcpy3D(&parms, stream, true);
if (status != hipSuccess) {
*failIdx = i;
break;
}
}
HIP_RETURN(status);
}
hipError_t packFillMemoryCommand(amd::Command*& command, amd::Memory* memory, size_t offset,
int64_t value, size_t valueSize, size_t sizeBytes,
hip::Stream* stream) {
@@ -3117,7 +3165,7 @@ hipError_t ihipMemset3DCommand(std::vector<amd::Command*> &commands, hipPitchedP
hipError_t ihipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent,
hipStream_t stream, bool isAsync = false) {
hipStream_t stream, bool isAsync = false, size_t elementSize = 1) {
auto sizeBytes = extent.width * extent.height * extent.depth;
if (sizeBytes == 0) {
@@ -3142,7 +3190,7 @@ hipError_t ihipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent
}
hip::Stream* hip_stream = hip::getStream(stream);
std::vector<amd::Command*> commands;
status = ihipMemset3DCommand(commands, pitchedDevPtr, value, extent, hip_stream);
status = ihipMemset3DCommand(commands, pitchedDevPtr, value, extent, hip_stream, elementSize);
if (status != hipSuccess) {
return status;
}
@@ -3157,9 +3205,10 @@ hipError_t ihipMemset3D(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent
}
hipError_t hipMemset2D_common(void* dst, size_t pitch, int value, size_t width,
size_t height, hipStream_t stream=nullptr) {
size_t height, hipStream_t stream=nullptr, size_t elementSize = 1) {
CHECK_STREAM_CAPTURING();
return ihipMemset3D({dst, pitch, width, height}, value, {width, height, 1}, stream);
return ihipMemset3D({dst, pitch, width, height}, value, {width, height, 1}, stream, false,
elementSize);
}
hipError_t hipMemset2D_spt(void* dst, size_t pitch, int value, size_t width, size_t height) {
@@ -3173,11 +3222,11 @@ hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t
HIP_RETURN(hipMemset2D_common(dst, pitch, value, width, height));
}
hipError_t hipMemset2DAsync_common(void* dst, size_t pitch, int value,
size_t width, size_t height, hipStream_t stream) {
hipError_t hipMemset2DAsync_common(void* dst, size_t pitch, int value, size_t width, size_t height,
hipStream_t stream, size_t elementSize = 1) {
STREAM_CAPTURE(hipMemset2DAsync, stream, dst, pitch, value, width, height);
return ihipMemset3D({dst, pitch, width, height}, value, {width, height, 1}, stream, true);
return ihipMemset3D({dst, pitch, width, height}, value, {width, height, 1}, stream, true,
elementSize);
}
hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value,
@@ -3193,6 +3242,48 @@ hipError_t hipMemset2DAsync_spt(void* dst, size_t pitch, int value,
HIP_RETURN(hipMemset2DAsync_common(dst, pitch, value, width, height, stream));
}
hipError_t hipMemsetD2D8(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, size_t width,
size_t height) {
HIP_INIT_API(hipMemsetD2D8, dst, dstPitch, value, width, height);
HIP_RETURN(hipMemset2D_common(dst, dstPitch, value, width, height, nullptr,
sizeof(unsigned char)));
}
hipError_t hipMemsetD2D8Async(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, size_t width,
size_t height, hipStream_t stream) {
HIP_INIT_API(hipMemsetD2D8Async, dst, dstPitch, value, width, height, stream);
HIP_RETURN(hipMemset2DAsync_common(dst, dstPitch, value, width, height, stream,
sizeof(unsigned char)));
}
hipError_t hipMemsetD2D16(hipDeviceptr_t dst, size_t dstPitch, unsigned short value, size_t width,
size_t height) {
HIP_INIT_API(hipMemsetD2D16, dst, dstPitch, value, width, height);
HIP_RETURN(hipMemset2D_common(dst, dstPitch, value, width, height, nullptr,
sizeof(unsigned short)));
}
hipError_t hipMemsetD2D16Async(hipDeviceptr_t dst, size_t dstPitch, unsigned short value,
size_t width, size_t height, hipStream_t stream) {
HIP_INIT_API(hipMemsetD2D16Async, dst, dstPitch, value, width, height, stream);
HIP_RETURN(hipMemset2DAsync_common(dst, dstPitch, value, width, height, stream,
sizeof(unsigned short)));
}
hipError_t hipMemsetD2D32(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, size_t width,
size_t height) {
HIP_INIT_API(hipMemsetD2D32, dst, dstPitch, value, width, height);
HIP_RETURN(hipMemset2D_common(dst, dstPitch, value, width, height, nullptr,
sizeof(unsigned int)));
}
hipError_t hipMemsetD2D32Async(hipDeviceptr_t dst, size_t dstPitch, unsigned int value,
size_t width, size_t height, hipStream_t stream) {
HIP_INIT_API(hipMemsetD2D32Async, dst, dstPitch, value, width, height, stream);
HIP_RETURN(hipMemset2DAsync_common(dst, dstPitch, value, width, height, stream,
sizeof(unsigned int)));
}
// ================================================================================================
hipError_t hipMemset3D_common(hipPitchedPtr pitchedDevPtr, int value, hipExtent extent,
hipStream_t stream=nullptr) {
@@ -4310,7 +4401,7 @@ hipError_t hipExternalMemoryGetMappedMipmappedArray(
(size_t)mipmapDesc->offset, buf));
}
hipError_t hipMemGetHandleForAddressRange(void* handle, hipDeviceptr_t dptr, size_t size,
hipError_t hipMemGetHandleForAddressRange(void* handle, hipDeviceptr_t dptr, size_t size,
hipMemRangeHandleType handleType,
unsigned long long flags) {
HIP_INIT_API(hipMemGetHandleForAddressRange, handle, dptr, size, handleType, flags);
projects/clr/hipamd/src/hip_module.cpp
+15
مشاهده پرونده
@@ -48,6 +48,11 @@ hipError_t hipModuleUnload(hipModule_t hmod) {
CHECK_STREAM_CAPTURE_SUPPORTED();
HIP_RETURN(PlatformState::instance().unloadModule(hmod));
}
hipError_t hipModuleLoadFatBinary(hipModule_t* module, const void* fatbin) {
HIP_INIT_API(hipModuleLoadFatBinary, module, fatbin);
HIP_RETURN(PlatformState::instance().loadModule(module, 0, fatbin));
HIP_RETURN(hipSuccess);
}
hipError_t hipModuleLoad(hipModule_t* module, const char* fname) {
HIP_INIT_API(hipModuleLoad, module, fname);
@@ -89,6 +94,15 @@ hipError_t hipModuleGetFunction(hipFunction_t* hfunc, hipModule_t hmod, const ch
HIP_RETURN(hipSuccess);
}
hipError_t hipModuleGetFunctionCount(unsigned int* count, hipModule_t mod) {
HIP_INIT_API(hipModuleGetFunctionCount, count, mod);
if (mod == nullptr) {
HIP_RETURN(hipErrorInvalidResourceHandle);
}
HIP_RETURN(PlatformState::instance().getFuncCount(count, mod););
}
hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod,
const char* name) {
HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name);
@@ -371,6 +385,7 @@ hipError_t ihipLaunchKernelCommand(amd::Command*& command, hipFunction_t f,
assert(kernelParams == nullptr);
if (extra[0] != HIP_LAUNCH_PARAM_BUFFER_POINTER || extra[2] != HIP_LAUNCH_PARAM_BUFFER_SIZE ||
extra[4] != HIP_LAUNCH_PARAM_END) {
kernelCommand->release();
return hipErrorInvalidValue;
}
kernargs = reinterpret_cast<address>(extra[1]);
projects/clr/hipamd/src/hip_peer.cpp
+35
مشاهده پرونده
@@ -21,6 +21,8 @@
#include <hip/hip_runtime.h>
#include "hip_internal.hpp"
#include "hip_conversions.hpp"
namespace hip {
@@ -217,6 +219,39 @@ hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int src
HIP_RETURN(ihipMemcpy(dst, src, sizeBytes, hipMemcpyDeviceToDevice, *hip_stream, true, true));
}
hipError_t hipMemcpy3DPeer(hipMemcpy3DPeerParms *p) {
HIP_INIT_API(hipMemcpy3DPeer, p);
if (p == NULL) {
HIP_RETURN(hipErrorInvalidValue);
}
if (p->srcDevice >= static_cast<int>(g_devices.size()) ||
p->dstDevice >= static_cast<int>(g_devices.size()) || p->srcDevice < 0 || p->dstDevice < 0) {
HIP_RETURN(hipErrorInvalidDevice);
}
hipMemcpy3DParms copyParms = getMemcpy3DParms(*p);
HIP_RETURN(ihipMemcpy3D(&copyParms, nullptr));
}
hipError_t hipMemcpy3DPeerAsync(hipMemcpy3DPeerParms *p, hipStream_t stream) {
HIP_INIT_API(hipMemcpy3DPeerAsync, p, stream);
if (p == NULL) {
HIP_RETURN(hipErrorInvalidValue);
}
if (p->srcDevice >= static_cast<int>(g_devices.size()) ||
p->dstDevice >= static_cast<int>(g_devices.size()) || p->srcDevice < 0 || p->dstDevice < 0) {
HIP_RETURN(hipErrorInvalidDevice);
}
getStreamPerThread(stream);
hip::Stream* hip_stream = hip::getStream(stream);
if (hip_stream == nullptr) {
return hipErrorInvalidValue;
}
hipMemcpy3DParms copyParms = getMemcpy3DParms(*p);
HIP_RETURN(ihipMemcpy3D(&copyParms), stream, true);
}
hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) {
HIP_INIT_API(hipCtxEnablePeerAccess, peerCtx, flags);
projects/clr/hipamd/src/hip_platform.cpp
+11
مشاهده پرونده
@@ -834,6 +834,17 @@ hipError_t PlatformState::getDynFunc(hipFunction_t* hfunc, hipModule_t hmod,
return it->second->getDynFunc(hfunc, func_name);
}
hipError_t PlatformState::getFuncCount(unsigned int* count, hipModule_t hmod) {
amd::ScopedLock lock(lock_);
auto it = dynCO_map_.find(hmod);
if (it == dynCO_map_.end()) {
LogPrintfError("Cannot find the module: 0x%x", hmod);
return hipErrorNotFound;
}
return it->second->getFuncCount(count);
}
bool PlatformState::isValidDynFunc(const void* hfunc) {
amd::ScopedLock lock(lock_);
return std::any_of(dynCO_map_.begin(), dynCO_map_.end(),
projects/clr/hipamd/src/hip_platform.hpp
+1
مشاهده پرونده
@@ -63,6 +63,7 @@ class PlatformState {
hipError_t unloadModule(hipModule_t hmod);
bool isValidDynFunc(const void* hfunc);
hipError_t getDynFunc(hipFunction_t* hfunc, hipModule_t hmod, const char* func_name);
hipError_t getFuncCount(unsigned int* count, hipModule_t hmod);
hipError_t getDynGlobalVar(const char* hostVar, hipModule_t hmod, hipDeviceptr_t* dev_ptr,
size_t* size_ptr);
hipError_t getDynTexRef(const char* hostVar, hipModule_t hmod, textureReference** texRef);
projects/clr/hipamd/src/hip_stream.cpp
+72 -1
مشاهده پرونده
@@ -560,7 +560,6 @@ hipError_t hipStreamQuery_spt(hipStream_t stream) {
hipError_t streamCallback_common(hipStream_t stream, StreamCallback* cbo, void* userData) {
getStreamPerThread(stream);
CHECK_SUPPORTED_DURING_CAPTURE();
hip::Stream* hip_stream = hip::getStream(stream);
amd::Command* last_command = hip_stream->getLastQueuedCommand(true);
amd::Command::EventWaitList eventWaitList;
@@ -821,4 +820,76 @@ hipError_t hipStreamGetDevice(hipStream_t stream, hipDevice_t* device) {
HIP_RETURN(hipSuccess);
}
// ================================================================================================
hipError_t hipStreamSetAttribute(hipStream_t stream, hipStreamAttrID attr,
const hipStreamAttrValue *value) {
HIP_INIT_API(hipStreamSetAttribute, stream, attr, value);
hipError_t status = hipSuccess;
if (value == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
if (!hip::isValid(stream)) {
HIP_RETURN(hipErrorInvalidResourceHandle);
}
getStreamPerThread(stream);
// if stream is capturing, don't allow changing stream attributes
if (hip::Stream::StreamCaptureOngoing(stream) == true) {
HIP_RETURN(hipErrorStreamCaptureUnsupported);
}
hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
switch (attr) {
case hipStreamAttributeSynchronizationPolicy: {
hipSynchronizationPolicy syncPolicy = value->syncPolicy;
// validate sync policy
if (syncPolicy < hipSyncPolicyAuto || syncPolicy > hipSyncPolicyBlockingSync) {
HIP_RETURN(hipErrorInvalidValue);
}
s->SetSyncPolicy(static_cast<amd::SyncPolicy>(syncPolicy));
break;
}
default: {
HIP_RETURN(hipErrorInvalidValue);
}
}
HIP_RETURN(hipSuccess);
}
hipError_t hipStreamGetAttribute(hipStream_t stream, hipStreamAttrID attr,
hipStreamAttrValue *value_out) {
HIP_INIT_API(hipStreamGetAttribute, stream, attr, value_out);
if (value_out == nullptr) {
return hipErrorInvalidValue;
}
if (!hip::isValid(stream)) {
HIP_RETURN(hipErrorInvalidResourceHandle);
}
getStreamPerThread(stream);
hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
switch(attr) {
case hipStreamAttributeSynchronizationPolicy: {
value_out->syncPolicy = static_cast<hipSynchronizationPolicy>(s->GetSyncPolicy());
break;
}
case hipStreamAttributePriority: {
value_out->priority = s->GetPriority();
break;
}
default: {
HIP_RETURN(hipErrorInvalidValue);
}
}
HIP_RETURN(hipSuccess);
}
} // hip namespace
projects/clr/hipamd/src/hip_table_interface.cpp
+78 -1
مشاهده پرونده
@@ -419,6 +419,16 @@ extern "C" hipError_t hipGetDevicePropertiesR0600(hipDeviceProp_tR0600* prop, in
extern "C" hipError_t hipGetDevicePropertiesR0000(hipDeviceProp_tR0000* prop, int device) {
return hip::GetHipDispatchTable()->hipGetDevicePropertiesR0000_fn(prop, device);
}
hipError_t hipGetDriverEntryPoint(const char* symbol, void** funcPtr, unsigned long long flags,
hipDriverEntryPointQueryResult* status) {
return hip::GetHipDispatchTable()->hipGetDriverEntryPoint_fn(symbol, funcPtr, flags,
status);
}
hipError_t hipGetDriverEntryPoint_spt(const char* symbol, void** funcPtr, unsigned long long flags,
hipDriverEntryPointQueryResult* status) {
return hip::GetHipDispatchTable()->hipGetDriverEntryPoint_spt_fn(symbol, funcPtr, flags,
status);
}
const char* hipGetErrorName(hipError_t hip_error) {
return hip::GetHipDispatchTable()->hipGetErrorName_fn(hip_error);
}
@@ -926,6 +936,10 @@ hipError_t hipMemAddressReserve(void** ptr, size_t size, size_t alignment, void*
hipError_t hipMemAdvise(const void* dev_ptr, size_t count, hipMemoryAdvise advice, int device) {
return hip::GetHipDispatchTable()->hipMemAdvise_fn(dev_ptr, count, advice, device);
}
hipError_t hipMemAdvise_v2(const void* dev_ptr, size_t count, hipMemoryAdvise advice,
hipMemLocation location) {
return hip::GetHipDispatchTable()->hipMemAdvise_v2_fn(dev_ptr, count, advice, location);
}
hipError_t hipMemAllocHost(void** ptr, size_t size) {
return hip::GetHipDispatchTable()->hipMemAllocHost_fn(ptr, size);
}
@@ -1020,6 +1034,11 @@ hipError_t hipMemPoolTrimTo(hipMemPool_t mem_pool, size_t min_bytes_to_hold) {
hipError_t hipMemPrefetchAsync(const void* dev_ptr, size_t count, int device, hipStream_t stream) {
return hip::GetHipDispatchTable()->hipMemPrefetchAsync_fn(dev_ptr, count, device, stream);
}
hipError_t hipMemPrefetchAsync_v2(const void* dev_ptr, size_t count, hipMemLocation location,
unsigned int flags, hipStream_t stream) {
return hip::GetHipDispatchTable()->hipMemPrefetchAsync_v2_fn(dev_ptr, count, location, flags,
stream);
}
hipError_t hipMemPtrGetInfo(void* ptr, size_t* size) {
return hip::GetHipDispatchTable()->hipMemPtrGetInfo_fn(ptr, size);
}
@@ -1220,6 +1239,9 @@ hipError_t hipMipmappedArrayGetLevel(hipArray_t* pLevelArray, hipMipmappedArray_
hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, const char* kname) {
return hip::GetHipDispatchTable()->hipModuleGetFunction_fn(function, module, kname);
}
hipError_t hipModuleGetFunctionCount(unsigned int* count, hipModule_t mod) {
return hip::GetHipDispatchTable()->hipModuleGetFunctionCount_fn(count, mod);
}
hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod,
const char* name) {
return hip::GetHipDispatchTable()->hipModuleGetGlobal_fn(dptr, bytes, hmod, name);
@@ -1251,6 +1273,9 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigne
f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, stream,
kernelParams, extra);
}
hipError_t hipModuleLoadFatBinary(hipModule_t* module, const void* fatbin) {
return hip::GetHipDispatchTable()->hipModuleLoadFatBinary_fn(module, fatbin);
}
hipError_t hipModuleLoad(hipModule_t* module, const char* fname) {
return hip::GetHipDispatchTable()->hipModuleLoad_fn(module, fname);
}
@@ -1880,9 +1905,61 @@ hipError_t hipDrvLaunchKernelEx(const HIP_LAUNCH_CONFIG* config, hipFunction_t f
return hip::GetHipDispatchTable()->hipDrvLaunchKernelEx_fn(config, f, kernel, extra);
}
hipError_t hipMemGetHandleForAddressRange(void* handle, hipDeviceptr_t dptr, size_t size,
hipError_t hipMemGetHandleForAddressRange(void* handle, hipDeviceptr_t dptr, size_t size,
hipMemRangeHandleType handleType,
unsigned long long flags) {
return hip::GetHipDispatchTable()->hipMemGetHandleForAddressRange_fn(handle, dptr, size,
handleType, flags);
}
hipError_t hipMemsetD2D8(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, size_t width,
size_t height) {
return hip::GetHipDispatchTable()->hipMemsetD2D8_fn(dst, dstPitch, value, width, height);
}
hipError_t hipMemsetD2D8Async(hipDeviceptr_t dst, size_t dstPitch, unsigned char value, size_t width,
size_t height, hipStream_t stream) {
return hip::GetHipDispatchTable()->hipMemsetD2D8Async_fn(dst, dstPitch, value, width, height,
stream);
}
hipError_t hipMemsetD2D16(hipDeviceptr_t dst, size_t dstPitch, unsigned short value, size_t width,
size_t height) {
return hip::GetHipDispatchTable()->hipMemsetD2D16_fn(dst, dstPitch, value, width, height);
}
hipError_t hipMemsetD2D16Async(hipDeviceptr_t dst, size_t dstPitch, unsigned short value, size_t width,
size_t height, hipStream_t stream) {
return hip::GetHipDispatchTable()->hipMemsetD2D16Async_fn(dst, dstPitch, value, width, height,
stream);
}
hipError_t hipMemsetD2D32(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, size_t width,
size_t height) {
return hip::GetHipDispatchTable()->hipMemsetD2D32_fn(dst, dstPitch, value, width, height);
}
hipError_t hipMemsetD2D32Async(hipDeviceptr_t dst, size_t dstPitch, unsigned int value, size_t width,
size_t height, hipStream_t stream) {
return hip::GetHipDispatchTable()->hipMemsetD2D32Async_fn(dst, dstPitch, value, width, height,
stream);
}
hipError_t hipStreamSetAttribute(hipStream_t stream, hipStreamAttrID attr,
const hipStreamAttrValue *value) {
return hip::GetHipDispatchTable()->hipStreamSetAttribute_fn(stream, attr, value);
}
hipError_t hipStreamGetAttribute(hipStream_t stream, hipStreamAttrID attr,
hipStreamAttrValue *value) {
return hip::GetHipDispatchTable()->hipStreamGetAttribute_fn(stream, attr, value);
}
hipError_t hipMemcpyBatchAsync(void **dsts, void **srcs, size_t *sizes, size_t count,
hipMemcpyAttributes *attrs, size_t *attrsIdxs, size_t numAttrs,
size_t *failIdx, hipStream_t stream) {
return hip::GetHipDispatchTable()->hipMemcpyBatchAsync_fn(dsts, srcs, sizes, count, attrs,
attrsIdxs, numAttrs, failIdx, stream);
}
hipError_t hipMemcpy3DBatchAsync(size_t numOps, struct hipMemcpy3DBatchOp *opList, size_t *failIdx,
unsigned long long flags, hipStream_t stream) {
return hip::GetHipDispatchTable()->hipMemcpy3DBatchAsync_fn(numOps, opList, failIdx, flags,
stream);
}
hipError_t hipMemcpy3DPeer(hipMemcpy3DPeerParms *p) {
return hip::GetHipDispatchTable()->hipMemcpy3DPeer_fn(p);
}
hipError_t hipMemcpy3DPeerAsync(hipMemcpy3DPeerParms *p, hipStream_t stream) {
return hip::GetHipDispatchTable()->hipMemcpy3DPeerAsync_fn(p, stream);
}
projects/clr/hipamd/src/hip_texture.cpp
+4 -1
مشاهده پرونده
@@ -123,7 +123,10 @@ hipError_t ihipCreateTextureObject(hipTextureObject_t* pTexObject,
LogInfo("Mipmap not supported on the device");
return hipErrorNotSupported;
}
if (pResDesc->res.mipmap.mipmap == nullptr || pTexDesc->normalizedCoords == 0) {
if (pResDesc->res.mipmap.mipmap == nullptr) {
return hipErrorInvalidValue;
}
if (pTexDesc->normalizedCoords == 0) {
return hipErrorInvalidChannelDescriptor;
}
}
projects/clr/hipamd/src/hip_vm.cpp
+11 -3
مشاهده پرونده
@@ -85,11 +85,15 @@ hipError_t hipMemCreate(hipMemGenericAllocationHandle_t* handle, size_t size,
HIP_INIT_API(hipMemCreate, handle, size, prop, flags);
// Currently we do not support Pinned memory
if (handle == nullptr || size == 0 || flags != 0 || prop == nullptr ||
if (handle == nullptr || size == 0 || prop == nullptr ||
prop->type != hipMemAllocationTypePinned || prop->location.type != hipMemLocationTypeDevice) {
HIP_RETURN(hipErrorInvalidValue);
}
if (flags != hipDeviceMallocUncached && flags != 0) {
HIP_RETURN(hipErrorInvalidValue);
}
if (prop->location.id < 0 || prop->location.id >= g_devices.size()) {
HIP_RETURN(hipErrorInvalidDevice);
}
@@ -112,8 +116,12 @@ hipError_t hipMemCreate(hipMemGenericAllocationHandle_t* handle, size_t size,
amd::Context* amdContext = g_devices[prop->location.id]->asContext();
// When ROCCLR_MEM_PHYMEM is set, ROCr impl gets and stores unique hsa handle. Flag no-op on PAL.
void* ptr = amd::SvmBuffer::malloc(*amdContext, ROCCLR_MEM_PHYMEM, size,
dev_info.memBaseAddrAlign_, nullptr);
uint64_t ihipFlags = ROCCLR_MEM_PHYMEM;
if (flags == hipDeviceMallocUncached) {
ihipFlags |= ROCCLR_MEM_HSA_UNCACHED | CL_MEM_SVM_ATOMICS;
}
void* ptr =
amd::SvmBuffer::malloc(*amdContext, ihipFlags, size, dev_info.memBaseAddrAlign_, nullptr);
// Handle out of memory cases,
if (ptr == nullptr) {
projects/clr/hipamd/src/hiprtc/CMakeLists.txt
-1
مشاهده پرونده
@@ -114,7 +114,6 @@ else()
target_include_directories(hiprtc PRIVATE $<TARGET_PROPERTY:rocclr,INCLUDE_DIRECTORIES>)
endif()
target_include_directories(hiprtc PRIVATE ${LLVM_INCLUDE_DIRS})
target_compile_definitions(hiprtc PUBLIC __HIP_PLATFORM_AMD__)
add_to_config(_versionInfo HIP_PACKAGING_VERSION_PATCH)
projects/clr/opencl/tools/cltrace/cltrace.cpp
+2 -2
مشاهده پرونده
@@ -1033,8 +1033,8 @@ getStringString(const char* src)
std::string str(src);
if (str.length() > 60) {
str = str.substr(0, 60).append("...");
if (str.length() > 6000) {
str = str.substr(0, 6000).append("...");
}
size_t found = 0;
projects/clr/rocclr/device/devhostcall.cpp
+2
مشاهده پرونده
@@ -94,6 +94,7 @@ static void handlePayload(MessageHandler& messages, uint32_t service, uint64_t*
if (payload[0]) {
amd::Memory* mem = amd::MemObjMap::FindMemObj(reinterpret_cast<void*>(payload[0]));
if (mem) {
const_cast<amd::Device*>(&dev)->RemoveHostcallMemory(mem);
amd::MemObjMap::RemoveMemObj(reinterpret_cast<void*>(payload[0]));
mem->release();
} else {
@@ -110,6 +111,7 @@ static void handlePayload(MessageHandler& messages, uint32_t service, uint64_t*
device::Memory* dm = buf->getDeviceMemory(dev);
va = dm->virtualAddress();
amd::MemObjMap::AddMemObj(reinterpret_cast<void*>(va), buf);
const_cast<amd::Device*>(&dev)->TrackHostcallMemory(buf);
} else {
buf->release();
}
projects/clr/rocclr/device/device.cpp
+82
مشاهده پرونده
@@ -81,6 +81,9 @@ bool VirtualDevice::ActiveWait() const {
return device_().ActiveWait();
}
#if defined(USE_COMGR_LIBRARY)
extern amd_comgr_status_t getMetaBuf(const amd_comgr_metadata_node_t meta, std::string* str);
#endif
}
static_assert(static_cast<uint32_t>(device::Memory::MemAccess::kMemAccessNone)
@@ -328,6 +331,61 @@ const Isa* Isa::end() {
return supportedIsas().second;
}
#if defined(USE_COMGR_LIBRARY)
void Isa::setAvailableSgprVgprCached() const {
std::call_once(setSgprVgprFlag, [this]() {
std::string buf;
amd_comgr_metadata_node_t isaMeta;
amd_comgr_metadata_node_t sgprMeta;
amd_comgr_metadata_node_t vgprMeta;
bool hasIsaMeta = false;
bool hasSgprMeta = false;
bool hasVgprMeta = false;
amd_comgr_status_t status = amd::Comgr::get_isa_metadata(isaName().c_str(), &isaMeta);
if (status == AMD_COMGR_STATUS_SUCCESS) {
hasIsaMeta = true;
status = amd::Comgr::metadata_lookup(isaMeta, "AddressableNumSGPRs", &sgprMeta);
}
if (status == AMD_COMGR_STATUS_SUCCESS) {
hasSgprMeta = true;
status = amd::device::getMetaBuf(sgprMeta, &buf);
}
sgprPerWavefront_ = (status == AMD_COMGR_STATUS_SUCCESS) ? atoi(buf.c_str()) : 0;
if (status == AMD_COMGR_STATUS_SUCCESS) {
status = amd::Comgr::metadata_lookup(isaMeta, "AddressableNumVGPRs", &vgprMeta);
}
if (status == AMD_COMGR_STATUS_SUCCESS) {
hasVgprMeta = true;
status = amd::device::getMetaBuf(vgprMeta, &buf);
}
vgprPerWavefront_ = (status == AMD_COMGR_STATUS_SUCCESS) ? atoi(buf.c_str()) : 0;
if (hasVgprMeta) {
amd::Comgr::destroy_metadata(vgprMeta);
}
if (hasSgprMeta) {
amd::Comgr::destroy_metadata(sgprMeta);
}
if (hasIsaMeta) {
amd::Comgr::destroy_metadata(isaMeta);
}
if (status != AMD_COMGR_STATUS_SUCCESS) {
DevLogPrintfError("Failed to set SGPR/VGPR for ISA: %s", isaName().c_str());
}
});
}
#endif
std::vector<Device*>* Device::devices_ = nullptr;
AppProfile Device::appProfile_;
@@ -749,6 +807,16 @@ Device::~Device() {
delete vaCacheMap_;
}
for (auto memory : hostcall_allocated_memories_) {
if (memory != nullptr) {
amd::MemObjMap::RemoveMemObj(
reinterpret_cast<void*>(memory->getDeviceMemory(*this, false)->virtualAddress()));
memory->release();
}
}
hostcall_allocated_memories_.clear();
delete vaCacheAccess_;
delete settings_;
delete[] info_.extensions_;
@@ -1147,6 +1215,20 @@ bool Device::GetHandleForAddressRange(void* dev_ptr, size_t size, void* handle)
return dev_mem->GetFDHandleForMem(dev_ptr, size, VmmPtr, handle);
}
// ================================================================================================
void Device::TrackHostcallMemory(amd::Memory* memory) {
hostcall_allocated_memories_.push_back(memory);
}
// ================================================================================================
void Device::RemoveHostcallMemory(amd::Memory* memory) {
auto it =
std::find(hostcall_allocated_memories_.begin(), hostcall_allocated_memories_.end(), memory);
if (it != hostcall_allocated_memories_.end()) {
hostcall_allocated_memories_.erase(it);
}
}
} // namespace amd
namespace amd::device {
projects/clr/rocclr/device/device.hpp
+50 -8
مشاهده پرونده
@@ -1383,6 +1383,14 @@ class VirtualDevice : public amd::HeapObject {
namespace amd {
/*! IHIP IPC MEMORY Structure */
#define AMD_IPC_MEM_HANDLE_SIZE 32
typedef enum SyncPolicy {
Auto = 1,
Spin = 2,
Yield = 3,
Blocking = 4
} SyncPolicy;
//! MemoryObject map lookup class
class MemObjMap : public AllStatic {
public:
@@ -1565,6 +1573,20 @@ class Isa {
return localMemBanks_;
}
#if defined(USE_COMGR_LIBRARY)
/// @returns This Isa's available sgprs per wavefront
size_t sgprPerWavefront() const {
setAvailableSgprVgprCached();
return sgprPerWavefront_;
}
/// @returns This Isa's available vgprs per wavefront
size_t vgprPerWavefront() const {
setAvailableSgprVgprCached();
return vgprPerWavefront_;
}
#endif
/// @returns True if @p codeObjectIsa and @p agentIsa are compatible,
/// false otherwise.
static bool isCompatible(const Isa &codeObjectIsa, const Isa &agentIsa);
@@ -1604,11 +1626,19 @@ class Isa {
simdInstructionWidth_(simdInstructionWidth),
memChannelBankWidth_(memChannelBankWidth),
localMemSizePerCU_(localMemSizePerCU),
localMemBanks_(localMemBanks) {}
localMemBanks_(localMemBanks),
sgprPerWavefront_(0),
vgprPerWavefront_(0) {}
// @brief Returns the begin and end iterators for the suppported ISAs.
static std::pair<const Isa*, const Isa*> supportedIsas();
#if defined(USE_COMGR_LIBRARY)
// @brief Populate this Isa's available sgprs/vgprs per wavefront from comgr.
// Only called once per Isa.
void setAvailableSgprVgprCached() const;
#endif
// @brief Isa's target ID name. Used for LLVM COde Object Manager
// compilations.
const char* targetId_;
@@ -1631,6 +1661,10 @@ class Isa {
uint32_t memChannelBankWidth_; //!< Memory channel bank width.
uint32_t localMemSizePerCU_; //!< Local memory size per CU.
uint32_t localMemBanks_; //!< Number of banks of local memory.
mutable size_t sgprPerWavefront_; //!< Number of sgpr per wavefront.
mutable size_t vgprPerWavefront_; //!< Number of vgpr per wavefront.
mutable std::once_flag setSgprVgprFlag; //!< Once flag for sgpr and vgpr retrieval.
}; // class Isa
/*! \addtogroup Runtime
@@ -1989,8 +2023,8 @@ class Device : public RuntimeObject {
/**
* @return True if the device successfully applied the SVM attributes in HMM for device memory
*/
virtual bool SetSvmAttributes(const void* dev_ptr, size_t count,
amd::MemoryAdvice advice, bool use_cpu = false) const {
virtual bool SetSvmAttributes(const void* dev_ptr, size_t count, amd::MemoryAdvice advice,
bool use_cpu = false, int numa_id = kDefaultNumaNode) const {
ShouldNotCallThis();
return false;
}
@@ -2024,9 +2058,9 @@ class Device : public RuntimeObject {
// Returns the status of HW event, associated with amd::Event
virtual bool IsHwEventReady(
const amd::Event& event, //!< AMD event for HW status validation
bool wait = false, //!< If true then forces the event completion
uint32_t hip_event_flags = 0 //!< flags associated with the event. 0 = hipEventDefault
const amd::Event& event, //!< AMD event for HW status validation
bool wait = false, //!< If true then forces the event completion
amd::SyncPolicy policy = amd::SyncPolicy::Auto
) const {
return false;
};
@@ -2200,7 +2234,12 @@ class Device : public RuntimeObject {
bool GetHandleForAddressRange(void* dev_ptr, size_t size, void* handle);
protected:
// Registers a memory object allocated via hostcall for later cleanup.
void TrackHostcallMemory(amd::Memory* memory);
// Removes a memory object from hostcall tracking.
void RemoveHostcallMemory(amd::Memory* memory);
//! Enable the specified extension
char* getExtensionString();
@@ -2247,7 +2286,10 @@ class Device : public RuntimeObject {
static amd::Monitor lockP2P_;
Monitor* vaCacheAccess_; //!< Lock to serialize VA caching access
std::map<uintptr_t, device::Memory*>* vaCacheMap_; //!< VA cache map
uint32_t index_; //!< Unique device index
uint32_t index_; //!< Unique device index
static constexpr int kDefaultNumaNode = -1; //! Default NUMA node value for SVM operations
// Tracks all amd::Memory objects allocated via hostcall for this device.
std::vector<amd::Memory*> hostcall_allocated_memories_;
};
/*! @}
projects/clr/rocclr/device/devkernel.cpp
-50
مشاهده پرونده
@@ -1155,56 +1155,6 @@ bool Kernel::GetAttrCodePropMetadata() {
return true;
}
bool Kernel::SetAvailableSgprVgpr() {
std::string buf;
amd_comgr_metadata_node_t isaMeta;
amd_comgr_metadata_node_t sgprMeta;
amd_comgr_metadata_node_t vgprMeta;
bool hasIsaMeta = false;
bool hasSgprMeta = false;
bool hasVgprMeta = false;
amd_comgr_status_t status = amd::Comgr::get_isa_metadata(
prog().device().isa().isaName().c_str(), &isaMeta);
if (status == AMD_COMGR_STATUS_SUCCESS) {
hasIsaMeta = true;
status = amd::Comgr::metadata_lookup(isaMeta, "AddressableNumSGPRs", &sgprMeta);
}
if (status == AMD_COMGR_STATUS_SUCCESS) {
hasSgprMeta = true;
status = getMetaBuf(sgprMeta, &buf);
}
workGroupInfo_.availableSGPRs_ = (status == AMD_COMGR_STATUS_SUCCESS) ? atoi(buf.c_str()) : 0;
if (status == AMD_COMGR_STATUS_SUCCESS) {
status = amd::Comgr::metadata_lookup(isaMeta, "AddressableNumVGPRs", &vgprMeta);
}
if (status == AMD_COMGR_STATUS_SUCCESS) {
hasVgprMeta = true;
status = getMetaBuf(vgprMeta, &buf);
}
workGroupInfo_.availableVGPRs_ = (status == AMD_COMGR_STATUS_SUCCESS) ? atoi(buf.c_str()) : 0;
if (hasVgprMeta) {
amd::Comgr::destroy_metadata(vgprMeta);
}
if (hasSgprMeta) {
amd::Comgr::destroy_metadata(sgprMeta);
}
if (hasIsaMeta) {
amd::Comgr::destroy_metadata(isaMeta);
}
return (status == AMD_COMGR_STATUS_SUCCESS);
}
bool Kernel::GetPrintfStr(std::vector<std::string>* printfStr) {
const amd_comgr_metadata_node_t programMD = prog().metadata();
amd_comgr_metadata_node_t printfMeta;
projects/clr/rocclr/device/devkernel.hpp
-3
مشاهده پرونده
@@ -367,9 +367,6 @@ class Kernel : public amd::HeapObject {
//! Retrieve kernel attribute and code properties metadata
bool GetAttrCodePropMetadata();
//! Retrieve the available SGPRs and VGPRs
bool SetAvailableSgprVgpr();
//! Retrieve the printf string metadata
bool GetPrintfStr(std::vector<std::string>* printfStr);
projects/clr/rocclr/device/rocm/mesa_glinterop.h
+3 -1
مشاهده پرونده
@@ -239,7 +239,7 @@ struct mesa_glinterop_export_out {
/* Structure version 1 ends here. */
};
#if IS_LINUX
/**
* Query device information.
*
@@ -301,6 +301,8 @@ typedef int (PFNMESAGLINTEROPEGLEXPORTOBJECTPROC)(EGLDisplay dpy, EGLContext con
struct mesa_glinterop_export_in *in,
struct mesa_glinterop_export_out *out);
#endif // IS_LINUX
#ifdef __cplusplus
}
#endif
projects/clr/rocclr/device/rocm/rocblit.cpp
+1 -1
مشاهده پرونده
@@ -2750,7 +2750,7 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM,
sp->child_queue = reinterpret_cast<uint64_t>(schedulerQueue);
sp->complete_signal = gpu().Barriers().ActiveSignal(kInitSignalValueOne, nullptr);
sp->vqueue_header = vqVM;
sp->parentAQL = reinterpret_cast<uint64_t>(aql_wrap);
sp->parentAQL = aql_wrap;
if (dev().info().maxEngineClockFrequency_ > 0) {
sp->eng_clk = (1000 * 1024) / dev().info().maxEngineClockFrequency_;
projects/clr/rocclr/device/rocm/rocdevice.cpp
+21 -28
مشاهده پرونده
@@ -56,7 +56,6 @@
#include <iostream>
#include <iomanip>
#include <memory>
#include <sys/resource.h>
#ifdef ROCCLR_SUPPORT_NUMA_POLICY
#include <numa.h>
#include <numaif.h>
@@ -1017,7 +1016,7 @@ bool Sampler::create(const amd::Sampler& owner) {
return false;
}
hwSrd_ = reinterpret_cast<uint64_t>(hsa_sampler.handle);
hwSrd_ = hsa_sampler.handle;
hwState_ = reinterpret_cast<address>(hsa_sampler.handle);
return true;
@@ -1272,8 +1271,7 @@ bool Device::populateOCLDeviceConstants() {
assert(alloc_granularity_ > 0);
} else {
// We suppose half of physical memory can be used by GPU in APU system
info_.globalMemSize_ =
uint64_t(sysconf(_SC_PAGESIZE)) * uint64_t(sysconf(_SC_PHYS_PAGES)) / 2;
info_.globalMemSize_ = amd::Os::hostTotalPhysicalMemory() / 2;
info_.globalMemSize_ = std::max(info_.globalMemSize_, uint64_t(1 * Gi));
info_.globalMemSize_ = (static_cast<uint64_t>(std::min(GPU_MAX_HEAP_SIZE, 100u)) *
static_cast<uint64_t>(info_.globalMemSize_)) / 100u;
@@ -2288,22 +2286,22 @@ void Device::updateFreeMemory(size_t size, bool free) {
void* Device::svmAlloc(amd::Context& context, size_t size, size_t alignment, cl_svm_mem_flags flags,
void* svmPtr) const {
amd::Memory* mem = nullptr;
void* svmPtrUsed = reinterpret_cast<void*>(amd::Memory::MemoryType::kSvmMemoryPtr);
void* svmPtrUsed = reinterpret_cast<void*>(amd::Memory::MemoryType::kSvmMemoryPtr);
if (nullptr != svmPtr) {
// Find the existing amd::mem object
mem = amd::MemObjMap::FindMemObj(svmPtr);
if (mem != nullptr) {
return mem->getSvmPtr();
}
}
if (flags & CL_MEM_USE_HOST_PTR ) {
svmPtrUsed = svmPtr;
} else {
DevLogPrintfError("Cannot find svm_ptr: 0x%x \n", svmPtr);
return nullptr;
}
}
}
// create a hidden buffer, which will allocated on the device later
mem = new (context) amd::Buffer(context, flags, size, svmPtrUsed);
if (mem == nullptr) {
@@ -2468,8 +2466,8 @@ amd::Memory* Device::ImportShareableVMMHandle(void* osHandle) {
}
// ================================================================================================
bool Device::SetSvmAttributesInt(const void* dev_ptr, size_t count,
amd::MemoryAdvice advice, bool first_alloc, bool use_cpu) const {
bool Device::SetSvmAttributesInt(const void* dev_ptr, size_t count, amd::MemoryAdvice advice,
bool first_alloc, bool use_cpu, int numa_id) const {
if ((settings().hmmFlags_ & Settings::Hmm::EnableSvmTracking) && !first_alloc) {
amd::Memory* svm_mem = amd::MemObjMap::FindMemObj(dev_ptr);
if ((nullptr == svm_mem) || ((svm_mem->getMemFlags() & CL_MEM_ALLOC_HOST_PTR) == 0) ||
@@ -2492,7 +2490,7 @@ bool Device::SetSvmAttributesInt(const void* dev_ptr, size_t count,
break;
case amd::MemoryAdvice::SetPreferredLocation:
if (use_cpu) {
attr.push_back({HSA_AMD_SVM_ATTRIB_PREFERRED_LOCATION, getCpuAgent().handle});
attr.push_back({HSA_AMD_SVM_ATTRIB_PREFERRED_LOCATION, getCpuAgent(numa_id).handle});
} else {
attr.push_back({HSA_AMD_SVM_ATTRIB_PREFERRED_LOCATION, getBackendDevice().handle});
}
@@ -2552,7 +2550,7 @@ bool Device::SetSvmAttributesInt(const void* dev_ptr, size_t count,
// ================================================================================================
bool Device::SetSvmAttributes(const void* dev_ptr, size_t count,
amd::MemoryAdvice advice, bool use_cpu) const {
amd::MemoryAdvice advice, bool use_cpu, int numa_id) const {
constexpr bool kFirstAlloc = false;
return SetSvmAttributesInt(dev_ptr, count, advice, kFirstAlloc, use_cpu);
}
@@ -2840,7 +2838,7 @@ bool Device::SetClockMode(const cl_set_device_clock_mode_input_amd setClockModeI
}
// ================================================================================================
bool Device::IsHwEventReady(const amd::Event& event, bool wait, uint32_t hip_event_flags) const {
bool Device::IsHwEventReady(const amd::Event& event, bool wait, amd::SyncPolicy policy) const {
void* hw_event =
(event.NotifyEvent() != nullptr) ? event.NotifyEvent()->HwEvent() : event.HwEvent();
if (hw_event == nullptr) {
@@ -2851,8 +2849,10 @@ bool Device::IsHwEventReady(const amd::Event& event, bool wait, uint32_t hip_eve
// when set the CPU gives up host thread for other work
// when not set the CPU enters a busy-wait on the event to occur
constexpr int kHipEventBlockingSync = 0x1;
bool active_wait = !(hip_event_flags & kHipEventBlockingSync) && ActiveWait();
return WaitForSignal(reinterpret_cast<ProfilingSignal*>(hw_event)->signal_, active_wait);
bool active_wait = !((policy == amd::SyncPolicy::Blocking) & kHipEventBlockingSync) &&
ActiveWait();
bool yield = (policy == amd::SyncPolicy::Yield);
return WaitForSignal(reinterpret_cast<ProfilingSignal*>(hw_event)->signal_, active_wait, yield);
}
auto signal = reinterpret_cast<ProfilingSignal*>(hw_event)->signal_;
@@ -3159,19 +3159,16 @@ void Device::releaseQueue(hsa_queue_t* queue, const std::vector<uint32_t>& cuMas
void* Device::getOrCreateHostcallBuffer(hsa_queue_t* queue, bool coop_queue,
const std::vector<uint32_t>& cuMask) {
decltype(queuePool_)::value_type::iterator qIter;
bool found = false;
if (!coop_queue) {
for (auto &it : cuMask.size() == 0 ? queuePool_ : queueWithCUMaskPool_) {
qIter = it.find(queue);
if (qIter != it.end()) {
found = true;
break;
}
}
if (cuMask.size() == 0) {
assert(qIter != queuePool_[QueuePriority::High].end());
} else {
assert(qIter != queueWithCUMaskPool_[QueuePriority::High].end());
}
assert(found && "Couldn't find queue");
if (qIter->second.hostcallBuffer_) {
return qIter->second.hostcallBuffer_;
@@ -3406,9 +3403,7 @@ hsa_status_t Device::BackendErrorCallBackHandler(const hsa_amd_event_t* event, v
}
// Execute the default handler if a GPU core file should be generated ...
struct rlimit rlimit;
if ((getrlimit(RLIMIT_CORE, &rlimit) == 0 && rlimit.rlim_cur != 0) ||
!HIP_SKIP_ABORT_ON_GPU_ERROR) {
if (amd::Os::DumpCoreFile() || !HIP_SKIP_ABORT_ON_GPU_ERROR) {
return HSA_STATUS_ERROR;
}
@@ -3654,9 +3649,7 @@ void callbackQueue(hsa_status_t status, hsa_queue_t* queue, void* data) {
errorMsg, status);
}
struct rlimit rlimit;
if ((getrlimit(RLIMIT_CORE, &rlimit) == 0 && rlimit.rlim_cur != 0) ||
!HIP_SKIP_ABORT_ON_GPU_ERROR) {
if (amd::Os::DumpCoreFile() || !HIP_SKIP_ABORT_ON_GPU_ERROR) {
abort();
}
amd::Device::gpu_error_ = ConvertHSAErrorIntoCLError(status);
projects/clr/rocclr/device/rocm/rocdevice.hpp
+19 -5
مشاهده پرونده
@@ -291,7 +291,8 @@ class NullDevice : public amd::Device {
}
bool IsHwEventReady(const amd::Event& event, bool wait = false,
uint32_t hip_event_flags = 0) const override {
amd::SyncPolicy policy = amd::SyncPolicy::Auto)
const override {
return false;
}
@@ -344,6 +345,15 @@ class Device : public NullDevice {
//! Get the CPU agent with the least NUMA distance to this GPU
const hsa_agent_t &getCpuAgent() const { return cpu_agent_info_->agent; }
//! Get the CPU agent that is in a 'index' NUMA node
const hsa_agent_t getCpuAgent(int index) const {
if ((index < 0) || (index >= cpu_agents_.size())) {
// Return default CPU agent
return cpu_agent_info_->agent;
}
return cpu_agents_[index].agent;
}
void setupCpuAgent(); // Setup the CPU agent which has the least NUMA distance to this GPU
void checkAtomicSupport(); //!< Check the support for pcie atomics
@@ -424,8 +434,8 @@ class Device : public NullDevice {
virtual void svmFree(void* ptr) const;
virtual bool SetSvmAttributes(const void* dev_ptr, size_t count,
amd::MemoryAdvice advice, bool use_cpu = false) const;
virtual bool SetSvmAttributes(const void* dev_ptr, size_t count, amd::MemoryAdvice advice,
bool use_cpu = false, int numa_id = kDefaultNumaNode) const;
virtual bool GetSvmAttributes(void** data, size_t* data_sizes, int* attributes,
size_t num_attributes, const void* dev_ptr, size_t count) const;
virtual size_t ScratchLimitCurrent() const final;
@@ -447,7 +457,8 @@ class Device : public NullDevice {
cl_set_device_clock_mode_output_amd* pSetClockModeOutput);
virtual bool IsHwEventReady(const amd::Event& event, bool wait = false,
uint32_t hip_event_flags = 0) const;
amd::SyncPolicy policy = amd::SyncPolicy::Auto)
const;
virtual void getHwEventTime(const amd::Event& event, uint64_t* start, uint64_t* end) const;
virtual void ReleaseGlobalSignal(void* signal) const;
virtual bool CreateUserEvent(amd::UserEvent* event) const;
@@ -576,8 +587,11 @@ class Device : public NullDevice {
//! Construct a new physical HSA device
Device(hsa_agent_t bkendDevice);
static constexpr int kDefaultNumaNode = -1;
bool SetSvmAttributesInt(const void* dev_ptr, size_t count, amd::MemoryAdvice advice,
bool first_alloc = false, bool use_cpu = false) const;
bool first_alloc = false, bool use_cpu = false,
int numa_id = kDefaultNumaNode) const;
static constexpr hsa_signal_value_t InitSignalValue = 1;
static hsa_ven_amd_loader_1_00_pfn_t amd_loader_ext_table;
projects/clr/rocclr/device/rocm/rocglinterop.hpp
+6
مشاهده پرونده
@@ -26,8 +26,14 @@
#include <GL/glx.h>
#include <EGL/egl.h>
#else
#include <windows.h>
#include <GL/gl.h>
#include <GL/glext.h>
#include <EGL/egl.h>
#ifndef GLX_H
struct _XDisplay;
struct __GLXcontextRec;
#endif
typedef _XDisplay Display;
typedef __GLXcontextRec* GLXContext;
#endif
projects/clr/rocclr/device/rocm/rockernel.cpp
+3 -28
مشاهده پرونده
@@ -47,11 +47,6 @@ bool Kernel::postLoad() {
workGroupInfo_.availableLDSSize_ = device().info().localMemSizePerCU_;
assert(workGroupInfo_.availableLDSSize_ > 0);
if (!SetAvailableSgprVgpr()) {
DevLogError("Cannot set available SGPR/VGPR\n");
return false;
}
// Get the kernel code handle
hsa_status_t hsaStatus;
hsa_executable_symbol_t symbol;
@@ -145,35 +140,15 @@ bool Kernel::postLoad() {
}
assert(wavefront_size > 0);
size_t const_size_bytes = 0;
hsa_executable_iterate_symbols(
program()->hsaExecutable(),
[](hsa_executable_t executable, hsa_executable_symbol_t symbol,
void *const_size_bytes) -> hsa_status_t {
bool variable_is_const = false;
hsa_status_t hsaStat = hsa_executable_symbol_get_info(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_IS_CONST, &variable_is_const);
if (hsaStat == HSA_STATUS_SUCCESS && variable_is_const) {
uint32_t variable_size = 0;
if (hsa_executable_symbol_get_info(
symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE,
&variable_size) == HSA_STATUS_SUCCESS) {
*(static_cast<size_t *>(const_size_bytes)) += variable_size;
}
}
return HSA_STATUS_SUCCESS;
},
&const_size_bytes);
workGroupInfo_.availableVGPRs_ = device().isa().vgprPerWavefront();
workGroupInfo_.availableSGPRs_ = device().isa().sgprPerWavefront();
workGroupInfo_.privateMemSize_ = workitemPrivateSegmentByteSize_;
workGroupInfo_.localMemSize_ = workgroupGroupSegmentByteSize_;
workGroupInfo_.usedLDSSize_ = workgroupGroupSegmentByteSize_;
workGroupInfo_.preferredSizeMultiple_ = wavefront_size;
workGroupInfo_.usedStackSize_ = kernelHasDynamicCallStack_;
workGroupInfo_.wavefrontPerSIMD_ = program()->rocDevice().info().maxWorkItemSizes_[0] / wavefront_size;
workGroupInfo_.constMemSize_ = const_size_bytes;
workGroupInfo_.constMemSize_ = 0;
workGroupInfo_.maxDynamicSharedSizeBytes_ = static_cast<int>(workGroupInfo_.availableLDSSize_ -
workGroupInfo_.localMemSize_);
if (workGroupInfo_.size_ == 0) {
projects/clr/rocclr/device/rocm/rockernel.hpp
+1 -5
مشاهده پرونده
@@ -21,7 +21,6 @@
#pragma once
#include <memory>
#include <cxxabi.h>
#include "rocprogram.hpp"
#include "top.hpp"
#include "rocprintf.hpp"
@@ -61,10 +60,7 @@ class Kernel : public device::Kernel {
private:
void initDemangledName() {
if (demangled_name_.empty()) {
int status = 0;
char* demangled = abi::__cxa_demangle(name().c_str(), nullptr, nullptr, &status);
demangled_name_ = (status == 0 && demangled != nullptr) ? demangled : name().c_str();
free(demangled);
amd::Os::CxaDemangle(name(), &demangled_name_);
}
}
projects/clr/rocclr/device/rocm/rocmemory.cpp
+9 -4
مشاهده پرونده
@@ -204,11 +204,17 @@ void Memory::cpuUnmap(device::VirtualDevice& vDev) {
}
// ================================================================================================
hsa_status_t Memory::interopMapBuffer(int fd) {
hsa_status_t Memory::interopMapBuffer(amd::Os::FileDesc fdn) {
hsa_agent_t agent = dev().getBackendDevice();
size_t size;
size_t metadata_size = 0;
void* metadata;
#if IS_WINDOWS
int fd = 0;
assert(!"Unimplemented");
#else
auto fd = fdn;
#endif
hsa_status_t status = hsa_amd_interop_map_buffer(
1, &agent, fd, 0, &size, &interop_deviceMemory_,
&metadata_size, (const void**)&metadata);
@@ -232,7 +238,7 @@ hsa_status_t Memory::interopMapBuffer(int fd) {
// Setup an interop buffer (dmabuf handle) as an OpenCL buffer
// ================================================================================================
bool Memory::createInteropBuffer(GLenum targetType, int miplevel) {
#if defined(_WIN32)
#if IS_WINDOWS
return false;
#else
assert(owner()->isInterop() && "Object is not an interop object.");
@@ -851,8 +857,7 @@ bool Buffer::create(bool alloc_local) {
return false;
}
deviceMemory_ = const_cast<long int*>(signalValuePtr); // conversion to void * is
// implicit
deviceMemory_ = const_cast<void*>(reinterpret_cast<volatile void*>(signalValuePtr));
// Disable host access to force blit path for memeory writes.
flags_ &= ~HostMemoryDirectAccess;
projects/clr/rocclr/device/rocm/rocmemory.hpp
+1 -1
مشاهده پرونده
@@ -128,7 +128,7 @@ class Memory : public device::Memory {
// Free / deregister device memory.
virtual void destroy() = 0;
hsa_status_t interopMapBuffer(int fd);
hsa_status_t interopMapBuffer(amd::Os::FileDesc fdn);
// Place interop object into HSA's flat address space
bool createInteropBuffer(GLenum targetType, int miplevel);
projects/clr/rocclr/device/rocm/rocprintf.hpp
-11
مشاهده پرونده
@@ -23,17 +23,6 @@
/*! \addtogroup GPU GPU Device Implementation
* @{
*/
#ifndef isinf
#ifdef _MSC_VER
#define isinf(X) (!_finite(X) && !_isnan(X))
#endif //_MSC_VER
#endif // isinf
#ifndef isnan
#ifdef _MSC_VER
#define isnan(X) (_isnan(X))
#endif //_MSC_VER
#endif // isnan
#ifndef copysign
#ifdef _MSC_VER
projects/clr/rocclr/device/rocm/rocvirtual.cpp
+26 -11
مشاهده پرونده
@@ -44,6 +44,8 @@
#include <string>
#include <thread>
#include <vector>
#include <atomic>
#include <cinttypes>
#if defined(__AVX__)
#if defined(__MINGW64__)
@@ -925,7 +927,12 @@ uint64_t VirtualGPU::getQueueID() {
// ================================================================================================
static inline void packet_store_release(uint32_t* packet, uint16_t header, uint16_t rest) {
__atomic_store_n(packet, header | (rest << 16), __ATOMIC_RELEASE);
#if IS_WINDOWS
std::atomic_ref<uint32_t> atomic_header(*packet);
atomic_header.store(header | (rest << 16), std::memory_order_release);
#else
__atomic_store_n(packet, header | (rest << 16), __ATOMIC_RELEASE);
#endif
}
// ================================================================================================
@@ -968,12 +975,12 @@ void VirtualGPU::AnalyzeAqlQueue() const {
} else {
printf("VGPU(%p) Queue(%p). Couldn't find kernel\n", this, gpu_queue_);
}
printf("VGPU=%p SWq=%p, HWq=%p, id=%ld\n\tDispatch Header = "
printf("VGPU=%p SWq=%p, HWq=%p, id=%" PRIu64 "\n\tDispatch Header ="
"0x%x (type=%d, barrier=%d, acquire=%d, release=%d), "
"setup=%d\n\tgrid=[%u, %u, %u], workgroup=[%u, %u, %u]\n\tprivate_seg_size=%u, "
"group_seg_size=%u\n\tkernel_obj=0x%lx, "
"kernarg_address=0x%p\n\tcompletion_signal=0x%lx, "
"correlation_id=%lu\n\trptr=%lu, wptr=%lu\n",
"group_seg_size=%u\n\tkernel_obj=0x%" PRIx64 ", "
"kernarg_address=0x%p\n\tcompletion_signal=0x%" PRIx64 ", "
"correlation_id=%" PRIu64 "\n\trptr=%" PRIu64 ", wptr=%" PRIu64 "\n ",
this, gpu_queue_, gpu_queue_->base_address, gpu_queue_->id, header,
extractAqlBits(header, HSA_PACKET_HEADER_TYPE, HSA_PACKET_HEADER_WIDTH_TYPE),
extractAqlBits(header, HSA_PACKET_HEADER_BARRIER, HSA_PACKET_HEADER_WIDTH_BARRIER),
@@ -987,8 +994,8 @@ void VirtualGPU::AnalyzeAqlQueue() const {
packet.kernarg_address, packet.completion_signal.handle, packet.reserved2,
read, index);
} else {
printf("VGPU(%p) Queue(%p) rptr=%lu, wptr=%lu. A barrier packet in the queue!\n",
this, gpu_queue_, read, index);
printf("VGPU(%p) Queue(%p) rptr=%" PRIu64 ", wptr=%" PRIu64
". A barrier packet in the queue!\n", this, gpu_queue_, read, index);
}
} else {
printf("VGPU(%p) Queue(%p) is idle\n", this, gpu_queue_);
@@ -1261,7 +1268,7 @@ void VirtualGPU::dispatchBarrierPacket(uint16_t packetHeader, bool skipSignal,
hsa_barrier_and_packet_t* aql_loc =
&(reinterpret_cast<hsa_barrier_and_packet_t*>(gpu_queue_->base_address))[index & queueMask];
*aql_loc = barrier_packet_;
__atomic_store_n(reinterpret_cast<uint32_t*>(aql_loc), packetHeader, __ATOMIC_RELEASE);
packet_store_release(reinterpret_cast<uint32_t*>(aql_loc), packetHeader, 0);
hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index);
ClPrint(amd::LOG_DEBUG, amd::LOG_AQL,
@@ -2109,9 +2116,10 @@ void VirtualGPU::submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) {
hsa_signal_t active = Barriers().ActiveSignal(kInitSignalValueOne, timestamp_);
// Find the requested agent for the transfer
hsa_agent_t agent = (cmd.cpu_access() ||
(dev().settings().hmmFlags_ & Settings::Hmm::EnableSystemMemory)) ?
dev().getCpuAgent() : (static_cast<const roc::Device*>(cmd.device()))->getBackendDevice();
hsa_agent_t agent =
(cmd.cpu_access() || (dev().settings().hmmFlags_ & Settings::Hmm::EnableSystemMemory))
? dev().getCpuAgent(cmd.numa_id())
: (static_cast<const roc::Device*>(cmd.device()))->getBackendDevice();
// Initiate a prefetch command
hsa_status_t status = hsa_amd_svm_prefetch_async(
@@ -3223,6 +3231,7 @@ bool VirtualGPU::createVirtualQueue(uint deviceQueueSize)
}
// ================================================================================================
#if IS_LINUX
__attribute__((optimize("unroll-all-loops"), always_inline))
static inline void nontemporalMemcpy(
void* __restrict dst, const void* __restrict src, size_t size) {
@@ -3270,6 +3279,12 @@ static inline void nontemporalMemcpy(
std::memcpy(dst, src, size);
#endif
}
#else
static inline void nontemporalMemcpy(void* __restrict dst, const void* __restrict src,
size_t size) {
std::memcpy(dst, src, size);
}
#endif
void VirtualGPU::HiddenHeapInit() { const_cast<Device&>(dev()).HiddenHeapInit(*this); }
projects/clr/rocclr/device/rocm/rocvirtual.hpp
+5 -1
مشاهده پرونده
@@ -32,6 +32,7 @@
#include "hsa/hsa_ven_amd_aqlprofile.h"
#include "rocsched.hpp"
#include "device/device.hpp"
#include "os/os.hpp"
#include <stack>
namespace amd::roc {
@@ -49,7 +50,7 @@ constexpr static uint64_t kUnlimitedWait = std::numeric_limits<uint64_t>::max();
constexpr static uint64_t kTimeout4Secs = 4 * M;
inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false) {
inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false, bool yield = false) {
hsa_wait_state_t wait_state = HSA_WAIT_STATE_BLOCKED;
if (active_wait) {
@@ -81,6 +82,9 @@ inline bool WaitForSignal(hsa_signal_t signal, bool active_wait = false) {
"(0x%lx) for %d ns", signal.handle, kTimeout4Secs);
return true;
}
if (yield && wait_state == HSA_WAIT_STATE_ACTIVE) {
amd::Os::yield();
}
}
}
projects/clr/rocclr/elf/elf.hpp
+2
مشاهده پرونده
@@ -26,6 +26,8 @@
#include "top.hpp"
#include "elfio/elfio.hpp"
#include <sstream>
using amd::ELFIO::Elf64_Ehdr;
using amd::ELFIO::Elf64_Shdr;
// Not sure where to put these in the libelf
#define AMD_BIF2 2 // AMD BIF Version 2.0
projects/clr/rocclr/os/os.cpp
+6 -5
مشاهده پرونده
@@ -37,16 +37,17 @@
namespace amd {
// ================================================================================================
bool Os::isValidFileDesc(const amd::Os::FileDesc& desc) {
#if defined(_WIN32)
return reinterpret_cast<int>(desc) > 0;
#if IS_WINDOWS
return desc != nullptr;
#else
return static_cast<int>(desc) > 0;
return desc > 0;
#endif
return false;
}
// ================================================================================================
void* Os::loadLibrary(const char* libraryname) {
void* handle;
@@ -57,7 +58,7 @@ void* Os::loadLibrary(const char* libraryname) {
namestart = (namestart != std::string::npos) ? namestart + 1 : 0;
if (namestart == 0) {
#if defined(ATI_OS_WIN)
#if IS_WINDOWS
// Try with the path of the current loaded dll(OCL runtime) first
HMODULE hm = NULL;
if (!GetModuleHandleExA(
projects/clr/rocclr/os/os.hpp
+7 -1
مشاهده پرونده
@@ -332,8 +332,14 @@ class Os : AllStatic {
//! Return the current process id
static int getProcessId();
// Prints the location of the currently loaded library (shared object or DLL)
//! Prints the location of the currently loaded library (shared object or DLL)
static void PrintLibraryLocation();
//! Checks if a core dump must be generated (rocgdb detection). Returns false in Windows
static bool DumpCoreFile();
//! Demangle a C++ name. The function will return the same name if couldn't demangle
static void CxaDemangle(const std::string& name, std::string* demangle);
};
/*@}*/
projects/clr/rocclr/os/os_posix.cpp
+35 -26
مشاهده پرونده
@@ -39,8 +39,10 @@
#include <pthread.h>
#include <dlfcn.h>
#include <signal.h>
#include <cxxabi.h>
#include <sys/prctl.h>
#include <sys/resource.h>
#include <link.h>
#include <time.h>
@@ -238,7 +240,8 @@ address Os::reserveMemory(address start, size_t size, size_t alignment, MemProt
if (size >= kLargePageSize) {
int status = madvise(aligned, size, MADV_HUGEPAGE);
if (status) {
ClPrint(amd::LOG_DEBUG, amd::LOG_CODE, "madvise with advice MADV_HUGEPAGE"
ClPrint(amd::LOG_DEBUG, amd::LOG_CODE,
"madvise with advice MADV_HUGEPAGE"
" starting at address %p and page size 0x%zx, returned %d, errno: %s",
aligned, size, status, strerror(errno));
}
@@ -335,7 +338,7 @@ void Os::setPreferredNumaNode(uint32_t node) {
numa_free_cpumask(bm);
}
#endif //ROCCLR_SUPPORT_NUMA_POLICY
#endif // ROCCLR_SUPPORT_NUMA_POLICY
}
void* Thread::entry(Thread* thread) {
@@ -736,8 +739,7 @@ void Os::getAppPathAndFileName(std::string& appName, std::string& appPathAndName
// Get filename without path and extension.
appName = std::string(basename(buff.get()));
appPathAndName = std::string(buff.get());
}
else {
} else {
appName = "";
appPathAndName = "";
}
@@ -747,9 +749,8 @@ void Os::getAppPathAndFileName(std::string& appName, std::string& appPathAndName
bool Os::GetURIFromMemory(const void* image, size_t image_size, std::string& uri) {
pid_t pid = getpid();
std::ostringstream uri_stream;
//Create a unique resource indicator to the memory address
uri_stream << "memory://" << pid
<< "#offset=0x" << std::hex << (uintptr_t)image << std::dec
// Create a unique resource indicator to the memory address
uri_stream << "memory://" << pid << "#offset=0x" << std::hex << (uintptr_t)image << std::dec
<< "&size=" << image_size;
uri = uri_stream.str();
return true;
@@ -757,7 +758,7 @@ bool Os::GetURIFromMemory(const void* image, size_t image_size, std::string& uri
bool Os::CloseFileHandle(FileDesc fdesc) {
// Return false if close system call fails
if(close(fdesc) < 0) {
if (close(fdesc) < 0) {
return false;
}
@@ -776,7 +777,7 @@ bool Os::GetFileHandle(const char* fname, FileDesc* fd_ptr, size_t* sz_ptr) {
return false;
}
//Retrieve stat info and size
// Retrieve stat info and size
if (fstat(*fd_ptr, &stat_buf) != 0) {
close(*fd_ptr);
return false;
@@ -789,7 +790,6 @@ bool Os::GetFileHandle(const char* fname, FileDesc* fd_ptr, size_t* sz_ptr) {
bool amd::Os::FindFileNameFromAddress(const void* image, std::string* fname_ptr,
size_t* foffset_ptr) {
// Get the list of mapped file list
bool ret_value = false;
std::ifstream proc_maps;
@@ -804,9 +804,7 @@ bool amd::Os::FindFileNameFromAddress(const void* image, std::string* fname_ptr,
char dash;
std::stringstream tokens(line);
uintptr_t low_address, high_address;
tokens >> std::hex >> low_address >> std::dec
>> dash
>> std::hex >> high_address >> std::dec;
tokens >> std::hex >> low_address >> std::dec >> dash >> std::hex >> high_address >> std::dec;
if (dash != '-') {
continue;
}
@@ -818,10 +816,7 @@ bool amd::Os::FindFileNameFromAddress(const void* image, std::string* fname_ptr,
std::string permissions, device, uri_file_path;
size_t offset;
uint64_t inode;
tokens >> permissions
>> std::hex >> offset >> std::dec
>> device
>> inode;
tokens >> permissions >> std::hex >> offset >> std::dec >> device >> inode;
std::getline(tokens >> std::ws, uri_file_path);
if (inode == 0 || uri_file_path.empty()) {
@@ -870,7 +865,7 @@ bool Os::MemoryMapFile(const char* fname, const void** mmap_ptr, size_t* mmap_si
struct stat stat_buf;
int fd = open(fname, O_RDONLY);
if (fd < 0 ) {
if (fd < 0) {
return false;
}
@@ -897,15 +892,15 @@ bool Os::MemoryMapFileTruncated(const char* fname, const void** mmap_ptr, size_t
}
struct stat stat_buf;
int fd = shm_open(fname, O_RDWR|O_CREAT, S_IRWXU|S_IRWXG|S_IRWXO);
if (fd < 0 ) {
int fd = shm_open(fname, O_RDWR | O_CREAT, S_IRWXU | S_IRWXG | S_IRWXO);
if (fd < 0) {
return false;
}
if (ftruncate(fd, mmap_size) != 0) {
return false;
}
*mmap_ptr = mmap(NULL, mmap_size, PROT_READ|PROT_WRITE, MAP_SHARED, fd, 0);
*mmap_ptr = mmap(NULL, mmap_size, PROT_READ | PROT_WRITE, MAP_SHARED, fd, 0);
close(fd);
@@ -916,13 +911,11 @@ bool Os::MemoryMapFileTruncated(const char* fname, const void** mmap_ptr, size_t
return true;
}
int Os::getProcessId() {
return ::getpid();
}
int Os::getProcessId() { return ::getpid(); }
// ================================================================================================
void* Os::CreateIpcMemory(const char* fname, size_t size, FileDesc* desc) {
*desc = shm_open(fname, O_RDWR | O_CREAT, S_IRWXU|S_IRWXG|S_IRWXO);
*desc = shm_open(fname, O_RDWR | O_CREAT, S_IRWXU | S_IRWXG | S_IRWXO);
if (*desc < 0) {
return nullptr;
}
@@ -940,7 +933,7 @@ void* Os::CreateIpcMemory(const char* fname, size_t size, FileDesc* desc) {
void* Os::OpenIpcMemory(const char* fname, const FileDesc desc, size_t size) {
FileDesc handle = desc;
if (fname != nullptr) {
handle = shm_open(fname, O_RDWR, S_IRWXU|S_IRWXG|S_IRWXO);
handle = shm_open(fname, O_RDWR, S_IRWXU | S_IRWXG | S_IRWXO);
}
if (handle < 0) {
@@ -961,6 +954,7 @@ void Os::CloseIpcMemory(const FileDesc desc, const void* ptr, size_t size) {
}
}
// ================================================================================================
void Os::PrintLibraryLocation() {
Dl_info dl_info;
if (dladdr(reinterpret_cast<void*>(Os::loadLibrary), &dl_info) && dl_info.dli_fname) {
@@ -970,6 +964,21 @@ void Os::PrintLibraryLocation() {
}
}
// ================================================================================================
bool Os::DumpCoreFile() {
// Execute the default handler if a GPU core file should be generated ...
struct rlimit rlimit;
return (getrlimit(RLIMIT_CORE, &rlimit) == 0 && rlimit.rlim_cur != 0);
}
// ================================================================================================
void Os::CxaDemangle(const std::string& name, std::string* result) {
int status = 0;
char* demangled = abi::__cxa_demangle(name.c_str(), nullptr, nullptr, &status);
*result = (status == 0 && demangled != nullptr) ? demangled : name;
free(demangled);
}
} // namespace amd
#endif // !defined(_WIN32) && !defined(__CYGWIN__)
projects/clr/rocclr/os/os_win32.cpp
+7
مشاهده پرونده
@@ -735,6 +735,7 @@ void Os::CloseIpcMemory(const FileDesc desc, const void* ptr, size_t size) {
}
}
// ================================================================================================
void Os::PrintLibraryLocation() {
HMODULE hm = NULL;
if (GetModuleHandleExA(
@@ -749,6 +750,12 @@ void Os::PrintLibraryLocation() {
ClPrint(amd::LOG_INFO, amd::LOG_INIT, "HIP Library Path: <unknown>");
}
// ================================================================================================
bool Os::DumpCoreFile() { return false; }
// ================================================================================================
void Os::CxaDemangle(const std::string& name, std::string* result) { *result = name; }
} // namespace amd
#endif // _WIN32 || __CYGWIN__
projects/clr/rocclr/platform/command.cpp
+15 -14
مشاهده پرونده
@@ -421,20 +421,21 @@ const Context& Command::context() const { return queue_->context(); }
NDRangeKernelCommand::NDRangeKernelCommand(HostQueue& queue, const EventWaitList& eventWaitList,
Kernel& kernel, const NDRangeContainer& sizes,
uint32_t sharedMemBytes, uint32_t extraParam,
uint32_t gridId, uint32_t numGrids,
uint64_t prevGridSum, uint64_t allGridSum,
uint32_t firstDevice, bool forceProfiling) :
Command(queue, CL_COMMAND_NDRANGE_KERNEL, eventWaitList, AMD_SERIALIZE_KERNEL |
(HIP_LAUNCH_BLOCKING << 1)),
kernel_(kernel),
sizes_(sizes),
sharedMemBytes_(sharedMemBytes),
extraParam_(extraParam),
gridId_(gridId),
numGrids_(numGrids),
prevGridSum_(prevGridSum),
allGridSum_(allGridSum),
firstDevice_(firstDevice) {
uint32_t gridId, uint32_t numGrids, uint64_t prevGridSum,
uint64_t allGridSum, uint32_t firstDevice,
bool forceProfiling)
: Command(queue, CL_COMMAND_NDRANGE_KERNEL, eventWaitList,
AMD_SERIALIZE_KERNEL | (HIP_LAUNCH_BLOCKING << 1)),
kernel_(kernel),
sizes_(sizes),
sharedMemBytes_(sharedMemBytes),
extraParam_(extraParam),
gridId_(gridId),
numGrids_(numGrids),
prevGridSum_(prevGridSum),
allGridSum_(allGridSum),
firstDevice_(firstDevice),
parameters_(nullptr) {
auto& device = queue.device();
auto devKernel = const_cast<device::Kernel*>(kernel.getDeviceKernel(device));
if (cooperativeGroups()) {
projects/clr/rocclr/platform/command.hpp
+5 -2
مشاهده پرونده
@@ -1856,13 +1856,15 @@ class SvmPrefetchAsyncCommand : public Command {
const void* dev_ptr_; //!< Device pointer to memory for prefetch
size_t count_; //!< the size for prefetch
bool cpu_access_; //!< Prefetch data into CPU location
int numa_id_; //!< Host NUMA node id
amd::Device* dev_; //!< Destination device to prefetch to
public:
SvmPrefetchAsyncCommand(HostQueue& queue, const EventWaitList& eventWaitList,
const void* dev_ptr, size_t count, amd::Device* dev, bool cpu_access)
const void* dev_ptr, size_t count, amd::Device* dev,
bool cpu_access, int numa_id)
: Command(queue, 1, eventWaitList), dev_ptr_(dev_ptr), count_(count),
cpu_access_(cpu_access), dev_(dev) {}
cpu_access_(cpu_access), dev_(dev), numa_id_(numa_id) {}
virtual void submit(device::VirtualDevice& device) { device.submitSvmPrefetchAsync(*this); }
@@ -1872,6 +1874,7 @@ class SvmPrefetchAsyncCommand : public Command {
size_t count() const { return count_; }
amd::Device* device() const { return dev_; }
size_t cpu_access() const { return cpu_access_; }
int numa_id() const { return numa_id_; }
};
/*! \brief A virtual map memory command.
projects/clr/rocclr/platform/commandqueue.cpp
+13 -10
مشاهده پرونده
@@ -41,7 +41,8 @@ HostQueue::HostQueue(Context& context, Device& device, cl_command_queue_properti
lastEnqueueCommand_(nullptr),
head_(nullptr),
tail_(nullptr),
isActive_(false) {
isActive_(false),
sync_policy_(amd::SyncPolicy::Auto) {
if (GPU_FORCE_QUEUE_PROFILING) {
properties().set(CL_QUEUE_PROFILING_ENABLE);
}
@@ -72,20 +73,21 @@ bool HostQueue::terminate() {
if (lastCommand != nullptr) {
// Check if CPU batch wasn't flushed for completion with the last command
if (GetSubmissionBatch() != nullptr) {
auto command = new Marker(*this, false);
if (command != nullptr) {
ClPrint(LOG_DEBUG, LOG_CMD, "Marker queued to ensure finish");
command->enqueue();
lastCommand = command;
}
auto command = new Marker(*this, false);
if (command != nullptr) {
ClPrint(LOG_DEBUG, LOG_CMD, "Marker queued to ensure finish");
command->enqueue();
lastCommand->release();
lastCommand = command;
}
}
if (device_.gpu_error_ == CL_SUCCESS) {
lastCommand->awaitCompletion();
lastCommand->awaitCompletion();
}
// Note that if lastCommand isn't a marker, it may not be lastEnqueueCommand_ now
// after lastCommand->awaitCompletion() is called.
if (lastEnqueueCommand_ != nullptr) {
lastEnqueueCommand_ ->release(); // lastEnqueueCommand_ should be a marker
lastEnqueueCommand_->release(); // lastEnqueueCommand_ should be a marker
lastEnqueueCommand_ = nullptr;
}
lastCommand->release();
@@ -197,9 +199,10 @@ void HostQueue::finish(bool cpu_wait) {
}
command->enqueue();
}
// Check HW status of the ROCcrl event. Note: not all ROCclr modes support HW status
static constexpr bool kWaitCompletion = true;
if (cpu_wait || !device().IsHwEventReady(command->event(), kWaitCompletion)) {
if (cpu_wait || !device().IsHwEventReady(command->event(), kWaitCompletion, GetSyncPolicy())) {
ClPrint(LOG_DEBUG, LOG_CMD,
"No HW event or batch size is less than %zu, "
"await command completion",
projects/clr/rocclr/platform/commandqueue.hpp
+9
مشاهده پرونده
@@ -307,6 +307,13 @@ class HostQueue : public CommandQueue {
return thread_.vdev()->getQueueID();
}
//! Returns Synchronization Policy for the current stream
amd::SyncPolicy GetSyncPolicy() const { return sync_policy_; }
//! Set Synchronization Policy used by Queue
void SetSyncPolicy(amd::SyncPolicy value) {
sync_policy_ = value;
}
private:
Command* head_; //!< Head of the batch list
Command* tail_; //!< Tail of the batch list
@@ -315,6 +322,8 @@ private:
//! True if this command queue is active
bool isActive_;
bool forceDestroy_ = false; //!< Destroy the queue in the current state
amd::SyncPolicy sync_policy_; //!< Used for controlling stream synchronization
};
class DeviceQueue : public CommandQueue {
projects/clr/rocclr/utils/flags.cpp
+5 -1
مشاهده پرونده
@@ -151,8 +151,12 @@ bool Flag::init() {
if (!flagIsDefault(AMD_LOG_LEVEL)) {
if (!flagIsDefault(AMD_LOG_LEVEL_FILE)) {
std::string fileName = AMD_LOG_LEVEL_FILE;
fileName = fileName + "_" + std::to_string(Os::getProcessId());
std::string pid = std::to_string(Os::getProcessId());
fileName = fileName + "_" + pid;
outFile = fopen(fileName.c_str(), "a");
if (outFile == NULL) {
outFile = fopen(("clr_logs_" + pid + ".txt").c_str(), "a");
}
}
}