diff --git a/include/cocl/llvm_dump.h b/include/cocl/llvm_dump.h
new file mode 100644
index 00000000..147d6e9c
--- /dev/null
+++ b/include/cocl/llvm_dump.h
@@ -0,0 +1,29 @@
+// Copyright Hugh Perkins 2016, 2017
+
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+
+//     http://www.apache.org/licenses/LICENSE-2.0
+
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+/// \author Aksel Alpay
+
+#ifndef COCL_LLVM_DUMP_H
+#define COCL_LLVM_DUMP_H
+
+#include <llvm/Support/Debug.h>
+
+#if LLVM_VERSION_MAJOR > 4
+#define COCL_LLVM_DUMP(entity) (entity)->print(llvm::dbgs(), true)
+#else
+#define COCL_LLVM_DUMP(entity) ((entity)->dump())
+#endif
+
+#endif
+
diff --git a/include/cocl/new_instruction_dumper.h b/include/cocl/new_instruction_dumper.h
index 22d88500..36755e37 100644
--- a/include/cocl/new_instruction_dumper.h
+++ b/include/cocl/new_instruction_dumper.h
@@ -17,7 +17,7 @@
 #include "cocl/LocalValueInfo.h"
 #include "cocl/InstructionDumper.h"
 #include "cocl/shims.h"
-
+#include "cocl/llvm_dump.h"
 #include <string>
 #include <stdexcept>
 
@@ -29,7 +29,7 @@ class NeedValueDependencyException : public std::exception {
     }
     virtual const char* what() const throw()
     {
-        value->dump();
+        COCL_LLVM_DUMP(value);
         return "Need dependent value";
     }
     llvm::Value *value;
diff --git a/include/cocl/patch_hostside.h b/include/cocl/patch_hostside.h
index 0702be85..06e3ae5c 100644
--- a/include/cocl/patch_hostside.h
+++ b/include/cocl/patch_hostside.h
@@ -58,6 +58,8 @@ Ok, so the doc is mostly below, inside the class declaration, at the bottom of t
 #include <vector>
 #include <memory>
 
+#include "cocl/llvm_dump.h"
+
 #include "llvm/IR/Module.h"
 // #include "llvm/IR/Verifier.h"
 #include "llvm/IR/Type.h"
@@ -148,7 +150,7 @@ class GenericCallInst_Call : public GenericCallInst {
     virtual llvm::Value *getOperand(int idx) override { return inst->getArgOperand(idx); }
     virtual llvm::Module *getModule() override { return inst->getModule(); }
     virtual llvm::Instruction *getInst() override { return inst; }
-    virtual void dump() override { inst->dump(); }
+    virtual void dump() override { COCL_LLVM_DUMP(inst); }
 };
 
 class GenericCallInst_Invoke : public GenericCallInst {
@@ -159,7 +161,7 @@ class GenericCallInst_Invoke : public GenericCallInst {
     virtual llvm::Value *getOperand(int idx) override { return inst->getArgOperand(idx); }
     virtual llvm::Module *getModule() override { return inst->getModule(); }
     virtual llvm::Instruction *getInst() override { return inst; }
-    virtual void dump() override { inst->dump(); }
+    virtual void dump() override { COCL_LLVM_DUMP(inst); }
 };
 
 class PatchHostside {
diff --git a/src/ClWriter.cpp b/src/ClWriter.cpp
index 937f26fe..56cb608c 100644
--- a/src/ClWriter.cpp
+++ b/src/ClWriter.cpp
@@ -24,6 +24,8 @@
 #include "llvm/IR/Instructions.h"
 #include "llvm/IR/Constants.h"
 
+#include "cocl/llvm_dump.h"
+
 #include <iostream>
 using namespace std;
 using namespace llvm;
@@ -74,7 +76,7 @@ void ClWriter::writeInlineCl(std::string indent, std::ostream &os) { // writes a
     }
     if(localValueInfo->toBeDeclared) {
         if(!localValueInfo->expressionValid) {
-            localValueInfo->value->dump();
+            COCL_LLVM_DUMP(localValueInfo->value);
             cout << "expression for " << localValueInfo->name + " not defined" << endl;
             throw runtime_error("expression for " + localValueInfo->name + " not defined");
         }
@@ -156,7 +158,7 @@ void AllocaClWriter::writeDeclaration(std::string indent, TypeDumper *typeDumper
             throw runtime_error("not implemented: alloca for count != 1");
         }
     } else {
-        alloca->dump();
+        COCL_LLVM_DUMP(alloca);
         throw runtime_error("dumpalloca not implemented for non pointer type");
     }
 }
@@ -244,7 +246,7 @@ void SharedClWriter::writeDeclaration(std::string indent, TypeDumper *typeDumper
             // primitiveType->dump();
         } else {
             cout << "ERROR: sharedclwriter::writedlecaraiotn, not implemneted for:" << endl;
-            value->dump();
+            COCL_LLVM_DUMP(value);
             cout << endl;
             // TODO: uncomment this line FIXME
             return;
@@ -256,7 +258,7 @@ void SharedClWriter::writeDeclaration(std::string indent, TypeDumper *typeDumper
         os << indent << "local " << typeDumper->dumpType(primitiveType) << " " << localValueInfo->name << "[" << numElements << "];\n";
     } else {
         cout << "sharedclwriter writedeclaration not implmeneted for htis type:" << endl;
-        value->dump();
+        COCL_LLVM_DUMP(value);
         cout << endl;
         throw runtime_error("not implemented: sharedclwriter for this type");
     }
@@ -311,7 +313,7 @@ void CallClWriter::writeInlineCl(std::string indent, std::ostream &os) { // writ
     }
     if(localValueInfo->toBeDeclared) {
         if(!localValueInfo->expressionValid) {
-            localValueInfo->value->dump();
+            COCL_LLVM_DUMP(localValueInfo->value);
             cout << "expression for " << localValueInfo->name + " not defined" << endl;
             throw runtime_error("expression for " + localValueInfo->name + " not defined");
         }
diff --git a/src/GlobalNames.cpp b/src/GlobalNames.cpp
index 8d1b281f..1ce2fd81 100644
--- a/src/GlobalNames.cpp
+++ b/src/GlobalNames.cpp
@@ -13,6 +13,7 @@
 // limitations under the License.
 
 #include "cocl/GlobalNames.h"
+#include "cocl/llvm_dump.h"
 
 #include "EasyCL/util/easycl_stringhelper.h"
 
@@ -35,7 +36,7 @@ std::string GlobalNames::getName(Value *value) {
     auto it = nameByValue.find(value);
     if(it == nameByValue.end()) {
         cout << "this value not found in global name map:" << endl;
-        value->dump();
+        COCL_LLVM_DUMP(value);
         cout << endl;
         throw runtime_error("value not found in global name map");
     }
@@ -99,7 +100,7 @@ std::string GlobalNames::createName(Value *value, std::string name) {
 std::string GlobalNames::getName(Type *type) {
     auto it = nameByType.find(type);
     if(it == nameByType.end()) {
-        type->dump();
+        COCL_LLVM_DUMP(type);
         cout << endl;
         throw runtime_error("type not found in name map");
     }
diff --git a/src/LocalNames.cpp b/src/LocalNames.cpp
index 646042a1..73a07484 100644
--- a/src/LocalNames.cpp
+++ b/src/LocalNames.cpp
@@ -1,4 +1,5 @@
 #include "cocl/LocalNames.h"
+#include "cocl/llvm_dump.h"
 
 #include "EasyCL/util/easycl_stringhelper.h"
 
@@ -14,7 +15,7 @@ std::string LocalNames::getName(Value *value) {
     auto it = nameByValue.find(value);
     if(it == nameByValue.end()) {
         cout << "this value not found in local names:" << endl;
-        value->dump();
+        COCL_LLVM_DUMP(value);
         cout << endl;
         throw runtime_error("value not found in local name map");
     }
@@ -67,7 +68,7 @@ bool LocalNames::hasValue(llvm::Value *value) {
 
 std::string LocalNames::createName(Value *value, std::string name) {
     if(valueByName.find(name) != valueByName.end()) {
-        valueByName[name]->dump();
+        COCL_LLVM_DUMP(valueByName[name]);
         cout << endl;
         throw runtime_error("duplicate name " + name);
     }
diff --git a/src/basicblockdumper.cpp b/src/basicblockdumper.cpp
index 05df31a9..e9fc8707 100644
--- a/src/basicblockdumper.cpp
+++ b/src/basicblockdumper.cpp
@@ -13,6 +13,7 @@
 // limitations under the License.
 
 #include "cocl/basicblockdumper.h"
+#include "cocl/llvm_dump.h"
 
 #include <iostream>
 #include <stdexcept>
@@ -98,7 +99,7 @@ bool BasicBlockDumper::runGeneration(const std::map<llvm::Function *, llvm::Type
             }
         } catch(runtime_error &e) {
             cout << "basicblockdumper.runGeneration got exception whilst processing:" << endl;
-            inst->dump();
+            COCL_LLVM_DUMP(inst);
             cout << endl;
             throw e;
         }
diff --git a/src/function_dumper.cpp b/src/function_dumper.cpp
index 2257fcfe..0d8d5a88 100644
--- a/src/function_dumper.cpp
+++ b/src/function_dumper.cpp
@@ -22,6 +22,8 @@
 
 #include "llvm/IR/Function.h"
 
+#include "cocl/llvm_dump.h"
+
 #include <sstream>
 
 using namespace std;
@@ -349,7 +351,7 @@ std::string FunctionDumper::dumpTerminator(Type **pReturnType, Instruction *term
         terminatorCl = dumpBranch(branch);
     } else {
         cout << "unhandled terminator type:";
-        terminator->dump();
+        COCL_LLVM_DUMP(terminator);
         throw runtime_error("unhandled terminator type");
     }
     return terminatorCl;
diff --git a/src/new_instruction_dumper.cpp b/src/new_instruction_dumper.cpp
index 56efe63b..0b9a2f64 100644
--- a/src/new_instruction_dumper.cpp
+++ b/src/new_instruction_dumper.cpp
@@ -25,6 +25,8 @@
 #include "cocl/readIR.h"
 #include "EasyCL/util/easycl_stringhelper.h"
 
+#include "cocl/llvm_dump.h"
+
 #include "llvm/IR/Instructions.h"
 #include "llvm/IR/Constants.h"
 #include "llvm/Transforms/Utils/Cloning.h"
@@ -125,7 +127,7 @@ LocalValueInfo *NewInstructionDumper::dumpConstant(llvm::Constant *constant) {
         return constantInfo;
     } else {
         cout << "dumpconstant, unhandled valuetype valueTy " << valueTy << endl;
-        constant->dump();
+        COCL_LLVM_DUMP(constant);
         cout << endl;
         throw runtime_error("unknown constnat type");
     }
@@ -475,7 +477,7 @@ void NewInstructionDumper::dumpGetElementPtr(cocl::LocalValueInfo *localValueInf
             }
         } else {
             cout << "type unimplemeneted in gep:" << endl;
-            currentType->dump();
+            COCL_LLVM_DUMP(currentType);
             cout << endl;
             cout << "isa pointer " << isa<PointerType>(currentType) << endl;
             cout << "isa struct " << isa<StructType>(currentType) << endl;
@@ -620,7 +622,7 @@ void NewInstructionDumper::dumpExtractValue(cocl::LocalValueInfo *localValueInfo
             }
         } else {
             cout << "NewInstructionDumper::dumpExtractValue unimplemented type" << endl;
-            currentType->dump();
+            COCL_LLVM_DUMP(currentType);
             throw runtime_error("type not implemented in extractvalue");
         }
         currentType = newType;
@@ -682,7 +684,7 @@ void NewInstructionDumper::dumpInsertValue(cocl::LocalValueInfo *localValueInfo)
                 newType = elementType;
             }
         } else {
-            currentType->dump();
+            COCL_LLVM_DUMP(currentType);
             throw runtime_error("type not implemented in insertvalue");
         }
         currentType = newType;
diff --git a/src/patch_hostside.cpp b/src/patch_hostside.cpp
index c6132359..33408e75 100644
--- a/src/patch_hostside.cpp
+++ b/src/patch_hostside.cpp
@@ -17,6 +17,7 @@
 #include "cocl/patch_hostside.h"
 
 #include "cocl/cocl_logging.h"
+#include "cocl/llvm_dump.h"
 
 #include "cocl/mutations.h"
 #include "argparsecpp/argparsecpp.h"
@@ -121,11 +122,17 @@ llvm::Instruction *PatchHostside::addSetKernelArgInst_int(llvm::Instruction *las
     } else {
         throw std::runtime_error("bitlength " + easycl::toString(bitLength) + " not implemented");
     }
+#if LLVM_VERSION_MAJOR > 4
     Function *setKernelArgInt = cast<Function>(M->getOrInsertFunction(
         mangledName,
         Type::getVoidTy(context),
-        IntegerType::get(context, bitLength),
-        NULL));
+        IntegerType::get(context, bitLength)));
+#else
+    Function *setKernelArgInt = cast<Function>(M->getOrInsertFunction(
+        mangledName,
+        Type::getVoidTy(context),
+        IntegerType::get(context, bitLength), NULL));
+#endif
     CallInst *call = CallInst::Create(setKernelArgInt, value);
     call->insertAfter(lastInst);
     lastInst = call;
@@ -145,11 +152,17 @@ llvm::Instruction *PatchHostside::addSetKernelArgInst_float(llvm::Instruction *l
         throw runtime_error("Executing functions with doubles as kernel parameters is not supported");
     }
 
+#if LLVM_VERSION_MAJOR > 4
     Function *setKernelArgFloat = cast<Function>(M->getOrInsertFunction(
         "setKernelArgFloat",
         Type::getVoidTy(context),
-        Type::getFloatTy(context),
-        NULL));
+        Type::getFloatTy(context)));
+#else
+    Function *setKernelArgFloat = cast<Function>(M->getOrInsertFunction(
+        "setKernelArgFloat",
+        Type::getVoidTy(context),
+        Type::getFloatTy(context), NULL));
+#endif
     CallInst *call = CallInst::Create(setKernelArgFloat, value);
     call->insertAfter(lastInst);
     return call;
@@ -176,12 +189,20 @@ llvm::Instruction *PatchHostside::addSetKernelArgInst_pointer(llvm::Instruction
     int allocSize = dataLayout->getTypeAllocSize(elementType);
     int32_t elementSize = allocSize;
 
+#if LLVM_VERSION_MAJOR > 4
     Function *setKernelArgGpuBuffer = cast<Function>(M->getOrInsertFunction(
         "setKernelArgGpuBuffer",
         Type::getVoidTy(context),
         PointerType::get(IntegerType::get(context, 8), 0),
-        IntegerType::get(context, 32),
-        NULL));
+        IntegerType::get(context, 32)));
+#else
+    Function *setKernelArgGpuBuffer = cast<Function>(M->getOrInsertFunction(
+        "setKernelArgGpuBuffer",
+        Type::getVoidTy(context),
+        PointerType::get(IntegerType::get(context, 8), 0),
+        IntegerType::get(context, 32), NULL));
+#endif
+
     Value *args[] = {bitcast, createInt32Constant(&context, elementSize)};
     CallInst *call = CallInst::Create(setKernelArgGpuBuffer, ArrayRef<Value *>(args));
     call->insertAfter(lastInst);
@@ -208,7 +229,7 @@ llvm::Instruction *PatchHostside::addSetKernelArgInst_pointerstruct(llvm::Instru
         cout << "ERROR: Coriander currently forbids pointers inside gpu-side structs, passed into kernels" << endl;
         cout << "If you need this, please create an issue at https://github.com/hughperkins/Coriander/issues" << endl;
         cout << "Note that it's pretty hard to do though" << endl;
-        structPointer->dump();
+        COCL_LLVM_DUMP(structPointer);
         cout << endl;
         throw std::runtime_error("ERROR: Coriander currently forbids pointers inside gpu-side structs, passed into kernels");
     }
@@ -229,7 +250,7 @@ llvm::Instruction *PatchHostside::addSetKernelArgInst_byvaluevector(llvm::Instru
     int elementSizeBytes = elementType->getPrimitiveSizeInBits() / 8;
     if(elementType->getPrimitiveSizeInBits() == 0) {
         cout << endl;
-        vectorType->dump();
+        COCL_LLVM_DUMP(vectorType);
         cout << endl;
         throw runtime_error("PatchHostside::addSetKernelArgInst_byvaluevector: not implemented for non-primitive types");
     }
@@ -246,12 +267,19 @@ llvm::Instruction *PatchHostside::addSetKernelArgInst_byvaluevector(llvm::Instru
     args[0] = bitcast;
     args[1] = createInt32Constant(&context, allocSize);
 
+#if LLVM_VERSION_MAJOR > 4
     Function *setKernelArgHostsideBuffer = cast<Function>(M->getOrInsertFunction(
         "setKernelArgHostsideBuffer",
         Type::getVoidTy(context),
         PointerType::get(IntegerType::get(context, 8), 0),
-        IntegerType::get(context, 32),
-        NULL));
+        IntegerType::get(context, 32)));
+#else
+    Function *setKernelArgHostsideBuffer = cast<Function>(M->getOrInsertFunction(
+        "setKernelArgHostsideBuffer",
+        Type::getVoidTy(context),
+        PointerType::get(IntegerType::get(context, 8), 0),
+        IntegerType::get(context, 32), NULL));
+#endif
 
     CallInst *call = CallInst::Create(setKernelArgHostsideBuffer, ArrayRef<Value *>(args));
     call->insertAfter(lastInst);
@@ -303,7 +331,12 @@ llvm::Instruction *PatchHostside::addSetKernelArgInst_byvaluestruct(llvm::Instru
 
     Value *sourceStruct = 0;
     if(structHasPointers) {
+#if LLVM_VERSION_MAJOR > 4
+        // Assume address space 0
+        Instruction *alloca = new AllocaInst(newType, 0, "newalloca");
+#else
         Instruction *alloca = new AllocaInst(newType, "newalloca");
+#endif
         alloca->insertAfter(lastInst);
         lastInst = alloca;
 
@@ -321,12 +354,19 @@ llvm::Instruction *PatchHostside::addSetKernelArgInst_byvaluestruct(llvm::Instru
     args[0] = bitcast;
     args[1] = createInt32Constant(&context, allocSize);
 
+#if LLVM_VERSION_MAJOR > 4
     Function *setKernelArgHostsideBuffer = cast<Function>(M->getOrInsertFunction(
         "setKernelArgHostsideBuffer",
         Type::getVoidTy(context),
         PointerType::get(IntegerType::get(context, 8), 0),
-        IntegerType::get(context, 32),
-        NULL));
+        IntegerType::get(context, 32)));
+#else
+    Function *setKernelArgHostsideBuffer = cast<Function>(M->getOrInsertFunction(
+        "setKernelArgHostsideBuffer",
+        Type::getVoidTy(context),
+        PointerType::get(IntegerType::get(context, 8), 0),
+        IntegerType::get(context, 32), NULL));
+#endif
 
     CallInst *call = CallInst::Create(setKernelArgHostsideBuffer, ArrayRef<Value *>(args));
     call->insertAfter(lastInst);
@@ -396,7 +436,7 @@ llvm::Instruction *PatchHostside::addSetKernelArgInst(llvm::Instruction *lastIns
         // so send it as a hostside buffer?
         lastInst = PatchHostside::addSetKernelArgInst_byvaluevector(lastInst, valueAsPointerInstr);
     } else {
-        value->dump();
+        COCL_LLVM_DUMP(value);
         outs() << "\n";
         throw std::runtime_error("kernel arg type type not implemented " + typeDumper.dumpType(value->getType()));
     }
@@ -423,7 +463,7 @@ void PatchHostside::getLaunchArgValue(GenericCallInst *inst, LaunchCallInfo *inf
         outs() << "getlaunchvalue, first operatnd of inst is not an instruction..." << "\n";
         inst->dump();
         outs() << "\n";
-        inst->getOperand(0)->dump();
+        COCL_LLVM_DUMP(inst->getOperand(0));
         outs() << "\n";
         throw runtime_error("getlaunchvalue, first operatnd of inst is not an instruction...");
     }
@@ -518,6 +558,15 @@ void PatchHostside::patchCudaLaunch(
     Instruction *llSourcecodeValue = addStringInstrExistingGlobal(M, devicellcode_stringname);
     llSourcecodeValue->insertBefore(inst->getInst());
 
+#if LLVM_VERSION_MAJOR > 4
+    Function *configureKernel = cast<Function>(F->getParent()->getOrInsertFunction(
+        "configureKernel",
+        Type::getVoidTy(context),
+        PointerType::get(IntegerType::get(context, 8), 0),
+        PointerType::get(IntegerType::get(context, 8), 0)
+        // PointerType::get(IntegerType::get(context, 8), 0),
+        ));
+#else
     Function *configureKernel = cast<Function>(F->getParent()->getOrInsertFunction(
         "configureKernel",
         Type::getVoidTy(context),
@@ -525,6 +574,8 @@ void PatchHostside::patchCudaLaunch(
         PointerType::get(IntegerType::get(context, 8), 0),
         // PointerType::get(IntegerType::get(context, 8), 0),
         NULL));
+#endif
+
     Value *args[] = {kernelNameValue, llSourcecodeValue};
     CallInst *callConfigureKernel = CallInst::Create(configureKernel, ArrayRef<Value *>(&args[0], &args[2]));
     callConfigureKernel->insertBefore(inst->getInst());
diff --git a/src/readIR.cpp b/src/readIR.cpp
index 87d36a99..33e1ac34 100644
--- a/src/readIR.cpp
+++ b/src/readIR.cpp
@@ -13,6 +13,7 @@
 // limitations under the License.
 
 #include "cocl/readIR.h"
+#include "cocl/llvm_dump.h"
 
 #include "llvm/Support/raw_os_ostream.h"
 #include "llvm/IR/Constants.h"
@@ -39,7 +40,7 @@ uint32_t ReadIR::readIntConstant_uint32(ConstantInt *constant) {
 
 std::string ReadIR::getName(StructType *type) {
     if(!type->hasName()) {
-        type->dump();
+        COCL_LLVM_DUMP(type);
         throw runtime_error("type doesnt have name");
     }
     return type->getName();
@@ -47,7 +48,8 @@ std::string ReadIR::getName(StructType *type) {
 
 std::string ReadIR::getName(Function *type) {
     if(!type->hasName()) {
-        type->dump();
+        type->print(llvm::dbgs());
+        //COCL_LLVM_DUMP(type);
         throw runtime_error("function doesnt have name");
     }
     return type->getName();
@@ -55,7 +57,7 @@ std::string ReadIR::getName(Function *type) {
 
 std::string ReadIR::getName(Value *value) {
     if(!value->hasName()) {
-        value->dump();
+        COCL_LLVM_DUMP(value);
         throw runtime_error("value doesnt have name");
     }
     return value->getName();
diff --git a/src/struct_clone.cpp b/src/struct_clone.cpp
index 7fea9efd..8a70b67c 100644
--- a/src/struct_clone.cpp
+++ b/src/struct_clone.cpp
@@ -18,6 +18,8 @@
 #include "EasyCL/util/easycl_stringhelper.h"
 #include "cocl/mutations.h"
 
+#include "cocl/llvm_dump.h"
+
 #include "llvm/Support/raw_os_ostream.h"
 #include "llvm/IR/Constants.h"
 #include "llvm/IR/Module.h"
@@ -168,7 +170,7 @@ std::string StructCloner::writeClCopyToDevicesideStruct(llvm::StructType *ptrful
             dstidx++;
         } else {
             outs() << "unhandled type\n";
-            childType->dump();
+            COCL_LLVM_DUMP(childType);
             outs() << "\n";
             throw runtime_error("unhandled type");
             srcidx++;
@@ -251,7 +253,7 @@ llvm::Instruction *StructCloner::writeHostsideIrCopyToMarshallingStruct(
                 }
             } else {
                 outs() << "unhandled type:\n";
-                childType->dump();
+                COCL_LLVM_DUMP(childType);
                 outs() << "\n";
                 throw runtime_error("structcloner unhandled type");
             }
@@ -260,7 +262,7 @@ llvm::Instruction *StructCloner::writeHostsideIrCopyToMarshallingStruct(
         }
     } else {
         outs() << "skipping type:\n";
-        structType->dump();
+        COCL_LLVM_DUMP(structType);
         outs() << "\n";
     }
     return lastInst;
@@ -297,7 +299,7 @@ void StructCloner::walkType(Module *M, StructInfo *structInfo, int level, int of
     } else if(isa<IntegerType>(type)) {
     } else if(type->getPrimitiveSizeInBits() != 0) {
     } else {
-        type->dump();
+        COCL_LLVM_DUMP(type);
         throw runtime_error("walktype type not handled");
     }
 }
diff --git a/src/type_dumper.cpp b/src/type_dumper.cpp
index 3ebe5094..67b886a9 100644
--- a/src/type_dumper.cpp
+++ b/src/type_dumper.cpp
@@ -17,6 +17,8 @@
 #include "cocl/mutations.h"
 #include "EasyCL/util/easycl_stringhelper.h"
 
+#include "cocl/llvm_dump.h"
+
 #include "llvm/IR/Instructions.h"
 #include "llvm/Support/raw_ostream.h"
 
@@ -187,7 +189,7 @@ std::string TypeDumper::addStructToGlobalNames(StructType *type) {
             }
         }
     } else {
-        type->dump();
+        COCL_LLVM_DUMP(type);
         throw runtime_error("not implemented: anonymous struct types");
     }
 }
@@ -217,7 +219,7 @@ std::string TypeDumper::dumpVectorType(VectorType *vectorType, bool decayArraysT
     Type *elementType = vectorType->getElementType();
     if(elementType->getPrimitiveSizeInBits() == 0) {
         cout << endl;
-        vectorType->dump();
+        COCL_LLVM_DUMP(vectorType);
         cout << endl;
         throw runtime_error("TypeDumper::dumpVectorType: not implemented for non-primitive types");
     }
@@ -286,7 +288,7 @@ std::string TypeDumper::dumpType(Type *type, bool decayArraysToPointer) {
         default:
             outs() << "type id " << typeID << "\n";
             cout << endl;
-            type->dump();
+            COCL_LLVM_DUMP(type);
             cout << endl;
             throw runtime_error("TypeDumper::dumpType(...): unrecognized type");
     }
diff --git a/test/gtest/test_block_dumper.cpp b/test/gtest/test_block_dumper.cpp
index 0239c650..f1b364c4 100644
--- a/test/gtest/test_block_dumper.cpp
+++ b/test/gtest/test_block_dumper.cpp
@@ -18,6 +18,8 @@
 #include "cocl/GlobalNames.h"
 #include "cocl/LocalNames.h"
 
+#include "cocl/llvm_dump.h"
+
 #include "llvm/IRReader/IRReader.h"
 #include "llvm/IR/Module.h"
 #include "llvm/Support/MemoryBuffer.h"
@@ -97,7 +99,7 @@ class LocalWrapper {
         for(auto it=F->arg_begin(); it != F->arg_end(); it++) {
             Argument *arg = &*it;
             // sring name = localNames.getOrCreateName(arg, arg->getName().str());
-            arg->dump();
+            COCL_LLVM_DUMP(arg);
             LocalValueInfo *localValueInfo = LocalValueInfo::getOrCreate(
                 &localNames, &localValueInfos, arg, arg->getName().str());
             localValueInfo->setExpression(localValueInfo->name);
@@ -226,7 +228,7 @@ TEST(test_block_dumper, basic2) {
     }
     ASSERT_EQ(1u, blockDumper->neededFunctions.size());
     Function *neededFunction = *blockDumper->neededFunctions.begin();
-    neededFunction->dump();
+    COCL_LLVM_DUMP(neededFunction);
     cout << endl;
     cout << neededFunction->getName().str() << endl;
     ASSERT_EQ("someFunc_gp", neededFunction->getName().str());