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

SWDEV-299127 - Merge 'develop' into 'amd-staging'

Browse Source 
Change-Id: I0862a538ce6c78294201437108649761de636406


[ROCm/hip commit: 38912f4ec5]
This commit is contained in:
Maneesh Gupta
2023-07-06 09:07:38 +00:00
parent a1dc3c51b6 fce17f9804
commit 9b81b1f256
7 changed files with 246 additions and 255 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/cmake/FindHIP.cmake
+20 -9
View File
@@ -52,17 +52,28 @@ if(NOT HIP_CXX_COMPILER)
set(HIP_CXX_COMPILER ${CMAKE_CXX_COMPILER})
endif()
# Approach: To find HIP_CLANG_PATH using HIP_CXX_COMPILER options
# Using --version option of HIP_CXX_COMPILER get the Install Directory
# This install directory is set as the HIP_CLANG_PATH.
# IF not successful in getting the Install Directory using HIP_CXX_COMPILER
# Approach: To find HIP_CLANG_PATH for HIP_CXX_COMPILER types
# For HIP_CXX_COMPILER as *hipcc use hipconfig -l option to get the clang path.
# For HIP_CXX_COMPILER as *clang use real path of HIP_CXX_COMPILER
# IF not successful in getting the CLANG_PATH using HIP_CXX_COMPILER
# fallback to Old Method to find HIP_CLANG_PATH from ENV Vars, ROCMPATH, HIPPATH etc.
execute_process(COMMAND ${HIP_CXX_COMPILER} --version
OUTPUT_STRIP_TRAILING_WHITESPACE
OUTPUT_VARIABLE HIP_CXX_COMPILER_VERSION_OUTPUT)
if(HIP_CXX_COMPILER MATCHES ".*hipcc")
if(HIP_CXX_COMPILER_VERSION_OUTPUT MATCHES "InstalledDir:[ \t]*([^\n]*)")
get_filename_component(HIP_CLANG_INSTALL_DIR "${CMAKE_MATCH_1}" DIRECTORY)
get_filename_component(HIPCC_PATH "${HIP_CXX_COMPILER}" DIRECTORY)
set(_HIPCONFIG_EXECUTABLE "${HIPCC_PATH}/hipconfig")
execute_process(COMMAND ${_HIPCONFIG_EXECUTABLE} -l
OUTPUT_VARIABLE _HIP_CLANG_INSTALL_PATH
OUTPUT_STRIP_TRAILING_WHITESPACE
ERROR_VARIABLE _HIPCONFIG_EXE_ERROR
ERROR_STRIP_TRAILING_WHITESPACE
RESULT_VARIABLE _HIPCONFIG_EXE_EXIT_CODE)
if( _HIPCONFIG_EXE_ERROR )
message( STATUS "hipconfig -l option failed with error: ${_HIPCONFIG_EXE_ERROR}" )
else() #IF hipconfig -l executed with no error
if( "${_HIPCONFIG_EXE_EXIT_CODE}" STREQUAL "0" )
set(HIP_CLANG_PATH "${_HIP_CLANG_INSTALL_PATH}")
else()
message( STATUS "${_HIPCONFIG_EXECUTABLE} Failed with Exit code: ${_HIPCONFIG_EXE_EXIT_CODE}" )
endif()
endif()
elseif (HIP_CXX_COMPILER MATCHES ".*clang\\+\\+")
get_filename_component(_HIP_CLANG_REAL_PATH "${HIP_CXX_COMPILER}" REALPATH)
projects/hip/include/hip/hip_ext.h
+3 -2
View File
@@ -49,6 +49,9 @@ THE SOFTWARE.
* @param [in] kernelParams pointer to kernel parameters.
* @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and
* must be in the memory layout and alignment expected by the kernel.
* All passed arguments must be naturally aligned according to their type. The memory address of each
* argument should be a multiple of its size in bytes. Please refer to hip_porting_driver_api.md
* for sample usage.
* @param [in] startEvent If non-null, specified event will be updated to track the start time of
* the kernel launch. The event must be created before calling this API.
* @param [in] stopEvent If non-null, specified event will be updated to track the stop time of
@@ -57,8 +60,6 @@ THE SOFTWARE.
* launched in any order.
* @returns #hipSuccess, #hipInvalidDeviceId, #hipErrorNotInitialized, #hipErrorInvalidValue.
*
* @warning kernellParams argument is not yet implemented in HIP, use extra instead.
* Please refer to hip_porting_driver_api.md for sample usage.
* HIP/ROCm actually updates the start event when the associated kernel completes.
* Currently, timing between startEvent and stopEvent does not include the time it takes to perform
* a system scope release/cache flush - only the time it takes to issues writes to cache.
projects/hip/include/hip/hip_runtime_api.h
+122 -170
View File
@@ -41,7 +41,16 @@ enum {
HIP_ERROR_NOT_INITIALIZED,
HIP_ERROR_LAUNCH_OUT_OF_RESOURCES
};
// hack to get these to show up in Doxygen:
/**
* @defgroup GlobalDefs Global enum and defines
* @{
*
*/
/**
* hipDeviceArch_t
*
*/
typedef struct {
// 32-bit Atomics
unsigned hasGlobalInt32Atomics : 1; ///< 32-bit integer atomics for global memory.
@@ -152,11 +161,9 @@ typedef struct hipDeviceProp_t {
int pageableMemoryAccessUsesHostPageTables; ///< Device accesses pageable memory via the host's page tables
} hipDeviceProp_t;
/*
* @brief HIP Memory type (for pointer attributes)
* @enum
* @ingroup Enumerations
/**
* hipMemoryType (for pointer attributes)
*
*/
typedef enum hipMemoryType {
hipMemoryTypeHost = 0, ///< Memory is physically located on host
@@ -186,14 +193,6 @@ typedef struct hipPointerAttribute_t {
/* peers? */
} hipPointerAttribute_t;
// hack to get these to show up in Doxygen:
/**
* @defgroup GlobalDefs Global enum and defines
* @{
*
*/
// Ignoring error-code return values from hip APIs is discouraged. On C++17,
// we can make that yield a warning
#if __cplusplus >= 201703L
@@ -202,10 +201,9 @@ typedef struct hipPointerAttribute_t {
#define __HIP_NODISCARD
#endif
/*
* @brief hipError_t
* @enum
* @ingroup Enumerations
/**
* HIP error type
*
*/
// Developer note - when updating these, update the hipErrorName and hipErrorString functions in
// NVCC and HCC paths Also update the hipCUDAErrorTohipError function in NVCC path.
@@ -214,29 +212,29 @@ typedef enum __HIP_NODISCARD hipError_t {
hipSuccess = 0, ///< Successful completion.
hipErrorInvalidValue = 1, ///< One or more of the parameters passed to the API call is NULL
///< or not in an acceptable range.
hipErrorOutOfMemory = 2,
hipErrorOutOfMemory = 2, ///< out of memory range.
// Deprecated
hipErrorMemoryAllocation = 2, ///< Memory allocation error.
hipErrorNotInitialized = 3,
hipErrorNotInitialized = 3, ///< Invalid not initialized
// Deprecated
hipErrorInitializationError = 3,
hipErrorDeinitialized = 4,
hipErrorDeinitialized = 4, ///< Deinitialized
hipErrorProfilerDisabled = 5,
hipErrorProfilerNotInitialized = 6,
hipErrorProfilerAlreadyStarted = 7,
hipErrorProfilerAlreadyStopped = 8,
hipErrorInvalidConfiguration = 9,
hipErrorInvalidPitchValue = 12,
hipErrorInvalidSymbol = 13,
hipErrorInvalidConfiguration = 9, ///< Invalide configuration
hipErrorInvalidPitchValue = 12, ///< Invalid pitch value
hipErrorInvalidSymbol = 13, ///< Invalid symbol
hipErrorInvalidDevicePointer = 17, ///< Invalid Device Pointer
hipErrorInvalidMemcpyDirection = 21, ///< Invalid memory copy direction
hipErrorInsufficientDriver = 35,
hipErrorMissingConfiguration = 52,
hipErrorPriorLaunchFailure = 53,
hipErrorInvalidDeviceFunction = 98,
hipErrorInvalidDeviceFunction = 98, ///< Invalid device function
hipErrorNoDevice = 100, ///< Call to hipGetDeviceCount returned 0 devices
hipErrorInvalidDevice = 101, ///< DeviceID must be in range 0...#compute-devices.
hipErrorInvalidImage = 200,
hipErrorInvalidDevice = 101, ///< DeviceID must be in range from 0 to compute-devices.
hipErrorInvalidImage = 200, ///< Invalid image
hipErrorInvalidContext = 201, ///< Produced when input context is invalid.
hipErrorContextAlreadyCurrent = 202,
hipErrorMapFailed = 205,
@@ -251,34 +249,34 @@ typedef enum __HIP_NODISCARD hipError_t {
hipErrorNotMappedAsArray = 212,
hipErrorNotMappedAsPointer = 213,
hipErrorECCNotCorrectable = 214,
hipErrorUnsupportedLimit = 215,
hipErrorContextAlreadyInUse = 216,
hipErrorUnsupportedLimit = 215, ///< Unsupported limit
hipErrorContextAlreadyInUse = 216, ///< The context is already in use
hipErrorPeerAccessUnsupported = 217,
hipErrorInvalidKernelFile = 218, ///< In CUDA DRV, it is CUDA_ERROR_INVALID_PTX
hipErrorInvalidGraphicsContext = 219,
hipErrorInvalidSource = 300,
hipErrorFileNotFound = 301,
hipErrorInvalidSource = 300, ///< Invalid source.
hipErrorFileNotFound = 301, ///< the file is not found.
hipErrorSharedObjectSymbolNotFound = 302,
hipErrorSharedObjectInitFailed = 303,
hipErrorOperatingSystem = 304,
hipErrorInvalidHandle = 400,
hipErrorSharedObjectInitFailed = 303, ///< Failed to initialize shared object.
hipErrorOperatingSystem = 304, ///< Not the correct operating system
hipErrorInvalidHandle = 400, ///< Invalide handle
// Deprecated
hipErrorInvalidResourceHandle = 400, ///< Resource handle (hipEvent_t or hipStream_t) invalid.
hipErrorIllegalState = 401, ///< Resource required is not in a valid state to perform operation.
hipErrorNotFound = 500,
hipErrorNotFound = 500, ///< Not found
hipErrorNotReady = 600, ///< Indicates that asynchronous operations enqueued earlier are not
///< ready. This is not actually an error, but is used to distinguish
///< from hipSuccess (which indicates completion). APIs that return
///< this error include hipEventQuery and hipStreamQuery.
hipErrorIllegalAddress = 700,
hipErrorLaunchOutOfResources = 701, ///< Out of resources error.
hipErrorLaunchTimeOut = 702,
hipErrorLaunchTimeOut = 702, ///< Timeout for the launch.
hipErrorPeerAccessAlreadyEnabled =
704, ///< Peer access was already enabled from the current device.
hipErrorPeerAccessNotEnabled =
705, ///< Peer access was never enabled from the current device.
hipErrorSetOnActiveProcess = 708,
hipErrorContextIsDestroyed = 709,
hipErrorSetOnActiveProcess = 708, ///< The process is active.
hipErrorContextIsDestroyed = 709, ///< The context is already destroyed
hipErrorAssert = 710, ///< Produced when the kernel calls assert.
hipErrorHostMemoryAlreadyRegistered =
712, ///< Produced when trying to lock a page-locked memory.
@@ -317,7 +315,7 @@ typedef enum __HIP_NODISCARD hipError_t {
///< not performed because it included changes which
///< violated constraintsspecific to instantiated graph
///< update.
hipErrorUnknown = 999, //< Unknown error.
hipErrorUnknown = 999, ///< Unknown error.
// HSA Runtime Error Codes start here.
hipErrorRuntimeMemory = 1052, ///< HSA runtime memory call returned error. Typically not seen
///< in production systems.
@@ -328,10 +326,9 @@ typedef enum __HIP_NODISCARD hipError_t {
#undef __HIP_NODISCARD
/*
* @brief hipDeviceAttribute_t
* @enum
* @ingroup Enumerations
/**
* hipDeviceAttribute_t
*
*/
typedef enum hipDeviceAttribute_t {
hipDeviceAttributeCudaCompatibleBegin = 0,
@@ -473,10 +470,6 @@ enum hipComputeMode {
hipComputeModeExclusiveProcess = 3
};
/**
* @}
*/
#if (defined(__HIP_PLATFORM_HCC__) || defined(__HIP_PLATFORM_AMD__)) && !(defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_PLATFORM_NVIDIA__))
#include <stdint.h>
@@ -556,14 +549,13 @@ typedef struct hipFuncAttributes {
} 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
hipLimitStackSize = 0x0, ///< limit of stack size in bytes on the current device
hipLimitPrintfFifoSize = 0x01, ///< size limit in bytes of fifo used by printf call on the device
hipLimitMallocHeapSize = 0x02, ///< limit of heap size in bytes on the current device
hipLimitRange ///< supported limit range
};
/**
* @addtogroup GlobalDefs More
* @{
* Flags that can be used with hipStreamCreateWithFlags.
*/
//Flags that can be used with hipStreamCreateWithFlags.
/** Default stream creation flags. These are used with hipStreamCreate().*/
@@ -711,10 +703,10 @@ enum hipLimit_t {
// Indicates that the external memory object is a dedicated resource
#define hipExternalMemoryDedicated 0x1
/*
* @brief HIP Memory Advise values
* @enum
* @ingroup Enumerations
/**
* HIP Memory Advise values
*
* @note This memory advise enumeration is used on Linux, not Windows.
*/
typedef enum hipMemoryAdvise {
hipMemAdviseSetReadMostly = 1, ///< Data will mostly be read and only occassionally
@@ -723,7 +715,7 @@ typedef enum hipMemoryAdvise {
hipMemAdviseSetPreferredLocation = 3, ///< Set the preferred location for the data as
///< the specified device
hipMemAdviseUnsetPreferredLocation = 4, ///< Clear the preferred location for the data
hipMemAdviseSetAccessedBy = 5, ///< Data will be accessed by the specified device,
hipMemAdviseSetAccessedBy = 5, ///< Data will be accessed by the specified device
///< so prevent page faults as much as possible
hipMemAdviseUnsetAccessedBy = 6, ///< Let HIP to decide on the page faulting policy
///< for the specified device
@@ -734,10 +726,8 @@ typedef enum hipMemoryAdvise {
///< boundaries for better performance
hipMemAdviseUnsetCoarseGrain = 101 ///< Restores cache coherency policy back to fine-grain
} hipMemoryAdvise;
/*
* @brief HIP Coherency Mode
* @enum
* @ingroup Enumerations
/**
* HIP Coherency Mode
*/
typedef enum hipMemRangeCoherencyMode {
hipMemRangeCoherencyModeFineGrain = 0, ///< Updates to memory with this attribute can be
@@ -748,10 +738,8 @@ typedef enum hipMemRangeCoherencyMode {
///< both hipMemRangeCoherencyModeFineGrain and
///< hipMemRangeCoherencyModeCoarseGrain attributes
} hipMemRangeCoherencyMode;
/*
* @brief HIP range attributes
* @enum
* @ingroup Enumerations
/**
* HIP range attributes
*/
typedef enum hipMemRangeAttribute {
hipMemRangeAttributeReadMostly = 1, ///< Whether the range will mostly be read and
@@ -766,9 +754,7 @@ typedef enum hipMemRangeAttribute {
} hipMemRangeAttribute;
/**
* @brief HIP memory pool attributes
* @enum
* @ingroup Enumerations
* HIP memory pool attributes
*/
typedef enum hipMemPoolAttr
{
@@ -827,9 +813,7 @@ typedef enum hipMemPoolAttr
hipMemPoolAttrUsedMemHigh = 0x8
} hipMemPoolAttr;
/**
* @brief Specifies the type of location
* @enum
* @ingroup Enumerations
* Specifies the type of location
*/
typedef enum hipMemLocationType {
hipMemLocationTypeInvalid = 0,
@@ -845,9 +829,8 @@ typedef struct hipMemLocation {
int id; ///< Identifier for the provided location type @p hipMemLocationType
} hipMemLocation;
/**
* @brief Specifies the memory protection flags for mapping
* @enum
* @ingroup Enumerations
* Specifies the memory protection flags for mapping
*
*/
typedef enum hipMemAccessFlags {
hipMemAccessFlagsProtNone = 0, ///< Default, make the address range not accessible
@@ -862,9 +845,7 @@ typedef struct hipMemAccessDesc {
hipMemAccessFlags flags; ///< Accessibility flags to set
} hipMemAccessDesc;
/**
* @brief Defines the allocation types
* @enum
* @ingroup Enumerations
* Defines the allocation types
*/
typedef enum hipMemAllocationType {
hipMemAllocationTypeInvalid = 0x0,
@@ -875,9 +856,8 @@ typedef enum hipMemAllocationType {
hipMemAllocationTypeMax = 0x7FFFFFFF
} hipMemAllocationType;
/**
* @brief Flags for specifying handle types for memory pool allocations
* @enum
* @ingroup Enumerations
* Flags for specifying handle types for memory pool allocations
*
*/
typedef enum hipMemAllocationHandleType {
hipMemHandleTypeNone = 0x0, ///< Does not allow any export mechanism
@@ -905,10 +885,8 @@ typedef struct hipMemPoolPtrExportData {
unsigned char reserved[64];
} hipMemPoolPtrExportData;
/*
* @brief hipJitOption
* @enum
* @ingroup Enumerations
/**
* hipJitOption
*/
typedef enum hipJitOption {
hipJitOptionMaxRegisters = 0,
@@ -959,7 +937,6 @@ typedef enum hipSharedMemConfig {
} hipSharedMemConfig;
/**
* Struct for data in 3D
*
*/
typedef struct dim3 {
uint32_t x; ///< x
@@ -969,6 +946,9 @@ typedef struct dim3 {
constexpr __host__ __device__ dim3(uint32_t _x = 1, uint32_t _y = 1, uint32_t _z = 1) : x(_x), y(_y), z(_z){};
#endif
} dim3;
/**
* struct hipLaunchParams_t
*/
typedef struct hipLaunchParams_t {
void* func; ///< Device function symbol
dim3 gridDim; ///< Grid dimentions
@@ -977,6 +957,9 @@ typedef struct hipLaunchParams_t {
size_t sharedMem; ///< Shared memory
hipStream_t stream; ///< Stream identifier
} hipLaunchParams;
/**
* struct hipFunctionLaunchParams_t
*/
typedef struct hipFunctionLaunchParams_t {
hipFunction_t function; ///< Kernel to launch
unsigned int gridDimX; ///< Width(X) of grid in blocks
@@ -1069,16 +1052,13 @@ typedef struct hipExternalSemaphoreWaitParams_st {
/**
* Internal use only. This API may change in the future
* Pre-Compiled header for online compilation
*
*/
void __hipGetPCH(const char** pch, unsigned int*size);
#endif
/*
* @brief HIP Devices used by current OpenGL Context.
* @enum
* @ingroup Enumerations
*/
/**
* HIP Devices used by current OpenGL Context.
*/
typedef enum hipGLDeviceList {
hipGLDeviceListAll = 1, ///< All hip devices used by current OpenGL context.
hipGLDeviceListCurrentFrame = 2, ///< Hip devices used by current OpenGL context in current
@@ -1087,11 +1067,9 @@ typedef enum hipGLDeviceList {
///< frame.
} hipGLDeviceList;
/*
* @brief HIP Access falgs for Interop resources.
* @enum
* @ingroup Enumerations
*/
/**
* HIP Access falgs for Interop resources.
*/
typedef enum hipGraphicsRegisterFlags {
hipGraphicsRegisterFlagsNone = 0,
hipGraphicsRegisterFlagsReadOnly = 1, ///< HIP will not write to this registered resource
@@ -1126,9 +1104,7 @@ typedef struct hipUserObject* hipUserObject_t;
/**
* @brief hipGraphNodeType
* @enum
*
* hipGraphNodeType
*/
typedef enum hipGraphNodeType {
hipGraphNodeTypeKernel = 0, ///< GPU kernel node
@@ -1181,9 +1157,7 @@ typedef struct hipMemAllocNodeParams {
} hipMemAllocNodeParams;
/**
* @brief hipKernelNodeAttrID
* @enum
*
* Kernel node attributeID
*/
typedef enum hipKernelNodeAttrID {
hipKernelNodeAttributeAccessPolicyWindow = 1,
@@ -1207,9 +1181,7 @@ typedef union hipKernelNodeAttrValue {
} hipKernelNodeAttrValue;
/**
* @brief hipGraphExecUpdateResult
* @enum
*
* Graph execution update result
*/
typedef enum hipGraphExecUpdateResult {
hipGraphExecUpdateSuccess = 0x0, ///< The update succeeded
@@ -1307,9 +1279,7 @@ typedef struct hipMemAllocationProp {
typedef struct ihipMemGenericAllocationHandle* hipMemGenericAllocationHandle_t;
/**
* @brief Flags for granularity
* @enum
* @ingroup Enumerations
* Flags for granularity
*/
typedef enum hipMemAllocationGranularity_flags {
hipMemAllocationGranularityMinimum = 0x0, ///< Minimum granularity
@@ -1317,18 +1287,14 @@ typedef enum hipMemAllocationGranularity_flags {
} hipMemAllocationGranularity_flags;
/**
* @brief Memory handle type
* @enum
* @ingroup Enumerations
* Memory handle type
*/
typedef enum hipMemHandleType {
hipMemHandleTypeGeneric = 0x0 ///< Generic handle type
} hipMemHandleType;
/**
* @brief Memory operation types
* @enum
* @ingroup Enumerations
* Memory operation types
*/
typedef enum hipMemOperationType {
hipMemOperationTypeMap = 0x1, ///< Map operation
@@ -1336,9 +1302,7 @@ typedef enum hipMemOperationType {
} hipMemOperationType;
/**
* @brief Subresource types for sparse arrays
* @enum
* @ingroup Enumerations
* Subresource types for sparse arrays
*/
typedef enum hipArraySparseSubresourceType {
hipArraySparseSubresourceTypeSparseLevel = 0x0, ///< Sparse level
@@ -1383,18 +1347,9 @@ typedef struct hipArrayMapInfo {
unsigned int reserved[2]; ///< Reserved for future use, must be zero now.
} 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;
/**
* @}
*/
/**
* @defgroup API HIP API
* @{
@@ -5042,15 +4997,15 @@ hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned
* @param [in] kernelParams
* @param [in] extra Pointer to kernel arguments. These are passed directly to the kernel and
* must be in the memory layout and alignment expected by the kernel.
* All passed arguments must be naturally aligned according to their type. The memory address of each
* argument should be a multiple of its size in bytes. Please refer to hip_porting_driver_api.md
* for sample usage.
*
* Please note, HIP does not support kernel launch with total work items defined in dimension with
* size gridDim x blockDim >= 2^32. So gridDim.x * blockDim.x, gridDim.y * blockDim.y
* and gridDim.z * blockDim.z are always less than 2^32.
*
* @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue
*
* @warning kernellParams argument is not yet implemented in HIP. Please use extra instead. Please
* refer to hip_porting_driver_api.md for sample usage.
*/
hipError_t hipModuleLaunchKernel(hipFunction_t f, unsigned int gridDimX, unsigned int gridDimY,
unsigned int gridDimZ, unsigned int blockDimX,
@@ -7691,7 +7646,7 @@ hipError_t hipGraphicsUnmapResources(int count, hipGraphicsResource_t* resources
* @brief Unregisters a graphics resource.
*
* @param [in] resource - Graphics resources to unregister.
*
*
* @returns #hipSuccess
*
*/
@@ -7701,7 +7656,40 @@ hipError_t hipGraphicsUnregisterResource(hipGraphicsResource_t resource);
* @}
*/
/**
*-------------------------------------------------------------------------------------------------
*-------------------------------------------------------------------------------------------------
* @defgroup Surface Surface Object
* @{
*
* This section describes surface object functions of HIP runtime API.
*
* @note APIs in this section are under development.
*
*/
/**
* @brief Create a surface object.
*
* @param [out] pSurfObject Pointer of surface object to be created.
* @param [in] pResDesc Pointer of suface object descriptor.
*
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, const hipResourceDesc* pResDesc);
/**
* @brief Destroy a surface object.
*
* @param [in] surfaceObject Surface object to be destroyed.
*
* @returns #hipSuccess, #hipErrorInvalidValue
*/
hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject);
// end of surface
/**
* @}
*/
#ifdef __cplusplus
} /* extern "c" */
#endif
@@ -8058,42 +8046,6 @@ inline hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchPara
unsigned int numDevices, unsigned int flags = 0) {
return hipExtLaunchMultiKernelMultiDevice(launchParamsList, numDevices, flags);
}
/**
*-------------------------------------------------------------------------------------------------
*-------------------------------------------------------------------------------------------------
* @defgroup Surface Surface Object
* @{
*
* This section describes surface object functions of HIP runtime API.
*
* @note APIs in this section are under development.
*
*/
/**
* @brief Create a surface object.
*
* @param [out] pSurfObject Pointer of surface object to be created.
* @param [in] pResDesc Pointer of suface object descriptor.
*
* @returns #hipSuccess, #hipErrorInvalidValue
*
*/
hipError_t hipCreateSurfaceObject(hipSurfaceObject_t* pSurfObject, const hipResourceDesc* pResDesc);
/**
* @brief Destroy a surface object.
*
* @param [in] surfaceObject Surface object to be destroyed.
*
* @returns #hipSuccess, #hipErrorInvalidValue
*/
hipError_t hipDestroySurfaceObject(hipSurfaceObject_t surfaceObject);
// end of surface
/**
* @}
*/
/**
* @brief Binds a memory area to a texture.
*
projects/hip/include/hip/hiprtc.h
+23 -20
View File
@@ -96,8 +96,8 @@ typedef enum hiprtcJIT_option {
HIPRTC_JIT_PREC_SQRT, ///< PREC_SQRT
HIPRTC_JIT_FMA, ///< FMA
HIPRTC_JIT_NUM_OPTIONS, ///< Number of options
HIPRTC_JIT_IR_TO_ISA_OPT_EXT = 10000, //< AMD only. Linker options to be passed on to
HIPRTC_JIT_IR_TO_ISA_OPT_COUNT_EXT, //< AMD only. Count of linker options
HIPRTC_JIT_IR_TO_ISA_OPT_EXT = 10000, ///< AMD only. Linker options to be passed on to
HIPRTC_JIT_IR_TO_ISA_OPT_COUNT_EXT, ///< AMD only. Count of linker options
} hiprtcJIT_option;
/**
@@ -131,7 +131,7 @@ typedef struct ihiprtcLinkState* hiprtcLinkState;
* @brief Returns text string message to explain the error which occurred
*
* @param [in] result code to convert to string.
* @return const char pointer to the NULL-terminated error string
* @returns const char pointer to the NULL-terminated error string
*
* @warning In HIP, this function returns the name of the error,
* if the hiprtc result is defined, it will return "Invalid HIPRTC error code"
@@ -147,8 +147,11 @@ const char* hiprtcGetErrorString(hiprtcResult result);
* @param [out] major HIP Runtime Compilation major version.
* @param [out] minor HIP Runtime Compilation minor version.
*
* @returns #HIPRTC_ERROR_INVALID_INPUT, #HIPRTC_SUCCESS
*
*/
hiprtcResult hiprtcVersion(int* major, int* minor);
/**
* hiprtc program
*
@@ -161,7 +164,7 @@ typedef struct _hiprtcProgram* hiprtcProgram;
*
* @param [in] prog runtime compilation program instance.
* @param [in] name_expression const char pointer to the name expression.
* @return #HIPRTC_SUCCESS
* @returns #HIPRTC_SUCCESS
*
* If const char pointer is NULL, it will return #HIPRTC_ERROR_INVALID_INPUT.
*
@@ -171,13 +174,13 @@ hiprtcResult hiprtcAddNameExpression(hiprtcProgram prog,
const char* name_expression);
/**
* @ingroup Runtime
* @ingroup Runtime
* @brief Compiles the given runtime compilation program.
*
* @param [in] prog runtime compilation program instance.
* @param [in] numOptions number of compiler options.
* @param [in] options compiler options as const array of strins.
* @return #HIPRTC_SUCCESS
* @returns #HIPRTC_SUCCESS
*
* If the compiler failed to build the runtime compilation program,
* it will return #HIPRTC_ERROR_COMPILATION.
@@ -199,7 +202,7 @@ hiprtcResult hiprtcCompileProgram(hiprtcProgram prog,
* @param [in] numHeaders number of headers.
* @param [in] headers array of strings pointing to headers.
* @param [in] includeNames array of strings pointing to names included in program source.
* @return #HIPRTC_SUCCESS
* @returns #HIPRTC_SUCCESS
*
* Any invalide input parameter, it will return #HIPRTC_ERROR_INVALID_INPUT
* or #HIPRTC_ERROR_INVALID_PROGRAM.
@@ -219,7 +222,7 @@ hiprtcResult hiprtcCreateProgram(hiprtcProgram* prog,
* @brief Destroys an instance of given hiprtcProgram.
* @ingroup Runtime
* @param [in] prog runtime compilation program instance.
* @return #HIPRTC_SUCCESS
* @returns #HIPRTC_SUCCESS
*
* If prog is NULL, it will return #HIPRTC_ERROR_INVALID_INPUT.
*
@@ -234,7 +237,7 @@ hiprtcResult hiprtcDestroyProgram(hiprtcProgram* prog);
* @param [in] prog runtime compilation program instance.
* @param [in] name_expression const char pointer to the name expression.
* @param [in, out] lowered_name const char array to the lowered (mangled) name.
* @return #HIPRTC_SUCCESS
* @returns #HIPRTC_SUCCESS
*
* If any invalide nullptr input parameters, it will return #HIPRTC_ERROR_INVALID_INPUT
*
@@ -253,7 +256,7 @@ hiprtcResult hiprtcGetLoweredName(hiprtcProgram prog,
* @ingroup Runtime
* @param [in] prog runtime compilation program instance.
* @param [out] log memory pointer to the generated log.
* @return HIPRTC_SUCCESS
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
@@ -264,7 +267,7 @@ hiprtcResult hiprtcGetProgramLog(hiprtcProgram prog, char* log);
*
* @param [in] prog runtime compilation program instance.
* @param [out] logSizeRet size of generated log.
* @return HIPRTC_SUCCESS
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
@@ -276,7 +279,7 @@ hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram prog,
* @ingroup Runtime
* @param [in] prog runtime compilation program instance.
* @param [out] code char pointer to binary.
* @return HIPRTC_SUCCESS
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
@@ -287,7 +290,7 @@ hiprtcResult hiprtcGetCode(hiprtcProgram prog, char* code);
* @ingroup Runtime
* @param [in] prog runtime compilation program instance.
* @param [out] codeSizeRet the size of binary.
* @return HIPRTC_SUCCESS
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
@@ -310,7 +313,7 @@ hiprtcResult hiprtcGetBitcode(hiprtcProgram prog, char* bitcode);
*
* @param [in] prog runtime compilation program instance.
* @param [out] bitcode_size the size of bitcode.
* @return #HIPRTC_SUCCESS
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
@@ -323,8 +326,8 @@ hiprtcResult hiprtcGetBitcodeSize(hiprtcProgram prog, size_t* bitcode_size);
* @param [in] option_ptr Array of options
* @param [in] option_vals_pptr Array of option values cast to void*
* @param [out] hip_link_state_ptr hiprtc link state created upon success
*
* @return #HIPRTC_SUCCESS, #HIPRTC_ERROR_INVALID_INPUT, #HIPRTC_ERROR_INVALID_OPTION
*
* @returns #HIPRTC_SUCCESS, #HIPRTC_ERROR_INVALID_INPUT, #HIPRTC_ERROR_INVALID_OPTION
*
* @see hiprtcResult
*/
@@ -341,7 +344,7 @@ hiprtcResult hiprtcLinkCreate(unsigned int num_options, hiprtcJIT_option* option
* @param [in] options_ptr Array of options applied to this input
* @param [in] option_values Array of option values cast to void*
*
* @return #HIPRTC_SUCCESS
* @returns #HIPRTC_SUCCESS
*
* If input values are invalid, it will
* @return #HIPRTC_ERROR_INVALID_INPUT
@@ -365,7 +368,7 @@ hiprtcResult hiprtcLinkAddFile(hiprtcLinkState hip_link_state, hiprtcJITInputTyp
* @param [in] options_ptr Array of options applied to this input
* @param [in] option_values Array of option values cast to void*
*
* @return #HIPRTC_SUCCESS, #HIPRTC_ERROR_INVALID_INPUT
* @returns #HIPRTC_SUCCESS, #HIPRTC_ERROR_INVALID_INPUT
*
* If adding the file fails, it will
* @return #HIPRTC_ERROR_PROGRAM_CREATION_FAILURE
@@ -385,7 +388,7 @@ hiprtcResult hiprtcLinkAddData(hiprtcLinkState hip_link_state, hiprtcJITInputTyp
* @param [out] bin_out Upon success, points to the output binary
* @param [out] size_out Size of the binary is stored (optional)
*
* @return #HIPRTC_SUCCESS
* @returns #HIPRTC_SUCCESS
*
* If adding the data fails, it will
* @return #HIPRTC_ERROR_LINKING
@@ -399,7 +402,7 @@ hiprtcResult hiprtcLinkComplete(hiprtcLinkState hip_link_state, void** bin_out,
* @ingroup Runtime
* @param [in] hip_link_state link state instance
*
* @return #HIPRTC_SUCCESS
* @returns #HIPRTC_SUCCESS
*
* @see hiprtcResult
*/
projects/hip/tests/src/runtimeApi/cooperativeGrps/cooperative_streams.cpp
+66 -43
View File
@@ -200,8 +200,8 @@ template<typename T>
bool verifyFullCapacity(T& single_kernel_time, T& double_kernel_time, T& triple_kernel_time)
{
// Test that the two cooperative kernels took roughly twice as long as the one
if (double_kernel_time < 1.8 * single_kernel_time ||
double_kernel_time > 2.2 * single_kernel_time ) {
if (double_kernel_time < 1.7 * single_kernel_time ||
double_kernel_time > 2.3 * single_kernel_time ) {
std::cerr << "ERROR!" << std::endl;
std::cerr << "Two cooperative kernels launched at the same ";
std::cerr << "time did not take roughly twice as long as a single ";
@@ -209,14 +209,14 @@ bool verifyFullCapacity(T& single_kernel_time, T& double_kernel_time, T& triple_
return false;
}
// Test that the three kernels together took roughly 1.6 times as long as the two
// Test that the three kernels together took roughly 1.9 times as long as the two
// cooperative kernels. If the first 2 kernels run very fast, the third
// won't share much time with the second kernel.
if (triple_kernel_time > 1.7 * double_kernel_time) {
if (triple_kernel_time > 1.9 * double_kernel_time) {
std::cerr << "ERROR!" << std::endl;
std::cerr << "Launching a normal kernel in parallel with two ";
std::cerr << "back-to-back cooperative kernels still ended up taking ";
std::cerr << "more than 70% longer than the two cooperative kernels ";
std::cerr << "more than 90% longer than the two cooperative kernels ";
std::cerr << "alone." << std::endl;
return false;
}
@@ -331,8 +331,8 @@ int main(int argc, char** argv) {
for (int i = 0; i < 3; i++) {
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&dev_array[i]),
warp_size * sizeof(long long)));
HIPCHECK(hipMemsetAsync(dev_array[i], 0, warp_size * sizeof(long long),
max_active_blocks * warp_size * sizeof(long long)));
HIPCHECK(hipMemsetAsync(dev_array[i], 0, max_active_blocks * warp_size * sizeof(long long),
streams[i]));
}
@@ -347,6 +347,20 @@ int main(int argc, char** argv) {
coop_params[i][2] = reinterpret_cast<void*>(&totalTicks);
}
hipEvent_t single_start0, single_end0;
hipEvent_t single_start, single_end;
hipEvent_t double_start, double_end;
hipEvent_t triple_start, triple_end;
HIPCHECK(hipEventCreate(&single_start0));
HIPCHECK(hipEventCreate(&single_end0));
HIPCHECK(hipEventCreate(&single_start));
HIPCHECK(hipEventCreate(&single_end));
HIPCHECK(hipEventCreate(&double_start));
HIPCHECK(hipEventCreate(&double_end));
HIPCHECK(hipEventCreate(&triple_start));
HIPCHECK(hipEventCreate(&triple_end));
// Verify over capacity
HIPCHECK_API(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel_used),
max_active_blocks + 1, warp_size,
@@ -355,81 +369,90 @@ int main(int argc, char** argv) {
std::cout << "Launching an initial single cooperative kernel..." << std::endl;
// We need exclude the the initial launching as it will need time to load code obj.
auto single_start0 = std::chrono::system_clock::now();
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel),
HIPCHECK(hipEventRecord(single_start0, 0));
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel_used),
max_active_blocks, warp_size,
coop_params[0], 0, streams[0]));
HIPCHECK(hipEventRecord(single_end0, 0));
HIPCHECK(hipDeviceSynchronize());
auto single_end0 = std::chrono::system_clock::now();
std::cout << "Launching a single cooperative kernel..." << std::endl;
auto single_start = std::chrono::system_clock::now();
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel),
HIPCHECK(hipEventRecord(single_start, 0));
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel_used),
coop_blocks, warp_size,
coop_params[0], 0, streams[0]));
HIPCHECK(hipEventRecord(single_end, 0));
HIPCHECK(hipDeviceSynchronize());
auto single_end = std::chrono::system_clock::now();
std::cout << "Launching 2 cooperative kernels to different streams...";
std::cout << std::endl;
auto double_start = std::chrono::system_clock::now();
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel),
HIPCHECK(hipEventRecord(double_start, 0));
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel_used),
coop_blocks, warp_size,
coop_params[0], 0, streams[0]));
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel),
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel_used),
coop_blocks, warp_size,
coop_params[1], 0, streams[1]));
HIPCHECK(hipEventRecord(double_end, 0));
HIPCHECK(hipDeviceSynchronize());
auto double_end = std::chrono::system_clock::now();
std::cout << "Launching 2 cooperative kernels and 1 normal kernel...";
std::cout << std::endl;
auto triple_start = std::chrono::system_clock::now();
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel),
HIPCHECK(hipEventRecord(triple_start, 0));
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel_used),
coop_blocks, warp_size,
coop_params[0], 0, streams[0]));
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel),
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel_used),
coop_blocks, warp_size,
coop_params[1], 0, streams[1]));
hipLaunchKernelGGL(test_kernel, dim3(reg_blocks), dim3(warp_size),
hipLaunchKernelGGL(test_kernel_used, dim3(reg_blocks), dim3(warp_size),
0, streams[2], loops, dev_array[2], totalTicks);
HIPCHECK(hipEventRecord(triple_end, 0));
HIPCHECK(hipDeviceSynchronize());
auto triple_end = std::chrono::system_clock::now();
std::chrono::duration<double> single_kernel_time0 =
(single_end0 - single_start0);
std::chrono::duration<double> single_kernel_time =
(single_end - single_start);
std::chrono::duration<double> double_kernel_time =
(double_end - double_start);
std::chrono::duration<double> triple_kernel_time =
(triple_end - triple_start);
float single_kernel_time0;
float single_kernel_time;
float double_kernel_time;
float triple_kernel_time;
HIPCHECK(hipEventElapsedTime(&single_kernel_time0, single_start0, single_end0));
HIPCHECK(hipEventElapsedTime(&single_kernel_time, single_start, single_end));
HIPCHECK(hipEventElapsedTime(&double_kernel_time, double_start, double_end));
HIPCHECK(hipEventElapsedTime(&triple_kernel_time, triple_start, triple_end));
std::cout << "Initial single kernel took:" << std::endl;
std::cout << " " << single_kernel_time0.count();
std::cout << " seconds" << std::endl;
std::cout << " " << (int)single_kernel_time0;
std::cout << " milli-seconds" << std::endl;
std::cout << std::endl;
std::cout << "A single kernel took:" << std::endl;
std::cout << " " << single_kernel_time.count();
std::cout << " seconds" << std::endl;
std::cout << " " << (int)single_kernel_time;
std::cout << " milli-seconds" << std::endl;
std::cout << std::endl;
std::cout << "Two cooperative kernels that could run together took:";
std::cout << std::endl;
std::cout << " " << double_kernel_time.count();
std::cout << " seconds" << std::endl;
std::cout << " " << (int)double_kernel_time;
std::cout << " milli-seconds" << std::endl;
std::cout << std::endl;
std::cout << "Two coop kernels and a third regular kernel took:";
std::cout << std::endl << " ";
std::cout << triple_kernel_time.count();
std::cout << " seconds" << std::endl;
std::cout << (int)triple_kernel_time;
std::cout << " milli-seconds" << std::endl;
std::cout << "Testing whether these times make sense.." << std::endl;
FailFlag = !verify(p_tests, single_kernel_time,
double_kernel_time, triple_kernel_time);
HIPCHECK(hipEventDestroy(single_start0));
HIPCHECK(hipEventDestroy(single_end0));
HIPCHECK(hipEventDestroy(single_start));
HIPCHECK(hipEventDestroy(single_end));
HIPCHECK(hipEventDestroy(double_start));
HIPCHECK(hipEventDestroy(double_end));
HIPCHECK(hipEventDestroy(triple_start));
HIPCHECK(hipEventDestroy(triple_end));
for (int k = 0; k < 3; ++k) {
HIPCHECK(hipFree(dev_array[k]));
HIPCHECK(hipStreamDestroy(streams[k]));
projects/hip/tests/src/runtimeApi/module/hipModule.cpp
+11 -5
View File
@@ -126,10 +126,16 @@ bool isRocmPathSet() {
bool testMultiTargArchCodeObj() {
bool btestPassed = true;
#ifdef __linux__
#if defined(__linux__) && defined(__HIP_PLATFORM_AMD__)
char command[COMMAND_LEN];
hipDeviceProp_t props;
hipGetDeviceProperties(&props, 0);
// Extract the base GPU arch name excluding any feature
std::string arch = std::string(props.gcnArchName);
auto pos = arch.find(":");
if (pos != std::string::npos)
arch = arch.substr(0, pos);
// Hardcoding the codeobject lines in multiple string to avoid cpplint warning
std::string CodeObjL1 = "#include \"hip/hip_runtime.h\"\n";
std::string CodeObjL2 =
@@ -161,8 +167,8 @@ bool testMultiTargArchCodeObj() {
const char* genco_option = "--offload-arch";
const char* input_codeobj = "/tmp/vcpy_kernel.cpp";
snprintf(command, COMMAND_LEN,
"%s --genco %s=gfx801,gfx802,gfx803,gfx900,gfx908,%s %s -o %s",
hipcc_path, genco_option, props.gcnArchName, input_codeobj,
"unset HIP_PATH;%s --genco %s=gfx801,gfx802,gfx803,gfx900,gfx908,%s %s -o %s",
hipcc_path, genco_option, arch.c_str(), input_codeobj,
CODE_OBJ_MULTIARCH);
printf("command = %s\n", command);
@@ -173,11 +179,11 @@ bool testMultiTargArchCodeObj() {
if (access(command, F_OK) == -1) {
printf("Code Object File not found \n");
return true;
return false;
}
btestPassed = testCodeObjFile(CODE_OBJ_MULTIARCH);
#else
printf("This test is skipped due to non linux environment.\n");
printf("This test is skipped due to non linux or non AMD environment.\n");
#endif
return btestPassed;
}
projects/hip/tests/src/runtimeApi/module/hipModuleLaunchKernel.cpp
+1 -6
View File
@@ -246,15 +246,10 @@ bool Module_GridBlock_Corner_Tests() {
unsigned int maxblockX = deviceProp.maxThreadsDim[0];
unsigned int maxblockY = deviceProp.maxThreadsDim[1];
unsigned int maxblockZ = deviceProp.maxThreadsDim[2];
#ifdef __HIP_PLATFORM_NVCC__
unsigned int maxgridX = deviceProp.maxGridSize[0];
unsigned int maxgridY = deviceProp.maxGridSize[1];
unsigned int maxgridZ = deviceProp.maxGridSize[2];
#else
unsigned int maxgridX = INT32_MAX;
unsigned int maxgridY = UINT16_MAX;
unsigned int maxgridZ = UINT16_MAX;
#endif
struct gridblockDim test[6] = {{1, 1, 1, maxblockX, 1, 1},
{1, 1, 1, 1, maxblockY, 1},
{1, 1, 1, 1, 1, maxblockZ},