hip_runtime_api.h

Go to the documentation of this 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


#include <string.h>  // for getDeviceProp
#include <hip/hip_version.h>
#include <hip/hip_common.h>

enum {
    HIP_SUCCESS = 0,
    HIP_ERROR_INVALID_VALUE,
    HIP_ERROR_NOT_INITIALIZED,
    HIP_ERROR_LAUNCH_OUT_OF_RESOURCES
};

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 HCC paths:

typedef struct hipDeviceProp_t {
    char name[256];            
    size_t totalGlobalMem;     
    size_t sharedMemPerBlock;  
    int regsPerBlock;          
    int warpSize;              
    int maxThreadsPerBlock;    
    int maxThreadsDim[3];      
    int maxGridSize[3];        
    int clockRate;             
    int memoryClockRate;       
    int memoryBusWidth;        
    size_t totalConstMem;      
    int major;  
    int minor;  
    int multiProcessorCount;          
    int l2CacheSize;                  
    int maxThreadsPerMultiProcessor;  
    int computeMode;                  
    int clockInstructionRate;  
    hipDeviceArch_t arch;      
    int concurrentKernels;     
    int pciDomainID;           
    int pciBusID;              
    int pciDeviceID;           
    size_t maxSharedMemoryPerMultiProcessor;  
    int isMultiGpuBoard;                      
    int canMapHostMemory;                     
    int gcnArch;                              
    char gcnArchName[256];                    
    int integrated;            
    int cooperativeLaunch;            
    int cooperativeMultiDeviceLaunch; 
    int maxTexture1DLinear;    
    int maxTexture1D;          
    int maxTexture2D[2];       
    int maxTexture3D[3];       
    unsigned int* hdpMemFlushCntl;      
    unsigned int* hdpRegFlushCntl;      
    size_t memPitch;                 
    size_t textureAlignment;         
    size_t texturePitchAlignment;    
    int kernelExecTimeoutEnabled;    
    int ECCEnabled;                  
    int tccDriver;                   
    int cooperativeMultiDeviceUnmatchedFunc;        
    int cooperativeMultiDeviceUnmatchedGridDim;     
    int cooperativeMultiDeviceUnmatchedBlockDim;    
    int cooperativeMultiDeviceUnmatchedSharedMem;   
    int isLargeBar;                  
    int asicRevision;                
    int managedMemory;               
    int directManagedMemAccessFromHost; 
    int concurrentManagedAccess;     
    int pageableMemoryAccess;        
    int pageableMemoryAccessUsesHostPageTables; 
} hipDeviceProp_t;


 /*
 * @brief HIP Memory type (for pointer attributes)
 * @enum
 * @ingroup Enumerations
 */
typedef enum hipMemoryType {
    hipMemoryTypeHost = 0,    
    hipMemoryTypeDevice = 1,  
    hipMemoryTypeArray = 2,   
    hipMemoryTypeUnified = 3, 
    hipMemoryTypeManaged = 4  
} hipMemoryType;

typedef struct hipPointerAttribute_t {
  union {
      // Deprecated, use instead type
      enum hipMemoryType memoryType;
      enum hipMemoryType type;
    };
    int device;
    void* devicePointer;
    void* hostPointer;
    int isManaged;
    unsigned allocationFlags; /* flags specified when memory was allocated*/
    /* peers? */
} hipPointerAttribute_t;


// hack to get these to show up in Doxygen:
// 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

/*
 * @brief hipError_t
 * @enum
 * @ingroup Enumerations
 */
// Developer note - when updating these, update the hipErrorName and hipErrorString functions in
// NVCC and HCC 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,  
    hipErrorUnknown = 999,  //< Unknown error.
    // HSA Runtime Error Codes start here.
    hipErrorRuntimeMemory = 1052,  
    hipErrorRuntimeOther = 1053,  
    hipErrorTbd  
} hipError_t;

#undef __HIP_NODISCARD

/*
 * @brief hipDeviceAttribute_t
 * @enum
 * @ingroup Enumerations
 */
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,              
    hipDeviceAttributeName,                             
    hipDeviceAttributePageableMemoryAccess,             
    hipDeviceAttributePageableMemoryAccessUsesHostPageTables, 
    hipDeviceAttributePciBusId,                         
    hipDeviceAttributePciDeviceId,                      
    hipDeviceAttributePciDomainID,                      
    hipDeviceAttributePersistingL2CacheMaxSize,         
    hipDeviceAttributeMaxRegistersPerBlock,             
    hipDeviceAttributeMaxRegistersPerMultiprocessor,    
    hipDeviceAttributeReservedSharedMemPerBlock,        
    hipDeviceAttributeMaxSharedMemoryPerBlock,          
    hipDeviceAttributeSharedMemPerBlockOptin,           
    hipDeviceAttributeSharedMemPerMultiprocessor,       
    hipDeviceAttributeSingleToDoublePrecisionPerfRatio, 
    hipDeviceAttributeStreamPrioritiesSupported,        
    hipDeviceAttributeSurfaceAlignment,                 
    hipDeviceAttributeTccDriver,                        
    hipDeviceAttributeTextureAlignment,                 
    hipDeviceAttributeTexturePitchAlignment,            
    hipDeviceAttributeTotalConstantMemory,              
    hipDeviceAttributeTotalGlobalMem,                   
    hipDeviceAttributeUnifiedAddressing,                
    hipDeviceAttributeUuid,                             
    hipDeviceAttributeWarpSize,                         
    hipDeviceAttributeMemoryPoolsSupported,             
    hipDeviceAttributeVirtualMemoryManagementSupported, 

    hipDeviceAttributeCudaCompatibleEnd = 9999,
    hipDeviceAttributeAmdSpecificBegin = 10000,

    hipDeviceAttributeClockInstructionRate = hipDeviceAttributeAmdSpecificBegin,  
    hipDeviceAttributeArch,                                     
    hipDeviceAttributeMaxSharedMemoryPerMultiprocessor,         
    hipDeviceAttributeGcnArch,                                  
    hipDeviceAttributeGcnArchName,                              
    hipDeviceAttributeHdpMemFlushCntl,                          
    hipDeviceAttributeHdpRegFlushCntl,                          
    hipDeviceAttributeCooperativeMultiDeviceUnmatchedFunc,      
    hipDeviceAttributeCooperativeMultiDeviceUnmatchedGridDim,   
    hipDeviceAttributeCooperativeMultiDeviceUnmatchedBlockDim,  
    hipDeviceAttributeCooperativeMultiDeviceUnmatchedSharedMem, 
    hipDeviceAttributeIsLargeBar,                               
    hipDeviceAttributeAsicRevision,                             
    hipDeviceAttributeCanUseStreamWaitValue,                    
    hipDeviceAttributeImageSupport,                             
    hipDeviceAttributePhysicalMultiProcessorCount,              
    hipDeviceAttributeFineGrainSupport,                         
    hipDeviceAttributeWallClockRate,                            

    hipDeviceAttributeAmdSpecificEnd = 19999,
    hipDeviceAttributeVendorSpecificBegin = 20000,
    // Extended attributes for vendors
} hipDeviceAttribute_t;

enum hipComputeMode {
    hipComputeModeDefault = 0,
    hipComputeModeExclusive = 1,
    hipComputeModeProhibited = 2,
    hipComputeModeExclusiveProcess = 3
};

#if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))

#include <stdint.h>
#include <stddef.h>
#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 DEPRECATED(msg) __declspec(deprecated(msg))
#else // !defined(_MSC_VER)
#define DEPRECATED(msg) __attribute__ ((deprecated(msg)))
#endif // !defined(_MSC_VER)
#define DEPRECATED_MSG "This API is marked as deprecated and may not be supported in future releases. For more details please refer https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_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 0
#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 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,        // limit device stack size
    hipLimitPrintfFifoSize = 0x01,  // limit printf fifo size
    hipLimitMallocHeapSize = 0x02,  // limit heap size
    hipLimitRange                   // supported limit range
};
//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

#define hipEventDisableSystemFence 0x20000000

#define hipEventReleaseToDevice  0x40000000

#define hipEventReleaseToSystem  0x80000000

//Flags that can be used with hipHostMalloc.
#define hipHostMallocDefault 0x0

#define hipHostMallocPortable 0x1

#define hipHostMallocMapped  0x2

#define hipHostMallocWriteCombined 0x4

#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

//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 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
// Stream per thread
#define hipStreamPerThread ((hipStream_t)2)

// Indicates that the external memory object is a dedicated resource
#define hipExternalMemoryDedicated 0x1
/*
 * @brief HIP Memory Advise values
 * @enum
 * @ingroup Enumerations
 */
typedef enum hipMemoryAdvise {
    hipMemAdviseSetReadMostly = 1,          
    hipMemAdviseUnsetReadMostly = 2,        
    hipMemAdviseSetPreferredLocation = 3,   
    hipMemAdviseUnsetPreferredLocation = 4, 
    hipMemAdviseSetAccessedBy = 5,          
    hipMemAdviseUnsetAccessedBy = 6,        
    hipMemAdviseSetCoarseGrain = 100,       
    hipMemAdviseUnsetCoarseGrain = 101      
} hipMemoryAdvise;
/*
 * @brief HIP Coherency Mode
 * @enum
 * @ingroup Enumerations
 */
typedef enum hipMemRangeCoherencyMode {
    hipMemRangeCoherencyModeFineGrain = 0,      
    hipMemRangeCoherencyModeCoarseGrain = 1,    
    hipMemRangeCoherencyModeIndeterminate = 2   
} hipMemRangeCoherencyMode;
/*
 * @brief HIP range attributes
 * @enum
 * @ingroup Enumerations
 */
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,
    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;
    unsigned char               reserved[64]; 
} hipMemPoolProps;
typedef struct hipMemPoolPtrExportData {
    unsigned char reserved[64];
} hipMemPoolPtrExportData;

/*
 * @brief hipJitOption
 * @enum
 * @ingroup Enumerations
 */
typedef enum hipJitOption {
    hipJitOptionMaxRegisters = 0,
    hipJitOptionThreadsPerBlock,
    hipJitOptionWallTime,
    hipJitOptionInfoLogBuffer,
    hipJitOptionInfoLogBufferSizeBytes,
    hipJitOptionErrorLogBuffer,
    hipJitOptionErrorLogBufferSizeBytes,
    hipJitOptionOptimizationLevel,
    hipJitOptionTargetFromContext,
    hipJitOptionTarget,
    hipJitOptionFallbackStrategy,
    hipJitOptionGenerateDebugInfo,
    hipJitOptionLogVerbose,
    hipJitOptionGenerateLineInfo,
    hipJitOptionCacheMode,
    hipJitOptionSm3xOpt,
    hipJitOptionFastCompile,
    hipJitOptionNumOptions
} hipJitOption;
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,
} hipExternalMemoryHandleType;
typedef struct hipExternalMemoryHandleDesc_st {
  hipExternalMemoryHandleType type;
  union {
    int fd;
    struct {
      void *handle;
      const void *name;
    } win32;
  } handle;
  unsigned long long size;
  unsigned int flags;
} hipExternalMemoryHandleDesc;
typedef struct hipExternalMemoryBufferDesc_st {
  unsigned long long offset;
  unsigned long long size;
  unsigned int flags;
} hipExternalMemoryBufferDesc;
typedef void* hipExternalMemory_t;
typedef enum hipExternalSemaphoreHandleType_enum {
  hipExternalSemaphoreHandleTypeOpaqueFd = 1,
  hipExternalSemaphoreHandleTypeOpaqueWin32 = 2,
  hipExternalSemaphoreHandleTypeOpaqueWin32Kmt = 3,
  hipExternalSemaphoreHandleTypeD3D12Fence = 4
} hipExternalSemaphoreHandleType;
typedef struct hipExternalSemaphoreHandleDesc_st {
  hipExternalSemaphoreHandleType type;
  union {
    int fd;
    struct {
      void* handle;
      const void* name;
    } win32;
  } handle;
  unsigned int flags;
} hipExternalSemaphoreHandleDesc;
typedef void* hipExternalSemaphore_t;
typedef struct hipExternalSemaphoreSignalParams_st {
  struct {
    struct {
      unsigned long long value;
    } fence;
    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;
    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

/*
    * @brief HIP Devices used by current OpenGL Context.
    * @enum
    * @ingroup Enumerations
    */
typedef enum hipGLDeviceList {
    hipGLDeviceListAll = 1,           
    hipGLDeviceListCurrentFrame = 2,  
    hipGLDeviceListNextFrame = 3      
} hipGLDeviceList;

/*
    * @brief HIP Access falgs for Interop resources.
    * @enum
    * @ingroup Enumerations
    */
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,    
  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 hipKernelNodeAttrID {
    hipKernelNodeAttributeAccessPolicyWindow = 1,
    hipKernelNodeAttributeCooperative = 2,
} hipKernelNodeAttrID;
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 union hipKernelNodeAttrValue {
    hipAccessPolicyWindow accessPolicyWindow;
    int cooperative;
} hipKernelNodeAttrValue;

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 struct hipMemAllocationProp {
    hipMemAllocationType type;                       
    hipMemAllocationHandleType requestedHandleType;  
    hipMemLocation location;                         
    void* win32HandleMetaData;                       
    struct {
        unsigned char compressionType;               
        unsigned char gpuDirectRDMACapable;          
        unsigned short usage;                        
    } allocFlags;
} hipMemAllocationProp;

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;
// Doxygen end group GlobalDefs
//-------------------------------------------------------------------------------------------------
// The handle allows the async commands to use the stream even if the parent hipStream_t goes
// out-of-scope.
// typedef class ihipStream_t * hipStream_t;
/*
 * Opaque structure allows the true event (pointed at by the handle) to remain "live" even if the
 * surrounding hipEvent_t goes out-of-scope. This is handy for cases where the hipEvent_t goes
 * out-of-scope but the true event is being written by some async queue or device */
// typedef struct hipEvent_t {
//    struct ihipEvent_t *_handle;
//} hipEvent_t;
// 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 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 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 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);
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);
// end doxygen Stream Memory Operations
hipError_t hipEventCreateWithFlags(hipEvent_t* event, unsigned flags);
hipError_t hipEventCreate(hipEvent_t* event);
#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);
// end of external resource
hipError_t hipMalloc(void** ptr, size_t size);
hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags);
DEPRECATED("use hipHostMalloc instead")
hipError_t hipMallocHost(void** ptr, size_t size);
DEPRECATED("use hipHostMalloc instead")
hipError_t hipMemAllocHost(void** ptr, size_t size);
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
DEPRECATED("use hipHostMalloc instead")
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);
DEPRECATED("use hipHostFree instead")
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, 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 hipMemcpyHtoDAsync(hipDeviceptr_t dst, 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 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 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** array, const hipChannelFormatDesc* desc, size_t width,
                          size_t height __dparm(0), unsigned int flags __dparm(hipArrayDefault));
hipError_t hipArrayCreate(hipArray** pHandle, const HIP_ARRAY_DESCRIPTOR* pAllocateArray);
hipError_t hipArrayDestroy(hipArray* array);
hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAllocateArray);
hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent);
hipError_t hipFreeArray(hipArray* array);
hipError_t hipFreeMipmappedArray(hipMipmappedArray_t mipmappedArray);
hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* desc,
                            struct hipExtent extent, unsigned int flags);
hipError_t hipMallocMipmappedArray(
    hipMipmappedArray_t *mipmappedArray,
    const struct hipChannelFormatDesc* desc,
    struct hipExtent extent,
    unsigned int numLevels,
    unsigned int flags __dparm(0));
hipError_t hipGetMipmappedArrayLevel(
    hipArray_t *levelArray,
    hipMipmappedArray_const_t mipmappedArray,
    unsigned int level);
hipError_t hipArrayGetInfo(hipChannelFormatDesc* desc, hipExtent* extent, unsigned int* flags,
                           hipArray* array);
hipError_t hipArrayGetDescriptor(HIP_ARRAY_DESCRIPTOR* pArrayDescriptor, hipArray* array);
hipError_t hipArray3DGetDescriptor(HIP_ARRAY3D_DESCRIPTOR* pArrayDescriptor, hipArray* 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* 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* 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));
DEPRECATED(DEPRECATED_MSG)
hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src,
                            size_t count, hipMemcpyKind kind);
DEPRECATED(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* srcArray, size_t srcOffset, size_t count);
hipError_t hipMemcpyHtoA(hipArray* 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);
// 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 hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr);
#ifndef USE_PEER_NON_UNIFIED
#define USE_PEER_NON_UNIFIED 1
#endif
#if USE_PEER_NON_UNIFIED == 1

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));
#endif
// doxygen end PeerToPeer
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxDestroy(hipCtx_t ctx);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxPopCurrent(hipCtx_t* ctx);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxPushCurrent(hipCtx_t ctx);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxSetCurrent(hipCtx_t ctx);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxGetCurrent(hipCtx_t* ctx);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxGetDevice(hipDevice_t* device);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxSynchronize(void);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxGetFlags(unsigned int* flags);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx);
// doxygen end Context deprecated
hipError_t hipDevicePrimaryCtxGetState(hipDevice_t dev, unsigned int* flags, int* active);
hipError_t hipDevicePrimaryCtxRelease(hipDevice_t dev);
hipError_t hipDevicePrimaryCtxRetain(hipCtx_t* pctx, hipDevice_t dev);
hipError_t hipDevicePrimaryCtxReset(hipDevice_t dev);
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 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 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);
hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
                                              int  numDevices, unsigned int  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:
DEPRECATED("use roctracer/rocTX instead")
hipError_t hipProfilerStart();
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 hipBindTextureToMipmappedArray(
    const textureReference* tex,
    hipMipmappedArray_const_t mipmappedArray,
    const hipChannelFormatDesc* desc);

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);

DEPRECATED(DEPRECATED_MSG)
 hipError_t hipGetTextureReference(
    const textureReference** texref,
    const void* symbol);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetAddressMode(
    textureReference* texRef,
    int dim,
    enum hipTextureAddressMode am);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetArray(
    textureReference* tex,
    hipArray_const_t array,
    unsigned int flags);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetFilterMode(
    textureReference* texRef,
    enum hipTextureFilterMode fm);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetFlags(
    textureReference* texRef,
    unsigned int Flags);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetFormat(
    textureReference* texRef,
    hipArray_Format fmt,
    int NumPackedComponents);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipBindTexture(
    size_t* offset,
    const textureReference* tex,
    const void* devPtr,
    const hipChannelFormatDesc* desc,
    size_t size __dparm(UINT_MAX));
DEPRECATED(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);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipBindTextureToArray(
    const textureReference* tex,
    hipArray_const_t array,
    const hipChannelFormatDesc* desc);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipGetTextureAlignmentOffset(
    size_t* offset,
    const textureReference* texref);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipUnbindTexture(const textureReference* tex);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefGetAddress(
    hipDeviceptr_t* dev_ptr,
    const textureReference* texRef);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefGetAddressMode(
    enum hipTextureAddressMode* pam,
    const textureReference* texRef,
    int dim);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefGetFilterMode(
    enum hipTextureFilterMode* pfm,
    const textureReference* texRef);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefGetFlags(
    unsigned int* pFlags,
    const textureReference* texRef);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefGetFormat(
    hipArray_Format* pFormat,
    int* pNumChannels,
    const textureReference* texRef);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefGetMaxAnisotropy(
    int* pmaxAnsio,
    const textureReference* texRef);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefGetMipmapFilterMode(
    enum hipTextureFilterMode* pfm,
    const textureReference* texRef);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefGetMipmapLevelBias(
    float* pbias,
    const textureReference* texRef);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefGetMipmapLevelClamp(
    float* pminMipmapLevelClamp,
    float* pmaxMipmapLevelClamp,
    const textureReference* texRef);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefGetMipMappedArray(
    hipMipmappedArray_t* pArray,
    const textureReference* texRef);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetAddress(
    size_t* ByteOffset,
    textureReference* texRef,
    hipDeviceptr_t dptr,
    size_t bytes);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetAddress2D(
    textureReference* texRef,
    const HIP_ARRAY_DESCRIPTOR* desc,
    hipDeviceptr_t dptr,
    size_t Pitch);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetMaxAnisotropy(
    textureReference* texRef,
    unsigned int maxAniso);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetBorderColor(
    textureReference* texRef,
    float* pBorderColor);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetMipmapFilterMode(
    textureReference* texRef,
    enum hipTextureFilterMode fm);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetMipmapLevelBias(
    textureReference* texRef,
    float bias);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetMipmapLevelClamp(
    textureReference* texRef,
    float minMipMapLevelClamp,
    float maxMipMapLevelClamp);
DEPRECATED(DEPRECATED_MSG)
hipError_t hipTexRefSetMipmappedArray(
    textureReference* texRef,
    struct hipMipmappedArray* mipmappedArray,
    unsigned int Flags);

// doxygen end deprecated texture management
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);

// 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 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 hipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream);

hipError_t hipGraphUpload(hipGraphExec_t graphExec, hipStream_t stream);

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 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);

// 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
typedef unsigned int GLuint;
typedef unsigned int GLenum;

hipError_t hipGLGetDevices(unsigned int* pHipDeviceCount, int* pHipDevices,
                           unsigned int hipDeviceCount, hipGLDeviceList deviceList);
hipError_t hipGraphicsGLRegisterBuffer(hipGraphicsResource** resource, GLuint buffer,
                                       unsigned int flags);
hipError_t hipGraphicsGLRegisterImage(hipGraphicsResource** resource, GLuint image,
                                      GLenum target, unsigned int flags);
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 GL Interop
#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);
}

hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, const hipResourceDesc* pResDesc);
hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject);
// end of surface
template <class T, int dim, enum hipTextureReadMode readMode>
DEPRECATED(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>
DEPRECATED(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>
DEPRECATED(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>
DEPRECATED(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>
DEPRECATED(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>
DEPRECATED(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>
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>
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>
DEPRECATED(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);
}
#endif // __cplusplus

#ifdef __GNUC__
#pragma GCC visibility pop
#endif


#elif !(defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && (defined(__HIP_PLATFORM_NVCC__) || 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 hipHostMalloc(T** ptr, size_t size,
                                       unsigned int flags = hipHostMallocDefault) {
    return hipHostMalloc((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