portelli/Grid
1
0
Fork 0
mirror of https://github.com/paboyle/Grid.git synced 2025-04-04 19:25:56 +01:00
Code Releases Activity

Proper offload (accelerator access) will require a mutable copy lambda.

Browse Source
This commit is contained in:
paboyle 2018-02-02 11:38:19 +00:00
parent 51eb2c5dfc
commit b9b5bdfc3a
1 changed files with 8 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

10
lib/threads/Pragmas.h
View File

@ -76,13 +76,19 @@ Author: paboyle <paboyle@ph.ed.ac.uk>
// Accelerator primitives; fall back to threading // Accelerator primitives; fall back to threading
////////////////////////////////////////////////////////////////////////////////// //////////////////////////////////////////////////////////////////////////////////
#ifdef GRID_NVCC #ifdef GRID_NVCC
#define accelerator_exec( ... ) \
auto lambda = [=] accelerator (void) { \
__VA_ARGS__; \
}; \
lambda();
#define accelerator __host__ __device__ #define accelerator __host__ __device__
#define accelerator_inline __host__ __device__ inline #define accelerator_inline __host__ __device__ inline
// FIXME ; need to make this a CUDA kernel call // FIXME ; need to make this a CUDA kernel call
#define accelerator_loop( iterator, range, ... ) \ #define accelerator_loop( iterator, range, ... ) \
typedef decltype(range.begin()) Iterator; \ typedef decltype(range.begin()) Iterator; \
auto lambda = [&] (Iterator iterator) { \ auto lambda = [=] accelerator (Iterator iterator) mutable { \
__VA_ARGS__; \ __VA_ARGS__; \
}; \ }; \
for(auto it=range.begin();it<range.end();it++){ \ for(auto it=range.begin();it<range.end();it++){ \
lambda(it); \ lambda(it); \