Commit Graph - Grid - DiRAC Tursa git server

portelli/Grid

0 Fork 0

mirror of https://github.com/paboyle/Grid.git synced 2026-11-07 04:33:19 +00:00

Commit Graph

Select branches

Hide Pull Requests

3c67d626ba05

FgridStaggered

KS_shifted

XmlIO

bugfix/dminus

chulwoo-dec12-2015

ckelly-dec12-2015

debug-crusher

develop

feature/BCG

feature/CG_repro

feature/Kpipi-masaaki-offload

feature/Lanczos

feature/S2xR

feature/a2a-integration

feature/a2a-offload

feature/adaptive_wflow

feature/benchiotimings

feature/benchmark-io-update

feature/bgq-asm

feature/block-cg

feature/block_lanczos

feature/block_lanczos22

feature/boosted

feature/cache-fix

feature/ckelly-gparity-merge

feature/clover

feature/conjugate-bc-dirs

feature/contractor

feature/cpu-threaded-smp

feature/ddhmc

feature/deprecate-uvm

feature/dirichlet

feature/dirichlet-gparity

feature/dwf-multirhs

feature/dwf-preconditioning

feature/eigen-relative

feature/feature/staggered-comms

feature/felix-mm-debug

feature/felix-slice-sum-fast

feature/fermion-laplace

feature/fermtoprop

feature/fthmc

feature/fthmc-optimise

feature/fugaku-optimisation

feature/gparity-merge-april24

feature/gparity_HMC

feature/gpu-port

feature/hadrons-a2a

feature/hadrons-twist

feature/half-prec-comms

feature/hdcr

feature/hw-multigrid

feature/knl-stats

feature/lanczos-reorg

feature/lanczos-simplify

feature/luchang-rng

feature/mixedCG

feature/mpi3

feature/multi-communicator

feature/normHP

feature/npr

feature/parallelio

feature/reduction-reorganisation

feature/resilient-io

feature/rmhmc

feature/scidacRNG

feature/serialisation-update

feature/shGordon

feature/shm-chmod

feature/sitmo-skipahead

feature/splitGridCGBatch

feature/staggered-comms-compute

feature/staggered-hdcg

feature/staggered-merge

feature/staggering

feature/sumd-npr

feature/summit-lanczos

feature/summit-multigrid

feature/sycl-simt

feature/verify-checksum-host

fix/HOST_NAME_MAX

gauge-group-covariance

gh-pages

hotfix/nvcc-warnings

hotfix/unwind-aarch64

hotfix/virtual-dtor

master

release/0.8.0

release/0.8.2

release/0.9.1

release/dirac-ITT

release/v0.7.0

rmhmc_merge2

specflow

syck

sycl

sycl-linking-hack

#1

#110

#114

#116

#118

#120

#121

#123

#124

#128

#129

#13

#13/merge

#130

#133

#135

#136

#138

#14

#145

#148

#148/merge

#150

#151

#152

#155

#157

#158

#159

#16

#163

#166

#168

#171

#172

#173

#174

#176

#177

#178

#180

#182

#183

#185

#186

#187

#188

#189

#19

#19/merge

#190

#191

#193

#195

#197

#199

#200

#201

#202

#203

#204

#205

#206

#207

#208

#209

#21

#210

#211

#212

#213

#214

#215

#216

#217

#218

#219

#22

#220

#221

#222

#224

#225

#226

#229

#231

#233

#235

#236

#238

#239

#24

#240

#241

#242

#243

#244

#246

#247

#248

#249

#250

#251

#252

#253

#254

#255

#256

#257

#259

#260

#261

#262

#263

#266

#267

#268

#269

#27

#27/merge

#270

#271

#272

#273

#274

#276

#277

#278

#279

#28

#280

#281

#282

#283

#287

#288

#29

#291

#292

#298

#304

#305

#307

#308

#309

#31

#310

#311

#312

#314

#316

#317

#318

#32

#325

#327

#328

#329

#33

#330

#332

#333

#335

#336

#337

#338

#339

#34

#340

#341

#342

#344

#347

#348

#349

#35

#354

#355

#356

#357

#358

#359

#360

#361

#363

#364

#365

#366

#367

#368

#369

#37

#370

#371

#375

#377

#379

#379/merge

#382

#384

#385

#387

#388

#389

#390

#391

#392

#394

#396

#397

#399

#4

#400

#401

#402

#403

#404

#405

#406

#407

#409

#410

#411

#412

#414

#415

#416

#417

#418

#419

#42

#42/merge

#421

#422

#424

#425

#426

#427

#428

#43

#431

#432

#433

#435

#437

#439

#44

#44/merge

#440

#441

#442

#444

#448

#449

#449/merge

#451

#453

#454

#455

#457

#458

#459

#461

#461/merge

#462

#462/merge

#465

#466

#469

#470

#471

#472

#473

#474

#474/merge

#475

#476

#476/merge

#477

#478

#479

#479/merge

#480

#481

#482

#483

#484

#485

#486

#488

#490

#491

#491/merge

#492

#493

#494

#494/merge

#495

#495/merge

#496

#496/merge

#497

#497/merge

#5

#50

#61

#62

#64

#65

#68

#69

#7

#72

#72/merge

#74

#75

#75/merge

#79

#80

#81

#84

#86

#87

#88

#89

#90

#91

#92

#92/merge

#93

0.8.1

0.8.2

DIRAC-ITT-2020

DiRAC-ITT-2020-UCX-WORKAROUND

ISC-freeze-1

ISC-freeze-2

dirac-ITT

dirac-ITT-fix

dirac-ITT-fix1

v0.5.0

v0.5.1

v0.6.0

v0.7.0

v0.8.0

d6ffadb33b Coalesced write Peter Boyle 2019-07-02 17:25:13 +01:00
ae3abbe53d Added the ability for Perambulator module to save unsmeared sinks through the addition of two optional parameters: UnsmearedSinkFileName: If present, specifies the filename to write to UnsmearedSinkMultiFile: defaults to true to write each sink vector to a different file, but can be set to 0 for a single file Michael Marshall 2019-07-01 17:28:27 +01:00
5fc0188205 started saving sinks Felix Erben 2019-07-01 14:51:59 +01:00
4c3225412b Drop 5dVEC Peter Boyle 2019-07-01 07:31:26 +01:00
b8f7bfbb26 Dont stream as poor perf in some cases Peter Boyle 2019-07-01 07:30:25 +01:00
7b7c470917 Accelerator loop Peter Boyle 2019-07-01 07:29:51 +01:00
532e226b22 cuda 9.1 fixes Peter Boyle 2019-07-01 07:29:22 +01:00
6a13731818 Move GPU cuda call earlier Peter Boyle 2019-07-01 07:28:41 +01:00
67690df3bd Changes nedded to have a current insertion on every second time slice - avoids unnecessary contractions fionnoh 2019-06-28 15:18:28 +08:00
1059189abf Bugfix for A2ALoop module fionnoh 2019-06-27 13:49:55 +08:00
ce29b18dc9 New modules for loading in MFs as diskvectors and producing propagaotrs from 4 quark contractions fionnoh 2019-06-27 13:46:06 +08:00
421a0a8a36 Changes to A2Autils, A2AMatirx and DiskVector code that is needed for Hadrons 4 quark contraction module fionnoh 2019-06-27 13:45:20 +08:00
ac530636ca A2Aloop bugfix fionnoh 2019-06-27 13:44:47 +08:00
2d940a598c Inserted four extra parameters just to make this test compile. Needs to be fixed properly Michael Marshall 2019-06-19 10:37:50 +01:00
c28c5fc61b Inserted four extra parameters just to make this test compile. Needs to be fixed properly Michael Marshall 2019-06-19 10:31:41 +01:00
015340d60c Elided superfluous copy on write Michael Marshall 2019-06-19 09:37:03 +01:00
1cd4ee0706 Thrust used on GPU builds Peter Boyle 2019-06-18 12:50:35 +01:00
b8f71b6777 Fix NVCC warning unused variable Peter Boyle 2019-06-17 13:58:45 +01:00
703dc20377 Compile tests fix Peter Boyle 2019-06-16 13:59:29 +01:00
d976e5c514 Pow is being awkward in thrust for reasons I don't understand. Possible thrust bug. Peter Boyle 2019-06-16 12:05:11 +01:00
d7b3efe893 Compile fix Peter Boyle 2019-06-15 17:03:15 +01:00
f710d7bd45 TODO list update Peter Boyle 2019-06-15 12:54:27 +01:00
cb336aa8f8 Thread loop constructs changing a little Peter Boyle 2019-06-15 12:54:11 +01:00
462900b48d Modified entire test directory to suit new GPU constructs for looping Peter Boyle 2019-06-15 12:53:27 +01:00
0561c2edeb Benchmarks modified for new GPU constructs Peter Boyle 2019-06-15 12:52:56 +01:00
0184719216 Change to predicate type Peter Boyle 2019-06-15 12:52:26 +01:00
24202dbc51 Thread loop construct change Peter Boyle 2019-06-15 12:52:07 +01:00
d763c303c5 Clean acceleerator barrier Peter Boyle 2019-06-15 12:51:45 +01:00
8e394d3bf9 New loop construct Peter Boyle 2019-06-15 12:51:15 +01:00
b881d5489b Move SchurDiagTwoKappa to Algorithms Peter Boyle 2019-06-15 12:50:45 +01:00
82306913a8 Move Schur operator into correct place Peter Boyle 2019-06-15 12:49:22 +01:00
49f90cc7eb use pragma once Peter Boyle 2019-06-15 12:45:22 +01:00
b77af0210b Thread loop. Probably deprecate this impl Peter Boyle 2019-06-15 12:44:56 +01:00
5254ede2d8 New loops. Revisit as accelerator loop in future audit Peter Boyle 2019-06-15 12:44:29 +01:00
16e5d7945e Hard to make 5D vec work with GPU code Peter Boyle 2019-06-15 12:43:43 +01:00
decc99ca76 Accelerator version Peter Boyle 2019-06-15 12:43:00 +01:00
464cd65931 Still to test this fully Peter Boyle 2019-06-15 12:35:14 +01:00
a1ec2f4723 Still to test this routine fully Peter Boyle 2019-06-15 12:33:55 +01:00
ea9662ec85 Thread loop changes Peter Boyle 2019-06-15 09:09:57 +01:00
52c74f1cac Thread loop changes Peter Boyle 2019-06-15 09:08:16 +01:00
9a13d2992c lean up Peter Boyle 2019-06-15 09:05:16 +01:00
b0449ae270 Thread loop changes Peter Boyle 2019-06-15 09:04:19 +01:00
1299225105 Accelerator loop changes Peter Boyle 2019-06-15 09:03:46 +01:00
5925e7f405 Thread for changes Peter Boyle 2019-06-15 09:01:30 +01:00
be1fd4930f Template instantiation make happy changes Peter Boyle 2019-06-15 08:37:34 +01:00
377fa5dec1 looping construct Peter Boyle 2019-06-15 08:36:48 +01:00
e8b78f596e Looping construct changes Peter Boyle 2019-06-15 08:35:57 +01:00
09720c40cd Coalesced loops Peter Boyle 2019-06-15 08:35:26 +01:00
bb024dd114 Loop construct changed Peter Boyle 2019-06-15 08:30:05 +01:00
52456b9ec7 New loop construct Peter Boyle 2019-06-15 08:28:45 +01:00
b285138be4 Better checking on types Peter Boyle 2019-06-15 08:27:48 +01:00
c7dbf4c87e Scalar support for GPU threads Peter Boyle 2019-06-15 08:25:43 +01:00
1e889c93b8 Insert a GPU synchronise Peter Boyle 2019-06-15 08:23:26 +01:00
7379047482 Threading and acceleration primitives further changes. accelerator_barrier() needed and used Peter Boyle 2019-06-15 08:22:48 +01:00
d836ce3b78 Clean up of acceleration and threading primitives Peter Boyle 2019-06-15 08:14:21 +01:00
cefaacbc07 Changing accelerator loop. Still have work to do for multi-GPU code Peter Boyle 2019-06-15 08:10:24 +01:00
0074ef7f69 thread loops Peter Boyle 2019-06-15 08:04:29 +01:00
20359ca15f Coalesced loops. Peter Boyle 2019-06-15 08:03:57 +01:00
736358b0cb Coalesced loops Peter Boyle 2019-06-15 08:03:13 +01:00
6b692aa726 Thread loops Peter Boyle 2019-06-15 08:02:26 +01:00
7f99e1cd3b Coalesced loops Peter Boyle 2019-06-15 08:01:39 +01:00
f3c89df948 Thread loop changes Peter Boyle 2019-06-15 08:00:37 +01:00
b7e6d111d7 Thread loop changes. Need to offload this file Peter Boyle 2019-06-15 07:59:10 +01:00
f39cf69c33 Accelerator loop change Peter Boyle 2019-06-15 07:58:23 +01:00
8e27338df2 Rationalise number of loop macros Peter Boyle 2019-06-15 07:57:40 +01:00
bcbb5e9d26 Remove assembly tests Peter Boyle 2019-06-15 07:57:05 +01:00
0ea7f5279d Accelerator loop changes Peter Boyle 2019-06-15 07:56:14 +01:00
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. Peter Boyle 2019-06-15 07:53:58 +01:00
e896d81235 Accelerator loop redefine. Coalesce most accesses, but ET engine still to go clean. Peter Boyle 2019-06-15 07:52:44 +01:00
7b8ccff4f4 Accelerated coalesced loops in most cases Peter Boyle 2019-06-15 07:48:00 +01:00
68541606ab Thread loop changes. Soon try these with accelerator loops and benchmark Peter Boyle 2019-06-15 07:46:42 +01:00
339ea10cc7 First touch only on CPU code Peter Boyle 2019-06-15 07:45:43 +01:00
d0d8dc8042 Thread loop changes Peter Boyle 2019-06-15 07:45:09 +01:00
81eb1fd9f2 Accelerator loop changes for coalesced access Peter Boyle 2019-06-15 07:44:47 +01:00
cb93d32cd9 Thread loop changes Peter Boyle 2019-06-15 07:44:08 +01:00
8f223962ff Thread loop changed Peter Boyle 2019-06-15 07:43:42 +01:00
9a8a63467e BC2 now runs. setup() runs twice, which had resulted in doubling up of momenta. Also fixed initialisation of momentum phases. Michael Marshall 2019-06-12 15:25:59 +01:00
36f06555a2 Simplify Impl Peter Boyle 2019-06-09 22:26:27 +01:00
d6c0e0756d Remove GPU version Peter Boyle 2019-06-09 11:23:42 +01:00
3e41b1055c Remove Gpu only kernels. Peter Boyle 2019-06-09 11:20:01 +01:00
9fbcfe612c Update TODO list Peter Boyle 2019-06-09 11:19:38 +01:00
e78a5e7838 ASM instantiation without link errors Peter Boyle 2019-06-09 01:25:21 +01:00
da8d87e9da Cuda switch off Peter Boyle 2019-06-08 17:11:38 +01:00
8e3a05d89b Moving the instantiation into a cleaner structure Peter Boyle 2019-06-08 13:48:33 +01:00
8adc5da7dd Testig out approaches to kernel writing introducing SIMT_loop temporarily Peter Boyle 2019-06-08 13:47:04 +01:00
29a244e423 Test of using a lane variable instead of repeated reference to threadIdx.y Peter Boyle 2019-06-08 13:46:26 +01:00
18cbfecf02 Use symlinks in find command Peter Boyle 2019-06-08 13:45:46 +01:00
c933ac2248 Temporarily introduce a SIMT_loop to test out approaches prior to making a global change to accelerator_loop Peter Boyle 2019-06-08 13:44:27 +01:00
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 Peter Boyle 2019-06-08 13:43:12 +01:00
86e7fb6e86 Instantiation relocation Peter Boyle 2019-06-08 13:42:46 +01:00
fb91dda7be Hand instantiation moved location Peter Boyle 2019-06-08 13:42:26 +01:00
82cf7bc5ab Move instantiation into fermion/instantiation Peter Boyle 2019-06-08 13:41:46 +01:00
e452cc0a22 Move static variables into instantiation .cc file Peter Boyle 2019-06-08 13:41:20 +01:00
4d2b938166 Remove explict instantiation from here Peter Boyle 2019-06-08 13:41:01 +01:00
10d16ab76c Remove explict instantiation from here Peter Boyle 2019-06-08 13:40:32 +01:00
1f997fa484 Instantiate via explict .cc files for parallel make. Peter Boyle 2019-06-08 13:39:51 +01:00
dc5024e88c The GPU reduction was not working for me and causing errors. Need to revisit. Gianluca is working on deterministic reduction/ Peter Boyle 2019-06-08 13:39:11 +01:00
6d77941990 Drop the 5D vec actions Peter Boyle 2019-06-08 13:38:05 +01:00
0ee6e77cbc Compiles GPU and CPU, still gives good performance on CPU Peter Boyle 2019-06-05 13:28:16 +01:00
18d3cde29a Compile on GPU workd Peter Boyle 2019-06-05 00:14:58 +01:00

... 34 35 36 37 38 ...