Remove @llvm.lifetime.* intrinsics

The corresponding OpLifetimeStart and OpLifetimeEnd instructions
require Kernel capability, so we can't use them in Vulkan shaders
anyway.  Just remove them.

Fixes google#142

lib/ReplaceLLVMIntrinsicsPass.cpp

@@ -12,15 +12,15 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include <llvm/IR/Constants.h>
-#include <llvm/IR/IRBuilder.h>
-#include <llvm/IR/Instructions.h>
-#include <llvm/IR/Module.h>
-#include <llvm/Pass.h>
-#include <llvm/Support/raw_ostream.h>
-#include <llvm/Transforms/Utils/Cloning.h>
+#include "llvm/IR/Constants.h"
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/Instructions.h"
+#include "llvm/IR/Module.h"
+#include "llvm/Pass.h"
+#include "llvm/Support/raw_ostream.h"
+#include "llvm/Transforms/Utils/Cloning.h"
 
-#include <spirv/1.0/spirv.hpp>
+#include "spirv/1.0/spirv.hpp"
 
 using namespace llvm;
 
@@ -34,6 +34,7 @@ struct ReplaceLLVMIntrinsicsPass final : public ModulePass {
   bool runOnModule(Module &M) override;
   bool replaceMemset(Module &M);
   bool replaceMemcpy(Module &M);
+  bool removeLifetimeDeclarations(Module &M);
 };
 }
 
@@ -50,6 +51,9 @@ ModulePass *createReplaceLLVMIntrinsicsPass() {
 bool ReplaceLLVMIntrinsicsPass::runOnModule(Module &M) {
   bool Changed = false;
 
+  // Remove lifetime annotations first.  They coulud be using memset
+  // and memcpy calls.
+  Changed |= removeLifetimeDeclarations(M);
   Changed |= replaceMemset(M);
   Changed |= replaceMemcpy(M);
 
@@ -332,3 +336,29 @@ bool ReplaceLLVMIntrinsicsPass::replaceMemcpy(Module &M) {
 
   return Changed;
 }
+
+bool ReplaceLLVMIntrinsicsPass::removeLifetimeDeclarations(Module &M) {
+  // SPIR-V OpLifetimeStart and OpLifetimeEnd require Kernel capability.
+  // Vulkan doesn't support that, so remove all lifteime bounds declarations.
+
+  bool Changed = false;
+
+  SmallVector<Function *, 2> WorkList;
+  for (auto &F : M) {
+    if (F.getName().startswith("llvm.lifetime.")) {
+      WorkList.push_back(&F);
+    }
+  }
+
+  for (auto *F : WorkList) {
+    Changed = true;
+    for (auto U : F->users()) {
+      if (auto *CI = dyn_cast<CallInst>(U)) {
+        CI->eraseFromParent();
+      }
+    }
+    F->eraseFromParent();
+  }
+
+  return Changed;
+}

test/LLVMIntrinsics/lifetime_start.cl

@@ -0,0 +1,34 @@
+// Remove @llvm.lifetime.start.*
+// Fixes https://github.com/google/clspv/issues/142
+
+
+// RUN: clspv %s -S -o %t.spvasm -cluster-pod-kernel-args
+// RUN: FileCheck %s < %t.spvasm
+// RUN: clspv %s -o %t.spv -cluster-pod-kernel-args
+// RUN: spirv-dis -o %t2.spvasm %t.spv
+// RUN: FileCheck %s < %t2.spvasm
+// RUN: spirv-val --target-env vulkan1.0 %t.spv
+
+// Just check that the compiler works at all.
+
+// CHECK: ; SPIR-V
+// CHECK: ; Version: 1.0
+// CHECK: OpEntryPoint
+// CHECK: OpFunctionEnd
+
+#define CHUNK_SIZE 32
+
+kernel void cfc(global const int *in, global int *out, int limit) {
+  size_t x = get_global_id(0);
+
+  int temp[CHUNK_SIZE];
+  for (int i = 0; i < CHUNK_SIZE; ++i) {
+    temp[i] = in[i];
+  }
+
+  if (x < limit) {
+    out[x] = x;
+  } else {
+    out[x] = temp[x % CHUNK_SIZE];
+  }
+}