intel · nbpatel · Aug 2, 2022 · Jul 15, 2022 · Jul 15, 2022 · Jul 19, 2022
diff --git a/.github/workflows/build.yml b/.github/workflows/build.yml
@@ -49,7 +49,7 @@ jobs:
 
       - name: Install CMake and Ninja
         run: |
-          conda install cmake ninja
+          conda install cmake ninja conda-forge::lit
 
       - name: Checkout repo
         uses: actions/checkout@v2
@@ -76,7 +76,7 @@ jobs:
         id: cache-llvm-mlir
         uses: actions/cache@v3
         env:
-          CACHE_NUMBER: 0  # Increase to reset cache
+          CACHE_NUMBER: 1  # Increase to reset cache
         with:
           path: |
             /home/runner/work/llvm-mlir/_mlir_install/**
@@ -152,12 +152,17 @@ jobs:
             -DCMAKE_INSTALL_PREFIX=/home/runner/work/llvm-mlir/_mlir_install || exit 1
           cmake --build . -j ${np} || exit 1
           cmake --install . || exit 1
+          cp bin/FileCheck /home/runner/work/llvm-mlir/_mlir_install/bin/
+          cp bin/count /home/runner/work/llvm-mlir/_mlir_install/bin/
+          cp bin/not /home/runner/work/llvm-mlir/_mlir_install/bin/
           popd
 
       - name: Build IMEX
         shell: bash -l {0}
 
         run: |
+          external_lit=`which lit`
+          echo ${external_lit}
           # make the llvm-mlir working dir
           mkdir -p /home/runner/work/tmpdir
           export TBB_PATH=/home/runner/work/tbb/oneapi-tbb-${TBB_VER}
@@ -168,7 +173,16 @@ jobs:
             --llvm-install /home/runner/work/llvm-mlir/_mlir_install           \
             --imex-enable-igpu                                                 \
             --imex-enable-tests                                                \
-            --imex-tbb-dir /${TBB_PATH}/lib/cmake/tbb
+            --imex-tbb-dir /${TBB_PATH}/lib/cmake/tbb                          \
+            --external-lit ${external_lit}
+
+
+      - name: Run IMEX core test cases
+        shell: bash -l {0}
+
+        run: |
+          cd _build || exit 1
+          cmake --build . --target check-l0-runner || exit 1
 
       - name: Build Python frontend
         shell: bash -l {0}

diff --git a/mlir/test/Conversion/gpu-to-spirv.mlir b/mlir/test/Conversion/gpu-to-spirv.mlir
@@ -0,0 +1,58 @@
+// RUN: level-zero-opt --gpu-to-spirv %s | FileCheck %s
+
+module attributes {gpu.container_module, spv.target_env = #spv.target_env<#spv.vce<v1.0, [Addresses, Float16Buffer, Int64, Int16, Int8, Kernel, Linkage, Vector16, GenericPointer, Groups, Float16, Float64, AtomicFloat32AddEXT, ExpectAssumeKHR], [SPV_EXT_shader_atomic_float_add, SPV_KHR_expect_assume]>, #spv.resource_limits<>>} {
+  func.func @main() {
+    %c8 = arith.constant 8 : index
+    %c1 = arith.constant 1 : index
+    %cst = arith.constant 2.200000e+00 : f32
+    %cst_0 = arith.constant 1.100000e+00 : f32
+    %cst_1 = arith.constant 0.000000e+00 : f32
+    %memref = gpu.alloc  () {gpu.alloc_shared} : memref<8xf32>
+    %memref_2 = gpu.alloc  () {gpu.alloc_shared} : memref<8xf32>
+    %memref_3 = gpu.alloc  () {gpu.alloc_shared} : memref<8xf32>
+    %0 = memref.cast %memref : memref<8xf32> to memref<?xf32>
+    %1 = memref.cast %memref_2 : memref<8xf32> to memref<?xf32>
+    %2 = memref.cast %memref_3 : memref<8xf32> to memref<?xf32>
+    call @fillResource1DFloat(%0, %cst_0) : (memref<?xf32>, f32) -> ()
+    call @fillResource1DFloat(%1, %cst) : (memref<?xf32>, f32) -> ()
+    call @fillResource1DFloat(%2, %cst_1) : (memref<?xf32>, f32) -> ()
+    gpu.launch_func  @main_kernel::@main_kernel blocks in (%c8, %c1, %c1) threads in (%c1, %c1, %c1) args(%memref : memref<8xf32>, %memref_2 : memref<8xf32>, %memref_3 : memref<8xf32>)
+    %3 = memref.cast %memref_3 : memref<8xf32> to memref<*xf32>
+    call @printMemrefF32(%3) : (memref<*xf32>) -> ()
+    return
+  }
+  gpu.module @main_kernel {
+    gpu.func @main_kernel(%arg0: memref<8xf32>, %arg1: memref<8xf32>, %arg2: memref<8xf32>) kernel attributes {spv.entry_point_abi = #spv.entry_point_abi<>} {
+      cf.br ^bb1
+    ^bb1:  // pred: ^bb0
+      %0 = gpu.block_id  x
+      %1 = memref.load %arg0[%0] : memref<8xf32>
+      %2 = memref.load %arg1[%0] : memref<8xf32>
+      %3 = arith.addf %1, %2 : f32
+      memref.store %3, %arg2[%0] : memref<8xf32>
+      gpu.return
+    }
+
+    // CHECK: spv.module @__spv__main_kernel Physical64 OpenCL {
+    // CHECK: spv.GlobalVariable @__builtin_var_WorkgroupId__ built_in("WorkgroupId") : !spv.ptr<vector<3xi64>, Input>
+    // CHECK: spv.func @main_kernel(%arg0: !spv.ptr<f32, CrossWorkgroup>, %arg1: !spv.ptr<f32, CrossWorkgroup>, %arg2: !spv.ptr<f32, CrossWorkgroup>) "None" attributes {spv.entry_point_abi = #spv.entry_point_abi<>, workgroup_attributions = 0 : i64} {
+    // CHECK: spv.Branch ^bb1
+    // CHECK: ^bb1:  // pred: ^bb0
+    // CHECK: %__builtin_var_WorkgroupId___addr = spv.mlir.addressof @__builtin_var_WorkgroupId__ : !spv.ptr<vector<3xi64>, Input>
+    // CHECK: %[[VAR0:.*]] = spv.Load "Input" %__builtin_var_WorkgroupId___addr : vector<3xi64>
+    // CHECK: %[[VAR1:.*]] = spv.CompositeExtract %[[VAR0:.*]][0 : i32] : vector<3xi64>
+    // CHECK: %[[VAR2:.*]] = spv.InBoundsPtrAccessChain %arg0[%[[VAR1:.*]]] : !spv.ptr<f32, CrossWorkgroup>, i64
+    // CHECK: %[[VAR3:.*]] = spv.Load "CrossWorkgroup" %[[VAR2:.*]] ["Aligned", 4] : f32
+    // CHECK: %[[VAR4:.*]] = spv.InBoundsPtrAccessChain %arg1[%[[VAR1:.*]]] : !spv.ptr<f32, CrossWorkgroup>, i64
+    // CHECK: %[[VAR5:.*]] = spv.Load "CrossWorkgroup" %[[VAR4:.*]] ["Aligned", 4] : f32
+    // CHECK:  %[[VAR6:.*]] = spv.FAdd %[[VAR3:.*]], %[[VAR5:.*]] : f32
+    // CHECK:  %[[VAR7:.*]] = spv.InBoundsPtrAccessChain %arg2[%[[VAR1:.*]]] : !spv.ptr<f32, CrossWorkgroup>, i64
+    // CHECK:   spv.Store "CrossWorkgroup" %[[VAR7:.*]], %[[VAR6:.*]] ["Aligned", 4] : f32
+    // CHECK:   spv.Return
+    // CHECK:  }
+    // CHECK: }
+
+  }
+  func.func private @fillResource1DFloat(memref<?xf32>, f32)
+  func.func private @printMemrefF32(memref<*xf32>)
+}
diff --git a/mlir/test/L0_RUNNER/lit.local.cfg b/mlir/test/L0_RUNNER/lit.local.cfg
@@ -0,0 +1 @@
+config.unsupported = True
diff --git a/mlir/test/Transforms/memref-alloc-to-gpu-alloc.mlir b/mlir/test/Transforms/memref-alloc-to-gpu-alloc.mlir
@@ -0,0 +1,29 @@
+// RUN: level-zero-opt --insert-gpu-alloc %s | FileCheck %s
+
+func.func @main() {
+  %c8 = arith.constant 8 : index
+  %c1 = arith.constant 1 : index
+  // CHECK: func.func @main()
+  %0 = memref.alloc() : memref<8xf32>
+  %1 = memref.alloc() : memref<8xf32>
+  %2 = memref.alloc() : memref<8xf32>
+  // CHECK: %[[MEMREF0:.*]]= gpu.alloc  () : memref<8xf32>
+  // CHECK: %[[MEMREF1:.*]]= gpu.alloc  () : memref<8xf32>
+  // CHECK: %[[MEMREF2:.*]]= gpu.alloc  () : memref<8xf32>
+  gpu.launch blocks(%arg0, %arg1, %arg2) in (%arg6 = %c8, %arg7 = %c1, %arg8 = %c1) threads(%arg3, %arg4, %arg5) in (%arg9 = %c1, %arg10 = %c1, %arg11 = %c1) {
+    // CHECK: gpu.launch {{.*}}
+    %7 = gpu.block_id  x
+    // CHECK: [[VAR0:.*]] = gpu.block_id  x
+    %8 = memref.load %0[%7] : memref<8xf32>
+    %9 = memref.load %1[%7] : memref<8xf32>
+
+    // CHECK: [[VAR1:.*]] = memref.load %[[MEMREF0:.*]][[[VAR0:.*]]] : memref<8xf32>
+    // CHECK: [[VAR2:.*]] = memref.load %[[MEMREF1:.*]][[[VAR0:.*]]] : memref<8xf32>
+    %10 = arith.addf %8, %9 : f32
+    memref.store %10, %2[%7] : memref<8xf32>
+    gpu.terminator
+    // CHECK: gpu.terminator
+  }
+  %6 = memref.cast %2 : memref<8xf32> to memref<*xf32>
+  return
+}
diff --git a/mlir/test/lit.cfg.py b/mlir/test/lit.cfg.py
@@ -49,13 +49,15 @@
 # test_exec_root: The root path where tests should be run.
 config.test_exec_root = os.path.join(config.imex_obj_root, 'mlir', 'test')
 config.imex_tools_dir = os.path.join(config.imex_obj_root, 'mlir', 'tools', 'level_zero_runner')
+config.imex_level_zero_opt_dir = os.path.join(config.imex_obj_root, 'mlir', 'tools', 'level-zero-opt')
 
 # Tweak the PATH to include the tools dir.
 llvm_config.with_environment('PATH', config.llvm_tools_dir, append_path=True)
 
-tool_dirs = [config.imex_tools_dir, config.llvm_tools_dir]
+tool_dirs = [config.imex_tools_dir, config.llvm_tools_dir, config.imex_level_zero_opt_dir]
 tools = [
-    'level_zero_runner'
+    'level_zero_runner',
+    'level-zero-opt'
 ]
 
 llvm_config.add_tool_substitutions(tools, tool_dirs)
diff --git a/mlir/tools/CMakeLists.txt b/mlir/tools/CMakeLists.txt
@@ -13,4 +13,5 @@
 # limitations under the License.
 
 add_subdirectory(level_zero_runner)
+add_subdirectory(level-zero-opt)
 
diff --git a/mlir/tools/level-zero-opt/CMakeLists.txt b/mlir/tools/level-zero-opt/CMakeLists.txt
@@ -0,0 +1,58 @@
+# Copyright 2022 Intel Corporation
+#
+# 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.
+
+project(level-zero-opt LANGUAGES CXX C)
+
+find_package(LLVM REQUIRED CONFIG)
+find_package(MLIR REQUIRED CONFIG)
+
+list(APPEND CMAKE_MODULE_PATH "${MLIR_CMAKE_DIR}")
+list(APPEND CMAKE_MODULE_PATH "${LLVM_CMAKE_DIR}")
+include(AddLLVM)
+include(AddMLIR)
+include(HandleLLVMOptions)
+
+get_property(dialect_libs GLOBAL PROPERTY MLIR_DIALECT_LIBS)
+get_property(conversion_libs GLOBAL PROPERTY MLIR_CONVERSION_LIBS)
+
+set(SOURCES_LIST
+    main.cpp
+    rewrite-wrapper.cpp
+    )
+set(HEADERS_LIST
+    )
+
+add_executable(${PROJECT_NAME} ${SOURCES_LIST} ${HEADERS_LIST})
+
+apply_llvm_compile_flags(${PROJECT_NAME})
+
+target_link_libraries(${PROJECT_NAME} PRIVATE
+    mlir-extensions
+    ${dialect_libs}
+    ${conversion_libs}
+    MLIROptLib
+    MLIRPass
+    MLIRSupport
+    )
+
+target_include_directories(${PROJECT_NAME} PRIVATE
+    ${LLVM_INCLUDE_DIRS}
+    ${MLIR_INCLUDE_DIRS}
+    )
+
+set(CMAKE_INSTALL_BINDIR ./tools/level-zero-opt)
+
+install(TARGETS level-zero-opt
+        PERMISSIONS OWNER_WRITE OWNER_READ OWNER_EXECUTE GROUP_READ GROUP_EXECUTE WORLD_READ WORLD_EXECUTE
+        )
diff --git a/mlir/tools/level-zero-opt/main.cpp b/mlir/tools/level-zero-opt/main.cpp
@@ -0,0 +1,27 @@
+// Copyright 2022 Intel Corporation
+//
+// 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.
+
+#include <mlir/IR/MLIRContext.h>
+#include <mlir/InitAllDialects.h>
+#include <mlir/InitAllPasses.h>
+#include <mlir/Tools/mlir-opt/MlirOptMain.h>
+
+int main(int argc, char **argv) {
+  mlir::registerAllPasses();
+  mlir::DialectRegistry registry;
+  mlir::registerAllDialects(registry);
+  return mlir::failed(MlirOptMain(argc, argv,
+                                  "L0 modular optimizer driver\n", registry,
+                                  /*preloadDialectsInContext=*/false));
+}
diff --git a/mlir/tools/level-zero-opt/rewrite-wrapper.cpp b/mlir/tools/level-zero-opt/rewrite-wrapper.cpp
@@ -0,0 +1,89 @@
+// Copyright 2022 Intel Corporation
+//
+// 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.
+
+#include <mlir/IR/PatternMatch.h>
+#include <mlir/Pass/Pass.h>
+#include <mlir/Pass/PassManager.h>
+#include <mlir/Pass/PassRegistry.h>
+#include <mlir/Transforms/GreedyPatternRewriteDriver.h>
+
+#include <mlir/Dialect/Func/IR/FuncOps.h>
+#include <mlir/Dialect/SCF/IR/SCF.h>
+
+#include "mlir/Dialect/GPU/Transforms/Passes.h"
+#include "mlir-extensions/Conversion/gpu_to_gpu_runtime.hpp"
+#include "mlir-extensions/Conversion/gpu_runtime_to_llvm.hpp"
+#include "mlir-extensions/Transforms/rewrite_wrapper.hpp"
+
+static mlir::PassPipelineRegistration<> ParallelLoopToGpu(
+    "parallel-loop-to-gpu", "Maps scf parallel loop to gpu",
+    [](mlir::OpPassManager &pm) {
+      pm.addNestedPass<mlir::func::FuncOp>(gpu_runtime::createParallelLoopGPUMappingPass());
+    });
+
+static mlir::PassPipelineRegistration<> InsertGpuAlloc(
+    "insert-gpu-alloc", "Converts memref alloc to gpu alloc",
+    [](mlir::OpPassManager &pm) {
+      pm.addNestedPass<mlir::func::FuncOp>(gpu_runtime::createInsertGPUAllocsPass());
+    });
+
+static mlir::PassPipelineRegistration<> UnstrideMemrefPass(
+    "unstride-memref", "Used to flatten 2D to 1D",
+    [](mlir::OpPassManager &pm) {
+      pm.addNestedPass<mlir::func::FuncOp>(gpu_runtime::createUnstrideMemrefsPass());
+    });
+
+static mlir::PassPipelineRegistration<> AbiAttrsPass(
+    "abi-attrs", "Create AbiAttrs Pass",
+    [](mlir::OpPassManager &pm) {
+      pm.addNestedPass<mlir::gpu::GPUModuleOp>(
+      gpu_runtime::createAbiAttrsPass());
+    });
+
+static mlir::PassPipelineRegistration<> SetSpirvCapabalities(
+    "set-spirv-capablilities", "Sets Spirv capabilities",
+    [](mlir::OpPassManager &pm) {
+      pm.addPass(gpu_runtime::createGPUToSpirvPass());
+    });
+
+static mlir::PassPipelineRegistration<> GpuToSpirv(
+    "gpu-to-spirv", "Converts Gpu to spirv module",
+    [](mlir::OpPassManager &pm) {
+      pm.addPass(gpu_runtime::createGPUToSpirvPass());
+    });
+
+static mlir::PassPipelineRegistration<> SerializeSpirv(
+    "serialize-spirv", "Serializes the spir-v binary",
+    [](mlir::OpPassManager &pm) {
+      pm.addPass(gpu_runtime::createSerializeSPIRVPass());
+    });
+
+ static mlir::PassPipelineRegistration<> GpuToGpuRuntime(
+    "gpu-to-gpu-runtime", "Converts Gpu ops to gpu runteim",
+    [](mlir::OpPassManager &pm) {
+      pm.addNestedPass<mlir::func::FuncOp>(gpu_runtime::createGPUExPass());
+    });
+
+ static mlir::PassPipelineRegistration<> EnumerateEvents(
+    "enumerate-events", "Adds event dependency",
+    [](mlir::OpPassManager &pm) {
+      pm.addPass(gpu_runtime::createEnumerateEventsPass());
+    });
+
+  static mlir::PassPipelineRegistration<> GpuToLlvm(
+    "convert-gpu-to-llvm", "Converts Gpu runtime dialect to llvm runtime calls",
+    [](mlir::OpPassManager &pm) {
+      pm.addPass(gpu_runtime::createGPUToLLVMPass());
+    });
+
diff --git a/scripts/build_locally.py b/scripts/build_locally.py
@@ -310,6 +310,7 @@ def _build_imex(
     enable_numba_fe=False,
     with_tbb=None,
     enable_numba_hotfix=False,
+    external_lit=None,
 ):
     """Builds Intel MLIR extensions (IMEX).
 
@@ -400,6 +401,9 @@ def _build_imex(
             cmake_config_args.append("-DIMEX_ENABLE_TBB_SUPPORT=ON")
         else:
             warnings.warn("Provided TBB directory path does not exist.")
+
+    if external_lit is not None:
+        cmake_config_args.append("-DLLVM_EXTERNAL_LIT=" + external_lit)
 
     build_dir = os.path.abspath(build_dir)
     # Configure
@@ -513,6 +517,12 @@ def _build_imex(
         action="store_true",
         help="Removes any existing build directory when building IMEX",
     )
+    imex_builder.add_argument(
+        "--external-lit",
+        type=str,
+        help="Path to a lit executable",
+        dest="external_lit",
+    )
 
     args = parser.parse_args()
 
@@ -629,6 +639,7 @@ def _build_imex(
         enable_igpu=args.imex_enable_igpu,
         with_tbb=g_tbb_dir,
         enable_tests=args.imex_enable_tests,
+        external_lit=args.external_lit,
     )
 
     # TODO
Original file line number	Diff line number	Diff line change
Expand Up		@@ -13,4 +13,5 @@
		# limitations under the License.

		add_subdirectory(level_zero_runner)
		add_subdirectory(level-zero-opt)