From 449e2c931bb7282b516f358391ae0290859eb566 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. --- samples/1_Utils/hipInfo/hipInfo.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/samples/1_Utils/hipInfo/hipInfo.cpp b/samples/1_Utils/hipInfo/hipInfo.cpp index e17f19675a..14faa7671b 100644 --- a/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/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 5144ee8c0402ec168bec3c234945aa53ef7b880c 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) --- samples/0_Intro/module_api/defaultDriver.cpp | 4 ++-- samples/0_Intro/module_api/launchKernelHcc.cpp | 4 ++-- samples/0_Intro/module_api/runKernel.cpp | 4 ++-- samples/0_Intro/module_api_global/runKernel.cpp | 4 ++-- 4 files changed, 8 insertions(+), 8 deletions(-) diff --git a/samples/0_Intro/module_api/defaultDriver.cpp b/samples/0_Intro/module_api/defaultDriver.cpp index ea36aabcf4..af8b413ac2 100644 --- a/samples/0_Intro/module_api/defaultDriver.cpp +++ b/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/samples/0_Intro/module_api/launchKernelHcc.cpp b/samples/0_Intro/module_api/launchKernelHcc.cpp index 38cf0d414c..90e569c5bc 100644 --- a/samples/0_Intro/module_api/launchKernelHcc.cpp +++ b/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/samples/0_Intro/module_api/runKernel.cpp b/samples/0_Intro/module_api/runKernel.cpp index a011b42666..1093b0dd54 100644 --- a/samples/0_Intro/module_api/runKernel.cpp +++ b/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/samples/0_Intro/module_api_global/runKernel.cpp b/samples/0_Intro/module_api_global/runKernel.cpp index 3a2804b7a2..4a2d49144c 100644 --- a/samples/0_Intro/module_api_global/runKernel.cpp +++ b/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 177457e54c9c874467ca5676c10d62e9ed6f2846 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) --- samples/2_Cookbook/13_occupancy/occupancy.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/samples/2_Cookbook/13_occupancy/occupancy.cpp b/samples/2_Cookbook/13_occupancy/occupancy.cpp index a9f4e198b0..01fa7aafed 100644 --- a/samples/2_Cookbook/13_occupancy/occupancy.cpp +++ b/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 a93ae7cbedca6821f8f5bdb7e4a98cded613a9a7 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) --- samples/0_Intro/square/README.md | 2 +- samples/2_Cookbook/0_MatrixTranspose/Readme.md | 2 +- samples/2_Cookbook/10_inline_asm/Readme.md | 2 +- samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md | 2 +- samples/2_Cookbook/1_hipEvent/Readme.md | 2 +- samples/2_Cookbook/3_shared_memory/Readme.md | 2 +- samples/2_Cookbook/4_shfl/Readme.md | 2 +- samples/2_Cookbook/5_2dshfl/Readme.md | 2 +- samples/2_Cookbook/6_dynamic_shared/Readme.md | 2 +- samples/2_Cookbook/7_streams/Readme.md | 2 +- samples/2_Cookbook/9_unroll/Readme.md | 2 +- 11 files changed, 11 insertions(+), 11 deletions(-) diff --git a/samples/0_Intro/square/README.md b/samples/0_Intro/square/README.md index 7a9e04fc5f..c185903993 100644 --- a/samples/0_Intro/square/README.md +++ b/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/samples/2_Cookbook/0_MatrixTranspose/Readme.md b/samples/2_Cookbook/0_MatrixTranspose/Readme.md index 9e1a342a07..432f9180dc 100644 --- a/samples/2_Cookbook/0_MatrixTranspose/Readme.md +++ b/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/samples/2_Cookbook/10_inline_asm/Readme.md b/samples/2_Cookbook/10_inline_asm/Readme.md index f65bbdcf20..e86085b648 100644 --- a/samples/2_Cookbook/10_inline_asm/Readme.md +++ b/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/samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md b/samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md index 1430e58ecc..937da30af0 100644 --- a/samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md +++ b/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/samples/2_Cookbook/1_hipEvent/Readme.md b/samples/2_Cookbook/1_hipEvent/Readme.md index c12c76e701..2bd389e25e 100644 --- a/samples/2_Cookbook/1_hipEvent/Readme.md +++ b/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/samples/2_Cookbook/3_shared_memory/Readme.md b/samples/2_Cookbook/3_shared_memory/Readme.md index ad23d58f73..756cb6e7f2 100644 --- a/samples/2_Cookbook/3_shared_memory/Readme.md +++ b/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/samples/2_Cookbook/4_shfl/Readme.md b/samples/2_Cookbook/4_shfl/Readme.md index 6adc98fb4e..ac5dff9292 100644 --- a/samples/2_Cookbook/4_shfl/Readme.md +++ b/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/samples/2_Cookbook/5_2dshfl/Readme.md b/samples/2_Cookbook/5_2dshfl/Readme.md index cc9484377b..fa10c71d6c 100644 --- a/samples/2_Cookbook/5_2dshfl/Readme.md +++ b/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/samples/2_Cookbook/6_dynamic_shared/Readme.md b/samples/2_Cookbook/6_dynamic_shared/Readme.md index 047cc94278..68782807bf 100644 --- a/samples/2_Cookbook/6_dynamic_shared/Readme.md +++ b/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/samples/2_Cookbook/7_streams/Readme.md b/samples/2_Cookbook/7_streams/Readme.md index 1c9186791c..14b6a9762a 100644 --- a/samples/2_Cookbook/7_streams/Readme.md +++ b/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/samples/2_Cookbook/9_unroll/Readme.md b/samples/2_Cookbook/9_unroll/Readme.md index c6b8a8cf35..6fad55e3c9 100644 --- a/samples/2_Cookbook/9_unroll/Readme.md +++ b/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 22e9bcf4247fec7f7bbc4c7e0f5602e76605612c 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 --- samples/0_Intro/module_api/defaultDriver.cpp | 4 ++-- samples/0_Intro/module_api/launchKernelHcc.cpp | 4 ++-- samples/0_Intro/module_api/runKernel.cpp | 4 ++-- samples/0_Intro/module_api_global/runKernel.cpp | 4 ++-- samples/1_Utils/hipInfo/hipInfo.cpp | 3 ++- 5 files changed, 10 insertions(+), 9 deletions(-) diff --git a/samples/0_Intro/module_api/defaultDriver.cpp b/samples/0_Intro/module_api/defaultDriver.cpp index ea36aabcf4..af8b413ac2 100644 --- a/samples/0_Intro/module_api/defaultDriver.cpp +++ b/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/samples/0_Intro/module_api/launchKernelHcc.cpp b/samples/0_Intro/module_api/launchKernelHcc.cpp index 38cf0d414c..90e569c5bc 100644 --- a/samples/0_Intro/module_api/launchKernelHcc.cpp +++ b/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/samples/0_Intro/module_api/runKernel.cpp b/samples/0_Intro/module_api/runKernel.cpp index a011b42666..1093b0dd54 100644 --- a/samples/0_Intro/module_api/runKernel.cpp +++ b/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/samples/0_Intro/module_api_global/runKernel.cpp b/samples/0_Intro/module_api_global/runKernel.cpp index 3a2804b7a2..4a2d49144c 100644 --- a/samples/0_Intro/module_api_global/runKernel.cpp +++ b/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/samples/1_Utils/hipInfo/hipInfo.cpp b/samples/1_Utils/hipInfo/hipInfo.cpp index e17f19675a..14faa7671b 100644 --- a/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/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 ed3b0eb391f92685756367a5f102f7abda2397d7 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 --- samples/0_Intro/bit_extract/Makefile | 7 ++++++- samples/0_Intro/square/Makefile | 7 +++++-- 2 files changed, 11 insertions(+), 3 deletions(-) diff --git a/samples/0_Intro/bit_extract/Makefile b/samples/0_Intro/bit_extract/Makefile index 08bca6e642..4a3a0bb4fe 100644 --- a/samples/0_Intro/bit_extract/Makefile +++ b/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/samples/0_Intro/square/Makefile b/samples/0_Intro/square/Makefile index 9bb0dd8205..aa046eeaaa 100644 --- a/samples/0_Intro/square/Makefile +++ b/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 af5a4ca38b340faa819dee11e94b31c8f32440ec 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 ed3b0eb391f92685756367a5f102f7abda2397d7. 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 --- samples/0_Intro/bit_extract/Makefile | 7 +------ samples/0_Intro/square/Makefile | 7 ++----- 2 files changed, 3 insertions(+), 11 deletions(-) diff --git a/samples/0_Intro/bit_extract/Makefile b/samples/0_Intro/bit_extract/Makefile index 4a3a0bb4fe..08bca6e642 100644 --- a/samples/0_Intro/bit_extract/Makefile +++ b/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/samples/0_Intro/square/Makefile b/samples/0_Intro/square/Makefile index aa046eeaaa..9bb0dd8205 100644 --- a/samples/0_Intro/square/Makefile +++ b/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 6c0a08ba4f93634784892ed4d9f6fd3dea770c38 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 --- samples/0_Intro/bit_extract/Makefile | 7 ++++++- samples/0_Intro/square/Makefile | 7 +++++-- 2 files changed, 11 insertions(+), 3 deletions(-) diff --git a/samples/0_Intro/bit_extract/Makefile b/samples/0_Intro/bit_extract/Makefile index 08bca6e642..4a3a0bb4fe 100644 --- a/samples/0_Intro/bit_extract/Makefile +++ b/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/samples/0_Intro/square/Makefile b/samples/0_Intro/square/Makefile index 9bb0dd8205..aa046eeaaa 100644 --- a/samples/0_Intro/square/Makefile +++ b/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 8ec28891de57cf1501d4b0ebd0ebb2106bedd257 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 --- samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp b/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp index 6fd49fdb0f..5831da0e9d 100644 --- a/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp +++ b/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 e46a4e3e48c1f36013532b3ec5f8b47b3a3556a5 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) --- .../hipDispatchLatency/hipDispatchLatency.cpp | 2 +- samples/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 samples/2_Cookbook/2_Profiler/Makefile create mode 100644 samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp create mode 100644 samples/2_Cookbook/2_Profiler/Readme.md diff --git a/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp b/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp index 625d8cd742..b528b0c75d 100644 --- a/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp +++ b/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/samples/2_Cookbook/13_occupancy/occupancy.cpp b/samples/2_Cookbook/13_occupancy/occupancy.cpp index 01fa7aafed..e772e82b1d 100644 --- a/samples/2_Cookbook/13_occupancy/occupancy.cpp +++ b/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/samples/2_Cookbook/2_Profiler/Readme.md b/samples/2_Cookbook/2_Profiler/Readme.md new file mode 100644 index 0000000000..8b32beb348 --- /dev/null +++ b/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 2ab06b307d53a8e05974981e15f6414d2173ef4f 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 --- samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) mode change 100755 => 100644 samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp diff --git a/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp old mode 100755 new mode 100644 index b42ac86ad1..01729222ee --- a/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp +++ b/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;