From 297a20eac75487a8bbc685c4b1b6ec51ae99e5d9 Mon Sep 17 00:00:00 2001 From: amd-lthakur Date: Thu, 24 Oct 2019 15:06:28 +0530 Subject: [PATCH 1/7] Adding a directed test case for hipExtModuleLaunchKernel() api. [ROCm/hip commit: 8b496e4715efed5bb3e1dbd8d657d8c239cea94e] --- .../module/hipExtModuleLaunchKernel.cpp | 143 ++++++++++++++++++ .../tests/src/runtimeApi/module/matmul.cpp | 47 ++++++ 2 files changed, 190 insertions(+) create mode 100755 projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp create mode 100755 projects/hip/tests/src/runtimeApi/module/matmul.cpp diff --git a/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp b/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp new file mode 100755 index 0000000000..456bd53282 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp @@ -0,0 +1,143 @@ +/* +Copyright (c) 2015-2019 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD_CMD: matmul.code %hc --genco %S/matmul.cpp -o matmul.code + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * TEST: %t + * HIT_END + */ + + + + +#include +#include "hip/hip_runtime.h" +#include "hip/hip_hcc.h" +#include "test_common.h" +using namespace std::chrono; + +#define fileName "matmul.code" +#define kernel_name1 "matmulK" +#define kernel_name2 "WaitKernel" + +//hipDeviceReset() call is used to clear all the allocations in case of any api failure. +#define HIP_CHECK(status) \ + if (status != hipSuccess) { \ + std::cout << "Got Status: " << hipGetErrorString(status) << " at Line: " << __LINE__ << std::endl; \ + hipDeviceReset(); \ + exit(0); \ + } + +int main() +{ + +int N=16384; +int SIZE=N*N; + +int *A=new int[N*N*sizeof(int)]; +int *B=new int[N*N*sizeof(int)]; +int *C; + +hipDeviceptr_t *Ad,*Bd; + +for(int i=0;i(stop-start); + + start=high_resolution_clock::now(); +HIP_CHECK(hipExtModuleLaunchKernel(Function2, 1,1, 1, 1,1 ,1 , 0, stream1, NULL, (void**)&config2, NULL, NULL,1 )); +HIP_CHECK(hipExtModuleLaunchKernel(Function1, N,N, 1, 32,32 ,1 , 0, stream1, NULL, (void**)&config1, NULL, NULL,1 )); +HIP_CHECK(hipStreamSynchronize(stream1)); + + stop=high_resolution_clock::now(); +auto duration2=duration_cast(stop-start); + +bool TEST_STATUS = true; + +if(! (duration2.count() < duration1.count())){ + std::cout<<"Test failed as there was no time gain observed when two kernels were launched using hipExtModuleLaunchKernel() with flag 1."< Date: Fri, 25 Oct 2019 09:19:49 +0530 Subject: [PATCH 2/7] Update hipExtModuleLaunchKernel.cpp [ROCm/hip commit: cd2514922514d3d3ec84fabfbf19a7c710146360] --- .../runtimeApi/module/hipExtModuleLaunchKernel.cpp | 14 +++++--------- 1 file changed, 5 insertions(+), 9 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp b/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp index 456bd53282..9a6719d392 100755 --- a/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp +++ b/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2019 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2019-Present Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights @@ -19,14 +19,11 @@ THE SOFTWARE. /* HIT_START * BUILD_CMD: matmul.code %hc --genco %S/matmul.cpp -o matmul.code - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11 + * BUILD: %t %s ../../test_common.cpp * TEST: %t * HIT_END */ - - - #include #include "hip/hip_runtime.h" #include "hip/hip_hcc.h" @@ -108,12 +105,12 @@ HIP_CHECK(hipStreamSynchronize(stream1)); auto stop=high_resolution_clock::now(); auto duration1=duration_cast(stop-start); - start=high_resolution_clock::now(); +start=high_resolution_clock::now(); HIP_CHECK(hipExtModuleLaunchKernel(Function2, 1,1, 1, 1,1 ,1 , 0, stream1, NULL, (void**)&config2, NULL, NULL,1 )); HIP_CHECK(hipExtModuleLaunchKernel(Function1, N,N, 1, 32,32 ,1 , 0, stream1, NULL, (void**)&config1, NULL, NULL,1 )); HIP_CHECK(hipStreamSynchronize(stream1)); - stop=high_resolution_clock::now(); +stop=high_resolution_clock::now(); auto duration2=duration_cast(stop-start); bool TEST_STATUS = true; @@ -127,8 +124,7 @@ for(int i=0;i Date: Fri, 25 Oct 2019 09:22:07 +0530 Subject: [PATCH 3/7] Update matmul.cpp [ROCm/hip commit: 318df5c36ba57d7178e5898448c3dafc14f0b30d] --- projects/hip/tests/src/runtimeApi/module/matmul.cpp | 6 +----- 1 file changed, 1 insertion(+), 5 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/module/matmul.cpp b/projects/hip/tests/src/runtimeApi/module/matmul.cpp index f6769c9d5a..bff103ebdd 100755 --- a/projects/hip/tests/src/runtimeApi/module/matmul.cpp +++ b/projects/hip/tests/src/runtimeApi/module/matmul.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2019 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2019 -> Present Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights @@ -17,12 +17,8 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ - - #include"hip/hip_runtime.h" - - extern "C" __global__ void matmulK(int* A, int* B, int* C, int N) { int ROW = blockIdx.y*blockDim.y+threadIdx.y; From 158cab3bb764360db1cace633b24717082139a94 Mon Sep 17 00:00:00 2001 From: amd-lthakur Date: Fri, 25 Oct 2019 10:44:38 +0530 Subject: [PATCH 4/7] Refactored the file as suggested [ROCm/hip commit: 564418c3087a3e8e6b62094905ed6ed1cd9e24c8] --- .../module/hipExtModuleLaunchKernel.cpp | 84 +++++++++---------- 1 file changed, 38 insertions(+), 46 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp b/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp index 9a6719d392..c040d7320c 100755 --- a/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp +++ b/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2019-Present Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2019 -> Present Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights @@ -34,14 +34,6 @@ using namespace std::chrono; #define kernel_name1 "matmulK" #define kernel_name2 "WaitKernel" -//hipDeviceReset() call is used to clear all the allocations in case of any api failure. -#define HIP_CHECK(status) \ - if (status != hipSuccess) { \ - std::cout << "Got Status: " << hipGetErrorString(status) << " at Line: " << __LINE__ << std::endl; \ - hipDeviceReset(); \ - exit(0); \ - } - int main() { @@ -55,32 +47,32 @@ int *C; hipDeviceptr_t *Ad,*Bd; for(int i=0;i(stop-start); start=high_resolution_clock::now(); -HIP_CHECK(hipExtModuleLaunchKernel(Function2, 1,1, 1, 1,1 ,1 , 0, stream1, NULL, (void**)&config2, NULL, NULL,1 )); -HIP_CHECK(hipExtModuleLaunchKernel(Function1, N,N, 1, 32,32 ,1 , 0, stream1, NULL, (void**)&config1, NULL, NULL,1 )); -HIP_CHECK(hipStreamSynchronize(stream1)); +HIPCHECK(hipExtModuleLaunchKernel(Function2, 1,1, 1, 1,1 ,1 , 0, stream1, NULL, (void**)&config2, NULL, NULL,1 )); +HIPCHECK(hipExtModuleLaunchKernel(Function1, N,N, 1, 32,32 ,1 , 0, stream1, NULL, (void**)&config1, NULL, NULL,1 )); +HIPCHECK(hipStreamSynchronize(stream1)); stop=high_resolution_clock::now(); auto duration2=duration_cast(stop-start); bool TEST_STATUS = true; -if(! (duration2.count() < duration1.count())){ +if(! (duration2.count() < duration1.count())) { std::cout<<"Test failed as there was no time gain observed when two kernels were launched using hipExtModuleLaunchKernel() with flag 1."< Date: Fri, 25 Oct 2019 15:52:11 +0530 Subject: [PATCH 5/7] Excluded the test case for nvcc platform [ROCm/hip commit: 4239c94fe5c86011d2897d4edf38827bf07f576c] --- .../runtimeApi/module/hipExtModuleLaunchKernel.cpp | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp b/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp index c040d7320c..54ddc63dde 100755 --- a/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp +++ b/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp @@ -18,8 +18,8 @@ THE SOFTWARE. */ /* HIT_START - * BUILD_CMD: matmul.code %hc --genco %S/matmul.cpp -o matmul.code - * BUILD: %t %s ../../test_common.cpp + * BUILD_CMD: matmul.code %hc --genco %S/matmul.cpp -o matmul.code EXCLUDE_HIP_PLATFORM nvcc + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc * TEST: %t * HIT_END */ @@ -108,18 +108,18 @@ auto duration2=duration_cast(stop-start); bool TEST_STATUS = true; if(! (duration2.count() < duration1.count())) { - std::cout<<"Test failed as there was no time gain observed when two kernels were launched using hipExtModuleLaunchKernel() with flag 1."< Date: Wed, 30 Oct 2019 18:09:21 -0700 Subject: [PATCH 6/7] Formatting changes ,variable name and check update [ROCm/hip commit: 4ab71216b486db0bd65d25579f8a8242663173fa] --- .../module/hipExtModuleLaunchKernel.cpp | 151 +++++++++--------- 1 file changed, 76 insertions(+), 75 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp b/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp index 54ddc63dde..f4c72ca1c5 100755 --- a/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp +++ b/projects/hip/tests/src/runtimeApi/module/hipExtModuleLaunchKernel.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2019 -> Present Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights @@ -36,96 +36,97 @@ using namespace std::chrono; int main() { - -int N=16384; -int SIZE=N*N; + int N=16384; + int SIZE=N*N; -int *A=new int[N*N*sizeof(int)]; -int *B=new int[N*N*sizeof(int)]; -int *C; + int *A=new int[N*N*sizeof(int)]; + int *B=new int[N*N*sizeof(int)]; + int *C; -hipDeviceptr_t *Ad,*Bd; + hipDeviceptr_t *Ad,*Bd; -for(int i=0;i(stop-start); + auto stop=high_resolution_clock::now(); + auto duration1=duration_cast(stop-start); -start=high_resolution_clock::now(); -HIPCHECK(hipExtModuleLaunchKernel(Function2, 1,1, 1, 1,1 ,1 , 0, stream1, NULL, (void**)&config2, NULL, NULL,1 )); -HIPCHECK(hipExtModuleLaunchKernel(Function1, N,N, 1, 32,32 ,1 , 0, stream1, NULL, (void**)&config1, NULL, NULL,1 )); -HIPCHECK(hipStreamSynchronize(stream1)); + start=high_resolution_clock::now(); + HIPCHECK(hipExtModuleLaunchKernel(Function2, 1,1, 1, 1,1 ,1 , 0, stream1, NULL, (void**)&config2, NULL, NULL,1 )); + HIPCHECK(hipExtModuleLaunchKernel(Function1, N,N, 1, 32,32 ,1 , 0, stream1, NULL, (void**)&config1, NULL, NULL,1 )); + HIPCHECK(hipStreamSynchronize(stream1)); -stop=high_resolution_clock::now(); -auto duration2=duration_cast(stop-start); + stop=high_resolution_clock::now(); + auto duration2=duration_cast(stop-start); -bool TEST_STATUS = true; + bool testStatus = true; -if(! (duration2.count() < duration1.count())) { - std::cout<<"Test failed as there was no time gain observed when two kernels were launched using hipExtModuleLaunchKernel() with flag 1."< Date: Wed, 30 Oct 2019 18:12:51 -0700 Subject: [PATCH 7/7] Formatting changes [ROCm/hip commit: 55f2a38120b2cfa7f690d1cd561bd5568341e773] --- .../tests/src/runtimeApi/module/matmul.cpp | 24 ++++++++----------- 1 file changed, 10 insertions(+), 14 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/module/matmul.cpp b/projects/hip/tests/src/runtimeApi/module/matmul.cpp index bff103ebdd..4f66a6820b 100755 --- a/projects/hip/tests/src/runtimeApi/module/matmul.cpp +++ b/projects/hip/tests/src/runtimeApi/module/matmul.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2019 -> Present Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2019 - present Advanced Micro Devices, Inc. All rights reserved. Permission is hereby granted, free of charge, to any person obtaining a copy of this software and associated documentation files (the "Software"), to deal in the Software without restriction, including without limitation the rights @@ -20,24 +20,20 @@ THE SOFTWARE. #include"hip/hip_runtime.h" extern "C" __global__ void matmulK(int* A, int* B, int* C, int N) { - int ROW = blockIdx.y*blockDim.y+threadIdx.y; int COL = blockIdx.x*blockDim.x+threadIdx.x; int tmpSum = 0; - - if (ROW < N && COL < N) { - // each thread computes one element of the block sub-matrix - for (int i = 0; i < N; i++) { - tmpSum += A[ROW * N + i] * B[i * N + COL]; - } - } - C[ROW * N + COL] =tmpSum; + if (ROW < N && COL < N) { + // each thread computes one element of the block sub-matrix + for (int i = 0; i < N; i++) { + tmpSum += A[ROW * N + i] * B[i * N + COL]; + } + } + C[ROW * N + COL] =tmpSum; } extern "C" __global__ void WaitKernel() { - unsigned long long int wait_t=32000000000,start=clock64(),cur; - do{cur=clock64()-start;} - while(cur