Skip to content

Commit 9899121

Browse files
whitneywhtsangwhitneywhtsang
committed
[SYCL-MLIR] Fix merge
Signed-off-by: Tsang, Whitney <[email protected]>
1 parent a43a776 commit 9899121

File tree

9 files changed

+43
-34
lines changed

9 files changed

+43
-34
lines changed

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -261,7 +261,7 @@ extern "C" SYCL_EXTERNAL void cons_9(const sycl::vec<sycl::cl_char, 3>::vector_t
261261
// CHECK: func.func @[[VEC_INITLIST_VEC_CTR]](%{{.*}}: memref<?x!sycl_vec_i64_16_, 4> {{{.*}}}, %{{.*}}: memref<?x!sycl_vec_i64_8_, 4> {{{.*}}}, %{{.*}}: memref<?x!sycl_vec_i64_4_, 4> {{{.*}}}, %{{.*}}: memref<?x!sycl_vec_i64_2_, 4> {{{.*}}}, %{{.*}}: memref<?xi64, 4> {{{.*}}}, %{{.*}}: memref<?xi64, 4> {{{.*}}}) attributes {[[SPIR_FUNCCC]], [[LINKONCE]], {{.*}}}
262262

263263
// CHECK-LLVM-LABEL: define spir_func void @cons_10(
264-
// CHECK-LLVM-SAME: %"class.sycl::_V1::vec.5" addrspace(4)* noundef align 64 %[[ARG0:.*]], %"class.sycl::_V1::vec.6" addrspace(4)* noundef align 32 %[[ARG1:.*]], %"class.sycl::_V1::vec.7" addrspace(4)* noundef align 16 %[[ARG2:.*]], i64 noundef %{{.*}}, i64 noundef %{{.*}}) #1 {
264+
// CHECK-LLVM-SAME: %"class.sycl::_V1::vec.5" addrspace(4)* noundef align 64 %[[ARG0:.*]], %"class.sycl::_V1::vec.6" addrspace(4)* noundef align 32 %[[ARG1:.*]], %"class.sycl::_V1::vec.7" addrspace(4)* noundef align 16 %[[ARG2:.*]], i64 noundef %{{.*}}, i64 noundef %{{.*}}) #[[FUNCATTRS]] {
265265
// CHECK-LLVM: call spir_func void @[[VEC_INITLIST_VEC_CTR:.*]](%"class.sycl::_V1::vec.4" addrspace(4)* %{{.*}}, %"class.sycl::_V1::vec.5" addrspace(4)* %[[ARG0]], %"class.sycl::_V1::vec.6" addrspace(4)* %[[ARG1]], %"class.sycl::_V1::vec.7" addrspace(4)* %[[ARG2]], i64 addrspace(4)* %{{.*}}, i64 addrspace(4)* %{{.*}})
266266
// CHECK-LLVM: define linkonce_odr spir_func void @[[VEC_INITLIST_VEC_CTR]](%"class.sycl::_V1::vec.4" addrspace(4)* noundef align 128 %{{.*}}, %"class.sycl::_V1::vec.5" addrspace(4)* noundef align 64 %{{.*}}, %"class.sycl::_V1::vec.6" addrspace(4)* noundef align 32 %{{.*}}, %"class.sycl::_V1::vec.7" addrspace(4)* noundef align 16 %{{.*}}, i64 addrspace(4)* noundef align 8 %{{.*}}, i64 addrspace(4)* noundef align 8 %{{.*}}) #[[FUNCATTRS]] {
267267

@@ -294,4 +294,4 @@ extern "C" SYCL_EXTERNAL void cons_11() {
294294
}
295295

296296
// Keep at the end.
297-
// CHECK-LLVM: attributes #[[FUNCATTRS]] = { convergent mustprogress norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="{{.*}}/Test/Verification/sycl/constructors.cpp" }
297+
// CHECK-LLVM: attributes #[[FUNCATTRS]] = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="{{.*}}/Test/Verification/sycl/constructors.cpp" }

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -165,7 +165,7 @@ SYCL_EXTERNAL void range_size(sycl::range<2> r) {
165165
// CHECK-LLVM: %{{.*}} = call spir_func %"class.sycl::_V1::range.2" @_ZNK4sycl3_V18nd_rangeILi2EE16get_global_rangeEv(%"class.sycl::_V1::nd_range.2" addrspace(4)* %{{.*}})
166166

167167
// VAL_3 has incorrect type. Issue 7972 open in GitHub to address this.
168-
// CHECK-LLVM: define linkonce_odr spir_func %"class.sycl::_V1::range.2" @_ZNK4sycl3_V18nd_rangeILi2EE16get_global_rangeEv(%"class.sycl::_V1::nd_range.2" addrspace(4)* noundef align 8 [[VAL_0:%.*]]) #1 {
168+
// CHECK-LLVM: define linkonce_odr spir_func %"class.sycl::_V1::range.2" @_ZNK4sycl3_V18nd_rangeILi2EE16get_global_rangeEv(%"class.sycl::_V1::nd_range.2" addrspace(4)* noundef align 8 [[VAL_0:%.*]]) #[[FUNCATTRS]] {
169169
// CHECK-LLVM-NEXT: [[VAL_2:%.*]] = alloca %"class.sycl::_V1::range.2", align 8
170170
// CHECK-LLVM-NEXT: [[VAL_3:%.*]] = getelementptr %"class.sycl::_V1::nd_range.2", %"class.sycl::_V1::nd_range.2" addrspace(4)* [[VAL_0]], i32 0, i32 0
171171
// CHECK-LLVM-NEXT: [[VAL_4:%.*]] = addrspacecast %"class.sycl::_V1::range.2"* [[VAL_2]] to %"class.sycl::_V1::range.2" addrspace(4)*
@@ -728,4 +728,4 @@ SYCL_EXTERNAL void static_1(sycl::id<2> a, sycl::id<2> b) {
728728
auto abs = sycl::abs(a.get(0) + a.get(1));
729729
}
730730

731-
// CHECK-LLVM: attributes #[[FUNCATTRS]] = { convergent mustprogress norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="{{.*}}polygeist/tools/cgeist/Test/Verification/sycl/functions.cpp" }
731+
// CHECK-LLVM: attributes #[[FUNCATTRS]] = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="{{.*}}polygeist/tools/cgeist/Test/Verification/sycl/functions.cpp" }

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

Lines changed: 25 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
// RUN: clang++ -fsycl -fsycl-device-only -O0 -w -emit-mlir %s -o - | FileCheck %s --check-prefix=CHECK-MLIR
2-
// RUN: clang++ -fsycl -fsycl-device-only -O3 -w -S -emit-llvm -fsycl-targets=spir64-unknown-unknown-syclmlir %s -o - | FileCheck %s --check-prefix=CHECK-LLVM
2+
// RUN: clang++ -fsycl -fsycl-device-only -O0 -w -S -emit-llvm -fsycl-targets=spir64-unknown-unknown-syclmlir %s -o - | FileCheck %s --check-prefix=CHECK-LLVM
33

44
#include <sycl/sycl.hpp>
55

@@ -31,26 +31,30 @@
3131

3232
// CHECK-LLVM-LABEL: define weak_odr spir_kernel void @_ZTSN4sycl3_V16detail18RoundedRangeKernelINS0_4itemILi1ELb1EEELi1EZ4testRNS0_5queueEEUlNS0_2idILi1EEEE_EE
3333
// CHECK-LLVM-SAME: (%"class.sycl::_V1::range.1"* noundef byval(%"class.sycl::_V1::range.1") align 8 %0,
34-
// CHECK-LLVM-SAME: { i32 addrspace(1)* }* noundef byval({ i32 addrspace(1)* }) align 8 %1) local_unnamed_addr #0 {
35-
// CHECK-LLVM-NEXT: %3 = addrspacecast %"class.sycl::_V1::range.1"* %0 to %"class.sycl::_V1::range.1" addrspace(4)*
36-
// CHECK-LLVM-NEXT: %4 = getelementptr %"class.sycl::_V1::range.1", %"class.sycl::_V1::range.1" addrspace(4)* %3, i64 0, i32 0, i32 0, i64 0
37-
// CHECK-LLVM-NEXT: %5 = load i64, i64 addrspace(4)* %4, align 8
38-
// CHECK-LLVM-NEXT: %6 = bitcast { i32 addrspace(1)* }* %1 to i32 addrspace(4)**
39-
// CHECK-LLVM-NEXT: %7 = addrspacecast i32 addrspace(4)** %6 to i32 addrspace(4)* addrspace(4)*
40-
// CHECK-LLVM-NEXT: %.val = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* %7, align 8
41-
// CHECK-LLVM-NEXT: %8 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32
42-
// CHECK-LLVM-NEXT: %9 = extractelement <3 x i64> %8, i64 0
43-
// CHECK-LLVM-NEXT: %10 = icmp slt i64 %9, 2147483648
44-
// CHECK-LLVM-NEXT: tail call void @llvm.assume(i1 %10)
45-
// CHECK-LLVM-NEXT: %.not = icmp slt i64 %9, %5
46-
// CHECK-LLVM-NEXT: br i1 %.not, label %11, label %14
47-
// CHECK-LLVM: 11:
48-
// CHECK-LLVM-NEXT: %12 = trunc i64 %9 to i32
49-
// CHECK-LLVM-NEXT: %13 = getelementptr i32, i32 addrspace(4)* %.val, i64 %9
50-
// CHECK-LLVM-NEXT: store i32 %12, i32 addrspace(4)* %13, align 4
51-
// CHECK-LLVM-NEXT: br label %14
52-
// CHECK-LLVM: 14:
53-
// CHECK-LLVM-NEXT: ret void
34+
// CHECK-LLVM-SAME: { i32 addrspace(1)* }* noundef byval({ i32 addrspace(1)* }) align 8 %1) #0 {
35+
// CHECK-LLVM-NEXT: %3 = alloca %"class.sycl::_V1::item.1.true", align 8
36+
// CHECK-LLVM-NEXT: %4 = alloca { i32 addrspace(4)* }, i64 1, align 8
37+
// CHECK-LLVM-NEXT: %5 = alloca %"class.sycl::_V1::range.1", align 8
38+
// CHECK-LLVM-NEXT: %6 = alloca { %"class.sycl::_V1::range.1", { i32 addrspace(4)* } }, align 8
39+
// CHECK-LLVM-NEXT: %7 = getelementptr { %"class.sycl::_V1::range.1", { i32 addrspace(4)* } }, { %"class.sycl::_V1::range.1", { i32 addrspace(4)* } }* %6, i32 0, i32 0
40+
// CHECK-LLVM-NEXT: %8 = addrspacecast %"class.sycl::_V1::range.1"* %5 to %"class.sycl::_V1::range.1" addrspace(4)*
41+
// CHECK-LLVM-NEXT: %9 = addrspacecast %"class.sycl::_V1::range.1"* %0 to %"class.sycl::_V1::range.1" addrspace(4)*
42+
// CHECK-LLVM-NEXT: call spir_func void @_ZN4sycl3_V15rangeILi1EEC1ERKS2_(%"class.sycl::_V1::range.1" addrspace(4)* %8, %"class.sycl::_V1::range.1" addrspace(4)* %9)
43+
// CHECK-LLVM-NEXT: %10 = load %"class.sycl::_V1::range.1", %"class.sycl::_V1::range.1"* %5, align 8
44+
// CHECK-LLVM-NEXT: store %"class.sycl::_V1::range.1" %10, %"class.sycl::_V1::range.1"* %7, align 8
45+
// CHECK-LLVM-NEXT: %11 = getelementptr { %"class.sycl::_V1::range.1", { i32 addrspace(4)* } }, { %"class.sycl::_V1::range.1", { i32 addrspace(4)* } }* %6, i32 0, i32 1
46+
// CHECK-LLVM-NEXT: %12 = bitcast { i32 addrspace(1)* }* %1 to { i32 addrspace(4)* }*
47+
// CHECK-LLVM-NEXT: %13 = addrspacecast { i32 addrspace(4)* }* %4 to { i32 addrspace(4)* } addrspace(4)*
48+
// CHECK-LLVM-NEXT: %14 = addrspacecast { i32 addrspace(4)* }* %12 to { i32 addrspace(4)* } addrspace(4)*
49+
// CHECK-LLVM-NEXT: call spir_func void @_ZZ4testRN4sycl3_V15queueEENUlNS0_2idILi1EEEE_C1ERKS5_({ i32 addrspace(4)* } addrspace(4)* %13, { i32 addrspace(4)* } addrspace(4)* %14)
50+
// CHECK-LLVM-NEXT: %15 = load { i32 addrspace(4)* }, { i32 addrspace(4)* }* %4, align 8
51+
// CHECK-LLVM-NEXT: store { i32 addrspace(4)* } %15, { i32 addrspace(4)* }* %11, align 8
52+
// CHECK-LLVM-NEXT: %16 = call spir_func %"class.sycl::_V1::item.1.true" addrspace(4)* @_ZN4sycl3_V16detail7declptrINS0_4itemILi1ELb1EEEEEPT_v()
53+
// CHECK-LLVM-NEXT: %17 = call spir_func %"class.sycl::_V1::item.1.true" @_ZN4sycl3_V16detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE(%"class.sycl::_V1::item.1.true" addrspace(4)* %16)
54+
// CHECK-LLVM-NEXT: %18 = addrspacecast { %"class.sycl::_V1::range.1", { i32 addrspace(4)* } }* %6 to { %"class.sycl::_V1::range.1", { i32 addrspace(4)* } } addrspace(4)*
55+
// CHECK-LLVM-NEXT: store %"class.sycl::_V1::item.1.true" %17, %"class.sycl::_V1::item.1.true"* %3, align 8
56+
// CHECK-LLVM-NEXT: call spir_func void @_ZNK4sycl3_V16detail18RoundedRangeKernelINS0_4itemILi1ELb1EEELi1EZ4testRNS0_5queueEEUlNS0_2idILi1EEEE_EclES4_({ %"class.sycl::_V1::range.1", { i32 addrspace(4)* } } addrspace(4)* %18, %"class.sycl::_V1::item.1.true"* %3)
57+
// CHECK-LLVM-NEXT: ret void
5458

5559
int test(sycl::queue &q) {
5660
int *x = sycl::malloc_device<int>(10, q);

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -82,4 +82,4 @@ SYCL_EXTERNAL void function_1(sycl::item<2, true> item) {
8282
}
8383

8484
// Keep at the end of the file.
85-
// CHECK-LLVM: attributes #[[FUNCATTRS]] = { convergent mustprogress norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="{{.*}}/polygeist/tools/cgeist/Test/Verification/sycl/kernels.cpp" }
85+
// CHECK-LLVM: attributes #[[FUNCATTRS]] = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="{{.*}}/polygeist/tools/cgeist/Test/Verification/sycl/kernels.cpp" }

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@
77
//
88
// CHECK-MLIR-DAG: gpu.func @_ZTS8kernel_1
99
// CHECK-MLIR-SAME: kernel attributes {llvm.cconv = #llvm.cconv<spir_kernelcc>, llvm.linkage = #llvm.linkage<weak_odr>,
10-
// CHECK-MLIR-SAME: [[PASSTHROUGH:passthrough = \["convergent", "mustprogress", "norecurse", "nounwind", \["frame-pointer", "all"\], \["no-trapping-math", "true"\], \["stack-protector-buffer-size", "8"\], \["sycl-module-id", ".*/polygeist/tools/cgeist/Test/Verification/sycl/kernels_funcs.cpp"\]\]]]} {
10+
// CHECK-MLIR-SAME: [[PASSTHROUGH:passthrough = \["convergent", "mustprogress", "noinline", "norecurse", "nounwind", "optnone", \["frame-pointer", "all"\], \["no-trapping-math", "true"\], \["stack-protector-buffer-size", "8"\], \["sycl-module-id", ".*/polygeist/tools/cgeist/Test/Verification/sycl/kernels_funcs.cpp"\]\]]]} {
1111
// CHECK-MLIR-DAG: gpu.func @_ZTSZZ6host_2vENKUlRN4sycl3_V17handlerEE_clES2_E8kernel_2
1212
// CHECK-MLIR-SAME: kernel attributes {llvm.cconv = #llvm.cconv<spir_kernelcc>, llvm.linkage = #llvm.linkage<weak_odr>, [[PASSTHROUGH]]} {
1313
// CHECK-MLIR-DAG: func.func @_ZN12StoreWrapperIiLi1ELN4sycl3_V16access4modeE1026EEC1ENS1_8accessorIiLi1ELS3_1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS1_3ext6oneapi22accessor_property_listIJEEEEENS1_2idILi1EEERKi
@@ -20,7 +20,7 @@
2020
// CHECK-LLVM-DAG: define weak_odr spir_kernel void @_ZTS8kernel_1({{.*}}) #[[FUNCATTRS1:[0-9]+]]
2121
// CHECK-LLVM-DAG: define weak_odr spir_kernel void @_ZTSZZ6host_2vENKUlRN4sycl3_V17handlerEE_clES2_E8kernel_2({{.*}}) #[[FUNCATTRS1]]
2222

23-
// CHECK-LLVM-DAG: attributes #[[FUNCATTRS1]] = { convergent mustprogress norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="{{.*}}/polygeist/tools/cgeist/Test/Verification/sycl/kernels_funcs.cpp" }
23+
// CHECK-LLVM-DAG: attributes #[[FUNCATTRS1]] = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="{{.*}}/polygeist/tools/cgeist/Test/Verification/sycl/kernels_funcs.cpp" }
2424

2525
template <typename DataT,
2626
int Dimensions = 1,

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -43,4 +43,4 @@ void host_single_task() {
4343
}
4444

4545
// Keep at the end of the file.
46-
// CHECK-LLVM: attributes #[[FUNCATTRS1]] = { convergent mustprogress norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="{{.*}}/polygeist/tools/cgeist/Test/Verification/sycl/sycl_extern_func.cpp" }
46+
// CHECK-LLVM: attributes #[[FUNCATTRS1]] = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="{{.*}}/polygeist/tools/cgeist/Test/Verification/sycl/sycl_extern_func.cpp" }

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -49,7 +49,7 @@
4949
// CHECK-LABEL: func.func @_Z10accessor_1N4sycl3_V18accessorIiLi1ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE(
5050
// CHECK: %arg0: memref<?x!sycl_accessor_1_i32_rw_gb> {llvm.align = 8 : i64, llvm.byval = !sycl_accessor_1_i32_rw_gb, llvm.noundef})
5151
// CHECK-SAME: attributes {[[SPIR_FUNCCC:llvm.cconv = #llvm.cconv<spir_funccc>]], [[LINKEXT:llvm.linkage = #llvm.linkage<external>]],
52-
// CHECK-SAME: [[PASSTHROUGH:passthrough = \["convergent", "mustprogress", "norecurse", "nounwind", \["frame-pointer", "all"\], \["no-trapping-math", "true"\], \["stack-protector-buffer-size", "8"\], \["sycl-module-id", ".*/polygeist/tools/cgeist/Test/Verification/sycl/types.cpp"\]\]]]} {
52+
// CHECK-SAME: [[PASSTHROUGH:passthrough = \["convergent", "mustprogress", "noinline", "norecurse", "nounwind", "optnone", \["frame-pointer", "all"\], \["no-trapping-math", "true"\], \["stack-protector-buffer-size", "8"\], \["sycl-module-id", ".*/polygeist/tools/cgeist/Test/Verification/sycl/types.cpp"\]\]]]} {
5353
SYCL_EXTERNAL void accessor_1(sycl::accessor<sycl::cl_int, 1, sycl::access::mode::read_write, sycl::access::target::global_buffer>) {}
5454

5555
// CHECK-LABEL: func.func @_Z10accessor_2N4sycl3_V18accessorIiLi2ELNS0_6access4modeE1026ELNS2_6targetE2014ELNS2_11placeholderE0ENS0_3ext6oneapi22accessor_property_listIJEEEEE(

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

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -16,13 +16,13 @@
1616
// CHECK-MLIR-NEXT: [[RESULT:%.*]] = arith.addf [[V1]], [[V2]] : f32
1717
// CHECK-MLIR: affine.store [[RESULT]], {{.*}}[0] : memref<?xf32, 4>
1818

19-
// CHECK-LLVM: define internal spir_func void [[FUNC:@_ZZZ21vec_add_device_simpleRSt5arrayIfLm32EES1_S1_ENKUlRN4sycl3_V17handlerEE_clES5_ENKUlNS3_2idILi1EEEE_clES8_]]({{.*}}) #[[FUNCATTRS2:[0-9]+]]
19+
// CHECK-LLVM: define internal spir_func void [[FUNC:@_ZZZ21vec_add_device_simpleRSt5arrayIfLm32EES1_S1_ENKUlRN4sycl3_V17handlerEE_clES5_ENKUlNS3_2idILi1EEEE_clES8_]]({{.*}}) #[[FUNCATTRS:[0-9]+]]
2020
// CHECK-LLVM-DAG: [[V1:%.*]] = load float, float addrspace(4)* {{.*}}, align 4
2121
// CHECK-LLVM-DAG: [[V2:%.*]] = load float, float addrspace(4)* {{.*}}, align 4
2222
// CHECK-LLVM: [[RESULT:%.*]] = fadd float [[V1]], [[V2]]
2323
// CHECK-LLVM: store float [[RESULT]], float addrspace(4)* {{.*}}, align 4
2424

25-
// CHECK-LLVM: define weak_odr spir_kernel void @{{.*}}vec_add_simple({{.*}}) #[[FUNCATTRS1:[0-9]+]]
25+
// CHECK-LLVM: define weak_odr spir_kernel void @{{.*}}vec_add_simple({{.*}}) #[[FUNCATTRS]]
2626
// CHECK-LLVM: call spir_func void [[FUNC]]
2727

2828
void vec_add_device_simple(std::array<float, N> &VA, std::array<float, N> &VB,
@@ -95,5 +95,4 @@ int main() {
9595
}
9696

9797
// Keep at the end of the file.
98-
// CHECK-LLVM-DAG: attributes #[[FUNCATTRS1]] = { convergent mustprogress norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="{{.*}}/polygeist/tools/cgeist/Test/Verification/sycl/vec_add.cpp" }
99-
// CHECK-LLVM-DAG: attributes #[[FUNCATTRS2]] = { convergent inlinehint mustprogress norecurse nounwind "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="{{.*}}/polygeist/tools/cgeist/Test/Verification/sycl/vec_add.cpp" }
98+
// CHECK-LLVM: attributes #[[FUNCATTRS]] = { convergent mustprogress noinline norecurse nounwind optnone "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="{{.*}}/polygeist/tools/cgeist/Test/Verification/sycl/vec_add.cpp" }

‎polygeist/tools/cgeist/driver.cc

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -916,6 +916,7 @@ splitCommandLineOptions(int argc, char **argv,
916916
SmallVector<const char *> &MLIRArgs) {
917917
bool SYCLIsDevice = false;
918918
bool LinkOnly = false;
919+
bool ClangOption = false;
919920
llvm::OptimizationLevel OptimizationLevel =
920921
CgeistOptions::getDefaultOptimizationLevel();
921922

@@ -946,9 +947,14 @@ splitCommandLineOptions(int argc, char **argv,
946947
} else if (Ref == "-fsycl-is-device") {
947948
SYCLIsDevice = true;
948949
MLIRArgs.push_back(argv[I]);
950+
} else if (Ref == "--args") {
951+
ClangOption = true;
952+
MLIRArgs.push_back(argv[I]);
949953
} else if (Ref.consume_front("-O") || Ref.consume_front("--optimize")) {
950954
// If several flags are passed, we keep the last one.
951955
OptimizationLevel = ExitOnErr(parseOptimizationLevel(Ref));
956+
if (ClangOption)
957+
MLIRArgs.push_back(argv[I]);
952958
LinkageArgs.push_back(argv[I]);
953959
} else if (Ref == "-g")
954960
LinkageArgs.push_back(argv[I]);

0 commit comments

Comments
 (0)