[Polly] [PPCGCodeGeneration] Add managed memory support to GPU code generation.

[bollu]

Summary

This needs changes to GPURuntime to expose synchronization between host
and device.

1. Needs better function naming, I want a better name than

"getOrCreateManagedDeviceArray"

2. DeviceAllocations is used by both the managed memory and the

non-managed memory path. This exploits the fact that the two code paths
are never run together. I'm not sure if this is the best design decision

[PhilippSchaad]

Hi Siddharth

I really like this idea! Seems like this could be helpful indeed.

However, I do understand it correctly that this should work effortlessly with for example the PolyBench/C kernels, out of the box?
Compiling a number of different kernels, the assertion

865:  assert(false && "multiple allocations for same scop");

gets triggered and the compilation fails with the following error:
clang -O3 -mllvm -polly -mllvm -polly-target=gpu -mllvm -polly-acc-codegen-managed-memory -ldl -lGPURuntime -I utilities -I linear-algebra/solvers/trisolv utilities/polybench.c linear-algebra/solvers/trisolv/trisolv.c -DPOLYBENCH_TIME

{P7990, lines=25}

[bollu]

• remove debug-time assert that was leftover from debugging

[bollu]

Yes, I realised this and pushed it to my git. Looks like I forgot to update phab. Thanks for the catch!
That was an assert error that was left out when I was trying to understand different cases.

Also, note that it will not work (fully) out of the box, you would need to change malloc/free to cudaMallocManaged/cudaFree, unless you're running on CUDA 8, in which case it should just work with regular malloc + kernel patches. But yes, the change is minimal

I did modify polybench to run this: [Click link - raw: https://github.com/bollu/polybench-c ](https://github.com/bollu/polybench-c)

[bollu]

(there is a makefile which can be run to create benchmarks)

[PhilippSchaad]

Using the modified Polybench/C I can confirm that it's working. When respecting those necessary changes, this patch looks good to me.

[grosser]

Hey guys,

thanks for pushing this forward. The patch indeed looks already very good. Here, some last some stylistic comments from me.

lib/CodeGen/PPCGCodeGeneration.cpp

511

Only use braces in case there is more than one statement.

654

Missing comment in assert.

674

Start with uppercase.

694

Missing comment.

717

Missing comment.

738

Missing comment.

863

Start with uppercase.

868

use llvm_unreachable

892

Start with uppercase.

954

Don't use braces

965

Here, braces are OK.

[bollu]

• [NFC with respect to branch] fixed one-line if and assert comments

[bollu]

lib/CodeGen/PPCGCodeGeneration.cpp

868

This should be reachable. I hadn't updated the differential with the commit.

965

quick style question because it doesn't seem to be there in the LLVM style guide either.

How do I style

question.cpp

if (predicate) {
    then();
} else {
    else1();
    else2();
}

do I transform this to:

remove-brace-from-then-clause

if (predicate)
    then();
else {
    else1();
    else2();
}

[bollu]

@grosser, @Meinersbur: ping. I've updated comments and fixed if {} styling.

[bollu]

[NoChange] Trying to use LLVM Phab instructions, need to have this commit track arc
patch.

[bollu]

Closed by 0cdb52faa664af4427e627b30872d660f177b1d9. I messed up the differential revision number. It should point to D32226 (this revision). Commit message points to [D32225](https://reviews.llvm.org/D32225)

lib/CodeGen/PPCGCodeGeneration.cpp

static cl::opt<bool> SharedMemory("polly-acc-use-shared",
                                   cl::desc("Use shared memory"), cl::Hidden,
                                   cl::init(false), cl::ZeroOrMore,
                                   cl::cat(PollyCategory));
 static cl::opt<bool> PrivateMemory("polly-acc-use-private",
                                    cl::desc("Use private memory"), cl::Hidden,
                                    cl::init(false), cl::ZeroOrMore,
                                    cl::cat(PollyCategory));
 
+static cl::opt<bool> ManagedMemory("polly-acc-codegen-managed-memory",
+                                   cl::desc("Generate Host kernel code assuming"
+                                            " that all memory has been"
+                                            " declared as managed memory"),
+                                   cl::Hidden, cl::init(false), 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<int>
     MinCompute("polly-acc-mincompute",
                cl::desc("Minimal number of compute statements to run on GPU."),
private:
 
   /// Compute the sizes of the execution grid for a given kernel.
   ///
   /// @param Kernel The kernel to compute grid sizes for.
   ///
   /// @returns A tuple with grid sizes for X and Y dimension
   std::tuple<Value *, Value *> getGridSizes(ppcg_kernel *Kernel);
 
+  /// Creates a array that can be sent to the kernel on the device using a
+  /// host pointer. This is required for managed memory, when we directly send
+  /// host pointers to the device.
+  /// \note
+  /// This is to be used only with managed memory
+  Value *getOrCreateManagedDeviceArray(gpu_array_info *Array,
+                                       ScopArrayInfo *ArrayInfo);
+
   /// Compute the sizes of the thread blocks for a given kernel.
   ///
   /// @param Kernel The kernel to compute thread block sizes for.
   ///
   /// @returns A tuple with thread block sizes for X, Y, and Z dimensions.
   std::tuple<Value *, Value *, Value *> getBlockSizes(ppcg_kernel *Kernel);
 
   /// Create kernel launch parameters.
private:
 
   /// Create a call to copy data from device to host.
   ///
   /// @param DevicePtr A pointer to the device data that should be copied.
   /// @param HostPtr A host pointer specifying the location to copy to.
   void createCallCopyFromDeviceToHost(Value *DevicePtr, Value *HostPtr,
                                       Value *Size);
 
+  /// 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 Entry  The name of the kernel function to call.
   ///
   /// @returns A pointer to a kernel object
   Value *createCallGetKernel(Value *Buffer, Value *Entry);
 

 
 void GPUNodeBuilder::initializeAfterRTH() {
   BasicBlock *NewBB = SplitBlock(Builder.GetInsertBlock(),
                                  &*Builder.GetInsertPoint(), &DT, &LI);
   NewBB->setName("polly.acc.initialize");
   Builder.SetInsertPoint(&NewBB->front());
 
   GPUContext = createCallInitContext();
+
+  if (!ManagedMemory)
     allocateDeviceArrays();
 }
 
 void GPUNodeBuilder::finalize() {
+  if (!ManagedMemory)
     freeDeviceArrays();
+
   createCallFreeContext(GPUContext);
   IslNodeBuilder::finalize();
 }
 
 void GPUNodeBuilder::allocateDeviceArrays() {
+  assert(!ManagedMemory && "Managed memory will directly send host pointers "
+                           "to the kernel. There is no need for device arrays");
   isl_ast_build *Build = isl_ast_build_from_context(S.getContext());
 
   for (int i = 0; i < Prog->n_array; ++i) {
     gpu_array_info *Array = &Prog->array[i];
     auto *ScopArray = (ScopArrayInfo *)Array->user;
     std::string DevArrayName("p_dev_array_");
     DevArrayName.append(Array->name);
 
Metadata *Elements[] = {
         ValueAsMetadata::get(V[2]),
     };
     MDNode *Node = MDNode::get(M->getContext(), Elements);
     AnnotationNode->addOperand(Node);
   }
 }
 
 void GPUNodeBuilder::freeDeviceArrays() {
+  assert(!ManagedMemory && "Managed memory does not use device arrays");
   for (auto &Array : DeviceAllocations)
     createCallFreeDeviceMemory(Array.second);
 }
 
 Value *GPUNodeBuilder::createCallGetKernel(Value *Buffer, Value *Entry) {
   const char *Name = "polly_getKernel";
   Module *M = Builder.GetInsertBlock()->getParent()->getParent();
   Function *F = M->getFunction(Name);
if (!F) {
     FunctionType *Ty = FunctionType::get(Builder.getVoidTy(), Args, false);
     F = Function::Create(Ty, Linkage, Name, M);
   }
 
   Builder.CreateCall(F, {GPUKernel});
 }
 
 void GPUNodeBuilder::createCallFreeDeviceMemory(Value *Array) {
+  assert(!ManagedMemory && "Managed memory does not allocate or free memory "
+                           "for device");
   const char *Name = "polly_freeDeviceMemory";
   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.getVoidTy(), Args, false);
     F = Function::Create(Ty, Linkage, Name, M);
   }
 
   Builder.CreateCall(F, {Array});
 }
 
 Value *GPUNodeBuilder::createCallAllocateMemoryForDevice(Value *Size) {
+  assert(!ManagedMemory && "Managed memory does not allocate or free memory "
+                           "for device");
   const char *Name = "polly_allocateMemoryForDevice";
   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.getInt64Ty());
     FunctionType *Ty = FunctionType::get(Builder.getInt8PtrTy(), Args, false);
     F = Function::Create(Ty, Linkage, Name, M);
   }
 
   return Builder.CreateCall(F, {Size});
 }
 
 void GPUNodeBuilder::createCallCopyFromHostToDevice(Value *HostData,
                                                     Value *DeviceData,
                                                     Value *Size) {
+  assert(!ManagedMemory && "Managed memory does not transfer memory between "
+                           "device and host");
   const char *Name = "polly_copyFromHostToDevice";
   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.getInt8PtrTy());
     Args.push_back(Builder.getInt64Ty());
     FunctionType *Ty = FunctionType::get(Builder.getVoidTy(), Args, false);
     F = Function::Create(Ty, Linkage, Name, M);
   }
 
   Builder.CreateCall(F, {HostData, DeviceData, Size});
 }
 
 void GPUNodeBuilder::createCallCopyFromDeviceToHost(Value *DeviceData,
                                                     Value *HostData,
                                                     Value *Size) {
+  assert(!ManagedMemory && "Managed memory does not transfer memory between "
+                           "device and host");
   const char *Name = "polly_copyFromDeviceToHost";
   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.getInt8PtrTy());
     Args.push_back(Builder.getInt64Ty());
     FunctionType *Ty = FunctionType::get(Builder.getVoidTy(), Args, false);
     F = Function::Create(Ty, Linkage, Name, M);
   }
 
   Builder.CreateCall(F, {DeviceData, HostData, Size});
 }
 
+void GPUNodeBuilder::createCallSynchronizeDevice() {
+  assert(ManagedMemory && "explicit synchronization is only necessary for "
+                          "managed memory");
+  const char *Name = "polly_synchronizeDevice";
+  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;
+    FunctionType *Ty = FunctionType::get(Builder.getVoidTy(), false);
+    F = Function::Create(Ty, Linkage, Name, M);
+  }
+
+  Builder.CreateCall(F);
+}
+
 Value *GPUNodeBuilder::createCallInitContext() {
   const char *Name = "polly_initContext";
   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;
Value *GPUNodeBuilder::getArrayOffset(gpu_array_info *Array) {
 
   Value *ResultValue = ExprBuilder.create(Result);
   isl_set_free(Min);
   isl_ast_build_free(Build);
 
   return ResultValue;
 }
 
+Value *GPUNodeBuilder::getOrCreateManagedDeviceArray(gpu_array_info *Array,
+                                                     ScopArrayInfo *ArrayInfo) {
+
+  assert(ManagedMemory && "Only used when you wish to get a host "
+                          "pointer for sending data to the kernel, "
+                          "with managed memory");
+  std::map<ScopArrayInfo *, Value *>::iterator it;
+  if ((it = DeviceAllocations.find(ArrayInfo)) != DeviceAllocations.end()) {
+    return it->second;
+  } else {
+    Value *HostPtr;
+
+    if (gpu_array_is_scalar(Array))
+      HostPtr = BlockGen.getOrCreateAlloca(ArrayInfo);
+    else
+      HostPtr = ArrayInfo->getBasePtr();
+
+    Value *Offset = getArrayOffset(Array);
+    if (Offset) {
+      HostPtr = Builder.CreatePointerCast(
+          HostPtr, ArrayInfo->getElementType()->getPointerTo());
+      HostPtr = Builder.CreateGEP(HostPtr, Offset);
+    }
+
+    HostPtr = Builder.CreatePointerCast(HostPtr, Builder.getInt8PtrTy());
+    DeviceAllocations[ArrayInfo] = HostPtr;
+    return HostPtr;
+  }
+}
+
 void GPUNodeBuilder::createDataTransfer(__isl_take isl_ast_node *TransferStmt,
                                         enum DataDirection Direction) {
+  assert(!ManagedMemory && "Managed memory needs no data transfers");
   isl_ast_expr *Expr = isl_ast_node_user_get_expr(TransferStmt);
   isl_ast_expr *Arg = isl_ast_expr_get_op_arg(Expr, 0);
   isl_id *Id = isl_ast_expr_get_id(Arg);
   auto Array = (gpu_array_info *)isl_id_get_user(Id);
   auto ScopArray = (ScopArrayInfo *)(Array->user);
 
   Value *Size = getArraySize(Array);
   Value *Offset = getArrayOffset(Array);
void GPUNodeBuilder::createUser(__isl_take isl_ast_node *UserStmt) {
   const char *Str = isl_id_get_name(Id);
   if (!strcmp(Str, "kernel")) {
     createKernel(UserStmt);
     isl_ast_expr_free(Expr);
     return;
   }
 
   if (isPrefix(Str, "to_device")) {
+    if (!ManagedMemory)
       createDataTransfer(UserStmt, HOST_TO_DEVICE);
+    else
+      isl_ast_node_free(UserStmt);
+
     isl_ast_expr_free(Expr);
     return;
   }
 
   if (isPrefix(Str, "from_device")) {
+    if (!ManagedMemory) {
       createDataTransfer(UserStmt, DEVICE_TO_HOST);
+    } else {
+      createCallSynchronizeDevice();
+      isl_ast_node_free(UserStmt);
+    }
     isl_ast_expr_free(Expr);
     return;
   }
 
   isl_id *Anno = isl_ast_node_get_annotation(UserStmt);
   struct ppcg_kernel_stmt *KernelStmt =
       (struct ppcg_kernel_stmt *)isl_id_get_user(Anno);
   isl_id_free(Anno);
GPUNodeBuilder::createLaunchParameters(ppcg_kernel *Kernel, Function *F,
   int Index = 0;
   for (long i = 0; i < Prog->n_array; i++) {
     if (!ppcg_kernel_requires_array_argument(Kernel, i))
       continue;
 
     isl_id *Id = isl_space_get_tuple_id(Prog->array[i].space, isl_dim_set);
     const ScopArrayInfo *SAI = ScopArrayInfo::getFromId(Id);
 
-    Value *DevArray = DeviceAllocations[const_cast<ScopArrayInfo *>(SAI)];
+    Value *DevArray = nullptr;
+    if (ManagedMemory) {
+      DevArray = getOrCreateManagedDeviceArray(
+          &Prog->array[i], const_cast<ScopArrayInfo *>(SAI));
+    } else {
+      DevArray = DeviceAllocations[const_cast<ScopArrayInfo *>(SAI)];
       DevArray = createCallGetDevicePtr(DevArray);
-
+    }
+    assert(DevArray != nullptr && "Array to be offloaded to device not "
+                                  "initialized");
     Value *Offset = getArrayOffset(&Prog->array[i]);
 
     if (Offset) {
       DevArray = Builder.CreatePointerCast(
           DevArray, SAI->getElementType()->getPointerTo());
       DevArray = Builder.CreateGEP(DevArray, Builder.CreateNeg(Offset));
       DevArray = Builder.CreatePointerCast(DevArray, Builder.getInt8PtrTy());
     }
     Value *Slot = Builder.CreateGEP(
         Parameters, {Builder.getInt64(0), Builder.getInt64(Index)});
 
     if (gpu_array_is_read_only_scalar(&Prog->array[i])) {
-      Value *ValPtr = BlockGen.getOrCreateAlloca(SAI);
+      Value *ValPtr = nullptr;
+      if (ManagedMemory)
+        ValPtr = DevArray;
+      else
+        ValPtr = BlockGen.getOrCreateAlloca(SAI);
+
+      assert(ValPtr != nullptr && "ValPtr that should point to a valid object"
+                                  " to be stored into Parameters");
       Value *ValPtrCast =
           Builder.CreatePointerCast(ValPtr, Builder.getInt8PtrTy());
       Builder.CreateStore(ValPtrCast, Slot);
     } else {
       Instruction *Param =
           new AllocaInst(Builder.getInt8PtrTy(), AddressSpace,
                          Launch + "_param_" + std::to_string(Index),
                          EntryBlock->getTerminator());

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

+; RUN: opt %loadPolly -S  -polly-process-unprofitable -polly-acc-mincompute=0 -polly-target=gpu  -polly-codegen-ppcg -polly-acc-codegen-managed-memory < %s | \
+; RUN: FileCheck %s
+
+; REQUIRES: pollyacc
+
+;
+;    #include <cuda_runtime.h>
+;
+;    static const int N = 45;
+;
+;    void copy(int *R, int *A) {
+;      for (int i = 0; i < N; i++) {
+;        R[i] = A[i] * 10;
+;      }
+;    }
+;
+;    int main() {
+;      int *A, *R;
+;
+;      cudaMallocManaged((void **)(&A), sizeof(int) * N, cudaMemAttachGlobal);
+;      cudaMallocManaged((void **)(&R), sizeof(int) * N, cudaMemAttachGlobal);
+;
+;      for (int i = 0; i < N; i++) {
+;        A[i] = i;
+;        R[i] = 0;
+;      }
+;      copy(R, A);
+;
+;      return 0;
+;    }
+;
+
+; CHECK-NOT: polly_copyFromHostToDevice
+; CHECK-NOT: polly_copyFromDeviceToHost
+; CHECK-NOT: polly_freeDeviceMemory
+; CHECK-NOT: polly_allocateMemoryForDevice
+
+; CHECK:       %13 = call i8* @polly_initContext()
+; CHECK-NEXT:  %14 = bitcast i32* %A to i8*
+; CHECK-NEXT:  %15 = 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:  %16 = bitcast i8** %polly_launch_0_param_0 to i8*
+; CHECK-NEXT:  store i8* %16, i8** %15
+; CHECK-NEXT:  %17 = bitcast i32* %R to i8*
+; CHECK-NEXT:  %18 = getelementptr [2 x i8*], [2 x i8*]* %polly_launch_0_params, i64 0, i64 1
+; CHECK-NEXT:  store i8* %17, 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 ([820 x i8], [820 x i8]* @kernel_0, i32 0, i32 0), i8* getelementptr inbounds ([9 x i8], [9 x i8]* @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:  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:
+  br label %for.cond
+
+for.cond:                                         ; preds = %for.inc, %entry
+  %indvars.iv = phi i64 [ %indvars.iv.next, %for.inc ], [ 0, %entry ]
+  %exitcond = icmp ne i64 %indvars.iv, 45
+  br i1 %exitcond, label %for.body, label %for.end
+
+for.body:                                         ; preds = %for.cond
+  %arrayidx = getelementptr inbounds i32, i32* %A, i64 %indvars.iv
+  %tmp = load i32, i32* %arrayidx, align 4
+  %mul = mul nsw i32 %tmp, 10
+  %arrayidx2 = getelementptr inbounds i32, i32* %R, i64 %indvars.iv
+  store i32 %mul, i32* %arrayidx2, align 4
+  br label %for.inc
+
+for.inc:                                          ; preds = %for.body
+  %indvars.iv.next = add nuw nsw i64 %indvars.iv, 1
+  br label %for.cond
+
+for.end:                                          ; preds = %for.cond
+  ret void
+}
+
+define i32 @main() {
+entry:
+  %A = alloca i32*, align 8
+  %R = alloca i32*, align 8
+  %tmp = bitcast i32** %A to i8**
+  %call = call i32 @cudaMallocManaged(i8** nonnull %tmp, i64 180, i32 1) #2
+  %tmp1 = bitcast i32** %R to i8**
+  %call1 = call i32 @cudaMallocManaged(i8** nonnull %tmp1, i64 180, i32 1) #2
+  br label %for.cond
+
+for.cond:                                         ; preds = %for.inc, %entry
+  %indvars.iv = phi i64 [ %indvars.iv.next, %for.inc ], [ 0, %entry ]
+  %exitcond = icmp ne i64 %indvars.iv, 45
+  br i1 %exitcond, label %for.body, label %for.end
+
+for.body:                                         ; preds = %for.cond
+  %tmp2 = load i32*, i32** %A, align 8
+  %arrayidx = getelementptr inbounds i32, i32* %tmp2, i64 %indvars.iv
+  %tmp3 = trunc i64 %indvars.iv to i32
+  store i32 %tmp3, i32* %arrayidx, align 4
+  %tmp4 = load i32*, i32** %R, align 8
+  %arrayidx3 = getelementptr inbounds i32, i32* %tmp4, i64 %indvars.iv
+  store i32 0, i32* %arrayidx3, align 4
+  br label %for.inc
+
+for.inc:                                          ; preds = %for.body
+  %indvars.iv.next = add nuw nsw i64 %indvars.iv, 1
+  br label %for.cond
+
+for.end:                                          ; preds = %for.cond
+  %tmp5 = load i32*, i32** %R, align 8
+  %tmp6 = load i32*, i32** %A, align 8
+  call void @copy(i32* %tmp5, i32* %tmp6)
+  ret i32 0
+}
+
+declare i32 @cudaMallocManaged(i8**, i64, i32) #1

tools/GPURuntime/GPUJIT.h

 PollyGPUContext *polly_initContext();
 PollyGPUFunction *polly_getKernel(const char *PTXBuffer,
                                   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 polly_freeDeviceMemory(PollyGPUDevicePtr *Allocation);
 void polly_freeContext(PollyGPUContext *Context);
 #endif /* GPUJIT_H_ */

tools/GPURuntime/GPUJIT.c

 
 typedef CUresult CUDAAPI CuLinkCompleteFcnTy(CUlinkState state, void **cubinOut,
                                              size_t *sizeOut);
 static CuLinkCompleteFcnTy *CuLinkCompleteFcnPtr;
 
 typedef CUresult CUDAAPI CuLinkDestroyFcnTy(CUlinkState state);
 static CuLinkDestroyFcnTy *CuLinkDestroyFcnPtr;
 
+typedef CUresult CUDAAPI CuCtxSynchronizeFcnTy();
+static CuCtxSynchronizeFcnTy *CuCtxSynchronizeFcnPtr;
+
 /* Type-defines of function pointer ot CUDA runtime APIs. */
 typedef cudaError_t CUDARTAPI CudaThreadSynchronizeFcnTy(void);
 static CudaThreadSynchronizeFcnTy *CudaThreadSynchronizeFcnPtr;
 
 static void *getAPIHandle(void *Handle, const char *FuncName) {
   char *Err;
   void *FuncPtr;
   dlerror();
CuLinkCreateFcnPtr =
       (CuLinkCreateFcnTy *)getAPIHandle(HandleCuda, "cuLinkCreate");
 
   CuLinkCompleteFcnPtr =
       (CuLinkCompleteFcnTy *)getAPIHandle(HandleCuda, "cuLinkComplete");
 
   CuLinkDestroyFcnPtr =
       (CuLinkDestroyFcnTy *)getAPIHandle(HandleCuda, "cuLinkDestroy");
 
+  CuCtxSynchronizeFcnPtr =
+      (CuCtxSynchronizeFcnTy *)getAPIHandle(HandleCuda, "cuCtxSynchronize");
+
   /* Get function pointer to CUDA Runtime APIs. */
   CudaThreadSynchronizeFcnPtr = (CudaThreadSynchronizeFcnTy *)getAPIHandle(
       HandleCudaRT, "cudaThreadSynchronize");
 
   return 1;
 }
 
 PollyGPUContext *polly_initContext() {
void polly_copyFromDeviceToHost(PollyGPUDevicePtr *DevData, void *HostData,
                                 long MemSize) {
   dump_function();
 
   if (CuMemcpyDtoHFcnPtr(HostData, DevData->Cuda, MemSize) != CUDA_SUCCESS) {
     fprintf(stdout, "Copying results from device to host memory failed.\n");
     exit(-1);
   }
 }
+void polly_synchronizeDevice() {
+  dump_function();
+  if (CuCtxSynchronizeFcnPtr() != CUDA_SUCCESS) {
+    fprintf(stdout, "Synchronizing device and host memory failed.\n");
+    exit(-1);
+  }
+}
 
 void polly_launchKernel(PollyGPUFunction *Kernel, unsigned int GridDimX,
                         unsigned int GridDimY, unsigned int BlockDimX,
                         unsigned int BlockDimY, unsigned int BlockDimZ,
                         void **Parameters) {
   dump_function();
 
   unsigned GridDimZ = 1;