From 7ba0720f7131ecbc5ca9b394e6823b21c085adeb 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 [ROCm/hip-tests commit: 70716bee425f1268b67ef8ac84f1d1a3188ed6de] --- .../samples/2_Cookbook/8_peer2peer/Makefile | 13 +++++++------ .../samples/7_Advanced/hipblas_saxpy/Makefile | 2 +- 2 files changed, 8 insertions(+), 7 deletions(-) diff --git a/projects/hip-tests/samples/2_Cookbook/8_peer2peer/Makefile b/projects/hip-tests/samples/2_Cookbook/8_peer2peer/Makefile index 5cb7473921..0bf9e6f93e 100644 --- a/projects/hip-tests/samples/2_Cookbook/8_peer2peer/Makefile +++ b/projects/hip-tests/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/projects/hip-tests/samples/7_Advanced/hipblas_saxpy/Makefile b/projects/hip-tests/samples/7_Advanced/hipblas_saxpy/Makefile index ed88be2dd0..8586e75d25 100644 --- a/projects/hip-tests/samples/7_Advanced/hipblas_saxpy/Makefile +++ b/projects/hip-tests/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 7ace03b518aae78f5696b877faa22ba282027e07 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 [ROCm/hip-tests commit: 054fc61f6e3beb7cd87c21fadf43ff4395bad410] --- .../samples/1_Utils/hipCommander/TODO | 50 ------------------- 1 file changed, 50 deletions(-) delete mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/TODO diff --git a/projects/hip-tests/samples/1_Utils/hipCommander/TODO b/projects/hip-tests/samples/1_Utils/hipCommander/TODO deleted file mode 100644 index 4c835cfced..0000000000 --- a/projects/hip-tests/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 257d60a3859a50fc84b71d326bf4afadaf3627e4 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 [ROCm/hip-tests commit: cee24a20f2edb40eea955e223729dd922944dac0] --- .../samples/1_Utils/hipInfo/hipInfo.cpp | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) diff --git a/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp b/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp index 42a879e732..cf4660eae7 100644 --- a/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/projects/hip-tests/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 b9ba4f6ca148183266fed272f59d907c092b28f1 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 [ROCm/hip-tests commit: ba8fe1675fe087bf44c6d37d7108246286bcd333] --- .../hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp b/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp index 9c07d066ba..0add1ce3e3 100644 --- a/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp +++ b/projects/hip-tests/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 09ede45f88a1d907a843eaa2606bb5293a8e78a3 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 [ROCm/hip-tests commit: 0390b12175aea6190b021c3758a675a4dffe4ad4] --- projects/hip-tests/samples/0_Intro/square/square.hipref.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp b/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp index 3c863b8b76..0073c1399a 100644 --- a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp +++ b/projects/hip-tests/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 78c27a5c51d4948fb51a021a9d26711a8a1d3304 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 [ROCm/hip-tests commit: dc2f5fcc7d9969d322bc4c7ae55556afc6234dc0] --- .../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 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_h2d.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_kernel.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_h2d.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_kernel.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_kernel.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_kernel.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_h2d.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_h2d.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_sync.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2d2h_wosync.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel_wosync.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_wosync.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2kernels.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2kernels_wosync.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_h2d.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_kernel.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_h2d.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_kernel.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_kernel.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_kernel.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_h2d.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_h2d.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_sync.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3d2h_wosync.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3h2d.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3h2d_wosync.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3kernels.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3kernels_wosync.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_4kernels.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_5kernels.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_6kernels.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_h2d.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_kernel.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_h2d.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_kernel.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_10.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h_wosync.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_wosync.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync_d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_barrier.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_h2d.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_d2h.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_h2d.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_streamcreate.hcm create mode 100644 projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_sync.hcm diff --git a/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_h2d.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_h2d.hcm new file mode 100644 index 0000000000..640bb2be79 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_kernel.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_kernel.hcm new file mode 100644 index 0000000000..c1bc0f6702 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_h2d.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_h2d.hcm new file mode 100644 index 0000000000..0e787f9bd0 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_kernel.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_d2h_sync_kernel.hcm new file mode 100644 index 0000000000..8d7fddc146 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_d2h.hcm new file mode 100644 index 0000000000..7d845d03a4 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_kernel.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_kernel.hcm new file mode 100644 index 0000000000..49c0d77146 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_d2h.hcm new file mode 100644 index 0000000000..fe1f14bee5 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_kernel.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_h2d_sync_kernel.hcm new file mode 100644 index 0000000000..0762001daa --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_d2h.hcm new file mode 100644 index 0000000000..88003ba476 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_h2d.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_h2d.hcm new file mode 100644 index 0000000000..01913f8481 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_d2h.hcm new file mode 100644 index 0000000000..530eb8f68e --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_h2d.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_kernel_sync_h2d.hcm new file mode 100644 index 0000000000..6d83ee87c9 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_sync.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2_sync.hcm new file mode 100644 index 0000000000..8b9e233a9e --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2d2h.hcm new file mode 100644 index 0000000000..83cdc4ff75 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2d2h_wosync.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2d2h_wosync.hcm new file mode 100644 index 0000000000..4b91403582 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d.hcm new file mode 100644 index 0000000000..a2e4311bf6 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel.hcm new file mode 100644 index 0000000000..0c622614cc --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel_wosync.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_kernel_wosync.hcm new file mode 100644 index 0000000000..d73467da10 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_wosync.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2h2d_wosync.hcm new file mode 100644 index 0000000000..35f5e68522 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2kernels.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2kernels.hcm new file mode 100644 index 0000000000..3b85c6bef8 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2kernels_wosync.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_2kernels_wosync.hcm new file mode 100644 index 0000000000..584d6b8021 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_h2d.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_h2d.hcm new file mode 100644 index 0000000000..7f0fce96c5 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_kernel.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_kernel.hcm new file mode 100644 index 0000000000..a384439b5c --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_h2d.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_h2d.hcm new file mode 100644 index 0000000000..1cab6ff0d2 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_kernel.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_d2h_sync_kernel.hcm new file mode 100644 index 0000000000..ff5b09a3dc --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_d2h.hcm new file mode 100644 index 0000000000..d8921a64e7 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_kernel.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_kernel.hcm new file mode 100644 index 0000000000..4ccbf9a83c --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_d2h.hcm new file mode 100644 index 0000000000..a3d9a282f5 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_kernel.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_h2d_sync_kernel.hcm new file mode 100644 index 0000000000..56554d15fd --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_d2h.hcm new file mode 100644 index 0000000000..a6e3a683d2 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_h2d.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_h2d.hcm new file mode 100644 index 0000000000..eae3eadc5b --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_d2h.hcm new file mode 100644 index 0000000000..9e21709b0b --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_h2d.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_kernel_sync_h2d.hcm new file mode 100644 index 0000000000..b1ef7ef9f4 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_sync.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3_sync.hcm new file mode 100644 index 0000000000..bc8d21c594 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3d2h.hcm new file mode 100644 index 0000000000..4e07574b99 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3d2h_wosync.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3d2h_wosync.hcm new file mode 100644 index 0000000000..e96707fed9 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3h2d.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3h2d.hcm new file mode 100644 index 0000000000..82151adb8b --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3h2d_wosync.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3h2d_wosync.hcm new file mode 100644 index 0000000000..7d96bfcfab --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3kernels.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3kernels.hcm new file mode 100644 index 0000000000..2e8306dfde --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3kernels_wosync.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_3kernels_wosync.hcm new file mode 100644 index 0000000000..85cd0dd4d2 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_4kernels.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_4kernels.hcm new file mode 100644 index 0000000000..48a8223626 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_5kernels.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_5kernels.hcm new file mode 100644 index 0000000000..70ad00c248 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_6kernels.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_6kernels.hcm new file mode 100644 index 0000000000..1bbb5694b1 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_d2h.hcm new file mode 100644 index 0000000000..54f06a3481 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_h2d.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_h2d.hcm new file mode 100644 index 0000000000..6667ba95fa --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_kernel.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_kernel.hcm new file mode 100644 index 0000000000..fe770c5e9d --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_h2d.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_h2d.hcm new file mode 100644 index 0000000000..20ec951509 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_kernel.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_d2h_sync_kernel.hcm new file mode 100644 index 0000000000..77e483b3df --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d.hcm new file mode 100644 index 0000000000..f5642bfdf0 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_10.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_10.hcm new file mode 100644 index 0000000000..05452b9c87 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_d2h.hcm new file mode 100644 index 0000000000..dad9fc7437 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel.hcm new file mode 100644 index 0000000000..1b60640b9e --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h.hcm new file mode 100644 index 0000000000..6e4e9f3544 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h_wosync.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_d2h_wosync.hcm new file mode 100644 index 0000000000..4e94a26ebf --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_wosync.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_kernel_wosync.hcm new file mode 100644 index 0000000000..b3b40d3190 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_d2h.hcm new file mode 100644 index 0000000000..030213d1b3 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync.hcm new file mode 100644 index 0000000000..146c74bcae --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync_d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_h2d_sync_kernel_sync_d2h.hcm new file mode 100644 index 0000000000..366d04f469 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel.hcm new file mode 100644 index 0000000000..027d89aad0 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_barrier.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_barrier.hcm new file mode 100644 index 0000000000..fb6a867e7f --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_d2h.hcm new file mode 100644 index 0000000000..2e64472dbd --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_h2d.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_h2d.hcm new file mode 100644 index 0000000000..b220a69c68 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_d2h.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_d2h.hcm new file mode 100644 index 0000000000..48b332b1c3 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_h2d.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_kernel_sync_h2d.hcm new file mode 100644 index 0000000000..5a45d55376 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_streamcreate.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_streamcreate.hcm new file mode 100644 index 0000000000..1e6aef5dc1 --- /dev/null +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_sync.hcm b/projects/hip-tests/samples/1_Utils/hipCommander/perf/scripts/latency_sync.hcm new file mode 100644 index 0000000000..a784013b1a --- /dev/null +++ b/projects/hip-tests/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 8e522c123c6bc3f38861d57cb9b61a6a18aca3f5 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 [ROCm/hip-tests commit: b5e3f593773648236a3c61ea6540b075cbf1ab67] --- .../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 projects/hip-tests/samples/7_Advanced/hipblas_saxpy/Makefile delete mode 100644 projects/hip-tests/samples/7_Advanced/hipblas_saxpy/saxpy.cublas.cpp delete mode 100644 projects/hip-tests/samples/7_Advanced/hipblas_saxpy/saxpy.hipblasref.cpp diff --git a/projects/hip-tests/samples/7_Advanced/hipblas_saxpy/Makefile b/projects/hip-tests/samples/7_Advanced/hipblas_saxpy/Makefile deleted file mode 100644 index 8586e75d25..0000000000 --- a/projects/hip-tests/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/projects/hip-tests/samples/7_Advanced/hipblas_saxpy/saxpy.cublas.cpp b/projects/hip-tests/samples/7_Advanced/hipblas_saxpy/saxpy.cublas.cpp deleted file mode 100644 index 03a38f3fb1..0000000000 --- a/projects/hip-tests/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/projects/hip-tests/samples/7_Advanced/hipblas_saxpy/saxpy.hipblasref.cpp b/projects/hip-tests/samples/7_Advanced/hipblas_saxpy/saxpy.hipblasref.cpp deleted file mode 100644 index 4610a612d4..0000000000 --- a/projects/hip-tests/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 e2b0e9c8ea3f94b58850d644d43b913e56fff3d9 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 [ROCm/hip-tests commit: af22699ec601e3942849b016cf3a088228d4d6c7] --- .../hip-tests/samples/0_Intro/square/square.hipref.cpp | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp b/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp index 0073c1399a..e694bfb8a4 100644 --- a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp +++ b/projects/hip-tests/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 [ROCm/hip-tests commit: a8fb90d9a924535a33a4cbfcc68cec0672a8adde] --- projects/hip-tests/samples/2_Cookbook/4_shfl/Makefile | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/projects/hip-tests/samples/2_Cookbook/4_shfl/Makefile b/projects/hip-tests/samples/2_Cookbook/4_shfl/Makefile index 3383cf2bf5..21c0e93959 100644 --- a/projects/hip-tests/samples/2_Cookbook/4_shfl/Makefile +++ b/projects/hip-tests/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 80820f361574885e0527f8ea9aa9b5f3af5dae3b 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 [ROCm/hip-tests commit: 5eb39f1c6b26bc8c92bed2ae27df8cf3828647c9] --- projects/hip-tests/samples/2_Cookbook/5_2dshfl/Makefile | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Makefile b/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Makefile index b742bbf80a..6abaf658b1 100644 --- a/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Makefile +++ b/projects/hip-tests/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 7dd5113bdf616d9fad8df36c003c00c75fa4fb2b 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. [ROCm/hip-tests commit: b25691cb87f8bb1010cb37a0973cbe9e5b572986] --- .../hipBusBandwidth/hipBusBandwidth.cpp | 163 +++++++++++------- 1 file changed, 104 insertions(+), 59 deletions(-) diff --git a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index 7cb3e7908e..a1b2fd1705 100644 --- a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/projects/hip-tests/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. [ROCm/hip-tests commit: 705ab93664add412e6d1d21bdb266d024c9991e2] --- .../hipBusBandwidth/hipBusBandwidth.cpp | 52 ++++++++++++++++++- 1 file changed, 51 insertions(+), 1 deletion(-) diff --git a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp b/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp index a1b2fd1705..09f78543c9 100644 --- a/projects/hip-tests/samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp +++ b/projects/hip-tests/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 [ROCm/hip-tests commit: 8d6cb1f5a37bee4e693623c81978b1b88a799875] --- projects/hip-tests/samples/2_Cookbook/4_shfl/Makefile | 6 +++++- projects/hip-tests/samples/2_Cookbook/5_2dshfl/Makefile | 6 +++++- 2 files changed, 10 insertions(+), 2 deletions(-) diff --git a/projects/hip-tests/samples/2_Cookbook/4_shfl/Makefile b/projects/hip-tests/samples/2_Cookbook/4_shfl/Makefile index 21c0e93959..56f54d9518 100644 --- a/projects/hip-tests/samples/2_Cookbook/4_shfl/Makefile +++ b/projects/hip-tests/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/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Makefile b/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Makefile index 6abaf658b1..cfadb1a311 100644 --- a/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Makefile +++ b/projects/hip-tests/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 1ec758282082f64288407b5cf4ad3bb302d3a134 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 [ROCm/hip-tests commit: 219343027fdf45bcb86a0070665cb58f54502c7f] --- .../samples/0_Intro/module_api/Makefile | 8 +- .../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 projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp create mode 100644 projects/hip-tests/samples/0_Intro/module_api/test.cl create mode 100755 projects/hip-tests/samples/0_Intro/module_api/test.co diff --git a/projects/hip-tests/samples/0_Intro/module_api/Makefile b/projects/hip-tests/samples/0_Intro/module_api/Makefile index f2c0ce555a..632a8d3e70 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/Makefile +++ b/projects/hip-tests/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/projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp b/projects/hip-tests/samples/0_Intro/module_api/defaultDriver.cpp new file mode 100644 index 0000000000..a29271a0dc --- /dev/null +++ b/projects/hip-tests/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 d33885889d3a17bb68853e13aa1fa391ddcbd642 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 [ROCm/hip-tests commit: 323807d02beb026bcb282059d8579f339af66127] --- projects/hip-tests/samples/0_Intro/square/square.hipref.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp b/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp index e694bfb8a4..118f8acf13 100644 --- a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp +++ b/projects/hip-tests/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 d59dbb3fd6c3a8c90c904c66dcc927b3bd9eafd7 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 [ROCm/hip-tests commit: c4c4d95db66100423477b8109601159642d15713] --- projects/hip-tests/samples/0_Intro/square/square.hipref.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp b/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp index 118f8acf13..963ab63260 100644 --- a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp +++ b/projects/hip-tests/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 60636db8ee44f2e3d3b89b8b524a16242fd59152 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 [ROCm/hip-tests commit: 7735b454a133547a23b189f617ba2062666f4ade] --- .../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/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp b/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp index 201892a4a1..e7d54beb54 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp +++ b/projects/hip-tests/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 [ROCm/hip-tests commit: 8bc80debe45e5442740cfb2d2b79bc1d43deecff] --- projects/hip-tests/samples/0_Intro/bit_extract/Makefile | 1 - .../hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp | 6 +++--- 2 files changed, 3 insertions(+), 4 deletions(-) diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/Makefile b/projects/hip-tests/samples/0_Intro/bit_extract/Makefile index 0965ae7296..78f6a2faa8 100644 --- a/projects/hip-tests/samples/0_Intro/bit_extract/Makefile +++ b/projects/hip-tests/samples/0_Intro/bit_extract/Makefile @@ -24,4 +24,3 @@ $(EXE): bit_extract.cpp clean: rm -f *.o $(EXE) - diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp b/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp index 1535d2bd98..a30f2d052d 100644 --- a/projects/hip-tests/samples/0_Intro/bit_extract/bit_extract.cpp +++ b/projects/hip-tests/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 [ROCm/hip-tests commit: 93a0b55616853e4a47e717dc93b5bfef0b42da78] --- .../hip-tests/samples/0_Intro/module_api/runKernel.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp b/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp index e7d54beb54..355f0bf5da 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/runKernel.cpp +++ b/projects/hip-tests/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 [ROCm/hip-tests commit: 8e2b7147a51bbe5b3a1d22b2a635970b439503e1] --- .../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 projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp diff --git a/projects/hip-tests/samples/0_Intro/module_api/Makefile b/projects/hip-tests/samples/0_Intro/module_api/Makefile index 632a8d3e70..38bd00a6a6 100644 --- a/projects/hip-tests/samples/0_Intro/module_api/Makefile +++ b/projects/hip-tests/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/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp b/projects/hip-tests/samples/0_Intro/module_api/launchKernelHcc.cpp new file mode 100644 index 0000000000..e86e44cb24 --- /dev/null +++ b/projects/hip-tests/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) [ROCm/hip-tests commit: 9fcc03e2b64dde04a82935ce2fd559a262498108] --- projects/hip-tests/samples/0_Intro/bit_extract/Makefile | 4 ---- projects/hip-tests/samples/1_Utils/hipCommander/Makefile | 3 --- .../hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp | 1 - 3 files changed, 8 deletions(-) diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/Makefile b/projects/hip-tests/samples/0_Intro/bit_extract/Makefile index 78f6a2faa8..08bca6e642 100644 --- a/projects/hip-tests/samples/0_Intro/bit_extract/Makefile +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/Makefile b/projects/hip-tests/samples/1_Utils/hipCommander/Makefile index e770c636a4..a411763b7f 100644 --- a/projects/hip-tests/samples/1_Utils/hipCommander/Makefile +++ b/projects/hip-tests/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/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp b/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp index 0add1ce3e3..4b93180b18 100644 --- a/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp +++ b/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp @@ -11,7 +11,6 @@ #include #include #include -#include #endif #include From 182d5261b3b72821f0254f4c916e46cf1f525a86 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) [ROCm/hip-tests commit: b8fd2f159a4003d204f60f24186e3627afef6c3a] --- projects/hip-tests/samples/0_Intro/square/Makefile | 1 + projects/hip-tests/samples/0_Intro/square/square.hipref.cpp | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/hip-tests/samples/0_Intro/square/Makefile b/projects/hip-tests/samples/0_Intro/square/Makefile index 1e8cdba080..aa48cc5864 100644 --- a/projects/hip-tests/samples/0_Intro/square/Makefile +++ b/projects/hip-tests/samples/0_Intro/square/Makefile @@ -15,5 +15,6 @@ square.hip.out: square.hipref.cpp + clean: rm -f *.o *.out diff --git a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp b/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp index 963ab63260..e694bfb8a4 100644 --- a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp +++ b/projects/hip-tests/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));