Skip to content

Commit cd73d78

Browse files
committed
Combine reduction kernels
1 parent a98fb0c commit cd73d78

File tree

1 file changed

+7
-18
lines changed

1 file changed

+7
-18
lines changed

external/Stats.cpp

Lines changed: 7 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -44,27 +44,16 @@ void updateStats(sycl::queue myQueue, sycl::buffer<Actor> actorBuf,
4444

4545
auto forceSumReduction =
4646
sycl::reduction(forceSumBuf, cgh, sycl::plus<float>());
47-
48-
cgh.parallel_for(sycl::range<1>{actorAcc.size()}, forceSumReduction,
49-
[=](sycl::id<1> index, auto &sum) {
50-
if (!actorAcc[index].getAtDestination()) {
51-
sum += actorAcc[index].getForce();
52-
}
53-
});
54-
});
55-
myQueue.throw_asynchronous();
56-
57-
myQueue.submit([&](sycl::handler &cgh) {
58-
auto actorAcc = actorBuf.get_access<sycl::access::mode::read>(cgh);
59-
47+
6048
auto activeActorsReduction =
6149
sycl::reduction(activeActorsBuf, cgh, sycl::plus<int>());
6250

63-
cgh.parallel_for(sycl::range<1>{actorAcc.size()},
64-
activeActorsReduction,
65-
[=](sycl::id<1> index, auto &sum) {
66-
if (!actorAcc[index].getAtDestination()) {
67-
sum += 1;
51+
int globalSize = (std::floor(actorAcc.size()/16) + 1) * 16;
52+
cgh.parallel_for(sycl::nd_range<1>{globalSize, 16}, forceSumReduction, activeActorsReduction,
53+
[=](sycl::nd_item<1> index, auto &forceSum, auto& actorSum) {
54+
if (!actorAcc[index.get_global_id()].getAtDestination()) {
55+
forceSum += actorAcc[index.get_global_id()].getForce();
56+
actorSum += 1;
6857
}
6958
});
7059
});

0 commit comments

Comments
 (0)