Skip to content

Commit

Permalink
[NFC][SYCL] Improve kernel metadata test (#1610)
Browse files Browse the repository at this point in the history
There was a bug in kernel_arg_addr_space metadata generation
(SYCL addr spaces were ignored), that was fixed in ac94b1e.
Expand the test accordingly.

Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>
  • Loading branch information
MrSidims authored Apr 30, 2020
1 parent c8d86a8 commit 4dd874b
Showing 1 changed file with 12 additions and 8 deletions.
20 changes: 12 additions & 8 deletions clang/test/CodeGenSYCL/kernel-metadata.cpp
Original file line number Diff line number Diff line change
@@ -1,14 +1,18 @@
// 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]] = !{!"", !"", !"", !""}

template <typename name, typename Func>
__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
kernelFunc();
}
#include "sycl.hpp"

int main() {
kernel_single_task<class kernel_function>([]() {});
cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write> accessorA;
cl::sycl::kernel_single_task<class kernel_function>(
[=]() {
accessorA.use();
});
return 0;
}

0 comments on commit 4dd874b

Please sign in to comment.