diff --git a/projects/hip/cmake/FindHIP.cmake b/projects/hip/cmake/FindHIP.cmake index 24e475d35b..29abcd7870 100644 --- a/projects/hip/cmake/FindHIP.cmake +++ b/projects/hip/cmake/FindHIP.cmake @@ -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) diff --git a/projects/hip/include/hip/hip_ext.h b/projects/hip/include/hip/hip_ext.h index 2548cbe1e5..09e9cf8b95 100644 --- a/projects/hip/include/hip/hip_ext.h +++ b/projects/hip/include/hip/hip_ext.h @@ -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. diff --git a/projects/hip/include/hip/hip_runtime_api.h b/projects/hip/include/hip/hip_runtime_api.h index a1f3822366..e55cdff4a8 100644 --- a/projects/hip/include/hip/hip_runtime_api.h +++ b/projects/hip/include/hip/hip_runtime_api.h @@ -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 @@ -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. * diff --git a/projects/hip/include/hip/hiprtc.h b/projects/hip/include/hip/hiprtc.h index cfc3a862ba..6f764bfa4a 100644 --- a/projects/hip/include/hip/hiprtc.h +++ b/projects/hip/include/hip/hiprtc.h @@ -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 */ diff --git a/projects/hip/tests/src/runtimeApi/cooperativeGrps/cooperative_streams.cpp b/projects/hip/tests/src/runtimeApi/cooperativeGrps/cooperative_streams.cpp index 70ce8b6a80..aa32bb4a3e 100644 --- a/projects/hip/tests/src/runtimeApi/cooperativeGrps/cooperative_streams.cpp +++ b/projects/hip/tests/src/runtimeApi/cooperativeGrps/cooperative_streams.cpp @@ -200,8 +200,8 @@ template 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(&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(&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(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(test_kernel), + HIPCHECK(hipEventRecord(single_start0, 0)); + HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(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(test_kernel), + HIPCHECK(hipEventRecord(single_start, 0)); + HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(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(test_kernel), + HIPCHECK(hipEventRecord(double_start, 0)); + HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel_used), coop_blocks, warp_size, coop_params[0], 0, streams[0])); - HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel), + HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(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(test_kernel), + HIPCHECK(hipEventRecord(triple_start, 0)); + HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel_used), coop_blocks, warp_size, coop_params[0], 0, streams[0])); - HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel), + HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(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 single_kernel_time0 = - (single_end0 - single_start0); - std::chrono::duration single_kernel_time = - (single_end - single_start); - std::chrono::duration double_kernel_time = - (double_end - double_start); - std::chrono::duration 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])); diff --git a/projects/hip/tests/src/runtimeApi/module/hipModule.cpp b/projects/hip/tests/src/runtimeApi/module/hipModule.cpp index e29c56185f..ed7ff62d79 100644 --- a/projects/hip/tests/src/runtimeApi/module/hipModule.cpp +++ b/projects/hip/tests/src/runtimeApi/module/hipModule.cpp @@ -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; } diff --git a/projects/hip/tests/src/runtimeApi/module/hipModuleLaunchKernel.cpp b/projects/hip/tests/src/runtimeApi/module/hipModuleLaunchKernel.cpp index d317405799..324ebd4fdf 100644 --- a/projects/hip/tests/src/runtimeApi/module/hipModuleLaunchKernel.cpp +++ b/projects/hip/tests/src/runtimeApi/module/hipModuleLaunchKernel.cpp @@ -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},