From 25c3e36bede0ee9f47a3c17169085d820ba5bcc5 Mon Sep 17 00:00:00 2001 From: Joseph Greathouse Date: Tue, 17 Mar 2020 03:30:51 -0500 Subject: [PATCH 01/11] Fix maxSharedMemoryPerMultiProcessor attribute (#1927) The maxSharedMemoryPerMultiProcessor attribute is meant to describe the number of bytes of shared memory (LDS space in AMD terminology) in each SM (CU in AMD terminology). For instance, on AMD GPUs this is often 64KB per CU, and some Nvidia GPUs it's 96KB per SM. This shared memory is a different address space from the normal global memory. However, the current HIP-HCC properties fill this in with a size that matches the totalGlboalMem property. This gives a drastically too-high calculation for the amount of LDS space that each CU has -- tens of GBs vs. 10s of KBs. This patch fixes this by pulling the maxSharedMemoryPerMultiProcessor property from the HSA pool that describes how much workgroup-local space is available on each CU. The HSA runtime eventually pulls this from the topology information about LDSSizeInKB, defined as "Size of Local Data Store in Kilobytes per SIMD". Previously, this HSA query was used to fill in the value of the sharedMemPerBlock property. On today's AMD GPUs, we know that the amount of LDS avaialble to the workgroup is identical to the amount of LDS space in the CU. However, in the future this may differ. As such, this patch changes around the order and fills in the "PerMultiProcessor" property from the HSA query (since what's what the query is defined to return), and then separately fills in the "PerBlock" property as we know it. [ROCm/hip-tests commit: 449e2c931bb7282b516f358391ae0290859eb566] --- projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp b/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp index e17f19675a..14faa7671b 100644 --- a/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp @@ -56,6 +56,7 @@ void printCompilerInfo() { #endif } +double bytesToKB(size_t s) { return (double)s / (1024.0); } double bytesToGB(size_t s) { return (double)s / (1024.0 * 1024.0 * 1024.0); } #define printLimit(w1, limit, units) \ @@ -97,7 +98,7 @@ void printDeviceProp(int deviceId) { cout << setw(w1) << "totalGlobalMem: " << fixed << setprecision(2) << bytesToGB(props.totalGlobalMem) << " GB" << endl; cout << setw(w1) << "maxSharedMemoryPerMultiProcessor: " << fixed << setprecision(2) - << bytesToGB(props.maxSharedMemoryPerMultiProcessor) << " GB" << endl; + << bytesToKB(props.maxSharedMemoryPerMultiProcessor) << " KB" << endl; cout << setw(w1) << "totalConstMem: " << props.totalConstMem << endl; cout << setw(w1) << "sharedMemPerBlock: " << (float)props.sharedMemPerBlock / 1024.0 << " KB" << endl; From d709d913cb9015ab215f9c8cedc83af4bbeb105b Mon Sep 17 00:00:00 2001 From: Sarbojit2019 <52527887+SarbojitAMD@users.noreply.github.com> Date: Fri, 27 Mar 2020 14:08:30 +0530 Subject: [PATCH 02/11] Fix few memory leaks in HIP (#1969) [ROCm/hip-tests commit: 5144ee8c0402ec168bec3c234945aa53ef7b880c] --- .../hip-tests/samples/0_Intro/module_api/defaultDriver.cpp | 4 ++-- .../hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp | 4 ++-- projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp | 4 ++-- .../hip-tests/samples/0_Intro/module_api_global/runKernel.cpp | 4 ++-- 4 files changed, 8 insertions(+), 8 deletions(-) diff --git a/projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp b/projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp index ea36aabcf4..af8b413ac2 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp +++ b/projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp @@ -80,8 +80,8 @@ int main() { hipFree(Ad); hipFree(Bd); - delete A; - delete B; + delete[] A; + delete[] B; hipCtxDestroy(context); return 0; } diff --git a/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp b/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp index 38cf0d414c..90e569c5bc 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp +++ b/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp @@ -107,8 +107,8 @@ int main() { hipFree(Ad); hipFree(Bd); - delete A; - delete B; + delete[] A; + delete[] B; hipCtxDestroy(context); return 0; } diff --git a/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp b/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp index a011b42666..1093b0dd54 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp +++ b/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp @@ -99,8 +99,8 @@ int main() { hipFree(Ad); hipFree(Bd); - delete A; - delete B; + delete[] A; + delete[] B; hipCtxDestroy(context); return 0; } diff --git a/projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp b/projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp index 3a2804b7a2..4a2d49144c 100644 --- a/projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp +++ b/projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp @@ -154,8 +154,8 @@ int main() { hipFree(Ad); hipFree(Bd); - delete A; - delete B; + delete[] A; + delete[] B; hipCtxDestroy(context); return 0; } From e9ffe1fe87d13d3784407562ebe554a775d0ccc4 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 7 Apr 2020 01:32:52 -0700 Subject: [PATCH 03/11] Rename hipDrvOccupancy to hipModuleOccupancy and match CUDA syntax (#1943) [ROCm/hip-tests commit: 177457e54c9c874467ca5676c10d62e9ed6f2846] --- .../samples/2_Cookbook/13_occupancy/occupancy.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/hip-tests/samples/2_Cookbook/13_occupancy/occupancy.cpp b/projects/hip-tests/samples/2_Cookbook/13_occupancy/occupancy.cpp index a9f4e198b0..01fa7aafed 100644 --- a/projects/hip-tests/samples/2_Cookbook/13_occupancy/occupancy.cpp +++ b/projects/hip-tests/samples/2_Cookbook/13_occupancy/occupancy.cpp @@ -56,9 +56,9 @@ void launchKernel(float* C, float* A, float* B, bool manual){ const unsigned threadsperblock = 32; const unsigned blocks = (NUM/threadsperblock)+1; - uint32_t mingridSize = 0; - uint32_t gridSize = 0; - uint32_t blockSize = 0; + int mingridSize = 0; + int gridSize = 0; + int blockSize = 0; if (manual){ blockSize = threadsperblock; @@ -86,7 +86,7 @@ void launchKernel(float* C, float* A, float* B, bool manual){ printf("kernel Execution time = %6.3fms\n", eventMs); //Calculate Occupancy - uint32_t numBlock = 0; + int numBlock = 0; HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, multiply, blockSize, 0)); if(devProp.maxThreadsPerMultiProcessor){ From a6c7b1ca12add9a9a324051054f80604e7e751dc Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 7 Apr 2020 11:33:19 +0300 Subject: [PATCH 04/11] [HIP][doc] Update docs due to moving of HIPIFY to a separate repo (#2001) [ROCm/hip-tests commit: a93ae7cbedca6821f8f5bdb7e4a98cded613a9a7] --- projects/hip-tests/samples/0_Intro/square/README.md | 2 +- .../hip-tests/samples/2_Cookbook/0_MatrixTranspose/Readme.md | 2 +- projects/hip-tests/samples/2_Cookbook/10_inline_asm/Readme.md | 2 +- .../samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md | 2 +- projects/hip-tests/samples/2_Cookbook/1_hipEvent/Readme.md | 2 +- projects/hip-tests/samples/2_Cookbook/3_shared_memory/Readme.md | 2 +- projects/hip-tests/samples/2_Cookbook/4_shfl/Readme.md | 2 +- projects/hip-tests/samples/2_Cookbook/5_2dshfl/Readme.md | 2 +- .../hip-tests/samples/2_Cookbook/6_dynamic_shared/Readme.md | 2 +- projects/hip-tests/samples/2_Cookbook/7_streams/Readme.md | 2 +- projects/hip-tests/samples/2_Cookbook/9_unroll/Readme.md | 2 +- 11 files changed, 11 insertions(+), 11 deletions(-) diff --git a/projects/hip-tests/samples/0_Intro/square/README.md b/projects/hip-tests/samples/0_Intro/square/README.md index 7a9e04fc5f..c185903993 100644 --- a/projects/hip-tests/samples/0_Intro/square/README.md +++ b/projects/hip-tests/samples/0_Intro/square/README.md @@ -1,6 +1,6 @@ # Square.md -Simple test which shows how to use hipify to port CUDA code to HIP. +Simple test which shows how to use hipify-perl to port CUDA code to HIP. See related [blog](http://gpuopen.com/hip-to-be-squared-an-introductory-hip-tutorial) that explains the example. Now it is even simpler and requires no manual modification to the hipified source code - just hipify and compile: diff --git a/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/Readme.md b/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/Readme.md index 9e1a342a07..432f9180dc 100644 --- a/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/Readme.md @@ -96,6 +96,6 @@ Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia - [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) - [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_porting_guide.md) - [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/README.md) +- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/README.md) - [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/blob/master/CONTRIBUTING.md) - [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/blob/master/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Readme.md b/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Readme.md index f65bbdcf20..e86085b648 100644 --- a/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Readme.md @@ -55,6 +55,6 @@ Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia - [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) - [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_porting_guide.md) - [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/README.md) +- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/README.md) - [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/blob/master/CONTRIBUTING.md) - [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/blob/master/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md b/projects/hip-tests/samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md index 1430e58ecc..937da30af0 100644 --- a/projects/hip-tests/samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md @@ -48,6 +48,6 @@ make - [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) - [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_porting_guide.md) - [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/README.md) +- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/README.md) - [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/blob/master/CONTRIBUTING.md) - [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/blob/master/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/1_hipEvent/Readme.md b/projects/hip-tests/samples/2_Cookbook/1_hipEvent/Readme.md index c12c76e701..2bd389e25e 100644 --- a/projects/hip-tests/samples/2_Cookbook/1_hipEvent/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/1_hipEvent/Readme.md @@ -75,6 +75,6 @@ Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia - [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) - [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_porting_guide.md) - [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/README.md) +- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/README.md) - [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/blob/master/CONTRIBUTING.md) - [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/blob/master/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/3_shared_memory/Readme.md b/projects/hip-tests/samples/2_Cookbook/3_shared_memory/Readme.md index ad23d58f73..756cb6e7f2 100644 --- a/projects/hip-tests/samples/2_Cookbook/3_shared_memory/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/3_shared_memory/Readme.md @@ -37,6 +37,6 @@ Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia - [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) - [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_porting_guide.md) - [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/README.md) +- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/README.md) - [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/blob/master/CONTRIBUTING.md) - [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/blob/master/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/4_shfl/Readme.md b/projects/hip-tests/samples/2_Cookbook/4_shfl/Readme.md index 6adc98fb4e..ac5dff9292 100644 --- a/projects/hip-tests/samples/2_Cookbook/4_shfl/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/4_shfl/Readme.md @@ -48,6 +48,6 @@ please make sure you have a 3.0 or higher compute capable device in order to use - [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) - [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_porting_guide.md) - [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/README.md) +- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/README.md) - [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/blob/master/CONTRIBUTING.md) - [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/blob/master/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Readme.md b/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Readme.md index cc9484377b..fa10c71d6c 100644 --- a/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Readme.md @@ -50,6 +50,6 @@ please make sure you have a 3.0 or higher compute capable device in order to use - [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) - [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_porting_guide.md) - [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/README.md) +- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/README.md) - [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/blob/master/CONTRIBUTING.md) - [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/blob/master/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/Readme.md b/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/Readme.md index 047cc94278..68782807bf 100644 --- a/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/Readme.md @@ -44,6 +44,6 @@ Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia - [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) - [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_porting_guide.md) - [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/README.md) +- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/README.md) - [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/blob/master/CONTRIBUTING.md) - [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/blob/master/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/7_streams/Readme.md b/projects/hip-tests/samples/2_Cookbook/7_streams/Readme.md index 1c9186791c..14b6a9762a 100644 --- a/projects/hip-tests/samples/2_Cookbook/7_streams/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/7_streams/Readme.md @@ -58,6 +58,6 @@ Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia - [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) - [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_porting_guide.md) - [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/README.md) +- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/README.md) - [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/blob/master/CONTRIBUTING.md) - [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/blob/master/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/9_unroll/Readme.md b/projects/hip-tests/samples/2_Cookbook/9_unroll/Readme.md index c6b8a8cf35..6fad55e3c9 100644 --- a/projects/hip-tests/samples/2_Cookbook/9_unroll/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/9_unroll/Readme.md @@ -43,6 +43,6 @@ please make sure you have a 3.0 or higher compute capable device in order to use - [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) - [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_porting_guide.md) - [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/README.md) +- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIPIFY/blob/master/README.md) - [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/blob/master/CONTRIBUTING.md) - [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/blob/master/RELEASE.md) From 4247ec7d9d32dd1d09c89079e47d68cfedf77f7f Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 7 Apr 2020 06:57:42 -0400 Subject: [PATCH 05/11] Merge branch 'amd-master' into amd-master-next Change-Id: I3094c15008093f2072bcd38aca4ea90aeae2d97b [ROCm/hip-tests commit: 22e9bcf4247fec7f7bbc4c7e0f5602e76605612c] --- .../hip-tests/samples/0_Intro/module_api/defaultDriver.cpp | 4 ++-- .../hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp | 4 ++-- projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp | 4 ++-- .../hip-tests/samples/0_Intro/module_api_global/runKernel.cpp | 4 ++-- projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp | 3 ++- 5 files changed, 10 insertions(+), 9 deletions(-) diff --git a/projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp b/projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp index ea36aabcf4..af8b413ac2 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp +++ b/projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp @@ -80,8 +80,8 @@ int main() { hipFree(Ad); hipFree(Bd); - delete A; - delete B; + delete[] A; + delete[] B; hipCtxDestroy(context); return 0; } diff --git a/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp b/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp index 38cf0d414c..90e569c5bc 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp +++ b/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp @@ -107,8 +107,8 @@ int main() { hipFree(Ad); hipFree(Bd); - delete A; - delete B; + delete[] A; + delete[] B; hipCtxDestroy(context); return 0; } diff --git a/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp b/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp index a011b42666..1093b0dd54 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp +++ b/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp @@ -99,8 +99,8 @@ int main() { hipFree(Ad); hipFree(Bd); - delete A; - delete B; + delete[] A; + delete[] B; hipCtxDestroy(context); return 0; } diff --git a/projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp b/projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp index 3a2804b7a2..4a2d49144c 100644 --- a/projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp +++ b/projects/hip-tests/samples/0_Intro/module_api_global/runKernel.cpp @@ -154,8 +154,8 @@ int main() { hipFree(Ad); hipFree(Bd); - delete A; - delete B; + delete[] A; + delete[] B; hipCtxDestroy(context); return 0; } diff --git a/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp b/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp index e17f19675a..14faa7671b 100644 --- a/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp @@ -56,6 +56,7 @@ void printCompilerInfo() { #endif } +double bytesToKB(size_t s) { return (double)s / (1024.0); } double bytesToGB(size_t s) { return (double)s / (1024.0 * 1024.0 * 1024.0); } #define printLimit(w1, limit, units) \ @@ -97,7 +98,7 @@ void printDeviceProp(int deviceId) { cout << setw(w1) << "totalGlobalMem: " << fixed << setprecision(2) << bytesToGB(props.totalGlobalMem) << " GB" << endl; cout << setw(w1) << "maxSharedMemoryPerMultiProcessor: " << fixed << setprecision(2) - << bytesToGB(props.maxSharedMemoryPerMultiProcessor) << " GB" << endl; + << bytesToKB(props.maxSharedMemoryPerMultiProcessor) << " KB" << endl; cout << setw(w1) << "totalConstMem: " << props.totalConstMem << endl; cout << setw(w1) << "sharedMemPerBlock: " << (float)props.sharedMemPerBlock / 1024.0 << " KB" << endl; From 0ab2c685bd99f38a33c7941b274618e6f2062e1b Mon Sep 17 00:00:00 2001 From: Tao Sang Date: Mon, 6 Apr 2020 09:58:35 -0400 Subject: [PATCH 06/11] Solve issues with hip-vdi runtime static lib 1.Combine libamdhip64_static_base.a and libamdvdi_static.a into libamdhip64_static.a. 2.Let hipcc use -use-staticlib to link libamdhip64_static.a. 3.Add some samples for static lib. 4.Fix compiling failure of code object. Change-Id: Ia2333622a8d05639b90974c4c5d3d85654ba0138 [ROCm/hip-tests commit: ed3b0eb391f92685756367a5f102f7abda2397d7] --- projects/hip-tests/samples/0_Intro/bit_extract/Makefile | 7 ++++++- projects/hip-tests/samples/0_Intro/square/Makefile | 7 +++++-- 2 files changed, 11 insertions(+), 3 deletions(-) diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/Makefile b/projects/hip-tests/samples/0_Intro/bit_extract/Makefile index 08bca6e642..4a3a0bb4fe 100644 --- a/projects/hip-tests/samples/0_Intro/bit_extract/Makefile +++ b/projects/hip-tests/samples/0_Intro/bit_extract/Makefile @@ -13,10 +13,15 @@ ifeq (${HIP_PLATFORM}, nvcc) endif EXE=bit_extract +EXE_STATIC=bit_extract_static $(EXE): bit_extract.cpp $(HIPCC) $(HIPCC_FLAGS) $< -o $@ +$(EXE_STATIC): bit_extract.cpp + $(HIPCC) -use-staticlib $(HIPCC_FLAGS) $< -o $@ + +all: $(EXE) $(EXE_STATIC) clean: - rm -f *.o $(EXE) + rm -f *.o $(EXE) $(EXE_STATIC) diff --git a/projects/hip-tests/samples/0_Intro/square/Makefile b/projects/hip-tests/samples/0_Intro/square/Makefile index 9bb0dd8205..aa046eeaaa 100644 --- a/projects/hip-tests/samples/0_Intro/square/Makefile +++ b/projects/hip-tests/samples/0_Intro/square/Makefile @@ -11,7 +11,7 @@ else SOURCES=square.cpp endif -all: square.out +all: square.out square.out.static # Step square.cpp: square.cu @@ -20,5 +20,8 @@ square.cpp: square.cu square.out: $(SOURCES) $(HIPCC) $(CXXFLAGS) $(SOURCES) -o $@ +square.out.static: $(SOURCES) + $(HIPCC) -use-staticlib $(CXXFLAGS) $(SOURCES) -o $@ + clean: - rm -f *.o *.out square.cpp + rm -f *.o *.out *.out.static square.cpp From 161c8414026470cc239a5f0dd4f4a23010855567 Mon Sep 17 00:00:00 2001 From: Tao Sang Date: Fri, 17 Apr 2020 10:13:08 -0500 Subject: [PATCH 07/11] Revert "Solve issues with hip-vdi runtime static lib" This reverts commit 0ab2c685bd99f38a33c7941b274618e6f2062e1b. Reason for revert: It is causing dkms-no-npi-hipclang broken. It is top priority to maintain dkms-no-npi-hipclang build, otherwise we lose track of regression analysis. So revert the change for now and recommit it after fixing it. Change-Id: Ia5136e888baecb6148c6c18eedbf37066fcb1eaa [ROCm/hip-tests commit: af5a4ca38b340faa819dee11e94b31c8f32440ec] --- projects/hip-tests/samples/0_Intro/bit_extract/Makefile | 7 +------ projects/hip-tests/samples/0_Intro/square/Makefile | 7 ++----- 2 files changed, 3 insertions(+), 11 deletions(-) diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/Makefile b/projects/hip-tests/samples/0_Intro/bit_extract/Makefile index 4a3a0bb4fe..08bca6e642 100644 --- a/projects/hip-tests/samples/0_Intro/bit_extract/Makefile +++ b/projects/hip-tests/samples/0_Intro/bit_extract/Makefile @@ -13,15 +13,10 @@ ifeq (${HIP_PLATFORM}, nvcc) endif EXE=bit_extract -EXE_STATIC=bit_extract_static $(EXE): bit_extract.cpp $(HIPCC) $(HIPCC_FLAGS) $< -o $@ -$(EXE_STATIC): bit_extract.cpp - $(HIPCC) -use-staticlib $(HIPCC_FLAGS) $< -o $@ - -all: $(EXE) $(EXE_STATIC) clean: - rm -f *.o $(EXE) $(EXE_STATIC) + rm -f *.o $(EXE) diff --git a/projects/hip-tests/samples/0_Intro/square/Makefile b/projects/hip-tests/samples/0_Intro/square/Makefile index aa046eeaaa..9bb0dd8205 100644 --- a/projects/hip-tests/samples/0_Intro/square/Makefile +++ b/projects/hip-tests/samples/0_Intro/square/Makefile @@ -11,7 +11,7 @@ else SOURCES=square.cpp endif -all: square.out square.out.static +all: square.out # Step square.cpp: square.cu @@ -20,8 +20,5 @@ square.cpp: square.cu square.out: $(SOURCES) $(HIPCC) $(CXXFLAGS) $(SOURCES) -o $@ -square.out.static: $(SOURCES) - $(HIPCC) -use-staticlib $(CXXFLAGS) $(SOURCES) -o $@ - clean: - rm -f *.o *.out *.out.static square.cpp + rm -f *.o *.out square.cpp From 28a2dbed5fdb0c32956676b32008b59f5bed1113 Mon Sep 17 00:00:00 2001 From: Tao Sang Date: Mon, 6 Apr 2020 09:58:35 -0400 Subject: [PATCH 08/11] Solve issues with hip-vdi runtime static lib 1.Combine libamdhip64_static_base.a and libamdvdi_static.a into libamdhip64_static.a. 2.Let hipcc use -use-staticlib to link libamdhip64_static.a. 3.Add some samples for static lib. 4.Fix compiling failure of code object. Change-Id: Ic8c95228eb139058da8b5d66ba8439486154ca6f [ROCm/hip-tests commit: 6c0a08ba4f93634784892ed4d9f6fd3dea770c38] --- projects/hip-tests/samples/0_Intro/bit_extract/Makefile | 7 ++++++- projects/hip-tests/samples/0_Intro/square/Makefile | 7 +++++-- 2 files changed, 11 insertions(+), 3 deletions(-) diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/Makefile b/projects/hip-tests/samples/0_Intro/bit_extract/Makefile index 08bca6e642..4a3a0bb4fe 100644 --- a/projects/hip-tests/samples/0_Intro/bit_extract/Makefile +++ b/projects/hip-tests/samples/0_Intro/bit_extract/Makefile @@ -13,10 +13,15 @@ ifeq (${HIP_PLATFORM}, nvcc) endif EXE=bit_extract +EXE_STATIC=bit_extract_static $(EXE): bit_extract.cpp $(HIPCC) $(HIPCC_FLAGS) $< -o $@ +$(EXE_STATIC): bit_extract.cpp + $(HIPCC) -use-staticlib $(HIPCC_FLAGS) $< -o $@ + +all: $(EXE) $(EXE_STATIC) clean: - rm -f *.o $(EXE) + rm -f *.o $(EXE) $(EXE_STATIC) diff --git a/projects/hip-tests/samples/0_Intro/square/Makefile b/projects/hip-tests/samples/0_Intro/square/Makefile index 9bb0dd8205..aa046eeaaa 100644 --- a/projects/hip-tests/samples/0_Intro/square/Makefile +++ b/projects/hip-tests/samples/0_Intro/square/Makefile @@ -11,7 +11,7 @@ else SOURCES=square.cpp endif -all: square.out +all: square.out square.out.static # Step square.cpp: square.cu @@ -20,5 +20,8 @@ square.cpp: square.cu square.out: $(SOURCES) $(HIPCC) $(CXXFLAGS) $(SOURCES) -o $@ +square.out.static: $(SOURCES) + $(HIPCC) -use-staticlib $(CXXFLAGS) $(SOURCES) -o $@ + clean: - rm -f *.o *.out square.cpp + rm -f *.o *.out *.out.static square.cpp From 6174d0f0dda57e962a7d453b5086940cf506d3eb Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Mon, 6 Apr 2020 10:57:03 -0400 Subject: [PATCH 09/11] [vdi] Refactor texture/surface reference support. Change-Id: I8014d82aae7139ef5f95e4b50c4fc6da200dbc9d [ROCm/hip-tests commit: 8ec28891de57cf1501d4b0ebd0ebb2106bedd257] --- .../samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/projects/hip-tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp b/projects/hip-tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp index 6fd49fdb0f..5831da0e9d 100644 --- a/projects/hip-tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp +++ b/projects/hip-tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp @@ -21,11 +21,7 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" -#if __HIP__ -__hip_pinned_shadow__ -#else -extern -#endif + texture tex; extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) { From 900fe14ea2d5642c1b8191b3e48730b03cdf595d Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Thu, 23 Apr 2020 21:42:06 +0530 Subject: [PATCH 10/11] Merge in the rocclr based hip runtime (#2032) * Merge master-next changes in master (include vdi development in master branch) [ROCm/hip-tests commit: e46a4e3e48c1f36013532b3ec5f8b47b3a3556a5] --- .../hipDispatchLatency/hipDispatchLatency.cpp | 2 +- .../2_Cookbook/13_occupancy/occupancy.cpp | 8 + .../samples/2_Cookbook/2_Profiler/Makefile | 53 +++++ .../2_Cookbook/2_Profiler/MatrixTranspose.cpp | 219 ++++++++++++++++++ .../samples/2_Cookbook/2_Profiler/Readme.md | 47 ++++ 5 files changed, 328 insertions(+), 1 deletion(-) create mode 100644 projects/hip-tests/samples/2_Cookbook/2_Profiler/Makefile create mode 100644 projects/hip-tests/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp create mode 100644 projects/hip-tests/samples/2_Cookbook/2_Profiler/Readme.md diff --git a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp index 625d8cd742..b528b0c75d 100644 --- a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp +++ b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp @@ -109,7 +109,7 @@ int main() { /***********************************************************************************/ //Timing directly the dispatch -#ifdef __HIP_PLATFORM_HCC__ +#if defined(__HIP_PLATFORM_HCC__) && GENERIC_GRID_LAUNCH == 1 && defined(__HCC__) for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) { hipExtLaunchKernelGGL((EmptyKernel), dim3(NUM_GROUPS), dim3(GROUP_SIZE), 0, stream0, start, stop, 0); hipEventSynchronize(stop); diff --git a/projects/hip-tests/samples/2_Cookbook/13_occupancy/occupancy.cpp b/projects/hip-tests/samples/2_Cookbook/13_occupancy/occupancy.cpp index 01fa7aafed..e772e82b1d 100644 --- a/projects/hip-tests/samples/2_Cookbook/13_occupancy/occupancy.cpp +++ b/projects/hip-tests/samples/2_Cookbook/13_occupancy/occupancy.cpp @@ -44,6 +44,8 @@ void multiplyCPU(float* C, float* A, float* B, int N){ } } +#if defined(__HIP_PLATFORM_HCC__) && GENERIC_GRID_LAUNCH == 1 && defined(__HCC__) + void launchKernel(float* C, float* A, float* B, bool manual){ hipDeviceProp_t devProp; @@ -93,8 +95,10 @@ void launchKernel(float* C, float* A, float* B, bool manual){ std::cout << "Theoretical Occupancy is " << (double)numBlock* blockSize/devProp.maxThreadsPerMultiProcessor * 100 << "%" << std::endl; } } +#endif int main() { +#if defined(__HIP_PLATFORM_HCC__) && GENERIC_GRID_LAUNCH == 1 && defined(__HCC__) float *A, *B, *C0, *C1, *cpuC; float *Ad, *Bd, *C0d, *C1d; int errors=0; @@ -173,4 +177,8 @@ int main() { free(C0); free(C1); free(cpuC); +#else + std::cout <<"hipOccupancyMaxPotentialBlockSize template not support for Clang compiler"< + +// hip header file +#include "hip/hip_runtime.h" +#include "hip/hip_profile.h" + +#define WIDTH 1024 + +#define NUM (WIDTH * WIDTH) + +#define THREADS_PER_BLOCK_X 4 +#define THREADS_PER_BLOCK_Y 4 +#define THREADS_PER_BLOCK_Z 1 + +#define ITERATIONS 10 + +// Cmdline parms to control start and stop triggers +int startTriggerIteration = -1; +int stopTriggerIteration = -1; + +// Device (Kernel) function, it must be void +__global__ void matrixTranspose(float* out, float* in, const int width) { + int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + out[y * width + x] = in[x * width + y]; +} + +// CPU implementation of matrix transpose +void matrixTransposeCPUReference(float* output, float* input, const unsigned int width) { + for (unsigned int j = 0; j < width; j++) { + for (unsigned int i = 0; i < width; i++) { + output[i * width + j] = input[j * width + i]; + } + } +} + + +// Use a separate function to demonstrate how to use function name as part of scoped marker: +void runGPU(float* Matrix, float* TransposeMatrix, float* gpuMatrix, float* gpuTransposeMatrix) { + // __func__ is a standard C++ macro which expands to the name of the function, in this case + // "runGPU" + HIP_SCOPED_MARKER(__func__, "MyGroup"); + + for (int i = 0; i < ITERATIONS; i++) { + if (i == startTriggerIteration) { + hipProfilerStart(); + } + if (i == stopTriggerIteration) { + hipProfilerStop(); + } + + float eventMs = 0.0f; + + hipEvent_t start, stop; + hipEventCreate(&start); + hipEventCreate(&stop); + + + // Record the start event + hipEventRecord(start, NULL); + + // Memory transfer from host to device + hipMemcpy(gpuMatrix, Matrix, NUM * sizeof(float), hipMemcpyHostToDevice); + + // Record the stop event + hipEventRecord(stop, NULL); + hipEventSynchronize(stop); + + hipEventElapsedTime(&eventMs, start, stop); + + printf("hipMemcpyHostToDevice time taken = %6.3fms\n", eventMs); + + // Record the start event + hipEventRecord(start, NULL); + + // Lauching kernel from host + hipLaunchKernelGGL(matrixTranspose, + dim3(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0, 0, gpuTransposeMatrix, + gpuMatrix, WIDTH); + + // Record the stop event + hipEventRecord(stop, NULL); + hipEventSynchronize(stop); + hipEventElapsedTime(&eventMs, start, stop); + + printf("kernel Execution time = %6.3fms\n", eventMs); + + // Record the start event + hipEventRecord(start, NULL); + + // Memory transfer from device to host + hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + + // Record the stop event + hipEventRecord(stop, NULL); + hipEventSynchronize(stop); + + hipEventElapsedTime(&eventMs, start, stop); + + printf("hipMemcpyDeviceToHost time taken = %6.3fms\n", eventMs); + } +}; + + +int main(int argc, char* argv[]) { + if (argc >= 2) { + startTriggerIteration = atoi(argv[1]); + printf("info : will start tracing at iteration:%d\n", startTriggerIteration); + } + if (argc >= 3) { + stopTriggerIteration = atoi(argv[2]); + printf("info : will stop tracing at iteration:%d\n", stopTriggerIteration); + } + + float* Matrix; + float* TransposeMatrix; + float* cpuTransposeMatrix; + + float* gpuMatrix; + float* gpuTransposeMatrix; + + hipDeviceProp_t devProp; + hipGetDeviceProperties(&devProp, 0); + + std::cout << "Device name " << devProp.name << std::endl; + + { + // Show example of how to create a "scoped marker". + // The scoped marker records the time spent inside the { scope } of the marker - the begin + // timestamp is at the beginning of the code scope, and the end is recorded when the SCOPE + // exits. This can be viewed in CodeXL timeline relative to other GPU and CPU events. This + // marker captures the time spent in setup including host allocation, initialization, and + // device memory allocation. + HIP_SCOPED_MARKER("Setup", "MyGroup"); + + + Matrix = (float*)malloc(NUM * sizeof(float)); + TransposeMatrix = (float*)malloc(NUM * sizeof(float)); + cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float)); + + // initialize the input data + for (int i = 0; i < NUM; i++) { + Matrix[i] = (float)i * 10.0f; + } + + + // allocate the memory on the device side + hipMalloc((void**)&gpuMatrix, NUM * sizeof(float)); + hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float)); + + // FYI, the scoped-marker will be destroyed here when the scope exits, and will record its + // "end" timestamp. + } + + runGPU(Matrix, TransposeMatrix, gpuMatrix, gpuTransposeMatrix); + + + // show how to use explicit begin/end markers: + // We begin the timed region with HIP_BEGIN_MARKER, passing in the markerName and group: + // The region will stop when HIP_END_MARKER is called + // This is another way to mark begin/end - as an alternative to scoped markers. + HIP_BEGIN_MARKER("Check&TearDown", "MyGroup"); + + int errors = 0; + + // CPU MatrixTranspose computation + matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH); + + // verify the results + double eps = 1.0E-6; + for (int i = 0; i < NUM; i++) { + if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps) { + errors++; + } + } + if (errors != 0) { + printf("FAILED: %d errors\n", errors); + } else { + printf("PASSED!\n"); + } + + // free the resources on device side + hipFree(gpuMatrix); + hipFree(gpuTransposeMatrix); + + // free the resources on host side + free(Matrix); + free(TransposeMatrix); + free(cpuTransposeMatrix); + + // This ends the last marker started in this thread, in this case "Check&TearDown" + HIP_END_MARKER(); + + return errors; +} diff --git a/projects/hip-tests/samples/2_Cookbook/2_Profiler/Readme.md b/projects/hip-tests/samples/2_Cookbook/2_Profiler/Readme.md new file mode 100644 index 0000000000..8b32beb348 --- /dev/null +++ b/projects/hip-tests/samples/2_Cookbook/2_Profiler/Readme.md @@ -0,0 +1,47 @@ +## Using hipEvents to measure performance ### + +This tutorial is follow-up of the previous two tutorial where we learn how to write our first hip program, in which we compute Matrix Transpose and in second one, we added feature to measure time taken for memory transfer and kernel execution. In this tutorial, we'll explain how to use the codexl/rocm-profiler for hip timeline tracing. Also, we will augment the source code with additional markers so we can see the high-level application flow alongside the information that CodeXL automatically collects. + + +## Introduction: + +CodeXL and rocm-profiler are the tool used for profiling the application, which is of prominent use in optimizing the application by means of finding the memory bottlenecks and etc. + +## Requirement: +[CodeXL Installation](http://gpuopen.com/compute-product/codexl/) + +## prerequiste knowledge: + +Programmers familiar with CUDA, OpenCL will be able to quickly learn and start coding with the HIP API. In case you are not, don't worry. You choose to start with the best one. We'll be explaining everything assuming you are completely new to gpgpu programming. + +## Simple Matrix Transpose + +We will be using the Simple Matrix Transpose source code from the previous tutorial as it is. + +## Using CodeXL markers for HIP Functions + +HIP can generate markers at function being/end which are displayed on the CodeXL timeline view. To do this, you need to install ROCm-Profiler and enable HIP to generate the markers: + +1. Install ROCm-Profiler Installing HIP from the rocm pre-built packages, installs the ROCm-Profiler as well. Alternatively, you can build ROCm-Profiler using the instructions given below. + + +2. Run with profiler enabled to generate ATP file. +(These steps are also captured in the Makefile) +The HIP_PROFILE_API enables display of the HIP APIs on the CodeXL trimeline view. +`/opt/rocm/bin/rocm-profiler -o -A -e HIP_PROFILE_API=1 ` + +##Using HIP_TRACE_API + +You can also print the HIP function strings to stderr using HIP_TRACE_API environment variable. This can also be combined with the more detailed debug information provided by the HIP_DB switch. For example: +`HIP_TRACE_API=1 HIP_DB=0x2 ./myHipApp` +Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors. + +## More Info: +- [HIP FAQ](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_faq.md) +- [HIP Kernel Language](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_kernel_language.md) +- [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) +- [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_porting_guide.md) +- [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/blob/master/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) +- [HIPIFY](https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/README.md) +- [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/blob/master/CONTRIBUTING.md) +- [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/blob/master/RELEASE.md) From b38734a50b57e478a7eefffad37b7e3680ebcec9 Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Thu, 23 Apr 2020 23:05:28 -0400 Subject: [PATCH 11/11] [vdi] Fix texture reference sample. - The driver code should not re-define `tex` again as it's already defined in the kernel code. Eventually, the driver code should be as regular C++ code instad of HIP code. Change-Id: I8c7cab204b98990619d6e7109b990d7089ea9261 [ROCm/hip-tests commit: 2ab06b307d53a8e05974981e15f6414d2173ef4f] --- .../samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) mode change 100755 => 100644 projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp diff --git a/projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp old mode 100755 new mode 100644 index b42ac86ad1..01729222ee --- a/projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp +++ b/projects/hip-tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp @@ -27,7 +27,6 @@ THE SOFTWARE. #define fileName "tex2dKernel.code" -texture tex; bool testResult = true; #define HIP_CHECK(cmd) \ @@ -122,7 +121,7 @@ bool runTest(int argc, char** argv) { } } } - HIP_CHECK(hipUnbindTexture(tex)); + HIP_CHECK(hipUnbindTexture(texref)); HIP_CHECK(hipFree(dData)); HIP_CHECK(hipFreeArray(array)); return testResult;