@@ -293,138 +293,137 @@ sycl::info::event_command_status eventStatus(Event ev)
293
293
}
294
294
#endif
295
295
296
- void StarComplexityGen::fillStarMap (unsigned maxStars )
296
+ void StarComplexityGen::fillStarMap (unsigned numStars )
297
297
{
298
298
if (elemStars.empty ()) return ;
299
- // insert the single star graph
300
- starMap.emplace (elemStars[0 ],1 );
301
- counts.emplace (elemStars[0 ],1 );
299
+ if (numStars==1 )
300
+ {
301
+ // insert the single star graph
302
+ starMap.emplace (elemStars[0 ],1 );
303
+ counts.emplace (elemStars[0 ],1 );
304
+ return ;
305
+ }
302
306
303
- // unsigned numStars=maxStars;
304
- for (unsigned numStars=2 ; numStars<=maxStars; ++numStars)
305
- {
306
- ecolab::DeviceType<BlockEvaluator> block (blockSize, numStars,elemStars);
307
- Event resultSwapped, alreadySeenSwapped, starMapPopulated,backedResultsReset;
308
- atomic<bool > populating=false ;
309
- atomic<unsigned > resultBufferConsumed=0 ;
310
- const unsigned numOps=1 <<(numStars-1 );
311
-
312
- bool blown=false ;
313
- auto populateStarMap=[&]() {
314
- for (auto & j: block->getResults ())
315
- {
316
- if (j.blown ())
317
- {
318
- blown=true ;
319
- return ;
320
- }
321
- for (auto k: j)
322
- {
323
- auto res=starMap.emplace (k, numStars);
324
- if (res.first ->second ==numStars) // only count least star operations
325
- counts[k]++;
326
- }
327
- }
328
- };
307
+ ecolab::DeviceType<BlockEvaluator> block (blockSize, numStars,elemStars);
308
+ Event resultSwapped, alreadySeenSwapped, starMapPopulated,backedResultsReset;
309
+ atomic<bool > populating=false ;
310
+ atomic<unsigned > resultBufferConsumed=0 ;
311
+ const unsigned numOps=1 <<(numStars-1 );
329
312
330
- // int numLoops=3000;
331
- do
313
+ bool blown=false ;
314
+ auto populateStarMap=[&]() {
315
+ for (auto & j: block->getResults ())
332
316
{
333
- auto start=time (nullptr );
317
+ if (j.blown ())
318
+ {
319
+ blown=true ;
320
+ return ;
321
+ }
322
+ for (auto k: j)
323
+ {
324
+ auto res=starMap.emplace (k, numStars);
325
+ if (res.first ->second ==numStars) // only count least star operations
326
+ counts[k]++;
327
+ }
328
+ }
329
+ };
334
330
331
+ // int numLoops=3000;
332
+ do
333
+ {
334
+ auto start=time (nullptr );
335
+
335
336
#ifdef SYCL_LANGUAGE_VERSION
336
- /* Using SYCL dependency graph structure, to run the compute
337
- side simultaneously with accumulating results in the
338
- starMap on the host. sycl::event is used to coordinate host
339
- and device threads. */
340
- vector<Event> compute;
341
-
342
- auto consumeResults=[&](bool lastLoop) {
343
- // retrieve results to host by swapping with backing buffers
344
- resultSwapped=syclQ ().single_task
345
- (compute,[block=&*block](){
346
- block->result .swap (block->backedResult );
347
- });
348
- backedResultsReset=syclQ ().parallel_for
349
- (blockSize,resultSwapped,[block=&*block](auto i){block->backedResult [i].reset ();});
350
-
351
- // accumulate starMap results on separate host thread
352
- starMapPopulated=syclQ ().submit ([&](auto & handler) {
353
- handler.depends_on (resultSwapped);
354
- handler.host_task ([&]() {
355
- resultBufferConsumed=0 ;
356
- populateStarMap ();
357
- // if (!lastLoop)
358
- // {
359
- // block->alreadySeenBacked.clear();
360
- // for (auto& k: starMap) block->alreadySeenBacked.push_back(k.first);
361
- // // push update already seen vector to device
362
- // alreadySeenSwapped=syclQ().single_task
363
- // (compute,
364
- // [=,block=&*block]{
365
- // block->alreadySeen.swap(block->alreadySeenBacked);
366
- // });
367
- // }
368
- populating=false ;
369
- });
337
+ /* Using SYCL dependency graph structure, to run the compute
338
+ side simultaneously with accumulating results in the
339
+ starMap on the host. sycl::event is used to coordinate host
340
+ and device threads. */
341
+ vector<Event> compute;
342
+
343
+ auto consumeResults=[&](bool lastLoop) {
344
+ // retrieve results to host by swapping with backing buffers
345
+ resultSwapped=syclQ ().single_task
346
+ (compute,[block=&*block](){
347
+ block->result .swap (block->backedResult );
348
+ });
349
+ backedResultsReset=syclQ ().parallel_for
350
+ (blockSize,resultSwapped,[block=&*block](auto i){block->backedResult [i].reset ();});
351
+
352
+ // accumulate starMap results on separate host thread
353
+ starMapPopulated=syclQ ().submit ([&](auto & handler) {
354
+ handler.depends_on (resultSwapped);
355
+ handler.host_task ([&]() {
356
+ resultBufferConsumed=0 ;
357
+ populateStarMap ();
358
+ // if (!lastLoop)
359
+ // {
360
+ // block->alreadySeenBacked.clear();
361
+ // for (auto& k: starMap) block->alreadySeenBacked.push_back(k.first);
362
+ // // push update already seen vector to device
363
+ // alreadySeenSwapped=syclQ().single_task
364
+ // (compute,
365
+ // [=,block=&*block]{
366
+ // block->alreadySeen.swap(block->alreadySeenBacked);
367
+ // });
368
+ // }
369
+ populating=false ;
370
370
});
371
+ });
371
372
372
- Event::wait ({resultSwapped,backedResultsReset});
373
- resultBufferConsumed=0 ;
374
- };
373
+ Event::wait ({resultSwapped,backedResultsReset});
374
+ resultBufferConsumed=0 ;
375
+ };
375
376
376
- for (unsigned i=0 ; i<block->numGraphs ; i+=blockSize)
377
- {
378
- // if (resultBufferConsumed+numOps>=OutputBuffer::maxQ)
379
- // {
380
- // consumeResults(i+blockSize>=block->numGraphs);
381
- // cout<<"pausing..."<<endl;
382
- // backedResultsReset.wait();
383
- // // while (resultBufferConsumed)
384
- // // usleep(1000); // pause until results consumed
385
- // }
386
-
387
- // compute a block of graphs
388
- compute.push_back
389
- (syclQ ().parallel_for (blockSize, {resultSwapped, alreadySeenSwapped},
390
- [=,block=&*block](auto j){block->eval (i,j);}));
391
- resultBufferConsumed+=numOps;
392
-
393
- if (eventStatus (starMapPopulated)==sycl::info::event_command_status::complete)
394
- // throttle consuming thread
395
- consumeResults (i+blockSize>=block->numGraphs );
396
- }
397
- // wait until all compute threads finish before bumping pos
398
- Event::wait (compute);
377
+ for (unsigned i=0 ; i<block->numGraphs ; i+=blockSize)
378
+ {
379
+ // if (resultBufferConsumed+numOps>=OutputBuffer::maxQ)
380
+ // {
381
+ // consumeResults(i+blockSize>=block->numGraphs);
382
+ // cout<<"pausing..."<<endl;
383
+ // backedResultsReset.wait();
384
+ // // while (resultBufferConsumed)
385
+ // // usleep(1000); // pause until results consumed
386
+ // }
387
+
388
+ // compute a block of graphs
389
+ compute.push_back
390
+ (syclQ ().parallel_for (blockSize, {resultSwapped, alreadySeenSwapped},
391
+ [=,block=&*block](auto j){block->eval (i,j);}));
392
+ resultBufferConsumed+=numOps;
393
+
394
+ if (eventStatus (starMapPopulated)==sycl::info::event_command_status::complete)
395
+ // throttle consuming thread
396
+ consumeResults (i+blockSize>=block->numGraphs );
397
+ }
398
+ // wait until all compute threads finish before bumping pos
399
+ Event::wait (compute);
399
400
#else
400
- for (unsigned i=0 ; i<block->numGraphs ; i+=block->size ())
401
- for (unsigned j=0 ; j<block->size (); ++j)
402
- block->eval (i,j);
403
- populateStarMap ();
401
+ for (unsigned i=0 ; i<block->numGraphs ; i+=block->size ())
402
+ for (unsigned j=0 ; j<block->size (); ++j)
403
+ block->eval (i,j);
404
+ populateStarMap ();
404
405
#endif
405
- if (blown)
406
- {
407
- cout<<" buffer blown, throwing" <<endl;
406
+ if (blown)
407
+ {
408
+ cout<<" buffer blown, throwing" <<endl;
408
409
#ifdef SYCL_LANGUAGE_VERSION
409
- syclQ ().wait (); // flush queue before destructors called
410
+ syclQ ().wait (); // flush queue before destructors called
410
411
#endif
411
- throw runtime_error (" Output buffer blown" );
412
- }
412
+ throw runtime_error (" Output buffer blown" );
413
+ }
413
414
414
- // cout<<"loops rem="<<numLoops<<" resultBufferConsumed:"<<unsigned(resultBufferConsumed)<<" "<<(time(nullptr)-start)<<"secs\n";
415
- } while (block->pos .next () /* && --numLoops>0*/ );
415
+ // cout<<"loops rem="<<numLoops<<" resultBufferConsumed:"<<unsigned(resultBufferConsumed)<<" "<<(time(nullptr)-start)<<"secs\n";
416
+ } while (block->pos .next () /* && --numLoops>0*/ );
416
417
#ifdef SYCL_LANGUAGE_VERSION
417
- auto start=time (nullptr );
418
- resultSwapped=syclQ ().single_task
419
- ([block=&*block](){
420
- block->result .swap (block->backedResult );
421
- });
422
- populateStarMap ();
423
- syclQ ().wait_and_throw (); // flush queue before destructors called
424
- cout<<" numStars done:" <<numStars<<endl;
425
- // cout<<"Final wait:"<<(time(nullptr)-start)<<"secs\n";
418
+ auto start=time (nullptr );
419
+ resultSwapped=syclQ ().single_task
420
+ ([block=&*block](){
421
+ block->result .swap (block->backedResult );
422
+ });
423
+ populateStarMap ();
424
+ syclQ ().wait_and_throw (); // flush queue before destructors called
426
425
#endif
427
- }
426
+
428
427
}
429
428
430
429
NautyRep toNautyRep (linkRep g, unsigned nodes)
0 commit comments