Skip to content

[Test][SYCL] Add test to verify SPIR kernel argument names #5819

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 2 commits into from
Mar 17, 2022
Merged

[Test][SYCL] Add test to verify SPIR kernel argument names #5819

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
21cfc12
[NFC][SYCL] Add test to verify openCL kernel argument names
elizabethandrews Mar 15, 2022
e3a1dc1
Implement review comments
elizabethandrews Mar 16, 2022
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
74 changes: 74 additions & 0 deletions clang/test/CodeGenSYCL/save-user-names.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,74 @@
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s

// Test to verify that user specified names are retained in SPIR
// kernel argument names.

#include "sycl.hpp"

struct NestedSimple {
int NestedSimpleField;
};

struct NestedComplex {
int NestedComplexField;
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> NestedAccField;
};

struct KernelFunctor {
int IntField;
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> AccField1;
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> AccField2;
NestedSimple NestedSimpleObj;
NestedComplex NestedComplexObj;
void operator()() const {}
};

int main() {
cl::sycl::queue q;

q.submit([&](cl::sycl::handler &cgh) {
KernelFunctor FunctorObj;
cgh.single_task<class Kernel1>(FunctorObj);
});

q.submit([&](cl::sycl::handler &cgh) {
int Data;
NestedSimple NestedSimpleObj;
NestedComplex NestedComplexObj;
cl::sycl::accessor<char, 1, cl::sycl::access::mode::read> CapturedAcc1, CapturedAcc2;
cgh.single_task<class Kernel2>([=]() {
Data;
CapturedAcc1;
CapturedAcc2;
NestedSimpleObj;
NestedComplexObj;
});
});

return 0;
}

// Check kernel parameters generated when kernel is defined as Functor

// NOTE: Accessor fields have 4 corresponding openCL kernel arguments. When
// the compiler generates the openCL kernel arguments, they are generated
// with the same name (i.e the user-specified name for accessor). Since LLVM
// IR cannot have same names, a number is appended in IR.
//
// CHECK: define {{.*}}spir_kernel void @{{.*}}Kernel1
// CHECK-SAME: %_arg_IntField
// CHECK-SAME: %_arg_AccField1{{.*}}%_arg_AccField11{{.*}}%_arg_AccField12{{.*}}%_arg_AccField13
// CHECK-SAME: %_arg_AccField2{{.*}}%_arg_AccField24{{.*}}%_arg_AccField25{{.*}}%_arg_AccField26
// CHECK-SAME: %_arg_NestedSimpleObj
// CHECK-SAME: %_arg_NestedComplexField
// CHECK-SAME: %_arg_NestedAccField{{.*}}%_arg_NestedAccField7{{.*}}%_arg_NestedAccField8{{.*}}%_arg_NestedAccField9

// Check kernel parameters generated when kernel is defined as Lambda
//
// CHECK: define {{.*}}spir_kernel void @{{.*}}Kernel2
// CHECK-SAME: %_arg_Data
// CHECK-SAME: %_arg_CapturedAcc1{{.*}}%_arg_CapturedAcc11{{.*}}%_arg_CapturedAcc12{{.*}}%_arg_CapturedAcc13
// CHECK-SAME: %_arg_CapturedAcc2{{.*}}%_arg_CapturedAcc24{{.*}}%_arg_CapturedAcc25{{.*}}%_arg_CapturedAcc26
// CHECK-SAME: %_arg_NestedSimpleObj
// CHECK-SAME: %_arg_NestedComplexField
// CHECK-SAME: %_arg_NestedAccField{{.*}}%_arg_NestedAccField7{{.*}}%_arg_NestedAccField8{{.*}}%_arg_NestedAccField9