portelli/Grid
1
0
Fork 0
mirror of https://github.com/paboyle/Grid.git synced 2026-07-17 15:43:27 +01:00
Code Releases Activity
4,965 Commits 95 Branches 14 Tags
f710d7bd45b7ad2e59953006bf4bbfc0b5055416
Commit Graph
100 Commits
This Branch
This Branch
All Branches
Author SHA1 Message Date
Peter Boyle f710d7bd45 TODO list update 2019-06-15 12:54:27 +01:00
Peter Boyle cb336aa8f8 Thread loop constructs changing a little 2019-06-15 12:54:11 +01:00
Peter Boyle 462900b48d Modified entire test directory to suit new GPU constructs for looping 2019-06-15 12:53:27 +01:00
Peter Boyle 0561c2edeb Benchmarks modified for new GPU constructs 2019-06-15 12:52:56 +01:00
Peter Boyle 0184719216 Change to predicate type 2019-06-15 12:52:26 +01:00
Peter Boyle 24202dbc51 Thread loop construct change 2019-06-15 12:52:07 +01:00
Peter Boyle d763c303c5 Clean acceleerator barrier 2019-06-15 12:51:45 +01:00
Peter Boyle 8e394d3bf9 New loop construct 2019-06-15 12:51:15 +01:00
Peter Boyle b881d5489b Move SchurDiagTwoKappa to Algorithms 2019-06-15 12:50:45 +01:00
Peter Boyle 82306913a8 Move Schur operator into correct place 2019-06-15 12:49:22 +01:00
Peter Boyle 49f90cc7eb use pragma once 2019-06-15 12:45:22 +01:00
Peter Boyle b77af0210b Thread loop. Probably deprecate this impl 2019-06-15 12:44:56 +01:00
Peter Boyle 5254ede2d8 New loops. Revisit as accelerator loop in future audit 2019-06-15 12:44:29 +01:00
Peter Boyle 16e5d7945e Hard to make 5D vec work with GPU code 2019-06-15 12:43:43 +01:00
Peter Boyle decc99ca76 Accelerator version 2019-06-15 12:43:00 +01:00
Peter Boyle 464cd65931 Still to test this fully 2019-06-15 12:35:14 +01:00
Peter Boyle a1ec2f4723 Still to test this routine fully 2019-06-15 12:33:55 +01:00
Peter Boyle ea9662ec85 Thread loop changes 2019-06-15 09:09:57 +01:00
Peter Boyle 52c74f1cac Thread loop changes 2019-06-15 09:08:16 +01:00
Peter Boyle 9a13d2992c lean up 2019-06-15 09:05:16 +01:00
Peter Boyle b0449ae270 Thread loop changes 2019-06-15 09:04:19 +01:00
Peter Boyle 1299225105 Accelerator loop changes 2019-06-15 09:03:46 +01:00
Peter Boyle 5925e7f405 Thread for changes 2019-06-15 09:01:30 +01:00
Peter Boyle be1fd4930f Template instantiation make happy changes 2019-06-15 08:37:34 +01:00
Peter Boyle 377fa5dec1 looping construct 2019-06-15 08:36:48 +01:00
Peter Boyle e8b78f596e Looping construct changes 2019-06-15 08:35:57 +01:00
Peter Boyle 09720c40cd Coalesced loops 2019-06-15 08:35:26 +01:00
Peter Boyle bb024dd114 Loop construct changed 2019-06-15 08:30:05 +01:00
Peter Boyle 52456b9ec7 New loop construct 2019-06-15 08:28:45 +01:00
Peter Boyle b285138be4 Better checking on types 2019-06-15 08:27:48 +01:00
Peter Boyle c7dbf4c87e Scalar support for GPU threads 2019-06-15 08:25:43 +01:00
Peter Boyle 1e889c93b8 Insert a GPU synchronise 2019-06-15 08:23:26 +01:00
Peter Boyle 7379047482 Threading and acceleration primitives further changes. accelerator_barrier() needed and used 2019-06-15 08:22:48 +01:00
Peter Boyle d836ce3b78 Clean up of acceleration and threading primitives 2019-06-15 08:14:21 +01:00
Peter Boyle cefaacbc07 Changing accelerator loop. Still have work to do for multi-GPU code 2019-06-15 08:10:24 +01:00
Peter Boyle 0074ef7f69 thread loops 2019-06-15 08:04:29 +01:00
Peter Boyle 20359ca15f Coalesced loops. 2019-06-15 08:03:57 +01:00
Peter Boyle 736358b0cb Coalesced loops 2019-06-15 08:03:13 +01:00
Peter Boyle 6b692aa726 Thread loops 2019-06-15 08:02:26 +01:00
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