SYCL implementation #44

jeffhammond · 2020-06-27T16:51:22Z

This is the SYCL/DPC++ port. It currently depends on two features that are not widely available:

USM (unified shared memory), which serves the same purpose as cudaMallocManaged. It is part of SYCL 2020 but only Intel DPC++ on Intel back-ends supports it today. CodePlay ComputeCpp has started implementing it but I don't think it is finished yet and I didn't test it.
sycl::intel::experimental::printf, which is, as one might expect, an Intel extension to support printf. There is an alternative extension in CodePlay ComputeCpp but I didn't bother with that because of the previous issue.

The following output appears to be identical to that of GCC OpenMP, but please let me know what other verification I need to do.

Today, the Intel DPC++ implementation is working with the host and Gen9 GPU devices, but not the CPU device because of an Intel OpenCL issue that is known and in the process of being fixed.

jrhammon@jrhammon-nuc:~/QUICKSILVER/src$ QS_DEVICE=GPU ./qs 
Copyright (c) 2016
Lawrence Livermore National Security, LLC
All Rights Reserved
Quicksilver Version     : 2020-Feb-4-22:35:56
Quicksilver Git Hash    : af27b3dcce08933786cb526e2f8a0bbe99d99b07
MPI Version             : 3.0
Number of MPI ranks     : 1
Number of OpenMP Threads: 1
Number of OpenMP CPUs   : 1

Simulation:
   dt: 1e-08
   fMax: 0.1
   inputFile: 
   energySpectrum: 
   boundaryCondition: reflect
   loadBalance: 0
   cycleTimers: 0
   debugThreads: 0
   lx: 100
   ly: 100
   lz: 100
   nParticles: 1000000
   batchSize: 0
   nBatches: 10
   nSteps: 10
   nx: 10
   ny: 10
   nz: 10
   seed: 1029384756
   xDom: 0
   yDom: 0
   zDom: 0
   eMax: 20
   eMin: 1e-09
   nGroups: 230
   lowWeightCutoff: 0.001
   bTally: 1
   fTally: 1
   cTally: 1
   coralBenchmark: 0
   crossSectionsOut:

Geometry:
   material: sourceMaterial
   shape: brick
   xMax: 100
   xMin: 0
   yMax: 100
   yMin: 0
   zMax: 100
   zMin: 0

Material:
   name: sourceMaterial
   mass: 1000
   nIsotopes: 10
   nReactions: 9
   sourceRate: 1e+10
   totalCrossSection: 1
   absorptionCrossSection: flat
   fissionCrossSection: flat
   scatteringCrossSection: flat
   absorptionCrossSectionRatio: 1
   fissionCrossSectionRatio: 0.1
   scatteringCrossSectionRatio: 1

CrossSection:
   name: flat
   A: 0
   B: 0
   C: 0
   D: 0
   E: 1
   nuBar: 2.4
is gpu
Building partition 0
Building partition 1
Building partition 2
Building partition 3
done building
Building MC_Domain 0
Building MC_Domain 1
Building MC_Domain 2
Building MC_Domain 3
Starting Consistency Check
Finished Consistency Check
Finished initMesh
Using SYCL device
cycle           start       source           rr        split       absorb      scatter      fission      produce      collisn       escape       census      num_seg   scalar_flux      cycleInit  cycleTracking  cycleFinalize
       0            0       100000            0       900000      1078182      1076792       107133       257364      2262107            0        72049      2670386  2.264064e+08   4.442000e-02   6.090415e+00   0.000000e+00
       1        72049       100000            0       828008      1107255      1106235       110306       264657      2323796            0        47153      2719702  2.438830e+08   4.589800e-02   3.454263e+00   0.000000e+00
       2        47153       100000            0       852712      1086097      1088696       108334       259738      2283127            0        65172      2687840  2.435394e+08   4.463300e-02   3.493513e+00   0.000000e+00
       3        65172       100000        68015       834785      1017555      1018659       101778       244593      2137992            0        57202      2517378  2.450517e+08   4.562600e-02   3.246188e+00   0.000000e+00
       4        57202       100000        62214       842934      1020418      1019522       101687       244038      2141627            0        59855      2522163  2.434017e+08   4.571000e-02   3.285100e+00   0.000000e+00
       5        59855       100000        56726       840345      1029994      1029682       103183       247672      2162859            0        57969      2545713  2.451216e+08   4.450600e-02   3.400606e+00   0.000000e+00
       6        57969       100000        52439       841925      1032190      1032180       102801       246877      2167171            0        59341      2551468  2.446226e+08   4.578000e-02   3.692880e+00   0.000000e+00
       7        59341       100000        59663       840635      1023444      1022593       102792       246649      2148829            0        60726      2531066  2.441845e+08   4.598600e-02   3.360731e+00   0.000000e+00
       8        60726       100000        68187       839357      1013501      1014287       101238       243112      2129026            0        60269      2508491  2.440307e+08   4.525700e-02   3.606579e+00   0.000000e+00
       9        60269       100000        71159       839953      1012439      1011892       101368       243262      2125699            0        58518      2500968  2.444142e+08   4.538600e-02   3.415050e+00   0.000000e+00

Timer                       Cumulative   Cumulative   Cumulative   Cumulative   Cumulative   Cumulative
Name                            number    microSecs    microSecs    microSecs    microSecs   Efficiency
                              of calls          min          avg          max       stddev       Rating
main                                 1    3.750e+07    3.750e+07    3.750e+07    0.000e+00       100.00
cycleInit                           10    4.532e+05    4.532e+05    4.532e+05    0.000e+00       100.00
cycleTracking                       10    3.705e+07    3.705e+07    3.705e+07    0.000e+00       100.00
cycleTracking_Kernel               992    3.699e+07    3.699e+07    3.699e+07    0.000e+00       100.00
cycleTracking_MPI                 1083    5.392e+04    5.392e+04    5.392e+04    0.000e+00       100.00
cycleTracking_Test_Done              0    0.000e+00    0.000e+00    0.000e+00    0.000e+00         0.00
cycleFinalize                       20    1.564e+03    1.564e+03    1.564e+03    0.000e+00       100.00
Figure Of Merit              6.952e+05 [Num Segments / Cycle Tracking Time]

look for HAVE_SYCL and SYCL_EXTERNAL for most of the changes. this works on GPU and HOST but not CPU (OpenCL issue). Signed-off-by: Jeff Hammond <[email protected]>

Signed-off-by: Jeff Hammond <[email protected]>

jeffhammond · 2020-06-27T16:54:17Z

Makefile is set to use DPC++ so we will want to change that before merging into the main branch, but I'm leaving it that way for testing purposes.

I am also leaving in a bit of debug code that I think is useful until we figure out the right way to control SYCL device dispatch. Because SYCL is a pluripotent back-end, it isn't obvious how to do this. Intel DPC++ allows one to set the default device with an environment variable, but I wanted to push control in QS for debugging purposes.

jeffhammond · 2020-10-09T16:38:35Z

I have verified that the SYCL implementation also runs correctly on NVIDIA (Pascal).

There is a compiler bug with printf on the device that requires me to disable 4 source instances of device printf but otherwise the code is identical.

Compiler

clang version 12.0.0 (https://github.com/intel/llvm.git b1cf776e91a0f8f99397a5c3668cceda19b1b000)

Hardware

$ nvidia-smi
Fri Oct  9 12:35:04 2020       
+-----------------------------------------------------------------------------+
| NVIDIA-SMI 450.51.06    Driver Version: 450.51.06    CUDA Version: 11.0     |
|-------------------------------+----------------------+----------------------+
| GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
| Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
|                               |                      |               MIG M. |
|===============================+======================+======================|
|   0  GeForce GTX 1080    On   | 00000000:3B:00.0 Off |                  N/A |
| 24%   38C    P8     8W / 180W |      7MiB /  8119MiB |      0%      Default |
|                               |                      |                  N/A |
+-------------------------------+----------------------+----------------------+
                                                                               
+-----------------------------------------------------------------------------+
| Processes:                                                                  |
|  GPU   GI   CI        PID   Type   Process name                  GPU Memory |
|        ID   ID                                                   Usage      |
|=============================================================================|
|    0   N/A  N/A      3309      G   /usr/lib/xorg/Xorg                  4MiB |
+-----------------------------------------------------------------------------+

Host execution

$ QS_DEVICE=HOST ./qs 
Copyright (c) 2016
Lawrence Livermore National Security, LLC
All Rights Reserved
Quicksilver Version     : 2020-Jun-27-09:44:01
Quicksilver Git Hash    : 7794a66f2a3d1ea6446519615dcadb7a15c2a707
MPI Version             : 3.0
Number of MPI ranks     : 1
Number of OpenMP Threads: 1
Number of OpenMP CPUs   : 1

Simulation:
   dt: 1e-08
   fMax: 0.1
   inputFile: 
   energySpectrum: 
   boundaryCondition: reflect
   loadBalance: 0
   cycleTimers: 0
   debugThreads: 0
   lx: 100
   ly: 100
   lz: 100
   nParticles: 1000000
   batchSize: 0
   nBatches: 10
   nSteps: 10
   nx: 10
   ny: 10
   nz: 10
   seed: 1029384756
   xDom: 0
   yDom: 0
   zDom: 0
   eMax: 20
   eMin: 1e-09
   nGroups: 230
   lowWeightCutoff: 0.001
   bTally: 1
   fTally: 1
   cTally: 1
   coralBenchmark: 0
   crossSectionsOut:

Geometry:
   material: sourceMaterial
   shape: brick
   xMax: 100
   xMin: 0
   yMax: 100
   yMin: 0
   zMax: 100
   zMin: 0

Material:
   name: sourceMaterial
   mass: 1000
   nIsotopes: 10
   nReactions: 9
   sourceRate: 1e+10
   totalCrossSection: 1
   absorptionCrossSection: flat
   fissionCrossSection: flat
   scatteringCrossSection: flat
   absorptionCrossSectionRatio: 1
   fissionCrossSectionRatio: 0.1
   scatteringCrossSectionRatio: 1

CrossSection:
   name: flat
   A: 0
   B: 0
   C: 0
   D: 0
   E: 1
   nuBar: 2.4
is host
Building partition 0
Building partition 1
Building partition 2
Building partition 3
done building
Building MC_Domain 0
Building MC_Domain 1
Building MC_Domain 2
Building MC_Domain 3
Starting Consistency Check
Finished Consistency Check
Finished initMesh
Using SYCL device
cycle           start       source           rr        split       absorb      scatter      fission      produce      collisn       escape       census      num_seg   scalar_flux      cycleInit  cycleTracking  cycleFinalize
       0            0       100000            0       900000      1078182      1076792       107133       257364      2262107            0        72049      2670386  2.264064e+08   8.100800e-02   2.926682e+00   0.000000e+00
       1        72049       100000            0       828008      1107255      1106235       110306       264657      2323796            0        47153      2719702  2.438830e+08   1.334280e-01   3.022620e+00   0.000000e+00
       2        47153       100000            0       852712      1086097      1088696       108334       259738      2283127            0        65172      2687840  2.435394e+08   1.117280e-01   2.942065e+00   0.000000e+00
       3        65172       100000        68015       834785      1017555      1018659       101778       244593      2137992            0        57202      2517378  2.450517e+08   1.087580e-01   2.757061e+00   0.000000e+00
       4        57202       100000        62214       842934      1020418      1019522       101687       244038      2141627            0        59855      2522163  2.434017e+08   1.091730e-01   2.750167e+00   0.000000e+00
       5        59855       100000        56726       840345      1029994      1029682       103183       247672      2162859            0        57969      2545713  2.451216e+08   8.949400e-02   2.795324e+00   0.000000e+00
       6        57969       100000        52439       841925      1032190      1032180       102801       246877      2167171            0        59341      2551468  2.446226e+08   8.452000e-02   2.786467e+00   0.000000e+00
       7        59341       100000        59663       840635      1023444      1022593       102792       246649      2148829            0        60726      2531066  2.441845e+08   8.383200e-02   2.782119e+00   1.000000e-06
       8        60726       100000        68187       839357      1013501      1014287       101238       243112      2129026            0        60269      2508491  2.440307e+08   8.301600e-02   2.735566e+00   1.000000e-06
       9        60269       100000        71159       839953      1012439      1011892       101368       243262      2125699            0        58518      2500968  2.444142e+08   8.994500e-02   2.722267e+00   1.000000e-06

Timer                       Cumulative   Cumulative   Cumulative   Cumulative   Cumulative   Cumulative
Name                            number    microSecs    microSecs    microSecs    microSecs   Efficiency
                              of calls          min          avg          max       stddev       Rating
main                                 1    2.920e+07    2.920e+07    2.920e+07    0.000e+00       100.00
cycleInit                           10    9.749e+05    9.749e+05    9.749e+05    0.000e+00       100.00
cycleTracking                       10    2.822e+07    2.822e+07    2.822e+07    0.000e+00       100.00
cycleTracking_Kernel               992    2.810e+07    2.810e+07    2.810e+07    0.000e+00       100.00
cycleTracking_MPI                 1083    1.171e+05    1.171e+05    1.171e+05    0.000e+00       100.00
cycleTracking_Test_Done              0    0.000e+00    0.000e+00    0.000e+00    0.000e+00         0.00
cycleFinalize                       20    1.387e+03    1.387e+03    1.387e+03    0.000e+00       100.00
Figure Of Merit              9.126e+05 [Num Segments / Cycle Tracking Time]

NVIDIA execution

$ QS_DEVICE=GPU /usr/local/cuda-11.0/bin/nvprof ./qs 
==44046== NVPROF is profiling process 44046, command: ./qs
Copyright (c) 2016
Lawrence Livermore National Security, LLC
All Rights Reserved
Quicksilver Version     : 2020-Jun-27-09:44:01
Quicksilver Git Hash    : 7794a66f2a3d1ea6446519615dcadb7a15c2a707
MPI Version             : 3.0
Number of MPI ranks     : 1
Number of OpenMP Threads: 1
Number of OpenMP CPUs   : 1

Simulation:
   dt: 1e-08
   fMax: 0.1
   inputFile: 
   energySpectrum: 
   boundaryCondition: reflect
   loadBalance: 0
   cycleTimers: 0
   debugThreads: 0
   lx: 100
   ly: 100
   lz: 100
   nParticles: 1000000
   batchSize: 0
   nBatches: 10
   nSteps: 10
   nx: 10
   ny: 10
   nz: 10
   seed: 1029384756
   xDom: 0
   yDom: 0
   zDom: 0
   eMax: 20
   eMin: 1e-09
   nGroups: 230
   lowWeightCutoff: 0.001
   bTally: 1
   fTally: 1
   cTally: 1
   coralBenchmark: 0
   crossSectionsOut:

Geometry:
   material: sourceMaterial
   shape: brick
   xMax: 100
   xMin: 0
   yMax: 100
   yMin: 0
   zMax: 100
   zMin: 0

Material:
   name: sourceMaterial
   mass: 1000
   nIsotopes: 10
   nReactions: 9
   sourceRate: 1e+10
   totalCrossSection: 1
   absorptionCrossSection: flat
   fissionCrossSection: flat
   scatteringCrossSection: flat
   absorptionCrossSectionRatio: 1
   fissionCrossSectionRatio: 0.1
   scatteringCrossSectionRatio: 1

CrossSection:
   name: flat
   A: 0
   B: 0
   C: 0
   D: 0
   E: 1
   nuBar: 2.4
is gpu
Building partition 0
Building partition 1
Building partition 2
Building partition 3
done building
Building MC_Domain 0
Building MC_Domain 1
Building MC_Domain 2
Building MC_Domain 3
Starting Consistency Check
Finished Consistency Check
Finished initMesh
Using SYCL device
cycle           start       source           rr        split       absorb      scatter      fission      produce      collisn       escape       census      num_seg   scalar_flux      cycleInit  cycleTracking  cycleFinalize
       0            0       100000            0       900000      1078182      1076792       107133       257364      2262107            0        72049      2670386  2.264064e+08   8.140800e-02   3.494830e-01   4.400000e-05
       1        72049       100000            0       828008      1107254      1106235       110306       264657      2323795            0        47154      2719703  2.438830e+08   1.319410e-01   3.669860e-01   4.200000e-05
       2        47154       100000            0       852712      1086097      1088693       108332       259734      2283122            0        65171      2687830  2.435392e+08   1.292360e-01   3.583770e-01   4.100000e-05
       3        65171       100000        68010       834788      1017561      1018651       101777       244591      2137989            0        57202      2517372  2.450488e+08   1.285070e-01   3.940940e-01   4.400000e-05
       4        57202       100000        62214       842933      1020417      1019520       101687       244038      2141624            0        59855      2522155  2.434013e+08   1.233870e-01   3.430330e-01   4.200000e-05
       5        59855       100000        56726       840345      1029994      1029682       103183       247672      2162859            0        57969      2545713  2.451217e+08   1.251520e-01   9.301870e-01   4.900000e-05
       6        57969       100000        52439       841925      1032190      1032180       102801       246877      2167171            0        59341      2551472  2.446224e+08   1.254430e-01   4.115140e-01   4.300000e-05
       7        59341       100000        59663       840635      1023444      1022593       102792       246649      2148829            0        60726      2531066  2.441844e+08   1.263230e-01   3.548820e-01   4.100000e-05
       8        60726       100000        68187       839357      1013501      1014287       101238       243112      2129026            0        60269      2508492  2.440307e+08   1.245810e-01   1.525691e+00   4.200000e-05
       9        60269       100000        71159       839953      1012439      1011892       101368       243262      2125699            0        58518      2500968  2.444142e+08   1.244990e-01   3.747170e-01   4.200000e-05

Timer                       Cumulative   Cumulative   Cumulative   Cumulative   Cumulative   Cumulative
Name                            number    microSecs    microSecs    microSecs    microSecs   Efficiency
                              of calls          min          avg          max       stddev       Rating
main                                 1    6.647e+06    6.647e+06    6.647e+06    0.000e+00       100.00
cycleInit                           10    1.220e+06    1.220e+06    1.220e+06    0.000e+00       100.00
cycleTracking                       10    5.409e+06    5.409e+06    5.409e+06    0.000e+00       100.00
cycleTracking_Kernel               992    5.028e+06    5.028e+06    5.028e+06    0.000e+00       100.00
cycleTracking_MPI                 1083    3.471e+05    3.471e+05    3.471e+05    0.000e+00       100.00
cycleTracking_Test_Done              0    0.000e+00    0.000e+00    0.000e+00    0.000e+00         0.00
cycleFinalize                       20    3.918e+03    3.918e+03    3.918e+03    0.000e+00       100.00
Figure Of Merit              4.762e+06 [Num Segments / Cycle Tracking Time]
==44046== Profiling application: ./qs
==44046== Profiling result:
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:  100.00%  5.00802s       195  25.682ms  175.84us  1.15108s  _ZTSZZ13cycleTrackingP10MonteCarloENKUlRN2cl4sycl7handlerEE276_52clES4_EUlNS2_4itemILi1ELb1EEEE277_66
      API calls:   88.88%  5.00789s       195  25.681ms  175.10us  1.15110s  cuEventSynchronize
                    7.76%  437.11ms         2  218.55ms  205.62ms  231.49ms  cuCtxCreate
                    2.40%  135.24ms         2  67.618ms  61.238ms  73.998ms  cuCtxDestroy
                    0.45%  25.116ms       393  63.908us  4.7490us  20.552ms  cuMemAllocManaged
                    0.39%  21.747ms       393  55.335us  8.9590us  950.74us  cuMemFree
                    0.06%  3.3578ms       195  17.219us  8.1610us  101.59us  cuLaunchKernel
                    0.03%  1.9149ms         1  1.9149ms  1.9149ms  1.9149ms  cuModuleLoadDataEx
                    0.01%  495.29us         1  495.29us  495.29us  495.29us  cuModuleUnload
                    0.01%  351.42us      1579     222ns     162ns  3.0640us  cuCtxGetCurrent
                    0.01%  326.05us       783     416ns     174ns  13.519us  cuDeviceGetAttribute
                    0.00%  272.17us       197  1.3810us     590ns  12.086us  cuEventCreate
                    0.00%  224.69us       197  1.1400us     708ns  5.0690us  cuEventRecord
                    0.00%  156.83us       197     796ns     412ns  7.6580us  cuEventDestroy
                    0.00%  154.58us       393     393ns     230ns  7.6360us  cuPointerGetAttribute
                    0.00%  25.972us         2  12.986us  11.569us  14.403us  cuStreamCreate
                    0.00%  24.368us         2  12.184us  7.1720us  17.196us  cuStreamDestroy
                    0.00%  8.0820us         2  4.0410us  2.6660us  5.4160us  cuStreamSynchronize
                    0.00%  4.0600us         2  2.0300us  1.4350us  2.6250us  cuCtxSynchronize
                    0.00%  2.7130us         1  2.7130us  2.7130us  2.7130us  cuDeviceGetPCIBusId
                    0.00%  2.1270us         4     531ns     307ns     748ns  cuCtxSetCurrent
                    0.00%  2.0830us         2  1.0410us     612ns  1.4710us  cuModuleGetFunction
                    0.00%  1.5310us         3     510ns     191ns  1.0640us  cuDeviceGetCount
                    0.00%  1.1070us         2     553ns     461ns     646ns  cuCtxPopCurrent
                    0.00%     733ns         2     366ns     199ns     534ns  cuDeviceGet

==44046== Unified Memory profiling result:
Device "GeForce GTX 1080 (0)"
   Count  Avg Size  Min Size  Max Size  Total Size  Total Time  Name
   28880  75.789KB  4.0000KB  0.9961MB  2.087399GB  214.8492ms  Host To Device
   13520  151.98KB  4.0000KB  0.9961MB  1.959534GB  169.6310ms  Device To Host
    9274         -         -         -           -  860.9489ms  Gpu page fault groups
     366  4.0000KB  4.0000KB  4.0000KB  1.429688MB           -  Memory thrashes
Total CPU Page faults: 8063
Total CPU thrashes: 366

CUDA execution

For reference, here is the execution with CUDA 11.0 and the master branch. There are some issues with preprocessor logic that breaks the sycl-clean branch when compiling for CUDA. I do not know what is wrong there yet.

$ QS_DEVICE=GPU /usr/local/cuda-11.0/bin/nvprof ./qs 
Copyright (c) 2016
Lawrence Livermore National Security, LLC
All Rights Reserved
Quicksilver Version     : 2020-Feb-4-22:35:56
Quicksilver Git Hash    : af27b3dcce08933786cb526e2f8a0bbe99d99b07
MPI Version             : 3.0
Number of MPI ranks     : 1
Number of OpenMP Threads: 1
Number of OpenMP CPUs   : 1

Simulation:
   dt: 1e-08
   fMax: 0.1
   inputFile: 
   energySpectrum: 
   boundaryCondition: reflect
   loadBalance: 0
   cycleTimers: 0
   debugThreads: 0
   lx: 100
   ly: 100
   lz: 100
   nParticles: 1000000
   batchSize: 0
   nBatches: 10
   nSteps: 10
   nx: 10
   ny: 10
   nz: 10
   seed: 1029384756
   xDom: 0
   yDom: 0
   zDom: 0
   eMax: 20
   eMin: 1e-09
   nGroups: 230
   lowWeightCutoff: 0.001
   bTally: 1
   fTally: 1
   cTally: 1
   coralBenchmark: 0
   crossSectionsOut:

Geometry:
   material: sourceMaterial
   shape: brick
   xMax: 100
   xMin: 0
   yMax: 100
   yMin: 0
   zMax: 100
   zMin: 0

Material:
   name: sourceMaterial
   mass: 1000
   nIsotopes: 10
   nReactions: 9
   sourceRate: 1e+10
   totalCrossSection: 1
   absorptionCrossSection: flat
   fissionCrossSection: flat
   scatteringCrossSection: flat
   absorptionCrossSectionRatio: 1
   fissionCrossSectionRatio: 0.1
   scatteringCrossSectionRatio: 1

CrossSection:
   name: flat
   A: 0
   B: 0
   C: 0
   D: 0
   E: 1
   nuBar: 2.4
==71046== NVPROF is profiling process 71046, command: ./qs
Building partition 0
Building partition 1
Building partition 2
Building partition 3
done building
Building MC_Domain 0
Building MC_Domain 1
Building MC_Domain 2
Building MC_Domain 3
Starting Consistency Check
Finished Consistency Check
Finished initMesh
cycle           start       source           rr        split       absorb      scatter      fission      produce      collisn       escape       census      num_seg   scalar_flux      cycleInit  cycleTracking  cycleFinalize
       0            0       100000            0       900000      1078182      1076792       107133       257364      2262107            0        72049      2670386  2.264064e+08   9.041900e-02   4.416690e-01   4.400000e-05
       1        72049       100000            0       828008      1107255      1106235       110306       264657      2323796            0        47153      2719702  2.438830e+08   1.405880e-01   4.555290e-01   4.300000e-05
       2        47153       100000            0       852712      1086097      1088696       108334       259738      2283127            0        65172      2687840  2.435394e+08   1.361020e-01   4.307970e-01   4.100000e-05
       3        65172       100000        68015       834785      1017555      1018659       101778       244593      2137992            0        57202      2517378  2.450517e+08   1.369190e-01   4.161950e-01   4.200000e-05
       4        57202       100000        62214       842934      1020418      1019522       101687       244038      2141627            0        59855      2522163  2.434017e+08   1.338490e-01   4.129190e-01   4.200000e-05
       5        59855       100000        56726       840345      1029994      1029682       103183       247672      2162859            0        57969      2545713  2.451216e+08   1.335090e-01   4.229580e-01   4.200000e-05
       6        57969       100000        52439       841925      1032190      1032180       102801       246877      2167171            0        59341      2551468  2.446226e+08   1.345140e-01   4.111950e-01   4.300000e-05
       7        59341       100000        59663       840635      1023444      1022593       102792       246649      2148829            0        60726      2531066  2.441845e+08   1.343080e-01   4.148660e-01   4.200000e-05
       8        60726       100000        68187       839357      1013501      1014287       101238       243112      2129026            0        60269      2508491  2.440307e+08   1.345430e-01   4.114840e-01   4.400000e-05
       9        60269       100000        71159       839953      1012439      1011892       101368       243262      2125699            0        58518      2500968  2.444142e+08   1.341150e-01   4.092590e-01   4.200000e-05

Timer                       Cumulative   Cumulative   Cumulative   Cumulative   Cumulative   Cumulative
Name                            number    microSecs    microSecs    microSecs    microSecs   Efficiency
                              of calls          min          avg          max       stddev       Rating
main                                 1    5.552e+06    5.552e+06    5.552e+06    0.000e+00       100.00
cycleInit                           10    1.309e+06    1.309e+06    1.309e+06    0.000e+00       100.00
cycleTracking                       10    4.227e+06    4.227e+06    4.227e+06    0.000e+00       100.00
cycleTracking_Kernel               992    3.889e+06    3.889e+06    3.889e+06    0.000e+00       100.00
cycleTracking_MPI                 1083    3.375e+05    3.375e+05    3.375e+05    0.000e+00       100.00
cycleTracking_Test_Done              0    0.000e+00    0.000e+00    0.000e+00    0.000e+00         0.00
cycleFinalize                       20    3.602e+03    3.602e+03    3.602e+03    0.000e+00       100.00
Figure Of Merit              6.093e+06 [Num Segments / Cycle Tracking Time]
==71046== Profiling application: ./qs
==71046== Profiling result:
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
 GPU activities:  100.00%  3.84139s       195  19.699ms  254.05us  38.108ms  CycleTrackingKernel(MonteCarlo*, int, ParticleVault*, ParticleVault*)
                    0.00%  1.3760us         1  1.3760us  1.3760us  1.3760us  _GLOBAL__N__48_tmpxft_00011231_00000000_7_cudaFunctions_cpp1_ii_2bb9853e::WarmUpKernel(void)
      API calls:   92.43%  3.84205s       196  19.602ms  10.971us  38.107ms  cudaDeviceSynchronize
                    6.96%  289.13ms       393  735.69us  5.5720us  284.39ms  cudaMallocManaged
                    0.50%  20.979ms       393  53.380us  9.7610us  742.77us  cudaFree
                    0.09%  3.8650ms       196  19.719us  8.9830us  202.60us  cudaLaunchKernel
                    0.01%  272.25us       101  2.6950us     183ns  170.40us  cuDeviceGetAttribute
                    0.01%  236.49us         1  236.49us  236.49us  236.49us  cuDeviceTotalMem
                    0.00%  63.885us       195     327ns     253ns  3.6680us  cudaPeekAtLastError
                    0.00%  26.808us         1  26.808us  26.808us  26.808us  cuDeviceGetName
                    0.00%  2.8380us         1  2.8380us  2.8380us  2.8380us  cuDeviceGetPCIBusId
                    0.00%  2.8210us         1  2.8210us  2.8210us  2.8210us  cudaSetDevice
                    0.00%  1.5900us         3     530ns     236ns  1.0870us  cuDeviceGetCount
                    0.00%     873ns         2     436ns     192ns     681ns  cuDeviceGet
                    0.00%     447ns         1     447ns     447ns     447ns  cudaGetDeviceCount
                    0.00%     371ns         1     371ns     371ns     371ns  cuDeviceGetUuid

==71046== Unified Memory profiling result:
Device "GeForce GTX 1080 (0)"
   Count  Avg Size  Min Size  Max Size  Total Size  Total Time  Name
   31624  69.160KB  4.0000KB  0.9961MB  2.085815GB  217.8466ms  Host To Device
   13544  151.74KB  4.0000KB  0.9961MB  1.960014GB  169.6492ms  Device To Host
    8236         -         -         -           -  916.3543ms  Gpu page fault groups
     258  4.0000KB  4.0000KB  4.0000KB  1.007813MB           -  Memory thrashes
Total CPU Page faults: 8087
Total CPU thrashes: 258

Signed-off-by: Jeff Hammond <[email protected]>

Jeff Hammond added 7 commits June 3, 2020 13:38

port to SYCL/DPC++

451f9ba

look for HAVE_SYCL and SYCL_EXTERNAL for most of the changes. this works on GPU and HOST but not CPU (OpenCL issue). Signed-off-by: Jeff Hammond <[email protected]>

better compiler flags

516be3e

implement SYCL atomics

7fa9b9d

re-enable printf in device code using Intel extension

d14bbb5

Signed-off-by: Jeff Hammond <[email protected]>

re-enable printf in device code using Intel extension

c6f20a6

Signed-off-by: Jeff Hammond <[email protected]>

re-enable printf in device code using Intel extension

a50696b

Signed-off-by: Jeff Hammond <[email protected]>

globally rename q->sycl_device_queue

7794a66

Signed-off-by: Jeff Hammond <[email protected]>

jeffhammond and others added 5 commits October 9, 2020 13:06

add DPC++ PTX stuff

96f8368

Signed-off-by: Jeff Hammond <[email protected]>

deal with broken printf w/ DPC++ NVIDIA backend

5512778

Signed-off-by: Jeff Hammond <[email protected]>

build system changes

5d992a5

flag workaround compiler bug

3ae32fa

disable unrelated flags

0b9c632

jeffhammond mentioned this pull request Oct 22, 2023

Quicksilver SYCL version #65

Open

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

SYCL implementation #44

SYCL implementation #44

jeffhammond commented Jun 27, 2020

jeffhammond commented Jun 27, 2020

jeffhammond commented Oct 9, 2020 •

edited

Loading

SYCL implementation #44

Are you sure you want to change the base?

SYCL implementation #44

Conversation

jeffhammond commented Jun 27, 2020

jeffhammond commented Jun 27, 2020

jeffhammond commented Oct 9, 2020 • edited Loading

Compiler

Hardware

Host execution

NVIDIA execution

CUDA execution

jeffhammond commented Oct 9, 2020 •

edited

Loading