Cycles: Expose image image extension mapping to the image manager
[blender.git] / intern / cycles / kernel / kernel_queues.h
1 /*
2  * Copyright 2011-2015 Blender Foundation
3  *
4  * Licensed under the Apache License, Version 2.0 (the "License");
5  * you may not use this file except in compliance with the License.
6  * You may obtain a copy of the License at
7  *
8  * http://www.apache.org/licenses/LICENSE-2.0
9  *
10  * Unless required by applicable law or agreed to in writing, software
11  * distributed under the License is distributed on an "AS IS" BASIS,
12  * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13  * See the License for the specific language governing permissions and
14  * limitations under the License.
15  */
16
17 #ifndef __KERNEL_QUEUE_H__
18 #define __KERNEL_QUEUE_H__
19
20 /*
21  * Queue utility functions for split kernel
22  */
23
24 #pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
25 #pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable
26
27 /*
28  * Enqueue ray index into the queue
29  */
30 ccl_device void enqueue_ray_index(
31         int ray_index,                /* Ray index to be enqueued. */
32         int queue_number,             /* Queue in which the ray index should be enqueued. */
33         ccl_global int *queues,       /* Buffer of all queues. */
34         int queue_size,               /* Size of each queue. */
35         ccl_global int *queue_index) /* Array of size num_queues; Used for atomic increment. */
36 {
37         /* This thread's queue index. */
38         int my_queue_index = atomic_inc(&queue_index[queue_number]) + (queue_number * queue_size);
39         queues[my_queue_index] = ray_index;
40 }
41
42 /*
43  * Get the ray index for this thread
44  * Returns a positive ray_index for threads that have to do some work;
45  * Returns 'QUEUE_EMPTY_SLOT' for threads that don't have any work
46  * i.e All ray's in the queue has been successfully allocated and there
47  * is no more ray to allocate to other threads.
48  */
49 ccl_device int get_ray_index(
50         int thread_index,       /* Global thread index. */
51         int queue_number,       /* Queue to operate on. */
52         ccl_global int *queues, /* Buffer of all queues. */
53         int queuesize,          /* Size of a queue. */
54         int empty_queue)        /* Empty the queue slot as soon as we fetch the ray index. */
55 {
56         int ray_index = queues[queue_number * queuesize + thread_index];
57         if(empty_queue && ray_index != QUEUE_EMPTY_SLOT) {
58                 queues[queue_number * queuesize + thread_index] = QUEUE_EMPTY_SLOT;
59         }
60         return ray_index;
61 }
62
63 /* The following functions are to realize Local memory variant of enqueue ray index function. */
64
65 /* All threads should call this function. */
66 ccl_device void enqueue_ray_index_local(
67         int ray_index,                               /* Ray index to enqueue. */
68         int queue_number,                            /* Queue in which to enqueue ray index. */
69         char enqueue_flag,                           /* True for threads whose ray index has to be enqueued. */
70         int queuesize,                               /* queue size. */
71         ccl_local unsigned int *local_queue_atomics,   /* To to local queue atomics. */
72         ccl_global int *Queue_data,                  /* Queues. */
73         ccl_global int *Queue_index)                 /* To do global queue atomics. */
74 {
75         int lidx = get_local_id(1) * get_local_size(0) + get_local_id(0);
76
77         /* Get local queue id .*/
78         unsigned int lqidx;
79         if(enqueue_flag) {
80                 lqidx = atomic_inc(local_queue_atomics);
81         }
82         barrier(CLK_LOCAL_MEM_FENCE);
83
84         /* Get global queue offset. */
85         if(lidx == 0) {
86                 *local_queue_atomics = atomic_add(&Queue_index[queue_number], *local_queue_atomics);
87         }
88         barrier(CLK_LOCAL_MEM_FENCE);
89
90         /* Get global queue index and enqueue ray. */
91         if(enqueue_flag) {
92                 unsigned int my_gqidx = queue_number * queuesize + (*local_queue_atomics) + lqidx;
93                 Queue_data[my_gqidx] = ray_index;
94         }
95 }
96
97 ccl_device unsigned int get_local_queue_index(
98         int queue_number, /* Queue in which to enqueue the ray; -1 if no queue */
99         ccl_local unsigned int *local_queue_atomics)
100 {
101         int my_lqidx = atomic_inc(&local_queue_atomics[queue_number]);
102         return my_lqidx;
103 }
104
105 ccl_device unsigned int get_global_per_queue_offset(
106         int queue_number,
107         ccl_local unsigned int *local_queue_atomics,
108         ccl_global int* global_queue_atomics)
109 {
110         unsigned int queue_offset = atomic_add(&global_queue_atomics[queue_number],
111                                                local_queue_atomics[queue_number]);
112         return queue_offset;
113 }
114
115 ccl_device unsigned int get_global_queue_index(
116     int queue_number,
117     int queuesize,
118     unsigned int lqidx,
119     ccl_local unsigned int * global_per_queue_offset)
120 {
121         int my_gqidx = queuesize * queue_number + lqidx + global_per_queue_offset[queue_number];
122         return my_gqidx;
123 }
124
125 #endif // __KERNEL_QUEUE_H__