Add initial context implementation.

APIs: hipInit, hipCtxCreate.
Track TLS default ctx.  Set deviceID now changes the ctx.
Add first context test.

Change-Id: If1cb9989b5a04a36147e25e84904336c7b6f3d88


[ROCm/clr commit: f19f2248bf]

projects/clr/hipamd/CMakeLists.txt

@@ -133,6 +133,7 @@ if(HIP_PLATFORM STREQUAL "hcc")
 
     set(SOURCE_FILES src/device_util.cpp
                      src/hip_hcc.cpp
+                     src/hip_context.cpp
                      src/hip_device.cpp
                      src/hip_error.cpp
                      src/hip_event.cpp

projects/clr/hipamd/include/hcc_detail/hip_hcc.h

@@ -66,8 +66,16 @@ extern int HIP_VISIBLE_DEVICES; /* Contains a comma-separated sequence of GPU id
 extern int HIP_DISABLE_HW_KERNEL_DEP;
 extern int HIP_DISABLE_HW_COPY_DEP;
 
-extern thread_local int tls_defaultDevice;
+//---
+//Extern tls
+extern thread_local int tls_defaultDeviceId;
+extern thread_local ihipCtx_t *tls_defaultCtx;  
+
 extern thread_local hipError_t tls_lastHipError;
+
+
+//---
+//Forward defs:
 class ihipStream_t;
 class ihipDevice_t;
 class ihipCtx_t;

projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h

@@ -43,7 +43,13 @@ THE SOFTWARE.
 extern "C" {
 #endif
 
-typedef struct ihipCtx_t hipCtx_t;
+//---
+//API-visible structures
+typedef struct ihipCtx_t    *hipCtx_t;
+
+// Note many APIs also use integer deviceIds as an alternative to the device pointer:
+typedef struct ihipDevice_t *hipDevice_t;
+
 typedef struct ihipStream_t *hipStream_t;
 typedef struct hipEvent_t {
     struct ihipEvent_t *_handle;
@@ -1023,16 +1029,29 @@ hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int src
  * @}
  */
 
-
-
 /**
  *-------------------------------------------------------------------------------------------------
  *-------------------------------------------------------------------------------------------------
- *  @defgroup Version Management
+ *  @defgroup Driver Initialization and Version 
  *  @{
  *
  */
 
+/**
+ * @brief Explicitly initializes 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.
+hipError_t hipInit(unsigned int flags) ;
+
+
+
+// TODO-ctx
+hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags,  hipDevice_t device);
+
+
 /**
  * @brief Returns the approximate HIP driver version.
  *

projects/clr/hipamd/src/hip_context.cpp

@@ -0,0 +1,89 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+//---
+// Driver initialization and reporting:
+
+#include <stack>
+
+#include "hip_runtime.h"
+#include "hcc_detail/hip_hcc.h"
+#include "hcc_detail/trace_helper.h"
+
+// 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);
+}
+
+
+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 
+    hipError_t e = hipSuccess;
+
+    *ctx = new ihipCtx_t(device, g_deviceCnt, flags);
+    tls_defaultCtx = *ctx;
+    tls_ctxStack.push(*ctx);
+
+    return ihipLogStatus(e);
+}
+
+
+hipError_t hipDeviceGet(hipDevice_t *device, int deviceId)
+{
+    HIP_INIT_API(device, deviceId); // FIXME - review if we want to init 
+
+    *device = ihipGetDevice(deviceId);
+
+    hipError_t e = hipSuccess;
+    if (*device == NULL) {
+        e = hipErrorInvalidDevice;
+    }
+
+    return ihipLogStatus(e);
+};
+
+
+/**
+ * @return #hipSuccess
+ */
+//---
+hipError_t hipDriverGetVersion(int *driverVersion)
+{
+    HIP_INIT_API(driverVersion);
+
+    if (driverVersion) {
+        *driverVersion = 4;
+    }
+
+    return ihipLogStatus(hipSuccess);
+}

projects/clr/hipamd/src/hip_device.cpp

@@ -28,11 +28,11 @@ THE SOFTWARE.
 /**
  * @return  #hipSuccess
  */
-hipError_t hipGetDevice(int *device)
+hipError_t hipGetDevice(int *deviceId)
 {
-    HIP_INIT_API(device);
+    HIP_INIT_API(deviceId);
 
-    *device = tls_defaultDevice;
+    *deviceId = tls_defaultDeviceId;
     return ihipLogStatus(hipSuccess);
 }
 
@@ -130,13 +130,14 @@ hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig )
 /**
  * @return #hipSuccess, #hipErrorInvalidDevice
  */
-hipError_t hipSetDevice(int device)
+hipError_t hipSetDevice(int deviceId)
 {
-    HIP_INIT_API(device);
-    if ((device < 0) || (device >= g_deviceCnt)) {
+    HIP_INIT_API(deviceId);
+    if ((deviceId < 0) || (deviceId >= g_deviceCnt)) {
         return ihipLogStatus(hipErrorInvalidDevice);
     } else {
-        tls_defaultDevice = device;
+        tls_defaultDeviceId = deviceId;
+        tls_defaultCtx = ihipGetPrimaryCtx(deviceId);
         return ihipLogStatus(hipSuccess);
     }
 }

projects/clr/hipamd/src/hip_hcc.cpp

@@ -79,8 +79,7 @@ int HIP_VISIBLE_DEVICES = 0; /* Contains a comma-separated sequence of GPU ident
 int HIP_DISABLE_HW_KERNEL_DEP = 0;
 int HIP_DISABLE_HW_COPY_DEP = 0;
 
-thread_local int tls_defaultDevice = 0;
-thread_local hipError_t tls_lastHipError = hipSuccess;
+
 
 
 
@@ -101,8 +100,22 @@ hsa_agent_t gpu_agent_;
 hsa_amd_memory_pool_t gpu_pool_;
 
 //=================================================================================================
-// "free" functions:
+// Thread-local storage:
+//=================================================================================================
+thread_local int tls_defaultDeviceId = 0;
 
+// This is the implicit context used by all HIP commands.
+// It can be set by hipSetDevice or by the CTX manipulation commands:
+thread_local ihipCtx_t *tls_defaultCtx;  
+
+thread_local hipError_t tls_lastHipError = hipSuccess;
+
+
+
+
+//=================================================================================================
+// Top-level "free" functions:
+//=================================================================================================
 static inline bool ihipIsValidDevice(unsigned deviceIndex)
 {
     // deviceIndex is unsigned so always > 0
@@ -134,9 +147,9 @@ ihipCtx_t *ihipGetTlsDefaultCtx()
     // If this is invalid, the TLS state is corrupt.
     // This can fire if called before devices are initialized.
     // TODO - consider replacing assert with error code
-    assert (ihipIsValidDevice(tls_defaultDevice));
+    assert (ihipIsValidDevice(tls_defaultDeviceId));
 
-    return ihipGetPrimaryCtx(tls_defaultDevice);
+    return ihipGetPrimaryCtx(tls_defaultDeviceId);
 }
 
 

projects/clr/hipamd/src/hip_peer.cpp

@@ -36,7 +36,7 @@ THE SOFTWARE.
  * HCC returns 0 in *canAccessPeer ; Need to update this function when RT supports P2P
  */
 //---
-hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t *thisCtx, hipCtx_t *peerCtx)
+hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx)
 {
     HIP_INIT_API(canAccessPeer, thisCtx, peerCtx);
 
@@ -67,7 +67,7 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t *thisCtx, hipCtx
 //---
 // Disable visibility of this device into memory allocated on peer device.
 // Remove this device from peer device peerlist.
-hipError_t hipDeviceDisablePeerAccess (hipCtx_t *peerCtx)
+hipError_t hipDeviceDisablePeerAccess (hipCtx_t peerCtx)
 {
     HIP_INIT_API(peerCtx);
 
@@ -109,7 +109,7 @@ hipError_t hipDeviceDisablePeerAccess (hipCtx_t *peerCtx)
 //---
 // Allow the current device to see all memory allocated on peerDevice.
 // This should add this device to the peer-device peer list.
-hipError_t hipDeviceEnablePeerAccess (hipCtx_t *peerCtx, unsigned int flags)
+hipError_t hipDeviceEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags)
 {
     HIP_INIT_API(peerCtx, flags);
 
@@ -140,7 +140,7 @@ hipError_t hipDeviceEnablePeerAccess (hipCtx_t *peerCtx, unsigned int flags)
 
 
 //---
-hipError_t hipMemcpyPeer (void* dst, hipCtx_t *dstCtx, const void* src, hipCtx_t *srcCtx, size_t sizeBytes)
+hipError_t hipMemcpyPeer (void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t srcCtx, size_t sizeBytes)
 {
     HIP_INIT_API(dst, dstCtx, src, srcCtx, sizeBytes);
 
@@ -150,7 +150,7 @@ hipError_t hipMemcpyPeer (void* dst, hipCtx_t *dstCtx, const void* src, hipCtx_t
 
 
 //---
-hipError_t hipMemcpyPeerAsync (void* dst, hipCtx_t *dstDevice, const void* src, hipCtx_t *srcDevice, size_t sizeBytes, hipStream_t stream)
+hipError_t hipMemcpyPeerAsync (void* dst, hipCtx_t dstDevice, const void* src, hipCtx_t srcDevice, size_t sizeBytes, hipStream_t stream)
 {
     HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes, stream);
     // HCC has a unified memory architecture so device specifiers are not required.
@@ -201,19 +201,5 @@ hipError_t hipMemcpyPeerAsync (void* dst, int  dstDevice, const void* src, int
 }
 
 
-/**
- * @return #hipSuccess
- */
-//---
-hipError_t hipDriverGetVersion(int *driverVersion)
-{
-    HIP_INIT_API(driverVersion);
-
-    if (driverVersion) {
-        *driverVersion = 4;
-    }
-
-    return ihipLogStatus(hipSuccess);
-}
 
 

projects/clr/hipamd/tests/src/context/CMakeLists.txt

@@ -0,0 +1,9 @@
+cmake_minimum_required (VERSION 2.6)
+
+# Functions for kernel attributes (grid_launch, __launch_bounds__, etc)
+project (kernel)
+
+include_directories( ${HIPTEST_SOURCE_DIR} )
+
+build_hip_executable_libcpp (hipCtx_simple hipCtx_simple.cpp)
+make_test(hipCtx_simple " " )

projects/clr/hipamd/tests/src/context/hipCtx_simple.cpp

@@ -0,0 +1,39 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "hip_runtime.h"
+#include "test_common.h"
+
+int main(int argc, char *argv[])
+{
+    HipTest::parseStandardArguments(argc, argv, true);
+
+    HIPCHECK(hipInit());
+
+    hipDevice_t device;
+    hipCtx_t    ctx;
+
+    HIPCHECK(hipDeviceGet(&device, 1));
+    HIPCHECK(hipCtxCreate(&ctx, 0, device));
+
+    passed();
+};