Skip to content

Commit d850c76

Browse files
whitneywhtsangetiotto
whitneywhtsang
authored and
etiotto
committed
untying polygeist from the inteon (#16)
This is a patch sent to use by CodePlay to allow compilation of the cgesit SYCL tests
1 parent e590edf commit d850c76

File tree

7 files changed

+33
-32
lines changed

7 files changed

+33
-32
lines changed

‎polygeist/tools/cgeist/Lib/clang-mlir.cc

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -4829,9 +4829,9 @@ MLIRASTConsumer::GetOrCreateMLIRFunction(const FunctionDecl *FD,
48294829
NamedAttrList attrs(function->getAttrDictionary());
48304830
attrs.set("llvm.linkage",
48314831
mlir::LLVM::LinkageAttr::get(builder.getContext(), lnk));
4832-
//if (FD->hasAttr<SYCLHalideAttr>() && FD->hasAttr<SYCLKernelAttr>()) {
4833-
// attrs.set("SYCLKernel", mlir::StringAttr::get(builder.getContext(), name));
4834-
//}
4832+
if (/*FD->hasAttr<SYCLHalideAttr>() && */ FD->hasAttr<SYCLKernelAttr>()) {
4833+
attrs.set("SYCLKernel", mlir::StringAttr::get(builder.getContext(), name));
4834+
}
48354835
function->setAttrs(attrs.getDictionary(builder.getContext()));
48364836

48374837
functions[name] = function;
@@ -5059,7 +5059,8 @@ bool MLIRASTConsumer::HandleTopLevelDecl(DeclGroupRef dg) {
50595059

50605060
if ((emitIfFound.count("*") && name != "fpclassify" && !fd->isStatic() &&
50615061
externLinkage) ||
5062-
emitIfFound.count(name) /*|| fd->hasAttr<SYCLHalideAttr>()*/) {
5062+
emitIfFound.count(name) || fd->hasAttr<OpenCLKernelAttr>() ||
5063+
fd->hasAttr<SYCLDeviceAttr>()) {
50635064
functionsToEmit.push_back(fd);
50645065
} else {
50655066
}

‎polygeist/tools/cgeist/Test/Verification/sycl/constructors.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -28,7 +28,7 @@
2828
// CHECK-NEXT: return
2929
// CHECK-NEXT: }
3030

31-
[[intel::halide]] void cons_1() {
31+
SYCL_EXTERNAL void cons_1() {
3232
auto id = sycl::id<2>{};
3333
}
3434

@@ -39,7 +39,7 @@
3939
// CHECK-NEXT: return
4040
// CHECK-NEXT: }
4141

42-
[[intel::halide]] void cons_2(size_t a, size_t b) {
42+
SYCL_EXTERNAL void cons_2(size_t a, size_t b) {
4343
auto id = sycl::id<2>{a, b};
4444
}
4545

@@ -53,7 +53,7 @@
5353
// CHECK-NEXT: return
5454
// CHECK-NEXT: }
5555

56-
[[intel::halide]] void cons_3(sycl::item<2, true> val) {
56+
SYCL_EXTERNAL void cons_3(sycl::item<2, true> val) {
5757
auto id = sycl::id<2>{val};
5858
}
5959

@@ -67,6 +67,6 @@
6767
// CHECK-NEXT: return
6868
// CHECK-NEXT: }
6969

70-
[[intel::halide]] void cons_4(sycl::id<2> val) {
70+
SYCL_EXTERNAL void cons_4(sycl::id<2> val) {
7171
auto id = sycl::id<2>{val};
7272
}

‎polygeist/tools/cgeist/Test/Verification/sycl/functions.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -25,7 +25,7 @@
2525
// CHECK-NEXT: return
2626
// CHECK-NEXT: }
2727

28-
[[intel::halide]] void method_1(sycl::item<2, true> item) {
28+
SYCL_EXTERNAL void method_1(sycl::item<2, true> item) {
2929
auto id = item.get_id(0);
3030
}
3131

@@ -37,7 +37,7 @@
3737
// CHECK-NEXT: return
3838
// CHECK-NEXT: }
3939

40-
[[intel::halide]] void method_2(sycl::item<2, true> item) {
40+
SYCL_EXTERNAL void method_2(sycl::item<2, true> item) {
4141
auto id = item.operator==(item);
4242
}
4343

@@ -54,7 +54,7 @@
5454
// CHECK-NEXT: return
5555
// CHECK-NEXT: }
5656

57-
[[intel::halide]] void op_1(sycl::id<2> a, sycl::id<2> b) {
57+
SYCL_EXTERNAL void op_1(sycl::id<2> a, sycl::id<2> b) {
5858
auto id = a == b;
5959
}
6060

@@ -72,6 +72,6 @@
7272
// CHECK-NEXT: return
7373
// CHECK-NEXT: }
7474

75-
[[intel::halide]] void static_1(sycl::id<2> a, sycl::id<2> b) {
75+
SYCL_EXTERNAL void static_1(sycl::id<2> a, sycl::id<2> b) {
7676
auto abs = sycl::abs(a.get(0) + a.get(1));
7777
}

‎polygeist/tools/cgeist/Test/Verification/sycl/kernels.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,7 @@ class kernel_1 {
2727
public:
2828
kernel_1(sycl::accessor<cl::sycl::cl_int, 1, sycl::access::mode::read_write> A) : A(A) {}
2929

30-
[[intel::halide]] void operator()(sycl::id<1> id) const {
30+
void operator()(sycl::id<1> id) const {
3131
A[id] = 42;
3232
}
3333
};
@@ -57,14 +57,14 @@ void host_2() {
5757
auto buf = sycl::buffer<int, 1>{nullptr, range};
5858
q.submit([&](sycl::handler &cgh) {
5959
auto A = buf.get_access<sycl::access::mode::read_write>(cgh);
60-
cgh.parallel_for<class kernel_2>(range, [=](sycl::id<1> id) [[intel::halide]] {
60+
cgh.parallel_for<class kernel_2>(range, [=](sycl::id<1> id) {
6161
A[id] = 42;
6262
});
6363
});
6464
}
6565
}
6666

6767
// CHECK-NOT: SYCLKernel =
68-
[[intel::halide]] void function_1(sycl::item<2, true> item) {
68+
SYCL_EXTERNAL void function_1(sycl::item<2, true> item) {
6969
auto id = item.get_id(0);
7070
}

‎polygeist/tools/cgeist/Test/Verification/sycl/types.cpp

Lines changed: 14 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -29,56 +29,56 @@
2929

3030
// CHECK: func.func @_Z4id_1N2cl4sycl2idILi1EEE(%arg0: !sycl_id_1_) attributes {llvm.linkage = #llvm.linkage<external>}
3131

32-
[[intel::halide]] void id_1(sycl::id<1> id) {}
32+
SYCL_EXTERNAL void id_1(sycl::id<1> id) {}
3333

3434
// CHECK: func.func @_Z4id_2N2cl4sycl2idILi2EEE(%arg0: !sycl_id_2_) attributes {llvm.linkage = #llvm.linkage<external>}
3535

36-
[[intel::halide]] void id_2(sycl::id<2> id) {}
36+
SYCL_EXTERNAL void id_2(sycl::id<2> id) {}
3737

3838
// CHECK: func.func @_Z5acc_1N2cl4sycl8accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE(%arg0: !sycl_accessor_1_i32_read_write_global_buffer) attributes {llvm.linkage = #llvm.linkage<external>}
3939

40-
[[intel::halide]] void acc_1(sycl::accessor<cl::sycl::cl_int, 1, sycl::access::mode::read_write>) {}
40+
SYCL_EXTERNAL void acc_1(sycl::accessor<cl::sycl::cl_int, 1, sycl::access::mode::read_write>) {}
4141

4242
// CHECK: func.func @_Z5acc_2N2cl4sycl8accessorIiLi2ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE(%arg0: !sycl_accessor_2_i32_read_write_global_buffer) attributes {llvm.linkage = #llvm.linkage<external>}
4343

44-
[[intel::halide]] void acc_2(sycl::accessor<cl::sycl::cl_int, 2, sycl::access::mode::read_write>) {}
44+
SYCL_EXTERNAL void acc_2(sycl::accessor<cl::sycl::cl_int, 2, sycl::access::mode::read_write>) {}
4545

4646
// CHECK: func.func @_Z7range_1N2cl4sycl5rangeILi1EEE(%arg0: !sycl_range_1_) attributes {llvm.linkage = #llvm.linkage<external>}
4747

48-
[[intel::halide]] void range_1(sycl::range<1> range) {}
48+
SYCL_EXTERNAL void range_1(sycl::range<1> range) {}
4949

5050
// CHECK: func.func @_Z7range_2N2cl4sycl5rangeILi2EEE(%arg0: !sycl_range_2_) attributes {llvm.linkage = #llvm.linkage<external>}
5151

52-
[[intel::halide]] void range_2(sycl::range<2> range) {}
52+
SYCL_EXTERNAL void range_2(sycl::range<2> range) {}
5353

5454
// CHECK: func.func @_Z5arr_1N2cl4sycl6detail5arrayILi1EEE(%arg0: !sycl_array_1_) attributes {llvm.linkage = #llvm.linkage<external>}
5555

56-
[[intel::halide]] void arr_1(sycl::detail::array<1> arr) {}
56+
SYCL_EXTERNAL void arr_1(sycl::detail::array<1> arr) {}
5757

5858
// CHECK: func.func @_Z5arr_2N2cl4sycl6detail5arrayILi2EEE(%arg0: !sycl_array_2_) attributes {llvm.linkage = #llvm.linkage<external>}
5959

60-
[[intel::halide]] void arr_2(sycl::detail::array<2> arr) {}
60+
SYCL_EXTERNAL void arr_2(sycl::detail::array<2> arr) {}
6161

6262
// CHECK: func.func @_Z11item_1_trueN2cl4sycl4itemILi1ELb1EEE(%arg0: !sycl_item_1_1_) attributes {llvm.linkage = #llvm.linkage<external>}
6363

64-
[[intel::halide]] void item_1_true(sycl::item<1, true> item) {}
64+
SYCL_EXTERNAL void item_1_true(sycl::item<1, true> item) {}
6565

6666
// CHECK: func.func @_Z12item_2_falseN2cl4sycl4itemILi2ELb0EEE(%arg0: !sycl_item_2_0_) attributes {llvm.linkage = #llvm.linkage<external>}
6767

68-
[[intel::halide]] void item_2_false(sycl::item<2, false> item) {}
68+
SYCL_EXTERNAL void item_2_false(sycl::item<2, false> item) {}
6969

7070
// CHECK: func.func @_Z9nd_item_1N2cl4sycl7nd_itemILi1EEE(%arg0: !sycl_nd_item_1_) attributes {llvm.linkage = #llvm.linkage<external>}
7171

72-
[[intel::halide]] void nd_item_1(sycl::nd_item<1> nd_item) {}
72+
SYCL_EXTERNAL void nd_item_1(sycl::nd_item<1> nd_item) {}
7373

7474
// CHECK: func.func @_Z9nd_item_2N2cl4sycl7nd_itemILi2EEE(%arg0: !sycl_nd_item_2_) attributes {llvm.linkage = #llvm.linkage<external>}
7575

76-
[[intel::halide]] void nd_item_2(sycl::nd_item<2> nd_item) {}
76+
SYCL_EXTERNAL void nd_item_2(sycl::nd_item<2> nd_item) {}
7777

7878
// CHECK: func.func @_Z7group_1N2cl4sycl5groupILi1EEE(%arg0: !sycl_group_1_) attributes {llvm.linkage = #llvm.linkage<external>}
7979

80-
[[intel::halide]] void group_1(sycl::group<1> group) {}
80+
SYCL_EXTERNAL void group_1(sycl::group<1> group) {}
8181

8282
// CHECK: func.func @_Z7group_2N2cl4sycl5groupILi2EEE(%arg0: !sycl_group_2_) attributes {llvm.linkage = #llvm.linkage<external>}
8383

84-
[[intel::halide]] void group_2(sycl::group<2> group) {}
84+
SYCL_EXTERNAL void group_2(sycl::group<2> group) {}

‎polygeist/tools/cgeist/driver.cc

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -645,7 +645,7 @@ int main(int argc, char **argv) {
645645
} else if (ref.startswith("-I")) {
646646
MLIRArgs.push_back("-I");
647647
MLIRArgs.push_back(&argv[i][2]);
648-
} else if (ref == "-fintel-halide") {
648+
} else if (ref == "-fsycl-is-device") {
649649
syclKernelsOnly = true;
650650
MLIRArgs.push_back(argv[i]);
651651
} else if (ref == "-g") {

‎polygeist/tools/cgeist/sycl-clang.py.in

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -22,11 +22,11 @@ def main():
2222
mlir_tool = bin_path + "/cgeist"
2323
clang_tool = bin_path + "/clang++"
2424
arg_file = sys.argv[1]
25-
clang_args = [clang_tool, "-###", "-fsycl", "-fintel-halide"]
25+
clang_args = [clang_tool, "-###", "-fsycl", "-fsycl-device-only", "-D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__"]
2626
clang_args.extend(sys.argv[1:])
2727
clang_res = subprocess.Popen(clang_args, stderr=subprocess.PIPE)
2828
output = clang_res.stderr.readlines()
29-
expanded_clang_args = output[6].decode("utf-8")
29+
expanded_clang_args = output[5].decode("utf-8")
3030
split_output = shlex.split(expanded_clang_args)
3131

3232
mlir_args = [mlir_tool, "-S", "--function=main", arg_file, "--args"]

0 commit comments

Comments
 (0)