intern/cycles/kernel/kernel_compat_cuda.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_CUDA_H__
  18 #define __KERNEL_COMPAT_CUDA_H__
  19
  20 #define __KERNEL_GPU__
  21 #define __KERNEL_CUDA__
  22 #define CCL_NAMESPACE_BEGIN
  23 #define CCL_NAMESPACE_END
  24
  25 /* Selective nodes compilation. */
  26 #ifndef __NODES_MAX_GROUP__
  27 #  define __NODES_MAX_GROUP__ NODE_GROUP_LEVEL_MAX
  28 #endif
  29 #ifndef __NODES_FEATURES__
  30 #  define __NODES_FEATURES__ NODE_FEATURE_ALL
  31 #endif
  32
  33 /* Manual definitions so we can compile without CUDA toolkit. */
  34
  35 typedef unsigned int uint32_t;
  36 typedef unsigned long long uint64_t;
  37 typedef unsigned short half;
  38 typedef unsigned long long CUtexObject;
  39
  40 #define FLT_MIN 1.175494350822287507969e-38f
  41 #define FLT_MAX 340282346638528859811704183484516925440.0f
  42
  43 __device__ half __float2half(const float f)
  44 {
  45        half val;
  46        asm("{  cvt.rn.f16.f32 %0, %1;}\n" : "=h"(val) : "f"(f));
  47        return val;
  48 }
  49
  50 /* Qualifier wrappers for different names on different devices */
  51
  52 #define ccl_device  __device__ __inline__
  53 #if __CUDA_ARCH__ < 500
  54 #  define ccl_device_inline  __device__ __forceinline__
  55 #  define ccl_device_forceinline  __device__ __forceinline__
  56 #else
  57 #  define ccl_device_inline  __device__ __inline__
  58 #  define ccl_device_forceinline  __device__ __forceinline__
  59 #endif
  60 #define ccl_device_noinline  __device__ __noinline__
  61 #define ccl_global
  62 #define ccl_static_constant __constant__
  63 #define ccl_constant const
  64 #define ccl_local __shared__
  65 #define ccl_local_param
  66 #define ccl_private
  67 #define ccl_may_alias
  68 #define ccl_addr_space
  69 #define ccl_restrict __restrict__
  70 /* TODO(sergey): In theory we might use references with CUDA, however
  71  * performance impact yet to be investigated.
  72  */
  73 #define ccl_ref
  74 #define ccl_align(n) __align__(n)
  75
  76 #define ATTR_FALLTHROUGH
  77
  78 #define CCL_MAX_LOCAL_SIZE (CUDA_THREADS_BLOCK_WIDTH*CUDA_THREADS_BLOCK_WIDTH)
  79
  80
  81 /* No assert supported for CUDA */
  82
  83 #define kernel_assert(cond)
  84
  85 /* Types */
  86
  87 #include "util/util_half.h"
  88 #include "util/util_types.h"
  89
  90 /* Work item functions */
  91
  92 ccl_device_inline uint ccl_local_id(uint d)
  93 {
  94         switch(d) {
  95                 case 0: return threadIdx.x;
  96                 case 1: return threadIdx.y;
  97                 case 2: return threadIdx.z;
  98                 default: return 0;
  99         }
 100 }
 101
 102 #define ccl_global_id(d) (ccl_group_id(d) * ccl_local_size(d) + ccl_local_id(d))
 103
 104 ccl_device_inline uint ccl_local_size(uint d)
 105 {
 106         switch(d) {
 107                 case 0: return blockDim.x;
 108                 case 1: return blockDim.y;
 109                 case 2: return blockDim.z;
 110                 default: return 0;
 111         }
 112 }
 113
 114 #define ccl_global_size(d) (ccl_num_groups(d) * ccl_local_size(d))
 115
 116 ccl_device_inline uint ccl_group_id(uint d)
 117 {
 118         switch(d) {
 119                 case 0: return blockIdx.x;
 120                 case 1: return blockIdx.y;
 121                 case 2: return blockIdx.z;
 122                 default: return 0;
 123         }
 124 }
 125
 126 ccl_device_inline uint ccl_num_groups(uint d)
 127 {
 128         switch(d) {
 129                 case 0: return gridDim.x;
 130                 case 1: return gridDim.y;
 131                 case 2: return gridDim.z;
 132                 default: return 0;
 133         }
 134 }
 135
 136 /* Textures */
 137
 138 /* Use arrays for regular data. This is a little slower than textures on Fermi,
 139  * but allows for cleaner code and we will stop supporting Fermi soon. */
 140 #define kernel_tex_fetch(t, index) t[(index)]
 141
 142 /* On Kepler (6xx) and above, we use Bindless Textures for images.
 143  * On Fermi cards (4xx and 5xx), we have to use regular textures. */
 144 #if __CUDA_ARCH__ < 300
 145 typedef texture<float4, 2> texture_image_float4;
 146 typedef texture<float4, 3> texture_image3d_float4;
 147 typedef texture<uchar4, 2, cudaReadModeNormalizedFloat> texture_image_uchar4;
 148 #endif
 149
 150 #define kernel_data __data
 151
 152 /* Use fast math functions */
 153
 154 #define cosf(x) __cosf(((float)(x)))
 155 #define sinf(x) __sinf(((float)(x)))
 156 #define powf(x, y) __powf(((float)(x)), ((float)(y)))
 157 #define tanf(x) __tanf(((float)(x)))
 158 #define logf(x) __logf(((float)(x)))
 159 #define expf(x) __expf(((float)(x)))
 160
 161 #endif /* __KERNEL_COMPAT_CUDA_H__ */