 Peter Boyle
|
7f99e1cd3b
|
Coalesced loops
|
2019-06-15 08:01:39 +01:00 |
|
|
Peter Boyle
|
f3c89df948
|
Thread loop changes
|
2019-06-15 08:00:37 +01:00 |
|
|
Peter Boyle
|
b7e6d111d7
|
Thread loop changes. Need to offload this file
|
2019-06-15 07:59:10 +01:00 |
|
|
Peter Boyle
|
f39cf69c33
|
Accelerator loop change
|
2019-06-15 07:58:23 +01:00 |
|
|
Peter Boyle
|
8e27338df2
|
Rationalise number of loop macros
|
2019-06-15 07:57:40 +01:00 |
|
|
Peter Boyle
|
bcbb5e9d26
|
Remove assembly tests
|
2019-06-15 07:57:05 +01:00 |
|
|
Peter Boyle
|
0ea7f5279d
|
Accelerator loop changes
|
2019-06-15 07:56:14 +01:00 |
|
|
Peter Boyle
|
18e5de426d
|
There is a stray use of predicatedWhere introduced by Andrew Lawson in the conserve currents.
The conserved currents need rewritten using data parallel operations.
|
2019-06-15 07:53:58 +01:00 |
|
|
Peter Boyle
|
e896d81235
|
Accelerator loop redefine. Coalesce most accesses, but ET engine still to go clean.
|
2019-06-15 07:52:44 +01:00 |
|
|
Peter Boyle
|
7b8ccff4f4
|
Accelerated coalesced loops in most cases
|
2019-06-15 07:48:00 +01:00 |
|
|
Peter Boyle
|
68541606ab
|
Thread loop changes. Soon try these with accelerator loops and benchmark
|
2019-06-15 07:46:42 +01:00 |
|
|
Peter Boyle
|
339ea10cc7
|
First touch only on CPU code
|
2019-06-15 07:45:43 +01:00 |
|
|
Peter Boyle
|
d0d8dc8042
|
Thread loop changes
|
2019-06-15 07:45:09 +01:00 |
|
|
Peter Boyle
|
81eb1fd9f2
|
Accelerator loop changes for coalesced access
|
2019-06-15 07:44:47 +01:00 |
|
|
Peter Boyle
|
cb93d32cd9
|
Thread loop changes
|
2019-06-15 07:44:08 +01:00 |
|
|
Peter Boyle
|
8f223962ff
|
Thread loop changed
|
2019-06-15 07:43:42 +01:00 |
|
|
Peter Boyle
|
36f06555a2
|
Simplify Impl
|
2019-06-09 22:26:27 +01:00 |
|
|
Peter Boyle
|
d6c0e0756d
|
Remove GPU version
|
2019-06-09 11:23:42 +01:00 |
|
|
Peter Boyle
|
3e41b1055c
|
Remove Gpu only kernels.
|
2019-06-09 11:20:01 +01:00 |
|
|
Peter Boyle
|
9fbcfe612c
|
Update TODO list
|
2019-06-09 11:19:38 +01:00 |
|
|
Peter Boyle
|
e78a5e7838
|
ASM instantiation without link errors
|
2019-06-09 01:25:21 +01:00 |
|
|
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 |
|