Add type converter from sycl (id and range) type to llvm type (#34)

Type conversions are still work in progress.
This PR adds the support of type converters of `sycl::IDType` and `sycl::RangeType` to LLVM types.

Signed-off-by: Tsang, Whitney <whitney.tsang@intel.com>

Author: whitneywhtsang

mlir-sycl/include/SYCL/SYCLToLLVM.h

@@ -0,0 +1,27 @@
+// Copyright (C) Intel
+
+//===--- SYCLToLLVM.h -----------------------------------------------------===//
+//
+// MLIR-SYCL is under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef SYCL_SYCLTOLLVM_H_
+#define SYCL_SYCLTOLLVM_H_
+
+namespace mlir {
+class LLVMTypeConverter;
+class RewritePatternSet;
+
+namespace sycl {
+/// Collect the patterns to convert from the SYCL dialect to LLVM dialect. The
+/// conversion patterns capture the LLVMTypeConverter by reference meaning the
+/// references have to remain alive during the entire pattern lifetime.
+void populateSYCLToLLVMConversionPatterns(LLVMTypeConverter &converter,
+                                          RewritePatternSet &patterns);
+} // namespace sycl
+} // namespace mlir
+
+#endif // SYCL_SYCLTOLLVM_H_

mlir-sycl/lib/SYCL/CMakeLists.txt

@@ -39,6 +39,7 @@ add_mlir_dialect_library(MLIRSYCLDialect
   SYCLOpsDialect.cpp
   SYCLOpsTypes.cpp
   SYCLOpsAlias.cpp
+  SYCLToLLVM.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${PROJECT_SOURCE_DIR}/include/SYCL

mlir-sycl/lib/SYCL/SYCLToLLVM.cpp

@@ -0,0 +1,70 @@
+// Copyright (C) Intel
+
+//===--- SYCLToLLVM.cpp ---------------------------------------------------===//
+//
+// MLIR-SYCL is under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file implements a pass to convert SYCL dialects types to their
+// corresponding LLVM dialect types.
+//
+//===----------------------------------------------------------------------===//
+
+#include "SYCL/SYCLToLLVM.h"
+#include "SYCL/SYCLOpsTypes.h"
+#include "mlir/Conversion/LLVMCommon/Pattern.h"
+#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
+
+using namespace mlir;
+
+namespace {
+// Get the LLVM type of "class.cl::sycl::detail::array" with number of
+// dimentions \p dimNum and element type \p type.
+static Type getSYCLArrayTy(MLIRContext &context, unsigned dimNum, Type type) {
+  assert((dimNum == 1 || dimNum == 2 || dimNum == 3) &&
+         "Expecting number of dimensions to be 1, 2, or 3.");
+  auto structTy = LLVM::LLVMStructType::getIdentified(
+      &context, "class.cl::sycl::detail::array." + std::to_string(dimNum));
+  if (!structTy.isInitialized()) {
+    auto arrayTy = LLVM::LLVMArrayType::get(type, dimNum);
+    auto res = structTy.setBody({arrayTy}, /*isPacked=*/false);
+    assert(succeeded(res) &&
+           "Unexpected failure from LLVMStructType::setBody.");
+  }
+  return structTy;
+}
+
+// Get the LLVM type of a SYCL range or id type, given \p type - the type in
+// SYCL, \p name - the expected LLVM type name, \p converter - LLVM type
+// converter.
+template <typename T>
+static Type getSYCLRangeOrIDTy(T type, StringRef name,
+                               LLVMTypeConverter &converter) {
+  unsigned dimNum = type.getDimension();
+  auto structTy = LLVM::LLVMStructType::getIdentified(
+      &converter.getContext(), name.str() + "." + std::to_string(dimNum));
+  if (!structTy.isInitialized()) {
+    auto res = structTy.setBody(getSYCLArrayTy(converter.getContext(), dimNum,
+                                               converter.getIndexType()),
+                                /*isPacked=*/false);
+    assert(succeeded(res) &&
+           "Unexpected failure from LLVMStructType::setBody.");
+  }
+  return LLVM::LLVMPointerType::get(structTy);
+}
+} // namespace
+
+void mlir::sycl::populateSYCLToLLVMConversionPatterns(
+    LLVMTypeConverter &converter, RewritePatternSet &patterns) {
+  converter.addConversion([&](mlir::sycl::IDType type) {
+    return getSYCLRangeOrIDTy<mlir::sycl::IDType>(type, "class.cl::sycl::id",
+                                                  converter);
+  });
+  converter.addConversion([&](mlir::sycl::RangeType type) {
+    return getSYCLRangeOrIDTy<mlir::sycl::RangeType>(
+        type, "class.cl::sycl::range", converter);
+  });
+}

polygeist/lib/polygeist/Passes/CMakeLists.txt

@@ -36,5 +36,6 @@ add_mlir_dialect_library(MLIRPolygeistTransforms
   MLIRPolygeist
   MLIRSideEffectInterfaces
   MLIRSCFToControlFlow
+  MLIRSYCLDialect
   MLIRTransformUtils
   )

polygeist/lib/polygeist/Passes/ConvertPolygeistToLLVM.cpp

@@ -11,6 +11,7 @@
 //===----------------------------------------------------------------------===//
 #include "PassDetails.h"
 
+#include "SYCL/SYCLToLLVM.h"
 #include "mlir/Analysis/DataLayoutAnalysis.h"
 #include "mlir/Conversion/ArithmeticToLLVM/ArithmeticToLLVM.h"
 #include "mlir/Conversion/ControlFlowToLLVM/ControlFlowToLLVM.h"
@@ -630,6 +631,7 @@ struct ConvertPolygeistToLLVMPass
 
       LLVMTypeConverter converter(&getContext(), options, &dataLayoutAnalysis);
       RewritePatternSet patterns(&getContext());
+      sycl::populateSYCLToLLVMConversionPatterns(converter, patterns);
       populatePolygeistToLLVMConversionPatterns(converter, patterns);
       populateSCFToControlFlowConversionPatterns(patterns);
       cf::populateControlFlowToLLVMConversionPatterns(converter, patterns);

polygeist/tools/cgeist/Test/Verification/sycl/SYCLToLLVM/types.mlir

@@ -0,0 +1,32 @@
+// Copyright (C) Intel
+
+//===--- types.mlir ------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+// RUN: cgeist %s -S -emit-llvm --args -fsycl-is-device | FileCheck %s
+
+// CHECK-DAG: [[ARRAY_1:.*]] = type { [1 x i64] }
+// CHECK-DAG: [[ARRAY_2:.*]] = type { [2 x i64] }
+// CHECK-DAG: [[ID:.*]] = type { [[ARRAY_1]] }
+// CHECK-DAG: [[RANGE_1:.*]] = type { [[ARRAY_1]] }
+// CHECK-DAG: [[RANGE_2:.*]] = type { [[ARRAY_2]] }
+// CHECK: define void @test_id([[ID]]* %0, [[ID]]* %1)
+// CHECK: define void @test_range.1([[RANGE_1]]* %0)
+// CHECK: define void @test_range.2([[RANGE_2]]* %0)
+
+module {
+  func.func @test_id(%arg0: !sycl.id<1>, %arg1: !sycl.id<1>) {
+    return
+  }
+  func.func @test_range.1(%arg0: !sycl.range<1>) {
+    return
+  }
+  func.func @test_range.2(%arg0: !sycl.range<2>) {
+    return
+  }
+}

polygeist/tools/cgeist/Test/lit.cfg

@@ -17,7 +17,7 @@ config.test_format = lit.formats.ShTest(not llvm_config.use_lit_shell)
 
 # suffixes: A list of file extensions to treat as test files. This is overriden
 # by individual lit.local.cfg files in the test subdirectories.
-config.suffixes = ['.c', '.cpp', '.cu']
+config.suffixes = ['.c', '.cpp', '.cu', '.mlir']
 
 # excludes: A list of directories or files to exclude from the testsuite even
 # if they match the suffixes pattern.