first commit

README.md

@@ -0,0 +1 @@
+# uvm_async_bench

data.zip

@@ -0,0 +1,3 @@
+version https://git-lfs.github.com/spec/v1
+oid sha256:24dcfa78fe0a79ebf4b189d3b2c8449842f15a9f30b7b781c400c96ba525756a
+size 8411674489

env.sh

@@ -0,0 +1 @@
+export UVMAsyncBench_BASE=$(pwd)

workloads/common/Makefile

@@ -0,0 +1,12 @@
+include ./make.config
+
+NVCC = $(CUDA_DIR)/bin/nvcc
+NVCC_FLAGS = -I$(CUDA_DIR)/include -I$(CUPTI_INCLUDE)
+
+all: cpu_timestapms.o cupti_add.o
+
+cpu_timestapms.o: cpu_timestamps.cpp
+	$(NVCC) $(NVCC_FLAGS) -c cpu_timestamps.cpp 
+
+cupti_add.o: cupti_add.cpp
+	$(NVCC) $(NVCC_FLAGS) -c cupti_add.cpp 

workloads/common/cpu_timestamps.cpp

@@ -0,0 +1,38 @@
+#include "cpu_timestamps.h"
+
+void startCPU() {
+  struct timespec tv;
+  if(clock_gettime(CLOCK_REALTIME, &tv))
+    printf("error clock_gettime\n");
+  startCPUTime = (tv.tv_sec * 1.0e9 + tv.tv_nsec);
+}
+
+
+
+void endCPU() {
+  struct timespec tv;
+  if(clock_gettime(CLOCK_REALTIME, &tv))
+    printf("error clock_gettime\n");
+
+  endCPUTime =  (tv.tv_sec * 1.0e9 + tv.tv_nsec);
+  //endCPUTimestamp1 = std::chrono::system_clock::now();
+  printf("CPU_Times,%lu,%lu,%lu\n", startCPUTime, endCPUTime, endCPUTime-startCPUTime);
+  printf("Overlap_Times,%lu,%lu,%lu\n", overlapStartCPUTime, overlapEndCPUTime, overlapEndCPUTime - overlapStartCPUTime);
+}
+
+void overlapStartCPU()
+{
+  struct timespec tv;
+  if (clock_gettime(CLOCK_REALTIME, &tv))
+    printf("error clock_gettime\n");
+  overlapStartCPUTime = (tv.tv_sec * 1.0e9 + tv.tv_nsec);
+}
+
+void overlapEndCPU()
+{
+  struct timespec tv;
+  if (clock_gettime(CLOCK_REALTIME, &tv))
+    printf("error clock_gettime\n");
+
+  overlapEndCPUTime = (tv.tv_sec * 1.0e9 + tv.tv_nsec);
+}

workloads/common/cpu_timestamps.h

@@ -0,0 +1,22 @@
+#ifndef CPU_TIMESTAMP_
+#define CPU_TIMESTAMP_
+
+#include <time.h>
+#include <sys/time.h>
+#include <stdio.h>
+#include <stdint.h>
+#include <error.h>
+
+static uint64_t startCPUTime;
+static uint64_t endCPUTime;
+
+static uint64_t overlapStartCPUTime = 0;
+static uint64_t overlapEndCPUTime = 0;
+
+void startCPU();
+void endCPU();
+
+void overlapStartCPU();
+void overlapEndCPU();
+
+#endif

workloads/common/cupti_add.cpp

@@ -0,0 +1,242 @@
+#include "cupti_add.h"
+
+static const char *
+getMemcpyKindString(CUpti_ActivityMemcpyKind kind)
+{
+    switch (kind)
+    {
+    case CUPTI_ACTIVITY_MEMCPY_KIND_HTOD:
+        return "HtoD";
+    case CUPTI_ACTIVITY_MEMCPY_KIND_DTOH:
+        return "DtoH";
+    case CUPTI_ACTIVITY_MEMCPY_KIND_HTOA:
+        return "HtoA";
+    case CUPTI_ACTIVITY_MEMCPY_KIND_ATOH:
+        return "AtoH";
+    case CUPTI_ACTIVITY_MEMCPY_KIND_ATOA:
+        return "AtoA";
+    case CUPTI_ACTIVITY_MEMCPY_KIND_ATOD:
+        return "AtoD";
+    case CUPTI_ACTIVITY_MEMCPY_KIND_DTOA:
+        return "DtoA";
+    case CUPTI_ACTIVITY_MEMCPY_KIND_DTOD:
+        return "DtoD";
+    case CUPTI_ACTIVITY_MEMCPY_KIND_HTOH:
+        return "HtoH";
+    default:
+        break;
+    }
+
+    return "<unknown>";
+}
+
+static const char *
+getUvmCounterKindString(CUpti_ActivityUnifiedMemoryCounterKind kind)
+{
+    switch (kind)
+    {
+    case CUPTI_ACTIVITY_UNIFIED_MEMORY_COUNTER_KIND_BYTES_TRANSFER_HTOD:
+        return "BYTES_TRANSFER_HTOD";
+    case CUPTI_ACTIVITY_UNIFIED_MEMORY_COUNTER_KIND_BYTES_TRANSFER_DTOH:
+        return "BYTES_TRANSFER_DTOH";
+    case CUPTI_ACTIVITY_UNIFIED_MEMORY_COUNTER_KIND_CPU_PAGE_FAULT_COUNT:
+        return "CPU_PAGE_FAULTS";
+    case CUPTI_ACTIVITY_UNIFIED_MEMORY_COUNTER_KIND_GPU_PAGE_FAULT:
+        return "GPU_PAGE_FAULTS";
+    default:
+        break;
+    }
+    return "<unknown>";
+}
+
+static void
+printActivity(CUpti_Activity *record)
+{
+    switch (record->kind)
+    {
+        case CUPTI_ACTIVITY_KIND_KERNEL:
+        {
+            int status;
+            CUpti_ActivityKernel4 *kernel = (CUpti_ActivityKernel4 *)record;
+            printf("KERNEL %s, %llu, %llu, %llu\n",
+                abi::__cxa_demangle(kernel->name, 0, 0, &status),
+                (unsigned long long)(kernel->start),
+                (unsigned long long)(kernel->end),
+                (unsigned long long)(kernel->end) - (kernel->start));
+            break;
+        }
+        case CUPTI_ACTIVITY_KIND_RUNTIME:
+        {
+            CUpti_ActivityAPI *api = (CUpti_ActivityAPI *)record;
+            const char *callback_name;
+            cuptiGetCallbackName(CUPTI_CB_DOMAIN_RUNTIME_API, api->cbid, &callback_name);
+            // printf("RUNTIME %s (cbid=%u) [ %llu - %llu ] process %u, thread %u, correlation %u\n",
+            //        callback_name, api->cbid,
+            //        (unsigned long long)(api->start - startTimestamp),
+            //        (unsigned long long)(api->end - startTimestamp),
+            //        api->processId, api->threadId, api->correlationId);
+            printf("RUNTIME %s (cbid=%u), %llu,%llu,%llu, process %u, thread %u, correlation %u\n",
+                callback_name, api->cbid,
+                (unsigned long long)(api->start),
+                (unsigned long long)(api->end),
+                (unsigned long long)(api->end - api->start),
+                api->processId, api->threadId, api->correlationId);
+            break;
+        }
+        case CUPTI_ACTIVITY_KIND_MEMCPY:
+        {
+            CUpti_ActivityMemcpy4 *memcpy = (CUpti_ActivityMemcpy4 *)record;
+            printf("MEMCPY %s, size %llu, %llu, %llu, %llu\n",
+                   getMemcpyKindString((CUpti_ActivityMemcpyKind)memcpy->copyKind),
+                   (unsigned long long)memcpy->bytes,
+                   (unsigned long long)(memcpy->start),
+                   (unsigned long long)(memcpy->end),
+                   (unsigned long long)(memcpy->end) - (memcpy->start));
+            break;
+        }
+        case CUPTI_ACTIVITY_KIND_UNIFIED_MEMORY_COUNTER:
+        {
+            CUpti_ActivityUnifiedMemoryCounter2 *uvm = (CUpti_ActivityUnifiedMemoryCounter2 *)record;
+            printf("UVM MEMCPY %s, size %llu, %llu, %llu, %llu \n",
+                   getUvmCounterKindString(uvm->counterKind),
+                   (unsigned long long)uvm->value,
+                   (unsigned long long)(uvm->start),
+                   (unsigned long long)(uvm->end),
+                   (unsigned long long)(uvm->end - uvm->start));
+            break;
+        }
+    }
+}
+
+void CUPTIAPI bufferRequested(uint8_t **buffer, size_t *size, size_t *maxNumRecords)
+{
+    uint8_t *bfr = (uint8_t *)malloc(BUF_SIZE + ALIGN_SIZE);
+    if (bfr == NULL)
+    {
+        printf("Error: out of memory\n");
+        exit(-1);
+    }
+
+    *size = BUF_SIZE;
+    *buffer = ALIGN_BUFFER(bfr, ALIGN_SIZE);
+    *maxNumRecords = 0;
+}
+
+void CUPTIAPI bufferCompleted(CUcontext ctx, uint32_t streamId, uint8_t *buffer, size_t size, size_t validSize)
+{
+    CUptiResult status;
+    CUpti_Activity *record = NULL;
+    if (validSize > 0)
+    {
+        do
+        {
+            status = cuptiActivityGetNextRecord(buffer, validSize, &record);
+            if (status == CUPTI_SUCCESS)
+            {
+                printActivity(record);
+            }
+            else if (status == CUPTI_ERROR_MAX_LIMIT_REACHED)
+                break;
+            else
+            {
+                CUPTI_CALL(status);
+            }
+        } while (1);
+
+        // report any records dropped from the queue
+        size_t dropped;
+        CUPTI_CALL(cuptiActivityGetNumDroppedRecords(ctx, streamId, &dropped));
+        if (dropped != 0)
+        {
+            printf("Dropped %u activity records\n", (unsigned int)dropped);
+        }
+    }
+
+    free(buffer);
+}
+
+// void initTrace() {
+//     return;
+// }
+
+// void finiTrace() {
+//     return;
+// }
+
+
+void initTrace()
+{
+    size_t attrValue = 0, attrValueSize = sizeof(size_t);
+
+    CUpti_ActivityUnifiedMemoryCounterConfig config[2];
+
+    // configure unified memory counters
+    config[0].scope = CUPTI_ACTIVITY_UNIFIED_MEMORY_COUNTER_SCOPE_PROCESS_SINGLE_DEVICE;
+    config[0].kind = CUPTI_ACTIVITY_UNIFIED_MEMORY_COUNTER_KIND_BYTES_TRANSFER_HTOD;
+    config[0].deviceId = 0;
+    config[0].enable = 1;
+
+    config[1].scope = CUPTI_ACTIVITY_UNIFIED_MEMORY_COUNTER_SCOPE_PROCESS_SINGLE_DEVICE;
+    config[1].kind = CUPTI_ACTIVITY_UNIFIED_MEMORY_COUNTER_KIND_BYTES_TRANSFER_DTOH;
+    config[1].deviceId = 0;
+    config[1].enable = 1;
+
+    CUptiResult res = cuptiActivityConfigureUnifiedMemoryCounter(config, 2);
+    if (res == CUPTI_ERROR_UM_PROFILING_NOT_SUPPORTED)
+    {
+        printf("Test is waived, unified memory is not supported on the underlying platform.\n");
+    }
+    else if (res == CUPTI_ERROR_UM_PROFILING_NOT_SUPPORTED_ON_DEVICE)
+    {
+        printf("Test is waived, unified memory is not supported on the device.\n");
+    }
+    else if (res == CUPTI_ERROR_UM_PROFILING_NOT_SUPPORTED_ON_NON_P2P_DEVICES)
+    {
+        printf("Test is waived, unified memory is not supported on the non-P2P multi-gpu setup.\n");
+    }
+    else
+    {
+        CUPTI_CALL(res);
+    }
+
+    CUPTI_CALL(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_KERNEL));
+    CUPTI_CALL(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_RUNTIME));
+    CUPTI_CALL(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_MEMCPY));
+    CUPTI_CALL(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_UNIFIED_MEMORY_COUNTER));
+    // CUPTI_CALL(cuptiActivityEnable(CUPTI_ACTIVITY_UNIFIED_MEMORY_COUNTER_KIND_COUNT));
+
+    // CUPTI_ACTIVITY_UNIFIED_MEMORY_COUNTER_KIND_GPU_PAGE_FAULT 
+    // CUPTI_ACTIVITY_UNIFIED_MEMORY_COUNTER_KIND_CPU_PAGE_FAULT_COUNT
+
+    // Register callbacks for buffer requests and for buffers completed by CUPTI.
+    CUPTI_CALL(cuptiActivityRegisterCallbacks(bufferRequested, bufferCompleted));
+
+    // Optionally get and set activity attributes.
+    // Attributes can be set by the CUPTI client to change behavior of the activity API.
+    // Some attributes require to be set before any CUDA context is created to be effective,
+    // e.g. to be applied to all device buffer allocations (see documentation).
+    CUPTI_CALL(cuptiActivityGetAttribute(CUPTI_ACTIVITY_ATTR_DEVICE_BUFFER_SIZE, &attrValueSize, &attrValue));
+    printf("%s = %llu B\n", "CUPTI_ACTIVITY_ATTR_DEVICE_BUFFER_SIZE", (long long unsigned)attrValue);
+    attrValue *= 2;
+    CUPTI_CALL(cuptiActivitySetAttribute(CUPTI_ACTIVITY_ATTR_DEVICE_BUFFER_SIZE, &attrValueSize, &attrValue));
+
+    CUPTI_CALL(cuptiActivityGetAttribute(CUPTI_ACTIVITY_ATTR_DEVICE_BUFFER_POOL_LIMIT, &attrValueSize, &attrValue));
+    printf("%s = %llu\n", "CUPTI_ACTIVITY_ATTR_DEVICE_BUFFER_POOL_LIMIT", (long long unsigned)attrValue);
+    attrValue *= 2;
+    CUPTI_CALL(cuptiActivitySetAttribute(CUPTI_ACTIVITY_ATTR_DEVICE_BUFFER_POOL_LIMIT, &attrValueSize, &attrValue));
+
+    CUPTI_CALL(cuptiGetTimestamp(&startTimestamp));
+}
+
+void finiTrace()
+{
+    // Force flush any remaining activity buffers before termination of the application
+    CUPTI_CALL(cuptiActivityFlushAll(1));
+}
+
+void GPU_argv_init() {
+  cudaDeviceProp deviceProp;
+  cudaGetDeviceProperties(&deviceProp, GPU_DEVICE);
+  printf("setting device %d with name %s\n", GPU_DEVICE, deviceProp.name);
+  cudaSetDevice(GPU_DEVICE);
+}

workloads/common/cupti_add.h

@@ -0,0 +1,56 @@
+#include <cupti.h>
+#include <stdio.h>
+#include <cxxabi.h>
+#include <time.h>
+
+#define GPU_DEVICE 7
+
+#define BUF_SIZE (32 * 1024)
+#define ALIGN_SIZE (8)
+#define ALIGN_BUFFER(buffer, align)                                            \
+  (((uintptr_t) (buffer) & ((align)-1)) ? ((buffer) + (align) - ((uintptr_t) (buffer) & ((align)-1))) : (buffer))
+
+static uint64_t startTimestamp;
+// Timestamp at trace initialization time. Used to normalized other
+// timestamps
+
+#define CUPTI_CALL(call)                                                \
+  do {                                                                  \
+    CUptiResult _status = call;                                         \
+    if (_status != CUPTI_SUCCESS) {                                     \
+      const char *errstr;                                               \
+      cuptiGetResultString(_status, &errstr);                           \
+      fprintf(stderr, "%s:%d: error: function %s failed with error %s.\n", \
+              __FILE__, __LINE__, #call, errstr);                       \
+      if(_status == CUPTI_ERROR_LEGACY_PROFILER_NOT_SUPPORTED)          \
+          exit(0);                                                      \
+      else                                                              \
+          exit(-1);                                                     \
+    }                                                                   \
+  } while (0)
+
+
+extern inline __attribute__((always_inline)) unsigned long rdtsc() {
+  unsigned long a, d;
+
+  __asm__ volatile("rdtsc" : "=a"(a), "=d"(d));
+
+  return (a | (d << 32));
+}
+
+extern inline __attribute__((always_inline)) unsigned long rdtsp() {
+  struct timespec tms;
+  if (clock_gettime(CLOCK_REALTIME, &tms)) {
+    return -1;
+  }
+  unsigned long ns = tms.tv_sec * 1000000000;
+  ns += tms.tv_nsec;
+  return ns;
+}
+
+void initTrace();
+void finiTrace();
+void GPU_argv_init();
+void CUPTIAPI bufferRequested(uint8_t **buffer, size_t *size, size_t *maxNumRecords);
+void CUPTIAPI bufferCompleted(CUcontext ctx, uint32_t streamId, uint8_t *buffer, size_t size, size_t validSize);
+static void printActivity(CUpti_Activity *record);

workloads/common/make.config

@@ -0,0 +1,10 @@
+CUDA_DIR = /apps/cuda-11.4
+
+COMPUTE = compute_80 #61
+SM_ARCH = sm_80 #61
+
+CUDA_LIB_DIR := $(CUDA_DIR)/lib64
+CUPTI_LIB_DIR := $(CUDA_DIR)/extras/CUPTI/lib64/
+CUPTI_INCLUDE := $(CUDA_DIR)/extras/CUPTI/include/
+
+CUPTI_ADD_COMMON = $(UVMAsyncBench_BASE)/workloads/common/

workloads/micro/async/2DCONV/2DConvolution.cu

@@ -0,0 +1,392 @@
+/**
+ * 2DConvolution.cu: This file is part of the PolyBench/GPU 1.0 test suite.
+ *
+ *
+ * Contact: Scott Grauer-Gray <sgrauerg@gmail.com>
+ * Louis-Noel Pouchet <pouchet@cse.ohio-state.edu>
+ * Web address: http://www.cse.ohio-state.edu/~pouchet/software/polybench/GPU
+ */
+
+#include <unistd.h>
+#include <stdio.h>
+#include <time.h>
+#include <sys/time.h>
+#include <stdlib.h>
+#include <stdarg.h>
+#include <string.h>
+#include <cuda.h>
+
+#include "../../../common/cupti_add.h"
+#include "../../../common/cpu_timestamps.h"
+
+#include <cooperative_groups.h>
+#include <cooperative_groups/memcpy_async.h>
+
+using namespace nvcuda::experimental;
+
+#define PREFETCH_COUNT 2
+
+#define SMALL_FLOAT_VAL 0.00000001f
+
+double rtclock()
+{
+	struct timezone Tzp;
+	struct timeval Tp;
+	uint64_t stat;
+	stat = gettimeofday(&Tp, &Tzp);
+	if (stat != 0)
+		printf("Error return from gettimeofday: %d", stat);
+	return (Tp.tv_sec + Tp.tv_usec * 1.0e-6);
+}
+
+float absVal(float a)
+{
+	if (a < 0)
+	{
+		return (a * -1);
+	}
+	else
+	{
+		return a;
+	}
+}
+
+float percentDiff(double val1, double val2)
+{
+	if ((absVal(val1) < 0.01) && (absVal(val2) < 0.01))
+	{
+		return 0.0f;
+	}
+
+	else
+	{
+		return 100.0f * (absVal(absVal(val1 - val2) / absVal(val1 + SMALL_FLOAT_VAL)));
+	}
+}
+
+//define the error threshold for the results "not matching"
+#define PERCENT_DIFF_ERROR_THRESHOLD 0.05
+
+/* Problem size */
+#define SIZE 4096
+#define NBLOCKS 32
+#define BATCH_SIZE 4
+
+uint64_t NI;
+uint64_t NJ;
+uint64_t nblocks;
+
+
+/* Thread block dimensions */
+#define KERNEL 3
+#define DIM_THREAD_BLOCK 8 
+
+/* Can switch DATA_TYPE between float and double */
+typedef float DATA_TYPE;
+
+void conv2D(DATA_TYPE* A, DATA_TYPE* B)
+{
+	uint64_t i, j;
+	DATA_TYPE c11, c12, c13, c21, c22, c23, c31, c32, c33;
+
+	c11 = +0.2;  c21 = +0.5;  c31 = -0.8;
+	c12 = -0.3;  c22 = +0.6;  c32 = -0.9;
+	c13 = +0.4;  c23 = +0.7;  c33 = +0.10;
+
+	for (i = 1; i < NI - 1; ++i) // 0
+	{
+		for (j = 1; j < NJ - 1; ++j) // 1
+		{
+			B[i*NJ + j] = c11 * A[(i - 1)*NJ + (j - 1)]  +  c12 * A[(i + 0)*NJ + (j - 1)]  +  c13 * A[(i + 1)*NJ + (j - 1)]
+				+ c21 * A[(i - 1)*NJ + (j + 0)]  +  c22 * A[(i + 0)*NJ + (j + 0)]  +  c23 * A[(i + 1)*NJ + (j + 0)] 
+				+ c31 * A[(i - 1)*NJ + (j + 1)]  +  c32 * A[(i + 0)*NJ + (j + 1)]  +  c33 * A[(i + 1)*NJ + (j + 1)];
+		}
+	}
+}
+
+
+void initGPU(DATA_TYPE* A_gpu)
+{
+	uint64_t i, j;
+
+	for (i = 0; i < NI; ++i) {
+		for (j = 0; j < NJ; ++j) {
+			A_gpu[i * NJ + j] = ((DATA_TYPE)i * j) / NI;
+		}
+    }
+}
+
+void initCPU(DATA_TYPE* A)
+{
+	uint64_t i, j;
+
+	for (i = 0; i < NI; ++i) {
+		for (j = 0; j < NJ; ++j) {
+			A[i * NJ + j] = ((DATA_TYPE)i * j) / NI;
+		}
+    }
+}
+
+
+void compareResults(DATA_TYPE* B, DATA_TYPE* B_outputFromGpu)
+{
+	uint64_t i, j, fail;
+	fail = 0;
+	
+	// Compare a and b
+	for (i=1; i < (NI-1); i++) 
+	{
+		for (j=1; j < (NJ-1); j++) 
+		{
+			if (percentDiff(B[i*NJ + j], B_outputFromGpu[i*NJ + j]) > PERCENT_DIFF_ERROR_THRESHOLD) 
+			{
+				printf("%d, %d, CPU is %f, GPU is %f.\n", i, j, B[i * NJ + j], B_outputFromGpu[i * NJ + j]);
+				fail++;
+			}
+		}
+	}
+	
+	// Print results
+	printf("Non-Matching CPU-GPU Outputs Beyond Error Threshold of %4.2f Percent: %d\n", PERCENT_DIFF_ERROR_THRESHOLD, fail);
+
+}
+
+__global__ void Convolution2D_kernel(DATA_TYPE *A, DATA_TYPE *B, uint64_t NI, uint64_t NJ, uint64_t block_size)
+{
+	cooperative_groups::thread_block block = cooperative_groups::this_thread_block();
+	pipeline pipe;
+	uint64_t tile_dim_x = (NJ + DIM_THREAD_BLOCK - 1) / (DIM_THREAD_BLOCK * BATCH_SIZE);
+
+	__shared__ DATA_TYPE tmp_A[PREFETCH_COUNT][DIM_THREAD_BLOCK * BATCH_SIZE + KERNEL - 1][DIM_THREAD_BLOCK * BATCH_SIZE + KERNEL - 1];
+	__shared__ DATA_TYPE tmp_B[DIM_THREAD_BLOCK * BATCH_SIZE][DIM_THREAD_BLOCK * BATCH_SIZE];
+
+	uint64_t total_tiles = tile_dim_x * tile_dim_x;
+
+	uint64_t tiles_this_block_x = (block_size / (DIM_THREAD_BLOCK * BATCH_SIZE));
+	uint64_t tiles_this_block = tiles_this_block_x * tiles_this_block_x;
+
+	// DATA_TYPE c11, c12, c13, c21, c22, c23, c31, c32, c33;
+
+	// c11 = +0.2;  c21 = +0.5;  c31 = -0.8;
+	// c12 = -0.3;  c22 = +0.6;  c32 = -0.9;
+	// c13 = +0.4;  c23 = +0.7;  c33 = +0.10;
+
+	DATA_TYPE c[KERNEL][KERNEL];
+
+	c[0][0] = +0.2;
+	c[1][0] = +0.5;
+	c[2][0] = -0.8;
+	c[0][1] = -0.3;
+	c[1][1] = +0.6;
+	c[2][1] = -0.9;
+	c[0][2] = +0.4;
+	c[1][2] = +0.7;
+	c[2][2] = +0.10;
+
+	uint64_t base_tile = (blockIdx.y * gridDim.x + blockIdx.x) * tiles_this_block;
+	uint64_t fetch = base_tile;
+	uint64_t end_tile = fetch + tiles_this_block;
+
+	for (uint64_t compute = fetch; compute < end_tile; compute++)
+	{
+		for (; fetch < end_tile && fetch < compute + PREFETCH_COUNT; fetch++)
+		{
+			// block id
+			uint64_t offset = fetch - base_tile;
+			uint64_t block_id = fetch / tiles_this_block;
+			uint64_t bx = block_id % gridDim.x * tiles_this_block_x + offset % tiles_this_block_x;
+			uint64_t by = block_id / gridDim.x * tiles_this_block_x + offset / tiles_this_block_x;
+
+			uint64_t batch_size = DIM_THREAD_BLOCK * BATCH_SIZE;
+
+			// thread id
+			uint64_t tx = threadIdx.x;
+			uint64_t ty = threadIdx.y;
+
+			uint64_t index_A_y = DIM_THREAD_BLOCK * BATCH_SIZE * by + BATCH_SIZE * ty;
+			uint64_t index_A_x = DIM_THREAD_BLOCK * BATCH_SIZE * bx + BATCH_SIZE * tx;
+
+			uint64_t index_A_y_start = DIM_THREAD_BLOCK * BATCH_SIZE * by;
+			uint64_t index_A_x_start = DIM_THREAD_BLOCK * BATCH_SIZE * bx;
+
+			uint64_t index_A_y_bound = DIM_THREAD_BLOCK * BATCH_SIZE * by + BATCH_SIZE * DIM_THREAD_BLOCK;
+			uint64_t index_A_x_bound = DIM_THREAD_BLOCK * BATCH_SIZE * bx + BATCH_SIZE * DIM_THREAD_BLOCK;
+
+			// fetch A
+			for (uint64_t i = 0; i < BATCH_SIZE; i++)
+			{
+				for (uint64_t j = 0; j < BATCH_SIZE; j++)
+				{
+					if ((index_A_y + i) < NI && (index_A_x + j) < NJ)
+					{
+						memcpy_async(tmp_A[fetch % PREFETCH_COUNT][ty * BATCH_SIZE + i][tx * BATCH_SIZE + j], A[(index_A_y + i) * NJ + index_A_x + j], pipe);
+						tmp_B[ty * BATCH_SIZE + i][tx * BATCH_SIZE + j] = 0;
+					}
+				}
+			}
+
+			// fetch A -- padding
+			for (uint64_t i = 0; i < KERNEL - 1; i++)
+			{
+				for (uint64_t j = 0; j < BATCH_SIZE * DIM_THREAD_BLOCK + KERNEL - 1; j++)
+				{
+					if ((index_A_y_bound + i) < NI && (index_A_x_start + j) < NJ)
+					{
+						memcpy_async(tmp_A[fetch % PREFETCH_COUNT][DIM_THREAD_BLOCK * BATCH_SIZE + i][j], A[(index_A_y_bound + i) * NJ + index_A_x_start + j], pipe);
+					}
+				}
+			}
+
+			// fetch A -- padding
+			for (uint64_t i = 0; i < BATCH_SIZE * DIM_THREAD_BLOCK + KERNEL - 1; i++)
+			{
+				for (uint64_t j = 0; j < KERNEL - 1; j++)
+				{
+					if ((index_A_y_start + i) < NI && (index_A_x_bound + j) < NJ)
+					{
+						memcpy_async(tmp_A[fetch % PREFETCH_COUNT][i][DIM_THREAD_BLOCK * BATCH_SIZE + j], A[(index_A_y_start + i) * NJ + index_A_x_bound + j], pipe);
+					}
+				}
+			}
+			pipe.commit();
+		}
+		if (fetch == end_tile)
+		{
+			for (uint64_t i = 0; i < PREFETCH_COUNT - 1; ++i)
+			{
+				pipe.commit();
+			}
+			++fetch;
+		}
+		pipe.wait_prior<PREFETCH_COUNT - 1>();
+		block.sync();
+
+		// block id
+		uint64_t offset = compute - base_tile;
+		uint64_t block_id = compute / tiles_this_block;
+		uint64_t bx = block_id % gridDim.x * tiles_this_block_x + offset % tiles_this_block_x;
+		uint64_t by = block_id / gridDim.x * tiles_this_block_x + offset / tiles_this_block_x;
+
+		// thread id
+		uint64_t tx = threadIdx.x;
+		uint64_t ty = threadIdx.y;
+
+		uint64_t index_B_y = DIM_THREAD_BLOCK * BATCH_SIZE * by + BATCH_SIZE * ty + 1;
+		uint64_t index_B_x = DIM_THREAD_BLOCK * BATCH_SIZE * bx + BATCH_SIZE * tx + 1;
+
+		// Computation
+		for (uint64_t i = 0; i < BATCH_SIZE; i++)
+		{
+			for (uint64_t j = 0; j < BATCH_SIZE; j++)
+			{
+				tmp_B[ty * BATCH_SIZE + i][tx * BATCH_SIZE + j] = 0;
+			}
+		}
+		block.sync();
+
+		for (uint64_t i = 0; i < BATCH_SIZE; i++)
+		{
+			for (uint64_t j = 0; j < BATCH_SIZE; j++)
+			{
+				for (uint64_t m = 0; m < KERNEL; m++)
+				{
+					for (uint64_t n = 0; n < KERNEL; n++)
+					{
+						tmp_B[ty * BATCH_SIZE + i][tx * BATCH_SIZE + j] += tmp_A[compute % PREFETCH_COUNT][ty * BATCH_SIZE + i + m][tx * BATCH_SIZE + j + n] * c[n][m];
+					}
+				}
+			}
+		}
+		block.sync();
+
+		// Store B
+		for (uint64_t i = 0; i < BATCH_SIZE; i++)
+		{
+			for (uint64_t j = 0; j < BATCH_SIZE; j++)
+			{
+				if ((index_B_y + i) < NI && (index_B_x + j) < NJ)
+				{
+					B[(index_B_y + i) * NJ + index_B_x + j] = tmp_B[ty * BATCH_SIZE + i][tx * BATCH_SIZE + j];
+				}
+			}
+		}
+		block.sync();
+	}
+}
+
+void convolution2DCuda(DATA_TYPE *A, DATA_TYPE *B, DATA_TYPE *A_gpu, DATA_TYPE *B_gpu)
+{
+	double t_start, t_end;
+
+	uint64_t output_width = NI - KERNEL + 1;
+	uint64_t output_height = NJ - KERNEL + 1;
+
+	dim3 block(DIM_THREAD_BLOCK, DIM_THREAD_BLOCK);
+	dim3 grid(nblocks, nblocks);
+
+	uint64_t block_size = (NJ + (nblocks - 1)) / nblocks;
+
+	// t_start = rtclock();
+	
+	cudaMemcpy(A_gpu, A, sizeof(DATA_TYPE) * NI * NJ, cudaMemcpyHostToDevice);
+	Convolution2D_kernel<<<grid,block>>>(A_gpu, B_gpu, NI, NJ, block_size);
+	cudaDeviceSynchronize();
+	cudaMemcpy(B, B_gpu, sizeof(DATA_TYPE) * NI * NJ, cudaMemcpyDeviceToHost);
+	
+	// t_end = rtclock();
+	// fprintf(stdout, "GPU Runtime: %0.6lfs\n", t_end - t_start);//);
+}
+
+int main(int argc, char *argv[])
+{
+	uint64_t start_tsc = rdtsc();
+	uint64_t start_tsp = rdtsp();
+	printf("start_tsc %lu start_tsp %lu\n", start_tsc, start_tsp);
+	if (argc >= 4) {
+		NI = atoll(argv[1]);
+		NJ = atoll(argv[2]);
+		nblocks = atoi(argv[3]);
+	} else {
+		NI = SIZE;
+		NJ = SIZE;
+		nblocks = NBLOCKS;
+	}
+	double t_start, t_end;
+
+	DATA_TYPE* A;
+	DATA_TYPE* B;
+	DATA_TYPE *B_ref;
+	DATA_TYPE *A_gpu;
+	DATA_TYPE *B_gpu;
+
+	A = (DATA_TYPE*)malloc(NI*NJ*sizeof(DATA_TYPE));
+	B = (DATA_TYPE*)malloc(NI*NJ*sizeof(DATA_TYPE));
+	B_ref = (DATA_TYPE *)malloc(NI * NJ * sizeof(DATA_TYPE));
+	initCPU(A);
+	GPU_argv_init();
+
+	initTrace();
+	startCPU();
+
+	cudaMalloc(&A_gpu, sizeof(DATA_TYPE) * NI * NJ);
+	cudaMalloc(&B_gpu, sizeof(DATA_TYPE) * NI * NJ);
+	// B_outputFromGpu = (DATA_TYPE*)malloc(NI*NJ*sizeof(DATA_TYPE));
+
+	convolution2DCuda(A, B, A_gpu, B_gpu);
+
+	cudaFree(A_gpu);
+	cudaFree(B_gpu);
+	endCPU();
+	finiTrace();
+	
+	// t_start = rtclock();
+	// conv2D(A, B_ref);
+	// t_end = rtclock();
+	// fprintf(stdout, "CPU Runtime: %0.6lfs\n", t_end - t_start);
+
+	// compareResults(B, B_ref);
+	free(A);
+	free(B);
+	
+	return 0;
+}

workloads/micro/async/2DCONV/Makefile

@@ -0,0 +1,12 @@
+include ../../../common/make.config
+
+NVCCCFLAGS = -I$(CUPTI_INCLUDE) -L$(CUPTI_LIB_DIR) -std=c++11 -lcuda -lcupti -arch=sm_80 -O3
+NVCC = $(CUDA_DIR)/bin/nvcc
+
+EXECUTABLE := 2DConvolution
+CUFILES := 2DConvolution.cu $(CUPTI_ADD_COMMON)/cupti_add.cpp $(CUPTI_ADD_COMMON)/cpu_timestamps.cpp
+
+all:
+	$(NVCC) ${NVCCCFLAGS} ${CUFILES} ${DEF} -o ${EXECUTABLE}
+clean:
+	rm -f *.o 2DConvolution

workloads/micro/async/2DCONV/run.sh

@@ -0,0 +1,2 @@
+# ./2DConvolution 16384 16384 32
+./2DConvolution 32768 32768 32

workloads/micro/async/2DCONV/run_large.sh

@@ -0,0 +1 @@
+./2DConvolution 8192 8192 32

workloads/micro/async/2DCONV/run_medium.sh

@@ -0,0 +1 @@
+./2DConvolution 4096 4096 32

workloads/micro/async/2DCONV/run_mega.sh

@@ -0,0 +1 @@
+./2DConvolution 65536 65536 32

workloads/micro/async/2DCONV/run_small.sh

@@ -0,0 +1 @@
+./2DConvolution 1024 1024 8

workloads/micro/async/2DCONV/run_super.sh

@@ -0,0 +1 @@
+./2DConvolution 32768 32768 32

workloads/micro/async/2DCONV/run_tiny.sh

@@ -0,0 +1 @@
+./2DConvolution 512 512 4

workloads/micro/async/3DCONV/3DConvolution.cu

@@ -0,0 +1,449 @@
+/**
+ * 3DConvolution.cu: This file is part of the PolyBench/GPU 1.0 test suite.
+ *
+ *
+ * Contact: Scott Grauer-Gray <sgrauerg@gmail.com>
+ * Louis-Noel Pouchet <pouchet@cse.ohio-state.edu>
+ * Web address: http://www.cse.ohio-state.edu/~pouchet/software/polybench/GPU
+ */
+
+#include <unistd.h>
+#include <stdio.h>
+#include <time.h>
+#include <sys/time.h>
+#include <stdlib.h>
+#include <stdarg.h>
+#include <string.h>
+#include <cuda.h>
+
+#include "../../../common/cupti_add.h"
+#include "../../../common/cpu_timestamps.h"
+
+#include <cooperative_groups.h>
+#include <cooperative_groups/memcpy_async.h>
+
+using namespace nvcuda::experimental;
+
+#define PREFETCH_COUNT 2
+
+#define SMALL_FLOAT_VAL 0.00000001f
+
+double rtclock()
+{
+	struct timezone Tzp;
+	struct timeval Tp;
+	uint64_t stat;
+	stat = gettimeofday(&Tp, &Tzp);
+	if (stat != 0)
+		printf("Error return from gettimeofday: %d", stat);
+	return (Tp.tv_sec + Tp.tv_usec * 1.0e-6);
+}
+
+float absVal(float a)
+{
+	if (a < 0)
+	{
+		return (a * -1);
+	}
+	else
+	{
+		return a;
+	}
+}
+
+float percentDiff(double val1, double val2)
+{
+	if ((absVal(val1) < 0.01) && (absVal(val2) < 0.01))
+	{
+		return 0.0f;
+	}
+
+	else
+	{
+		return 100.0f * (absVal(absVal(val1 - val2) / absVal(val1 + SMALL_FLOAT_VAL)));
+	}
+}
+
+// define the error threshold for the results "not matching"
+#define PERCENT_DIFF_ERROR_THRESHOLD 0.05
+
+/* Problem size */
+#define SIZE 4096
+#define NBLOCKS 2
+#define BATCH_SIZE 3
+
+uint64_t NI;
+uint64_t NJ;
+uint64_t NK;
+uint64_t nblocks;
+
+/* Thread block dimensions */
+#define DIM_THREAD_BLOCK 4
+
+#define KERNEL 3
+
+/* Can switch DATA_TYPE between float and double */
+typedef float DATA_TYPE;
+
+
+
+void conv3D(DATA_TYPE* A, DATA_TYPE* B)
+{
+	uint64_t i, j, k;
+	DATA_TYPE c11, c12, c13, c21, c22, c23, c31, c32, c33;
+
+	c11 = +2;  c21 = +5;  c31 = -8;
+	c12 = -3;  c22 = +6;  c32 = -9;
+	c13 = +4;  c23 = +7;  c33 = +10;
+
+	for (i = 1; i < NI - 1; ++i) // 0
+	{
+		for (j = 1; j < NJ - 1; ++j) // 1
+		{
+			for (k = 1; k < NK -1; ++k) // 2
+			{
+				B[i*(NK * NJ) + j*NK + k] = c11 * A[(i - 1)*(NK * NJ) + (j - 1)*NK + (k - 1)]  +  c13 * A[(i + 1)*(NK * NJ) + (j - 1)*NK + (k - 1)]
+					     +   c21 * A[(i - 1)*(NK * NJ) + (j - 1)*NK + (k - 1)]  +  c23 * A[(i + 1)*(NK * NJ) + (j - 1)*NK + (k - 1)]
+					     +   c31 * A[(i - 1)*(NK * NJ) + (j - 1)*NK + (k - 1)]  +  c33 * A[(i + 1)*(NK * NJ) + (j - 1)*NK + (k - 1)]
+					     +   c12 * A[(i + 0)*(NK * NJ) + (j - 1)*NK + (k + 0)]  +  c22 * A[(i + 0)*(NK * NJ) + (j + 0)*NK + (k + 0)]   
+					     +   c32 * A[(i + 0)*(NK * NJ) + (j + 1)*NK + (k + 0)]  +  c11 * A[(i - 1)*(NK * NJ) + (j - 1)*NK + (k + 1)]  
+					     +   c13 * A[(i + 1)*(NK * NJ) + (j - 1)*NK + (k + 1)]  +  c21 * A[(i - 1)*(NK * NJ) + (j + 0)*NK + (k + 1)]  
+					     +   c23 * A[(i + 1)*(NK * NJ) + (j + 0)*NK + (k + 1)]  +  c31 * A[(i - 1)*(NK * NJ) + (j + 1)*NK + (k + 1)]  
+					     +   c33 * A[(i + 1)*(NK * NJ) + (j + 1)*NK + (k + 1)];
+			}
+		}
+	}
+}
+
+void initGPU(DATA_TYPE *A_gpu)
+{
+	uint64_t i, j, k;
+
+	for (i = 0; i < NI; ++i)
+	{
+		for (j = 0; j < NJ; ++j)
+		{
+			for (k = 0; k < NK; ++k)
+			{
+				A_gpu[i * (NK * NJ) + j * NK + k] = i % 12 + 2 * (j % 7) + 3 * (k % 13);
+			}
+		}
+	}
+}
+
+void initCPU(DATA_TYPE *A)
+{
+	uint64_t i, j, k;
+
+	for (i = 0; i < NI; ++i)
+    {
+		for (j = 0; j < NJ; ++j)
+		{
+			for (k = 0; k < NK; ++k)
+			{
+				A[i*(NK * NJ) + j*NK + k] = i % 12 + 2 * (j % 7) + 3 * (k % 13);
+			}
+		}
+	}
+}
+
+
+void compareResults(DATA_TYPE* B, DATA_TYPE* B_outputFromGpu)
+{
+	uint64_t i, j, k, fail;
+	fail = 0;
+	
+	// Compare result from cpu and gpu...
+	for (i = 1; i < NI - 1; ++i) // 0
+	{
+		for (j = 1; j < NJ - 1; ++j) // 1
+		{
+			for (k = 1; k < NK - 1; ++k) // 2
+			{
+				if (percentDiff(B[i*(NK * NJ) + j*NK + k], B_outputFromGpu[i*(NK * NJ) + j*NK + k]) > PERCENT_DIFF_ERROR_THRESHOLD)
+				{
+					printf("%d, %d, %d, CPU is %f, GPU is %f.\n", i, j, k, B[i * (NK * NJ) + j * NK + k], B_outputFromGpu[i * (NK * NJ) + j * NK + k]);
+					fail++;
+				}
+			}	
+		}
+	}
+	
+	// Print results
+	printf("Non-Matching CPU-GPU Outputs Beyond Error Threshold of %4.2f Percent: %d\n", PERCENT_DIFF_ERROR_THRESHOLD, fail);
+}
+
+__global__ void convolution3D_kernel(DATA_TYPE *A, DATA_TYPE *B, uint64_t NI, uint64_t NJ, uint64_t NK, uint64_t block_size)
+{
+	cooperative_groups::thread_block block = cooperative_groups::this_thread_block();
+	pipeline pipe;
+	DATA_TYPE c11, c12, c13, c21, c22, c23, c31, c32, c33;
+
+	c11 = +2;
+	c21 = +5;
+	c31 = -8;
+	c12 = -3;
+	c22 = +6;
+	c32 = -9;
+	c13 = +4;
+	c23 = +7;
+	c33 = +10;
+
+	uint64_t tile_dim_x = (NJ + DIM_THREAD_BLOCK - 1) / (DIM_THREAD_BLOCK * BATCH_SIZE);
+
+	__shared__ DATA_TYPE tmp_A[PREFETCH_COUNT][DIM_THREAD_BLOCK * BATCH_SIZE + KERNEL - 1][DIM_THREAD_BLOCK * BATCH_SIZE + KERNEL - 1][DIM_THREAD_BLOCK * BATCH_SIZE + KERNEL - 1];
+	__shared__ DATA_TYPE tmp_B[DIM_THREAD_BLOCK * BATCH_SIZE][DIM_THREAD_BLOCK * BATCH_SIZE][DIM_THREAD_BLOCK * BATCH_SIZE];
+
+	// uint64_t total_tiles = tile_dim_x * tile_dim_x * tile_dim_x;
+
+	uint64_t tiles_this_block_x = (block_size / (DIM_THREAD_BLOCK * BATCH_SIZE));
+	uint64_t tiles_this_block = tiles_this_block_x * tiles_this_block_x * tiles_this_block_x;
+
+	uint64_t base_tile = (blockIdx.z * gridDim.y * gridDim.x + blockIdx.y * gridDim.x + blockIdx.x) * tiles_this_block;
+	uint64_t fetch = base_tile;
+	uint64_t end_tile = fetch + tiles_this_block;
+
+	// printf("block_size is %d, tile_dim_x is %d, tiles_this_block_x is %d.\n", block_size, tile_dim_x, tiles_this_block_x);
+
+	for (uint64_t compute = fetch; compute < end_tile; compute++)
+	{
+		for (; fetch < end_tile && fetch < compute + PREFETCH_COUNT; fetch++)
+		{
+			// block id
+			uint64_t offset = fetch - base_tile;
+			uint64_t block_id = fetch / tiles_this_block;
+
+			uint64_t bz = block_id / (gridDim.y * gridDim.x) * tiles_this_block_x + offset / (tiles_this_block_x * tiles_this_block_x);
+			uint64_t by = block_id % (gridDim.y * gridDim.x) / gridDim.x * tiles_this_block_x + offset % (tiles_this_block_x * tiles_this_block_x) / tiles_this_block_x;
+			uint64_t bx = block_id % (gridDim.y * gridDim.x) % gridDim.x * tiles_this_block_x + offset % (tiles_this_block_x * tiles_this_block_x) % tiles_this_block_x;
+
+			// thread id
+			uint64_t tx = threadIdx.x;
+			uint64_t ty = threadIdx.y;
+			uint64_t tz = threadIdx.z;
+
+			uint64_t index_A_z = DIM_THREAD_BLOCK * BATCH_SIZE * bz + BATCH_SIZE * tz;
+			uint64_t index_A_y = DIM_THREAD_BLOCK * BATCH_SIZE * by + BATCH_SIZE * ty;
+			uint64_t index_A_x = DIM_THREAD_BLOCK * BATCH_SIZE * bx + BATCH_SIZE * tx;
+
+			uint64_t index_A_z_start = DIM_THREAD_BLOCK * BATCH_SIZE * bz;
+			uint64_t index_A_y_start = DIM_THREAD_BLOCK * BATCH_SIZE * by;
+			uint64_t index_A_x_start = DIM_THREAD_BLOCK * BATCH_SIZE * bx;
+
+			uint64_t index_A_z_bound = DIM_THREAD_BLOCK * BATCH_SIZE * bz + BATCH_SIZE * DIM_THREAD_BLOCK;
+			uint64_t index_A_y_bound = DIM_THREAD_BLOCK * BATCH_SIZE * by + BATCH_SIZE * DIM_THREAD_BLOCK;
+			uint64_t index_A_x_bound = DIM_THREAD_BLOCK * BATCH_SIZE * bx + BATCH_SIZE * DIM_THREAD_BLOCK;
+
+			// fetch A
+			for (uint64_t i = 0; i < BATCH_SIZE; i++)
+			{
+				for (uint64_t j = 0; j < BATCH_SIZE; j++)
+				{
+					for (uint64_t k = 0; k < BATCH_SIZE; k++)
+					{
+						if ((index_A_z + i) < NI && (index_A_y + j) < NJ && (index_A_x + k) < NK)
+						{
+							memcpy_async(tmp_A[fetch % PREFETCH_COUNT][tz * BATCH_SIZE + i][ty * BATCH_SIZE + j][tx * BATCH_SIZE + k], A[(index_A_z + i) * NJ * NK + (index_A_y + j) * NK + index_A_x + k], pipe);
+						}
+					}
+				}
+			}
+
+			// fetch A -- padding
+			for (uint64_t i = 0; i < KERNEL - 1; i++)
+			{
+				for (uint64_t j = 0; j < BATCH_SIZE * DIM_THREAD_BLOCK + KERNEL - 1; j++)
+				{
+					for (uint64_t k = 0; k < BATCH_SIZE * DIM_THREAD_BLOCK + KERNEL - 1; k++)
+					{
+						if ((index_A_z_bound + i) < NI && (index_A_y_start + j) < NJ && (index_A_x_start + k) < NK)
+						{
+							memcpy_async(tmp_A[fetch % PREFETCH_COUNT][DIM_THREAD_BLOCK * BATCH_SIZE + i][j][k], A[(index_A_z_bound + i) * NJ * NK + (index_A_y_start + j) * NK + index_A_x_start + k], pipe);
+						}
+					}
+				}
+			}
+
+			// fetch A -- padding
+			for (uint64_t i = 0; i < BATCH_SIZE * DIM_THREAD_BLOCK + KERNEL - 1; i++)
+			{
+				for (uint64_t j = 0; j < KERNEL - 1; j++)
+				{
+					for (uint64_t k = 0; k < BATCH_SIZE * DIM_THREAD_BLOCK + KERNEL - 1; k++)
+					{
+						if ((index_A_z_start + i) < NI && (index_A_y_bound + j) < NJ && (index_A_x_start + k) < NK)
+						{
+							memcpy_async(tmp_A[fetch % PREFETCH_COUNT][i][DIM_THREAD_BLOCK * BATCH_SIZE + j][k], A[(index_A_z_start + i) * NJ * NK + (index_A_y_bound + j) * NK + index_A_x_start + k], pipe);
+						}
+					}
+				}
+			}
+
+			// fetch A -- padding
+			for (uint64_t i = 0; i < BATCH_SIZE * DIM_THREAD_BLOCK + KERNEL - 1; i++)
+			{
+				for (uint64_t j = 0; j < BATCH_SIZE * DIM_THREAD_BLOCK + KERNEL - 1; j++)
+				{
+					for (uint64_t k = 0; k < KERNEL - 1; k++)
+					{
+						if ((index_A_z_start + i) < NI && (index_A_y_start + j) < NJ && (index_A_x_bound + k) < NK)
+						{
+							memcpy_async(tmp_A[fetch % PREFETCH_COUNT][i][j][DIM_THREAD_BLOCK * BATCH_SIZE + k], A[(index_A_z_start + i) * NJ * NK + (index_A_y_start + j) * NK + index_A_x_bound + k], pipe);
+						}
+					}
+				}
+			}
+			pipe.commit();
+		}
+		if (fetch == end_tile)
+		{
+			for (uint64_t i = 0; i < PREFETCH_COUNT - 1; ++i)
+			{
+				pipe.commit();
+			}
+			++fetch;
+		}
+		pipe.wait_prior<PREFETCH_COUNT - 1>();
+		block.sync();
+
+		// block id
+		uint64_t offset = compute - base_tile;
+		uint64_t block_id = compute / tiles_this_block;
+
+		uint64_t bz = block_id / (gridDim.y * gridDim.x) * tiles_this_block_x + offset / (tiles_this_block_x * tiles_this_block_x);
+		uint64_t by = block_id % (gridDim.y * gridDim.x) / gridDim.x * tiles_this_block_x + offset % (tiles_this_block_x * tiles_this_block_x) / tiles_this_block_x;
+		uint64_t bx = block_id % (gridDim.y * gridDim.x) % gridDim.x * tiles_this_block_x + offset % (tiles_this_block_x * tiles_this_block_x) % tiles_this_block_x;
+
+		// thread id
+		uint64_t tx = threadIdx.x;
+		uint64_t ty = threadIdx.y;
+		uint64_t tz = threadIdx.z;
+
+		uint64_t index_B_z = DIM_THREAD_BLOCK * BATCH_SIZE * bz + BATCH_SIZE * tz + 1;
+		uint64_t index_B_y = DIM_THREAD_BLOCK * BATCH_SIZE * by + BATCH_SIZE * ty + 1;
+		uint64_t index_B_x = DIM_THREAD_BLOCK * BATCH_SIZE * bx + BATCH_SIZE * tx + 1;
+
+							
+		// Computation
+		for (uint64_t i = 0; i < BATCH_SIZE; i++)
+		{
+			for (uint64_t j = 0; j < BATCH_SIZE; j++)
+			{
+				for (uint64_t k = 0; k < BATCH_SIZE; k++)
+				{
+					tmp_B[tz * BATCH_SIZE + i][ty * BATCH_SIZE + j][tx * BATCH_SIZE + k] = 0;
+				}
+			}
+		}
+		block.sync();
+
+		for (uint64_t i = 0; i < BATCH_SIZE; i++)
+		{
+			for (uint64_t j = 0; j < BATCH_SIZE; j++)
+			{
+				for (uint64_t k = 0; k < BATCH_SIZE; k++)
+				{
+					tmp_B[tz * BATCH_SIZE + i][ty * BATCH_SIZE + j][tx * BATCH_SIZE + k] =
+						c11 * tmp_A[compute % PREFETCH_COUNT][tz * BATCH_SIZE + i][ty * BATCH_SIZE + j][tx * BATCH_SIZE + k] + c13 * tmp_A[compute % PREFETCH_COUNT][tz * BATCH_SIZE + i + 2][ty * BATCH_SIZE + j][tx * BATCH_SIZE + k] + c21 * tmp_A[compute % PREFETCH_COUNT][tz * BATCH_SIZE + i][ty * BATCH_SIZE + j][tx * BATCH_SIZE + k] + c23 * tmp_A[compute % PREFETCH_COUNT][tz * BATCH_SIZE + i + 2][ty * BATCH_SIZE + j][tx * BATCH_SIZE + k] + c31 * tmp_A[compute % PREFETCH_COUNT][tz * BATCH_SIZE + i][ty * BATCH_SIZE + j][tx * BATCH_SIZE + k] + c33 * tmp_A[compute % PREFETCH_COUNT][tz * BATCH_SIZE + i + 2][ty * BATCH_SIZE + j][tx * BATCH_SIZE + k] + c12 * tmp_A[compute % PREFETCH_COUNT][tz * BATCH_SIZE + i + 1][ty * BATCH_SIZE + j][tx * BATCH_SIZE + k + 1] + c22 * tmp_A[compute % PREFETCH_COUNT][tz * BATCH_SIZE + i + 1][ty * BATCH_SIZE + j + 1][tx * BATCH_SIZE + k + 1] + c32 * tmp_A[compute % PREFETCH_COUNT][tz * BATCH_SIZE + i + 1][ty * BATCH_SIZE + j + 2][tx * BATCH_SIZE + k + 1] + c11 * tmp_A[compute % PREFETCH_COUNT][tz * BATCH_SIZE + i][ty * BATCH_SIZE + j][tx * BATCH_SIZE + k + 2] + c13 * tmp_A[compute % PREFETCH_COUNT][tz * BATCH_SIZE + i + 2][ty * BATCH_SIZE + j][tx * BATCH_SIZE + k + 2] + c21 * tmp_A[compute % PREFETCH_COUNT][tz * BATCH_SIZE + i][ty * BATCH_SIZE + j + 1][tx * BATCH_SIZE + k + 2] + c23 * tmp_A[compute % PREFETCH_COUNT][tz * BATCH_SIZE + i + 2][ty * BATCH_SIZE + j + 1][tx * BATCH_SIZE + k + 2] + c31 * tmp_A[compute % PREFETCH_COUNT][tz * BATCH_SIZE + i][ty * BATCH_SIZE + j + 2][tx * BATCH_SIZE + k + 2] + c33 * tmp_A[compute % PREFETCH_COUNT][tz * BATCH_SIZE + i + 2][ty * BATCH_SIZE + j + 2][tx * BATCH_SIZE + k + 2];
+				}
+			}
+		}
+		block.sync();
+
+		// Store B
+		for (uint64_t i = 0; i < BATCH_SIZE; i++)
+		{
+			for (uint64_t j = 0; j < BATCH_SIZE; j++)
+			{
+				for (uint64_t k = 0; k < BATCH_SIZE; k++)
+				{
+					if ((index_B_z + i + 1) < NI && (index_B_y + j + 1) < NJ && (index_B_x + k + 1) < NK)
+					{
+						B[(index_B_z + i) * NJ * NK + (index_B_y + j) * NK + index_B_x + k] = tmp_B[tz * BATCH_SIZE + i][ty * BATCH_SIZE + j][tx * BATCH_SIZE + k];
+					}
+				}
+			}
+		}
+		block.sync();
+	}
+}
+
+
+
+void convolution3DCuda(DATA_TYPE* A, DATA_TYPE* B, DATA_TYPE* A_gpu, DATA_TYPE* B_gpu)
+{
+	double t_start, t_end;
+
+	dim3 block(DIM_THREAD_BLOCK, DIM_THREAD_BLOCK, DIM_THREAD_BLOCK);
+	dim3 grid(nblocks, nblocks, nblocks);
+
+	uint64_t block_size = (NI + (nblocks - 1)) / nblocks;
+
+	// t_start = rtclock();
+
+	cudaMemcpy(A_gpu, A, sizeof(DATA_TYPE) * NI * NJ * NK, cudaMemcpyHostToDevice);
+	convolution3D_kernel<<<grid, block>>>(A_gpu, B_gpu, NI, NJ, NK, block_size);
+	cudaDeviceSynchronize();
+	cudaMemcpy(B, B_gpu, sizeof(DATA_TYPE) * NI * NJ * NK, cudaMemcpyDeviceToHost);
+
+	// t_end = rtclock();
+	// fprintf(stdout, "GPU Runtime: %0.6lfs\n", t_end - t_start);
+
+}
+
+int main(int argc, char *argv[])
+{
+	uint64_t start_tsc = rdtsc();
+	uint64_t start_tsp = rdtsp();
+	printf("start_tsc %lu start_tsp %lu\n", start_tsc, start_tsp);
+	if (argc >= 5) {
+		NI = atoll(argv[1]);
+		NJ = atoll(argv[2]);
+		NK = atoll(argv[3]);
+		nblocks = atoi(argv[4]);
+	} else {
+		NI = SIZE;
+		NJ = SIZE;
+		NK = SIZE;
+		nblocks = NBLOCKS;
+	}
+	double t_start, t_end;
+
+	DATA_TYPE* A;
+	DATA_TYPE* B;
+	DATA_TYPE *B_ref;
+	DATA_TYPE *A_gpu;
+	DATA_TYPE *B_gpu;
+
+	A = (DATA_TYPE*)malloc(NI*NJ*NK*sizeof(DATA_TYPE));
+	B = (DATA_TYPE*)malloc(NI*NJ*NK*sizeof(DATA_TYPE));
+	B_ref = (DATA_TYPE*)malloc(NI*NJ*NK*sizeof(DATA_TYPE));
+	initCPU(A);
+	GPU_argv_init();
+
+	initTrace();
+	startCPU();
+
+	cudaMalloc(&A_gpu, sizeof(DATA_TYPE) * NI * NJ * NK);
+	cudaMalloc(&B_gpu, sizeof(DATA_TYPE) * NI * NJ * NK);
+	
+	// initGPU(A_gpu);
+
+	convolution3DCuda(A, B, A_gpu, B_gpu);
+
+	cudaFree(A_gpu);
+	cudaFree(B_gpu);
+	endCPU();
+	finiTrace();
+
+	// t_start = rtclock();
+	// conv3D(A, B_ref);
+	// t_end = rtclock();
+	// fprintf(stdout, "CPU Runtime: %0.6lfs\n", t_end - t_start);
+	
+	// compareResults(B, B_ref);
+	free(A);
+	free(B);
+
+    return 0;
+}

workloads/micro/async/3DCONV/Makefile

@@ -0,0 +1,12 @@
+include ../../../common/make.config
+
+NVCCCFLAGS = -I$(CUPTI_INCLUDE) -L$(CUPTI_LIB_DIR) -std=c++11 -lcuda -lcupti -arch=sm_80 -O3
+NVCC = $(CUDA_DIR)/bin/nvcc
+
+EXECUTABLE := 3DConvolution
+CUFILES := 3DConvolution.cu $(CUPTI_ADD_COMMON)/cupti_add.cpp $(CUPTI_ADD_COMMON)/cpu_timestamps.cpp
+
+all:
+	$(NVCC) ${NVCCCFLAGS} ${CUFILES} ${DEF} -o ${EXECUTABLE}
+clean:
+	rm -f *.o 3DConvolution

workloads/micro/async/3DCONV/run.sh

@@ -0,0 +1,2 @@
+#./3DConvolution 768 768 768 8
+./3DConvolution 1536 1536 1536 8

workloads/micro/async/3DCONV/run_large.sh

@@ -0,0 +1 @@
+./3DConvolution 384 384 384 8

workloads/micro/async/3DCONV/run_medium.sh

@@ -0,0 +1 @@
+./3DConvolution 192 192 192 4

workloads/micro/async/3DCONV/run_mega.sh

@@ -0,0 +1 @@
+./3DConvolution 1536 1536 1536 8

workloads/micro/async/3DCONV/run_small.sh

@@ -0,0 +1 @@
+./3DConvolution 96 96 96 4

workloads/micro/async/3DCONV/run_super.sh

@@ -0,0 +1 @@
+./3DConvolution 768 768 768 8

workloads/micro/async/3DCONV/run_tiny.sh

@@ -0,0 +1 @@
+./3DConvolution 48 48 48 2

workloads/micro/async/gemm/Makefile

@@ -0,0 +1,12 @@
+include ../../../common/make.config
+
+NVCCCFLAGS = -I$(CUPTI_INCLUDE) -L$(CUPTI_LIB_DIR) -std=c++11 -lcuda -lcupti -arch=sm_80 -O3
+NVCC = $(CUDA_DIR)/bin/nvcc
+
+EXECUTABLE := gemm
+CUFILES := gemm.cu $(CUPTI_ADD_COMMON)/cupti_add.cpp $(CUPTI_ADD_COMMON)/cpu_timestamps.cpp
+
+all:
+	$(NVCC) ${NVCCCFLAGS} ${CUFILES} ${DEF} -o ${EXECUTABLE}
+clean:
+	rm -f *.o gemm

workloads/micro/async/gemm/gemm.cu

@@ -0,0 +1,277 @@
+/**
+ * gemm.cu: This file is part of the PolyBench/GPU 1.0 test suite.
+ *
+ *
+ * Contact: Scott Grauer-Gray <sgrauerg@gmail.com>
+ * Louis-Noel Pouchet <pouchet@cse.ohio-state.edu>
+ * Web address: http://www.cse.ohio-state.edu/~pouchet/software/polybench/GPU
+ */
+
+#include <unistd.h>
+#include <stdio.h>
+#include <time.h>
+#include <sys/time.h>
+#include <stdlib.h>
+#include <stdarg.h>
+#include <string.h>
+#include <cuda.h>
+#include "../../../common/cupti_add.h"
+#include "../../../common/cpu_timestamps.h"
+
+#include <cooperative_groups.h>
+#include <cooperative_groups/memcpy_async.h>
+
+using namespace nvcuda::experimental;
+
+#define PREFETCH_COUNT 2
+
+#define SMALL_FLOAT_VAL 0.00000001f
+
+double rtclock()
+{
+	struct timezone Tzp;
+	struct timeval Tp;
+	uint64_t stat;
+	stat = gettimeofday(&Tp, &Tzp);
+	if (stat != 0)
+		printf("Error return from gettimeofday: %d", stat);
+	return (Tp.tv_sec + Tp.tv_usec * 1.0e-6);
+}
+
+float absVal(float a)
+{
+	if (a < 0)
+	{
+		return (a * -1);
+	}
+	else
+	{
+		return a;
+	}
+}
+
+float percentDiff(double val1, double val2)
+{
+	if ((absVal(val1) < 0.01) && (absVal(val2) < 0.01))
+	{
+		return 0.0f;
+	}
+
+	else
+	{
+		return 100.0f * (absVal(absVal(val1 - val2) / absVal(val1 + SMALL_FLOAT_VAL)));
+	}
+}
+
+//define the error threshold for the results "not matching"
+#define PERCENT_DIFF_ERROR_THRESHOLD 0.05
+
+/* Problem size */
+#define SIZE 4096
+uint64_t NI;
+uint64_t NJ;
+uint64_t NK;
+
+/* Thread block dimensions */
+#define DIM_THREAD_BLOCK_X 32
+#define DIM_THREAD_BLOCK_Y 32
+
+
+/* Declared constant values for ALPHA and BETA (same as values in PolyBench 2.0) */
+#define ALPHA 1.1f
+#define BETA 1.1f
+
+/* Can switch DATA_TYPE between float and double */
+typedef float DATA_TYPE;
+// typedef uint64_t DATA_TYPE;
+
+void gemm(DATA_TYPE *A, DATA_TYPE *B, DATA_TYPE *C)
+{
+	uint64_t i,j,k;
+	
+	for (i = 0; i < NI; i++) {
+    	for (j = 0; j < NJ; j++) {
+			C[i*NJ + j] *= BETA;
+			for (k = 0; k < NK; ++k) {
+	  			C[i*NJ + j] += ALPHA * A[i*NK + k] * B[k*NJ + j];
+			}
+      	}
+	}
+}
+
+void init(DATA_TYPE *A, DATA_TYPE *B, DATA_TYPE *C, DATA_TYPE *C_ref)
+{
+	uint64_t i, j;
+
+  	for (i = 0; i < NI; i++)
+    	for (j = 0; j < NK; j++)
+			A[i*NK + j] = ((DATA_TYPE) i*j) / NI;
+
+  	for (i = 0; i < NK; i++)
+    	for (j = 0; j < NJ; j++)
+			B[i*NJ + j] = ((DATA_TYPE) i*j + 1) / NJ;
+
+  	for (i = 0; i < NI; i++) {
+		for (j = 0; j < NJ; j++) {
+			C[i * NJ + j] = ((DATA_TYPE)i * j + 2) / NJ;
+			C_ref[i * NJ + j] = ((DATA_TYPE)i * j + 2) / NJ;
+		}
+	} 
+    	
+}
+
+
+void compareResults(DATA_TYPE* C, DATA_TYPE* C_outputFromGpu)
+{
+	uint64_t i, j, fail;
+	fail = 0;
+	
+	// Compare C1 and C2
+	for (i=0; i < NI; i++) 
+	{
+		for (j=0; j < NJ; j++) 
+		{
+			// printf("%d, %d, GPU is %f, CPU is %f.\n", i, j, C[i*NJ + j], C_outputFromGpu[i*NJ + j]);
+			if (percentDiff(C[i*NJ + j], C_outputFromGpu[i*NJ + j]) > PERCENT_DIFF_ERROR_THRESHOLD) 
+			{
+				printf("%d, %d, GPU is %f, CPU is %f.\n", i, j, C[i*NJ + j], C_outputFromGpu[i*NJ + j]);
+				fail++;
+			}
+		}
+	}
+	
+	// Print results
+	printf("Non-Matching CPU-GPU Outputs Beyond Error Threshold of %4.2f Percent: %d\n", PERCENT_DIFF_ERROR_THRESHOLD, fail);
+}
+
+__global__ void gemm_kernel(DATA_TYPE *a, DATA_TYPE *b, DATA_TYPE *c, uint64_t NI, uint64_t NK, uint64_t NJ)
+{
+	cooperative_groups::thread_block block = cooperative_groups::this_thread_block();
+	pipeline pipe;
+
+	uint64_t row = blockIdx.y * blockDim.y + threadIdx.y;
+	uint64_t col = blockIdx.x * blockDim.x + threadIdx.x;
+
+	__shared__ DATA_TYPE s_a[DIM_THREAD_BLOCK_X * DIM_THREAD_BLOCK_Y * PREFETCH_COUNT];
+	__shared__ DATA_TYPE s_b[DIM_THREAD_BLOCK_X * DIM_THREAD_BLOCK_Y * PREFETCH_COUNT];
+
+	DATA_TYPE tmp = BETA * c[row * NJ + col];
+
+	uint64_t base_tiles = 0;
+	uint64_t end_tile = base_tiles + NK / blockDim.x;
+
+	uint64_t fetch = base_tiles;
+	uint64_t tile_size = DIM_THREAD_BLOCK_X;
+	uint64_t mem_size = DIM_THREAD_BLOCK_X * DIM_THREAD_BLOCK_Y;
+
+	for (uint64_t compute = fetch; compute < end_tile; compute++)
+	{
+		for (; fetch < end_tile && fetch < compute + PREFETCH_COUNT; fetch++)
+		{
+			memcpy_async(s_a[(fetch % PREFETCH_COUNT) * mem_size + (threadIdx.y * blockDim.x + threadIdx.x)], a[row * NK + fetch * tile_size + threadIdx.x], pipe);
+			memcpy_async(s_b[(fetch % PREFETCH_COUNT) * mem_size + (threadIdx.y * blockDim.x + threadIdx.x)], b[(fetch * tile_size + threadIdx.y) * NJ + col], pipe);
+
+			pipe.commit();
+		}
+		if (fetch == end_tile) {
+			for (uint64_t i = 0; i < PREFETCH_COUNT-1; ++i) { pipe.commit(); }
+      		++fetch;
+		}
+		pipe.wait_prior<PREFETCH_COUNT - 1>();
+		block.sync();
+
+		for (uint64_t k = 0; k < blockDim.x; k++)
+		{
+			tmp += ALPHA * s_a[(compute % PREFETCH_COUNT) * mem_size + (threadIdx.y * blockDim.x + k)] * s_b[(compute % PREFETCH_COUNT) * mem_size + (k * blockDim.x + threadIdx.x)];
+		}
+		block.sync();
+	}
+
+	c[row * NJ + col] = tmp;
+}
+
+void gemmCuda(DATA_TYPE *A, DATA_TYPE *B, DATA_TYPE *C, DATA_TYPE *A_gpu, DATA_TYPE *B_gpu, DATA_TYPE *C_gpu)
+{
+	double t_start, t_end;
+
+	dim3 block(DIM_THREAD_BLOCK_X, DIM_THREAD_BLOCK_Y);
+	dim3 grid((size_t)(ceil( ((float)NI)/ ((float)block.x) )),(size_t)(ceil( ((float)NJ)/ ((float)block.y) )));
+
+	//t_start = rtclock();
+	cudaMemcpy(A_gpu, A, sizeof(DATA_TYPE) * NI * NK, cudaMemcpyHostToDevice);
+	cudaMemcpy(B_gpu, B, sizeof(DATA_TYPE) * NK * NJ, cudaMemcpyHostToDevice);
+	cudaMemcpy(C_gpu, C, sizeof(DATA_TYPE) * NI * NJ, cudaMemcpyHostToDevice);
+	gemm_kernel<<< grid, block >>>(A_gpu, B_gpu, C_gpu, NI, NK, NJ);
+	cudaDeviceSynchronize();
+	cudaMemcpy(C, C_gpu, sizeof(DATA_TYPE) * NI * NJ, cudaMemcpyDeviceToHost);
+	//t_end = rtclock();
+
+	//fprintf(stdout, "GPU Runtime: %0.6lfs\n", t_end - t_start);   
+}
+
+int main(int argc, char *argv[])
+{
+	uint64_t start_tsc = rdtsc();
+	uint64_t start_tsp = rdtsp();
+	printf("start_tsc %lu start_tsp %lu\n", start_tsc, start_tsp);
+	if (argc >= 4) {
+		NI = atoll(argv[1]);
+		NK = atoll(argv[2]);
+		NJ = atoll(argv[3]);
+	} else {
+		NI = SIZE;
+		NK = SIZE;
+		NJ = SIZE;
+	}
+
+	double t_start, t_end;
+
+	DATA_TYPE* A;
+	DATA_TYPE* B;  
+	DATA_TYPE* C;
+	DATA_TYPE *C_ref;
+
+	DATA_TYPE *A_gpu;
+	DATA_TYPE *B_gpu;
+	DATA_TYPE *C_gpu; 
+
+	A = (DATA_TYPE*)malloc(NI*NK*sizeof(DATA_TYPE)); 
+	B = (DATA_TYPE*)malloc(NK*NJ*sizeof(DATA_TYPE));   
+	C = (DATA_TYPE*)malloc(NI*NJ*sizeof(DATA_TYPE));
+	C_ref = (DATA_TYPE *)malloc(NI * NJ * sizeof(DATA_TYPE));
+
+	//cudaMallocManaged(&A_gpu, sizeof(DATA_TYPE) * NI * NK);
+	//cudaMallocManaged(&B_gpu, sizeof(DATA_TYPE) * NK * NJ);
+	//cudaMallocManaged(&C_gpu, sizeof(DATA_TYPE) * NI * NJ);
+	init(A, B, C, C_ref);
+
+	GPU_argv_init();
+	
+	initTrace();
+	startCPU();
+	
+	cudaMalloc(&A_gpu, sizeof(DATA_TYPE) * NI * NK);
+	cudaMalloc(&B_gpu, sizeof(DATA_TYPE) * NK * NJ);
+	cudaMalloc(&C_gpu, sizeof(DATA_TYPE) * NI * NJ);
+
+	gemmCuda(A, B, C, A_gpu, B_gpu, C_gpu);
+
+	cudaFree(A_gpu);
+	cudaFree(B_gpu);
+	cudaFree(C_gpu);
+	endCPU();
+	finiTrace();
+
+	// t_start = rtclock();	
+	// gemm(A, B, C_ref);
+	// t_end = rtclock();
+	// fprintf(stdout, "CPU Runtime: %0.6lfs\n", t_end - t_start);
+	
+	// compareResults(C, C_ref);
+	free(A);
+	free(B);  
+	free(C);
+	free(C_ref);
+    return 0;
+}
+

workloads/micro/async/gemm/run.sh

@@ -0,0 +1,3 @@
+#./gemm 1024 1024 1024
+./gemm 32768 32768 32768
+#./gemm 512 512 512

workloads/micro/async/gemm/run_large.sh

@@ -0,0 +1 @@
+./gemm 8192 8192 8192

workloads/micro/async/gemm/run_medium.sh

@@ -0,0 +1 @@
+./gemm 4096 4096 4096

workloads/micro/async/gemm/run_mega.sh

@@ -0,0 +1 @@
+./gemm 32768 32768 32768

workloads/micro/async/gemm/run_small.sh

@@ -0,0 +1 @@
+./gemm 1024 1024 1024

workloads/micro/async/gemm/run_super.sh

@@ -0,0 +1 @@
+./gemm 16384 16384 16384

workloads/micro/async/gemm/run_tiny.sh

@@ -0,0 +1 @@
+./gemm 512 512 512

workloads/micro/async/gemm_perf/Makefile

@@ -0,0 +1,12 @@
+include ../../../common/make.config
+
+NVCCCFLAGS = -I$(CUPTI_INCLUDE) -L$(CUPTI_LIB_DIR) -std=c++11 -lcuda -lcupti -arch=sm_80 -O3
+NVCC = $(CUDA_DIR)/bin/nvcc
+
+EXECUTABLE := gemm
+CUFILES := gemm.cu $(CUPTI_ADD_COMMON)/cupti_add.cpp $(CUPTI_ADD_COMMON)/cpu_timestamps.cpp
+
+all:
+	$(NVCC) ${NVCCCFLAGS} ${CUFILES} ${DEF} -o ${EXECUTABLE}
+clean:
+	rm -f *.o gemm

workloads/micro/async/gemm_perf/gemm.cu

@@ -0,0 +1,277 @@
+/**
+ * gemm.cu: This file is part of the PolyBench/GPU 1.0 test suite.
+ *
+ *
+ * Contact: Scott Grauer-Gray <sgrauerg@gmail.com>
+ * Louis-Noel Pouchet <pouchet@cse.ohio-state.edu>
+ * Web address: http://www.cse.ohio-state.edu/~pouchet/software/polybench/GPU
+ */
+
+#include <unistd.h>
+#include <stdio.h>
+#include <time.h>
+#include <sys/time.h>
+#include <stdlib.h>
+#include <stdarg.h>
+#include <string.h>
+#include <cuda.h>
+#include "../../../common/cupti_add.h"
+#include "../../../common/cpu_timestamps.h"
+
+#include <cooperative_groups.h>
+#include <cooperative_groups/memcpy_async.h>
+
+using namespace nvcuda::experimental;
+
+#define PREFETCH_COUNT 2
+
+#define SMALL_FLOAT_VAL 0.00000001f
+
+double rtclock()
+{
+	struct timezone Tzp;
+	struct timeval Tp;
+	uint64_t stat;
+	stat = gettimeofday(&Tp, &Tzp);
+	if (stat != 0)
+		printf("Error return from gettimeofday: %d", stat);
+	return (Tp.tv_sec + Tp.tv_usec * 1.0e-6);
+}
+
+float absVal(float a)
+{
+	if (a < 0)
+	{
+		return (a * -1);
+	}
+	else
+	{
+		return a;
+	}
+}
+
+float percentDiff(double val1, double val2)
+{
+	if ((absVal(val1) < 0.01) && (absVal(val2) < 0.01))
+	{
+		return 0.0f;
+	}
+
+	else
+	{
+		return 100.0f * (absVal(absVal(val1 - val2) / absVal(val1 + SMALL_FLOAT_VAL)));
+	}
+}
+
+//define the error threshold for the results "not matching"
+#define PERCENT_DIFF_ERROR_THRESHOLD 0.05
+
+/* Problem size */
+#define SIZE 4096
+uint64_t NI;
+uint64_t NJ;
+uint64_t NK;
+
+/* Thread block dimensions */
+#define DIM_THREAD_BLOCK_X 32
+#define DIM_THREAD_BLOCK_Y 32
+
+
+/* Declared constant values for ALPHA and BETA (same as values in PolyBench 2.0) */
+#define ALPHA 1.1f
+#define BETA 1.1f
+
+/* Can switch DATA_TYPE between float and double */
+typedef float DATA_TYPE;
+// typedef uint64_t DATA_TYPE;
+
+void gemm(DATA_TYPE *A, DATA_TYPE *B, DATA_TYPE *C)
+{
+	uint64_t i,j,k;
+	
+	for (i = 0; i < NI; i++) {
+    	for (j = 0; j < NJ; j++) {
+			C[i*NJ + j] *= BETA;
+			for (k = 0; k < NK; ++k) {
+	  			C[i*NJ + j] += ALPHA * A[i*NK + k] * B[k*NJ + j];
+			}
+      	}
+	}
+}
+
+void init(DATA_TYPE *A, DATA_TYPE *B, DATA_TYPE *C, DATA_TYPE *C_ref)
+{
+	uint64_t i, j;
+
+  	for (i = 0; i < NI; i++)
+    	for (j = 0; j < NK; j++)
+			A[i*NK + j] = ((DATA_TYPE) i*j) / NI;
+
+  	for (i = 0; i < NK; i++)
+    	for (j = 0; j < NJ; j++)
+			B[i*NJ + j] = ((DATA_TYPE) i*j + 1) / NJ;
+
+  	for (i = 0; i < NI; i++) {
+		for (j = 0; j < NJ; j++) {
+			C[i * NJ + j] = ((DATA_TYPE)i * j + 2) / NJ;
+			C_ref[i * NJ + j] = ((DATA_TYPE)i * j + 2) / NJ;
+		}
+	} 
+    	
+}
+
+
+void compareResults(DATA_TYPE* C, DATA_TYPE* C_outputFromGpu)
+{
+	uint64_t i, j, fail;
+	fail = 0;
+	
+	// Compare C1 and C2
+	for (i=0; i < NI; i++) 
+	{
+		for (j=0; j < NJ; j++) 
+		{
+			// printf("%d, %d, GPU is %f, CPU is %f.\n", i, j, C[i*NJ + j], C_outputFromGpu[i*NJ + j]);
+			if (percentDiff(C[i*NJ + j], C_outputFromGpu[i*NJ + j]) > PERCENT_DIFF_ERROR_THRESHOLD) 
+			{
+				printf("%d, %d, GPU is %f, CPU is %f.\n", i, j, C[i*NJ + j], C_outputFromGpu[i*NJ + j]);
+				fail++;
+			}
+		}
+	}
+	
+	// Print results
+	printf("Non-Matching CPU-GPU Outputs Beyond Error Threshold of %4.2f Percent: %d\n", PERCENT_DIFF_ERROR_THRESHOLD, fail);
+}
+
+__global__ void gemm_kernel(DATA_TYPE *a, DATA_TYPE *b, DATA_TYPE *c, uint64_t NI, uint64_t NK, uint64_t NJ)
+{
+	cooperative_groups::thread_block block = cooperative_groups::this_thread_block();
+	pipeline pipe;
+
+	uint64_t row = blockIdx.y * blockDim.y + threadIdx.y;
+	uint64_t col = blockIdx.x * blockDim.x + threadIdx.x;
+
+	__shared__ DATA_TYPE s_a[DIM_THREAD_BLOCK_X * DIM_THREAD_BLOCK_Y * PREFETCH_COUNT];
+	__shared__ DATA_TYPE s_b[DIM_THREAD_BLOCK_X * DIM_THREAD_BLOCK_Y * PREFETCH_COUNT];
+
+	DATA_TYPE tmp = BETA * c[row * NJ + col];
+
+	uint64_t base_tiles = 0;
+	uint64_t end_tile = base_tiles + NK / blockDim.x;
+
+	uint64_t fetch = base_tiles;
+	uint64_t tile_size = DIM_THREAD_BLOCK_X;
+	uint64_t mem_size = DIM_THREAD_BLOCK_X * DIM_THREAD_BLOCK_Y;
+
+	for (uint64_t compute = fetch; compute < end_tile; compute++)
+	{
+		for (; fetch < end_tile && fetch < compute + PREFETCH_COUNT; fetch++)
+		{
+			memcpy_async(s_a[(fetch % PREFETCH_COUNT) * mem_size + (threadIdx.y * blockDim.x + threadIdx.x)], a[row * NK + fetch * tile_size + threadIdx.x], pipe);
+			memcpy_async(s_b[(fetch % PREFETCH_COUNT) * mem_size + (threadIdx.y * blockDim.x + threadIdx.x)], b[(fetch * tile_size + threadIdx.y) * NJ + col], pipe);
+
+			pipe.commit();
+		}
+		if (fetch == end_tile) {
+			for (uint64_t i = 0; i < PREFETCH_COUNT-1; ++i) { pipe.commit(); }
+      		++fetch;
+		}
+		pipe.wait_prior<PREFETCH_COUNT - 1>();
+		block.sync();
+
+		for (uint64_t k = 0; k < blockDim.x; k++)
+		{
+			tmp += ALPHA * s_a[(compute % PREFETCH_COUNT) * mem_size + (threadIdx.y * blockDim.x + k)] * s_b[(compute % PREFETCH_COUNT) * mem_size + (k * blockDim.x + threadIdx.x)];
+		}
+		block.sync();
+	}
+
+	c[row * NJ + col] = tmp;
+}
+
+void gemmCuda(DATA_TYPE *A, DATA_TYPE *B, DATA_TYPE *C, DATA_TYPE *A_gpu, DATA_TYPE *B_gpu, DATA_TYPE *C_gpu)
+{
+	double t_start, t_end;
+
+	dim3 block(DIM_THREAD_BLOCK_X, DIM_THREAD_BLOCK_Y);
+	dim3 grid((size_t)(ceil( ((float)NI)/ ((float)block.x) )),(size_t)(ceil( ((float)NJ)/ ((float)block.y) )));
+
+	//t_start = rtclock();
+	cudaMemcpy(A_gpu, A, sizeof(DATA_TYPE) * NI * NK, cudaMemcpyHostToDevice);
+	cudaMemcpy(B_gpu, B, sizeof(DATA_TYPE) * NK * NJ, cudaMemcpyHostToDevice);
+	cudaMemcpy(C_gpu, C, sizeof(DATA_TYPE) * NI * NJ, cudaMemcpyHostToDevice);
+	gemm_kernel<<< grid, block >>>(A_gpu, B_gpu, C_gpu, NI, NK, NJ);
+	cudaDeviceSynchronize();
+	cudaMemcpy(C, C_gpu, sizeof(DATA_TYPE) * NI * NJ, cudaMemcpyDeviceToHost);
+	//t_end = rtclock();
+
+	//fprintf(stdout, "GPU Runtime: %0.6lfs\n", t_end - t_start);   
+}
+
+int main(int argc, char *argv[])
+{
+	uint64_t start_tsc = rdtsc();
+	uint64_t start_tsp = rdtsp();
+	printf("start_tsc %lu start_tsp %lu\n", start_tsc, start_tsp);
+	if (argc >= 4) {
+		NI = atoll(argv[1]);
+		NK = atoll(argv[2]);
+		NJ = atoll(argv[3]);
+	} else {
+		NI = SIZE;
+		NK = SIZE;
+		NJ = SIZE;
+	}
+
+	double t_start, t_end;
+
+	DATA_TYPE* A;
+	DATA_TYPE* B;  
+	DATA_TYPE* C;
+	DATA_TYPE *C_ref;
+
+	DATA_TYPE *A_gpu;
+	DATA_TYPE *B_gpu;
+	DATA_TYPE *C_gpu; 
+
+	A = (DATA_TYPE*)malloc(NI*NK*sizeof(DATA_TYPE)); 
+	B = (DATA_TYPE*)malloc(NK*NJ*sizeof(DATA_TYPE));   
+	C = (DATA_TYPE*)malloc(NI*NJ*sizeof(DATA_TYPE));
+	C_ref = (DATA_TYPE *)malloc(NI * NJ * sizeof(DATA_TYPE));
+
+	//cudaMallocManaged(&A_gpu, sizeof(DATA_TYPE) * NI * NK);
+	//cudaMallocManaged(&B_gpu, sizeof(DATA_TYPE) * NK * NJ);
+	//cudaMallocManaged(&C_gpu, sizeof(DATA_TYPE) * NI * NJ);
+	init(A, B, C, C_ref);
+
+	GPU_argv_init();
+	
+	//initTrace();
+	startCPU();
+	
+	cudaMalloc(&A_gpu, sizeof(DATA_TYPE) * NI * NK);
+	cudaMalloc(&B_gpu, sizeof(DATA_TYPE) * NK * NJ);
+	cudaMalloc(&C_gpu, sizeof(DATA_TYPE) * NI * NJ);
+
+	gemmCuda(A, B, C, A_gpu, B_gpu, C_gpu);
+
+	cudaFree(A_gpu);
+	cudaFree(B_gpu);
+	cudaFree(C_gpu);
+	endCPU();
+	//finiTrace();
+
+	// t_start = rtclock();	
+	// gemm(A, B, C_ref);
+	// t_end = rtclock();
+	// fprintf(stdout, "CPU Runtime: %0.6lfs\n", t_end - t_start);
+	
+	// compareResults(C, C_ref);
+	free(A);
+	free(B);  
+	free(C);
+	free(C_ref);
+    return 0;
+}
+

workloads/micro/async/gemm_perf/run.sh

@@ -0,0 +1,3 @@
+#./gemm 1024 1024 1024
+./gemm 32768 32768 32768
+#./gemm 512 512 512

workloads/micro/async/gemm_perf/run_large.sh

@@ -0,0 +1 @@
+./gemm 8192 8192 8192

workloads/micro/async/gemm_perf/run_medium.sh

@@ -0,0 +1 @@
+./gemm 4096 4096 4096

workloads/micro/async/gemm_perf/run_mega.sh

@@ -0,0 +1 @@
+./gemm 32768 32768 32768

workloads/micro/async/gemm_perf/run_small.sh

@@ -0,0 +1 @@
+./gemm 1024 1024 1024

workloads/micro/async/gemm_perf/run_super.sh

@@ -0,0 +1 @@
+./gemm 16384 16384 16384

workloads/micro/async/gemm_perf/run_tiny.sh

@@ -0,0 +1 @@
+./gemm 512 512 512

workloads/micro/async/gemv/Makefile

@@ -0,0 +1,12 @@
+include ../../../common/make.config
+
+NVCCCFLAGS = -I$(CUPTI_INCLUDE) -L$(CUPTI_LIB_DIR) -std=c++11 -lcuda -lcupti -arch=sm_80 -O3
+NVCC = $(CUDA_DIR)/bin/nvcc
+
+EXECUTABLE := gemv
+CUFILES := gemv.cu $(CUPTI_ADD_COMMON)/cupti_add.cpp $(CUPTI_ADD_COMMON)/cpu_timestamps.cpp
+
+all:
+	$(NVCC) ${NVCCCFLAGS} ${CUFILES} ${DEF} -o ${EXECUTABLE}
+clean:
+	rm -f *.o gemv

workloads/micro/async/gemv/gemv.cu

@@ -0,0 +1,269 @@
+/**
+ * gemm.cu: This file is part of the PolyBench/GPU 1.0 test suite.
+ *
+ *
+ * Contact: Scott Grauer-Gray <sgrauerg@gmail.com>
+ * Louis-Noel Pouchet <pouchet@cse.ohio-state.edu>
+ * Web address: http://www.cse.ohio-state.edu/~pouchet/software/polybench/GPU
+ */
+
+#include <unistd.h>
+#include <stdio.h>
+#include <time.h>
+#include <sys/time.h>
+#include <stdlib.h>
+#include <stdarg.h>
+#include <string.h>
+#include <cuda.h>
+#include "../../../common/cupti_add.h"
+#include "../../../common/cpu_timestamps.h"
+
+#include <cooperative_groups.h>
+#include <cooperative_groups/memcpy_async.h>
+
+using namespace nvcuda::experimental;
+
+#define PREFETCH_COUNT 2
+
+#define SMALL_FLOAT_VAL 0.00000001f
+
+double rtclock()
+{
+	struct timezone Tzp;
+	struct timeval Tp;
+	uint64_t stat;
+	stat = gettimeofday(&Tp, &Tzp);
+	if (stat != 0)
+		printf("Error return from gettimeofday: %d", stat);
+	return (Tp.tv_sec + Tp.tv_usec * 1.0e-6);
+}
+
+float absVal(float a)
+{
+	if (a < 0)
+	{
+		return (a * -1);
+	}
+	else
+	{
+		return a;
+	}
+}
+
+float percentDiff(double val1, double val2)
+{
+	if ((absVal(val1) < 0.01) && (absVal(val2) < 0.01))
+	{
+		return 0.0f;
+	}
+
+	else
+	{
+		return 100.0f * (absVal(absVal(val1 - val2) / absVal(val1 + SMALL_FLOAT_VAL)));
+	}
+}
+
+// define the error threshold for the results "not matching"
+#define PERCENT_DIFF_ERROR_THRESHOLD 0.05
+
+/* Problem size */
+#define SIZE 40960
+uint64_t NI;
+uint64_t NJ;
+
+/* Thread block dimensions */
+#define DIM_THREAD_BLOCK 256
+
+#define BATCH_SIZE 16
+
+/* Declared constant values for ALPHA and BETA (same as values in PolyBench 2.0) */
+#define ALPHA 1.1f
+#define BETA 1.1f
+
+/* Can switch DATA_TYPE between float and double */
+typedef float DATA_TYPE;
+// typedef uint64_t DATA_TYPE;
+
+void gemv(DATA_TYPE *A, DATA_TYPE *B, DATA_TYPE *C)
+{
+	uint64_t i, j;
+
+	for (i = 0; i < NI; i++)
+	{
+		C[i] *= BETA;
+		for (j = 0; j < NJ; j++)
+		{
+			C[i] += ALPHA * A[i * NJ + j] * B[j];
+		}
+	}
+}
+
+void init(DATA_TYPE *A, DATA_TYPE *B, DATA_TYPE *C, DATA_TYPE *C_ref)
+{
+	uint64_t i, j;
+
+	for (i = 0; i < NI; i++)
+		for (j = 0; j < NJ; j++)
+			A[i * NJ + j] = ((DATA_TYPE)i * j) / NI;
+
+	for (j = 0; j < NJ; j++)
+		B[j] = ((DATA_TYPE)j + 1) / NJ;
+
+	for (i = 0; i < NI; i++)
+	{
+		C[i] = ((DATA_TYPE)i + 2) / NI;
+		C_ref[i] = ((DATA_TYPE)i + 2) / NI;
+	}
+}
+
+void compareResults(DATA_TYPE *C, DATA_TYPE *C_outputFromGpu)
+{
+	uint64_t i, fail;
+	fail = 0;
+
+	// Compare C1 and C2
+	for (i = 0; i < NI; i++)
+	{
+		if (percentDiff(C[i], C_outputFromGpu[i]) > PERCENT_DIFF_ERROR_THRESHOLD)
+		{
+			fail++;
+			printf("%d, GPU is %f, CPU is %f.\n", i, C[i], C_outputFromGpu[i]);
+		}
+	}
+
+	// Print results
+	printf("Non-Matching CPU-GPU Outputs Beyond Error Threshold of %4.2f Percent: %d\n", PERCENT_DIFF_ERROR_THRESHOLD, fail);
+}
+
+__global__ void gemv_kernel(DATA_TYPE *a, DATA_TYPE *b, DATA_TYPE *c, uint64_t NI, uint64_t NJ)
+{
+	cooperative_groups::thread_block block = cooperative_groups::this_thread_block();
+	pipeline pipe;
+	uint64_t row = blockIdx.x * blockDim.x + threadIdx.x;
+	uint64_t tx = threadIdx.x;
+
+	__shared__ DATA_TYPE s_b[PREFETCH_COUNT][DIM_THREAD_BLOCK][BATCH_SIZE];
+
+	DATA_TYPE tmp = BETA * c[row];
+	__syncthreads();
+
+	uint64_t fetch = 0;
+	uint64_t end_tile = NJ / BATCH_SIZE;
+
+	for (uint64_t compute = fetch; compute < end_tile; compute++)
+	{
+		for (; fetch < end_tile && fetch < compute + PREFETCH_COUNT; fetch++) 
+		{
+			uint64_t base_index = fetch * BATCH_SIZE;
+			for (uint64_t k = 0; k < BATCH_SIZE; k++)
+			{
+				memcpy_async(s_b[fetch % PREFETCH_COUNT][tx][k], b[base_index + k], pipe);
+			}
+			pipe.commit();
+		}
+		if (fetch == end_tile)
+		{
+			for (uint64_t i = 0; i < PREFETCH_COUNT - 1; ++i)
+			{
+				pipe.commit();
+			}
+			++fetch;
+		}
+		pipe.wait_prior<PREFETCH_COUNT - 1>();
+		block.sync();
+
+		uint64_t base_index = compute * BATCH_SIZE;
+		for (uint64_t k = 0; k < BATCH_SIZE; k++)
+		{
+			tmp += ALPHA * a[row * NJ + base_index + k] * s_b[compute % PREFETCH_COUNT][tx][k];
+		}
+		block.sync();
+	}
+	c[row] = tmp;
+}
+
+void gemvCuda(DATA_TYPE *A, DATA_TYPE *B, DATA_TYPE *C, DATA_TYPE *A_gpu, DATA_TYPE *B_gpu, DATA_TYPE *C_gpu)
+{
+	double t_start, t_end;
+
+	dim3 block(DIM_THREAD_BLOCK);
+	dim3 grid(NI / (DIM_THREAD_BLOCK));
+
+	// t_start = rtclock();
+	cudaMemcpy(A_gpu, A, sizeof(DATA_TYPE) * NI * NJ, cudaMemcpyHostToDevice);
+	cudaMemcpy(B_gpu, B, sizeof(DATA_TYPE) * NJ, cudaMemcpyHostToDevice);
+	cudaMemcpy(C_gpu, C, sizeof(DATA_TYPE) * NI, cudaMemcpyHostToDevice);
+	gemv_kernel<<<grid, block>>>(A_gpu, B_gpu, C_gpu, NI, NJ);
+	cudaDeviceSynchronize();
+	cudaMemcpy(C, C_gpu, sizeof(DATA_TYPE) * NI, cudaMemcpyDeviceToHost);
+	// t_end = rtclock();
+
+	// fprintf(stdout, "GPU Runtime: %0.6lfs\n", t_end - t_start);
+}
+
+int main(int argc, char *argv[])
+{
+	uint64_t start_tsc = rdtsc();
+	uint64_t start_tsp = rdtsp();
+	printf("start_tsc %lu start_tsp %lu\n", start_tsc, start_tsp);
+	if (argc >= 3)
+	{
+		NI = atoll(argv[1]);
+		NJ = atoll(argv[2]);
+	}
+	else
+	{
+		NI = SIZE;
+		NJ = SIZE;
+	}
+
+	double t_start, t_end;
+
+	DATA_TYPE *A;
+	DATA_TYPE *B;
+	DATA_TYPE *C;
+	DATA_TYPE *C_ref;
+
+	DATA_TYPE *A_gpu;
+	DATA_TYPE *B_gpu;
+	DATA_TYPE *C_gpu;
+
+	A = (DATA_TYPE *)malloc(NI * NJ * sizeof(DATA_TYPE));
+	B = (DATA_TYPE *)malloc(NJ * sizeof(DATA_TYPE));
+	C = (DATA_TYPE *)malloc(NI * sizeof(DATA_TYPE));
+	C_ref = (DATA_TYPE *)malloc(NI * sizeof(DATA_TYPE));
+
+	// cudaMallocManaged(&A_gpu, sizeof(DATA_TYPE) * NI * NK);
+	// cudaMallocManaged(&B_gpu, sizeof(DATA_TYPE) * NK * NJ);
+	// cudaMallocManaged(&C_gpu, sizeof(DATA_TYPE) * NI * NJ);
+
+	init(A, B, C, C_ref);
+
+	GPU_argv_init();
+	initTrace();
+	startCPU();
+
+	cudaMalloc(&A_gpu, sizeof(DATA_TYPE) * NI * NJ);
+	cudaMalloc(&B_gpu, sizeof(DATA_TYPE) * NJ);
+	cudaMalloc(&C_gpu, sizeof(DATA_TYPE) * NI);
+
+	gemvCuda(A, B, C, A_gpu, B_gpu, C_gpu);
+
+	cudaFree(A_gpu);
+	cudaFree(B_gpu);
+	cudaFree(C_gpu);
+	endCPU();
+	finiTrace();
+
+	// t_start = rtclock();
+	// gemv(A, B, C_ref);
+	// t_end = rtclock();
+	// fprintf(stdout, "CPU Runtime: %0.6lfs\n", t_end - t_start);
+
+	// compareResults(C, C_ref);
+	free(A);
+	free(B);
+	free(C);
+	free(C_ref);
+	return 0;
+}