 Peter Boyle
|
da8d87e9da
|
Cuda switch off
|
2019-06-08 17:11:38 +01:00 |
|
|
Peter Boyle
|
8e3a05d89b
|
Moving the instantiation into a cleaner structure
|
2019-06-08 13:48:33 +01:00 |
|
|
Peter Boyle
|
8adc5da7dd
|
Testig out approaches to kernel writing introducing SIMT_loop temporarily
|
2019-06-08 13:47:04 +01:00 |
|
|
Peter Boyle
|
29a244e423
|
Test of using a lane variable instead of repeated reference to threadIdx.y
|
2019-06-08 13:46:26 +01:00 |
|
|
Peter Boyle
|
18cbfecf02
|
Use symlinks in find command
|
2019-06-08 13:45:46 +01:00 |
|
|
Peter Boyle
|
c933ac2248
|
Temporarily introduce a SIMT_loop to test out approaches prior to making a global change to
accelerator_loop
|
2019-06-08 13:44:27 +01:00 |
|
|
Peter Boyle
|
ad2c433574
|
Instantiations move. Tried using Gianluca's suggestion about avoiding threadIdx but doesn't
seem to make a difference. Will revisit this and probably remove the lane parameter from the coalescedRead
|
2019-06-08 13:43:12 +01:00 |
|
|
Peter Boyle
|
86e7fb6e86
|
Instantiation relocation
|
2019-06-08 13:42:46 +01:00 |
|
|
Peter Boyle
|
fb91dda7be
|
Hand instantiation moved location
|
2019-06-08 13:42:26 +01:00 |
|
|
Peter Boyle
|
82cf7bc5ab
|
Move instantiation into fermion/instantiation
|
2019-06-08 13:41:46 +01:00 |
|
|
Peter Boyle
|
e452cc0a22
|
Move static variables into instantiation .cc file
|
2019-06-08 13:41:20 +01:00 |
|
|
Peter Boyle
|
4d2b938166
|
Remove explict instantiation from here
|
2019-06-08 13:41:01 +01:00 |
|
|
Peter Boyle
|
10d16ab76c
|
Remove explict instantiation from here
|
2019-06-08 13:40:32 +01:00 |
|
|
Peter Boyle
|
1f997fa484
|
Instantiate via explict .cc files for parallel make.
|
2019-06-08 13:39:51 +01:00 |
|
|
Peter Boyle
|
dc5024e88c
|
The GPU reduction was not working for me and causing errors. Need to revisit.
Gianluca is working on deterministic reduction/
|
2019-06-08 13:39:11 +01:00 |
|
|
Peter Boyle
|
6d77941990
|
Drop the 5D vec actions
|
2019-06-08 13:38:05 +01:00 |
|
|
Peter Boyle
|
0ee6e77cbc
|
Compiles GPU and CPU, still gives good performance on CPU
|
2019-06-05 13:28:16 +01:00 |
|
|
Peter Boyle
|
18d3cde29a
|
Compile on GPU workd
|
2019-06-05 00:14:58 +01:00 |
|
|
Peter Boyle
|
7323099966
|
Instatiation fix
|
2019-06-05 00:14:38 +01:00 |
|
|
Peter Boyle
|
6379651cdd
|
Generic or GPU ready for benchmark test on GPU
|
2019-06-05 00:13:52 +01:00 |
|
|
Peter Boyle
|
ba4fd756b9
|
Fix signature, but deprecating this loops style
|
2019-06-05 00:12:36 +01:00 |
|
|
Peter Boyle
|
d185fc1ebf
|
clean up instantiation
|
2019-06-05 00:11:52 +01:00 |
|
|
Peter Boyle
|
96b36d8367
|
Instantiation clean up
|
2019-06-05 00:11:27 +01:00 |
|
|
Peter Boyle
|
899f8b5065
|
Instantiation clean up 5d vec removal
|
2019-06-05 00:11:05 +01:00 |
|
|
Peter Boyle
|
c8d0483fe9
|
Remove 5d vectorisation
|
2019-06-05 00:10:37 +01:00 |
|
|
Peter Boyle
|
0f214e5f76
|
Clean up instantiation
|
2019-06-05 00:10:13 +01:00 |
|
|
Peter Boyle
|
8eea568426
|
GPU loop ; presently differentiated with ifdef, find a way to unify.
|
2019-06-05 00:09:28 +01:00 |
|
|
Peter Boyle
|
9636324069
|
GPU happy code
|
2019-06-05 00:08:54 +01:00 |
|
|
Peter Boyle
|
8a5489d9e6
|
Move the loop into a central kernel call.
|
2019-06-05 00:08:13 +01:00 |
|
|
Peter Boyle
|
8113845f9c
|
coalesce loop. Need to rationalise this file
|
2019-06-04 23:49:29 +01:00 |
|
|
Peter Boyle
|
b47f73c222
|
GPU happy
|
2019-06-04 21:30:39 +01:00 |
|
|
Peter Boyle
|
5720ced0fd
|
Simplifying
|
2019-06-04 21:30:08 +01:00 |
|
|
Peter Boyle
|
2c87b56b53
|
Making GPU happier
|
2019-06-04 21:29:44 +01:00 |
|
|
Peter Boyle
|
dbad48d802
|
Remove Ls vectorised DWF
|
2019-06-04 21:27:40 +01:00 |
|
|
Peter Boyle
|
4557a1365a
|
Remove Ls vectorised DWF
|
2019-06-04 20:59:59 +01:00 |
|
|
Peter Boyle
|
16e9b87d98
|
Remove Ls vectorised DWF as unused and hard to maintain
|
2019-06-04 20:59:01 +01:00 |
|
|
Peter Boyle
|
685eea3d0f
|
Small cosmetic
|
2019-06-04 20:58:14 +01:00 |
|
|
Peter Boyle
|
65b48831fb
|
Simplify code
|
2019-06-04 20:56:30 +01:00 |
|
|
Peter Boyle
|
57396fc595
|
Simplify code
|
2019-06-04 20:56:23 +01:00 |
|
|
Peter Boyle
|
a2e199df50
|
Simplifying Cayley cases.
|
2019-06-04 20:54:52 +01:00 |
|
|
Peter Boyle
|
020346c848
|
WOrk list. Will have to clean up Fermion sector.
|
2019-06-04 20:54:00 +01:00 |
|
|
Peter Boyle
|
c2625a127e
|
Non blocking loop. Want to change the naming here.
|
2019-06-04 20:52:59 +01:00 |
|
|
Peter Boyle
|
8794d35c78
|
GPU
|
2019-06-04 20:52:27 +01:00 |
|
|
Peter Boyle
|
24bff6dbe6
|
Minor improvements
|
2019-06-04 20:51:48 +01:00 |
|
|
Peter Boyle
|
45b15d10d3
|
GPU happy changes
|
2019-06-04 20:49:16 +01:00 |
|
|
Peter Boyle
|
33d6bbe32b
|
GPU must use accelerator vectors
|
2019-06-04 20:48:52 +01:00 |
|
|
Peter Boyle
|
7a1569bd46
|
Annoying, cannot rely on equivalence of Grid ComplexD adn Eigen Complex type on GPU.
Solve with ComplexD typecasts but must be a better way
|
2019-06-04 20:47:49 +01:00 |
|
|
Peter Boyle
|
6e2e904a0e
|
NVCC compiles happy. Start to develop strategy for writing generic
code for GPU kernels and CPU kernels.
|
2019-06-04 20:46:35 +01:00 |
|
|
Peter Boyle
|
d92a17f359
|
Suppress NVCC warnings in pugixml with pragma
|
2019-06-04 20:45:53 +01:00 |
|
|
Peter Boyle
|
47c063f984
|
Remove Ls Vec cases from benchmarks
|
2019-06-04 20:45:35 +01:00 |
|
|
Peter Boyle
|
7e27a5213a
|
Tests builds clean.
|
2019-06-04 20:45:20 +01:00 |
|
|
Peter Boyle
|
ade4a126da
|
Getting closer on the GPU port, but will start deleting 5th dim vectorised variants
for code maintainability
|
2019-06-04 11:53:44 +01:00 |
|
|
Peter Boyle
|
7b59ab5bd7
|
Compiling after reorganisation
|
2019-06-03 15:46:26 +01:00 |
|
|
Peter Boyle
|
fcd8cfe257
|
Gparity in
|
2019-06-03 15:45:09 +01:00 |
|
|
Peter Boyle
|
b4b53812cb
|
Move implementation to specific implementation headers
|
2019-06-03 15:43:01 +01:00 |
|
|
Peter Boyle
|
085cac583f
|
Implementation in header
|
2019-06-03 15:42:36 +01:00 |
|
|
Peter Boyle
|
25e3b8640c
|
Move to header
|
2019-06-03 15:42:05 +01:00 |
|
|
Peter Boyle
|
44bbec50b0
|
Making GPU compile happy
|
2019-06-03 14:57:04 +01:00 |
|
|
Peter Boyle
|
ec68b67d5d
|
Attempt at unified GPU and CPU kernel
|
2019-06-03 14:55:51 +01:00 |
|
|
Peter Boyle
|
778450e0c8
|
Move to implementation subdir
|
2019-06-03 14:53:56 +01:00 |
|
|
Peter Boyle
|
567aa5f366
|
Move to implementation subdir
|
2019-06-03 14:53:33 +01:00 |
|
|
Peter Boyle
|
2ab7e2b175
|
Force instantiation in .cc files.
Eventually move into multiple files
|
2019-06-03 14:52:59 +01:00 |
|
|
Peter Boyle
|
6f61be044d
|
Dont instantiate in header
|
2019-06-03 14:52:01 +01:00 |
|
|
Peter Boyle
|
269e00509e
|
Don't instantiate in header
|
2019-06-03 14:51:24 +01:00 |
|
|
Peter Boyle
|
a5e90b0ddc
|
Making the kernels more GPU happy
|
2019-06-03 14:50:54 +01:00 |
|
|
Peter Boyle
|
5622faf226
|
pragma once ifdef guard
|
2019-06-03 14:50:26 +01:00 |
|
|
Peter Boyle
|
82ecd520c7
|
Macos happy fix under nvcc
|
2019-06-03 14:48:50 +01:00 |
|
|
Peter Boyle
|
ffde81f22a
|
Nsimd() and coalesced support
|
2019-05-25 12:44:07 +01:00 |
|
|
Peter Boyle
|
d8098f1ecd
|
coalesced support
|
2019-05-25 12:43:31 +01:00 |
|
|
Peter Boyle
|
aca788cf4f
|
Move coalesced read into tensors
|
2019-05-25 12:43:00 +01:00 |
|
|
Peter Boyle
|
a0e9f3b0a0
|
Plan for GPU port
|
2019-05-20 09:46:19 +01:00 |
|
|
Peter Boyle
|
a9342c6ae5
|
Udpdate TODO afer gianluc marge
|
2019-05-18 22:58:25 +01:00 |
|
|
Peter Boyle
|
4e9df9e93c
|
GPU patches
|
2019-05-18 17:43:11 +01:00 |
|
|
Peter Boyle
|
9fe68857a9
|
Runs multiGPU with coalesced access on tesseract
|
2019-05-18 17:42:41 +01:00 |
|
|
Peter Boyle
|
37336c9e0c
|
Allow compress to be either vector or scalar types
|
2019-05-18 17:41:13 +01:00 |
|
|
Peter Boyle
|
6c4da3bbc7
|
Stencil now runs with coalesced accesses
|
2019-05-18 17:40:35 +01:00 |
|
|
Peter Boyle
|
a584b16c4a
|
Adding a non-blocking kernel launch
|
2019-05-18 17:39:54 +01:00 |
|
|
Peter Boyle
|
60330e05a3
|
NVCC wacky compiler options frozen. Possibly Cuda 9.2 specific
|
2019-04-28 07:39:33 +01:00 |
|
|
Peter Boyle
|
f9b8c0cccf
|
Vector changes for UVM
|
2019-04-28 07:38:57 +01:00 |
|
|
Peter Boyle
|
3cad67e569
|
Compile on tesseract
|
2019-04-28 07:38:09 +01:00 |
|
|
Peter Boyle
|
170ba4e619
|
Ensure different MPI ranks use different GPUs. The mapping works on Tesseract.
|
2019-04-28 07:32:30 +01:00 |
|
|
Peter Boyle
|
204a090497
|
Inner product is not working on GPU. Why?
|
2019-04-28 07:31:56 +01:00 |
|
|
Peter Boyle
|
3c717c47ef
|
GPU no compile on Wilson Multigrid fixed
|
2019-04-28 07:31:19 +01:00 |
|
|
Peter Boyle
|
c5e081d69c
|
Re-Merge branch 'develop' into feature/gpu-port
Pull in Regensburg MultiGrid pull request
|
2019-01-03 01:50:16 +00:00 |
|
|
Peter Boyle
|
535a6aaf05
|
Update todo list
|
2019-01-02 22:07:51 +00:00 |
|
|
Peter Boyle
|
91a7fe247b
|
Merge branch 'DanielRichtmann-feature/wilsonmg' into develop
|
2019-01-02 14:40:31 +00:00 |
|
|
Peter Boyle
|
8a1be021d3
|
Merge branch 'feature/wilsonmg' of https://github.com/DanielRichtmann/Grid into DanielRichtmann-feature/wilsonmg
|
2019-01-02 14:39:59 +00:00 |
|
|
Peter Boyle
|
e73b909a48
|
Make tests running past nvcc. Different NVCC versions proving tricky to keep happy. This is 9.2
|
2019-01-02 12:05:30 +00:00 |
|
|
Peter Boyle
|
a4d9200293
|
Fixing AVX 512 instantiation error. Need to move to extern templates urgently.
|
2019-01-02 00:27:07 +00:00 |
|
|
Peter Boyle
|
350508bdb3
|
pugixml problem
|
2019-01-01 16:38:54 +00:00 |
|
|
Peter Boyle
|
38852737e4
|
No compile fix on clang
|
2019-01-01 15:55:13 +00:00 |
|
|
Peter Boyle
|
802404c78c
|
Remove warnings under NVCC and move parallel_for to thread-loop
|
2019-01-01 15:08:09 +00:00 |
|
|
Peter Boyle
|
0e9b591c1c
|
NVCC warning suppression
|
2019-01-01 15:07:47 +00:00 |
|
|
Peter Boyle
|
c43a2b599a
|
GPU support
|
2019-01-01 15:07:29 +00:00 |
|
|
Peter Boyle
|
8c91e82ee8
|
GPU clean up, remove parallel_for. Split into accelerator_loop, thread_loop
cases, and collides with parallel_for in thrust
|
2019-01-01 15:06:46 +00:00 |
|
|
Peter Boyle
|
9d866d062a
|
GPU support improvements
|
2019-01-01 15:05:03 +00:00 |
|
|
Peter Boyle
|
3a4e397e72
|
Deprecating JSON, too hard to support under NVCC
|
2019-01-01 15:04:33 +00:00 |
|
|
Peter Boyle
|
2b6cfe555f
|
Disable JSON on NVCC. Maybe unsupport JSON full stop. XML and JSON is too many formats in my view.
|
2019-01-01 15:03:50 +00:00 |
|
|
Peter Boyle
|
7df58dd883
|
Photon syntax gave problems with NVCC
|
2019-01-01 15:03:29 +00:00 |
|
|
Peter Boyle
|
4bf86ae60a
|
NVCC clean up
|
2019-01-01 15:02:50 +00:00 |
|