From 0c8ca4b37dadaaa17af0703c71fb7e7aa4238326 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 8 Aug 2016 17:49:02 -0500 Subject: [PATCH] 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: f19f2248bfb0e22dddec47be12b2654509d80cc5] --- projects/clr/hipamd/CMakeLists.txt | 1 + .../clr/hipamd/include/hcc_detail/hip_hcc.h | 10 ++- .../include/hcc_detail/hip_runtime_api.h | 27 +++++- projects/clr/hipamd/src/hip_context.cpp | 89 +++++++++++++++++++ projects/clr/hipamd/src/hip_device.cpp | 15 ++-- projects/clr/hipamd/src/hip_hcc.cpp | 23 +++-- projects/clr/hipamd/src/hip_peer.cpp | 24 ++--- .../hipamd/tests/src/context/CMakeLists.txt | 9 ++ .../tests/src/context/hipCtx_simple.cpp | 39 ++++++++ 9 files changed, 201 insertions(+), 36 deletions(-) create mode 100644 projects/clr/hipamd/src/hip_context.cpp create mode 100644 projects/clr/hipamd/tests/src/context/CMakeLists.txt create mode 100644 projects/clr/hipamd/tests/src/context/hipCtx_simple.cpp diff --git a/projects/clr/hipamd/CMakeLists.txt b/projects/clr/hipamd/CMakeLists.txt index a26e848985..58046c069f 100644 --- a/projects/clr/hipamd/CMakeLists.txt +++ b/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 diff --git a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h b/projects/clr/hipamd/include/hcc_detail/hip_hcc.h index a197e90899..933cff16e2 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_hcc.h +++ b/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; 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 d85a6f6800..273dc760ff 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/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. * diff --git a/projects/clr/hipamd/src/hip_context.cpp b/projects/clr/hipamd/src/hip_context.cpp new file mode 100644 index 0000000000..9c7392aaae --- /dev/null +++ b/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 + +#include "hip_runtime.h" +#include "hcc_detail/hip_hcc.h" +#include "hcc_detail/trace_helper.h" + +// 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); +} + + +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); +} diff --git a/projects/clr/hipamd/src/hip_device.cpp b/projects/clr/hipamd/src/hip_device.cpp index dff736bb94..c0f52a7a5a 100644 --- a/projects/clr/hipamd/src/hip_device.cpp +++ b/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); } } diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index 363272775c..6cfc73d037 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/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); } diff --git a/projects/clr/hipamd/src/hip_peer.cpp b/projects/clr/hipamd/src/hip_peer.cpp index f5f9287baa..fd51599815 100644 --- a/projects/clr/hipamd/src/hip_peer.cpp +++ b/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); -} diff --git a/projects/clr/hipamd/tests/src/context/CMakeLists.txt b/projects/clr/hipamd/tests/src/context/CMakeLists.txt new file mode 100644 index 0000000000..d04985d001 --- /dev/null +++ b/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 " " ) diff --git a/projects/clr/hipamd/tests/src/context/hipCtx_simple.cpp b/projects/clr/hipamd/tests/src/context/hipCtx_simple.cpp new file mode 100644 index 0000000000..d54e4b67ca --- /dev/null +++ b/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(); +};