[Polly] [Fortran Support] Add pattern matching to find Fortran array descriptor

[bollu]

Summary

Dragonegg genereates Fortran arrays that have been created using Allocate(..),
as well as globals in a specific way. This allows us to acces their internal
structure and query for properties such as dimension.

This patch adds preliminary support for detection and storage of the Fortran
array descriptor structures. For now, the tests simply check that Fortran
arrays are being detected.

We can build up on this to emit run-time checks for array dimensions. That will
be a separate patch

[Meinersbur]

Can you explain why you need to call ConstantExpr::getAsInstruction(). The patch could be significantly simpler we do not "convert" a ConstantExpr to Instruction. Am I missing something?

DraggonEgg is a de-facto abandoned project (The test cases were generated by LLVM 3.3.1). To use DraggonEgg with Polly today, one has to emit the code with an old release of LLVM with DraggenEgg to a bitcode file, then import it into a newer LLVM with Polly. To compile Fortran, there is flang. Is there still a need to support DraggonEgg?

include/polly/ScopBuilder.h

90

Please use proper capitalization: @see polly::FortranArrayDescriptor

92

Please use the same capitalization as the argument name: @param Inst ...

include/polly/ScopInfo.h

413

"dumb struct" is also called a Plain Old Data (POD). See http://www.cplusplus.com/reference/type_traits/is_pod/.

It is not, btw. std::unique_ptr and AssertingVH have destructors.

I think you should just remove that line, as it is misleading.

432

Do you have other means of explaining the delete-issue than pointing to an unofficial mirror url?

452

This line is only used for asserts. You could put that part with this asserts into a #ifndef NDEBUG block.

468

Did you consider making this method const?

This seems to be the only public accessor of this class. DescriptorOwningInstruction is used to not prematurely free the instruction. To do need to store a reference to Descriptor or would its name be enough?

689

I don't know if Doxygen processes @see if not on its own line.

690

Did you consider moving this fields where the other fields of this class are defined?

lib/Analysis/ScopBuilder.cpp

1

Mistake?

117–130

Please don't just copy the IR from somewhere, only include things that are relevant. For instance !tbaa !5 doesn't mean anything without the MDNode.

Please also explain what the code you are matching is supposed to do.

181–182

I don't undersant why we need to call ConstantExpr::getAsInstruction? Can't we just also patch the ConstantExpr?

201

Can we avoid matching by name? The name should have no influence on the semantics.

252–253

Again, why the need for ConstantExpr::getAsInstruction()?

lib/Analysis/ScopInfo.cpp

1026

Did you consider having FAD as an argument of `MemoryAccesses ctor. It seems to be required only for initialization.

[bollu]

I'm not super familiar with the LLVM API to manipulate IR, so apologies for any API abuse!

Can you explain why you need to call ConstantExpr::getAsInstruction(). The patch could be significantly simpler we do not "convert" a ConstantExpr to Instruction. Am I missing something?

From what I understand, the docs seem to imply that ConstantExpr is a black-box. I wish to pattern match on the GetElementPtr constant expression.

• [lvm::GetElementPtrConstantExpr](http://llvm.org/docs/doxygen/html/classllvm_1_1GetElementPtrConstantExpr.html) mentions that "...GetElementPtrConstantExpr - This class is private to Constants.cpp, and is used behind the s.cenes to implement getelementpr constant exprs".
• Similarly, I wish to destructure a ConstantExpr into a Bitcast, since I wanted to access the array that is being bitcasted.

The only API that seemed to expose the internals of which ContantExpr is being exposed is getAsInstruction(). If this is a misuse of the API, I'd like to know what the correct pattern is. Yes, if there's a clean way to get this information, the ugliness goes away.

[bollu]

include/polly/ScopBuilder.h

92

Thanks for the catch, will fix.

include/polly/ScopInfo.h

413

right. I meant "dumb" in the sense of "contains no dynamic behaviour" / "is used as a data aggregate", not "has no complex behaviour on destruction".

Do you have a better name for that sort of thing? I agree, "dumb" is misleading.

432

Everywhere else the method getAsInstruction is used, they attach to a basic-block that LLVM frees.

I can point to the official LLVM sources instead of the mirror, but the point is the same, is it not? It is used in a test within LLVM to free the chunk of memory.

468

yes, this should be const. Thanks for the catch.

I think I will need the Descriptor later on, to generate things such as run-time checks. I will need access to the LLVM IR Value to generate Load/Store instructions from the Descriptor.

lib/Analysis/ScopBuilder.cpp

201

do you want me to match based on the internal structure of the struct? I could do that, but again, there could be another struct with the same layout.

Would be OK with the solution where I pattern match on both the name and the structure?

lib/Analysis/ScopInfo.cpp

1026

I considered it. However, in that case, construction of the MemoryAccess becomes more complex (and tied in with the Fortran stuff). I wanted to keep the changes to MemoryAccess as isolated as possible.

Currently, computing the array descriptor is somewhat orthogonal from the computation of the memory access. I didn't want to tightly couple them for the reasons mentioned above.

[philip.pfaffe]

If this is a misuse of the API, I'd like to know what the correct pattern is. Yes, if there's a clean way to get this information, the ugliness goes away.

I wouldn't call it a misuse, but it isn't beautiful either. On the other hand, there might not be a beautiful way.

What you can do is check the Opcode of the ConstantExpr CE->getOpcode() == Instruction::GetElementPtr, and manually inspect the Operands, e.g. check if they're all ConstantInts.

[Meinersbur]

In D32639#740731, @bollu wrote:

I'm not super familiar with the LLVM API to manipulate IR, so apologies for any API abuse!

Can you explain why you need to call ConstantExpr::getAsInstruction(). The patch could be significantly simpler we do not "convert" a ConstantExpr to Instruction. Am I missing something?

From what I understand, the docs seem to imply that ConstantExpr is a black-box. I wish to pattern match on the GetElementPtr constant expression.

• [lvm::GetElementPtrConstantExpr](http://llvm.org/docs/doxygen/html/classllvm_1_1GetElementPtrConstantExpr.html) mentions that "...GetElementPtrConstantExpr - This class is private to Constants.cpp, and is used behind the s.cenes to implement getelementpr constant exprs".
• Similarly, I wish to destructure a ConstantExpr into a Bitcast, since I wanted to access the array that is being bitcasted.

The only API that seemed to expose the internals of which ContantExpr is being exposed is getAsInstruction(). If this is a misuse of the API, I'd like to know what the correct pattern is. Yes, if there's a clean way to get this information, the ugliness goes away.

Have you tried llvm::GEPOperator?

[bollu]

• [NFC wrt patch] fix capitalization, comment order and style, comment phrasing.

• Fix mistake comment
• Change order of comments to separate out Fortran parts
• Wrap Fortran pattern matching in ScopBuilder.h in a ..@
• Document pattern matching functions properly, remove extra un-necessary IR cruft.
• remove "dumb struct" comment. Still looking for a way to describe that this struct should not have behaviour. Rather, it is used as an aggregate.

[bollu]

[NFC wrt patch] Don't return StringRef for name, clone & return a std::string

[Meinersbur]

Fails for msvc with:

C:\Users\Meinersbur\src\llvm\tools\polly\lib\Analysis\ScopBuilder.cpp(218): error C2668: 'llvm::make_unique': ambiguous call to overloaded function
C:\Users\Meinersbur\src\llvm\include\llvm/ADT/STLExtras.h(927): note: could be 'std::unique_ptr<polly::FortranArrayDescriptor,std::default_delete<_Ty>> llvm::make_unique<polly::FortranArrayDescriptor,llvm::GlobalValue*&,std::unique_ptr<llvm::Instruction,std::default_delete<llvm::Instruction>>>(llvm::GlobalValue *&,std::unique_ptr<llvm::Instruction,std::default_delete<llvm::Instruction>> &&)'
        with
        [
            _Ty=polly::FortranArrayDescriptor
        ]
C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\INCLUDE\memory(1628): note: or       'std::unique_ptr<polly::FortranArrayDescriptor,std::default_delete<_Ty>> std::make_unique<polly::FortranArrayDescriptor,llvm::GlobalValue*&,std::unique_ptr<llvm::Instruction,std::default_delete<llvm::Instruction>>>(llvm::GlobalValue *&,std::unique_ptr<llvm::Instruction,std::default_delete<llvm::Instruction>> &&)' [found using argument-dependent lookup]
        with
        [
            _Ty=polly::FortranArrayDescriptor
        ]
C:\Users\Meinersbur\src\llvm\tools\polly\lib\Analysis\ScopBuilder.cpp(218): note: while trying to match the argument list '(llvm::GlobalValue *, std::unique_ptr<llvm::Instruction,std::default_delete<_Ty>>)'
        with
        [
            _Ty=llvm::Instruction
        ]
C:\Users\Meinersbur\src\llvm\tools\polly\lib\Analysis\ScopBuilder.cpp(276): error C2668: 'llvm::make_unique': ambiguous call to overloaded function
C:\Users\Meinersbur\src\llvm\include\llvm/ADT/STLExtras.h(927): note: could be 'std::unique_ptr<polly::FortranArrayDescriptor,std::default_delete<_Ty>> llvm::make_unique<polly::FortranArrayDescriptor,llvm::GlobalValue*&,std::unique_ptr<llvm::Instruction,std::default_delete<llvm::Instruction>>>(llvm::GlobalValue *&,std::unique_ptr<llvm::Instruction,std::default_delete<llvm::Instruction>> &&)'
        with
        [
            _Ty=polly::FortranArrayDescriptor
        ]
C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\INCLUDE\memory(1628): note: or       'std::unique_ptr<polly::FortranArrayDescriptor,std::default_delete<_Ty>> std::make_unique<polly::FortranArrayDescriptor,llvm::GlobalValue*&,std::unique_ptr<llvm::Instruction,std::default_delete<llvm::Instruction>>>(llvm::GlobalValue *&,std::unique_ptr<llvm::Instruction,std::default_delete<llvm::Instruction>> &&)' [found using argument-dependent lookup]
        with
        [
            _Ty=polly::FortranArrayDescriptor
        ]
C:\Users\Meinersbur\src\llvm\tools\polly\lib\Analysis\ScopBuilder.cpp(276): note: while trying to match the argument list '(llvm::GlobalValue *, std::unique_ptr<llvm::Instruction,std::default_delete<_Ty>>)'
        with
        [
            _Ty=llvm::Instruction
        ]

lib/Analysis/ScopBuilder.cpp

181–182

With this your test cases still work:

diff --git a/lib/Analysis/ScopBuilder.cpp b/lib/Analysis/ScopBuilder.cpp
index c399e0d2..9d9b381d 100644
--- a/lib/Analysis/ScopBuilder.cpp
+++ b/lib/Analysis/ScopBuilder.cpp
@@ -183,16 +183,7 @@ ScopBuilder::findFortranArrayDescriptorForAllocArrayAccess(MemAccInst Inst) {
     if (!DescriptorGEP)
       continue;

-    // this needs to be deleted, on the basis of:
-    // https://github.com/llvm-mirror/llvm/blob/93e6e5414ded14bcbb233baaaa5567132fee9a0c/unittests/IR/ConstantsTest.cpp#L186
-    std::unique_ptr<Instruction> DescriptorGEPRawInst(
-        DescriptorGEP->getAsInstruction());
-
-    // We want an instruction and not a ConstantExpr since this should
-    // give us more flexibility to inspect the GEP.
-    auto *DescriptorGEPInst =
-        dyn_cast<GetElementPtrInst>(DescriptorGEPRawInst.get());
-
+    auto *DescriptorGEPInst = dyn_cast<GEPOperator>(DescriptorGEP);
     if (!(DescriptorGEPInst && DescriptorGEPInst->hasAllConstantIndices())) {
       continue;
     }
@@ -214,8 +205,7 @@ ScopBuilder::findFortranArrayDescriptorForAllocArrayAccess(MemAccInst Inst) {
     if (!Descriptor) {
       continue;
     }
-    return make_unique<FortranArrayDescriptor>(Descriptor,
-                                               std::move(DescriptorGEPRawInst));
+    return make_unique<FortranArrayDescriptor>(Descriptor, nullptr);
   }

   return nullptr;

278

The ; is unnecessary

[etherzhhb]

lib/Analysis/ScopBuilder.cpp

140

Maybe functions from llvm/IR/PatternMatch.h can help you to simplify this function

685–691

Actually findFortranArrayDescriptorForAllocArrayAccess and findFortranArrayDescriptorForNonAllocArrayAccess do not need to return a std::unique_ptr.

If setFortranArrayDescriptor is written as:

void MemoryAccess::setFortranArrayDescriptor(FortranArrayDescriptor *FAD) {
  this->FAD.reset(FAD);

  // TODO: write checks to make sure it looks _exactly_ like a Fortran array
  // descriptor
};

[bollu]

@Meinersbur: Are there windows buildbots for Polly whose outputs I can look at? I think the ambiguity is between llvm::make_unique and std::make_unique.

[bollu]

• [NFC wrt patch] clarified which make_unique to pick. Pick llvm::make_unique

[Meinersbur]

If you are not going to use dyn_cast<GEPOperator>, could you please state the reason why?

[bollu]

• replace getAsInstruction with dyn_cast<[GEP|BitCast]Operator>. Thanks @Meinersbur!

[bollu]

@Meinersbur:

Fixed, I'm using dyn_cast<GEPOperator>. I wanted to explore the other error, not that I was against using GEPOperator. Thanks for pointing me to it!

Could you point me to how I could have discovered that GEPOperator works in this case on my own? I later read that *Operator acts as a unifying API to manipulate *Instruction and *ConstantExpr, but is there someplace in the documentation that this relationship is expressed? It's described in llvm::Operator, but not in llvm::Instruction or llvm::ConstantExpr from what I've seen.

@etherzhhb:

Maybe functions from llvm/IR/PatternMatch.h can help you to simplify this function

I looked at PatternMatch.h, and I did not find a way to pattern match against GEP instructions. Am I missing it or does the functionality not exist?

Actually findFortranArrayDescriptorForAllocArrayAccess and findFortranArrayDescriptorForNonAllocArrayAccess do not need to return a std::unique_ptr.

I would prefer to return a unique_ptr with make_unique at the site of the allocation itself since it prevents a whole class of bugs. Could you tell me the rationale for not wishing to
take a unique_ptr as a parameter to MemoryAccess::setFortranArrayDescriptor?

[Meinersbur]

I hoped to get rid of FortranaArrayDescriptor entirely. All the it is used for is getName(). ScopArrayInfo already has a BaseName property that you could use for this purpose.

In D32639#743129, @bollu wrote:

Could you point me to how I could have discovered that GEPOperator works in this case on my own? I later read that *Operator acts as a unifying API to manipulate *Instruction and *ConstantExpr, but is there someplace in the documentation that this relationship is expressed? It's described in llvm::Operator, but not in llvm::Instruction or llvm::ConstantExpr from what I've seen.

I look how other code in LLVM does things.

In this case I was looking for uses of GetElementPtrConstantExpr as I did not believe that no other code in LLVM was expecting the structure of constants. I found non-Constants.cpp uses in GEPOperator, which is used in all of LLVM.

I could not find a non-Doxygen documentation for this.

@etherzhhb:

Maybe functions from llvm/IR/PatternMatch.h can help you to simplify this function

I looked at PatternMatch.h, and I did not find a way to pattern match against GEP instructions. Am I missing it or does the functionality not exist?

I cannot find a matcher for GEP either.

Actually findFortranArrayDescriptorForAllocArrayAccess and findFortranArrayDescriptorForNonAllocArrayAccess do not need to return a std::unique_ptr.

I would prefer to return a unique_ptr with make_unique at the site of the allocation itself since it prevents a whole class of bugs. Could you tell me the rationale for not wishing to
take a unique_ptr as a parameter to MemoryAccess::setFortranArrayDescriptor?

If FortranArrayDescriptor is a type that has to be allocated on the heap (because it is identified by memory location; compared by its address) then I find unique_ptr appropriate. If it is used as a container only ("dumb struct"), then you can also pass-by-value, especially if it has just one member. The bug classes "use-after-free" don't "memory leak" don't even exist for pass-by-value.

With one member only, I don't even see a need to introduce that struct. You can just pass the GlobalValue.

include/polly/ScopInfo.h

437–438

This is not necessary anymore because with NDEBUG not set, the asserts will use the variable.

lib/Analysis/ScopBuilder.cpp

1

There is still a change in the diff.

201

Mmmh, not sure.

Ideally, DraggonEgg would add an annotation saying this is an array (access) so we don't have to rely on name and/or structure and on other compilers not generating the same name/structure for some other purpose. Would that be possible?

If not, let's keep the detection by name.

226–228

The order is confusing to me. Could you consider

///  2. [%slot = getelementptr inbounds i8, i8* %mem, i64 <index>]
///      %slot is optional because if you are writing to the 0th index, we don't need a GEP.

230–231

The implementation does not check for alignment.

253–258

Shorten to

auto *BitcastOperator = dyn_cast<BitCastOperator>(MemLoad->getPointerOperand());
if (!BitcastOperator)
  return nullptr;

?

[bollu]

I hoped to get rid of FortranArrayDescriptor entirely. All the it is used for is getName(). ScopArrayInfo already has a BaseName property that you could use for this purpose.

Ah yes, that's fair. I was considering keeping FortranArrayDescriptor around in case I wish to describe more properties, but for now I suppose I could kill the struct.

I'm writing a patch currently which will use the GlobalValue in the FortranArrayDescriptor. I suppose I could make this a memory of the ScopArrayInfo, and pull it out later on if need be. Does that sound OK?

[Meinersbur]

In D32639#744277, @bollu wrote:

I suppose I could make this a memory of the ScopArrayInfo, and pull it out later on if need be. Does that sound OK?

Yes, sounds great!

[bollu]

• [NFC wrt patch] nuke all references to FortranArayDescriptor struct
• [NFC wrt patch] Make changes suggested by @Meinersbur
• [NFW wrt patch] removed braces for one-line if/else

[bollu]

• Check for alignment of store/load

[grosser]

For some reason your diff contains parts of Philip's changes :(

lib/Analysis/ScopBuilder.cpp

1

Unrelated change.

[bollu]

• [NFC wrt patch] fixing diff of patch to diff against upstream/master

[bollu]

• [NFC wrt patch] reorder comment to make it easier to read
• [NFC wrt patch] add period at comment end

[bollu]

@Meinersbur: I restructured the code as you requested. Could you please take another look?

[Meinersbur]

Accepting, because I don't see more significant code improvements.

However, I have some concerns remaining:

• The detection code is too general. findFortranArrayDescriptorForNonAllocArrayAccess is a load from a global pointer, GEP and an access to it. That is very common code and we will get false positives. The prefix "struct.array" also does not seem to be unique to draggonegg's Fortran-style arrays. At the moment there is no effect of array descriptors, but let me suggest to add a command-line switch that enables detection of these arrays and can be enabled manually when compiling Fortran code (i.e. off by default)
• I still don't know that the Fortran Array Descriptor will be used for, hence I cannot evaluate whether the code is useful at all.

[bollu]

• Add better pattern matching checks for type of fortran array descriptor

[bollu]

lib/Analysis/ScopBuilder.cpp

136

@Meinersbur: I added more rigorous checks to make sure that a Global that we believe to be a fortran array is almost surely one. I would appreciate it if you take a look at it.

I am somewhat against adding a flag, for two reasons:

1. Not being behind a flag allows us to catch false positives for this sort of code, which is valuable

2. Polly already has a lot of options, and I don't want to add to the combinatorial explosion :)

If you think that even these pattern matches are too less, I could actually regex match against the name.

However, I am of the opinion that since this is a structural and a name based match, plus incredibly specific nature of "struct.descriptor_dimension" should make this foolproof.

[Meinersbur]

lib/Analysis/ScopBuilder.cpp

136

There is still nothing unique that only DragonEgg would generate. The name "struct.descriptor_dimension" make false positives less likely, but not I can even craft one using clang:

struct descriptor_dimension {int x,y,z};

clang -emit-llvm -S -o -

%struct.descriptor_dimension = type { i32, i32, i32 }

I wouldn't bet on that no code ever written or will be written in the future does not contain a struct of 3 ints called "descriptor_dimension" or "array1_real".

1. Not being behind a flag allows us to catch false positives for this sort of code, which is valuable

If you want to test for false positves, just use the flag. Being off by default doesn't make it impossible to test your code. I'd rather not use our users as guinea pigs. I'd also prefer not having to tell someone in a bug report that they were using the wrong variable/type name.

Once we compile something different depending on whether a Fortran array is detected, we don't even know whether a miscompilation is caused by a false positve since the only difference you can observe is the miscompilation (no error message or the like). Since there isn't even a flag to switch it off to see whether it works then, I cannot even check whether a false positive was the cause of a miscompilation.

Sorry, I have a rather strong opinion about this.

2. Polly already has a lot of options, and I don't want to add to the combinatorial explosion.

Tobias seems to have no hesitation to add more flags (we already have so many). Most are hidden, i.e. not meant for end users.

Telling the compiler what kind of input we are handling is a perfect reason for a new command line switch.

I can offer the alternative to modify your LLVM+DragonEgg such that it generates a unique flag/attribute/metadata that identifies a Fortran array. DragonEgg is not developed anymore, and you are probably the only one using it with Polly. Since DraggonEgg is in the LLVM repository, you can even upstream it.

[bollu]

I'd rather not use our users as guinea pigs. I'd also prefer not having to tell someone in a bug report that they were using the wrong variable/type name.

I understand. I didn't consider the ramifications of this generating false positives, I'm sorry.

Very well, I'll put this behind a flag for now. If I choose to modify dragonegg such that we will generate no false positives, then we can perhaps remove the flag. I hope that sounds okay?

[Meinersbur]

In D32639#749650, @bollu wrote:

Very well, I'll put this behind a flag for now. If I choose to modify dragonegg such that we will generate no false positives, then we can perhaps remove the flag. I hope that sounds okay?

Yes, this is what I had in mind. Thank you.

[bollu]

• hide fortran array detection behind option that is disabled by default

[bollu]

@Meinersbur: Done, I added an option to hide this behind.

[bollu]

@Meinersbur I would like it if you re-check this in its current state, after which I'll be comfortable upstreaming this. Thanks!

[Meinersbur]

Yes, looks great.

lib/Analysis/ScopBuilder.cpp

40

Capitalization of first latter.

687

If there are only two functions to call, the answer is: no.

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/OpenCL, as GPU code generation test cases can be run
-  # without a CUDA/OpenCL library.
+  # Do not require CUDA, as GPU code generation test cases can be run without
+  # a cuda 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)
-  add_definitions(-DHAS_LIBCUDART)
   INCLUDE_DIRECTORIES( ${CUDALIB_INCLUDE_DIR} )
 endif(CUDALIB_FOUND)
-if (OpenCL_FOUND)
-  add_definitions(-DHAS_LIBOPENCL)
-  INCLUDE_DIRECTORIES( ${OpenCL_INCLUDE_DIR} )
-endif(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/CodeGen/PPCGCodeGeneration.h

-//===--- polly/PPCGCodeGeneration.h - Polly Accelerator Code Generation. --===//
-//
-//                     The LLVM Compiler Infrastructure
-//
-// This file is distributed under the University of Illinois Open Source
-// License. See LICENSE.TXT for details.
-//
-//===----------------------------------------------------------------------===//
-//
-// Take a scop created by ScopInfo and map it to GPU code using the ppcg
-// GPU mapping strategy.
-//
-//===----------------------------------------------------------------------===//
-
-#ifndef POLLY_PPCGCODEGENERATION_H
-#define POLLY_PPCGCODEGENERATION_H
-
-/// The GPU Architecture to target.
-enum GPUArch { NVPTX64 };
-
-/// The GPU Runtime implementation to use.
-enum GPURuntime { CUDA, OpenCL };
-
-#endif // POLLY_PPCGCODEGENERATION_H

include/polly/LinkAllPasses.h

 // This header file pulls in all transformation and analysis passes for tools
 // like opt and bugpoint that need this functionality.
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef POLLY_LINKALLPASSES_H
 #define POLLY_LINKALLPASSES_H
 
-#include "polly/CodeGen/PPCGCodeGeneration.h"
 #include "polly/Config/config.h"
 #include "polly/PruneUnprofitable.h"
 #include "polly/Simplify.h"
 #include "polly/Support/DumpModulePass.h"
 #include "llvm/ADT/StringRef.h"
 #include <cstdlib>
 
 namespace llvm {

 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(GPUArch Arch = GPUArch::NVPTX64,
-                                         GPURuntime Runtime = GPURuntime::CUDA);
+llvm::Pass *createPPCGCodeGenerationPass();
 #endif
 llvm::Pass *createIslScheduleOptimizerPass();
 llvm::Pass *createFlattenSchedulePass();
 llvm::Pass *createDeLICMPass();
 
 extern char &CodePreparationID;
 } // namespace polly
 

include/polly/ScopBuilder.h

class ScopBuilder {
   SmallVector<Instruction *, 16> GlobalReads;
 
   /// Set of all accessed array base pointers.
   SmallSetVector<Value *, 16> ArrayBasePointers;
 
   // The Scop
   std::unique_ptr<Scop> scop;
 
+  // Methods for pattern matching against Fortran code generated by dragonegg.
+  // @{
+
+  /// Try to pattern match and find the array descriptor structure in case of a
+  /// fortran array accesss. succeeds on load/store into a fortran array that
+  /// has been allocated.
+  ///
+  /// @see polly::FortranArrayDescriptor
+  ///
+  /// @param Inst The load/store instruction that access the memory.
+  ///
+  /// @note assumes -polly-canonicalize has been run.
+  GlobalValue *findFortranArrayDescriptorForAllocArrayAccess(MemAccInst Inst);
+
+  /// Try to pattern match and find the array descriptor structure in case of a
+  /// fortran array accesss. succeeds on load/store into a fortran array that
+  /// has been allocated.
+  ///
+  /// @see polly::FortranArrayDescriptor
+  ///
+  /// @param Inst The load/store instruction that access the memory.
+  ///
+  /// @note assumes -polly-canonicalize has been run.
+  GlobalValue *
+  findFortranArrayDescriptorForNonAllocArrayAccess(MemAccInst Inst);
+  // @}
+
   // Build the SCoP for Region @p R.
   void buildScop(Region &R, AssumptionCache &AC);
 
   /// Try to build a multi-dimensional fixed sized MemoryAccess from the
   /// Load/Store instruction.
   ///
   /// @param Inst       The Load/Store instruction that access the memory
   /// @param Stmt       The parent statement of the instruction
   ///
   /// @returns True if the access could be built, False otherwise.
   bool buildAccessMultiDimFixed(MemAccInst Inst, ScopStmt *Stmt);
 
   /// Try to build a multi-dimensional parameteric sized MemoryAccess.
   ///        from the Load/Store instruction.
   ///
   /// @param Inst       The Load/Store instruction that access the memory
   /// @param Stmt       The parent statement of the instruction

include/polly/ScopInfo.h

private:
   /// The data layout of the module.
   const DataLayout &DL;
 
   /// The scop this SAI object belongs to.
   Scop &S;
 };
 
 /// Represent memory accesses in statements.
 class MemoryAccess {
   friend class Scop;
   friend class ScopStmt;
 
 public:
   /// The access type of a memory access
   ///
   /// There are three kind of access types:
   ///
   /// * A read access
   ///
   /// A certain set of memory locations are read and may be used for internal
   /// calculations.
   ///
   /// * A must-write access
   ///
   /// A certain set of memory locations is definitely written. The old value is
   /// replaced by a newly calculated value. The old value is not read or used at
   /// all.
   ///
   /// * A may-write access
   ///
   /// A certain set of memory locations may be written. The memory location may
   /// contain a new value if there is actually a write or the old value may
   /// remain, if no write happens.
   enum AccessType {
     READ = 0x1,
     MUST_WRITE = 0x2,
     MAY_WRITE = 0x3,
   };
 
   /// Reduction access type
   ///
   /// Commutative and associative binary operations suitable for reductions
   enum ReductionType {
     RT_NONE, ///< Indicate no reduction at all
     RT_ADD,  ///< Addition
     RT_MUL,  ///< Multiplication
     RT_BOR,  ///< Bitwise Or
     RT_BXOR, ///< Bitwise XOr
     RT_BAND, ///< Bitwise And
   };
 
 private:
   MemoryAccess(const MemoryAccess &) = delete;
   const MemoryAccess &operator=(const MemoryAccess &) = delete;
 
   /// A unique identifier for this memory access.
   ///
   /// The identifier is unique between all memory accesses belonging to the same
   /// scop statement.
   isl_id *Id;
 
   /// What is modeled by this MemoryAccess.
   /// @see MemoryKind
   MemoryKind Kind;
 
   /// Whether it a reading or writing access, and if writing, whether it
   /// is conditional (MAY_WRITE).
   enum AccessType AccType;
 
   /// Reduction type for reduction like accesses, RT_NONE otherwise
   ///
   /// An access is reduction like if it is part of a load-store chain in which
private:
   ///      for i
   ///    S:     A[i] = *((double*)&A[4 * i]);
   ///
   ///    => { S[i] -> A[i]; S[i] -> A[o] : 4i <= o <= 4i + 3 }
   isl_map *AccessRelation;
 
   /// Updated access relation read from JSCOP file.
   isl_map *NewAccessRelation;
+
+  /// Fortran arrays that are created using "Allocate" are stored in terms
+  /// of a descriptor struct. This maintains a raw pointer to the memory,
+  /// along with auxiliary fields with information such as dimensions.
+  /// We hold a reference to the descriptor corresponding to a MemoryAccess
+  /// into a Fortran array. FAD for "Fortran Array Descriptor"
+  AssertingVH<GlobalValue> FAD;
   // @}
 
   __isl_give isl_basic_map *createBasicAccessMap(ScopStmt *Statement);
 
   void assumeNoOutOfBound();
 
   /// Compute bounds on an over approximated  access relation.
   ///
private:
   /// constant size.
   ///
   /// Example:
   ///   float (*A)[4];
   ///   A[1][6] -> A[2][2]
   void wrapConstantDimensions();
 
 public:
   /// Create a new MemoryAccess.
   ///
   /// @param Stmt       The parent statement.
   /// @param AccessInst The instruction doing the access.
   /// @param BaseAddr   The accessed array's address.
   /// @param ElemType   The type of the accessed array elements.
   /// @param AccType    Whether read or write access.
   /// @param IsAffine   Whether the subscripts are affine expressions.
   /// @param Kind       The kind of memory accessed.
   /// @param Subscripts Subscipt expressions
public:
   bool isAnyPHIKind() const { return isOriginalAnyPHIKind(); }
 
   /// Get the statement that contains this memory access.
   ScopStmt *getStatement() const { return Statement; }
 
   /// Get the reduction type of this access
   ReductionType getReductionType() const { return RedType; }
 
+  /// Set the array descriptor corresponding to the Array on which the
+  /// memory access is performed.
+  void setFortranArrayDescriptor(GlobalValue *FAD);
+
   /// Update the original access relation.
   ///
   /// We need to update the original access relation during scop construction,
   /// when unifying the memory accesses that access the same scop array info
   /// object. After the scop has been constructed, the original access relation
   /// should not be changed any more. Instead setNewAccessRelation should
   /// be called.
   void setAccessRelation(__isl_take isl_map *AccessRelation);

lib/Analysis/ScopBuilder.cpp

 //===- ScopBuilder.cpp ---------------------------------------------------===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 //

 STATISTIC(InfeasibleScops,
           "Number of SCoPs with statically infeasible context.");
 
 static cl::opt<bool> ModelReadOnlyScalars(
     "polly-analyze-read-only-scalars",
     cl::desc("Model read-only scalar values in the scop description"),
     cl::Hidden, cl::ZeroOrMore, cl::init(true), cl::cat(PollyCategory));
 
 void ScopBuilder::buildPHIAccesses(PHINode *PHI, Region *NonAffineSubRegion,
                                    bool IsExitBlock) {
 
   // PHI nodes that are in the exit block of the region, hence if IsExitBlock is
   // true, are not modeled as ordinary PHI nodes as they are not part of the
   // region. However, we model the operands in the predecessor blocks that are
   // part of the region as regular scalar accesses.
 
   // If we can synthesize a PHI we can skip it, however only if it is in
if (!scop->contains(UseParent) ||
          scop->hasSingleExitEdge())) {
       // At least one escaping use found.
       ensureValueWrite(Inst);
       break;
     }
   }
 }
 
+/// This is matching against code generated by dragonegg after simplifier
+/// passes have been run.
+///
+/// This is trying to match against "@globaldescriptor", the descriptor
+/// of the Fortran array that is being accessed at load/store. This style
+/// of code is generated for arrays that have been allocated using "Allocate"
+/// in the same module
+///
+/// Pattern Match:
+/// 1. %mallocmem = i8* @malloc(i64 40)
+///
+/// 5. store i8* %mallocmem, i8** getelementptr inbounds
+/// (%"struct.array1_real(kind=8)", %"struct.array1_real(kind=8)"*
+/// @globaldescriptor, i64 0, i32 0), align 32
+///
+/// 2. %typedmem = bitcast i8* %mallocmem to <memtype>*
+///
+/// 3 is optional because if you are writing to the 0th index, you don't
+//    need a GEP.
+/// 3. [%slot = getelementptr inbounds i8, i8* %typedmem, i64 <index>]
+///
+/// 4.1 store/load <memtype> <val>, <memtype>* %typedmem, align 8
+/// 4.2 store/load <memtype> <val>, <memtype>* %slot, align 8
+GlobalValue *
+ScopBuilder::findFortranArrayDescriptorForAllocArrayAccess(MemAccInst Inst) {
+  // match: 4.1 & 4.2 store/load
+  if (!isa<LoadInst>(Inst) && !isa<StoreInst>(Inst))
+    return nullptr;
+
+  Value *Address = Inst.getPointerOperand();
+
+  const BitCastInst *Bitcast = nullptr;
+  // [match: 3]
+  if (auto *Slot = dyn_cast<GetElementPtrInst>(Address)) {
+    Value *TypedMem = Slot->getPointerOperand();
+    // match: 2
+    Bitcast = dyn_cast<BitCastInst>(TypedMem);
+  } else {
+    // match: 2
+    Bitcast = dyn_cast<BitCastInst>(Address);
+  }
+
+  if (!Bitcast)
+    return nullptr;
+
+  auto *MallocMem = Bitcast->getOperand(0);
+
+  // match: 1
+  auto *MallocCall = dyn_cast<CallInst>(MallocMem);
+  if (!MallocCall)
+    return nullptr;
+
+  Function *MallocFn = MallocCall->getCalledFunction();
+  if (!(MallocFn && MallocFn->hasName() && MallocFn->getName() == "malloc"))
+    return nullptr;
+
+  // Find all uses the malloc'd memory.
+  // We are looking for a "store" into a struct with the type being the Fortran
+  // descriptor type
+  for (auto user : MallocMem->users()) {
+
+    /// match: 5
+    auto *MallocStore = dyn_cast<StoreInst>(user);
+    if (!MallocStore)
+      continue;
+
+    auto *DescriptorGEP =
+        dyn_cast<GEPOperator>(MallocStore->getPointerOperand());
+    if (!DescriptorGEP)
+      continue;
+
+    // match: 5
+    auto DescriptorType =
+        dyn_cast<StructType>(DescriptorGEP->getSourceElementType());
+    if (!(DescriptorType && DescriptorType->hasName()))
+      continue;
+
+    // name does not match expected name
+    if (!DescriptorType->getName().startswith("struct.array"))
+      continue;
+
+    GlobalValue *Descriptor =
+        dyn_cast<GlobalValue>(DescriptorGEP->getPointerOperand());
+
+    if (!Descriptor)
+      continue;
+
+    return Descriptor;
+  }
+
+  return nullptr;
+}
+
+/// This is matching against code generated by dragonegg after simplifier
+/// passes have been run.
+///
+/// This is trying to match against "@globaldescriptor", the descriptor
+/// of the Fortran array that is being accessed at load/store. This style
+/// of code is generated for arrays that have been declared global, and
+/// are being accessed across modules
+///
+/// Pattern Match:
+///  1. %mem = load double*, double** bitcast (%"struct.array1_real(kind=8)"*
+///    @globaldescriptor to double**), align 32
+///
+///  2 is optional because if you are writing to the 0th index, you don't
+///     need a GEP.
+///  2. [%slot = getelementptr inbounds i8, i8* %mem, i64 <index>]
+///
+///  3.1 store/load <memtype> <val>, <memtype>* %slot, align 8
+///  3.2 store/load <memtype> <val>, <memtype>* %mem, align 8
+GlobalValue *
+ScopBuilder::findFortranArrayDescriptorForNonAllocArrayAccess(MemAccInst Inst) {
+  // match: 3
+  if (!isa<LoadInst>(Inst) && !isa<StoreInst>(Inst)) {
+    return nullptr;
+  }
+  Value *Slot = Inst.getPointerOperand();
+
+  LoadInst *MemLoad = nullptr;
+  // [match: 2]
+  if (auto *SlotGEP = dyn_cast<GetElementPtrInst>(Slot)) {
+    // match: 1
+    MemLoad = dyn_cast<LoadInst>(SlotGEP->getPointerOperand());
+  } else {
+    // match: 1
+    MemLoad = dyn_cast<LoadInst>(Slot);
+  }
+
+  if (!MemLoad)
+    return nullptr;
+
+  auto *BitcastOperator =
+      dyn_cast<BitCastOperator>(MemLoad->getPointerOperand());
+  if (!BitcastOperator)
+    return nullptr;
+
+  GlobalValue *Descriptor =
+      dyn_cast<GlobalValue>(BitcastOperator->getOperand(0));
+  if (!Descriptor)
+    return nullptr;
+
+  return Descriptor;
+}
+
 bool ScopBuilder::buildAccessMultiDimFixed(MemAccInst Inst, ScopStmt *Stmt) {
   Value *Val = Inst.getValueOperand();
   Type *ElementType = Val->getType();
   Value *Address = Inst.getPointerOperand();
   const SCEV *AccessFunction =
       SE.getSCEVAtScope(Address, LI.getLoopFor(Inst->getParent()));
   const SCEVUnknown *BasePointer =
       dyn_cast<SCEVUnknown>(SE.getPointerBase(AccessFunction));
   enum MemoryAccess::AccessType AccType =
       isa<LoadInst>(Inst) ? MemoryAccess::READ : MemoryAccess::MUST_WRITE;
 
   if (auto *BitCast = dyn_cast<BitCastInst>(Address)) {
     auto *Src = BitCast->getOperand(0);
     auto *SrcTy = Src->getType();
     auto *DstTy = BitCast->getType();
     // Do not try to delinearize non-sized (opaque) pointers.
     if ((SrcTy->isPointerTy() && !SrcTy->getPointerElementType()->isSized()) ||
         (DstTy->isPointerTy() && !DstTy->getPointerElementType()->isSized())) {
       return false;
     }
     if (SrcTy->isPointerTy() && DstTy->isPointerTy() &&
         DL.getTypeAllocSize(SrcTy->getPointerElementType()) ==
             DL.getTypeAllocSize(DstTy->getPointerElementType()))
       Address = Src;
   }
 
MemoryAccess *ScopBuilder::addMemoryAccess(
   return Access;
 }
 
 void ScopBuilder::addArrayAccess(
     MemAccInst MemAccInst, MemoryAccess::AccessType AccType, Value *BaseAddress,
     Type *ElementType, bool IsAffine, ArrayRef<const SCEV *> Subscripts,
     ArrayRef<const SCEV *> Sizes, Value *AccessValue) {
   ArrayBasePointers.insert(BaseAddress);
-  addMemoryAccess(MemAccInst->getParent(), MemAccInst, AccType, BaseAddress,
-                  ElementType, IsAffine, AccessValue, Subscripts, Sizes,
-                  MemoryKind::Array);
+  auto *MemAccess = addMemoryAccess(
+      MemAccInst->getParent(), MemAccInst, AccType, BaseAddress, ElementType,
+      IsAffine, AccessValue, Subscripts, Sizes, MemoryKind::Array);
+
+  // TODO: change to loop of function pointers?
+  if (GlobalValue *FAD =
+          findFortranArrayDescriptorForAllocArrayAccess(MemAccInst))
+    MemAccess->setFortranArrayDescriptor(FAD);
+  else if (GlobalValue *FAD =
+               findFortranArrayDescriptorForNonAllocArrayAccess(MemAccInst))
+    MemAccess->setFortranArrayDescriptor(FAD);
 }
 
 void ScopBuilder::ensureValueWrite(Instruction *Inst) {
   ScopStmt *Stmt = scop->getStmtFor(Inst);
 
   // Inst not defined within this SCoP.
   if (!Stmt)
     return;
 

lib/Analysis/ScopInfo.cpp

MemoryAccess::MemoryAccess(ScopStmt *Stmt, Instruction *AccessInst,
                            ArrayRef<const SCEV *> Subscripts,
                            ArrayRef<const SCEV *> Sizes, Value *AccessValue,
                            MemoryKind Kind)
     : Kind(Kind), AccType(AccType), RedType(RT_NONE), Statement(Stmt),
       InvalidDomain(nullptr), BaseAddr(BaseAddress), ElementType(ElementType),
       Sizes(Sizes.begin(), Sizes.end()), AccessInstruction(AccessInst),
       AccessValue(AccessValue), IsAffine(Affine),
       Subscripts(Subscripts.begin(), Subscripts.end()), AccessRelation(nullptr),
-      NewAccessRelation(nullptr) {
+      NewAccessRelation(nullptr), FAD(nullptr) {
   static const std::string TypeStrings[] = {"", "_Read", "_Write", "_MayWrite"};
   const std::string Access = TypeStrings[AccType] + utostr(Stmt->size());
 
   std::string IdName = Stmt->getBaseName() + Access;
   Id = isl_id_alloc(Stmt->getParent()->getIslCtx(), IdName.c_str(), this);
 }
 
 MemoryAccess::MemoryAccess(ScopStmt *Stmt, AccessType AccType,
                            __isl_take isl_map *AccRel)
     : Kind(MemoryKind::Array), AccType(AccType), RedType(RT_NONE),
       Statement(Stmt), InvalidDomain(nullptr), AccessInstruction(nullptr),
-      IsAffine(true), AccessRelation(nullptr), NewAccessRelation(AccRel) {
+      IsAffine(true), AccessRelation(nullptr), NewAccessRelation(AccRel),
+      FAD(nullptr) {
   auto *ArrayInfoId = isl_map_get_tuple_id(NewAccessRelation, isl_dim_out);
   auto *SAI = ScopArrayInfo::getFromId(ArrayInfoId);
   Sizes.push_back(nullptr);
   for (unsigned i = 1; i < SAI->getNumberOfDimensions(); i++)
     Sizes.push_back(SAI->getDimensionSize(i));
   ElementType = SAI->getElementType();
   BaseAddr = SAI->getBasePtr();
   static const std::string TypeStrings[] = {"", "_Read", "_Write", "_MayWrite"};
raw_ostream &polly::operator<<(raw_ostream &OS,
                                MemoryAccess::ReductionType RT) {
   if (RT == MemoryAccess::RT_NONE)
     OS << "NONE";
   else
     OS << MemoryAccess::getReductionOperatorStr(RT);
   return OS;
 }
 
+void MemoryAccess::setFortranArrayDescriptor(GlobalValue *FAD) {
+  this->FAD = FAD;
+
+// TODO: write checks to make sure it looks _exactly_ like a Fortran array
+// descriptor
+#ifdef NDEBUG
+  StructType *ty = dyn_cast<StructType>(Descriptor->getValueType());
+  assert(ty && "expected value of type Fortran array descriptor");
+  assert(ty->hasName() && ty->getName().startswith("struct.array") &&
+         "expected global to follow Fortran array descriptor type naming "
+         "convention");
+  assert(ty->getNumElements() == 4 &&
+         "expected layout to be like Fortran array descriptor type");
+#endif
+}
+
 void MemoryAccess::print(raw_ostream &OS) const {
   switch (AccType) {
   case READ:
     OS.indent(12) << "ReadAccess :=\t";
     break;
   case MUST_WRITE:
     OS.indent(12) << "MustWriteAccess :=\t";
     break;
   case MAY_WRITE:
     OS.indent(12) << "MayWriteAccess :=\t";
     break;
   }
+
   OS << "[Reduction Type: " << getReductionType() << "] ";
+
+  if (FAD) {
+    OS << "[Fortran array descriptor: " << FAD->getName();
+    OS << "] ";
+  };
+
   OS << "[Scalar: " << isScalarKind() << "]\n";
   OS.indent(16) << getOriginalAccessRelationStr() << ";\n";
   if (hasNewAccessRelation())
     OS.indent(11) << "new: " << getNewAccessRelationStr() << ";\n";
 }
 
 void MemoryAccess::dump() const { print(errs()); }
 

lib/CodeGen/PPCGCodeGeneration.cpp

 //===------ PPCGCodeGeneration.cpp - Polly Accelerator Code Generation. ---===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 //
 // Take a scop created by ScopInfo and map it to GPU code using the ppcg
 // GPU mapping strategy.
 //
 //===----------------------------------------------------------------------===//
 
-#include "polly/CodeGen/PPCGCodeGeneration.h"
 #include "polly/CodeGen/IslAst.h"
 #include "polly/CodeGen/IslNodeBuilder.h"
 #include "polly/CodeGen/Utils.h"
 #include "polly/DependenceInfo.h"
 #include "polly/LinkAllPasses.h"
 #include "polly/Options.h"
 #include "polly/ScopDetection.h"
 #include "polly/ScopInfo.h"

 /// 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, GPURuntime Runtime, GPUArch Arch)
+                 gpu_prog *Prog)
       : IslNodeBuilder(Builder, Annotator, DL, LI, SE, DT, S, StartBlock),
-        Prog(Prog), Runtime(Runtime), Arch(Arch) {
+        Prog(Prog) {
     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).
-  GPURuntime Runtime;
-
-  /// The GPU Architecture to target.
-  GPUArch Arch;
-
   /// 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(), false);
     F = Function::Create(Ty, Linkage, Name, M);
   }
 
   Builder.CreateCall(F);
 }
 
 Value *GPUNodeBuilder::createCallInitContext() {
-  const char *Name;
-
-  switch (Runtime) {
-  case GPURuntime::CUDA:
-    Name = "polly_initContextCUDA";
-    break;
-  case GPURuntime::OpenCL:
-    Name = "polly_initContextCL";
-    break;
-  }
-
+  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;
     std::vector<Type *> Args;
     FunctionType *Ty = FunctionType::get(Builder.getInt8PtrTy(), Args, false);
void GPUNodeBuilder::createScopStmt(isl_ast_expr *Expr,
   if (Stmt->isBlockStmt())
     BlockGen.copyStmt(*Stmt, LTS, Indexes);
   else
     RegionGen.copyStmt(*Stmt, LTS, Indexes);
 }
 
 void GPUNodeBuilder::createKernelSync() {
   Module *M = Builder.GetInsertBlock()->getParent()->getParent();
-
-  Function *Sync;
-
-  switch (Arch) {
-  case GPUArch::NVPTX64:
-    Sync = Intrinsic::getDeclaration(M, Intrinsic::nvvm_barrier0);
-    break;
-  }
-
+  auto *Sync = Intrinsic::getDeclaration(M, Intrinsic::nvvm_barrier0);
   Builder.CreateCall(Sync, {});
 }
 
 /// Collect llvm::Values referenced from @p Node
 ///
 /// This function only applies to isl_ast_nodes that are user_nodes referring
 /// to a ScopStmt. All other node types are ignore.
 ///
GPUNodeBuilder::createKernelFunctionDecl(ppcg_kernel *Kernel,
   }
 
   for (auto *V : SubtreeValues)
     Args.push_back(V->getType());
 
   auto *FT = FunctionType::get(Builder.getVoidTy(), Args, false);
   auto *FN = Function::Create(FT, Function::ExternalLinkage, Identifier,
                               GPUModule.get());
-
-  switch (Arch) {
-  case GPUArch::NVPTX64:
-    FN->setCallingConv(CallingConv::PTX_Kernel);
-    break;
-  }
 
   auto Arg = FN->arg_begin();
   for (long i = 0; i < Kernel->n_array; i++) {
     if (!ppcg_kernel_requires_array_argument(Kernel, i))
       continue;
 
     Arg->setName(Kernel->array[i].array->name);
 
for (auto *V : SubtreeValues) {
     ValueMap[V] = &*Arg;
     Arg++;
   }
 
   return FN;
 }
 
 void GPUNodeBuilder::insertKernelIntrinsics(ppcg_kernel *Kernel) {
-  Intrinsic::ID IntrinsicsBID[2];
-  Intrinsic::ID IntrinsicsTID[3];
+  Intrinsic::ID IntrinsicsBID[] = {Intrinsic::nvvm_read_ptx_sreg_ctaid_x,
+                                   Intrinsic::nvvm_read_ptx_sreg_ctaid_y};
 
-  switch (Arch) {
-  case GPUArch::NVPTX64:
-    IntrinsicsBID[0] = Intrinsic::nvvm_read_ptx_sreg_ctaid_x;
-    IntrinsicsBID[1] = Intrinsic::nvvm_read_ptx_sreg_ctaid_y;
-
-    IntrinsicsTID[0] = Intrinsic::nvvm_read_ptx_sreg_tid_x;
-    IntrinsicsTID[1] = Intrinsic::nvvm_read_ptx_sreg_tid_y;
-    IntrinsicsTID[2] = Intrinsic::nvvm_read_ptx_sreg_tid_z;
-    break;
-  }
+  Intrinsic::ID IntrinsicsTID[] = {Intrinsic::nvvm_read_ptx_sreg_tid_x,
+                                   Intrinsic::nvvm_read_ptx_sreg_tid_y,
+                                   Intrinsic::nvvm_read_ptx_sreg_tid_z};
 
   auto addId = [this](__isl_take isl_id *Id, Intrinsic::ID Intr) mutable {
     std::string Name = isl_id_get_name(Id);
     Module *M = Builder.GetInsertBlock()->getParent()->getParent();
     Function *IntrinsicFn = Intrinsic::getDeclaration(M, Intr);
     Value *Val = Builder.CreateCall(IntrinsicFn, {});
     Val = Builder.CreateIntCast(Val, Builder.getInt64Ty(), false, Name);
     IDToValue[Id] = Val;
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) {
+
   std::string Identifier = "kernel_" + std::to_string(Kernel->id);
   GPUModule.reset(new Module(Identifier, Builder.getContext()));
-
-  switch (Arch) {
-  case GPUArch::NVPTX64:
-    if (Runtime == GPURuntime::CUDA)
-      GPUModule->setTargetTriple(Triple::normalize("nvptx64-nvidia-cuda"));
-    else if (Runtime == GPURuntime::OpenCL)
-      GPUModule->setTargetTriple(Triple::normalize("nvptx64-nvidia-nvcl"));
-    GPUModule->setDataLayout(computeNVPTXDataLayout(true /* is64Bit */));
-    break;
-  }
 
   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;
-
-  switch (Arch) {
-  case GPUArch::NVPTX64:
-    switch (Runtime) {
-    case GPURuntime::CUDA:
-      GPUTriple = llvm::Triple(Triple::normalize("nvptx64-nvidia-cuda"));
-      break;
-    case GPURuntime::OpenCL:
-      GPUTriple = llvm::Triple(Triple::normalize("nvptx64-nvidia-nvcl"));
-      break;
-    }
-    break;
-  }
-
+  llvm::Triple GPUTriple(Triple::normalize("nvptx64-nvidia-cuda"));
   std::string ErrMsg;
   auto GPUTarget = TargetRegistry::lookupTarget(GPUTriple.getTriple(), ErrMsg);
 
   if (!GPUTarget) {
     errs() << ErrMsg << "\n";
     return "";
   }
 
   TargetOptions Options;
   Options.UnsafeFPMath = FastMath;
-
-  std::string subtarget;
-
-  switch (Arch) {
-  case GPUArch::NVPTX64:
-    subtarget = CudaVersion;
-    break;
-  }
-
-  std::unique_ptr<TargetMachine> TargetM(GPUTarget->createTargetMachine(
-      GPUTriple.getTriple(), subtarget, "", Options, Optional<Reloc::Model>()));
+  std::unique_ptr<TargetMachine> TargetM(
+      GPUTarget->createTargetMachine(GPUTriple.getTriple(), CudaVersion, "",
+                                     Options, Optional<Reloc::Model>()));
 
   SmallString<0> ASMString;
   raw_svector_ostream ASMStream(ASMString);
   llvm::legacy::PassManager PM;
 
   PM.add(createTargetTransformInfoWrapperPass(TargetM->getTargetIRAnalysis()));
 
   if (TargetM->addPassesToEmitFile(
std::string GPUNodeBuilder::finalizeKernelFunction() {
   return Assembly;
 }
 
 namespace {
 class PPCGCodeGeneration : public ScopPass {
 public:
   static char ID;
 
-  GPURuntime Runtime = GPURuntime::CUDA;
-
-  GPUArch Architecture = GPUArch::NVPTX64;
-
   /// The scop that is currently processed.
   Scop *S;
 
   LoopInfo *LI;
   DominatorTree *DT;
   ScalarEvolution *SE;
   const DataLayout *DL;
   RegionInfo *RI;
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, Runtime, Architecture);
+                               StartBlock, Prog);
 
     // 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 {
     AU.addPreserved<RegionInfoPass>();
     AU.addPreserved<ScopInfoRegionPass>();
   }
 };
 } // namespace
 
 char PPCGCodeGeneration::ID = 1;
 
-Pass *polly::createPPCGCodeGenerationPass(GPUArch Arch, GPURuntime Runtime) {
-  PPCGCodeGeneration *generator = new PPCGCodeGeneration();
-  generator->Runtime = Runtime;
-  generator->Architecture = Arch;
-  return generator;
-}
+Pass *polly::createPPCGCodeGenerationPass() { 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

 // changed, but that the flag '-polly' provided at optimization level '-O3'
 // enables additional polyhedral optimizations.
 //===----------------------------------------------------------------------===//
 
 #include "polly/RegisterPasses.h"
 #include "polly/Canonicalization.h"
 #include "polly/CodeGen/CodeGeneration.h"
 #include "polly/CodeGen/CodegenCleanup.h"
-#include "polly/CodeGen/PPCGCodeGeneration.h"
 #include "polly/DeLICM.h"
 #include "polly/DependenceInfo.h"
 #include "polly/FlattenSchedule.h"
 #include "polly/LinkAllPasses.h"
 #include "polly/Options.h"
 #include "polly/PolyhedralInfo.h"
 #include "polly/ScopDetection.h"
 #include "polly/ScopInfo.h"
Target("polly-target", cl::desc("The hardware to target"),
            cl::values(clEnumValN(TARGET_CPU, "cpu", "generate CPU code")
 #ifdef GPU_CODEGEN
                           ,
                       clEnumValN(TARGET_GPU, "gpu", "generate GPU code")
 #endif
                           ),
            cl::init(TARGET_CPU), cl::ZeroOrMore, cl::cat(PollyCategory));
 
-#ifdef GPU_CODEGEN
-static cl::opt<GPURuntime> GPURuntimeChoice(
-    "polly-gpu-runtime", cl::desc("The GPU Runtime API to target"),
-    cl::values(clEnumValN(GPURuntime::CUDA, "libcudart",
-                          "use the CUDA Runtime API"),
-               clEnumValN(GPURuntime::OpenCL, "libopencl",
-                          "use the OpenCL Runtime API")),
-    cl::init(GPURuntime::CUDA), cl::ZeroOrMore, cl::cat(PollyCategory));
-
-static cl::opt<GPUArch>
-    GPUArchChoice("polly-gpu-arch", cl::desc("The GPU Architecture to target"),
-                  cl::values(clEnumValN(GPUArch::NVPTX64, "nvptx64",
-                                        "target NVIDIA 64-bit architecture")),
-                  cl::init(GPUArch::NVPTX64), cl::ZeroOrMore,
-                  cl::cat(PollyCategory));
-#endif
-
 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"),
         clEnumValN(
if (Target == TARGET_GPU) {
     }
   }
 
   if (ExportJScop)
     PM.add(polly::createJSONExporterPass());
 
   if (Target == TARGET_GPU) {
 #ifdef GPU_CODEGEN
-    PM.add(
-        polly::createPPCGCodeGenerationPass(GPUArchChoice, GPURuntimeChoice));
+    PM.add(polly::createPPCGCodeGenerationPass());
 #endif
   } else {
     switch (CodeGeneration) {
     case CODEGEN_AST:
       PM.add(polly::createIslAstInfoPass());
       break;
     case CODEGEN_FULL:
       PM.add(polly::createCodeGenerationPass());

test/FortranDetection/global-malloc-nonvectored.ll

+; RUN: opt -S  -analyze -polly-process-unprofitable  -polly-remarks-minimal \
+; RUN: -polly-canonicalize -polly-scops -polly-dependences \
+; RUN: -debug-only=polly-dependence -polly-canonicalize -polly-allow-nonaffine \
+; RUN: -polly-ignore-aliasing   -polly-invariant-load-hoisting \
+; RUN: < %s| FileCheck %s
+;
+; MODULE src_soil
+; USE data_parameters, ONLY :   &
+;     wp,        & ! KIND-type parameter for real variables
+;     iintegers    ! KIND-type parameter for standard integer variables
+; IMPLICIT NONE
+; REAL (KIND = wp),     ALLOCATABLE, PRIVATE  :: &
+;   xdzs     (:)
+; CONTAINS
+; SUBROUTINE terra1(n)
+;   INTEGER, intent(in) :: n
+;   INTEGER (KIND=iintegers) ::  &
+;     j
+;   Allocate(xdzs(n));
+;    DO j = 2, n
+;         xdzs(j) = xdzs(j) * xdzs(j) + xdzs(j - 1)
+;   END DO
+; END SUBROUTINE terra1
+; END MODULE src_soil
+
+target datalayout = "e-p:64:64:64-S128-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f16:16:16-f32:32:32-f64:64:64-f128:128:128-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64"
+target triple = "x86_64-unknown-linux-gnu"
+
+module asm "\09.ident\09\22GCC: (GNU) 4.6.4 LLVM: 3.3.1\22"
+
+%"struct.array1_real(kind=8)" = type { i8*, i64, i64, [1 x %struct.descriptor_dimension] }
+%struct.descriptor_dimension = type { i64, i64, i64 }
+
+@__src_soil_MOD_xdzs = unnamed_addr global %"struct.array1_real(kind=8)" zeroinitializer, align 32
+@.cst = private unnamed_addr constant [67 x i8] c"Integer overflow when calculating the amount of memory to allocate\00", align 64
+@.cst1 = private unnamed_addr constant [37 x i8] c"Allocation would exceed memory limit\00", align 64
+@.cst2 = private unnamed_addr constant [93 x i8] c"At line 23 of file /home/siddhart/cosmo-self-installation/cosmo-pompa/cosmo/src/src_soil.f90\00", align 64
+@.cst3 = private unnamed_addr constant [55 x i8] c"Attempting to allocate already allocated variable '%s'\00", align 64
+@.cst4 = private unnamed_addr constant [5 x i8] c"xdzs\00", align 8
+
+; Function Attrs: nounwind uwtable
+define void @__src_soil_MOD_terra1(i32* noalias nocapture %n) unnamed_addr #0 {
+entry:
+  store i64 537, i64* getelementptr inbounds (%"struct.array1_real(kind=8)", %"struct.array1_real(kind=8)"* @__src_soil_MOD_xdzs, i64 0, i32 2), align 16, !tbaa !0
+  store i64 1, i64* getelementptr inbounds (%"struct.array1_real(kind=8)", %"struct.array1_real(kind=8)"* @__src_soil_MOD_xdzs, i64 0, i32 3, i64 0, i32 1), align 8, !tbaa !0
+  %0 = load i32, i32* %n, align 4, !tbaa !3
+  %1 = sext i32 %0 to i64
+  store i64 %1, i64* getelementptr inbounds (%"struct.array1_real(kind=8)", %"struct.array1_real(kind=8)"* @__src_soil_MOD_xdzs, i64 0, i32 3, i64 0, i32 2), align 8, !tbaa !0
+  store i64 1, i64* getelementptr inbounds (%"struct.array1_real(kind=8)", %"struct.array1_real(kind=8)"* @__src_soil_MOD_xdzs, i64 0, i32 3, i64 0, i32 0), align 8, !tbaa !0
+  %2 = icmp slt i32 %0, 0
+  %3 = select i1 %2, i64 0, i64 %1
+  %4 = icmp eq i64 %3, 0
+  br i1 %4, label %"16", label %"8"
+
+"8":                                              ; preds = %entry
+  %5 = sdiv i64 9223372036854775807, %1
+  %6 = icmp slt i64 %5, 1
+  %7 = icmp slt i32 %0, 1
+  %8 = shl nsw i64 %3, 3
+  %.2 = select i1 %7, i64 0, i64 %8
+  br i1 %6, label %"15", label %"16"
+
+"15":                                             ; preds = %"8"
+
+  unreachable
+
+"16":                                             ; preds = %"8", %entry
+  %.24 = phi i64 [ %.2, %"8" ], [ 0, %entry ]
+  %9 = load i8*, i8** getelementptr inbounds (%"struct.array1_real(kind=8)", %"struct.array1_real(kind=8)"* @__src_soil_MOD_xdzs, i64 0, i32 0), align 32, !tbaa !5
+  %10 = icmp eq i8* %9, null
+  br i1 %10, label %"17", label %"20"
+
+"17":                                             ; preds = %"16"
+  %11 = icmp ne i64 %.24, 0
+  %12 = select i1 %11, i64 %.24, i64 1
+  %13 = tail call noalias i8* @malloc(i64 %12) #2 ;<= 1. malloc
+  %14 = icmp eq i8* %13, null
+  br i1 %14, label %"18", label %"19"
+
+"18":                                             ; preds = %"17"
+  unreachable
+
+"19":                                             ; preds = %"17"
+  store i8* %13, i8** getelementptr inbounds (%"struct.array1_real(kind=8)", %"struct.array1_real(kind=8)"* @__src_soil_MOD_xdzs, i64 0, i32 0), align 32, !tbaa !5
+  store i64 -1, i64* getelementptr inbounds (%"struct.array1_real(kind=8)", %"struct.array1_real(kind=8)"* @__src_soil_MOD_xdzs, i64 0, i32 1), align 8, !tbaa !0
+  %15 = icmp sgt i32 %0, 1
+  br i1 %15, label %"21.preheader", label %return
+
+"21.preheader":                                   ; preds = %"19"
+  %16 = bitcast i8* %13 to double* ;<= 2. bitcast to double*
+  %17 = add i32 %0, 1
+  br label %"21"
+
+"20":                                             ; preds = %"16"
+  unreachable
+
+"21":                                             ; preds = %"21", %"21.preheader"
+  %18 = phi double [ undef, %"21.preheader" ], [ %23, %"21" ]
+  %indvars.iv = phi i64 [ 2, %"21.preheader" ], [ %indvars.iv.next, %"21" ]
+  %19 = add nsw i64 %indvars.iv, -1
+  %20 = getelementptr inbounds double, double* %16, i64 %19 ;<= 3. GEP
+  %21 = load double, double* %20, align 8, !tbaa !7
+  %22 = fmul double %21, %21
+  %23 = fadd double %22, %18
+  store double %23, double* %20, align 8, !tbaa !7 ;<= 4. store
+  %indvars.iv.next = add i64 %indvars.iv, 1
+  %lftr.wideiv = trunc i64 %indvars.iv.next to i32
+  %exitcond = icmp eq i32 %lftr.wideiv, %17
+  br i1 %exitcond, label %return, label %"21"
+
+return:                                           ; preds = %"21", %"19"
+  ret void
+}
+
+; Function Attrs: noreturn
+declare void @_gfortran_runtime_error(i8*, ...) #1
+
+; Function Attrs: nounwind
+declare noalias i8* @malloc(i64) #2
+
+; Function Attrs: noreturn
+declare void @_gfortran_os_error(i8*) #1
+
+; Function Attrs: noreturn
+declare void @_gfortran_runtime_error_at(i8*, i8*, ...) #1
+
+attributes #0 = { nounwind uwtable }
+attributes #1 = { noreturn }
+attributes #2 = { nounwind }
+attributes #3 = { noreturn nounwind }
+
+!0 = !{!1, !1, i64 0}
+!1 = !{!"alias set 4: integer(kind=8)", !2}
+!2 = distinct !{!2}
+!3 = !{!4, !4, i64 0}
+!4 = !{!"alias set 11: integer(kind=4)", !2}
+!5 = !{!6, !6, i64 0}
+!6 = !{!"alias set 3: void*", !2}
+!7 = !{!8, !8, i64 0}
+!8 = !{!"alias set 18: real(kind=8)", !2}
+
+; CHECK: ReadAccess :=  [Reduction Type: NONE] [Fortran array descriptor: __src_soil_MOD_xdzs] [Scalar: 0]
+; CHECK: MustWriteAccess := [Reduction Type: NONE] [Fortran array descriptor: __src_soil_MOD_xdzs] [Scalar: 0]

test/FortranDetection/global-nonmalloc-nonvectored.ll

+; RUN: opt -S  -analyze -polly-process-unprofitable  -polly-remarks-minimal \
+; RUN: -polly-canonicalize -polly-scops -polly-dependences \
+; RUN: -debug-only=polly-dependence -polly-canonicalize -polly-allow-nonaffine \
+; RUN: -polly-ignore-aliasing   -polly-invariant-load-hoisting \
+; RUN: < %s| FileCheck %s
+;
+; MODULE src_soil
+; USE data_parameters, ONLY :   &
+;     wp,        & ! KIND-type parameter for real variables
+;     iintegers    ! KIND-type parameter for standard integer variables
+; IMPLICIT NONE
+; REAL (KIND = wp),     ALLOCATABLE, PRIVATE  :: &
+;   xdzs     (:)
+; CONTAINS
+;
+; SUBROUTINE terra1(n)
+;   INTEGER, intent(in) :: n
+;
+;   INTEGER (KIND=iintegers) ::  &
+;     j
+;
+;    DO j = 22, n
+;         xdzs(j) = xdzs(j) * xdzs(j) + xdzs(j - 1)
+;   END DO
+; END SUBROUTINE terra1
+; END MODULE src_soil
+
+target datalayout = "e-p:64:64:64-S128-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f16:16:16-f32:32:32-f64:64:64-f128:128:128-v64:64:64-v128:128:128-a0:0:64-s0:64:64-f80:128:128-n8:16:32:64"
+target triple = "x86_64-unknown-linux-gnu"
+
+module asm "\09.ident\09\22GCC: (GNU) 4.6.4 LLVM: 3.3.1\22"
+
+%"struct.array1_real(kind=8)" = type { i8*, i64, i64, [1 x %struct.descriptor_dimension] }
+%struct.descriptor_dimension = type { i64, i64, i64 }
+
+@__src_soil_MOD_xdzs = unnamed_addr global %"struct.array1_real(kind=8)" zeroinitializer, align 32
+
+; Function Attrs: nounwind uwtable
+define void @__src_soil_MOD_terra1(i32* noalias nocapture %n) unnamed_addr #0 {
+entry:
+  %0 = load i32, i32* %n, align 4, !tbaa !0
+  %1 = icmp sgt i32 %0, 21
+  br i1 %1, label %"3.preheader", label %return
+
+"3.preheader":                                    ; preds = %entry
+  %2 = load i64, i64* getelementptr inbounds (%"struct.array1_real(kind=8)", %"struct.array1_real(kind=8)"* @__src_soil_MOD_xdzs, i64 0, i32 1), align 8, !tbaa !3
+  %3 = load i8*, i8** getelementptr inbounds (%"struct.array1_real(kind=8)", %"struct.array1_real(kind=8)"* @__src_soil_MOD_xdzs, i64 0, i32 0), align 32, !tbaa !5
+  %4 = bitcast i8* %3 to double*
+  %5 = add i32 %0, 1
+  br label %"3"
+
+"3":                                              ; preds = %"3", %"3.preheader"
+  %indvars.iv = phi i64 [ 22, %"3.preheader" ], [ %indvars.iv.next, %"3" ]
+  %6 = add nsw i64 %indvars.iv, %2
+  %7 = getelementptr inbounds double, double* %4, i64 %6
+  %8 = load double, double* %7, align 8, !tbaa !7
+  %9 = fmul double %8, %8
+  %10 = add nsw i64 %indvars.iv, -1
+  %11 = add nsw i64 %10, %2
+  %12 = getelementptr inbounds double, double* %4, i64 %11
+  %13 = load double, double* %12, align 8, !tbaa !7
+  %14 = fadd double %9, %13
+  store double %14, double* %7, align 8, !tbaa !7
+  %indvars.iv.next = add i64 %indvars.iv, 1
+  %lftr.wideiv = trunc i64 %indvars.iv.next to i32
+  %exitcond = icmp eq i32 %lftr.wideiv, %5
+  br i1 %exitcond, label %return, label %"3"
+
+return:                                           ; preds = %"3", %entry
+  ret void
+}
+
+attributes #0 = { nounwind uwtable }
+
+!0 = !{!1, !1, i64 0}
+!1 = !{!"alias set 11: integer(kind=4)", !2}
+!2 = distinct !{!2}
+!3 = !{!4, !4, i64 0}
+!4 = !{!"alias set 4: integer(kind=8)", !2}
+!5 = !{!6, !6, i64 0}
+!6 = !{!"alias set 3: void*", !2}
+!7 = !{!8, !8, i64 0}
+!8 = !{!"alias set 18: real(kind=8)", !2}
+
+; CHECK: ReadAccess :=  [Reduction Type: NONE] [Fortran array descriptor: __src_soil_MOD_xdzs] [Scalar: 0]
+; CHECK: ReadAccess :=  [Reduction Type: NONE] [Fortran array descriptor: __src_soil_MOD_xdzs] [Scalar: 0]
+; CHECK: MustWriteAccess := [Reduction Type: NONE] [Fortran array descriptor: __src_soil_MOD_xdzs] [Scalar: 0]

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

 ;    }
 ;
 
 ; CHECK-NOT: polly_copyFromHostToDevice
 ; CHECK-NOT: polly_copyFromDeviceToHost
 ; CHECK-NOT: polly_freeDeviceMemory
 ; CHECK-NOT: polly_allocateMemoryForDevice
 
-; CHECK:       %13 = call i8* @polly_initContextCUDA()
+; 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 ([750 x i8], [750 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: %20 = call i8* @polly_getKernel(i8* getelementptr inbounds ([750 x i8], [750 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) {

test/GPGPU/size-cast.ll

 ; CODE:   cudaCheckReturn(cudaMemcpy(MemRef_arg2, dev_MemRef_arg2, (arg) * sizeof(double), cudaMemcpyDeviceToHost));
 ; CODE-NEXT: }
 
 ; CODE: # kernel0
 ; CODE-NEXT: for (int c0 = 0; c0 <= (arg - 32 * b0 - 1) / 1048576; c0 += 1)
 ; CODE-NEXT:   if (arg >= 32 * b0 + t0 + 1048576 * c0 + 1)
 ; CODE-NEXT:     Stmt_bb6(0, 32 * b0 + t0 + 1048576 * c0);
 
-; IR:        call i8* @polly_initContextCUDA()
+; IR:        call i8* @polly_initContext()
 ; IR-NEXT:   sext i32 %arg to i64
 ; IR-NEXT:   mul i64
 ; IR-NEXT:   @polly_allocateMemoryForDevice
 
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 target triple = "x86_64-unknown-linux-gnu"
 
 define void @hoge(i32 %arg, i32 %arg1, [1000 x double]* %arg2, double* %arg3) {

tools/CMakeLists.txt

-if (CUDALIB_FOUND OR OpenCL_FOUND)
+if (CUDALIB_FOUND)
   add_subdirectory(GPURuntime)
-endif (CUDALIB_FOUND OR OpenCL_FOUND)
+endif (CUDALIB_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;
 
-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 *BinaryBuffer,
+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

 /******************** 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"
-
-#ifdef HAS_LIBCUDART
 #include <cuda.h>
 #include <cuda_runtime.h>
-#endif /* HAS_LIBCUDART */
-
-#ifdef HAS_LIBOPENCL
-#ifdef __APPLE__
-#include <OpenCL/opencl.h>
-#else
-#include <CL/cl.h>
-#endif
-#endif /* HAS_LIBOPENCL */
-
 #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 KERNEL_CACHE_SIZE 10
-
-static void err_runtime() {
-  fprintf(stderr, "Runtime not correctly initialized.\n");
-  exit(-1);
-}
-
+/* Define Polly's GPGPU data types. */
 struct PollyGPUContextT {
-  void *Context;
-};
-
-struct PollyGPUFunctionT {
-  void *Kernel;
-};
-
-struct PollyGPUDevicePtrT {
-  void *DevicePtr;
-};
-
-/******************************************************************************/
-/*                                  OpenCL                                    */
-/******************************************************************************/
-#ifdef HAS_LIBOPENCL
-
-struct OpenCLContextT {
-  cl_context Context;
-  cl_command_queue CommandQueue;
-};
-
-struct OpenCLKernelT {
-  cl_kernel Kernel;
-  cl_program Program;
-  const char *BinaryString;
-};
-
-struct OpenCLDevicePtrT {
-  cl_mem MemObj;
-};
-
-/* 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 NumEntries,
-                                     cl_platform_id *Platforms,
-                                     cl_uint *NumPlatforms);
-static clGetPlatformIDsFcnTy *clGetPlatformIDsFcnPtr;
-
-typedef cl_int clGetDeviceIDsFcnTy(cl_platform_id Platform,
-                                   cl_device_type DeviceType,
-                                   cl_uint NumEntries, cl_device_id *Devices,
-                                   cl_uint *NumDevices);
-static clGetDeviceIDsFcnTy *clGetDeviceIDsFcnPtr;
-
-typedef cl_int clGetDeviceInfoFcnTy(cl_device_id Device,
-                                    cl_device_info ParamName,
-                                    size_t ParamValueSize, void *ParamValue,
-                                    size_t *ParamValueSizeRet);
-static clGetDeviceInfoFcnTy *clGetDeviceInfoFcnPtr;
-
-typedef cl_int clGetKernelInfoFcnTy(cl_kernel Kernel, cl_kernel_info ParamName,
-                                    size_t ParamValueSize, void *ParamValue,
-                                    size_t *ParamValueSizeRet);
-static clGetKernelInfoFcnTy *clGetKernelInfoFcnPtr;
-
-typedef cl_context clCreateContextFcnTy(
-    const cl_context_properties *Properties, cl_uint NumDevices,
-    const cl_device_id *Devices,
-    void CL_CALLBACK *pfn_notify(const char *Errinfo, const void *PrivateInfo,
-                                 size_t CB, void *UserData),
-    void *UserData, cl_int *ErrcodeRet);
-static clCreateContextFcnTy *clCreateContextFcnPtr;
-
-typedef cl_command_queue
-clCreateCommandQueueFcnTy(cl_context Context, cl_device_id Device,
-                          cl_command_queue_properties Properties,
-                          cl_int *ErrcodeRet);
-static clCreateCommandQueueFcnTy *clCreateCommandQueueFcnPtr;
-
-typedef cl_mem clCreateBufferFcnTy(cl_context Context, cl_mem_flags Flags,
-                                   size_t Size, void *HostPtr,
-                                   cl_int *ErrcodeRet);
-static clCreateBufferFcnTy *clCreateBufferFcnPtr;
-
-typedef cl_int
-clEnqueueWriteBufferFcnTy(cl_command_queue CommandQueue, cl_mem Buffer,
-                          cl_bool BlockingWrite, size_t Offset, size_t Size,
-                          const void *Ptr, cl_uint NumEventsInWaitList,
-                          const cl_event *EventWaitList, cl_event *Event);
-static clEnqueueWriteBufferFcnTy *clEnqueueWriteBufferFcnPtr;
-
-typedef cl_program clCreateProgramWithBinaryFcnTy(
-    cl_context Context, cl_uint NumDevices, const cl_device_id *DeviceList,
-    const size_t *Lengths, const unsigned char **Binaries, cl_int *BinaryStatus,
-    cl_int *ErrcodeRet);
-static clCreateProgramWithBinaryFcnTy *clCreateProgramWithBinaryFcnPtr;
-
-typedef cl_int clBuildProgramFcnTy(
-    cl_program Program, cl_uint NumDevices, const cl_device_id *DeviceList,
-    const char *Options,
-    void(CL_CALLBACK *pfn_notify)(cl_program Program, void *UserData),
-    void *UserData);
-static clBuildProgramFcnTy *clBuildProgramFcnPtr;
-
-typedef cl_kernel clCreateKernelFcnTy(cl_program Program,
-                                      const char *KernelName,
-                                      cl_int *ErrcodeRet);
-static clCreateKernelFcnTy *clCreateKernelFcnPtr;
-
-typedef cl_int clSetKernelArgFcnTy(cl_kernel Kernel, cl_uint ArgIndex,
-                                   size_t ArgSize, const void *ArgValue);
-static clSetKernelArgFcnTy *clSetKernelArgFcnPtr;
-
-typedef cl_int clEnqueueNDRangeKernelFcnTy(
-    cl_command_queue CommandQueue, cl_kernel Kernel, cl_uint WorkDim,
-    const size_t *GlobalWorkOffset, const size_t *GlobalWorkSize,
-    const size_t *LocalWorkSize, cl_uint NumEventsInWaitList,
-    const cl_event *EventWaitList, cl_event *Event);
-static clEnqueueNDRangeKernelFcnTy *clEnqueueNDRangeKernelFcnPtr;
-
-typedef cl_int clEnqueueReadBufferFcnTy(cl_command_queue CommandQueue,
-                                        cl_mem Buffer, cl_bool BlockingRead,
-                                        size_t Offset, size_t Size, void *Ptr,
-                                        cl_uint NumEventsInWaitList,
-                                        const cl_event *EventWaitList,
-                                        cl_event *Event);
-static clEnqueueReadBufferFcnTy *clEnqueueReadBufferFcnPtr;
-
-typedef cl_int clFlushFcnTy(cl_command_queue CommandQueue);
-static clFlushFcnTy *clFlushFcnPtr;
-
-typedef cl_int clFinishFcnTy(cl_command_queue CommandQueue);
-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 CommandQueue);
-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(stderr, "Load OpenCL Runtime API failed: %s. \n", Err);
-    return 0;
-  }
-  return FuncPtr;
-}
-
-static int initialDeviceAPILibrariesCL() {
-  HandleOpenCL = dlopen("libOpenCL.so", RTLD_LAZY);
-  if (!HandleOpenCL) {
-    fprintf(stderr, "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. */
-static void printOpenCLError(int Error);
-
-static void checkOpenCLError(int Ret, const char *format, ...) {
-  if (Ret == CL_SUCCESS)
-    return;
-
-  printOpenCLError(Ret);
-  va_list args;
-  va_start(args, format);
-  vfprintf(stderr, format, args);
-  va_end(args);
-  exit(-1);
-}
-
-static PollyGPUContext *initContextCL() {
-  dump_function();
-
-  PollyGPUContext *Context;
-
-  cl_platform_id PlatformID = NULL;
-  cl_device_id DeviceID = NULL;
-  cl_uint NumDevicesRet;
-  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(stderr, "Getting the \"handle\" for the OpenCL Runtime failed.\n");
-    exit(-1);
-  }
-
-  /* Get number of devices that support OpenCL. */
-  static const int NumberOfPlatforms = 1;
-  Ret = clGetPlatformIDsFcnPtr(NumberOfPlatforms, &PlatformID, NULL);
-  checkOpenCLError(Ret, "Failed to get platform IDs.\n");
-  // TODO: Extend to CL_DEVICE_TYPE_ALL?
-  static const int NumberOfDevices = 1;
-  Ret = clGetDeviceIDsFcnPtr(PlatformID, CL_DEVICE_TYPE_GPU, NumberOfDevices,
-                             &DeviceID, &NumDevicesRet);
-  checkOpenCLError(Ret, "Failed to get device IDs.\n");
-
-  GlobalDeviceID = DeviceID;
-  if (NumDevicesRet == 0) {
-    fprintf(stderr, "There is no device supporting OpenCL.\n");
-    exit(-1);
-  }
-
-  /* Get device revision. */
-  Ret =
-      clGetDeviceInfoFcnPtr(DeviceID, CL_DEVICE_VERSION, sizeof(DeviceRevision),
-                            DeviceRevision, &DeviceRevisionRetSize);
-  checkOpenCLError(Ret, "Failed to fetch device revision.\n");
-
-  /* Get device name. */
-  Ret = clGetDeviceInfoFcnPtr(DeviceID, CL_DEVICE_NAME, sizeof(DeviceName),
-                              DeviceName, &DeviceNameRetSize);
-  checkOpenCLError(Ret, "Failed to fetch device name.\n");
-
-  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(stderr, "Allocate memory for Polly GPU context failed.\n");
-    exit(-1);
-  }
-  Context->Context = (OpenCLContext *)malloc(sizeof(OpenCLContext));
-  if (Context->Context == 0) {
-    fprintf(stderr, "Allocate memory for Polly OpenCL context failed.\n");
-    exit(-1);
-  }
-  ((OpenCLContext *)Context->Context)->Context =
-      clCreateContextFcnPtr(NULL, NumDevicesRet, &DeviceID, NULL, NULL, &Ret);
-  checkOpenCLError(Ret, "Failed to create context.\n");
-
-  static const int ExtraProperties = 0;
-  ((OpenCLContext *)Context->Context)->CommandQueue =
-      clCreateCommandQueueFcnPtr(((OpenCLContext *)Context->Context)->Context,
-                                 DeviceID, ExtraProperties, &Ret);
-  checkOpenCLError(Ret, "Failed to create command queue.\n");
-
-  if (CacheMode)
-    CurrentContext = Context;
-
-  GlobalContext = Context;
-  return Context;
-}
-
-static void freeKernelCL(PollyGPUFunction *Kernel) {
-  dump_function();
-
-  if (CacheMode)
-    return;
-
-  if (!GlobalContext) {
-    fprintf(stderr, "GPGPU-code generation not correctly initialized.\n");
-    exit(-1);
-  }
-
-  cl_int Ret;
-  Ret = clFlushFcnPtr(((OpenCLContext *)GlobalContext->Context)->CommandQueue);
-  checkOpenCLError(Ret, "Failed to flush command queue.\n");
-  Ret = clFinishFcnPtr(((OpenCLContext *)GlobalContext->Context)->CommandQueue);
-  checkOpenCLError(Ret, "Failed to finish command queue.\n");
-
-  if (((OpenCLKernel *)Kernel->Kernel)->Kernel) {
-    cl_int Ret =
-        clReleaseKernelFcnPtr(((OpenCLKernel *)Kernel->Kernel)->Kernel);
-    checkOpenCLError(Ret, "Failed to release kernel.\n");
-  }
-
-  if (((OpenCLKernel *)Kernel->Kernel)->Program) {
-    cl_int Ret =
-        clReleaseProgramFcnPtr(((OpenCLKernel *)Kernel->Kernel)->Program);
-    checkOpenCLError(Ret, "Failed to release program.\n");
-  }
-
-  if (Kernel->Kernel)
-    free((OpenCLKernel *)Kernel->Kernel);
-
-  if (Kernel)
-    free(Kernel);
-}
-
-static PollyGPUFunction *getKernelCL(const char *BinaryBuffer,
-                                     const char *KernelName) {
-  dump_function();
-
-  if (!GlobalContext) {
-    fprintf(stderr, "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)->BinaryString ==
-            BinaryBuffer) {
-      debug_print("  -> using cached kernel\n");
-      return KernelCache[i];
-    }
-  }
-
-  PollyGPUFunction *Function = malloc(sizeof(PollyGPUFunction));
-  if (Function == 0) {
-    fprintf(stderr, "Allocate memory for Polly GPU function failed.\n");
-    exit(-1);
-  }
-  Function->Kernel = (OpenCLKernel *)malloc(sizeof(OpenCLKernel));
-  if (Function->Kernel == 0) {
-    fprintf(stderr, "Allocate memory for Polly OpenCL kernel failed.\n");
-    exit(-1);
-  }
-
-  if (!GlobalDeviceID) {
-    fprintf(stderr, "GPGPU-code generation not initialized correctly.\n");
-    exit(-1);
-  }
-
-  cl_int Ret;
-  size_t BinarySize = strlen(BinaryBuffer);
-  ((OpenCLKernel *)Function->Kernel)->Program = clCreateProgramWithBinaryFcnPtr(
-      ((OpenCLContext *)GlobalContext->Context)->Context, 1, &GlobalDeviceID,
-      (const size_t *)&BinarySize, (const unsigned char **)&BinaryBuffer, NULL,
-      &Ret);
-  checkOpenCLError(Ret, "Failed to create program from binary.\n");
-
-  Ret = clBuildProgramFcnPtr(((OpenCLKernel *)Function->Kernel)->Program, 1,
-                             &GlobalDeviceID, NULL, NULL, NULL);
-  checkOpenCLError(Ret, "Failed to build program.\n");
-
-  ((OpenCLKernel *)Function->Kernel)->Kernel = clCreateKernelFcnPtr(
-      ((OpenCLKernel *)Function->Kernel)->Program, KernelName, &Ret);
-  checkOpenCLError(Ret, "Failed to create kernel.\n");
-
-  ((OpenCLKernel *)Function->Kernel)->BinaryString = BinaryBuffer;
-
-  if (CacheMode) {
-    if (KernelCache[NextCacheItem])
-      freeKernelCL(KernelCache[NextCacheItem]);
-
-    KernelCache[NextCacheItem] = Function;
-
-    NextCacheItem = (NextCacheItem + 1) % KERNEL_CACHE_SIZE;
-  }
-
-  return Function;
-}
-
-static void copyFromHostToDeviceCL(void *HostData, PollyGPUDevicePtr *DevData,
-                                   long MemSize) {
-  dump_function();
-
-  if (!GlobalContext) {
-    fprintf(stderr, "GPGPU-code generation not correctly initialized.\n");
-    exit(-1);
-  }
-
-  cl_int Ret;
-  Ret = clEnqueueWriteBufferFcnPtr(
-      ((OpenCLContext *)GlobalContext->Context)->CommandQueue,
-      ((OpenCLDevicePtr *)DevData->DevicePtr)->MemObj, CL_TRUE, 0, MemSize,
-      HostData, 0, NULL, NULL);
-  checkOpenCLError(Ret, "Copying data from host memory to device failed.\n");
-}
-
-static void copyFromDeviceToHostCL(PollyGPUDevicePtr *DevData, void *HostData,
-                                   long MemSize) {
-  dump_function();
-
-  if (!GlobalContext) {
-    fprintf(stderr, "GPGPU-code generation not correctly initialized.\n");
-    exit(-1);
-  }
-
-  cl_int Ret;
-  Ret = clEnqueueReadBufferFcnPtr(
-      ((OpenCLContext *)GlobalContext->Context)->CommandQueue,
-      ((OpenCLDevicePtr *)DevData->DevicePtr)->MemObj, CL_TRUE, 0, MemSize,
-      HostData, 0, NULL, NULL);
-  checkOpenCLError(Ret, "Copying results from device to host memory failed.\n");
-}
-
-static void launchKernelCL(PollyGPUFunction *Kernel, unsigned int GridDimX,
-                           unsigned int GridDimY, unsigned int BlockDimX,
-                           unsigned int BlockDimY, unsigned int BlockDimZ,
-                           void **Parameters) {
-  dump_function();
-
-  cl_int Ret;
-  cl_uint NumArgs;
-
-  if (!GlobalContext) {
-    fprintf(stderr, "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), &NumArgs, NULL);
-  checkOpenCLError(Ret, "Failed to get number of kernel arguments.\n");
-
-  // TODO: Pass the size of the kernel arguments in to launchKernelCL, along
-  // with the arguments themselves. This is a dirty workaround that can be
-  // broken.
-  for (cl_uint i = 0; i < NumArgs; i++) {
-    Ret = clSetKernelArgFcnPtr(CLKernel->Kernel, i, 8, (void *)Parameters[i]);
-    if (Ret == CL_INVALID_ARG_SIZE) {
-      Ret = clSetKernelArgFcnPtr(CLKernel->Kernel, i, 4, (void *)Parameters[i]);
-      if (Ret == CL_INVALID_ARG_SIZE) {
-        Ret =
-            clSetKernelArgFcnPtr(CLKernel->Kernel, i, 2, (void *)Parameters[i]);
-        if (Ret == CL_INVALID_ARG_SIZE) {
-          Ret = clSetKernelArgFcnPtr(CLKernel->Kernel, i, 1,
-                                     (void *)Parameters[i]);
-          checkOpenCLError(Ret, "Failed to set Kernel argument %d.\n", i);
-        }
-      }
-    }
-    if (Ret != CL_SUCCESS && Ret != CL_INVALID_ARG_SIZE) {
-      fprintf(stderr, "Failed to set Kernel argument.\n");
-      printOpenCLError(Ret);
-      exit(-1);
-    }
-  }
-
-  unsigned int GridDimZ = 1;
-  size_t GlobalWorkSize[3] = {BlockDimX * GridDimX, BlockDimY * GridDimY,
-                              BlockDimZ * GridDimZ};
-  size_t LocalWorkSize[3] = {BlockDimX, BlockDimY, BlockDimZ};
-
-  static const int WorkDim = 3;
-  OpenCLContext *CLContext = (OpenCLContext *)GlobalContext->Context;
-  Ret = clEnqueueNDRangeKernelFcnPtr(CLContext->CommandQueue, CLKernel->Kernel,
-                                     WorkDim, NULL, GlobalWorkSize,
-                                     LocalWorkSize, 0, NULL, NULL);
-  checkOpenCLError(Ret, "Launching OpenCL kernel failed.\n");
-}
-
-static void freeDeviceMemoryCL(PollyGPUDevicePtr *Allocation) {
-  dump_function();
-
-  OpenCLDevicePtr *DevPtr = (OpenCLDevicePtr *)Allocation->DevicePtr;
-  cl_int Ret = clReleaseMemObjectFcnPtr((cl_mem)DevPtr->MemObj);
-  checkOpenCLError(Ret, "Failed to free device memory.\n");
-
-  free(DevPtr);
-  free(Allocation);
-}
-
-static PollyGPUDevicePtr *allocateMemoryForDeviceCL(long MemSize) {
-  dump_function();
-
-  if (!GlobalContext) {
-    fprintf(stderr, "GPGPU-code generation not correctly initialized.\n");
-    exit(-1);
-  }
-
-  PollyGPUDevicePtr *DevData = malloc(sizeof(PollyGPUDevicePtr));
-  if (DevData == 0) {
-    fprintf(stderr, "Allocate memory for GPU device memory pointer failed.\n");
-    exit(-1);
-  }
-  DevData->DevicePtr = (OpenCLDevicePtr *)malloc(sizeof(OpenCLDevicePtr));
-  if (DevData->DevicePtr == 0) {
-    fprintf(stderr, "Allocate memory for GPU device memory pointer failed.\n");
-    exit(-1);
-  }
-
-  cl_int Ret;
-  ((OpenCLDevicePtr *)DevData->DevicePtr)->MemObj =
-      clCreateBufferFcnPtr(((OpenCLContext *)GlobalContext->Context)->Context,
-                           CL_MEM_READ_WRITE, MemSize, NULL, &Ret);
-  checkOpenCLError(Ret,
-                   "Allocate memory for GPU device memory pointer failed.\n");
-
-  return DevData;
-}
-
-static void *getDevicePtrCL(PollyGPUDevicePtr *Allocation) {
-  dump_function();
-
-  OpenCLDevicePtr *DevPtr = (OpenCLDevicePtr *)Allocation->DevicePtr;
-  return (void *)DevPtr->MemObj;
-}
-
-static void synchronizeDeviceCL() {
-  dump_function();
-
-  if (!GlobalContext) {
-    fprintf(stderr, "GPGPU-code generation not correctly initialized.\n");
-    exit(-1);
-  }
-
-  if (clFinishFcnPtr(((OpenCLContext *)GlobalContext->Context)->CommandQueue) !=
-      CL_SUCCESS) {
-    fprintf(stderr, "Synchronizing device and host memory failed.\n");
-    exit(-1);
-  }
-}
-
-static void freeContextCL(PollyGPUContext *Context) {
-  dump_function();
-
-  cl_int Ret;
-
-  GlobalContext = NULL;
-
-  OpenCLContext *Ctx = (OpenCLContext *)Context->Context;
-  if (Ctx->CommandQueue) {
-    Ret = clReleaseCommandQueueFcnPtr(Ctx->CommandQueue);
-    checkOpenCLError(Ret, "Could not release command queue.\n");
-  }
-
-  if (Ctx->Context) {
-    Ret = clReleaseContextFcnPtr(Ctx->Context);
-    checkOpenCLError(Ret, "Could not release context.\n");
-  }
-
-  free(Ctx);
-  free(Context);
-}
-
-static void printOpenCLError(int Error) {
-
-  switch (Error) {
-  case CL_SUCCESS:
-    // Success, don't print an error.
-    break;
-
-  // JIT/Runtime errors.
-  case CL_DEVICE_NOT_FOUND:
-    fprintf(stderr, "Device not found.\n");
-    break;
-  case CL_DEVICE_NOT_AVAILABLE:
-    fprintf(stderr, "Device not available.\n");
-    break;
-  case CL_COMPILER_NOT_AVAILABLE:
-    fprintf(stderr, "Compiler not available.\n");
-    break;
-  case CL_MEM_OBJECT_ALLOCATION_FAILURE:
-    fprintf(stderr, "Mem object allocation failure.\n");
-    break;
-  case CL_OUT_OF_RESOURCES:
-    fprintf(stderr, "Out of resources.\n");
-    break;
-  case CL_OUT_OF_HOST_MEMORY:
-    fprintf(stderr, "Out of host memory.\n");
-    break;
-  case CL_PROFILING_INFO_NOT_AVAILABLE:
-    fprintf(stderr, "Profiling info not available.\n");
-    break;
-  case CL_MEM_COPY_OVERLAP:
-    fprintf(stderr, "Mem copy overlap.\n");
-    break;
-  case CL_IMAGE_FORMAT_MISMATCH:
-    fprintf(stderr, "Image format mismatch.\n");
-    break;
-  case CL_IMAGE_FORMAT_NOT_SUPPORTED:
-    fprintf(stderr, "Image format not supported.\n");
-    break;
-  case CL_BUILD_PROGRAM_FAILURE:
-    fprintf(stderr, "Build program failure.\n");
-    break;
-  case CL_MAP_FAILURE:
-    fprintf(stderr, "Map failure.\n");
-    break;
-  case CL_MISALIGNED_SUB_BUFFER_OFFSET:
-    fprintf(stderr, "Misaligned sub buffer offset.\n");
-    break;
-  case CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST:
-    fprintf(stderr, "Exec status error for events in wait list.\n");
-    break;
-  case CL_COMPILE_PROGRAM_FAILURE:
-    fprintf(stderr, "Compile program failure.\n");
-    break;
-  case CL_LINKER_NOT_AVAILABLE:
-    fprintf(stderr, "Linker not available.\n");
-    break;
-  case CL_LINK_PROGRAM_FAILURE:
-    fprintf(stderr, "Link program failure.\n");
-    break;
-  case CL_DEVICE_PARTITION_FAILED:
-    fprintf(stderr, "Device partition failed.\n");
-    break;
-  case CL_KERNEL_ARG_INFO_NOT_AVAILABLE:
-    fprintf(stderr, "Kernel arg info not available.\n");
-    break;
-
-  // Compiler errors.
-  case CL_INVALID_VALUE:
-    fprintf(stderr, "Invalid value.\n");
-    break;
-  case CL_INVALID_DEVICE_TYPE:
-    fprintf(stderr, "Invalid device type.\n");
-    break;
-  case CL_INVALID_PLATFORM:
-    fprintf(stderr, "Invalid platform.\n");
-    break;
-  case CL_INVALID_DEVICE:
-    fprintf(stderr, "Invalid device.\n");
-    break;
-  case CL_INVALID_CONTEXT:
-    fprintf(stderr, "Invalid context.\n");
-    break;
-  case CL_INVALID_QUEUE_PROPERTIES:
-    fprintf(stderr, "Invalid queue properties.\n");
-    break;
-  case CL_INVALID_COMMAND_QUEUE:
-    fprintf(stderr, "Invalid command queue.\n");
-    break;
-  case CL_INVALID_HOST_PTR:
-    fprintf(stderr, "Invalid host pointer.\n");
-    break;
-  case CL_INVALID_MEM_OBJECT:
-    fprintf(stderr, "Invalid memory object.\n");
-    break;
-  case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:
-    fprintf(stderr, "Invalid image format descriptor.\n");
-    break;
-  case CL_INVALID_IMAGE_SIZE:
-    fprintf(stderr, "Invalid image size.\n");
-    break;
-  case CL_INVALID_SAMPLER:
-    fprintf(stderr, "Invalid sampler.\n");
-    break;
-  case CL_INVALID_BINARY:
-    fprintf(stderr, "Invalid binary.\n");
-    break;
-  case CL_INVALID_BUILD_OPTIONS:
-    fprintf(stderr, "Invalid build options.\n");
-    break;
-  case CL_INVALID_PROGRAM:
-    fprintf(stderr, "Invalid program.\n");
-    break;
-  case CL_INVALID_PROGRAM_EXECUTABLE:
-    fprintf(stderr, "Invalid program executable.\n");
-    break;
-  case CL_INVALID_KERNEL_NAME:
-    fprintf(stderr, "Invalid kernel name.\n");
-    break;
-  case CL_INVALID_KERNEL_DEFINITION:
-    fprintf(stderr, "Invalid kernel definition.\n");
-    break;
-  case CL_INVALID_KERNEL:
-    fprintf(stderr, "Invalid kernel.\n");
-    break;
-  case CL_INVALID_ARG_INDEX:
-    fprintf(stderr, "Invalid arg index.\n");
-    break;
-  case CL_INVALID_ARG_VALUE:
-    fprintf(stderr, "Invalid arg value.\n");
-    break;
-  case CL_INVALID_ARG_SIZE:
-    fprintf(stderr, "Invalid arg size.\n");
-    break;
-  case CL_INVALID_KERNEL_ARGS:
-    fprintf(stderr, "Invalid kernel args.\n");
-    break;
-  case CL_INVALID_WORK_DIMENSION:
-    fprintf(stderr, "Invalid work dimension.\n");
-    break;
-  case CL_INVALID_WORK_GROUP_SIZE:
-    fprintf(stderr, "Invalid work group size.\n");
-    break;
-  case CL_INVALID_WORK_ITEM_SIZE:
-    fprintf(stderr, "Invalid work item size.\n");
-    break;
-  case CL_INVALID_GLOBAL_OFFSET:
-    fprintf(stderr, "Invalid global offset.\n");
-    break;
-  case CL_INVALID_EVENT_WAIT_LIST:
-    fprintf(stderr, "Invalid event wait list.\n");
-    break;
-  case CL_INVALID_EVENT:
-    fprintf(stderr, "Invalid event.\n");
-    break;
-  case CL_INVALID_OPERATION:
-    fprintf(stderr, "Invalid operation.\n");
-    break;
-  case CL_INVALID_GL_OBJECT:
-    fprintf(stderr, "Invalid GL object.\n");
-    break;
-  case CL_INVALID_BUFFER_SIZE:
-    fprintf(stderr, "Invalid buffer size.\n");
-    break;
-  case CL_INVALID_MIP_LEVEL:
-    fprintf(stderr, "Invalid mip level.\n");
-    break;
-  case CL_INVALID_GLOBAL_WORK_SIZE:
-    fprintf(stderr, "Invalid global work size.\n");
-    break;
-  case CL_INVALID_PROPERTY:
-    fprintf(stderr, "Invalid property.\n");
-    break;
-  case CL_INVALID_IMAGE_DESCRIPTOR:
-    fprintf(stderr, "Invalid image descriptor.\n");
-    break;
-  case CL_INVALID_COMPILER_OPTIONS:
-    fprintf(stderr, "Invalid compiler options.\n");
-    break;
-  case CL_INVALID_LINKER_OPTIONS:
-    fprintf(stderr, "Invalid linker options.\n");
-    break;
-  case CL_INVALID_DEVICE_PARTITION_COUNT:
-    fprintf(stderr, "Invalid device partition count.\n");
-    break;
-  case CL_INVALID_PIPE_SIZE:
-    fprintf(stderr, "Invalid pipe size.\n");
-    break;
-  case CL_INVALID_DEVICE_QUEUE:
-    fprintf(stderr, "Invalid device queue.\n");
-    break;
-
-  // NVIDIA specific error.
-  case -9999:
-    fprintf(stderr, "NVIDIA invalid read or write buffer.\n");
-    break;
-
-  default:
-    fprintf(stderr, "Unknown error code!\n");
-    break;
-  }
-}
-
-#endif /* HAS_LIBOPENCL */
-/******************************************************************************/
-/*                                   CUDA                                     */
-/******************************************************************************/
-#ifdef HAS_LIBCUDART
-
-struct CUDAContextT {
   CUcontext Cuda;
 };
 
-struct CUDAKernelT {
+struct PollyGPUFunctionT {
   CUfunction Cuda;
   CUmodule CudaModule;
-  const char *BinaryString;
+  const char *PTXString;
 };
 
-struct CUDADevicePtrT {
+struct PollyGPUDevicePtrT {
   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 CuMemAllocFcnTy(CUdeviceptr *, size_t);
 static CuMemAllocFcnTy *CuMemAllocFcnPtr;
 
 typedef CUresult CUDAAPI CuLaunchKernelFcnTy(
-    CUfunction F, unsigned int GridDimX, unsigned int GridDimY,
-    unsigned int gridDimZ, unsigned int blockDimX, unsigned int BlockDimY,
-    unsigned int BlockDimZ, unsigned int SharedMemBytes, CUstream HStream,
-    void **KernelParams, void **Extra);
+    CUfunction f, unsigned int gridDimX, unsigned int gridDimY,
+    unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY,
+    unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream,
+    void **kernelParams, void **extra);
 static CuLaunchKernelFcnTy *CuLaunchKernelFcnPtr;
 
 typedef CUresult CUDAAPI CuMemcpyDtoHFcnTy(void *, CUdeviceptr, size_t);
 static CuMemcpyDtoHFcnTy *CuMemcpyDtoHFcnPtr;
 
 typedef CUresult CUDAAPI CuMemcpyHtoDFcnTy(CUdeviceptr, const void *, size_t);
 static CuMemcpyHtoDFcnTy *CuMemcpyHtoDFcnPtr;
 

 typedef CUresult CUDAAPI CuDeviceGetFcnTy(CUdevice *, int);
 static CuDeviceGetFcnTy *CuDeviceGetFcnPtr;
 
 typedef CUresult CUDAAPI CuModuleLoadDataExFcnTy(CUmodule *, const void *,
                                                  unsigned int, CUjit_option *,
                                                  void **);
 static CuModuleLoadDataExFcnTy *CuModuleLoadDataExFcnPtr;
 
-typedef CUresult CUDAAPI CuModuleLoadDataFcnTy(CUmodule *Module,
-                                               const void *Image);
+typedef CUresult CUDAAPI CuModuleLoadDataFcnTy(CUmodule *module,
+                                               const void *image);
 static CuModuleLoadDataFcnTy *CuModuleLoadDataFcnPtr;
 
 typedef CUresult CUDAAPI CuModuleGetFunctionFcnTy(CUfunction *, CUmodule,
                                                   const char *);
 static CuModuleGetFunctionFcnTy *CuModuleGetFunctionFcnPtr;
 
 typedef CUresult CUDAAPI CuDeviceComputeCapabilityFcnTy(int *, int *, CUdevice);
 static CuDeviceComputeCapabilityFcnTy *CuDeviceComputeCapabilityFcnPtr;
 
 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);
+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);
+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);
+typedef CUresult CUDAAPI CuLinkCompleteFcnTy(CUlinkState state, void **cubinOut,
+                                             size_t *sizeOut);
 static CuLinkCompleteFcnTy *CuLinkCompleteFcnPtr;
 
-typedef CUresult CUDAAPI CuLinkDestroyFcnTy(CUlinkState State);
+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 *getAPIHandleCUDA(void *Handle, const char *FuncName) {
+static void *getAPIHandle(void *Handle, const char *FuncName) {
   char *Err;
   void *FuncPtr;
   dlerror();
   FuncPtr = dlsym(Handle, FuncName);
   if ((Err = dlerror()) != 0) {
-    fprintf(stderr, "Load CUDA driver API failed: %s. \n", Err);
+    fprintf(stdout, "Load CUDA driver API failed: %s. \n", Err);
     return 0;
   }
   return FuncPtr;
 }
 
-static int initialDeviceAPILibrariesCUDA() {
+static int initialDeviceAPILibraries() {
   HandleCuda = dlopen("libcuda.so", RTLD_LAZY);
   if (!HandleCuda) {
-    fprintf(stderr, "Cannot open library: %s. \n", dlerror());
+    printf("Cannot open library: %s. \n", dlerror());
     return 0;
   }
 
   HandleCudaRT = dlopen("libcudart.so", RTLD_LAZY);
   if (!HandleCudaRT) {
-    fprintf(stderr, "Cannot open library: %s. \n", dlerror());
+    printf("Cannot open library: %s. \n", dlerror());
     return 0;
   }
 
   return 1;
 }
 
-static int initialDeviceAPIsCUDA() {
-  if (initialDeviceAPILibrariesCUDA() == 0)
+static int initialDeviceAPIs() {
+  if (initialDeviceAPILibraries() == 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 *)getAPIHandleCUDA(HandleCuda, "cuLaunchKernel");
+      (CuLaunchKernelFcnTy *)getAPIHandle(HandleCuda, "cuLaunchKernel");
 
   CuMemAllocFcnPtr =
-      (CuMemAllocFcnTy *)getAPIHandleCUDA(HandleCuda, "cuMemAlloc_v2");
+      (CuMemAllocFcnTy *)getAPIHandle(HandleCuda, "cuMemAlloc_v2");
 
-  CuMemFreeFcnPtr =
-      (CuMemFreeFcnTy *)getAPIHandleCUDA(HandleCuda, "cuMemFree_v2");
+  CuMemFreeFcnPtr = (CuMemFreeFcnTy *)getAPIHandle(HandleCuda, "cuMemFree_v2");
 
   CuMemcpyDtoHFcnPtr =
-      (CuMemcpyDtoHFcnTy *)getAPIHandleCUDA(HandleCuda, "cuMemcpyDtoH_v2");
+      (CuMemcpyDtoHFcnTy *)getAPIHandle(HandleCuda, "cuMemcpyDtoH_v2");
 
   CuMemcpyHtoDFcnPtr =
-      (CuMemcpyHtoDFcnTy *)getAPIHandleCUDA(HandleCuda, "cuMemcpyHtoD_v2");
+      (CuMemcpyHtoDFcnTy *)getAPIHandle(HandleCuda, "cuMemcpyHtoD_v2");
 
   CuModuleUnloadFcnPtr =
-      (CuModuleUnloadFcnTy *)getAPIHandleCUDA(HandleCuda, "cuModuleUnload");
+      (CuModuleUnloadFcnTy *)getAPIHandle(HandleCuda, "cuModuleUnload");
 
   CuCtxDestroyFcnPtr =
-      (CuCtxDestroyFcnTy *)getAPIHandleCUDA(HandleCuda, "cuCtxDestroy");
+      (CuCtxDestroyFcnTy *)getAPIHandle(HandleCuda, "cuCtxDestroy");
 
-  CuInitFcnPtr = (CuInitFcnTy *)getAPIHandleCUDA(HandleCuda, "cuInit");
+  CuInitFcnPtr = (CuInitFcnTy *)getAPIHandle(HandleCuda, "cuInit");
 
   CuDeviceGetCountFcnPtr =
-      (CuDeviceGetCountFcnTy *)getAPIHandleCUDA(HandleCuda, "cuDeviceGetCount");
+      (CuDeviceGetCountFcnTy *)getAPIHandle(HandleCuda, "cuDeviceGetCount");
 
   CuDeviceGetFcnPtr =
-      (CuDeviceGetFcnTy *)getAPIHandleCUDA(HandleCuda, "cuDeviceGet");
+      (CuDeviceGetFcnTy *)getAPIHandle(HandleCuda, "cuDeviceGet");
 
   CuCtxCreateFcnPtr =
-      (CuCtxCreateFcnTy *)getAPIHandleCUDA(HandleCuda, "cuCtxCreate_v2");
+      (CuCtxCreateFcnTy *)getAPIHandle(HandleCuda, "cuCtxCreate_v2");
 
-  CuModuleLoadDataExFcnPtr = (CuModuleLoadDataExFcnTy *)getAPIHandleCUDA(
-      HandleCuda, "cuModuleLoadDataEx");
+  CuModuleLoadDataExFcnPtr =
+      (CuModuleLoadDataExFcnTy *)getAPIHandle(HandleCuda, "cuModuleLoadDataEx");
 
   CuModuleLoadDataFcnPtr =
-      (CuModuleLoadDataFcnTy *)getAPIHandleCUDA(HandleCuda, "cuModuleLoadData");
+      (CuModuleLoadDataFcnTy *)getAPIHandle(HandleCuda, "cuModuleLoadData");
 
-  CuModuleGetFunctionFcnPtr = (CuModuleGetFunctionFcnTy *)getAPIHandleCUDA(
+  CuModuleGetFunctionFcnPtr = (CuModuleGetFunctionFcnTy *)getAPIHandle(
       HandleCuda, "cuModuleGetFunction");
 
   CuDeviceComputeCapabilityFcnPtr =
-      (CuDeviceComputeCapabilityFcnTy *)getAPIHandleCUDA(
+      (CuDeviceComputeCapabilityFcnTy *)getAPIHandle(
           HandleCuda, "cuDeviceComputeCapability");
 
   CuDeviceGetNameFcnPtr =
-      (CuDeviceGetNameFcnTy *)getAPIHandleCUDA(HandleCuda, "cuDeviceGetName");
+      (CuDeviceGetNameFcnTy *)getAPIHandle(HandleCuda, "cuDeviceGetName");
 
   CuLinkAddDataFcnPtr =
-      (CuLinkAddDataFcnTy *)getAPIHandleCUDA(HandleCuda, "cuLinkAddData");
+      (CuLinkAddDataFcnTy *)getAPIHandle(HandleCuda, "cuLinkAddData");
 
   CuLinkCreateFcnPtr =
-      (CuLinkCreateFcnTy *)getAPIHandleCUDA(HandleCuda, "cuLinkCreate");
+      (CuLinkCreateFcnTy *)getAPIHandle(HandleCuda, "cuLinkCreate");
 
   CuLinkCompleteFcnPtr =
-      (CuLinkCompleteFcnTy *)getAPIHandleCUDA(HandleCuda, "cuLinkComplete");
+      (CuLinkCompleteFcnTy *)getAPIHandle(HandleCuda, "cuLinkComplete");
 
   CuLinkDestroyFcnPtr =
-      (CuLinkDestroyFcnTy *)getAPIHandleCUDA(HandleCuda, "cuLinkDestroy");
+      (CuLinkDestroyFcnTy *)getAPIHandle(HandleCuda, "cuLinkDestroy");
 
   CuCtxSynchronizeFcnPtr =
-      (CuCtxSynchronizeFcnTy *)getAPIHandleCUDA(HandleCuda, "cuCtxSynchronize");
+      (CuCtxSynchronizeFcnTy *)getAPIHandle(HandleCuda, "cuCtxSynchronize");
 
   /* Get function pointer to CUDA Runtime APIs. */
-  CudaThreadSynchronizeFcnPtr = (CudaThreadSynchronizeFcnTy *)getAPIHandleCUDA(
+  CudaThreadSynchronizeFcnPtr = (CudaThreadSynchronizeFcnTy *)getAPIHandle(
       HandleCudaRT, "cudaThreadSynchronize");
 
   return 1;
 }
 
-static PollyGPUContext *initContextCUDA() {
+PollyGPUContext *polly_initContext() {
+  DebugMode = getenv("POLLY_DEBUG") != 0;
+
   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 (initialDeviceAPIsCUDA() == 0) {
-    fprintf(stderr, "Getting the \"handle\" for the CUDA driver API failed.\n");
+  if (initialDeviceAPIs() == 0) {
+    fprintf(stdout, "Getting the \"handle\" for the CUDA driver API failed.\n");
     exit(-1);
   }
 
   if (CuInitFcnPtr(0) != CUDA_SUCCESS) {
-    fprintf(stderr, "Initializing the CUDA driver API failed.\n");
+    fprintf(stdout, "Initializing the CUDA driver API failed.\n");
     exit(-1);
   }
 
   /* Get number of devices that supports CUDA. */
   CuDeviceGetCountFcnPtr(&DeviceCount);
   if (DeviceCount == 0) {
-    fprintf(stderr, "There is no device supporting CUDA.\n");
+    fprintf(stdout, "There is no device supporting CUDA.\n");
     exit(-1);
   }
 
   CuDeviceGetFcnPtr(&Device, 0);
 
   /* Get compute capabilities and the device name. */
   CuDeviceComputeCapabilityFcnPtr(&Major, &Minor, Device);
   CuDeviceGetNameFcnPtr(DeviceName, 256, Device);
   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(stderr, "Allocate memory for Polly GPU context failed.\n");
+    fprintf(stdout, "Allocate memory for Polly GPU context failed.\n");
     exit(-1);
   }
-  Context->Context = malloc(sizeof(CUDAContext));
-  if (Context->Context == 0) {
-    fprintf(stderr, "Allocate memory for Polly CUDA context failed.\n");
-    exit(-1);
-  }
-  CuCtxCreateFcnPtr(&(((CUDAContext *)Context->Context)->Cuda), 0, Device);
+  CuCtxCreateFcnPtr(&(Context->Cuda), 0, Device);
+
+  CacheMode = getenv("POLLY_NOCACHE") == 0;
 
   if (CacheMode)
     CurrentContext = Context;
 
   return Context;
 }
 
-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);
+static void freeKernel(PollyGPUFunction *Kernel) {
+  if (Kernel->CudaModule)
+    CuModuleUnloadFcnPtr(Kernel->CudaModule);
 
   if (Kernel)
     free(Kernel);
 }
 
-static PollyGPUFunction *getKernelCUDA(const char *BinaryBuffer,
-                                       const char *KernelName) {
+#define KERNEL_CACHE_SIZE 10
+
+PollyGPUFunction *polly_getKernel(const char *PTXBuffer,
   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] &&
-        ((CUDAKernel *)KernelCache[i]->Kernel)->BinaryString == BinaryBuffer) {
+    if (KernelCache[i] && KernelCache[i]->PTXString == PTXBuffer) {
       debug_print("  -> using cached kernel\n");
       return KernelCache[i];
     }
   }
 
   PollyGPUFunction *Function = malloc(sizeof(PollyGPUFunction));
+
   if (Function == 0) {
-    fprintf(stderr, "Allocate memory for Polly GPU function failed.\n");
-    exit(-1);
-  }
-  Function->Kernel = (CUDAKernel *)malloc(sizeof(CUDAKernel));
-  if (Function->Kernel == 0) {
-    fprintf(stderr, "Allocate memory for Polly CUDA function failed.\n");
+    fprintf(stdout, "Allocate memory for Polly GPU function failed.\n");
     exit(-1);
   }
 
   CUresult Res;
   CUlinkState LState;
   CUjit_option Options[6];
   void *OptionVals[6];
   float Walltime = 0;
PollyGPUFunction *polly_getKernel(const char *PTXBuffer,
   OptionVals[4] = (void *)LogSize;
   // Make the linker verbose
   Options[5] = CU_JIT_LOG_VERBOSE;
   OptionVals[5] = (void *)1;
 
   memset(ErrorLog, 0, sizeof(ErrorLog));
 
   CuLinkCreateFcnPtr(6, Options, OptionVals, &LState);
-  Res = CuLinkAddDataFcnPtr(LState, CU_JIT_INPUT_PTX, (void *)BinaryBuffer,
-                            strlen(BinaryBuffer) + 1, 0, 0, 0, 0);
+  Res = CuLinkAddDataFcnPtr(LState, CU_JIT_INPUT_PTX, (void *)PTXBuffer,
+                            strlen(PTXBuffer) + 1, 0, 0, 0, 0);
   if (Res != CUDA_SUCCESS) {
-    fprintf(stderr, "PTX Linker Error:\n%s\n%s", ErrorLog, InfoLog);
+    fprintf(stdout, "PTX Linker Error:\n%s\n%s", ErrorLog, InfoLog);
     exit(-1);
   }
 
   Res = CuLinkCompleteFcnPtr(LState, &CuOut, &OutSize);
   if (Res != CUDA_SUCCESS) {
-    fprintf(stderr, "Complete ptx linker step failed.\n");
-    fprintf(stderr, "\n%s\n", ErrorLog);
+    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(&(((CUDAKernel *)Function->Kernel)->CudaModule),
-                               CuOut);
+  Res = CuModuleLoadDataFcnPtr(&(Function->CudaModule), CuOut);
   if (Res != CUDA_SUCCESS) {
-    fprintf(stderr, "Loading ptx assembly text failed.\n");
+    fprintf(stdout, "Loading ptx assembly text failed.\n");
     exit(-1);
   }
 
-  Res = CuModuleGetFunctionFcnPtr(&(((CUDAKernel *)Function->Kernel)->Cuda),
-                                  ((CUDAKernel *)Function->Kernel)->CudaModule,
+  Res = CuModuleGetFunctionFcnPtr(&(Function->Cuda), Function->CudaModule,
                                   KernelName);
   if (Res != CUDA_SUCCESS) {
-    fprintf(stderr, "Loading kernel function failed.\n");
+    fprintf(stdout, "Loading kernel function failed.\n");
     exit(-1);
   }
 
   CuLinkDestroyFcnPtr(LState);
 
-  ((CUDAKernel *)Function->Kernel)->BinaryString = BinaryBuffer;
+  Function->PTXString = PTXBuffer;
 
   if (CacheMode) {
     if (KernelCache[NextCacheItem])
-      freeKernelCUDA(KernelCache[NextCacheItem]);
+      freeKernel(KernelCache[NextCacheItem]);
 
     KernelCache[NextCacheItem] = Function;
 
     NextCacheItem = (NextCacheItem + 1) % KERNEL_CACHE_SIZE;
   }
 
   return Function;
 }
 
-static void synchronizeDeviceCUDA() {
+void polly_freeKernel(PollyGPUFunction *Kernel) {
   dump_function();
-  if (CuCtxSynchronizeFcnPtr() != CUDA_SUCCESS) {
-    fprintf(stderr, "Synchronizing device and host memory failed.\n");
-    exit(-1);
-  }
+
+  if (CacheMode)
+    return;
+
+  freeKernel(Kernel);
 }
 
-static void copyFromHostToDeviceCUDA(void *HostData, PollyGPUDevicePtr *DevData,
-                                     long MemSize) {
+void polly_copyFromHostToDevice(void *HostData, PollyGPUDevicePtr *DevData,
   dump_function();
 
-  CUdeviceptr CuDevData = ((CUDADevicePtr *)DevData->DevicePtr)->Cuda;
+  CUdeviceptr CuDevData = DevData->Cuda;
   CuMemcpyHtoDFcnPtr(CuDevData, HostData, MemSize);
 }
 
-static void copyFromDeviceToHostCUDA(PollyGPUDevicePtr *DevData, void *HostData,
-                                     long MemSize) {
+void polly_copyFromDeviceToHost(PollyGPUDevicePtr *DevData, void *HostData,
   dump_function();
 
-  if (CuMemcpyDtoHFcnPtr(HostData, ((CUDADevicePtr *)DevData->DevicePtr)->Cuda,
-                         MemSize) != CUDA_SUCCESS) {
-    fprintf(stderr, "Copying results from device to host memory failed.\n");
+  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);
   }
 }
 
-static void launchKernelCUDA(PollyGPUFunction *Kernel, unsigned int GridDimX,
-                             unsigned int GridDimY, unsigned int BlockDimX,
-                             unsigned int BlockDimY, unsigned int BlockDimZ,
-                             void **Parameters) {
+void polly_launchKernel(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(((CUDAKernel *)Kernel->Kernel)->Cuda, GridDimX,
-                           GridDimY, GridDimZ, BlockDimX, BlockDimY, BlockDimZ,
-                           SharedMemBytes, Stream, Parameters, Extra);
+  Res = CuLaunchKernelFcnPtr(Kernel->Cuda, GridDimX, GridDimY, GridDimZ,
+                             BlockDimX, BlockDimY, BlockDimZ, SharedMemBytes,
+                             Stream, Parameters, Extra);
   if (Res != CUDA_SUCCESS) {
-    fprintf(stderr, "Launching CUDA kernel failed.\n");
+    fprintf(stdout, "Launching CUDA kernel failed.\n");
     exit(-1);
   }
 }
 
-static void freeDeviceMemoryCUDA(PollyGPUDevicePtr *Allocation) {
+void polly_freeDeviceMemory(PollyGPUDevicePtr *Allocation) {
   dump_function();
-  CUDADevicePtr *DevPtr = (CUDADevicePtr *)Allocation->DevicePtr;
-  CuMemFreeFcnPtr((CUdeviceptr)DevPtr->Cuda);
-  free(DevPtr);
+  CuMemFreeFcnPtr((CUdeviceptr)Allocation->Cuda);
   free(Allocation);
 }
 
-static PollyGPUDevicePtr *allocateMemoryForDeviceCUDA(long MemSize) {
+PollyGPUDevicePtr *polly_allocateMemoryForDevice(long MemSize) {
   dump_function();
 
   PollyGPUDevicePtr *DevData = malloc(sizeof(PollyGPUDevicePtr));
+
   if (DevData == 0) {
-    fprintf(stderr, "Allocate memory for GPU device memory pointer failed.\n");
-    exit(-1);
-  }
-  DevData->DevicePtr = (CUDADevicePtr *)malloc(sizeof(CUDADevicePtr));
-  if (DevData->DevicePtr == 0) {
-    fprintf(stderr, "Allocate memory for GPU device memory pointer failed.\n");
+    fprintf(stdout, "Allocate memory for GPU device memory pointer failed.\n");
     exit(-1);
   }
 
-  CUresult Res =
-      CuMemAllocFcnPtr(&(((CUDADevicePtr *)DevData->DevicePtr)->Cuda), MemSize);
+  CUresult Res = CuMemAllocFcnPtr(&(DevData->Cuda), MemSize);
 
   if (Res != CUDA_SUCCESS) {
-    fprintf(stderr, "Allocate memory for GPU device memory pointer failed.\n");
+    fprintf(stdout, "Allocate memory for GPU device memory pointer failed.\n");
     exit(-1);
   }
 
   return DevData;
 }
 
-static void *getDevicePtrCUDA(PollyGPUDevicePtr *Allocation) {
-  dump_function();
-
-  CUDADevicePtr *DevPtr = (CUDADevicePtr *)Allocation->DevicePtr;
-  return (void *)DevPtr->Cuda;
-}
-
-static 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);
-}
-
-#endif /* HAS_LIBCUDART */
-/******************************************************************************/
-/*                                    API                                     */
-/******************************************************************************/
-
-PollyGPUContext *polly_initContext() {
-  DebugMode = getenv("POLLY_DEBUG") != 0;
-  CacheMode = getenv("POLLY_NOCACHE") == 0;
-
-  dump_function();
-
-  PollyGPUContext *Context;
-
-  switch (Runtime) {
-#ifdef HAS_LIBCUDART
-  case RUNTIME_CUDA:
-    Context = initContextCUDA();
-    break;
-#endif /* HAS_LIBCUDART */
-#ifdef HAS_LIBOPENCL
-  case RUNTIME_CL:
-    Context = initContextCL();
-    break;
-#endif /* HAS_LIBOPENCL */
-  default:
-    err_runtime();
-  }
-
-  return Context;
-}
-
-void polly_freeKernel(PollyGPUFunction *Kernel) {
-  dump_function();
-
-  switch (Runtime) {
-#ifdef HAS_LIBCUDART
-  case RUNTIME_CUDA:
-    freeKernelCUDA(Kernel);
-    break;
-#endif /* HAS_LIBCUDART */
-#ifdef HAS_LIBOPENCL
-  case RUNTIME_CL:
-    freeKernelCL(Kernel);
-    break;
-#endif /* HAS_LIBOPENCL */
-  default:
-    err_runtime();
-  }
-}
-
-PollyGPUFunction *polly_getKernel(const char *BinaryBuffer,
-                                  const char *KernelName) {
-  dump_function();
-
-  PollyGPUFunction *Function;
-
-  switch (Runtime) {
-#ifdef HAS_LIBCUDART
-  case RUNTIME_CUDA:
-    Function = getKernelCUDA(BinaryBuffer, KernelName);
-    break;
-#endif /* HAS_LIBCUDART */
-#ifdef HAS_LIBOPENCL
-  case RUNTIME_CL:
-    Function = getKernelCL(BinaryBuffer, KernelName);
-    break;
-#endif /* HAS_LIBOPENCL */
-  default:
-    err_runtime();
-  }
-
-  return Function;
-}
-
-void polly_copyFromHostToDevice(void *HostData, PollyGPUDevicePtr *DevData,
-                                long MemSize) {
-  dump_function();
-
-  switch (Runtime) {
-#ifdef HAS_LIBCUDART
-  case RUNTIME_CUDA:
-    copyFromHostToDeviceCUDA(HostData, DevData, MemSize);
-    break;
-#endif /* HAS_LIBCUDART */
-#ifdef HAS_LIBOPENCL
-  case RUNTIME_CL:
-    copyFromHostToDeviceCL(HostData, DevData, MemSize);
-    break;
-#endif /* HAS_LIBOPENCL */
-  default:
-    err_runtime();
-  }
-}
-
-void polly_copyFromDeviceToHost(PollyGPUDevicePtr *DevData, void *HostData,
-                                long MemSize) {
-  dump_function();
-
-  switch (Runtime) {
-#ifdef HAS_LIBCUDART
-  case RUNTIME_CUDA:
-    copyFromDeviceToHostCUDA(DevData, HostData, MemSize);
-    break;
-#endif /* HAS_LIBCUDART */
-#ifdef HAS_LIBOPENCL
-  case RUNTIME_CL:
-    copyFromDeviceToHostCL(DevData, HostData, MemSize);
-    break;
-#endif /* HAS_LIBOPENCL */
-  default:
-    err_runtime();
-  }
-}
-
-void polly_launchKernel(PollyGPUFunction *Kernel, unsigned int GridDimX,
-                        unsigned int GridDimY, unsigned int BlockDimX,
-                        unsigned int BlockDimY, unsigned int BlockDimZ,
-                        void **Parameters) {
-  dump_function();
-
-  switch (Runtime) {
-#ifdef HAS_LIBCUDART
-  case RUNTIME_CUDA:
-    launchKernelCUDA(Kernel, GridDimX, GridDimY, BlockDimX, BlockDimY,
-                     BlockDimZ, Parameters);
-    break;
-#endif /* HAS_LIBCUDART */
-#ifdef HAS_LIBOPENCL
-  case RUNTIME_CL:
-    launchKernelCL(Kernel, GridDimX, GridDimY, BlockDimX, BlockDimY, BlockDimZ,
-                   Parameters);
-    break;
-#endif /* HAS_LIBOPENCL */
-  default:
-    err_runtime();
-  }
-}
-
-void polly_freeDeviceMemory(PollyGPUDevicePtr *Allocation) {
-  dump_function();
-
-  switch (Runtime) {
-#ifdef HAS_LIBCUDART
-  case RUNTIME_CUDA:
-    freeDeviceMemoryCUDA(Allocation);
-    break;
-#endif /* HAS_LIBCUDART */
-#ifdef HAS_LIBOPENCL
-  case RUNTIME_CL:
-    freeDeviceMemoryCL(Allocation);
-    break;
-#endif /* HAS_LIBOPENCL */
-  default:
-    err_runtime();
-  }
-}
-
-PollyGPUDevicePtr *polly_allocateMemoryForDevice(long MemSize) {
-  dump_function();
-
-  PollyGPUDevicePtr *DevData;
-
-  switch (Runtime) {
-#ifdef HAS_LIBCUDART
-  case RUNTIME_CUDA:
-    DevData = allocateMemoryForDeviceCUDA(MemSize);
-    break;
-#endif /* HAS_LIBCUDART */
-#ifdef HAS_LIBOPENCL
-  case RUNTIME_CL:
-    DevData = allocateMemoryForDeviceCL(MemSize);
-    break;
-#endif /* HAS_LIBOPENCL */
-  default:
-    err_runtime();
-  }
-
-  return DevData;
-}
-
 void *polly_getDevicePtr(PollyGPUDevicePtr *Allocation) {
   dump_function();
 
-  void *DevPtr;
-
-  switch (Runtime) {
-#ifdef HAS_LIBCUDART
-  case RUNTIME_CUDA:
-    DevPtr = getDevicePtrCUDA(Allocation);
-    break;
-#endif /* HAS_LIBCUDART */
-#ifdef HAS_LIBOPENCL
-  case RUNTIME_CL:
-    DevPtr = getDevicePtrCL(Allocation);
-    break;
-#endif /* HAS_LIBOPENCL */
-  default:
-    err_runtime();
-  }
-
-  return DevPtr;
-}
-
-void polly_synchronizeDevice() {
-  dump_function();
-
-  switch (Runtime) {
-#ifdef HAS_LIBCUDART
-  case RUNTIME_CUDA:
-    synchronizeDeviceCUDA();
-    break;
-#endif /* HAS_LIBCUDART */
-#ifdef HAS_LIBOPENCL
-  case RUNTIME_CL:
-    synchronizeDeviceCL();
-    break;
-#endif /* HAS_LIBOPENCL */
-  default:
-    err_runtime();
-  }
+  return (void *)Allocation->Cuda;
 }
 
 void polly_freeContext(PollyGPUContext *Context) {
   dump_function();
 
   if (CacheMode)
     return;
 
-  switch (Runtime) {
-#ifdef HAS_LIBCUDART
-  case RUNTIME_CUDA:
-    freeContextCUDA(Context);
-    break;
-#endif /* HAS_LIBCUDART */
-#ifdef HAS_LIBOPENCL
-  case RUNTIME_CL:
-    freeContextCL(Context);
-    break;
-#endif /* HAS_LIBOPENCL */
-  default:
-    err_runtime();
-  }
-}
-
-/* Initialize GPUJIT with CUDA as runtime library. */
-PollyGPUContext *polly_initContextCUDA() {
-#ifdef HAS_LIBCUDART
-  Runtime = RUNTIME_CUDA;
-  return polly_initContext();
-#else
-  fprintf(stderr, "GPU Runtime was built without CUDA support.\n");
-  exit(-1);
-#endif /* HAS_LIBCUDART */
+  if (Context->Cuda) {
+    CuCtxDestroyFcnPtr(Context->Cuda);
+    free(Context);
   }
 
-/* Initialize GPUJIT with OpenCL as runtime library. */
-PollyGPUContext *polly_initContextCL() {
-#ifdef HAS_LIBOPENCL
-  Runtime = RUNTIME_CL;
-  return polly_initContext();
-#else
-  fprintf(stderr, "GPU Runtime was built without OpenCL support.\n");
-  exit(-1);
-#endif /* HAS_LIBOPENCL */
+  dlclose(HandleCuda);
+  dlclose(HandleCudaRT);
 }