Merge from 'sycl' to 'sycl-web' (#1)

Ignored the revert 1a8a660 (cherry-pick of 0c28a7c). It was done
for sycl-web => sycl pulldown. The revert has already happened once in
sycl-web (as 0c28a7c), followed by a re-commit 4c4093e.

The conflict was because 1a8a660 was trying to apply on top of
the re-commit 4c4093e.

  CONFLICT (modify/delete): llvm/test/Transforms/InstCombine/fpclass.ll deleted in 9f41bf6 and modified in HEAD. Version HEAD of llvm/test/Transforms/InstCombine/fpclass.ll left in tree.
  CONFLICT (modify/delete): llvm/test/CodeGen/PowerPC/ppc-fpclass.ll deleted in 9f41bf6 and modified in HEAD. Version HEAD of llvm/test/CodeGen/PowerPC/ppc-fpclass.ll left in tree.
  CONFLICT (content): Merge conflict in llvm/lib/Transforms/InstCombine/InstCombineCalls.cpp
  CONFLICT (content): Merge conflict in llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp

Author: abhinavgaba

clang/test/Driver/sycl-offload-intelfpga.cpp

@@ -13,13 +13,13 @@
 /// -fintelfpga implies -g and -MMD
 // RUN:   %clangxx -### -target x86_64-unknown-linux-gnu -fsycl -fintelfpga -Xshardware %s 2>&1 \
 // RUN:   | FileCheck -check-prefix=CHK-TOOLS-INTELFPGA %s
-// CHK-TOOLS-INTELFPGA: clang{{.*}} "-debug-info-kind=limited" {{.*}} "-dependency-file"
+// CHK-TOOLS-INTELFPGA: clang{{.*}} "-debug-info-kind=constructor" {{.*}} "-dependency-file"
 // CHK-TOOLS-INTELFPGA: aoc{{.*}} "-dep-files={{.*}}"
 
 /// -fintelfpga implies -g but -g0 should override
 // RUN:   %clangxx -### -target x86_64-unknown-linux-gnu -g0 -fsycl -fintelfpga %s 2>&1 \
 // RUN:   | FileCheck -check-prefix=CHK-TOOLS-INTELFPGA-G0 %s
-// CHK-TOOLS-INTELFPGA-G0-NOT: clang{{.*}} "-debug-info-kind=limited"
+// CHK-TOOLS-INTELFPGA-G0-NOT: clang{{.*}} "-debug-info-kind=constructor"
 
 /// FPGA target implies -fsycl-disable-range-rounding
 // RUN:   %clangxx -### -target x86_64-unknown-linux-gnu -fsycl -fintelfpga %s 2>&1 \

clang/test/Driver/sycl-offload-nvptx.cpp

@@ -12,7 +12,7 @@
 // CHK-ACTIONS: "-cc1" "-triple" "nvptx64-nvidia-cuda" "-aux-triple" "x86_64-unknown-linux-gnu"{{.*}} "-fsycl-is-device"{{.*}} "-Wno-sycl-strict"{{.*}} "-sycl-std=2020" {{.*}} "-emit-llvm-bc" {{.*}} "-internal-isystem" "{{.*}}bin{{[/\\]+}}..{{[/\\]+}}include{{[/\\]+}}sycl"{{.*}} "-mlink-builtin-bitcode" "{{.*}}libspirv.bc"{{.*}} "-mlink-builtin-bitcode" "{{.*}}libdevice{{.*}}.10.bc"{{.*}} "-target-feature" "+ptx42"{{.*}} "-target-sdk-version=[[CUDA_VERSION:[0-9.]+]]"{{.*}} "-target-cpu" "sm_50"{{.*}} "-std=c++11"{{.*}}
 // CHK-ACTIONS: sycl-post-link{{.*}} "-split=auto"
 // CHK-ACTIONS: file-table-tform" "-extract=Code" "-drop_titles"
-// CHK-ACTIONS: llvm-foreach" {{.*}} "--" "{{.*}}clang-13"
+// CHK-ACTIONS: llvm-foreach" {{.*}} "--" "{{.*}}clang-{{[0-9]+}}"
 // CHK-ACTIONS: llvm-foreach" {{.*}} "--" "{{.*}}ptxas"
 // CHK-ACTIONS: llvm-foreach" {{.*}} "--" "{{.*}}fatbinary"
 // CHK-ACTIONS: file-table-tform" "-replace=Code,Code"

llvm-spirv/.github/workflows/check-code-style.yml

@@ -24,7 +24,7 @@ env:
   # We need compile command database in order to perform clang-tidy check. So,
   # in order to perform configure step we need to setup llvm-dev package. This
   # env variable used to specify desired version of it
-  LLVM_VERSION: 13
+  LLVM_VERSION: 14
 
 jobs:
   clang-format-and-tidy:

llvm-spirv/.github/workflows/check-in-tree-build.yml

@@ -32,7 +32,7 @@ on:
     - cron: 0 0 * * *
 
 env:
-  LLVM_VERSION: 13
+  LLVM_VERSION: 14
 
 jobs:
   build_and_test_linux:
@@ -70,6 +70,16 @@ jobs:
         uses: actions/checkout@v2
         with:
           path: llvm-project/llvm/projects/SPIRV-LLVM-Translator
+      - name:  Get tag for SPIR-V Headers
+        id: spirv-headers-tag
+        run: |
+          echo "spirv_headers_tag=$(cat llvm-project/llvm/projects/SPIRV-LLVM-Translator/spirv-headers-tag.conf)" >> $GITHUB_ENV
+      - name:  Checkout SPIR-V Headers
+        uses: actions/checkout@v2
+        with:
+          repository: KhronosGroup/SPIRV-Headers
+          ref: ${{ env.spirv_headers_tag }}
+          path: llvm-project/llvm/projects/SPIRV-Headers
       - name: Configure
         run: |
           mkdir build && cd build
@@ -115,6 +125,16 @@ jobs:
         uses: actions/checkout@v2
         with:
           path: llvm-project\\llvm\\projects\\SPIRV-LLVM-Translator
+      - name:  Get tag for SPIR-V Headers
+        id: spirv-headers-tag
+        run: |
+          echo "spirv_headers_tag=$(type llvm-project\\llvm\\projects\\SPIRV-LLVM-Translator\\spirv-headers-tag.conf)" >> $GITHUB_ENV
+      - name:  Checkout SPIR-V Headers
+        uses: actions/checkout@v2
+        with:
+          repository: KhronosGroup/SPIRV-Headers
+          ref: ${{ env.spirv_headers_tag }}
+          path: llvm-project\\llvm\\projects\\SPIRV-Headers
       - name: Configure
         shell: bash
         run: |
@@ -157,6 +177,16 @@ jobs:
         uses: actions/checkout@v2
         with:
           path: llvm-project/llvm/projects/SPIRV-LLVM-Translator
+      - name:  Get tag for SPIR-V Headers
+        id: spirv-headers-tag
+        run: |
+          echo "spirv_headers_tag=$(cat llvm-project/llvm/projects/SPIRV-LLVM-Translator/spirv-headers-tag.conf)" >> $GITHUB_ENV
+      - name:  Checkout SPIR-V Headers
+        uses: actions/checkout@v2
+        with:
+          repository: KhronosGroup/SPIRV-Headers
+          ref: ${{ env.spirv_headers_tag }}
+          path: llvm-project/llvm/projects/SPIRV-Headers
       - name: Configure
         run: |
           mkdir build && cd build

llvm-spirv/.github/workflows/check-out-of-tree-build.yml

@@ -29,7 +29,7 @@ on:
     - cron: 0 0 * * *
 
 env:
-  LLVM_VERSION: 13
+  LLVM_VERSION: 14
 
 jobs:
   build_and_test:
@@ -58,14 +58,28 @@ jobs:
           sudo update-alternatives --install /usr/bin/clang clang /usr/bin/clang-${{ env.LLVM_VERSION }} 1000
       - name:  Checkout the translator sources
         uses: actions/checkout@v2
+        with:
+          path: SPIRV-LLVM-Translator
+      - name:  Get tag for SPIR-V Headers
+        id: spirv-headers-tag
+        run: |
+          echo "spirv_headers_tag=$(cat SPIRV-LLVM-Translator/spirv-headers-tag.conf)" >> $GITHUB_ENV
+      - name:  Checkout SPIR-V Headers
+        uses: actions/checkout@v2
+        with:
+          repository: KhronosGroup/SPIRV-Headers
+          ref: ${{ env.spirv_headers_tag }}
+          path: SPIRV-Headers
       - name: Configure
         run: |
           mkdir build && cd build
-          cmake ${{ github.workspace }} \
+          cmake ${{ github.workspace }}/SPIRV-LLVM-Translator \
             -DCMAKE_BUILD_TYPE=${{ matrix.build_type }} \
             -DCMAKE_CXX_FLAGS="-Werror" \
             -DLLVM_INCLUDE_TESTS=ON \
             -DLLVM_EXTERNAL_LIT="/usr/lib/llvm-${{ env.LLVM_VERSION }}/build/utils/lit/lit.py" \
+            -DLLVM_EXTERNAL_PROJECTS="SPIRV-Headers" \
+            -DLLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR=${{ github.workspace }}/SPIRV-Headers \
             -G "Unix Makefiles"
       - name: Build
         run: |

llvm-spirv/.travis.yml

@@ -42,7 +42,7 @@ compiler:
 env:
   global:
     - MAKEFLAGS="-j2"
-    - LLVM_VERSION=13
+    - LLVM_VERSION=14
   matrix:
     - BUILD_TYPE=Release BUILD_EXTERNAL=1 SHARED_LIBS=ON  MAKE_TARGETS=""           MAKE_TEST_TARGET="test"
     - BUILD_TYPE=Debug   BUILD_EXTERNAL=1 SHARED_LIBS=ON  MAKE_TARGETS=""           MAKE_TEST_TARGET="test"

llvm-spirv/CMakeLists.txt

@@ -1,8 +1,10 @@
-cmake_minimum_required(VERSION 3.3)
+cmake_minimum_required(VERSION 3.13.4)
 
-set (BASE_LLVM_VERSION 13.0.0)
+set (BASE_LLVM_VERSION 14.0.0)
 set(LLVM_SPIRV_VERSION ${BASE_LLVM_VERSION}.0)
 
+include(FetchContent)
+
 option(LLVM_SPIRV_INCLUDE_TESTS
   "Generate build targets for the llvm-spirv lit tests."
   ${LLVM_INCLUDE_TESTS})
@@ -14,6 +16,42 @@ if (NOT DEFINED LLVM_SPIRV_BUILD_EXTERNAL)
   endif(CMAKE_SOURCE_DIR STREQUAL CMAKE_CURRENT_SOURCE_DIR)
 endif (NOT DEFINED LLVM_SPIRV_BUILD_EXTERNAL)
 
+# Download spirv.hpp from the official SPIRV-Headers repository.
+# One can skip this step by manually setting
+# LLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR path.
+if(NOT DEFINED LLVM_TOOL_SPIRV_HEADERS_BUILD AND
+   NOT DEFINED LLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR)
+  set(LLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR
+    "${CMAKE_CURRENT_BINARY_DIR}/SPIRV-Headers")
+  message(STATUS "SPIR-V Headers location is not specified. Will try to download
+          spirv.hpp from https://github.com/KhronosGroup/SPIRV-Headers into
+          ${LLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR}")
+  file(READ spirv-headers-tag.conf SPIRV_HEADERS_TAG)
+  # Strip the potential trailing newline from tag
+  string(STRIP "${SPIRV_HEADERS_TAG}" SPIRV_HEADERS_TAG)
+  FetchContent_Declare(spirv-headers
+    GIT_REPOSITORY    https://github.com/KhronosGroup/SPIRV-Headers.git
+    GIT_TAG           ${SPIRV_HEADERS_TAG}
+    SOURCE_DIR ${LLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR}
+  )
+  FetchContent_MakeAvailable(spirv-headers)
+else()
+  if(NOT DEFINED LLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR)
+    # This means LLVM_TOOL_SPIRV_HEADERS_BUILD is defined, therefore
+    # SPIRV-Headers exist as a subproject.
+    set(LLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR
+      "${CMAKE_SOURCE_DIR}/projects/SPIRV-Headers")
+    if(NOT EXISTS ${LLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR})
+      message(FATAL_ERROR "No location specified for SPIRV-Headers.
+              Try setting the LLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR
+              path or put the project into the llvm/projects folder
+              under the name 'SPIRV-Headers'")
+    endif()
+  endif()
+  message(STATUS "Using SPIR-V Headers from
+          ${LLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR}")
+endif()
+
 if(LLVM_SPIRV_BUILD_EXTERNAL)
   project(LLVM_SPIRV
     VERSION

llvm-spirv/CONTRIBUTING.md

@@ -36,6 +36,29 @@ Coding Standards]. Compliance of your code is checked automatically using
 Travis CI. See [clang-format] and [clang-tidy] configs for more details about
 coding standards.
 
+## How to add an extension
+
+First of all please make sure you have added a link to the
+specification for the extension in your PR. Then to add definitions of
+new Op Codes you shall modify [spirv.hpp], which is an external
+dependency for this project. To do so, you should add new definitions
+to [json grammar file], rebuild the header following the
+[instructions] in [SPIR-V Headers repository] and push your changes
+for review, i.e. make a PR. Once the PR is merged, a new spirv.hpp
+will have to be downloaded during build of the translator; make sure
+to update the hash for SPIRV-Headers in [spirv-headers-tag.conf]
+so that tokens from your extension can be visible to the translator
+build.
+
+It's highly recommended to add the definitions to [SPIR-V Headers repository]
+first, but if you don't want to bring it there yet, you can define new Op Codes
+in the [internal SPIR-V header file].
+
+For local testing you can copy your spirv.hpp variant to
+`<PATH_TO_SPIRV_HEADERS>/include/spirv/unified1` and/or modify it
+there. See [README.md](README.md#configuring-spir-v-headers) for build
+instructions that should be employed with such modifications.
+
 ### Conditions to merge a PR
 
 In order to get your PR merged, the following conditions must be met:
@@ -79,5 +102,10 @@ it can be merged:
 [LLVM Coding Standards]: https://llvm.org/docs/CodingStandards.html
 [clang-format]: [.clang-format]
 [clang-tidy]: [.clang-tidy]
+[spirv.hpp]: https://github.com/KhronosGroup/SPIRV-Headers/blob/master/include/spirv/unified1/spirv.hpp
+[json grammar file]: https://github.com/KhronosGroup/SPIRV-Headers/blob/master/include/spirv/unified1/spirv.core.grammar.json
+[instructions]: https://github.com/KhronosGroup/SPIRV-Headers#generating-headers-from-the-json-grammar-for-the-spir-v-core-instruction-set
+[SPIR-V Headers repository]: https://github.com/KhronosGroup/SPIRV-Headers
+[internal SPIR-V header file]: https://github.com/KhronosGroup/SPIRV-LLVM-Translator/blob/master/lib/SPIRV/libSPIRV/spirv_internal.hpp
 [Contributor License Agreement]: https://cla-assistant.io/KhronosGroup/SPIRV-LLVM-Translator
 [Travis CI testing]: https://travis-ci.org/KhronosGroup/SPIRV-LLVM-Translator

llvm-spirv/README.md

@@ -29,7 +29,7 @@ The translator can be built with the latest(nightly) package of LLVM. For Ubuntu
 wget -O - https://apt.llvm.org/llvm-snapshot.gpg.key | sudo apt-key add -
 sudo add-apt-repository "deb http://apt.llvm.org/xenial/ llvm-toolchain-xenial main"
 sudo apt-get update
-sudo apt-get install llvm-13-dev llvm-13-tools clang-13 libclang-13-dev
+sudo apt-get install llvm-14-dev llvm-14-tools clang-14 libclang-14-dev
 ```
 The installed version of LLVM will be used by default for out-of-tree build of the translator.
 ```
@@ -93,6 +93,26 @@ Building clang from sources takes time and resources and it can be avoided:
   moment, see [KhronosGroup/SPIRV-LLVM-Translator#477](https://github.com/KhronosGroup/SPIRV-LLVM-Translator/issues/477)
   to track progress, discuss and contribute.
 
+## Configuring SPIR-V Headers
+
+The translator build is dependent on the official Khronos header file
+`spirv.hpp` that maps SPIR-V extensions, decorations, instructions,
+etc. onto numeric tokens. The official header version is available at
+[KhronosGroup/SPIRV-Headers](https://github.com/KhronosGroup/SPIRV-Headers).
+There are several options for accessing the header file:
+- By default, the header file repository will be downloaded from
+  Khronos Group GitHub and put into `<build_dir>/SPIRV-Headers`.
+- If you are building the translator in-tree, you can manually
+  download the SPIR-V Headers repo into `llvm/projects` - this
+  location will be automatically picked up by the LLVM build
+  scripts. Make sure the folder retains its default naming in
+  that of `SPIRV-Headers`.
+- Any build type can also use an external installation of SPIR-V
+  Headers - if you have the headers downloaded somewhere in your
+  system and want to use that version, simply extend your CMake
+  command with `-DLLVM_EXTERNAL_PROJECTS="SPIRV-Headers"
+  -DLLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR=</path/to/headers_dir>`.
+
 ## Test instructions
 
 All tests related to the translator are placed in the [test](test) directory. A number of the tests require spirv-as (part of SPIR-V Tools) to run, but the remainder of the tests can still be run without this. Optionally the tests can make use of spirv-val (part of SPIRV-Tools) in order to validate the generated SPIR-V against the official SPIR-V specification.
@@ -105,7 +125,7 @@ make test
 ```
 This requires that the `-DLLVM_SPIRV_INCLUDE_TESTS=ON` argument is
 passed to CMake during the build step. Additionally,
-`-DLLVM_EXTERNAL_LIT="/usr/lib/llvm-13/build/utils/lit/lit.py"` is
+`-DLLVM_EXTERNAL_LIT="/usr/lib/llvm-14/build/utils/lit/lit.py"` is
 needed when building with a pre-installed version of LLVM.
 
 The translator test suite can be disabled by passing

llvm-spirv/lib/SPIRV/CMakeLists.txt

@@ -51,6 +51,10 @@ target_include_directories(LLVMSPIRVLib
   PRIVATE
     ${LLVM_INCLUDE_DIRS}
     ${LLVM_SPIRV_INCLUDE_DIRS}
+    # TODO: Consider using SPIRV-Headers' as a header-only INTERFACE
+    # instead. Right now this runs into exporting issues with
+    # the LLVM in-tree builds.
+    ${LLVM_EXTERNAL_SPIRV_HEADERS_SOURCE_DIR}/include
     ${CMAKE_CURRENT_SOURCE_DIR}
     ${CMAKE_CURRENT_SOURCE_DIR}/libSPIRV
     ${CMAKE_CURRENT_SOURCE_DIR}/Mangler

llvm-spirv/lib/SPIRV/OCLToSPIRV.cpp

@@ -108,9 +108,6 @@ class OCLToSPIRVBase : public InstVisitor<OCLToSPIRVBase> {
   /// Transform OCL builtin function to SPIR-V builtin function.
   void transBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info);
 
-  /// Transform OCL work item builtin functions to SPIR-V builtin variables.
-  void transWorkItemBuiltinsToVariables();
-
   /// Transform atomic_work_item_fence/mem_fence to __spirv_MemoryBarrier.
   /// func(flag, order, scope) =>
   ///   __spirv_MemoryBarrier(map(scope), map(flag)|map(order))
@@ -367,8 +364,6 @@ bool OCLToSPIRVBase::runOCLToSPIRV(Module &Module) {
 
   LLVM_DEBUG(dbgs() << "Enter OCLToSPIRV:\n");
 
-  transWorkItemBuiltinsToVariables();
-
   visit(*M);
 
   for (auto &I : ValuesToDelete)
@@ -1061,6 +1056,7 @@ void OCLToSPIRVBase::transBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info) {
   AttributeList Attrs = CI->getCalledFunction()->getAttributes();
   Op OC = OpNop;
   unsigned ExtOp = ~0U;
+  SPIRVBuiltinVariableKind BVKind = BuiltInMax;
   if (StringRef(Info.UniqName).startswith(kSPIRVName::Prefix))
     return;
   if (OCLSPIRVBuiltinMap::find(Info.UniqName, &OC)) {
@@ -1081,7 +1077,11 @@ void OCLToSPIRVBase::transBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info) {
     }
   } else if ((ExtOp = getExtOp(Info.MangledName, Info.UniqName)) != ~0U)
     Info.UniqName = getSPIRVExtFuncName(SPIRVEIS_OpenCL, ExtOp);
-  else
+  else if (SPIRSPIRVBuiltinVariableMap::find(Info.UniqName, &BVKind)) {
+    // Map OCL work item builtins to SPV-IR work item builtins.
+    // e.g. get_global_id() --> __spirv_BuiltinGlobalInvocationId()
+    Info.UniqName = getSPIRVFuncName(BVKind);
+  } else
     return;
   if (!Info.RetTy)
     mutateCallInstSPIRV(
@@ -1273,60 +1273,6 @@ void OCLToSPIRVBase::visitCallBuiltinSimple(CallInst *CI, StringRef MangledName,
   transBuiltin(CI, Info);
 }
 
-/// Translates OCL work-item builtin functions to SPIRV builtin variables.
-/// Function like get_global_id(i) -> x = load GlobalInvocationId; extract x, i
-/// Function like get_work_dim() -> load WorkDim
-void OCLToSPIRVBase::transWorkItemBuiltinsToVariables() {
-  LLVM_DEBUG(dbgs() << "Enter transWorkItemBuiltinsToVariables\n");
-  std::vector<Function *> WorkList;
-  for (auto &I : *M) {
-    StringRef DemangledName;
-    if (!oclIsBuiltin(I.getName(), DemangledName))
-      continue;
-    LLVM_DEBUG(dbgs() << "Function demangled name: " << DemangledName << '\n');
-    std::string BuiltinVarName;
-    SPIRVBuiltinVariableKind BVKind;
-    if (!SPIRSPIRVBuiltinVariableMap::find(DemangledName.str(), &BVKind))
-      continue;
-    BuiltinVarName =
-        std::string(kSPIRVName::Prefix) + SPIRVBuiltInNameMap::map(BVKind);
-    LLVM_DEBUG(dbgs() << "builtin variable name: " << BuiltinVarName << '\n');
-    bool IsVec = I.getFunctionType()->getNumParams() > 0;
-    Type *GVType =
-        IsVec ? FixedVectorType::get(I.getReturnType(), 3) : I.getReturnType();
-    auto BV = new GlobalVariable(*M, GVType, true, GlobalValue::ExternalLinkage,
-                                 nullptr, BuiltinVarName, 0,
-                                 GlobalVariable::NotThreadLocal, SPIRAS_Input);
-    std::vector<Instruction *> InstList;
-    for (auto UI = I.user_begin(), UE = I.user_end(); UI != UE; ++UI) {
-      auto CI = dyn_cast<CallInst>(*UI);
-      assert(CI && "invalid instruction");
-      const DebugLoc &DLoc = CI->getDebugLoc();
-      Instruction *NewValue = new LoadInst(GVType, BV, "", CI);
-      if (DLoc)
-        NewValue->setDebugLoc(DLoc);
-      LLVM_DEBUG(dbgs() << "Transform: " << *CI << " => " << *NewValue << '\n');
-      if (IsVec) {
-        NewValue =
-            ExtractElementInst::Create(NewValue, CI->getArgOperand(0), "", CI);
-        if (DLoc)
-          NewValue->setDebugLoc(DLoc);
-        LLVM_DEBUG(dbgs() << *NewValue << '\n');
-      }
-      NewValue->takeName(CI);
-      CI->replaceAllUsesWith(NewValue);
-      InstList.push_back(CI);
-    }
-    for (auto &Inst : InstList) {
-      Inst->eraseFromParent();
-    }
-    WorkList.push_back(&I);
-  }
-  for (auto &I : WorkList) {
-    I->eraseFromParent();
-  }
-}
-
 void OCLToSPIRVBase::visitCallReadWriteImage(CallInst *CI,
                                              StringRef DemangledName) {
   OCLBuiltinTransInfo Info;