From 23dc2f02cd361bfc3d0d410b8493f037c06ed81e Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Wed, 26 Oct 2016 10:30:42 -0500
Subject: [PATCH 1/8] Rename HIP_ATP_MARKER and profiling vars

HIP_PROFILE_API
HIP_DB_START_API
HIP_DB_STOP_API

Change-Id: I6c4da67212ff8217e6356a2622d4c6278a188c34
---
 samples/2_Cookbook/{2_HIP_ATP_MARKER => 2_CodeXL_ATP}/Makefile    | 0
 .../{2_HIP_ATP_MARKER => 2_CodeXL_ATP}/MatrixTranspose.cpp        | 0
 samples/2_Cookbook/{2_HIP_ATP_MARKER => 2_CodeXL_ATP}/Readme.md   | 0
 3 files changed, 0 insertions(+), 0 deletions(-)
 rename samples/2_Cookbook/{2_HIP_ATP_MARKER => 2_CodeXL_ATP}/Makefile (100%)
 rename samples/2_Cookbook/{2_HIP_ATP_MARKER => 2_CodeXL_ATP}/MatrixTranspose.cpp (100%)
 rename samples/2_Cookbook/{2_HIP_ATP_MARKER => 2_CodeXL_ATP}/Readme.md (100%)

diff --git a/samples/2_Cookbook/2_HIP_ATP_MARKER/Makefile b/samples/2_Cookbook/2_CodeXL_ATP/Makefile
similarity index 100%
rename from samples/2_Cookbook/2_HIP_ATP_MARKER/Makefile
rename to samples/2_Cookbook/2_CodeXL_ATP/Makefile
diff --git a/samples/2_Cookbook/2_HIP_ATP_MARKER/MatrixTranspose.cpp b/samples/2_Cookbook/2_CodeXL_ATP/MatrixTranspose.cpp
similarity index 100%
rename from samples/2_Cookbook/2_HIP_ATP_MARKER/MatrixTranspose.cpp
rename to samples/2_Cookbook/2_CodeXL_ATP/MatrixTranspose.cpp
diff --git a/samples/2_Cookbook/2_HIP_ATP_MARKER/Readme.md b/samples/2_Cookbook/2_CodeXL_ATP/Readme.md
similarity index 100%
rename from samples/2_Cookbook/2_HIP_ATP_MARKER/Readme.md
rename to samples/2_Cookbook/2_CodeXL_ATP/Readme.md

From f8c4fa982a7ccb60967a8ab383fa71931c58bb4e Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Thu, 27 Oct 2016 20:37:46 -0500
Subject: [PATCH 2/8] Add new hipdemangleatp and snapshot sample update for new
 functionality

Change-Id: Ie19c683b2b0bdfeb0c3fcf89444c2e21b7f606e7
---
 samples/0_Intro/square/Makefile               |   6 +-
 .../2_CodeXL_ATP/MatrixTranspose.cpp          | 172 ---------------
 .../{2_CodeXL_ATP => 2_Profiler}/Makefile     |  22 +-
 .../2_Cookbook/2_Profiler/MatrixTranspose.cpp | 195 ++++++++++++++++++
 .../{2_CodeXL_ATP => 2_Profiler}/Readme.md    |  14 +-
 5 files changed, 218 insertions(+), 191 deletions(-)
 delete mode 100644 samples/2_Cookbook/2_CodeXL_ATP/MatrixTranspose.cpp
 rename samples/2_Cookbook/{2_CodeXL_ATP => 2_Profiler}/Makefile (52%)
 create mode 100644 samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp
 rename samples/2_Cookbook/{2_CodeXL_ATP => 2_Profiler}/Readme.md (83%)

diff --git a/samples/0_Intro/square/Makefile b/samples/0_Intro/square/Makefile
index 89921c2072..1e8cdba080 100644
--- a/samples/0_Intro/square/Makefile
+++ b/samples/0_Intro/square/Makefile
@@ -1,7 +1,4 @@
 HIP_PATH?= $(wildcard /opt/rocm/hip)
-ifeq (,$(HIP_PATH))
-	HIP_PATH=../../..
-endif
 HIPCC=$(HIP_PATH)/bin/hipcc
 
 all: square.hip.out
@@ -11,9 +8,10 @@ square.cuda.out : square.cu
 
 #hipify square.cu > square.cpp
 # Then review & finish port in square.cpp
+#
 
 square.hip.out: square.hipref.cpp
-	$(HIPCC) square.hipref.cpp -o $@
+	$(HIPCC) $(CXXFLAGS) square.hipref.cpp -o $@
 
 
 
diff --git a/samples/2_Cookbook/2_CodeXL_ATP/MatrixTranspose.cpp b/samples/2_Cookbook/2_CodeXL_ATP/MatrixTranspose.cpp
deleted file mode 100644
index f2aea146e4..0000000000
--- a/samples/2_Cookbook/2_CodeXL_ATP/MatrixTranspose.cpp
+++ /dev/null
@@ -1,172 +0,0 @@
-/*
-Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
-
-Permission is hereby granted, free of charge, to any person obtaining a copy
-of this software and associated documentation files (the "Software"), to deal
-in the Software without restriction, including without limitation the rights
-to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
-copies of the Software, and to permit persons to whom the Software is
-furnished to do so, subject to the following conditions:
-
-The above copyright notice and this permission notice shall be included in
-all copies or substantial portions of the Software.
-
-THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
-IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
-FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
-AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
-LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
-OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
-THE SOFTWARE.
-*/
-
-#include<iostream>
-
-// hip header file
-#include "hip/hip_runtime.h"
-
-#define WIDTH     1024
-
-#define NUM       (WIDTH*WIDTH)
-
-#define THREADS_PER_BLOCK_X  4
-#define THREADS_PER_BLOCK_Y  4
-#define THREADS_PER_BLOCK_Z  1
-
-// Device (Kernel) function, it must be void
-// hipLaunchParm provides the execution configuration
-__global__ void matrixTranspose(hipLaunchParm lp,
-                                float *out,
-                                float *in,
-                                const int width)
-{
-    int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
-    int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
-
-    out[y * width + x] = in[x * width + y];
-}
-
-// CPU implementation of matrix transpose
-void matrixTransposeCPUReference(
-    float * output,
-    float * input,
-    const unsigned int width)
-{
-    for(unsigned int j=0; j < width; j++)
-    {
-        for(unsigned int i=0; i < width; i++)
-        {
-            output[i*width + j] = input[j*width + i];
-        }
-    }
-}
-
-int main() {
-
-  float* Matrix;
-  float* TransposeMatrix;
-  float* cpuTransposeMatrix;
-
-  float* gpuMatrix;
-  float* gpuTransposeMatrix;
-
-  hipDeviceProp_t devProp;
-  hipGetDeviceProperties(&devProp, 0);
-
-  std::cout << "Device name " << devProp.name << std::endl;
-
-  hipEvent_t start, stop;
-  hipEventCreate(&start);
-  hipEventCreate(&stop);
-  float eventMs = 1.0f;
-
-  int i;
-  int errors;
-
-  Matrix = (float*)malloc(NUM * sizeof(float));
-  TransposeMatrix = (float*)malloc(NUM * sizeof(float));
-  cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
-
-  // initialize the input data
-  for (i = 0; i < NUM; i++) {
-    Matrix[i] = (float)i*10.0f;
-  }
-
-  // allocate the memory on the device side
-  hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
-  hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
-
-  // Record the start event
-  hipEventRecord(start, NULL);
-
-  // Memory transfer from host to device
-  hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);
-
-  // Record the stop event
-  hipEventRecord(stop, NULL);
-  hipEventSynchronize(stop);
-
-  hipEventElapsedTime(&eventMs, start, stop);
-
-  printf ("hipMemcpyHostToDevice time taken  = %6.3fms\n", eventMs);
-
-  // Record the start event
-  hipEventRecord(start, NULL);
-
-  // Lauching kernel from host
-  hipLaunchKernel(matrixTranspose,
-                  dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
-                  dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
-                  0, 0,
-                  gpuTransposeMatrix , gpuMatrix, WIDTH);
-
-  // Record the stop event
-  hipEventRecord(stop, NULL);
-  hipEventSynchronize(stop);
-
-  hipEventElapsedTime(&eventMs, start, stop);
-
-  printf ("kernel Execution time             = %6.3fms\n", eventMs);
-
-  // Record the start event
-  hipEventRecord(start, NULL);
-
-  // Memory transfer from device to host
-  hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost);
-
-  // Record the stop event
-  hipEventRecord(stop, NULL);
-  hipEventSynchronize(stop);
-
-  hipEventElapsedTime(&eventMs, start, stop);
-
-  printf ("hipMemcpyDeviceToHost time taken  = %6.3fms\n", eventMs);
-
-  // CPU MatrixTranspose computation
-  matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
-
-  // verify the results
-  errors = 0;
-  double eps = 1.0E-6;
-  for (i = 0; i < NUM; i++) {
-    if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) {
-      errors++;
-    }
-  }
-  if (errors!=0) {
-    printf("FAILED: %d errors\n",errors);
-  } else {
-    printf ("PASSED!\n");
-  }
-
-  //free the resources on device side
-  hipFree(gpuMatrix);
-  hipFree(gpuTransposeMatrix);
-
-  //free the resources on host side
-  free(Matrix);
-  free(TransposeMatrix);
-  free(cpuTransposeMatrix);
-
-  return errors;
-}
diff --git a/samples/2_Cookbook/2_CodeXL_ATP/Makefile b/samples/2_Cookbook/2_Profiler/Makefile
similarity index 52%
rename from samples/2_Cookbook/2_CodeXL_ATP/Makefile
rename to samples/2_Cookbook/2_Profiler/Makefile
index d3630a1c19..4b9a063f38 100644
--- a/samples/2_Cookbook/2_CodeXL_ATP/Makefile
+++ b/samples/2_Cookbook/2_Profiler/Makefile
@@ -1,10 +1,12 @@
 HIP_PATH?= $(wildcard /opt/rocm/hip)
-ifeq (,$(HIP_PATH))
-	HIP_PATH=../../..
-endif
 
 HIPCC=$(HIP_PATH)/bin/hipcc
 
+
+HIPPROFILER=/opt/rocm/bin/rocm-profiler
+PROFILER_OPT=-A -o MT.atp -e HIP_PROFILE_API=1
+HIPPROFILER_POST_CMD=$(HIP_PATH)/bin/hipdemangleatp MT.atp
+
 TARGET=hcc
 
 SOURCES = MatrixTranspose.cpp
@@ -15,9 +17,12 @@ EXECUTABLE=./MatrixTranspose
 .PHONY: test
 
 
-all: $(EXECUTABLE) test
+all: $(EXECUTABLE) profile
 
-CXXFLAGS =-g
+
+
+OPT =-g 
+CXXFLAGS =$(OPT) 
 CXX=$(HIPCC)
 
 
@@ -25,7 +30,12 @@ $(EXECUTABLE): $(OBJECTS)
 	$(HIPCC) $(OBJECTS) -o $@
 
 
-test: $(EXECUTABLE)
+profile: $(EXECUTABLE)
+	$(HIPPROFILER) $(PROFILER_OPT) $(EXECUTABLE)
+	$(HIPPROFILER_POST_CMD)
+
+
+run: $(EXECUTABLE)
 	$(EXECUTABLE)
 
 
diff --git a/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp b/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp
new file mode 100644
index 0000000000..7500957dfc
--- /dev/null
+++ b/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp
@@ -0,0 +1,195 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include<iostream>
+
+// hip header file
+#include "hip/hip_runtime.h"
+#include "hip/hip_profile.h"
+
+#define WIDTH     1024
+
+#define NUM       (WIDTH*WIDTH)
+
+#define THREADS_PER_BLOCK_X  4
+#define THREADS_PER_BLOCK_Y  4
+#define THREADS_PER_BLOCK_Z  1
+
+// Device (Kernel) function, it must be void
+// hipLaunchParm provides the execution configuration
+__global__ void matrixTranspose(hipLaunchParm lp,
+                                float *out,
+                                float *in,
+                                const int width)
+{
+    int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+    int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+
+    out[y * width + x] = in[x * width + y];
+}
+
+// CPU implementation of matrix transpose
+void matrixTransposeCPUReference(
+    float * output,
+    float * input,
+    const unsigned int width)
+{
+    for(unsigned int j=0; j < width; j++)
+    {
+        for(unsigned int i=0; i < width; i++)
+        {
+            output[i*width + j] = input[j*width + i];
+        }
+    }
+}
+
+int main() {
+
+  //HIP_SCOPED_MARKER(__func__, "MainFunc");
+  HIP_BEGIN_MARKER(__func__, "MainFunc");
+
+  float* Matrix;
+  float* TransposeMatrix;
+  float* cpuTransposeMatrix;
+
+  float* gpuMatrix;
+  float* gpuTransposeMatrix;
+
+  hipDeviceProp_t devProp;
+  hipGetDeviceProperties(&devProp, 0);
+
+  std::cout << "Device name " << devProp.name << std::endl;
+
+  hipEvent_t start, stop;
+  float eventMs = 1.0f;
+  {
+      // Show example of how to create a "scoped marker".  
+      // The scoped marker records the time spent inside the { scope } of the marker - the begin timestamp is at the
+      // beginning of the code scope, and the end is recorded when the SCOPE exits.  This can be viewed in CodeXL
+      // timeline relative to other GPU and CPU events.
+      // This marker captures the time spent in setup including host allocation, initialization, and device memory allocation.
+      HIP_SCOPED_MARKER("Setup", "App");
+
+      hipEventCreate(&start);
+      hipEventCreate(&stop);
+
+
+      Matrix = (float*)malloc(NUM * sizeof(float));
+      TransposeMatrix = (float*)malloc(NUM * sizeof(float));
+      cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
+
+      // initialize the input data
+      for (int i = 0; i < NUM; i++) {
+        Matrix[i] = (float)i*10.0f;
+      }
+
+
+
+      // allocate the memory on the device side
+      hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
+      hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
+  }
+
+  {
+      HIP_SCOPED_MARKER("Loop", "App");
+
+      // Record the start event
+      hipEventRecord(start, NULL);
+
+      // Memory transfer from host to device
+      hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);
+
+      // Record the stop event
+      hipEventRecord(stop, NULL);
+      hipEventSynchronize(stop);
+
+      hipEventElapsedTime(&eventMs, start, stop);
+
+      printf ("hipMemcpyHostToDevice time taken  = %6.3fms\n", eventMs);
+
+      // Record the start event
+      hipEventRecord(start, NULL);
+
+      // Lauching kernel from host
+      hipLaunchKernel(matrixTranspose,
+                      dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
+                      dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
+                      0, 0,
+                      gpuTransposeMatrix , gpuMatrix, WIDTH);
+
+      // Record the stop event
+      hipEventRecord(stop, NULL);
+      hipEventSynchronize(stop);
+      hipEventElapsedTime(&eventMs, start, stop);
+  }
+
+  int errors = 0;
+  {
+      HIP_SCOPED_MARKER("Teardown", "App");
+
+
+      printf ("kernel Execution time             = %6.3fms\n", eventMs);
+
+      // Record the start event
+      hipEventRecord(start, NULL);
+
+      // Memory transfer from device to host
+      hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost);
+
+      // Record the stop event
+      hipEventRecord(stop, NULL);
+      hipEventSynchronize(stop);
+
+      hipEventElapsedTime(&eventMs, start, stop);
+
+      printf ("hipMemcpyDeviceToHost time taken  = %6.3fms\n", eventMs);
+
+      // CPU MatrixTranspose computation
+      matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
+
+      // verify the results
+      double eps = 1.0E-6;
+      for (int i = 0; i < NUM; i++) {
+        if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) {
+          errors++;
+        }
+      }
+      if (errors!=0) {
+        printf("FAILED: %d errors\n",errors);
+      } else {
+        printf ("PASSED!\n");
+      }
+
+      //free the resources on device side
+      hipFree(gpuMatrix);
+      hipFree(gpuTransposeMatrix);
+
+      //free the resources on host side
+      free(Matrix);
+      free(TransposeMatrix);
+      free(cpuTransposeMatrix);
+  }
+
+  HIP_END_MARKER();
+  
+  return errors;
+}
diff --git a/samples/2_Cookbook/2_CodeXL_ATP/Readme.md b/samples/2_Cookbook/2_Profiler/Readme.md
similarity index 83%
rename from samples/2_Cookbook/2_CodeXL_ATP/Readme.md
rename to samples/2_Cookbook/2_Profiler/Readme.md
index de1a800572..92a8be228e 100644
--- a/samples/2_Cookbook/2_CodeXL_ATP/Readme.md
+++ b/samples/2_Cookbook/2_Profiler/Readme.md
@@ -1,6 +1,6 @@
 ## Using hipEvents to measure performance ###
 
-This tutorial is follow-up of the previous two tutorial where we learn how to write our first hip program, in which we compute Matrix Transpose and in second one, we added feature to measure time taken for memory transfer and kernel execution. In this tutorial, we won't make amy changes to the source code. We'll explain how to use the  codexl/rocm-profiler for hip timeline tracing.
+This tutorial is follow-up of the previous two tutorial where we learn how to write our first hip program, in which we compute Matrix Transpose and in second one, we added feature to measure time taken for memory transfer and kernel execution. In this tutorial, we'll explain how to use the codexl/rocm-profiler for hip timeline tracing.  Also, we will augment the source code with additional markers so we can see the high-level application flow alongside the information that CodeXL automatically collects.
 
 
 ## Introduction:
@@ -24,15 +24,11 @@ HIP can generate markers at function being/end which are displayed on the CodeXL
 
 1. Install ROCm-Profiler Installing HIP from the rocm pre-built packages, installs the ROCm-Profiler as well. Alternatively, you can build ROCm-Profiler using the instructions given below.
 
-2. Build HIP with ATP markers enabled HIP pre-built packages are enabled with ATP marker support by default. To enable ATP marker support when building HIP from source, use the option -DCOMPILE_HIP_ATP_MARKER=1 during the cmake configure step.
 
-3. Set HIP_ATP_MARKER
-`export HIP_ATP_MARKER=1`
-
-4. Recompile the target application
-
-5. Run with profiler enabled to generate ATP file.
-`/opt/rocm/bin/rocm-profiler -o <outputATPFileName> -A <applicationName> <applicationArguments>`
+2. Run with profiler enabled to generate ATP file.
+(These steps are also captured in the Makefile)
+The HIP_PROFILE_API enables display of the HIP APIs on the CodeXL trimeline view.
+`/opt/rocm/bin/rocm-profiler -o <outputATPFileName> -A <applicationName> -e HIP_PROFILE_API=1 <applicationArguments>`
 
 ##Using HIP_TRACE_API
 

From e2fbab109d117de9aecf809bf515f0f1f38d5a15 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Thu, 27 Oct 2016 21:26:28 -0500
Subject: [PATCH 3/8] show how to use variety of HIP_PROFILE features

Change-Id: I6edd66ac4c068b64e1dc3787d7f1f69ab3238469
---
 .../2_Cookbook/2_Profiler/MatrixTranspose.cpp | 177 ++++++++++--------
 1 file changed, 96 insertions(+), 81 deletions(-)

diff --git a/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp b/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp
index 7500957dfc..b6a6b141d2 100644
--- a/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp
+++ b/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp
@@ -34,6 +34,8 @@ THE SOFTWARE.
 #define THREADS_PER_BLOCK_Y  4
 #define THREADS_PER_BLOCK_Z  1
 
+#define ITERATIONS 10
+
 // Device (Kernel) function, it must be void
 // hipLaunchParm provides the execution configuration
 __global__ void matrixTranspose(hipLaunchParm lp,
@@ -62,10 +64,72 @@ void matrixTransposeCPUReference(
     }
 }
 
-int main() {
 
-  //HIP_SCOPED_MARKER(__func__, "MainFunc");
-  HIP_BEGIN_MARKER(__func__, "MainFunc");
+// Use a separate function to demonstrate how to use function name as part of scoped marker:
+void runGPU(float *Matrix, float *TransposeMatrix, 
+            float* gpuMatrix, float* gpuTransposeMatrix)  {
+
+  // __func__ is a standard C++ macro which expands to the name of the function, in this case "runGPU"
+  HIP_SCOPED_MARKER(__func__, "MyGroup");
+
+  for (int i=0; i<ITERATIONS; i++) {
+
+    float eventMs = 0.0f;
+
+    hipEvent_t start, stop;
+    hipEventCreate(&start);
+    hipEventCreate(&stop);
+
+
+    // Record the start event
+    hipEventRecord(start, NULL);
+
+    // Memory transfer from host to device
+    hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);
+
+    // Record the stop event
+    hipEventRecord(stop, NULL);
+    hipEventSynchronize(stop);
+
+    hipEventElapsedTime(&eventMs, start, stop);
+
+    printf ("hipMemcpyHostToDevice time taken  = %6.3fms\n", eventMs);
+
+    // Record the start event
+    hipEventRecord(start, NULL);
+
+    // Lauching kernel from host
+    hipLaunchKernel(matrixTranspose,
+                    dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
+                    dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
+                    0, 0,
+                    gpuTransposeMatrix , gpuMatrix, WIDTH);
+
+    // Record the stop event
+    hipEventRecord(stop, NULL);
+    hipEventSynchronize(stop);
+    hipEventElapsedTime(&eventMs, start, stop);
+
+    printf ("kernel Execution time             = %6.3fms\n", eventMs);
+
+    // Record the start event
+    hipEventRecord(start, NULL);
+
+    // Memory transfer from device to host
+    hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost);
+
+    // Record the stop event
+    hipEventRecord(stop, NULL);
+    hipEventSynchronize(stop);
+
+    hipEventElapsedTime(&eventMs, start, stop);
+
+    printf ("hipMemcpyDeviceToHost time taken  = %6.3fms\n", eventMs);
+  }
+};
+
+
+int main() {
 
   float* Matrix;
   float* TransposeMatrix;
@@ -79,18 +143,14 @@ int main() {
 
   std::cout << "Device name " << devProp.name << std::endl;
 
-  hipEvent_t start, stop;
-  float eventMs = 1.0f;
   {
       // Show example of how to create a "scoped marker".  
       // The scoped marker records the time spent inside the { scope } of the marker - the begin timestamp is at the
       // beginning of the code scope, and the end is recorded when the SCOPE exits.  This can be viewed in CodeXL
       // timeline relative to other GPU and CPU events.
       // This marker captures the time spent in setup including host allocation, initialization, and device memory allocation.
-      HIP_SCOPED_MARKER("Setup", "App");
+      HIP_SCOPED_MARKER("Setup", "MyGroup");
 
-      hipEventCreate(&start);
-      hipEventCreate(&stop);
 
 
       Matrix = (float*)malloc(NUM * sizeof(float));
@@ -103,91 +163,46 @@ int main() {
       }
 
 
-
       // allocate the memory on the device side
       hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
       hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
   }
 
-  {
-      HIP_SCOPED_MARKER("Loop", "App");
+  runGPU(Matrix, TransposeMatrix, gpuMatrix, gpuTransposeMatrix);
 
-      // Record the start event
-      hipEventRecord(start, NULL);
 
-      // Memory transfer from host to device
-      hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);
-
-      // Record the stop event
-      hipEventRecord(stop, NULL);
-      hipEventSynchronize(stop);
-
-      hipEventElapsedTime(&eventMs, start, stop);
-
-      printf ("hipMemcpyHostToDevice time taken  = %6.3fms\n", eventMs);
-
-      // Record the start event
-      hipEventRecord(start, NULL);
-
-      // Lauching kernel from host
-      hipLaunchKernel(matrixTranspose,
-                      dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
-                      dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
-                      0, 0,
-                      gpuTransposeMatrix , gpuMatrix, WIDTH);
-
-      // Record the stop event
-      hipEventRecord(stop, NULL);
-      hipEventSynchronize(stop);
-      hipEventElapsedTime(&eventMs, start, stop);
-  }
+  // show how to use explicit begin/end markers:
+  // We begin the timed region with HIP_BEGIN_MARKER, passing in the markerName and group:
+  // The region will stop when HIP_END_MARKER is called
+  // This is another way to mark begin/end - as an alternative to scoped markers.
+  HIP_BEGIN_MARKER("Check&TearDown", "MyGroup");
 
   int errors = 0;
-  {
-      HIP_SCOPED_MARKER("Teardown", "App");
 
+  // CPU MatrixTranspose computation
+  matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
 
-      printf ("kernel Execution time             = %6.3fms\n", eventMs);
-
-      // Record the start event
-      hipEventRecord(start, NULL);
-
-      // Memory transfer from device to host
-      hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost);
-
-      // Record the stop event
-      hipEventRecord(stop, NULL);
-      hipEventSynchronize(stop);
-
-      hipEventElapsedTime(&eventMs, start, stop);
-
-      printf ("hipMemcpyDeviceToHost time taken  = %6.3fms\n", eventMs);
-
-      // CPU MatrixTranspose computation
-      matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
-
-      // verify the results
-      double eps = 1.0E-6;
-      for (int i = 0; i < NUM; i++) {
-        if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) {
-          errors++;
-        }
-      }
-      if (errors!=0) {
-        printf("FAILED: %d errors\n",errors);
-      } else {
-        printf ("PASSED!\n");
-      }
-
-      //free the resources on device side
-      hipFree(gpuMatrix);
-      hipFree(gpuTransposeMatrix);
-
-      //free the resources on host side
-      free(Matrix);
-      free(TransposeMatrix);
-      free(cpuTransposeMatrix);
+  // verify the results
+  double eps = 1.0E-6;
+  for (int i = 0; i < NUM; i++) {
+    if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) {
+      errors++;
+    }
   }
+  if (errors!=0) {
+    printf("FAILED: %d errors\n",errors);
+  } else {
+    printf ("PASSED!\n");
+  }
+
+  //free the resources on device side
+  hipFree(gpuMatrix);
+  hipFree(gpuTransposeMatrix);
+
+  //free the resources on host side
+  free(Matrix);
+  free(TransposeMatrix);
+  free(cpuTransposeMatrix);
 
   HIP_END_MARKER();
   

From b5b7663c51013d58fff9334ee21e433d0d18faeb Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Thu, 27 Oct 2016 22:05:52 -0500
Subject: [PATCH 4/8] Add initial hipProfileStart/Stop

And modify sample to show how to use.
Still needs some work to understand interaction with CXL.

Change-Id: I2579824d2dd7863ea23874d34f0dabb3cb305d3e
---
 samples/2_Cookbook/2_Profiler/Makefile        |  7 +++++
 .../2_Cookbook/2_Profiler/MatrixTranspose.cpp | 27 +++++++++++++++++--
 2 files changed, 32 insertions(+), 2 deletions(-)

diff --git a/samples/2_Cookbook/2_Profiler/Makefile b/samples/2_Cookbook/2_Profiler/Makefile
index 4b9a063f38..db2d008182 100644
--- a/samples/2_Cookbook/2_Profiler/Makefile
+++ b/samples/2_Cookbook/2_Profiler/Makefile
@@ -35,6 +35,13 @@ profile: $(EXECUTABLE)
 	$(HIPPROFILER_POST_CMD)
 
 
+# Pass option to control start and stop iterations for profiling - see MatrixTranspose.cpp for implementation:
+# Note we start profiler in --startdisabled mode - no timing collected until app enabled it via hipProfilerStart()
+profile_trigger: $(EXECUTABLE)
+	$(HIPPROFILER) $(PROFILER_OPT)  --startdisabled   $(EXECUTABLE) 3 6
+	$(HIPPROFILER_POST_CMD)
+
+
 run: $(EXECUTABLE)
 	$(EXECUTABLE)
 
diff --git a/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp b/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp
index b6a6b141d2..3747bb4ec5 100644
--- a/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp
+++ b/samples/2_Cookbook/2_Profiler/MatrixTranspose.cpp
@@ -36,6 +36,10 @@ THE SOFTWARE.
 
 #define ITERATIONS 10
 
+// Cmdline parms to control start and stop triggers
+int startTriggerIteration=-1;
+int stopTriggerIteration=-1;
+
 // Device (Kernel) function, it must be void
 // hipLaunchParm provides the execution configuration
 __global__ void matrixTranspose(hipLaunchParm lp,
@@ -74,6 +78,13 @@ void runGPU(float *Matrix, float *TransposeMatrix,
 
   for (int i=0; i<ITERATIONS; i++) {
 
+    if (i==startTriggerIteration) {
+      hipProfilerStart();
+    }
+    if (i==stopTriggerIteration) {
+      hipProfilerStop();
+    }
+
     float eventMs = 0.0f;
 
     hipEvent_t start, stop;
@@ -129,7 +140,16 @@ void runGPU(float *Matrix, float *TransposeMatrix,
 };
 
 
-int main() {
+int main(int argc, char *argv[]) {
+
+  if (argc >= 2) {
+    startTriggerIteration = atoi(argv[1]);
+    printf ("info : will start tracing at iteration:%d\n", startTriggerIteration);
+  } 
+  if (argc >= 3) {
+    stopTriggerIteration = atoi(argv[2]);
+    printf ("info : will stop tracing at iteration:%d\n", stopTriggerIteration);
+  }
 
   float* Matrix;
   float* TransposeMatrix;
@@ -166,6 +186,8 @@ int main() {
       // allocate the memory on the device side
       hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
       hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
+
+      // FYI, the scoped-marker will be destroyed here when the scope exits, and will record its "end" timestamp.
   }
 
   runGPU(Matrix, TransposeMatrix, gpuMatrix, gpuTransposeMatrix);
@@ -204,7 +226,8 @@ int main() {
   free(TransposeMatrix);
   free(cpuTransposeMatrix);
 
-  HIP_END_MARKER();
+  // This ends the last marker started in this thread, in this case "Check&TearDown"
+  HIP_END_MARKER();  
   
   return errors;
 }

From ab73e76987c41f5ad0ab5d0bba732b385df697e8 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Fri, 4 Nov 2016 06:34:07 -0500
Subject: [PATCH 5/8] Print non-peers too

Change-Id: I2a6905edcdf144aa732ae3120c17780477f232ac
---
 samples/1_Utils/hipInfo/hipInfo.cpp | 9 +++++++++
 1 file changed, 9 insertions(+)

diff --git a/samples/1_Utils/hipInfo/hipInfo.cpp b/samples/1_Utils/hipInfo/hipInfo.cpp
index 0403162bd1..42a879e732 100644
--- a/samples/1_Utils/hipInfo/hipInfo.cpp
+++ b/samples/1_Utils/hipInfo/hipInfo.cpp
@@ -133,6 +133,15 @@ void printDeviceProp (int deviceId)
         }
     }
     cout << endl;
+    cout << setw(w1) << "non-peers: ";
+    for (int i=0; i<deviceCnt; i++) {
+        int isPeer;
+        hipDeviceCanAccessPeer(&isPeer, i, deviceId);
+        if (!isPeer) {
+            cout << "device#" << i << " ";
+        }
+    }
+    cout << endl;
 
 
 

From 997dbf0d6a85e3cee7e8f4258cbe72e4714c6bf5 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Sun, 6 Nov 2016 04:07:51 -0600
Subject: [PATCH 6/8] Update gitignore for some common output files

Change-Id: I9cd60f042af4dba07fe0fdbd2ee442936ff8c7bd
---
 samples/0_Intro/hcc_dialects/.gitignore | 5 +++++
 samples/0_Intro/module_api/.gitignore   | 5 +++++
 2 files changed, 10 insertions(+)
 create mode 100644 samples/0_Intro/hcc_dialects/.gitignore
 create mode 100644 samples/0_Intro/module_api/.gitignore

diff --git a/samples/0_Intro/hcc_dialects/.gitignore b/samples/0_Intro/hcc_dialects/.gitignore
new file mode 100644
index 0000000000..bce1cdf193
--- /dev/null
+++ b/samples/0_Intro/hcc_dialects/.gitignore
@@ -0,0 +1,5 @@
+vadd_amp_arrayview
+vadd_hc_am
+vadd_hc_array
+vadd_hc_arrayview
+vadd_hip
diff --git a/samples/0_Intro/module_api/.gitignore b/samples/0_Intro/module_api/.gitignore
new file mode 100644
index 0000000000..c1d81e043f
--- /dev/null
+++ b/samples/0_Intro/module_api/.gitignore
@@ -0,0 +1,5 @@
+runKernel.hip.out
+vcpy_isa.code
+vcpy_isa.hsaco
+vcpy_kernel.co
+vcpy_kernel.code

From aaf1547bff0a29418d8b5f5a77815a5be6ac6bc7 Mon Sep 17 00:00:00 2001
From: Maneesh Gupta <maneesh.gupta@amd.com>
Date: Thu, 10 Nov 2016 11:27:28 +0530
Subject: [PATCH 7/8] hcc_dialects/Makefile: use clamp-config

Change-Id: I86df82f75b75125825e22d0545209a19386d9936
---
 samples/0_Intro/hcc_dialects/Makefile | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/samples/0_Intro/hcc_dialects/Makefile b/samples/0_Intro/hcc_dialects/Makefile
index 3b5ceca7f0..4a514b6691 100644
--- a/samples/0_Intro/hcc_dialects/Makefile
+++ b/samples/0_Intro/hcc_dialects/Makefile
@@ -5,8 +5,8 @@ OPT=-O2
 HCC_CFLAGS= `$(HCC_HOME)/bin/hcc-config --cxxflags` ${OPT}
 HCC_LDFLAGS= `$(HCC_HOME)/bin/hcc-config --ldflags` ${OPT}
 
-CPPAMP_CFLAGS= -std=c++amp -stdlib=libc++ -I$(HCC_HOME)/include
-CPPAMP_LDFLAGS=  -std=c++amp -L$(HCC_HOME)/lib -Wl,--rpath=$(HCC_HOME)/lib -lc++ -lc++abi -ldl -lpthread -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive 
+CPPAMP_CFLAGS= `$(HCC_HOME)/bin/clamp-config --cxxflags`
+CPPAMP_LDFLAGS= `$(HCC_HOME)/bin/clamp-config --ldflags`
 
 HIP_PATH?= $(wildcard /opt/rocm/hip)
 ifeq (,$(HIP_PATH))

From 8829e2626c2b37c87a690e783dd0559971338ffe Mon Sep 17 00:00:00 2001
From: Sandeep Kumar <Sandeep.Kumar@amd.com>
Date: Wed, 9 Nov 2016 12:06:45 +0530
Subject: [PATCH 8/8] Add p2p for cookbook

Change-Id: Id2e77ab31123ef95885d665efe34bc0d4596733a
(cherry picked from commit 6fbd0352713ca36e399b1ed4f17c486207a53875)
---
 samples/2_Cookbook/8_peer2peer/Makefile      |  36 +++
 samples/2_Cookbook/8_peer2peer/peer2peer.cpp | 241 +++++++++++++++++++
 2 files changed, 277 insertions(+)
 create mode 100644 samples/2_Cookbook/8_peer2peer/Makefile
 create mode 100644 samples/2_Cookbook/8_peer2peer/peer2peer.cpp

diff --git a/samples/2_Cookbook/8_peer2peer/Makefile b/samples/2_Cookbook/8_peer2peer/Makefile
new file mode 100644
index 0000000000..a1dad7d1da
--- /dev/null
+++ b/samples/2_Cookbook/8_peer2peer/Makefile
@@ -0,0 +1,36 @@
+HIP_PATH?= $(wildcard /opt/rocm/hip)
+ifeq (,$(HIP_PATH))
+	HIP_PATH=../../..
+endif
+
+HIPCC=$(HIP_PATH)/bin/hipcc
+
+TARGET=hcc
+
+SOURCES = peer2peer.cpp
+OBJECTS = $(SOURCES:.cpp=.o)
+
+EXECUTABLE=./peer2peer
+
+.PHONY: test
+
+
+all: $(EXECUTABLE) test
+
+CXXFLAGS =-g
+CXX=$(HIPCC)
+
+
+$(EXECUTABLE): $(OBJECTS)
+	$(HIPCC) $(OBJECTS) -o $@
+
+
+test: $(EXECUTABLE)
+	$(EXECUTABLE)
+
+
+clean:
+	rm -f $(EXECUTABLE)
+	rm -f $(OBJECTS)
+	rm -f $(HIP_PATH)/src/*.o
+
diff --git a/samples/2_Cookbook/8_peer2peer/peer2peer.cpp b/samples/2_Cookbook/8_peer2peer/peer2peer.cpp
new file mode 100644
index 0000000000..624de56cb0
--- /dev/null
+++ b/samples/2_Cookbook/8_peer2peer/peer2peer.cpp
@@ -0,0 +1,241 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANUMTY OF ANY KIND, EXPRESS OR
+IMPLIED, INUMCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNUMESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANUMY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INUM AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INUM CONUMECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <iostream>
+#include <hip/hip_runtime.h>
+#include <assert.h>
+#define WIDTH     32
+
+#define NUM       (WIDTH*WIDTH)
+
+#define THREADS_PER_BLOCK_X  4
+#define THREADS_PER_BLOCK_Y  4
+#define THREADS_PER_BLOCK_Z  1
+
+using namespace std;
+
+#define KNRM  "\x1B[0m"
+#define KRED  "\x1B[31m"
+
+#define failed(...) \
+    printf ("%serror: ", KRED);\
+    printf (__VA_ARGS__);\
+    printf ("\n");\
+    printf ("error: TEST FAILED\n%s", KNRM );\
+    abort();
+
+#define HIPCHECK(error) \
+{\
+    hipError_t localError = error; \
+    if (localError != hipSuccess) { \
+        printf("%serror: '%s'(%d) from %s at %s:%d%s\n", \
+        KRED, hipGetErrorString(localError), localError,\
+        #error,__FILE__, __LINE__, KNRM); \
+        failed("API returned error code.");\
+    }\
+}
+
+void checkPeer2PeerSupport()
+{
+    int gpuCount;
+    int canAccessPeer;
+    int p2pCapableDeviceCount=0;
+
+    HIPCHECK(hipGetDeviceCount(&gpuCount));
+
+    if (gpuCount < 2)
+        printf("Peer2Peer application requires atleast 2 gpu devices");
+
+    for (int currentGpu=0; currentGpu<gpuCount; currentGpu++)
+    {
+        HIPCHECK(hipSetDevice(currentGpu));
+
+        for (int peerGpu=0; peerGpu<currentGpu; peerGpu++)
+        {
+            if (currentGpu!=peerGpu)
+            {
+                HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu));
+                printf ("currentGpu#%d canAccessPeer: peerGpu#%d=%d\n", currentGpu, peerGpu, canAccessPeer);
+            }
+
+            HIPCHECK(hipSetDevice(peerGpu));
+            HIPCHECK(hipDeviceReset());
+        }
+    HIPCHECK(hipSetDevice(currentGpu));
+    HIPCHECK(hipDeviceReset());
+    }
+}
+
+void enablePeer2Peer(int currentGpu, int peerGpu)
+{
+    int canAccessPeer;
+
+    // Must be on a multi-gpu system:
+    assert (currentGpu != peerGpu);
+
+    HIPCHECK(hipSetDevice(currentGpu));
+    hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu);
+
+    if(canAccessPeer==1){
+        HIPCHECK(hipDeviceEnablePeerAccess(peerGpu, 0));
+    }
+    else
+    printf("peer2peer transfer not possible between the selected gpu devices");
+}
+
+void disablePeer2Peer(int currentGpu, int peerGpu)
+{
+    int canAccessPeer;
+
+    // Must be on a multi-gpu system:
+    assert (currentGpu != peerGpu);
+
+    HIPCHECK(hipSetDevice(currentGpu));
+    hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu);
+	
+    if(canAccessPeer==1){
+        HIPCHECK(hipDeviceDisablePeerAccess(peerGpu));
+    }
+    else
+    printf("peer2peer disable not required");
+}
+
+
+__global__ void matrixTranspose_static_shared(hipLaunchParm lp,
+                                float *out,
+                                float *in,
+                                const int width)
+{
+    __shared__ float sharedMem[WIDTH*WIDTH];
+
+    int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+    int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+
+    sharedMem[y * width + x] = in[x * width + y];
+
+    __syncthreads();
+
+    out[y * width + x] = sharedMem[y * width + x];
+}
+
+__global__ void matrixTranspose_dynamic_shared(hipLaunchParm lp,
+                                float *out,
+                                float *in,
+                                const int width)
+{
+    // declare dynamic shared memory
+    HIP_DYNAMIC_SHARED(float, sharedMem)
+
+    int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+    int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
+
+    sharedMem[y * width + x] = in[x * width + y];
+
+    __syncthreads();
+
+    out[y * width + x] = sharedMem[y * width + x];
+}
+
+int main(){
+
+    checkPeer2PeerSupport();
+    
+    int gpuCount;
+    int currentGpu, peerGpu;
+
+    HIPCHECK(hipGetDeviceCount(&gpuCount));
+
+    currentGpu = 0;
+    peerGpu = (currentGpu + 1);
+
+    printf ("currentGpu=%d peerGpu=%d (Total no. of gpu = %d)\n", currentGpu, peerGpu, gpuCount);
+
+    float *data[2], *TransposeMatrix[2], *gpuTransposeMatrix[2], *randArray;
+
+    int width = WIDTH;
+
+    randArray = (float*)malloc(NUM * sizeof(float));
+
+    for(int i = 0; i < NUM; i++)
+    {
+        randArray[i] = (float)i*1.0f;
+    }
+
+    enablePeer2Peer(currentGpu, peerGpu);
+
+    HIPCHECK(hipSetDevice(currentGpu));
+    TransposeMatrix[0] = (float*)malloc(NUM * sizeof(float));
+    hipMalloc((void**)&gpuTransposeMatrix[0], NUM * sizeof(float));
+    hipMalloc((void**)&data[0], NUM * sizeof(float));
+    hipMemcpy(data[0], randArray, NUM * sizeof(float), hipMemcpyHostToDevice);
+
+    hipLaunchKernel(matrixTranspose_static_shared,
+                    dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
+                    dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
+                    0, 0,
+                    gpuTransposeMatrix[0], data[0], width);
+
+	HIPCHECK(hipSetDevice(peerGpu));
+	TransposeMatrix[1] = (float*)malloc(NUM * sizeof(float));
+	hipMalloc((void**)&gpuTransposeMatrix[1], NUM * sizeof(float));
+	hipMalloc((void**)&data[1], NUM * sizeof(float));
+    hipMemcpy(data[1], gpuTransposeMatrix[0], NUM * sizeof(float), hipMemcpyDeviceToDevice);
+
+    hipLaunchKernel(matrixTranspose_dynamic_shared,
+                    dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
+                    dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
+                    sizeof(float)*WIDTH*WIDTH, 0,
+                    gpuTransposeMatrix[1], data[1], width);
+
+    hipMemcpy(TransposeMatrix[1], gpuTransposeMatrix[1], NUM*sizeof(float), hipMemcpyDeviceToHost);
+
+    hipDeviceSynchronize();
+
+    disablePeer2Peer(currentGpu, peerGpu);
+
+    // verify the results
+    int errors = 0;
+    double eps = 1.0E-6;
+    for (int i = 0; i < NUM; i++) {
+        if (std::abs(randArray[i] - TransposeMatrix[1][i]) > eps ) {
+        printf("%d cpu: %f gpu peered data  %f\n",i,randArray[i],TransposeMatrix[1][i]);
+        errors++;
+        }
+    }
+    if (errors!=0) {
+        printf("FAILED: %d errors\n",errors);
+    } else {
+        printf ("Peer2Peer PASSED!\n");
+    }
+
+    free(randArray);
+    for(int i=0;i<2;i++){
+       hipFree(data[i]);
+       hipFree(gpuTransposeMatrix[i]);
+       free(TransposeMatrix[i]);
+    }
+
+    HIPCHECK(hipSetDevice(peerGpu));
+    HIPCHECK(hipDeviceReset());
+
+    HIPCHECK(hipSetDevice(currentGpu));
+    HIPCHECK(hipDeviceReset());
+
+    return 0;
+}