Skip to content

Commit 502d3b9

Browse files
etiottoetiotto
committed
Create infrastructure for SYCLToLLVM conversion (#37)
Signed-off-by: Tiotto, Ettore <ettore.tiotto@intel.com>
1 parent 800f288 commit 502d3b9

31 files changed

+290
-94
lines changed

‎mlir-sycl/CMakeLists.txt

+1-1
Original file line numberDiff line numberDiff line change
@@ -67,7 +67,7 @@ set(MLIR_SYCL_TOOLS_DIR ${CMAKE_BINARY_DIR}/bin)
6767
set(LLVM_LIT_ARGS "-sv" CACHE STRING "lit default options")
6868
list(APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake/modules")
6969

70-
add_subdirectory(include)
70+
add_subdirectory(include/mlir)
7171
add_subdirectory(lib)
7272
add_subdirectory(test)
7373

‎mlir-sycl/include/CMakeLists.txt

-11
This file was deleted.

‎mlir-sycl/include/SYCL/SYCLToLLVM.h

-27
This file was deleted.

‎mlir-sycl/include/mlir/CMakeLists.txt

+2
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,2 @@
1+
add_subdirectory(Conversion)
2+
add_subdirectory(Dialect)

‎mlir-sycl/include/mlir/Conversion/CMakeLists.txt

+3
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,3 @@
1+
set(LLVM_TARGET_DEFINITIONS SYCLPasses.td)
2+
mlir_tablegen(SYCLPasses.h.inc -gen-pass-decls -name Conversion)
3+
add_public_tablegen_target(MLIRSYCLConversionPassIncGen)

‎mlir-sycl/include/mlir/Conversion/SYCLPasses.h

+22
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,22 @@
1+
//===- SYCLPasses.h - Conversion Pass Construction and Registration -----------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#ifndef MLIR_CONVERSION_SYCLPASSES_H
10+
#define MLIR_CONVERSION_SYCLPASSES_H
11+
12+
#include "mlir/Conversion/SYCLToLLVM/SYCLToLLVM.h"
13+
14+
namespace mlir {
15+
16+
/// Generate the code for registering conversion passes.
17+
#define GEN_PASS_REGISTRATION
18+
#include "mlir/Conversion/SYCLPasses.h.inc"
19+
20+
} // namespace mlir
21+
22+
#endif // MLIR_CONVERSION_SYCLPASSES_H

‎mlir-sycl/include/mlir/Conversion/SYCLPasses.td

+28
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,28 @@
1+
//===-- SYCLPasses.td - Conversion pass definition file ----*- tablegen -*-===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#ifndef MLIR_CONVERSION_SYCLPASSES
10+
#define MLIR_CONVERSION_SYCLPASSES
11+
12+
include "mlir/Pass/PassBase.td"
13+
14+
//===----------------------------------------------------------------------===//
15+
// SYCLToLLVM
16+
//===----------------------------------------------------------------------===//
17+
18+
def ConvertSYCLToLLVM : Pass<"convert-sycl-to-llvm", "ModuleOp"> {
19+
let summary = "Convert SYCL dialect to LLVM dialect";
20+
let description = [{
21+
See docs/SYCLToLLVMDialectConversion/ for more details.
22+
TODO: add docs referenced above.
23+
}];
24+
let constructor = "mlir::createConvertSYCLToLLVMPass()";
25+
let dependentDialects = ["LLVM::LLVMDialect"];
26+
}
27+
28+
#endif // MLIR_CONVERSION_SYCLPASSES

‎mlir-sycl/include/mlir/Conversion/SYCLToLLVM/SYCLToLLVM.h

+46
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,46 @@
1+
//===- SYCLToLLVM.h - SYCL to LLVM Patterns ---------------------*- C++ -*-===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
//
9+
// Provides patterns to convert SYCL dialect to LLVM dialect.
10+
//
11+
//===----------------------------------------------------------------------===//
12+
13+
#ifndef MLIR_CONVERSION_SYCLTOLLVM_SYCLTOLLVM_H
14+
#define MLIR_CONVERSION_SYCLTOLLVM_SYCLTOLLVM_H
15+
16+
#include "mlir/Transforms/DialectConversion.h"
17+
18+
namespace mlir {
19+
class LLVMTypeConverter;
20+
class MLIRContext;
21+
class ModuleOp;
22+
23+
namespace sycl {
24+
template <typename SYCLOp>
25+
class SYCLToLLVMConversion : public OpConversionPattern<SYCLOp> {
26+
public:
27+
SYCLToLLVMConversion(MLIRContext *context, LLVMTypeConverter &typeConverter,
28+
PatternBenefit benefit = 1)
29+
: OpConversionPattern<SYCLOp>(typeConverter, context, benefit),
30+
typeConverter(typeConverter) {}
31+
32+
protected:
33+
LLVMTypeConverter &typeConverter;
34+
};
35+
36+
/// Populates type conversions with additional SYCL types.
37+
void populateSYCLToLLVMTypeConversion(LLVMTypeConverter &typeConverter);
38+
39+
/// Populates the given list with patterns that convert from SYCL to LLVM.
40+
void populateSYCLToLLVMConversionPatterns(LLVMTypeConverter &typeConverter,
41+
RewritePatternSet &patterns);
42+
43+
} // namespace sycl
44+
} // namespace mlir
45+
46+
#endif // MLIR_CONVERSION_SYCLTOLLVM_SYCLTOLLVM_H

‎mlir-sycl/include/mlir/Conversion/SYCLToLLVM/SYCLToLLVMPass.h

+27
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,27 @@
1+
//===- SYCLToLLVMPass.h - SYCL to LLVM Passes -------------------*- C++ -*-===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
//
9+
// Provides passes to convert SYCL dialect to LLVM dialect.
10+
//
11+
//===----------------------------------------------------------------------===//
12+
13+
#ifndef MLIR_CONVERSION_SYCLTOLLVM_SYCLTOLLVMPASS_H
14+
#define MLIR_CONVERSION_SYCLTOLLVM_SYCLTOLLVMPASS_H
15+
16+
#include <memory>
17+
18+
namespace mlir {
19+
class ModuleOp;
20+
template <typename T> class OperationPass;
21+
22+
/// Creates a pass to convert SYCL operations to the LLVMIR dialect.
23+
std::unique_ptr<OperationPass<ModuleOp>> createConvertSYCLToLLVMPass();
24+
25+
} // namespace mlir
26+
27+
#endif // MLIR_CONVERSION_SYCLTOLLVM_SYCLTOLLVMPASS_H

‎mlir-sycl/include/mlir/Dialect/CMakeLists.txt

+1
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
add_subdirectory(SYCL)

‎mlir-sycl/include/mlir/Dialect/SYCL/CMakeLists.txt

+1
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
add_subdirectory(IR)

‎mlir-sycl/include/SYCL/CMakeLists.txt ‎mlir-sycl/include/mlir/Dialect/SYCL/IR/CMakeLists.txt

File renamed without changes.

‎mlir-sycl/include/SYCL/SYCLOps.td ‎mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOps.td

File renamed without changes.

‎mlir-sycl/include/SYCL/SYCLOpsAlias.h ‎mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOpsAlias.h

File renamed without changes.

‎mlir-sycl/include/SYCL/SYCLOpsDialect.h ‎mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOpsDialect.h

+2-2
Original file line numberDiff line numberDiff line change
@@ -19,11 +19,11 @@
1919

2020
/// Include the auto-generated header file containing the declaration of the
2121
/// sycl dialect.
22-
#include "SYCL/SYCLOpsDialect.h.inc"
22+
#include "mlir/Dialect/SYCL/IR/SYCLOpsDialect.h.inc"
2323

2424
/// Include the auto-generated header file containing the declarations of the
2525
/// sycl operations.
2626
#define GET_OP_CLASSES
27-
#include "SYCL/SYCLOps.h.inc"
27+
#include "mlir/Dialect/SYCL/IR/SYCLOps.h.inc"
2828

2929
#endif // MLIR_SYCL_OPS_DIALECT_H_

‎mlir-sycl/include/SYCL/SYCLOpsTypes.h ‎mlir-sycl/include/mlir/Dialect/SYCL/IR/SYCLOpsTypes.h

File renamed without changes.

‎mlir-sycl/lib/CMakeLists.txt

+4-10
Original file line numberDiff line numberDiff line change
@@ -1,11 +1,5 @@
1-
# Copyright (C) Codeplay Software Limited
1+
# Enable errors for any global constructors.
2+
add_flag_if_supported("-Werror=global-constructors" WERROR_GLOBAL_CONSTRUCTOR)
23

3-
#===--- CMakeLists.txt -----------------------------------------------------===#
4-
#
5-
# MLIR-SYCL is under the Apache License v2.0 with LLVM Exceptions.
6-
# See https://llvm.org/LICENSE.txt for license information.
7-
# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
8-
#
9-
#===------------------------------------------------------------------------===#
10-
11-
add_subdirectory(SYCL)
4+
add_subdirectory(Conversion)
5+
add_subdirectory(Dialect)

‎mlir-sycl/lib/Conversion/CMakeLists.txt

+1
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
add_subdirectory(SYCLToLLVM)

‎mlir-sycl/lib/Conversion/PassDetail.h

+31
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,31 @@
1+
//===- PassDetail.h - Conversion Pass class details -------------*- C++ -*-===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#ifndef SYCL_CONVERSION_PASSDETAIL_H_
10+
#define SYCL_CONVERSION_PASSDETAIL_H_
11+
12+
#include "mlir/Pass/Pass.h"
13+
#include "mlir/IR/BuiltinOps.h"
14+
#include "mlir/IR/FunctionInterfaces.h"
15+
16+
namespace mlir {
17+
18+
namespace LLVM {
19+
class LLVMDialect;
20+
}
21+
22+
namespace sycl {
23+
class SYCLDialect;
24+
}
25+
26+
#define GEN_PASS_CLASSES
27+
#include "mlir/Conversion/SYCLPasses.h.inc"
28+
29+
} // namespace mlir
30+
31+
#endif // SYCL_CONVERSION_PASSDETAIL_H_

‎mlir-sycl/lib/Conversion/SYCLToLLVM/CMakeLists.txt

+20
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,20 @@
1+
add_mlir_conversion_library(MLIRSYCLToLLVM
2+
SYCLToLLVM.cpp
3+
SYCLToLLVMPass.cpp
4+
5+
ADDITIONAL_HEADER_DIRS
6+
${MLIR_SYCL_SOURCE_DIR}/mlir/Conversion/SYCLToLLVM
7+
8+
DEPENDS
9+
MLIRSYCLConversionPassIncGen
10+
11+
LINK_LIBS PUBLIC
12+
MLIRArithmeticToLLVM
13+
MLIRFuncToLLVM
14+
MLIRIR
15+
MLIRLLVMCommonConversion
16+
MLIRLLVMDialect
17+
MLIRMemRefToLLVM
18+
MLIRSYCLDialect
19+
MLIRTransforms
20+
)

‎mlir-sycl/lib/SYCL/SYCLToLLVM.cpp ‎mlir-sycl/lib/Conversion/SYCLToLLVM/SYCLToLLVM.cpp

+33-16
Original file line numberDiff line numberDiff line change
@@ -1,26 +1,35 @@
1-
// Copyright (C) Intel
2-
3-
//===--- SYCLToLLVM.cpp ---------------------------------------------------===//
1+
//===- SYCLToLLVM.cpp - SYCL to LLVM Patterns -----------------------------===//
42
//
5-
// MLIR-SYCL is under the Apache License v2.0 with LLVM Exceptions.
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
64
// See https://llvm.org/LICENSE.txt for license information.
75
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
86
//
97
//===----------------------------------------------------------------------===//
108
//
11-
// This file implements a pass to convert SYCL dialects types to their
12-
// corresponding LLVM dialect types.
9+
// This file implements patterns to convert SYCL dialect to LLVM dialect.
1310
//
1411
//===----------------------------------------------------------------------===//
1512

16-
#include "SYCL/SYCLToLLVM.h"
17-
#include "SYCL/SYCLOpsTypes.h"
13+
#include "mlir/Conversion/SYCLToLLVM/SYCLToLLVM.h"
1814
#include "mlir/Conversion/LLVMCommon/Pattern.h"
15+
#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
1916
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
17+
#include "mlir/Dialect/SYCL/IR/SYCLOpsDialect.h"
18+
#include "mlir/Dialect/SYCL/IR/SYCLOpsTypes.h"
19+
#include "mlir/IR/BuiltinOps.h"
20+
#include "mlir/IR/PatternMatch.h"
21+
#include "mlir/Support/LogicalResult.h"
22+
#include "mlir/Transforms/DialectConversion.h"
23+
#include "llvm/Support/Debug.h"
24+
25+
#define DEBUG_TYPE "sycl-to-llvm-pattern"
2026

2127
using namespace mlir;
2228

23-
namespace {
29+
//===----------------------------------------------------------------------===//
30+
// Utility functions
31+
//===----------------------------------------------------------------------===//
32+
2433
// Get the LLVM type of "class.cl::sycl::detail::array" with number of
2534
// dimentions \p dimNum and element type \p type.
2635
static Type getSYCLArrayTy(MLIRContext &context, unsigned dimNum, Type type) {
@@ -55,16 +64,24 @@ static Type getSYCLRangeOrIDTy(T type, StringRef name,
5564
}
5665
return LLVM::LLVMPointerType::get(structTy);
5766
}
58-
} // namespace
5967

60-
void mlir::sycl::populateSYCLToLLVMConversionPatterns(
61-
LLVMTypeConverter &converter, RewritePatternSet &patterns) {
62-
converter.addConversion([&](mlir::sycl::IDType type) {
68+
//===----------------------------------------------------------------------===//
69+
// Pattern population
70+
//===----------------------------------------------------------------------===//
71+
72+
void mlir::sycl::populateSYCLToLLVMTypeConversion(
73+
LLVMTypeConverter &typeConverter) {
74+
typeConverter.addConversion([&](mlir::sycl::IDType type) {
6375
return getSYCLRangeOrIDTy<mlir::sycl::IDType>(type, "class.cl::sycl::id",
64-
converter);
76+
typeConverter);
6577
});
66-
converter.addConversion([&](mlir::sycl::RangeType type) {
78+
typeConverter.addConversion([&](mlir::sycl::RangeType type) {
6779
return getSYCLRangeOrIDTy<mlir::sycl::RangeType>(
68-
type, "class.cl::sycl::range", converter);
80+
type, "class.cl::sycl::range", typeConverter);
6981
});
7082
}
83+
84+
void mlir::sycl::populateSYCLToLLVMConversionPatterns(
85+
LLVMTypeConverter &typeConverter, RewritePatternSet &patterns) {
86+
populateSYCLToLLVMTypeConversion(typeConverter);
87+
}

‎mlir-sycl/lib/Conversion/SYCLToLLVM/SYCLToLLVMPass.cpp

+50
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,50 @@
1+
//===- SYCLToLLVMPass.cpp - SYCL to LLVM Passes ---------------------------===//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
//
9+
// This file implements a pass to convert MLIR SYCL ops into LLVM ops
10+
//
11+
//===----------------------------------------------------------------------===//
12+
13+
#include "mlir/Conversion/SYCLToLLVM/SYCLToLLVMPass.h"
14+
#include "../PassDetail.h"
15+
#include "mlir/Conversion/LLVMCommon/TypeConverter.h"
16+
#include "mlir/Conversion/SYCLToLLVM/SYCLToLLVM.h"
17+
#include "mlir/Dialect/LLVMIR/LLVMDialect.h"
18+
#include "mlir/Dialect/SYCL/IR/SYCLOpsDialect.h"
19+
20+
using namespace mlir;
21+
22+
namespace {
23+
/// A pass converting MLIR SYCL operations into LLVM dialect.
24+
class ConvertSYCLToLLVMPass
25+
: public ConvertSYCLToLLVMBase<ConvertSYCLToLLVMPass> {
26+
void runOnOperation() override;
27+
};
28+
} // namespace
29+
30+
void ConvertSYCLToLLVMPass::runOnOperation() {
31+
MLIRContext *context = &getContext();
32+
ModuleOp module = getOperation();
33+
LLVMTypeConverter converter(&getContext());
34+
35+
RewritePatternSet patterns(context);
36+
37+
sycl::populateSYCLToLLVMConversionPatterns(converter, patterns);
38+
39+
ConversionTarget target(*context);
40+
target.addIllegalDialect<sycl::SYCLDialect>();
41+
target.addLegalDialect<LLVM::LLVMDialect>();
42+
43+
target.addLegalOp<ModuleOp>();
44+
if (failed(applyPartialConversion(module, target, std::move(patterns))))
45+
signalPassFailure();
46+
}
47+
48+
std::unique_ptr<OperationPass<ModuleOp>> mlir::createConvertSYCLToLLVMPass() {
49+
return std::make_unique<ConvertSYCLToLLVMPass>();
50+
}

0 commit comments

Comments
 (0)