diff --git a/MICRO_AE_example/anti-coalescing-transformation/hist.cu b/MICRO_AE_example/anti-coalescing-transformation/hist.cu new file mode 100644 index 0000000..5518c35 --- /dev/null +++ b/MICRO_AE_example/anti-coalescing-transformation/hist.cu @@ -0,0 +1,104 @@ +#include +#include +#include +#include + +__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; +} diff --git a/MICRO_AE_example/anti-coalescing-transformation/run.sh b/MICRO_AE_example/anti-coalescing-transformation/run.sh new file mode 100644 index 0000000..a19df79 --- /dev/null +++ b/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 diff --git a/MICRO_AE_example/block-size-invariant-analysis/run.sh b/MICRO_AE_example/block-size-invariant-analysis/run.sh new file mode 100644 index 0000000..7b552df --- /dev/null +++ b/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 diff --git a/MICRO_AE_example/block-size-invariant-analysis/vecadd.cu b/MICRO_AE_example/block-size-invariant-analysis/vecadd.cu new file mode 100644 index 0000000..af5725e --- /dev/null +++ b/MICRO_AE_example/block-size-invariant-analysis/vecadd.cu @@ -0,0 +1,50 @@ +#include +#include +#include +#include + +__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<<>>(d_a, d_b, d_c, size); + vecadd<<>>(d_a, d_b, d_c, size); + vecadd<<>>(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; +} diff --git a/MICRO_AE_example/tail-block-adaptive-synchronization/run.sh b/MICRO_AE_example/tail-block-adaptive-synchronization/run.sh new file mode 100644 index 0000000..7b552df --- /dev/null +++ b/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 diff --git a/MICRO_AE_example/tail-block-adaptive-synchronization/vecadd.cu b/MICRO_AE_example/tail-block-adaptive-synchronization/vecadd.cu new file mode 100644 index 0000000..ce576b0 --- /dev/null +++ b/MICRO_AE_example/tail-block-adaptive-synchronization/vecadd.cu @@ -0,0 +1,47 @@ +#include +#include +#include +#include + +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<<>>(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; +} diff --git a/README.md b/README.md index 36e558a..f7ab2fa 100644 --- a/README.md +++ b/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]() +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. diff --git a/compilation/CMakeLists.txt b/compilation/CMakeLists.txt index e7fec26..e0f8bbc 100644 --- a/compilation/CMakeLists.txt +++ b/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) diff --git a/compilation/KernelTranslation.cpp b/compilation/KernelTranslation.cpp index 2dd318b..a8b9f10 100644 --- a/compilation/KernelTranslation.cpp +++ b/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; diff --git a/compilation/KernelTranslation/include/x86/generate_x86_format.h b/compilation/KernelTranslation/include/x86/generate_x86_format.h index 747b1b1..c568b0f 100644 --- a/compilation/KernelTranslation/include/x86/generate_x86_format.h +++ b/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); diff --git a/compilation/KernelTranslation/include/x86/global_mem_coalescing_optimization.h b/compilation/KernelTranslation/include/x86/global_mem_coalescing_optimization.h new file mode 100644 index 0000000..c2fe13e --- /dev/null +++ b/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 \ No newline at end of file diff --git a/compilation/KernelTranslation/include/x86/performance.h b/compilation/KernelTranslation/include/x86/performance.h index bb9bf88..308faec 100644 --- a/compilation/KernelTranslation/include/x86/performance.h +++ b/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 \ No newline at end of file diff --git a/compilation/KernelTranslation/include/x86/tail_block_adaptive_sync_optimization.h b/compilation/KernelTranslation/include/x86/tail_block_adaptive_sync_optimization.h new file mode 100644 index 0000000..aa7c3b8 --- /dev/null +++ b/compilation/KernelTranslation/include/x86/tail_block_adaptive_sync_optimization.h @@ -0,0 +1,32 @@ +#ifndef __NVVM2x86_TAIL_BLOCK_ADAPTIVE_SYNC_OPTIMIZATION__ +#define __NVVM2x86_TAIL_BLOCK_ADAPTIVE_SYNC_OPTIMIZATION__ + +#include "llvm/IR/Function.h" + +/* +Before transformation: +__global__ void k1(...) { + int gid = bid * bdim + tid; + // only the last block can be false + if (gid < N) + out[gid] = in[gid]; +} +After transformation: +__global__ void k1(...) { + int gid = bid * bdim + tid; + if (blockIdx.x == blockDim.x - 1) { + if (gid < N) { + out[gid] = in[gid]; + __syncthreads(); + } + else { + out[gid] = in[gid]; + __syncthreads(); + } +} +*/ + +void tail_block_adaptive_sync_optimization(llvm::Module *kernel_module, + llvm::Module *host_module); + +#endif diff --git a/compilation/KernelTranslation/include/x86/tool.h b/compilation/KernelTranslation/include/x86/tool.h index b675eed..f7c2fa0 100644 --- a/compilation/KernelTranslation/include/x86/tool.h +++ b/compilation/KernelTranslation/include/x86/tool.h @@ -5,8 +5,10 @@ #include "llvm/IR/Instructions.h" #include "llvm/IR/Module.h" -llvm::Module *LoadModuleFromFilr(char *file_name); -void DumpModule(llvm::Module *M, char *file_name); +#include + +llvm::Module *LoadModuleFromFilr(const char *file_name); +void DumpModule(llvm::Module *M, const char *file_name); bool isKernelFunction(llvm::Module *M, llvm::Function *F); void replace_block(llvm::Function *F, llvm::BasicBlock *before, llvm::BasicBlock *after); @@ -31,4 +33,37 @@ llvm::Value *createInBoundsGEP(llvm::IRBuilder<> &B, llvm::Value *ptr, llvm::ArrayRef idxlist); llvm::Value *createGEP(llvm::IRBuilder<> &B, llvm::Value *ptr, llvm::ArrayRef idxlist); + +// Used to represent the CUDA grid/block size configuration +// This is an auxiliary structure to store the possible block sizes. +struct Dim3SizeConfig { + int _x, _y, _z; + Dim3SizeConfig(int x, int y, int z) : _x(x), _y(y), _z(z) {} + // This function is used to convert the block size configuration to a + // string. + std::string toString() const { + std::ostringstream oss; + oss << _x << "_" << _y << "_" << _z; + return oss.str(); + } + bool operator==(const Dim3SizeConfig &other) const { + return (_x == other._x && _y == other._y && _z == other._z); + } + bool operator<(const Dim3SizeConfig &other) const { + if (_x != other._x) + return _x < other._x; + if (_y != other._y) + return _y < other._y; + return _z < other._z; + } +}; + +std::set +get_possible_grid_or_block_size(llvm::Module *host_module, bool getBlockSize); + +llvm::Instruction *find_thread_idx_x(llvm::BasicBlock *BB); +llvm::Instruction *find_block_idx_x(llvm::BasicBlock *BB); +llvm::Instruction *find_block_dim_x(llvm::BasicBlock *BB); +llvm::Instruction *find_grid_dim_x(llvm::BasicBlock *BB); + #endif diff --git a/compilation/KernelTranslation/src/x86/generate_x86_format.cpp b/compilation/KernelTranslation/src/x86/generate_x86_format.cpp index 23c97d0..19e7ab0 100644 --- a/compilation/KernelTranslation/src/x86/generate_x86_format.cpp +++ b/compilation/KernelTranslation/src/x86/generate_x86_format.cpp @@ -30,9 +30,20 @@ void set_meta_data(llvm::Module *M) { "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"); } -// as pthread only accept a single void* for input +// As pthread only accept a single void* for input // we have to decode this input inside the kernel -void decode_input(llvm::Module *M) { +void generate_wrapper_func(llvm::Module *M, llvm::Module *host_module) { + + std::set possible_block_size_list = + get_possible_grid_or_block_size(host_module, /*getBlockSize=*/true); + if (possible_block_size_list.size() != 0) { + printf("possible block sizes: \n"); + for (auto block_size : possible_block_size_list) { + printf("x: %d y: %d z: %d\n", block_size._x, block_size._y, + block_size._z); + } + printf("\n"); + } std::set need_remove; @@ -54,8 +65,7 @@ void decode_input(llvm::Module *M) { if (!isKernelFunction(M, F)) continue; auto func_name = F->getName().str(); - // remove mangle prefix - // remove _Z24 + // remove mangle prefix: _Z24 for (int pos = 2; pos < func_name.length(); pos++) { if (func_name[pos] >= '0' && func_name[pos] <= '9') continue; @@ -68,6 +78,7 @@ void decode_input(llvm::Module *M) { M->getOrInsertFunction(func_name + "_wrapper", LauncherFuncT); Function *WorkGroup = dyn_cast(fc.getCallee()); + // create the entry block BasicBlock *Block = BasicBlock::Create(M->getContext(), "", WorkGroup); Builder.SetInsertPoint(Block); @@ -137,8 +148,68 @@ void decode_input(llvm::Module *M) { ++idx; } - CallInst *c = Builder.CreateCall(F, ArrayRef(Arguments)); - Builder.CreateRetVoid(); + // replace the block_size to constant in the kernel functions, + // and use switch to choose to call which function + + BasicBlock *Exit_Block = BasicBlock::Create(M->getContext(), "", WorkGroup); + { + llvm::IRBuilder<> Builder3(M->getContext()); + Builder3.SetInsertPoint(Exit_Block); + Builder3.CreateRetVoid(); + } + + auto block_size_global = M->getGlobalVariable("block_size"); + auto loaded_block_size = Builder.CreateLoad( + block_size_global->getType()->getElementType(), block_size_global); + BasicBlock *default_call_block = + BasicBlock::Create(*C, "default_block_size", WorkGroup); + auto default_call_inst = llvm::CallInst::Create( + F, ArrayRef(Arguments), "", default_call_block); + llvm::BranchInst::Create(Exit_Block, default_call_block); + auto switchInst = + Builder.CreateSwitch(loaded_block_size, default_call_block); + + for (auto block_size_config : possible_block_size_list) { + // Clone a new function + ValueToValueMapTy EmptyMap; + Function *Clone = CloneFunction(F, EmptyMap); + Clone->setName(F->getName() + "_block_size_" + + block_size_config.toString()); + // replace all reference to block_size to constants + auto replace_global_variable_with_constant = + [&](llvm::GlobalVariable *global_variable, int constant) { + ConstantInt *constant_int = + dyn_cast(ConstantInt::get(Int32T, constant, true)); + std::vector Users(global_variable->user_begin(), + global_variable->user_end()); + for (auto *U : Users) { + if (auto *loadInst = llvm::dyn_cast(U)) + if (loadInst->getParent()->getParent() == Clone) { + loadInst->replaceAllUsesWith(constant_int); + } + } + }; + + int total_block_size = + block_size_config._x * block_size_config._y * block_size_config._z; + replace_global_variable_with_constant(M->getGlobalVariable("block_size"), + total_block_size); + replace_global_variable_with_constant( + M->getGlobalVariable("block_size_x"), block_size_config._x); + replace_global_variable_with_constant( + M->getGlobalVariable("block_size_y"), block_size_config._y); + replace_global_variable_with_constant( + M->getGlobalVariable("block_size_z"), block_size_config._z); + BasicBlock *possible_call_block = BasicBlock::Create( + *C, "block_size_" + block_size_config.toString(), WorkGroup); + auto call_inst = llvm::CallInst::Create( + Clone, ArrayRef(Arguments), "", possible_call_block); + auto branch_inst = + llvm::BranchInst::Create(Exit_Block, possible_call_block); + switchInst->addCase(dyn_cast( + ConstantInt::get(Int32T, total_block_size, true)), + possible_call_block); + } } for (auto f : need_remove) { f->dropAllReferences(); @@ -173,12 +244,12 @@ void remove_useless_var(llvm::Module *M) { M->getGlobalVariable("inter_warp_index")->eraseFromParent(); } -void generate_x86_format(llvm::Module *M) { +void generate_x86_format(llvm::Module *M, llvm::Module *host_module) { DEBUG_INFO("generate x86 format\n"); // change metadata set_meta_data(M); - // decode argument - decode_input(M); + // generate wrapper func for the kernel functions + generate_wrapper_func(M, host_module); // remove barrier remove_barrier(M); // remove useless func/variable diff --git a/compilation/KernelTranslation/src/x86/global_mem_coalescing_optimization.cpp b/compilation/KernelTranslation/src/x86/global_mem_coalescing_optimization.cpp new file mode 100644 index 0000000..7d6abff --- /dev/null +++ b/compilation/KernelTranslation/src/x86/global_mem_coalescing_optimization.cpp @@ -0,0 +1,218 @@ +#include "global_mem_coalescing_optimization.h" +#include "debug.hpp" +#include "performance.h" +#include "tool.h" +#include "llvm/Analysis/LoopInfo.h" +#include "llvm/Analysis/LoopPass.h" +#include "llvm/IR/IRBuilder.h" +#include "llvm/IR/LegacyPassManager.h" +#include "llvm/InitializePasses.h" +#include "llvm/PassInfo.h" +#include "llvm/PassRegistry.h" +#include "llvm/Transforms/Utils/BasicBlockUtils.h" +using namespace llvm; + +// insert a barrier before the for-loop which contains global memory coalescing +struct MemCoalescingTransformation : public llvm::FunctionPass { + +public: + static char ID; + + MemCoalescingTransformation() : FunctionPass(ID) {} + + void getAnalysisUsage(AnalysisUsage &AU) const { + AU.addRequiredTransitive(); + } + + bool runOnFunction(Function &F) { + auto M = F.getParent(); + if (!isKernelFunction(M, &F)) + return 0; + // check whether this loop has barrier + // if the loop contains barrier, we do not need to implement optimizations + // for memory coalescing + std::set mem_coalescing_loop; + LoopInfo &LI = getAnalysis().getLoopInfo(); + SmallVector LoopStack(LI.begin(), LI.end()); + for (auto L : LoopStack) { + bool contains_barrier = 0; + for (Loop::block_iterator i = L->block_begin(), e = L->block_end(); + i != e; ++i) { + for (BasicBlock::iterator j = (*i)->begin(), e = (*i)->end(); j != e; + ++j) { + if (auto Call = dyn_cast(j)) { + if (Call->isInlineAsm()) + continue; + auto func_name = Call->getCalledFunction()->getName().str(); + if (func_name == "llvm.nvvm.barrier0" || + func_name == "llvm.nvvm.bar.warp.sync" || + func_name == "llvm.nvvm.barrier.sync") { + contains_barrier = true; + break; + } + } + } + } + if (contains_barrier) + continue; + + // check whether this loop contains global memory coalescing + if (!loop_contains_global_memory_coalescing(L)) + continue; + if (!L->getExitBlock()) + continue; + mem_coalescing_loop.insert(L); + } + // implement transformation + for (auto L : mem_coalescing_loop) { + LLVMContext &context = M->getContext(); + auto I8Ptr = llvm::Type::getInt8PtrTy(context); + auto I8 = llvm::Type::getInt8Ty(context); + IRBuilder<> builder(context); + // create basic blocks + BasicBlock *do_while_latch = BasicBlock::Create( + context, "do_while_latch", L->getHeader()->getParent()); + BasicBlock *do_while_preheader = BasicBlock::Create( + context, "do_while_preheader", L->getHeader()->getParent()); + BasicBlock *do_while_header = BasicBlock::Create( + context, "do_while_header", L->getHeader()->getParent()); + // create do while latch + builder.SetInsertPoint(do_while_latch); + llvm::Instruction *has_activated_thread = createLoad( + builder, M->getGlobalVariable("has_activated_thread_addr")); + auto branch_var = builder.CreateICmpNE(has_activated_thread, + ConstantInt::get(I8, 0, true)); + builder.CreateCondBr(branch_var, do_while_header, L->getExitBlock()); + // Part0: change the loop preheader, to jump to do while preheader + auto preheader_br = + dyn_cast(L->getLoopPreheader()->getTerminator()); + assert(preheader_br->isUnconditional()); + preheader_br->setSuccessor(0, do_while_preheader); + // Part1: do while preheader + // create a variable thread_activated, to record whether a thread is + // activated or not + builder.SetInsertPoint(do_while_preheader); + llvm::Instruction *thread_activated_addr = + builder.CreateAlloca(I8, 0, "thread_activated_addr"); + builder.CreateStore(ConstantInt::get(I8, 1, true), thread_activated_addr); + builder.CreateBr(do_while_header); + // Part2: do while header + // set has_activated_thread to false + builder.SetInsertPoint(do_while_header); + Instruction *last_inst = builder.CreateStore( + ConstantInt::get(I8, 0), + M->getGlobalVariable("has_activated_thread_addr")); + // get condition ins the origial loop + // (TODO): currently, we assume the loop header contains/calculates the + // loop condition, and the cond is the instruction before condition + llvm::BasicBlock *loop_body; + llvm::BasicBlock *loop_cond = L->getHeader(); + for (auto b_iter = loop_cond->begin(); b_iter != loop_cond->end(); + ++b_iter) { + Instruction *org_inst = dyn_cast(&*b_iter); + if (isa(org_inst)) { + auto br = dyn_cast(org_inst); + loop_body = br->getSuccessor(0); + break; + } + auto new_inst = org_inst->clone(); + new_inst->insertAfter(last_inst); + org_inst->replaceAllUsesWith(new_inst); + last_inst = new_inst; + } + if (!isa(last_inst)) + return 0; + CreateInterWarpBarrier(last_inst); + // set thread_activated = thread_activated & (cond) + auto thread_activated = new LoadInst(I8, thread_activated_addr, + "thread_activated", do_while_header); + last_inst = llvm::CastInst::CreateIntegerCast(last_inst, I8, false, "", + do_while_header); + auto and_result = BinaryOperator::Create( + Instruction::And, last_inst, thread_activated, "", do_while_header); + new StoreInst(and_result, thread_activated_addr, do_while_header); + // has_activated_thread |= and_result + has_activated_thread = + new LoadInst(I8, M->getGlobalVariable("has_activated_thread_addr"), + "has_activated_thread", do_while_header); + auto or_result = + BinaryOperator::Create(Instruction::Or, has_activated_thread, + and_result, "", do_while_header); + new StoreInst(or_result, + M->getGlobalVariable("has_activated_thread_addr"), + do_while_header); + // create branch + auto branch_res = + new ICmpInst(*do_while_header, llvm::CmpInst::Predicate::ICMP_NE, + and_result, ConstantInt::get(I8, 0)); + llvm::BranchInst::Create(loop_body, do_while_latch, branch_res, + do_while_header); + // Part3: replace original loop's latches target + if (auto latch = L->getLoopLatch()) { + auto t = dyn_cast(latch->getTerminator()); + assert(t->isUnconditional()); + t->setSuccessor(0, do_while_latch); + } + // Part4: remove useless block: loop header, as it has been copied to + // do_while_header + DeleteDeadBlocks(loop_cond); + } + return 1; + } +}; + +char MemCoalescingTransformation::ID = 0; + +namespace { +static RegisterPass insert_mem_coalescing_barrier( + "mem-coalescing-opt", + "Apply transformation on CUDA memory coalescing pattern"); +} // namespace + +/* + *The CUDA global memory coalescing optimization will result to low cache hit + *rate on CPU. Thus, we need to implement transformation. + * For example: + * Input CUDA: + * uint32_t index = tid; + * while (index < num_pixels) { + * uint32_t color = pixels[index]; + * priv_hist[color]++; + * index += gsize; + * } + * Output CUDA: + * uint32_t index = tid; + * __shared__ bool has_activated_thread; + * bool thread_activated = true; + * do { + * has_activated_thread = false; + * __syncthreads(); + * thread_activated = thread_activated & (index < num_pixels); + * has_activated_thread |= thread_activated; + * if (thread_activated) { + * uint32_t color = pixels[index]; + * priv_hist[color]++; + * index += gsize; + * } + * } while (has_activated_thread); + * __syncthreads(); + */ +void global_mem_coalescing_optimization(llvm::Module *M) { + DEBUG_INFO("global memory coalescing optimization\n"); + auto Registry = PassRegistry::getPassRegistry(); + + llvm::legacy::PassManager Passes; + + std::vector passes; + passes.push_back("mem-coalescing-opt"); + for (auto pass : passes) { + const PassInfo *PIs = Registry->getPassInfo(StringRef(pass)); + if (PIs) { + Pass *thispass = PIs->createPass(); + Passes.add(thispass); + } else { + assert(0 && "Pass not found\n"); + } + } + Passes.run(*M); +} diff --git a/compilation/KernelTranslation/src/x86/init.cpp b/compilation/KernelTranslation/src/x86/init.cpp index 049d4b2..7e313f3 100644 --- a/compilation/KernelTranslation/src/x86/init.cpp +++ b/compilation/KernelTranslation/src/x86/init.cpp @@ -167,6 +167,11 @@ void create_global_variable(llvm::Module *M) { *M, VoteArrayType, false, llvm::GlobalValue::ExternalLinkage, NULL, "warp_vote", NULL, llvm::GlobalValue::GeneralDynamicTLSModel, 0, false); warp_vote->setAlignment(llvm::MaybeAlign(32)); + // TLS variable used for global memory coalescing optimization + new llvm::GlobalVariable(*M, I8, false, llvm::GlobalValue::ExternalLinkage, + llvm::ConstantInt::get(I8, 0, true), + "has_activated_thread_addr", NULL, + llvm::GlobalValue::GeneralDynamicTLSModel, 0, false); } void remove_metadata(llvm::Module *M) { @@ -302,7 +307,6 @@ bool lower_constant_expr(llvm::Module *M) { } void replace_cuda_math_built_in(llvm::Module *M) { - // replace _ZL3expd, just delete its body for (Module::iterator i = M->begin(), e = M->end(); i != e; ++i) { Function *F = &(*i); auto func_name = F->getName().str(); diff --git a/compilation/KernelTranslation/src/x86/insert_warp_loop.cpp b/compilation/KernelTranslation/src/x86/insert_warp_loop.cpp index 0819dcf..bdb2399 100644 --- a/compilation/KernelTranslation/src/x86/insert_warp_loop.cpp +++ b/compilation/KernelTranslation/src/x86/insert_warp_loop.cpp @@ -319,6 +319,7 @@ void handle_local_variable_intra_warp(std::vector PRs, // we should handle allocation generated by PHI { std::vector instruction_to_fix; + std::vector instruction_to_move; auto F = PRs[0].start_block->getParent(); for (auto bb = F->begin(); bb != F->end(); bb++) { for (auto ii = bb->begin(); ii != bb->end(); ii++) { @@ -343,6 +344,7 @@ void handle_local_variable_intra_warp(std::vector PRs, } } if (allStoreNonDivergence) { + instruction_to_move.push_back(alloc); continue; } // Do not duplicate var used outside PRs @@ -366,6 +368,8 @@ void handle_local_variable_intra_warp(std::vector PRs, } if (!used_in_non_PR) { instruction_to_fix.push_back(alloc); + } else { + instruction_to_move.push_back(alloc); } } } @@ -373,6 +377,18 @@ void handle_local_variable_intra_warp(std::vector PRs, for (auto inst : instruction_to_fix) { AddContextSaveRestore(inst, intra_warp_loop); } + for (auto alloc : instruction_to_move) { + // need to move all allocInst to the entry basic block + IRBuilder<> builder(&*(alloc->getParent() + ->getParent() + ->getEntryBlock() + .getFirstInsertionPt())); + auto newAllocInst = + builder.CreateAlloca(alloc->getType()->getElementType(), + alloc->getArraySize(), alloc->getName()); + alloc->replaceAllUsesWith(newAllocInst); + alloc->removeFromParent(); + } } for (auto parallel_regions : PRs) { @@ -681,8 +697,6 @@ class InsertWarpLoopPass : public llvm::FunctionPass { is_single_conditional_branch_block = 1; } else { // generate by replicate local variable - printf( - "[WARNING] match single conditional branch with HARD CODE\n"); bool branch_to_intra_init = false; for (unsigned suc = 0; suc < br->getNumSuccessors(); ++suc) { llvm::BasicBlock *entryCandidate = br->getSuccessor(suc); diff --git a/compilation/KernelTranslation/src/x86/performance.cpp b/compilation/KernelTranslation/src/x86/performance.cpp index 2d79384..6a95239 100644 --- a/compilation/KernelTranslation/src/x86/performance.cpp +++ b/compilation/KernelTranslation/src/x86/performance.cpp @@ -6,7 +6,6 @@ #include "llvm/ADT/Triple.h" #include "llvm/Analysis/LoopInfo.h" #include "llvm/Analysis/LoopPass.h" -#include "llvm/Analysis/PostDominators.h" #include "llvm/Analysis/TargetLibraryInfo.h" #include "llvm/Analysis/TargetTransformInfo.h" #include "llvm/CodeGen/MachineModuleInfo.h" @@ -41,8 +40,66 @@ using namespace llvm; +void eliminate_redudant_threadIdx_computation(llvm::Module *M) { + for (Module::iterator f = M->begin(); f != M->end(); ++f) { + Function *F = &(*f); + if (!isKernelFunction(M, F)) + continue; + // check whether we have inter_warp + bool has_inter_warp = false; + for (Function::iterator b = F->begin(); b != F->end(); ++b) { + BasicBlock *B = &(*b); + if (B->getName().startswith("inter_warp_cond")) { + has_inter_warp = true; + break; + } + } + if (!has_inter_warp) { + // since there is not inter_warp, the inter_warp_index is always 0, + // threadIdx = 32*inter_warp_index+intra_warp_index --> threadIdx = + // intra_warp_index + for (Function::iterator b = F->begin(); b != F->end(); ++b) { + BasicBlock *B = &(*b); + LoadInst *load_intra_warp_index = NULL; + // if there is store to local_intra_warp_idx, it is not safe to replace + bool safe_to_replace = true; + std::set need_to_replace; + for (BasicBlock::iterator i = B->begin(); i != B->end(); ++i) { + Instruction *inst = &(*i); + if (auto loadInst = dyn_cast(inst)) { + if (loadInst->getOperand(0)->getName().equals( + "local_intra_warp_idx") && + !load_intra_warp_index) + load_intra_warp_index = loadInst; + } else if (auto storeInst = dyn_cast(inst)) { + if (storeInst->getOperand(1)->getName().equals( + "local_intra_warp_idx")) { + safe_to_replace = false; + break; + } + } else if (auto binaryInst = dyn_cast(inst)) { + if (binaryInst->getName().startswith("thread_idx") && + binaryInst->getOpcode() == Instruction::Add) { + need_to_replace.insert(binaryInst); + } + } + } + if (safe_to_replace) { + for (auto inst : need_to_replace) { + inst->replaceAllUsesWith(load_intra_warp_index); + inst->eraseFromParent(); + } + } + } + } + } +} + void performance_optimization(llvm::Module *M) { DEBUG_INFO("performance optimization\n"); + // remove useless load + eliminate_redudant_threadIdx_computation(M); + // we assume no alias in arguments for (auto F = M->begin(); F != M->end(); F++) { for (auto I = F->arg_begin(); I != F->arg_end(); ++I) { if (I->getType()->isPointerTy()) { @@ -62,7 +119,7 @@ void performance_optimization(llvm::Module *M) { Options.FloatABIType = FloatABI::Hard; TargetMachine *TM = TheTarget->createTargetMachine( - triple.getTriple(), llvm::sys::getHostCPUName().str(), StringRef("+m,+f"), + triple.getTriple(), llvm::sys::getHostCPUName().str(), StringRef(""), Options, Reloc::PIC_, CodeModel::Small, CodeGenOpt::Aggressive); assert(TM && "No Machine Information\n"); @@ -80,9 +137,132 @@ void performance_optimization(llvm::Module *M) { Builder.LoopVectorize = true; Builder.SLPVectorize = true; - Builder.VerifyInput = true; - Builder.VerifyOutput = true; + // Builder.VerifyInput = true; + // Builder.VerifyOutput = true; Builder.populateModulePassManager(Passes); Passes.run(*M); } + +/* + * Check whether the value of inst is linear with blockDim + */ +bool linear_with_blockDim(llvm::Instruction *inst, + std::set visited) { + + bool result = false; + if (llvm::CallBase *callInst = llvm::dyn_cast(inst)) { + if (Function *calledFunction = callInst->getCalledFunction()) { + if (calledFunction->getName().startswith( + "llvm.nvvm.read.ptx.sreg.ntid")) { + result = true; + } + } else { + result = false; + } + } else if (auto binOp = llvm::dyn_cast(inst)) { + if (auto lhs = dyn_cast(binOp->getOperand(0))) + if (visited.find(lhs) == visited.end()) { + visited.insert(lhs); + result |= linear_with_blockDim(lhs, visited); + } + + if (auto rhs = dyn_cast(binOp->getOperand(1))) + if (visited.find(rhs) == visited.end()) { + visited.insert(rhs); + result |= linear_with_blockDim(rhs, visited); + } + } else if (auto loadInst = llvm::dyn_cast(inst)) { + // find all store that related to this load + auto address = loadInst->getOperand(0); + bool all_linear_with_blockDim = true; + for (auto U : address->users()) { + if (auto store = dyn_cast(U)) { + all_linear_with_blockDim &= linear_with_blockDim(store, visited); + } else if (!isa(U) && U != loadInst) { + all_linear_with_blockDim = false; + } + } + result = all_linear_with_blockDim; + } else if (auto storeInst = llvm::dyn_cast(inst)) { + if (isa(storeInst->getOperand(0))) + result = linear_with_blockDim( + dyn_cast(storeInst->getOperand(0)), visited); + } + return result; +} +/* + * check whether a variable des is linear related with variable src + */ +bool linear_related(llvm::Instruction *src, llvm::Instruction *des) { + if (src == des) + return true; + if (auto cast = dyn_cast(des)) { + if (auto cast_var = dyn_cast(cast->getOperand(0))) + return linear_related(src, cast_var); + } else if (auto load = dyn_cast(des)) { + if (auto load_var = dyn_cast(load->getOperand(0))) + return linear_related(src, load_var); + } + return false; +} +/* + * check whether the loop contains global memory coalescing optimizations + * we identify the memory coalescing according to that: + * 1) the loop iteration variable is linearly related with blockDim; + * 2) there are memory access, where the index is: induction_variable + + * threadIdx and the index is the -1 dimension; + */ +bool loop_contains_global_memory_coalescing(llvm::Loop *L) { + printf("check loop with header: %s\n", + L->getHeader()->getName().str().c_str()); + // find iteration variable + auto loop_latch = L->getLoopLatch(); + auto F = loop_latch->getParent(); + if (!loop_latch) { + return false; + } + llvm::Instruction *iteration_var = NULL; + llvm::Instruction *inc_inst = NULL; + for (BasicBlock::reverse_iterator i = loop_latch->rbegin(), + e = loop_latch->rend(); + i != e; ++i) { + if (auto Store = dyn_cast(&*i)) { + if (isa(Store->getOperand(1))) { + iteration_var = dyn_cast(Store->getOperand(1)); + inc_inst = dyn_cast(Store->getOperand(0)); + + break; + } + } + } + if (!iteration_var) { + // cannot find iteration variable + exit(1); + } + + // check whether the stride is a linear function of blockDim + std::set visited; + if (!linear_with_blockDim(inc_inst, visited)) { + // the stride is not linear with blockDim + printf("not linear with block\n"); + return false; + } + // check whether the loop contains GEP insturction, with iteration_var as the + // last dimension + for (Loop::block_iterator i = L->block_begin(), e = L->block_end(); i != e; + ++i) { + for (BasicBlock::iterator j = (*i)->begin(), e = (*i)->end(); j != e; ++j) { + if (auto GEP = dyn_cast(j)) { + auto last_dim = GEP->getOperand(GEP->getNumIndices()); + if (auto last_dim_var = dyn_cast(last_dim)) { + if (linear_related(iteration_var, last_dim_var)) { + printf("Find global memory coalescing\n"); + return true; + } + } + } + } + } + return false; +} diff --git a/compilation/KernelTranslation/src/x86/tail_block_adaptive_sync.cpp b/compilation/KernelTranslation/src/x86/tail_block_adaptive_sync.cpp new file mode 100644 index 0000000..8bce89a --- /dev/null +++ b/compilation/KernelTranslation/src/x86/tail_block_adaptive_sync.cpp @@ -0,0 +1,243 @@ +#include "debug.hpp" +#include "global_mem_coalescing_optimization.h" +#include "performance.h" +#include "tool.h" +#include "llvm/Analysis/LoopInfo.h" +#include "llvm/Analysis/LoopPass.h" +#include "llvm/Analysis/ScalarEvolution.h" +#include "llvm/Analysis/ScalarEvolutionExpressions.h" +#include "llvm/IR/Function.h" +#include "llvm/IR/IRBuilder.h" +#include "llvm/IR/Instructions.h" +#include "llvm/IR/LegacyPassManager.h" +#include "llvm/IR/Module.h" +#include "llvm/IR/PassManager.h" +#include "llvm/InitializePasses.h" +#include "llvm/PassInfo.h" +#include "llvm/PassRegistry.h" +#include "llvm/Passes/PassBuilder.h" +#include "llvm/Passes/PassPlugin.h" +#include "llvm/Support/CommandLine.h" +#include "llvm/Transforms/Scalar/IndVarSimplify.h" +#include "llvm/Transforms/Utils/BasicBlockUtils.h" +#include "llvm/Transforms/Utils/Cloning.h" +using namespace llvm; + +std::set possible_block_size_list; +std::set possible_grid_size_list; + +int tail_block_threshold; + +struct TailBlockSyncTransformation : public llvm::FunctionPass { + +public: + static char ID; + + TailBlockSyncTransformation() : FunctionPass(ID) {} + + void getAnalysisUsage(AnalysisUsage &AU) const { + AU.addRequired(); + AU.addRequired(); + AU.addRequired(); + } + + void apply_tail_block_adaptive_sync(Function *F, DominatorTree *DT, + LoopInfo *LI) { + // Step 1: Split the header by the condition + BasicBlock *header_block = &*F->begin(); + Instruction *thread_idx_x = find_thread_idx_x(header_block); + Instruction *block_idx_x = find_block_idx_x(header_block); + Instruction *block_dim_x = find_block_dim_x(header_block); + Instruction *grid_dim_x = find_grid_dim_x(header_block); + + BasicBlock *body_block = &*std::next(F->begin()); + BasicBlock *ret_block = &*std::next(std::next(F->begin())); + + Instruction *splitPoint = header_block->getTerminator(); + BasicBlock *tail_block_start = + SplitBlock(header_block, splitPoint, DT, LI, nullptr, "tail_block"); + CreateInterWarpBarrier(tail_block_start->getTerminator()); + // Step 2: Construct new blocks for non-tail GPU blocks + ValueToValueMapTy VMap; + BasicBlock *non_tail_block_start = + CloneBasicBlock(body_block, VMap, "non_tail_block", F); + + for (Instruction &I : *non_tail_block_start) { + RemapInstruction(&I, VMap, + RF_NoModuleLevelChanges | RF_IgnoreMissingLocals); + } + CreateInterWarpBarrier(non_tail_block_start->getTerminator()); + // Step 3: Build a new condition instruction at end of the header block + IRBuilder<> Builder(header_block->getTerminator()); + + // Check whether the block is the last block + Value *is_tail_block = Builder.CreateCmp( + CmpInst::ICMP_EQ, block_idx_x, + Builder.CreateSub(block_dim_x, + ConstantInt::get(block_dim_x->getType(), 1))); + Instruction *new_branch_for_header = Builder.CreateCondBr( + is_tail_block, tail_block_start, non_tail_block_start); + + // Remove the original unconditional branch + Instruction *unconditional_branch = header_block->getTerminator(); + unconditional_branch->eraseFromParent(); + } + + bool isSingleConditionFormat(Function *F) { + // The idle format should be like: + // void func() { + // int gid = blockIdx.x * blockDim.x + threadIdx.x; + // if (gid < N) { + // // kernel body + // } + // } + //} + // + // check the length of blocks + if (F->size() != 3) { + return false; + } + auto header_block = F->begin(); + auto ret_block = std::next(std::next(F->begin())); + // Step1: Check the return block + if (ret_block->size() != 1) { + // The return block should only contains a single ret instruction + return false; + } + // Step2: Check the header block + auto branch_inst = dyn_cast(header_block->getTerminator()); + if (!branch_inst || branch_inst->isUnconditional()) { + return false; + } + auto cmp_inst = dyn_cast(branch_inst->getCondition()); + if (!cmp_inst || cmp_inst->getPredicate() != CmpInst::ICMP_SLT) { + return false; + } + // The left operand should be a variable + auto lhr = dyn_cast(cmp_inst->getOperand(0)); + // The right operand should be a constant + auto rhr = dyn_cast(cmp_inst->getOperand(1)); + if (!lhr || !rhr) { + return false; + } + tail_block_threshold = rhr->getSExtValue(); + // Check the left operand, it should be blockIdx.x * blockDim.x + tid.x + for (auto inst : lhr->getPointerOperand()->users()) { + if (auto store_inst = dyn_cast(inst)) { + auto add_inst = dyn_cast(store_inst->getValueOperand()); + if (!add_inst || add_inst->getOpcode() != Instruction::Add) { + return false; + } + // The second operand should be threadIdx.x + auto tid_inst = dyn_cast(add_inst->getOperand(1)); + if (!tid_inst || tid_inst->getCalledFunction()->getName() != + "llvm.nvvm.read.ptx.sreg.tid.x") { + return false; + } + // The first operand should be blockIdx.x * blockDim.x + if (auto mul_inst = dyn_cast(add_inst->getOperand(0))) { + if (mul_inst->getOpcode() == Instruction::Mul) { + auto block_idx_inst = dyn_cast(mul_inst->getOperand(0)); + auto grid_size_inst = dyn_cast(mul_inst->getOperand(1)); + + if (!block_idx_inst || !grid_size_inst) { + return false; + } + if (block_idx_inst->getCalledFunction()->getName() == + "llvm.nvvm.read.ptx.sreg.ctaid.x" && + grid_size_inst->getCalledFunction()->getName() == + "llvm.nvvm.read.ptx.sreg.ntid.x") { + return true; + } + } + } + } + } + return false; + } + bool runOnFunction(Function &F) { + auto M = F.getParent(); + if (!isKernelFunction(M, &F)) + return 0; + // Check whether the function body can be optimized. + if (!isSingleConditionFormat(&F)) { + printf("the function does not contain tail block\n"); + return 0; + } + auto &SE = getAnalysis().getSE(); + // Check whether the condition is true for all but the last blocks. + APInt condition_constant(32, tail_block_threshold); + + ConstantRange ThreadIdxRange = ConstantRange( + APInt(32, 0), APInt(32, possible_block_size_list.begin()->_x + 1)); + ConstantRange BlockIdxRange = ConstantRange( + APInt(32, 0), APInt(32, possible_grid_size_list.begin()->_x)); + + ConstantRange MulRange = BlockIdxRange.multiply(ThreadIdxRange); + + MulRange.print(llvm::outs()); + llvm::outs() << "\n"; + if (!MulRange.icmp(CmpInst::ICMP_SLT, condition_constant)) { + printf("Condition is not always true.\n"); + return 0; + } + printf("The function can be optimized by tail block adaptive sync\n"); + // Apply transformation + DominatorTree *DT = &getAnalysis().getDomTree(); + LoopInfo *LI = &getAnalysis().getLoopInfo(); + apply_tail_block_adaptive_sync(&F, DT, LI); + return 1; + } +}; + +char TailBlockSyncTransformation::ID = 0; + +namespace { +static RegisterPass tail_block_adaptive_sync( + "tail-block-adaptive-sync", + "Apply tail block adaptive synchronization transformation"); +} // namespace + +void tail_block_adaptive_sync_optimization(llvm::Module *kernel_module, + llvm::Module *host_module) { + DEBUG_INFO("tail block adaptive synchronization optimization\n"); + + // Try to analyze the host module to get the range of grid/block size + + possible_block_size_list = + get_possible_grid_or_block_size(host_module, /*getBlockSize=*/true); + possible_grid_size_list = + get_possible_grid_or_block_size(host_module, /*getBlockSize=*/false); + + // We only optimize the cases when there is single grid/block configuration + if (possible_block_size_list.size() != 1 || + possible_grid_size_list.size() != 1) { + DEBUG_INFO("Not a single static grid/block configuration. Skip.\n"); + return; + } + + // The block size and grid size should be 1 dimension + if (possible_block_size_list.begin()->_y != 1 || + possible_block_size_list.begin()->_z != 1 || + possible_grid_size_list.begin()->_y != 1 || + possible_grid_size_list.begin()->_z != 1) { + DEBUG_INFO("Not a single dimension grid/block configuration. Skip.\n"); + return; + } + auto Registry = PassRegistry::getPassRegistry(); + + llvm::legacy::PassManager Passes; + + std::vector passes; + passes.push_back("tail-block-adaptive-sync"); + for (auto pass : passes) { + const PassInfo *PIs = Registry->getPassInfo(StringRef(pass)); + if (PIs) { + Pass *thispass = PIs->createPass(); + Passes.add(thispass); + } else { + assert(0 && "Pass not found\n"); + } + } + Passes.run(*kernel_module); +} diff --git a/compilation/KernelTranslation/src/x86/tool.cpp b/compilation/KernelTranslation/src/x86/tool.cpp index d3024ac..2d08b12 100644 --- a/compilation/KernelTranslation/src/x86/tool.cpp +++ b/compilation/KernelTranslation/src/x86/tool.cpp @@ -28,7 +28,7 @@ using namespace llvm; -llvm::Module *LoadModuleFromFilr(char *file_name) { +llvm::Module *LoadModuleFromFilr(const char *file_name) { llvm::SMDiagnostic Err; llvm::LLVMContext *globalContext = new llvm::LLVMContext; auto program = parseIRFile(file_name, Err, *globalContext).release(); @@ -46,7 +46,7 @@ void VerifyModule(llvm::Module *program) { llvm::report_fatal_error(os.str().c_str()); } -void DumpModule(llvm::Module *M, char *file_name) { +void DumpModule(llvm::Module *M, const char *file_name) { std::string msg; llvm::raw_string_ostream os(msg); std::error_code EC; @@ -487,11 +487,13 @@ void replace_built_in_function(llvm::Module *M) { func_name == "__nv_fast_powf" || func_name == "__nv_powf" || func_name == "__nv_logf" || func_name == "__nv_expf" || func_name == "__nv_fabsf" || - func_name == "__nv_log10f" || + func_name == "__nv_log10f" || func_name == "__nv_log" || func_name == "__nv_fmodf" || func_name == "__nv_sqrt" || func_name == "__nv_sqrtf" || func_name == "__nv_exp" || - func_name == "__nv_isnanf" || + func_name == "__nv_floorf" || func_name == "__nv_cosf" || + func_name == "__nv_isnanf" || func_name == "__nv_sinf" || func_name == "__nv_isinff" || func_name == "__nv_powi" || + func_name == "__nvvm_fabs_f" || func_name == "__nv_powif") { Call->getCalledFunction()->deleteBody(); } else if (func_name == "llvm.nvvm.fma.rn.d") { @@ -685,3 +687,166 @@ Value *createGEP(IRBuilder<> &B, Value *ptr, ArrayRef idxlist) { return B.CreateGEP(ptr->getType()->getScalarType()->getPointerElementType(), ptr, idxlist); } + +// The following functions are used to analyze the host module and extract the +// grid/block size information. +// This function checkes the depdency between two values recursively. +// Although there is off-the-shelf def-use analysis, it cannot handle bitcast +// and memcpy. +bool isDependentRecursively(llvm::Value *V1, llvm::Value *V2, + std::set &visited) { + // We also need to handle the bit cast + if (llvm::BitCastInst *BitCastInst = llvm::dyn_cast(V1)) { + if (isDependentRecursively(BitCastInst->getOperand(0), V2, visited)) { + return true; + } + } + if (llvm::CallInst *CallInst = llvm::dyn_cast(V1)) { + if (CallInst->getCalledFunction() && + CallInst->getCalledFunction()->getName() == + "llvm.memcpy.p0i8.p0i8.i64") { + if (isDependentRecursively(CallInst->getArgOperand(0), V2, visited)) { + return true; + } + } + return false; + } + + if (visited.count(V1)) { + return false; + } + visited.insert(V1); + for (auto *User : V1->users()) { + if (User == V2) { + return true; + } + if (llvm::Instruction *Inst = llvm::dyn_cast(User)) { + if (isDependentRecursively(Inst, V2, visited)) { + return true; + } + } + } + return false; +} + +// This function is used to find the block size configuration for the +// __cudaPushCallConfiguration instruction. Specifically, this function +// accepts a grid configuration and a block configuration, and they have the +// same data type: struct dim3. Thus, we need to write a function to find the +// corresponding block configuration for a given kernel launch. This function +// is an auxiliary function for the block size invariant analysis. +llvm::CallInst * +findDim3SizeForKernelLaunch(llvm::CallInst *cudaPushCallConfiguration, + bool getBlockSize) { + // The 3rd and 4th arguments are related with block size + llvm::Value *block_arg = cudaPushCallConfiguration->getArgOperand(3); + llvm::Value *grid_arg = cudaPushCallConfiguration->getArgOperand(1); + + llvm::Value *required_arg = getBlockSize ? block_arg : grid_arg; + + for (auto &F : *cudaPushCallConfiguration->getFunction()->getParent()) { + for (auto &BB : F) { + for (auto &I : BB) { + if (auto *callInst = llvm::dyn_cast(&I)) { + llvm::Function *calledFunc = callInst->getCalledFunction(); + if (calledFunc && calledFunc->getName() == "_ZN4dim3C2Ejjj") { + std::set visited; + // In CUDA host module, all grid/block configurations are + // generated by _ZN4dim3C2Ejjj. For example: call void + // @_ZN4dim3C2Ejjj(%struct.dim3* %15, i32 42, i32 1, i32 1) will + // generate a grid/block configuration with size: 42 * 1 * 1. + if (isDependentRecursively(callInst->getArgOperand(0), required_arg, + visited)) { + return callInst; + } + } + } + } + } + } + return nullptr; +} + +// This function analyzes the kernel launch instructions in the host module, +// and get all possible block sizes. +std::set +get_possible_grid_or_block_size(llvm::Module *host_module, bool getBlockSize) { + std::set possible_size_set; + for (auto &F : *host_module) { + for (auto &BB : F) { + for (auto &I : BB) { + if (auto *CI = dyn_cast(&I)) { + if (CI->getCalledFunction()->getName().str() == + "__cudaPushCallConfiguration") { + auto dim3_size_initialization = + findDim3SizeForKernelLaunch(CI, getBlockSize); + + if (dim3_size_initialization) { + auto dim3_size_x = dyn_cast( + dim3_size_initialization->getArgOperand(1)); + auto dim3_size_y = dyn_cast( + dim3_size_initialization->getArgOperand(2)); + auto dim3_size_z = dyn_cast( + dim3_size_initialization->getArgOperand(3)); + if (dim3_size_x && dim3_size_y && dim3_size_z) { + Dim3SizeConfig dim3_size(dim3_size_x->getSExtValue(), + dim3_size_y->getSExtValue(), + dim3_size_z->getSExtValue()); + possible_size_set.insert(dim3_size); + } + } + } + } + } + } + } + return possible_size_set; +} + +Instruction *find_thread_idx_x(BasicBlock *BB) { + for (auto &I : *BB) { + if (auto *CI = dyn_cast(&I)) { + if (CI->getCalledFunction()->getName() == + "llvm.nvvm.read.ptx.sreg.tid.x") { + return CI; + } + } + } + return nullptr; +} +Instruction *find_block_idx_x(BasicBlock *BB) { + for (auto &I : *BB) { + if (auto *CI = dyn_cast(&I)) { + if (CI->getCalledFunction()->getName() == + "llvm.nvvm.read.ptx.sreg.ctaid.x") { + return CI; + } + } + } + return nullptr; +} + +Instruction *find_block_dim_x(BasicBlock *BB) { + for (auto &I : *BB) { + if (auto *CI = dyn_cast(&I)) { + if (CI->getCalledFunction()->getName() == + "llvm.nvvm.read.ptx.sreg.ntid.x") { + return CI; + } + } + } + + return nullptr; +} + +Instruction *find_grid_dim_x(BasicBlock *BB) { + for (auto &I : *BB) { + if (auto *CI = dyn_cast(&I)) { + if (CI->getCalledFunction()->getName() == + "llvm.nvvm.read.ptx.sreg.nctaid.x") { + return CI; + } + } + } + return nullptr; +} diff --git a/runtime/include/x86/cudaKernelImpl.h b/runtime/include/x86/cudaKernelImpl.h index e65d0c1..4f376e7 100644 --- a/runtime/include/x86/cudaKernelImpl.h +++ b/runtime/include/x86/cudaKernelImpl.h @@ -6,8 +6,12 @@ extern "C" { double __nv_exp(double); double __nv_sqrt(double); float __nv_sqrtf(float); +float __nvvm_fabs_f(float); float __nv_powif(float, int); float __nv_logf(float); +float __nv_floorf(float); +float __nv_sinf(float); +float __nv_cosf(float); float __nv_expf(float); float __nv_log10f(float); float __nv_fast_log2f(float); @@ -24,5 +28,6 @@ double __nv_fmaxd(double, double); int __nvvm_mul24_i(int, int); double _ZL3expd(double); double _ZL8copysigndd(double, double); +double __nv_log(double); } #endif diff --git a/runtime/src/x86/cudaKernelImpl.cpp b/runtime/src/x86/cudaKernelImpl.cpp index 36eedf1..d947595 100644 --- a/runtime/src/x86/cudaKernelImpl.cpp +++ b/runtime/src/x86/cudaKernelImpl.cpp @@ -5,9 +5,14 @@ double __nv_sqrt(double v) { return sqrt(v); } float __nv_sqrtf(float v) { return sqrt(v); } float __nv_powif(float base, int exp) { return pow(base, exp); } float __nv_logf(float v) { return logf(v); } +double __nv_log(double v) { return log(v); } float __nv_expf(float v) { return expf(v); } float __nv_log10f(float v) { return log10f(v); } float __nv_fast_log2f(float v) { return log2f(v); } +float __nv_floorf(float v) { return floor(v); } +float __nv_sinf(float v) { return sin(v); }; +float __nv_cosf(float v) { return cos(v); }; +float __nvvm_fabs_f(float v) { return abs(v); } double __nv_powi(double base, int exp) { return pow(base, exp); } float __nv_powf(float base, float exp) { return pow(base, exp); } float __nv_fast_powf(float base, float exp) { return pow(base, exp); } diff --git a/runtime/threadPool/src/x86/api.cpp b/runtime/threadPool/src/x86/api.cpp index 4a319fc..33a13d4 100644 --- a/runtime/threadPool/src/x86/api.cpp +++ b/runtime/threadPool/src/x86/api.cpp @@ -23,8 +23,6 @@ int init_device() { return C_ERROR_MEMALLOC; device->max_compute_units = std::thread::hardware_concurrency(); - DEBUG_INFO("%d concurrent threads are supported.\n", - device->max_compute_units); device_max_compute_units = device->max_compute_units; // initialize scheduler