/*
- * Copyright 2011, Blender Foundation.
+ * Copyright 2011-2013 Blender Foundation
*
- * This program is free software; you can redistribute it and/or
- * modify it under the terms of the GNU General Public License
- * as published by the Free Software Foundation; either version 2
- * of the License, or (at your option) any later version.
+ * 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
*
- * This program is distributed in the hope that it will be useful,
- * but WITHOUT ANY WARRANTY; without even the implied warranty of
- * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
- * GNU General Public License for more details.
+ * http://www.apache.org/licenses/LICENSE-2.0
*
- * You should have received a copy of the GNU General Public License
- * along with this program; if not, write to the Free Software Foundation,
- * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+ * 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
+ * limitations under the License.
*/
/* Constant Globals */
-#ifdef __KERNEL_CPU__
-
-#ifdef __OSL__
-#include "osl_globals.h"
-#endif
-
-#endif
-
CCL_NAMESPACE_BEGIN
/* On the CPU, we pass along the struct KernelGlobals to nearly everywhere in
#ifdef __KERNEL_CPU__
-#define MAX_BYTE_IMAGES 512
-#define MAX_FLOAT_IMAGES 5
+# ifdef __OSL__
+struct OSLGlobals;
+struct OSLThreadData;
+struct OSLShadingSystem;
+# endif
typedef struct KernelGlobals {
- texture_image_uchar4 texture_byte_images[MAX_BYTE_IMAGES];
- texture_image_float4 texture_float_images[MAX_FLOAT_IMAGES];
+ texture_image_uchar4 texture_byte4_images[TEX_NUM_BYTE4_IMAGES_CPU];
+ texture_image_float4 texture_float4_images[TEX_NUM_FLOAT4_IMAGES_CPU];
-#define KERNEL_TEX(type, ttype, name) ttype name;
-#define KERNEL_IMAGE_TEX(type, ttype, name)
-#include "kernel_textures.h"
+# define KERNEL_TEX(type, ttype, name) ttype name;
+# define KERNEL_IMAGE_TEX(type, ttype, name)
+# include "kernel_textures.h"
KernelData __data;
-#ifdef __OSL__
+# ifdef __OSL__
/* On the CPU, we also have the OSL globals here. Most data structures are shared
* with SVM, the difference is in the shaders and object/mesh attributes. */
- OSLGlobals osl;
-#endif
+ OSLGlobals *osl;
+ OSLShadingSystem *osl_ss;
+ OSLThreadData *osl_tdata;
+# endif
} KernelGlobals;
-#endif
+#endif /* __KERNEL_CPU__ */
/* For CUDA, constant memory textures must be globals, so we can't put them
* into a struct. As a result we don't actually use this struct and use actual
__constant__ KernelData __data;
typedef struct KernelGlobals {} KernelGlobals;
-#define KERNEL_TEX(type, ttype, name) ttype name;
-#define KERNEL_IMAGE_TEX(type, ttype, name) ttype name;
-#include "kernel_textures.h"
+# ifdef __KERNEL_CUDA_TEX_STORAGE__
+# define KERNEL_TEX(type, ttype, name) ttype name;
+# else
+# define KERNEL_TEX(type, ttype, name) const __constant__ __device__ type *name;
+# endif
+# define KERNEL_IMAGE_TEX(type, ttype, name) ttype name;
+# include "kernel_textures.h"
-#endif
+#endif /* __KERNEL_CUDA__ */
/* OpenCL */
#ifdef __KERNEL_OPENCL__
-typedef struct KernelGlobals {
- __constant KernelData *data;
+typedef ccl_addr_space struct KernelGlobals {
+ ccl_constant KernelData *data;
+
+# define KERNEL_TEX(type, ttype, name) \
+ ccl_global type *name;
+# include "kernel_textures.h"
-#define KERNEL_TEX(type, ttype, name) \
- __global type *name;
-#include "kernel_textures.h"
+# ifdef __SPLIT_KERNEL__
+ ShaderData *sd_input;
+ Intersection *isect_shadow;
+# endif
} KernelGlobals;
-#endif
+#endif /* __KERNEL_OPENCL__ */
+
+/* Interpolated lookup table access */
+
+ccl_device float lookup_table_read(KernelGlobals *kg, float x, int offset, int size)
+{
+ x = saturate(x)*(size-1);
+
+ int index = min(float_to_int(x), size-1);
+ int nindex = min(index+1, size-1);
+ float t = x - index;
+
+ float data0 = kernel_tex_fetch(__lookup_table, index + offset);
+ if(t == 0.0f)
+ return data0;
+
+ float data1 = kernel_tex_fetch(__lookup_table, nindex + offset);
+ return (1.0f - t)*data0 + t*data1;
+}
+
+ccl_device float lookup_table_read_2D(KernelGlobals *kg, float x, float y, int offset, int xsize, int ysize)
+{
+ y = saturate(y)*(ysize-1);
+
+ int index = min(float_to_int(y), ysize-1);
+ int nindex = min(index+1, ysize-1);
+ float t = y - index;
+
+ float data0 = lookup_table_read(kg, x, offset + xsize*index, xsize);
+ if(t == 0.0f)
+ return data0;
+
+ float data1 = lookup_table_read(kg, x, offset + xsize*nindex, xsize);
+ return (1.0f - t)*data0 + t*data1;
+}
CCL_NAMESPACE_END
projects / blender-staging.git / blobdiff