Skip to content

[SYCL] Disable inlining in 'ID queries fit in int' test #3751

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 10 commits into from
May 27, 2021
+7 −13
Merged

[SYCL] Disable inlining in 'ID queries fit in int' test #3751

Changes from 8 commits
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
e1ab75d
[SYCL] Disable inlining in 'ID queries fit in int' test
May 13, 2021
652dc2d
Fix typo
May 17, 2021
46723d3
Leave -O0
May 18, 2021
e541be3
Replace -fno-inline-functions with -disable-llvm-passes
May 18, 2021
b4eccf6
Remove redundant argument
May 19, 2021
8bdf821
Merge remote-tracking branch 'public/sycl' into fix-test
May 26, 2021
c44a25c
Make code the device one
May 26, 2021
6a81ac2
Fix style issue
May 26, 2021
af1f0ef
Eliminate use of runtime queue and macro to disable parallel for rang…
May 26, 2021
1b5ca99
Removed main
May 26, 2021
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
81 changes: 43 additions & 38 deletions sycl/test/check_device_code/id_queries_fit_int.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,51 +1,56 @@
// RUN: %clangxx -fsycl -fsycl-id-queries-fit-in-int -O0 -c -S -emit-llvm -o %t.ll %s
// RUN: %clangxx -D__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ -fsycl-device-only -fsycl-id-queries-fit-in-int -fno-sycl-early-optimizations -S -o %t.ll %s
// RUN: FileCheck %s --input-file %t.ll

#include <CL/sycl.hpp>

using namespace sycl;

// CHECK: define dso_local i32 @main() {{.*}} {
int main() {
item<1, true> TestItem = detail::Builder::createItem<1, true>({3}, {2}, {1});
// CHECK: call void @llvm.assume(i1 {{.*}})
int Id = TestItem.get_id(0);
// CHECK: call void @llvm.assume(i1 {{.*}})
int Range = TestItem.get_range(0);
// CHECK: call void @llvm.assume(i1 {{.*}})
int LinearId = TestItem.get_linear_id();
queue Q;

cl::sycl::nd_item<1> TestNDItem =
detail::Builder::createNDItem<1>(detail::Builder::createItem<1, false>({4}, {2}),
detail::Builder::createItem<1, false>({2}, {0}),
detail::Builder::createGroup<1>({4}, {2}, {1}));
// CHECK: define {{.*}}dso_local spir_kernel void @"{{.*}}Test"()
Q.submit([&](handler &H) {
H.parallel_for<class Test>(range<1>{3}, [=](item<1> TestItem) {
// CHECK: call void @llvm.assume(i1 {{.*}})
int Id = TestItem.get_id(0);
// CHECK: call void @llvm.assume(i1 {{.*}})
int Range = TestItem.get_range(0);
// CHECK: call void @llvm.assume(i1 {{.*}})
int LinearId = TestItem.get_linear_id();
});
});

// CHECK: call void @llvm.assume(i1 {{.*}})
int GlobalId = TestNDItem.get_global_id(0);
// CHECK: call void @llvm.assume(i1 {{.*}})
int GlobalLinearId = TestNDItem.get_global_linear_id();
// CHECK: call void @llvm.assume(i1 {{.*}})
int LocalId = TestNDItem.get_local_id(0);
// CHECK: call void @llvm.assume(i1 {{.*}})
int LocalLinearId = TestNDItem.get_local_linear_id();
// CHECK: call void @llvm.assume(i1 {{.*}})
int GroupRange = TestNDItem.get_group_range(0);
// CHECK: call void @llvm.assume(i1 {{.*}})
int GroupId = TestNDItem.get_group(0);
// CHECK: call void @llvm.assume(i1 {{.*}})
int GroupLinearId = TestNDItem.get_group_linear_id();
// CHECK: call void @llvm.assume(i1 {{.*}})
int GlobalRange = TestNDItem.get_global_range(0);
// CHECK: call void @llvm.assume(i1 {{.*}})
int LocalRange = TestNDItem.get_local_range(0);
// CHECK: define {{.*}}dso_local spir_kernel void @"{{.*}}NDTest"()
Q.submit([&](handler &H) {
H.parallel_for<class NDTest>(
nd_range<1>{range<1>{3}, range<1>{1}}, [=](nd_item<1> TestNDItem) {
// CHECK: call void @llvm.assume(i1 {{.*}})
int GlobalId = TestNDItem.get_global_id(0);
// CHECK: call void @llvm.assume(i1 {{.*}})
int GlobalLinearId = TestNDItem.get_global_linear_id();
// CHECK: call void @llvm.assume(i1 {{.*}})
int LocalId = TestNDItem.get_local_id(0);
// CHECK: call void @llvm.assume(i1 {{.*}})
int LocalLinearId = TestNDItem.get_local_linear_id();
// CHECK: call void @llvm.assume(i1 {{.*}})
int GroupRange = TestNDItem.get_group_range(0);
// CHECK: call void @llvm.assume(i1 {{.*}})
int GroupId = TestNDItem.get_group(0);
// CHECK: call void @llvm.assume(i1 {{.*}})
int GroupLinearId = TestNDItem.get_group_linear_id();
// CHECK: call void @llvm.assume(i1 {{.*}})
int GlobalRange = TestNDItem.get_global_range(0);
// CHECK: call void @llvm.assume(i1 {{.*}})
int LocalRange = TestNDItem.get_local_range(0);

int GlobalIdConverted = TestNDItem.get_global_id();
// CHECK: call void @llvm.assume(i1 {{.*}})
int LocalIdConverted = TestNDItem.get_local_id();
// CHECK: call void @llvm.assume(i1 {{.*}})
int OffsetConferted = TestNDItem.get_offset();
// CHECK: call void @llvm.assume(i1 {{.*}})
// CHECK: call void @llvm.assume(i1 {{.*}})
int GlobalIdConverted = TestNDItem.get_global_id();
// CHECK: call void @llvm.assume(i1 {{.*}})
int LocalIdConverted = TestNDItem.get_local_id();
// CHECK: call void @llvm.assume(i1 {{.*}})
int OffsetConferted = TestNDItem.get_offset();
});
});

return 0;
}
// CHECK: }