portelli/Grid
1
0
Fork 0
mirror of https://github.com/paboyle/Grid.git synced 2025-04-26 05:35:55 +01:00
Code Releases Activity

Loop for s and xyzt offlow

Browse Source
This commit is contained in:
paboyle 2018-06-27 21:49:57 +01:00
parent 3e947527cb
commit f8e880b445
1 changed files with 19 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

19
lib/threads/Pragmas.h
View File

@ -106,6 +106,21 @@ void LambdaApply(uint64_t base, uint64_t Num, lambda Lambda)
exit(0); \
}
#define accelerator_loopN( iterator, num, ... ) \
typedef decltype(num) Iterator; \
auto lambda = [=] accelerator (Iterator iterator) mutable { \
__VA_ARGS__; \
}; \
Iterator base = 0; \
Iterator num_block = (num+gpu_threads-1)/gpu_threads; \
LambdaApply<<<num_block,gpu_threads>>>(base,num,lambda); \
cudaDeviceSynchronize(); \
cudaError err = cudaGetLastError(); \
if ( cudaSuccess != err ) { \
printf("Cuda error %s\n",cudaGetErrorString( err )); \
exit(0); \
}
#define cpu_loop( iterator, range, ... ) thread_loop( (auto iterator = range.begin();iterator<range.end();iterator++), { __VA_ARGS__ });
template<typename lambda> __global__
@ -141,6 +156,10 @@ void LambdaApply2D(uint64_t Osites, uint64_t Isites, lambda Lambda)
#define accelerator_inline strong_inline
#define accelerator_loop( iterator, range, ... ) \
thread_loop( (auto iterator = range.begin();iterator<range.end();iterator++), { __VA_ARGS__ });
#define accelerator_loopN( iterator, num, ... ) \
thread_loop( (auto iterator = 0;iterator<num;iterator++), { __VA_ARGS__ });
#define cpu_loop( iterator, range, ... ) \
thread_loop( (auto iterator = range.begin();iterator<range.end();iterator++), { __VA_ARGS__ });