From a039dbf2dd7518e84a043f0d7533712ee29c58ea Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 20 Apr 2018 17:34:48 +0530 Subject: [PATCH 01/12] Mark hipCtx APIs as deprecated --- docs/markdown/hip_deprecated_api_list.md | 22 +++++++++++++++ include/hip/hcc_detail/hip_runtime_api.h | 35 +++++++++++++----------- 2 files changed, 41 insertions(+), 16 deletions(-) create mode 100644 docs/markdown/hip_deprecated_api_list.md diff --git a/docs/markdown/hip_deprecated_api_list.md b/docs/markdown/hip_deprecated_api_list.md new file mode 100644 index 0000000000..a96f7f4d3f --- /dev/null +++ b/docs/markdown/hip_deprecated_api_list.md @@ -0,0 +1,22 @@ +# HIP Deprecated API List + +## HIP Context API + +CUDA supports cuCtx API, the Driver API that defines "Context" and "Devices" as separate entities. Contexts contain a single device, and a device can theoretically have multiple contexts. HIP initially added limited support for these API to facilitate easy porting from existing driver codes. These API are marked as deprecated now since there are better alternate interface (such as hipSetDevice or the stream API) to achieve the required functions. + +###hipCtxCreate +###hipCtxDestroy +###hipCtxPopCurrent +###hipCtxPushCurrent +###hipCtxSetCurrent +###hipCtxGetCurrent +###hipCtxGetDevice +###hipCtxGetApiVersion +###hipCtxGetCacheConfig +###hipCtxSetCacheConfig +###hipCtxSetSharedMemConfig +###hipCtxGetSharedMemConfig +###hipCtxSynchronize +###hipCtxGetFlags +###hipCtxEnablePeerAccess +###hipCtxDisablePeerAccess diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 8674824c48..f7a5718894 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -42,6 +42,9 @@ THE SOFTWARE. #include #include +#define DEPRECATED(msg, func) func __attribute__ ((deprecated(msg))) +#define DEPRECATED_MSG "This API is marked as deprecated and may not be supported in future releases.For more details please refer https://github.com/ROCm-Developer-Tools/HIP/tree/master/docs/markdown/hip_deprecated_api_list" + #if defined(__HCC__) && (__hcc_workweek__ < 16155) #error("This version of HIP requires a newer version of HCC."); #endif @@ -1828,7 +1831,7 @@ hipError_t hipInit(unsigned int flags); * @see hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxPushCurrent, * hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device)); /** * @brief Destroy a HIP context. @@ -1840,7 +1843,7 @@ hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device); * @see hipCtxCreate, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent,hipCtxSetCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize , hipCtxGetDevice */ -hipError_t hipCtxDestroy(hipCtx_t ctx); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxDestroy(hipCtx_t ctx)); /** * @brief Pop the current/default context and return the popped context. @@ -1852,7 +1855,7 @@ hipError_t hipCtxDestroy(hipCtx_t ctx); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxSetCurrent, hipCtxGetCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxPopCurrent(hipCtx_t* ctx); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxPopCurrent(hipCtx_t* ctx)); /** * @brief Push the context to be set as current/ default context @@ -1864,7 +1867,7 @@ hipError_t hipCtxPopCurrent(hipCtx_t* ctx); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize , hipCtxGetDevice */ -hipError_t hipCtxPushCurrent(hipCtx_t ctx); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxPushCurrent(hipCtx_t ctx)); /** * @brief Set the passed context as current/default @@ -1876,7 +1879,7 @@ hipError_t hipCtxPushCurrent(hipCtx_t ctx); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize , hipCtxGetDevice */ -hipError_t hipCtxSetCurrent(hipCtx_t ctx); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxSetCurrent(hipCtx_t ctx)); /** * @brief Get the handle of the current/ default context @@ -1888,7 +1891,7 @@ hipError_t hipCtxSetCurrent(hipCtx_t ctx); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetDevice, hipCtxGetFlags, hipCtxPopCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxGetCurrent(hipCtx_t* ctx); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetCurrent(hipCtx_t* ctx)); /** * @brief Get the handle of the device associated with current/default context @@ -1901,7 +1904,7 @@ hipError_t hipCtxGetCurrent(hipCtx_t* ctx); * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize */ -hipError_t hipCtxGetDevice(hipDevice_t* device); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetDevice(hipDevice_t* device)); /** * @brief Returns the approximate HIP api version. @@ -1920,7 +1923,7 @@ hipError_t hipCtxGetDevice(hipDevice_t* device); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetDevice, hipCtxGetFlags, hipCtxPopCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion)); /** * @brief Set Cache configuration for a specific function @@ -1935,7 +1938,7 @@ hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig)); /** * @brief Set L1/Shared cache partition. @@ -1950,7 +1953,7 @@ hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig)); /** * @brief Set Shared memory bank configuration. @@ -1965,7 +1968,7 @@ hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config)); /** * @brief Get Shared memory bank configuration. @@ -1980,7 +1983,7 @@ hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig)); /** * @brief Blocks until the default context has completed all preceding requested tasks. @@ -1993,7 +1996,7 @@ hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxGetDevice */ -hipError_t hipCtxSynchronize(void); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxSynchronize(void)); /** * @brief Return flags used for creating default context. @@ -2005,7 +2008,7 @@ hipError_t hipCtxSynchronize(void); * @see hipCtxCreate, hipCtxDestroy, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -hipError_t hipCtxGetFlags(unsigned int* flags); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetFlags(unsigned int* flags)); /** * @brief Enables direct access to memory allocations in a peer context. @@ -2026,7 +2029,7 @@ hipError_t hipCtxGetFlags(unsigned int* flags); * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice * @warning PeerToPeer support is experimental. */ -hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags)); /** * @brief Disable direct access from current context's virtual address space to memory allocations @@ -2044,7 +2047,7 @@ hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags); * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice * @warning PeerToPeer support is experimental. */ -hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx); +DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx)); /** * @brief Get the state of the primary context. From 19cb6c61abb84a107932f689be6a71435ec54d61 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 20 Apr 2018 21:57:23 +0530 Subject: [PATCH 02/12] Moved deprecated macro to the end --- include/hip/hcc_detail/hip_runtime_api.h | 34 ++++++++++++------------ 1 file changed, 17 insertions(+), 17 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index f7a5718894..86678f90ef 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -42,7 +42,7 @@ THE SOFTWARE. #include #include -#define DEPRECATED(msg, func) func __attribute__ ((deprecated(msg))) +#define DEPRECATED(msg) __attribute__ ((deprecated(msg))) #define DEPRECATED_MSG "This API is marked as deprecated and may not be supported in future releases.For more details please refer https://github.com/ROCm-Developer-Tools/HIP/tree/master/docs/markdown/hip_deprecated_api_list" #if defined(__HCC__) && (__hcc_workweek__ < 16155) @@ -1831,7 +1831,7 @@ hipError_t hipInit(unsigned int flags); * @see hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxPushCurrent, * hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device)); +hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device)DEPRECATED(DEPRECATED_MSG); /** * @brief Destroy a HIP context. @@ -1843,7 +1843,7 @@ DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int fl * @see hipCtxCreate, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent,hipCtxSetCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize , hipCtxGetDevice */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxDestroy(hipCtx_t ctx)); +hipError_t hipCtxDestroy(hipCtx_t ctx) DEPRECATED(DEPRECATED_MSG); /** * @brief Pop the current/default context and return the popped context. @@ -1855,7 +1855,7 @@ DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxDestroy(hipCtx_t ctx)); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxSetCurrent, hipCtxGetCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxPopCurrent(hipCtx_t* ctx)); +hipError_t hipCtxPopCurrent(hipCtx_t* ctx) DEPRECATED(DEPRECATED_MSG); /** * @brief Push the context to be set as current/ default context @@ -1867,7 +1867,7 @@ DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxPopCurrent(hipCtx_t* ctx)); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize , hipCtxGetDevice */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxPushCurrent(hipCtx_t ctx)); +hipError_t hipCtxPushCurrent(hipCtx_t ctx) DEPRECATED(DEPRECATED_MSG); /** * @brief Set the passed context as current/default @@ -1879,7 +1879,7 @@ DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxPushCurrent(hipCtx_t ctx)); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize , hipCtxGetDevice */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxSetCurrent(hipCtx_t ctx)); +hipError_t hipCtxSetCurrent(hipCtx_t ctx) DEPRECATED(DEPRECATED_MSG); /** * @brief Get the handle of the current/ default context @@ -1891,7 +1891,7 @@ DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxSetCurrent(hipCtx_t ctx)); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetDevice, hipCtxGetFlags, hipCtxPopCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetCurrent(hipCtx_t* ctx)); +hipError_t hipCtxGetCurrent(hipCtx_t* ctx) DEPRECATED(DEPRECATED_MSG); /** * @brief Get the handle of the device associated with current/default context @@ -1904,7 +1904,7 @@ DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetCurrent(hipCtx_t* ctx)); * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetDevice(hipDevice_t* device)); +hipError_t hipCtxGetDevice(hipDevice_t* device) DEPRECATED(DEPRECATED_MSG); /** * @brief Returns the approximate HIP api version. @@ -1923,7 +1923,7 @@ DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetDevice(hipDevice_t* device)); * @see hipCtxCreate, hipCtxDestroy, hipCtxGetDevice, hipCtxGetFlags, hipCtxPopCurrent, * hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion)); +hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiVersion) DEPRECATED(DEPRECATED_MSG); /** * @brief Set Cache configuration for a specific function @@ -1938,7 +1938,7 @@ DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetApiVersion(hipCtx_t ctx, int* apiV * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig)); +hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheConfig) DEPRECATED(DEPRECATED_MSG); /** * @brief Set L1/Shared cache partition. @@ -1953,7 +1953,7 @@ DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetCacheConfig(hipFuncCache_t* cacheC * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig)); +hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheConfig) DEPRECATED(DEPRECATED_MSG); /** * @brief Set Shared memory bank configuration. @@ -1968,7 +1968,7 @@ DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxSetCacheConfig(hipFuncCache_t cacheCo * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config)); +hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig config) DEPRECATED(DEPRECATED_MSG); /** * @brief Get Shared memory bank configuration. @@ -1983,7 +1983,7 @@ DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxSetSharedMemConfig(hipSharedMemConfig * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig)); +hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) DEPRECATED(DEPRECATED_MSG); /** * @brief Blocks until the default context has completed all preceding requested tasks. @@ -1996,7 +1996,7 @@ DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxGetDevice */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxSynchronize(void)); +hipError_t hipCtxSynchronize(void) DEPRECATED(DEPRECATED_MSG); /** * @brief Return flags used for creating default context. @@ -2008,7 +2008,7 @@ DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxSynchronize(void)); * @see hipCtxCreate, hipCtxDestroy, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxGetCurrent, * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetFlags(unsigned int* flags)); +hipError_t hipCtxGetFlags(unsigned int* flags) DEPRECATED(DEPRECATED_MSG); /** * @brief Enables direct access to memory allocations in a peer context. @@ -2029,7 +2029,7 @@ DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxGetFlags(unsigned int* flags)); * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice * @warning PeerToPeer support is experimental. */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags)); +hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) DEPRECATED(DEPRECATED_MSG); /** * @brief Disable direct access from current context's virtual address space to memory allocations @@ -2047,7 +2047,7 @@ DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, un * hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice * @warning PeerToPeer support is experimental. */ -DEPRECATED(DEPRECATED_MSG,hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx)); +hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx) DEPRECATED(DEPRECATED_MSG); /** * @brief Get the state of the primary context. From 39b5b07dc30bc2903ed8f662332acc48e8874474 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 20 Apr 2018 22:43:58 +0530 Subject: [PATCH 03/12] Added deprecation note in cuCtx section --- docs/markdown/hip_porting_driver_api.md | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/docs/markdown/hip_porting_driver_api.md b/docs/markdown/hip_porting_driver_api.md index 47cb7fb009..e0489d5c16 100644 --- a/docs/markdown/hip_porting_driver_api.md +++ b/docs/markdown/hip_porting_driver_api.md @@ -39,7 +39,8 @@ HIP as well as a recent versions of CUDA Runtime provide other mechanisms to acc The CUDA Runtime API unifies the Context API with the Device API. This simplifies the APIs and has little loss of functionality since each Context can contain a single device, and the benefits of multiple contexts has been replaced with other interfaces. HIP provides a context API to facilitate easy porting from existing Driver codes. In HIP, the Ctx functions largely provide an alternate syntax for changing the active device. -Most new applications will prefer to use `hipSetDevice` or the stream APIs. + +Most new applications will prefer to use `hipSetDevice` or the stream APIs , therefore HIP has marked hipCtx APIs as **deprecated**. Support for these APIs may not be available in future releases. For more details on deprecated APIs please refer [HIP deprecated APIs](https://github.com/ROCm-Developer-Tools/HIP/tree/master/docs/markdown/hip_deprecated_api_list.md). ## HIP Module and Ctx APIs From 433866eb60f0d8df1376c40882f535adb1364b3a Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Mon, 23 Apr 2018 15:25:35 -0400 Subject: [PATCH 04/12] Remove unsupported targets gfx801 and gfx802 These architectures are no longer supported. Should not compile for them. --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index e9a3451222..075b916a4d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -201,7 +201,7 @@ if(HIP_PLATFORM STREQUAL "hcc") execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS) set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS} -Wl,-Bsymbolic") - set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx801 --amdgpu-target=gfx802 --amdgpu-target=gfx803 --amdgpu-target=gfx900") + set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx803 --amdgpu-target=gfx900") if(COMPILE_HIP_ATP_MARKER) set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -L/opt/rocm/profiler/CXLActivityLogger/bin/x86_64 -lCXLActivityLogger") endif() From af586bbbf292d01dd6d305bc637aeefcaf2e0c86 Mon Sep 17 00:00:00 2001 From: Deven Desai Date: Tue, 24 Apr 2018 18:10:07 +0000 Subject: [PATCH 05/12] Checkin to fix bugs in math functions. This change fixes the following bugs that were discovered while debuggnig TF unit test failures (cwise_ops_test) 1. __hisinf and __hisnan routines Both had incorrect implementations. 2. abs A "long long" (64bit int) version was missing, resulting in the 32bit version being used for 64bit ints (which resulted in incorrect results, when the value passed in was outside the 32bit int range) 3. lgamma We seemed to have a custom version for the 'double' datatype (which was giving incorrect results). Replaced it with a call to the 'double' version of the underlying 'hc::precision_math::lgamma' --- include/hip/hcc_detail/math_functions.h | 1 + src/hip_fp16.cpp | 24 ++++++++++++++++-------- src/math_functions.cpp | 12 ++++-------- 3 files changed, 21 insertions(+), 16 deletions(-) diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h index 5482f34093..e717df07c1 100644 --- a/include/hip/hcc_detail/math_functions.h +++ b/include/hip/hcc_detail/math_functions.h @@ -57,6 +57,7 @@ __device__ float exp2f(float x); __device__ float expf(float x); __device__ float expm1f(float x); __device__ int abs(int x); +__device__ long long abs(long long x); __device__ float fabsf(float x); __device__ float fdimf(float x, float y); __device__ float fdividef(float x, float y); diff --git a/src/hip_fp16.cpp b/src/hip_fp16.cpp index e4c1b43786..3004b7805d 100644 --- a/src/hip_fp16.cpp +++ b/src/hip_fp16.cpp @@ -29,10 +29,6 @@ struct hipHalfHolder { }; }; -#define HINF 65504 - -__device__ static struct hipHalfHolder __hInfValue = {HINF}; - __device__ __half __hadd(__half a, __half b) { return a + b; } __device__ __half __hadd_sat(__half a, __half b) { return a + b; } @@ -63,9 +59,21 @@ __device__ bool __hge(__half a, __half b) { return a >= b ? true : false; } __device__ bool __hgt(__half a, __half b) { return a > b ? true : false; } -__device__ bool __hisinf(__half a) { return a == HINF ? true : false; } +__device__ bool __hisinf(__half a) { + hipHalfHolder hH; + hH.h = a; + // mask with 0x7fff to drop the sign bit + // 0x7c00 is bit pattern for inf (exp = 11111, significand = 0) + return ((hH.s & 0x7fff) == 0x7c00) ? true : false; +} -__device__ bool __hisnan(__half a) { return a > HINF ? true : false; } +__device__ bool __hisnan(__half a) { + hipHalfHolder hH; + hH.h = a; + // mask with 0x7fff to drop the sign bit + // 0x7cXX is bit pattern for inf (exp = 11111, significand = 0) + return ((hH.s & 0x7fff) > 0x7c00) ? true : false; +} __device__ bool __hle(__half a, __half b) { return a <= b ? true : false; } @@ -124,8 +132,8 @@ __device__ __half2 __hgt2(__half2 a, __half2 b) { __device__ __half2 __hisnan2(__half2 a) { __half2 c; - c.x = (a.x > HINF) ? (__half)1 : (__half)0; - c.y = (a.y > HINF) ? (__half)1 : (__half)0; + c.x = (__hisnan(a.x)) ? (__half)1 : (__half)0; + c.y = (__hisnan(a.y)) ? (__half)1 : (__half)0; return c; } diff --git a/src/math_functions.cpp b/src/math_functions.cpp index 3c0a7f6541..dedc40f2ae 100644 --- a/src/math_functions.cpp +++ b/src/math_functions.cpp @@ -56,6 +56,9 @@ __device__ float expm1f(float x) { return hc::precise_math::expm1f(x); } __device__ int abs(int x) { return x >= 0 ? x : -x; // TODO - optimize with OCML } +__device__ long long abs(long long x) { + return x >= 0 ? x : -x; +} __device__ float fabsf(float x) { return hc::precise_math::fabsf(x); } __device__ float fdimf(float x, float y) { return hc::precise_math::fdimf(x, y); } __device__ float fdividef(float x, float y) { return x / y; } @@ -220,14 +223,7 @@ __device__ double j0(double x) { return __hip_j0(x); } __device__ double j1(double x) { return __hip_j1(x); } __device__ double jn(int n, double x) { return __hip_jn(n, x); } __device__ double ldexp(double x, int exp) { return hc::precise_math::ldexp(x, exp); } -__device__ double lgamma(double x) { - double val = 0.0; - double y = x - 1; - while (y > 0) { - val += log(y--); - } - return val; -} +__device__ double lgamma(double x) { return hc::precise_math::lgamma(x); } __device__ long long int llrint(double x) { long long int y = hc::precise_math::round(x); return y; From 639001933cc468fd9beb99bd6bdc7d518195d304 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 25 Apr 2018 12:30:11 +0530 Subject: [PATCH 06/12] Fix gcc build issues for host code --- include/hip/hcc_detail/driver_types.h | 52 +++++++++++++----------- include/hip/hcc_detail/hip_runtime_api.h | 43 ++++---------------- include/hip/hcc_detail/texture_types.h | 8 ++-- include/hip/hip_runtime_api.h | 4 +- 4 files changed, 43 insertions(+), 64 deletions(-) diff --git a/include/hip/hcc_detail/driver_types.h b/include/hip/hcc_detail/driver_types.h index c25053ec32..661af64cd0 100644 --- a/include/hip/hcc_detail/driver_types.h +++ b/include/hip/hcc_detail/driver_types.h @@ -23,27 +23,31 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_DRIVER_TYPES_H #define HIP_INCLUDE_HIP_HCC_DETAIL_DRIVER_TYPES_H +#ifndef __cplusplus +#include +#endif + typedef void* hipDeviceptr_t; -enum hipChannelFormatKind { +typedef enum hipChannelFormatKind { hipChannelFormatKindSigned = 0, hipChannelFormatKindUnsigned = 1, hipChannelFormatKindFloat = 2, hipChannelFormatKindNone = 3 -}; +}hipChannelFormatKind; -struct hipChannelFormatDesc { +typedef struct hipChannelFormatDesc { int x; int y; int z; int w; enum hipChannelFormatKind f; -}; +}hipChannelFormatDesc; #define HIP_TRSF_NORMALIZED_COORDINATES 0x01 #define HIP_TRSF_READ_AS_INTEGER 0x00 #define HIP_TRSA_OVERRIDE_FORMAT 0x01 -enum hipArray_Format { +typedef enum hipArray_Format { HIP_AD_FORMAT_UNSIGNED_INT8 = 0x01, HIP_AD_FORMAT_UNSIGNED_INT16 = 0x02, HIP_AD_FORMAT_UNSIGNED_INT32 = 0x03, @@ -52,18 +56,18 @@ enum hipArray_Format { HIP_AD_FORMAT_SIGNED_INT32 = 0x0a, HIP_AD_FORMAT_HALF = 0x10, HIP_AD_FORMAT_FLOAT = 0x20 -}; +}hipArray_Format; -struct HIP_ARRAY_DESCRIPTOR { +typedef struct HIP_ARRAY_DESCRIPTOR { enum hipArray_Format format; unsigned int numChannels; size_t width; size_t height; unsigned int flags; size_t depth; -}; +}HIP_ARRAY_DESCRIPTOR; -struct hipArray { +typedef struct hipArray { void* data; // FIXME: generalize this struct hipChannelFormatDesc desc; unsigned int type; @@ -73,7 +77,7 @@ struct hipArray { struct HIP_ARRAY_DESCRIPTOR drvDesc; bool isDrv; unsigned int textureType; -}; +}hipArray; typedef struct hip_Memcpy2D { size_t height; @@ -115,17 +119,17 @@ typedef const struct hipMipmappedArray* hipMipmappedArray_const_t; /** * hip resource types */ -enum hipResourceType { +typedef enum hipResourceType { hipResourceTypeArray = 0x00, hipResourceTypeMipmappedArray = 0x01, hipResourceTypeLinear = 0x02, hipResourceTypePitch2D = 0x03 -}; +}hipResourceType; /** * hip texture resource view formats */ -enum hipResourceViewFormat { +typedef enum hipResourceViewFormat { hipResViewFormatNone = 0x00, hipResViewFormatUnsignedChar1 = 0x01, hipResViewFormatUnsignedChar2 = 0x02, @@ -161,12 +165,12 @@ enum hipResourceViewFormat { hipResViewFormatUnsignedBlockCompressed6H = 0x20, hipResViewFormatSignedBlockCompressed6H = 0x21, hipResViewFormatUnsignedBlockCompressed7 = 0x22 -}; +}hipResourceViewFormat; /** * HIP resource descriptor */ -struct hipResourceDesc { +typedef struct hipResourceDesc { enum hipResourceType resType; union { @@ -189,7 +193,7 @@ struct hipResourceDesc { size_t pitchInBytes; } pitch2D; } res; -}; +}hipResourceDesc; /** * hip resource view descriptor @@ -218,27 +222,27 @@ typedef enum hipMemcpyKind { 4 ///< Runtime will automatically determine copy-kind based on virtual addresses. } hipMemcpyKind; -struct hipPitchedPtr { +typedef struct hipPitchedPtr { void* ptr; size_t pitch; size_t xsize; size_t ysize; -}; +}hipPitchedPtr; -struct hipExtent { +typedef struct hipExtent { size_t width; // Width in elements when referring to array memory, in bytes when referring to // linear memory size_t height; size_t depth; -}; +}hipExtent; -struct hipPos { +typedef struct hipPos { size_t x; size_t y; size_t z; -}; +}hipPos; -struct hipMemcpy3DParms { +typedef struct hipMemcpy3DParms { hipArray_t srcArray; struct hipPos srcPos; struct hipPitchedPtr srcPtr; @@ -273,7 +277,7 @@ struct hipMemcpy3DParms { size_t srcXInBytes; size_t srcY; size_t srcZ; -}; +}hipMemcpy3DParms; static __inline__ struct hipPitchedPtr make_hipPitchedPtr(void* d, size_t p, size_t xsz, size_t ysz) { diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 8674824c48..e78ccb02c4 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -30,7 +30,6 @@ THE SOFTWARE. */ #include #include -#include #ifndef GENERIC_GRID_LAUNCH #define GENERIC_GRID_LAUNCH 1 @@ -1314,7 +1313,7 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz * hipMemcpyFromSymbolAsync */ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t sizeBytes, - size_t offset = 0, hipMemcpyKind kind = hipMemcpyHostToDevice); + size_t offset __dparm(0), hipMemcpyKind kind __dparm(hipMemcpyHostToDevice)); /** @@ -1343,13 +1342,13 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t siz * hipMemcpyFromSymbolAsync */ hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_t sizeBytes, - size_t offset, hipMemcpyKind kind, hipStream_t stream = 0); + size_t offset, hipMemcpyKind kind, hipStream_t stream __dparm(0)); hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t sizeBytes, - size_t offset = 0, hipMemcpyKind kind = hipMemcpyDeviceToHost); + size_t offset __dparm(0), hipMemcpyKind kind __dparm( hipMemcpyDeviceToHost )); hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t sizeBytes, - size_t offset, hipMemcpyKind kind, hipStream_t stream = 0); + size_t offset, hipMemcpyKind kind, hipStream_t stream __dparm(0)); /** * @brief Copy data from src to dst asynchronously. @@ -1379,13 +1378,8 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t si * hipMemcpyFromArrayAsync, hipMemcpy2DFromArrayAsync, hipMemcpyToSymbolAsync, * hipMemcpyFromSymbolAsync */ -#if __cplusplus hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, - hipStream_t stream = 0); -#else -hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, - hipStream_t stream); -#endif + hipStream_t stream __dparm(0)); /** * @brief Fills the first sizeBytes bytes of the memory area pointed to by dest with the constant @@ -1424,11 +1418,7 @@ hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, size_t sizeByte * @param[in] stream - Stream identifier * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorMemoryFree */ -#if __cplusplus -hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream = 0); -#else -hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream); -#endif +hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream __dparm(0)); /** * @brief Fills the memory area pointed to by dst with the constant value. @@ -1484,13 +1474,8 @@ hipError_t hipMemPtrGetInfo(void* ptr, size_t* size); * * @see hipMalloc, hipMallocPitch, hipFree, hipFreeArray, hipHostMalloc, hipHostFree */ -#if __cplusplus hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width, - size_t height = 0, unsigned int flags = hipArrayDefault); -#else -hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, size_t width, - size_t height, unsigned int flags); -#endif + size_t height __dparm(0), unsigned int flags __dparm(hipArrayDefault)); hipError_t hipArrayCreate(hipArray** pHandle, const HIP_ARRAY_DESCRIPTOR* pAllocateArray); hipError_t hipArray3DCreate(hipArray_t* array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray); @@ -1558,13 +1543,8 @@ hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy); * @see hipMemcpy, hipMemcpyToArray, hipMemcpy2DToArray, hipMemcpyFromArray, hipMemcpyToSymbol, * hipMemcpyAsync */ -#if __cplusplus hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, - size_t height, hipMemcpyKind kind, hipStream_t stream = 0); -#else -hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, - size_t height, hipMemcpyKind kind, hipStream_t stream); -#endif + size_t height, hipMemcpyKind kind, hipStream_t stream __dparm(0)); /** * @brief Copies data between host and device. @@ -1776,13 +1756,8 @@ hipError_t hipMemcpyPeer(void* dst, int dstDeviceId, const void* src, int srcDev * * @returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDevice */ -#if __cplusplus hipError_t hipMemcpyPeerAsync(void* dst, int dstDeviceId, const void* src, int srcDevice, - size_t sizeBytes, hipStream_t stream = 0); -#else -hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int srcDevice, - size_t sizeBytes, hipStream_t stream); -#endif + size_t sizeBytes, hipStream_t stream __dparm(0)); #endif diff --git a/include/hip/hcc_detail/texture_types.h b/include/hip/hcc_detail/texture_types.h index 3d0d44bd55..bc334de248 100644 --- a/include/hip/hcc_detail/texture_types.h +++ b/include/hip/hcc_detail/texture_types.h @@ -70,7 +70,7 @@ enum hipTextureReadMode { hipReadModeElementType = 0, hipReadModeNormalizedFloat /** * hip texture reference */ -struct textureReference { +typedef struct textureReference { int normalized; enum hipTextureFilterMode filterMode; enum hipTextureAddressMode addressMode[3]; // Texture address mode for up to 3 dimensions @@ -85,12 +85,12 @@ struct textureReference { hipTextureObject_t textureObject; int numChannels; enum hipArray_Format format; -}; +}textureReference; /** * hip texture descriptor */ -struct hipTextureDesc { +typedef struct hipTextureDesc { enum hipTextureAddressMode addressMode[3]; // Texture address mode for up to 3 dimensions enum hipTextureFilterMode filterMode; enum hipTextureReadMode readMode; @@ -102,6 +102,6 @@ struct hipTextureDesc { float mipmapLevelBias; float minMipmapLevelClamp; float maxMipmapLevelClamp; -}; +}hipTextureDesc; #endif diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index 8d27724ad6..bfde2e942d 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -120,14 +120,14 @@ typedef struct hipDeviceProp_t { /** * Memory type (for pointer attributes) */ -enum hipMemoryType { +typedef enum hipMemoryType { hipMemoryTypeHost, ///< Memory is physically located on host hipMemoryTypeDevice, ///< Memory is physically located on device. (see deviceId for specific ///< device) hipMemoryTypeArray, ///< Array memory, physically located on device. (see deviceId for specific ///< device) hipMemoryTypeUnified ///< Not used currently -}; +}hipMemoryType; /** From 4be4cf644f576aa3b5f0455b7e38add755dbecb7 Mon Sep 17 00:00:00 2001 From: Deven Desai Date: Wed, 25 Apr 2018 17:53:21 +0000 Subject: [PATCH 07/12] Checkin to add unit tests for fixes in my previous commit This change adds unit tests for 1. __hisinf 2. __hisnan 2. abs(long long) 3. lgamma(double) --- tests/src/deviceLib/hipMathFunctions.cpp | 138 +++++++++++++++++++++ tests/src/deviceLib/hipTestHalf.cpp | 145 +++++++++++++++++++++++ 2 files changed, 283 insertions(+) diff --git a/tests/src/deviceLib/hipMathFunctions.cpp b/tests/src/deviceLib/hipMathFunctions.cpp index ba3707bc59..ff69087f82 100644 --- a/tests/src/deviceLib/hipMathFunctions.cpp +++ b/tests/src/deviceLib/hipMathFunctions.cpp @@ -20,12 +20,150 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * RUN: %t + * HIT_END + */ #include "hip/hip_runtime.h" #include "test_common.h" +#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ + +__global__ void kernel_abs_int64(hipLaunchParm lp, long long *input, long long *output) { + int tx = threadIdx.x; + output[tx] = abs(input[tx]); +} + +__global__ void kernel_lgamma_double(hipLaunchParm lp, double *input, double *output) { + int tx = threadIdx.x; + output[tx] = lgamma(input[tx]); +} + +#endif + +#define CHECK_LGAMMA_DOUBLE(IN, OUT, EXP) \ + { \ + if (OUT != EXP) { \ + failed("check_abs_int64 failed on %f (output = %f, expected = %fd)\n", IN, OUT, EXP); \ + } \ + } + +#define CHECK_ABS_INT64(IN, OUT, EXP) \ + { \ + if (OUT != EXP) { \ + failed("check_abs_int64 failed on %lld (output = %lld, expected = %lld)\n", IN, OUT, EXP); \ + } \ + } + +void check_lgamma_double() { + + using datatype_t = double; + + const int NUM_INPUTS = 8; + auto memsize = NUM_INPUTS * sizeof(datatype_t); + + // allocate memories + datatype_t *inputCPU = (datatype_t *) malloc(memsize); + datatype_t *outputCPU = (datatype_t *) malloc(memsize); + datatype_t *inputGPU = nullptr; hipMalloc((void**)&inputGPU, memsize); + datatype_t *outputGPU = nullptr; hipMalloc((void**)&outputGPU, memsize); + + // populate input + inputCPU[0] = -3.5; + inputCPU[0] = -2.5; + inputCPU[0] = -1.5; + inputCPU[0] = -0.5; + inputCPU[0] = 0.5; + inputCPU[0] = 1.5; + inputCPU[0] = 2.5; + inputCPU[0] = 3.5; + + // copy inputs to device + hipMemcpy(inputGPU, inputCPU, memsize, hipMemcpyHostToDevice); + + // launch kernel + hipLaunchKernel(kernel_lgamma_double, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU); + + // copy outputs from device + hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost); + + // check outputs + for (int i=0; i #include #include "hip/hip_runtime.h" @@ -59,8 +65,143 @@ __global__ void __half2Math(hipLaunchParm lp, __half2* A, __half2* B, __half2* C c = __hmul2_sat(b, c); } +__global__ void kernel_hisnan(hipLaunchParm lp, __half* input, int* output) { + int tx = threadIdx.x; + output[tx] = __hisnan(input[tx]); +} + +__global__ void kernel_hisinf(hipLaunchParm lp, __half* input, int* output) { + int tx = threadIdx.x; + output[tx] = __hisinf(input[tx]); +} + #endif + +__half host_ushort_as_half(unsigned short s) { + union {__half h; unsigned short s; } converter; + converter.s = s; + return converter.h; +} + + +void check_hisnan(int NUM_INPUTS, __half* inputCPU, __half* inputGPU) { + + // allocate memory + auto memsize = NUM_INPUTS * sizeof(int); + int* outputGPU = nullptr; + hipMalloc((void**)&outputGPU, memsize); + + // launch the kernel + hipLaunchKernel(kernel_hisnan, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU); + + // copy output from device + int* outputCPU = (int*) malloc(memsize); + hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost); + + // check output + for (int i=0; i Date: Fri, 27 Apr 2018 12:59:51 +0000 Subject: [PATCH 08/12] Fixing a copy/paste error in my previous checkin --- tests/src/deviceLib/hipMathFunctions.cpp | 11 +++-------- 1 file changed, 3 insertions(+), 8 deletions(-) diff --git a/tests/src/deviceLib/hipMathFunctions.cpp b/tests/src/deviceLib/hipMathFunctions.cpp index ff69087f82..7fe0003672 100644 --- a/tests/src/deviceLib/hipMathFunctions.cpp +++ b/tests/src/deviceLib/hipMathFunctions.cpp @@ -71,14 +71,9 @@ void check_lgamma_double() { datatype_t *outputGPU = nullptr; hipMalloc((void**)&outputGPU, memsize); // populate input - inputCPU[0] = -3.5; - inputCPU[0] = -2.5; - inputCPU[0] = -1.5; - inputCPU[0] = -0.5; - inputCPU[0] = 0.5; - inputCPU[0] = 1.5; - inputCPU[0] = 2.5; - inputCPU[0] = 3.5; + for (int i=0; i Date: Mon, 30 Apr 2018 09:36:37 -0500 Subject: [PATCH 09/12] Update hip_bugs.md Remove several out-date links in hip_bugs.md document --- docs/markdown/hip_bugs.md | 4 ---- 1 file changed, 4 deletions(-) diff --git a/docs/markdown/hip_bugs.md b/docs/markdown/hip_bugs.md index ef6fd5decb..46dfa6d0d2 100644 --- a/docs/markdown/hip_bugs.md +++ b/docs/markdown/hip_bugs.md @@ -1,10 +1,6 @@ # HIP Bugs - -- [Can't find kernels inside dynamic linked library](#cant-find-kernels-inside-dynamic-linked-library) -- [What is the current limitation of HIP Generic Grid Launch method?](#what-is-the-current-limitation-of-hip-generic-grid-launch-method) -- [Errors related to `no matching constructor`](#errors-related-to-no-matching-constructor) - [HIP is more restrictive in enforcing restrictions](#hip-is-more-restrictive-in-enforcing-restrictions) From 9de5f23d54fcb2ba528837af057210f775730c12 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 2 May 2018 11:56:37 +0530 Subject: [PATCH 10/12] Fix texture 3D for HIP/NVCC --- include/hip/hcc_detail/hip_runtime_api.h | 4 ++-- include/hip/nvcc_detail/hip_runtime_api.h | 19 ++++++++++++++++++- src/hip_memory.cpp | 4 ++-- 3 files changed, 22 insertions(+), 5 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 0bed1ea239..8becab3c9c 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -1481,7 +1481,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, si size_t height __dparm(0), unsigned int flags __dparm(hipArrayDefault)); hipError_t hipArrayCreate(hipArray** pHandle, const HIP_ARRAY_DESCRIPTOR* pAllocateArray); -hipError_t hipArray3DCreate(hipArray_t* array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray); +hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray); hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent); @@ -1507,7 +1507,7 @@ hipError_t hipFreeArray(hipArray* array); * @see hipMalloc, hipMallocPitch, hipFree, hipFreeArray, hipHostMalloc, hipHostFree */ -hipError_t hipMalloc3DArray(hipArray_t* array, const struct hipChannelFormatDesc* desc, +hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* desc, struct hipExtent extent, unsigned int flags); /** * @brief Copies data between host and device. diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index d391c31db8..7f64845fbe 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -150,13 +150,20 @@ typedef CUfunction hipFunction_t; typedef CUdeviceptr hipDeviceptr_t; typedef struct cudaArray hipArray; typedef struct cudaArray* hipArray_const_t; +typedef cudaMemcpy3DParms hipMemcpy3DParms; +//#define hipMemcpy3DParms cudaMemcpy3DParms #define hipArrayDefault cudaArrayDefault typedef cudaTextureObject_t hipTextureObject_t; typedef cudaSurfaceObject_t hipSurfaceObject_t; -#define hipTextureType2D cudaTextureType2D; +#define hipTextureType2D cudaTextureType2D +#define hipTextureType3D cudaTextureType3D #define hipDeviceMapHost cudaDeviceMapHost +#define hipExtent cudaExtent +#define make_hipExtent make_cudaExtent +#define make_hipPos make_cudaPos +#define make_hipPitchedPtr make_cudaPitchedPtr // Flags that can be used with hipStreamCreateWithFlags #define hipStreamDefault cudaStreamDefault #define hipStreamNonBlocking cudaStreamNonBlocking @@ -399,6 +406,11 @@ inline static hipError_t hipMallocArray(hipArray** array, const struct hipChanne return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); } +inline static hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* desc, + struct hipExtent extent, unsigned int flags) { + return hipCUDAErrorTohipError(cudaMalloc3DArray(array, desc, extent, flags)); +} + inline static hipError_t hipFreeArray(hipArray* array) { return hipCUDAErrorTohipError(cudaFreeArray(array)); } @@ -531,6 +543,11 @@ inline static hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, cudaMemcpy2D(dst, dpitch, src, spitch, width, height, hipMemcpyKindToCudaMemcpyKind(kind))); } +inline static hipError_t hipMemcpy3D(const struct hipMemcpy3DParms *p) +{ + return hipCUDAErrorTohipError(cudaMemcpy3D(p)); +} + inline static hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index fbbba027de..3219dac3db 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -641,7 +641,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, si return ihipLogStatus(hip_status); } -hipError_t hipArray3DCreate(hipArray_t* array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) { +hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) { HIP_INIT_SPECIAL_API((TRACE_MEM), array, pAllocateArray); hipError_t hip_status = hipSuccess; @@ -764,7 +764,7 @@ hipError_t hipArray3DCreate(hipArray_t* array, const HIP_ARRAY_DESCRIPTOR* pAllo return ihipLogStatus(hip_status); } -hipError_t hipMalloc3DArray(hipArray_t* array, const struct hipChannelFormatDesc* desc, +hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc* desc, struct hipExtent extent, unsigned int flags) { From b4585a2a90ab5a55e702425036f9ca6dc55cf8c4 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 2 May 2018 12:54:17 +0530 Subject: [PATCH 11/12] Cleaned up --- include/hip/nvcc_detail/hip_runtime_api.h | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/include/hip/nvcc_detail/hip_runtime_api.h b/include/hip/nvcc_detail/hip_runtime_api.h index 7f64845fbe..6f222648a2 100644 --- a/include/hip/nvcc_detail/hip_runtime_api.h +++ b/include/hip/nvcc_detail/hip_runtime_api.h @@ -150,8 +150,7 @@ typedef CUfunction hipFunction_t; typedef CUdeviceptr hipDeviceptr_t; typedef struct cudaArray hipArray; typedef struct cudaArray* hipArray_const_t; -typedef cudaMemcpy3DParms hipMemcpy3DParms; -//#define hipMemcpy3DParms cudaMemcpy3DParms +#define hipMemcpy3DParms cudaMemcpy3DParms #define hipArrayDefault cudaArrayDefault typedef cudaTextureObject_t hipTextureObject_t; From 6411ca1f6d0e24d41ef182614da71e378bd14dff Mon Sep 17 00:00:00 2001 From: Lakhan Singh Date: Thu, 3 May 2018 09:27:50 +0530 Subject: [PATCH 12/12] Null checks added for hipmallocpitch and hipmemcpy apis --- src/hip_memory.cpp | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 3219dac3db..48a8da81ca 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -339,6 +339,11 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) { // width in bytes hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height, size_t depth) { hipError_t hip_status = hipSuccess; + if(ptr==NULL) + { + hip_status=hipErrorInvalidValue; + return hip_status; + } // hardcoded 128 bytes *pitch = ((((int)width - 1) / 128) + 1) * 128; const size_t sizeBytes = (*pitch) * height; @@ -1132,7 +1137,11 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind hc::completion_future marker; hipError_t e = hipSuccess; - + if(dst==NULL || src==NULL) + { + e=hipErrorInvalidValue; + return e; + } try { stream->locked_copySync(dst, src, sizeBytes, kind); } catch (ihipException& ex) {