Lines Matching refs:tid
50 inline int calc_lut(__local int* smem, int val, int tid)
52 smem[tid] = val;
55 if (tid == 0)
60 return smem[tid];
64 inline void reduce(volatile __local int* smem, int val, int tid)
66 smem[tid] = val;
69 if (tid < 128)
70 smem[tid] = val += smem[tid + 128];
73 if (tid < 64)
74 smem[tid] = val += smem[tid + 64];
77 if (tid < 32)
78 smem[tid] += smem[tid + 32];
81 if (tid < 16)
82 smem[tid] += smem[tid + 16];
85 if (tid < 8)
86 smem[tid] += smem[tid + 8];
89 if (tid < 4)
90 smem[tid] += smem[tid + 4];
93 if (tid < 2)
94 smem[tid] += smem[tid + 2];
97 if (tid < 1)
98 smem[256] = smem[tid] + smem[tid + 1];
104 inline void reduce(__local volatile int* smem, int val, int tid)
106 smem[tid] = val;
109 if (tid < 128)
110 smem[tid] = val += smem[tid + 128];
113 if (tid < 64)
114 smem[tid] = val += smem[tid + 64];
117 if (tid < 32)
119 smem[tid] += smem[tid + 32];
123 if (tid < 16)
126 smem[tid] += smem[tid + 16];
131 if (tid < 8)
134 smem[tid] += smem[tid + 8];
135 smem[tid] += smem[tid + 4];
136 smem[tid] += smem[tid + 2];
137 smem[tid] += smem[tid + 1];
152 int tid = get_local_id(1) * get_local_size(0)
154 smem[tid] = 0;
168 int tHistVal = smem[tid];
182 reduce(smem, clipped, tid);
194 if (tid == 0)
204 if (tid < residual)
208 const int lutVal = calc_lut(smem, tHistVal, tid);
210 lut[(ty * tilesX + tx) * dstStep + tid + dst_offset] =