Context update.

- Remove tls_deviceID.
- Add first passing test.

Change-Id: If3e2f254abf589028cfe4f9e6369745f04160de0

include/hcc_detail/hip_hcc.h

@@ -68,9 +68,6 @@ extern int HIP_DISABLE_HW_COPY_DEP;
 
 //---
 //Extern tls
-extern thread_local int tls_defaultDeviceId;
-extern thread_local ihipCtx_t *tls_defaultCtx;
-
 extern thread_local hipError_t tls_lastHipError;
 
 
@@ -653,6 +650,7 @@ extern hsa_agent_t g_cpu_agent ;   // the CPU agent.
 extern void ihipInit();
 extern const char *ihipErrorString(hipError_t);
 extern ihipCtx_t    *ihipGetTlsDefaultCtx();
+extern void          ihipSetTlsDefaultCtx(ihipCtx_t *ctx);
 
 extern ihipDevice_t *ihipGetDevice(int);
 ihipCtx_t * ihipGetPrimaryCtx(unsigned deviceIndex);

include/hcc_detail/hip_runtime_api.h

@@ -244,6 +244,8 @@ hipError_t hipSetDevice(int deviceId);
  * hipGetDevice returns in * @p device the default device for the calling host thread.
  *
  * @see hipSetDevice, hipGetDevicesizeBytes
+ *
+ * @returns hipSuccess, hipErrorInvalidDevice
  */
 hipError_t hipGetDevice(int *deviceId);
 
@@ -1052,6 +1054,13 @@ hipError_t hipInit(unsigned int flags) ;
 hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags,  hipDevice_t device);
 
 
+// TODO-ctx
+/**
+ * @return hipSuccess, hipErrorInvalidDevice
+ */
+hipError_t hipDeviceGetFromId(hipDevice_t *device, int deviceId);
+
+
 /**
  * @brief Returns the approximate HIP driver version.
  *

src/hip_context.cpp

@@ -51,7 +51,7 @@ hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags,  hipDevice_t device)
     hipError_t e = hipSuccess;
 
     *ctx = new ihipCtx_t(device, g_deviceCnt, flags);
-    tls_defaultCtx = *ctx;
+    ihipSetTlsDefaultCtx(*ctx);
     tls_ctxStack.push(*ctx);
 
     return ihipLogStatus(e);

src/hip_device.cpp

@@ -26,14 +26,25 @@ THE SOFTWARE.
 //-------------------------------------------------------------------------------------------------
 //---
 /**
- * @return  #hipSuccess
+ * @return  #hipSuccess, hipErrorInvalidDevice
  */
+// TODO - does this initialize HIP runtime?
 hipError_t hipGetDevice(int *deviceId)
 {
     HIP_INIT_API(deviceId);
 
-    *deviceId = tls_defaultDeviceId;
-    return ihipLogStatus(hipSuccess);
+    hipError_t e = hipSuccess;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+
+    if (ctx == nullptr) {
+        e = hipErrorInvalidDevice; // TODO, check error code.
+        *deviceId = -1;
+    } else {
+        *deviceId = ctx->getDevice()->_deviceId;
+    }
+
+    return ihipLogStatus(e);
 }
 
 
@@ -41,6 +52,7 @@ hipError_t hipGetDevice(int *deviceId)
 /**
  * @return  #hipSuccess, #hipErrorNoDevice
  */
+// TODO - does this initialize HIP runtime?
 hipError_t hipGetDeviceCount(int *count)
 {
     HIP_INIT_API(count);
@@ -136,8 +148,7 @@ hipError_t hipSetDevice(int deviceId)
     if ((deviceId < 0) || (deviceId >= g_deviceCnt)) {
         return ihipLogStatus(hipErrorInvalidDevice);
     } else {
-        tls_defaultDeviceId = deviceId;
-        tls_defaultCtx = ihipGetPrimaryCtx(deviceId);
+        ihipSetTlsDefaultCtx(ihipGetPrimaryCtx(deviceId));
         return ihipLogStatus(hipSuccess);
     }
 }
@@ -299,3 +310,20 @@ hipError_t hipSetDeviceFlags( unsigned int flags)
 };
 
 
+
+
+hipError_t hipDeviceGetFromId(hipDevice_t *device, int deviceId)
+{
+    HIP_INIT_API(device, deviceId);
+
+    hipError_t e = hipSuccess;
+
+    *device = ihipGetDevice(deviceId);
+
+    if (device == nullptr) {
+        e = hipErrorInvalidDevice;
+    }
+
+
+    return ihipLogStatus(e);
+}

src/hip_hcc.cpp

@@ -102,11 +102,9 @@ hsa_amd_memory_pool_t gpu_pool_;
 //=================================================================================================
 // 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;
 
@@ -139,17 +137,25 @@ ihipCtx_t * ihipGetPrimaryCtx(unsigned deviceIndex)
 };
 
 
+static thread_local ihipCtx_t *tls_defaultCtx = nullptr;  
+void ihipSetTlsDefaultCtx(ihipCtx_t *ctx)
+{
+    tls_defaultCtx = ctx;
+}
+
 
 //---
-//FIXME - this needs to return the active context for this CPU thread - not primary for device.
+//TODO - review the context creation strategy here.  Really should be:
+//  - first "non-device" runtime call creates the context for this thread.  Allowed to call setDevice first.
+//  - hipDeviceReset destroys the primary context for device?
+//  - Then context is created again for next usage.
 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_defaultDeviceId));
-
-    return ihipGetPrimaryCtx(tls_defaultDeviceId);
+    // Per-thread initialization of the TLS:
+    if ((tls_defaultCtx == nullptr) && (g_deviceCnt>0)) {
+        ihipSetTlsDefaultCtx(ihipGetPrimaryCtx(0));
+    }
+    return tls_defaultCtx;
 }
 
 
@@ -1221,6 +1227,7 @@ void ihipInit()
         assert(deviceCnt == g_deviceCnt);
     }
 
+
     tprintf(DB_SYNC, "pid=%u %-30s\n", getpid(), "<ihipInit>");
 }
 

tests/src/CMakeLists.txt

@@ -219,6 +219,7 @@ make_hipify_test(specialFunc.cu )
 #make_test(hipDynamicShared " ")
 
 # Add subdirs here:
+add_subdirectory(context)
 add_subdirectory(deviceLib)
 add_subdirectory(runtimeApi)
 add_subdirectory(kernel)

tests/src/context/hipCtx_simple.cpp

@@ -27,12 +27,12 @@ int main(int argc, char *argv[])
 {
     HipTest::parseStandardArguments(argc, argv, true);
 
-    HIPCHECK(hipInit());
+    HIPCHECK(hipInit(0));
 
     hipDevice_t device;
     hipCtx_t    ctx;
 
-    HIPCHECK(hipDeviceGet(&device, 1));
+    HIPCHECK(hipDeviceGetFromId(&device, 0));
     HIPCHECK(hipCtxCreate(&ctx, 0, device));
 
     passed();

tests/src/hipHostAlloc.cpp

@@ -1,24 +1,24 @@
 /*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+   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:
+   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 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.
-*/
+   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"test_common.h"
 
@@ -26,43 +26,43 @@ THE SOFTWARE.
 #define SIZE LEN*sizeof(float)
 
 __global__ void Add(hipLaunchParm lp, float *Ad, float *Bd, float *Cd){
-int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
-Cd[tx] = Ad[tx] + Bd[tx];
+    int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
+    Cd[tx] = Ad[tx] + Bd[tx];
 }
 
 int main(){
-float *A, *B, *C;
-float *Ad, *Bd, *Cd;
+    float *A, *B, *C;
+    float *Ad, *Bd, *Cd;
 
-hipDeviceProp_t prop;
-int device;
-HIPCHECK(hipGetDevice(&device));
-HIPCHECK(hipGetDeviceProperties(&prop, device));
-if(prop.canMapHostMemory != 1){
-std::cout<<"Exiting..."<<std::endl;
-failed("Does support HostPinned Memory");
-}
+    hipDeviceProp_t prop;
+    int device;
+    HIPCHECK(hipGetDevice(&device));
+    HIPCHECK(hipGetDeviceProperties(&prop, device));
+    if(prop.canMapHostMemory != 1){
+        std::cout<<"Exiting..."<<std::endl;
+        failed("Does support HostPinned Memory");
+    }
 
-HIPCHECK(hipHostMalloc((void**)&A, SIZE, hipHostMallocWriteCombined | hipHostMallocMapped));
-HIPCHECK(hipHostMalloc((void**)&B, SIZE, hipHostMallocDefault));
-HIPCHECK(hipHostMalloc((void**)&C, SIZE, hipHostMallocMapped));
+    HIPCHECK(hipHostMalloc((void**)&A, SIZE, hipHostMallocWriteCombined | hipHostMallocMapped));
+    HIPCHECK(hipHostMalloc((void**)&B, SIZE, hipHostMallocDefault));
+    HIPCHECK(hipHostMalloc((void**)&C, SIZE, hipHostMallocMapped));
 
-HIPCHECK(hipHostGetDevicePointer((void**)&Ad, A, 0));
-HIPCHECK(hipHostGetDevicePointer((void**)&Cd, C, 0));
+    HIPCHECK(hipHostGetDevicePointer((void**)&Ad, A, 0));
+    HIPCHECK(hipHostGetDevicePointer((void**)&Cd, C, 0));
 
-for(int i=0;i<LEN;i++){
-A[i] = 1.0f;
-B[i] = 2.0f;
-}
+    for(int i=0;i<LEN;i++){
+        A[i] = 1.0f;
+        B[i] = 2.0f;
+    }
 
-HIPCHECK(hipMalloc((void**)&Bd, SIZE));
-HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
+    HIPCHECK(hipMalloc((void**)&Bd, SIZE));
+    HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
 
-dim3 dimGrid(LEN/512,1,1);
-dim3 dimBlock(512,1,1);
+    dim3 dimGrid(LEN/512,1,1);
+    dim3 dimBlock(512,1,1);
 
-hipLaunchKernel(HIP_KERNEL_NAME(Add), dimGrid, dimBlock, 0, 0, Ad, Bd, Cd);
+    hipLaunchKernel(HIP_KERNEL_NAME(Add), dimGrid, dimBlock, 0, 0, Ad, Bd, Cd);
 
-passed();
+    passed();
 
 }