Blender  V3.3
cuda/compat.h
Go to the documentation of this file.
1 /* SPDX-License-Identifier: Apache-2.0
2  * Copyright 2011-2022 Blender Foundation */
3 
4 #pragma once
5 
6 #define __KERNEL_GPU__
7 #define __KERNEL_CUDA__
8 #define CCL_NAMESPACE_BEGIN
9 #define CCL_NAMESPACE_END
10 
11 #ifndef ATTR_FALLTHROUGH
12 # define ATTR_FALLTHROUGH
13 #endif
14 
15 /* Manual definitions so we can compile without CUDA toolkit. */
16 
17 #ifdef __CUDACC_RTC__
18 typedef unsigned int uint32_t;
19 typedef unsigned long long uint64_t;
20 #else
21 # include <stdint.h>
22 #endif
23 
24 #ifdef CYCLES_CUBIN_CC
25 # define FLT_MIN 1.175494350822287507969e-38f
26 # define FLT_MAX 340282346638528859811704183484516925440.0f
27 # define FLT_EPSILON 1.192092896e-07F
28 #endif
29 
30 /* Qualifiers */
31 
32 #define ccl_device __device__ __inline__
33 #if __CUDA_ARCH__ < 500
34 # define ccl_device_inline __device__ __forceinline__
35 # define ccl_device_forceinline __device__ __forceinline__
36 #else
37 # define ccl_device_inline __device__ __inline__
38 # define ccl_device_forceinline __device__ __forceinline__
39 #endif
40 #define ccl_device_noinline __device__ __noinline__
41 #define ccl_device_noinline_cpu ccl_device
42 #define ccl_device_inline_method ccl_device
43 #define ccl_global
44 #define ccl_inline_constant __constant__
45 #define ccl_device_constant __constant__ __device__
46 #define ccl_constant const
47 #define ccl_gpu_shared __shared__
48 #define ccl_private
49 #define ccl_may_alias
50 #define ccl_restrict __restrict__
51 #define ccl_loop_no_unroll
52 #define ccl_align(n) __align__(n)
53 #define ccl_optional_struct_init
54 
55 /* No assert supported for CUDA */
56 
57 #define kernel_assert(cond)
58 
59 /* GPU thread, block, grid size and index */
60 
61 #define ccl_gpu_thread_idx_x (threadIdx.x)
62 #define ccl_gpu_block_dim_x (blockDim.x)
63 #define ccl_gpu_block_idx_x (blockIdx.x)
64 #define ccl_gpu_grid_dim_x (gridDim.x)
65 #define ccl_gpu_warp_size (warpSize)
66 #define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))
67 
68 #define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x)
69 #define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x)
70 
71 /* GPU warp synchronization. */
72 
73 #define ccl_gpu_syncthreads() __syncthreads()
74 #define ccl_gpu_ballot(predicate) __ballot_sync(0xFFFFFFFF, predicate)
75 
76 /* GPU texture objects */
77 
78 typedef unsigned long long CUtexObject;
81 
82 template<typename T>
84  const float x,
85  const float y)
86 {
87  return tex2D<T>(texobj, x, y);
88 }
89 
90 template<typename T>
92  const float x,
93  const float y,
94  const float z)
95 {
96  return tex3D<T>(texobj, x, y, z);
97 }
98 
99 /* Use fast math functions */
100 
101 #define cosf(x) __cosf(((float)(x)))
102 #define sinf(x) __sinf(((float)(x)))
103 #define powf(x, y) __powf(((float)(x)), ((float)(y)))
104 #define tanf(x) __tanf(((float)(x)))
105 #define logf(x) __logf(((float)(x)))
106 #define expf(x) __expf(((float)(x)))
107 
108 /* Half */
109 
110 typedef unsigned short half;
111 
113 {
114  half val;
115  asm("{ cvt.rn.f16.f32 %0, %1;}\n" : "=h"(val) : "f"(f));
116  return val;
117 }
118 
120 {
121  float val;
122  asm("{ cvt.f32.f16 %0, %1;}\n" : "=f"(val) : "h"(h));
123  return val;
124 }
125 
126 /* Types */
127 
128 #include "util/half.h"
129 #include "util/types.h"
_GL_VOID GLfloat value _GL_VOID_RET _GL_VOID const GLuint GLboolean *residences _GL_BOOL_RET _GL_VOID GLsizei GLfloat GLfloat GLfloat GLfloat const GLubyte *bitmap _GL_VOID_RET _GL_VOID GLenum const void *lists _GL_VOID_RET _GL_VOID const GLdouble *equation _GL_VOID_RET _GL_VOID GLdouble GLdouble blue _GL_VOID_RET _GL_VOID GLfloat GLfloat blue _GL_VOID_RET _GL_VOID GLint GLint blue _GL_VOID_RET _GL_VOID GLshort GLshort blue _GL_VOID_RET _GL_VOID GLubyte GLubyte blue _GL_VOID_RET _GL_VOID GLuint GLuint blue _GL_VOID_RET _GL_VOID GLushort GLushort blue _GL_VOID_RET _GL_VOID GLbyte GLbyte GLbyte alpha _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble alpha _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat alpha _GL_VOID_RET _GL_VOID GLint GLint GLint alpha _GL_VOID_RET _GL_VOID GLshort GLshort GLshort alpha _GL_VOID_RET _GL_VOID GLubyte GLubyte GLubyte alpha _GL_VOID_RET _GL_VOID GLuint GLuint GLuint alpha _GL_VOID_RET _GL_VOID GLushort GLushort GLushort alpha _GL_VOID_RET _GL_VOID GLenum mode _GL_VOID_RET _GL_VOID GLint GLsizei GLsizei GLenum type _GL_VOID_RET _GL_VOID GLsizei GLenum GLenum const void *pixels _GL_VOID_RET _GL_VOID const void *pointer _GL_VOID_RET _GL_VOID GLdouble v _GL_VOID_RET _GL_VOID GLfloat v _GL_VOID_RET _GL_VOID GLint GLint i2 _GL_VOID_RET _GL_VOID GLint j _GL_VOID_RET _GL_VOID GLfloat param _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble GLdouble GLdouble zFar _GL_VOID_RET _GL_UINT GLdouble *equation _GL_VOID_RET _GL_VOID GLenum GLint *params _GL_VOID_RET _GL_VOID GLenum GLfloat *v _GL_VOID_RET _GL_VOID GLenum GLfloat *params _GL_VOID_RET _GL_VOID GLfloat *values _GL_VOID_RET _GL_VOID GLushort *values _GL_VOID_RET _GL_VOID GLenum GLfloat *params _GL_VOID_RET _GL_VOID GLenum GLdouble *params _GL_VOID_RET _GL_VOID GLenum GLint *params _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_BOOL GLfloat param _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID GLenum GLfloat param _GL_VOID_RET _GL_VOID GLenum GLint param _GL_VOID_RET _GL_VOID GLushort pattern _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLint const GLdouble *points _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLint GLdouble GLdouble GLint GLint const GLdouble *points _GL_VOID_RET _GL_VOID GLdouble GLdouble u2 _GL_VOID_RET _GL_VOID GLdouble GLdouble GLint GLdouble GLdouble v2 _GL_VOID_RET _GL_VOID GLenum GLfloat param _GL_VOID_RET _GL_VOID GLenum GLint param _GL_VOID_RET _GL_VOID GLenum mode _GL_VOID_RET _GL_VOID GLdouble GLdouble nz _GL_VOID_RET _GL_VOID GLfloat GLfloat nz _GL_VOID_RET _GL_VOID GLint GLint nz _GL_VOID_RET _GL_VOID GLshort GLshort nz _GL_VOID_RET _GL_VOID GLsizei const void *pointer _GL_VOID_RET _GL_VOID GLsizei const GLfloat *values _GL_VOID_RET _GL_VOID GLsizei const GLushort *values _GL_VOID_RET _GL_VOID GLint param _GL_VOID_RET _GL_VOID const GLuint const GLclampf *priorities _GL_VOID_RET _GL_VOID GLdouble y _GL_VOID_RET _GL_VOID GLfloat y _GL_VOID_RET _GL_VOID GLint y _GL_VOID_RET _GL_VOID GLshort y _GL_VOID_RET _GL_VOID GLdouble GLdouble z _GL_VOID_RET _GL_VOID GLfloat GLfloat z _GL_VOID_RET _GL_VOID GLint GLint z _GL_VOID_RET _GL_VOID GLshort GLshort z _GL_VOID_RET _GL_VOID GLdouble GLdouble z
_GL_VOID GLfloat value _GL_VOID_RET _GL_VOID const GLuint GLboolean *residences _GL_BOOL_RET _GL_VOID GLsizei GLfloat GLfloat GLfloat GLfloat const GLubyte *bitmap _GL_VOID_RET _GL_VOID GLenum const void *lists _GL_VOID_RET _GL_VOID const GLdouble *equation _GL_VOID_RET _GL_VOID GLdouble GLdouble blue _GL_VOID_RET _GL_VOID GLfloat GLfloat blue _GL_VOID_RET _GL_VOID GLint GLint blue _GL_VOID_RET _GL_VOID GLshort GLshort blue _GL_VOID_RET _GL_VOID GLubyte GLubyte blue _GL_VOID_RET _GL_VOID GLuint GLuint blue _GL_VOID_RET _GL_VOID GLushort GLushort blue _GL_VOID_RET _GL_VOID GLbyte GLbyte GLbyte alpha _GL_VOID_RET _GL_VOID GLdouble GLdouble GLdouble alpha _GL_VOID_RET _GL_VOID GLfloat GLfloat GLfloat alpha _GL_VOID_RET _GL_VOID GLint GLint GLint alpha _GL_VOID_RET _GL_VOID GLshort GLshort GLshort alpha _GL_VOID_RET _GL_VOID GLubyte GLubyte GLubyte alpha _GL_VOID_RET _GL_VOID GLuint GLuint GLuint alpha _GL_VOID_RET _GL_VOID GLushort GLushort GLushort alpha _GL_VOID_RET _GL_VOID GLenum mode _GL_VOID_RET _GL_VOID GLint y
Definition: half.h:41
ccl_device_forceinline T ccl_gpu_tex_object_read_2D(const ccl_gpu_tex_object_2D texobj, const float x, const float y)
Definition: cuda/compat.h:83
#define ccl_device_forceinline
Definition: cuda/compat.h:35
__device__ half __float2half(const float f)
Definition: cuda/compat.h:112
CUtexObject ccl_gpu_tex_object_3D
Definition: cuda/compat.h:80
unsigned short half
Definition: cuda/compat.h:110
__device__ float __half2float(const half h)
Definition: cuda/compat.h:119
ccl_device_forceinline T ccl_gpu_tex_object_read_3D(const ccl_gpu_tex_object_3D texobj, const float x, const float y, const float z)
Definition: cuda/compat.h:91
unsigned long long CUtexObject
Definition: cuda/compat.h:78
CUtexObject ccl_gpu_tex_object_2D
Definition: cuda/compat.h:79
#define T
#define __device__
Definition: metal/compat.h:248
unsigned int uint32_t
Definition: stdint.h:80
unsigned __int64 uint64_t
Definition: stdint.h:90