include/hip/hip_runtime_api.h Source File

include/hip/hip_runtime_api.h Source File#

HIP Runtime API Reference: include/hip/hip_runtime_api.h Source File
/*
Copyright (c) 2015 - 2023 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.
*/
 
#ifndef HIP_INCLUDE_HIP_HIP_RUNTIME_API_H
#define HIP_INCLUDE_HIP_HIP_RUNTIME_API_H
 
#if __cplusplus
#include <climits>
#include <cstdint>
#include <cstdlib>
#else
#include <limits.h>
#include <stdint.h>
#include <stdlib.h>
#endif
 
#include <hip/hip_version.h>
#include <hip/hip_common.h>
#include <hip/linker_types.h>
 
enum {
    HIP_SUCCESS = 0,
    HIP_ERROR_INVALID_VALUE,
    HIP_ERROR_NOT_INITIALIZED,
    HIP_ERROR_LAUNCH_OUT_OF_RESOURCES
};
// hack to get these to show up in Doxygen:
typedef struct {
    // 32-bit Atomics
    unsigned hasGlobalInt32Atomics : 1;     
    unsigned hasGlobalFloatAtomicExch : 1;  
    unsigned hasSharedInt32Atomics : 1;     
    unsigned hasSharedFloatAtomicExch : 1;  
    unsigned hasFloatAtomicAdd : 1;  
 
    // 64-bit Atomics
    unsigned hasGlobalInt64Atomics : 1;  
    unsigned hasSharedInt64Atomics : 1;  
 
    // Doubles
    unsigned hasDoubles : 1;  
 
    // Warp cross-lane operations
    unsigned hasWarpVote : 1;     
    unsigned hasWarpBallot : 1;   
    unsigned hasWarpShuffle : 1;  
    unsigned hasFunnelShift : 1;  
 
    // Sync
    unsigned hasThreadFenceSystem : 1;  
    unsigned hasSyncThreadsExt : 1;     
 
    // Misc
    unsigned hasSurfaceFuncs : 1;        
    unsigned has3dGrid : 1;              
    unsigned hasDynamicParallelism : 1;  
} hipDeviceArch_t;
 
typedef struct hipUUID_t {
    char bytes[16];
} hipUUID;
 
//---
// Common headers for both NVCC and HIP-Clang paths:
 
#define hipGetDeviceProperties hipGetDevicePropertiesR0600
#define hipDeviceProp_t hipDeviceProp_tR0600
#define hipChooseDevice hipChooseDeviceR0600
 
typedef struct hipDeviceProp_t {
    char name[256];                   
    hipUUID uuid;                     
    char luid[8];                     
    unsigned int luidDeviceNodeMask;  
    size_t totalGlobalMem;            
    size_t sharedMemPerBlock;         
    int regsPerBlock;                 
    int warpSize;                     
    size_t memPitch;                  
    int maxThreadsPerBlock;           
    int maxThreadsDim[3];             
    int maxGridSize[3];               
    int clockRate;                    
    size_t totalConstMem;             
    int major;  
    int minor;  
    size_t textureAlignment;       
    size_t texturePitchAlignment;  
    int deviceOverlap;             
    int multiProcessorCount;       
    int kernelExecTimeoutEnabled;  
    int integrated;                
    int canMapHostMemory;          
    int computeMode;               
    int maxTexture1D;              
    int maxTexture1DMipmap;        
    int maxTexture1DLinear;        
    int maxTexture2D[2];  
    int maxTexture2DMipmap[2];  
    int maxTexture2DLinear[3];  
    int maxTexture2DGather[2];  
    int maxTexture3D[3];  
    int maxTexture3DAlt[3];           
    int maxTextureCubemap;            
    int maxTexture1DLayered[2];       
    int maxTexture2DLayered[3];       
    int maxTextureCubemapLayered[2];  
    int maxSurface1D;                 
    int maxSurface2D[2];              
    int maxSurface3D[3];              
    int maxSurface1DLayered[2];       
    int maxSurface2DLayered[3];       
    int maxSurfaceCubemap;            
    int maxSurfaceCubemapLayered[2];  
    size_t surfaceAlignment;          
    int concurrentKernels;         
    int ECCEnabled;                
    int pciBusID;                  
    int pciDeviceID;               
    int pciDomainID;               
    int tccDriver;                 
    int asyncEngineCount;          
    int unifiedAddressing;         
    int memoryClockRate;           
    int memoryBusWidth;            
    int l2CacheSize;               
    int persistingL2CacheMaxSize;  
    int maxThreadsPerMultiProcessor;    
    int streamPrioritiesSupported;      
    int globalL1CacheSupported;         
    int localL1CacheSupported;          
    size_t sharedMemPerMultiprocessor;  
    int regsPerMultiprocessor;          
    int managedMemory;         
    int isMultiGpuBoard;       
    int multiGpuBoardGroupID;  
    int hostNativeAtomicSupported;         
    int singleToDoublePrecisionPerfRatio;  
    int pageableMemoryAccess;              
    int concurrentManagedAccess;  
    int computePreemptionSupported;         
    int canUseHostPointerForRegisteredMem;  
    int cooperativeLaunch;                  
    int cooperativeMultiDeviceLaunch;       
    size_t
        sharedMemPerBlockOptin;  
    int pageableMemoryAccessUsesHostPageTables;  
    int directManagedMemAccessFromHost;  
    int maxBlocksPerMultiProcessor;      
    int accessPolicyMaxWindowSize;       
    size_t reservedSharedMemPerBlock;    
    int hostRegisterSupported;           
    int sparseHipArraySupported;         
    int hostRegisterReadOnlySupported;   
    int timelineSemaphoreInteropSupported;  
    int memoryPoolsSupported;  
    int gpuDirectRDMASupported;                    
    unsigned int gpuDirectRDMAFlushWritesOptions;  
    int gpuDirectRDMAWritesOrdering;               
    unsigned int
        memoryPoolSupportedHandleTypes;  
    int deferredMappingHipArraySupported;  
    int ipcEventSupported;                 
    int clusterLaunch;                     
    int unifiedFunctionPointers;           
    int reserved[63];                      
 
    int hipReserved[32];  
 
    /* HIP Only struct members */
    char gcnArchName[256];                    
    size_t maxSharedMemoryPerMultiProcessor;  
    int clockInstructionRate;  
    hipDeviceArch_t arch;      
    unsigned int* hdpMemFlushCntl;            
    unsigned int* hdpRegFlushCntl;            
    int cooperativeMultiDeviceUnmatchedFunc;  
    int cooperativeMultiDeviceUnmatchedGridDim;    
    int cooperativeMultiDeviceUnmatchedBlockDim;   
    int cooperativeMultiDeviceUnmatchedSharedMem;  
    int isLargeBar;                                
    int asicRevision;                              
} hipDeviceProp_t;
 
typedef enum hipMemoryType {
    hipMemoryTypeUnregistered = 0,  
    hipMemoryTypeHost         = 1,  
    hipMemoryTypeDevice       = 2,  
    hipMemoryTypeManaged      = 3,  
    hipMemoryTypeArray        = 10, 
    hipMemoryTypeUnified      = 11  
 
} hipMemoryType;
 
typedef struct hipPointerAttribute_t {
    enum hipMemoryType type;
    int device;
    void* devicePointer;
    void* hostPointer;
    int isManaged;
    unsigned allocationFlags; /* flags specified when memory was allocated*/
    /* peers? */
} hipPointerAttribute_t;
 
// Ignoring error-code return values from hip APIs is discouraged. On C++17,
// we can make that yield a warning
#if __cplusplus >= 201703L
#define __HIP_NODISCARD [[nodiscard]]
#else
#define __HIP_NODISCARD
#endif
 
// Developer note - when updating these, update the hipErrorName and hipErrorString functions in
// NVCC and HIP-Clang paths Also update the hipCUDAErrorTohipError function in NVCC path.
 
typedef enum __HIP_NODISCARD hipError_t {
    hipSuccess = 0,  
    hipErrorInvalidValue = 1,  
    hipErrorOutOfMemory = 2,   
    // Deprecated
    hipErrorMemoryAllocation = 2,  
    hipErrorNotInitialized = 3,    
    // Deprecated
    hipErrorInitializationError = 3,
    hipErrorDeinitialized = 4,      
    hipErrorProfilerDisabled = 5,
    hipErrorProfilerNotInitialized = 6,
    hipErrorProfilerAlreadyStarted = 7,
    hipErrorProfilerAlreadyStopped = 8,
    hipErrorInvalidConfiguration = 9,  
    hipErrorInvalidPitchValue = 12,   
    hipErrorInvalidSymbol = 13,   
    hipErrorInvalidDevicePointer = 17,  
    hipErrorInvalidMemcpyDirection = 21,  
    hipErrorInsufficientDriver = 35,
    hipErrorMissingConfiguration = 52,
    hipErrorPriorLaunchFailure = 53,
    hipErrorInvalidDeviceFunction = 98,  
    hipErrorNoDevice = 100,  
    hipErrorInvalidDevice = 101,  
    hipErrorInvalidImage = 200,   
    hipErrorInvalidContext = 201,  
    hipErrorContextAlreadyCurrent = 202,
    hipErrorMapFailed = 205,
    // Deprecated
    hipErrorMapBufferObjectFailed = 205,  
    hipErrorUnmapFailed = 206,
    hipErrorArrayIsMapped = 207,
    hipErrorAlreadyMapped = 208,
    hipErrorNoBinaryForGpu = 209,
    hipErrorAlreadyAcquired = 210,
    hipErrorNotMapped = 211,
    hipErrorNotMappedAsArray = 212,
    hipErrorNotMappedAsPointer = 213,
    hipErrorECCNotCorrectable = 214,
    hipErrorUnsupportedLimit = 215,   
    hipErrorContextAlreadyInUse = 216,   
    hipErrorPeerAccessUnsupported = 217,
    hipErrorInvalidKernelFile = 218,  
    hipErrorInvalidGraphicsContext = 219,
    hipErrorInvalidSource = 300,   
    hipErrorFileNotFound = 301,   
    hipErrorSharedObjectSymbolNotFound = 302,
    hipErrorSharedObjectInitFailed = 303,   
    hipErrorOperatingSystem = 304,   
    hipErrorInvalidHandle = 400,  
    // Deprecated
    hipErrorInvalidResourceHandle = 400,  
    hipErrorIllegalState = 401, 
    hipErrorNotFound = 500,   
    hipErrorNotReady = 600,  
    hipErrorIllegalAddress = 700,
    hipErrorLaunchOutOfResources = 701,  
    hipErrorLaunchTimeOut = 702,   
    hipErrorPeerAccessAlreadyEnabled = 704,  
    hipErrorPeerAccessNotEnabled = 705,  
    hipErrorSetOnActiveProcess = 708,   
    hipErrorContextIsDestroyed = 709,   
    hipErrorAssert = 710,  
    hipErrorHostMemoryAlreadyRegistered = 712,  
    hipErrorHostMemoryNotRegistered = 713,  
    hipErrorLaunchFailure = 719,  
    hipErrorCooperativeLaunchTooLarge = 720,  
    hipErrorNotSupported = 801,  
    hipErrorStreamCaptureUnsupported = 900,  
    hipErrorStreamCaptureInvalidated = 901,  
    hipErrorStreamCaptureMerge = 902,  
    hipErrorStreamCaptureUnmatched = 903,  
    hipErrorStreamCaptureUnjoined = 904,  
    hipErrorStreamCaptureIsolation = 905,  
    hipErrorStreamCaptureImplicit = 906,  
    hipErrorCapturedEvent = 907,  
    hipErrorStreamCaptureWrongThread = 908,  
    hipErrorGraphExecUpdateFailure = 910,  
    hipErrorInvalidChannelDescriptor = 911,  
    hipErrorInvalidTexture = 912,  
    hipErrorUnknown = 999,  
    // HSA Runtime Error Codes start here.
    hipErrorRuntimeMemory = 1052,  
    hipErrorRuntimeOther = 1053,  
    hipErrorTbd  
} hipError_t;
 
#undef __HIP_NODISCARD
 
typedef enum hipDeviceAttribute_t {
    hipDeviceAttributeCudaCompatibleBegin = 0,
 
    hipDeviceAttributeEccEnabled = hipDeviceAttributeCudaCompatibleBegin, 
    hipDeviceAttributeAccessPolicyMaxWindowSize,        
    hipDeviceAttributeAsyncEngineCount,                 
    hipDeviceAttributeCanMapHostMemory,                 
    hipDeviceAttributeCanUseHostPointerForRegisteredMem,
    hipDeviceAttributeClockRate,                        
    hipDeviceAttributeComputeMode,                      
    hipDeviceAttributeComputePreemptionSupported,       
    hipDeviceAttributeConcurrentKernels,                
    hipDeviceAttributeConcurrentManagedAccess,          
    hipDeviceAttributeCooperativeLaunch,                
    hipDeviceAttributeCooperativeMultiDeviceLaunch,     
    hipDeviceAttributeDeviceOverlap,                    
    hipDeviceAttributeDirectManagedMemAccessFromHost,   
    hipDeviceAttributeGlobalL1CacheSupported,           
    hipDeviceAttributeHostNativeAtomicSupported,        
    hipDeviceAttributeIntegrated,                       
    hipDeviceAttributeIsMultiGpuBoard,                  
    hipDeviceAttributeKernelExecTimeout,                
    hipDeviceAttributeL2CacheSize,                      
    hipDeviceAttributeLocalL1CacheSupported,            
    hipDeviceAttributeLuid,                             
    hipDeviceAttributeLuidDeviceNodeMask,               
    hipDeviceAttributeComputeCapabilityMajor,           
    hipDeviceAttributeManagedMemory,                    
    hipDeviceAttributeMaxBlocksPerMultiProcessor,       
    hipDeviceAttributeMaxBlockDimX,                     
    hipDeviceAttributeMaxBlockDimY,                     
    hipDeviceAttributeMaxBlockDimZ,                     
    hipDeviceAttributeMaxGridDimX,                      
    hipDeviceAttributeMaxGridDimY,                      
    hipDeviceAttributeMaxGridDimZ,                      
    hipDeviceAttributeMaxSurface1D,                     
    hipDeviceAttributeMaxSurface1DLayered,              
    hipDeviceAttributeMaxSurface2D,                     
    hipDeviceAttributeMaxSurface2DLayered,              
    hipDeviceAttributeMaxSurface3D,                     
    hipDeviceAttributeMaxSurfaceCubemap,                
    hipDeviceAttributeMaxSurfaceCubemapLayered,         
    hipDeviceAttributeMaxTexture1DWidth,                
    hipDeviceAttributeMaxTexture1DLayered,              
    hipDeviceAttributeMaxTexture1DLinear,               
    hipDeviceAttributeMaxTexture1DMipmap,               
    hipDeviceAttributeMaxTexture2DWidth,                
    hipDeviceAttributeMaxTexture2DHeight,               
    hipDeviceAttributeMaxTexture2DGather,               
    hipDeviceAttributeMaxTexture2DLayered,              
    hipDeviceAttributeMaxTexture2DLinear,               
    hipDeviceAttributeMaxTexture2DMipmap,               
    hipDeviceAttributeMaxTexture3DWidth,                
    hipDeviceAttributeMaxTexture3DHeight,               
    hipDeviceAttributeMaxTexture3DDepth,                
    hipDeviceAttributeMaxTexture3DAlt,                  
    hipDeviceAttributeMaxTextureCubemap,                
    hipDeviceAttributeMaxTextureCubemapLayered,         
    hipDeviceAttributeMaxThreadsDim,                    
    hipDeviceAttributeMaxThreadsPerBlock,               
    hipDeviceAttributeMaxThreadsPerMultiProcessor,      
    hipDeviceAttributeMaxPitch,                         
    hipDeviceAttributeMemoryBusWidth,                   
    hipDeviceAttributeMemoryClockRate,                  
    hipDeviceAttributeComputeCapabilityMinor,           
    hipDeviceAttributeMultiGpuBoardGroupID,             
    hipDeviceAttributeMultiprocessorCount,              
    hipDeviceAttributeUnused1,                          
    hipDeviceAttributePageableMemoryAccess,             
    hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 
    hipDeviceAttributePciBusId,                         
    hipDeviceAttributePciDeviceId,                      
    hipDeviceAttributePciDomainId,                      
    hipDeviceAttributePciDomainID = hipDeviceAttributePciDomainId,  
    hipDeviceAttributePersistingL2CacheMaxSize,         
    hipDeviceAttributeMaxRegistersPerBlock,             
    hipDeviceAttributeMaxRegistersPerMultiprocessor,    
    hipDeviceAttributeReservedSharedMemPerBlock,        
    hipDeviceAttributeMaxSharedMemoryPerBlock,          
    hipDeviceAttributeSharedMemPerBlockOptin,           
    hipDeviceAttributeSharedMemPerMultiprocessor,       
    hipDeviceAttributeSingleToDoublePrecisionPerfRatio, 
    hipDeviceAttributeStreamPrioritiesSupported,        
    hipDeviceAttributeSurfaceAlignment,                 
    hipDeviceAttributeTccDriver,                        
    hipDeviceAttributeTextureAlignment,                 
    hipDeviceAttributeTexturePitchAlignment,            
    hipDeviceAttributeTotalConstantMemory,              
    hipDeviceAttributeTotalGlobalMem,                   
    hipDeviceAttributeUnifiedAddressing,                
    hipDeviceAttributeUnused2,                          
    hipDeviceAttributeWarpSize,                         
    hipDeviceAttributeMemoryPoolsSupported,             
    hipDeviceAttributeVirtualMemoryManagementSupported, 
    hipDeviceAttributeHostRegisterSupported,            
    hipDeviceAttributeMemoryPoolSupportedHandleTypes,   
 
    hipDeviceAttributeCudaCompatibleEnd = 9999,
    hipDeviceAttributeAmdSpecificBegin = 10000,
 
    hipDeviceAttributeClockInstructionRate = hipDeviceAttributeAmdSpecificBegin,  
    hipDeviceAttributeUnused3,                                  
    hipDeviceAttributeMaxSharedMemoryPerMultiprocessor,         
    hipDeviceAttributeUnused4,                                  
    hipDeviceAttributeUnused5,                                  
    hipDeviceAttributeHdpMemFlushCntl,                          
    hipDeviceAttributeHdpRegFlushCntl,                          
    hipDeviceAttributeCooperativeMultiDeviceUnmatchedFunc,      
    hipDeviceAttributeCooperativeMultiDeviceUnmatchedGridDim,   
    hipDeviceAttributeCooperativeMultiDeviceUnmatchedBlockDim,  
    hipDeviceAttributeCooperativeMultiDeviceUnmatchedSharedMem, 
    hipDeviceAttributeIsLargeBar,                               
    hipDeviceAttributeAsicRevision,                             
    hipDeviceAttributeCanUseStreamWaitValue,                    
    hipDeviceAttributeImageSupport,                             
    hipDeviceAttributePhysicalMultiProcessorCount,              
    hipDeviceAttributeFineGrainSupport,                         
    hipDeviceAttributeWallClockRate,                            
    hipDeviceAttributeNumberOfXccs,                             
    hipDeviceAttributePciChipId,                                
 
    hipDeviceAttributeAmdSpecificEnd = 19999,
    hipDeviceAttributeVendorSpecificBegin = 20000,
    // Extended attributes for vendors
} hipDeviceAttribute_t;
 
typedef enum hipDriverProcAddressQueryResult {
    HIP_GET_PROC_ADDRESS_SUCCESS = 0,
    HIP_GET_PROC_ADDRESS_SYMBOL_NOT_FOUND = 1,
    HIP_GET_PROC_ADDRESS_VERSION_NOT_SUFFICIENT = 2
} hipDriverProcAddressQueryResult;
 
enum hipComputeMode {
    hipComputeModeDefault = 0,
    hipComputeModeExclusive = 1,
    hipComputeModeProhibited = 2,
    hipComputeModeExclusiveProcess = 3
};
 
enum hipFlushGPUDirectRDMAWritesOptions {
  hipFlushGPUDirectRDMAWritesOptionHost = 1 << 0,
  hipFlushGPUDirectRDMAWritesOptionMemOps = 1 << 1
};
 
enum hipGPUDirectRDMAWritesOrdering {
  hipGPUDirectRDMAWritesOrderingNone = 0,
  hipGPUDirectRDMAWritesOrderingOwner = 100,
  hipGPUDirectRDMAWritesOrderingAllDevices = 200
};
 
#if defined(__HIP_PLATFORM_AMD__) && !defined(__HIP_PLATFORM_NVIDIA__)
 
#ifndef GENERIC_GRID_LAUNCH
#define GENERIC_GRID_LAUNCH 1
#endif
#include <hip/amd_detail/host_defines.h>
#include <hip/driver_types.h>
#include <hip/texture_types.h>
#include <hip/surface_types.h>
#if defined(_MSC_VER)
#define HIP_DEPRECATED(msg) __declspec(deprecated(msg))
#else // !defined(_MSC_VER)
#define HIP_DEPRECATED(msg) __attribute__ ((deprecated(msg)))
#endif // !defined(_MSC_VER)
#define HIP_DEPRECATED_MSG "This API is marked as deprecated and might not be supported in future releases. For more details please refer https://github.com/ROCm/HIP/blob/develop/docs/reference/deprecated_api_list.md"
#define HIP_LAUNCH_PARAM_BUFFER_POINTER ((void*)0x01)
#define HIP_LAUNCH_PARAM_BUFFER_SIZE ((void*)0x02)
#define HIP_LAUNCH_PARAM_END ((void*)0x03)
#ifdef __cplusplus
  #define __dparm(x) \
          = x
#else
  #define __dparm(x)
#endif
#ifdef __GNUC__
#pragma GCC visibility push (default)
#endif
#ifdef __cplusplus
namespace hip_impl {
hipError_t hip_init();
}  // namespace hip_impl
#endif
// Structure definitions:
#ifdef __cplusplus
extern "C" {
#endif
//---
// API-visible structures
typedef struct ihipCtx_t* hipCtx_t;
// Note many APIs also use integer deviceIds as an alternative to the device pointer:
typedef int hipDevice_t;
typedef enum hipDeviceP2PAttr {
  hipDevP2PAttrPerformanceRank = 0,
  hipDevP2PAttrAccessSupported,
  hipDevP2PAttrNativeAtomicSupported,
  hipDevP2PAttrHipArrayAccessSupported
} hipDeviceP2PAttr;
typedef struct ihipStream_t* hipStream_t;
#define hipIpcMemLazyEnablePeerAccess 0x01
#define HIP_IPC_HANDLE_SIZE 64
typedef struct hipIpcMemHandle_st {
    char reserved[HIP_IPC_HANDLE_SIZE];
} hipIpcMemHandle_t;
typedef struct hipIpcEventHandle_st {
    char reserved[HIP_IPC_HANDLE_SIZE];
} hipIpcEventHandle_t;
typedef struct ihipModule_t* hipModule_t;
typedef struct ihipModuleSymbol_t* hipFunction_t;
typedef struct ihipLinkState_t* hipLinkState_t;
typedef struct ihipMemPoolHandle_t* hipMemPool_t;
 
typedef struct hipFuncAttributes {
    int binaryVersion;
    int cacheModeCA;
    size_t constSizeBytes;
    size_t localSizeBytes;
    int maxDynamicSharedSizeBytes;
    int maxThreadsPerBlock;
    int numRegs;
    int preferredShmemCarveout;
    int ptxVersion;
    size_t sharedSizeBytes;
} hipFuncAttributes;
typedef struct ihipEvent_t* hipEvent_t;
 
enum hipLimit_t {
    hipLimitStackSize = 0x0,        
    hipLimitPrintfFifoSize = 0x01,  
    hipLimitMallocHeapSize = 0x02,  
    hipExtLimitScratchMin = 0x1000, 
    hipExtLimitScratchMax = 0x1001, 
    hipExtLimitScratchCurrent = 0x1002, 
    hipLimitRange                   
};
 
//Flags that can be used with hipStreamCreateWithFlags.
#define hipStreamDefault  0x00
 
#define hipStreamNonBlocking 0x01
 
//Flags that can be used with hipEventCreateWithFlags.
#define hipEventDefault 0x0
 
#define hipEventBlockingSync 0x1
 
#define hipEventDisableTiming  0x2
 
#define hipEventInterprocess 0x4
 
//Flags that can be used with hipEventRecordWithFlags.
#define hipEventRecordDefault 0x00
 
#define hipEventRecordExternal 0x01
 
//Flags that can be used with hipStreamWaitEvent.
#define hipEventWaitDefault 0x00
 
#define hipEventWaitExternal 0x01
 
#define hipEventDisableSystemFence 0x20000000
 
#define hipEventReleaseToDevice  0x40000000
 
#define hipEventReleaseToSystem  0x80000000
 
//Flags that can be used with hipHostMalloc/hipHostAlloc.
#define hipHostAllocDefault 0x0
 
#define hipHostMallocDefault 0x0
 
#define hipHostAllocPortable 0x1
 
#define hipHostMallocPortable 0x1
 
#define hipHostAllocMapped  0x2
 
#define hipHostMallocMapped  0x2
 
#define hipHostAllocWriteCombined 0x4
 
#define hipHostMallocWriteCombined 0x4
 
#define hipHostMallocUncached  0x10000000
#define hipHostAllocUncached   hipHostMallocUncached
 
#define hipHostMallocNumaUser  0x20000000
 
#define hipHostMallocCoherent  0x40000000
 
#define hipHostMallocNonCoherent  0x80000000
 
#define hipMemAttachGlobal  0x01
 
#define hipMemAttachHost    0x02
 
#define hipMemAttachSingle  0x04
 
#define hipDeviceMallocDefault 0x0
 
#define hipDeviceMallocFinegrained 0x1
 
#define hipMallocSignalMemory 0x2
 
#define hipDeviceMallocUncached 0x3
 
#define hipDeviceMallocContiguous 0x4
 
//Flags that can be used with hipHostRegister.
#define hipHostRegisterDefault 0x0
 
#define hipHostRegisterPortable 0x1
 
#define hipHostRegisterMapped  0x2
 
#define hipHostRegisterIoMemory 0x4
 
#define hipHostRegisterReadOnly 0x08
 
#define hipExtHostRegisterCoarseGrained 0x8
 
#define hipExtHostRegisterUncached 0x80000000
 
#define hipDeviceScheduleAuto 0x0
 
#define hipDeviceScheduleSpin  0x1
 
#define hipDeviceScheduleYield  0x2
#define hipDeviceScheduleBlockingSync 0x4
#define hipDeviceScheduleMask 0x7
#define hipDeviceMapHost 0x8
#define hipDeviceLmemResizeToMax 0x10
#define hipArrayDefault 0x00
#define hipArrayLayered 0x01
#define hipArraySurfaceLoadStore 0x02
#define hipArrayCubemap 0x04
#define hipArrayTextureGather 0x08
#define hipOccupancyDefault 0x00
#define hipOccupancyDisableCachingOverride 0x01
#define hipCooperativeLaunchMultiDeviceNoPreSync 0x01
#define hipCooperativeLaunchMultiDeviceNoPostSync 0x02
#define hipCpuDeviceId ((int)-1)
#define hipInvalidDeviceId ((int)-2)
//Flags that can be used with hipExtLaunch Set of APIs.
#define hipExtAnyOrderLaunch 0x01
// Flags to be used with hipStreamWaitValue32 and hipStreamWaitValue64.
#define hipStreamWaitValueGte 0x0
#define hipStreamWaitValueEq 0x1
#define hipStreamWaitValueAnd 0x2
#define hipStreamWaitValueNor 0x3
 
typedef enum hipStreamBatchMemOpType {
  hipStreamMemOpWaitValue32 = 0x1,
  hipStreamMemOpWriteValue32 = 0x2,
  hipStreamMemOpWaitValue64 = 0x4,
  hipStreamMemOpWriteValue64 = 0x5,
  hipStreamMemOpBarrier = 0x6,            
  hipStreamMemOpFlushRemoteWrites = 0x3   
} hipStreamBatchMemOpType;
 
typedef union hipStreamBatchMemOpParams_union {
  hipStreamBatchMemOpType operation;
  struct hipStreamMemOpWaitValueParams_t{
    hipStreamBatchMemOpType operation;
    hipDeviceptr_t address;
    union {
      uint32_t value;
      uint64_t value64;
    };
    unsigned int flags;
    hipDeviceptr_t alias;   
  } waitValue;
  struct hipStreamMemOpWriteValueParams_t{
    hipStreamBatchMemOpType operation;
    hipDeviceptr_t address;
    union {
      uint32_t value;
      uint64_t value64;
    };
    unsigned int flags;
    hipDeviceptr_t alias;   
    } writeValue;
    struct hipStreamMemOpFlushRemoteWritesParams_t{
      hipStreamBatchMemOpType operation;
      unsigned int flags;
    } flushRemoteWrites;    
    struct hipStreamMemOpMemoryBarrierParams_t{
      hipStreamBatchMemOpType operation;
      unsigned int flags;
    } memoryBarrier;        
    uint64_t pad[6];
} hipStreamBatchMemOpParams;
 
typedef struct hipBatchMemOpNodeParams {
    hipCtx_t ctx;
    unsigned int count;
    hipStreamBatchMemOpParams *paramArray;
    unsigned int flags;
} hipBatchMemOpNodeParams;
 
// Stream per thread
#define hipStreamPerThread ((hipStream_t)2)
 
#define hipStreamLegacy ((hipStream_t)1)
 
// Indicates that the external memory object is a dedicated resource
#define hipExternalMemoryDedicated 0x1
typedef enum hipMemoryAdvise {
    hipMemAdviseSetReadMostly = 1,          
    hipMemAdviseUnsetReadMostly = 2,        
    hipMemAdviseSetPreferredLocation = 3,   
    hipMemAdviseUnsetPreferredLocation = 4, 
    hipMemAdviseSetAccessedBy = 5,          
    hipMemAdviseUnsetAccessedBy = 6,        
    hipMemAdviseSetCoarseGrain = 100,       
    hipMemAdviseUnsetCoarseGrain = 101      
} hipMemoryAdvise;
typedef enum hipMemRangeCoherencyMode {
    hipMemRangeCoherencyModeFineGrain = 0,      
    hipMemRangeCoherencyModeCoarseGrain = 1,    
    hipMemRangeCoherencyModeIndeterminate = 2   
} hipMemRangeCoherencyMode;
typedef enum hipMemRangeAttribute {
    hipMemRangeAttributeReadMostly = 1,         
    hipMemRangeAttributePreferredLocation = 2,  
    hipMemRangeAttributeAccessedBy = 3,         
    hipMemRangeAttributeLastPrefetchLocation = 4,
    hipMemRangeAttributeCoherencyMode = 100,    
} hipMemRangeAttribute;
 
typedef enum hipMemPoolAttr
{
    hipMemPoolReuseFollowEventDependencies   = 0x1,
    hipMemPoolReuseAllowOpportunistic        = 0x2,
    hipMemPoolReuseAllowInternalDependencies = 0x3,
    hipMemPoolAttrReleaseThreshold           = 0x4,
    hipMemPoolAttrReservedMemCurrent         = 0x5,
    hipMemPoolAttrReservedMemHigh            = 0x6,
    hipMemPoolAttrUsedMemCurrent             = 0x7,
    hipMemPoolAttrUsedMemHigh                = 0x8
} hipMemPoolAttr;
 typedef enum hipMemLocationType {
    hipMemLocationTypeInvalid = 0,
    hipMemLocationTypeDevice = 1    
} hipMemLocationType;
typedef struct hipMemLocation {
    hipMemLocationType type;  
    int id;                   
} hipMemLocation;
typedef enum hipMemAccessFlags {
    hipMemAccessFlagsProtNone      = 0,  
    hipMemAccessFlagsProtRead      = 1,  
    hipMemAccessFlagsProtReadWrite = 3   
} hipMemAccessFlags;
typedef struct hipMemAccessDesc {
    hipMemLocation      location; 
    hipMemAccessFlags   flags;    
} hipMemAccessDesc;
typedef enum hipMemAllocationType {
    hipMemAllocationTypeInvalid = 0x0,
    hipMemAllocationTypePinned  = 0x1,
    hipMemAllocationTypeUncached = 0x40000000,
    hipMemAllocationTypeMax     = 0x7FFFFFFF
} hipMemAllocationType;
typedef enum hipMemAllocationHandleType {
    hipMemHandleTypeNone                    = 0x0,  
    hipMemHandleTypePosixFileDescriptor     = 0x1,  
    hipMemHandleTypeWin32                   = 0x2,  
    hipMemHandleTypeWin32Kmt                = 0x4   
} hipMemAllocationHandleType;
typedef struct hipMemPoolProps {
    hipMemAllocationType       allocType;   
    hipMemAllocationHandleType handleTypes; 
    hipMemLocation             location;    
    void*                       win32SecurityAttributes;
    size_t                      maxSize;  
    unsigned char               reserved[56];  
} hipMemPoolProps;
typedef struct hipMemPoolPtrExportData {
    unsigned char reserved[64];
} hipMemPoolPtrExportData;
 
typedef enum hipFuncAttribute {
  hipFuncAttributeMaxDynamicSharedMemorySize = 8,          
  hipFuncAttributePreferredSharedMemoryCarveout = 9,       
  hipFuncAttributeMax
} hipFuncAttribute;
typedef enum hipFuncCache_t {
    hipFuncCachePreferNone,    
    hipFuncCachePreferShared,  
    hipFuncCachePreferL1,      
    hipFuncCachePreferEqual,   
} hipFuncCache_t;
typedef enum hipSharedMemConfig {
    hipSharedMemBankSizeDefault,  
    hipSharedMemBankSizeFourByte,  
    hipSharedMemBankSizeEightByte  
} hipSharedMemConfig;
typedef struct dim3 {
    uint32_t x;  
    uint32_t y;  
    uint32_t z;  
#ifdef __cplusplus
    constexpr __host__ __device__ dim3(uint32_t _x = 1, uint32_t _y = 1, uint32_t _z = 1) : x(_x), y(_y), z(_z){};
#endif
} dim3;
typedef struct hipLaunchParams_t {
    void* func;             
    dim3 gridDim;           
    dim3 blockDim;          
    void **args;            
    size_t sharedMem;       
    hipStream_t stream;     
} hipLaunchParams;
typedef struct hipFunctionLaunchParams_t {
    hipFunction_t function;      
    unsigned int gridDimX;       
    unsigned int gridDimY;       
    unsigned int gridDimZ;       
    unsigned int blockDimX;      
    unsigned int blockDimY;      
    unsigned int blockDimZ;      
    unsigned int sharedMemBytes; 
    hipStream_t hStream;         
    void **kernelParams;         
} hipFunctionLaunchParams;
typedef enum hipExternalMemoryHandleType_enum {
  hipExternalMemoryHandleTypeOpaqueFd = 1,
  hipExternalMemoryHandleTypeOpaqueWin32 = 2,
  hipExternalMemoryHandleTypeOpaqueWin32Kmt = 3,
  hipExternalMemoryHandleTypeD3D12Heap = 4,
  hipExternalMemoryHandleTypeD3D12Resource = 5,
  hipExternalMemoryHandleTypeD3D11Resource = 6,
  hipExternalMemoryHandleTypeD3D11ResourceKmt = 7,
  hipExternalMemoryHandleTypeNvSciBuf         = 8
} hipExternalMemoryHandleType;
typedef struct hipExternalMemoryHandleDesc_st {
  hipExternalMemoryHandleType type;
  union {
    int fd;
    struct {
      void *handle;
      const void *name;
    } win32;
    const void *nvSciBufObject;
  } handle;
  unsigned long long size;
  unsigned int flags;
  unsigned int reserved[16];
} hipExternalMemoryHandleDesc;
typedef struct hipExternalMemoryBufferDesc_st {
  unsigned long long offset;
  unsigned long long size;
  unsigned int flags;
  unsigned int reserved[16];
} hipExternalMemoryBufferDesc;
typedef struct hipExternalMemoryMipmappedArrayDesc_st {
  unsigned long long offset;
  hipChannelFormatDesc formatDesc;
  hipExtent extent;
  unsigned int flags;
  unsigned int numLevels;
} hipExternalMemoryMipmappedArrayDesc;
typedef void* hipExternalMemory_t;
typedef enum hipExternalSemaphoreHandleType_enum {
  hipExternalSemaphoreHandleTypeOpaqueFd = 1,
  hipExternalSemaphoreHandleTypeOpaqueWin32 = 2,
  hipExternalSemaphoreHandleTypeOpaqueWin32Kmt = 3,
  hipExternalSemaphoreHandleTypeD3D12Fence = 4,
  hipExternalSemaphoreHandleTypeD3D11Fence = 5,
  hipExternalSemaphoreHandleTypeNvSciSync = 6,
  hipExternalSemaphoreHandleTypeKeyedMutex = 7,
  hipExternalSemaphoreHandleTypeKeyedMutexKmt = 8,
  hipExternalSemaphoreHandleTypeTimelineSemaphoreFd = 9,
  hipExternalSemaphoreHandleTypeTimelineSemaphoreWin32 = 10
} hipExternalSemaphoreHandleType;
typedef struct hipExternalSemaphoreHandleDesc_st {
  hipExternalSemaphoreHandleType type;
  union {
    int fd;
    struct {
      void* handle;
      const void* name;
    } win32;
    const void* NvSciSyncObj;
  } handle;
  unsigned int flags;
  unsigned int reserved[16];
} hipExternalSemaphoreHandleDesc;
typedef void* hipExternalSemaphore_t;
typedef struct hipExternalSemaphoreSignalParams_st {
  struct {
    struct {
      unsigned long long value;
    } fence;
    union {
      void *fence;
      unsigned long long reserved;
    } nvSciSync;
    struct {
      unsigned long long key;
    } keyedMutex;
    unsigned int reserved[12];
  } params;
  unsigned int flags;
  unsigned int reserved[16];
} hipExternalSemaphoreSignalParams;
typedef struct hipExternalSemaphoreWaitParams_st {
  struct {
    struct {
      unsigned long long value;
    } fence;
    union {
      void *fence;
      unsigned long long reserved;
    } nvSciSync;
    struct {
      unsigned long long key;
      unsigned int timeoutMs;
    } keyedMutex;
    unsigned int reserved[10];
  } params;
  unsigned int flags;
  unsigned int reserved[16];
} hipExternalSemaphoreWaitParams;
 
#if __HIP_HAS_GET_PCH
    void __hipGetPCH(const char** pch, unsigned int*size);
#endif
 
typedef enum hipGraphicsRegisterFlags {
    hipGraphicsRegisterFlagsNone = 0,
    hipGraphicsRegisterFlagsReadOnly = 1,  
    hipGraphicsRegisterFlagsWriteDiscard =
        2,  
    hipGraphicsRegisterFlagsSurfaceLoadStore = 4,  
    hipGraphicsRegisterFlagsTextureGather =
        8  
} hipGraphicsRegisterFlags;
 
typedef struct _hipGraphicsResource hipGraphicsResource;
 
typedef hipGraphicsResource* hipGraphicsResource_t;
 
typedef struct ihipGraph* hipGraph_t;
typedef struct hipGraphNode* hipGraphNode_t;
typedef struct hipGraphExec* hipGraphExec_t;
 
typedef struct hipUserObject* hipUserObject_t;
 
 
typedef enum hipGraphNodeType {
  hipGraphNodeTypeKernel = 0,             
  hipGraphNodeTypeMemcpy = 1,             
  hipGraphNodeTypeMemset = 2,             
  hipGraphNodeTypeHost = 3,               
  hipGraphNodeTypeGraph = 4,              
  hipGraphNodeTypeEmpty = 5,              
  hipGraphNodeTypeWaitEvent = 6,          
  hipGraphNodeTypeEventRecord = 7,        
  hipGraphNodeTypeExtSemaphoreSignal = 8, 
  hipGraphNodeTypeExtSemaphoreWait = 9,   
  hipGraphNodeTypeMemAlloc = 10,          
  hipGraphNodeTypeMemFree = 11,           
  hipGraphNodeTypeMemcpyFromSymbol = 12,  
  hipGraphNodeTypeMemcpyToSymbol = 13,    
  hipGraphNodeTypeBatchMemOp = 14,        
  hipGraphNodeTypeCount
} hipGraphNodeType;
 
typedef void (*hipHostFn_t)(void* userData);
typedef struct hipHostNodeParams {
  hipHostFn_t fn;
  void* userData;
} hipHostNodeParams;
typedef struct hipKernelNodeParams {
  dim3 blockDim;
  void** extra;
  void* func;
  dim3 gridDim;
  void** kernelParams;
  unsigned int sharedMemBytes;
} hipKernelNodeParams;
typedef struct hipMemsetParams {
  void* dst;
  unsigned int elementSize;
  size_t height;
  size_t pitch;
  unsigned int value;
  size_t width;
} hipMemsetParams;
 
typedef struct hipMemAllocNodeParams {
    hipMemPoolProps poolProps;          
    const hipMemAccessDesc* accessDescs;
    size_t  accessDescCount;            
    size_t  bytesize;                   
    void*   dptr;                       
} hipMemAllocNodeParams;
 
 
typedef enum hipAccessProperty {
    hipAccessPropertyNormal = 0,
    hipAccessPropertyStreaming  = 1,
    hipAccessPropertyPersisting = 2,
} hipAccessProperty;
typedef struct hipAccessPolicyWindow {
    void* base_ptr;
    hipAccessProperty hitProp;
    float hitRatio;
    hipAccessProperty missProp;
    size_t num_bytes;
} hipAccessPolicyWindow;
 
typedef enum hipLaunchAttributeID {
  hipLaunchAttributeAccessPolicyWindow = 1,                
  hipLaunchAttributeCooperative = 2,                       
  hipLaunchAttributePriority = 8,                          
  hipLaunchAttributeMax
} hipLaunchAttributeID;
 
 
typedef union hipLaunchAttributeValue {
    char pad[64];                              
    hipAccessPolicyWindow accessPolicyWindow;  
    int cooperative;                           
    int priority;                              
} hipLaunchAttributeValue;
 
#define hipKernelNodeAttrID hipLaunchAttributeID
#define hipKernelNodeAttributeAccessPolicyWindow hipLaunchAttributeAccessPolicyWindow
#define hipKernelNodeAttributeCooperative hipLaunchAttributeCooperative
#define hipKernelNodeAttributePriority hipLaunchAttributePriority
 
#define hipKernelNodeAttrValue hipLaunchAttributeValue
 
#define hipDrvLaunchAttributeCooperative  hipLaunchAttributeCooperative
 
#define hipDrvLaunchAttributeID hipLaunchAttributeID
#define hipDrvLaunchAttributeValue hipLaunchAttributeValue
#define hipDrvLaunchAttribute hipLaunchAttribute
 
typedef enum hipGraphExecUpdateResult {
  hipGraphExecUpdateSuccess = 0x0,  
  hipGraphExecUpdateError = 0x1,  
  hipGraphExecUpdateErrorTopologyChanged = 0x2,  
  hipGraphExecUpdateErrorNodeTypeChanged = 0x3,  
  hipGraphExecUpdateErrorFunctionChanged =
      0x4,  
  hipGraphExecUpdateErrorParametersChanged =
      0x5,  
  hipGraphExecUpdateErrorNotSupported =
      0x6,  
  hipGraphExecUpdateErrorUnsupportedFunctionChange = 0x7
} hipGraphExecUpdateResult;
 
typedef enum hipStreamCaptureMode {
  hipStreamCaptureModeGlobal = 0,
  hipStreamCaptureModeThreadLocal,
  hipStreamCaptureModeRelaxed
} hipStreamCaptureMode;
typedef enum hipStreamCaptureStatus {
  hipStreamCaptureStatusNone = 0,    
  hipStreamCaptureStatusActive,      
  hipStreamCaptureStatusInvalidated  
} hipStreamCaptureStatus;
 
typedef enum hipStreamUpdateCaptureDependenciesFlags {
  hipStreamAddCaptureDependencies = 0,  
  hipStreamSetCaptureDependencies,      
} hipStreamUpdateCaptureDependenciesFlags;
 
typedef enum hipGraphMemAttributeType {
  hipGraphMemAttrUsedMemCurrent = 0, 
  hipGraphMemAttrUsedMemHigh,        
  hipGraphMemAttrReservedMemCurrent, 
  hipGraphMemAttrReservedMemHigh,    
}hipGraphMemAttributeType;
typedef enum hipUserObjectFlags {
  hipUserObjectNoDestructorSync = 0x1, 
} hipUserObjectFlags;
 
typedef enum hipUserObjectRetainFlags {
  hipGraphUserObjectMove = 0x1, 
} hipUserObjectRetainFlags;
 
typedef enum hipGraphInstantiateFlags {
  hipGraphInstantiateFlagAutoFreeOnLaunch =
      1,  
  hipGraphInstantiateFlagUpload =
      2, 
  hipGraphInstantiateFlagDeviceLaunch  =
      4, 
  hipGraphInstantiateFlagUseNodePriority =
      8, 
} hipGraphInstantiateFlags;
 
enum hipGraphDebugDotFlags {
  hipGraphDebugDotFlagsVerbose = 1
      << 0, 
  hipGraphDebugDotFlagsKernelNodeParams = 1 << 2, 
  hipGraphDebugDotFlagsMemcpyNodeParams = 1 << 3, 
  hipGraphDebugDotFlagsMemsetNodeParams = 1 << 4, 
  hipGraphDebugDotFlagsHostNodeParams = 1 << 5,   
  hipGraphDebugDotFlagsEventNodeParams = 1
      << 6, 
  hipGraphDebugDotFlagsExtSemasSignalNodeParams = 1
      << 7, 
  hipGraphDebugDotFlagsExtSemasWaitNodeParams = 1
      << 8, 
  hipGraphDebugDotFlagsKernelNodeAttributes = 1
      << 9, 
  hipGraphDebugDotFlagsHandles = 1
      << 10 
};
 
typedef enum hipGraphInstantiateResult {
    hipGraphInstantiateSuccess = 0, 
    hipGraphInstantiateError = 1, 
    hipGraphInstantiateInvalidStructure = 2, 
    hipGraphInstantiateNodeOperationNotSupported = 3, 
    hipGraphInstantiateMultipleDevicesNotSupported = 4, 
}hipGraphInstantiateResult;
 
typedef struct hipGraphInstantiateParams {
    hipGraphNode_t errNode_out; 
    unsigned long long flags; 
    hipGraphInstantiateResult result_out; 
    hipStream_t uploadStream; 
} hipGraphInstantiateParams;
 
 
typedef struct hipMemAllocationProp {
    hipMemAllocationType type;                       
    union {
      hipMemAllocationHandleType requestedHandleType;   
      hipMemAllocationHandleType requestedHandleTypes;  
    };
    hipMemLocation location;                         
    void* win32HandleMetaData;                       
    struct {
        unsigned char compressionType;               
        unsigned char gpuDirectRDMACapable;          
        unsigned short usage;                        
    } allocFlags;
} hipMemAllocationProp;
 
typedef struct hipExternalSemaphoreSignalNodeParams {
    hipExternalSemaphore_t* extSemArray;
    const hipExternalSemaphoreSignalParams* paramsArray;
    unsigned int numExtSems;
} hipExternalSemaphoreSignalNodeParams;
 
typedef struct hipExternalSemaphoreWaitNodeParams {
    hipExternalSemaphore_t* extSemArray;
    const hipExternalSemaphoreWaitParams* paramsArray;
    unsigned int numExtSems;
} hipExternalSemaphoreWaitNodeParams;
 
typedef struct ihipMemGenericAllocationHandle* hipMemGenericAllocationHandle_t;
 
typedef enum hipMemAllocationGranularity_flags {
    hipMemAllocationGranularityMinimum     = 0x0,  
    hipMemAllocationGranularityRecommended = 0x1   
} hipMemAllocationGranularity_flags;
 
typedef enum hipMemHandleType {
    hipMemHandleTypeGeneric = 0x0  
} hipMemHandleType;
 
typedef enum hipMemOperationType {
    hipMemOperationTypeMap   = 0x1,   
    hipMemOperationTypeUnmap = 0x2    
} hipMemOperationType;
 
typedef enum hipArraySparseSubresourceType {
    hipArraySparseSubresourceTypeSparseLevel = 0x0, 
    hipArraySparseSubresourceTypeMiptail     = 0x1  
} hipArraySparseSubresourceType;
 
typedef struct hipArrayMapInfo {
     hipResourceType resourceType;                   
     union {
         hipMipmappedArray mipmap;
         hipArray_t array;
     } resource;
     hipArraySparseSubresourceType subresourceType;  
     union {
         struct {
             unsigned int level;   
             unsigned int layer;   
             unsigned int offsetX;                   
             unsigned int offsetY;                   
             unsigned int offsetZ;                   
             unsigned int extentWidth;               
             unsigned int extentHeight;              
             unsigned int extentDepth;               
         } sparseLevel;
         struct {
             unsigned int layer;   
             unsigned long long offset;              
             unsigned long long size;                
         } miptail;
     } subresource;
     hipMemOperationType memOperationType;           
     hipMemHandleType memHandleType;                 
     union {
         hipMemGenericAllocationHandle_t memHandle;
     } memHandle;
     unsigned long long offset;                      
     unsigned int deviceBitMask;                     
     unsigned int flags;                             
     unsigned int reserved[2];                       
} hipArrayMapInfo;
 
typedef struct hipMemcpyNodeParams {
    int flags;                     
    int reserved[3];               
    hipMemcpy3DParms copyParams;   
} hipMemcpyNodeParams;
 
typedef struct hipChildGraphNodeParams {
    hipGraph_t graph; 
} hipChildGraphNodeParams;
 
typedef struct hipEventWaitNodeParams {
    hipEvent_t event; 
} hipEventWaitNodeParams;
 
typedef struct hipEventRecordNodeParams {
    hipEvent_t event; 
} hipEventRecordNodeParams;
 
typedef struct hipMemFreeNodeParams {
    void *dptr; 
} hipMemFreeNodeParams;
 
typedef struct hipGraphNodeParams {
    hipGraphNodeType type;
    int reserved0[3];
    union {
        long long                            reserved1[29];
        hipKernelNodeParams                  kernel;
        hipMemcpyNodeParams                  memcpy;
        hipMemsetParams                      memset;
        hipHostNodeParams                    host;
        hipChildGraphNodeParams              graph;
        hipEventWaitNodeParams               eventWait;
        hipEventRecordNodeParams             eventRecord;
        hipExternalSemaphoreSignalNodeParams extSemSignal;
        hipExternalSemaphoreWaitNodeParams   extSemWait;
        hipMemAllocNodeParams                alloc;
        hipMemFreeNodeParams                 free;
    };
 
    long long reserved2;
} hipGraphNodeParams;
 
#define hipGraphKernelNodePortDefault 0
 
#define hipGraphKernelNodePortLaunchCompletion 2
 
#define hipGraphKernelNodePortProgrammatic 1
 
typedef enum hipGraphDependencyType {
  hipGraphDependencyTypeDefault = 0,
  hipGraphDependencyTypeProgrammatic = 1
}hipGraphDependencyType;
 
typedef struct hipGraphEdgeData {
  unsigned char
      from_port;  
  unsigned char reserved[5];  
  unsigned char
      to_port;  
  unsigned char type;  
} hipGraphEdgeData;
 
 
typedef struct hipLaunchAttribute_st {
  hipLaunchAttributeID id;                        
  char pad[8 - sizeof(hipLaunchAttributeID)];     
  union {
    hipLaunchAttributeValue val;                  
    hipLaunchAttributeValue value;                
  };
} hipLaunchAttribute;
 
typedef struct hipLaunchConfig_st {
  dim3 gridDim;              
  dim3 blockDim;             
  size_t dynamicSmemBytes;   
  hipStream_t stream;        
  hipLaunchAttribute* attrs; 
  unsigned int numAttrs;     
} hipLaunchConfig_t;
 
typedef struct HIP_LAUNCH_CONFIG_st {
  unsigned int gridDimX;        
  unsigned int gridDimY;        
  unsigned int gridDimZ;        
  unsigned int blockDimX;       
  unsigned int blockDimY;       
  unsigned int blockDimZ;       
  unsigned int sharedMemBytes;  
  hipStream_t hStream;          
  hipLaunchAttribute* attrs;    
  unsigned int numAttrs;        
} HIP_LAUNCH_CONFIG;
 
typedef enum hipMemRangeHandleType {
  hipMemRangeHandleTypeDmaBufFd = 0x1,
  hipMemRangeHandleTypeMax = 0x7fffffff
} hipMemRangeHandleType;
 
typedef enum hipMemRangeFlags {
  hipMemRangeFlagDmaBufMappingTypePcie = 0x1,
  hipMemRangeFlagsMax = 0x7fffffff
} hipMemRangeFlags;
 
// Doxygen end group GlobalDefs
// TODO-ctx - more description on error codes.
hipError_t hipInit(unsigned int flags);
 
hipError_t hipDriverGetVersion(int* driverVersion);
hipError_t hipRuntimeGetVersion(int* runtimeVersion);
hipError_t hipDeviceGet(hipDevice_t* device, int ordinal);
 
hipError_t hipDeviceComputeCapability(int* major, int* minor, hipDevice_t device);
hipError_t hipDeviceGetName(char* name, int len, hipDevice_t device);
hipError_t hipDeviceGetUuid(hipUUID* uuid, hipDevice_t device);
hipError_t hipDeviceGetP2PAttribute(int* value, hipDeviceP2PAttr attr,
                                    int srcDevice, int dstDevice);
hipError_t hipDeviceGetPCIBusId(char* pciBusId, int len, int device);
hipError_t hipDeviceGetByPCIBusId(int* device, const char* pciBusId);
hipError_t hipDeviceTotalMem(size_t* bytes, hipDevice_t device);
// doxygen end initialization
hipError_t hipDeviceSynchronize(void);
hipError_t hipDeviceReset(void);
hipError_t hipSetDevice(int deviceId);
hipError_t hipSetValidDevices(int* device_arr, int len);
hipError_t hipGetDevice(int* deviceId);
hipError_t hipGetDeviceCount(int* count);
hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int deviceId);
hipError_t hipDeviceGetDefaultMemPool(hipMemPool_t* mem_pool, int device);
hipError_t hipDeviceSetMemPool(int device, hipMemPool_t mem_pool);
hipError_t hipDeviceGetMemPool(hipMemPool_t* mem_pool, int device);
hipError_t hipGetDeviceProperties(hipDeviceProp_t* prop, int deviceId);
hipError_t hipDeviceGetTexture1DLinearMaxWidth(size_t* max_width,
                                               const hipChannelFormatDesc* desc,
                                               int device);
hipError_t hipDeviceSetCacheConfig(hipFuncCache_t cacheConfig);
hipError_t hipDeviceGetCacheConfig(hipFuncCache_t* cacheConfig);
hipError_t hipDeviceGetLimit(size_t* pValue, enum hipLimit_t limit);
hipError_t hipDeviceSetLimit ( enum hipLimit_t limit, size_t value );
hipError_t hipDeviceGetSharedMemConfig(hipSharedMemConfig* pConfig);
hipError_t hipGetDeviceFlags(unsigned int* flags);
hipError_t hipDeviceSetSharedMemConfig(hipSharedMemConfig config);
hipError_t hipSetDeviceFlags(unsigned flags);
hipError_t hipChooseDevice(int* device, const hipDeviceProp_t* prop);
hipError_t hipExtGetLinkTypeAndHopCount(int device1, int device2, uint32_t* linktype, uint32_t* hopcount);
// TODO: implement IPC apis
hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr);
hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags);
hipError_t hipIpcCloseMemHandle(void* devPtr);
 
hipError_t hipIpcGetEventHandle(hipIpcEventHandle_t* handle, hipEvent_t event);
 
hipError_t hipIpcOpenEventHandle(hipEvent_t* event, hipIpcEventHandle_t handle);
 
// end doxygen Device
hipError_t hipFuncSetAttribute(const void* func, hipFuncAttribute attr, int value);
hipError_t hipFuncSetCacheConfig(const void* func, hipFuncCache_t config);
hipError_t hipFuncSetSharedMemConfig(const void* func, hipSharedMemConfig config);
//doxygen end execution
hipError_t hipGetLastError(void);
 
hipError_t hipExtGetLastError(void);
 
hipError_t hipPeekAtLastError(void);
const char* hipGetErrorName(hipError_t hip_error);
const char* hipGetErrorString(hipError_t hipError);
hipError_t hipDrvGetErrorName(hipError_t hipError, const char** errorString);
hipError_t hipDrvGetErrorString(hipError_t hipError, const char** errorString);
// end doxygen Error
hipError_t hipStreamCreate(hipStream_t* stream);
hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags);
hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, int priority);
hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority);
hipError_t hipStreamDestroy(hipStream_t stream);
hipError_t hipStreamQuery(hipStream_t stream);
hipError_t hipStreamSynchronize(hipStream_t stream);
hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags __dparm(0));
hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags);
hipError_t hipStreamGetPriority(hipStream_t stream, int* priority);
hipError_t hipStreamGetDevice(hipStream_t stream, hipDevice_t* device);
hipError_t hipExtStreamCreateWithCUMask(hipStream_t* stream, uint32_t cuMaskSize, const uint32_t* cuMask);
hipError_t hipExtStreamGetCUMask(hipStream_t stream, uint32_t cuMaskSize, uint32_t* cuMask);
typedef void (*hipStreamCallback_t)(hipStream_t stream, hipError_t status, void* userData);
hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData,
                                unsigned int flags);
// end doxygen Stream
hipError_t hipStreamWaitValue32(hipStream_t stream, void* ptr, uint32_t value, unsigned int flags,
                                uint32_t mask __dparm(0xFFFFFFFF));
 
hipError_t hipStreamWaitValue64(hipStream_t stream, void* ptr, uint64_t value, unsigned int flags,
                                uint64_t mask __dparm(0xFFFFFFFFFFFFFFFF));
 
hipError_t hipStreamWriteValue32(hipStream_t stream, void* ptr, uint32_t value, unsigned int flags);
hipError_t hipStreamWriteValue64(hipStream_t stream, void* ptr, uint64_t value, unsigned int flags);
 
hipError_t hipStreamBatchMemOp(hipStream_t stream, unsigned int count,
                               hipStreamBatchMemOpParams* paramArray, unsigned int flags);
 
hipError_t hipGraphAddBatchMemOpNode(hipGraphNode_t *phGraphNode, hipGraph_t hGraph,
                                     const hipGraphNode_t *dependencies, size_t numDependencies,
                                     const hipBatchMemOpNodeParams* nodeParams);
 
hipError_t hipGraphBatchMemOpNodeGetParams(hipGraphNode_t hNode,
                                           hipBatchMemOpNodeParams* nodeParams_out);
 
hipError_t hipGraphBatchMemOpNodeSetParams(hipGraphNode_t hNode,
                                           hipBatchMemOpNodeParams* nodeParams);
 
hipError_t hipGraphExecBatchMemOpNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t hNode,
                                               const hipBatchMemOpNodeParams* nodeParams);
 
// end doxygen Stream Memory Operations
hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags);
hipError_t hipEventCreate(hipEvent_t* event);
hipError_t hipEventRecordWithFlags(hipEvent_t event, hipStream_t stream __dparm(0),
                                   unsigned int flags __dparm(0));
#ifdef __cplusplus
hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream = NULL);
#else
hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream);
#endif
hipError_t hipEventDestroy(hipEvent_t event);
hipError_t hipEventSynchronize(hipEvent_t event);
hipError_t hipEventElapsedTime(float* ms, hipEvent_t start, hipEvent_t stop);
hipError_t hipEventQuery(hipEvent_t event);
// end doxygen Events
hipError_t hipPointerSetAttribute(const void* value, hipPointer_attribute attribute,
                                  hipDeviceptr_t ptr);
 
 
hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void* ptr);
hipError_t hipPointerGetAttribute(void* data, hipPointer_attribute attribute,
                                  hipDeviceptr_t ptr);
hipError_t hipDrvPointerGetAttributes(unsigned int numAttributes, hipPointer_attribute* attributes,
                                      void** data, hipDeviceptr_t ptr);
hipError_t hipImportExternalSemaphore(hipExternalSemaphore_t* extSem_out,
                                      const hipExternalSemaphoreHandleDesc* semHandleDesc);
hipError_t hipSignalExternalSemaphoresAsync(const hipExternalSemaphore_t* extSemArray,
                                            const hipExternalSemaphoreSignalParams* paramsArray,
                                            unsigned int numExtSems, hipStream_t stream);
hipError_t hipWaitExternalSemaphoresAsync(const hipExternalSemaphore_t* extSemArray,
                                              const hipExternalSemaphoreWaitParams* paramsArray,
                                              unsigned int numExtSems, hipStream_t stream);
hipError_t hipDestroyExternalSemaphore(hipExternalSemaphore_t extSem);
 
hipError_t hipImportExternalMemory(hipExternalMemory_t* extMem_out, const hipExternalMemoryHandleDesc* memHandleDesc);
hipError_t hipExternalMemoryGetMappedBuffer(void **devPtr, hipExternalMemory_t extMem, const hipExternalMemoryBufferDesc *bufferDesc);
hipError_t hipDestroyExternalMemory(hipExternalMemory_t extMem);
hipError_t hipExternalMemoryGetMappedMipmappedArray(hipMipmappedArray_t* mipmap, hipExternalMemory_t extMem,
    const hipExternalMemoryMipmappedArrayDesc* mipmapDesc);
 // end of external resource
hipError_t hipMalloc(void** ptr, size_t size);
hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags);
 
 
HIP_DEPRECATED("use hipHostMalloc instead")
hipError_t hipMallocHost(void** ptr, size_t size);
HIP_DEPRECATED("use hipHostMalloc instead")
hipError_t hipMemAllocHost(void** ptr, size_t size);
// end doxygen deprecated management memory
hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags);
hipError_t hipMallocManaged(void** dev_ptr,
                            size_t size,
                            unsigned int flags __dparm(hipMemAttachGlobal));
hipError_t hipMemPrefetchAsync(const void* dev_ptr,
                               size_t count,
                               int device,
                               hipStream_t stream __dparm(0));
hipError_t hipMemAdvise(const void* dev_ptr,
                        size_t count,
                        hipMemoryAdvise advice,
                        int device);
hipError_t hipMemRangeGetAttribute(void* data,
                                   size_t data_size,
                                   hipMemRangeAttribute attribute,
                                   const void* dev_ptr,
                                   size_t count);
hipError_t hipMemRangeGetAttributes(void** data,
                                    size_t* data_sizes,
                                    hipMemRangeAttribute* attributes,
                                    size_t num_attributes,
                                    const void* dev_ptr,
                                    size_t count);
hipError_t hipStreamAttachMemAsync(hipStream_t stream,
                                   void* dev_ptr,
                                   size_t length __dparm(0),
                                   unsigned int flags __dparm(hipMemAttachSingle));
// end doxygen Managed Memory
hipError_t hipMallocAsync(void** dev_ptr, size_t size, hipStream_t stream);
hipError_t hipFreeAsync(void* dev_ptr, hipStream_t stream);
hipError_t hipMemPoolTrimTo(hipMemPool_t mem_pool, size_t min_bytes_to_hold);
hipError_t hipMemPoolSetAttribute(hipMemPool_t mem_pool, hipMemPoolAttr attr, void* value);
hipError_t hipMemPoolGetAttribute(hipMemPool_t mem_pool, hipMemPoolAttr attr, void* value);
hipError_t hipMemPoolSetAccess(hipMemPool_t mem_pool, const hipMemAccessDesc* desc_list, size_t count);
hipError_t hipMemPoolGetAccess(hipMemAccessFlags* flags, hipMemPool_t mem_pool, hipMemLocation* location);
hipError_t hipMemPoolCreate(hipMemPool_t* mem_pool, const hipMemPoolProps* pool_props);
hipError_t hipMemPoolDestroy(hipMemPool_t mem_pool);
hipError_t hipMallocFromPoolAsync(void** dev_ptr, size_t size, hipMemPool_t mem_pool, hipStream_t stream);
hipError_t hipMemPoolExportToShareableHandle(
    void*                      shared_handle,
    hipMemPool_t               mem_pool,
    hipMemAllocationHandleType handle_type,
    unsigned int               flags);
hipError_t hipMemPoolImportFromShareableHandle(
    hipMemPool_t*              mem_pool,
    void*                      shared_handle,
    hipMemAllocationHandleType handle_type,
    unsigned int               flags);
hipError_t hipMemPoolExportPointer(hipMemPoolPtrExportData* export_data, void* dev_ptr);
hipError_t hipMemPoolImportPointer(
    void**                   dev_ptr,
    hipMemPool_t             mem_pool,
    hipMemPoolPtrExportData* export_data);
// Doxygen end of ordered memory allocator
hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags);
hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, unsigned int flags);
hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr);
hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags);
hipError_t hipHostUnregister(void* hostPtr);
hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height);
hipError_t hipMemAllocPitch(hipDeviceptr_t* dptr, size_t* pitch, size_t widthInBytes, size_t height,
   unsigned int elementSizeBytes);
hipError_t hipFree(void* ptr);
hipError_t hipFreeHost(void* ptr);
hipError_t hipHostFree(void* ptr);
hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind);
hipError_t hipMemcpyWithStream(void* dst, const void* src, size_t sizeBytes,
                               hipMemcpyKind kind, hipStream_t stream);
hipError_t hipMemcpyHtoD(hipDeviceptr_t dst, const void* src, size_t sizeBytes);
hipError_t hipMemcpyDtoH(void* dst, hipDeviceptr_t src, size_t sizeBytes);
hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes);
hipError_t hipMemcpyAtoD(hipDeviceptr_t dstDevice, hipArray_t srcArray, size_t srcOffset,
                         size_t ByteCount);
hipError_t hipMemcpyDtoA(hipArray_t dstArray, size_t dstOffset, hipDeviceptr_t srcDevice,
                         size_t ByteCount);
 
hipError_t hipMemcpyAtoA(hipArray_t dstArray, size_t dstOffset, hipArray_t srcArray,
                         size_t srcOffset, size_t ByteCount);
hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, const void* src, size_t sizeBytes,
                              hipStream_t stream);
hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, hipStream_t stream);
hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeBytes,
                              hipStream_t stream);
hipError_t hipMemcpyAtoHAsync(void* dstHost, hipArray_t srcArray, size_t srcOffset,
                              size_t ByteCount, hipStream_t stream);
hipError_t hipMemcpyHtoAAsync(hipArray_t dstArray, size_t dstOffset, const void* srcHost,
                              size_t ByteCount, hipStream_t stream);
hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
    hipModule_t hmod, const char* name);
 
hipError_t hipGetSymbolAddress(void** devPtr, const void* symbol);
 
 
 
hipError_t hipGetSymbolSize(size_t* size, const void* symbol);
 
hipError_t hipGetProcAddress(const char* symbol, void** pfn, int  hipVersion, uint64_t flags,
                             hipDriverProcAddressQueryResult* symbolStatus);
 
hipError_t hipMemcpyToSymbol(const void* symbol, const void* src,
                             size_t sizeBytes, size_t offset __dparm(0),
                             hipMemcpyKind kind __dparm(hipMemcpyHostToDevice));
 
hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src,
                                  size_t sizeBytes, size_t offset,
                                  hipMemcpyKind kind, hipStream_t stream __dparm(0));
 
hipError_t hipMemcpyFromSymbol(void* dst, const void* symbol,
                               size_t sizeBytes, size_t offset __dparm(0),
                               hipMemcpyKind kind __dparm(hipMemcpyDeviceToHost));
 
hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbol,
                                    size_t sizeBytes, size_t offset,
                                    hipMemcpyKind kind,
                                    hipStream_t stream __dparm(0));
hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
                          hipStream_t stream __dparm(0));
hipError_t hipMemset(void* dst, int value, size_t sizeBytes);
hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t count);
hipError_t hipMemsetD8Async(hipDeviceptr_t dest, unsigned char value, size_t count, hipStream_t stream __dparm(0));
hipError_t hipMemsetD16(hipDeviceptr_t dest, unsigned short value, size_t count);
hipError_t hipMemsetD16Async(hipDeviceptr_t dest, unsigned short value, size_t count, hipStream_t stream __dparm(0));
hipError_t hipMemsetD32(hipDeviceptr_t dest, int value, size_t count);
hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream __dparm(0));
hipError_t hipMemsetD32Async(hipDeviceptr_t dst, int value, size_t count,
                             hipStream_t stream __dparm(0));
hipError_t hipMemset2D(void* dst, size_t pitch, int value, size_t width, size_t height);
hipError_t hipMemset2DAsync(void* dst, size_t pitch, int value, size_t width, size_t height,hipStream_t stream __dparm(0));
hipError_t hipMemset3D(hipPitchedPtr pitchedDevPtr, int  value, hipExtent extent );
hipError_t hipMemset3DAsync(hipPitchedPtr pitchedDevPtr, int  value, hipExtent extent ,hipStream_t stream __dparm(0));
hipError_t hipMemGetInfo(size_t* free, size_t* total);
 
hipError_t hipMemPtrGetInfo(void* ptr, size_t* size);
hipError_t hipMallocArray(hipArray_t* array, const hipChannelFormatDesc* desc, size_t width,
                          size_t height __dparm(0), unsigned int flags __dparm(hipArrayDefault));
hipError_t hipArrayCreate(hipArray_t* pHandle, const HIP_ARRAY_DESCRIPTOR* pAllocateArray);
hipError_t hipArrayDestroy(hipArray_t array);
hipError_t hipArray3DCreate(hipArray_t* array, const HIP_ARRAY3D_DESCRIPTOR* pAllocateArray);
hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent);
hipError_t hipFreeArray(hipArray_t array);
hipError_t hipMalloc3DArray(hipArray_t* array, const struct hipChannelFormatDesc* desc,
                            struct hipExtent extent, unsigned int flags);
hipError_t hipArrayGetInfo(hipChannelFormatDesc* desc, hipExtent* extent, unsigned int* flags,
                           hipArray_t array);
hipError_t hipArrayGetDescriptor(HIP_ARRAY_DESCRIPTOR* pArrayDescriptor, hipArray_t array);
hipError_t hipArray3DGetDescriptor(HIP_ARRAY3D_DESCRIPTOR* pArrayDescriptor, hipArray_t array);
hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
                       size_t height, hipMemcpyKind kind);
hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy);
hipError_t hipMemcpyParam2DAsync(const hip_Memcpy2D* pCopy, hipStream_t stream __dparm(0));
hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width,
                            size_t height, hipMemcpyKind kind, hipStream_t stream __dparm(0));
hipError_t hipMemcpy2DToArray(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src,
                              size_t spitch, size_t width, size_t height, hipMemcpyKind kind);
hipError_t hipMemcpy2DToArrayAsync(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src,
                                   size_t spitch, size_t width, size_t height, hipMemcpyKind kind,
                                   hipStream_t stream __dparm(0));
hipError_t hipMemcpy2DArrayToArray(hipArray_t dst, size_t wOffsetDst, size_t hOffsetDst,
                                   hipArray_const_t src, size_t wOffsetSrc, size_t hOffsetSrc,
                                   size_t width, size_t height, hipMemcpyKind kind);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipMemcpyToArray(hipArray_t dst, size_t wOffset, size_t hOffset, const void* src,
                            size_t count, hipMemcpyKind kind);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipMemcpyFromArray(void* dst, hipArray_const_t srcArray, size_t wOffset, size_t hOffset,
                              size_t count, hipMemcpyKind kind);
hipError_t hipMemcpy2DFromArray( void* dst, size_t dpitch, hipArray_const_t src, size_t wOffset, size_t hOffset, size_t width, size_t height, hipMemcpyKind kind);
hipError_t hipMemcpy2DFromArrayAsync( void* dst, size_t dpitch, hipArray_const_t src, size_t wOffset, size_t hOffset, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream __dparm(0));
hipError_t hipMemcpyAtoH(void* dst, hipArray_t srcArray, size_t srcOffset, size_t count);
hipError_t hipMemcpyHtoA(hipArray_t dstArray, size_t dstOffset, const void* srcHost, size_t count);
hipError_t hipMemcpy3D(const struct hipMemcpy3DParms* p);
hipError_t hipMemcpy3DAsync(const struct hipMemcpy3DParms* p, hipStream_t stream __dparm(0));
hipError_t hipDrvMemcpy3D(const HIP_MEMCPY3D* pCopy);
hipError_t hipDrvMemcpy3DAsync(const HIP_MEMCPY3D* pCopy, hipStream_t stream);
hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr);
// doxygen end Memory
hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, int deviceId, int peerDeviceId);
hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags);
hipError_t hipDeviceDisablePeerAccess(int peerDeviceId);
 
hipError_t hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDeviceId,
                         size_t sizeBytes);
hipError_t hipMemcpyPeerAsync(void* dst, int dstDeviceId, const void* src, int srcDevice,
                              size_t sizeBytes, hipStream_t stream __dparm(0));
 
// doxygen end PeerToPeer
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxDestroy(hipCtx_t ctx);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxPopCurrent(hipCtx_t* ctx);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxPushCurrent(hipCtx_t ctx);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxSetCurrent(hipCtx_t ctx);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxGetCurrent(hipCtx_t* ctx);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxGetDevice(hipDevice_t* device);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxGetApiVersion(hipCtx_t ctx, unsigned int* apiVersion);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxSynchronize(void);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxGetFlags(unsigned int* flags);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx);
 
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int* flags, int* active);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipDevicePrimaryCtxRelease(hipDevice_t dev);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipDevicePrimaryCtxRetain(hipCtx_t* pctx, hipDevice_t dev);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipDevicePrimaryCtxReset(hipDevice_t dev);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipDevicePrimaryCtxSetFlags(hipDevice_t dev, unsigned int flags);
// doxygen end Context Management
hipError_t hipModuleLoad(hipModule_t* module, const char* fname);
hipError_t hipModuleUnload(hipModule_t module);
hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, const char* kname);
hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func);
hipError_t hipFuncGetAttribute(int* value, hipFunction_attribute attrib, hipFunction_t hfunc);
hipError_t hipGetFuncBySymbol(hipFunction_t* functionPtr, const void* symbolPtr);
hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name);
hipError_t hipModuleLoadData(hipModule_t* module, const void* image);
hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned int numOptions,
                               hipJitOption* options, void** optionValues);
hipError_t hipLinkAddData(hipLinkState_t state, hipJitInputType type, void* data, size_t size,
                                const char* name, unsigned int numOptions, hipJitOption* options,
                                void** optionValues);
 
hipError_t hipLinkAddFile(hipLinkState_t state, hipJitInputType type, const char* path, unsigned int numOptions,
                            hipJitOption* options, void** optionValues);
 
hipError_t hipLinkComplete(hipLinkState_t state, void** hipBinOut, size_t* sizeOut);
 
hipError_t hipLinkCreate(unsigned int numOptions, hipJitOption* options,
                            void** optionValues, hipLinkState_t* stateOut);
hipError_t hipLinkDestroy(hipLinkState_t state);
 
hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY,
                                 unsigned int gridDimZ, unsigned int blockDimX,
                                 unsigned int blockDimY, unsigned int blockDimZ,
                                 unsigned int sharedMemBytes, hipStream_t stream,
                                 void** kernelParams, void** extra);
hipError_t hipModuleLaunchCooperativeKernel(hipFunction_t f, unsigned int gridDimX,
                                            unsigned int gridDimY, unsigned int gridDimZ,
                                            unsigned int blockDimX, unsigned int blockDimY,
                                            unsigned int blockDimZ, unsigned int sharedMemBytes,
                                            hipStream_t stream, void** kernelParams);
hipError_t hipModuleLaunchCooperativeKernelMultiDevice(hipFunctionLaunchParams* launchParamsList,
                                                       unsigned int numDevices,
                                                       unsigned int flags);
hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDimX,
                                      void** kernelParams, unsigned int sharedMemBytes,
                                      hipStream_t stream);
hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList,
                                                 int  numDevices, unsigned int  flags);
 
// Doxygen end group ModuleCooperativeG
hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
                                              int  numDevices, unsigned int  flags);
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,
                                           hipMemRangeHandleType handleType,
                                           unsigned long long flags);
// doxygen end Module
//TODO - Match CUoccupancyB2DSize
hipError_t hipModuleOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize,
                                             hipFunction_t f, size_t dynSharedMemPerBlk,
                                             int blockSizeLimit);
//TODO - Match CUoccupancyB2DSize
hipError_t hipModuleOccupancyMaxPotentialBlockSizeWithFlags(int* gridSize, int* blockSize,
                                             hipFunction_t f, size_t dynSharedMemPerBlk,
                                             int blockSizeLimit, unsigned int  flags);
hipError_t hipModuleOccupancyMaxActiveBlocksPerMultiprocessor(
   int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk);
hipError_t hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
   int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk, unsigned int flags);
hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
   int* numBlocks, const void* f, int blockSize, size_t dynSharedMemPerBlk);
hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
   int* numBlocks, const void* f, int blockSize, size_t dynSharedMemPerBlk, unsigned int flags __dparm(hipOccupancyDefault));
hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize,
                                             const void* f, size_t dynSharedMemPerBlk,
                                             int blockSizeLimit);
// doxygen end Occupancy
// TODO - expand descriptions:
HIP_DEPRECATED("use roctracer/rocTX instead")
hipError_t hipProfilerStart();
HIP_DEPRECATED("use roctracer/rocTX instead")
hipError_t hipProfilerStop();
// doxygen end profiler
hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem __dparm(0), hipStream_t stream __dparm(0));
hipError_t hipSetupArgument(const void* arg, size_t size, size_t offset);
hipError_t hipLaunchByPtr(const void* func);
hipError_t __hipPushCallConfiguration(dim3 gridDim,
                                      dim3 blockDim,
                                      size_t sharedMem __dparm(0),
                                      hipStream_t stream __dparm(0));
hipError_t __hipPopCallConfiguration(dim3 *gridDim,
                                     dim3 *blockDim,
                                     size_t *sharedMem,
                                     hipStream_t *stream);
hipError_t hipLaunchKernel(const void* function_address,
                           dim3 numBlocks,
                           dim3 dimBlocks,
                           void** args,
                           size_t sharedMemBytes __dparm(0),
                           hipStream_t stream __dparm(0));
 
hipError_t hipLaunchHostFunc(hipStream_t stream, hipHostFn_t fn, void* userData);
 
hipError_t hipDrvMemcpy2DUnaligned(const hip_Memcpy2D* pCopy);
//TODO: Move this to hip_ext.h
hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, dim3 dimBlocks,
                              void** args, size_t sharedMemBytes, hipStream_t stream,
                              hipEvent_t startEvent, hipEvent_t stopEvent, int flags);
// doxygen end Clang launch
hipError_t hipCreateTextureObject(
    hipTextureObject_t* pTexObject,
    const hipResourceDesc* pResDesc,
    const hipTextureDesc* pTexDesc,
    const struct hipResourceViewDesc* pResViewDesc);
 
hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject);
 
hipError_t hipGetChannelDesc(
    hipChannelFormatDesc* desc,
    hipArray_const_t array);
 
hipError_t hipGetTextureObjectResourceDesc(
    hipResourceDesc* pResDesc,
    hipTextureObject_t textureObject);
 
hipError_t hipGetTextureObjectResourceViewDesc(
    struct hipResourceViewDesc* pResViewDesc,
    hipTextureObject_t textureObject);
 
hipError_t hipGetTextureObjectTextureDesc(
    hipTextureDesc* pTexDesc,
    hipTextureObject_t textureObject);
 
hipError_t hipTexObjectCreate(
    hipTextureObject_t* pTexObject,
    const HIP_RESOURCE_DESC* pResDesc,
    const HIP_TEXTURE_DESC* pTexDesc,
    const HIP_RESOURCE_VIEW_DESC* pResViewDesc);
 
hipError_t hipTexObjectDestroy(
    hipTextureObject_t texObject);
 
hipError_t hipTexObjectGetResourceDesc(
    HIP_RESOURCE_DESC* pResDesc,
    hipTextureObject_t texObject);
 
hipError_t hipTexObjectGetResourceViewDesc(
    HIP_RESOURCE_VIEW_DESC* pResViewDesc,
    hipTextureObject_t texObject);
 
hipError_t hipTexObjectGetTextureDesc(
    HIP_TEXTURE_DESC* pTexDesc,
    hipTextureObject_t texObject);
 
hipError_t hipMallocMipmappedArray(
    hipMipmappedArray_t *mipmappedArray,
    const struct hipChannelFormatDesc* desc,
    struct hipExtent extent,
    unsigned int numLevels,
    unsigned int flags __dparm(0));
 
hipError_t hipFreeMipmappedArray(hipMipmappedArray_t mipmappedArray);
 
hipError_t hipGetMipmappedArrayLevel(
    hipArray_t *levelArray,
    hipMipmappedArray_const_t mipmappedArray,
    unsigned int level);
 
hipError_t hipMipmappedArrayCreate(
    hipMipmappedArray_t* pHandle,
    HIP_ARRAY3D_DESCRIPTOR* pMipmappedArrayDesc,
    unsigned int numMipmapLevels);
 
hipError_t hipMipmappedArrayDestroy(hipMipmappedArray_t hMipmappedArray);
 
hipError_t hipMipmappedArrayGetLevel(
    hipArray_t* pLevelArray,
    hipMipmappedArray_t hMipMappedArray,
    unsigned int level);
 
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipBindTextureToMipmappedArray(
    const textureReference* tex,
    hipMipmappedArray_const_t mipmappedArray,
    const hipChannelFormatDesc* desc);
 
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipGetTextureReference(
    const textureReference** texref,
    const void* symbol);
 
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefGetBorderColor(float* pBorderColor, const textureReference* texRef);
 
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefGetArray(hipArray_t* pArray, const textureReference* texRef);
 
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefSetAddressMode(
    textureReference* texRef,
    int dim,
    enum hipTextureAddressMode am);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefSetArray(
    textureReference* tex,
    hipArray_const_t array,
    unsigned int flags);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefSetFilterMode(
    textureReference* texRef,
    enum hipTextureFilterMode fm);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefSetFlags(
    textureReference* texRef,
    unsigned int Flags);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefSetFormat(
    textureReference* texRef,
    hipArray_Format fmt,
    int NumPackedComponents);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipBindTexture(
    size_t* offset,
    const textureReference* tex,
    const void* devPtr,
    const hipChannelFormatDesc* desc,
    size_t size __dparm(UINT_MAX));
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipBindTexture2D(
    size_t* offset,
    const textureReference* tex,
    const void* devPtr,
    const hipChannelFormatDesc* desc,
    size_t width,
    size_t height,
    size_t pitch);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipBindTextureToArray(
    const textureReference* tex,
    hipArray_const_t array,
    const hipChannelFormatDesc* desc);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipGetTextureAlignmentOffset(
    size_t* offset,
    const textureReference* texref);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipUnbindTexture(const textureReference* tex);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefGetAddress(
    hipDeviceptr_t* dev_ptr,
    const textureReference* texRef);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefGetAddressMode(
    enum hipTextureAddressMode* pam,
    const textureReference* texRef,
    int dim);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefGetFilterMode(
    enum hipTextureFilterMode* pfm,
    const textureReference* texRef);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefGetFlags(
    unsigned int* pFlags,
    const textureReference* texRef);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefGetFormat(
    hipArray_Format* pFormat,
    int* pNumChannels,
    const textureReference* texRef);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefGetMaxAnisotropy(
    int* pmaxAnsio,
    const textureReference* texRef);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefGetMipmapFilterMode(
    enum hipTextureFilterMode* pfm,
    const textureReference* texRef);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefGetMipmapLevelBias(
    float* pbias,
    const textureReference* texRef);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefGetMipmapLevelClamp(
    float* pminMipmapLevelClamp,
    float* pmaxMipmapLevelClamp,
    const textureReference* texRef);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefGetMipMappedArray(
    hipMipmappedArray_t* pArray,
    const textureReference* texRef);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefSetAddress(
    size_t* ByteOffset,
    textureReference* texRef,
    hipDeviceptr_t dptr,
    size_t bytes);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefSetAddress2D(
    textureReference* texRef,
    const HIP_ARRAY_DESCRIPTOR* desc,
    hipDeviceptr_t dptr,
    size_t Pitch);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefSetMaxAnisotropy(
    textureReference* texRef,
    unsigned int maxAniso);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefSetBorderColor(
    textureReference* texRef,
    float* pBorderColor);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefSetMipmapFilterMode(
    textureReference* texRef,
    enum hipTextureFilterMode fm);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefSetMipmapLevelBias(
    textureReference* texRef,
    float bias);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefSetMipmapLevelClamp(
    textureReference* texRef,
    float minMipMapLevelClamp,
    float maxMipMapLevelClamp);
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
hipError_t hipTexRefSetMipmappedArray(
    textureReference* texRef,
    struct hipMipmappedArray* mipmappedArray,
    unsigned int Flags);
 
// doxygen end deprecated texture management
// doxygen end Texture management
// This group is for HIPrtc
 
// doxygen end Runtime
const char* hipApiName(uint32_t id);
const char* hipKernelNameRef(const hipFunction_t f);
const char* hipKernelNameRefByPtr(const void* hostFunction, hipStream_t stream);
int hipGetStreamDeviceId(hipStream_t stream);
 
// doxygen end Callback
hipError_t hipStreamBeginCapture(hipStream_t stream, hipStreamCaptureMode mode);
 
hipError_t hipStreamBeginCaptureToGraph(hipStream_t stream, hipGraph_t graph,
                                        const hipGraphNode_t* dependencies,
                                        const hipGraphEdgeData* dependencyData,
                                        size_t numDependencies, hipStreamCaptureMode mode);
 
hipError_t hipStreamEndCapture(hipStream_t stream, hipGraph_t* pGraph);
 
hipError_t hipStreamGetCaptureInfo(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus,
                                   unsigned long long* pId);
 
hipError_t hipStreamGetCaptureInfo_v2(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out,
                                      unsigned long long* id_out __dparm(0),
                                      hipGraph_t* graph_out __dparm(0),
                                      const hipGraphNode_t** dependencies_out __dparm(0),
                                      size_t* numDependencies_out __dparm(0));
 
hipError_t hipStreamIsCapturing(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus);
 
hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t* dependencies,
                                              size_t numDependencies,
                                              unsigned int flags __dparm(0));
 
hipError_t hipThreadExchangeStreamCaptureMode(hipStreamCaptureMode* mode);
 
hipError_t hipGraphCreate(hipGraph_t* pGraph, unsigned int flags);
 
hipError_t hipGraphDestroy(hipGraph_t graph);
 
hipError_t hipGraphAddDependencies(hipGraph_t graph, const hipGraphNode_t* from,
                                   const hipGraphNode_t* to, size_t numDependencies);
 
hipError_t hipGraphRemoveDependencies(hipGraph_t graph, const hipGraphNode_t* from,
                                      const hipGraphNode_t* to, size_t numDependencies);
 
hipError_t hipGraphGetEdges(hipGraph_t graph, hipGraphNode_t* from, hipGraphNode_t* to,
                            size_t* numEdges);
 
hipError_t hipGraphGetNodes(hipGraph_t graph, hipGraphNode_t* nodes, size_t* numNodes);
 
hipError_t hipGraphGetRootNodes(hipGraph_t graph, hipGraphNode_t* pRootNodes,
                                size_t* pNumRootNodes);
 
hipError_t hipGraphNodeGetDependencies(hipGraphNode_t node, hipGraphNode_t* pDependencies,
                                       size_t* pNumDependencies);
 
hipError_t hipGraphNodeGetDependentNodes(hipGraphNode_t node, hipGraphNode_t* pDependentNodes,
                                         size_t* pNumDependentNodes);
 
hipError_t hipGraphNodeGetType(hipGraphNode_t node, hipGraphNodeType* pType);
 
hipError_t hipGraphDestroyNode(hipGraphNode_t node);
 
hipError_t hipGraphClone(hipGraph_t* pGraphClone, hipGraph_t originalGraph);
 
hipError_t hipGraphNodeFindInClone(hipGraphNode_t* pNode, hipGraphNode_t originalNode,
                                   hipGraph_t clonedGraph);
 
hipError_t hipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph,
                               hipGraphNode_t* pErrorNode, char* pLogBuffer, size_t bufferSize);
 
hipError_t hipGraphInstantiateWithFlags(hipGraphExec_t* pGraphExec, hipGraph_t graph,
                                        unsigned long long flags);
 
hipError_t hipGraphInstantiateWithParams(hipGraphExec_t* pGraphExec, hipGraph_t graph,
                                        hipGraphInstantiateParams *instantiateParams);
hipError_t hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream);
 
hipError_t hipGraphUpload(hipGraphExec_t graphExec, hipStream_t stream);
 
hipError_t hipGraphAddNode(hipGraphNode_t *pGraphNode, hipGraph_t graph,
                           const hipGraphNode_t *pDependencies, size_t numDependencies,
                           hipGraphNodeParams *nodeParams);
 
hipError_t hipGraphExecGetFlags(hipGraphExec_t graphExec, unsigned long long* flags);
 
hipError_t hipGraphNodeSetParams(hipGraphNode_t node, hipGraphNodeParams *nodeParams);
 
hipError_t hipGraphExecNodeSetParams(hipGraphExec_t graphExec, hipGraphNode_t node, hipGraphNodeParams* nodeParams);
 
hipError_t hipGraphExecDestroy(hipGraphExec_t graphExec);
 
// Check whether an executable graph can be updated with a graph and perform the update if possible.
hipError_t hipGraphExecUpdate(hipGraphExec_t hGraphExec, hipGraph_t hGraph,
                              hipGraphNode_t* hErrorNode_out,
                              hipGraphExecUpdateResult* updateResult_out);
 
hipError_t hipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                 const hipGraphNode_t* pDependencies, size_t numDependencies,
                                 const hipKernelNodeParams* pNodeParams);
 
hipError_t hipGraphKernelNodeGetParams(hipGraphNode_t node, hipKernelNodeParams* pNodeParams);
 
hipError_t hipGraphKernelNodeSetParams(hipGraphNode_t node, const hipKernelNodeParams* pNodeParams);
 
hipError_t hipGraphExecKernelNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node,
                                           const hipKernelNodeParams* pNodeParams);
 
hipError_t hipDrvGraphAddMemcpyNode(hipGraphNode_t* phGraphNode, hipGraph_t hGraph,
                                    const hipGraphNode_t* dependencies,
                                    size_t numDependencies,
                                    const HIP_MEMCPY3D* copyParams, hipCtx_t ctx);
hipError_t hipGraphAddMemcpyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                 const hipGraphNode_t* pDependencies, size_t numDependencies,
                                 const hipMemcpy3DParms* pCopyParams);
hipError_t hipGraphMemcpyNodeGetParams(hipGraphNode_t node, hipMemcpy3DParms* pNodeParams);
 
hipError_t hipGraphMemcpyNodeSetParams(hipGraphNode_t node, const hipMemcpy3DParms* pNodeParams);
 
hipError_t hipGraphKernelNodeSetAttribute(hipGraphNode_t hNode, hipKernelNodeAttrID attr,
                                          const hipKernelNodeAttrValue* value);
hipError_t hipGraphKernelNodeGetAttribute(hipGraphNode_t hNode, hipKernelNodeAttrID attr,
                                          hipKernelNodeAttrValue* value);
hipError_t hipGraphExecMemcpyNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node,
                                           hipMemcpy3DParms* pNodeParams);
 
hipError_t hipGraphAddMemcpyNode1D(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                   const hipGraphNode_t* pDependencies, size_t numDependencies,
                                   void* dst, const void* src, size_t count, hipMemcpyKind kind);
 
hipError_t hipGraphMemcpyNodeSetParams1D(hipGraphNode_t node, void* dst, const void* src,
                                         size_t count, hipMemcpyKind kind);
 
hipError_t hipGraphExecMemcpyNodeSetParams1D(hipGraphExec_t hGraphExec, hipGraphNode_t node,
                                             void* dst, const void* src, size_t count,
                                             hipMemcpyKind kind);
 
hipError_t hipGraphAddMemcpyNodeFromSymbol(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                           const hipGraphNode_t* pDependencies,
                                           size_t numDependencies, void* dst, const void* symbol,
                                           size_t count, size_t offset, hipMemcpyKind kind);
 
hipError_t hipGraphMemcpyNodeSetParamsFromSymbol(hipGraphNode_t node, void* dst, const void* symbol,
                                                 size_t count, size_t offset, hipMemcpyKind kind);
 
hipError_t hipGraphExecMemcpyNodeSetParamsFromSymbol(hipGraphExec_t hGraphExec, hipGraphNode_t node,
                                                     void* dst, const void* symbol, size_t count,
                                                     size_t offset, hipMemcpyKind kind);
 
hipError_t hipGraphAddMemcpyNodeToSymbol(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                         const hipGraphNode_t* pDependencies,
                                         size_t numDependencies, const void* symbol,
                                         const void* src, size_t count, size_t offset,
                                         hipMemcpyKind kind);
 
hipError_t hipGraphMemcpyNodeSetParamsToSymbol(hipGraphNode_t node, const void* symbol,
                                               const void* src, size_t count, size_t offset,
                                               hipMemcpyKind kind);
 
 
hipError_t hipGraphExecMemcpyNodeSetParamsToSymbol(hipGraphExec_t hGraphExec, hipGraphNode_t node,
                                                   const void* symbol, const void* src,
                                                   size_t count, size_t offset, hipMemcpyKind kind);
 
hipError_t hipGraphAddMemsetNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                 const hipGraphNode_t* pDependencies, size_t numDependencies,
                                 const hipMemsetParams* pMemsetParams);
 
hipError_t hipGraphMemsetNodeGetParams(hipGraphNode_t node, hipMemsetParams* pNodeParams);
 
hipError_t hipGraphMemsetNodeSetParams(hipGraphNode_t node, const hipMemsetParams* pNodeParams);
 
hipError_t hipGraphExecMemsetNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node,
                                           const hipMemsetParams* pNodeParams);
 
hipError_t hipGraphAddHostNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                               const hipGraphNode_t* pDependencies, size_t numDependencies,
                               const hipHostNodeParams* pNodeParams);
 
hipError_t hipGraphHostNodeGetParams(hipGraphNode_t node, hipHostNodeParams* pNodeParams);
 
hipError_t hipGraphHostNodeSetParams(hipGraphNode_t node, const hipHostNodeParams* pNodeParams);
 
hipError_t hipGraphExecHostNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node,
                                         const hipHostNodeParams* pNodeParams);
 
hipError_t hipGraphAddChildGraphNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                     const hipGraphNode_t* pDependencies, size_t numDependencies,
                                     hipGraph_t childGraph);
 
hipError_t hipGraphChildGraphNodeGetGraph(hipGraphNode_t node, hipGraph_t* pGraph);
 
hipError_t hipGraphExecChildGraphNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node,
                                               hipGraph_t childGraph);
 
hipError_t hipGraphAddEmptyNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                const hipGraphNode_t* pDependencies, size_t numDependencies);
 
 
hipError_t hipGraphAddEventRecordNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                      const hipGraphNode_t* pDependencies, size_t numDependencies,
                                      hipEvent_t event);
 
hipError_t hipGraphEventRecordNodeGetEvent(hipGraphNode_t node, hipEvent_t* event_out);
 
hipError_t hipGraphEventRecordNodeSetEvent(hipGraphNode_t node, hipEvent_t event);
 
hipError_t hipGraphExecEventRecordNodeSetEvent(hipGraphExec_t hGraphExec, hipGraphNode_t hNode,
                                               hipEvent_t event);
 
hipError_t hipGraphAddEventWaitNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                                    const hipGraphNode_t* pDependencies, size_t numDependencies,
                                    hipEvent_t event);
 
 
hipError_t hipGraphEventWaitNodeGetEvent(hipGraphNode_t node, hipEvent_t* event_out);
 
hipError_t hipGraphEventWaitNodeSetEvent(hipGraphNode_t node, hipEvent_t event);
 
hipError_t hipGraphExecEventWaitNodeSetEvent(hipGraphExec_t hGraphExec, hipGraphNode_t hNode,
                                             hipEvent_t event);
 
hipError_t hipGraphAddMemAllocNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
    const hipGraphNode_t* pDependencies, size_t numDependencies, hipMemAllocNodeParams* pNodeParams);
 
hipError_t hipGraphMemAllocNodeGetParams(hipGraphNode_t node, hipMemAllocNodeParams* pNodeParams);
 
hipError_t hipGraphAddMemFreeNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
    const hipGraphNode_t* pDependencies, size_t numDependencies, void* dev_ptr);
 
hipError_t hipGraphMemFreeNodeGetParams(hipGraphNode_t node, void* dev_ptr);
 
hipError_t hipDeviceGetGraphMemAttribute(int device, hipGraphMemAttributeType attr, void* value);
 
hipError_t hipDeviceSetGraphMemAttribute(int device, hipGraphMemAttributeType attr, void* value);
 
hipError_t hipDeviceGraphMemTrim(int device);
 
hipError_t hipUserObjectCreate(hipUserObject_t* object_out, void* ptr, hipHostFn_t destroy,
                               unsigned int initialRefcount, unsigned int flags);
 
hipError_t hipUserObjectRelease(hipUserObject_t object, unsigned int count __dparm(1));
 
hipError_t hipUserObjectRetain(hipUserObject_t object, unsigned int count __dparm(1));
 
hipError_t hipGraphRetainUserObject(hipGraph_t graph, hipUserObject_t object,
                                    unsigned int count __dparm(1), unsigned int flags __dparm(0));
 
hipError_t hipGraphReleaseUserObject(hipGraph_t graph, hipUserObject_t object,
                                     unsigned int count __dparm(1));
 
hipError_t hipGraphDebugDotPrint(hipGraph_t graph, const char* path, unsigned int flags);
 
hipError_t hipGraphKernelNodeCopyAttributes(hipGraphNode_t hSrc, hipGraphNode_t hDst);
 
hipError_t hipGraphNodeSetEnabled(hipGraphExec_t hGraphExec, hipGraphNode_t hNode,
                                  unsigned int isEnabled);
hipError_t hipGraphNodeGetEnabled(hipGraphExec_t hGraphExec, hipGraphNode_t hNode,
                                  unsigned int* isEnabled);
 
hipError_t hipGraphAddExternalSemaphoresWaitNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                               const hipGraphNode_t* pDependencies, size_t numDependencies,
                               const hipExternalSemaphoreWaitNodeParams* nodeParams);
 
hipError_t hipGraphAddExternalSemaphoresSignalNode(hipGraphNode_t* pGraphNode, hipGraph_t graph,
                               const hipGraphNode_t* pDependencies, size_t numDependencies,
                               const hipExternalSemaphoreSignalNodeParams* nodeParams);
hipError_t hipGraphExternalSemaphoresSignalNodeSetParams(hipGraphNode_t hNode,
                                                         const hipExternalSemaphoreSignalNodeParams* nodeParams);
hipError_t hipGraphExternalSemaphoresWaitNodeSetParams(hipGraphNode_t hNode,
                                                       const hipExternalSemaphoreWaitNodeParams* nodeParams);
hipError_t hipGraphExternalSemaphoresSignalNodeGetParams(hipGraphNode_t hNode,
                                                         hipExternalSemaphoreSignalNodeParams* params_out);
hipError_t hipGraphExternalSemaphoresWaitNodeGetParams(hipGraphNode_t hNode,
                                                       hipExternalSemaphoreWaitNodeParams* params_out);
hipError_t hipGraphExecExternalSemaphoresSignalNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t hNode,
                                                             const hipExternalSemaphoreSignalNodeParams* nodeParams);
hipError_t hipGraphExecExternalSemaphoresWaitNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t hNode,
                                                           const hipExternalSemaphoreWaitNodeParams* nodeParams);
 
hipError_t hipDrvGraphMemcpyNodeGetParams(hipGraphNode_t hNode, HIP_MEMCPY3D* nodeParams);
 
hipError_t hipDrvGraphMemcpyNodeSetParams(hipGraphNode_t hNode, const HIP_MEMCPY3D* nodeParams);
 
hipError_t hipDrvGraphAddMemsetNode(hipGraphNode_t* phGraphNode, hipGraph_t hGraph,
                                 const hipGraphNode_t* dependencies, size_t numDependencies,
                                 const hipMemsetParams* memsetParams, hipCtx_t ctx);
 
hipError_t hipDrvGraphAddMemFreeNode(hipGraphNode_t* phGraphNode, hipGraph_t hGraph,
                                  const hipGraphNode_t* dependencies, size_t numDependencies,
                                  hipDeviceptr_t dptr);
 
hipError_t hipDrvGraphExecMemcpyNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t hNode,
                                   const HIP_MEMCPY3D* copyParams, hipCtx_t ctx);
 
hipError_t hipDrvGraphExecMemsetNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t hNode,
                                   const hipMemsetParams* memsetParams, hipCtx_t ctx);
 
// doxygen end graph API
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 hipMemCreate(hipMemGenericAllocationHandle_t* handle, size_t size, const hipMemAllocationProp* prop, unsigned long long flags);
 
hipError_t hipMemExportToShareableHandle(void* shareableHandle, hipMemGenericAllocationHandle_t handle, hipMemAllocationHandleType handleType, unsigned long long flags);
 
hipError_t hipMemGetAccess(unsigned long long* flags, const hipMemLocation* location, void* ptr);
 
hipError_t hipMemGetAllocationGranularity(size_t* granularity, const hipMemAllocationProp* prop, hipMemAllocationGranularity_flags option);
 
hipError_t hipMemGetAllocationPropertiesFromHandle(hipMemAllocationProp* prop, hipMemGenericAllocationHandle_t handle);
 
hipError_t hipMemImportFromShareableHandle(hipMemGenericAllocationHandle_t* handle, void* osHandle, hipMemAllocationHandleType shHandleType);
 
hipError_t hipMemMap(void* ptr, size_t size, size_t offset, hipMemGenericAllocationHandle_t handle, unsigned long long flags);
 
hipError_t hipMemMapArrayAsync(hipArrayMapInfo* mapInfoList, unsigned int  count, hipStream_t stream);
 
hipError_t hipMemRelease(hipMemGenericAllocationHandle_t handle);
 
hipError_t hipMemRetainAllocationHandle(hipMemGenericAllocationHandle_t* handle, void* addr);
 
hipError_t hipMemSetAccess(void* ptr, size_t size, const hipMemAccessDesc* desc, size_t count);
 
hipError_t hipMemUnmap(void* ptr, size_t size);
 
// doxygen end virtual memory management API
hipError_t hipGraphicsMapResources(int count, hipGraphicsResource_t* resources,
                                   hipStream_t stream  __dparm(0) );
hipError_t hipGraphicsSubResourceGetMappedArray(hipArray_t* array, hipGraphicsResource_t resource,
                                                unsigned int arrayIndex, unsigned int mipLevel);
hipError_t hipGraphicsResourceGetMappedPointer(void** devPtr, size_t* size,
                                               hipGraphicsResource_t resource);
hipError_t hipGraphicsUnmapResources(int count, hipGraphicsResource_t* resources,
                                     hipStream_t stream __dparm(0));
hipError_t hipGraphicsUnregisterResource(hipGraphicsResource_t resource);
// doxygen end GraphicsInterop
hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, const hipResourceDesc* pResDesc);
hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject);
// end of surface
#ifdef __cplusplus
} /* extern "c" */
#endif
#ifdef __cplusplus
#if defined(__clang__) && defined(__HIP__)
template <typename T>
static hipError_t __host__ inline hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize,
    T f, size_t dynSharedMemPerBlk = 0, int blockSizeLimit = 0) {
    return hipOccupancyMaxPotentialBlockSize(gridSize, blockSize, reinterpret_cast<const void*>(f),dynSharedMemPerBlk,blockSizeLimit);
}
template <typename T>
static hipError_t __host__ inline hipOccupancyMaxPotentialBlockSizeWithFlags(int* gridSize, int* blockSize,
    T f, size_t dynSharedMemPerBlk = 0, int blockSizeLimit = 0, unsigned int  flags = 0 ) {
    return hipOccupancyMaxPotentialBlockSize(gridSize, blockSize, reinterpret_cast<const void*>(f),dynSharedMemPerBlk,blockSizeLimit);
}
#endif // defined(__clang__) && defined(__HIP__)
 
template <typename T>
hipError_t hipGetSymbolAddress(void** devPtr, const T &symbol) {
  return ::hipGetSymbolAddress(devPtr, (const void *)&symbol);
}
template <typename T>
hipError_t hipGetSymbolSize(size_t* size, const T &symbol) {
  return ::hipGetSymbolSize(size, (const void *)&symbol);
}
 
template <typename T>
hipError_t hipMemcpyToSymbol(const T& symbol, const void* src, size_t sizeBytes,
                             size_t offset __dparm(0),
                             hipMemcpyKind kind __dparm(hipMemcpyHostToDevice)) {
  return ::hipMemcpyToSymbol((const void*)&symbol, src, sizeBytes, offset, kind);
}
template <typename T>
hipError_t hipMemcpyToSymbolAsync(const T& symbol, const void* src, size_t sizeBytes, size_t offset,
                                  hipMemcpyKind kind, hipStream_t stream __dparm(0)) {
  return ::hipMemcpyToSymbolAsync((const void*)&symbol, src, sizeBytes, offset, kind, stream);
}
template <typename T>
hipError_t hipMemcpyFromSymbol(void* dst, const T &symbol,
                               size_t sizeBytes, size_t offset __dparm(0),
                               hipMemcpyKind kind __dparm(hipMemcpyDeviceToHost)) {
  return ::hipMemcpyFromSymbol(dst, (const void*)&symbol, sizeBytes, offset, kind);
}
template <typename T>
hipError_t hipMemcpyFromSymbolAsync(void* dst, const T& symbol, size_t sizeBytes, size_t offset,
                                    hipMemcpyKind kind, hipStream_t stream __dparm(0)) {
  return ::hipMemcpyFromSymbolAsync(dst, (const void*)&symbol, sizeBytes, offset, kind, stream);
}
 
template <class T>
inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
    int* numBlocks, T f, int blockSize, size_t dynSharedMemPerBlk) {
    return hipOccupancyMaxActiveBlocksPerMultiprocessor(
        numBlocks, reinterpret_cast<const void*>(f), blockSize, dynSharedMemPerBlk);
}
template <class T>
inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
    int* numBlocks, T f, int blockSize, size_t dynSharedMemPerBlk, unsigned int flags) {
    return hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
        numBlocks, reinterpret_cast<const void*>(f), blockSize, dynSharedMemPerBlk, flags);
}
template<typename UnaryFunction, class T>
static hipError_t __host__ inline hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(
    int*          min_grid_size,
    int*          block_size,
    T             func,
    UnaryFunction block_size_to_dynamic_smem_size,
    int           block_size_limit = 0,
    unsigned int  flags = 0) {
  if (min_grid_size == nullptr || block_size == nullptr ||
      reinterpret_cast<const void*>(func) == nullptr) {
    return hipErrorInvalidValue;
  }
 
  int dev;
  hipError_t status;
  if ((status = hipGetDevice(&dev)) != hipSuccess) {
    return status;
  }
 
  int max_threads_per_cu;
  if ((status = hipDeviceGetAttribute(&max_threads_per_cu,
      hipDeviceAttributeMaxThreadsPerMultiProcessor, dev)) != hipSuccess) {
    return status;
  }
 
  int warp_size;
  if ((status = hipDeviceGetAttribute(&warp_size,
      hipDeviceAttributeWarpSize, dev)) != hipSuccess) {
    return status;
  }
 
  int max_cu_count;
  if ((status = hipDeviceGetAttribute(&max_cu_count,
      hipDeviceAttributeMultiprocessorCount, dev)) != hipSuccess) {
    return status;
  }
 
  struct hipFuncAttributes attr;
  if ((status = hipFuncGetAttributes(&attr, reinterpret_cast<const void*>(func))) != hipSuccess) {
    return status;
  }
 
  // Initial limits for the execution
  const int func_max_threads_per_block = attr.maxThreadsPerBlock;
  if (block_size_limit == 0) {
    block_size_limit = func_max_threads_per_block;
  }
 
  if (func_max_threads_per_block < block_size_limit) {
    block_size_limit = func_max_threads_per_block;
  }
 
  const int block_size_limit_aligned =
    ((block_size_limit + (warp_size - 1)) / warp_size) * warp_size;
 
  // For maximum search
  int max_threads = 0;
  int max_block_size{};
  int max_num_blocks{};
  for (int block_size_check_aligned = block_size_limit_aligned;
       block_size_check_aligned > 0;
       block_size_check_aligned -= warp_size) {
    // Make sure the logic uses the requested limit and not aligned
    int block_size_check = (block_size_limit < block_size_check_aligned) ?
        block_size_limit : block_size_check_aligned;
 
    size_t dyn_smem_size = block_size_to_dynamic_smem_size(block_size_check);
    int optimal_blocks;
    if ((status = hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
        &optimal_blocks, func, block_size_check, dyn_smem_size, flags)) != hipSuccess) {
      return status;
    }
 
    int total_threads = block_size_check * optimal_blocks;
    if (total_threads > max_threads) {
      max_block_size = block_size_check;
      max_num_blocks = optimal_blocks;
      max_threads = total_threads;
    }
 
    // Break if the logic reached possible maximum
    if (max_threads_per_cu == max_threads) {
      break;
    }
  }
 
  // Grid size is the number of blocks per CU * CU count
  *min_grid_size = max_num_blocks * max_cu_count;
  *block_size = max_block_size;
 
  return status;
}
 
template<typename UnaryFunction, class T>
static hipError_t __host__ inline hipOccupancyMaxPotentialBlockSizeVariableSMem(
    int*          min_grid_size,
    int*          block_size,
    T             func,
    UnaryFunction block_size_to_dynamic_smem_size,
    int           block_size_limit = 0)
{
    return hipOccupancyMaxPotentialBlockSizeVariableSMemWithFlags(min_grid_size, block_size, func,
      block_size_to_dynamic_smem_size, block_size_limit);
}
template <typename F>
inline hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize,
                                                    F kernel, size_t dynSharedMemPerBlk, uint32_t blockSizeLimit) {
return hipOccupancyMaxPotentialBlockSize(gridSize, blockSize,(hipFunction_t)kernel, dynSharedMemPerBlk, blockSizeLimit);
}
template <class T>
inline hipError_t hipLaunchCooperativeKernel(T f, dim3 gridDim, dim3 blockDim,
                                             void** kernelParams, unsigned int sharedMemBytes, hipStream_t stream) {
    return hipLaunchCooperativeKernel(reinterpret_cast<const void*>(f), gridDim,
                                      blockDim, kernelParams, sharedMemBytes, stream);
}
template <class T>
inline hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList,
                                                        unsigned int  numDevices, unsigned int  flags = 0) {
    return hipLaunchCooperativeKernelMultiDevice(launchParamsList, numDevices, flags);
}
template <class T>
inline hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
                                                     unsigned int  numDevices, unsigned int  flags = 0) {
    return hipExtLaunchMultiKernelMultiDevice(launchParamsList, numDevices, flags);
}
template <class T, int dim, enum hipTextureReadMode readMode>
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
static inline hipError_t hipBindTexture(size_t* offset, const struct texture<T, dim, readMode>& tex,
                                        const void* devPtr, size_t size = UINT_MAX) {
    return hipBindTexture(offset, &tex, devPtr, &tex.channelDesc, size);
}
template <class T, int dim, enum hipTextureReadMode readMode>
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
static inline hipError_t
    hipBindTexture(size_t* offset, const struct texture<T, dim, readMode>& tex, const void* devPtr,
                   const struct hipChannelFormatDesc& desc, size_t size = UINT_MAX) {
    return hipBindTexture(offset, &tex, devPtr, &desc, size);
}
template<class T, int dim, enum hipTextureReadMode readMode>
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
static inline hipError_t hipBindTexture2D(
    size_t *offset,
    const struct texture<T, dim, readMode> &tex,
    const void *devPtr,
    size_t width,
    size_t height,
    size_t pitch)
{
    return hipBindTexture2D(offset, &tex, devPtr, &tex.channelDesc, width, height, pitch);
}
template<class T, int dim, enum hipTextureReadMode readMode>
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
static inline hipError_t hipBindTexture2D(
  size_t *offset,
  const struct texture<T, dim, readMode> &tex,
  const void *devPtr,
  const struct hipChannelFormatDesc &desc,
  size_t width,
  size_t height,
  size_t pitch)
{
  return hipBindTexture2D(offset, &tex, devPtr, &desc, width, height, pitch);
}
template<class T, int dim, enum hipTextureReadMode readMode>
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
static inline hipError_t hipBindTextureToArray(
    const struct texture<T, dim, readMode> &tex,
    hipArray_const_t array)
{
    struct hipChannelFormatDesc desc;
    hipError_t err = hipGetChannelDesc(&desc, array);
    return (err == hipSuccess) ? hipBindTextureToArray(&tex, array, &desc) : err;
}
template<class T, int dim, enum hipTextureReadMode readMode>
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
static inline hipError_t hipBindTextureToArray(
    const struct texture<T, dim, readMode> &tex,
    hipArray_const_t array,
    const struct hipChannelFormatDesc &desc)
{
    return hipBindTextureToArray(&tex, array, &desc);
}
template<class T, int dim, enum hipTextureReadMode readMode>
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
static inline hipError_t hipBindTextureToMipmappedArray(
    const struct texture<T, dim, readMode> &tex,
    hipMipmappedArray_const_t mipmappedArray)
{
    struct hipChannelFormatDesc desc;
    hipArray_t levelArray;
    hipError_t err = hipGetMipmappedArrayLevel(&levelArray, mipmappedArray, 0);
    if (err != hipSuccess) {
        return err;
    }
    err = hipGetChannelDesc(&desc, levelArray);
    return (err == hipSuccess) ? hipBindTextureToMipmappedArray(&tex, mipmappedArray, &desc) : err;
}
template<class T, int dim, enum hipTextureReadMode readMode>
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
static inline hipError_t hipBindTextureToMipmappedArray(
    const struct texture<T, dim, readMode> &tex,
    hipMipmappedArray_const_t mipmappedArray,
    const struct hipChannelFormatDesc &desc)
{
    return hipBindTextureToMipmappedArray(&tex, mipmappedArray, &desc);
}
template<class T, int dim, enum hipTextureReadMode readMode>
HIP_DEPRECATED(HIP_DEPRECATED_MSG)
static inline hipError_t hipUnbindTexture(
    const struct texture<T, dim, readMode> &tex)
{
    return hipUnbindTexture(&tex);
}
static inline hipError_t hipMallocAsync(
  void**        dev_ptr,
  size_t        size,
  hipMemPool_t  mem_pool,
  hipStream_t   stream) {
  return hipMallocFromPoolAsync(dev_ptr, size, mem_pool, stream);
}
template<class T>
static inline hipError_t hipMallocAsync(
  T**           dev_ptr,
  size_t        size,
  hipMemPool_t  mem_pool,
  hipStream_t   stream) {
  return hipMallocFromPoolAsync(reinterpret_cast<void**>(dev_ptr), size, mem_pool, stream);
}
template<class T>
static inline hipError_t hipMallocAsync(
  T**           dev_ptr,
  size_t        size,
  hipStream_t   stream) {
  return hipMallocAsync(reinterpret_cast<void**>(dev_ptr), size, stream);
}
template<class T>
static inline hipError_t hipMallocFromPoolAsync(
  T**           dev_ptr,
  size_t        size,
  hipMemPool_t  mem_pool,
  hipStream_t   stream) {
  return hipMallocFromPoolAsync(reinterpret_cast<void**>(dev_ptr), size, mem_pool, stream);
}
template <typename... KernelArgs, typename... Params>
static inline __host__ hipError_t hipLaunchKernelEx(const hipLaunchConfig_t* config,
                                                    void (*kernel)(KernelArgs...),
                                                    Params&&... args) {
  return [&](KernelArgs... convertedArgs) {
    void* pArgs[] = {&convertedArgs...};
    return ::hipLaunchKernelExC(config, reinterpret_cast<void*>(kernel), pArgs);
  }(std::forward<Params>(args)...);
}
#endif // __cplusplus
 
#ifdef __GNUC__
#pragma GCC visibility pop
#endif
 
 
#elif !defined(__HIP_PLATFORM_AMD__) && defined(__HIP_PLATFORM_NVIDIA__)
#include "hip/nvidia_detail/nvidia_hip_runtime_api.h"
#else
#error("Must define exactly one of __HIP_PLATFORM_AMD__ or __HIP_PLATFORM_NVIDIA__");
#endif
 
 
#if defined(__cplusplus) && !defined(__HIP_DISABLE_CPP_FUNCTIONS__)
template <class T>
static inline hipError_t hipMalloc(T** devPtr, size_t size) {
    return hipMalloc((void**)devPtr, size);
}
template <class T>
static inline hipError_t hipMallocPitch(T** devPtr, size_t* pitch, size_t width, size_t height) {
    return hipMallocPitch((void**)devPtr, pitch, width, height);
}
template <class T>
static inline hipError_t hipHostMalloc(T** ptr, size_t size,
                                       unsigned int flags = hipHostMallocDefault) {
    return hipHostMalloc((void**)ptr, size, flags);
}
template <class T>
static inline hipError_t hipHostAlloc(T** ptr, size_t size,
                                       unsigned int flags = hipHostAllocDefault) {
    return hipHostAlloc((void**)ptr, size, flags);
}
template <class T>
static inline hipError_t hipMallocManaged(T** devPtr, size_t size,
                                       unsigned int flags = hipMemAttachGlobal) {
    return hipMallocManaged((void**)devPtr, size, flags);
}
 
 
#endif
#endif
// doxygen end HIP API
#include <hip/amd_detail/amd_hip_runtime_pt_api.h>
 
#if USE_PROF_API
#include <hip/amd_detail/hip_prof_str.h>
#endif