From d70a1d4bececc7bbede343f445dae7a6540f43a8 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Fri, 24 Apr 2020 12:26:34 +0300 Subject: [PATCH 1/3] [NFC][SYCL] Improve kernel metadata test There was a bug in kernel_arg_addr_space metadata generation (SYCL addr spaces were ignored), that was fixed in ac94b1eb537. Expand the test accordingly. Signed-off-by: Dmitry Sidorov --- clang/test/CodeGenSYCL/kernel-metadata.cpp | 17 +++++++++++++---- 1 file changed, 13 insertions(+), 4 deletions(-) diff --git a/clang/test/CodeGenSYCL/kernel-metadata.cpp b/clang/test/CodeGenSYCL/kernel-metadata.cpp index 7e07220663868..c6224066753ae 100644 --- a/clang/test/CodeGenSYCL/kernel-metadata.cpp +++ b/clang/test/CodeGenSYCL/kernel-metadata.cpp @@ -1,7 +1,12 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -emit-llvm %s -o - | FileCheck %s -// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function() {{[^{]+}} !kernel_arg_addr_space ![[MD:[0-9]+]] !kernel_arg_access_qual ![[MD]] !kernel_arg_type ![[MD]] !kernel_arg_base_type ![[MD]] !kernel_arg_type_qual ![[MD]] -// CHECK: ![[MD]] = !{} +// CHECK: define {{.*}}spir_kernel void @_ZTSZ4mainE15kernel_function{{.*}} !kernel_arg_addr_space ![[MDAS:[0-9]+]] !kernel_arg_access_qual ![[MDAC:[0-9]+]] !kernel_arg_type ![[MDAT:[0-9]+]] !kernel_arg_base_type ![[MDAT:[0-9]+]] !kernel_arg_type_qual ![[MDATQ:[0-9]+]] +// CHECK: ![[MDAS]] = !{i32 1, i32 0, i32 0, i32 0} +// CHECK: ![[MDAC]] = !{!"none", !"none", !"none", !"none"} +// CHECK: ![[MDAT]] = !{!"int*", !"cl::sycl::range<1>", !"cl::sycl::range<1>", !"cl::sycl::id<1>"} +// CHECK: ![[MDATQ]] = !{!"", !"", !"", !""} + +#include "sycl.hpp" template __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { @@ -9,6 +14,10 @@ __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { } int main() { - kernel_single_task([]() {}); + cl::sycl::accessor accessorA; + kernel_single_task( + [=]() { + accessorA.use(); + }); return 0; } From 42bb3042db1fa4ca8c299e1e7a7eb776e4ce8442 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Wed, 29 Apr 2020 13:51:24 +0300 Subject: [PATCH 2/3] Fix clang-format concerns Signed-off-by: Dmitry Sidorov --- clang/test/CodeGenSYCL/kernel-metadata.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/test/CodeGenSYCL/kernel-metadata.cpp b/clang/test/CodeGenSYCL/kernel-metadata.cpp index c6224066753ae..e9a4d7b6f934f 100644 --- a/clang/test/CodeGenSYCL/kernel-metadata.cpp +++ b/clang/test/CodeGenSYCL/kernel-metadata.cpp @@ -16,8 +16,8 @@ __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { int main() { cl::sycl::accessor accessorA; kernel_single_task( - [=]() { - accessorA.use(); - }); + [=]() { + accessorA.use(); + }); return 0; } From 0554515b09d13c82ae6a664d9e49264618546dc8 Mon Sep 17 00:00:00 2001 From: Dmitry Sidorov Date: Wed, 29 Apr 2020 13:54:21 +0300 Subject: [PATCH 3/3] Remove unnecessary kernel definition Signed-off-by: Dmitry Sidorov --- clang/test/CodeGenSYCL/kernel-metadata.cpp | 7 +------ 1 file changed, 1 insertion(+), 6 deletions(-) diff --git a/clang/test/CodeGenSYCL/kernel-metadata.cpp b/clang/test/CodeGenSYCL/kernel-metadata.cpp index e9a4d7b6f934f..dd502fa1dd844 100644 --- a/clang/test/CodeGenSYCL/kernel-metadata.cpp +++ b/clang/test/CodeGenSYCL/kernel-metadata.cpp @@ -8,14 +8,9 @@ #include "sycl.hpp" -template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { - kernelFunc(); -} - int main() { cl::sycl::accessor accessorA; - kernel_single_task( + cl::sycl::kernel_single_task( [=]() { accessorA.use(); });