Skip to content

[SYCL] Make tests insenstive to dso_local #3037

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Jan 17, 2021
Merged

[SYCL] Make tests insenstive to dso_local #3037

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/accessor_inheritance.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,7 @@ int main() {
// Check kernel parameters
// CHECK: %[[RANGE_TYPE:"struct.*cl::sycl::range"]]
// CHECK: %[[ID_TYPE:"struct.*cl::sycl::id"]]
// CHECK: define spir_kernel void @_ZTSZ4mainE6kernel
// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE6kernel
// CHECK-SAME: i32 [[ARG_A:%[a-zA-Z0-9_]+]],
// CHECK-SAME: i32 [[ARG_B:%[a-zA-Z0-9_]+]],
// CHECK-SAME: i8 addrspace(1)* [[ACC1_DATA:%[a-zA-Z0-9_]+]],
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/address-space-new.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -103,7 +103,7 @@ void test() {
//
Y yy;
baz(yy);
// CHECK: define spir_func void @{{.*}}baz{{.*}}
// CHECK: define {{.*}}spir_func void @{{.*}}baz{{.*}}
// CHECK: %[[FIRST:[a-zA-Z0-9]+]] = bitcast %struct.{{.*}}.Y addrspace(4)* %{{.*}} to i8 addrspace(4)*
// CHECK: %[[OFFSET:[a-zA-Z0-9]+]].ptr = getelementptr inbounds i8, i8 addrspace(4)* %[[FIRST]], i64 8
// CHECK: %[[SECOND:[a-zA-Z0-9]+]] = bitcast i8 addrspace(4)* %[[OFFSET]].ptr to %struct.{{.*}}.HasX addrspace(4)*
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/address-space-of-returns.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,7 @@ A ret_agg() {
A a;
return a;
}
// CHECK: define spir_func void @{{.*}}ret_agg{{.*}}(%struct.{{.*}}.A addrspace(4)* {{.*}} %agg.result)
// CHECK: define {{.*}}spir_func void @{{.*}}ret_agg{{.*}}(%struct.{{.*}}.A addrspace(4)* {{.*}} %agg.result)

template <typename name, typename Func>
__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
Expand Down
16 changes: 8 additions & 8 deletions clang/test/CodeGenSYCL/address-space-parameter-conversions.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,20 +1,20 @@
// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
void bar(int & Data) {}
// CHECK-DAG: define spir_func void @[[RAW_REF:[a-zA-Z0-9_]+]](i32 addrspace(4)* align 4 dereferenceable(4) %
// CHECK-DAG: define {{.*}}spir_func void @[[RAW_REF:[a-zA-Z0-9_]+]](i32 addrspace(4)* align 4 dereferenceable(4) %
void bar2(int & Data) {}
// CHECK-DAG: define spir_func void @[[RAW_REF2:[a-zA-Z0-9_]+]](i32 addrspace(4)* align 4 dereferenceable(4) %
// CHECK-DAG: define {{.*}}spir_func void @[[RAW_REF2:[a-zA-Z0-9_]+]](i32 addrspace(4)* align 4 dereferenceable(4) %
void bar(__attribute__((opencl_local)) int &Data) {}
// CHECK-DAG: define spir_func void [[LOC_REF:@[a-zA-Z0-9_]+]](i32 addrspace(3)* align 4 dereferenceable(4) %
// CHECK-DAG: define {{.*}}spir_func void [[LOC_REF:@[a-zA-Z0-9_]+]](i32 addrspace(3)* align 4 dereferenceable(4) %
void bar3(__attribute__((opencl_global)) int &Data) {}
// CHECK-DAG: define spir_func void @[[GLOB_REF:[a-zA-Z0-9_]+]](i32 addrspace(1)* align 4 dereferenceable(4) %
// CHECK-DAG: define {{.*}}spir_func void @[[GLOB_REF:[a-zA-Z0-9_]+]](i32 addrspace(1)* align 4 dereferenceable(4) %
void foo(int * Data) {}
// CHECK-DAG: define spir_func void @[[RAW_PTR:[a-zA-Z0-9_]+]](i32 addrspace(4)* %
// CHECK-DAG: define {{.*}}spir_func void @[[RAW_PTR:[a-zA-Z0-9_]+]](i32 addrspace(4)* %
void foo2(int * Data) {}
// CHECK-DAG: define spir_func void @[[RAW_PTR2:[a-zA-Z0-9_]+]](i32 addrspace(4)* %
// CHECK-DAG: define {{.*}}spir_func void @[[RAW_PTR2:[a-zA-Z0-9_]+]](i32 addrspace(4)* %
void foo(__attribute__((opencl_local)) int *Data) {}
// CHECK-DAG: define spir_func void [[LOC_PTR:@[a-zA-Z0-9_]+]](i32 addrspace(3)* %
// CHECK-DAG: define {{.*}}spir_func void [[LOC_PTR:@[a-zA-Z0-9_]+]](i32 addrspace(3)* %
void foo3(__attribute__((opencl_global)) int *Data) {}
// CHECK-DAG: define spir_func void @[[GLOB_PTR:[a-zA-Z0-9_]+]](i32 addrspace(1)* %
// CHECK-DAG: define {{.*}}spir_func void @[[GLOB_PTR:[a-zA-Z0-9_]+]](i32 addrspace(1)* %

template<typename T>
void tmpl(T t){}
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/basic-kernel-wrapper.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@ int main() {
return 0;
}

// CHECK: define spir_kernel void @{{.*}}kernel_function
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_function
// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG:%[a-zA-Z0-9_]+]],
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE:%[a-zA-Z0-9_]+_1]],
// CHECK-SAME: %"struct.{{.*}}.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE:%[a-zA-Z0-9_]+_2]],
Expand Down
4 changes: 2 additions & 2 deletions clang/test/CodeGenSYCL/device-functions.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,7 @@ int main() {
kernel_single_task<class fake_kernel>([]() { foo(); });
return 0;
}
// CHECK: define spir_kernel void @_ZTSZ4mainE11fake_kernel()
// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE11fake_kernel()
// CHECK: define internal spir_func void @"_ZZ4mainENK3$_0clEv"(%"class.{{.*}}.anon" addrspace(4)* {{[^,]*}} %this)
// CHECK: define spir_func void @_Z3foov()
// CHECK: define {{.*}}spir_func void @_Z3foov()
// CHECK: define linkonce_odr spir_func i32 @_Z3barIiET_S0_(i32 %arg)
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/device-indirectly-callable-attr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,18 +7,18 @@ void foo() {
helper();
}

// CHECK: define spir_func void @{{.*foo.*}}() #[[ATTRS_INDIR_CALL:[0-9]+]]
// CHECK: define {{.*}}spir_func void @{{.*foo.*}}() #[[ATTRS_INDIR_CALL:[0-9]+]]
// CHECK: call spir_func void @{{.*helper.*}}()
//
// CHECK: define spir_func void @{{.*helper.*}}() #[[ATTRS_NOT_INDIR_CALL:[0-9]+]]
// CHECK: define {{.*}}spir_func void @{{.*helper.*}}() #[[ATTRS_NOT_INDIR_CALL:[0-9]+]]
//

int bar20(int a) { return a + 20; }

class A {
public:
// CHECK-DAG: define linkonce_odr spir_func void @_ZN1A3fooEv{{.*}}#[[ATTRS_INDIR_CALL]]
// CHECK-DAG: define spir_func i32 @_Z5bar20{{.*}}#[[ATTRS_NOT_INDIR_CALL]]
// CHECK-DAG: define {{.*}}spir_func i32 @_Z5bar20{{.*}}#[[ATTRS_NOT_INDIR_CALL]]
[[intel::device_indirectly_callable]] void foo() { bar20(10); }

// CHECK-DAG: define linkonce_odr spir_func void @_ZN1AC1Ev{{.*}}#[[ATTRS_INDIR_CALL_1:[0-9]+]]
Expand Down
3 changes: 1 addition & 2 deletions clang/test/CodeGenSYCL/emit-kernel-in-virtual-func.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,5 +26,4 @@ int main() {
}

// Ensure that the SPIR-Kernel function is actually emitted.
// CHECK: define spir_kernel void @_ZTSZN7DERIVEDIiE10initializeEvE2FF

// CHECK: define {{.*}}spir_kernel void @_ZTSZN7DERIVEDIiE10initializeEvE2FF
3 changes: 1 addition & 2 deletions clang/test/CodeGenSYCL/emit-kernel-in-virtual-func2.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,5 +31,4 @@ int main() {
}

// Ensure that the SPIR-Kernel function is actually emitted.
// CHECK: define spir_kernel void @_ZTSZ2TTIiEvvE2PP

// CHECK: define {{.*}}spir_kernel void @_ZTSZ2TTIiEvvE2PP
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/esimd-accessor-ptr-md.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,7 @@ void test(int val) {
});

// --- Name
// CHECK-LABEL: define spir_kernel void @"_ZTSZZ4testiENK3$_0clERN2cl4sycl7handlerEE12esimd_kernel"(
// CHECK-LABEL: define {{.*}}spir_kernel void @"_ZTSZZ4testiENK3$_0clERN2cl4sycl7handlerEE12esimd_kernel"(
// --- Signature
// CHECK: i32 addrspace(1)* "VCArgumentDesc"="buffer_t" "VCArgumentIOKind"="0" "VCArgumentKind"="2" %_arg_,
// CHECK: i32 "VCArgumentDesc" "VCArgumentIOKind"="0" "VCArgumentKind"="0" %_arg_1,
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/esimd_metadata1.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,7 +15,7 @@ void kernel(const Func &f) __attribute__((sycl_kernel)) {

void bar() {
kernel<class MyKernel>([=]() __attribute__((sycl_explicit_simd)){});
// CHECK: define spir_kernel void @_ZTSZ3barvE8MyKernel() {{.*}} !sycl_explicit_simd ![[EMPTY:[0-9]+]] !intel_reqd_sub_group_size ![[REQD_SIZE:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @_ZTSZ3barvE8MyKernel() {{.*}} !sycl_explicit_simd ![[EMPTY:[0-9]+]] !intel_reqd_sub_group_size ![[REQD_SIZE:[0-9]+]]
}

// CHECK: !spirv.Source = !{[[LANG:![0-9]+]]}
Expand Down
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/esimd_metadata2.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -11,9 +11,9 @@ __attribute__((sycl_device)) void shared_func() { shared_func_decl(); }

__attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func() { shared_func(); }

// CHECK-ESIMD-DAG: define spir_kernel void @{{.*}}kernel_cm() #{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} {{.*}} !intel_reqd_sub_group_size ![[SGSIZE1:[0-9]+]] {{.*}}{
// CHECK-ESIMD-DAG: define spir_func void @{{.*}}esimd_funcv() #{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} {
// CHECK-ESIMD-DAG: define spir_func void @{{.*}}shared_funcv() #{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} {
// CHECK-ESIMD-DAG: define {{.*}}spir_kernel void @{{.*}}kernel_cm() #{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} {{.*}} !intel_reqd_sub_group_size ![[SGSIZE1:[0-9]+]] {{.*}}{
// CHECK-ESIMD-DAG: define {{.*}}spir_func void @{{.*}}esimd_funcv() #{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} {
// CHECK-ESIMD-DAG: define {{.*}}spir_func void @{{.*}}shared_funcv() #{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} {
// CHECK-ESIMD-DAG: define linkonce_odr spir_func void @_ZN12ESIMDFunctorclEv({{.*}}) #{{[0-9]+}} {{.*}} !sycl_explicit_simd !{{[0-9]+}} {
// CHECK-ESIMD-DAG: declare spir_func void @{{.*}}shared_func_declv() #{{[0-9]+}}

Expand Down
4 changes: 2 additions & 2 deletions clang/test/CodeGenSYCL/esimd_metadata3.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,5 +3,5 @@
__attribute__((sycl_device)) void funcWithSpirvIntrin() {}
__attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void standaloneCmFunc() { funcWithSpirvIntrin(); }

// CHECK-ESIMD-DAG: define spir_func void @{{.*}}funcWithSpirvIntrinv() #{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} {
// CHECK-ESIMD-DAG: define spir_func void @{{.*}}standaloneCmFuncv() #{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} {
// CHECK-ESIMD-DAG: define {{.*}}spir_func void @{{.*}}funcWithSpirvIntrinv() #{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} {
// CHECK-ESIMD-DAG: define {{.*}}spir_func void @{{.*}}standaloneCmFuncv() #{{[0-9]+}} !sycl_explicit_simd !{{[0-9]+}} {
12 changes: 6 additions & 6 deletions clang/test/CodeGenSYCL/image_accessor.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,27 +7,27 @@
// RUN: FileCheck < %t.ll --enable-var-scope %s --check-prefix=CHECK-3DWO
//
// CHECK-1DRO: %opencl.image1d_ro_t = type opaque
// CHECK-1DRO: define spir_kernel void @{{.*}}(%opencl.image1d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-1DRO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image1d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-1DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_ro_t addrspace(1)* %{{[0-9]+}})
//
// CHECK-2DRO: %opencl.image2d_ro_t = type opaque
// CHECK-2DRO: define spir_kernel void @{{.*}}(%opencl.image2d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-2DRO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image2d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-2DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_ro_t addrspace(1)* %{{[0-9]+}})
//
// CHECK-3DRO: %opencl.image3d_ro_t = type opaque
// CHECK-3DRO: define spir_kernel void @{{.*}}(%opencl.image3d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-3DRO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image3d_ro_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-3DRO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_ro_t addrspace(1)* %{{[0-9]+}})
//
// CHECK-1DWO: %opencl.image1d_wo_t = type opaque
// CHECK-1DWO: define spir_kernel void @{{.*}}(%opencl.image1d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-1DWO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image1d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-1DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image1d_wo_t addrspace(1)* %{{[0-9]+}})
//
// CHECK-2DWO: %opencl.image2d_wo_t = type opaque
// CHECK-2DWO: define spir_kernel void @{{.*}}(%opencl.image2d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-2DWO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image2d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-2DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image2d_wo_t addrspace(1)* %{{[0-9]+}})
//
// CHECK-3DWO: %opencl.image3d_wo_t = type opaque
// CHECK-3DWO: define spir_kernel void @{{.*}}(%opencl.image3d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-3DWO: define {{.*}}spir_kernel void @{{.*}}(%opencl.image3d_wo_t addrspace(1)* [[IMAGE_ARG:%[a-zA-Z0-9_]+]])
// CHECK-3DWO: call spir_func void @{{.*}}__init{{.*}}(%{{.*}}cl::sycl::accessor{{.*}} %{{[0-9]+}}, %opencl.image3d_wo_t addrspace(1)* %{{[0-9]+}})
//
// TODO: Add tests for the image_array opencl datatype support.
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/inheritance.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,7 +40,7 @@ int main() {
}

// Check kernel paramters
// CHECK: define spir_kernel void @{{.*}}derived(%struct.{{.*}}.base* byval(%struct.{{.*}}.base) align 4 %_arg__base, %struct.{{.*}}.__wrapper_class* byval(%struct.{{.*}}.__wrapper_class) align 8 %_arg_e, i32 %_arg_a)
// CHECK: define {{.*}}spir_kernel void @{{.*}}derived(%struct.{{.*}}.base* byval(%struct.{{.*}}.base) align 4 %_arg__base, %struct.{{.*}}.__wrapper_class* byval(%struct.{{.*}}.__wrapper_class) align 8 %_arg_e, i32 %_arg_a)

// Check alloca for kernel paramters
// CHECK: %[[ARG_A:[a-zA-Z0-9_.]+]] = alloca i32, align 4
Expand Down
18 changes: 9 additions & 9 deletions clang/test/CodeGenSYCL/intel-fpga-ivdep-array.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,7 @@

// Array-specific ivdep - annotate the correspondent GEPs only
//
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_array_no_safelenv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_array_no_safelenv()
void ivdep_array_no_safelen() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand All @@ -18,7 +18,7 @@ void ivdep_array_no_safelen() {
}

// Array-specific ivdep w/ safelen - annotate the correspondent GEPs only
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_array_with_safelenv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_array_with_safelenv()
void ivdep_array_with_safelen() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand All @@ -35,7 +35,7 @@ void ivdep_array_with_safelen() {

// Multiple array-specific ivdeps - annotate the correspondent GEPs
//
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_multiple_arraysv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_multiple_arraysv()
void ivdep_multiple_arrays() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand Down Expand Up @@ -63,7 +63,7 @@ void ivdep_multiple_arrays() {

// Global ivdep with INF safelen & array-specific ivdep with the same safelen
//
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_array_and_globalv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_array_and_globalv()
void ivdep_array_and_global() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand All @@ -81,7 +81,7 @@ void ivdep_array_and_global() {

// Global ivdep with INF safelen & array-specific ivdep with lesser safelen
//
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_array_and_inf_globalv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_array_and_inf_globalv()
void ivdep_array_and_inf_global() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand All @@ -99,7 +99,7 @@ void ivdep_array_and_inf_global() {

// Global ivdep with specified safelen & array-specific ivdep with lesser safelen
//
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_array_and_greater_globalv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_array_and_greater_globalv()
void ivdep_array_and_greater_global() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand All @@ -117,7 +117,7 @@ void ivdep_array_and_greater_global() {

// Global safelen, array-specific safelens
//
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_mul_arrays_and_globalv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_mul_arrays_and_globalv()
void ivdep_mul_arrays_and_global() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand All @@ -138,7 +138,7 @@ void ivdep_mul_arrays_and_global() {
}
}

// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_ptrv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_ptrv()
void ivdep_ptr() {
int *ptr;
// CHECK: %[[PTR:[0-9a-z]+]] = alloca i32 addrspace(4)*
Expand All @@ -149,7 +149,7 @@ void ivdep_ptr() {
// CHECK: br label %for.cond, !llvm.loop ![[MD_LOOP_PTR:[0-9]+]]
}

// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_structv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_structv()
void ivdep_struct() {
struct S {
int *ptr;
Expand Down
12 changes: 6 additions & 6 deletions clang/test/CodeGenSYCL/intel-fpga-ivdep-embedded-loops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,7 @@

// Accesses from the inner loop only, various global safelens for the outer and the inner loops.
//
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_inner_loop_accessv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_inner_loop_accessv()
void ivdep_inner_loop_access() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand All @@ -19,7 +19,7 @@ void ivdep_inner_loop_access() {

// Accesses from both inner and outer loop, same global (INF) safelen for both.
//
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_embedded_global_safelenv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_embedded_global_safelenv()
void ivdep_embedded_global_safelen() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand All @@ -39,7 +39,7 @@ void ivdep_embedded_global_safelen() {

// Accesses from both inner and outer loop, with various safelens per loop.
//
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_embedded_various_safelensv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_embedded_various_safelensv()
void ivdep_embedded_various_safelens() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand All @@ -61,7 +61,7 @@ void ivdep_embedded_various_safelens() {
// Outer loop: array-specific ivdeps for all arrays with various safelens
// Inner loop: global ivdep with its own safelen
//
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_embedded_multiple_arraysv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_embedded_multiple_arraysv()
void ivdep_embedded_multiple_arrays() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand Down Expand Up @@ -90,7 +90,7 @@ void ivdep_embedded_multiple_arrays() {
// As the outer loop's ivdep applies to a particular, other array(s) shouldn't be marked
// into any index group at the outer loop level
//
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_embedded_multiple_arrays_globalv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_embedded_multiple_arrays_globalv()
void ivdep_embedded_multiple_arrays_global() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand All @@ -114,7 +114,7 @@ void ivdep_embedded_multiple_arrays_global() {
}

// Accesses within each dimension of a multi-dimensional (n > 2) loop
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_embedded_multiple_dimensionsv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_embedded_multiple_dimensionsv()
void ivdep_embedded_multiple_dimensions() {
int a[10];
[[intel::ivdep]] for (int i = 0; i != 10; ++i) {
Expand Down
8 changes: 4 additions & 4 deletions clang/test/CodeGenSYCL/intel-fpga-ivdep-global.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,7 @@

// Global ivdep - annotate all GEPs
//
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_no_paramv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_no_paramv()
void ivdep_no_param() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand All @@ -20,7 +20,7 @@ void ivdep_no_param() {
// Global ivdep - annotate all GEPs
// Make sure that ALL of the relevant GEPs for an array are marked into the array's index groups
//
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_no_param_multiple_gepsv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_no_param_multiple_gepsv()
void ivdep_no_param_multiple_geps() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand All @@ -42,7 +42,7 @@ void ivdep_no_param_multiple_geps() {

// Global ivdep w/ safelen specified - annotate all GEPs
//
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_safelenv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_safelenv()
void ivdep_safelen() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand All @@ -59,7 +59,7 @@ void ivdep_safelen() {

// Global ivdep, albeit conflicting safelens - annotate all GEPs
//
// CHECK: define spir_func void @_Z{{[0-9]+}}ivdep_conflicting_safelenv()
// CHECK: define {{.*}}spir_func void @_Z{{[0-9]+}}ivdep_conflicting_safelenv()
void ivdep_conflicting_safelen() {
// CHECK: %[[ARRAY_A:[0-9a-z]+]] = alloca [10 x i32]
int a[10];
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/intel-fpga-mem-builtin.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,7 +12,7 @@ struct State {
// CHECK: [[ANN1:@.str[\.]*[0-9]*]] = {{.*}}{params:384}{cache-size:0}
// CHECK: [[ANN2:@.str[\.]*[0-9]*]] = {{.*}}{params:384}{cache-size:127}

// CHECK: define spir_func void @{{.*}}(float addrspace(4)* %A, i32 addrspace(4)* %B, [[STRUCT]] addrspace(4)* %C, [[STRUCT]] addrspace(4)*{{.*}}%D)
// CHECK: define {{.*}}spir_func void @{{.*}}(float addrspace(4)* %A, i32 addrspace(4)* %B, [[STRUCT]] addrspace(4)* %C, [[STRUCT]] addrspace(4)*{{.*}}%D)
void foo(float *A, int *B, State *C, State &D) {
float *x;
int *y;
Expand Down
Loading