[Polly] Added OpenCL Runtime to GPURuntime Library for GPGPU CodeGen

[PhilippSchaad]

Summary

When compiling for GPU, one can now choose to compile for OpenCL or CUDA,
with the corresponding polly-gpu-runtime flag (libopencl / libcudart). The
GPURuntime library (GPUJIT) has been extended with the OpenCL Runtime library
for that purpose, correctly choosing the corresponding library calls to the
option chosen when compiling (via different initialization calls).

Additionally, a specific GPU Target architecture can now be chosen with -polly-gpu-arch (only nvptx64 implemented thus far).

[PhilippSchaad]

• Replaced magic numbers, added assertions and fixed if-braces.

[Meinersbur]

I wrote a runtime with similar scope here: https://github.com/Meinersbur/prl . We were one discussing to use it for Polly as well. What's the status of that?

lib/CodeGen/PPCGCodeGeneration.cpp

56–58

Did you consider an enum?

1562

Is there some vendor-neutral triple?

1656

Why a static flag?

2530–2538

Did you consider

Pass *polly::createPPCGCodeGenerationPass(int Runtime);

?

lib/Support/RegisterPasses.cpp

312–319

A switch instead?

[PhilippSchaad]

In D32431#736600, @Meinersbur wrote:

I wrote a runtime with similar scope here: https://github.com/Meinersbur/prl . We were one discussing to use it for Polly as well. What's the status of that?

I have looked into it a tiny little bit about a month ago, but had then decided to write a basic OpenCL Runtime from scratch in GPUJIT. So to my knowledge, nothing has changed on that status yet.

Currently looking into the rest of your comment-mentioned points.

lib/CodeGen/PPCGCodeGeneration.cpp

1562

Do you mean like nvptx64-nvcl / nvptx64-cuda?

[PhilippSchaad]

lib/CodeGen/PPCGCodeGeneration.cpp

2530–2538

That seems reasonable, but I get a template-conflict for the LLVM Pass-Creation template when trying to change the pass-creation-method structure. I thought it might be easier this way?

[PhilippSchaad]

lib/CodeGen/PPCGCodeGeneration.cpp

2530–2538

Correction: looking at wrong function of course, you mean a different one :-)

[PhilippSchaad]

• Stylistic changes and switch to -polly-gpu-runtime=cuda/opencl compiler flag

[PhilippSchaad]

• Removed left over commented out macros

[Meinersbur]

lib/CodeGen/PPCGCodeGeneration.cpp

58–60

See http://llvm.org/docs/CodingStandards.html#name-types-functions-variables-and-enumerators-properly for LLVM's coding policy for enum members.

Nitpick: A "T" suffix is rather unusual.

153

Nitpick: No need to use an enum qualifier.

1562

I hoped that there might be some kind of triple that works for OpenCL in general, not only for nvidia (nvptx, nvcl). If the generated program only works for devices that support cuda anyway, I don't see where the benefit of such a backend is.

If there is indeed no backend that also works on non-nvidia devices, should we call the the runtime accordingly, e.g. "nvcl" then?

lib/Support/RegisterPasses.cpp

312–319

Now that createPPCGCodeGenerationPass takes an argument, you don't need a switch anymore.

tools/GPURuntime/GPUJIT.c

303–311

Consistent variable name style? What style do you intend to use in this file?

347–348

Replace the magic number 256 by sizeof(DeviceRevision)?

[etherzhhb]

include/polly/LinkAllPasses.h

51–52

is this Runtime supposed to be with type GPURuntimeT ? it is a little bit tricky here.
Maybe we need to introduce a PPCG header and define the runtime enum there, than include that runtime enum.

or we can declare the function as

llvm::Pass *createPPCGCodeGenerationPass(int Runtime = 0);

to at least avoid the magic number 0 in line 86.

lib/CodeGen/PPCGCodeGeneration.cpp

1562

for opencl, it can be "spir-unknown-unknown" or "spir64-unknown-unknown", but that may not work :)

[PhilippSchaad]

Looking into the rest of your comments.

include/polly/LinkAllPasses.h

51–52

Yes, it would be. The reason it's not is exactly the one you mentioned. I was considering adding a PPCG header, but refrained from it because I was hesitant about creating a header 'just for one enum'. If you agree that this is a good solution, I will indeed introduce a new header for PPCG and define the enum there, to get rid of magic numbers. The second option seems reasonable too though.

lib/CodeGen/PPCGCodeGeneration.cpp

1562

Looking into it. The next goal would be to add the AMDGPU backend to generate AMD ISA, which would then again utilize the same OpenCL Runtime implemented here. (I realize there will have to be some naming changes to make that clear in the GPUJIT, but as you pointed out, I have a naming-mess to fix there anyway.

[etherzhhb]

include/polly/LinkAllPasses.h

51–52

we could start from the second option if you think it is reasonable

[PhilippSchaad]

• Addressed consistency and naming concerns

[PhilippSchaad]

• Made CUDA Runtime default, fixed formatting, adapted test case

[grosser]

Hi Philip and others,

this already looks very cool. I also added some minor comments.

Best,
Tobias

lib/CodeGen/PPCGCodeGeneration.cpp

56–58

You can use C++11 enums ala

enum class GPURuntime { CUDA, OpenCL };

1562

Making OpenCL work for CUDA is just the first step. I expect that when adding AMDGPU support, we will use here different triples depending on which vendor to target. AMD will have a specific one, CUDA will have a specific one, and for Intel we likely use the generic SPIR-V comment. I assume this could then also work for Xilinx.

[PhilippSchaad]

• Fixed enum style to C++11

[Meinersbur]

lib/CodeGen/PPCGCodeGeneration.cpp

1562

At compile time, we don't know on which hardware it will run on, so we cannot specify a triple here.

Unless you think of a runtime dispatch system, then you need to generate all kernels at once. In that case, I still would like to select a single target only for when I know I will run only on that hardware and to keep the executable small.

[PhilippSchaad]

lib/CodeGen/PPCGCodeGeneration.cpp

1562

I thought the goal was to let the user compile for a specific target, i.e. providing something like -polly-gpu-arch=amd/nvidia/intel, and then choosing the correct target triple according to said selection. Meaning for example -polly-gpu-arch=amd would utilize the AMDGPU backend triple and feed that into the OpenCL runtime.

Am I misunderstanding something?

[Meinersbur]

lib/CodeGen/PPCGCodeGeneration.cpp

1562

I think we were miscommunicating. The -polly-gpu-arch switch is new to me and doesn't appear in this patch. I assumed a fat executable when you mentioned an AMD backend.

OpenCL claims to be hardware-independent with platform's driver translating OpenCL-C or SPIR(-V) to its proprietary format. In NVidia's terminology, CUDA is a platform of which CUDA C++, their OpenCL implementation, cudart (CUDA runtime) etc. are part of. We are still vendor-locked to CUDA, since it only works with CUDA's OpenCL runtime library. -polly-gpu-runtime=opencl therefore is misleading (at least it was to me), it it no alternative to CUDA.

It might resolve if it is indeed just the runtime library GPUJIT is linked to. If so, could you make it more clear? I suggest the following switches:

-polly-target=cpu/gpu

if -polly-target=gpu then
-polly-gpu-arch=nvptx64/hsa/spir/spir-v/opencl-c
(with -polly-gpu-arch=nvptx64 the only one implemented so far)

if -polly-gpu-arch=nvptx64 then there is a choice between
-polly-cuda-runtime=libcudart/libopencl

[PhilippSchaad]

• GPURuntime works on systems with just one of CUDA/OpenCL now.

[PhilippSchaad]

lib/CodeGen/PPCGCodeGeneration.cpp

1562

This change should address exactly this. The framework is now set to introduce new architectures and utilize eg. the AMDGPU backend instead of NVPTX etc.

[singam-sanjay]

lib/CodeGen/PPCGCodeGeneration.cpp

1592

Does nvptx64-nvidia-nvcl mean OpenCL code meant to be run on NVIDIA GPUs ?

[PhilippSchaad]

lib/CodeGen/PPCGCodeGeneration.cpp

1592

Yes, exactly. It generates a slightly different flavor of PTX, which can be used by OpenCL to generate a kernel from the PTX binary (on NVIDIA GPUs). If you were to use the standard CUDA PTX, OpenCL would complain because of wrong argument accesses.

[singam-sanjay]

lib/CodeGen/PPCGCodeGeneration.cpp

1592

Okay. From what you're saying, nvptx64-nvidia-nvcl indicates that backend must generate NVPTX code for a 64bit architecture for an NVIDIA GPU controlled by a OpenCL driver. Please correct me if I'm wrong.

[PhilippSchaad]

lib/CodeGen/PPCGCodeGeneration.cpp

1592

That is correct.

[singam-sanjay]

lib/CodeGen/PPCGCodeGeneration.cpp

1592

Thank you ! That was helpful.

[PhilippSchaad]

• Integrated D32226 - Managed memory support

[PhilippSchaad]

@grosser @Meinersbur ping

[PhilippSchaad]

• Fixed formatting and managed-memory test case (including pre-existing bug)

[Meinersbur]

I only consider the clSetKernelArg as a remaining bigger issue. Having only "polly_"-prefixed function non-static would also be great.

Tobias is the sole author of Polly-ACC. I think he should give the final LGTM.

include/polly/LinkAllPasses.h

51–52

I assume you kept the arguments of type int to not include header files here.

lib/CodeGen/PPCGCodeGeneration.cpp

712–715

Can you make this a switch so you get warned by the compiler when adding more runtimes?

1589–1592

Could also be a switch.

2539

We should check whether the arguments are valid. Such as:

switch (Runtime) {
case 1:
case 2:
  ...
default:
  llvm_unreachable("Invalid argument for Runtime");
}

2541–2545

Similarly:

switch (Arch) {
case 1:
  ...
default:
  llvm_unreachable("Invalid argument for Arch");
}

lib/Support/RegisterPasses.cpp

314–319

int Arch;
switch (GPUArch) {
case GPU_ARCH_NVPTX64;
  Arch = 1;
  break;
}

int Runtime;
switch (GPURuntime) {
case GPU_RUNTIME_CUDA:
  Runtime = 1;
  break;
case GPU_RUNTIME_OPENCL:
  Runtime = 2;
  breal
}

PM.add(polly::createPPCGCodeGenerationPass(Arch, Runtime))

With "you don't need a switch anymore", I was thinking about the like of:

PM.add(polly::createPPCGCodeGenerationPass(1, GPURuntime + 1));

Your choice.

It could be helpful to have createPPCGCodeGenerationPass accept the GPURuntime enum as arguments instead.

In your solution, I don't see the use of static const int local variables. If you want identifiers that give names to the accepted arguments, declare them as #define or static const int in the header file that also declares createPPCGCodeGenerationPass, so these can be used in the implementation of createPPCGCodeGenerationPass as well.

tools/GPURuntime/GPUJIT.c

392–396

Did you consider introducing a new function this sequence of code? It appears quite often.

605–613

Trying each argument size after the other and hoping one matches is not good. The caller must know the argument sizes. You probably have to pass the sizes in another argument to launchKernelCL that contains those sizes for each argument, generated by Polly.

Without this, the code will fail if you pass a struct (or vector) of size other than 8, 4, 2, or 1.

671

Shouldn't these print to stderr?

745

The function name does not follow the naming of other functions in this file. In C it is common have the public API functions prefixed with the library name (here: "polly") and everything else static. Don't choose the prefix of another library (here: "cl_"). This avoids symbol conflicts because multiple libraries happen to give the same name for a function.

[PhilippSchaad]

include/polly/LinkAllPasses.h

51–52

That is correct, yes.

lib/Support/RegisterPasses.cpp

314–319

With "you don't need a switch anymore", I was thinking about the like of:

PM.add(polly::createPPCGCodeGenerationPass(1, GPURuntime + 1));

Your choice.

Personally, I think the current solution is tiny little bit more 'documenting'. Both is fine though, good call.

It could be helpful to have createPPCGCodeGenerationPass accept the GPURuntime enum as arguments instead.

That is true, would mean having to provide a header with that GPURuntime enum instead though, right?

In your solution, I don't see the use of static const int local variables. If you want identifiers that give names to the accepted arguments, declare them as #define or static const int in the header file that also declares createPPCGCodeGenerationPass, so these can be used in the implementation of createPPCGCodeGenerationPass as well.

Would it maybe make sense to introduce a PPCGCodeGeneration header at this point?

tools/GPURuntime/GPUJIT.c

605–613

Yes, this is a priority issue still. The issue will have to be resolved at some point. This is basically a temporary way around some (probably) major argument handling changes in PPCG etc.

[PhilippSchaad]

• Addressed multiple issues pointed out in comment

[PhilippSchaad]

• Fixed formatting

[Meinersbur]

You changed stdout to stderr everywhere, which is better in my point of view, but logically is a different change. Sorry that I didn't realize that the libcudart also printed to stdout before so you tried to be consistent. Could you commit that change separately beforehand? (Maybe also the change in argument name capitalization)

I am accepting the patch proved that you are going to improve the clSetKernelArg situation later and add a TODO into the code about it. The other stuff is of stylistic nature only.

Please also wait for Tobias' approval.

lib/Support/RegisterPasses.cpp

314–319

Please remove at least the static keyword. It makes sense for global constants, but not for function-local ones.

The style

static const int ArgumentName = 0;
func(ArgumentName);

rather unusual in LLVM-style code (but not bad if applied consistenly, which is unfortunately not the case in Polly). I've seen

func(/* ArgumentName = */ 0);

much more often.

In this case I think UseOpenCLRuntime, UseCUDARuntime and TargetNVPTX64 should really be global constants declated close to the declaration of createPPCGCodeGenerationPass so it can be used by ever caller of that function.

Would it maybe make sense to introduce a PPCGCodeGeneration header at this point?

Yes, that sounds good to me as well.

tools/GPURuntime/GPUJIT.c

601–602

Thanks for the introduction of checkOpenCLError. You could also introduce one for these two lines. for instance:

if (!GlobalContext)
  handleError("GPGPU-code generation not correctly initialized.\n");

handleError could also be called by checkOpenCLError. It helps centralising the error handling, such that if we change some detail about it (e.g. the return code on exit, or some cleanup code), there is a single function for that.

605–613

Also note that I am not sure that OpenCL ICD's are required to check for correct CL_INVALID_ARG_SIZE. It might just trust the caller, or be a badly written one.

1015–1016

These are unrelated changes Tobias usually complains about. I personally don't care.

1144

Unrelated whitespace change?

[PhilippSchaad]

Addressed most of your concerns. @grosser it should be ready now, what do you think?

[PhilippSchaad]

• Introduced PPCGCodeGeneration header file for simplicity

[PhilippSchaad]

@Meinersbur the unrelated changes you mentioned have been added/moved to D32852 and D32854.

[grosser]

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

49 ↗

(On Diff #97804)

This change is unrelated.

[PhilippSchaad]

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

49 ↗

(On Diff #97804)

It is, but it got fixed in the meantime anyway. Removing it.

[PhilippSchaad]

Reopened for rebase

[PhilippSchaad]

Rebase

CMakeLists.txt

 
 # Add path for custom modules
 set(CMAKE_MODULE_PATH ${CMAKE_MODULE_PATH} "${POLLY_SOURCE_DIR}/cmake")
 
 SET(CMAKE_INSTALL_RPATH_USE_LINK_PATH TRUE)
 
 option(POLLY_ENABLE_GPGPU_CODEGEN "Enable GPGPU code generation feature" OFF)
 if (POLLY_ENABLE_GPGPU_CODEGEN)
-  # Do not require CUDA, as GPU code generation test cases can be run without
-  # a cuda library.
+  # Do not require CUDA/OpenCL, as GPU code generation test cases can be run
+  # without a CUDA/OpenCL library.
   FIND_PACKAGE(CUDA)
+  FIND_PACKAGE(OpenCL)
   set(GPU_CODEGEN TRUE)
 else(POLLY_ENABLE_GPGPU_CODEGEN)
   set(GPU_CODEGEN FALSE)
 endif(POLLY_ENABLE_GPGPU_CODEGEN)
 
 
 # Support GPGPU code generation if the library is available.
-if (CUDALIB_FOUND)
+if (CUDALIB_FOUND AND OpenCL_FOUND)
   INCLUDE_DIRECTORIES( ${CUDALIB_INCLUDE_DIR} )
-endif(CUDALIB_FOUND)
+  INCLUDE_DIRECTORIES( ${OpenCL_INCLUDE_DIR} )
+endif(CUDALIB_FOUND AND OpenCL_FOUND)
 
 option(POLLY_BUNDLED_ISL "Use the bundled version of libisl included in Polly" ON)
 if (NOT POLLY_BUNDLED_ISL)
   find_package(ISL MODULE REQUIRED)
   message(STATUS "Using external libisl ${ISL_VERSION} in: ${ISL_PREFIX}")
   set(ISL_TARGET ISL)
 else()
   set(ISL_INCLUDE_DIRS

include/polly/LinkAllPasses.h

 llvm::Pass *createPollyCanonicalizePass();
 llvm::Pass *createPolyhedralInfoPass();
 llvm::Pass *createScopDetectionPass();
 llvm::Pass *createScopInfoRegionPassPass();
 llvm::Pass *createScopInfoWrapperPassPass();
 llvm::Pass *createIslAstInfoPass();
 llvm::Pass *createCodeGenerationPass();
 #ifdef GPU_CODEGEN
-llvm::Pass *createPPCGCodeGenerationPass();
+llvm::Pass *createPPCGCodeGenerationPassCUDA();
+llvm::Pass *createPPCGCodeGenerationPassOpenCL();
 #endif
 llvm::Pass *createIslScheduleOptimizerPass();
 llvm::Pass *createFlattenSchedulePass();
 llvm::Pass *createDeLICMPass();
 
 extern char &CodePreparationID;
 } // namespace polly
 
PollyForcePassLinking() {
     polly::createJSONImporterPass();
     polly::createScopDetectionPass();
     polly::createScopInfoRegionPassPass();
     polly::createPollyCanonicalizePass();
     polly::createPolyhedralInfoPass();
     polly::createIslAstInfoPass();
     polly::createCodeGenerationPass();
 #ifdef GPU_CODEGEN
-    polly::createPPCGCodeGenerationPass();
+    polly::createPPCGCodeGenerationPassCUDA();
+    polly::createPPCGCodeGenerationPassOpenCL();
 #endif
     polly::createIslScheduleOptimizerPass();
     polly::createFlattenSchedulePass();
     polly::createDeLICMPass();
     polly::createDumpModulePass("", true);
     polly::createSimplifyPass();
     polly::createPruneUnprofitablePass();
   }

lib/CodeGen/PPCGCodeGeneration.cpp

 
 #include "llvm/Support/Debug.h"
 
 using namespace polly;
 using namespace llvm;
 
 #define DEBUG_TYPE "polly-codegen-ppcg"
 
+#define GPU_RUNTIME_NONE 0
+#define GPU_RUNTIME_CUDA 1
+#define GPU_RUNTIME_OPENCL 2
+
 static cl::opt<bool> DumpSchedule("polly-acc-dump-schedule",
                                   cl::desc("Dump the computed GPU Schedule"),
                                   cl::Hidden, cl::init(false), cl::ZeroOrMore,
                                   cl::cat(PollyCategory));
 
 static cl::opt<bool>
     DumpCode("polly-acc-dump-code",
              cl::desc("Dump C code describing the GPU mapping"), cl::Hidden,
              cl::init(false), cl::ZeroOrMore, cl::cat(PollyCategory));

 /// for generating GPU specific user nodes.
 ///
 /// @see GPUNodeBuilder::createUser
 class GPUNodeBuilder : public IslNodeBuilder {
 public:
   GPUNodeBuilder(PollyIRBuilder &Builder, ScopAnnotator &Annotator,
                  const DataLayout &DL, LoopInfo &LI, ScalarEvolution &SE,
                  DominatorTree &DT, Scop &S, BasicBlock *StartBlock,
-                 gpu_prog *Prog)
+                 gpu_prog *Prog, int GPURuntime)
       : IslNodeBuilder(Builder, Annotator, DL, LI, SE, DT, S, StartBlock),
-        Prog(Prog) {
+        Prog(Prog), Runtime(GPURuntime) {
     getExprBuilder().setIDToSAI(&IDToSAI);
   }
 
   /// Create after-run-time-check initialization code.
   void initializeAfterRTH();
 
   /// Finalize the generated scop.
   virtual void finalize();
private:
   /// A module containing GPU code.
   ///
   /// This pointer is only set in case we are currently generating GPU code.
   std::unique_ptr<Module> GPUModule;
 
   /// The GPU program we generate code for.
   gpu_prog *Prog;
 
+  /// The GPU Runtime implementation to use (OpenCL or CUDA).
+  int Runtime;
+
   /// Class to free isl_ids.
   class IslIdDeleter {
   public:
     void operator()(__isl_take isl_id *Id) { isl_id_free(Id); };
   };
 
   /// A set containing all isl_ids allocated in a GPU kernel.
   ///
if (!F) {
     FunctionType *Ty = FunctionType::get(Builder.getVoidTy(), Args, false);
     F = Function::Create(Ty, Linkage, Name, M);
   }
 
   Builder.CreateCall(F, {DeviceData, HostData, Size});
 }
 
 Value *GPUNodeBuilder::createCallInitContext() {
-  const char *Name = "polly_initContext";
+  assert(Runtime != GPU_RUNTIME_NONE && "GPU Runtime not set");
+  const char *Name;
+
+  if (Runtime == GPU_RUNTIME_CUDA)
+    Name = "polly_initContextCUDA";
+  else if (Runtime == GPU_RUNTIME_OPENCL)
+    Name = "polly_initContextCL";
+
   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;
     FunctionType *Ty = FunctionType::get(Builder.getInt8PtrTy(), Args, false);
for (int i = 0; i < Kernel->n_var; ++i) {
     LocalArrays.push_back(Allocation);
     KernelIds.push_back(Id);
     IDToSAI[Id] = SAI;
   }
 }
 
 void GPUNodeBuilder::createKernelFunction(ppcg_kernel *Kernel,
                                           SetVector<Value *> &SubtreeValues) {
+  assert(Runtime != GPU_RUNTIME_NONE && "GPU Runtime not set");
 
   std::string Identifier = "kernel_" + std::to_string(Kernel->id);
   GPUModule.reset(new Module(Identifier, Builder.getContext()));
+
+  if (GPUNodeBuilder::Runtime == GPU_RUNTIME_CUDA)
     GPUModule->setTargetTriple(Triple::normalize("nvptx64-nvidia-cuda"));
+  else if (GPUNodeBuilder::Runtime == GPU_RUNTIME_OPENCL)
+    GPUModule->setTargetTriple(Triple::normalize("nvptx64-nvidia-nvcl"));
+
   GPUModule->setDataLayout(computeNVPTXDataLayout(true /* is64Bit */));
 
   Function *FN = createKernelFunctionDecl(Kernel, SubtreeValues);
 
   BasicBlock *PrevBlock = Builder.GetInsertBlock();
   auto EntryBlock = BasicBlock::Create(Builder.getContext(), "entry", FN);
 
   DT.addNewBlock(EntryBlock, PrevBlock);
 
   Builder.SetInsertPoint(EntryBlock);
   Builder.CreateRetVoid();
   Builder.SetInsertPoint(EntryBlock, EntryBlock->begin());
 
   ScopDetection::markFunctionAsInvalid(FN);
 
   prepareKernelArguments(Kernel, FN);
   createKernelVariables(Kernel, FN);
   insertKernelIntrinsics(Kernel);
 }
 
 std::string GPUNodeBuilder::createKernelASM() {
-  llvm::Triple GPUTriple(Triple::normalize("nvptx64-nvidia-cuda"));
+  assert(Runtime != GPU_RUNTIME_NONE && "GPU Runtime not set");
+  llvm::Triple GPUTriple;
+
+  if (GPUNodeBuilder::Runtime == GPU_RUNTIME_CUDA)
+    GPUTriple = llvm::Triple(Triple::normalize("nvptx64-nvidia-cuda"));
+  else if (GPUNodeBuilder::Runtime == GPU_RUNTIME_OPENCL)
+    GPUTriple = llvm::Triple(Triple::normalize("nvptx64-nvidia-nvcl"));
+
   std::string ErrMsg;
   auto GPUTarget = TargetRegistry::lookupTarget(GPUTriple.getTriple(), ErrMsg);
 
   if (!GPUTarget) {
     errs() << ErrMsg << "\n";
     return "";
   }
 
std::string GPUNodeBuilder::finalizeKernelFunction() {
 
   return Assembly;
 }
 
 namespace {
 class PPCGCodeGeneration : public ScopPass {
 public:
   static char ID;
+  static int Runtime;
 
   /// The scop that is currently processed.
   Scop *S;
 
   LoopInfo *LI;
   DominatorTree *DT;
   ScalarEvolution *SE;
   const DataLayout *DL;
void generateCode(__isl_take isl_ast_node *Root, gpu_prog *Prog) {
     // branch will guard the original scop from new induction variables that
     // the SCEVExpander may introduce while code generating the parameters and
     // which may introduce scalar dependences that prevent us from correctly
     // code generating this scop.
     BasicBlock *StartBlock =
         executeScopConditionally(*S, Builder.getTrue(), *DT, *RI, *LI);
 
     GPUNodeBuilder NodeBuilder(Builder, Annotator, *DL, *LI, *SE, *DT, *S,
-                               StartBlock, Prog);
+                               StartBlock, Prog, Runtime);
 
     // TODO: Handle LICM
     auto SplitBlock = StartBlock->getSinglePredecessor();
     Builder.SetInsertPoint(SplitBlock->getTerminator());
     NodeBuilder.addParameters(S->getContext());
 
     isl_ast_build *Build = isl_ast_build_alloc(S->getIslCtx());
     isl_ast_expr *Condition = IslAst::buildRunCondition(S, Build);
void getAnalysisUsage(AnalysisUsage &AU) const override {
     //        region tree.
     AU.addPreserved<RegionInfoPass>();
     AU.addPreserved<ScopInfoRegionPass>();
   }
 };
 } // namespace
 
 char PPCGCodeGeneration::ID = 1;
+int PPCGCodeGeneration::Runtime = GPU_RUNTIME_NONE;
 
-Pass *polly::createPPCGCodeGenerationPass() { return new PPCGCodeGeneration(); }
+Pass *polly::createPPCGCodeGenerationPassCUDA() {
+  PPCGCodeGeneration::Runtime = GPU_RUNTIME_CUDA;
+  return new PPCGCodeGeneration();
+}
+
+Pass *polly::createPPCGCodeGenerationPassOpenCL() {
+  PPCGCodeGeneration::Runtime = GPU_RUNTIME_OPENCL;
+  return new PPCGCodeGeneration();
+}
 
 INITIALIZE_PASS_BEGIN(PPCGCodeGeneration, "polly-codegen-ppcg",
                       "Polly - Apply PPCG translation to SCOP", false, false)
 INITIALIZE_PASS_DEPENDENCY(DependenceInfo);
 INITIALIZE_PASS_DEPENDENCY(DominatorTreeWrapperPass);
 INITIALIZE_PASS_DEPENDENCY(LoopInfoWrapperPass);
 INITIALIZE_PASS_DEPENDENCY(RegionInfoPass);
 INITIALIZE_PASS_DEPENDENCY(ScalarEvolutionWrapperPass);
 INITIALIZE_PASS_DEPENDENCY(ScopDetection);
 INITIALIZE_PASS_END(PPCGCodeGeneration, "polly-codegen-ppcg",
                     "Polly - Apply PPCG translation to SCOP", false, false)

lib/Support/RegisterPasses.cpp

 enum CodeGenChoice { CODEGEN_FULL, CODEGEN_AST, CODEGEN_NONE };
 static cl::opt<CodeGenChoice> CodeGeneration(
     "polly-code-generation", cl::desc("How much code-generation to perform"),
     cl::values(clEnumValN(CODEGEN_FULL, "full", "AST and IR generation"),
                clEnumValN(CODEGEN_AST, "ast", "Only AST generation"),
                clEnumValN(CODEGEN_NONE, "none", "No code generation")),
     cl::Hidden, cl::init(CODEGEN_FULL), cl::ZeroOrMore, cl::cat(PollyCategory));
 
-enum TargetChoice { TARGET_CPU, TARGET_GPU };
-static cl::opt<TargetChoice>
-    Target("polly-target", cl::desc("The hardware to target"),
-           cl::values(clEnumValN(TARGET_CPU, "cpu", "generate CPU code")
+enum TargetChoice { TARGET_CPU, TARGET_GPU_OPENCL, TARGET_GPU_CUDA };
+static cl::opt<TargetChoice> Target(
+    "polly-target", cl::desc("The hardware to target"),
 #ifdef GPU_CODEGEN
-                          ,
-                      clEnumValN(TARGET_GPU, "gpu", "generate GPU code")
+               clEnumValN(TARGET_GPU_OPENCL, "gpu_opencl",
+                          "generate OpenCL GPU code"),
+               clEnumValN(TARGET_GPU_CUDA, "gpu_cuda", "generate CUDA GPU code")
 #endif
-                          ),
-           cl::init(TARGET_CPU), cl::ZeroOrMore, cl::cat(PollyCategory));
 
 VectorizerChoice polly::PollyVectorizerChoice;
 static cl::opt<polly::VectorizerChoice, true> Vectorizer(
     "polly-vectorizer", cl::desc("Select the vectorization strategy"),
     cl::values(
         clEnumValN(polly::VECTORIZER_NONE, "none", "No Vectorization"),
         clEnumValN(polly::VECTORIZER_POLLY, "polly",
                    "Polly internal vectorizer"),
if (ImportJScop)
     PM.add(polly::createJSONImporterPass());
 
   if (DeadCodeElim)
     PM.add(polly::createDeadCodeElimPass());
 
   if (EnablePruneUnprofitable)
     PM.add(polly::createPruneUnprofitablePass());
 
-  if (Target == TARGET_GPU) {
+  if (Target == TARGET_GPU_CUDA || Target == TARGET_GPU_OPENCL) {
     // GPU generation provides its own scheduling optimization strategy.
   } else {
     switch (Optimizer) {
     case OPTIMIZER_NONE:
       break; /* Do nothing */
 
     case OPTIMIZER_ISL:
       PM.add(polly::createIslScheduleOptimizerPass());
       break;
     }
   }
 
   if (ExportJScop)
     PM.add(polly::createJSONExporterPass());
 
-  if (Target == TARGET_GPU) {
+  if (Target == TARGET_GPU_CUDA) {
 #ifdef GPU_CODEGEN
-    PM.add(polly::createPPCGCodeGenerationPass());
+    PM.add(polly::createPPCGCodeGenerationPassCUDA());
+#endif
+  } else if (Target == TARGET_GPU_OPENCL) {
+#ifdef GPU_CODEGEN
+    PM.add(polly::createPPCGCodeGenerationPassOpenCL());
 #endif
   } else {
     switch (CodeGeneration) {
     case CODEGEN_AST:
       PM.add(polly::createIslAstInfoPass());
       break;
     case CODEGEN_FULL:
       PM.add(polly::createCodeGenerationPass());
       break;
#endif
   if (DumpAfter)
     PM.add(polly::createDumpModulePass("-after", true));
   for (auto &Filename : DumpAfterFile)
     PM.add(polly::createDumpModulePass(Filename, false));
 
   if (CFGPrinter)
     PM.add(llvm::createCFGPrinterLegacyPassPass());
 
-  if (Target == TARGET_GPU) {
+  if (Target == TARGET_GPU_CUDA || Target == TARGET_GPU_OPENCL) {
     // Invariant load hoisting not yet supported by GPU code generation.
     PollyInvariantLoadHoisting = false;
   }
 }
 
 static bool shouldEnablePolly() {
   if (PollyOnlyPrinter || PollyPrinter || PollyOnlyViewer || PollyViewer)
     PollyTrackFailures = true;

tools/CMakeLists.txt

-if (CUDALIB_FOUND)
+if (CUDALIB_FOUND AND OpenCL_FOUND)
   add_subdirectory(GPURuntime)
-endif (CUDALIB_FOUND)
+endif (CUDALIB_FOUND AND OpenCL_FOUND)
 
 set(LLVM_COMMON_DEPENDS ${LLVM_COMMON_DEPENDS} PARENT_SCOPE)

tools/GPURuntime/GPUJIT.h

  *   polly_copyFromDeviceToHost(HostData, DevData, MemSize);
  *   polly_freeKernel(Kernel);
  *   polly_freeDeviceMemory(DevArray);
  *   polly_freeContext(Context);
  * }
  *
  */
 
+typedef enum PollyGPURuntimeT {
+  RUNTIME_NONE,
+  RUNTIME_CUDA,
+  RUNTIME_CL
+} PollyGPURuntime;
+
 typedef struct PollyGPUContextT PollyGPUContext;
 typedef struct PollyGPUFunctionT PollyGPUFunction;
 typedef struct PollyGPUDevicePtrT PollyGPUDevicePtr;
 
-PollyGPUContext *polly_initContext();
+typedef struct OpenCLContextT OpenCLContext;
+typedef struct OpenCLKernelT OpenCLKernel;
+typedef struct OpenCLDevicePtrT OpenCLDevicePtr;
+
+typedef struct CUDAContextT CUDAContext;
+typedef struct CUDAKernelT CUDAKernel;
+typedef struct CUDADevicePtrT CUDADevicePtr;
+
+PollyGPUContext *polly_initContextCUDA();
+PollyGPUContext *polly_initContextCL();
 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_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

 /******************** GPUJIT.c - GPUJIT Execution Engine **********************/
 /*                                                                            */
 /*                     The LLVM Compiler Infrastructure                       */
 /*                                                                            */
 /* This file is dual licensed under the MIT and the University of Illinois    */
 /* Open Source License. See LICENSE.TXT for details.                          */
 /*                                                                            */
 /******************************************************************************/
 /*                                                                            */
 /*  This file implements GPUJIT, a ptx string execution engine for GPU.       */
 /*                                                                            */
 /******************************************************************************/
 
 #include "GPUJIT.h"
 #include <cuda.h>
 #include <cuda_runtime.h>
+
+#ifdef __APPLE__
+#include <OpenCL/opencl.h>
+#else
+#include <CL/cl.h>
+#endif
+
 #include <dlfcn.h>
 #include <stdarg.h>
 #include <stdio.h>
 #include <string.h>
 
 static int DebugMode;
 static int CacheMode;
 
+static PollyGPURuntime Runtime = RUNTIME_NONE;
+
 static void debug_print(const char *format, ...) {
   if (!DebugMode)
     return;
 
   va_list args;
   va_start(args, format);
   vfprintf(stderr, format, args);
   va_end(args);
 }
 #define dump_function() debug_print("-> %s\n", __func__)
 
-/* Define Polly's GPGPU data types. */
+#define KERNEL_CACHE_SIZE 10
+
+static void err_runtime() {
+  fprintf(stdout, "Runtime not correctly initialized.\n");
+  exit(-1);
+}
+
 struct PollyGPUContextT {
-  CUcontext Cuda;
+  void *Context;
 };
 
 struct PollyGPUFunctionT {
+  void *Kernel;
+};
+
+struct PollyGPUDevicePtrT {
+  void *DevicePtr;
+};
+
+/******************************************************************************/
+/*                                  OpenCL                                    */
+/******************************************************************************/
+
+struct OpenCLContextT {
+  cl_context context;
+  cl_command_queue command_queue;
+};
+
+struct OpenCLKernelT {
+  cl_kernel kernel;
+  cl_program program;
+  const char *PTXString;
+};
+
+struct OpenCLDevicePtrT {
+  cl_mem mem;
+};
+
+/* Dynamic library handles for the OpenCL runtime library. */
+static void *HandleOpenCL;
+
+/* Type-defines of function pointer to OpenCL Runtime API. */
+typedef cl_int clGetPlatformIDsFcnTy(cl_uint num_entries,
+                                     cl_platform_id *platforms,
+                                     cl_uint *num_platforms);
+static clGetPlatformIDsFcnTy *clGetPlatformIDsFcnPtr;
+
+typedef cl_int clGetDeviceIDsFcnTy(cl_platform_id platform,
+                                   cl_device_type device_type,
+                                   cl_uint num_entries, cl_device_id *devices,
+                                   cl_uint *num_devices);
+static clGetDeviceIDsFcnTy *clGetDeviceIDsFcnPtr;
+
+typedef cl_int clGetDeviceInfoFcnTy(cl_device_id device,
+                                    cl_device_info param_name,
+                                    size_t param_value_size, void *param_value,
+                                    size_t *param_value_size_ret);
+static clGetDeviceInfoFcnTy *clGetDeviceInfoFcnPtr;
+
+typedef cl_int clGetKernelInfoFcnTy(cl_kernel kernel, cl_kernel_info param_name,
+                                    size_t param_value_size, void *param_value,
+                                    size_t *param_value_size_ret);
+static clGetKernelInfoFcnTy *clGetKernelInfoFcnPtr;
+
+typedef cl_context clCreateContextFcnTy(
+    const cl_context_properties *properties, cl_uint num_devices,
+    const cl_device_id *devices,
+    void CL_CALLBACK *pfn_notify(const char *errinfo, const void *private_info,
+                                 size_t cb, void *user_data),
+    void *user_data, cl_int *errcode_ret);
+static clCreateContextFcnTy *clCreateContextFcnPtr;
+
+typedef cl_command_queue
+clCreateCommandQueueFcnTy(cl_context context, cl_device_id device,
+                          cl_command_queue_properties properties,
+                          cl_int *errcode_ret);
+static clCreateCommandQueueFcnTy *clCreateCommandQueueFcnPtr;
+
+typedef cl_mem clCreateBufferFcnTy(cl_context context, cl_mem_flags flags,
+                                   size_t size, void *host_ptr,
+                                   cl_int *errcode_ret);
+static clCreateBufferFcnTy *clCreateBufferFcnPtr;
+
+typedef cl_int
+clEnqueueWriteBufferFcnTy(cl_command_queue command_queue, cl_mem buffer,
+                          cl_bool blocking_write, size_t offset, size_t size,
+                          const void *ptr, cl_uint num_events_in_wait_list,
+                          const cl_event *event_wait_list, cl_event *event);
+static clEnqueueWriteBufferFcnTy *clEnqueueWriteBufferFcnPtr;
+
+typedef cl_program clCreateProgramWithBinaryFcnTy(
+    cl_context context, cl_uint num_devices, const cl_device_id *device_list,
+    const size_t *lengths, const unsigned char **binaries,
+    cl_int *binary_status, cl_int *errcode_ret);
+static clCreateProgramWithBinaryFcnTy *clCreateProgramWithBinaryFcnPtr;
+
+typedef cl_int clBuildProgramFcnTy(
+    cl_program program, cl_uint num_devices, const cl_device_id *device_list,
+    const char *options,
+    void(CL_CALLBACK *pfn_notify)(cl_program program, void *user_data),
+    void *user_data);
+static clBuildProgramFcnTy *clBuildProgramFcnPtr;
+
+typedef cl_kernel clCreateKernelFcnTy(cl_program program,
+                                      const char *kernel_name,
+                                      cl_int *errcode_ret);
+static clCreateKernelFcnTy *clCreateKernelFcnPtr;
+
+typedef cl_int clSetKernelArgFcnTy(cl_kernel kernel, cl_uint arg_index,
+                                   size_t arg_size, const void *arg_value);
+static clSetKernelArgFcnTy *clSetKernelArgFcnPtr;
+
+typedef cl_int clEnqueueNDRangeKernelFcnTy(
+    cl_command_queue command_queue, cl_kernel kernel, cl_uint work_dim,
+    const size_t *global_work_offset, const size_t *global_work_size,
+    const size_t *local_work_size, cl_uint num_events_in_wait_list,
+    const cl_event *event_wait_list, cl_event *event);
+static clEnqueueNDRangeKernelFcnTy *clEnqueueNDRangeKernelFcnPtr;
+
+typedef cl_int clEnqueueReadBufferFcnTy(cl_command_queue command_queue,
+                                        cl_mem buffer, cl_bool blocking_read,
+                                        size_t offset, size_t size, void *ptr,
+                                        cl_uint num_events_in_wait_list,
+                                        const cl_event *event_wait_list,
+                                        cl_event *event);
+static clEnqueueReadBufferFcnTy *clEnqueueReadBufferFcnPtr;
+
+typedef cl_int clFlushFcnTy(cl_command_queue command_queue);
+static clFlushFcnTy *clFlushFcnPtr;
+
+typedef cl_int clFinishFcnTy(cl_command_queue command_queue);
+static clFinishFcnTy *clFinishFcnPtr;
+
+typedef cl_int clReleaseKernelFcnTy(cl_kernel kernel);
+static clReleaseKernelFcnTy *clReleaseKernelFcnPtr;
+
+typedef cl_int clReleaseProgramFcnTy(cl_program program);
+static clReleaseProgramFcnTy *clReleaseProgramFcnPtr;
+
+typedef cl_int clReleaseMemObjectFcnTy(cl_mem memobject);
+static clReleaseMemObjectFcnTy *clReleaseMemObjectFcnPtr;
+
+typedef cl_int clReleaseCommandQueueFcnTy(cl_command_queue command_queue);
+static clReleaseCommandQueueFcnTy *clReleaseCommandQueueFcnPtr;
+
+typedef cl_int clReleaseContextFcnTy(cl_context context);
+static clReleaseContextFcnTy *clReleaseContextFcnPtr;
+
+static void *getAPIHandleCL(void *Handle, const char *FuncName) {
+  char *Err;
+  void *FuncPtr;
+  dlerror();
+  FuncPtr = dlsym(Handle, FuncName);
+  if ((Err = dlerror()) != 0) {
+    fprintf(stdout, "Load OpenCL Runtime API failed: %s. \n", Err);
+    return 0;
+  }
+  return FuncPtr;
+}
+
+static int initialDeviceAPILibrariesCL() {
+  HandleOpenCL = dlopen("libOpenCL.so", RTLD_LAZY);
+  if (!HandleOpenCL) {
+    printf("Cannot open library: %s. \n", dlerror());
+    return 0;
+  }
+  return 1;
+}
+
+static int initialDeviceAPIsCL() {
+  if (initialDeviceAPILibrariesCL() == 0)
+    return 0;
+
+  /* Get function pointer to OpenCL Runtime API.
+   *
+   * Note that compilers conforming to the ISO C standard are required to
+   * generate a warning if a conversion from a void * pointer to a function
+   * pointer is attempted as in the following statements. The warning
+   * of this kind of cast may not be emitted by clang and new versions of gcc
+   * as it is valid on POSIX 2008.
+   */
+  clGetPlatformIDsFcnPtr =
+      (clGetPlatformIDsFcnTy *)getAPIHandleCL(HandleOpenCL, "clGetPlatformIDs");
+
+  clGetDeviceIDsFcnPtr =
+      (clGetDeviceIDsFcnTy *)getAPIHandleCL(HandleOpenCL, "clGetDeviceIDs");
+
+  clGetDeviceInfoFcnPtr =
+      (clGetDeviceInfoFcnTy *)getAPIHandleCL(HandleOpenCL, "clGetDeviceInfo");
+
+  clGetKernelInfoFcnPtr =
+      (clGetKernelInfoFcnTy *)getAPIHandleCL(HandleOpenCL, "clGetKernelInfo");
+
+  clCreateContextFcnPtr =
+      (clCreateContextFcnTy *)getAPIHandleCL(HandleOpenCL, "clCreateContext");
+
+  clCreateCommandQueueFcnPtr = (clCreateCommandQueueFcnTy *)getAPIHandleCL(
+      HandleOpenCL, "clCreateCommandQueue");
+
+  clCreateBufferFcnPtr =
+      (clCreateBufferFcnTy *)getAPIHandleCL(HandleOpenCL, "clCreateBuffer");
+
+  clEnqueueWriteBufferFcnPtr = (clEnqueueWriteBufferFcnTy *)getAPIHandleCL(
+      HandleOpenCL, "clEnqueueWriteBuffer");
+
+  clCreateProgramWithBinaryFcnPtr =
+      (clCreateProgramWithBinaryFcnTy *)getAPIHandleCL(
+          HandleOpenCL, "clCreateProgramWithBinary");
+
+  clBuildProgramFcnPtr =
+      (clBuildProgramFcnTy *)getAPIHandleCL(HandleOpenCL, "clBuildProgram");
+
+  clCreateKernelFcnPtr =
+      (clCreateKernelFcnTy *)getAPIHandleCL(HandleOpenCL, "clCreateKernel");
+
+  clSetKernelArgFcnPtr =
+      (clSetKernelArgFcnTy *)getAPIHandleCL(HandleOpenCL, "clSetKernelArg");
+
+  clEnqueueNDRangeKernelFcnPtr = (clEnqueueNDRangeKernelFcnTy *)getAPIHandleCL(
+      HandleOpenCL, "clEnqueueNDRangeKernel");
+
+  clEnqueueReadBufferFcnPtr = (clEnqueueReadBufferFcnTy *)getAPIHandleCL(
+      HandleOpenCL, "clEnqueueReadBuffer");
+
+  clFlushFcnPtr = (clFlushFcnTy *)getAPIHandleCL(HandleOpenCL, "clFlush");
+
+  clFinishFcnPtr = (clFinishFcnTy *)getAPIHandleCL(HandleOpenCL, "clFinish");
+
+  clReleaseKernelFcnPtr =
+      (clReleaseKernelFcnTy *)getAPIHandleCL(HandleOpenCL, "clReleaseKernel");
+
+  clReleaseProgramFcnPtr =
+      (clReleaseProgramFcnTy *)getAPIHandleCL(HandleOpenCL, "clReleaseProgram");
+
+  clReleaseMemObjectFcnPtr = (clReleaseMemObjectFcnTy *)getAPIHandleCL(
+      HandleOpenCL, "clReleaseMemObject");
+
+  clReleaseCommandQueueFcnPtr = (clReleaseCommandQueueFcnTy *)getAPIHandleCL(
+      HandleOpenCL, "clReleaseCommandQueue");
+
+  clReleaseContextFcnPtr =
+      (clReleaseContextFcnTy *)getAPIHandleCL(HandleOpenCL, "clReleaseContext");
+
+  return 1;
+}
+
+/* Context and Device. */
+static PollyGPUContext *GlobalContext = NULL;
+static cl_device_id GlobalDeviceID = NULL;
+
+/* Fd-Decl: Print out OpenCL Error codes to human readable strings. */
+void cl_printError(int error);
+
+PollyGPUContext *initContextCL() {
+  dump_function();
+
+  PollyGPUContext *Context;
+
+  cl_platform_id platform_id = NULL;
+  cl_device_id device_id = NULL;
+  cl_uint num_platforms;
+  cl_uint num_devices;
+  cl_int ret;
+
+  char DeviceRevision[256];
+  char DeviceName[256];
+  size_t DeviceRevisionRetSize, DeviceNameRetSize;
+
+  static __thread PollyGPUContext *CurrentContext = NULL;
+
+  if (CurrentContext)
+    return CurrentContext;
+
+  /* Get API handles. */
+  if (initialDeviceAPIsCL() == 0) {
+    fprintf(stdout, "Getting the \"handle\" for the OpenCL Runtime failed.\n");
+    exit(-1);
+  }
+
+  /* Get number of devices that support OpenCL. */
+  ret = clGetPlatformIDsFcnPtr(1, &platform_id, &num_platforms);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Failed to get platform IDs.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+  // TODO: Extend to CL_DEVICE_TYPE_ALL?
+  ret = clGetDeviceIDsFcnPtr(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id,
+                             &num_devices);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Failed to get device IDs.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+
+  GlobalDeviceID = device_id;
+  if (num_devices == 0) {
+    fprintf(stdout, "There is no device supporting OpenCL.\n");
+    exit(-1);
+  }
+
+  /* Get device revision. */
+  ret = clGetDeviceInfoFcnPtr(device_id, CL_DEVICE_VERSION, 256, DeviceRevision,
+                              &DeviceRevisionRetSize);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Failed to fetch device revision.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+
+  /* Get device name. */
+  ret = clGetDeviceInfoFcnPtr(device_id, CL_DEVICE_NAME, 256, DeviceName,
+                              &DeviceNameRetSize);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Failed to fetch device name.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+
+  debug_print("> Running on GPU device %d : %s.\n", device_id, DeviceName);
+
+  /* Create context on the device. */
+  Context = (PollyGPUContext *)malloc(sizeof(PollyGPUContext));
+  if (Context == 0) {
+    fprintf(stdout, "Allocate memory for Polly GPU context failed.\n");
+    exit(-1);
+  }
+  Context->Context = (OpenCLContext *)malloc(sizeof(OpenCLContext));
+  if (Context->Context == 0) {
+    fprintf(stdout, "Allocate memory for Polly OpenCL context failed.\n");
+    exit(-1);
+  }
+  ((OpenCLContext *)Context->Context)->context =
+      clCreateContextFcnPtr(NULL, 1, &device_id, NULL, NULL, &ret);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Failed to create context.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+
+  ((OpenCLContext *)Context->Context)->command_queue =
+      clCreateCommandQueueFcnPtr(((OpenCLContext *)Context->Context)->context,
+                                 device_id, 0, &ret);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Failed to create command queue.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+
+  if (CacheMode)
+    CurrentContext = Context;
+
+  GlobalContext = Context;
+  return Context;
+}
+
+static void freeKernelCL(PollyGPUFunction *Kernel) {
+  dump_function();
+
+  if (CacheMode)
+    return;
+
+  if (!GlobalContext) {
+    fprintf(stdout, "GPGPU-code generation not correctly initialized.\n");
+    exit(-1);
+  }
+
+  cl_int ret;
+  ret = clFlushFcnPtr(((OpenCLContext *)GlobalContext->Context)->command_queue);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Failed to flush command queue.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+  ret =
+      clFinishFcnPtr(((OpenCLContext *)GlobalContext->Context)->command_queue);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Failed to flush command queue.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+
+  if (((OpenCLKernel *)Kernel->Kernel)->kernel) {
+    cl_int ret =
+        clReleaseKernelFcnPtr(((OpenCLKernel *)Kernel->Kernel)->kernel);
+    if (ret != CL_SUCCESS) {
+      fprintf(stdout, "Failed to release kernel.\n");
+      cl_printError(ret);
+      exit(-1);
+    }
+  }
+
+  if (((OpenCLKernel *)Kernel->Kernel)->program) {
+    cl_int ret =
+        clReleaseProgramFcnPtr(((OpenCLKernel *)Kernel->Kernel)->program);
+    if (ret != CL_SUCCESS) {
+      fprintf(stdout, "Failed to release program.\n");
+      cl_printError(ret);
+      exit(-1);
+    }
+  }
+
+  if (Kernel->Kernel)
+    free((OpenCLKernel *)Kernel->Kernel);
+
+  if (Kernel)
+    free(Kernel);
+}
+
+PollyGPUFunction *getKernelCL(const char *PTXBuffer, const char *KernelName) {
+  dump_function();
+
+  if (!GlobalContext) {
+    fprintf(stdout, "GPGPU-code generation not correctly initialized.\n");
+    exit(-1);
+  }
+
+  static __thread PollyGPUFunction *KernelCache[KERNEL_CACHE_SIZE];
+  static __thread int NextCacheItem = 0;
+
+  for (long i = 0; i < KERNEL_CACHE_SIZE; i++) {
+    // We exploit here the property that all Polly-ACC kernels are allocated
+    // as global constants, hence a pointer comparision is sufficient to
+    // determin equality.
+    if (KernelCache[i] &&
+        ((OpenCLKernel *)KernelCache[i]->Kernel)->PTXString == PTXBuffer) {
+      debug_print("  -> using cached kernel\n");
+      return KernelCache[i];
+    }
+  }
+
+  PollyGPUFunction *Function = malloc(sizeof(PollyGPUFunction));
+  if (Function == 0) {
+    fprintf(stdout, "Allocate memory for Polly GPU function failed.\n");
+    exit(-1);
+  }
+  Function->Kernel = (OpenCLKernel *)malloc(sizeof(OpenCLKernel));
+  if (Function->Kernel == 0) {
+    fprintf(stdout, "Allocate memory for Polly OpenCL kernel failed.\n");
+    exit(-1);
+  }
+
+  if (!GlobalDeviceID) {
+    fprintf(stdout, "GPGPU-code generation not initialized correctly.\n");
+    exit(-1);
+  }
+
+  cl_int bin_status;
+  cl_int ret;
+  size_t binary_size = strlen(PTXBuffer);
+  ((OpenCLKernel *)Function->Kernel)->program = clCreateProgramWithBinaryFcnPtr(
+      ((OpenCLContext *)GlobalContext->Context)->context, 1, &GlobalDeviceID,
+      (const size_t *)&binary_size, (const unsigned char **)&PTXBuffer,
+      &bin_status, &ret);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Failed to create program from binary.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+
+  ret = clBuildProgramFcnPtr(((OpenCLKernel *)Function->Kernel)->program, 1,
+                             &GlobalDeviceID, NULL, NULL, NULL);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Failed to build program.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+
+  ((OpenCLKernel *)Function->Kernel)->kernel = clCreateKernelFcnPtr(
+      ((OpenCLKernel *)Function->Kernel)->program, KernelName, &ret);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Failed to create kernel.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+
+  ((OpenCLKernel *)Function->Kernel)->PTXString = PTXBuffer;
+
+  if (CacheMode) {
+    if (KernelCache[NextCacheItem])
+      freeKernelCL(KernelCache[NextCacheItem]);
+
+    KernelCache[NextCacheItem] = Function;
+
+    NextCacheItem = (NextCacheItem + 1) % KERNEL_CACHE_SIZE;
+  }
+
+  return Function;
+}
+
+void copyFromHostToDeviceCL(void *HostData, PollyGPUDevicePtr *DevData,
+                            long MemSize) {
+  dump_function();
+
+  if (!GlobalContext) {
+    fprintf(stdout, "GPGPU-code generation not correctly initialized.\n");
+    exit(-1);
+  }
+
+  cl_int ret;
+  ret = clEnqueueWriteBufferFcnPtr(
+      ((OpenCLContext *)GlobalContext->Context)->command_queue,
+      ((OpenCLDevicePtr *)DevData->DevicePtr)->mem, CL_TRUE, 0, MemSize,
+      HostData, 0, NULL, NULL);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Copying data from host memory to device failed.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+}
+
+void copyFromDeviceToHostCL(PollyGPUDevicePtr *DevData, void *HostData,
+                            long MemSize) {
+  dump_function();
+
+  if (!GlobalContext) {
+    fprintf(stdout, "GPGPU-code generation not correctly initialized.\n");
+    exit(-1);
+  }
+
+  cl_int ret;
+  ret = clEnqueueReadBufferFcnPtr(
+      ((OpenCLContext *)GlobalContext->Context)->command_queue,
+      ((OpenCLDevicePtr *)DevData->DevicePtr)->mem, CL_TRUE, 0, MemSize,
+      HostData, 0, NULL, NULL);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Copying results from device to host memory failed.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+}
+
+void launchKernelCL(PollyGPUFunction *Kernel, unsigned int GridDimX,
+                    unsigned int GridDimY, unsigned int BlockDimX,
+                    unsigned int BlockDimY, unsigned int BlockDimZ,
+                    void **Parameters) {
+  dump_function();
+
+  cl_int ret;
+  cl_uint num_args;
+
+  if (!GlobalContext) {
+    fprintf(stdout, "GPGPU-code generation not correctly initialized.\n");
+    exit(-1);
+  }
+
+  OpenCLKernel *CLKernel = (OpenCLKernel *)Kernel->Kernel;
+  ret = clGetKernelInfoFcnPtr(CLKernel->kernel, CL_KERNEL_NUM_ARGS,
+                              sizeof(cl_uint), &num_args, NULL);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Failed to get number of kernel arguments.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+
+  for (cl_uint i = 0; i < num_args; i++) {
+    debug_print("Setting arg %d.\n", i);
+    ret = clSetKernelArgFcnPtr(CLKernel->kernel, i, 8, (void *)Parameters[i]);
+    if (ret == CL_INVALID_ARG_SIZE) {
+      debug_print("Arg %d with 8 bytes failed, trying 4.\n", i);
+      ret = clSetKernelArgFcnPtr(CLKernel->kernel, i, 4, (void *)Parameters[i]);
+      if (ret == CL_INVALID_ARG_SIZE) {
+        debug_print("Arg %d with 4 bytes failed, trying 2.\n", i);
+        ret =
+            clSetKernelArgFcnPtr(CLKernel->kernel, i, 2, (void *)Parameters[i]);
+        if (ret == CL_INVALID_ARG_SIZE) {
+          debug_print("Arg %d with 2 bytes failed, trying 1.\n", i);
+          ret = clSetKernelArgFcnPtr(CLKernel->kernel, i, 1,
+                                     (void *)Parameters[i]);
+          if (ret != CL_SUCCESS) {
+            fprintf(stdout, "Failed to set Kernel argument %d.\n", i);
+            cl_printError(ret);
+            exit(-1);
+          }
+        }
+      }
+    }
+    if (ret != CL_SUCCESS && ret != CL_INVALID_ARG_SIZE) {
+      fprintf(stdout, "Failed to set Kernel argument.\n");
+      cl_printError(ret);
+      exit(-1);
+    }
+  }
+
+  unsigned int GridDimZ = 1;
+  size_t global_work_size[3] = {BlockDimX * GridDimX, BlockDimY * GridDimY,
+                                BlockDimZ * GridDimZ};
+  size_t local_work_size[3] = {BlockDimX, BlockDimY, BlockDimZ};
+
+  OpenCLContext *CLContext = (OpenCLContext *)GlobalContext->Context;
+  ret = clEnqueueNDRangeKernelFcnPtr(CLContext->command_queue, CLKernel->kernel,
+                                     3, NULL, global_work_size, local_work_size,
+                                     0, NULL, NULL);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Launching OpenCL kernel failed.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+}
+
+void freeDeviceMemoryCL(PollyGPUDevicePtr *Allocation) {
+  dump_function();
+
+  OpenCLDevicePtr *DevPtr = (OpenCLDevicePtr *)Allocation->DevicePtr;
+  cl_int ret = clReleaseMemObjectFcnPtr((cl_mem)DevPtr->mem);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Failed to free device memory.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+
+  free(DevPtr);
+  free(Allocation);
+}
+
+PollyGPUDevicePtr *allocateMemoryForDeviceCL(long MemSize) {
+  dump_function();
+
+  if (!GlobalContext) {
+    fprintf(stdout, "GPGPU-code generation not correctly initialized.\n");
+    exit(-1);
+  }
+
+  PollyGPUDevicePtr *DevData = malloc(sizeof(PollyGPUDevicePtr));
+  if (DevData == 0) {
+    fprintf(stdout, "Allocate memory for GPU device memory pointer failed.\n");
+    exit(-1);
+  }
+  DevData->DevicePtr = (OpenCLDevicePtr *)malloc(sizeof(OpenCLDevicePtr));
+  if (DevData->DevicePtr == 0) {
+    fprintf(stdout, "Allocate memory for GPU device memory pointer failed.\n");
+    exit(-1);
+  }
+
+  cl_int ret;
+  ((OpenCLDevicePtr *)DevData->DevicePtr)->mem =
+      clCreateBufferFcnPtr(((OpenCLContext *)GlobalContext->Context)->context,
+                           CL_MEM_READ_WRITE, MemSize, NULL, &ret);
+  if (ret != CL_SUCCESS) {
+    fprintf(stdout, "Allocate memory for GPU device memory pointer failed.\n");
+    cl_printError(ret);
+    exit(-1);
+  }
+
+  return DevData;
+}
+
+void *getDevicePtrCL(PollyGPUDevicePtr *Allocation) {
+  dump_function();
+
+  OpenCLDevicePtr *DevPtr = (OpenCLDevicePtr *)Allocation->DevicePtr;
+  return (void *)DevPtr->mem;
+}
+
+void freeContextCL(PollyGPUContext *Context) {
+  dump_function();
+
+  cl_int ret;
+
+  GlobalContext = NULL;
+
+  OpenCLContext *Ctx = (OpenCLContext *)Context->Context;
+  if (Ctx->command_queue) {
+    ret = clReleaseCommandQueueFcnPtr(Ctx->command_queue);
+    if (ret != CL_SUCCESS) {
+      fprintf(stdout, "Could not release command queue.\n");
+      cl_printError(ret);
+      exit(-1);
+    }
+  }
+
+  if (Ctx->context) {
+    ret = clReleaseContextFcnPtr(Ctx->context);
+    if (ret != CL_SUCCESS) {
+      fprintf(stdout, "Could not release context.\n");
+      cl_printError(ret);
+      exit(-1);
+    }
+  }
+
+  free(Ctx);
+  free(Context);
+}
+
+void cl_printError(int error) {
+
+  switch (error) {
+  case CL_SUCCESS:
+    // Success, don't print an error.
+    break;
+
+  // JIT/Runtime errors.
+  case CL_DEVICE_NOT_FOUND:
+    fprintf(stdout, "Device not found.\n");
+    break;
+  case CL_DEVICE_NOT_AVAILABLE:
+    fprintf(stdout, "Device not available.\n");
+    break;
+  case CL_COMPILER_NOT_AVAILABLE:
+    fprintf(stdout, "Compiler not available.\n");
+    break;
+  case CL_MEM_OBJECT_ALLOCATION_FAILURE:
+    fprintf(stdout, "Mem object allocation failure.\n");
+    break;
+  case CL_OUT_OF_RESOURCES:
+    fprintf(stdout, "Out of resources.\n");
+    break;
+  case CL_OUT_OF_HOST_MEMORY:
+    fprintf(stdout, "Out of host memory.\n");
+    break;
+  case CL_PROFILING_INFO_NOT_AVAILABLE:
+    fprintf(stdout, "Profiling info not available.\n");
+    break;
+  case CL_MEM_COPY_OVERLAP:
+    fprintf(stdout, "Mem copy overlap.\n");
+    break;
+  case CL_IMAGE_FORMAT_MISMATCH:
+    fprintf(stdout, "Image format mismatch.\n");
+    break;
+  case CL_IMAGE_FORMAT_NOT_SUPPORTED:
+    fprintf(stdout, "Image format not supported.\n");
+    break;
+  case CL_BUILD_PROGRAM_FAILURE:
+    fprintf(stdout, "Build program failure.\n");
+    break;
+  case CL_MAP_FAILURE:
+    fprintf(stdout, "Map failure.\n");
+    break;
+  case CL_MISALIGNED_SUB_BUFFER_OFFSET:
+    fprintf(stdout, "Misaligned sub buffer offset.\n");
+    break;
+  case CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST:
+    fprintf(stdout, "Exec status error for events in wait list.\n");
+    break;
+  case CL_COMPILE_PROGRAM_FAILURE:
+    fprintf(stdout, "Compile program failure.\n");
+    break;
+  case CL_LINKER_NOT_AVAILABLE:
+    fprintf(stdout, "Linker not available.\n");
+    break;
+  case CL_LINK_PROGRAM_FAILURE:
+    fprintf(stdout, "Link program failure.\n");
+    break;
+  case CL_DEVICE_PARTITION_FAILED:
+    fprintf(stdout, "Device partition failed.\n");
+    break;
+  case CL_KERNEL_ARG_INFO_NOT_AVAILABLE:
+    fprintf(stdout, "Kernel arg info not available.\n");
+    break;
+
+  // Compiler errors.
+  case CL_INVALID_VALUE:
+    fprintf(stdout, "Invalid value.\n");
+    break;
+  case CL_INVALID_DEVICE_TYPE:
+    fprintf(stdout, "Invalid device type.\n");
+    break;
+  case CL_INVALID_PLATFORM:
+    fprintf(stdout, "Invalid platform.\n");
+    break;
+  case CL_INVALID_DEVICE:
+    fprintf(stdout, "Invalid device.\n");
+    break;
+  case CL_INVALID_CONTEXT:
+    fprintf(stdout, "Invalid context.\n");
+    break;
+  case CL_INVALID_QUEUE_PROPERTIES:
+    fprintf(stdout, "Invalid queue properties.\n");
+    break;
+  case CL_INVALID_COMMAND_QUEUE:
+    fprintf(stdout, "Invalid command queue.\n");
+    break;
+  case CL_INVALID_HOST_PTR:
+    fprintf(stdout, "Invalid host pointer.\n");
+    break;
+  case CL_INVALID_MEM_OBJECT:
+    fprintf(stdout, "Invalid memory object.\n");
+    break;
+  case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:
+    fprintf(stdout, "Invalid image format descriptor.\n");
+    break;
+  case CL_INVALID_IMAGE_SIZE:
+    fprintf(stdout, "Invalid image size.\n");
+    break;
+  case CL_INVALID_SAMPLER:
+    fprintf(stdout, "Invalid sampler.\n");
+    break;
+  case CL_INVALID_BINARY:
+    fprintf(stdout, "Invalid binary.\n");
+    break;
+  case CL_INVALID_BUILD_OPTIONS:
+    fprintf(stdout, "Invalid build options.\n");
+    break;
+  case CL_INVALID_PROGRAM:
+    fprintf(stdout, "Invalid program.\n");
+    break;
+  case CL_INVALID_PROGRAM_EXECUTABLE:
+    fprintf(stdout, "Invalid program executable.\n");
+    break;
+  case CL_INVALID_KERNEL_NAME:
+    fprintf(stdout, "Invalid kernel name.\n");
+    break;
+  case CL_INVALID_KERNEL_DEFINITION:
+    fprintf(stdout, "Invalid kernel definition.\n");
+    break;
+  case CL_INVALID_KERNEL:
+    fprintf(stdout, "Invalid kernel.\n");
+    break;
+  case CL_INVALID_ARG_INDEX:
+    fprintf(stdout, "Invalid arg index.\n");
+    break;
+  case CL_INVALID_ARG_VALUE:
+    fprintf(stdout, "Invalid arg value.\n");
+    break;
+  case CL_INVALID_ARG_SIZE:
+    fprintf(stdout, "Invalid arg size.\n");
+    break;
+  case CL_INVALID_KERNEL_ARGS:
+    fprintf(stdout, "Invalid kernel args.\n");
+    break;
+  case CL_INVALID_WORK_DIMENSION:
+    fprintf(stdout, "Invalid work dimension.\n");
+    break;
+  case CL_INVALID_WORK_GROUP_SIZE:
+    fprintf(stdout, "Invalid work group size.\n");
+    break;
+  case CL_INVALID_WORK_ITEM_SIZE:
+    fprintf(stdout, "Invalid work item size.\n");
+    break;
+  case CL_INVALID_GLOBAL_OFFSET:
+    fprintf(stdout, "Invalid global offset.\n");
+    break;
+  case CL_INVALID_EVENT_WAIT_LIST:
+    fprintf(stdout, "Invalid event wait list.\n");
+    break;
+  case CL_INVALID_EVENT:
+    fprintf(stdout, "Invalid event.\n");
+    break;
+  case CL_INVALID_OPERATION:
+    fprintf(stdout, "Invalid operation.\n");
+    break;
+  case CL_INVALID_GL_OBJECT:
+    fprintf(stdout, "Invalid GL object.\n");
+    break;
+  case CL_INVALID_BUFFER_SIZE:
+    fprintf(stdout, "Invalid buffer size.\n");
+    break;
+  case CL_INVALID_MIP_LEVEL:
+    fprintf(stdout, "Invalid mip level.\n");
+    break;
+  case CL_INVALID_GLOBAL_WORK_SIZE:
+    fprintf(stdout, "Invalid global work size.\n");
+    break;
+  case CL_INVALID_PROPERTY:
+    fprintf(stdout, "Invalid property.\n");
+    break;
+  case CL_INVALID_IMAGE_DESCRIPTOR:
+    fprintf(stdout, "Invalid image descriptor.\n");
+    break;
+  case CL_INVALID_COMPILER_OPTIONS:
+    fprintf(stdout, "Invalid compiler options.\n");
+    break;
+  case CL_INVALID_LINKER_OPTIONS:
+    fprintf(stdout, "Invalid linker options.\n");
+    break;
+  case CL_INVALID_DEVICE_PARTITION_COUNT:
+    fprintf(stdout, "Invalid device partition count.\n");
+    break;
+  case CL_INVALID_PIPE_SIZE:
+    fprintf(stdout, "Invalid pipe size.\n");
+    break;
+  case CL_INVALID_DEVICE_QUEUE:
+    fprintf(stdout, "Invalid device queue.\n");
+    break;
+
+  // NVIDIA specific error.
+  case -9999:
+    fprintf(stdout, "NVIDIA invalid read or write buffer.\n");
+    break;
+
+  default:
+    fprintf(stdout, "Unknown error code!\n");
+    break;
+  }
+}
+
+/******************************************************************************/
+/*                                   CUDA                                     */
+/******************************************************************************/
+
+struct CUDAContextT {
+  CUcontext Cuda;
+};
+
+struct CUDAKernelT {
   CUfunction Cuda;
   CUmodule CudaModule;
   const char *PTXString;
 };
 
-struct PollyGPUDevicePtrT {
+struct CUDADevicePtrT {
   CUdeviceptr Cuda;
 };
 
 /* Dynamic library handles for the CUDA and CUDA runtime library. */
 static void *HandleCuda;
 static void *HandleCudaRT;
 
 /* Type-defines of function pointer to CUDA driver APIs. */

 typedef CUresult CUDAAPI CuDeviceGetNameFcnTy(char *, int, CUdevice);
 static CuDeviceGetNameFcnTy *CuDeviceGetNameFcnPtr;
 
 typedef CUresult CUDAAPI CuLinkAddDataFcnTy(CUlinkState state,
                                             CUjitInputType type, void *data,
                                             size_t size, const char *name,
                                             unsigned int numOptions,
                                             CUjit_option *options,
                                             void **optionValues);
 static CuLinkAddDataFcnTy *CuLinkAddDataFcnPtr;
 
 typedef CUresult CUDAAPI CuLinkCreateFcnTy(unsigned int numOptions,
                                            CUjit_option *options,
                                            void **optionValues,
                                            CUlinkState *stateOut);
 static CuLinkCreateFcnTy *CuLinkCreateFcnPtr;
 
 typedef CUresult CUDAAPI CuLinkCompleteFcnTy(CUlinkState state, void **cubinOut,
                                              size_t *sizeOut);
 static CuLinkCompleteFcnTy *CuLinkCompleteFcnPtr;
 
 typedef CUresult CUDAAPI CuLinkDestroyFcnTy(CUlinkState state);
 static CuLinkDestroyFcnTy *CuLinkDestroyFcnPtr;
 
 /* 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) {
+static void *getAPIHandleCUDA(void *Handle, const char *FuncName) {
   char *Err;
   void *FuncPtr;
   dlerror();
   FuncPtr = dlsym(Handle, FuncName);
   if ((Err = dlerror()) != 0) {
     fprintf(stdout, "Load CUDA driver API failed: %s. \n", Err);
     return 0;
   }
   return FuncPtr;
 }
 
-static int initialDeviceAPILibraries() {
+static int initialDeviceAPILibrariesCUDA() {
   HandleCuda = dlopen("libcuda.so", RTLD_LAZY);
   if (!HandleCuda) {
     printf("Cannot open library: %s. \n", dlerror());
     return 0;
   }
 
   HandleCudaRT = dlopen("libcudart.so", RTLD_LAZY);
   if (!HandleCudaRT) {
     printf("Cannot open library: %s. \n", dlerror());
     return 0;
   }
 
   return 1;
 }
 
-static int initialDeviceAPIs() {
-  if (initialDeviceAPILibraries() == 0)
+static int initialDeviceAPIsCUDA() {
+  if (initialDeviceAPILibrariesCUDA() == 0)
     return 0;
 
   /* Get function pointer to CUDA Driver APIs.
    *
    * Note that compilers conforming to the ISO C standard are required to
    * generate a warning if a conversion from a void * pointer to a function
    * pointer is attempted as in the following statements. The warning
    * of this kind of cast may not be emitted by clang and new versions of gcc
    * as it is valid on POSIX 2008.
    */
   CuLaunchKernelFcnPtr =
-      (CuLaunchKernelFcnTy *)getAPIHandle(HandleCuda, "cuLaunchKernel");
+      (CuLaunchKernelFcnTy *)getAPIHandleCUDA(HandleCuda, "cuLaunchKernel");
 
   CuMemAllocFcnPtr =
-      (CuMemAllocFcnTy *)getAPIHandle(HandleCuda, "cuMemAlloc_v2");
+      (CuMemAllocFcnTy *)getAPIHandleCUDA(HandleCuda, "cuMemAlloc_v2");
 
-  CuMemFreeFcnPtr = (CuMemFreeFcnTy *)getAPIHandle(HandleCuda, "cuMemFree_v2");
+  CuMemFreeFcnPtr =
+      (CuMemFreeFcnTy *)getAPIHandleCUDA(HandleCuda, "cuMemFree_v2");
 
   CuMemcpyDtoHFcnPtr =
-      (CuMemcpyDtoHFcnTy *)getAPIHandle(HandleCuda, "cuMemcpyDtoH_v2");
+      (CuMemcpyDtoHFcnTy *)getAPIHandleCUDA(HandleCuda, "cuMemcpyDtoH_v2");
 
   CuMemcpyHtoDFcnPtr =
-      (CuMemcpyHtoDFcnTy *)getAPIHandle(HandleCuda, "cuMemcpyHtoD_v2");
+      (CuMemcpyHtoDFcnTy *)getAPIHandleCUDA(HandleCuda, "cuMemcpyHtoD_v2");
 
   CuModuleUnloadFcnPtr =
-      (CuModuleUnloadFcnTy *)getAPIHandle(HandleCuda, "cuModuleUnload");
+      (CuModuleUnloadFcnTy *)getAPIHandleCUDA(HandleCuda, "cuModuleUnload");
 
   CuCtxDestroyFcnPtr =
-      (CuCtxDestroyFcnTy *)getAPIHandle(HandleCuda, "cuCtxDestroy");
+      (CuCtxDestroyFcnTy *)getAPIHandleCUDA(HandleCuda, "cuCtxDestroy");
 
-  CuInitFcnPtr = (CuInitFcnTy *)getAPIHandle(HandleCuda, "cuInit");
+  CuInitFcnPtr = (CuInitFcnTy *)getAPIHandleCUDA(HandleCuda, "cuInit");
 
   CuDeviceGetCountFcnPtr =
-      (CuDeviceGetCountFcnTy *)getAPIHandle(HandleCuda, "cuDeviceGetCount");
+      (CuDeviceGetCountFcnTy *)getAPIHandleCUDA(HandleCuda, "cuDeviceGetCount");
 
   CuDeviceGetFcnPtr =
-      (CuDeviceGetFcnTy *)getAPIHandle(HandleCuda, "cuDeviceGet");
+      (CuDeviceGetFcnTy *)getAPIHandleCUDA(HandleCuda, "cuDeviceGet");
 
   CuCtxCreateFcnPtr =
-      (CuCtxCreateFcnTy *)getAPIHandle(HandleCuda, "cuCtxCreate_v2");
+      (CuCtxCreateFcnTy *)getAPIHandleCUDA(HandleCuda, "cuCtxCreate_v2");
 
-  CuModuleLoadDataExFcnPtr =
-      (CuModuleLoadDataExFcnTy *)getAPIHandle(HandleCuda, "cuModuleLoadDataEx");
+  CuModuleLoadDataExFcnPtr = (CuModuleLoadDataExFcnTy *)getAPIHandleCUDA(
+      HandleCuda, "cuModuleLoadDataEx");
 
   CuModuleLoadDataFcnPtr =
-      (CuModuleLoadDataFcnTy *)getAPIHandle(HandleCuda, "cuModuleLoadData");
+      (CuModuleLoadDataFcnTy *)getAPIHandleCUDA(HandleCuda, "cuModuleLoadData");
 
-  CuModuleGetFunctionFcnPtr = (CuModuleGetFunctionFcnTy *)getAPIHandle(
+  CuModuleGetFunctionFcnPtr = (CuModuleGetFunctionFcnTy *)getAPIHandleCUDA(
       HandleCuda, "cuModuleGetFunction");
 
   CuDeviceComputeCapabilityFcnPtr =
-      (CuDeviceComputeCapabilityFcnTy *)getAPIHandle(
+      (CuDeviceComputeCapabilityFcnTy *)getAPIHandleCUDA(
           HandleCuda, "cuDeviceComputeCapability");
 
   CuDeviceGetNameFcnPtr =
-      (CuDeviceGetNameFcnTy *)getAPIHandle(HandleCuda, "cuDeviceGetName");
+      (CuDeviceGetNameFcnTy *)getAPIHandleCUDA(HandleCuda, "cuDeviceGetName");
 
   CuLinkAddDataFcnPtr =
-      (CuLinkAddDataFcnTy *)getAPIHandle(HandleCuda, "cuLinkAddData");
+      (CuLinkAddDataFcnTy *)getAPIHandleCUDA(HandleCuda, "cuLinkAddData");
 
   CuLinkCreateFcnPtr =
-      (CuLinkCreateFcnTy *)getAPIHandle(HandleCuda, "cuLinkCreate");
+      (CuLinkCreateFcnTy *)getAPIHandleCUDA(HandleCuda, "cuLinkCreate");
 
   CuLinkCompleteFcnPtr =
-      (CuLinkCompleteFcnTy *)getAPIHandle(HandleCuda, "cuLinkComplete");
+      (CuLinkCompleteFcnTy *)getAPIHandleCUDA(HandleCuda, "cuLinkComplete");
 
   CuLinkDestroyFcnPtr =
-      (CuLinkDestroyFcnTy *)getAPIHandle(HandleCuda, "cuLinkDestroy");
+      (CuLinkDestroyFcnTy *)getAPIHandleCUDA(HandleCuda, "cuLinkDestroy");
 
   /* Get function pointer to CUDA Runtime APIs. */
-  CudaThreadSynchronizeFcnPtr = (CudaThreadSynchronizeFcnTy *)getAPIHandle(
+  CudaThreadSynchronizeFcnPtr = (CudaThreadSynchronizeFcnTy *)getAPIHandleCUDA(
       HandleCudaRT, "cudaThreadSynchronize");
 
   return 1;
 }
 
-PollyGPUContext *polly_initContext() {
-  DebugMode = getenv("POLLY_DEBUG") != 0;
-
+PollyGPUContext *initContextCUDA() {
   dump_function();
+
   PollyGPUContext *Context;
   CUdevice Device;
 
   int Major = 0, Minor = 0, DeviceID = 0;
   char DeviceName[256];
   int DeviceCount = 0;
 
   static __thread PollyGPUContext *CurrentContext = NULL;
 
   if (CurrentContext)
     return CurrentContext;
 
   /* Get API handles. */
-  if (initialDeviceAPIs() == 0) {
+  if (initialDeviceAPIsCUDA() == 0) {
     fprintf(stdout, "Getting the \"handle\" for the CUDA driver API failed.\n");
     exit(-1);
   }
 
   if (CuInitFcnPtr(0) != CUDA_SUCCESS) {
     fprintf(stdout, "Initializing the CUDA driver API failed.\n");
     exit(-1);
   }
PollyGPUContext *initContextCUDA() {
   debug_print("> Running on GPU device %d : %s.\n", DeviceID, DeviceName);
 
   /* Create context on the device. */
   Context = (PollyGPUContext *)malloc(sizeof(PollyGPUContext));
   if (Context == 0) {
     fprintf(stdout, "Allocate memory for Polly GPU context failed.\n");
     exit(-1);
   }
-  CuCtxCreateFcnPtr(&(Context->Cuda), 0, Device);
-
-  CacheMode = getenv("POLLY_NOCACHE") == 0;
+  Context->Context = malloc(sizeof(CUDAContext));
+  if (Context->Context == 0) {
+    fprintf(stdout, "Allocate memory for Polly CUDA context failed.\n");
+    exit(-1);
+  }
+  CuCtxCreateFcnPtr(&(((CUDAContext *)Context->Context)->Cuda), 0, Device);
 
   if (CacheMode)
     CurrentContext = Context;
 
   return Context;
 }
 
-static void freeKernel(PollyGPUFunction *Kernel) {
-  if (Kernel->CudaModule)
-    CuModuleUnloadFcnPtr(Kernel->CudaModule);
+static void freeKernelCUDA(PollyGPUFunction *Kernel) {
+  dump_function();
+
+  if (CacheMode)
+    return;
+
+  if (((CUDAKernel *)Kernel->Kernel)->CudaModule)
+    CuModuleUnloadFcnPtr(((CUDAKernel *)Kernel->Kernel)->CudaModule);
+
+  if (Kernel->Kernel)
+    free((CUDAKernel *)Kernel->Kernel);
 
   if (Kernel)
     free(Kernel);
 }
 
-#define KERNEL_CACHE_SIZE 10
-
-PollyGPUFunction *polly_getKernel(const char *PTXBuffer,
-                                  const char *KernelName) {
+PollyGPUFunction *getKernelCUDA(const char *PTXBuffer, const char *KernelName) {
   dump_function();
 
   static __thread PollyGPUFunction *KernelCache[KERNEL_CACHE_SIZE];
   static __thread int NextCacheItem = 0;
 
   for (long i = 0; i < KERNEL_CACHE_SIZE; i++) {
     // We exploit here the property that all Polly-ACC kernels are allocated
     // as global constants, hence a pointer comparision is sufficient to
     // determin equality.
-    if (KernelCache[i] && KernelCache[i]->PTXString == PTXBuffer) {
+    if (KernelCache[i] &&
+        ((CUDAKernel *)KernelCache[i]->Kernel)->PTXString == PTXBuffer) {
       debug_print("  -> using cached kernel\n");
       return KernelCache[i];
     }
   }
 
   PollyGPUFunction *Function = malloc(sizeof(PollyGPUFunction));
-
   if (Function == 0) {
     fprintf(stdout, "Allocate memory for Polly GPU function failed.\n");
     exit(-1);
   }
+  Function->Kernel = (CUDAKernel *)malloc(sizeof(CUDAKernel));
+  if (Function->Kernel == 0) {
+    fprintf(stdout, "Allocate memory for Polly CUDA function failed.\n");
+    exit(-1);
+  }
 
   CUresult Res;
   CUlinkState LState;
   CUjit_option Options[6];
   void *OptionVals[6];
   float Walltime = 0;
   unsigned long LogSize = 8192;
   char ErrorLog[8192], InfoLog[8192];
if (Res != CUDA_SUCCESS) {
     fprintf(stdout, "Complete ptx linker step failed.\n");
     fprintf(stdout, "\n%s\n", ErrorLog);
     exit(-1);
   }
 
   debug_print("CUDA Link Completed in %fms. Linker Output:\n%s\n", Walltime,
               InfoLog);
 
-  Res = CuModuleLoadDataFcnPtr(&(Function->CudaModule), CuOut);
+  Res = CuModuleLoadDataFcnPtr(&(((CUDAKernel *)Function->Kernel)->CudaModule),
+                               CuOut);
   if (Res != CUDA_SUCCESS) {
     fprintf(stdout, "Loading ptx assembly text failed.\n");
     exit(-1);
   }
 
-  Res = CuModuleGetFunctionFcnPtr(&(Function->Cuda), Function->CudaModule,
+  Res = CuModuleGetFunctionFcnPtr(&(((CUDAKernel *)Function->Kernel)->Cuda),
+                                  ((CUDAKernel *)Function->Kernel)->CudaModule,
                                   KernelName);
   if (Res != CUDA_SUCCESS) {
     fprintf(stdout, "Loading kernel function failed.\n");
     exit(-1);
   }
 
   CuLinkDestroyFcnPtr(LState);
 
-  Function->PTXString = PTXBuffer;
+  ((CUDAKernel *)Function->Kernel)->PTXString = PTXBuffer;
 
   if (CacheMode) {
     if (KernelCache[NextCacheItem])
-      freeKernel(KernelCache[NextCacheItem]);
+      freeKernelCUDA(KernelCache[NextCacheItem]);
 
     KernelCache[NextCacheItem] = Function;
 
     NextCacheItem = (NextCacheItem + 1) % KERNEL_CACHE_SIZE;
   }
 
   return Function;
 }
 
-void polly_freeKernel(PollyGPUFunction *Kernel) {
-  dump_function();
-
-  if (CacheMode)
-    return;
-
-  freeKernel(Kernel);
-}
-
-void polly_copyFromHostToDevice(void *HostData, PollyGPUDevicePtr *DevData,
-                                long MemSize) {
+void copyFromHostToDeviceCUDA(void *HostData, PollyGPUDevicePtr *DevData,
   dump_function();
 
-  CUdeviceptr CuDevData = DevData->Cuda;
+  CUdeviceptr CuDevData = ((CUDADevicePtr *)DevData->DevicePtr)->Cuda;
   CuMemcpyHtoDFcnPtr(CuDevData, HostData, MemSize);
 }
 
-void polly_copyFromDeviceToHost(PollyGPUDevicePtr *DevData, void *HostData,
-                                long MemSize) {
+void copyFromDeviceToHostCUDA(PollyGPUDevicePtr *DevData, void *HostData,
   dump_function();
 
-  if (CuMemcpyDtoHFcnPtr(HostData, DevData->Cuda, MemSize) != CUDA_SUCCESS) {
+  if (CuMemcpyDtoHFcnPtr(HostData, ((CUDADevicePtr *)DevData->DevicePtr)->Cuda,
+                         MemSize) != CUDA_SUCCESS) {
     fprintf(stdout, "Copying results from device to 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) {
+void launchKernelCUDA(PollyGPUFunction *Kernel, unsigned int GridDimX,
   dump_function();
 
   unsigned GridDimZ = 1;
   unsigned int SharedMemBytes = CU_SHARED_MEM_CONFIG_DEFAULT_BANK_SIZE;
   CUstream Stream = 0;
   void **Extra = 0;
 
   CUresult Res;
-  Res = CuLaunchKernelFcnPtr(Kernel->Cuda, GridDimX, GridDimY, GridDimZ,
-                             BlockDimX, BlockDimY, BlockDimZ, SharedMemBytes,
-                             Stream, Parameters, Extra);
+  Res =
+      CuLaunchKernelFcnPtr(((CUDAKernel *)Kernel->Kernel)->Cuda, GridDimX,
+                           GridDimY, GridDimZ, BlockDimX, BlockDimY, BlockDimZ,
+                           SharedMemBytes, Stream, Parameters, Extra);
   if (Res != CUDA_SUCCESS) {
     fprintf(stdout, "Launching CUDA kernel failed.\n");
     exit(-1);
   }
 }
 
-void polly_freeDeviceMemory(PollyGPUDevicePtr *Allocation) {
+void freeDeviceMemoryCUDA(PollyGPUDevicePtr *Allocation) {
   dump_function();
-  CuMemFreeFcnPtr((CUdeviceptr)Allocation->Cuda);
+  CUDADevicePtr *DevPtr = (CUDADevicePtr *)Allocation->DevicePtr;
+  CuMemFreeFcnPtr((CUdeviceptr)DevPtr->Cuda);
+  free(DevPtr);
   free(Allocation);
 }
 
-PollyGPUDevicePtr *polly_allocateMemoryForDevice(long MemSize) {
+PollyGPUDevicePtr *allocateMemoryForDeviceCUDA(long MemSize) {
   dump_function();
 
   PollyGPUDevicePtr *DevData = malloc(sizeof(PollyGPUDevicePtr));
-
   if (DevData == 0) {
     fprintf(stdout, "Allocate memory for GPU device memory pointer failed.\n");
     exit(-1);
   }
+  DevData->DevicePtr = (CUDADevicePtr *)malloc(sizeof(CUDADevicePtr));
+  if (DevData->DevicePtr == 0) {
+    fprintf(stdout, "Allocate memory for GPU device memory pointer failed.\n");
+    exit(-1);
+  }
 
-  CUresult Res = CuMemAllocFcnPtr(&(DevData->Cuda), MemSize);
+  CUresult Res =
+      CuMemAllocFcnPtr(&(((CUDADevicePtr *)DevData->DevicePtr)->Cuda), MemSize);
 
   if (Res != CUDA_SUCCESS) {
     fprintf(stdout, "Allocate memory for GPU device memory pointer failed.\n");
     exit(-1);
   }
 
   return DevData;
 }
 
+void *getDevicePtrCUDA(PollyGPUDevicePtr *Allocation) {
+  dump_function();
+
+  CUDADevicePtr *DevPtr = (CUDADevicePtr *)Allocation->DevicePtr;
+  return (void *)DevPtr->Cuda;
+}
+
+void freeContextCUDA(PollyGPUContext *Context) {
+  dump_function();
+
+  CUDAContext *Ctx = (CUDAContext *)Context->Context;
+  if (Ctx->Cuda) {
+    CuCtxDestroyFcnPtr(Ctx->Cuda);
+    free(Ctx);
+    free(Context);
+  }
+
+  dlclose(HandleCuda);
+  dlclose(HandleCudaRT);
+}
+
+/******************************************************************************/
+/*                                    API                                     */
+/******************************************************************************/
+
+PollyGPUContext *polly_initContext() {
+  DebugMode = getenv("POLLY_DEBUG") != 0;
+  CacheMode = getenv("POLLY_NOCACHE") == 0;
+
+  dump_function();
+
+  PollyGPUContext *Context;
+
+  switch (Runtime) {
+  case RUNTIME_CUDA:
+    Context = initContextCUDA();
+    break;
+  case RUNTIME_CL:
+    Context = initContextCL();
+    break;
+  default:
+    err_runtime();
+  }
+
+  return Context;
+}
+
+void polly_freeKernel(PollyGPUFunction *Kernel) {
+  dump_function();
+
+  switch (Runtime) {
+  case RUNTIME_CUDA:
+    freeKernelCUDA(Kernel);
+    break;
+  case RUNTIME_CL:
+    freeKernelCL(Kernel);
+    break;
+  default:
+    err_runtime();
+  }
+}
+
+PollyGPUFunction *polly_getKernel(const char *PTXBuffer,
+                                  const char *KernelName) {
+  dump_function();
+
+  PollyGPUFunction *Function;
+
+  switch (Runtime) {
+  case RUNTIME_CUDA:
+    Function = getKernelCUDA(PTXBuffer, KernelName);
+    break;
+  case RUNTIME_CL:
+    Function = getKernelCL(PTXBuffer, KernelName);
+    break;
+  default:
+    err_runtime();
+  }
+
+  return Function;
+}
+
+void polly_copyFromHostToDevice(void *HostData, PollyGPUDevicePtr *DevData,
+                                long MemSize) {
+  dump_function();
+
+  switch (Runtime) {
+  case RUNTIME_CUDA:
+    copyFromHostToDeviceCUDA(HostData, DevData, MemSize);
+    break;
+  case RUNTIME_CL:
+    copyFromHostToDeviceCL(HostData, DevData, MemSize);
+    break;
+  default:
+    err_runtime();
+  }
+}
+
+void polly_copyFromDeviceToHost(PollyGPUDevicePtr *DevData, void *HostData,
+                                long MemSize) {
+  dump_function();
+
+  switch (Runtime) {
+  case RUNTIME_CUDA:
+    copyFromDeviceToHostCUDA(DevData, HostData, MemSize);
+    break;
+  case RUNTIME_CL:
+    copyFromDeviceToHostCL(DevData, HostData, MemSize);
+    break;
+  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) {
+  dump_function();
+
+  switch (Runtime) {
+  case RUNTIME_CUDA:
+    launchKernelCUDA(Kernel, GridDimX, GridDimY, BlockDimX, BlockDimY,
+                     BlockDimZ, Parameters);
+    break;
+  case RUNTIME_CL:
+    launchKernelCL(Kernel, GridDimX, GridDimY, BlockDimX, BlockDimY, BlockDimZ,
+                   Parameters);
+    break;
+  default:
+    err_runtime();
+  }
+}
+
+void polly_freeDeviceMemory(PollyGPUDevicePtr *Allocation) {
+  dump_function();
+
+  switch (Runtime) {
+  case RUNTIME_CUDA:
+    freeDeviceMemoryCUDA(Allocation);
+    break;
+  case RUNTIME_CL:
+    freeDeviceMemoryCL(Allocation);
+    break;
+  default:
+    err_runtime();
+  }
+}
+
+PollyGPUDevicePtr *polly_allocateMemoryForDevice(long MemSize) {
+  dump_function();
+
+  PollyGPUDevicePtr *DevData;
+
+  switch (Runtime) {
+  case RUNTIME_CUDA:
+    DevData = allocateMemoryForDeviceCUDA(MemSize);
+    break;
+  case RUNTIME_CL:
+    DevData = allocateMemoryForDeviceCL(MemSize);
+    break;
+  default:
+    err_runtime();
+  }
+
+  return DevData;
+}
+
 void *polly_getDevicePtr(PollyGPUDevicePtr *Allocation) {
   dump_function();
 
-  return (void *)Allocation->Cuda;
+  void *DevPtr;
+
+  switch (Runtime) {
+  case RUNTIME_CUDA:
+    DevPtr = getDevicePtrCUDA(Allocation);
+    break;
+  case RUNTIME_CL:
+    DevPtr = getDevicePtrCL(Allocation);
+    break;
+  default:
+    err_runtime();
+  }
+
+  return DevPtr;
 }
 
 void polly_freeContext(PollyGPUContext *Context) {
   dump_function();
 
   if (CacheMode)
     return;
 
-  if (Context->Cuda) {
-    CuCtxDestroyFcnPtr(Context->Cuda);
-    free(Context);
-  }
+  switch (Runtime) {
+  case RUNTIME_CUDA:
+    freeContextCUDA(Context);
+    break;
+  case RUNTIME_CL:
+    freeContextCL(Context);
+    break;
+  default:
+    err_runtime();
+  }
 
-  dlclose(HandleCuda);
-  dlclose(HandleCudaRT);
+/* Initialize GPUJIT with CUDA as runtime library. */
+PollyGPUContext *polly_initContextCUDA() {
+  Runtime = RUNTIME_CUDA;
+  return polly_initContext();
+}
+
+/* Initialize GPUJIT with OpenCL as runtime library. */
+PollyGPUContext *polly_initContextCL() {
+  Runtime = RUNTIME_CL;
+  return polly_initContext();
 }