[Bf-blender-cvs] [3d6037b013] cycles-tiles-rework: Cycles: Add generic work item functions for CUDA

[Mai Lavelle]

Commit: 3d6037b013ed2cb24ddf2b19c18fe4e121a6a06d
Author: Mai Lavelle
Date:   Fri Jan 13 23:08:43 2017 -0500
Branches: cycles-tiles-rework
https://developer.blender.org/rB3d6037b013ed2cb24ddf2b19c18fe4e121a6a06d

Cycles: Add generic work item functions for CUDA

===================================================================

M	intern/cycles/kernel/kernel_compat_cuda.h

===================================================================

diff --git a/intern/cycles/kernel/kernel_compat_cuda.h b/intern/cycles/kernel/kernel_compat_cuda.h
index e0c7b17c6a..ec46d7f29f 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)