Cycles: Remove ccl_fetch and SOA
[blender.git] / intern / cycles / kernel / geom / geom_triangle_intersect.h
1 /*
2  * Copyright 2014, 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 /* Triangle/Ray intersections.
18  *
19  * For BVH ray intersection we use a precomputed triangle storage to accelerate
20  * intersection at the cost of more memory usage.
21  */
22
23 CCL_NAMESPACE_BEGIN
24
25 /* Workaround stupidness of CUDA/OpenCL which doesn't allow to access indexed
26  * component of float3 value.
27  */
28 #ifndef __KERNEL_CPU__
29 #  define IDX(vec, idx) \
30     ((idx == 0) ? ((vec).x) : ( (idx == 1) ? ((vec).y) : ((vec).z) ))
31 #else
32 #  define IDX(vec, idx) ((vec)[idx])
33 #endif
34
35 /* Ray-Triangle intersection for BVH traversal
36  *
37  * Sven Woop
38  * Watertight Ray/Triangle Intersection
39  *
40  * http://jcgt.org/published/0002/01/05/paper.pdf
41  */
42
43 /* Precalculated data for the ray->tri intersection. */
44 typedef struct IsectPrecalc {
45         /* Maximal dimension kz, and orthogonal dimensions. */
46         int kx, ky, kz;
47
48         /* Shear constants. */
49         float Sx, Sy, Sz;
50 } IsectPrecalc;
51
52 #if (defined(__KERNEL_OPENCL_APPLE__)) || \
53     (defined(__KERNEL_CUDA__) && (defined(i386) || defined(_M_IX86)))
54 ccl_device_noinline
55 #else
56 ccl_device_inline
57 #endif
58 void triangle_intersect_precalc(float3 dir,
59                                 IsectPrecalc *isect_precalc)
60 {
61         /* Calculate dimension where the ray direction is maximal. */
62 #ifndef __KERNEL_SSE__
63         int kz = util_max_axis(make_float3(fabsf(dir.x),
64                                            fabsf(dir.y),
65                                            fabsf(dir.z)));
66         int kx = kz + 1; if(kx == 3) kx = 0;
67         int ky = kx + 1; if(ky == 3) ky = 0;
68 #else
69         int kx, ky, kz;
70         /* Avoiding mispredicted branch on direction. */
71         kz = util_max_axis(fabs(dir));
72         static const char inc_xaxis[] = {1, 2, 0, 55};
73         static const char inc_yaxis[] = {2, 0, 1, 55};
74         kx = inc_xaxis[kz];
75         ky = inc_yaxis[kz];
76 #endif
77
78         float dir_kz = IDX(dir, kz);
79
80         /* Swap kx and ky dimensions to preserve winding direction of triangles. */
81         if(dir_kz < 0.0f) {
82                 int tmp = kx;
83                 kx = ky;
84                 ky = tmp;
85         }
86
87         /* Calculate the shear constants. */
88         float inv_dir_z = 1.0f / dir_kz;
89         isect_precalc->Sx = IDX(dir, kx) * inv_dir_z;
90         isect_precalc->Sy = IDX(dir, ky) * inv_dir_z;
91         isect_precalc->Sz = inv_dir_z;
92
93         /* Store the dimensions. */
94         isect_precalc->kx = kx;
95         isect_precalc->ky = ky;
96         isect_precalc->kz = kz;
97 }
98
99 /* TODO(sergey): Make it general utility function. */
100 ccl_device_inline float xor_signmask(float x, int y)
101 {
102         return __int_as_float(__float_as_int(x) ^ y);
103 }
104
105 ccl_device_inline bool triangle_intersect(KernelGlobals *kg,
106                                           const IsectPrecalc *isect_precalc,
107                                           Intersection *isect,
108                                           float3 P,
109                                           uint visibility,
110                                           int object,
111                                           int prim_addr)
112 {
113         const int kx = isect_precalc->kx;
114         const int ky = isect_precalc->ky;
115         const int kz = isect_precalc->kz;
116         const float Sx = isect_precalc->Sx;
117         const float Sy = isect_precalc->Sy;
118         const float Sz = isect_precalc->Sz;
119
120         /* Calculate vertices relative to ray origin. */
121         const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, prim_addr);
122
123 #if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__)
124         const avxf avxf_P(P.m128, P.m128);
125
126         const avxf tri_ab = kernel_tex_fetch_avxf(__prim_tri_verts, tri_vindex + 0);
127         const avxf tri_bc = kernel_tex_fetch_avxf(__prim_tri_verts, tri_vindex + 1);
128
129         const avxf AB = tri_ab - avxf_P;
130         const avxf BC = tri_bc - avxf_P;
131
132         const __m256i permute_mask = _mm256_set_epi32(0x3, kz, ky, kx, 0x3, kz, ky, kx);
133
134         const avxf AB_k = shuffle(AB, permute_mask);
135         const avxf BC_k = shuffle(BC, permute_mask);
136
137         /* Akz, Akz, Bkz, Bkz, Bkz, Bkz, Ckz, Ckz */
138         const avxf ABBC_kz = shuffle<2>(AB_k, BC_k);
139
140         /* Akx, Aky, Bkx, Bky, Bkx,Bky, Ckx, Cky */
141         const avxf ABBC_kxy = shuffle<0,1,0,1>(AB_k, BC_k);
142
143         const avxf Sxy(Sy, Sx, Sy, Sx);
144
145         /* Ax, Ay, Bx, By, Bx, By, Cx, Cy */
146         const avxf ABBC_xy = nmadd(ABBC_kz, Sxy, ABBC_kxy);
147
148         float ABBC_kz_array[8];
149         _mm256_storeu_ps((float*)&ABBC_kz_array, ABBC_kz);
150
151         const float A_kz = ABBC_kz_array[0];
152         const float B_kz = ABBC_kz_array[2];
153         const float C_kz = ABBC_kz_array[6];
154
155         /* By, Bx, Cy, Cx, By, Bx, Ay, Ax */
156         const avxf BCBA_yx = permute<3,2,7,6,3,2,1,0>(ABBC_xy);
157
158         const avxf neg_mask(0,0,0,0,0x80000000, 0x80000000, 0x80000000, 0x80000000);
159
160         /* W           U                             V
161          * (AxBy-AyBx) (BxCy-ByCx) XX XX (BxBy-ByBx) (CxAy-CyAx) XX XX
162          */
163         const avxf WUxxxxVxx_neg = _mm256_hsub_ps(ABBC_xy * BCBA_yx, neg_mask /* Dont care */);
164
165         const avxf WUVWnegWUVW = permute<0,1,5,0,0,1,5,0>(WUxxxxVxx_neg) ^ neg_mask;
166
167         /* Calculate scaled barycentric coordinates. */
168         float WUVW_array[4];
169         _mm_storeu_ps((float*)&WUVW_array, _mm256_castps256_ps128 (WUVWnegWUVW));
170
171         const float W = WUVW_array[0];
172         const float U = WUVW_array[1];
173         const float V = WUVW_array[2];
174
175         const int WUVW_mask = 0x7 & _mm256_movemask_ps(WUVWnegWUVW);
176         const int WUVW_zero = 0x7 & _mm256_movemask_ps(_mm256_cmp_ps(WUVWnegWUVW,
177                                                        _mm256_setzero_ps(), 0));
178
179         if(!((WUVW_mask == 7) || (WUVW_mask == 0)) && ((WUVW_mask | WUVW_zero) != 7)) {
180                 return false;
181         }
182 #else
183         const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex+0),
184                      tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex+1),
185                      tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex+2);
186         const float3 A = make_float3(tri_a.x - P.x, tri_a.y - P.y, tri_a.z - P.z);
187         const float3 B = make_float3(tri_b.x - P.x, tri_b.y - P.y, tri_b.z - P.z);
188         const float3 C = make_float3(tri_c.x - P.x, tri_c.y - P.y, tri_c.z - P.z);
189
190         const float A_kx = IDX(A, kx), A_ky = IDX(A, ky), A_kz = IDX(A, kz);
191         const float B_kx = IDX(B, kx), B_ky = IDX(B, ky), B_kz = IDX(B, kz);
192         const float C_kx = IDX(C, kx), C_ky = IDX(C, ky), C_kz = IDX(C, kz);
193
194         /* Perform shear and scale of vertices. */
195         const float Ax = A_kx - Sx * A_kz;
196         const float Ay = A_ky - Sy * A_kz;
197         const float Bx = B_kx - Sx * B_kz;
198         const float By = B_ky - Sy * B_kz;
199         const float Cx = C_kx - Sx * C_kz;
200         const float Cy = C_ky - Sy * C_kz;
201
202         /* Calculate scaled barycentric coordinates. */
203         float U = Cx * By - Cy * Bx;
204         float V = Ax * Cy - Ay * Cx;
205         float W = Bx * Ay - By * Ax;
206         if((U < 0.0f || V < 0.0f || W < 0.0f) &&
207            (U > 0.0f || V > 0.0f || W > 0.0f))
208         {
209                 return false;
210         }
211 #endif
212
213         /* Calculate determinant. */
214         float det = U + V + W;
215         if(UNLIKELY(det == 0.0f)) {
216                 return false;
217         }
218
219         /* Calculate scaled z-coordinates of vertices and use them to calculate
220          * the hit distance.
221          */
222         const float T = (U * A_kz + V * B_kz + W * C_kz) * Sz;
223         const int sign_det = (__float_as_int(det) & 0x80000000);
224         const float sign_T = xor_signmask(T, sign_det);
225         if((sign_T < 0.0f) ||
226            (sign_T > isect->t * xor_signmask(det, sign_det)))
227         {
228                 return false;
229         }
230
231 #ifdef __VISIBILITY_FLAG__
232         /* visibility flag test. we do it here under the assumption
233          * that most triangles are culled by node flags */
234         if(kernel_tex_fetch(__prim_visibility, prim_addr) & visibility)
235 #endif
236         {
237 #ifdef __KERNEL_CUDA__
238                 if(A == B && B == C) {
239                         return false;
240                 }
241 #endif
242                 /* Normalize U, V, W, and T. */
243                 const float inv_det = 1.0f / det;
244                 isect->prim = prim_addr;
245                 isect->object = object;
246                 isect->type = PRIMITIVE_TRIANGLE;
247                 isect->u = U * inv_det;
248                 isect->v = V * inv_det;
249                 isect->t = T * inv_det;
250                 return true;
251         }
252         return false;
253 }
254
255 /* Special ray intersection routines for subsurface scattering. In that case we
256  * only want to intersect with primitives in the same object, and if case of
257  * multiple hits we pick a single random primitive as the intersection point.
258  */
259
260 #ifdef __SUBSURFACE__
261 ccl_device_inline void triangle_intersect_subsurface(
262         KernelGlobals *kg,
263         const IsectPrecalc *isect_precalc,
264         SubsurfaceIntersection *ss_isect,
265         float3 P,
266         int object,
267         int prim_addr,
268         float tmax,
269         uint *lcg_state,
270         int max_hits)
271 {
272         const int kx = isect_precalc->kx;
273         const int ky = isect_precalc->ky;
274         const int kz = isect_precalc->kz;
275         const float Sx = isect_precalc->Sx;
276         const float Sy = isect_precalc->Sy;
277         const float Sz = isect_precalc->Sz;
278
279         /* Calculate vertices relative to ray origin. */
280         const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, prim_addr);
281         const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex+0),
282                      tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex+1),
283                      tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex+2);
284
285 #if defined(__KERNEL_AVX2__) && defined(__KERNEL_SSE__)
286         const avxf avxf_P(P.m128, P.m128);
287
288         const avxf tri_ab = kernel_tex_fetch_avxf(__prim_tri_verts, tri_vindex + 0);
289         const avxf tri_bc = kernel_tex_fetch_avxf(__prim_tri_verts, tri_vindex + 1);
290
291         const avxf AB = tri_ab - avxf_P;
292         const avxf BC = tri_bc - avxf_P;
293
294         const __m256i permuteMask = _mm256_set_epi32(0x3, kz, ky, kx, 0x3, kz, ky, kx);
295
296         const avxf AB_k = shuffle(AB, permuteMask);
297         const avxf BC_k = shuffle(BC, permuteMask);
298
299         /* Akz, Akz, Bkz, Bkz, Bkz, Bkz, Ckz, Ckz */
300         const avxf ABBC_kz = shuffle<2>(AB_k, BC_k);
301
302         /* Akx, Aky, Bkx, Bky, Bkx,Bky, Ckx, Cky */
303         const avxf ABBC_kxy = shuffle<0,1,0,1>(AB_k, BC_k);
304
305         const avxf Sxy(Sy, Sx, Sy, Sx);
306
307         /* Ax, Ay, Bx, By, Bx, By, Cx, Cy */
308         const avxf ABBC_xy = nmadd(ABBC_kz, Sxy, ABBC_kxy);
309
310         float ABBC_kz_array[8];
311         _mm256_storeu_ps((float*)&ABBC_kz_array, ABBC_kz);
312
313         const float A_kz = ABBC_kz_array[0];
314         const float B_kz = ABBC_kz_array[2];
315         const float C_kz = ABBC_kz_array[6];
316
317         /* By, Bx, Cy, Cx, By, Bx, Ay, Ax */
318         const avxf BCBA_yx = permute<3,2,7,6,3,2,1,0>(ABBC_xy);
319
320         const avxf negMask(0,0,0,0,0x80000000, 0x80000000, 0x80000000, 0x80000000);
321
322         /* W           U                             V
323          * (AxBy-AyBx) (BxCy-ByCx) XX XX (BxBy-ByBx) (CxAy-CyAx) XX XX
324          */
325         const avxf WUxxxxVxx_neg = _mm256_hsub_ps(ABBC_xy * BCBA_yx, negMask /* Dont care */);
326
327         const avxf WUVWnegWUVW = permute<0,1,5,0,0,1,5,0>(WUxxxxVxx_neg) ^ negMask;
328
329         /* Calculate scaled barycentric coordinates. */
330         float WUVW_array[4];
331         _mm_storeu_ps((float*)&WUVW_array, _mm256_castps256_ps128 (WUVWnegWUVW));
332
333         const float W = WUVW_array[0];
334         const float U = WUVW_array[1];
335         const float V = WUVW_array[2];
336
337         const int WUVW_mask = 0x7 & _mm256_movemask_ps(WUVWnegWUVW);
338         const int WUVW_zero = 0x7 & _mm256_movemask_ps(_mm256_cmp_ps(WUVWnegWUVW,
339                                                        _mm256_setzero_ps(), 0));
340
341         if(!((WUVW_mask == 7) || (WUVW_mask == 0)) && ((WUVW_mask | WUVW_zero) != 7)) {
342                 return;
343         }
344 #else
345         const float3 A = make_float3(tri_a.x - P.x, tri_a.y - P.y, tri_a.z - P.z);
346         const float3 B = make_float3(tri_b.x - P.x, tri_b.y - P.y, tri_b.z - P.z);
347         const float3 C = make_float3(tri_c.x - P.x, tri_c.y - P.y, tri_c.z - P.z);
348
349         const float A_kx = IDX(A, kx), A_ky = IDX(A, ky), A_kz = IDX(A, kz);
350         const float B_kx = IDX(B, kx), B_ky = IDX(B, ky), B_kz = IDX(B, kz);
351         const float C_kx = IDX(C, kx), C_ky = IDX(C, ky), C_kz = IDX(C, kz);
352
353         /* Perform shear and scale of vertices. */
354         const float Ax = A_kx - Sx * A_kz;
355         const float Ay = A_ky - Sy * A_kz;
356         const float Bx = B_kx - Sx * B_kz;
357         const float By = B_ky - Sy * B_kz;
358         const float Cx = C_kx - Sx * C_kz;
359         const float Cy = C_ky - Sy * C_kz;
360
361         /* Calculate scaled barycentric coordinates. */
362         float U = Cx * By - Cy * Bx;
363         float V = Ax * Cy - Ay * Cx;
364         float W = Bx * Ay - By * Ax;
365
366         if((U < 0.0f || V < 0.0f || W < 0.0f) &&
367            (U > 0.0f || V > 0.0f || W > 0.0f))
368         {
369                 return;
370         }
371 #endif
372
373         /* Calculate determinant. */
374         float det = U + V + W;
375         if(UNLIKELY(det == 0.0f)) {
376                 return;
377         }
378
379         /* Calculate scaled z−coordinates of vertices and use them to calculate
380          * the hit distance.
381          */
382         const int sign_det = (__float_as_int(det) & 0x80000000);
383         const float T = (U * A_kz + V * B_kz + W * C_kz) * Sz;
384         const float sign_T = xor_signmask(T, sign_det);
385         if((sign_T < 0.0f) ||
386            (sign_T > tmax * xor_signmask(det, sign_det)))
387         {
388                 return;
389         }
390
391         /* Normalize U, V, W, and T. */
392         const float inv_det = 1.0f / det;
393
394         const float t = T * inv_det;
395         for(int i = min(max_hits, ss_isect->num_hits) - 1; i >= 0; --i) {
396                 if(ss_isect->hits[i].t == t) {
397                         return;
398                 }
399         }
400
401         ss_isect->num_hits++;
402         int hit;
403
404         if(ss_isect->num_hits <= max_hits) {
405                 hit = ss_isect->num_hits - 1;
406         }
407         else {
408                 /* reservoir sampling: if we are at the maximum number of
409                  * hits, randomly replace element or skip it */
410                 hit = lcg_step_uint(lcg_state) % ss_isect->num_hits;
411
412                 if(hit >= max_hits)
413                         return;
414         }
415
416         /* record intersection */
417         Intersection *isect = &ss_isect->hits[hit];
418         isect->prim = prim_addr;
419         isect->object = object;
420         isect->type = PRIMITIVE_TRIANGLE;
421         isect->u = U * inv_det;
422         isect->v = V * inv_det;
423         isect->t = t;
424
425         /* Record geometric normal. */
426         /* TODO(sergey): Use float4_to_float3() on just an edges. */
427         const float3 v0 = float4_to_float3(tri_a);
428         const float3 v1 = float4_to_float3(tri_b);
429         const float3 v2 = float4_to_float3(tri_c);
430         ss_isect->Ng[hit] = normalize(cross(v1 - v0, v2 - v0));
431 }
432 #endif
433
434 /* Refine triangle intersection to more precise hit point. For rays that travel
435  * far the precision is often not so good, this reintersects the primitive from
436  * a closer distance. */
437
438 /* Reintersections uses the paper:
439  *
440  * Tomas Moeller
441  * Fast, minimum storage ray/triangle intersection
442  * http://www.cs.virginia.edu/~gfx/Courses/2003/ImageSynthesis/papers/Acceleration/Fast%20MinimumStorage%20RayTriangle%20Intersection.pdf
443  */
444
445 ccl_device_inline float3 triangle_refine(KernelGlobals *kg,
446                                          ShaderData *sd,
447                                          const Intersection *isect,
448                                          const Ray *ray)
449 {
450         float3 P = ray->P;
451         float3 D = ray->D;
452         float t = isect->t;
453
454 #ifdef __INTERSECTION_REFINE__
455         if(isect->object != OBJECT_NONE) {
456                 if(UNLIKELY(t == 0.0f)) {
457                         return P;
458                 }
459 #  ifdef __OBJECT_MOTION__
460                 Transform tfm = sd->ob_itfm;
461 #  else
462                 Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_INVERSE_TRANSFORM);
463 #  endif
464
465                 P = transform_point(&tfm, P);
466                 D = transform_direction(&tfm, D*t);
467                 D = normalize_len(D, &t);
468         }
469
470         P = P + D*t;
471
472         const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect->prim);
473         const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex+0),
474                      tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex+1),
475                      tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex+2);
476         float3 edge1 = make_float3(tri_a.x - tri_c.x, tri_a.y - tri_c.y, tri_a.z - tri_c.z);
477         float3 edge2 = make_float3(tri_b.x - tri_c.x, tri_b.y - tri_c.y, tri_b.z - tri_c.z);
478         float3 tvec = make_float3(P.x - tri_c.x, P.y - tri_c.y, P.z - tri_c.z);
479         float3 qvec = cross(tvec, edge1);
480         float3 pvec = cross(D, edge2);
481         float det = dot(edge1, pvec);
482         if(det != 0.0f) {
483                 /* If determinant is zero it means ray lies in the plane of
484                  * the triangle. It is possible in theory due to watertight
485                  * nature of triangle intersection. For such cases we simply
486                  * don't refine intersection hoping it'll go all fine.
487                  */
488                 float rt = dot(edge2, qvec) / det;
489                 P = P + D*rt;
490         }
491
492         if(isect->object != OBJECT_NONE) {
493 #  ifdef __OBJECT_MOTION__
494                 Transform tfm = sd->ob_tfm;
495 #  else
496                 Transform tfm = object_fetch_transform(kg, isect->object, OBJECT_TRANSFORM);
497 #  endif
498
499                 P = transform_point(&tfm, P);
500         }
501
502         return P;
503 #else
504         return P + D*t;
505 #endif
506 }
507
508 /* Same as above, except that isect->t is assumed to be in object space for
509  * instancing.
510  */
511 ccl_device_inline float3 triangle_refine_subsurface(KernelGlobals *kg,
512                                                     ShaderData *sd,
513                                                     const Intersection *isect,
514                                                     const Ray *ray)
515 {
516         float3 P = ray->P;
517         float3 D = ray->D;
518         float t = isect->t;
519
520         if(isect->object != OBJECT_NONE) {
521 #ifdef __OBJECT_MOTION__
522                 Transform tfm = sd->ob_itfm;
523 #else
524                 Transform tfm = object_fetch_transform(kg,
525                                                        isect->object,
526                                                        OBJECT_INVERSE_TRANSFORM);
527 #endif
528
529                 P = transform_point(&tfm, P);
530                 D = transform_direction(&tfm, D);
531                 D = normalize(D);
532         }
533
534         P = P + D*t;
535
536 #ifdef __INTERSECTION_REFINE__
537         const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect->prim);
538         const float4 tri_a = kernel_tex_fetch(__prim_tri_verts, tri_vindex+0),
539                      tri_b = kernel_tex_fetch(__prim_tri_verts, tri_vindex+1),
540                      tri_c = kernel_tex_fetch(__prim_tri_verts, tri_vindex+2);
541         float3 edge1 = make_float3(tri_a.x - tri_c.x, tri_a.y - tri_c.y, tri_a.z - tri_c.z);
542         float3 edge2 = make_float3(tri_b.x - tri_c.x, tri_b.y - tri_c.y, tri_b.z - tri_c.z);
543         float3 tvec = make_float3(P.x - tri_c.x, P.y - tri_c.y, P.z - tri_c.z);
544         float3 qvec = cross(tvec, edge1);
545         float3 pvec = cross(D, edge2);
546         float det = dot(edge1, pvec);
547         if(det != 0.0f) {
548                 /* If determinant is zero it means ray lies in the plane of
549                  * the triangle. It is possible in theory due to watertight
550                  * nature of triangle intersection. For such cases we simply
551                  * don't refine intersection hoping it'll go all fine.
552                  */
553                 float rt = dot(edge2, qvec) / det;
554                 P = P + D*rt;
555         }
556 #endif  /* __INTERSECTION_REFINE__ */
557
558         if(isect->object != OBJECT_NONE) {
559 #ifdef __OBJECT_MOTION__
560                 Transform tfm = sd->ob_tfm;
561 #else
562                 Transform tfm = object_fetch_transform(kg,
563                                                        isect->object,
564                                                        OBJECT_TRANSFORM);
565 #endif
566
567                 P = transform_point(&tfm, P);
568         }
569
570         return P;
571 }
572
573 #undef IDX
574
575 CCL_NAMESPACE_END