From e1ab75d0bebd1101664d5c1771fb3fa25f7bbc37 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Thu, 13 May 2021 16:49:40 +0300 Subject: [PATCH 1/9] [SYCL] Disable inlining in 'ID queries fit in int' test Signed-off-by: Sergey Kanaev --- sycl/test/check_device_code/id_queries_fit_int.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/check_device_code/id_queries_fit_int.cpp b/sycl/test/check_device_code/id_queries_fit_int.cpp index 6736ff6c24c7d..35d9fad4e9f49 100644 --- a/sycl/test/check_device_code/id_queries_fit_int.cpp +++ b/sycl/test/check_device_code/id_queries_fit_int.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl -fsycl-id-queries-fit-in-int -O0 -c -S -emit-llvm -o %t.ll %s +// RUN: %clangxx -fsycl -Xclang -fsycl-is-host -fsycl-id-queries-fit-in-int -fno-inline-functions -O1 -c -S -emit-llvm -o %t.ll %s // RUN: FileCheck %s --input-file %t.ll #include From 652dc2d17167b116ab0e8d138ad82f65c7ced3f0 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Mon, 17 May 2021 17:43:17 +0300 Subject: [PATCH 2/9] Fix typo Signed-off-by: Sergey Kanaev --- sycl/test/check_device_code/id_queries_fit_int.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/check_device_code/id_queries_fit_int.cpp b/sycl/test/check_device_code/id_queries_fit_int.cpp index 35d9fad4e9f49..27c19884d3dea 100644 --- a/sycl/test/check_device_code/id_queries_fit_int.cpp +++ b/sycl/test/check_device_code/id_queries_fit_int.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl -Xclang -fsycl-is-host -fsycl-id-queries-fit-in-int -fno-inline-functions -O1 -c -S -emit-llvm -o %t.ll %s +// RUN: %clangxx -fsycl -Xclang -fsycl-id-queries-fit-in-int -fno-inline-functions -O1 -c -S -emit-llvm -o %t.ll %s // RUN: FileCheck %s --input-file %t.ll #include From 46723d3b8842f6a4f67904401cc559c259325b43 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Tue, 18 May 2021 18:51:36 +0300 Subject: [PATCH 3/9] Leave -O0 Signed-off-by: Sergey Kanaev --- sycl/test/check_device_code/id_queries_fit_int.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/check_device_code/id_queries_fit_int.cpp b/sycl/test/check_device_code/id_queries_fit_int.cpp index 27c19884d3dea..0f630881d0d01 100644 --- a/sycl/test/check_device_code/id_queries_fit_int.cpp +++ b/sycl/test/check_device_code/id_queries_fit_int.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl -Xclang -fsycl-id-queries-fit-in-int -fno-inline-functions -O1 -c -S -emit-llvm -o %t.ll %s +// RUN: %clangxx -fsycl -Xclang -fsycl-id-queries-fit-in-int -fno-inline-functions -O0 -c -S -emit-llvm -o %t.ll %s // RUN: FileCheck %s --input-file %t.ll #include From e541be3381ea91411309a07dfaaae9502ad6a184 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Tue, 18 May 2021 18:52:01 +0300 Subject: [PATCH 4/9] Replace -fno-inline-functions with -disable-llvm-passes Signed-off-by: Sergey Kanaev --- sycl/test/check_device_code/id_queries_fit_int.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/check_device_code/id_queries_fit_int.cpp b/sycl/test/check_device_code/id_queries_fit_int.cpp index 0f630881d0d01..6612db1b1ae8c 100644 --- a/sycl/test/check_device_code/id_queries_fit_int.cpp +++ b/sycl/test/check_device_code/id_queries_fit_int.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl -Xclang -fsycl-id-queries-fit-in-int -fno-inline-functions -O0 -c -S -emit-llvm -o %t.ll %s +// RUN: %clangxx -fsycl -Xclang -fsycl-id-queries-fit-in-int -disable-llvm-passes -O0 -c -S -emit-llvm -o %t.ll %s // RUN: FileCheck %s --input-file %t.ll #include From b4eccf68c169aac4920a07cd204423d0455e9820 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 19 May 2021 12:57:01 +0300 Subject: [PATCH 5/9] Remove redundant argument Signed-off-by: Sergey Kanaev --- sycl/test/check_device_code/id_queries_fit_int.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/check_device_code/id_queries_fit_int.cpp b/sycl/test/check_device_code/id_queries_fit_int.cpp index 6612db1b1ae8c..f5f62c4004e2c 100644 --- a/sycl/test/check_device_code/id_queries_fit_int.cpp +++ b/sycl/test/check_device_code/id_queries_fit_int.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl -Xclang -fsycl-id-queries-fit-in-int -disable-llvm-passes -O0 -c -S -emit-llvm -o %t.ll %s +// RUN: %clangxx -fsycl -fsycl-id-queries-fit-in-int -disable-llvm-passes -O0 -c -S -emit-llvm -o %t.ll %s // RUN: FileCheck %s --input-file %t.ll #include From c44a25ca679e6e434107139295fdbaa40f03e8e4 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 26 May 2021 14:29:22 +0300 Subject: [PATCH 6/9] Make code the device one Signed-off-by: Sergey Kanaev --- .../check_device_code/id_queries_fit_int.cpp | 81 ++++++++++--------- 1 file changed, 43 insertions(+), 38 deletions(-) diff --git a/sycl/test/check_device_code/id_queries_fit_int.cpp b/sycl/test/check_device_code/id_queries_fit_int.cpp index f5f62c4004e2c..5b0f3dd66a468 100644 --- a/sycl/test/check_device_code/id_queries_fit_int.cpp +++ b/sycl/test/check_device_code/id_queries_fit_int.cpp @@ -1,51 +1,56 @@ -// RUN: %clangxx -fsycl -fsycl-id-queries-fit-in-int -disable-llvm-passes -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 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(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( + 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: } From 6a81ac2547c7165008c34472ee2905c5480c0c56 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 26 May 2021 14:37:05 +0300 Subject: [PATCH 7/9] Fix style issue Signed-off-by: Sergey Kanaev --- sycl/test/check_device_code/id_queries_fit_int.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/check_device_code/id_queries_fit_int.cpp b/sycl/test/check_device_code/id_queries_fit_int.cpp index 5b0f3dd66a468..a7d8f9b729a77 100644 --- a/sycl/test/check_device_code/id_queries_fit_int.cpp +++ b/sycl/test/check_device_code/id_queries_fit_int.cpp @@ -49,7 +49,7 @@ int main() { int LocalIdConverted = TestNDItem.get_local_id(); // CHECK: call void @llvm.assume(i1 {{.*}}) int OffsetConferted = TestNDItem.get_offset(); - }); + }); }); return 0; From af1f0ef8997c3de75d56d35da458caee4cdcb921 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 26 May 2021 14:52:31 +0300 Subject: [PATCH 8/9] Eliminate use of runtime queue and macro to disable parallel for range rounding Signed-off-by: Sergey Kanaev --- .../check_device_code/id_queries_fit_int.cpp | 85 +++++++++---------- 1 file changed, 38 insertions(+), 47 deletions(-) diff --git a/sycl/test/check_device_code/id_queries_fit_int.cpp b/sycl/test/check_device_code/id_queries_fit_int.cpp index a7d8f9b729a77..498e00eb99bcc 100644 --- a/sycl/test/check_device_code/id_queries_fit_int.cpp +++ b/sycl/test/check_device_code/id_queries_fit_int.cpp @@ -1,56 +1,47 @@ -// 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: %clangxx -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 using namespace sycl; -int main() { - queue Q; - - // CHECK: define {{.*}}dso_local spir_kernel void @"{{.*}}Test"() - Q.submit([&](handler &H) { - H.parallel_for(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: define {{.*}}dso_local spir_kernel void @"{{.*}}NDTest"() - Q.submit([&](handler &H) { - H.parallel_for( - 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); +// CHECK: define {{.*}}dso_local spir_func void @{{.*}}testItem{{.*}}(%"class.{{.*}}item"*{{.*}}%TestItem) +SYCL_EXTERNAL void testItem(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 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: define {{.*}}dso_local spir_func void @{{.*}}testNDItem{{.*}}(%"class.{{.*}}nd_item"*{{.*}}%TestNDItem) +SYCL_EXTERNAL void testNDItem(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); - return 0; + // 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(); } + +int main() { return 0; } From 1b5ca990204682ca98713c8547f881868faa1349 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 26 May 2021 15:55:32 +0300 Subject: [PATCH 9/9] Removed main Signed-off-by: Sergey Kanaev --- sycl/test/check_device_code/id_queries_fit_int.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/sycl/test/check_device_code/id_queries_fit_int.cpp b/sycl/test/check_device_code/id_queries_fit_int.cpp index 498e00eb99bcc..266219d0cb7bd 100644 --- a/sycl/test/check_device_code/id_queries_fit_int.cpp +++ b/sycl/test/check_device_code/id_queries_fit_int.cpp @@ -43,5 +43,3 @@ SYCL_EXTERNAL void testNDItem(nd_item<1> TestNDItem) { // CHECK: call void @llvm.assume(i1 {{.*}}) int OffsetConferted = TestNDItem.get_offset(); } - -int main() { return 0; }