From 70716bee425f1268b67ef8ac84f1d1a3188ed6de Mon Sep 17 00:00:00 2001 From: Sandeep Kumar Date: Wed, 14 Dec 2016 15:49:40 +0530 Subject: [PATCH 01/22] Fixes in Makefile of couple of samples - modified Makefile for hipblas_saxpy to replaced hcblas.so with hipblas.so as part of HCSWAP-100 - Resolved missing separator issue in peer2peer cookbook Makefile Change-Id: I678fea267eee1481f02da09379339ed78d3f95f2 --- samples/2_Cookbook/8_peer2peer/Makefile | 13 +++++++------ samples/7_Advanced/hipblas_saxpy/Makefile | 2 +- 2 files changed, 8 insertions(+), 7 deletions(-) diff --git a/samples/2_Cookbook/8_peer2peer/Makefile b/samples/2_Cookbook/8_peer2peer/Makefile index 5cb7473921..0bf9e6f93e 100644 --- a/samples/2_Cookbook/8_peer2peer/Makefile +++ b/samples/2_Cookbook/8_peer2peer/Makefile @@ -1,6 +1,6 @@ HIP_PATH?= $(wildcard /opt/rocm/hip) ifeq (,$(HIP_PATH)) - HIP_PATH=../../.. + HIP_PATH=../../.. endif HIPCC=$(HIP_PATH)/bin/hipcc @@ -22,14 +22,15 @@ CXX=$(HIPCC) $(EXECUTABLE): $(OBJECTS) - $(HIPCC) $(OBJECTS) -o $@ + $(HIPCC) $(OBJECTS) -o $@ + test: $(EXECUTABLE) - $(EXECUTABLE) + $(EXECUTABLE) clean: -rm -f $(EXECUTABLE) -rm -f $(OBJECTS) -rm -f $(HIP_PATH)/src/*.o + rm -f $(EXECUTABLE) + rm -f $(OBJECTS) + rm -f $(HIP_PATH)/src/*.o diff --git a/samples/7_Advanced/hipblas_saxpy/Makefile b/samples/7_Advanced/hipblas_saxpy/Makefile index ed88be2dd0..8586e75d25 100644 --- a/samples/7_Advanced/hipblas_saxpy/Makefile +++ b/samples/7_Advanced/hipblas_saxpy/Makefile @@ -12,7 +12,7 @@ endif ifeq (${HIP_PLATFORM}, hcc) HCBLAS_ROOT?= $(wildcard /opt/rocm/hcblas) HIPCC_FLAGS += -stdlib=libc++ -I$(HCBLAS_ROOT)/include - LIBS = -L$(HCBLAS_ROOT)/lib -lhcblas + LIBS = -L$(HCBLAS_ROOT)/lib -lhipblas -rpath $(HIP_PATH)/lib endif From 054fc61f6e3beb7cd87c21fadf43ff4395bad410 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 15 Dec 2016 14:41:27 -0600 Subject: [PATCH 02/22] remove TODO file --- samples/1_Utils/hipCommander/TODO | 50 ------------------------------- 1 file changed, 50 deletions(-) delete mode 100644 samples/1_Utils/hipCommander/TODO diff --git a/samples/1_Utils/hipCommander/TODO b/samples/1_Utils/hipCommander/TODO deleted file mode 100644 index 4c835cfced..0000000000 --- a/samples/1_Utils/hipCommander/TODO +++ /dev/null @@ -1,50 +0,0 @@ -_ Add AQL kernel. -_ Fix &*kernel command so the kernel name/type is an argument not a new command. - -_ Add command to parse only. -_ Add regression to parse all the hcm files. - -_ Partition HCC, HIP, HSA, OpenCL commands into separate files. - - -_ Show time for back-to-back copies. -_ Add variables. - %loopcnt - - ./hipCommander %loopcnt=4 - -_ Add datasize command. - - -_ Add ( ) to parsing. -_ Add argument parsing and checking. - -_ Add verbose option to print each step of setup. - - print deliniater between setup and run. Add run start message. - - - print sizes of all buffers. - - print each command before running. - - show start/stop of timer routine. - -_ -_ Clear documentation on what each oepration does. -_ Add time instrumentation for each command. -_ Add pcie atomic. - - -_ Add tests for negative cases, ie endloop w/o opening loop. - - -README tips ---- -- HIP_API_TRACE combined with -v is useful to track the exact commands generates by hipCommander. - - -Other ideas: ---- -[ ] Perf guide : stream creation very slow on HCC and should be avoided. - - -Scratch: - - From cee24a20f2edb40eea955e223729dd922944dac0 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 16 Dec 2016 08:55:11 -0600 Subject: [PATCH 03/22] Print limits on CUDA devices --- samples/1_Utils/hipInfo/hipInfo.cpp | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) diff --git a/samples/1_Utils/hipInfo/hipInfo.cpp b/samples/1_Utils/hipInfo/hipInfo.cpp index 42a879e732..cf4660eae7 100644 --- a/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/samples/1_Utils/hipInfo/hipInfo.cpp @@ -63,6 +63,14 @@ double bytesToGB(size_t s) return (double)s / (1024.0*1024.0*1024.0); } +#define printLimit(w1, limit, units) \ +{\ + size_t val;\ + cudaDeviceGetLimit(&val, limit);\ + std::cout << setw(w1) << #limit": " << val << " " << units << std::endl;\ +} + + void printDeviceProp (int deviceId) { using namespace std; @@ -144,6 +152,17 @@ void printDeviceProp (int deviceId) cout << endl; +#ifdef __HIP_PLATFORM_NVCC__ + // Limits: + cout << endl; + printLimit(w1, cudaLimitStackSize, "bytes/thread"); + printLimit(w1, cudaLimitPrintfFifoSize, "bytes/device"); + printLimit(w1, cudaLimitMallocHeapSize, "bytes/device"); + printLimit(w1, cudaLimitDevRuntimeSyncDepth, "grids"); + printLimit(w1, cudaLimitDevRuntimePendingLaunchCount, "launches"); +#endif + + cout << endl; From ba8fe1675fe087bf44c6d37d7108246286bcd333 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 19 Jan 2017 15:04:32 +0530 Subject: [PATCH 04/22] Fixed hipcommander default execution for HCSWAP-106 Change-Id: I9fbd10dfaeeb4928b2ec23ceed131b5200a658f9 --- samples/1_Utils/hipCommander/hipCommander.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/samples/1_Utils/hipCommander/hipCommander.cpp b/samples/1_Utils/hipCommander/hipCommander.cpp index 9c07d066ba..0add1ce3e3 100644 --- a/samples/1_Utils/hipCommander/hipCommander.cpp +++ b/samples/1_Utils/hipCommander/hipCommander.cpp @@ -23,7 +23,7 @@ bool g_printedTiming = false; // Cmdline parms: int p_device = 0; -const char* p_command = "H2D; NullKernel; D2H"; +const char* p_command = "setstream(1); H2D; NullKernel; D2H;"; const char* p_file = nullptr; unsigned p_verbose = 0x0; unsigned p_db = 0x0; From 0390b12175aea6190b021c3758a675a4dffe4ad4 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 19 Jan 2017 12:33:06 -0600 Subject: [PATCH 05/22] Doc update - describe debug techniques Also tweak sample to remove unneeded HIP_KERNEL_NAME. Comment update --- samples/0_Intro/square/square.hipref.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/samples/0_Intro/square/square.hipref.cpp b/samples/0_Intro/square/square.hipref.cpp index 3c863b8b76..0073c1399a 100644 --- a/samples/0_Intro/square/square.hipref.cpp +++ b/samples/0_Intro/square/square.hipref.cpp @@ -81,7 +81,7 @@ int main(int argc, char *argv[]) const unsigned threadsPerBlock = 256; printf ("info: launch 'vector_square' kernel\n"); - hipLaunchKernel(HIP_KERNEL_NAME(vector_square), dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N); + hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N); printf ("info: copy Device2Host\n"); CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); From dc2f5fcc7d9969d322bc4c7ae55556afc6234dc0 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Tue, 7 Feb 2017 15:03:46 +0530 Subject: [PATCH 06/22] Command scripts for latency measurements Change-Id: I8c28765a09fb0358447367939de524b12699a317 --- .../hipCommander/perf/scripts/latency_2_d2h_h2d.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_2_d2h_kernel.hcm | 9 +++++++++ .../perf/scripts/latency_2_d2h_sync_h2d.hcm | 9 +++++++++ .../perf/scripts/latency_2_d2h_sync_kernel.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_2_h2d_d2h.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_2_h2d_kernel.hcm | 9 +++++++++ .../perf/scripts/latency_2_h2d_sync_d2h.hcm | 9 +++++++++ .../perf/scripts/latency_2_h2d_sync_kernel.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_2_kernel_d2h.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_2_kernel_h2d.hcm | 9 +++++++++ .../perf/scripts/latency_2_kernel_sync_d2h.hcm | 9 +++++++++ .../perf/scripts/latency_2_kernel_sync_h2d.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_2_sync.hcm | 9 +++++++++ .../1_Utils/hipCommander/perf/scripts/latency_2d2h.hcm | 10 ++++++++++ .../hipCommander/perf/scripts/latency_2d2h_wosync.hcm | 10 ++++++++++ .../1_Utils/hipCommander/perf/scripts/latency_2h2d.hcm | 10 ++++++++++ .../hipCommander/perf/scripts/latency_2h2d_kernel.hcm | 9 +++++++++ .../perf/scripts/latency_2h2d_kernel_wosync.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_2h2d_wosync.hcm | 10 ++++++++++ .../hipCommander/perf/scripts/latency_2kernels.hcm | 10 ++++++++++ .../perf/scripts/latency_2kernels_wosync.hcm | 10 ++++++++++ .../hipCommander/perf/scripts/latency_3_d2h_h2d.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_3_d2h_kernel.hcm | 9 +++++++++ .../perf/scripts/latency_3_d2h_sync_h2d.hcm | 9 +++++++++ .../perf/scripts/latency_3_d2h_sync_kernel.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_3_h2d_d2h.hcm | 10 ++++++++++ .../hipCommander/perf/scripts/latency_3_h2d_kernel.hcm | 9 +++++++++ .../perf/scripts/latency_3_h2d_sync_d2h.hcm | 9 +++++++++ .../perf/scripts/latency_3_h2d_sync_kernel.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_3_kernel_d2h.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_3_kernel_h2d.hcm | 9 +++++++++ .../perf/scripts/latency_3_kernel_sync_d2h.hcm | 9 +++++++++ .../perf/scripts/latency_3_kernel_sync_h2d.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_3_sync.hcm | 9 +++++++++ .../1_Utils/hipCommander/perf/scripts/latency_3d2h.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_3d2h_wosync.hcm | 9 +++++++++ .../1_Utils/hipCommander/perf/scripts/latency_3h2d.hcm | 10 ++++++++++ .../hipCommander/perf/scripts/latency_3h2d_wosync.hcm | 10 ++++++++++ .../hipCommander/perf/scripts/latency_3kernels.hcm | 10 ++++++++++ .../perf/scripts/latency_3kernels_wosync.hcm | 10 ++++++++++ .../hipCommander/perf/scripts/latency_4kernels.hcm | 10 ++++++++++ .../hipCommander/perf/scripts/latency_5kernels.hcm | 10 ++++++++++ .../hipCommander/perf/scripts/latency_6kernels.hcm | 10 ++++++++++ .../1_Utils/hipCommander/perf/scripts/latency_d2h.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_d2h_h2d.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_d2h_kernel.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_d2h_sync_h2d.hcm | 9 +++++++++ .../perf/scripts/latency_d2h_sync_kernel.hcm | 9 +++++++++ .../1_Utils/hipCommander/perf/scripts/latency_h2d.hcm | 10 ++++++++++ .../hipCommander/perf/scripts/latency_h2d_10.hcm | 2 ++ .../hipCommander/perf/scripts/latency_h2d_d2h.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_h2d_kernel.hcm | 9 +++++++++ .../perf/scripts/latency_h2d_kernel_d2h.hcm | 9 +++++++++ .../perf/scripts/latency_h2d_kernel_d2h_wosync.hcm | 9 +++++++++ .../perf/scripts/latency_h2d_kernel_wosync.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_h2d_sync_d2h.hcm | 9 +++++++++ .../perf/scripts/latency_h2d_sync_kernel_sync.hcm | 9 +++++++++ .../perf/scripts/latency_h2d_sync_kernel_sync_d2h.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_kernel.hcm | 10 ++++++++++ .../perf/scripts/latency_kernel_barrier.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_kernel_d2h.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_kernel_h2d.hcm | 9 +++++++++ .../perf/scripts/latency_kernel_sync_d2h.hcm | 9 +++++++++ .../perf/scripts/latency_kernel_sync_h2d.hcm | 9 +++++++++ .../hipCommander/perf/scripts/latency_streamcreate.hcm | 2 ++ .../1_Utils/hipCommander/perf/scripts/latency_sync.hcm | 9 +++++++++ 66 files changed, 596 insertions(+) create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2_sync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2d2h_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2kernels.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_2kernels_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3_sync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3d2h_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3h2d_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3kernels.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_3kernels_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_4kernels.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_5kernels.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_6kernels.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_d2h_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_d2h_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_10.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_wosync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_kernel.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_kernel_barrier.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_kernel_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_kernel_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_d2h.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_h2d.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_streamcreate.hcm create mode 100644 samples/1_Utils/hipCommander/perf/scripts/latency_sync.hcm diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_h2d.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_h2d.hcm new file mode 100644 index 0000000000..640bb2be79 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_h2d.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1); +loop(10); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1); +loop(100); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1); +loop(100); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1); +loop(1000); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1); +loop(1000); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1); +loop(10000); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1); +loop(10000); D2H; H2D; streamsync;D2H; H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_kernel.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_kernel.hcm new file mode 100644 index 0000000000..c1bc0f6702 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_kernel.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1); +loop(10); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1); +loop(100); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1); +loop(100); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1); +loop(1000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1); +loop(1000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1); +loop(10000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1); +loop(10000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_h2d.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_h2d.hcm new file mode 100644 index 0000000000..0e787f9bd0 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_h2d.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1); +loop(10); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1); +loop(100); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1); +loop(100); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1); +loop(1000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1); +loop(1000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1); +loop(10000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1); +loop(10000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_kernel.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_kernel.hcm new file mode 100644 index 0000000000..8d7fddc146 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_kernel.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1); +loop(10); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1); +loop(100); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1); +loop(100); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1); +loop(1000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1); +loop(1000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1); +loop(10000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1); +loop(10000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_d2h.hcm new file mode 100644 index 0000000000..7d845d03a4 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_d2h.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1); +loop(10); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1); +loop(100); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1); +loop(100); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1); +loop(1000); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1); +loop(1000); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1); +loop(10000); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1); +loop(10000); H2D; D2H; streamsync;H2D; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_kernel.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_kernel.hcm new file mode 100644 index 0000000000..49c0d77146 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_kernel.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1); +loop(10); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync;endloop(1); +loop(100); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync;endloop(1); +loop(100); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync; endloop(1); +loop(1000); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync; endloop(1); +loop(1000); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync; endloop(1); +loop(10000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1); +loop(10000); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_d2h.hcm new file mode 100644 index 0000000000..fe1f14bee5 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_d2h.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1); +loop(10); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1); +loop(100); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1); +loop(100); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1); +loop(1000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1); +loop(1000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1); +loop(10000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1); +loop(10000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_kernel.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_kernel.hcm new file mode 100644 index 0000000000..0762001daa --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_kernel.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(10); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(100); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(100); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(1000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(1000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(10000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(10000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_d2h.hcm new file mode 100644 index 0000000000..88003ba476 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_d2h.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1); +loop(10); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1); +loop(100); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1); +loop(100); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1); +loop(1000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1); +loop(1000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1); +loop(10000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1); +loop(10000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_h2d.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_h2d.hcm new file mode 100644 index 0000000000..01913f8481 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_h2d.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1); +loop(10); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1); +loop(100); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1); +loop(100); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1); +loop(1000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1); +loop(1000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1); +loop(10000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1); +loop(10000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_d2h.hcm new file mode 100644 index 0000000000..530eb8f68e --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_d2h.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(10); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(100); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(100); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_h2d.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_h2d.hcm new file mode 100644 index 0000000000..6d83ee87c9 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_h2d.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(10); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(100); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(100); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2_sync.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2_sync.hcm new file mode 100644 index 0000000000..8b9e233a9e --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2_sync.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); streamsync; streamsync; endloop(1); +loop(10); streamsync; streamsync; endloop(1); +loop(100); streamsync; streamsync; endloop(1); +loop(100); streamsync; streamsync; endloop(1); +loop(1000); streamsync; streamsync; endloop(1); +loop(1000); streamsync; streamsync; endloop(1); +loop(10000); streamsync; streamsync; endloop(1); +loop(10000); streamsync; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2d2h.hcm new file mode 100644 index 0000000000..83cdc4ff75 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2d2h.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); D2H; streamsync; D2H; streamsync; endloop(1); +loop(10); D2H; streamsync; D2H; streamsync; endloop(1); +loop(100); D2H; streamsync; D2H; streamsync; endloop(1); +loop(100); D2H; streamsync; D2H; streamsync; endloop(1); +loop(1000); D2H;streamsync; D2H; streamsync; endloop(1); +loop(1000); D2H; streamsync; D2H; streamsync; endloop(1); +loop(1000); D2H; streamsync; D2H; streamsync; endloop(1); +loop(10000); D2H; streamsync; D2H; streamsync; endloop(1); +loop(10000); D2H; streamsync; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2d2h_wosync.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2d2h_wosync.hcm new file mode 100644 index 0000000000..4b91403582 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2d2h_wosync.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); D2H; D2H; streamsync; endloop(1); +loop(10); D2H; D2H; streamsync; endloop(1); +loop(100); D2H; D2H; streamsync; endloop(1); +loop(100); D2H; D2H; streamsync; endloop(1); +loop(1000); D2H; D2H; streamsync; endloop(1); +loop(1000); D2H; D2H; streamsync; endloop(1); +loop(1000); D2H; D2H; streamsync; endloop(1); +loop(10000); D2H; D2H; streamsync; endloop(1); +loop(10000); D2H; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d.hcm new file mode 100644 index 0000000000..a2e4311bf6 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); H2D; streamsync; H2D; streamsync; endloop(1); +loop(10); H2D; streamsync; H2D; streamsync; endloop(1); +loop(100); H2D; streamsync; H2D; streamsync; endloop(1); +loop(100); H2D; streamsync; H2D; streamsync; endloop(1); +loop(1000); H2D;streamsync; H2D; streamsync; endloop(1); +loop(1000); H2D; streamsync; H2D; streamsync; endloop(1); +loop(1000); H2D; streamsync; H2D; streamsync; endloop(1); +loop(10000); H2D; streamsync; H2D; streamsync; endloop(1); +loop(10000); H2D; streamsync; H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel.hcm new file mode 100644 index 0000000000..0c622614cc --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1); +loop(10); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(100); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1); +loop(100); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1); +loop(1000); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1); +loop(1000); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1); +loop(10000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(10000); H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel; streamsync;endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel_wosync.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel_wosync.hcm new file mode 100644 index 0000000000..d73467da10 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel_wosync.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1); +loop(10); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync;endloop(1); +loop(100); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync;endloop(1); +loop(100); H2D; NullKernel; streamsync; H2D; NullKernel; streamsync;endloop(1); +loop(1000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1); +loop(1000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1); +loop(10000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1); +loop(10000); H2D ; NullKernel; streamsync;H2D; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_wosync.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_wosync.hcm new file mode 100644 index 0000000000..35f5e68522 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_wosync.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); H2D; H2D; streamsync; endloop(1); +loop(10); H2D; H2D; streamsync; endloop(1); +loop(100); H2D; H2D; streamsync; endloop(1); +loop(100); H2D; H2D; streamsync; endloop(1); +loop(1000); H2D; H2D; streamsync; endloop(1); +loop(1000); H2D; H2D; streamsync; endloop(1); +loop(1000); H2D; H2D; streamsync; endloop(1); +loop(10000); H2D; H2D; streamsync; endloop(1); +loop(10000); H2D; H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2kernels.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2kernels.hcm new file mode 100644 index 0000000000..3b85c6bef8 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2kernels.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(10); NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(100); NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(100); NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_2kernels_wosync.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_2kernels_wosync.hcm new file mode 100644 index 0000000000..584d6b8021 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_2kernels_wosync.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); NullKernel; NullKernel; streamsync; endloop(1); +loop(10); NullKernel; NullKernel; streamsync; endloop(1); +loop(100); NullKernel; NullKernel; streamsync; endloop(1); +loop(100); NullKernel; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; NullKernel; streamsync; endloop(1); +loop(10000); NullKernel; NullKernel; streamsync; endloop(1); +loop(10000); NullKernel; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_h2d.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_h2d.hcm new file mode 100644 index 0000000000..7f0fce96c5 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_h2d.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1); +loop(10); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1); +loop(100); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1); +loop(100); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1); +loop(1000); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1); +loop(1000); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1); +loop(10000); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1); +loop(10000); D2H; H2D; streamsync;D2H; H2D; streamsync; D2H; H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_kernel.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_kernel.hcm new file mode 100644 index 0000000000..a384439b5c --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_kernel.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1); +loop(10); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1); +loop(100); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1); +loop(100); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1); +loop(1000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1); +loop(1000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1); +loop(10000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1); +loop(10000); D2H; NullKernel; streamsync;D2H; NullKernel; streamsync;streamsync; D2H; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_h2d.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_h2d.hcm new file mode 100644 index 0000000000..1cab6ff0d2 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_h2d.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1); +loop(10); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1); +loop(100); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1); +loop(100); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1); +loop(1000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1); +loop(1000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1); +loop(10000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1); +loop(10000); D2H; streamsync; H2D; streamsync;D2H; streamsync; H2D; streamsync; D2H; streamsync; H2D;streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_kernel.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_kernel.hcm new file mode 100644 index 0000000000..ff5b09a3dc --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_kernel.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1); +loop(10); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1); +loop(100); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1); +loop(100); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1); +loop(1000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1); +loop(1000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1); +loop(10000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1); +loop(10000); D2H; streamsync; NullKernel; streamsync;D2H; streamsync; NullKernel; streamsync; D2H; streamsync; NullKernel;streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_d2h.hcm new file mode 100644 index 0000000000..d8921a64e7 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_d2h.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1); +loop(10); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1); +loop(100); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1); +loop(100); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1); +loop(1000); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1); +loop(1000); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1); +loop(10000); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1); +loop(10000); H2D; D2H; streamsync;H2D; D2H; streamsync; H2D; D2H; streamsync; endloop(1); + diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_kernel.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_kernel.hcm new file mode 100644 index 0000000000..4ccbf9a83c --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_kernel.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1); +loop(10); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1); +loop(100); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1); +loop(100); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1); +loop(1000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1); +loop(1000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1); +loop(10000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1); +loop(10000); H2D; NullKernel; streamsync;H2D; NullKernel; streamsync;streamsync; H2D; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_d2h.hcm new file mode 100644 index 0000000000..a3d9a282f5 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_d2h.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1); +loop(10); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1); +loop(100); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1); +loop(100); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1); +loop(1000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1); +loop(1000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1); +loop(10000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1); +loop(10000); H2D; streamsync; D2H; streamsync;H2D; streamsync; D2H; streamsync; H2D; streamsync; D2H;streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_kernel.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_kernel.hcm new file mode 100644 index 0000000000..56554d15fd --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_kernel.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1); +loop(10); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1); +loop(100); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1); +loop(100); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1); +loop(1000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1); +loop(1000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1); +loop(10000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1); +loop(10000); H2D; streamsync; NullKernel; streamsync;H2D; streamsync; NullKernel; streamsync; H2D; streamsync; NullKernel;streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_d2h.hcm new file mode 100644 index 0000000000..a6e3a683d2 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_d2h.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1); +loop(10); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1); +loop(100); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1); +loop(100); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1); +loop(1000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1); +loop(1000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1); +loop(10000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1); +loop(10000); NullKernel; D2H; streamsync;NullKernel; D2H; streamsync; NullKernel; D2H; streamsync;endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_h2d.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_h2d.hcm new file mode 100644 index 0000000000..eae3eadc5b --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_h2d.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1); +loop(10); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1); +loop(100); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1); +loop(100); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1); +loop(1000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1); +loop(1000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1); +loop(10000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1); +loop(10000); NullKernel; H2D; streamsync;NullKernel; H2D; streamsync; NullKernel; H2D; streamsync;endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_d2h.hcm new file mode 100644 index 0000000000..9e21709b0b --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_d2h.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(10); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(100); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(100); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync;NullKernel; streamsync; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_h2d.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_h2d.hcm new file mode 100644 index 0000000000..b1ef7ef9f4 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_h2d.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(10); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(100); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(100); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync;NullKernel; streamsync; H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3_sync.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3_sync.hcm new file mode 100644 index 0000000000..bc8d21c594 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3_sync.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10);streamsync; streamsync; streamsync; endloop(1); +loop(10);streamsync; streamsync; streamsync; endloop(1); +loop(100);streamsync; streamsync; streamsync; endloop(1); +loop(100);streamsync; streamsync; streamsync; endloop(1); +loop(1000);streamsync; streamsync; streamsync; endloop(1); +loop(1000);streamsync; streamsync; streamsync; endloop(1); +loop(10000);streamsync; streamsync; streamsync; endloop(1); +loop(10000);streamsync; streamsync; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3d2h.hcm new file mode 100644 index 0000000000..4e07574b99 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3d2h.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1); +loop(10); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1); +loop(100); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1); +loop(100); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1); +loop(1000); D2H;streamsync; D2H; streamsync; D2H; streamsync; endloop(1); +loop(1000); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1); +loop(10000); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1); +loop(10000); D2H; streamsync; D2H; streamsync; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3d2h_wosync.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3d2h_wosync.hcm new file mode 100644 index 0000000000..e96707fed9 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3d2h_wosync.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); D2H; D2H; D2H; streamsync; endloop(1); +loop(10); D2H; D2H; D2H; streamsync; endloop(1); +loop(100); D2H; D2H; D2H; streamsync; endloop(1); +loop(100); D2H; D2H; D2H; streamsync; endloop(1); +loop(1000); D2H; D2H; D2H; streamsync; endloop(1); +loop(1000); D2H; D2H; D2H; streamsync; endloop(1); +loop(10000); D2H; D2H; D2H; streamsync; endloop(1); +loop(10000); D2H; D2H; D2H;streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3h2d.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3h2d.hcm new file mode 100644 index 0000000000..82151adb8b --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3h2d.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); H2D; streamsync;H2D;streamsync; H2D; streamsync; endloop(1); +loop(10); H2D; streamsync;H2D;streamsync; H2D; streamsync; endloop(1); +loop(100); H2D; streamsync;H2D; streamsync;H2D; streamsync; endloop(1); +loop(100); H2D;streamsync; H2D; streamsync;H2D; streamsync; endloop(1); +loop(1000); H2D;streamsync; H2D;streamsync; H2D; streamsync; endloop(1); +loop(1000); H2D;streamsync; H2D; streamsync;H2D; streamsync; endloop(1); +loop(1000); H2D;streamsync; H2D; streamsync;H2D; streamsync; endloop(1); +loop(10000); H2D;streamsync; H2D; streamsync;H2D; streamsync; endloop(1); +loop(10000); H2D;streamsync; H2D;streamsync; H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3h2d_wosync.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3h2d_wosync.hcm new file mode 100644 index 0000000000..7d96bfcfab --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3h2d_wosync.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); H2D; H2D; H2D; streamsync; endloop(1); +loop(10); H2D; H2D; H2D; streamsync; endloop(1); +loop(100); H2D; H2D; H2D; streamsync; endloop(1); +loop(100); H2D; H2D; H2D; streamsync; endloop(1); +loop(1000); H2D; H2D; H2D; streamsync; endloop(1); +loop(1000); H2D; H2D; H2D; streamsync; endloop(1); +loop(1000); H2D; H2D; H2D; streamsync; endloop(1); +loop(10000); H2D; H2D; H2D; streamsync; endloop(1); +loop(10000); H2D; H2D; H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3kernels.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3kernels.hcm new file mode 100644 index 0000000000..2e8306dfde --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3kernels.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(10); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(100); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(100); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; NullKernel; streamsync; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_3kernels_wosync.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_3kernels_wosync.hcm new file mode 100644 index 0000000000..85cd0dd4d2 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_3kernels_wosync.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(10); NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(100); NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(100); NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(10000); NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(10000); NullKernel; NullKernel; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_4kernels.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_4kernels.hcm new file mode 100644 index 0000000000..48a8223626 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_4kernels.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(10); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(100); NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1); +loop(100); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_5kernels.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_5kernels.hcm new file mode 100644 index 0000000000..70ad00c248 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_5kernels.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); NullKernel; NullKernel; NullKernel; NullKernel;NullKernel; streamsync; endloop(1); +loop(10); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(100); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1); +loop(100); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1); +loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_6kernels.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_6kernels.hcm new file mode 100644 index 0000000000..1bbb5694b1 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_6kernels.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); NullKernel; NullKernel; NullKernel; NullKernel;NullKernel;NullKernel; streamsync; endloop(1); +loop(10); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1); +loop(100); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;NullKernel;streamsync; endloop(1); +loop(100); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1); +loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1); +loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1); +loop(1000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1); +loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1); +loop(10000); NullKernel; NullKernel; NullKernel; NullKernel; NullKernel; NullKernel;streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_d2h.hcm new file mode 100644 index 0000000000..54f06a3481 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_d2h.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); D2H; streamsync; endloop(1); +loop(10); D2H; streamsync; endloop(1); +loop(100); D2H; streamsync; endloop(1); +loop(100); D2H; streamsync; endloop(1); +loop(1000); D2H; streamsync; endloop(1); +loop(1000); D2H; streamsync; endloop(1); +loop(10000); D2H; streamsync; endloop(1); +loop(10000); D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_h2d.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_h2d.hcm new file mode 100644 index 0000000000..6667ba95fa --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_h2d.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); D2H; H2D; streamsync; endloop(1); +loop(10); D2H; H2D; streamsync; endloop(1); +loop(100); D2H; H2D; streamsync; endloop(1); +loop(100); D2H; H2D; streamsync; endloop(1); +loop(1000); D2H; H2D; streamsync; endloop(1); +loop(1000); D2H; H2D; streamsync; endloop(1); +loop(10000); D2H; H2D; streamsync; endloop(1); +loop(10000); D2H; H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_kernel.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_kernel.hcm new file mode 100644 index 0000000000..fe770c5e9d --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_kernel.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); D2H; NullKernel; streamsync; endloop(1); +loop(10); D2H; NullKernel; streamsync; endloop(1); +loop(100); D2H; NullKernel; streamsync; endloop(1); +loop(100); D2H; NullKernel; streamsync; endloop(1); +loop(1000); D2H; NullKernel; streamsync; endloop(1); +loop(1000); D2H; NullKernel; streamsync; endloop(1); +loop(10000); D2H; NullKernel; streamsync; endloop(1); +loop(10000); D2H; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_h2d.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_h2d.hcm new file mode 100644 index 0000000000..20ec951509 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_h2d.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); D2H; streamsync; H2D; streamsync; endloop(1); +loop(10); D2H; streamsync; H2D; streamsync; endloop(1); +loop(100); D2H; streamsync; H2D; streamsync; endloop(1); +loop(100); D2H; streamsync; H2D; streamsync; endloop(1); +loop(1000); D2H; streamsync; H2D; streamsync; endloop(1); +loop(1000); D2H; streamsync; H2D; streamsync; endloop(1); +loop(10000); D2H; streamsync; H2D; streamsync; endloop(1); +loop(10000); D2H; streamsync; H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_kernel.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_kernel.hcm new file mode 100644 index 0000000000..77e483b3df --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_kernel.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); D2H; streamsync;NullKernel; streamsync; endloop(1); +loop(10); D2H; streamsync;NullKernel; streamsync; endloop(1); +loop(100); D2H; streamsync;NullKernel; streamsync; endloop(1); +loop(100); D2H; streamsync;NullKernel; streamsync; endloop(1); +loop(1000); D2H; streamsync;NullKernel; streamsync; endloop(1); +loop(1000); D2H; streamsync;NullKernel; streamsync; endloop(1); +loop(10000); D2H; streamsync;NullKernel; streamsync; endloop(1); +loop(10000); D2H; streamsync;NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_h2d.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d.hcm new file mode 100644 index 0000000000..f5642bfdf0 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); H2D; streamsync; endloop(1); +loop(10); H2D; streamsync; endloop(1); +loop(100); H2D; streamsync; endloop(1); +loop(100); H2D; streamsync; endloop(1); +loop(1000); H2D; streamsync; endloop(1); +loop(1000); H2D; streamsync; endloop(1); +loop(1000); H2D; streamsync; endloop(1); +loop(10000); H2D; streamsync; endloop(1); +loop(10000); H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_10.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_10.hcm new file mode 100644 index 0000000000..05452b9c87 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_10.hcm @@ -0,0 +1,2 @@ +setstream(1); +loop(10); H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_d2h.hcm new file mode 100644 index 0000000000..dad9fc7437 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_d2h.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); H2D; D2H; streamsync; endloop(1); +loop(10); H2D; D2H; streamsync; endloop(1); +loop(100); H2D; D2H; streamsync; endloop(1); +loop(100); H2D; D2H; streamsync; endloop(1); +loop(1000); H2D; D2H; streamsync; endloop(1); +loop(1000); H2D; D2H; streamsync; endloop(1); +loop(10000); H2D; D2H; streamsync; endloop(1); +loop(10000); H2D; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel.hcm new file mode 100644 index 0000000000..1b60640b9e --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); H2D; NullKernel; streamsync; endloop(1); +loop(10); H2D; NullKernel; streamsync; endloop(1); +loop(100); H2D; NullKernel; streamsync; endloop(1); +loop(100); H2D; NullKernel; streamsync; endloop(1); +loop(1000); H2D; NullKernel; streamsync; endloop(1); +loop(1000); H2D; NullKernel; streamsync; endloop(1); +loop(10000); H2D; NullKernel; streamsync; endloop(1); +loop(10000); H2D; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h.hcm new file mode 100644 index 0000000000..6e4e9f3544 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10);H2D; streamsync; NullKernel; D2H; streamsync;endloop(1); +loop(10); H2D; streamsync; NullKernel; D2H; streamsync; endloop(1); +loop(100); H2D; streamsync; NullKernel; D2H; streamsync; endloop(1); +loop(100); H2D; streamsync; NullKernel; D2H; streamsync; endloop(1); +loop(1000); H2D; streamsync; NullKernel; D2H; streamsync; endloop(1); +loop(1000); H2D; streamsync; NullKernel; D2H; streamsync; endloop(1); +loop(10000); H2D; streamsync; NullKernel; D2H; streamsync; endloop(1); +loop(10000); H2D; streamsync; NullKernel; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h_wosync.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h_wosync.hcm new file mode 100644 index 0000000000..4e94a26ebf --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h_wosync.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10);H2D; NullKernel; D2H; streamsync;endloop(1); +loop(10); H2D; NullKernel; D2H; streamsync; endloop(1); +loop(100); H2D; NullKernel; D2H; streamsync; endloop(1); +loop(100); H2D; NullKernel; D2H; streamsync; endloop(1); +loop(1000); H2D; NullKernel; D2H; streamsync; endloop(1); +loop(1000); H2D; NullKernel; D2H; streamsync; endloop(1); +loop(10000); H2D; NullKernel; D2H; streamsync; endloop(1); +loop(10000); H2D; NullKernel; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_wosync.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_wosync.hcm new file mode 100644 index 0000000000..b3b40d3190 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_wosync.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); H2D; NullKernel; streamsync; endloop(1); +loop(10); H2D; NullKernel; streamsync; endloop(1); +loop(100); H2D; NullKernel; streamsync; endloop(1); +loop(100); H2D; NullKernel; streamsync; endloop(1); +loop(1000); H2D; NullKernel; streamsync; endloop(1); +loop(1000); H2D; NullKernel; streamsync; endloop(1); +loop(10000); H2D; NullKernel; streamsync; endloop(1); +loop(10000); H2D ; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_d2h.hcm new file mode 100644 index 0000000000..030213d1b3 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_d2h.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); H2D; streamsync; D2H; streamsync; endloop(1); +loop(10); H2D; streamsync; D2H; streamsync; endloop(1); +loop(100); H2D; streamsync; D2H; streamsync; endloop(1); +loop(100); H2D; streamsync; D2H; streamsync; endloop(1); +loop(1000); H2D; streamsync; D2H; streamsync; endloop(1); +loop(1000); H2D; streamsync; D2H; streamsync; endloop(1); +loop(10000); H2D; streamsync; D2H; streamsync; endloop(1); +loop(10000); H2D; streamsync; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync.hcm new file mode 100644 index 0000000000..146c74bcae --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(10); H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(100); H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(100); H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(1000); H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(1000); H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(10000); H2D; streamsync; NullKernel; streamsync; endloop(1); +loop(10000); H2D; streamsync; NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync_d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync_d2h.hcm new file mode 100644 index 0000000000..366d04f469 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync_d2h.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10);H2D; streamsync; NullKernel;streamsync; D2H; streamsync;endloop(1); +loop(10); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(100); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(100); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(1000); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(1000); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(10000); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(10000); H2D; streamsync; NullKernel; streamsync; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_kernel.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_kernel.hcm new file mode 100644 index 0000000000..027d89aad0 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_kernel.hcm @@ -0,0 +1,10 @@ +setstream(1); +loop(10); NullKernel; streamsync; endloop(1); +loop(10); NullKernel; streamsync; endloop(1); +loop(100); NullKernel; streamsync; endloop(1); +loop(100); NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_barrier.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_barrier.hcm new file mode 100644 index 0000000000..fb6a867e7f --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_barrier.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); NullKernel; streamsync; streamsync; endloop(1); +loop(10); NullKernel; streamsync; streamsync; endloop(1); +loop(100); NullKernel; streamsync; streamsync; endloop(1); +loop(100); NullKernel; streamsync; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_d2h.hcm new file mode 100644 index 0000000000..2e64472dbd --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_d2h.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); NullKernel; D2H; streamsync; endloop(1); +loop(10); NullKernel; D2H; streamsync; endloop(1); +loop(100); NullKernel; D2H; streamsync; endloop(1); +loop(100); NullKernel; D2H; streamsync; endloop(1); +loop(1000); NullKernel; D2H; streamsync; endloop(1); +loop(1000); NullKernel; D2H; streamsync; endloop(1); +loop(10000); NullKernel; D2H; streamsync; endloop(1); +loop(10000); NullKernel; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_h2d.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_h2d.hcm new file mode 100644 index 0000000000..b220a69c68 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_h2d.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); NullKernel; H2D; streamsync; endloop(1); +loop(10); NullKernel; H2D; streamsync; endloop(1); +loop(100); NullKernel; H2D; streamsync; endloop(1); +loop(100); NullKernel; H2D; streamsync; endloop(1); +loop(1000); NullKernel; H2D; streamsync; endloop(1); +loop(1000); NullKernel; H2D; streamsync; endloop(1); +loop(10000); NullKernel; H2D; streamsync; endloop(1); +loop(10000); NullKernel; H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_d2h.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_d2h.hcm new file mode 100644 index 0000000000..48b332b1c3 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_d2h.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(10); NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(100); NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(100); NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; D2H; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; D2H; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_h2d.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_h2d.hcm new file mode 100644 index 0000000000..5a45d55376 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_h2d.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(10); NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(100); NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(100); NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(1000); NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; H2D; streamsync; endloop(1); +loop(10000); NullKernel; streamsync; H2D; streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_streamcreate.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_streamcreate.hcm new file mode 100644 index 0000000000..1e6aef5dc1 --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_streamcreate.hcm @@ -0,0 +1,2 @@ +setstream(1); +loop(10);setstream(1);setstream(2);setstream(3);setstream(4);setstream(5);streamsync; endloop(1); diff --git a/samples/1_Utils/hipCommander/perf/scripts/latency_sync.hcm b/samples/1_Utils/hipCommander/perf/scripts/latency_sync.hcm new file mode 100644 index 0000000000..a784013b1a --- /dev/null +++ b/samples/1_Utils/hipCommander/perf/scripts/latency_sync.hcm @@ -0,0 +1,9 @@ +setstream(1); +loop(10); streamsync; endloop(1); +loop(10); streamsync; endloop(1); +loop(100); streamsync; endloop(1); +loop(100); streamsync; endloop(1); +loop(1000); streamsync; endloop(1); +loop(1000); streamsync; endloop(1); +loop(10000); streamsync; endloop(1); +loop(10000); streamsync; endloop(1); From b5e3f593773648236a3c61ea6540b075cbf1ab67 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Fri, 17 Feb 2017 08:51:02 -0600 Subject: [PATCH 07/22] removed hipblas samples as it is not yet supported Change-Id: I354b710e652ce0d0413d670530ceb8b70f4993d5 --- samples/7_Advanced/hipblas_saxpy/Makefile | 34 ------- .../7_Advanced/hipblas_saxpy/saxpy.cublas.cpp | 94 ------------------- .../hipblas_saxpy/saxpy.hipblasref.cpp | 94 ------------------- 3 files changed, 222 deletions(-) delete mode 100644 samples/7_Advanced/hipblas_saxpy/Makefile delete mode 100644 samples/7_Advanced/hipblas_saxpy/saxpy.cublas.cpp delete mode 100644 samples/7_Advanced/hipblas_saxpy/saxpy.hipblasref.cpp diff --git a/samples/7_Advanced/hipblas_saxpy/Makefile b/samples/7_Advanced/hipblas_saxpy/Makefile deleted file mode 100644 index 8586e75d25..0000000000 --- a/samples/7_Advanced/hipblas_saxpy/Makefile +++ /dev/null @@ -1,34 +0,0 @@ -HIP_PATH?= $(wildcard /opt/rocm/hip) -ifeq (,$(HIP_PATH)) - HIP_PATH=../../.. -endif -HIPCC=$(HIP_PATH)/bin/hipcc - -HIPCC_FLAGS += -std=c++11 -HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform) -ifeq (${HIP_PLATFORM}, nvcc) - LIBS = -lcublas -endif -ifeq (${HIP_PLATFORM}, hcc) - HCBLAS_ROOT?= $(wildcard /opt/rocm/hcblas) - HIPCC_FLAGS += -stdlib=libc++ -I$(HCBLAS_ROOT)/include - LIBS = -L$(HCBLAS_ROOT)/lib -lhipblas -rpath $(HIP_PATH)/lib -endif - - -all: saxpy.hipblas.out - -saxpy.cublas.out : saxpy.cublas.cpp - nvcc -std=c++11 -I$(CUDA_HOME)/include saxpy.cublas.cpp -o $@ -L$(CUDA_HOME)/lib64 -lcublas - -# $HIPBLAS_ROOT/bin/hipifyblas ./saxpy.cublas.cpp > ./saxpy.hipblas.cpp -# Then review & finish port in saxpy.hipblas.cpp - -saxpy.hipblasref.o: saxpy.hipblasref.cpp - $(HIPCC) $(HIPCC_FLAGS) -c $< -o $@ - -saxpy.hipblas.out: saxpy.hipblasref.o - $(HIPCC) $< -o $@ $(LIBS) - -clean: - rm -f *.o *.out diff --git a/samples/7_Advanced/hipblas_saxpy/saxpy.cublas.cpp b/samples/7_Advanced/hipblas_saxpy/saxpy.cublas.cpp deleted file mode 100644 index 03a38f3fb1..0000000000 --- a/samples/7_Advanced/hipblas_saxpy/saxpy.cublas.cpp +++ /dev/null @@ -1,94 +0,0 @@ - -#include -#include -#include -#include - -// header file for the GPU API -#include -#include - -#define N (1024 * 500) - -#define CHECK(cmd) \ -{\ - cudaError_t error = cmd; \ - if (error != cudaSuccess) { \ - fprintf(stderr, "error: '%s'(%d) at %s:%d\n", cudaGetErrorString(error), error,__FILE__, __LINE__); \ - exit(EXIT_FAILURE);\ - }\ -} - -#define CHECK_BLAS(cmd) \ -{\ - cublasStatus_t error = cmd;\ - if (error != CUBLAS_STATUS_SUCCESS) { \ - fprintf(stderr, "error: (%d) at %s:%d\n", error,__FILE__, __LINE__); \ - exit(EXIT_FAILURE);\ - }\ -} - -int main() { - - const float a = 100.0f; - float x[N]; - float y[N], y_cpu_res[N], y_gpu_res[N]; - - // initialize the input data - std::default_random_engine random_gen; - std::uniform_real_distribution distribution(-N, N); - std::generate_n(x, N, [&]() { return distribution(random_gen); }); - std::generate_n(y, N, [&]() { return distribution(random_gen); }); - std::copy_n(y, N, y_cpu_res); - - // Explicit GPU code: - - size_t Nbytes = N*sizeof(float); - float *x_gpu, *y_gpu; - - cublasHandle_t handle; - - cudaDeviceProp props; - CHECK(cudaGetDeviceProperties(&props, 0/*deviceID*/)); - printf ("info: running on device %s\n", props.name); - - printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); - printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); - CHECK(cudaMalloc(&x_gpu, Nbytes)); - CHECK(cudaMalloc(&y_gpu, Nbytes)); - - // Initialize the blas library - CHECK_BLAS ( cublasCreate(&handle)); - - // copy n elements from a vector in host memory space to a vector in GPU memory space - printf ("info: copy Host2Device\n"); - CHECK_BLAS ( cublasSetVector(N, sizeof(*x), x, 1, x_gpu, 1)); - CHECK_BLAS ( cublasSetVector(N, sizeof(*y), y, 1, y_gpu, 1)); - - printf ("info: launch 'saxpy' kernel\n"); - CHECK_BLAS ( cublasSaxpy(handle, N, &a, x_gpu, 1, y_gpu, 1)); - - cudaDeviceSynchronize(); - - printf ("info: copy Device2Host\n"); - CHECK_BLAS ( cublasGetVector(N, sizeof(*y_gpu_res), y_gpu, 1, y_gpu_res, 1)); - - // CPU implementation of saxpy - for (int i = 0; i < N; i++) { - y_cpu_res[i] = a * x[i] + y[i]; - } - - // verify the results - int errors = 0; - for (int i = 0; i < N; i++) { - if (fabs(y_cpu_res[i] - y_gpu_res[i]) > fabs(y_cpu_res[i] * 0.0001f)) - errors++; - } - std::cout << errors << " errors" << std::endl; - - CHECK( cudaFree(x_gpu)); - CHECK( cudaFree(y_gpu)); - CHECK_BLAS( cublasDestroy(handle)); - - return errors; -} diff --git a/samples/7_Advanced/hipblas_saxpy/saxpy.hipblasref.cpp b/samples/7_Advanced/hipblas_saxpy/saxpy.hipblasref.cpp deleted file mode 100644 index 4610a612d4..0000000000 --- a/samples/7_Advanced/hipblas_saxpy/saxpy.hipblasref.cpp +++ /dev/null @@ -1,94 +0,0 @@ - -#include -#include -#include -#include - -// header file for the GPU API -#include "hip/hip_runtime.h" -#include - -#define N (1024 * 500) - -#define CHECK(cmd) \ -{\ - hipError_t error = cmd; \ - if (error != hipSuccess) { \ - fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \ - exit(EXIT_FAILURE);\ - }\ -} - -#define CHECK_BLAS(cmd) \ -{\ - hipblasStatus_t error = cmd;\ - if (error != HIPBLAS_STATUS_SUCCESS) { \ - fprintf(stderr, "error: (%d) at %s:%d\n", error,__FILE__, __LINE__); \ - exit(EXIT_FAILURE);\ - }\ -} - -int main() { - - const float a = 100.0f; - float x[N]; - float y[N], y_cpu_res[N], y_gpu_res[N]; - - // initialize the input data - std::default_random_engine random_gen; - std::uniform_real_distribution distribution(-N, N); - std::generate_n(x, N, [&]() { return distribution(random_gen); }); - std::generate_n(y, N, [&]() { return distribution(random_gen); }); - std::copy_n(y, N, y_cpu_res); - - // Explicit GPU code: - - size_t Nbytes = N*sizeof(float); - float *x_gpu, *y_gpu; - - hipblasHandle_t handle; - - hipDeviceProp_t props; - CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/)); - printf ("info: running on device %s\n", props.name); - - printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); - printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); - CHECK(hipMalloc(&x_gpu, Nbytes)); - CHECK(hipMalloc(&y_gpu, Nbytes)); - - // Initialize the blas library - CHECK_BLAS ( hipblasCreate(&handle)); - - // copy n elements from a vector in host memory space to a vector in GPU memory space - printf ("info: copy Host2Device\n"); - CHECK_BLAS ( hipblasSetVector(N, sizeof(*x), x, 1, x_gpu, 1)); - CHECK_BLAS ( hipblasSetVector(N, sizeof(*y), y, 1, y_gpu, 1)); - - printf ("info: launch 'saxpy' kernel\n"); - CHECK_BLAS ( hipblasSaxpy(handle, N, &a, x_gpu, 1, y_gpu, 1)); - - hipDeviceSynchronize(); - - printf ("info: copy Device2Host\n"); - CHECK_BLAS ( hipblasGetVector(N, sizeof(*y_gpu_res), y_gpu, 1, y_gpu_res, 1)); - - // CPU implementation of saxpy - for (int i = 0; i < N; i++) { - y_cpu_res[i] = a * x[i] + y[i]; - } - - // verify the results - int errors = 0; - for (int i = 0; i < N; i++) { - if (fabs(y_cpu_res[i] - y_gpu_res[i]) > fabs(y_cpu_res[i] * 0.0001f)) - errors++; - } - std::cout << errors << " errors" << std::endl; - - CHECK( hipFree(x_gpu)); - CHECK( hipFree(y_gpu)); - CHECK_BLAS( hipblasDestroy(handle)); - - return errors; -} From af22699ec601e3942849b016cf3a088228d4d6c7 Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Tue, 7 Mar 2017 11:24:32 -0600 Subject: [PATCH 08/22] added new field to hipDeviceProp_t structure gcnArch. 1. It is an integer containing gfx values 701, 801, 802, 803 2. On NV path, it is zero Change-Id: I2b4c7f48981d0214d8c6b1905d2cc85b16203419 --- samples/0_Intro/square/square.hipref.cpp | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/samples/0_Intro/square/square.hipref.cpp b/samples/0_Intro/square/square.hipref.cpp index 0073c1399a..e694bfb8a4 100644 --- a/samples/0_Intro/square/square.hipref.cpp +++ b/samples/0_Intro/square/square.hipref.cpp @@ -32,7 +32,7 @@ THE SOFTWARE. }\ } -/* +/* * Square each element in the array A and write to array C. */ template @@ -58,16 +58,18 @@ int main(int argc, char *argv[]) hipDeviceProp_t props; CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/)); printf ("info: running on device %s\n", props.name); - + #ifdef __HIP_PLATFORM_HCC__ + printf ("info: architecture on AMD GPU device is: %d\n",props.gcnArch); + #endif printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); A_h = (float*)malloc(Nbytes); CHECK(A_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); C_h = (float*)malloc(Nbytes); CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); // Fill with Phi + i - for (size_t i=0; i Date: Thu, 9 Mar 2017 08:52:50 -0600 Subject: [PATCH 09/22] make 4_shfl cookbook sample only for fiji 1. __shfl is not supported on hawaii gfx701 Change-Id: Iac09f5d30ee0674b8f58a6e74ec5c49b02be32ad --- samples/2_Cookbook/4_shfl/Makefile | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/samples/2_Cookbook/4_shfl/Makefile b/samples/2_Cookbook/4_shfl/Makefile index 3383cf2bf5..21c0e93959 100644 --- a/samples/2_Cookbook/4_shfl/Makefile +++ b/samples/2_Cookbook/4_shfl/Makefile @@ -22,7 +22,7 @@ CXX=$(HIPCC) $(EXECUTABLE): $(OBJECTS) - $(HIPCC) $(OBJECTS) -o $@ + $(HIPCC) --amdgpu-target=gfx803 $(OBJECTS) -o $@ test: $(EXECUTABLE) @@ -33,4 +33,3 @@ clean: rm -f $(EXECUTABLE) rm -f $(OBJECTS) rm -f $(HIP_PATH)/src/*.o - From 5eb39f1c6b26bc8c92bed2ae27df8cf3828647c9 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 10 Mar 2017 10:29:52 +0530 Subject: [PATCH 10/22] Fix for HCSWAP-128, make 5_2dshfl cookbook sample only for fiji Change-Id: I8869c28151bca1bd47a053a2808e93a801d16d00 --- samples/2_Cookbook/5_2dshfl/Makefile | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/samples/2_Cookbook/5_2dshfl/Makefile b/samples/2_Cookbook/5_2dshfl/Makefile index b742bbf80a..6abaf658b1 100644 --- a/samples/2_Cookbook/5_2dshfl/Makefile +++ b/samples/2_Cookbook/5_2dshfl/Makefile @@ -22,7 +22,7 @@ CXX=$(HIPCC) $(EXECUTABLE): $(OBJECTS) - $(HIPCC) $(OBJECTS) -o $@ + $(HIPCC) --amdgpu-target=gfx803 $(OBJECTS) -o $@ test: $(EXECUTABLE) From b25691cb87f8bb1010cb37a0973cbe9e5b572986 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 17 Feb 2017 17:14:55 -0600 Subject: [PATCH 11/22] Add first step to a "registerd" mode in hipBusBandwidth. --- .../hipBusBandwidth/hipBusBandwidth.cpp | 163 +++++++++++------- 1 file changed, 104 insertions(+), 59 deletions(-) diff --git a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index 7cb3e7908e..a1b2fd1705 100644 --- a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -6,9 +6,12 @@ #include "ResultDatabase.h" +enum MallocMode {MallocPinned, MallocUnpinned, MallocRegistered}; + // Cmdline parms: bool p_verbose = false; -bool p_pinned = true; +MallocMode p_malloc_mode = MallocPinned; +int p_numa_ctl = -1; int p_iterations = 10; int p_beatsperiteration=1; int p_device = 0; @@ -21,7 +24,7 @@ bool p_h2d = true; bool p_d2h = true; bool p_bidir = true; - +#define NO_CHECK #define CHECK_HIP_ERROR() \ @@ -36,6 +39,14 @@ bool p_bidir = true; } +std::string mallocModeString(int mallocMode) { + switch (mallocMode) { + case MallocPinned : return "pinned"; + case MallocUnpinned: return "unpinned"; + case MallocRegistered: return "registered"; + default: return "mallocmode-UNKNOWN"; + }; +}; // **************************************************************************** int sizeToBytes(int size) { @@ -106,7 +117,7 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) // Create some host memory pattern float *hostMem = NULL; - if (p_pinned) + if (p_malloc_mode == MallocPinned) { hipHostMalloc((void**)&hostMem, sizeof(float) * numMaxFloats); while (hipGetLastError() != hipSuccess) @@ -116,20 +127,29 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) --nSizes; if (nSizes < 1) { - std::cerr << "Error: Couldn't allocated any pinned buffer\n"; + std::cerr << "Error: Couldn't allocate any pinned buffer\n"; return; } numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; hipHostMalloc((void**)&hostMem, sizeof(float) * numMaxFloats); } } - else + else if (p_malloc_mode == MallocUnpinned) { if (p_alignedhost) { hostMem = (float*)aligned_alloc(p_alignedhost, numMaxFloats*sizeof(float)); } else { hostMem = new float[numMaxFloats]; } + } + else if (p_malloc_mode == MallocRegistered) + { + if (p_numa_ctl == -1) { + hostMem = (float*)malloc(numMaxFloats*sizeof(float)); + } + + hipHostRegister(hostMem, numMaxFloats * sizeof(float), 0); + CHECK_HIP_ERROR(); } for (int i = 0; i < numMaxFloats; i++) @@ -146,7 +166,7 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) --nSizes; if (nSizes < 1) { - std::cerr << "Error: Couldn't allocated any device buffer\n"; + std::cerr << "Error: Couldn't allocate any device buffer\n"; return; } numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; @@ -199,8 +219,8 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) } else { sprintf(sizeStr, "%9s", sizeToString(thisSize).c_str()); } - resultDB.AddResult(std::string("H2D_Bandwidth") + (p_pinned ? "_Pinned" : "_Unpinned"), sizeStr, "GB/sec", speed); - resultDB.AddResult(std::string("H2D_Time") + (p_pinned ? "_Pinned" : "_Unpinned"), sizeStr, "ms", t); + resultDB.AddResult(std::string("H2D_Bandwidth") + "_" + mallocModeString(p_malloc_mode), sizeStr, "GB/sec", speed); + resultDB.AddResult(std::string("H2D_Time") + mallocModeString(p_malloc_mode), sizeStr, "ms", t); if (p_onesize) { break; @@ -212,6 +232,8 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) numMaxFloats = sizeToBytes(p_onesize) / sizeof(float); } +#ifndef NO_CHECK + // Check. First reset the host memory, then copy-back result. Then compare against original ref value. for (int i = 0; i < numMaxFloats; i++) { @@ -225,24 +247,36 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) printf ("error: H2D. i=%d reference:%6.f != copyback:%6.2f\n", i, ref, hostMem[i]); } } +#endif // Cleanup hipFree((void*)device); CHECK_HIP_ERROR(); - if (p_pinned) - { + switch (p_malloc_mode) { + case MallocPinned: hipHostFree((void*)hostMem); CHECK_HIP_ERROR(); - } - else - { + break; + + case MallocUnpinned: if (p_alignedhost) { delete[] hostMem; } else { free(hostMem); } + break; + + case MallocRegistered: + hipHostUnregister(hostMem); + CHECK_HIP_ERROR(); + free(hostMem); + break; + default: + assert(0); } + + hipEventDestroy(start); hipEventDestroy(stop); } @@ -257,38 +291,40 @@ void RunBenchmark_D2H(ResultDatabase &resultDB) // Create some host memory pattern float *hostMem1; float *hostMem2; - if (p_pinned) + if (p_malloc_mode == MallocPinned) { hipHostMalloc((void**)&hostMem1, sizeof(float)*numMaxFloats); hipError_t err1 = hipGetLastError(); hipHostMalloc((void**)&hostMem2, sizeof(float)*numMaxFloats); hipError_t err2 = hipGetLastError(); - while (err1 != hipSuccess || err2 != hipSuccess) - { - // free the first buffer if only the second failed - if (err1 == hipSuccess) - hipHostFree((void*)hostMem1); + while (err1 != hipSuccess || err2 != hipSuccess) + { + // free the first buffer if only the second failed + if (err1 == hipSuccess) + hipHostFree((void*)hostMem1); - // drop the size and try again - if (p_verbose) std::cout << " - dropping size allocating pinned mem\n"; - --nSizes; - if (nSizes < 1) - { - std::cerr << "Error: Couldn't allocated any pinned buffer\n"; - return; - } - numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; - hipHostMalloc((void**)&hostMem1, sizeof(float)*numMaxFloats); - err1 = hipGetLastError(); - hipHostMalloc((void**)&hostMem2, sizeof(float)*numMaxFloats); - err2 = hipGetLastError(); - } - } - else + // drop the size and try again + if (p_verbose) std::cout << " - dropping size allocating pinned mem\n"; + --nSizes; + if (nSizes < 1) + { + std::cerr << "Error: Couldn't allocate any pinned buffer\n"; + return; + } + numMaxFloats = 1024 * (sizes[nSizes-1]) / 4; + hipHostMalloc((void**)&hostMem1, sizeof(float)*numMaxFloats); + err1 = hipGetLastError(); + hipHostMalloc((void**)&hostMem2, sizeof(float)*numMaxFloats); + err2 = hipGetLastError(); + } + } + else if (p_malloc_mode == MallocUnpinned) { hostMem1 = new float[numMaxFloats]; hostMem2 = new float[numMaxFloats]; } + + for (int i=0; i Date: Fri, 10 Mar 2017 15:04:46 -0600 Subject: [PATCH 12/22] Refactor registered memory calls. --- .../hipBusBandwidth/hipBusBandwidth.cpp | 52 ++++++++++++++++++- 1 file changed, 51 insertions(+), 1 deletion(-) diff --git a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index a1b2fd1705..09f78543c9 100644 --- a/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp @@ -24,7 +24,7 @@ bool p_h2d = true; bool p_d2h = true; bool p_bidir = true; -#define NO_CHECK +//#define NO_CHECK #define CHECK_HIP_ERROR() \ @@ -151,6 +151,10 @@ void RunBenchmark_H2D(ResultDatabase &resultDB) hipHostRegister(hostMem, numMaxFloats * sizeof(float), 0); CHECK_HIP_ERROR(); } + else + { + assert(0); + } for (int i = 0; i < numMaxFloats; i++) { @@ -323,6 +327,22 @@ void RunBenchmark_D2H(ResultDatabase &resultDB) hostMem1 = new float[numMaxFloats]; hostMem2 = new float[numMaxFloats]; } + else if (p_malloc_mode == MallocRegistered) + { + if (p_numa_ctl == -1) { + hostMem1 = (float*)malloc(numMaxFloats*sizeof(float)); + hostMem2 = (float*)malloc(numMaxFloats*sizeof(float)); + } + + hipHostRegister(hostMem1, numMaxFloats * sizeof(float), 0); + CHECK_HIP_ERROR(); + hipHostRegister(hostMem2, numMaxFloats * sizeof(float), 0); + CHECK_HIP_ERROR(); + } + else + { + assert(0); + } for (int i=0; i Date: Tue, 14 Mar 2017 15:56:18 +0530 Subject: [PATCH 13/22] 4_shfl and 5_2dshfl samples are unsupported on gfx701 Change-Id: I81eb880350f25e89573ba14c62b549c6c43f8c91 --- samples/2_Cookbook/4_shfl/Makefile | 6 +++++- samples/2_Cookbook/5_2dshfl/Makefile | 6 +++++- 2 files changed, 10 insertions(+), 2 deletions(-) diff --git a/samples/2_Cookbook/4_shfl/Makefile b/samples/2_Cookbook/4_shfl/Makefile index 21c0e93959..56f54d9518 100644 --- a/samples/2_Cookbook/4_shfl/Makefile +++ b/samples/2_Cookbook/4_shfl/Makefile @@ -3,6 +3,10 @@ ifeq (,$(HIP_PATH)) HIP_PATH=../../.. endif +ifeq (gfx701, $(findstring gfx701,$(HCC_AMDGPU_TARGET))) + $(error gfx701 is not a supported device for this sample) +endif + HIPCC=$(HIP_PATH)/bin/hipcc TARGET=hcc @@ -22,7 +26,7 @@ CXX=$(HIPCC) $(EXECUTABLE): $(OBJECTS) - $(HIPCC) --amdgpu-target=gfx803 $(OBJECTS) -o $@ + $(HIPCC) $(OBJECTS) -o $@ test: $(EXECUTABLE) diff --git a/samples/2_Cookbook/5_2dshfl/Makefile b/samples/2_Cookbook/5_2dshfl/Makefile index 6abaf658b1..cfadb1a311 100644 --- a/samples/2_Cookbook/5_2dshfl/Makefile +++ b/samples/2_Cookbook/5_2dshfl/Makefile @@ -3,6 +3,10 @@ ifeq (,$(HIP_PATH)) HIP_PATH=../../.. endif +ifeq (gfx701, $(findstring gfx701,$(HCC_AMDGPU_TARGET))) + $(error gfx701 is not a supported device for this sample) +endif + HIPCC=$(HIP_PATH)/bin/hipcc TARGET=hcc @@ -22,7 +26,7 @@ CXX=$(HIPCC) $(EXECUTABLE): $(OBJECTS) - $(HIPCC) --amdgpu-target=gfx803 $(OBJECTS) -o $@ + $(HIPCC) $(OBJECTS) -o $@ test: $(EXECUTABLE) From 219343027fdf45bcb86a0070665cb58f54502c7f Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Fri, 17 Mar 2017 13:11:34 -0500 Subject: [PATCH 14/22] Added default module launch api functionality 1. As in hipModuleLaunchKernel(..., kernelParams, nullptr); works with this commit 2. Added headers AMDGPUPTNote.h, AMDGPURuntimeMetadata.h to do code object meta data parsing 3. Changed CMake to look at llvm link libraries 4. HIP developer should set env variable LLVM_HOME to remove link errors 5. HIP depends on installed LLVM (not source, not build) 6. Added sample to test out the feature 7. Right now HCC does not support embedding metadata in code object. Use clang opencl 8. Changed HIPCC to read LLVM_HOME env var 9. New argument to CMake should be given -DLLVM_HOME= Change-Id: Iba38194aa872d97cc2c90a8e5ff746c48055c868 --- samples/0_Intro/module_api/Makefile | 8 +- samples/0_Intro/module_api/defaultDriver.cpp | 89 +++++++++++++++++++ samples/0_Intro/module_api/runKernel.cpp | 4 +- samples/0_Intro/module_api/test.cl | 12 +++ samples/0_Intro/module_api/test.co | Bin 0 -> 9824 bytes 5 files changed, 108 insertions(+), 5 deletions(-) create mode 100644 samples/0_Intro/module_api/defaultDriver.cpp create mode 100644 samples/0_Intro/module_api/test.cl create mode 100755 samples/0_Intro/module_api/test.co diff --git a/samples/0_Intro/module_api/Makefile b/samples/0_Intro/module_api/Makefile index f2c0ce555a..632a8d3e70 100644 --- a/samples/0_Intro/module_api/Makefile +++ b/samples/0_Intro/module_api/Makefile @@ -5,14 +5,16 @@ endif HIPCC=$(HIP_PATH)/bin/hipcc HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler) -all: vcpy_kernel.code runKernel.hip.out +all: vcpy_kernel.code runKernel.hip.out defaultDriver.hip.out runKernel.hip.out: runKernel.cpp $(HIPCC) $(HIPCC_FLAGS) $< -o $@ +defaultDriver.hip.out: defaultDriver.cpp + $(HIPCC) $(HIPCC_FLAGS) $< -o $@ + vcpy_kernel.code: vcpy_kernel.cpp - $(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@ + $(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@ clean: rm -f *.code *.out - diff --git a/samples/0_Intro/module_api/defaultDriver.cpp b/samples/0_Intro/module_api/defaultDriver.cpp new file mode 100644 index 0000000000..a29271a0dc --- /dev/null +++ b/samples/0_Intro/module_api/defaultDriver.cpp @@ -0,0 +1,89 @@ +/* +Copyright (c) 2015-2017 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/hip_runtime.h" +#include "hip/hip_runtime_api.h" +#include +#include +#include + +#define LEN 64 +#define SIZE LEN<<2 + +#define fileName "test.co" +#define kernel_name "vadd" + +int main(){ + float *A, *B, *C; + hipDeviceptr_t Ad, Bd, Cd; + A = new float[LEN]; + B = new float[LEN]; + C = new float[LEN]; + + for(uint32_t i=0;iH@ikcag#Q!>Xh`B3Zc@Bch`yO`hvZV z8c^LpLG*zK-~;#qzJcbkyp=xZp)W{%;2VH*c4oX@lU6MRTGTy~_spEjoHOT~nZ1l} z950_xGMPCdV1(Q!8En7E33D8g)kC5IXpY!q9IpGw1j#`^8__s_jw9f5*p3DxMB;Hl zM>6U?3@%S_j>OU*)Eh;X@NXF38R2>oOFrtMuY}?SfJWCoUKL%5rJe%xC()DhJNIko zMf0O#;32#!$kGn_hj143B7401*jKa_*+YjBhVX+iPL^249x?ii_Df3(ks7v<{wDk| zoJh}-u}BWOeHzdgmoIYJ?_qo#2W6+Hrzbw6tnTHP7FPUrr`Bu~>4#LGq8hzC6){eA z*r`UfS)^P}bT6?DzuIm#!t=FH{UWuDgmAg(vS7)tH`|+nZ<+-wA#(g;XW4JhUo6t3 z>Jqz#V6$>3okX>R~OsOZi|nutA?ErJHNPmt`r1Kcd^YH-2fcalGZOY zLl&Ir)={uiq_&zgQ}VpF-|4VQ;KxFm3Sq%ei4d&(#Yl1{(gwMRl3lQhSY$ zPXVM2Ao8e!B+eH4rJ%M3brBOz*BV}tikN}G3Jbdah0PWQy03+O?y9@a+I@1#bM>4vFQ~;}Q16q1nEP0f(G_FvEDp zdo-Z)5Dz0CW%36LY=e0o*ZEhoA|si)SgkcMk*J>)*g9D)BtEB_2fZULMh3;u?1Z9iIpA zhh6!I!6n3%-$4kj^8`KZq!CCXkVYVlz%vyA+{gJ3cYN5{#XkQqCrjiLfEcVMCSZnf+KJ%R4w0R2tOC@(z5D$wWS6@Ug44gVlun9{K%iLY{sdMXsioGy-V^ z(g>sxNF$I&AdNs8fiwbX1fJ;#$g}1LFK{;}g}3<5gAkpy16~%d`|_=Sf5a7M&d^pR zL(u!^Tf5VcuNGp1clS53#5wk_3hJ#SjG zOx18bdX(%o%#rmlY;}q=Gi$YQy<5q<&HBta=K0M=bZFVEx4NM}6KUo_glcZZa~$2Q zT2|Gr*sf{m%=dJM`F_PTD_*6l8LDBaN9p5@90VKnXkbG(w1VbY1&>twb`of+XcJoR8-Y44cjz*%hoJI_x&fHhV8hj zrMX^3ovr!>rn_pjU>H`V;5eq~!FhY7QgwYtJxa>8##O&vq~-DgPia5pyEX7T16EV{ zUlPV(!lv(9^Ao8&|P zS{U_Mz2=fU|3gFG?RYF?1Z%18H$VV1LskL2-U-{h#W6W{OuM2*SDg%=d|X-h;7n=h z)O^zUztO|EVT_`$N2o6w0?Yq%APWKiw=oE)i)&h{$#pLANg<4SXbZn>A;|TA0NS_* zNPW4^C8oebkbG%h;#Yx(=S7k)_YsNZK7wN)@H*jbydD%ZxlhS+DvA65^N}h;0~`W6 z>Y@)3&?H_JUFol!zxLu>A%ch8EE~v3t2?>h5knaG;9vrUjf=~`tM<20mLx81J_~wKLJkil%fCt literal 0 HcmV?d00001 From 323807d02beb026bcb282059d8579f339af66127 Mon Sep 17 00:00:00 2001 From: pensun Date: Fri, 17 Mar 2017 17:17:12 +0000 Subject: [PATCH 15/22] Initial integration with Alex' Generic Grid Launch Change-Id: I559afb80e9e39ec0d119bb3bf3b85ef9e448caf6 --- samples/0_Intro/square/square.hipref.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/samples/0_Intro/square/square.hipref.cpp b/samples/0_Intro/square/square.hipref.cpp index e694bfb8a4..118f8acf13 100644 --- a/samples/0_Intro/square/square.hipref.cpp +++ b/samples/0_Intro/square/square.hipref.cpp @@ -83,7 +83,7 @@ int main(int argc, char *argv[]) const unsigned threadsPerBlock = 256; printf ("info: launch 'vector_square' kernel\n"); - hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N); + hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, nullptr, C_d, A_d, N); printf ("info: copy Device2Host\n"); CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); From c4c4d95db66100423477b8109601159642d15713 Mon Sep 17 00:00:00 2001 From: "Sun, Peng" Date: Mon, 20 Mar 2017 17:03:21 -0500 Subject: [PATCH 16/22] revert workaround for square sample and update doc on GGL Change-Id: I731c68ca4111e7dc2e45bef51c4cad2c23fc81f8 --- samples/0_Intro/square/square.hipref.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/samples/0_Intro/square/square.hipref.cpp b/samples/0_Intro/square/square.hipref.cpp index 118f8acf13..963ab63260 100644 --- a/samples/0_Intro/square/square.hipref.cpp +++ b/samples/0_Intro/square/square.hipref.cpp @@ -83,7 +83,7 @@ int main(int argc, char *argv[]) const unsigned threadsPerBlock = 256; printf ("info: launch 'vector_square' kernel\n"); - hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, nullptr, C_d, A_d, N); + hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, nullptr, C_d, A_d, N); printf ("info: copy Device2Host\n"); CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); From 7735b454a133547a23b189f617ba2062666f4ade Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Fri, 31 Mar 2017 12:11:34 -0500 Subject: [PATCH 17/22] added new api hipHccModuleLaunchKernel 1. hipHccModuleLaunchKernel is same as hipModuleLaunchKernel with OpenCL workitem model 2. Added copy right 3. Fixed header naming Change-Id: I6a7c35a3566e2f8d3f5056613e34193775d4b236 --- samples/0_Intro/module_api/runKernel.cpp | 18 ++++++++++-------- samples/0_Intro/module_api/vcpy_kernel.cpp | 3 +-- 2 files changed, 11 insertions(+), 10 deletions(-) diff --git a/samples/0_Intro/module_api/runKernel.cpp b/samples/0_Intro/module_api/runKernel.cpp index 201892a4a1..e7d54beb54 100644 --- a/samples/0_Intro/module_api/runKernel.cpp +++ b/samples/0_Intro/module_api/runKernel.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015 - 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 @@ -22,9 +22,10 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include "hip/hip_runtime_api.h" -#include -#include -#include +#include +#include +#include +#include #define LEN 64 #define SIZE LEN<<2 @@ -43,10 +44,10 @@ int main(){ B[i] = 0.0f; } - hipInit(0); - hipDevice_t device; - hipCtx_t context; - hipDeviceGet(&device, 0); + hipInit(0); + hipDevice_t device; + hipCtx_t context; + hipDeviceGet(&device, 0); hipCtxCreate(&context, 0, device); hipMalloc((void**)&Ad, SIZE); @@ -101,6 +102,7 @@ int main(){ hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); hipMemcpyDtoH(B, Bd, SIZE); + int mismatchCount = 0; for(uint32_t i=0;i Date: Fri, 31 Mar 2017 13:35:26 -0500 Subject: [PATCH 18/22] Fixed bit_extract Change-Id: I92d7b7a302e3fa0db84889fb5dc6b612e6a53c73 --- samples/0_Intro/bit_extract/Makefile | 1 - samples/0_Intro/bit_extract/bit_extract.cpp | 6 +++--- 2 files changed, 3 insertions(+), 4 deletions(-) diff --git a/samples/0_Intro/bit_extract/Makefile b/samples/0_Intro/bit_extract/Makefile index 0965ae7296..78f6a2faa8 100644 --- a/samples/0_Intro/bit_extract/Makefile +++ b/samples/0_Intro/bit_extract/Makefile @@ -24,4 +24,3 @@ $(EXE): bit_extract.cpp clean: rm -f *.o $(EXE) - diff --git a/samples/0_Intro/bit_extract/bit_extract.cpp b/samples/0_Intro/bit_extract/bit_extract.cpp index 1535d2bd98..a30f2d052d 100644 --- a/samples/0_Intro/bit_extract/bit_extract.cpp +++ b/samples/0_Intro/bit_extract/bit_extract.cpp @@ -37,7 +37,7 @@ THE SOFTWARE. }\ } -void __global__ +__global__ void bit_extract_kernel(hipLaunchParm lp, uint32_t *C_d, const uint32_t *A_d, size_t N) { size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); @@ -45,7 +45,7 @@ bit_extract_kernel(hipLaunchParm lp, uint32_t *C_d, const uint32_t *A_d, size_t for (size_t i=offset; i> 8); #endif @@ -73,7 +73,7 @@ int main(int argc, char *argv[]) C_h = (uint32_t*)malloc(Nbytes); CHECK(C_h == 0 ? hipErrorMemoryAllocation : hipSuccess ); - for (size_t i=0; i Date: Fri, 31 Mar 2017 14:13:46 -0500 Subject: [PATCH 19/22] added debug support for HIP sample Change-Id: Ia7265234082039b68114f7421f4dbcb7149d4d2b --- samples/0_Intro/module_api/runKernel.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/samples/0_Intro/module_api/runKernel.cpp b/samples/0_Intro/module_api/runKernel.cpp index e7d54beb54..355f0bf5da 100644 --- a/samples/0_Intro/module_api/runKernel.cpp +++ b/samples/0_Intro/module_api/runKernel.cpp @@ -33,6 +33,9 @@ THE SOFTWARE. #define fileName "vcpy_kernel.code" #define kernel_name "hello_world" +#define HIP_CHECK(status) \ +if(status != hipSuccess) {std::cout<<"Got Status: "< Date: Fri, 31 Mar 2017 14:30:40 -0500 Subject: [PATCH 20/22] added module api sample which uses hipHccModuleLaunchKernel Change-Id: I7bce60b4480a3b5ff7ed69c3256078ded65a0945 --- samples/0_Intro/module_api/Makefile | 3 + .../0_Intro/module_api/launchKernelHcc.cpp | 112 ++++++++++++++++++ samples/0_Intro/module_api/runKernel.cpp | 4 - 3 files changed, 115 insertions(+), 4 deletions(-) create mode 100644 samples/0_Intro/module_api/launchKernelHcc.cpp diff --git a/samples/0_Intro/module_api/Makefile b/samples/0_Intro/module_api/Makefile index 632a8d3e70..38bd00a6a6 100644 --- a/samples/0_Intro/module_api/Makefile +++ b/samples/0_Intro/module_api/Makefile @@ -10,6 +10,9 @@ all: vcpy_kernel.code runKernel.hip.out defaultDriver.hip.out runKernel.hip.out: runKernel.cpp $(HIPCC) $(HIPCC_FLAGS) $< -o $@ +launchKernelHcc.hip.out: launchKernelHcc.cpp + $(HIPCC) $(HIPCC_FLAGS) $< -o $@ + defaultDriver.hip.out: defaultDriver.cpp $(HIPCC) $(HIPCC_FLAGS) $< -o $@ diff --git a/samples/0_Intro/module_api/launchKernelHcc.cpp b/samples/0_Intro/module_api/launchKernelHcc.cpp new file mode 100644 index 0000000000..e86e44cb24 --- /dev/null +++ b/samples/0_Intro/module_api/launchKernelHcc.cpp @@ -0,0 +1,112 @@ +/* +Copyright (c) 2015 - 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 +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/hip_runtime.h" +#include "hip/hip_runtime_api.h" +#include +#include +#include + +#ifdef __HIP_PLATFORM_HCC__ +#include +#endif + +#define LEN 64 +#define SIZE LEN<<2 + +#define fileName "vcpy_kernel.code" +#define kernel_name "hello_world" + +#define HIP_CHECK(status) \ +if(status != hipSuccess) {std::cout<<"Got Status: "< Date: Fri, 7 Apr 2017 15:40:09 +0530 Subject: [PATCH 21/22] Merge branch 'amd-develop' into amd-master Change-Id: I53d5a8916d769c4f0fe60d2ee3b240551da80b4f (cherry picked from commit 01c523f6c9fc5a66d5f6e60d5efe78c29c9fe317) --- samples/0_Intro/bit_extract/Makefile | 4 ---- samples/1_Utils/hipCommander/Makefile | 3 --- samples/1_Utils/hipCommander/hipCommander.cpp | 1 - 3 files changed, 8 deletions(-) diff --git a/samples/0_Intro/bit_extract/Makefile b/samples/0_Intro/bit_extract/Makefile index 78f6a2faa8..08bca6e642 100644 --- a/samples/0_Intro/bit_extract/Makefile +++ b/samples/0_Intro/bit_extract/Makefile @@ -11,10 +11,6 @@ HIPCC=$(HIP_PATH)/bin/hipcc ifeq (${HIP_PLATFORM}, nvcc) HIPCC_FLAGS = -gencode=arch=compute_20,code=sm_20 endif -ifeq (${HIP_PLATFORM}, hcc) - HIPCC_FLAGS = -stdlib=libc++ -endif - EXE=bit_extract diff --git a/samples/1_Utils/hipCommander/Makefile b/samples/1_Utils/hipCommander/Makefile index e770c636a4..a411763b7f 100644 --- a/samples/1_Utils/hipCommander/Makefile +++ b/samples/1_Utils/hipCommander/Makefile @@ -10,9 +10,6 @@ OPT=-O3 CXXFLAGS = $(OPT) --std=c++11 HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform) -ifeq (${HIP_PLATFORM}, hcc) - CXXFLAGS += " -stdlib=libc++" -endif CODE_OBJECTS=nullkernel.hsaco diff --git a/samples/1_Utils/hipCommander/hipCommander.cpp b/samples/1_Utils/hipCommander/hipCommander.cpp index 0add1ce3e3..4b93180b18 100644 --- a/samples/1_Utils/hipCommander/hipCommander.cpp +++ b/samples/1_Utils/hipCommander/hipCommander.cpp @@ -11,7 +11,6 @@ #include #include #include -#include #endif #include From b8fd2f159a4003d204f60f24186e3627afef6c3a Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 24 Apr 2017 08:50:43 +0530 Subject: [PATCH 22/22] Merge branch 'amd-develop' into amd-master Change-Id: I312fb9d1181733ef5160d1e993e2ae57ced0f6b3 (cherry picked from commit 88fb807af081f31314031d4549e98d9b621cfc41) --- samples/0_Intro/square/Makefile | 1 + samples/0_Intro/square/square.hipref.cpp | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/samples/0_Intro/square/Makefile b/samples/0_Intro/square/Makefile index 1e8cdba080..aa48cc5864 100644 --- a/samples/0_Intro/square/Makefile +++ b/samples/0_Intro/square/Makefile @@ -15,5 +15,6 @@ square.hip.out: square.hipref.cpp + clean: rm -f *.o *.out diff --git a/samples/0_Intro/square/square.hipref.cpp b/samples/0_Intro/square/square.hipref.cpp index 963ab63260..e694bfb8a4 100644 --- a/samples/0_Intro/square/square.hipref.cpp +++ b/samples/0_Intro/square/square.hipref.cpp @@ -83,7 +83,7 @@ int main(int argc, char *argv[]) const unsigned threadsPerBlock = 256; printf ("info: launch 'vector_square' kernel\n"); - hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, nullptr, C_d, A_d, N); + hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N); printf ("info: copy Device2Host\n"); CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));