archive/blender-archive
Archived
1
1
Fork 0
Code Pull Requests Packages Activity
This repository has been archived on 2023-10-09. You can view files and clone it, but cannot push or open issues or pull requests.
Files
f680c1b54a28a02fb86271bca649da0660542e9a
blender-archive/intern/cycles/kernel/kernel_compat_cuda.h

94 lines
2.6 KiB
C++
Raw Normal View History

Cycles render engine, initial commit. This is the engine itself, blender modifications and build instructions will follow later. Cycles uses code from some great open source projects, many thanks them: * BVH building and traversal code from NVidia's "Understanding the Efficiency of Ray Traversal on GPUs": http://code.google.com/p/understanding-the-efficiency-of-ray-traversal-on-gpus/ * Open Shading Language for a large part of the shading system: http://code.google.com/p/openshadinglanguage/ * Blender for procedural textures and a few other nodes. * Approximate Catmull Clark subdivision from NVidia Mesh tools: http://code.google.com/p/nvidia-mesh-tools/ * Sobol direction vectors from: http://web.maths.unsw.edu.au/~fkuo/sobol/ * Film response functions from: http://www.cs.columbia.edu/CAVE/software/softlib/dorf.php
2011-04-27 11:58:34 +00:00
/*
Cycles: relicense GNU GPL source code to Apache version 2.0. More information in this post: http://code.blender.org/ Thanks to all contributes for giving their permission!
2013-08-18 14:16:15 +00:00
* Copyright 2011-2013 Blender Foundation
Cycles render engine, initial commit. This is the engine itself, blender modifications and build instructions will follow later. Cycles uses code from some great open source projects, many thanks them: * BVH building and traversal code from NVidia's "Understanding the Efficiency of Ray Traversal on GPUs": http://code.google.com/p/understanding-the-efficiency-of-ray-traversal-on-gpus/ * Open Shading Language for a large part of the shading system: http://code.google.com/p/openshadinglanguage/ * Blender for procedural textures and a few other nodes. * Approximate Catmull Clark subdivision from NVidia Mesh tools: http://code.google.com/p/nvidia-mesh-tools/ * Sobol direction vectors from: http://web.maths.unsw.edu.au/~fkuo/sobol/ * Film response functions from: http://www.cs.columbia.edu/CAVE/software/softlib/dorf.php
2011-04-27 11:58:34 +00:00
*
Cycles: relicense GNU GPL source code to Apache version 2.0. More information in this post: http://code.blender.org/ Thanks to all contributes for giving their permission!
2013-08-18 14:16:15 +00:00
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
Cycles render engine, initial commit. This is the engine itself, blender modifications and build instructions will follow later. Cycles uses code from some great open source projects, many thanks them: * BVH building and traversal code from NVidia's "Understanding the Efficiency of Ray Traversal on GPUs": http://code.google.com/p/understanding-the-efficiency-of-ray-traversal-on-gpus/ * Open Shading Language for a large part of the shading system: http://code.google.com/p/openshadinglanguage/ * Blender for procedural textures and a few other nodes. * Approximate Catmull Clark subdivision from NVidia Mesh tools: http://code.google.com/p/nvidia-mesh-tools/ * Sobol direction vectors from: http://web.maths.unsw.edu.au/~fkuo/sobol/ * Film response functions from: http://www.cs.columbia.edu/CAVE/software/softlib/dorf.php
2011-04-27 11:58:34 +00:00
*
Cycles: relicense GNU GPL source code to Apache version 2.0. More information in this post: http://code.blender.org/ Thanks to all contributes for giving their permission!
2013-08-18 14:16:15 +00:00
* http://www.apache.org/licenses/LICENSE-2.0
Cycles render engine, initial commit. This is the engine itself, blender modifications and build instructions will follow later. Cycles uses code from some great open source projects, many thanks them: * BVH building and traversal code from NVidia's "Understanding the Efficiency of Ray Traversal on GPUs": http://code.google.com/p/understanding-the-efficiency-of-ray-traversal-on-gpus/ * Open Shading Language for a large part of the shading system: http://code.google.com/p/openshadinglanguage/ * Blender for procedural textures and a few other nodes. * Approximate Catmull Clark subdivision from NVidia Mesh tools: http://code.google.com/p/nvidia-mesh-tools/ * Sobol direction vectors from: http://web.maths.unsw.edu.au/~fkuo/sobol/ * Film response functions from: http://www.cs.columbia.edu/CAVE/software/softlib/dorf.php
2011-04-27 11:58:34 +00:00
*
Cycles: relicense GNU GPL source code to Apache version 2.0. More information in this post: http://code.blender.org/ Thanks to all contributes for giving their permission!
2013-08-18 14:16:15 +00:00
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
Cleanup: Fix Cycles Apache header. This was already mixed a bit, but the dot belongs there.
2014-12-25 02:50:24 +01:00
* limitations under the License.
Cycles render engine, initial commit. This is the engine itself, blender modifications and build instructions will follow later. Cycles uses code from some great open source projects, many thanks them: * BVH building and traversal code from NVidia's "Understanding the Efficiency of Ray Traversal on GPUs": http://code.google.com/p/understanding-the-efficiency-of-ray-traversal-on-gpus/ * Open Shading Language for a large part of the shading system: http://code.google.com/p/openshadinglanguage/ * Blender for procedural textures and a few other nodes. * Approximate Catmull Clark subdivision from NVidia Mesh tools: http://code.google.com/p/nvidia-mesh-tools/ * Sobol direction vectors from: http://web.maths.unsw.edu.au/~fkuo/sobol/ * Film response functions from: http://www.cs.columbia.edu/CAVE/software/softlib/dorf.php
2011-04-27 11:58:34 +00:00
*/
#ifndef __KERNEL_COMPAT_CUDA_H__
#define __KERNEL_COMPAT_CUDA_H__
#define __KERNEL_GPU__
#define __KERNEL_CUDA__
Cycles: * OpenCL now only uses GPU/Accelerator devices, it's only confusing if CPU device is used, easy to enable in the code for debugging. * OpenCL kernel binaries are now cached for faster startup after the first time compiling. * CUDA kernels can now be compiled and cached at runtime if the CUDA toolkit is installed. This means that even if the build does not have CUDA enabled, it's still possible to use it as long as you install the toolkit.
2011-09-09 12:04:39 +00:00
#define CCL_NAMESPACE_BEGIN
#define CCL_NAMESPACE_END
Cycles render engine, initial commit. This is the engine itself, blender modifications and build instructions will follow later. Cycles uses code from some great open source projects, many thanks them: * BVH building and traversal code from NVidia's "Understanding the Efficiency of Ray Traversal on GPUs": http://code.google.com/p/understanding-the-efficiency-of-ray-traversal-on-gpus/ * Open Shading Language for a large part of the shading system: http://code.google.com/p/openshadinglanguage/ * Blender for procedural textures and a few other nodes. * Approximate Catmull Clark subdivision from NVidia Mesh tools: http://code.google.com/p/nvidia-mesh-tools/ * Sobol direction vectors from: http://web.maths.unsw.edu.au/~fkuo/sobol/ * Film response functions from: http://www.cs.columbia.edu/CAVE/software/softlib/dorf.php
2011-04-27 11:58:34 +00:00
Cycles: Initial work towards selective nodes support compilation The goal is to be able to compile kernel with nodes which are actually needed to render current scene, hence improving performance of the kernel, The idea is: - Have few node groups, starting with a group which contains nodes are used really often, and then couple of groups which will be extension of this one. - Have feature-based nodes disabling, so it's possible to disable nodes related to features which are not used with the currently used nodes group. This commit only lays down needed routines for this approach, actual split will happen later after gathering statistics from bunch of production scenes.
2015-05-09 19:22:16 +05:00
/* Selective nodes compilation. */
#ifndef __NODES_MAX_GROUP__
# define __NODES_MAX_GROUP__ NODE_GROUP_LEVEL_MAX
#endif
#ifndef __NODES_FEATURES__
# define __NODES_FEATURES__ NODE_FEATURE_ALL
#endif
Cycles render engine, initial commit. This is the engine itself, blender modifications and build instructions will follow later. Cycles uses code from some great open source projects, many thanks them: * BVH building and traversal code from NVidia's "Understanding the Efficiency of Ray Traversal on GPUs": http://code.google.com/p/understanding-the-efficiency-of-ray-traversal-on-gpus/ * Open Shading Language for a large part of the shading system: http://code.google.com/p/openshadinglanguage/ * Blender for procedural textures and a few other nodes. * Approximate Catmull Clark subdivision from NVidia Mesh tools: http://code.google.com/p/nvidia-mesh-tools/ * Sobol direction vectors from: http://web.maths.unsw.edu.au/~fkuo/sobol/ * Film response functions from: http://www.cs.columbia.edu/CAVE/software/softlib/dorf.php
2011-04-27 11:58:34 +00:00
#include <cuda.h>
#include <float.h>
/* Qualifier wrappers for different names on different devices */
Cycles: change __device and similar qualifiers to ccl_device in kernel code. This to avoids build conflicts with libc++ on FreeBSD, these __ prefixed values are reserved for compilers. I apologize to anyone who has patches or branches and has to go through the pain of merging this change, it may be easiest to do these same replacements in your code and then apply/merge the patch. Ref T37477.
2013-11-16 00:17:10 +01:00
#define ccl_device __device__ __inline__
#define ccl_device_inline __device__ __inline__
#define ccl_device_noinline __device__ __noinline__
#define ccl_global
#define ccl_constant
#define ccl_may_alias
Cycles render engine, initial commit. This is the engine itself, blender modifications and build instructions will follow later. Cycles uses code from some great open source projects, many thanks them: * BVH building and traversal code from NVidia's "Understanding the Efficiency of Ray Traversal on GPUs": http://code.google.com/p/understanding-the-efficiency-of-ray-traversal-on-gpus/ * Open Shading Language for a large part of the shading system: http://code.google.com/p/openshadinglanguage/ * Blender for procedural textures and a few other nodes. * Approximate Catmull Clark subdivision from NVidia Mesh tools: http://code.google.com/p/nvidia-mesh-tools/ * Sobol direction vectors from: http://web.maths.unsw.edu.au/~fkuo/sobol/ * Film response functions from: http://www.cs.columbia.edu/CAVE/software/softlib/dorf.php
2011-04-27 11:58:34 +00:00
/* No assert supported for CUDA */
#define kernel_assert(cond)
Cycles: viewport render now takes scene color management settings into account, except for curves, that's still missing from the OpenColorIO GLSL shader. The pixels are stored in a half float texture, converterd from full float with native GPU instructions and SIMD on the CPU, so it should be pretty quick. Using a GLSL shader is useful for GPU render because it avoids a copy through CPU memory.
2013-08-30 23:49:38 +00:00
/* Types */
Code cleanup: move half float functions to separate header file.
2014-01-15 15:29:22 +01:00
#include "util_half.h"
Cycles: viewport render now takes scene color management settings into account, except for curves, that's still missing from the OpenColorIO GLSL shader. The pixels are stored in a half float texture, converterd from full float with native GPU instructions and SIMD on the CPU, so it should be pretty quick. Using a GLSL shader is useful for GPU render because it avoids a copy through CPU memory.
2013-08-30 23:49:38 +00:00
#include "util_types.h"
Cycles render engine, initial commit. This is the engine itself, blender modifications and build instructions will follow later. Cycles uses code from some great open source projects, many thanks them: * BVH building and traversal code from NVidia's "Understanding the Efficiency of Ray Traversal on GPUs": http://code.google.com/p/understanding-the-efficiency-of-ray-traversal-on-gpus/ * Open Shading Language for a large part of the shading system: http://code.google.com/p/openshadinglanguage/ * Blender for procedural textures and a few other nodes. * Approximate Catmull Clark subdivision from NVidia Mesh tools: http://code.google.com/p/nvidia-mesh-tools/ * Sobol direction vectors from: http://web.maths.unsw.edu.au/~fkuo/sobol/ * Film response functions from: http://www.cs.columbia.edu/CAVE/software/softlib/dorf.php
2011-04-27 11:58:34 +00:00
/* Textures */
typedef texture<float4, 1> texture_float4;
Sample as Lamp option for world shaders, to enable multiple importance sampling. By default lighting from the world is computed solely with indirect light sampling. However for more complex environment maps this can be too noisy, as sampling the BSDF may not easily find the highlights in the environment map image. By enabling this option, the world background will be sampled as a lamp, with lighter parts automatically given more samples. Map Resolution specifies the size of the importance map (res x res). Before rendering starts, an importance map is generated by "baking" a grayscale image from the world shader. This will then be used to determine which parts of the background are light and so should receive more samples than darker parts. Higher resolutions will result in more accurate sampling but take more setup time and memory. Patch by Mike Farnsworth, thanks!
2012-01-20 17:49:17 +00:00
typedef texture<float2, 1> texture_float2;
Cycles render engine, initial commit. This is the engine itself, blender modifications and build instructions will follow later. Cycles uses code from some great open source projects, many thanks them: * BVH building and traversal code from NVidia's "Understanding the Efficiency of Ray Traversal on GPUs": http://code.google.com/p/understanding-the-efficiency-of-ray-traversal-on-gpus/ * Open Shading Language for a large part of the shading system: http://code.google.com/p/openshadinglanguage/ * Blender for procedural textures and a few other nodes. * Approximate Catmull Clark subdivision from NVidia Mesh tools: http://code.google.com/p/nvidia-mesh-tools/ * Sobol direction vectors from: http://web.maths.unsw.edu.au/~fkuo/sobol/ * Film response functions from: http://www.cs.columbia.edu/CAVE/software/softlib/dorf.php
2011-04-27 11:58:34 +00:00
typedef texture<float, 1> texture_float;
typedef texture<uint, 1> texture_uint;
typedef texture<int, 1> texture_int;
typedef texture<uint4, 1> texture_uint4;
Cycles: OpenCL image texture support, fix an attribute node issue and refactor feature enabling #defines a bit.
2012-05-13 12:32:44 +00:00
typedef texture<uchar4, 1> texture_uchar4;
Cycles render engine, initial commit. This is the engine itself, blender modifications and build instructions will follow later. Cycles uses code from some great open source projects, many thanks them: * BVH building and traversal code from NVidia's "Understanding the Efficiency of Ray Traversal on GPUs": http://code.google.com/p/understanding-the-efficiency-of-ray-traversal-on-gpus/ * Open Shading Language for a large part of the shading system: http://code.google.com/p/openshadinglanguage/ * Blender for procedural textures and a few other nodes. * Approximate Catmull Clark subdivision from NVidia Mesh tools: http://code.google.com/p/nvidia-mesh-tools/ * Sobol direction vectors from: http://web.maths.unsw.edu.au/~fkuo/sobol/ * Film response functions from: http://www.cs.columbia.edu/CAVE/software/softlib/dorf.php
2011-04-27 11:58:34 +00:00
typedef texture<float4, 2> texture_image_float4;
typedef texture<uchar4, 2, cudaReadModeNormalizedFloat> texture_image_uchar4;
/* Macros to handle different memory storage on different devices */
Fix #35684: cycles unable to use full 6GB of memory on NVidia Titan GPU. We now use arrays instead of textures for general storage on this card (image textures are still stored as texture). Textures were found to be faster on older cards, but the limits on 1D texture size have not increased along with the memory size, which meant that the full 6 GB could not be used. The performance actually seems to be slightly better with arrays in some tests on Titan. For older cards there seems to be a bit of a mix, some are better and others not. We may change those to use arrays too, but more testing is needed, only Titan and Tesla K20 (sm_35) is changed for now. The fact that arrays are faster is a bit surprising, as others found textures to be faster on Kepler. However even if they were, the memory limitation is more important to solve anyway. https://research.nvidia.com/publication/understanding-efficiency-ray-traversal-gpus-kepler-and-fermi-addendum
2013-09-27 19:09:31 +00:00
/* In order to use full 6GB of memory on Titan cards, use arrays instead
* of textures. On earlier cards this seems slower, but on Titan it is
* actually slightly faster in tests. */
Cycles CUDA: make CUDA toolkit 6.0 the official supported version. This also updates the configurations to build kernels for compute capability 5.0 cards, when using and older CUDA toolkit version this will be skipped. Also includes tweaks to improve performance with this version: * Increase max registers on sm_30, sm_35 and sm_50 * No longer use texture storage on sm_30
2014-04-30 10:54:17 +02:00
#if __CUDA_ARCH__ < 300
Fix #35684: cycles unable to use full 6GB of memory on NVidia Titan GPU. We now use arrays instead of textures for general storage on this card (image textures are still stored as texture). Textures were found to be faster on older cards, but the limits on 1D texture size have not increased along with the memory size, which meant that the full 6 GB could not be used. The performance actually seems to be slightly better with arrays in some tests on Titan. For older cards there seems to be a bit of a mix, some are better and others not. We may change those to use arrays too, but more testing is needed, only Titan and Tesla K20 (sm_35) is changed for now. The fact that arrays are faster is a bit surprising, as others found textures to be faster on Kepler. However even if they were, the memory limitation is more important to solve anyway. https://research.nvidia.com/publication/understanding-efficiency-ray-traversal-gpus-kepler-and-fermi-addendum
2013-09-27 19:09:31 +00:00
#define __KERNEL_CUDA_TEX_STORAGE__
#endif
#ifdef __KERNEL_CUDA_TEX_STORAGE__
Cycles render engine, initial commit. This is the engine itself, blender modifications and build instructions will follow later. Cycles uses code from some great open source projects, many thanks them: * BVH building and traversal code from NVidia's "Understanding the Efficiency of Ray Traversal on GPUs": http://code.google.com/p/understanding-the-efficiency-of-ray-traversal-on-gpus/ * Open Shading Language for a large part of the shading system: http://code.google.com/p/openshadinglanguage/ * Blender for procedural textures and a few other nodes. * Approximate Catmull Clark subdivision from NVidia Mesh tools: http://code.google.com/p/nvidia-mesh-tools/ * Sobol direction vectors from: http://web.maths.unsw.edu.au/~fkuo/sobol/ * Film response functions from: http://www.cs.columbia.edu/CAVE/software/softlib/dorf.php
2011-04-27 11:58:34 +00:00
#define kernel_tex_fetch(t, index) tex1Dfetch(t, index)
Fix #35684: cycles unable to use full 6GB of memory on NVidia Titan GPU. We now use arrays instead of textures for general storage on this card (image textures are still stored as texture). Textures were found to be faster on older cards, but the limits on 1D texture size have not increased along with the memory size, which meant that the full 6 GB could not be used. The performance actually seems to be slightly better with arrays in some tests on Titan. For older cards there seems to be a bit of a mix, some are better and others not. We may change those to use arrays too, but more testing is needed, only Titan and Tesla K20 (sm_35) is changed for now. The fact that arrays are faster is a bit surprising, as others found textures to be faster on Kepler. However even if they were, the memory limitation is more important to solve anyway. https://research.nvidia.com/publication/understanding-efficiency-ray-traversal-gpus-kepler-and-fermi-addendum
2013-09-27 19:09:31 +00:00
#else
#define kernel_tex_fetch(t, index) t[(index)]
#endif
Cycles render engine, initial commit. This is the engine itself, blender modifications and build instructions will follow later. Cycles uses code from some great open source projects, many thanks them: * BVH building and traversal code from NVidia's "Understanding the Efficiency of Ray Traversal on GPUs": http://code.google.com/p/understanding-the-efficiency-of-ray-traversal-on-gpus/ * Open Shading Language for a large part of the shading system: http://code.google.com/p/openshadinglanguage/ * Blender for procedural textures and a few other nodes. * Approximate Catmull Clark subdivision from NVidia Mesh tools: http://code.google.com/p/nvidia-mesh-tools/ * Sobol direction vectors from: http://web.maths.unsw.edu.au/~fkuo/sobol/ * Film response functions from: http://www.cs.columbia.edu/CAVE/software/softlib/dorf.php
2011-04-27 11:58:34 +00:00
#define kernel_tex_image_interp(t, x, y) tex2D(t, x, y)
#define kernel_data __data
Cycles: fixes to make CUDA 4.2 work, compiling gave errors in shadows and other places, was mainly due to instancing not working, but also found issues in procedural textures. The problem was with --use_fast_math, this seems to now have way lower precision for some operations. Disabled this flag and selectively use fast math functions. Did not find performance regression on GTX 460 after doing this.
2012-05-28 19:21:13 +00:00
/* Use fast math functions */
Cycles: correct math wrappers include the parens around value before cast, in some cases was causing double/float promotion by only casting the left value.
2014-10-08 00:02:46 +02:00
#define cosf(x) __cosf(((float)(x)))
#define sinf(x) __sinf(((float)(x)))
#define powf(x, y) __powf(((float)(x)), ((float)(y)))
#define tanf(x) __tanf(((float)(x)))
#define logf(x) __logf(((float)(x)))
#define expf(x) __expf(((float)(x)))
Cycles: fixes to make CUDA 4.2 work, compiling gave errors in shadows and other places, was mainly due to instancing not working, but also found issues in procedural textures. The problem was with --use_fast_math, this seems to now have way lower precision for some operations. Disabled this flag and selectively use fast math functions. Did not find performance regression on GTX 460 after doing this.
2012-05-28 19:21:13 +00:00
Cycles render engine, initial commit. This is the engine itself, blender modifications and build instructions will follow later. Cycles uses code from some great open source projects, many thanks them: * BVH building and traversal code from NVidia's "Understanding the Efficiency of Ray Traversal on GPUs": http://code.google.com/p/understanding-the-efficiency-of-ray-traversal-on-gpus/ * Open Shading Language for a large part of the shading system: http://code.google.com/p/openshadinglanguage/ * Blender for procedural textures and a few other nodes. * Approximate Catmull Clark subdivision from NVidia Mesh tools: http://code.google.com/p/nvidia-mesh-tools/ * Sobol direction vectors from: http://web.maths.unsw.edu.au/~fkuo/sobol/ * Film response functions from: http://www.cs.columbia.edu/CAVE/software/softlib/dorf.php
2011-04-27 11:58:34 +00:00
#endif /* __KERNEL_COMPAT_CUDA_H__ */
Copy Permalink