[GPGPU][PPCGCodeGen][GPUJIT] Added AMD support to GPGPU code generation

[PhilippSchaad]

Summary

We can now use the AMDGPU back-end to generate valid AMD kernel code, which
gets linked by LLD to valid AMD executables. The OpenCL runtime can execute
said kernels. This requires the ROCm driver stack to be installed, supported
are any AMD GPUs from the Fiji family and up. The specific GPU family can be
chosen when compiling, default is Fiji.

The LLD pass is currently a bit of a hack with a temporary file, but sadly this
is the only way this is possible at the moment, since LLD is not available as a
from-code pass. The way the temp file is being handled will change, such that
the name is random/different between compilations.

[PhilippSchaad]

• Updated test cases to reflect changes and fixed formatting

[PhilippSchaad]

• Improved temporary file handling.
• Fixed error in test case

[philip.pfaffe]

I don't fully understand the reasoning behind this, and I dislike the LLD hack for a couple of reasons:

• Why invoke a program instead of linking against lld?
• What if lld is not installed? Why not default to ld?
• Generally I feel like this is a huge violation of layering.

Can you elaborate on why this is necessary?

Further, LLVM offers OS-independent abstractions for interfacing with the system in the llvm::sys and llvm::sys::fs namespaces. Please use those instead of POSIX.

[PhilippSchaad]

I agree, it is a really 'disgusting' hack at this point.

• So the reason lld is used instead of ld is simply that I do not know and have not been able to test if ld works. The AMDGPU backend documentation here specifically mentions lld, so I felt like that was the safest bet. But you are absolutely right, if ld works, that is the better default.
• As for the program call: I do not know of any other way of going about this currently. I've asked on the mailing list and was told that there is no way of reaching lld from code at this point in time.
• I will look into the elimination of POSIX specific system interactions, and instead try to adapt this hack to use llvm's own abstractions. Thank you for pointing that out, I was unaware of that!

include/polly/CodeGen/PPCGCodeGeneration.h

 // GPU mapping strategy.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef POLLY_PPCGCODEGENERATION_H
 #define POLLY_PPCGCODEGENERATION_H
 
 /// The GPU Architecture to target.
-enum GPUArch { NVPTX64, SPIR32, SPIR64 };
+enum GPUArch { NVPTX64, SPIR32, SPIR64, AMDGCN64 };
 
 /// The GPU Runtime implementation to use.
 enum GPURuntime { CUDA, OpenCL };
 
 namespace polly {
 extern bool PollyManagedMemory;
 }
 
 #endif // POLLY_PPCGCODEGENERATION_H

lib/CodeGen/PPCGCodeGeneration.cpp

 #include "ppcg/gpu.h"
 #include "ppcg/gpu_print.h"
 #include "ppcg/ppcg.h"
 #include "ppcg/schedule.h"
 }
 
 #include "llvm/Support/Debug.h"
 
+#include <fstream>
+
 using namespace polly;
 using namespace llvm;
 
 #define DEBUG_TYPE "polly-codegen-ppcg"
 
 static cl::opt<bool> DumpSchedule("polly-acc-dump-schedule",
                                   cl::desc("Dump the computed GPU Schedule"),
                                   cl::Hidden, cl::init(false), cl::ZeroOrMore,
static cl::opt<std::string> CUDALibDevice(
     cl::init("/usr/local/cuda/nvvm/libdevice/libdevice.compute_20.10.ll"),
     cl::ZeroOrMore, cl::cat(PollyCategory));
 
 static cl::opt<std::string>
     CudaVersion("polly-acc-cuda-version",
                 cl::desc("The CUDA version to compile for"), cl::Hidden,
                 cl::init("sm_30"), cl::ZeroOrMore, cl::cat(PollyCategory));
 
+static cl::opt<std::string>
+    AMDArch("polly-acc-amd-arch",
+            cl::desc("The AMD architecture version to compile for"), cl::Hidden,
+            cl::init("gfx803"), cl::ZeroOrMore, cl::cat(PollyCategory));
+
 static cl::opt<int>
     MinCompute("polly-acc-mincompute",
                cl::desc("Minimal number of compute statements to run on GPU."),
                cl::Hidden, cl::init(10 * 512 * 512));
 
 extern bool polly::PerfMonitoring;
 
 /// Return  a unique name for a Scop, which is the scop region with the
private:
   /// Create a call to synchronize Host & Device.
   /// \note
   /// This is to be used only with managed memory.
   void createCallSynchronizeDevice();
 
   /// Create a call to get a kernel from an assembly string.
   ///
   /// @param Buffer The string describing the kernel.
+  /// @param Size   The size of the string describing the kernel.
   /// @param Entry  The name of the kernel function to call.
   ///
   /// @returns A pointer to a kernel object
-  Value *createCallGetKernel(Value *Buffer, Value *Entry);
+  Value *createCallGetKernel(Value *Buffer, Value *Size, Value *Entry);
 
   /// Create a call to free a GPU kernel.
   ///
   /// @param GPUKernel THe kernel to free.
   void createCallFreeKernel(Value *GPUKernel);
 
   /// Create a call to launch a GPU kernel.
   ///
   /// @param GPUKernel  The kernel to launch.
   /// @param GridDimX   The size of the first grid dimension.
   /// @param GridDimY   The size of the second grid dimension.
   /// @param GridBlockX The size of the first block dimension.
   /// @param GridBlockY The size of the second block dimension.
   /// @param GridBlockZ The size of the third block dimension.
   /// @param Parameters A pointer to an array that contains itself pointers to
   ///                   the parameter values passed for each kernel argument.
+  /// @param CLUseLocalWorkSize A boolean dictating whether the OpenCL Runtime
+  ///                           should automatically determine the Local work
+  ///                           group size or use the provided dimensions.
   void createCallLaunchKernel(Value *GPUKernel, Value *GridDimX,
                               Value *GridDimY, Value *BlockDimX,
                               Value *BlockDimY, Value *BlockDimZ,
-                              Value *Parameters);
+                              Value *Parameters, Value *CLUseLocalWorkSize);
 };
 
 std::string GPUNodeBuilder::getKernelFuncName(int Kernel_id) {
   return "FUNC_" + S.getFunction().getName().str() + "_SCOP_" +
          std::to_string(S.getID()) + "_KERNEL_" + std::to_string(Kernel_id);
 }
 
 void GPUNodeBuilder::initializeAfterRTH() {

 }
 
 void GPUNodeBuilder::freeDeviceArrays() {
   assert(!PollyManagedMemory && "Managed memory does not use device arrays");
   for (auto &Array : DeviceAllocations)
     createCallFreeDeviceMemory(Array.second);
 }
 
-Value *GPUNodeBuilder::createCallGetKernel(Value *Buffer, Value *Entry) {
+Value *GPUNodeBuilder::createCallGetKernel(Value *Buffer, Value *Size,
+                                           Value *Entry) {
   const char *Name = "polly_getKernel";
   Module *M = Builder.GetInsertBlock()->getParent()->getParent();
   Function *F = M->getFunction(Name);
 
   // If F is not available, declare it.
   if (!F) {
     GlobalValue::LinkageTypes Linkage = Function::ExternalLinkage;
     std::vector<Type *> Args;
     Args.push_back(Builder.getInt8PtrTy());
+    Args.push_back(Builder.getInt64Ty());
     Args.push_back(Builder.getInt8PtrTy());
     FunctionType *Ty = FunctionType::get(Builder.getInt8PtrTy(), Args, false);
     F = Function::Create(Ty, Linkage, Name, M);
   }
 
-  return Builder.CreateCall(F, {Buffer, Entry});
+  return Builder.CreateCall(F, {Buffer, Size, Entry});
 }
 
 Value *GPUNodeBuilder::createCallGetDevicePtr(Value *Allocation) {
   const char *Name = "polly_getDevicePtr";
   Module *M = Builder.GetInsertBlock()->getParent()->getParent();
   Function *F = M->getFunction(Name);
 
   // If F is not available, declare it.
   if (!F) {
     GlobalValue::LinkageTypes Linkage = Function::ExternalLinkage;
     std::vector<Type *> Args;
     Args.push_back(Builder.getInt8PtrTy());
     FunctionType *Ty = FunctionType::get(Builder.getInt8PtrTy(), Args, false);
     F = Function::Create(Ty, Linkage, Name, M);
   }
 
   return Builder.CreateCall(F, {Allocation});
 }
 
 void GPUNodeBuilder::createCallLaunchKernel(Value *GPUKernel, Value *GridDimX,
                                             Value *GridDimY, Value *BlockDimX,
                                             Value *BlockDimY, Value *BlockDimZ,
-                                            Value *Parameters) {
+                                            Value *Parameters,
+                                            Value *CLUseLocalWorkSize) {
   const char *Name = "polly_launchKernel";
   Module *M = Builder.GetInsertBlock()->getParent()->getParent();
   Function *F = M->getFunction(Name);
 
   // If F is not available, declare it.
   if (!F) {
     GlobalValue::LinkageTypes Linkage = Function::ExternalLinkage;
     std::vector<Type *> Args;
     Args.push_back(Builder.getInt8PtrTy());
     Args.push_back(Builder.getInt32Ty());
     Args.push_back(Builder.getInt32Ty());
     Args.push_back(Builder.getInt32Ty());
     Args.push_back(Builder.getInt32Ty());
     Args.push_back(Builder.getInt32Ty());
     Args.push_back(Builder.getInt8PtrTy());
+    Args.push_back(Builder.getInt32Ty());
     FunctionType *Ty = FunctionType::get(Builder.getVoidTy(), Args, false);
     F = Function::Create(Ty, Linkage, Name, M);
   }
 
   Builder.CreateCall(F, {GPUKernel, GridDimX, GridDimY, BlockDimX, BlockDimY,
-                         BlockDimZ, Parameters});
+                         BlockDimZ, Parameters, CLUseLocalWorkSize});
 }
 
 void GPUNodeBuilder::createCallFreeKernel(Value *GPUKernel) {
   const char *Name = "polly_freeKernel";
   Module *M = Builder.GetInsertBlock()->getParent()->getParent();
   Function *F = M->getFunction(Name);
 
   // If F is not available, declare it.
if (!Sync) {
       FunctionType *Ty = FunctionType::get(Builder.getVoidTy(), Args, false);
       Sync = Function::Create(Ty, Linkage, SpirName, M);
       Sync->setCallingConv(CallingConv::SPIR_FUNC);
     }
     break;
   case GPUArch::NVPTX64:
     Sync = Intrinsic::getDeclaration(M, Intrinsic::nvvm_barrier0);
     break;
+  case GPUArch::AMDGCN64:
+    Sync = Intrinsic::getDeclaration(M, Intrinsic::amdgcn_s_barrier);
+    break;
   }
 
   Builder.CreateCall(Sync, {});
 }
 
 /// Collect llvm::Values referenced from @p Node
 ///
 /// This function only applies to isl_ast_nodes that are user_nodes referring
void GPUNodeBuilder::createKernel(__isl_take isl_ast_node *KernelStmt) {
   LocalArrays.clear();
 
   std::string ASMString = finalizeKernelFunction();
   Builder.SetInsertPoint(&HostInsertPoint);
   Value *Parameters = createLaunchParameters(Kernel, F, SubtreeValues);
 
   std::string Name = getKernelFuncName(Kernel->id);
   Value *KernelString = Builder.CreateGlobalStringPtr(ASMString, Name);
+  Value *KernelSize = Builder.getInt64(ASMString.length());
   Value *NameString = Builder.CreateGlobalStringPtr(Name, Name + "_name");
-  Value *GPUKernel = createCallGetKernel(KernelString, NameString);
+  Value *GPUKernel = createCallGetKernel(KernelString, KernelSize, NameString);
 
   Value *GridDimX, *GridDimY;
   std::tie(GridDimX, GridDimY) = getGridSizes(Kernel);
 
+  Value *CLUseLocalWorkSize;
+
+  switch (Arch) {
+  case GPUArch::SPIR32:
+  case GPUArch::SPIR64:
+    CLUseLocalWorkSize = Builder.getInt32(1);
+    break;
+  case GPUArch::AMDGCN64:
+  case GPUArch::NVPTX64:
+    CLUseLocalWorkSize = Builder.getInt32(0);
+    break;
+  }
+
   createCallLaunchKernel(GPUKernel, GridDimX, GridDimY, BlockDimX, BlockDimY,
-                         BlockDimZ, Parameters);
+                         BlockDimZ, Parameters, CLUseLocalWorkSize);
   createCallFreeKernel(GPUKernel);
 
   for (auto Id : KernelIds)
     isl_id_free(Id);
 
   KernelIds.clear();
 }
 
if (!is64Bit) {
     Ret += "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:"
            "64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:"
            "64-v128:128:128-n16:32:64";
   }
 
   return Ret;
 }
 
+/// Compute the DataLayout string for the AMDGPU backend.
+static std::string computeAMDGCNDataLayout() {
+  return "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32"
+         "-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128"
+         "-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64";
+}
+
 /// Compute the DataLayout string for a SPIR kernel.
 ///
 /// @param is64Bit Are we looking for a 64 bit architecture?
 static std::string computeSPIRDataLayout(bool is64Bit) {
   std::string Ret = "";
 
   if (!is64Bit) {
     Ret += "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:"
if (Arch == GPUArch::SPIR32 || Arch == GPUArch::SPIR64) {
     FN->setMetadata("kernel_arg_base_type",
                     MDNode::get(FN->getContext(), EmptyStrings));
   }
 
   switch (Arch) {
   case GPUArch::NVPTX64:
     FN->setCallingConv(CallingConv::PTX_Kernel);
     break;
+  case GPUArch::AMDGCN64:
+    FN->setCallingConv(CallingConv::AMDGPU_KERNEL);
+    break;
   case GPUArch::SPIR32:
   case GPUArch::SPIR64:
     FN->setCallingConv(CallingConv::SPIR_KERNEL);
     break;
   }
 
   auto Arg = FN->arg_begin();
   for (long i = 0; i < Kernel->n_array; i++) {
void GPUNodeBuilder::insertKernelIntrinsics(ppcg_kernel *Kernel) {
   case GPUArch::NVPTX64:
     IntrinsicsBID[0] = Intrinsic::nvvm_read_ptx_sreg_ctaid_x;
     IntrinsicsBID[1] = Intrinsic::nvvm_read_ptx_sreg_ctaid_y;
 
     IntrinsicsTID[0] = Intrinsic::nvvm_read_ptx_sreg_tid_x;
     IntrinsicsTID[1] = Intrinsic::nvvm_read_ptx_sreg_tid_y;
     IntrinsicsTID[2] = Intrinsic::nvvm_read_ptx_sreg_tid_z;
     break;
+  case GPUArch::AMDGCN64:
+    IntrinsicsBID[0] = Intrinsic::amdgcn_workitem_id_x;
+    IntrinsicsBID[1] = Intrinsic::amdgcn_workitem_id_y;
+
+    IntrinsicsTID[0] = Intrinsic::amdgcn_workgroup_id_x;
+    IntrinsicsTID[1] = Intrinsic::amdgcn_workgroup_id_y;
+    IntrinsicsTID[2] = Intrinsic::amdgcn_workgroup_id_z;
+    break;
   }
 
   auto addId = [this](__isl_take isl_id *Id, Intrinsic::ID Intr) mutable {
     std::string Name = isl_id_get_name(Id);
     Module *M = Builder.GetInsertBlock()->getParent()->getParent();
     Function *IntrinsicFn = Intrinsic::getDeclaration(M, Intr);
     Value *Val = Builder.CreateCall(IntrinsicFn, {});
     Val = Builder.CreateIntCast(Val, Builder.getInt64Ty(), false, Name);
void GPUNodeBuilder::createKernelFunction(
   switch (Arch) {
   case GPUArch::NVPTX64:
     if (Runtime == GPURuntime::CUDA)
       GPUModule->setTargetTriple(Triple::normalize("nvptx64-nvidia-cuda"));
     else if (Runtime == GPURuntime::OpenCL)
       GPUModule->setTargetTriple(Triple::normalize("nvptx64-nvidia-nvcl"));
     GPUModule->setDataLayout(computeNVPTXDataLayout(true /* is64Bit */));
     break;
+  case GPUArch::AMDGCN64:
+    if (Runtime == GPURuntime::CUDA)
+      llvm_unreachable("Cannot generate AMD code for CUDA runtime");
+    else if (Runtime == GPURuntime::OpenCL)
+      GPUModule->setTargetTriple(Triple::normalize("amdgcn-amd-amdhsa-opencl"));
+    GPUModule->setDataLayout(computeAMDGCNDataLayout());
+    break;
   case GPUArch::SPIR32:
     GPUModule->setTargetTriple(Triple::normalize("spir-unknown-unknown"));
     GPUModule->setDataLayout(computeSPIRDataLayout(false /* is64Bit */));
     break;
   case GPUArch::SPIR64:
     GPUModule->setTargetTriple(Triple::normalize("spir64-unknown-unknown"));
     GPUModule->setDataLayout(computeSPIRDataLayout(true /* is64Bit */));
     break;
void GPUNodeBuilder::createKernelFunction(
   Builder.SetInsertPoint(EntryBlock, EntryBlock->begin());
 
   ScopDetection::markFunctionAsInvalid(FN);
 
   prepareKernelArguments(Kernel, FN);
   createKernelVariables(Kernel, FN);
 
   switch (Arch) {
+  case GPUArch::AMDGCN64:
   case GPUArch::NVPTX64:
     insertKernelIntrinsics(Kernel);
     break;
   case GPUArch::SPIR32:
   case GPUArch::SPIR64:
     insertKernelCallsSPIR(Kernel);
     break;
   }
 }
 
 std::string GPUNodeBuilder::createKernelASM() {
   llvm::Triple GPUTriple;
 
   switch (Arch) {
   case GPUArch::NVPTX64:
     switch (Runtime) {
     case GPURuntime::CUDA:
       GPUTriple = llvm::Triple(Triple::normalize("nvptx64-nvidia-cuda"));
       break;
     case GPURuntime::OpenCL:
       GPUTriple = llvm::Triple(Triple::normalize("nvptx64-nvidia-nvcl"));
       break;
     }
     break;
+  case GPUArch::AMDGCN64:
+    switch (Runtime) {
+    case GPURuntime::CUDA:
+      llvm_unreachable("Cannot generate AMD code for CUDA runtime");
+      break;
+    case GPURuntime::OpenCL:
+      GPUTriple = llvm::Triple(Triple::normalize("amdgcn-amd-amdhsa-opencl"));
+      break;
+    }
+    break;
   case GPUArch::SPIR64:
   case GPUArch::SPIR32:
     std::string SPIRAssembly;
     raw_string_ostream IROstream(SPIRAssembly);
     IROstream << *GPUModule;
     IROstream.flush();
     return SPIRAssembly;
   }
std::string GPUNodeBuilder::createKernelASM() {
   Options.UnsafeFPMath = FastMath;
 
   std::string subtarget;
 
   switch (Arch) {
   case GPUArch::NVPTX64:
     subtarget = CudaVersion;
     break;
+  case GPUArch::AMDGCN64:
+    subtarget = AMDArch;
+    break;
   case GPUArch::SPIR32:
   case GPUArch::SPIR64:
     llvm_unreachable("No subtarget for SPIR architecture");
   }
 
   std::unique_ptr<TargetMachine> TargetM(GPUTarget->createTargetMachine(
       GPUTriple.getTriple(), subtarget, "", Options, Optional<Reloc::Model>()));
 
   SmallString<0> ASMString;
   raw_svector_ostream ASMStream(ASMString);
   llvm::legacy::PassManager PM;
 
   PM.add(createTargetTransformInfoWrapperPass(TargetM->getTargetIRAnalysis()));
 
+  if (Arch == GPUArch::AMDGCN64) {
     if (TargetM->addPassesToEmitFile(
-          PM, ASMStream, TargetMachine::CGFT_AssemblyFile, true /* verify */)) {
+            PM, ASMStream, TargetMachine::CGFT_ObjectFile, true /* verify */)) {
       errs() << "The target does not support generation of this file type!\n";
       return "";
     }
+  } else {
+    if (TargetM->addPassesToEmitFile(PM, ASMStream,
+                                     TargetMachine::CGFT_AssemblyFile,
+                                     true /* verify */)) {
+      errs() << "The target does not support generation of this file type!\n";
+      return "";
+    }
+  }
 
   PM.run(*GPUModule);
 
-  return ASMStream.str();
+  std::string Assembly = ASMStream.str();
+
+  if (Arch == GPUArch::AMDGCN64) {
+    std::string FileDir = "/tmp/";
+    std::string FileName = "polly_temp_kernel_amd.asm";
+    std::string OutName = "polly_temp_kernel_amd.bin";
+    std::ofstream ASMFile;
+    ASMFile.open(FileDir + FileName);
+    ASMFile << Assembly;
+    ASMFile.close();
+
+    std::string LLDCommand;
+    LLDCommand =
+        "ld.lld -shared " + FileDir + FileName + " -o " + FileDir + OutName;
+    system(LLDCommand.c_str());
+
+    std::ifstream ifs(FileDir + OutName);
+    Assembly.assign((std::istreambuf_iterator<char>(ifs)),
+                    (std::istreambuf_iterator<char>()));
+
+    remove((FileDir + FileName).c_str());
+    remove((FileDir + OutName).c_str());
+  }
+
+  return Assembly;
 }
 
 bool GPUNodeBuilder::requiresCUDALibDevice() {
   bool RequiresLibDevice = false;
   for (Function &F : GPUModule->functions()) {
     if (!F.isDeclaration())
       continue;
 

lib/Support/RegisterPasses.cpp

cl::values(clEnumValN(GPURuntime::CUDA, "libcudart",
                clEnumValN(GPURuntime::OpenCL, "libopencl",
                           "use the OpenCL Runtime API")),
     cl::init(GPURuntime::CUDA), cl::ZeroOrMore, cl::cat(PollyCategory));
 
 static cl::opt<GPUArch>
     GPUArchChoice("polly-gpu-arch", cl::desc("The GPU Architecture to target"),
                   cl::values(clEnumValN(GPUArch::NVPTX64, "nvptx64",
                                         "target NVIDIA 64-bit architecture"),
+                             clEnumValN(GPUArch::AMDGCN64, "amdgcn64",
+                                        "target AMD GCN 64-bit architecture"),
                              clEnumValN(GPUArch::SPIR32, "spir32",
                                         "target SPIR 32-bit architecture"),
                              clEnumValN(GPUArch::SPIR64, "spir64",
                                         "target SPIR 64-bit architecture")),
                   cl::init(GPUArch::NVPTX64), cl::ZeroOrMore,
                   cl::cat(PollyCategory));
 #endif
 

test/GPGPU/add-scalars-in-scop-to-kills.ll

 ; RUN: opt %loadPolly -analyze -polly-scops < %s | FileCheck %s -check-prefix=SCOP
 ; RUN: opt %loadPolly -S -polly-codegen-ppcg < %s | FileCheck %s -check-prefix=HOST-IR
 
 ; REQUIRES: pollyacc
 
 ; Check that we detect a scop.
 ; SCOP:       Function: checkScalarKill
 ; SCOP-NEXT: Region: %XLoopInit---%for.end
 ; SCOP-NEXT: Max Loop Depth:  1
 
 ; Check that we have a scalar that is not a phi node in the scop.
 ; SCOP: i32 MemRef_x_0; // Element size 4
 
 ; Check that kernel launch is generated in host IR.
 ; the declare would not be generated unless a call to a kernel exists.
-; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*)
+; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*, i32)
 
 ; Check that we add variables that are local to a scop into the kills that we
 ; pass to PPCG. This should enable PPCG to codegen this example.
 ; void checkScalarKill(int A[], int B[], int C[], const int control1, int control2) {
 ; int x;
 ; #pragma scop
 ;     for(int i = 0; i < 1000; i++) {
 ; XLoopInit:        x = 0;

test/GPGPU/bounds-construction-with-ignore-param-bounds.ll

 ; does not contain all parameters.
 ;
 ; We check that Polly does the right thing in this case and sets up the parameter
 ; dimensions correctly.
 
 
 ; Check that kernel launch is generated in host IR.
 ; the declare would not be generated unless a call to a kernel exists.
-; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*)
+; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*, i32)
 ; ModuleID = 'test/GPGPU/bounds-construction-with-ignore-param-bounds.ll'
 
 ; C pseudocode
 ; ------------
 ; void f(int *arr, long niters, long stride) {
 ;     for(int i = 0; i < niters; i++) {
 ;       arr[i * stride] = 1;
 ;     }

test/GPGPU/cuda-managed-memory-simple.ll

 ; CHECK-NEXT:  %16 = getelementptr [2 x i8*], [2 x i8*]* %polly_launch_0_params, i64 0, i64 0
 ; CHECK-NEXT:  store i8* %14, i8** %polly_launch_0_param_0
 ; CHECK-NEXT:  %17 = bitcast i8** %polly_launch_0_param_0 to i8*
 ; CHECK-NEXT:  store i8* %17, i8** %16
 ; CHECK-NEXT:  %18 = getelementptr [2 x i8*], [2 x i8*]* %polly_launch_0_params, i64 0, i64 1
 ; CHECK-NEXT:  store i8* %15, i8** %polly_launch_0_param_1
 ; CHECK-NEXT:  %19 = bitcast i8** %polly_launch_0_param_1 to i8*
 ; CHECK-NEXT:  store i8* %19, i8** %18
-; CHECK-NEXT:  %20 = call i8* @polly_getKernel(i8* getelementptr inbounds ([852 x i8], [852 x i8]* @FUNC_copy_SCOP_0_KERNEL_0, i32 0, i32 0), i8* getelementptr inbounds ([26 x i8], [26 x i8]* @FUNC_copy_SCOP_0_KERNEL_0_name, i32 0, i32 0))
-; CHECK-NEXT:  call void @polly_launchKernel(i8* %20, i32 2, i32 1, i32 32, i32 1, i32 1, i8* %polly_launch_0_params_i8ptr)
+; CHECK-NEXT:  %20 = call i8* @polly_getKernel(i8* getelementptr inbounds ([852 x i8], [852 x i8]* @FUNC_copy_SCOP_0_KERNEL_0, i32 0, i32 0), i64 851, i8* getelementptr inbounds ([26 x i8], [26 x i8]* @FUNC_copy_SCOP_0_KERNEL_0_name, i32 0, i32 0))
+; CHECK-NEXT:  call void @polly_launchKernel(i8* %20, i32 2, i32 1, i32 32, i32 1, i32 1, i8* %polly_launch_0_params_i8ptr, i32 0)
 ; CHECK-NEXT:  call void @polly_freeKernel(i8* %20)
 ; CHECK-NEXT:  call void @polly_synchronizeDevice()
 ; CHECK-NEXT:  call void @polly_freeContext(i8* %13)
 
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 
 define void @copy(i32* %R, i32* %A) {
 entry:

test/GPGPU/intrinsic-copied-into-kernel.ll

 ; Check that the intrinsic call is present in the kernel IR.
 ; KERNEL-IR:   %p_sqrt = tail call float @llvm.sqrt.f32(float %A.arr.i.val_p_scalar_)
 ; KERNEL-IR:   declare float @llvm.sqrt.f32(float)
 ; KERNEL-IR:   declare float @llvm.fabs.f32(float)
 ; KERNEL-IR:   declare float @llvm.powi.f32(float, i32)
 
 ; Check that kernel launch is generated in host IR.
 ; the declare would not be generated unless a call to a kernel exists.
-; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*)
+; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*, i32)
 
 
 ; void f(float *A, float *B, int N) {
 ;   for(int i = 0; i < N; i++) {
 ;       float tmp0 = A[i];
 ;       float tmp1 = sqrt(tmp1);
 ;       float tmp2 = fabs(tmp2);
 ;       float tmp3 = copysignf(tmp1, tmp2);

test/GPGPU/invariant-load-array-access.ll

 ; SCOP-NEXT:         Execution Context: [tmp] -> {  :  }
 ; SCOP-NEXT:         ReadAccess :=  [Reduction Type: NONE] [Scalar: 0]
 ; SCOP-NEXT:             [tmp] -> { Stmt_if_then[i0] -> MemRef_readarr[0] };
 ; SCOP-NEXT:         Execution Context: [tmp] -> {  : tmp >= 4 }
 ; SCOP-NEXT: }
 
 ; Check that kernel launch is generated in host IR.
 ; the declare would not be generated unless a call to a kernel exists.
-; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*)
+; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*, i32)
 
 ; This test makes sure that such an access pattern is handled correctly
 ; by PPCGCodeGeneration. It appears that not calling `preloadInvariantLoads`
 ; was the main reason that caused this test case to crash.
 ;
 ; void f(int *arr, const int *control, const int *readarr) {
 ;     for(int i = 0; i < 1000; i++) {
 ;         int t = 0;

test/GPGPU/invariant-load-hoisting-of-array.ll

 ; Check that we get back data from the kernel.
 ; HOST-IR: polly.acc.initialize:                             ; preds = %polly.start
 ; HOST-IR: [[FIRSTINDEX:%.+]] = getelementptr double, double* %polly.access.arr.of.ptrs.load, i64 1
 ; HOST-IR: [[BITCASTED:%.+]] = bitcast double* [[FIRSTINDEX]] to i8*
 ; HOST-IR: call void @polly_copyFromDeviceToHost(i8* %p_dev_array_MemRef_loaded_ptr, i8* [[BITCASTED]], i64 800)
 
 ; Check that the kernel launch is generated in the host IR.
 ; This declaration would not have been generated unless a kernel launch exists.
-; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*)
+; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*, i32)
 
 
 ; C pseudocode equivalent
 ; void f(double **arr_of_ptrs) {
 ;     double *loaded_ptr = arr_of_ptrs[0];
 ;     if (false) { return; }
 ;     else {
 ;         for(int i = 1; i < 100; i++) {

test/GPGPU/invariant-load-hoisting-with-variable-bounds.ll

 ; SCOP-NEXT:         ReadAccess :=  [Reduction Type: NONE] [Scalar: 0]
 ; SCOP-NEXT:             [tmp1, tmp4] -> { Stmt_for_body[i0] -> MemRef_end[0] };
 ; SCOP-NEXT:         Execution Context: [tmp1, tmp4] -> {  :  }
 ; SCOP-NEXT: }
 
 
 ; Check that the kernel launch is generated in the host IR.
 ; This declaration would not have been generated unless a kernel launch exists.
-; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*)
+; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*, i32)
 
 ;    void f(int *begin, int *end, int *arr) {
 ;      for (int i = *begin; i < *end; i++) {
 ;        arr[i] = 0;
 ;      }
 ;    }
 ;
 

test/GPGPU/invariant-load-hoisting-with-variable-lower-bound.ll

 ; SCOP-NEXT: Invariant Accesses: {
 ; SCOP-NEXT:         ReadAccess :=  [Reduction Type: NONE] [Scalar: 0]
 ; SCOP-NEXT:             [beginval] -> { Stmt_entry_split[] -> MemRef_begin[0] };
 ; SCOP-NEXT:         Execution Context: [beginval] -> {  :  }
 ; SCOP-NEXT: }
 
 ; Check that the kernel launch is generated in the host IR.
 ; This declaration would not have been generated unless a kernel launch exists.
-; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*)
+; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*, i32)
 
 ; 
 ; void f(int *begin, int *arr) {
 ;     for (int i = *begin; i < 100; i++) {
 ;         arr[i] = 0;
 ;     }
 ; }
 

test/GPGPU/invariant-load-hoisting-with-variable-upper-bound.ll

 ; SCOP-NEXT: Invariant Accesses: {
 ; SCOP-NEXT:         ReadAccess :=  [Reduction Type: NONE] [Scalar: 0]
 ; SCOP-NEXT:             [tmp2] -> { Stmt_for_body[i0] -> MemRef_idx[0] };
 ; SCOP-NEXT:         Execution Context: [tmp2] -> {  :  }
 ; SCOP-NEXT: }
 
 ; Check that kernel launch is generated in host IR.
 ; the declare would not be generated unless a call to a kernel exists.
-; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*)
+; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*, i32)
 
 ; Check if we generate GPU code for simple loop with variable upper bound.
 ; This always worked, but have this test to prevent regressions.
 ;    void f(int *idx, int *arr) {
 ;      for (int i = 0; i < *idx; i++) {
 ;        arr[i] = 0;
 ;      }
 ;    }

test/GPGPU/invariant-load-hoisting.ll

 ; SCOP:       Function: f
 ; SCOP-NEXT:  Region: %entry.split---%for.end26
 ; SCOP-NEXT:  Max Loop Depth:  3
 ; SCOP-NEXT:  Invariant Accesses: {
 ; SCOP-NEXT:          ReadAccess := [Reduction Type: NONE] [Scalar: 0]
 ; SCOP-NEXT:              [n, tmp12] -> { Stmt_for_body6[i0, i1, i2] -> MemRef_invariant[0] };
 ; SCOP-NEXT:          Execution Context: [n, tmp12] -> {  : n > 0 }
 ; SCOP-NEXT:  }
-; HOST-IR:      call void @polly_launchKernel(i8* %209, i32 %215, i32 1, i32 32, i32 1, i32 1, i8* %polly_launch_0_params_i8ptr)
+; HOST-IR:      call void @polly_launchKernel(i8* %209, i32 %215, i32 1, i32 32, i32 1, i32 1, i8* %polly_launch_0_params_i8ptr, i32 0)
 ; HOST-IR-NEXT: call void @polly_freeKernel(i8* %209)
 
 ; KERNEL-IR: define ptx_kernel void @FUNC_f_SCOP_0_KERNEL_0(i8 addrspace(1)* %MemRef_B, i8 addrspace(1)* %MemRef_A, i32 %n, i32 %tmp12, i32 %polly.preload.tmp21.merge)
 
 
 ; Check that we generate correct GPU code in case of invariant load hoisting.
 ;
 ;

test/GPGPU/libdevice-functions-copied-into-kernel.ll

 
 ; Check that the intrinsic call is present in the kernel IR.
 ; KERNEL-IR:   %p_expf = tail call float @__nv_expf(float %A.arr.i.val_p_scalar_)
 ; KERNEL-IR:   %p_cosf = tail call float @__nv_cosf(float %p_expf)
 ; KERNEL-IR:   %p_logf = tail call float @__nv_logf(float %p_cosf)
 
 ; Check that kernel launch is generated in host IR.
 ; the declare would not be generated unless a call to a kernel exists.
-; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*)
+; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*, i32)
 
 
 ; void f(float *A, float *B, int N) {
 ;   for(int i = 0; i < N; i++) {
 ;       float tmp0 = A[i];
 ;       float expf  = expf(tmp1);
 ;       cosf = cosf(expf);
 ;       logf = logf(cosf);

test/GPGPU/privatization-simple.ll

 ; RUN: opt %loadPolly -analyze -polly-scops < %s | FileCheck %s -check-prefix=SCOP
 ; RUN: opt %loadPolly -S -polly-codegen-ppcg < %s | FileCheck %s -check-prefix=HOST-IR
 
 ; REQUIRES: pollyacc
 
 ; SCOP:      Function: f
 ; SCOP-NEXT: Region: %for.body---%for.end
 ; SCOP-NEXT: Max Loop Depth:  1
 
 ; Check that kernel launch is generated in host IR.
 ; the declare would not be generated unless a call to a kernel exists.
-; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*)
+; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*, i32)
 
 ; void f(int A[], int B[], int control, int C[]) {
 ;     int x;
 ; #pragma scop
 ;     for(int i = 0; i < 1000; i ++) {
 ;         x = 0;
 ;         if(control) x = C[i];
 ;         B[i] = x * A[i];

test/GPGPU/privatization.ll

 ; RUN: opt %loadPolly -analyze -polly-scops < %s | FileCheck %s -check-prefix=SCOP
 ; RUN: opt %loadPolly -S -polly-codegen-ppcg < %s | FileCheck %s -check-prefix=HOST-IR
 
 ; REQUIRES: pollyacc
 
 ; SCOP:      Function: checkPrivatization
 ; SCOP-NEXT: Region: %for.body---%for.end
 ; SCOP-NEXT: Max Loop Depth:  1
 
 
 ; Check that kernel launch is generated in host IR.
 ; the declare would not be generated unless a call to a kernel exists.
-; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*)
+; HOST-IR: declare void @polly_launchKernel(i8*, i32, i32, i32, i32, i32, i8*, i32)
 
 ; 
 ;
 ;    void checkPrivatization(int A[], int B[], int C[], int control) {
 ;      int x;
 ;    #pragma scop
 ;      for (int i = 0; i < 1000; i++) {
 ;        x = 0;

tools/GPURuntime/GPUJIT.h

  * const char *Entry = "_Z8myKernelPi";
  *
  * int main() {
  *   PollyGPUFunction *Kernel;
  *   PollyGPUContext *Context;
  *   PollyGPUDevicePtr *DevArray;
  *   int *HostData;
  *   int MemSize;
+ *   size_t KernelSize = strlen(KernelString);
  *
  *   int GridX = 8;
  *   int GridY = 8;
  *
  *   int BlockX = 16;
  *   int BlockY = 16;
  *   int BlockZ = 1;
  *
  *   MemSize = 256*64*sizeof(int);
  *   Context = polly_initContext();
  *   DevArray = polly_allocateMemoryForDevice(MemSize);
- *   Kernel = polly_getKernel(KernelString, KernelName);
+ *   Kernel = polly_getKernel(KernelString, KernelSize, KernelName);
  *
  *   void *Params[1];
  *   void *DevPtr = polly_getDevicePtr(DevArray)
  *   Params[0] = &DevPtr;
  *
  *   polly_launchKernel(Kernel, GridX, GridY, BlockX, BlockY, BlockZ, Params);
  *
  *   polly_copyFromDeviceToHost(HostData, DevData, MemSize);

 
 typedef struct CUDAContextT CUDAContext;
 typedef struct CUDAKernelT CUDAKernel;
 typedef struct CUDADevicePtrT CUDADevicePtr;
 
 PollyGPUContext *polly_initContextCUDA();
 PollyGPUContext *polly_initContextCL();
 PollyGPUFunction *polly_getKernel(const char *BinaryBuffer,
+                                  const size_t BinarySize,
                                   const char *KernelName);
 void polly_freeKernel(PollyGPUFunction *Kernel);
 void polly_copyFromHostToDevice(void *HostData, PollyGPUDevicePtr *DevData,
                                 long MemSize);
 void polly_copyFromDeviceToHost(PollyGPUDevicePtr *DevData, void *HostData,
                                 long MemSize);
 void polly_synchronizeDevice();
 void polly_launchKernel(PollyGPUFunction *Kernel, unsigned int GridDimX,
                         unsigned int GridDimY, unsigned int BlockSizeX,
                         unsigned int BlockSizeY, unsigned int BlockSizeZ,
-                        void **Parameters);
+                        void **Parameters, unsigned int CLUseLocalWorkSize);
 void polly_freeDeviceMemory(PollyGPUDevicePtr *Allocation);
 void polly_freeContext(PollyGPUContext *Context);
 
 // Note that polly_{malloc/free}Managed are currently not used by Polly.
 // We use them in COSMO by replacing all malloc with polly_mallocManaged and all
 // frees with cudaFree, so we can get managed memory "automatically".
 // Needless to say, this is a hack.
 // Please make sure that this code is not present in Polly when 2018 rolls in.
 // If this is still present, ping Siddharth Bhat <[e-mail]>
 void *polly_mallocManaged(size_t size);
 void polly_freeManaged(void *mem);
 #endif /* GPUJIT_H_ */

tools/GPURuntime/GPUJIT.c

static void freeKernelCL(PollyGPUFunction *Kernel) {
   if (Kernel->Kernel)
     free((OpenCLKernel *)Kernel->Kernel);
 
   if (Kernel)
     free(Kernel);
 }
 
 static PollyGPUFunction *getKernelCL(const char *BinaryBuffer,
+                                     const size_t BinarySize,
                                      const char *KernelName) {
   dump_function();
 
   if (!GlobalContext) {
     fprintf(stderr, "GPGPU-code generation not correctly initialized.\n");
     exit(-1);
   }
 
if (HandleOpenCLBeignet) {
 
     ((OpenCLKernel *)Function->Kernel)->Program =
         clCreateProgramWithLLVMIntelFcnPtr(
             ((OpenCLContext *)GlobalContext->Context)->Context, 1,
             &GlobalDeviceID, "kernel.ll", &Ret);
     checkOpenCLError(Ret, "Failed to create program from llvm.\n");
     unlink("kernel.ll");
   } else {
-    size_t BinarySize = strlen(BinaryBuffer);
     ((OpenCLKernel *)Function->Kernel)->Program =
         clCreateProgramWithBinaryFcnPtr(
             ((OpenCLContext *)GlobalContext->Context)->Context, 1,
             &GlobalDeviceID, (const size_t *)&BinarySize,
             (const unsigned char **)&BinaryBuffer, NULL, &Ret);
     checkOpenCLError(Ret, "Failed to create program from binary.\n");
   }
 
Ret = clEnqueueReadBufferFcnPtr(
       ((OpenCLDevicePtr *)DevData->DevicePtr)->MemObj, CL_TRUE, 0, MemSize,
       HostData, 0, NULL, NULL);
   checkOpenCLError(Ret, "Copying results from device to host memory failed.\n");
 }
 
 static void launchKernelCL(PollyGPUFunction *Kernel, unsigned int GridDimX,
                            unsigned int GridDimY, unsigned int BlockDimX,
                            unsigned int BlockDimY, unsigned int BlockDimZ,
-                           void **Parameters) {
+                           void **Parameters, unsigned int UseLocalWorkSize) {
   dump_function();
 
   cl_int Ret;
   cl_uint NumArgs;
 
   if (!GlobalContext) {
     fprintf(stderr, "GPGPU-code generation not correctly initialized.\n");
     exit(-1);
static void launchKernelCL(PollyGPUFunction *Kernel, unsigned int GridDimX,
 
   unsigned int GridDimZ = 1;
   size_t GlobalWorkSize[3] = {BlockDimX * GridDimX, BlockDimY * GridDimY,
                               BlockDimZ * GridDimZ};
   size_t LocalWorkSize[3] = {BlockDimX, BlockDimY, BlockDimZ};
 
   static const int WorkDim = 3;
   OpenCLContext *CLContext = (OpenCLContext *)GlobalContext->Context;
-  Ret = clEnqueueNDRangeKernelFcnPtr(CLContext->CommandQueue, CLKernel->Kernel,
-                                     WorkDim, NULL, GlobalWorkSize,
-                                     LocalWorkSize, 0, NULL, NULL);
+  if (UseLocalWorkSize)
+    Ret = clEnqueueNDRangeKernelFcnPtr(
+        CLContext->CommandQueue, CLKernel->Kernel, WorkDim, NULL,
+        GlobalWorkSize, LocalWorkSize, 0, NULL, NULL);
+  else
+    Ret = clEnqueueNDRangeKernelFcnPtr(CLContext->CommandQueue,
+                                       CLKernel->Kernel, WorkDim, NULL,
+                                       GlobalWorkSize, NULL, 0, NULL, NULL);
   checkOpenCLError(Ret, "Launching OpenCL kernel failed.\n");
 }
 
 static void freeDeviceMemoryCL(PollyGPUDevicePtr *Allocation) {
   dump_function();
 
   OpenCLDevicePtr *DevPtr = (OpenCLDevicePtr *)Allocation->DevicePtr;
   cl_int Ret = clReleaseMemObjectFcnPtr((cl_mem)DevPtr->MemObj);
case RUNTIME_CL:
     break;
 #endif /* HAS_LIBOPENCL */
   default:
     err_runtime();
   }
 }
 
 PollyGPUFunction *polly_getKernel(const char *BinaryBuffer,
+                                  const size_t BinarySize,
                                   const char *KernelName) {
   dump_function();
 
   PollyGPUFunction *Function;
 
   switch (Runtime) {
 #ifdef HAS_LIBCUDART
   case RUNTIME_CUDA:
     Function = getKernelCUDA(BinaryBuffer, KernelName);
     break;
 #endif /* HAS_LIBCUDART */
 #ifdef HAS_LIBOPENCL
   case RUNTIME_CL:
-    Function = getKernelCL(BinaryBuffer, KernelName);
+    Function = getKernelCL(BinaryBuffer, BinarySize, KernelName);
     break;
 #endif /* HAS_LIBOPENCL */
   default:
     err_runtime();
   }
 
   return Function;
 }
#endif /* HAS_LIBOPENCL */
   default:
     err_runtime();
   }
 }
 
 void polly_launchKernel(PollyGPUFunction *Kernel, unsigned int GridDimX,
                         unsigned int GridDimY, unsigned int BlockDimX,
                         unsigned int BlockDimY, unsigned int BlockDimZ,
-                        void **Parameters) {
+                        void **Parameters, unsigned int CLUseLocalWorkSize) {
   dump_function();
 
   switch (Runtime) {
 #ifdef HAS_LIBCUDART
   case RUNTIME_CUDA:
     launchKernelCUDA(Kernel, GridDimX, GridDimY, BlockDimX, BlockDimY,
                      BlockDimZ, Parameters);
     break;
 #endif /* HAS_LIBCUDART */
 #ifdef HAS_LIBOPENCL
   case RUNTIME_CL:
     launchKernelCL(Kernel, GridDimX, GridDimY, BlockDimX, BlockDimY, BlockDimZ,
-                   Parameters);
+                   Parameters, CLUseLocalWorkSize);
     break;
 #endif /* HAS_LIBOPENCL */
   default:
     err_runtime();
   }
 }
 
 void polly_freeDeviceMemory(PollyGPUDevicePtr *Allocation) {