Added further hipCtxXXX Apis

Change-Id: I286d962a06cee656c1c652b3f6b45078587fbc41


[ROCm/clr commit: 96de030c2d]

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,

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<ihipCtx_t *>  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);
+}