From 4a273380c8ae37734d25d0294148bd0a327b42a9 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Sat, 30 Jun 2018 11:40:32 +0530 Subject: [PATCH 1/7] Updated indentation [ROCm/clr commit: cc14d6440f9cd6bebf6b0a4acc58c2f9079c853c] --- .../docs/markdown/hip_deprecated_api_list.md | 32 +++++++++---------- 1 file changed, 16 insertions(+), 16 deletions(-) diff --git a/projects/clr/hipamd/docs/markdown/hip_deprecated_api_list.md b/projects/clr/hipamd/docs/markdown/hip_deprecated_api_list.md index a96f7f4d3f..6a9ed48839 100644 --- a/projects/clr/hipamd/docs/markdown/hip_deprecated_api_list.md +++ b/projects/clr/hipamd/docs/markdown/hip_deprecated_api_list.md @@ -4,19 +4,19 @@ 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 +### hipCtxCreate +### hipCtxDestroy +### hipCtxPopCurrent +### hipCtxPushCurrent +### hipCtxSetCurrent +### hipCtxGetCurrent +### hipCtxGetDevice +### hipCtxGetApiVersion +### hipCtxGetCacheConfig +### hipCtxSetCacheConfig +### hipCtxSetSharedMemConfig +### hipCtxGetSharedMemConfig +### hipCtxSynchronize +### hipCtxGetFlags +### hipCtxEnablePeerAccess +### hipCtxDisablePeerAccess From c8da0cde2708d47ce6e7681867e11dd33171e457 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Sat, 30 Jun 2018 11:42:17 +0530 Subject: [PATCH 2/7] Updated heading [ROCm/clr commit: 805906d2ad19c5d58b02818193826459b3a88c96] --- projects/clr/hipamd/docs/markdown/hip_deprecated_api_list.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/docs/markdown/hip_deprecated_api_list.md b/projects/clr/hipamd/docs/markdown/hip_deprecated_api_list.md index 6a9ed48839..dfb202c8ee 100644 --- a/projects/clr/hipamd/docs/markdown/hip_deprecated_api_list.md +++ b/projects/clr/hipamd/docs/markdown/hip_deprecated_api_list.md @@ -1,4 +1,4 @@ -# HIP Deprecated API List +# HIP Deprecated APIs ## HIP Context API From a796eedde99e4a3eedcef9501b2d9d4f28e17727 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 2 Jul 2018 14:32:11 +0530 Subject: [PATCH 3/7] Revert "Use memcpy kernel for all pinned memory cases in hipMemcpy2DAsync" [ROCm/clr commit: f554e48db35e6e1bd4b2ee44a756526755e13609] --- projects/clr/hipamd/src/hip_memory.cpp | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 787e49683b..d6c04ae98c 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -1680,12 +1680,9 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp actualDest = pinnedPtr; } } -#if 0 if((width == dpitch) && (width == spitch)) { hip_internal::memcpyAsync(dst, src, width*height, kind, stream); - } else -#endif - { + } else { try { if(!isLocked){ for (int i = 0; i < height; ++i) From 676c2893104f48a7c41a62f13d4879efd7522808 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 3 Jul 2018 08:54:17 +0530 Subject: [PATCH 4/7] Fixed offset null check in bind texture functions [ROCm/clr commit: 77cbd4476de20270d577fdc54f2f7abd9b3c19b9] --- projects/clr/hipamd/src/hip_texture.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/src/hip_texture.cpp b/projects/clr/hipamd/src/hip_texture.cpp index 24c6eef3af..d6caf853de 100644 --- a/projects/clr/hipamd/src/hip_texture.cpp +++ b/projects/clr/hipamd/src/hip_texture.cpp @@ -389,7 +389,8 @@ hipError_t ihipBindTextureImpl(int dim, enum hipTextureReadMode readMode, size_t enum hipTextureFilterMode filterMode = tex->filterMode; int normalizedCoords = tex->normalized; hipTextureObject_t& textureObject = tex->textureObject; - *offset = 0; + if(offset != nullptr) + *offset = 0; auto ctx = ihipGetTlsDefaultCtx(); if (ctx) { hc::accelerator acc = ctx->getDevice()->_acc; @@ -459,7 +460,8 @@ hipError_t ihipBindTexture2DImpl(int dim, enum hipTextureReadMode readMode, size enum hipTextureFilterMode filterMode = tex->filterMode; int normalizedCoords = tex->normalized; hipTextureObject_t& textureObject = tex->textureObject; - *offset = 0; + if(offset != nullptr) + *offset = 0; auto ctx = ihipGetTlsDefaultCtx(); if (ctx) { hc::accelerator acc = ctx->getDevice()->_acc; From c548c3d83320517128b3ba4eab012a34a3f3c84a Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 5 Jul 2018 23:00:41 +0530 Subject: [PATCH 5/7] Added tex2dlayered mapping for HIP/NVCC [ROCm/clr commit: f080abe5a5b00f4f124ae10c1d13406d0eaeb72f] --- projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 3b8a3661f7..afe7f11f52 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -163,6 +163,7 @@ typedef cudaSurfaceObject_t hipSurfaceObject_t; #define hipTextureType1D cudaTextureType1D #define hipTextureType1DLayered cudaTextureType1DLayered #define hipTextureType2D cudaTextureType2D +#define hipTextureType2DLayered cudaTextureType2DLayered #define hipTextureType3D cudaTextureType3D #define hipDeviceMapHost cudaDeviceMapHost From 61b632944bc0c0ebcc5d893c901d5d39ee13f56a Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 5 Jul 2018 23:11:39 +0530 Subject: [PATCH 6/7] Corrected enum type [ROCm/clr commit: b15b08d0dba9a28ae8aa755236703691cad028f2] --- .../clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 3b8a3661f7..557b15e9a2 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -1168,20 +1168,20 @@ inline static hipError_t hipOccupancyMaxPotentialBlockSize(int* minGridSize, int return hipCUDAErrorTohipError(cerror); } -template +template inline static hipError_t hipBindTexture(size_t* offset, const struct texture& tex, const void* devPtr, size_t size = UINT_MAX) { return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, size)); } -template +template inline static hipError_t hipBindTexture(size_t* offset, struct texture& tex, const void* devPtr, const struct hipChannelFormatDesc& desc, size_t size = UINT_MAX) { return hipCUDAErrorTohipError(cudaBindTexture(offset, tex, devPtr, desc, size)); } -template +template inline static hipError_t hipUnbindTexture(struct texture* tex) { return hipCUDAErrorTohipError(cudaUnbindTexture(tex)); } @@ -1198,7 +1198,7 @@ inline static hipError_t hipBindTextureToArray(struct texture& return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array, desc)); } -template +template inline static hipError_t hipBindTextureToArray(struct texture& tex, hipArray_const_t array) { return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array)); From a6f74d51b944c7e07b09b093c2244a656fe1240d Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 10 Jul 2018 20:37:54 +0300 Subject: [PATCH 7/7] [HIPIFY][docs] Update README.md after testing with new LLVM releases 5.0.2 and 6.0.1 [ROCm/clr commit: 2806dff2be9e132f6011fe5ac9fcbe756e0be733] --- projects/clr/hipamd/hipify-clang/README.md | 23 +++++++++++++++------- 1 file changed, 16 insertions(+), 7 deletions(-) diff --git a/projects/clr/hipamd/hipify-clang/README.md b/projects/clr/hipamd/hipify-clang/README.md index c6f59a0251..e099085ca6 100644 --- a/projects/clr/hipamd/hipify-clang/README.md +++ b/projects/clr/hipamd/hipify-clang/README.md @@ -31,7 +31,9 @@ | 4.0.1 | 8.0 | | 5.0.0 | 8.0 | | 5.0.1 | 8.0 | +| 5.0.2 | 8.0 | | 6.0.0 | 9.0 | +| 6.0.1 | 9.0 | In most cases, you can get a suitable version of LLVM+CLANG with your package manager. @@ -80,7 +82,7 @@ To run it: cmake \ -DCMAKE_INSTALL_PREFIX=../dist \ -DLLVM_SOURCE_DIR=../llvm \ - -DCMAKE_BUILD_TYPE=Debug \ + -DCMAKE_BUILD_TYPE=Release \ -Thost=x64 \ ../llvm @@ -114,9 +116,9 @@ To run it: * Starting with LLVM 6.0.0 path to llvm-lit.py script should be specified by the `LLVM_EXTERNAL_LIT` option: - `-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.0/build/Debug/bin/llvm-lit.py`, + `-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.0/build/Release/bin/llvm-lit.py`, - where `f:/LLVM/6.0.0/build/Debug` is LLVM build directory. + where `f:/LLVM/6.0.0/build/Release` is LLVM build directory. 7. Build with the `HIPIFY_CLANG_TESTS` option turned on: -DHIPIFY_CLANG_TESTS=1. 8. `make test-hipify` @@ -124,9 +126,15 @@ To run it: ### Windows -On Windows the following tested configuration is recommended: +On Windows the following configurations are tested: -LLVM 6.0.0 (exact), CUDA 9.0 (exact), cudnn-9.0 (exact), Python 3.6 (min), cmake 3.10 (min), Visual Studio 15.5 2017 (min). +LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn-9.0 + +LLVM 5.0.0 - 5.0.2, CUDA 8.0, cudnn-8.0 + +Build system for the above configurations: + +Python 3.6 (min), cmake 3.10 (min), Visual Studio 15.5 2017 (min). Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 15 2017`: @@ -134,13 +142,14 @@ Here is an example of building `hipify-clang` with testing support on `Windows 1 cmake -G "Visual Studio 15 2017 Win64" \ -DHIPIFY_CLANG_TESTS=1 \ - -DCMAKE_BUILD_TYPE=Debug \ + -DCMAKE_BUILD_TYPE=Release \ -DCMAKE_INSTALL_PREFIX=../dist \ -DCMAKE_PREFIX_PATH=f:/LLVM/6.0.0/dist \ -DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0" \ -DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0" \ -DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.1 \ - -DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.0/build/Debug/bin/llvm-lit.py \ + -DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.0/build/Release/bin/llvm-lit.py \ + -Thost=x64 .. ``` A corresponding successful output: