Add OpenCL runtime (#191)

This introduces a OCL stream wrapper, and the implementation of gpux dialect's and upstream runtime with OCL.

IMEX runtime wrapper
If context and device is passed with nullptr, we will return a global stream object. The stream object is responsible for releasing the context. If both are given, we "borrow" the context and we don't release the context at the destructor of the stream object.

Upstream runtime wrapper
We use the same stream object class as used in IMEX wrapper. In upstream ROCM/Cuda/Sycl wrapper, they use raw queue pointer as the type of queue parameter in mgpu* APIS (like sycl::queue*).

Author: Menooker

README.md

@@ -58,7 +58,15 @@ cmake --build . --target gc-check
 Notes:
  * `/PATH/TO/llvm-project/llvm-install` should be the install path of LLVM. If you installed LLVM elsewhere by `-DCMAKE_INSTALL_PREFIX` option when building LLVM, you need to change the path in `-DMLIR_DIR` accordingly.
  *  The cmake option `-DLLVM_EXTERNAL_LIT` is for the tests of this project. It requires the `lit` tool to be installed in the system. You can install it via `pip install lit`. If you don't need to run the tests of this repo, you can omit this option in the command line.
- * If GPU components are on (`-DGC_USE_GPU=ON`), make sure the Level-zero runtime is installed in your system. Either install Level-zero runtime via system package managers (e.g. `apt`), or follow the instructions of [IMEX](https://github.com/intel/mlir-extensions).
+
+More notes if GPU components are on (`-DGC_USE_GPU=ON`):
+ * make sure the OpenCL runtime is installed in your system. You can either
+  install using OS-provided package (Ubuntu 22.04)
+```sh
+sudo apt install -y intel-opencl-icd opencl-c-headers
+```
+  Or, download and install package from: https://github.com/intel/compute-runtime/releases
+ * the LLVM codebase needs to be patched to support XeGPU lowering (from IMEX). Please follow instructions of [IMEX](https://github.com/intel/mlir-extensions) on patching LLVM.
 
 Graph Compiler supports the following build-time options.
 

cmake/imex.cmake

@@ -4,11 +4,11 @@ get_property(IMEX_INCLUDES GLOBAL PROPERTY IMEX_INCLUDES)
 if (NOT DEFINED IMEX_INCLUDES)
     include(functions)
     set(IMEX_CHECK_LLVM_VERSION ON)
-    set(IMEX_ENABLE_L0_RUNTIME 1)
+    set(IMEX_ENABLE_L0_RUNTIME 0)
     # TODO: Change to main https://github.com/oneapi-src/oneDNN.git when all the
     # required functionality is merged.
     gc_fetch_content(imex 496b240093b5e132b60c5ee69878300fe69be300 https://github.com/Menooker/mlir-extensions
-            CMAKE_ARGS "-DMLIR_DIR=${MLIR_DIR};-DIMEX_CHECK_LLVM_VERSION=ON;-DIMEX_ENABLE_L0_RUNTIME=1"
+            CMAKE_ARGS "-DMLIR_DIR=${MLIR_DIR};-DIMEX_CHECK_LLVM_VERSION=ON;-DIMEX_ENABLE_L0_RUNTIME=0"
     )
 
     set(IMEX_INCLUDES

include/gc/Transforms/Passes.td

@@ -32,7 +32,7 @@ def ConvertOneDNNGraphToLinalg : Pass<"convert-onednn-graph-to-linalg"> {
   ];
 }
 
-
+#ifdef GC_USE_GPU
 def LinalgToXeGPU : Pass<"linalg-to-xegpu", "func::FuncOp"> {
   let summary = "Convert linalg dialect to XeGPU dialect.";
   let description = [{
@@ -57,5 +57,6 @@ def LinalgToXeGPU : Pass<"linalg-to-xegpu", "func::FuncOp"> {
                "DPAS register block sizes MxNxK">,
   ];
 }
+#endif
 
 #endif // GC_DIALECT_GC_PASSES

lib/gc/CAPI/CMakeLists.txt

@@ -1,10 +1,16 @@
+set(GC_ALL_LIBS 
+  MLIROneDNNGraph
+  MLIRCPURuntimeDialect
+  GCPasses
+  MLIRCPURuntimeTransforms)
+
+if(GC_USE_GPU)
+  list(APPEND GC_ALL_LIBS GCGPUPasses)
+endif()
+
 add_mlir_public_c_api_library(GcCAPI
   Dialects.cpp
   Passes.cpp
   LINK_LIBS PUBLIC
-  MLIROneDNNGraph
-  MLIRCPURuntimeDialect
-  GCPasses
-  GCGPUPasses
-  MLIRCPURuntimeTransforms
+  ${GC_ALL_LIBS}
 )

lib/gc/ExecutionEngine/CMakeLists.txt

@@ -1,2 +1,5 @@
 add_subdirectory(CPURuntime)
-add_subdirectory(Driver)
+add_subdirectory(Driver)
+if(GC_USE_GPU)
+    add_subdirectory(OpenCLRuntime)
+endif()

lib/gc/ExecutionEngine/Driver/CMakeLists.txt

@@ -26,6 +26,11 @@ else()
   )
 endif()
 
+set(GC_PASSES GCPasses)
+if(GC_USE_GPU)
+  list(APPEND GC_PASSES GCGPUPasses)
+endif()
+
 add_mlir_library(GCJitWrapper
   Driver.cpp
 
@@ -35,8 +40,7 @@ add_mlir_library(GCJitWrapper
   LINK_LIBS PUBLIC
     ${MLIR_LINK_COMPONENTS}
     ${dialect_libs}
-    ${conversion_libs} 
-    GCPasses
-    GCGPUPasses
+    ${conversion_libs}
+    ${GC_PASSES}
   )
 

lib/gc/ExecutionEngine/OpenCLRuntime/CMakeLists.txt

@@ -0,0 +1,22 @@
+find_package(OpenCL REQUIRED)
+
+add_mlir_library(mlir_opencl_runtime
+    SHARED
+    OpenCLRuntimeWrappers.cpp
+
+    EXCLUDE_FROM_LIBMLIR
+  )
+
+check_cxx_compiler_flag("-frtti" CXX_HAS_FRTTI_FLAG)
+if(NOT CXX_HAS_FRTTI_FLAG)
+    message(FATAL_ERROR "CXX compiler does not accept flag -frtti")
+endif()
+target_compile_options (mlir_opencl_runtime PUBLIC -fexceptions -frtti)
+
+target_include_directories(mlir_opencl_runtime PRIVATE
+    ${MLIR_INCLUDE_DIRS}
+    ${OpenCL_INCLUDE_DIRS}
+    )
+
+message(STATUS "OpenCL Libraries: ${OpenCL_LIBRARIES}")
+target_link_libraries(mlir_opencl_runtime PUBLIC ${OpenCL_LIBRARIES})