intern/cycles/kernel/kernel_compat_cpu.h

   1 /*
   2  * Copyright 2011-2013 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_COMPAT_CPU_H__
  18 #define __KERNEL_COMPAT_CPU_H__
  19
  20 #define __KERNEL_CPU__
  21
  22 /* Release kernel has too much false-positive maybe-uninitialized warnings,
  23  * which makes it possible to miss actual warnings.
  24  */
  25 #if (defined(__GNUC__) && !defined(__clang__)) && defined(NDEBUG)
  26 #  pragma GCC diagnostic ignored "-Wmaybe-uninitialized"
  27 #  pragma GCC diagnostic ignored "-Wuninitialized"
  28 #endif
  29
  30 /* Selective nodes compilation. */
  31 #ifndef __NODES_MAX_GROUP__
  32 #  define __NODES_MAX_GROUP__ NODE_GROUP_LEVEL_MAX
  33 #endif
  34 #ifndef __NODES_FEATURES__
  35 #  define __NODES_FEATURES__ NODE_FEATURE_ALL
  36 #endif
  37
  38 #include "util/util_debug.h"
  39 #include "util/util_math.h"
  40 #include "util/util_simd.h"
  41 #include "util/util_half.h"
  42 #include "util/util_types.h"
  43 #include "util/util_texture.h"
  44
  45 #define ccl_addr_space
  46
  47 #define ccl_local_id(d) 0
  48 #define ccl_global_id(d) (kg->global_id[d])
  49
  50 #define ccl_local_size(d) 1
  51 #define ccl_global_size(d) (kg->global_size[d])
  52
  53 #define ccl_group_id(d) ccl_global_id(d)
  54 #define ccl_num_groups(d) ccl_global_size(d)
  55
  56 /* On x86_64, versions of glibc < 2.16 have an issue where expf is
  57  * much slower than the double version.  This was fixed in glibc 2.16.
  58  */
  59 #if !defined(__KERNEL_GPU__)  && defined(__x86_64__) && defined(__x86_64__) && \
  60      defined(__GNU_LIBRARY__) && defined(__GLIBC__ ) && defined(__GLIBC_MINOR__) && \
  61      (__GLIBC__ <= 2 && __GLIBC_MINOR__ < 16)
  62 #  define expf(x) ((float)exp((double)(x)))
  63 #endif
  64
  65 CCL_NAMESPACE_BEGIN
  66
  67 /* Assertions inside the kernel only work for the CPU device, so we wrap it in
  68  * a macro which is empty for other devices */
  69
  70 #define kernel_assert(cond) assert(cond)
  71
  72 /* Texture types to be compatible with CUDA textures. These are really just
  73  * simple arrays and after inlining fetch hopefully revert to being a simple
  74  * pointer lookup. */
  75
  76 template<typename T> struct texture  {
  77         ccl_always_inline T fetch(int index)
  78         {
  79                 kernel_assert(index >= 0 && index < width);
  80                 return data[index];
  81         }
  82
  83 #ifdef __KERNEL_AVX__
  84         /* Reads 256 bytes but indexes in blocks of 128 bytes to maintain
  85          * compatibility with existing indicies and data structures.
  86          */
  87         ccl_always_inline avxf fetch_avxf(const int index)
  88         {
  89                 kernel_assert(index >= 0 && (index+1) < width);
  90                 ssef *ssef_data = (ssef*)data;
  91                 ssef *ssef_node_data = &ssef_data[index];
  92                 return _mm256_loadu_ps((float *)ssef_node_data);
  93         }
  94
  95 #endif
  96
  97 #ifdef __KERNEL_SSE2__
  98         ccl_always_inline ssef fetch_ssef(int index)
  99         {
 100                 kernel_assert(index >= 0 && index < width);
 101                 return ((ssef*)data)[index];
 102         }
 103
 104         ccl_always_inline ssei fetch_ssei(int index)
 105         {
 106                 kernel_assert(index >= 0 && index < width);
 107                 return ((ssei*)data)[index];
 108         }
 109 #endif
 110
 111         T *data;
 112         int width;
 113 };
 114
 115 template<typename T> struct texture_image  {
 116 #define SET_CUBIC_SPLINE_WEIGHTS(u, t) \
 117         { \
 118                 u[0] = (((-1.0f/6.0f)* t + 0.5f) * t - 0.5f) * t + (1.0f/6.0f); \
 119                 u[1] =  ((      0.5f * t - 1.0f) * t       ) * t + (2.0f/3.0f); \
 120                 u[2] =  ((     -0.5f * t + 0.5f) * t + 0.5f) * t + (1.0f/6.0f); \
 121                 u[3] = (1.0f / 6.0f) * t * t * t; \
 122         } (void)0
 123
 124         ccl_always_inline float4 read(float4 r)
 125         {
 126                 return r;
 127         }
 128
 129         ccl_always_inline float4 read(uchar4 r)
 130         {
 131                 float f = 1.0f/255.0f;
 132                 return make_float4(r.x*f, r.y*f, r.z*f, r.w*f);
 133         }
 134
 135         ccl_always_inline float4 read(uchar r)
 136         {
 137                 float f = r*(1.0f/255.0f);
 138                 return make_float4(f, f, f, 1.0f);
 139         }
 140
 141         ccl_always_inline float4 read(float r)
 142         {
 143                 /* TODO(dingto): Optimize this, so interpolation
 144                  * happens on float instead of float4 */
 145                 return make_float4(r, r, r, 1.0f);
 146         }
 147
 148         ccl_always_inline float4 read(half4 r)
 149         {
 150                 return half4_to_float4(r);
 151         }
 152
 153         ccl_always_inline float4 read(half r)
 154         {
 155                 float f = half_to_float(r);
 156                 return make_float4(f, f, f, 1.0f);
 157         }
 158
 159         ccl_always_inline int wrap_periodic(int x, int width)
 160         {
 161                 x %= width;
 162                 if(x < 0)
 163                         x += width;
 164                 return x;
 165         }
 166
 167         ccl_always_inline int wrap_clamp(int x, int width)
 168         {
 169                 return clamp(x, 0, width-1);
 170         }
 171
 172         ccl_always_inline float frac(float x, int *ix)
 173         {
 174                 int i = float_to_int(x) - ((x < 0.0f)? 1: 0);
 175                 *ix = i;
 176                 return x - (float)i;
 177         }
 178
 179         ccl_always_inline float4 interp(float x, float y)
 180         {
 181                 if(UNLIKELY(!data))
 182                         return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 183
 184                 int ix, iy, nix, niy;
 185
 186                 if(interpolation == INTERPOLATION_CLOSEST) {
 187                         frac(x*(float)width, &ix);
 188                         frac(y*(float)height, &iy);
 189                         switch(extension) {
 190                                 case EXTENSION_REPEAT:
 191                                         ix = wrap_periodic(ix, width);
 192                                         iy = wrap_periodic(iy, height);
 193                                         break;
 194                                 case EXTENSION_CLIP:
 195                                         if(x < 0.0f || y < 0.0f || x > 1.0f || y > 1.0f) {
 196                                                 return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 197                                         }
 198                                         /* Fall through. */
 199                                 case EXTENSION_EXTEND:
 200                                         ix = wrap_clamp(ix, width);
 201                                         iy = wrap_clamp(iy, height);
 202                                         break;
 203                                 default:
 204                                         kernel_assert(0);
 205                                         return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 206                         }
 207                         return read(data[ix + iy*width]);
 208                 }
 209                 else if(interpolation == INTERPOLATION_LINEAR) {
 210                         float tx = frac(x*(float)width - 0.5f, &ix);
 211                         float ty = frac(y*(float)height - 0.5f, &iy);
 212
 213                         switch(extension) {
 214                                 case EXTENSION_REPEAT:
 215                                         ix = wrap_periodic(ix, width);
 216                                         iy = wrap_periodic(iy, height);
 217
 218                                         nix = wrap_periodic(ix+1, width);
 219                                         niy = wrap_periodic(iy+1, height);
 220                                         break;
 221                                 case EXTENSION_CLIP:
 222                                         if(x < 0.0f || y < 0.0f || x > 1.0f || y > 1.0f) {
 223                                                 return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 224                                         }
 225                                         /* Fall through. */
 226                                 case EXTENSION_EXTEND:
 227                                         nix = wrap_clamp(ix+1, width);
 228                                         niy = wrap_clamp(iy+1, height);
 229
 230                                         ix = wrap_clamp(ix, width);
 231                                         iy = wrap_clamp(iy, height);
 232                                         break;
 233                                 default:
 234                                         kernel_assert(0);
 235                                         return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 236                         }
 237
 238                         float4 r = (1.0f - ty)*(1.0f - tx)*read(data[ix + iy*width]);
 239                         r += (1.0f - ty)*tx*read(data[nix + iy*width]);
 240                         r += ty*(1.0f - tx)*read(data[ix + niy*width]);
 241                         r += ty*tx*read(data[nix + niy*width]);
 242
 243                         return r;
 244                 }
 245                 else {
 246                         /* Bicubic b-spline interpolation. */
 247                         float tx = frac(x*(float)width - 0.5f, &ix);
 248                         float ty = frac(y*(float)height - 0.5f, &iy);
 249                         int pix, piy, nnix, nniy;
 250                         switch(extension) {
 251                                 case EXTENSION_REPEAT:
 252                                         ix = wrap_periodic(ix, width);
 253                                         iy = wrap_periodic(iy, height);
 254
 255                                         pix = wrap_periodic(ix-1, width);
 256                                         piy = wrap_periodic(iy-1, height);
 257
 258                                         nix = wrap_periodic(ix+1, width);
 259                                         niy = wrap_periodic(iy+1, height);
 260
 261                                         nnix = wrap_periodic(ix+2, width);
 262                                         nniy = wrap_periodic(iy+2, height);
 263                                         break;
 264                                 case EXTENSION_CLIP:
 265                                         if(x < 0.0f || y < 0.0f || x > 1.0f || y > 1.0f) {
 266                                                 return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 267                                         }
 268                                         /* Fall through. */
 269                                 case EXTENSION_EXTEND:
 270                                         pix = wrap_clamp(ix-1, width);
 271                                         piy = wrap_clamp(iy-1, height);
 272
 273                                         nix = wrap_clamp(ix+1, width);
 274                                         niy = wrap_clamp(iy+1, height);
 275
 276                                         nnix = wrap_clamp(ix+2, width);
 277                                         nniy = wrap_clamp(iy+2, height);
 278
 279                                         ix = wrap_clamp(ix, width);
 280                                         iy = wrap_clamp(iy, height);
 281                                         break;
 282                                 default:
 283                                         kernel_assert(0);
 284                                         return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 285                         }
 286
 287                         const int xc[4] = {pix, ix, nix, nnix};
 288                         const int yc[4] = {width * piy,
 289                                            width * iy,
 290                                            width * niy,
 291                                            width * nniy};
 292                         float u[4], v[4];
 293                         /* Some helper macro to keep code reasonable size,
 294                          * let compiler to inline all the matrix multiplications.
 295                          */
 296 #define DATA(x, y) (read(data[xc[x] + yc[y]]))
 297 #define TERM(col) \
 298                         (v[col] * (u[0] * DATA(0, col) + \
 299                                    u[1] * DATA(1, col) + \
 300                                    u[2] * DATA(2, col) + \
 301                                    u[3] * DATA(3, col)))
 302
 303                         SET_CUBIC_SPLINE_WEIGHTS(u, tx);
 304                         SET_CUBIC_SPLINE_WEIGHTS(v, ty);
 305
 306                         /* Actual interpolation. */
 307                         return TERM(0) + TERM(1) + TERM(2) + TERM(3);
 308
 309 #undef TERM
 310 #undef DATA
 311                 }
 312         }
 313
 314         ccl_always_inline float4 interp_3d(float x, float y, float z)
 315         {
 316                 return interp_3d_ex(x, y, z, interpolation);
 317         }
 318
 319         ccl_always_inline float4 interp_3d_ex_closest(float x, float y, float z)
 320         {
 321                 int ix, iy, iz;
 322                 frac(x*(float)width, &ix);
 323                 frac(y*(float)height, &iy);
 324                 frac(z*(float)depth, &iz);
 325
 326                 switch(extension) {
 327                         case EXTENSION_REPEAT:
 328                                 ix = wrap_periodic(ix, width);
 329                                 iy = wrap_periodic(iy, height);
 330                                 iz = wrap_periodic(iz, depth);
 331                                 break;
 332                         case EXTENSION_CLIP:
 333                                 if(x < 0.0f || y < 0.0f || z < 0.0f ||
 334                                    x > 1.0f || y > 1.0f || z > 1.0f)
 335                                 {
 336                                         return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 337                                 }
 338                                 /* Fall through. */
 339                         case EXTENSION_EXTEND:
 340                                 ix = wrap_clamp(ix, width);
 341                                 iy = wrap_clamp(iy, height);
 342                                 iz = wrap_clamp(iz, depth);
 343                                 break;
 344                         default:
 345                                 kernel_assert(0);
 346                                 return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 347                 }
 348
 349                 return read(data[ix + iy*width + iz*width*height]);
 350         }
 351
 352         ccl_always_inline float4 interp_3d_ex_linear(float x, float y, float z)
 353         {
 354                 int ix, iy, iz;
 355                 int nix, niy, niz;
 356
 357                 float tx = frac(x*(float)width - 0.5f, &ix);
 358                 float ty = frac(y*(float)height - 0.5f, &iy);
 359                 float tz = frac(z*(float)depth - 0.5f, &iz);
 360
 361                 switch(extension) {
 362                         case EXTENSION_REPEAT:
 363                                 ix = wrap_periodic(ix, width);
 364                                 iy = wrap_periodic(iy, height);
 365                                 iz = wrap_periodic(iz, depth);
 366
 367                                 nix = wrap_periodic(ix+1, width);
 368                                 niy = wrap_periodic(iy+1, height);
 369                                 niz = wrap_periodic(iz+1, depth);
 370                                 break;
 371                         case EXTENSION_CLIP:
 372                                 if(x < 0.0f || y < 0.0f || z < 0.0f ||
 373                                    x > 1.0f || y > 1.0f || z > 1.0f)
 374                                 {
 375                                         return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 376                                 }
 377                                 /* Fall through. */
 378                         case EXTENSION_EXTEND:
 379                                 nix = wrap_clamp(ix+1, width);
 380                                 niy = wrap_clamp(iy+1, height);
 381                                 niz = wrap_clamp(iz+1, depth);
 382
 383                                 ix = wrap_clamp(ix, width);
 384                                 iy = wrap_clamp(iy, height);
 385                                 iz = wrap_clamp(iz, depth);
 386                                 break;
 387                         default:
 388                                 kernel_assert(0);
 389                                 return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 390                 }
 391
 392                 float4 r;
 393
 394                 r  = (1.0f - tz)*(1.0f - ty)*(1.0f - tx)*read(data[ix + iy*width + iz*width*height]);
 395                 r += (1.0f - tz)*(1.0f - ty)*tx*read(data[nix + iy*width + iz*width*height]);
 396                 r += (1.0f - tz)*ty*(1.0f - tx)*read(data[ix + niy*width + iz*width*height]);
 397                 r += (1.0f - tz)*ty*tx*read(data[nix + niy*width + iz*width*height]);
 398
 399                 r += tz*(1.0f - ty)*(1.0f - tx)*read(data[ix + iy*width + niz*width*height]);
 400                 r += tz*(1.0f - ty)*tx*read(data[nix + iy*width + niz*width*height]);
 401                 r += tz*ty*(1.0f - tx)*read(data[ix + niy*width + niz*width*height]);
 402                 r += tz*ty*tx*read(data[nix + niy*width + niz*width*height]);
 403
 404                 return r;
 405         }
 406
 407         /* TODO(sergey): For some unspeakable reason both GCC-6 and Clang-3.9 are
 408          * causing stack overflow issue in this function unless it is inlined.
 409          *
 410          * Only happens for AVX2 kernel and global __KERNEL_SSE__ vectorization
 411          * enabled.
 412          */
 413 #ifdef __GNUC__
 414         ccl_always_inline
 415 #else
 416         ccl_never_inline
 417 #endif
 418         float4 interp_3d_ex_tricubic(float x, float y, float z)
 419         {
 420                 int ix, iy, iz;
 421                 int nix, niy, niz;
 422                 /* Tricubic b-spline interpolation. */
 423                 const float tx = frac(x*(float)width - 0.5f, &ix);
 424                 const float ty = frac(y*(float)height - 0.5f, &iy);
 425                 const float tz = frac(z*(float)depth - 0.5f, &iz);
 426                 int pix, piy, piz, nnix, nniy, nniz;
 427
 428                 switch(extension) {
 429                         case EXTENSION_REPEAT:
 430                                 ix = wrap_periodic(ix, width);
 431                                 iy = wrap_periodic(iy, height);
 432                                 iz = wrap_periodic(iz, depth);
 433
 434                                 pix = wrap_periodic(ix-1, width);
 435                                 piy = wrap_periodic(iy-1, height);
 436                                 piz = wrap_periodic(iz-1, depth);
 437
 438                                 nix = wrap_periodic(ix+1, width);
 439                                 niy = wrap_periodic(iy+1, height);
 440                                 niz = wrap_periodic(iz+1, depth);
 441
 442                                 nnix = wrap_periodic(ix+2, width);
 443                                 nniy = wrap_periodic(iy+2, height);
 444                                 nniz = wrap_periodic(iz+2, depth);
 445                                 break;
 446                         case EXTENSION_CLIP:
 447                                 if(x < 0.0f || y < 0.0f || z < 0.0f ||
 448                                    x > 1.0f || y > 1.0f || z > 1.0f)
 449                                 {
 450                                         return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 451                                 }
 452                                 /* Fall through. */
 453                         case EXTENSION_EXTEND:
 454                                 pix = wrap_clamp(ix-1, width);
 455                                 piy = wrap_clamp(iy-1, height);
 456                                 piz = wrap_clamp(iz-1, depth);
 457
 458                                 nix = wrap_clamp(ix+1, width);
 459                                 niy = wrap_clamp(iy+1, height);
 460                                 niz = wrap_clamp(iz+1, depth);
 461
 462                                 nnix = wrap_clamp(ix+2, width);
 463                                 nniy = wrap_clamp(iy+2, height);
 464                                 nniz = wrap_clamp(iz+2, depth);
 465
 466                                 ix = wrap_clamp(ix, width);
 467                                 iy = wrap_clamp(iy, height);
 468                                 iz = wrap_clamp(iz, depth);
 469                                 break;
 470                         default:
 471                                 kernel_assert(0);
 472                                 return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 473                 }
 474
 475                 const int xc[4] = {pix, ix, nix, nnix};
 476                 const int yc[4] = {width * piy,
 477                                    width * iy,
 478                                    width * niy,
 479                                    width * nniy};
 480                 const int zc[4] = {width * height * piz,
 481                                    width * height * iz,
 482                                    width * height * niz,
 483                                    width * height * nniz};
 484                 float u[4], v[4], w[4];
 485
 486                 /* Some helper macro to keep code reasonable size,
 487                  * let compiler to inline all the matrix multiplications.
 488                  */
 489 #define DATA(x, y, z) (read(data[xc[x] + yc[y] + zc[z]]))
 490 #define COL_TERM(col, row) \
 491                 (v[col] * (u[0] * DATA(0, col, row) + \
 492                            u[1] * DATA(1, col, row) + \
 493                            u[2] * DATA(2, col, row) + \
 494                            u[3] * DATA(3, col, row)))
 495 #define ROW_TERM(row) \
 496                 (w[row] * (COL_TERM(0, row) + \
 497                            COL_TERM(1, row) + \
 498                            COL_TERM(2, row) + \
 499                            COL_TERM(3, row)))
 500
 501                 SET_CUBIC_SPLINE_WEIGHTS(u, tx);
 502                 SET_CUBIC_SPLINE_WEIGHTS(v, ty);
 503                 SET_CUBIC_SPLINE_WEIGHTS(w, tz);
 504
 505                 /* Actual interpolation. */
 506                 return ROW_TERM(0) + ROW_TERM(1) + ROW_TERM(2) + ROW_TERM(3);
 507
 508 #undef COL_TERM
 509 #undef ROW_TERM
 510 #undef DATA
 511         }
 512
 513         ccl_always_inline float4 interp_3d_ex(float x, float y, float z,
 514                                               int interpolation = INTERPOLATION_LINEAR)
 515         {
 516                 if(UNLIKELY(!data))
 517                         return make_float4(0.0f, 0.0f, 0.0f, 0.0f);
 518
 519                 switch(interpolation) {
 520                         case INTERPOLATION_CLOSEST:
 521                                 return interp_3d_ex_closest(x, y, z);
 522                         case INTERPOLATION_LINEAR:
 523                                 return interp_3d_ex_linear(x, y, z);
 524                         default:
 525                                 return interp_3d_ex_tricubic(x, y, z);
 526                 }
 527         }
 528
 529         ccl_always_inline void dimensions_set(int width_, int height_, int depth_)
 530         {
 531                 width = width_;
 532                 height = height_;
 533                 depth = depth_;
 534         }
 535
 536         T *data;
 537         int interpolation;
 538         ExtensionType extension;
 539         int width, height, depth;
 540 #undef SET_CUBIC_SPLINE_WEIGHTS
 541 };
 542
 543 typedef texture<float4> texture_float4;
 544 typedef texture<float2> texture_float2;
 545 typedef texture<float> texture_float;
 546 typedef texture<uint> texture_uint;
 547 typedef texture<int> texture_int;
 548 typedef texture<uint4> texture_uint4;
 549 typedef texture<uchar4> texture_uchar4;
 550 typedef texture<uchar> texture_uchar;
 551 typedef texture_image<float> texture_image_float;
 552 typedef texture_image<uchar> texture_image_uchar;
 553 typedef texture_image<half> texture_image_half;
 554 typedef texture_image<float4> texture_image_float4;
 555 typedef texture_image<uchar4> texture_image_uchar4;
 556 typedef texture_image<half4> texture_image_half4;
 557
 558 /* Macros to handle different memory storage on different devices */
 559
 560 #define kernel_tex_fetch(tex, index) (kg->tex.fetch(index))
 561 #define kernel_tex_fetch_avxf(tex, index) (kg->tex.fetch_avxf(index))
 562 #define kernel_tex_fetch_ssef(tex, index) (kg->tex.fetch_ssef(index))
 563 #define kernel_tex_fetch_ssei(tex, index) (kg->tex.fetch_ssei(index))
 564 #define kernel_tex_lookup(tex, t, offset, size) (kg->tex.lookup(t, offset, size))
 565
 566 #define kernel_tex_image_interp(tex,x,y) kernel_tex_image_interp_impl(kg,tex,x,y)
 567 #define kernel_tex_image_interp_3d(tex, x, y, z) kernel_tex_image_interp_3d_impl(kg,tex,x,y,z)
 568 #define kernel_tex_image_interp_3d_ex(tex, x, y, z, interpolation) kernel_tex_image_interp_3d_ex_impl(kg,tex, x, y, z, interpolation)
 569
 570 #define kernel_data (kg->__data)
 571
 572 #ifdef __KERNEL_SSE2__
 573 typedef vector3<sseb> sse3b;
 574 typedef vector3<ssef> sse3f;
 575 typedef vector3<ssei> sse3i;
 576
 577 ccl_device_inline void print_sse3b(const char *label, sse3b& a)
 578 {
 579         print_sseb(label, a.x);
 580         print_sseb(label, a.y);
 581         print_sseb(label, a.z);
 582 }
 583
 584 ccl_device_inline void print_sse3f(const char *label, sse3f& a)
 585 {
 586         print_ssef(label, a.x);
 587         print_ssef(label, a.y);
 588         print_ssef(label, a.z);
 589 }
 590
 591 ccl_device_inline void print_sse3i(const char *label, sse3i& a)
 592 {
 593         print_ssei(label, a.x);
 594         print_ssei(label, a.y);
 595         print_ssei(label, a.z);
 596 }
 597
 598 #endif
 599
 600 CCL_NAMESPACE_END
 601
 602 #endif /* __KERNEL_COMPAT_CPU_H__ */
 603