diff --git a/perftests/perfDispatch/hipPerfDispatchSpeed.cpp b/perftests/dispatch/hipPerfDispatchSpeed.cpp
similarity index 98%
rename from perftests/perfDispatch/hipPerfDispatchSpeed.cpp
rename to perftests/dispatch/hipPerfDispatchSpeed.cpp
index 84ba73c3aa..12999ebc33 100644
--- a/perftests/perfDispatch/hipPerfDispatchSpeed.cpp
+++ b/perftests/dispatch/hipPerfDispatchSpeed.cpp
@@ -7,7 +7,7 @@
 #include "test_common.h"
 
 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp timer.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */
diff --git a/perftests/perfDispatch/hipPerfBufferCopyRectSpeed.cpp b/perftests/memory/hipPerfBufferCopyRectSpeed.cpp
similarity index 98%
rename from perftests/perfDispatch/hipPerfBufferCopyRectSpeed.cpp
rename to perftests/memory/hipPerfBufferCopyRectSpeed.cpp
index 5000904af9..3cb3243e80 100644
--- a/perftests/perfDispatch/hipPerfBufferCopyRectSpeed.cpp
+++ b/perftests/memory/hipPerfBufferCopyRectSpeed.cpp
@@ -7,7 +7,7 @@
 #include "test_common.h"
 
 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp timer.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */
diff --git a/perftests/perfDispatch/hipPerfBufferCopySpeed.cpp b/perftests/memory/hipPerfBufferCopySpeed.cpp
similarity index 98%
rename from perftests/perfDispatch/hipPerfBufferCopySpeed.cpp
rename to perftests/memory/hipPerfBufferCopySpeed.cpp
index 6f284ae7fb..d9a2d443a2 100644
--- a/perftests/perfDispatch/hipPerfBufferCopySpeed.cpp
+++ b/perftests/memory/hipPerfBufferCopySpeed.cpp
@@ -7,7 +7,7 @@
 #include "test_common.h"
 
 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp timer.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD: %t %s ../../src/test_common.cpp ../../src/timer.cpp EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */
diff --git a/perftests/memory/hipPerfDevMemReadSpeed.cpp b/perftests/memory/hipPerfDevMemReadSpeed.cpp
new file mode 100644
index 0000000000..181cd37f24
--- /dev/null
+++ b/perftests/memory/hipPerfDevMemReadSpeed.cpp
@@ -0,0 +1,136 @@
+/*
+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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t
+ * HIT_END
+ */
+
+#include <iostream>
+#include <chrono>
+#include "test_common.h"
+
+using namespace std;
+
+#define arraySize 16
+
+typedef struct d_uint16 {
+  uint data[arraySize];
+} d_uint16;
+
+__global__ void read_kernel(d_uint16 *src, ulong N, uint *dst) {
+
+  size_t idx = (blockIdx.x * blockDim.x + threadIdx.x);
+  size_t stride = blockDim.x * gridDim.x ;
+
+  uint tmp = 0;
+  for (size_t i = idx; i < N; i += stride) {
+    for (size_t j = 0; j < arraySize; j++) {
+      tmp += src[i].data[j];
+    }
+  }
+
+  atomicAdd(dst, tmp);
+}
+
+int main(int argc, char* argv[]) {
+  d_uint16 *dSrc;
+  d_uint16 *hSrc;
+  uint *dDst;
+  uint *hDst;
+  hipStream_t stream;
+  ulong N = 4 * 1024 * 1024;
+  uint nBytes = N * sizeof(d_uint16);
+
+  int nGpu = 0;
+  HIPCHECK(hipGetDeviceCount(&nGpu));
+  if (nGpu < 1) {
+      cout << "info: didn't find any GPU! skipping the test!\n";
+      passed();
+      return 0;
+  }
+
+  static int device = 0;
+  HIPCHECK(hipSetDevice(device));
+  hipDeviceProp_t props;
+  HIPCHECK(hipGetDeviceProperties(&props, device));
+  cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name <<
+      " with " << props.multiProcessorCount << " CUs" << endl;
+
+  const unsigned threadsPerBlock = 64;
+  const unsigned blocks = props.multiProcessorCount * 4;
+
+  uint inputData = 0x1;
+  int nIter = 1000;
+
+  hSrc =  new d_uint16[nBytes];
+  HIPCHECK(hSrc == 0 ? hipErrorOutOfMemory : hipSuccess);
+  hDst =  new uint;
+  hDst[0] = 0;
+  HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess);
+  for (size_t i = 0; i < N; i++) {
+    for (int j = 0; j < arraySize; j++) {
+      hSrc[i].data[j] = inputData;
+    }
+  }
+
+  HIPCHECK(hipMalloc(&dSrc, nBytes));
+  HIPCHECK(hipMalloc(&dDst, sizeof(uint)));
+
+  HIPCHECK(hipStreamCreate(&stream));
+
+  HIPCHECK(hipMemcpy(dSrc, hSrc, nBytes, hipMemcpyHostToDevice));
+  HIPCHECK(hipMemcpy(dDst, hDst, sizeof(uint), hipMemcpyHostToDevice));
+
+  hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dSrc, N, dDst);
+  HIPCHECK(hipMemcpy(hDst, dDst, sizeof(uint), hipMemcpyDeviceToHost));
+  hipDeviceSynchronize();
+
+  if (hDst[0] != (nBytes / sizeof(uint))) {
+    cout << "info: Data validation failed for warm up run!" << endl;
+    cout << "info: expected " << nBytes / sizeof(uint) << " got " << hDst[0] << endl;
+    HIPCHECK(hipErrorUnknown);
+  }
+
+  // measure performance based on host time
+  auto all_start = chrono::steady_clock::now();
+
+  for(int i = 0; i < nIter; i++) {
+    hipLaunchKernelGGL(read_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dSrc, N, dDst);
+  }
+  hipDeviceSynchronize();
+
+  auto all_end = chrono::steady_clock::now();
+  chrono::duration<double> all_kernel_time = all_end - all_start;
+
+  // read speed in GB/s
+  double perf = ((double)nBytes * nIter * (double)(1e-09)) / all_kernel_time.count();
+
+  cout << "info: average read speed of " << perf << " GB/s " << "achieved for memory size of " <<
+      nBytes / (1024 * 1024) << " MB" << endl;
+
+  delete [] hSrc;
+  delete hDst;
+  hipFree(dSrc);
+  hipFree(dDst);
+  HIPCHECK(hipStreamDestroy(stream));
+
+  passed();
+}
diff --git a/perftests/memory/hipPerfDevMemWriteSpeed.cpp b/perftests/memory/hipPerfDevMemWriteSpeed.cpp
new file mode 100644
index 0000000000..4d706cdde9
--- /dev/null
+++ b/perftests/memory/hipPerfDevMemWriteSpeed.cpp
@@ -0,0 +1,126 @@
+/*
+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.
+*/
+
+/* HIT_START
+ * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t
+ * HIT_END
+ */
+
+#include <iostream>
+#include <chrono>
+#include "test_common.h"
+
+using namespace std;
+
+#define arraySize 16
+
+typedef struct d_uint16 {
+  uint data[arraySize];
+} d_uint16;
+
+__global__ void write_kernel(d_uint16 *dst, ulong N, d_uint16 pval) {
+    size_t idx = (blockIdx.x * blockDim.x + threadIdx.x);
+    size_t stride = blockDim.x * gridDim.x;
+    for (size_t i = idx; i < N; i += stride) {
+      dst[i] = pval;
+    }
+};
+
+int main(int argc, char* argv[]) {
+  d_uint16 *dDst;
+  d_uint16 *hDst;
+  hipStream_t stream;
+  ulong N = 4 * 1024 * 1024;
+  uint nBytes = N * sizeof(d_uint16);
+  d_uint16 pval;
+
+  for (int i = 0; i < arraySize; i++) {
+    pval.data[i] = 0xabababab;
+  }
+
+  int nGpu = 0;
+  HIPCHECK(hipGetDeviceCount(&nGpu));
+  if (nGpu < 1) {
+      cout << "info: didn't find any GPU! skipping the test!\n";
+      passed();
+      return 0;
+  }
+
+  static int device = 0;
+  HIPCHECK(hipSetDevice(device));
+  hipDeviceProp_t props;
+  HIPCHECK(hipGetDeviceProperties(&props, device));
+  cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name <<
+      " with " << props.multiProcessorCount << " CUs" << endl;
+
+  size_t threadsPerBlock = 64;
+  size_t blocks = props.multiProcessorCount * 4;
+
+  uint inputData = 0xabababab;
+  int nIter = 1000;
+
+  hDst =  new d_uint16[nBytes];
+  HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess);
+  for (size_t i = 0; i < N; i++) {
+    for (size_t j = 0; j < arraySize; j++) {
+      hDst[i].data[j] = 0;
+    }
+  }
+
+  HIPCHECK(hipMalloc(&dDst, nBytes));
+
+  HIPCHECK(hipStreamCreate(&stream));
+
+  hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval);
+  HIPCHECK(hipMemcpy(hDst, dDst, nBytes , hipMemcpyDeviceToHost));
+  hipDeviceSynchronize();
+
+  for (uint i = 0; i < N; i++) {
+    for (uint j = 0; j < arraySize; j++) {
+      if (hDst[i].data[j] != inputData) {
+        cout << "info: Data validation failed for warm up run! " << endl;
+        cout << "at index i: " << i << " element j: " << j << endl;
+        cout << hex << "expected 0x" << inputData << " but got 0x" << hDst[i].data[j] << endl;
+        HIPCHECK(hipErrorUnknown);
+      }
+    }
+  }
+
+  auto all_start = chrono::steady_clock::now();
+  for(int i = 0; i < nIter; i++) {
+    hipLaunchKernelGGL(write_kernel, dim3(blocks), dim3(threadsPerBlock), 0, stream, dDst, N, pval);
+  }
+  hipDeviceSynchronize();
+  auto all_end = chrono::steady_clock::now();
+  chrono::duration<double> all_kernel_time = all_end - all_start;
+
+  // read speed in GB/s
+  double perf = ((double)nBytes * nIter * (double)(1e-09)) / all_kernel_time.count();
+
+  cout << "info: average write speed of " << perf << " GB/s " << "achieved for memory size of " <<
+      nBytes / (1024 * 1024) << " MB" << endl;
+
+
+  delete [] hDst;
+  hipFree(dDst);
+  HIPCHECK(hipStreamDestroy(stream));
+
+  passed();
+}
diff --git a/perftests/memory/hipHostNumaAlloc.cpp b/perftests/memory/hipPerfHostNumaAlloc.cpp
similarity index 94%
rename from perftests/memory/hipHostNumaAlloc.cpp
rename to perftests/memory/hipPerfHostNumaAlloc.cpp
index 75d6edf0cf..a5e60c8549 100644
--- a/perftests/memory/hipHostNumaAlloc.cpp
+++ b/perftests/memory/hipPerfHostNumaAlloc.cpp
@@ -34,12 +34,13 @@ THE SOFTWARE.
 #include <array>
 #include "hip/hip_runtime.h"
 /* HIT_START
- * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * BUILD_CMD: hipPerfHostNumaAlloc %hc -I%S/../../src %S/%s %S/../../src/test_common.cpp -lnuma -o %T/%t EXCLUDE_HIP_PLATFORM nvcc
  * TEST: %t
  * HIT_END
  */
 
-// To run it correctly, we must not export HIP_VISIBLE_DEVICES
+// To run it correctly, we must not export HIP_VISIBLE_DEVICES.
+// And we must explicitly link libnuma because of numa api move_pages().
 #define NUM_PAGES 4
 char *h = nullptr;
 char *d_h = nullptr;
@@ -127,6 +128,7 @@ bool test(int cpuId, int gpuId, int numaMode, unsigned int hostMallocflags) {
   printf("\n");
 
   HIPCHECK(hipHostFree((void* )h));
+  hipHostUnregister(m);
   free(m);
 
   if (cpuId >= 0 && (numaMode == MPOL_BIND || numaMode == MPOL_PREFERRED)) {
@@ -149,8 +151,7 @@ bool runTest(const int &cpuCount, const int &gpuCount,
 
     for (int i = 0; i < cpuCount; i++) {
       for (int j = 0; j < gpuCount; j++) {
-        if (!test(i, j, mode[m],
-                  hipHostMallocDefault | hipHostMallocNumaUser)) {
+        if (!test(i, j, mode[m], hostMallocflags)) {
           return false;
         }
       }
diff --git a/perftests/memory/hipPerfSharedMemReadSpeed.cpp b/perftests/memory/hipPerfSharedMemReadSpeed.cpp
new file mode 100644
index 0000000000..86adc5f354
--- /dev/null
+++ b/perftests/memory/hipPerfSharedMemReadSpeed.cpp
@@ -0,0 +1,250 @@
+/*
+ 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.
+ */
+
+/* HIT_START
+ * BUILD: %t %s ../../src/test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t
+ * HIT_END
+ */
+
+#include <iostream>
+#include <chrono>
+#include "test_common.h"
+
+using namespace std;
+
+#define sharedMemSize1 2048
+#define sharedMemSize2 256
+
+__global__ void sharedMemReadSpeed1(float *outBuf, ulong N) {
+
+  size_t gid = (blockIdx.x * blockDim.x + threadIdx.x);
+  size_t lid = threadIdx.x;
+  __shared__ float local[sharedMemSize1];
+
+  float val1 = 0;
+  float val2 = 0;
+  float val3 = 0;
+  float val4 = 0;
+
+  for (int i = 0; i < (sharedMemSize1 / 64); i++) {
+    local[lid + i * 64] = lid;
+  }
+
+  __syncthreads();
+
+  val1 += local[lid];
+  val2 += local[lid + 64];
+  val3 += local[lid + 128];
+  val4 += local[lid + 192];
+  val1 += local[lid + 256];
+  val2 += local[lid + 320];
+  val3 += local[lid + 384];
+  val4 += local[lid + 448];
+  val1 += local[lid + 512];
+  val2 += local[lid + 576];
+  val3 += local[lid + 640];
+  val4 += local[lid + 704];
+  val1 += local[lid + 768];
+  val2 += local[lid + 832];
+  val3 += local[lid + 896];
+  val4 += local[lid + 960];
+  val1 += local[lid + 1024];
+  val2 += local[lid + 1088];
+  val3 += local[lid + 1152];
+  val4 += local[lid + 1216];
+  val1 += local[lid + 1280];
+  val2 += local[lid + 1344];
+  val3 += local[lid + 1408];
+  val4 += local[lid + 1472];
+  val1 += local[lid + 1536];
+  val2 += local[lid + 1600];
+  val3 += local[lid + 1664];
+  val4 += local[lid + 1728];
+  val1 += local[lid + 1792];
+  val2 += local[lid + 1856];
+  val3 += local[lid + 1920];
+  val4 += local[lid + 1984];
+
+  if (gid < N) {
+    outBuf[gid] = val1 + val2 + val3 + val4;
+  }
+};
+
+__global__ void sharedMemReadSpeed2(float *outBuf, ulong N) {
+  size_t gid = (blockIdx.x * blockDim.x + threadIdx.x);
+  size_t lid = threadIdx.x;
+  __shared__ float local[sharedMemSize2];
+
+  float val0 = 0.0f;
+  float val1 = 0.0f;
+
+  for (int i = 0; i < (sharedMemSize2 / 64); i++) {
+    local[lid + i * 64] = lid;
+  }
+
+  __syncthreads();
+
+#pragma nounroll
+  for (uint i = 0; i < 32; i++) {
+    val0 += local[8 * i + 0];
+    val1 += local[8 * i + 1];
+    val0 += local[8 * i + 2];
+    val1 += local[8 * i + 3];
+    val0 += local[8 * i + 4];
+    val1 += local[8 * i + 5];
+    val0 += local[8 * i + 6];
+    val1 += local[8 * i + 7];
+  }
+
+  if (gid < N) {
+    outBuf[gid] = val0 + val1;
+  }
+};
+
+int main(int argc, char *argv[]) {
+  float *dDst;
+  float *hDst;
+  hipStream_t stream;
+  constexpr uint numSizes = 4;
+  constexpr uint Sizes[numSizes] = {262144, 1048576, 4194304, 16777216};
+  uint numReads1 = 32;
+  uint numReads2 = 256;
+  uint sharedMemSizeBytes1 = sharedMemSize1 * sizeof(float);
+  uint sharedMemSizeBytes2 = sharedMemSize2 * sizeof(float);
+  int nIter = 1000;
+  const unsigned threadsPerBlock = 64;
+
+  int nGpu = 0;
+  HIPCHECK(hipGetDeviceCount(&nGpu));
+  if (nGpu < 1) {
+    cout << "info: didn't find any GPU! skipping the test!\n";
+    passed();
+    return 0;
+  }
+
+  static int device = 0;
+  HIPCHECK(hipSetDevice(device));
+  hipDeviceProp_t props;
+  HIPCHECK(hipGetDeviceProperties(&props, device));
+  cout << "info: running on bus " << "0x" << props.pciBusID << " " << props.name
+      << " with " << props.multiProcessorCount << " CUs" << endl;
+
+  HIPCHECK(hipStreamCreate(&stream));
+
+  for (int nTest = 0; nTest < numSizes; nTest++) {
+    uint nBytes = Sizes[nTest % numSizes];
+    ulong N = nBytes / sizeof(float);
+    const unsigned blocks = N / threadsPerBlock;
+
+    hDst = new float[nBytes];
+    HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess);
+    memset(hDst, 0, nBytes);
+
+    HIPCHECK(hipMalloc(&dDst, nBytes));
+    HIPCHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice));
+
+    hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks), dim3(threadsPerBlock),
+        0, stream, dDst, N);
+    HIPCHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost));
+    hipDeviceSynchronize();
+
+    int tmp = 0;
+    for (int i = 0; i < N; i++) {
+      if (i % threadsPerBlock == 0) {
+        tmp = 0;
+      }
+      if (hDst[i] != tmp) {
+        cout << "info: Data validation failed for warm up run!" << endl;
+        cout << "info: expected " << tmp << " got " << hDst[i] << endl;
+        HIPCHECK (hipErrorUnknown);
+      }
+      tmp += threadsPerBlock / 2;
+    }
+
+    auto all_start = chrono::steady_clock::now();
+    for (int i = 0; i < nIter; i++) {
+      hipLaunchKernelGGL(sharedMemReadSpeed1, dim3(blocks),
+          dim3(threadsPerBlock), 0, stream, dDst, N);
+    }
+    hipDeviceSynchronize();
+
+    auto all_end = chrono::steady_clock::now();
+    chrono::duration<double> all_kernel_time = all_end - all_start;
+
+    // read speed in GB/s
+    double perf = ((double) blocks * threadsPerBlock
+        * (numReads1 * sizeof(float) + sharedMemSizeBytes1 / 64) * nIter
+        * (double) (1e-09)) / all_kernel_time.count();
+
+    cout << "info: read speed = " << setw(8) << perf << " GB/s for "
+        << sharedMemSizeBytes1 / 1024 << " KB shared memory"
+            " with " << setw(8) << blocks * threadsPerBlock << " threads, "
+        << setw(4) << numReads1 << " reads in sharedMemReadSpeed1 kernel" << endl;
+
+    delete[] hDst;
+    hipFree(dDst);
+  }
+
+
+  for (int nTest = 0; nTest < numSizes; nTest++) {
+    uint nBytes = Sizes[nTest % numSizes];
+    ulong N = nBytes / sizeof(float);
+    const unsigned blocks = N / threadsPerBlock;
+
+    hDst = new float[nBytes];
+    HIPCHECK(hDst == 0 ? hipErrorOutOfMemory : hipSuccess);
+    memset(hDst, 0, nBytes);
+
+    HIPCHECK(hipMalloc(&dDst, nBytes));
+    HIPCHECK(hipMemcpy(dDst, hDst, nBytes, hipMemcpyHostToDevice));
+
+    hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks), dim3(threadsPerBlock),
+        0, stream, dDst, N);
+    HIPCHECK(hipMemcpy(hDst, dDst, nBytes, hipMemcpyDeviceToHost));
+    hipDeviceSynchronize();
+
+    auto all_start = chrono::steady_clock::now();
+    for (int i = 0; i < nIter; i++) {
+      hipLaunchKernelGGL(sharedMemReadSpeed2, dim3(blocks),
+          dim3(threadsPerBlock), 0, stream, dDst, N);
+    }
+    hipDeviceSynchronize();
+
+    auto all_end = chrono::steady_clock::now();
+    chrono::duration<double> all_kernel_time = all_end - all_start;
+
+    // read speed in GB/s
+    double perf = ((double) blocks * threadsPerBlock
+        * (numReads2 * sizeof(float) + sharedMemSizeBytes2 / 64) * nIter
+        * (double) (1e-09)) / all_kernel_time.count();
+
+    cout << "info: read speed = " << setw(8) << perf << " GB/s for "
+        << sharedMemSizeBytes2 / 1024 << " KB shared memory"
+            " with " << setw(8) << blocks * threadsPerBlock << " threads, "
+        << setw(4) << numReads2 << " reads in sharedMemReadSpeed2 kernel" << endl;
+
+    delete[] hDst;
+    hipFree(dDst);
+  }
+
+  HIPCHECK(hipStreamDestroy(stream));
+
+  passed();
+}
diff --git a/perftests/perfDispatch/timer.cpp b/perftests/perfDispatch/timer.cpp
deleted file mode 100644
index ea9c6ea1d9..0000000000
--- a/perftests/perfDispatch/timer.cpp
+++ /dev/null
@@ -1,116 +0,0 @@
-#include "timer.h"
-
-#include <stdlib.h>
-
-#ifdef _WIN32
-#define WIN32_LEAN_AND_MEAN
-#define VC_EXTRALEAN
-#include <windows.h>
-#pragma comment(lib, "user32")
-#endif
-
-#ifdef __linux__
-#include <time.h>
-#define NANOSECONDS_PER_SEC 1000000000
-#endif
-
-CPerfCounter::CPerfCounter() : _clocks(0), _start(0)
-{
-
-#ifdef _WIN32
-
-    QueryPerformanceFrequency((LARGE_INTEGER *)&_freq);
-
-#endif
-
-#ifdef __linux__
-    _freq = NANOSECONDS_PER_SEC;
-#endif
-
-}
-
-CPerfCounter::~CPerfCounter()
-{
-    // EMPTY!
-}
-
-void
-CPerfCounter::Start(void)
-{
-
-#ifdef _WIN32
-
-    if( _start )
-    {
-        MessageBox(NULL, "Bad Perf Counter Start", "Error", MB_OK);
-        exit(0);
-    }
-    QueryPerformanceCounter((LARGE_INTEGER *)&_start);
-
-#endif
-#ifdef __linux__
-
-    struct timespec s;
-    clock_gettime(CLOCK_MONOTONIC, &s);
-    _start = (i64)s.tv_sec * NANOSECONDS_PER_SEC + (i64)s.tv_nsec ;
-
-#endif
-
-}
-
-void
-CPerfCounter::Stop(void)
-{
-    i64 n;
-
-#ifdef _WIN32
-
-    if( !_start )
-    {
-        MessageBox(NULL, "Bad Perf Counter Stop", "Error", MB_OK);
-        exit(0);
-    }
-
-    QueryPerformanceCounter((LARGE_INTEGER *)&n);
-
-#endif
-#ifdef __linux__
-
-    struct timespec s;
-    clock_gettime(CLOCK_MONOTONIC, &s);
-    n = (i64)s.tv_sec * NANOSECONDS_PER_SEC + (i64)s.tv_nsec ;
-
-#endif
-
-    n -= _start;
-    _start = 0;
-    _clocks += n;
-}
-
-void
-CPerfCounter::Reset(void)
-{
-
-#ifdef _WIN32
-    if( _start )
-    {
-        MessageBox(NULL, "Bad Perf Counter Reset", "Error", MB_OK);
-        exit(0);
-    }
-#endif
-    _clocks = 0;
-}
-
-double
-CPerfCounter::GetElapsedTime(void)
-{
-#ifdef _WIN32
-    if( _start ) {
-        MessageBox(NULL, "Trying to get time while still running.", "Error", MB_OK);
-        exit(0);
-    }
-#endif
-
-    return (double)_clocks / (double)_freq;
-
-}
diff --git a/perftests/perfDispatch/timer.h b/perftests/perfDispatch/timer.h
deleted file mode 100644
index 28bfeff74b..0000000000
--- a/perftests/perfDispatch/timer.h
+++ /dev/null
@@ -1,28 +0,0 @@
-#ifndef _TIMER_H_
-#define _TIMER_H_
-
-#ifdef _WIN32
-typedef __int64 i64 ;
-#endif
-#ifdef __linux__
-typedef long long i64;
-#endif
-
-class CPerfCounter {
-
-public:
-    CPerfCounter();
-    ~CPerfCounter();
-    void Start(void);
-    void Stop(void);
-    void Reset(void);
-    double GetElapsedTime(void);
-
-private:
-
-    i64 _freq;
-    i64 _clocks;
-    i64 _start;
-};
-
-#endif // _TIMER_H_