diff options
author | Mai Lavelle <mai.lavelle@gmail.com> | 2017-01-14 07:08:43 +0300 |
---|---|---|
committer | Mai Lavelle <mai.lavelle@gmail.com> | 2017-01-14 07:08:43 +0300 |
commit | 3d6037b013ed2cb24ddf2b19c18fe4e121a6a06d (patch) | |
tree | b7e812774bc671e9e5f623c354df6cddeff23bec /intern | |
parent | aa50a5de171ab3cc3d0b2512ddc80f62c9e73da3 (diff) |
Cycles: Add generic work item functions for CUDA
Diffstat (limited to 'intern')
-rw-r--r-- | intern/cycles/kernel/kernel_compat_cuda.h | 44 |
1 files changed, 44 insertions, 0 deletions
diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h index e0c7b17c6a0..ec46d7f29f1 100644 --- a/intern/cycles/kernel/kernel_compat_cuda.h +++ b/intern/cycles/kernel/kernel_compat_cuda.h @@ -51,6 +51,50 @@ #define ccl_restrict __restrict__ #define ccl_align(n) __align__(n) +ccl_device_inline uint ccl_local_id(uint d) +{ + switch(d) { + case 0: return threadIdx.x; + case 1: return threadIdx.y; + case 2: return threadIdx.z; + default: return 0; + } +} + +#define ccl_global_id(d) (ccl_group_id(d) * ccl_local_size(d) + ccl_local_id(d)) + +ccl_device_inline uint ccl_local_size(uint d) +{ + switch(d) { + case 0: return blockDim.x; + case 1: return blockDim.y; + case 2: return blockDim.z; + default: return 0; + } +} + +#define ccl_global_size(d) (ccl_num_groups(d) * ccl_local_size(d)) + +ccl_device_inline uint ccl_group_id(uint d) +{ + switch(d) { + case 0: return blockIdx.x; + case 1: return blockIdx.y; + case 2: return blockIdx.z; + default: return 0; + } +} + +ccl_device_inline uint ccl_num_groups(uint d) +{ + switch(d) { + case 0: return gridDim.x; + case 1: return gridDim.y; + case 2: return gridDim.z; + default: return 0; + } +} + /* No assert supported for CUDA */ #define kernel_assert(cond) |