4
4
void updateStats (sycl::queue myQueue, sycl::buffer<Actor> actorBuf,
5
5
std::vector<float > &averageForces,
6
6
std::vector<std::array<int , 2 >> &destinationTimes,
7
- std::chrono::high_resolution_clock::time_point startPoint, int timestep) {
7
+ std::chrono::high_resolution_clock::time_point startPoint,
8
+ int timestep) {
8
9
try {
9
10
float forceSum = 0 ;
10
11
auto forceSumBuf = sycl::buffer<float >(&forceSum, 1 );
@@ -13,87 +14,83 @@ void updateStats(sycl::queue myQueue, sycl::buffer<Actor> actorBuf,
13
14
auto activeActorsBuf = sycl::buffer<int >(&activeActors, 1 );
14
15
15
16
// Calculate average force applied to actors this iteration
16
- myQueue
17
- .submit ([&](sycl::handler &cgh) {
18
- auto actorAcc =
19
- actorBuf.get_access <sycl::access::mode::read>(cgh);
20
-
21
- auto forceSumReduction =
22
- sycl::reduction (forceSumBuf, cgh, sycl::plus<float >());
23
-
24
- cgh.parallel_for (sycl::range<1 >{actorAcc.size ()},
25
- forceSumReduction,
26
- [=](sycl::id<1 > index, auto &sum) {
27
- if (!actorAcc[index].getAtDestination ()) {
28
- sum += actorAcc[index].getForce ();
29
- }
30
- });
31
- });
17
+ myQueue.submit ([&](sycl::handler &cgh) {
18
+ auto actorAcc = actorBuf.get_access <sycl::access::mode::read>(cgh);
19
+
20
+ auto forceSumReduction =
21
+ sycl::reduction (forceSumBuf, cgh, sycl::plus<float >());
22
+
23
+ cgh.parallel_for (sycl::range<1 >{actorAcc.size ()}, forceSumReduction,
24
+ [=](sycl::id<1 > index, auto &sum) {
25
+ if (!actorAcc[index].getAtDestination ()) {
26
+ sum += actorAcc[index].getForce ();
27
+ }
28
+ });
29
+ });
32
30
myQueue.throw_asynchronous ();
33
31
34
- myQueue
35
- .submit ([&](sycl::handler &cgh) {
36
- auto actorAcc =
37
- actorBuf.get_access <sycl::access::mode::read>(cgh);
38
-
39
- auto activeActorsReduction =
40
- sycl::reduction (activeActorsBuf, cgh, sycl::plus<int >());
41
-
42
- cgh.parallel_for (sycl::range<1 >{actorAcc.size ()},
43
- activeActorsReduction,
44
- [=](sycl::id<1 > index, auto &sum) {
45
- if (!actorAcc[index].getAtDestination ()) {
46
- sum += 1 ;
47
- }
48
- });
49
- });
32
+ myQueue.submit ([&](sycl::handler &cgh) {
33
+ auto actorAcc = actorBuf.get_access <sycl::access::mode::read>(cgh);
34
+
35
+ auto activeActorsReduction =
36
+ sycl::reduction (activeActorsBuf, cgh, sycl::plus<int >());
37
+
38
+ cgh.parallel_for (sycl::range<1 >{actorAcc.size ()},
39
+ activeActorsReduction,
40
+ [=](sycl::id<1 > index, auto &sum) {
41
+ if (!actorAcc[index].getAtDestination ()) {
42
+ sum += 1 ;
43
+ }
44
+ });
45
+ });
50
46
myQueue.throw_asynchronous ();
51
47
52
48
sycl::host_accessor<float , 1 , sycl::access::mode::read> forceSumHostAcc (
53
49
forceSumBuf);
54
50
sycl::host_accessor<int , 1 , sycl::access::mode::read>
55
51
activeActorsHostAcc (activeActorsBuf);
56
- averageForces.push_back (forceSumHostAcc[0 ] / float (activeActorsHostAcc[0 ]));
52
+ averageForces.push_back (forceSumHostAcc[0 ] /
53
+ float (activeActorsHostAcc[0 ]));
57
54
58
55
// Find actors which have reached their destination and record how long
59
56
// it took them
60
- auto destinationTimesBuf =
61
- sycl::buffer<std::array< int , 2 >>( destinationTimes.data (), destinationTimes.size ());
57
+ auto destinationTimesBuf = sycl::buffer<std::array< int , 2 >>(
58
+ destinationTimes.data (), destinationTimes.size ());
62
59
auto timestepBuf = sycl::buffer<int >(×tep, 1 );
63
60
64
- myQueue
65
- . submit ([&]( sycl::handler & cgh) {
66
- auto actorAcc =
67
- actorBuf .get_access <sycl::access::mode::read>(cgh);
68
- auto destinationTimesAcc =
69
- destinationTimesBuf
70
- .get_access <sycl::access::mode::read_write >(cgh);
71
- auto timestepAcc = timestepBuf. get_access <sycl::access::mode::read>(cgh);
72
-
73
- auto end = std::chrono::high_resolution_clock::now ();
74
- auto duration =
75
- std::chrono::duration_cast<std::chrono::milliseconds>(
76
- end - startPoint)
77
- . count ();
78
-
79
- cgh. parallel_for (sycl::range <1 >{destinationTimesAcc. size ()},
80
- [=](sycl::id< 1 > index) {
81
- if (actorAcc [index]. getAtDestination () &&
82
- destinationTimesAcc[index][ 0 ] == 0 ) {
83
- destinationTimesAcc[index] = { int (duration), timestepAcc[0 ]};
84
- }
85
- });
86
- });
61
+ myQueue. submit ([&](sycl::handler &cgh) {
62
+ auto actorAcc = actorBuf. get_access < sycl::access::mode::read>( cgh);
63
+ auto destinationTimesAcc =
64
+ destinationTimesBuf .get_access <sycl::access::mode::read_write>(
65
+ cgh);
66
+ auto timestepAcc =
67
+ timestepBuf .get_access <sycl::access::mode::read >(cgh);
68
+
69
+ auto end = std::chrono::high_resolution_clock::now ();
70
+ auto duration =
71
+ std::chrono::duration_cast<std::chrono::milliseconds>(
72
+ end - startPoint)
73
+ . count ();
74
+
75
+ cgh. parallel_for (sycl::range< 1 >{destinationTimesAcc. size ()},
76
+ [=] (sycl::id <1 > index) {
77
+ if (actorAcc[index]. getAtDestination () &&
78
+ destinationTimesAcc [index][ 0 ] == 0 ) {
79
+ destinationTimesAcc[index] = {
80
+ int (duration), timestepAcc[0 ]};
81
+ }
82
+ });
83
+ });
87
84
myQueue.throw_asynchronous ();
88
-
85
+
89
86
} catch (const sycl::exception &e) {
90
87
std::cout << " SYCL exception caught:\n "
91
88
<< e.what () << " \n [updateStats]" ;
92
89
}
93
90
}
94
91
95
92
void finalizeStats (sycl::queue myQueue, std::vector<float > averageForces,
96
- std::vector<std::array<int ,2 >> destinationTimes,
93
+ std::vector<std::array<int , 2 >> destinationTimes,
97
94
std::vector<int > kernelDurations, int numActors,
98
95
int totalExecutionTime) {
99
96
try {
@@ -106,21 +103,18 @@ void finalizeStats(sycl::queue myQueue, std::vector<float> averageForces,
106
103
sycl::buffer<int >(kernelDurations.data (), kernelDurations.size ());
107
104
108
105
// Calculate average kernel duration
109
- myQueue
110
- .submit ([&](sycl::handler &cgh) {
111
- auto durationAcc =
112
- kernelDurationsBuf.get_access <sycl::access::mode::read>(
113
- cgh);
114
-
115
- auto sumReduction =
116
- sycl::reduction (durationSumBuf, cgh, sycl::plus<int >());
117
-
118
- cgh.parallel_for (sycl::range<1 >{durationAcc.size ()},
119
- sumReduction,
120
- [=](sycl::id<1 > index, auto &sum) {
121
- sum += durationAcc[index];
122
- });
123
- });
106
+ myQueue.submit ([&](sycl::handler &cgh) {
107
+ auto durationAcc =
108
+ kernelDurationsBuf.get_access <sycl::access::mode::read>(cgh);
109
+
110
+ auto sumReduction =
111
+ sycl::reduction (durationSumBuf, cgh, sycl::plus<int >());
112
+
113
+ cgh.parallel_for (sycl::range<1 >{durationAcc.size ()}, sumReduction,
114
+ [=](sycl::id<1 > index, auto &sum) {
115
+ sum += durationAcc[index];
116
+ });
117
+ });
124
118
myQueue.throw_asynchronous ();
125
119
126
120
sycl::host_accessor<int , 1 , sycl::access::mode::read>
@@ -153,11 +147,13 @@ void finalizeStats(sycl::queue myQueue, std::vector<float> averageForces,
153
147
outputFile << std::setprecision (2 ) << std::fixed;
154
148
outputFile << std::setw (8 ) << x << " |" ;
155
149
if (destinationTimes[x][0 ] == 0 ) {
156
- outputFile << std::setw (31 ) << " NA" << " |" ;
150
+ outputFile << std::setw (31 ) << " NA"
151
+ << " |" ;
157
152
} else {
158
- outputFile << std::setw (31 ) << destinationTimes[x][0 ] << " |" ;
153
+ outputFile << std::setw (31 ) << destinationTimes[x][0 ]
154
+ << " |" ;
159
155
}
160
- if (destinationTimes[x][1 ] == 0 ) {
156
+ if (destinationTimes[x][1 ] == 0 ) {
161
157
outputFile << std::setw (8 ) << " NA" ;
162
158
} else {
163
159
outputFile << std::setw (8 ) << destinationTimes[x][1 ];
0 commit comments