Skip to content

Commit 8436827

Browse files
HPS-1HPS-1
committed
[SYCL] Revise VirtualFunctions Tests
Signed-off-by: Hu, Peisen <[email protected]>
1 parent 37dc53b commit 8436827

File tree

5 files changed

+89
-97
lines changed

5 files changed

+89
-97
lines changed

‎sycl/test-e2e/VirtualFunctions/2/1/1/missing-overrides.cpp‎

Lines changed: 18 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -66,6 +66,21 @@ void applyOp(int *DataPtr, Base *ObjPtr) {
6666
ObjPtr->multiply(DataPtr);
6767
}
6868

69+
template <typename T1, typename T2> struct KernelFunctor {
70+
T1 mStorageAcc;
71+
T2 mDataAcc;
72+
KernelFunctor(T1 StorageAcc, T2 DataAcc)
73+
: mStorageAcc(StorageAcc), mDataAcc(DataAcc) {}
74+
75+
void operator()() const {
76+
auto *Ptr = mStorageAcc[0].construct</* ret type = */ Base>(TestCase);
77+
applyOp(mDataAcc.get_multi_ptr<sycl::access::decorated::no>().get(), Ptr);
78+
}
79+
auto get(oneapi::properties_tag) const {
80+
return oneapi::properties{oneapi::assume_indirect_calls};
81+
}
82+
};
83+
6984
int main() try {
7085
using storage_t = obj_storage_t<IncrementBy1, IncrementBy1AndSubstractBy2,
7186
MultiplyBy2, MultiplyBy2AndIncrementBy8,
@@ -80,31 +95,15 @@ int main() try {
8095

8196
sycl::queue q(asyncHandler);
8297

83-
struct KernelFunctor {
84-
sycl::buffer<storage_t> mDeviceStorage;
85-
sycl::buffer<int> mDataStorage;
86-
sycl::handler mCGH;
87-
KernelFunctor(sycl::buffer<storage_t> DeviceStorage,
88-
sycl::buffer<int> DataStorage, sycl::handler CGH)
89-
: mDeviceStorage(DeviceStorage), mDataStorage(DataStorage), mCGH(CGH) {}
90-
91-
void operator()() const {
92-
sycl::accessor StorageAcc(mDeviceStorage, mCGH, sycl::write_only);
93-
sycl::accessor DataAcc(mDataStorage, mCGH, sycl::write_only);
94-
auto *Ptr = StorageAcc[0].construct</* ret type = */ Base>(TestCase);
95-
applyOp(DataAcc.get_multi_ptr<sycl::access::decorated::no>().get(), Ptr);
96-
}
97-
auto get(oneapi::properties_tag) const {
98-
return oneapi::properties{oneapi::assume_indirect_calls};
99-
}
100-
};
10198
for (unsigned TestCase = 0; TestCase < 6; ++TestCase) {
10299
int HostData = 42;
103100
int Data = HostData;
104101
sycl::buffer<int> DataStorage(&Data, sycl::range{1});
105102

106103
q.submit([&](sycl::handler &CGH) {
107-
CGH.single_task(KernelFunctor(DeviceStorage, DataStorage, CGH));
104+
sycl::accessor StorageAcc(DeviceStorage, CGH, sycl::write_only);
105+
sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only);
106+
CGH.single_task(KernelFunctor(StorageAcc, DataAcc));
108107
});
109108

110109
Base *Ptr = HostStorage.construct</* ret type = */ Base>(TestCase);

‎sycl/test-e2e/VirtualFunctions/2/1/1/more-complex-hierarchy.cpp‎

Lines changed: 18 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -45,6 +45,21 @@ class IncrementBy8 : public IncrementOp {
4545

4646
void applyOp(int *Data, AbstractOp *Obj) { Obj->applyOp(Data); }
4747

48+
template <typename T1, typename T2> struct KernelFunctor {
49+
T1 mStorageAcc;
50+
T2 mDataAcc;
51+
KernelFunctor(T1 &StorageAcc, T2 &DataAcc)
52+
: mStorageAcc(StorageAcc), mDataAcc(DataAcc) {}
53+
54+
void operator()() const {
55+
auto *Ptr = mStorageAcc[0].construct</* ret type = */ AbstractOp>(TestCase);
56+
applyOp(mDataAcc.get_multi_ptr<sycl::access::decorated::no>().get(), Ptr);
57+
}
58+
auto get(oneapi::properties_tag) const {
59+
return oneapi::properties{oneapi::assume_indirect_calls};
60+
}
61+
};
62+
4863
int main() try {
4964
using storage_t =
5065
obj_storage_t<IncrementBy1, IncrementBy2, IncrementBy4, IncrementBy8>;
@@ -59,32 +74,15 @@ int main() try {
5974

6075
sycl::queue q(asyncHandler);
6176

62-
struct KernelFunctor {
63-
sycl::buffer<storage_t> mDeviceStorage;
64-
sycl::buffer<int> mDataStorage;
65-
sycl::handler mCGH;
66-
KernelFunctor(sycl::buffer<storage_t> DeviceStorage,
67-
sycl::buffer<int> DataStorage, sycl::handler CGH)
68-
: mDeviceStorage(DeviceStorage), mDataStorage(DataStorage), mCGH(CGH) {}
69-
70-
void operator()() const {
71-
sycl::accessor StorageAcc(mDeviceStorage, mCGH, sycl::write_only);
72-
sycl::accessor DataAcc(mDataStorage, mCGH, sycl::write_only);
73-
auto *Ptr =
74-
StorageAcc[0].construct</* ret type = */ AbstractOp>(TestCase);
75-
applyOp(DataAcc.get_multi_ptr<sycl::access::decorated::no>().get(), Ptr);
76-
}
77-
auto get(oneapi::properties_tag) const {
78-
return oneapi::properties{oneapi::assume_indirect_calls};
79-
}
80-
};
8177
for (unsigned TestCase = 0; TestCase < 4; ++TestCase) {
8278
int HostData = 42;
8379
int Data = HostData;
8480
sycl::buffer<int> DataStorage(&Data, sycl::range{1});
8581

8682
q.submit([&](sycl::handler &CGH) {
87-
CGH.single_task(KernelFunctor(DeviceStorage, DataStorage, CGH));
83+
sycl::accessor StorageAcc(DeviceStorage, CGH, sycl::write_only);
84+
sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only);
85+
CGH.single_task(KernelFunctor(StorageAcc, DataAcc));
8886
});
8987

9088
auto *Ptr = HostStorage.construct</* ret type = */ AbstractOp>(TestCase);

‎sycl/test-e2e/VirtualFunctions/2/1/1/simple-hierarchy.cpp‎

Lines changed: 18 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -30,6 +30,21 @@ class IncrementBy8 : public BaseIncrement {
3030
void increment(int *Data) override { *Data += 8; }
3131
};
3232

33+
template <typename T1, typename T2> struct KernelFunctor {
34+
T1 mStorageAcc;
35+
T2 mDataAcc;
36+
KernelFunctor(T1 StorageAcc, T2 DataAcc)
37+
: mStorageAcc(StorageAcc), mDataAcc(DataAcc) {}
38+
void operator()() const {
39+
auto *Ptr =
40+
mStorageAcc[0].construct</* ret type = */ BaseIncrement>(TestCase);
41+
Ptr->increment(mDataAcc.get_multi_ptr<sycl::access::decorated::no>().get());
42+
}
43+
auto get(oneapi::properties_tag) const {
44+
return oneapi::properties{oneapi::assume_indirect_calls};
45+
}
46+
};
47+
3348
int main() try {
3449
using storage_t =
3550
obj_storage_t<BaseIncrement, IncrementBy2, IncrementBy4, IncrementBy8>;
@@ -44,32 +59,15 @@ int main() try {
4459

4560
sycl::queue q(asyncHandler);
4661

47-
struct KernelFunctor {
48-
sycl::buffer<storage_t> mDeviceStorage;
49-
sycl::buffer<int> mDataStorage;
50-
sycl::handler mCGH;
51-
KernelFunctor(sycl::buffer<storage_t> DeviceStorage,
52-
sycl::buffer<int> DataStorage, sycl::handler CGH)
53-
: mDeviceStorage(DeviceStorage), mDataStorage(DataStorage), mCGH(CGH) {}
54-
void operator()() const {
55-
sycl::accessor StorageAcc(mDeviceStorage, mCGH, sycl::write_only);
56-
sycl::accessor DataAcc(mDataStorage, mCGH, sycl::write_only);
57-
auto *Ptr =
58-
StorageAcc[0].construct</* ret type = */ BaseIncrement>(TestCase);
59-
Ptr->increment(
60-
DataAcc.get_multi_ptr<sycl::access::decorated::no>().get());
61-
}
62-
auto get(oneapi::properties_tag) const {
63-
return oneapi::properties{oneapi::assume_indirect_calls};
64-
}
65-
};
6662
for (unsigned TestCase = 0; TestCase < 4; ++TestCase) {
6763
int HostData = 42;
6864
int Data = HostData;
6965
sycl::buffer<int> DataStorage(&Data, sycl::range{1});
7066

7167
q.submit([&](sycl::handler &CGH) {
72-
CGH.single_task(KernelFunctor(DeviceStorage, DataStorage, CGH));
68+
sycl::accessor StorageAcc(DeviceStorage, CGH, sycl::write_only);
69+
sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only);
70+
CGH.single_task(KernelFunctor(StorageAcc, DataAcc));
7371
});
7472

7573
auto *Ptr = HostStorage.construct</* ret type = */ BaseIncrement>(TestCase);

‎sycl/test-e2e/VirtualFunctions/2/2/single-construct-single-use.cpp‎

Lines changed: 18 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -57,6 +57,21 @@ class IncrementBy16 : public BaseIncrement {
5757
void increment(int *Data) override { *Data += 16 + Mod; }
5858
};
5959

60+
template <typename T1, typename T2> struct KernelFunctor {
61+
T1 mStorageAcc;
62+
T2 mDataAcc;
63+
KernelFunctor(T1 StorageAcc, T2 DataAcc)
64+
: mStorageAcc(StorageAcc), mDataAcc(DataAcc) {}
65+
void operator()() const {
66+
auto *Ptr = mStorageAcc[0].getAs<BaseIncrement>();
67+
Ptr->increment(mDataAcc.get_multi_ptr<sycl::access::decorated::no>().get());
68+
}
69+
auto get(oneapi::properties_tag) const {
70+
return oneapi::properties{
71+
oneapi::assume_indirect_calls_to<void, SetIncBy16>};
72+
}
73+
};
74+
6075
int main() try {
6176
using storage_t = obj_storage_t<BaseIncrement, IncrementBy2, IncrementBy4,
6277
IncrementBy8, IncrementBy16>;
@@ -72,25 +87,6 @@ int main() try {
7287
sycl::queue q(asyncHandler);
7388

7489
// TODO: cover uses case when objects are passed through USM
75-
struct KernelFunctor {
76-
sycl::buffer<storage_t> mDeviceStorage;
77-
sycl::buffer<int> mDataStorage;
78-
sycl::handler mCGH;
79-
KernelFunctor(sycl::buffer<storage_t> DeviceStorage,
80-
sycl::buffer<int> DataStorage, sycl::handler CGH)
81-
: mDeviceStorage(DeviceStorage), mDataStorage(DataStorage), mCGH(CGH) {}
82-
void operator()() const {
83-
sycl::accessor StorageAcc(mDeviceStorage, mCGH, sycl::read_write);
84-
sycl::accessor DataAcc(mDataStorage, mCGH, sycl::write_only);
85-
auto *Ptr = StorageAcc[0].getAs<BaseIncrement>();
86-
Ptr->increment(
87-
DataAcc.get_multi_ptr<sycl::access::decorated::no>().get());
88-
}
89-
auto get(oneapi::properties_tag) const {
90-
return oneapi::properties{
91-
oneapi::assume_indirect_calls_to<void, SetIncBy16>};
92-
}
93-
};
9490
for (unsigned TestCase = 0; TestCase < 5; ++TestCase) {
9591
int HostData = 42;
9692
int Data = HostData;
@@ -105,7 +101,9 @@ int main() try {
105101
});
106102

107103
q.submit([&](sycl::handler &CGH) {
108-
CGH.single_task(KernelFunctor(DeviceStorage, DataStorage, CGH));
104+
sycl::accessor StorageAcc(DeviceStorage, CGH, sycl::read_write);
105+
sycl::accessor DataAcc(DataStorage, CGH, sycl::write_only);
106+
CGH.single_task(KernelFunctor(StorageAcc, DataAcc));
109107
});
110108

111109
auto *Ptr =

‎sycl/test-e2e/VirtualFunctions/misc/math.cpp‎

Lines changed: 17 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -40,6 +40,21 @@ class RoundOp : public BaseOp {
4040
virtual float apply(float V) { return sycl::round(V); }
4141
};
4242

43+
template <typename T1, typename T2> struct KernelFunctor {
44+
T1 mDataAcc;
45+
T2 mDeviceStorage;
46+
KernelFunctor(T1 &DataAcc, T2 &DeviceStorage)
47+
: mDataAcc(DataAcc), mDeviceStorage(DeviceStorage) {}
48+
49+
void operator()() const {
50+
auto *Ptr = mDeviceStorage->getAs<BaseOp>();
51+
mDataAcc[0] = Ptr->apply(mDataAcc[0]);
52+
}
53+
auto get(oneapi::properties_tag) const {
54+
return oneapi::properties{oneapi::assume_indirect_calls};
55+
}
56+
};
57+
4358
int main() try {
4459
using storage_t = obj_storage_t<FloorOp, CeilOp, RoundOp>;
4560

@@ -49,23 +64,6 @@ int main() try {
4964

5065
auto *DeviceStorage = sycl::malloc_shared<storage_t>(1, q);
5166

52-
template <typename T> struct KernelFunctor {
53-
T mDeviceStorage;
54-
sycl::buffer<int> mDataStorage;
55-
sycl::handler mCGH;
56-
KernelFunctor(T DeviceStorage, sycl::buffer<float> DataStorage,
57-
sycl::handler CGH)
58-
: mDeviceStorage(DeviceStorage), mDataStorage(DataStorage), mCGH(CGH) {}
59-
60-
void operator()() const {
61-
sycl::accessor DataAcc(mDataStorage, mCGH, sycl::read_write);
62-
auto *Ptr = DeviceStorage->getAs<BaseOp>();
63-
DataAcc[0] = Ptr->apply(DataAcc[0]);
64-
}
65-
auto get(oneapi::properties_tag) const {
66-
return oneapi::properties{oneapi::assume_indirect_calls};
67-
}
68-
};
6967
for (unsigned TestCase = 0; TestCase < 3; ++TestCase) {
7068
float HostData = 3.56;
7169
float Data = HostData;
@@ -78,7 +76,8 @@ int main() try {
7876
}).wait_and_throw();
7977

8078
q.submit([&](sycl::handler &CGH) {
81-
CGH.single_task(KernelFunctor(DeviceStorage, DataStorage, CGH));
79+
sycl::accessor DataAcc(DataStorage, CGH, sycl::read_write);
80+
CGH.single_task(KernelFunctor(DataAcc, DeviceStorage));
8281
});
8382

8483
auto *Ptr = HostStorage.construct</* ret type = */ BaseOp>(TestCase);

0 commit comments

Comments
 (0)