From e376b1baec93ad87c239053529628e9dd3476ebd Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 15 Mar 2016 13:39:15 -0500 Subject: [PATCH 1/4] v2 deprecating hipMallocHost with hipHostAlloc --- samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index 88d547e2c5..3d2aa659ae 100644 --- a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -103,7 +103,7 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) float *hostMem = NULL; if (p_pinned) { - hipMallocHost((void**)&hostMem, sizeof(float) * numMaxFloats); + hipHostAlloc((void**)&hostMem, sizeof(float) * numMaxFloats, hipHostAllocDefault); while (hipGetLastError() != hipSuccess) { // drop the size and try again @@ -115,7 +115,7 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) return; } numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; - hipMallocHost((void**)&hostMem, sizeof(float) * numMaxFloats); + hipHostAlloc((void**)&hostMem, sizeof(float) * numMaxFloats, hipHostAllocDefault); } } else @@ -217,9 +217,9 @@ void RunBenchmark_D2H(ResultDatabase &resultDB) float *hostMem2; if (p_pinned) { - hipMallocHost((void**)&hostMem1, sizeof(float)*numMaxFloats); + hipHostAlloc((void**)&hostMem1, sizeof(float)*numMaxFloats, hipHostAllocDefault); hipError_t err1 = hipGetLastError(); - hipMallocHost((void**)&hostMem2, sizeof(float)*numMaxFloats); + hipHostAlloc((void**)&hostMem2, sizeof(float)*numMaxFloats, hipHostAllocDefault); hipError_t err2 = hipGetLastError(); while (err1 != hipSuccess || err2 != hipSuccess) { From 93c30afe2f8cc84a70eec27cdb10eaafdb46eca4 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 15 Mar 2016 12:37:24 -0500 Subject: [PATCH 2/4] added performance metrics for kernel dispatch --- samples/1_Utils/hipDispatchLatency/Makefile | 15 +++ .../hipDispatchLatency/hipDispatchLatency.cpp | 120 ++++++++++++++++++ 2 files changed, 135 insertions(+) create mode 100644 samples/1_Utils/hipDispatchLatency/Makefile create mode 100644 samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp diff --git a/samples/1_Utils/hipDispatchLatency/Makefile b/samples/1_Utils/hipDispatchLatency/Makefile new file mode 100644 index 0000000000..b4543e6f99 --- /dev/null +++ b/samples/1_Utils/hipDispatchLatency/Makefile @@ -0,0 +1,15 @@ +HIP_PATH?=$(shell hipconfig -p) +HIPCC=$(HIP_PATH)/bin/hipcc + +EXE=hipDispatchLatency + +all: install + +$(EXE): hipDispatchLatency.cpp + $(HIPCC) hipDispatchLatency.cpp -o $@ + +install: $(EXE) + cp $(EXE) $(HIP_PATH)/bin + +clean: + rm -f *.o $(EXE) diff --git a/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp b/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp new file mode 100644 index 0000000000..fbdbced063 --- /dev/null +++ b/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp @@ -0,0 +1,120 @@ +/* +Copyright (c) 2015-2016 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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include"hip_runtime.h" +#include +#include + +#define check(msg, status) \ +if(status != hipSuccess){ \ + printf("%s failed.\n",#msg); \ + exit(1); \ +} + +#define LEN 1024*1024 +#define SIZE LEN * sizeof(float) +#define ITER 10000 + +__global__ void One(hipLaunchParm lp, float* Ad){ +} + +int main(){ + +hipError_t err; +float *A, *Ad; + +A = new float[LEN]; + +for(int i=0;i Date: Tue, 15 Mar 2016 21:05:15 -0500 Subject: [PATCH 3/4] Added single kernel launch to sample --- .../hipDispatchLatency/hipDispatchLatency.cpp | 169 ++++++++++-------- 1 file changed, 90 insertions(+), 79 deletions(-) diff --git a/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp b/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp index fbdbced063..a47931958f 100644 --- a/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp +++ b/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp @@ -36,85 +36,96 @@ __global__ void One(hipLaunchParm lp, float* Ad){ int main(){ -hipError_t err; -float *A, *Ad; + hipError_t err; + float *A, *Ad; -A = new float[LEN]; - -for(int i=0;i Date: Tue, 15 Mar 2016 14:22:00 -0500 Subject: [PATCH 4/4] corrected first and second kernel dispatch --- samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp b/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp index a47931958f..9827b94229 100644 --- a/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp +++ b/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp @@ -64,13 +64,13 @@ int main(){ hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad); hipEventRecord(stop); hipEventElapsedTime(&mS, start, stop); - std::cout<<"First Kernel Launch: \t\t"<