add hardcode mode for kernel launch

MICRO_AE_example/anti-coalescing-transformation/hist.cu

@@ -0,0 +1,104 @@
+#include <cstdio>
+#include <random>
+#include <stdint.h>
+#include <sys/time.h>
+
+__global__ void Histogram(uint32_t *pixels, uint32_t *histogram,
+                          uint32_t num_colors, uint32_t num_pixels) {
+  int tid = blockIdx.x * blockDim.x + threadIdx.x;
+  int gsize = gridDim.x * blockDim.x;
+
+  uint32_t priv_hist[256];
+  for (uint32_t i = 0; i < num_colors; i++) {
+    priv_hist[i] = 0;
+  }
+  uint32_t index = tid;
+  while (index < num_pixels) {
+    uint32_t color = pixels[index];
+    priv_hist[color]++;
+    index += gsize;
+  }
+
+  __syncthreads();
+  for (uint32_t i = 0; i < num_colors; i++) {
+    if (priv_hist[i] > 0) {
+      atomicAdd(histogram + i, priv_hist[i]);
+    }
+  }
+}
+
+__global__ void opt_Histogram(uint32_t *pixels, uint32_t *histogram,
+                              uint32_t num_colors, uint32_t num_pixels) {
+  int tid = blockIdx.x * blockDim.x + threadIdx.x;
+  int gsize = gridDim.x * blockDim.x;
+
+  uint32_t priv_hist[256];
+  for (uint32_t i = 0; i < num_colors; i++) {
+    priv_hist[i] = 0;
+  }
+  uint32_t index = tid;
+  __shared__ bool has_activated_thread;
+  bool activated = true;
+  do {
+    has_activated_thread = false;
+    __syncthreads();
+    activated = activated & (index < num_pixels);
+    has_activated_thread |= activated;
+    if (activated) {
+      uint32_t color = pixels[index];
+      priv_hist[color]++;
+      index += gsize;
+    }
+  } while (has_activated_thread);
+  __syncthreads();
+  for (uint32_t i = 0; i < num_colors; i++) {
+    if (priv_hist[i] > 0) {
+      atomicAdd(histogram + i, priv_hist[i]);
+    }
+  }
+}
+
+timeval time_start, time_end;
+unsigned int totalKernelTime;
+
+int main() {
+  int num_pixel_ = 7;
+  int num_color_ = 256;
+  uint32_t *pixels_ = new uint32_t[num_pixel_ * 2];
+  unsigned int seed = 42;
+  for (uint32_t i = 0; i < num_pixel_ * 2; i++) {
+    pixels_[i] = rand_r(&seed) % num_color_;
+  }
+
+  uint32_t *opt_histogram_ = new uint32_t[num_color_]();
+  uint32_t *no_opt_histogram_ = new uint32_t[num_color_]();
+
+  uint32_t *d_histogram;
+  uint32_t *d_pixels;
+  cudaMalloc(&d_pixels, num_pixel_ * 2 * sizeof(uint32_t));
+  cudaMemcpy(d_pixels, pixels_, num_pixel_ * 2 * sizeof(uint32_t),
+             cudaMemcpyHostToDevice);
+  // optimized
+  cudaMalloc(&d_histogram, num_color_ * sizeof(uint32_t));
+  opt_Histogram<<<8192 / 64, 64>>>(d_pixels, d_histogram, num_color_,
+                                   num_pixel_);
+  cudaDeviceSynchronize();
+  cudaMemcpy(opt_histogram_, d_histogram, num_color_ * sizeof(uint32_t),
+             cudaMemcpyDeviceToHost);
+
+  // unoptimized
+  cudaMalloc(&d_histogram, num_color_ * sizeof(uint32_t));
+  Histogram<<<8192 / 64, 64>>>(d_pixels, d_histogram, num_color_, num_pixel_);
+  cudaDeviceSynchronize();
+  cudaMemcpy(no_opt_histogram_, d_histogram, num_color_ * sizeof(uint32_t),
+             cudaMemcpyDeviceToHost);
+  for (int i = 0; i < num_color_; i++) {
+    if (no_opt_histogram_[i] != opt_histogram_[i]) {
+      printf("Error!\n");
+      printf("%d %d\n", no_opt_histogram_[i], opt_histogram_[i]);
+      exit(1);
+    }
+  }
+  printf("PASS\n");
+  return 0;
+}

MICRO_AE_example/anti-coalescing-transformation/run.sh

@@ -0,0 +1,17 @@
+#!/bin/bash
+set -e
+clang++ -std=c++11 hist.cu --cuda-path=$CUDA_PATH \
+    --cuda-gpu-arch=sm_50 -L$CUDA_PATH/lib64 \
+    -lcudart_static -ldl -lrt -pthread -save-temps -v
+
+kernelTranslator hist-cuda-nvptx64-nvidia-cuda-sm_50.bc hist-host-x86_64-unknown-linux-gnu.bc kernel.bc
+hostTranslator hist-host-x86_64-unknown-linux-gnu.bc host.bc
+
+llc --relocation-model=pic --filetype=obj  kernel.bc
+llc --relocation-model=pic --filetype=obj  host.bc
+
+g++ -o hist -fPIC -no-pie -L$CuPBoP_BUILD_PATH/runtime \
+  -L$CuPBoP_BUILD_PATH/runtime/threadPool \
+  host.o kernel.o -lpthread -lc -lx86Runtime -lthreadPool
+
+./hist

MICRO_AE_example/block-size-invariant-analysis/run.sh

@@ -0,0 +1,17 @@
+#!/bin/bash
+set -e
+clang++ -std=c++11 vecadd.cu --cuda-path=$CUDA_PATH \
+    --cuda-gpu-arch=sm_50 -L$CUDA_PATH/lib64 \
+    -lcudart_static -ldl -lrt -pthread -save-temps -v
+
+kernelTranslator vecadd-cuda-nvptx64-nvidia-cuda-sm_50.bc vecadd-host-x86_64-unknown-linux-gnu.bc kernel.bc
+hostTranslator vecadd-host-x86_64-unknown-linux-gnu.bc host.bc
+
+llc --relocation-model=pic --filetype=obj  kernel.bc
+llc --relocation-model=pic --filetype=obj  host.bc
+
+g++ -o vecadd -fPIC -no-pie -L$CuPBoP_BUILD_PATH/runtime \
+  -L$CuPBoP_BUILD_PATH/runtime/threadPool \
+  host.o kernel.o -lpthread -lc -lx86Runtime -lthreadPool
+
+./vecadd

MICRO_AE_example/block-size-invariant-analysis/vecadd.cu

@@ -0,0 +1,50 @@
+#include <cstdio>
+#include <random>
+#include <stdint.h>
+#include <sys/time.h>
+
+__global__ void vecadd(int *a, int *b, int *c, int n) {
+  int tid = blockIdx.x * blockDim.x + threadIdx.x;
+  if (tid < n) {
+    c[tid] += a[tid] + b[tid];
+  }
+}
+
+int main() {
+  int size = 512;
+  int *h_a, *h_b, *h_c;
+  h_a = (int *)malloc(size * sizeof(int));
+  h_b = (int *)malloc(size * sizeof(int));
+  h_c = (int *)malloc(size * sizeof(int));
+
+  for (int i = 0; i < size; i++) {
+    h_a[i] = i;
+    h_b[i] = 2 * i;
+    h_c[i] = 0;
+  }
+  int *d_a, *d_b, *d_c;
+  cudaMalloc(&d_a, size * sizeof(int));
+  cudaMalloc(&d_b, size * sizeof(int));
+  cudaMalloc(&d_c, size * sizeof(int));
+
+  cudaMemcpy(d_a, h_a, size * sizeof(int), cudaMemcpyHostToDevice);
+  cudaMemcpy(d_b, h_b, size * sizeof(int), cudaMemcpyHostToDevice);
+  cudaMemcpy(d_c, h_c, size * sizeof(int), cudaMemcpyHostToDevice);
+
+  // Launch with different block size
+  vecadd<<<std::ceil(size / 16.0), 16>>>(d_a, d_b, d_c, size);
+  vecadd<<<std::ceil(size / 32.0), 32>>>(d_a, d_b, d_c, size);
+  vecadd<<<std::ceil(size / 42.0), 42>>>(d_a, d_b, d_c, size);
+
+  cudaMemcpy(h_c, d_c, size * sizeof(int), cudaMemcpyDeviceToHost);
+
+  // Verify the result
+  for (int i = 0; i < size; i++) {
+    if (h_c[i] != 3 * (h_a[i] + h_b[i])) {
+      printf("Error at index %d\n", i);
+      return 1;
+    }
+  }
+  printf("PASS\n");
+  return 0;
+}

MICRO_AE_example/tail-block-adaptive-synchronization/run.sh

@@ -0,0 +1,17 @@
+#!/bin/bash
+set -e
+clang++ -std=c++11 vecadd.cu --cuda-path=$CUDA_PATH \
+    --cuda-gpu-arch=sm_50 -L$CUDA_PATH/lib64 \
+    -lcudart_static -ldl -lrt -pthread -save-temps -v
+
+kernelTranslator vecadd-cuda-nvptx64-nvidia-cuda-sm_50.bc vecadd-host-x86_64-unknown-linux-gnu.bc kernel.bc
+hostTranslator vecadd-host-x86_64-unknown-linux-gnu.bc host.bc
+
+llc --relocation-model=pic --filetype=obj  kernel.bc
+llc --relocation-model=pic --filetype=obj  host.bc
+
+g++ -o vecadd -fPIC -no-pie -L$CuPBoP_BUILD_PATH/runtime \
+  -L$CuPBoP_BUILD_PATH/runtime/threadPool \
+  host.o kernel.o -lpthread -lc -lx86Runtime -lthreadPool
+
+./vecadd

MICRO_AE_example/tail-block-adaptive-synchronization/vecadd.cu

@@ -0,0 +1,47 @@
+#include <cstdio>
+#include <random>
+#include <stdint.h>
+#include <sys/time.h>
+
+const int SIZE = 512;
+__global__ void vecadd(int *a, int *b, int *c) {
+  int tid = blockIdx.x * blockDim.x + threadIdx.x;
+  if (tid < SIZE) {
+    c[tid] += a[tid] + b[tid];
+  }
+}
+
+int main() {
+  int *h_a, *h_b, *h_c;
+  h_a = (int *)malloc(SIZE * sizeof(int));
+  h_b = (int *)malloc(SIZE * sizeof(int));
+  h_c = (int *)malloc(SIZE * sizeof(int));
+
+  for (int i = 0; i < SIZE; i++) {
+    h_a[i] = i;
+    h_b[i] = 2 * i;
+    h_c[i] = 0;
+  }
+  int *d_a, *d_b, *d_c;
+  cudaMalloc(&d_a, SIZE * sizeof(int));
+  cudaMalloc(&d_b, SIZE * sizeof(int));
+  cudaMalloc(&d_c, SIZE * sizeof(int));
+
+  cudaMemcpy(d_a, h_a, SIZE * sizeof(int), cudaMemcpyHostToDevice);
+  cudaMemcpy(d_b, h_b, SIZE * sizeof(int), cudaMemcpyHostToDevice);
+  cudaMemcpy(d_c, h_c, SIZE * sizeof(int), cudaMemcpyHostToDevice);
+
+  vecadd<<<SIZE/16, 16>>>(d_a, d_b, d_c);
+
+  cudaMemcpy(h_c, d_c, SIZE * sizeof(int), cudaMemcpyDeviceToHost);
+
+  // Verify the result
+  for (int i = 0; i < SIZE; i++) {
+    if (h_c[i] != (h_a[i] + h_b[i])) {
+      printf("Error at index %d\n", i);
+      return 1;
+    }
+  }
+  printf("PASS\n");
+  return 0;
+}

README.md

@@ -104,3 +104,19 @@ papers:
 - Haotian Sheng
 - Blaise Tine
 - [Hyesoon Kim](https://faculty.cc.gatech.edu/~hyesoon/)
+
+## Acknowledgements
+
+- [POCL](http://portablecl.org/) is an open-source
+OpenCL implementations that based on LLVM.
+We reuse some code from it
+(e.g., apply optimizations, load/store LLVM IRs).
+- [Hetero-Mark](https://github.com/NUCAR-DEV/Hetero-Mark)
+and [Rodinia Benchmark](https://github.com/yuhc/gpu-rodinia)
+are two benchmark suites
+for heterogeneous system computation.
+CuPBoP uses them as integrated test to verify the correctness.
+- [moodycamel::ConcurrentQueue](<https://github.com/cameron314/concurrentqueue/tree/master>)
+is a fast multi-producer,
+multi-consumer lock-free concurrent queue for C++11.
+CuPBoP uses it as the task queue for launching and executing kernels.

compilation/CMakeLists.txt

@@ -1,7 +1,7 @@
 cmake_minimum_required(VERSION 3.1 FATAL_ERROR)
 project(
   NVVM2X86
-  DESCRIPTION "Translate NVVM IR to LLVM IR for X86"
+  DESCRIPTION "Translate NVVM IR to LLVM IR for X86 backend"
   LANGUAGES CXX)
 
 set(CMAKE_VERBOSE_MAKEFILE ON)

compilation/KernelTranslation.cpp

@@ -1,9 +1,11 @@
 #include "generate_x86_format.h"
+#include "global_mem_coalescing_optimization.h"
 #include "handle_sync.h"
 #include "init.h"
 #include "insert_sync.h"
 #include "insert_warp_loop.h"
 #include "performance.h"
+#include "tail_block_adaptive_sync_optimization.h"
 #include "tool.h"
 #include "warp_func.h"
 #include "llvm/IR/Module.h"
@@ -20,42 +22,51 @@ using namespace llvm;
 std::string PATH = "kernel_meta.log";
 
 int main(int argc, char **argv) {
-  assert(argc == 3 && "incorrect number of arguments\n");
-  llvm::Module *program = LoadModuleFromFilr(argv[1]);
+  assert(argc == 4 && "incorrect number of arguments\n");
+  const char *input_kernel_module_path = argv[1];
+  const char *input_host_module_path = argv[2];
+  const char *output_path = argv[3];
+  llvm::Module *kernel_module = LoadModuleFromFilr(input_kernel_module_path);
+  llvm::Module *host_module = LoadModuleFromFilr(input_host_module_path);
 
   std::ofstream fout;
   fout.open(PATH);
 
   // inline, and create auxiliary global variables
-  init_block(program, fout);
+  init_block(kernel_module, fout);
+
+  // Apply tail block adaptive synchronization transformation
+  tail_block_adaptive_sync_optimization(kernel_module, host_module);
+  VerifyModule(kernel_module);
   // insert sync before each vote, and replace the
   // original vote function to warp vote
-  handle_warp_vote(program);
+  handle_warp_vote(kernel_module);
 
   // replace warp shuffle
-  handle_warp_shfl(program);
+  handle_warp_shfl(kernel_module);
+
+  // global memory access coalescing optimization
+  global_mem_coalescing_optimization(kernel_module);
+  VerifyModule(kernel_module);
 
   // insert sync
-  insert_sync(program);
+  insert_sync(kernel_module);
 
   // split block by sync
-  split_block_by_sync(program);
+  split_block_by_sync(kernel_module);
   // add loop for intra&intera thread
-  insert_warp_loop(program);
+  insert_warp_loop(kernel_module);
 
-  // (TODO): replace this patch
-  replace_built_in_function(program);
+  replace_built_in_function(kernel_module);
 
-  // TODO: replace with a more general function
-  // Not only for x86 backend
-  generate_x86_format(program);
+  generate_x86_format(kernel_module, host_module);
 
   // performance optimization
-  performance_optimization(program);
+  performance_optimization(kernel_module);
 
-  VerifyModule(program);
+  VerifyModule(kernel_module);
 
-  DumpModule(program, argv[2]);
+  DumpModule(kernel_module, output_path);
 
   fout.close();
   return 0;

compilation/KernelTranslation/include/x86/generate_x86_format.h

@@ -3,7 +3,7 @@
 
 #include "llvm/IR/Module.h"
 
-void generate_x86_format(llvm::Module *M);
+void generate_x86_format(llvm::Module *M, llvm::Module *host_module);
 
 void set_meta_data(llvm::Module *M);
 

compilation/KernelTranslation/include/x86/global_mem_coalescing_optimization.h

@@ -0,0 +1,8 @@
+#ifndef __NVVM2x86_GLOBAL_MEM_COALESCING_OPTIMIZATION__
+#define __NVVM2x86_GLOBAL_MEM_COALESCING_OPTIMIZATION__
+
+#include "llvm/IR/Function.h"
+
+void global_mem_coalescing_optimization(llvm::Module *M);
+
+#endif

compilation/KernelTranslation/include/x86/performance.h

@@ -1,7 +1,11 @@
 #ifndef __NVVM2x86_PERFORMANCE__
 #define __NVVM2x86_PERFORMANCE__
 
+#include "llvm/Analysis/LoopInfo.h"
+#include "llvm/Analysis/PostDominators.h"
 #include "llvm/IR/Module.h"
 
 void performance_optimization(llvm::Module *M);
-#endif
+
+bool loop_contains_global_memory_coalescing(llvm::Loop *L);
+#endif