diff --git a/lib/ReplaceLLVMIntrinsicsPass.cpp b/lib/ReplaceLLVMIntrinsicsPass.cpp index 651c04c12..a89405664 100644 --- a/lib/ReplaceLLVMIntrinsicsPass.cpp +++ b/lib/ReplaceLLVMIntrinsicsPass.cpp @@ -12,15 +12,15 @@ // See the License for the specific language governing permissions and // limitations under the License. -#include -#include -#include -#include -#include -#include -#include +#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 +#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 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(U)) { + CI->eraseFromParent(); + } + } + F->eraseFromParent(); + } + + return Changed; +} diff --git a/test/LLVMIntrinsics/lifetime_start.cl b/test/LLVMIntrinsics/lifetime_start.cl new file mode 100644 index 000000000..0236cf586 --- /dev/null +++ b/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]; + } +}