diff --git a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h index abdd820f3d..e606ebc471 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h @@ -427,7 +427,7 @@ const char *hipGetErrorString(hipError_t hipError); * @return #hipSuccess, #hipErrorInvalidValue * * Create a new asynchronous stream. @p stream returns an opaque handle that can be used to reference the newly - * created stream in subsequent hipStream* commands. The stream is allocated on the heap and will remain allocated + * created stream in subsequent hipStream* commands. The stream is allocated on the heap and will remain allocated * * even if the handle goes out-of-scope. To release the memory used by the stream, applicaiton must call hipStreamDestroy. * Flags controls behavior of the stream. See #hipStreamDefault, #hipStreamNonBlocking. @@ -444,13 +444,13 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags); * @return #hipSuccess, #hipErrorInvalidValue * * Create a new asynchronous stream. @p stream returns an opaque handle that can be used to reference the newly - * created stream in subsequent hipStream* commands. The stream is allocated on the heap and will remain allocated + * created stream in subsequent hipStream* commands. The stream is allocated on the heap and will remain allocated * even if the handle goes out-of-scope. To release the memory used by the stream, applicaiton must call hipStreamDestroy. - * + * * * @see hipStreamDestroy * - * @return + * @return * */ hipError_t hipStreamCreate(hipStream_t *stream); @@ -705,7 +705,7 @@ hipError_t hipMalloc(void** ptr, size_t size) ; hipError_t hipMallocHost(void** ptr, size_t size) __attribute__((deprecated("use hipHostMalloc instead"))) ; /** - * @brief Allocate device accessible page locked host memory + * @brief Allocate device accessible page locked host memory * * @param[out] ptr Pointer to the allocated host pinned memory * @param[in] size Requested memory size @@ -747,9 +747,9 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) ; * - #hipHostRegisterMapped Map the allocation into the address space for the current device. The device pointer can be obtained with #hipHostGetDevicePointer. * * - * After registering the memory, use #hipHostGetDevicePointer to obtain the mapped device pointer. + * After registering the memory, use #hipHostGetDevicePointer to obtain the mapped device pointer. * On many systems, the mapped device pointer will have a different value than the mapped host pointer. Applications - * must use the device pointer in device code, and the host pointer in device code. + * must use the device pointer in device code, and the host pointer in device code. * * On some systems, registered memory is pinned. On some systems, registered memory may not be actually be pinned * but uses OS or hardware facilities to all GPU access to the host memory. @@ -757,7 +757,7 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) ; * Developers are strongly encouraged to register memory blocks which are aligned to the host cache-line size. * (typically 64-bytes but can be obtains from the CPUID instruction). * - * If registering non-aligned pointers, the application must take care when register pointers from the same cache line + * If registering non-aligned pointers, the application must take care when register pointers from the same cache line * on different devices. HIP's coarse-grained synchronization model does not guarantee correct results if different * devices write to different parts of the same cache block - typically one of the writes will "win" and overwrite data * from the other registered memory region. @@ -795,7 +795,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height * If pointer is NULL, the hip runtime is initialized and hipSuccess is returned. * * @param[in] ptr Pointer to memory to be freed - * @return #hipSuccess + * @return #hipSuccess * @return #hipErrorInvalidDevicePointer (if pointer is invalid, including host pointers allocated with hipHostMalloc) */ hipError_t hipFree(void* ptr); @@ -816,7 +816,7 @@ hipError_t hipFreeHost(void* ptr) __attribute__((deprecated("use hipHostFree ins * If pointer is NULL, the hip runtime is initialized and hipSuccess is returned. * * @param[in] ptr Pointer to memory to be freed - * @return #hipSuccess, + * @return #hipSuccess, * #hipErrorInvalidValue (if pointer is invalid, including device pointers allocated with hipMalloc) */ hipError_t hipHostFree(void* ptr); @@ -832,7 +832,7 @@ hipError_t hipHostFree(void* ptr); * * For hipMemcpy, the copy is always performed by the current device (set by hipSetDevice). * For multi-gpu or peer-to-peer configurations, it is recommended to set the current device to the device where the src data is physically located. - * For optimal peer-to-peer copies, the copy device must be able to access the src and dst pointers (by calling hipDeviceEnablePeerAccess with copy agent as the + * For optimal peer-to-peer copies, the copy device must be able to access the src and dst pointers (by calling hipDeviceEnablePeerAccess with copy agent as the * current device and src/dest as the peerDevice argument. if this is not done, the hipMemcpy will still work, but will perform the copy using a staging buffer * on the host. * @@ -850,7 +850,7 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind * * The memory areas may not overlap. Symbol can either be a variable that resides in global or constant memory space, or it can be a character string, * naming a variable that resides in global or constant memory space. Kind can be either hipMemcpyHostToDevice or hipMemcpyDeviceToDevice - * TODO: cudaErrorInvalidSymbol and cudaErrorInvalidMemcpyDirection is not supported, use hipErrorUnknown for now. + * TODO: cudaErrorInvalidSymbol and cudaErrorInvalidMemcpyDirection is not supported, use hipErrorUnknown for now. * * @param[in] symbolName - Symbol destination on device * @param[in] src - Data being copy from @@ -871,7 +871,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t siz * For hipMemcpy, the copy is always performed by the device associated with the specified stream. * * For multi-gpu or peer-to-peer configurations, it is recommended to use a stream which is a attached to the device where the src data is physically located. - * For optimal peer-to-peer copies, the copy device must be able to access the src and dst pointers (by calling hipDeviceEnablePeerAccess with copy agent as the + * For optimal peer-to-peer copies, the copy device must be able to access the src and dst pointers (by calling hipDeviceEnablePeerAccess with copy agent as the * current device and src/dest as the peerDevice argument. if this is not done, the hipMemcpy will still work, but will perform the copy using a staging buffer * on the host. * @@ -958,7 +958,7 @@ hipError_t hipMemGetInfo (size_t * free, size_t * total) ; * * Returns "0" in @p canAccessPeer if deviceId == peerDeviceId, and both are valid devices : a device is not a peer of itself. * - * @returns #hipSuccess, + * @returns #hipSuccess, * @returns #hipErrorInvalidDevice if deviceId or peerDeviceId are not valid devices * @warning PeerToPeer support is experimental. */ @@ -966,7 +966,7 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDev /** - * @brief Enable direct access from current device's virtual address space to memory allocations physically located on a peer device. + * @brief Enable direct access from current device's virtual address space to memory allocations physically located on a peer device. * * Memory which already allocated on peer device will be mapped into the address space of the current device. In addition, all * future memory allocations on peerDeviceId will be mapped into the address space of the current device when the memory is allocated. @@ -976,7 +976,7 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDev * @param [in] peerDeviceId * @param [in] flags * - * Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, + * Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, * @returns #hipErrorPeerAccessAlreadyEnabled if peer access is already enabled for this device. * @warning PeerToPeer support is experimental. */ @@ -984,7 +984,7 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags); /** - * @brief Disable direct access from current device's virtual address space to memory allocations physically located on a peer device. + * @brief Disable direct access from current device's virtual address space to memory allocations physically located on a peer device. * * Returns hipErrorPeerAccessNotEnabled if direct access to memory on peerDevice has not yet been enabled from the current device. * @@ -1039,15 +1039,15 @@ hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int src /** *------------------------------------------------------------------------------------------------- *------------------------------------------------------------------------------------------------- - * @defgroup Driver Initialization and Version + * @defgroup Driver Initialization and Version * @{ * */ /** * @brief Explicitly initializes the HIP runtime. - * - * Most HIP APIs implicitly initialize the HIP runtime. + * + * Most HIP APIs implicitly initialize the HIP runtime. * This API provides control over the timing of the initialization. */ // TODO-ctx - more description on error codes. @@ -1070,6 +1070,16 @@ hipError_t hipCtxGetCurrent(hipCtx_t* ctx); hipError_t hipCtxGetDevice(hipDevice_t *device); +hipError_t hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion); + +hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig ); + +hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig ); + +hipError_t hipCtxSetSharedMemConfig ( hipSharedMemConfig config ); + +hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig ); + // TODO-ctx /** * @return hipSuccess, hipErrorInvalidDevice @@ -1094,7 +1104,7 @@ hipError_t hipModuleLoad(hipModule *module, const char *fname); hipError_t hipModuleGetFunction(hipFunction *function, hipModule module, const char *kname); -hipError_t hipDrvLaunchKernel(hipFunction f, +hipError_t hipDrvLaunchKernel(hipFunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, diff --git a/projects/clr/hipamd/src/hip_context.cpp b/projects/clr/hipamd/src/hip_context.cpp index 9eb65fae39..ee9e37a1a1 100644 --- a/projects/clr/hipamd/src/hip_context.cpp +++ b/projects/clr/hipamd/src/hip_context.cpp @@ -26,20 +26,20 @@ THE SOFTWARE. #include "hcc_detail/hip_hcc.h" #include "hcc_detail/trace_helper.h" -// Stack of contexts +// Stack of contexts thread_local std::stack tls_ctxStack; hipError_t hipInit(unsigned int flags) { HIP_INIT_API(flags); - + hipError_t e = hipSuccess; // Flags must be 0 if (flags != 0) { e = hipErrorInvalidValue; - } + } return ihipLogStatus(e); } @@ -47,7 +47,7 @@ hipError_t hipInit(unsigned int flags) hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device) { - HIP_INIT_API(ctx, flags, device); // FIXME - review if we want to init + HIP_INIT_API(ctx, flags, device); // FIXME - review if we want to init hipError_t e = hipSuccess; *ctx = new ihipCtx_t(device, g_deviceCnt, flags); @@ -60,7 +60,7 @@ hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device) hipError_t hipDeviceGet(hipDevice_t *device, int deviceId) { - HIP_INIT_API(device, deviceId); // FIXME - review if we want to init + HIP_INIT_API(device, deviceId); // FIXME - review if we want to init *device = ihipGetDevice(deviceId); @@ -103,15 +103,19 @@ hipError_t hipCtxDestroy(hipCtx_t ctx) hipError_t hipCtxPopCurrent(hipCtx_t* ctx) { hipError_t e = hipSuccess; - tls_ctxStack.pop(); + ihipCtx_t* tempCtx; + *ctx = ihipGetTlsDefaultCtx(); if(!tls_ctxStack.empty()) { - *ctx= tls_ctxStack.top(); + tls_ctxStack.pop(); } - else { - *ctx = nullptr; + if(!tls_ctxStack.empty()) { + tempCtx= tls_ctxStack.top(); } - - ihipSetTlsDefaultCtx(*ctx); //TOD0 - Shall check for NULL? + else { + tempCtx = nullptr; + } + + ihipSetTlsDefaultCtx(tempCtx); //TOD0 - Shall check for NULL? return ihipLogStatus(e); } @@ -166,3 +170,50 @@ hipError_t hipCtxGetDevice(hipDevice_t *device) } return ihipLogStatus(e); } + +hipError_t hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion) +{ + HIP_INIT_API(apiVersion); + + if (apiVersion) { + *apiVersion = 4; + } + + return ihipLogStatus(hipSuccess); +} + +hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig ) +{ + HIP_INIT_API(cacheConfig); + + *cacheConfig = hipFuncCachePreferNone; + + return ihipLogStatus(hipSuccess); +} + +hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig ) +{ + HIP_INIT_API(cacheConfig); + + // Nop, AMD does not support variable cache configs. + + return ihipLogStatus(hipSuccess); +} + +hipError_t hipCtxSetSharedMemConfig ( hipSharedMemConfig config ) +{ + HIP_INIT_API(config); + + // Nop, AMD does not support variable shared mem configs. + + return ihipLogStatus(hipSuccess); +} + +hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig ) +{ + HIP_INIT_API(pConfig); + + *pConfig = hipSharedMemBankSizeFourByte; + + return ihipLogStatus(hipSuccess); +} \ No newline at end of file