/*! [config] name: LuxRays slg tests build_options: -Werror clc_version_min: 10 [test] name: Test LuxRays slg4 InitFrameBuffer kernel_name: InitFrameBuffer dimensions: 1 global_size: 309444 0 0 arg_out: 0 buffer float[1237776] repeat 5.0 !*/ //global size is (WIDTH+2) * (HEIGHT+2) = 642 * 482 = 309444 //Framebuffer SIze is global_size * sizeof(Pixel) = global_size * 16 = 4951104 bytes = 1237776 floats #define PARAM_IMAGE_WIDTH 640 #define PARAM_IMAGE_HEIGHT 480 typedef struct { float r, g, b; } Spectrum; typedef struct { Spectrum c; float count; } Pixel; void VSTORE4F(const float4 v, __global float *p) { vstore4(v, 0, p); } __kernel __attribute__((work_group_size_hint(64, 1, 1))) void InitFrameBuffer( __global Pixel *frameBuffer ) { const size_t gid = get_global_id(0); if (gid >= (PARAM_IMAGE_WIDTH + 2) * (PARAM_IMAGE_HEIGHT + 2)) return; VSTORE4F(5.0f, &frameBuffer[gid].c.r); } #if 0 #define LUXRAYS_OPENCL_KERNEL #define SLG_OPENCL_KERNEL #define PARAM_TASK_COUNT 65536 #define PARAM_RAY_EPSILON_MIN 1.000000e-09f #define PARAM_RAY_EPSILON_MAX 1.000000e-01f #define PARAM_MAX_PATH_DEPTH 16 #define PARAM_RR_DEPTH 5 #define PARAM_RR_CAP 5.000000e-01f #define PARAM_ACCEL_QBVH #define PARAM_HAS_NORMALS_BUFFER #define PARAM_HAS_UVS_BUFFER #define PARAM_HAS_COLS_BUFFER #define PARAM_ENABLE_TEX_CONST_FLOAT #define PARAM_ENABLE_TEX_CONST_FLOAT3 #define PARAM_ENABLE_MAT_MATTE #define PARAM_ENABLE_MAT_GLASS #define PARAM_HAS_PASSTHROUGH #define PARAM_DIRECT_LIGHT_SAMPLING #define PARAM_DL_LIGHT_COUNT 4 #define PARAM_IMAGE_FILTER_TYPE 0 #define PARAM_SAMPLER_TYPE 0 #define PARAM_DEVICE_INDEX 0 #define PARAM_DEVICE_COUNT 1 #define PARAM_TONEMAP_LINEAR_SCALE 1.000000e+00f #define PARAM_GAMMA 2.200000e+00f #line 2 "luxrays_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ #define NULL_INDEX (0xffffffffu) #if defined(LUXRAYS_OPENCL_KERNEL) #if defined(__APPLE_CL__) float3 __OVERLOAD__ mix(float3 a, float3 b, float t) { return a + ( b - a ) * t; } #endif #if defined(__APPLE_FIX__) float2 VLOAD2F(const __global float *p) { return (float2)(p[0], p[1]); } void VSTORE2F(const float2 v, __global float *p) { p[0] = v.x; p[1] = v.y; } float3 VLOAD3F(const __global float *p) { return (float3)(p[0], p[1], p[2]); } void VSTORE3F(const float3 v, __global float *p) { p[0] = v.x; p[1] = v.y; p[2] = v.z; } float4 VLOAD4F(const __global float *p) { return (float4)(p[0], p[1], p[2], p[3]); } void VSTORE4F(const float4 v, __global float *p) { p[0] = v.x; p[1] = v.y; p[2] = v.z; p[3] = v.w; } #else float2 VLOAD2F(const __global float *p) { return vload2(0, p); } void VSTORE2F(const float2 v, __global float *p) { vstore2(v, 0, p); } float3 VLOAD3F(const __global float *p) { return vload3(0, p); } void VSTORE3F(const float3 v, __global float *p) { vstore3(v, 0, p); } float4 VLOAD4F(const __global float *p) { return vload4(0, p); } #endif #endif #line 2 "uv_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ #define ASSIGN_UV(a, b) { (a).u = (b).u; (a).v = (b).v; } typedef struct { float u,v; } UV; #line 2 "point_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ typedef struct { float x, y, z; } Point; #line 2 "vector_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ #define ASSIGN_VECTOR(a, b) { (a).x = (b).x; (a).y = (b).y; (a).z = (b).z; } typedef struct { float x, y, z; } Vector; #line 2 "normal_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ typedef struct { float x, y, z; } Normal; #line 2 "triangle_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ typedef struct { unsigned int v[3]; } Triangle; #line 2 "ray_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ typedef struct { Point o; Vector d; float mint, maxt, time; float pad[3]; } Ray; typedef struct { float t, b1, b2; unsigned int meshIndex, triangleIndex; } RayHit; #line 2 "epsilon_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ // NOTE: DEFAULT_EPSILON_MIN is very small. A plane passing exactly for the // origin will suffer of self shadow problems because the Ray class will use // MachineEpsilon(ray.o) as epsilon for the ray.mint. However it is pretty much // the only case where there is a problem so better to not change anything. // As workaround, moving the plane away from the origin is enough. #define DEFAULT_EPSILON_MIN 1e-9f #define DEFAULT_EPSILON_MAX 1e-1f #define DEFAULT_EPSILON_STATIC 1e-5f // An epsilon that can be used as threshold for cos(theta). For instance: // if (Dot(N, LightDir) < DEFAULT_COS_EPSILON_STATIC) return Spectrum(); #define DEFAULT_COS_EPSILON_STATIC 1e-4f // This is about 1e-5f for values near 1.f #define DEFAULT_EPSILON_DISTANCE_FROM_VALUE 0x80u #line 2 "spectrum_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ #if defined(SLG_OPENCL_KERNEL) #define BLACK ((float3)(0.f, 0.f, 0.f)) #define WHITE ((float3)(1.f, 1.f, 1.f)) #endif #define ASSIGN_SPECTRUM(c0, c1) { (c0).r = (c1).r; (c0).g = (c1).g; (c0).b = (c1).b; } #line 2 "frame_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ typedef struct { Vector X, Y, Z; } Frame; #line 2 "matrix4x4_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ typedef struct { float m[4][4]; } Matrix4x4; #line 2 "ray_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ typedef struct { Matrix4x4 m, mInv; } Transform; #line 2 "randomgen_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ typedef struct { unsigned int s1, s2, s3; } Seed; #line 2 "trianglemesh_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ typedef struct { // Vertex informaiton unsigned int vertsOffset; unsigned int normalsOffset; unsigned int uvsOffset; unsigned int colsOffset; unsigned int alphasOffset; // Triangle information unsigned int trisOffset; // Information used by MQBVH Transform trans; } Mesh; #line 2 "hitpoint_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ // This is defined only under OpenCL because of variable size structures #if defined(SLG_OPENCL_KERNEL) typedef struct { // The incoming direction. It is the eyeDir when fromLight = false and // lightDir when fromLight = true Vector fixedDir; Point p; UV uv; Normal geometryN; Normal shadeN; #if defined(PARAM_ENABLE_TEX_HITPOINTCOLOR) || defined(PARAM_ENABLE_TEX_HITPOINTGREY) Spectrum color; #endif #if defined(PARAM_ENABLE_TEX_HITPOINTALPHA) float alpha; #endif #if defined(PARAM_HAS_PASSTHROUGH) // passThroughEvent can be stored here in a path state even before of // BSDF initialization (while tracing the next path vertex ray) float passThroughEvent; #endif } HitPoint; #endif #line 2 "mapping_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ //------------------------------------------------------------------------------ // TextureMapping2D //------------------------------------------------------------------------------ typedef enum { UVMAPPING2D } TextureMapping2DType; typedef struct { float uScale, vScale, uDelta, vDelta; } UVMappingParam; typedef struct { TextureMapping2DType type; union { UVMappingParam uvMapping2D; }; } TextureMapping2D; //------------------------------------------------------------------------------ // TextureMapping3D //------------------------------------------------------------------------------ typedef enum { UVMAPPING3D, GLOBALMAPPING3D } TextureMapping3DType; typedef struct { TextureMapping3DType type; Transform worldToLocal; //union { // UVMapping3D has no parameters // GlobalMapping3D has no parameters //}; } TextureMapping3D; #line 2 "texture_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ #define DUDV_VALUE 0.001f typedef enum { CONST_FLOAT, CONST_FLOAT3, IMAGEMAP, SCALE_TEX, FRESNEL_APPROX_N, FRESNEL_APPROX_K, MIX_TEX, ADD_TEX, HITPOINTCOLOR, HITPOINTALPHA, HITPOINTGREY, // Procedural textures CHECKERBOARD2D, CHECKERBOARD3D, FBM_TEX, MARBLE, DOTS, BRICK, WINDY, WRINKLED, UV_TEX, BAND_TEX } TextureType; typedef struct { float value; } ConstFloatParam; typedef struct { Spectrum color; } ConstFloat3Param; typedef struct { unsigned int channelCount, width, height; unsigned int pageIndex, pixelsIndex; } ImageMap; typedef struct { TextureMapping2D mapping; float gain, Du, Dv; unsigned int imageMapIndex; } ImageMapTexParam; typedef struct { unsigned int tex1Index, tex2Index; } ScaleTexParam; typedef struct { unsigned int texIndex; } FresnelApproxNTexParam; typedef struct { unsigned int texIndex; } FresnelApproxKTexParam; typedef struct { TextureMapping2D mapping; unsigned int tex1Index, tex2Index; } CheckerBoard2DTexParam; typedef struct { TextureMapping3D mapping; unsigned int tex1Index, tex2Index; } CheckerBoard3DTexParam; typedef struct { unsigned int amountTexIndex, tex1Index, tex2Index; } MixTexParam; typedef struct { TextureMapping3D mapping; int octaves; float omega; } FBMTexParam; typedef struct { TextureMapping3D mapping; int octaves; float omega, scale, variation; } MarbleTexParam; typedef struct { TextureMapping2D mapping; unsigned int insideIndex, outsideIndex; } DotsTexParam; typedef enum { FLEMISH, RUNNING, ENGLISH, HERRINGBONE, BASKET, KETTING } MasonryBond; typedef struct { TextureMapping3D mapping; unsigned int tex1Index, tex2Index, tex3Index; MasonryBond bond; float offsetx, offsety, offsetz; float brickwidth, brickheight, brickdepth, mortarsize; float proportion, invproportion, run; float mortarwidth, mortarheight, mortardepth; float bevelwidth, bevelheight, beveldepth; int usebevel; } BrickTexParam; typedef struct { unsigned int tex1Index, tex2Index; } AddTexParam; typedef struct { TextureMapping3D mapping; int octaves; float omega; } WindyTexParam; typedef struct { TextureMapping3D mapping; int octaves; float omega; } WrinkledTexParam; typedef struct { TextureMapping2D mapping; } UVTexParam; #define BAND_TEX_MAX_SIZE 16 typedef struct { unsigned int amountTexIndex; unsigned int size; float offsets[BAND_TEX_MAX_SIZE]; Spectrum values[BAND_TEX_MAX_SIZE]; } BandTexParam; typedef struct { unsigned int channel; } HitPointGreyTexParam; typedef struct { TextureType type; union { ConstFloatParam constFloat; ConstFloat3Param constFloat3; ImageMapTexParam imageMapTex; ScaleTexParam scaleTex; FresnelApproxNTexParam fresnelApproxN; FresnelApproxKTexParam fresnelApproxK; CheckerBoard2DTexParam checkerBoard2D; CheckerBoard3DTexParam checkerBoard3D; MixTexParam mixTex; FBMTexParam fbm; MarbleTexParam marble; DotsTexParam dots; BrickTexParam brick; AddTexParam addTex; WindyTexParam windy; WrinkledTexParam wrinkled; UVTexParam uvTex; BandTexParam band; HitPointGreyTexParam hitPointGrey; }; } Texture; //------------------------------------------------------------------------------ // Some macro trick in order to have more readable code //------------------------------------------------------------------------------ #if defined(PARAM_HAS_IMAGEMAPS) #define IMAGEMAPS_PARAM_DECL , __global ImageMap *imageMapDescs, __global float **imageMapBuff #define IMAGEMAPS_PARAM , imageMapDescs, imageMapBuff #else #define IMAGEMAPS_PARAM_DECL #define IMAGEMAPS_PARAM #endif #define TEXTURES_PARAM_DECL , __global Texture *texs IMAGEMAPS_PARAM_DECL #define TEXTURES_PARAM , texs IMAGEMAPS_PARAM #line 2 "material_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ typedef enum { MATTE, MIRROR, GLASS, METAL, ARCHGLASS, MIX, NULLMAT, MATTETRANSLUCENT, GLOSSY2, METAL2, ROUGHGLASS } MaterialType; typedef struct { unsigned int kdTexIndex; } MatteParam; typedef struct { unsigned int krTexIndex; } MirrorParam; typedef struct { unsigned int krTexIndex; unsigned int ktTexIndex; unsigned int ousideIorTexIndex, iorTexIndex; } GlassParam; typedef struct { unsigned int krTexIndex; unsigned int expTexIndex; } MetalParam; typedef struct { unsigned int krTexIndex; unsigned int ktTexIndex; unsigned int ousideIorTexIndex, iorTexIndex; } ArchGlassParam; typedef struct { unsigned int matAIndex, matBIndex; unsigned int mixFactorTexIndex; } MixParam; typedef struct { unsigned int krTexIndex; unsigned int ktTexIndex; } MatteTranslucentParam; typedef struct { unsigned int kdTexIndex; unsigned int ksTexIndex; unsigned int nuTexIndex; unsigned int nvTexIndex; unsigned int kaTexIndex; unsigned int depthTexIndex; unsigned int indexTexIndex; int multibounce; } Glossy2Param; typedef struct { unsigned int nTexIndex; unsigned int kTexIndex; unsigned int nuTexIndex; unsigned int nvTexIndex; } Metal2Param; typedef struct { unsigned int krTexIndex; unsigned int ktTexIndex; unsigned int ousideIorTexIndex, iorTexIndex; unsigned int nuTexIndex; unsigned int nvTexIndex; } RoughGlassParam; typedef struct { MaterialType type; unsigned int emitTexIndex, bumpTexIndex, normalTexIndex; union { MatteParam matte; MirrorParam mirror; GlassParam glass; MetalParam metal; ArchGlassParam archglass; MixParam mix; // NULLMAT has no parameters MatteTranslucentParam matteTranslucent; Glossy2Param glossy2; Metal2Param metal2; RoughGlassParam roughglass; }; } Material; //------------------------------------------------------------------------------ // Some macro trick in order to have more readable code //------------------------------------------------------------------------------ #define MATERIALS_PARAM_DECL , __global Material *mats TEXTURES_PARAM_DECL #define MATERIALS_PARAM , mats TEXTURES_PARAM #line 2 "bsdf_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ typedef enum { NONE = 0, DIFFUSE = 1, GLOSSY = 2, SPECULAR = 4, REFLECT = 8, TRANSMIT = 16 } BSDFEventType; typedef int BSDFEvent; // This is defined only under OpenCL because of variable size structures #if defined(SLG_OPENCL_KERNEL) typedef struct { HitPoint hitPoint; unsigned int materialIndex; #if (PARAM_DL_LIGHT_COUNT > 0) unsigned int triangleLightSourceIndex; #endif Frame frame; } BSDF; #endif #line 2 "sampler_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ //------------------------------------------------------------------------------ // Indices of Sample related u[] array //------------------------------------------------------------------------------ #if defined(SLG_OPENCL_KERNEL) #define IDX_SCREEN_X 0 #define IDX_SCREEN_Y 1 #if defined(PARAM_CAMERA_HAS_DOF) && defined(PARAM_HAS_PASSTHROUGH) #define IDX_EYE_PASSTHROUGH 2 #define IDX_DOF_X 3 #define IDX_DOF_Y 4 #define IDX_BSDF_OFFSET 5 #elif defined(PARAM_CAMERA_HAS_DOF) #define IDX_DOF_X 2 #define IDX_DOF_Y 3 #define IDX_BSDF_OFFSET 4 #elif defined(PARAM_HAS_PASSTHROUGH) #define IDX_EYE_PASSTHROUGH 2 #define IDX_BSDF_OFFSET 3 #else #define IDX_BSDF_OFFSET 2 #endif // Relative to IDX_BSDF_OFFSET + PathDepth * VERTEX_SAMPLE_SIZE #if defined(PARAM_DIRECT_LIGHT_SAMPLING) && defined(PARAM_HAS_PASSTHROUGH) #define IDX_PASSTHROUGH 0 #define IDX_BSDF_X 1 #define IDX_BSDF_Y 2 #define IDX_DIRECTLIGHT_X 3 #define IDX_DIRECTLIGHT_Y 4 #define IDX_DIRECTLIGHT_Z 5 #define IDX_DIRECTLIGHT_W 6 #define IDX_DIRECTLIGHT_A 7 #define IDX_RR 8 #define VERTEX_SAMPLE_SIZE 9 #elif defined(PARAM_DIRECT_LIGHT_SAMPLING) #define IDX_BSDF_X 0 #define IDX_BSDF_Y 1 #define IDX_DIRECTLIGHT_X 2 #define IDX_DIRECTLIGHT_Y 3 #define IDX_DIRECTLIGHT_Z 4 #define IDX_DIRECTLIGHT_W 5 #define IDX_RR 6 #define VERTEX_SAMPLE_SIZE 7 #elif defined(PARAM_HAS_PASSTHROUGH) #define IDX_PASSTHROUGH 0 #define IDX_BSDF_X 1 #define IDX_BSDF_Y 2 #define IDX_RR 3 #define VERTEX_SAMPLE_SIZE 4 #else #define IDX_BSDF_X 0 #define IDX_BSDF_Y 1 #define IDX_RR 2 #define VERTEX_SAMPLE_SIZE 3 #endif #if (PARAM_SAMPLER_TYPE == 0) || (PARAM_SAMPLER_TYPE == 2) #define TOTAL_U_SIZE 2 #endif #if (PARAM_SAMPLER_TYPE == 1) #define TOTAL_U_SIZE (IDX_BSDF_OFFSET + PARAM_MAX_PATH_DEPTH * VERTEX_SAMPLE_SIZE) #endif #endif //------------------------------------------------------------------------------ // Sample data types //------------------------------------------------------------------------------ // This is defined only under OpenCL because of variable size structures #if defined(SLG_OPENCL_KERNEL) typedef struct { Spectrum radiance; #if defined(PARAM_ENABLE_ALPHA_CHANNEL) float alpha; #endif } RandomSample; typedef struct { Spectrum radiance; float alpha; float totalI; // Using ushort here totally freeze the ATI driver unsigned int largeMutationCount, smallMutationCount; unsigned int current, proposed, consecutiveRejects; float weight; Spectrum currentRadiance; #if defined(PARAM_ENABLE_ALPHA_CHANNEL) float currentAlpha; #endif } MetropolisSample; typedef struct { float rng0, rng1; unsigned int pixelIndex, pass; Spectrum radiance; #if defined(PARAM_ENABLE_ALPHA_CHANNEL) float alpha; #endif } SobolSample; #if (PARAM_SAMPLER_TYPE == 0) typedef RandomSample Sample; #endif #if (PARAM_SAMPLER_TYPE == 1) typedef MetropolisSample Sample; #endif #if (PARAM_SAMPLER_TYPE == 2) typedef SobolSample Sample; #endif #endif //------------------------------------------------------------------------------ // Sampler data types //------------------------------------------------------------------------------ typedef enum { RANDOM = 0, METROPOLIS = 1, SOBOL = 2 } SamplerType; typedef struct { SamplerType type; union { struct { float largeMutationProbability, imageMutationRange; unsigned int maxRejects; } metropolis; }; } Sampler; #define SOBOL_BITS 32 #define SOBOL_MAX_DIMENSIONS 21201 #line 2 "filter_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ //------------------------------------------------------------------------------ // Frame buffer data types //------------------------------------------------------------------------------ typedef struct { float alpha; } AlphaPixel; //------------------------------------------------------------------------------ // Filter data types //------------------------------------------------------------------------------ typedef enum { FILTER_NONE, FILTER_BOX, FILTER_GAUSSIAN, FILTER_MITCHELL } FilterType; typedef struct { FilterType type; union { struct { float widthX, widthY; } box; struct { float widthX, widthY; float alpha; } gaussian; struct { float widthX, widthY; float B, C; } mitchell; }; } Filter; #line 2 "camera_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ typedef struct { Transform rasterToCamera[2]; // 2 used for stereo rendering Transform cameraToWorld[2]; // 2 used for stereo rendering // Placed here for Transform memory alignement float lensRadius; float focalDistance; float yon, hither; } Camera; #line 2 "light_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ typedef enum { TYPE_IL, TYPE_IL_SKY, TYPE_SUN, TYPE_TRIANGLE } LightSourceType; typedef struct { Transform light2World; Spectrum gain; TextureMapping2D mapping; unsigned int imageMapIndex; } InfiniteLight; typedef struct { Transform light2World; Spectrum gain; float thetaS; float phiS; float zenith_Y, zenith_x, zenith_y; float perez_Y[6], perez_x[6], perez_y[6]; } SkyLight; typedef struct { Vector sunDir; Spectrum gain; float turbidity; float relSize; // XY Vectors for cone sampling Vector x, y; float cosThetaMax; Spectrum sunColor; } SunLight; typedef struct { Vector v0, v1, v2; UV uv0, uv1, uv2; float invArea; unsigned int materialIndex; } TriangleLight; #line 2 "epsilon_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ float MachineEpsilon_FloatAdvance(const float value) { return as_float(as_uint(value) + DEFAULT_EPSILON_DISTANCE_FROM_VALUE); } float MachineEpsilon_E(const float value) { const float epsilon = fabs(MachineEpsilon_FloatAdvance(value) - value); return clamp(epsilon, PARAM_RAY_EPSILON_MIN, PARAM_RAY_EPSILON_MAX); } float MachineEpsilon_E_Float3(const float3 v) { return fmax(MachineEpsilon_E(v.x), fmax(MachineEpsilon_E(v.y), MachineEpsilon_E(v.z))); } #line 2 "utils_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ int Mod(int a, int b) { if (b == 0) b = 1; a %= b; if (a < 0) a += b; return a; } int Floor2Int(const float val) { return (int)floor(val); } float Lerp(const float t, const float v1, const float v2) { return mix(v1, v2, t); } float3 Lerp3(const float t, const float3 v1, const float3 v2) { return mix(v1, v2, t); } float SmoothStep(const float min, const float max, const float value) { const float v = clamp((value - min) / (max - min), 0.f, 1.f); return v * v * (-2.f * v + 3.f); } float CosTheta(const float3 v) { return v.z; } float SinTheta2(const float3 w) { return fmax(0.f, 1.f - CosTheta(w) * CosTheta(w)); } #line 2 "vector_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ float SphericalTheta(const float3 v) { return acos(clamp(v.z, -1.f, 1.f)); } float SphericalPhi(const float3 v) { float p = atan2(v.y, v.x); return (p < 0.f) ? p + 2.f * M_PI_F : p; } void CoordinateSystem(const float3 v1, float3 *v2, float3 *v3) { if (fabs(v1.x) > fabs(v1.y)) { float invLen = 1.f / sqrt(v1.x * v1.x + v1.z * v1.z); (*v2).x = -v1.z * invLen; (*v2).y = 0.f; (*v2).z = v1.x * invLen; } else { float invLen = 1.f / sqrt(v1.y * v1.y + v1.z * v1.z); (*v2).x = 0.f; (*v2).y = v1.z * invLen; (*v2).z = -v1.y * invLen; } *v3 = cross(v1, *v2); } #line 2 "ray_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ void Ray_Init4(__global Ray *ray, const float3 orig, const float3 dir, const float mint, const float maxt) { VSTORE3F(orig, &ray->o.x); VSTORE3F(dir, &ray->d.x); ray->mint = mint; ray->maxt = maxt; } void Ray_Init3(__global Ray *ray, const float3 orig, const float3 dir, const float maxt) { VSTORE3F(orig, &ray->o.x); VSTORE3F(dir, &ray->d.x); ray->mint = MachineEpsilon_E_Float3(orig); ray->maxt = maxt; } void Ray_Init2(__global Ray *ray, const float3 orig, const float3 dir) { VSTORE3F(orig, &ray->o.x); VSTORE3F(dir, &ray->d.x); ray->mint = MachineEpsilon_E_Float3(orig); ray->maxt = INFINITY; } void Ray_ReadAligned4(__global Ray *ray, float3 *rayOrig, float3 *rayDir, float *mint, float *maxt) { __global float4 *basePtr =(__global float4 *)ray; const float4 data0 = (*basePtr++); const float4 data1 = (*basePtr); *rayOrig = (float3)(data0.x, data0.y, data0.z); *rayDir = (float3)(data0.w, data1.x, data1.y); *mint = data1.z; *maxt = data1.w; } #line 2 "specturm_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ bool Spectrum_IsEqual(const float3 a, const float3 b) { return all(isequal(a, b)); } bool Spectrum_IsBlack(const float3 a) { return Spectrum_IsEqual(a, BLACK); } float Spectrum_Filter(const float3 s) { return fmax(s.s0, fmax(s.s1, s.s2)); } float Spectrum_Y(const float3 s) { return 0.212671f * s.s0 + 0.715160f * s.s1 + 0.072169f * s.s2; } float3 Spectrum_Clamp(const float3 s) { return clamp(s, BLACK, WHITE); } float3 Spectrum_Exp(const float3 s) { return (float3)(exp(s.x), exp(s.y), exp(s.z)); } float3 Spectrum_Sqrt(const float3 s) { return (float3)(sqrt(s.x), sqrt(s.y), sqrt(s.z)); } #line 2 "matrix4x4_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ float3 Matrix4x4_ApplyPoint(__global Matrix4x4 *m, const float3 point) { const float4 point4 = (float4)(point.x, point.y, point.z, 1.f); const float4 row3 = VLOAD4F(&m->m[3][0]); const float iw = 1.f / dot(row3, point4); const float4 row0 = VLOAD4F(&m->m[0][0]); const float4 row1 = VLOAD4F(&m->m[1][0]); const float4 row2 = VLOAD4F(&m->m[2][0]); return (float3)( iw * dot(row0, point4), iw * dot(row1, point4), iw * dot(row2, point4) ); } float3 Matrix4x4_ApplyVector(__global Matrix4x4 *m, const float3 vector) { const float3 row0 = VLOAD3F(&m->m[0][0]); const float3 row1 = VLOAD3F(&m->m[1][0]); const float3 row2 = VLOAD3F(&m->m[2][0]); return (float3)( dot(row0, vector), dot(row1, vector), dot(row2, vector) ); } float3 Matrix4x4_ApplyNormal(__global Matrix4x4 *m, const float3 normal) { const float3 row0 = (float3)(m->m[0][0], m->m[1][0], m->m[2][0]); const float3 row1 = (float3)(m->m[0][1], m->m[1][1], m->m[2][1]); const float3 row2 = (float3)(m->m[0][2], m->m[1][2], m->m[2][2]); return (float3)( dot(row0, normal), dot(row1, normal), dot(row2, normal) ); } #line 2 "mc_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ void ConcentricSampleDisk(const float u0, const float u1, float *dx, float *dy) { float r, theta; // Map uniform random numbers to $[-1,1]^2$ float sx = 2.f * u0 - 1.f; float sy = 2.f * u1 - 1.f; // Map square to $(r,\theta)$ // Handle degeneracy at the origin if (sx == 0.f && sy == 0.f) { *dx = 0.f; *dy = 0.f; return; } if (sx >= -sy) { if (sx > sy) { // Handle first region of disk r = sx; if (sy > 0.f) theta = sy / r; else theta = 8.f + sy / r; } else { // Handle second region of disk r = sy; theta = 2.f - sx / r; } } else { if (sx <= sy) { // Handle third region of disk r = -sx; theta = 4.f - sy / r; } else { // Handle fourth region of disk r = -sy; theta = 6.f + sx / r; } } theta *= M_PI_F / 4.f; *dx = r * cos(theta); *dy = r * sin(theta); } float3 CosineSampleHemisphere(const float u0, const float u1) { float x, y; ConcentricSampleDisk(u0, u1, &x, &y); const float z = sqrt(fmax(0.f, 1.f - x * x - y * y)); return (float3)(x, y, z); } float3 CosineSampleHemisphereWithPdf(const float u0, const float u1, float *pdfW) { float x, y; ConcentricSampleDisk(u0, u1, &x, &y); const float z = sqrt(fmax(0.f, 1.f - x * x - y * y)); *pdfW = z * M_1_PI_F; return (float3)(x, y, z); } float3 UniformSampleCone(const float u0, const float u1, const float costhetamax, const float3 x, const float3 y, const float3 z) { const float costheta = mix(costhetamax, 1.f, u0); const float sintheta = sqrt(1.f - costheta * costheta); const float phi = u1 * 2.f * M_PI_F; const float kx = cos(phi) * sintheta; const float ky = sin(phi) * sintheta; const float kz = costheta; return (float3)(kx * x.x + ky * y.x + kz * z.x, kx * x.y + ky * y.y + kz * z.y, kx * x.z + ky * y.z + kz * z.z); } float UniformConePdf(const float costhetamax) { return 1.f / (2.f * M_PI_F * (1.f - costhetamax)); } float PowerHeuristic(const float fPdf, const float gPdf) { const float f = fPdf; const float g = gPdf; return (f * f) / (f * f + g * g); } float PdfWtoA(const float pdfW, const float dist, const float cosThere) { return pdfW * fabs(cosThere) / (dist * dist); } float PdfAtoW(const float pdfA, const float dist, const float cosThere) { return pdfA * dist * dist / fabs(cosThere); } #line 2 "frame_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ void Frame_SetFromZ(__global Frame *frame, const float3 Z) { float3 X, Y; CoordinateSystem(Z, &X, &Y); VSTORE3F(X, &frame->X.x); VSTORE3F(Y, &frame->Y.x); VSTORE3F(Z, &frame->Z.x); } float3 ToWorld(const float3 X, const float3 Y, const float3 Z, const float3 v) { return X * v.x + Y * v.y + Z * v.z; } float3 Frame_ToWorld(__global Frame *frame, const float3 v) { return ToWorld(VLOAD3F(&frame->X.x), VLOAD3F(&frame->Y.x), VLOAD3F(&frame->Z.x), v); } float3 ToLocal(const float3 X, const float3 Y, const float3 Z, const float3 a) { return (float3)(dot(a, X), dot(a, Y), dot(a, Z)); } float3 Frame_ToLocal(__global Frame *frame, const float3 v) { return ToLocal(VLOAD3F(&frame->X.x), VLOAD3F(&frame->Y.x), VLOAD3F(&frame->Z.x), v); } #line 2 "transform_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ float3 Transform_ApplyPoint(__global Transform *trans, const float3 point) { return Matrix4x4_ApplyPoint(&trans->m, point); } float3 Transform_ApplyVector(__global Transform *trans, const float3 vector) { return Matrix4x4_ApplyVector(&trans->m, vector); } float3 Transform_ApplyNormal(__global Transform *trans, const float3 normal) { return Matrix4x4_ApplyNormal(&trans->m, normal); } float3 Transform_InvApplyPoint(__global Transform *trans, const float3 point) { return Matrix4x4_ApplyPoint(&trans->mInv, point); } float3 Transform_InvApplyVector(__global Transform *trans, const float3 vector) { return Matrix4x4_ApplyVector(&trans->mInv, vector); } float3 Transform_InvApplyNormal(__global Transform *trans, const float3 normal) { return Matrix4x4_ApplyNormal(&trans->mInv, normal); } #line 2 "randomgen_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ //------------------------------------------------------------------------------ // Random number generator // maximally equidistributed combined Tausworthe generator //------------------------------------------------------------------------------ #define FLOATMASK 0x00ffffffu uint TAUSWORTHE(const uint s, const uint a, const uint b, const uint c, const uint d) { return ((s&c)<> b); } uint LCG(const uint x) { return x * 69069; } uint ValidSeed(const uint x, const uint m) { return (x < m) ? (x + m) : x; } void Rnd_Init(uint seed, Seed *s) { // Avoid 0 value seed = (seed == 0) ? (seed + 0xffffffu) : seed; s->s1 = ValidSeed(LCG(seed), 1); s->s2 = ValidSeed(LCG(s->s1), 7); s->s3 = ValidSeed(LCG(s->s2), 15); } unsigned long Rnd_UintValue(Seed *s) { s->s1 = TAUSWORTHE(s->s1, 13, 19, 4294967294UL, 12); s->s2 = TAUSWORTHE(s->s2, 2, 25, 4294967288UL, 4); s->s3 = TAUSWORTHE(s->s3, 3, 11, 4294967280UL, 17); return ((s->s1) ^ (s->s2) ^ (s->s3)); } float Rnd_FloatValue(Seed *s) { return (Rnd_UintValue(s) & FLOATMASK) * (1.f / (FLOATMASK + 1UL)); } #line 2 "triangle_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ void Triangle_UniformSample(const float u0, const float u1, float *b1, float *b2) { const float su1 = sqrt(u0); *b1 = 1.f - su1; *b2 = u1 * su1; } float3 Triangle_Sample(const float3 p0, const float3 p1, const float3 p2, const float u0, const float u1, float *b0, float *b1, float *b2) { Triangle_UniformSample(u0, u1, b1, b2); *b0 = 1.f - (*b1) - (*b2); return (*b0) * p0 + (*b1) * p1 + (*b2) * p2; } float3 Triangle_GetGeometryNormal(const float3 p0, const float3 p1, const float3 p2) { return normalize(cross(p1 - p0, p2 - p0)); } float3 Triangle_InterpolateNormal(const float3 n0, const float3 n1, const float3 n2, const float b0, const float b1, const float b2) { return normalize(b0 * n0 + b1 * n1 + b2 * n2); } float2 Triangle_InterpolateUV(const float2 uv0, const float2 uv1, const float2 uv2, const float b0, const float b1, const float b2) { return b0 * uv0 + b1 * uv1 + b2 * uv2; } float3 Triangle_InterpolateColor(const float3 rgb0, const float3 rgb1, const float3 rgb2, const float b0, const float b1, const float b2) { return b0 * rgb0 + b1 * rgb1 + b2 * rgb2; } float Triangle_InterpolateAlpha(const float a0, const float a1, const float a2, const float b0, const float b1, const float b2) { return b0 * a0 + b1 * a1 + b2 * a2; } void Triangle_Intersect( const float3 rayOrig, const float3 rayDir, const float mint, float *maxt, uint *hitMeshIndex, uint *hitTriangleIndex, float *hitB1, float *hitB2, const uint currentMeshIndex, const uint currentTriangleIndex, const float3 v0, const float3 v1, const float3 v2) { // Calculate intersection const float3 e1 = v1 - v0; const float3 e2 = v2 - v0; const float3 s1 = cross(rayDir, e2); const float divisor = dot(s1, e1); if (divisor == 0.f) return; const float invDivisor = 1.f / divisor; // Compute first barycentric coordinate const float3 d = rayOrig - v0; const float b1 = dot(d, s1) * invDivisor; if (b1 < 0.f) return; // Compute second barycentric coordinate const float3 s2 = cross(d, e1); const float b2 = dot(rayDir, s2) * invDivisor; if (b2 < 0.f) return; const float b0 = 1.f - b1 - b2; if (b0 < 0.f) return; // Compute _t_ to intersection point const float t = dot(e2, s2) * invDivisor; if (t < mint || t > *maxt) return; *maxt = t; *hitB1 = b1; *hitB2 = b2; *hitMeshIndex = currentMeshIndex; *hitTriangleIndex = currentTriangleIndex; } #line 2 "trianglemesh_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ float3 Mesh_GetGeometryNormal(__global Point *vertices, __global Triangle *triangles, const uint triIndex) { __global Triangle *tri = &triangles[triIndex]; const float3 p0 = VLOAD3F(&vertices[tri->v[0]].x); const float3 p1 = VLOAD3F(&vertices[tri->v[1]].x); const float3 p2 = VLOAD3F(&vertices[tri->v[2]].x); return Triangle_GetGeometryNormal(p0, p1, p2); } float3 Mesh_InterpolateNormal(__global Vector *normals, __global Triangle *triangles, const uint triIndex, const float b1, const float b2) { __global Triangle *tri = &triangles[triIndex]; const float3 n0 = VLOAD3F(&normals[tri->v[0]].x); const float3 n1 = VLOAD3F(&normals[tri->v[1]].x); const float3 n2 = VLOAD3F(&normals[tri->v[2]].x); const float b0 = 1.f - b1 - b2; return Triangle_InterpolateNormal(n0, n1, n2, b0, b1, b2); } float2 Mesh_InterpolateUV(__global UV *vertUVs, __global Triangle *triangles, const uint triIndex, const float b1, const float b2) { __global Triangle *tri = &triangles[triIndex]; const float2 uv0 = VLOAD2F(&vertUVs[tri->v[0]].u); const float2 uv1 = VLOAD2F(&vertUVs[tri->v[1]].u); const float2 uv2 = VLOAD2F(&vertUVs[tri->v[2]].u); const float b0 = 1.f - b1 - b2; return Triangle_InterpolateUV(uv0, uv1, uv2, b0, b1, b2); } float3 Mesh_InterpolateColor(__global Spectrum *vertCols, __global Triangle *triangles, const uint triIndex, const float b1, const float b2) { __global Triangle *tri = &triangles[triIndex]; const float3 rgb0 = VLOAD3F(&vertCols[tri->v[0]].r); const float3 rgb1 = VLOAD3F(&vertCols[tri->v[1]].r); const float3 rgb2 = VLOAD3F(&vertCols[tri->v[2]].r); const float b0 = 1.f - b1 - b2; return Triangle_InterpolateColor(rgb0, rgb1, rgb2, b0, b1, b2); } float Mesh_InterpolateAlpha(__global float *vertAlphas, __global Triangle *triangles, const uint triIndex, const float b1, const float b2) { __global Triangle *tri = &triangles[triIndex]; const float a0 = vertAlphas[tri->v[0]]; const float a1 = vertAlphas[tri->v[1]]; const float a2 = vertAlphas[tri->v[2]]; const float b0 = 1.f - b1 - b2; return Triangle_InterpolateAlpha(a0, a1, a2, b0, b1, b2); } #line 2 "mapping_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ float2 UVMapping2D_Map(__global TextureMapping2D *mapping, __global HitPoint *hitPoint) { const float2 scale = VLOAD2F(&mapping->uvMapping2D.uScale); const float2 delta = VLOAD2F(&mapping->uvMapping2D.uDelta); const float2 uv = VLOAD2F(&hitPoint->uv.u); return uv * scale + delta; } float3 UVMapping3D_Map(__global TextureMapping3D *mapping, __global HitPoint *hitPoint) { const float2 uv = VLOAD2F(&hitPoint->uv.u); return Transform_ApplyPoint(&mapping->worldToLocal, (float3)(uv.xy, 0.f)); } float3 GlobalMapping3D_Map(__global TextureMapping3D *mapping, __global HitPoint *hitPoint) { const float3 p = VLOAD3F(&hitPoint->p.x); return Transform_ApplyPoint(&mapping->worldToLocal, p); } float2 TextureMapping2D_Map(__global TextureMapping2D *mapping, __global HitPoint *hitPoint) { switch (mapping->type) { case UVMAPPING2D: return UVMapping2D_Map(mapping, hitPoint); default: return 0.f; } } float3 TextureMapping3D_Map(__global TextureMapping3D *mapping, __global HitPoint *hitPoint) { switch (mapping->type) { case UVMAPPING3D: return UVMapping3D_Map(mapping, hitPoint); case GLOBALMAPPING3D: return GlobalMapping3D_Map(mapping, hitPoint); default: return 0.f; } } #line 2 "texture_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ #define TEXTURE_STACK_SIZE 16 //------------------------------------------------------------------------------ // Texture utility functions //------------------------------------------------------------------------------ // Perlin Noise Data #define NOISE_PERM_SIZE 256 __constant int NoisePerm[2 * NOISE_PERM_SIZE] = { 151, 160, 137, 91, 90, 15, 131, 13, 201, 95, 96, 53, 194, 233, 7, 225, 140, 36, 103, 30, 69, 142, // Rest of noise permutation table 8, 99, 37, 240, 21, 10, 23, 190, 6, 148, 247, 120, 234, 75, 0, 26, 197, 62, 94, 252, 219, 203, 117, 35, 11, 32, 57, 177, 33, 88, 237, 149, 56, 87, 174, 20, 125, 136, 171, 168, 68, 175, 74, 165, 71, 134, 139, 48, 27, 166, 77, 146, 158, 231, 83, 111, 229, 122, 60, 211, 133, 230, 220, 105, 92, 41, 55, 46, 245, 40, 244, 102, 143, 54, 65, 25, 63, 161, 1, 216, 80, 73, 209, 76, 132, 187, 208, 89, 18, 169, 200, 196, 135, 130, 116, 188, 159, 86, 164, 100, 109, 198, 173, 186, 3, 64, 52, 217, 226, 250, 124, 123, 5, 202, 38, 147, 118, 126, 255, 82, 85, 212, 207, 206, 59, 227, 47, 16, 58, 17, 182, 189, 28, 42, 223, 183, 170, 213, 119, 248, 152, 2, 44, 154, 163, 70, 221, 153, 101, 155, 167, 43, 172, 9, 129, 22, 39, 253, 19, 98, 108, 110, 79, 113, 224, 232, 178, 185, 112, 104, 218, 246, 97, 228, 251, 34, 242, 193, 238, 210, 144, 12, 191, 179, 162, 241, 81, 51, 145, 235, 249, 14, 239, 107, 49, 192, 214, 31, 181, 199, 106, 157, 184, 84, 204, 176, 115, 121, 50, 45, 127, 4, 150, 254, 138, 236, 205, 93, 222, 114, 67, 29, 24, 72, 243, 141, 128, 195, 78, 66, 215, 61, 156, 180, 151, 160, 137, 91, 90, 15, 131, 13, 201, 95, 96, 53, 194, 233, 7, 225, 140, 36, 103, 30, 69, 142, 8, 99, 37, 240, 21, 10, 23, 190, 6, 148, 247, 120, 234, 75, 0, 26, 197, 62, 94, 252, 219, 203, 117, 35, 11, 32, 57, 177, 33, 88, 237, 149, 56, 87, 174, 20, 125, 136, 171, 168, 68, 175, 74, 165, 71, 134, 139, 48, 27, 166, 77, 146, 158, 231, 83, 111, 229, 122, 60, 211, 133, 230, 220, 105, 92, 41, 55, 46, 245, 40, 244, 102, 143, 54, 65, 25, 63, 161, 1, 216, 80, 73, 209, 76, 132, 187, 208, 89, 18, 169, 200, 196, 135, 130, 116, 188, 159, 86, 164, 100, 109, 198, 173, 186, 3, 64, 52, 217, 226, 250, 124, 123, 5, 202, 38, 147, 118, 126, 255, 82, 85, 212, 207, 206, 59, 227, 47, 16, 58, 17, 182, 189, 28, 42, 223, 183, 170, 213, 119, 248, 152, 2, 44, 154, 163, 70, 221, 153, 101, 155, 167, 43, 172, 9, 129, 22, 39, 253, 19, 98, 108, 110, 79, 113, 224, 232, 178, 185, 112, 104, 218, 246, 97, 228, 251, 34, 242, 193, 238, 210, 144, 12, 191, 179, 162, 241, 81, 51, 145, 235, 249, 14, 239, 107, 49, 192, 214, 31, 181, 199, 106, 157, 184, 84, 204, 176, 115, 121, 50, 45, 127, 4, 150, 254, 138, 236, 205, 93, 222, 114, 67, 29, 24, 72, 243, 141, 128, 195, 78, 66, 215, 61, 156, 180 }; float Grad(int x, int y, int z, float dx, float dy, float dz) { const int h = NoisePerm[NoisePerm[NoisePerm[x] + y] + z] & 15; const float u = h < 8 || h == 12 || h == 13 ? dx : dy; const float v = h < 4 || h == 12 || h == 13 ? dy : dz; return ((h & 1) ? -u : u) + ((h & 2) ? -v : v); } float NoiseWeight(float t) { const float t3 = t * t * t; const float t4 = t3 * t; return 6.f * t4 * t - 15.f * t4 + 10.f * t3; } float Noise(float x, float y, float z) { // Compute noise cell coordinates and offsets int ix = Floor2Int(x); int iy = Floor2Int(y); int iz = Floor2Int(z); const float dx = x - ix, dy = y - iy, dz = z - iz; // Compute gradient weights ix &= (NOISE_PERM_SIZE - 1); iy &= (NOISE_PERM_SIZE - 1); iz &= (NOISE_PERM_SIZE - 1); const float w000 = Grad(ix, iy, iz, dx, dy, dz); const float w100 = Grad(ix + 1, iy, iz, dx - 1, dy, dz); const float w010 = Grad(ix, iy + 1, iz, dx, dy - 1, dz); const float w110 = Grad(ix + 1, iy + 1, iz, dx - 1, dy - 1, dz); const float w001 = Grad(ix, iy, iz + 1, dx, dy, dz - 1); const float w101 = Grad(ix + 1, iy, iz + 1, dx - 1, dy, dz - 1); const float w011 = Grad(ix, iy + 1, iz + 1, dx, dy - 1, dz - 1); const float w111 = Grad(ix + 1, iy + 1, iz + 1, dx - 1, dy - 1, dz - 1); // Compute trilinear interpolation of weights const float wx = NoiseWeight(dx); const float wy = NoiseWeight(dy); const float wz = NoiseWeight(dz); const float x00 = Lerp(wx, w000, w100); const float x10 = Lerp(wx, w010, w110); const float x01 = Lerp(wx, w001, w101); const float x11 = Lerp(wx, w011, w111); const float y0 = Lerp(wy, x00, x10); const float y1 = Lerp(wy, x01, x11); return Lerp(wz, y0, y1); } float Noise3(const float3 P) { return Noise(P.x, P.y, P.z); } float FBm(const float3 P, const float omega, const int maxOctaves) { // Compute number of octaves for anti-aliased FBm const float foctaves = (float)maxOctaves; const int octaves = Floor2Int(foctaves); // Compute sum of octaves of noise for FBm float sum = 0.f, lambda = 1.f, o = 1.f; for (int i = 0; i < octaves; ++i) { sum += o * Noise3(lambda * P); lambda *= 1.99f; o *= omega; } const float partialOctave = foctaves - (float)octaves; sum += o * SmoothStep(.3f, .7f, partialOctave) * Noise3(lambda * P); return sum; } float Turbulence(const float3 P, const float omega, const int maxOctaves) { // Compute number of octaves for anti-aliased FBm const float foctaves = (float)maxOctaves; const int octaves = Floor2Int(foctaves); // Compute sum of octaves of noise for turbulence float sum = 0.f, lambda = 1.f, o = 1.f; for (int i = 0; i < octaves; ++i) { sum += o * fabs(Noise3(lambda * P)); lambda *= 1.99f; o *= omega; } const float partialOctave = foctaves - (float)(octaves); sum += o * SmoothStep(.3f, .7f, partialOctave) * fabs(Noise3(lambda * P)); // finally, add in value to account for average value of fabsf(Noise()) // (~0.2) for the remaining octaves... sum += (maxOctaves - foctaves) * 0.2f; return sum; } //------------------------------------------------------------------------------ // ImageMaps support //------------------------------------------------------------------------------ #if defined(PARAM_HAS_IMAGEMAPS) __global float *ImageMap_GetPixelsAddress(__global float **imageMapBuff, const uint page, const uint offset) { return &imageMapBuff[page][offset]; } float ImageMap_GetTexel_Float(__global float *pixels, const uint width, const uint height, const uint channelCount, const int s, const int t) { const uint u = Mod(s, width); const uint v = Mod(t, height); const uint index = channelCount * (v * width + u); return (channelCount == 1) ? pixels[index] : Spectrum_Y(VLOAD3F(&pixels[index])); } float3 ImageMap_GetTexel_Spectrum(__global float *pixels, const uint width, const uint height, const uint channelCount, const int s, const int t) { const uint u = Mod(s, width); const uint v = Mod(t, height); const uint index = channelCount * (v * width + u); return (channelCount == 1) ? pixels[index] : VLOAD3F(&pixels[index]); } float ImageMap_GetFloat(__global float *pixels, const uint width, const uint height, const uint channelCount, const float u, const float v) { const float s = u * width - 0.5f; const float t = v * height - 0.5f; const int s0 = Floor2Int(s); const int t0 = Floor2Int(t); const float ds = s - s0; const float dt = t - t0; const float ids = 1.f - ds; const float idt = 1.f - dt; const float c0 = ImageMap_GetTexel_Float(pixels, width, height, channelCount, s0, t0); const float c1 = ImageMap_GetTexel_Float(pixels, width, height, channelCount, s0, t0 + 1); const float c2 = ImageMap_GetTexel_Float(pixels, width, height, channelCount, s0 + 1, t0); const float c3 = ImageMap_GetTexel_Float(pixels, width, height, channelCount, s0 + 1, t0 + 1); const float k0 = ids * idt; const float k1 = ids * dt; const float k2 = ds * idt; const float k3 = ds * dt; return (k0 * c0 + k1 *c1 + k2 * c2 + k3 * c3); } float3 ImageMap_GetSpectrum(__global float *pixels, const uint width, const uint height, const uint channelCount, const float u, const float v) { const float s = u * width - 0.5f; const float t = v * height - 0.5f; const int s0 = Floor2Int(s); const int t0 = Floor2Int(t); const float ds = s - s0; const float dt = t - t0; const float ids = 1.f - ds; const float idt = 1.f - dt; const float3 c0 = ImageMap_GetTexel_Spectrum(pixels, width, height, channelCount, s0, t0); const float3 c1 = ImageMap_GetTexel_Spectrum(pixels, width, height, channelCount, s0, t0 + 1); const float3 c2 = ImageMap_GetTexel_Spectrum(pixels, width, height, channelCount, s0 + 1, t0); const float3 c3 = ImageMap_GetTexel_Spectrum(pixels, width, height, channelCount, s0 + 1, t0 + 1); const float k0 = ids * idt; const float k1 = ids * dt; const float k2 = ds * idt; const float k3 = ds * dt; return (k0 * c0 + k1 *c1 + k2 * c2 + k3 * c3); } #endif //------------------------------------------------------------------------------ // ConstFloat texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_TEX_CONST_FLOAT) void ConstFloatTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = texture->constFloat.value; } void ConstFloatTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = texture->constFloat.value; } void ConstFloatTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = 0.f; } #endif //------------------------------------------------------------------------------ // ConstFloat3 texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_TEX_CONST_FLOAT3) void ConstFloat3Texture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = Spectrum_Y(VLOAD3F(&texture->constFloat3.color.r)); } void ConstFloat3Texture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = VLOAD3F(&texture->constFloat3.color.r); } void ConstFloat3Texture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = 0.f; } #endif //------------------------------------------------------------------------------ // ImageMap texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_TEX_IMAGEMAP) void ImageMapTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize IMAGEMAPS_PARAM_DECL) { __global ImageMap *imageMap = &imageMapDescs[texture->imageMapTex.imageMapIndex]; __global float *pixels = ImageMap_GetPixelsAddress( imageMapBuff, imageMap->pageIndex, imageMap->pixelsIndex); const float2 uv = VLOAD2F(&hitPoint->uv.u); const float2 mapUV = TextureMapping2D_Map(&texture->imageMapTex.mapping, hitPoint); texValues[(*texValuesSize)++] = texture->imageMapTex.gain * ImageMap_GetFloat( pixels, imageMap->width, imageMap->height, imageMap->channelCount, mapUV.s0, mapUV.s1); } void ImageMapTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize IMAGEMAPS_PARAM_DECL) { __global ImageMap *imageMap = &imageMapDescs[texture->imageMapTex.imageMapIndex]; __global float *pixels = ImageMap_GetPixelsAddress( imageMapBuff, imageMap->pageIndex, imageMap->pixelsIndex); const float2 uv = VLOAD2F(&hitPoint->uv.u); const float2 mapUV = TextureMapping2D_Map(&texture->imageMapTex.mapping, hitPoint); texValues[(*texValuesSize)++] = texture->imageMapTex.gain * ImageMap_GetSpectrum( pixels, imageMap->width, imageMap->height, imageMap->channelCount, mapUV.s0, mapUV.s1); } void ImageMapTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = VLOAD2F(&texture->imageMapTex.Du); } #endif //------------------------------------------------------------------------------ // Scale texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_TEX_SCALE) void ScaleTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float value = texValues[--(*texValuesSize)] * texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = value; } void ScaleTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float3 value = texValues[--(*texValuesSize)] * texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = value; } void ScaleTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float2 dudv1 = texValues[--(*texValuesSize)]; const float2 dudv2 = texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = fmax(dudv1, dudv2); } #endif //------------------------------------------------------------------------------ // FresnelApproxN & FresnelApproxK texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_FRESNEL_APPROX_N) float FresnelApproxN(const float Fr) { const float sqrtReflectance = sqrt(clamp(Fr, 0.f, .999f)); return (1.f + sqrtReflectance) / (1.f - sqrtReflectance); } float3 FresnelApproxN3(const float3 Fr) { const float3 sqrtReflectance = Spectrum_Sqrt(clamp(Fr, 0.f, .999f)); return (WHITE + sqrtReflectance) / (WHITE - sqrtReflectance); } void FresnelApproxNTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float value = texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = FresnelApproxN(value); } void FresnelApproxNTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float3 value = texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = FresnelApproxN3(value); } void FresnelApproxNTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = texValues[--(*texValuesSize)]; } #endif #if defined (PARAM_ENABLE_FRESNEL_APPROX_K) float FresnelApproxK(const float Fr) { const float reflectance = clamp(Fr, 0.f, .999f); return 2.f * sqrt(reflectance / (1.f - reflectance)); } float3 FresnelApproxK3(const float3 Fr) { const float3 reflectance = clamp(Fr, 0.f, .999f); return 2.f * Spectrum_Sqrt(reflectance / (WHITE - reflectance)); } void FresnelApproxKTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float value = texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = FresnelApproxK(value); } void FresnelApproxKTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float3 value = texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = FresnelApproxK3(value); } void FresnelApproxKTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = texValues[--(*texValuesSize)]; } #endif //------------------------------------------------------------------------------ // CheckerBoard 2D & 3D texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_CHECKERBOARD2D) void CheckerBoard2DTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float value1 = texValues[--(*texValuesSize)]; const float value2 = texValues[--(*texValuesSize)]; const float2 uv = VLOAD2F(&hitPoint->uv.u); const float2 mapUV = TextureMapping2D_Map(&texture->checkerBoard2D.mapping, hitPoint); texValues[(*texValuesSize)++] = ((Floor2Int(mapUV.s0) + Floor2Int(mapUV.s1)) % 2 == 0) ? value1 : value2; } void CheckerBoard2DTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float3 value1 = texValues[--(*texValuesSize)]; const float3 value2 = texValues[--(*texValuesSize)]; const float2 uv = VLOAD2F(&hitPoint->uv.u); const float2 mapUV = TextureMapping2D_Map(&texture->checkerBoard2D.mapping, hitPoint); texValues[(*texValuesSize)++] = ((Floor2Int(mapUV.s0) + Floor2Int(mapUV.s1)) % 2 == 0) ? value1 : value2; } void CheckerBoard2DTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float2 dudv1 = texValues[--(*texValuesSize)]; const float2 dudv2 = texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = fmax(dudv1, dudv2); } #endif #if defined (PARAM_ENABLE_CHECKERBOARD3D) void CheckerBoard3DTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float value1 = texValues[--(*texValuesSize)]; const float value2 = texValues[--(*texValuesSize)]; const float3 mapP = TextureMapping3D_Map(&texture->checkerBoard3D.mapping, hitPoint); texValues[(*texValuesSize)++] = ((Floor2Int(mapP.x) + Floor2Int(mapP.y) + Floor2Int(mapP.z)) % 2 == 0) ? value1 : value2; } void CheckerBoard3DTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float3 value1 = texValues[--(*texValuesSize)]; const float3 value2 = texValues[--(*texValuesSize)]; const float3 mapP = TextureMapping3D_Map(&texture->checkerBoard3D.mapping, hitPoint); texValues[(*texValuesSize)++] = ((Floor2Int(mapP.x) + Floor2Int(mapP.y) + Floor2Int(mapP.z)) % 2 == 0) ? value1 : value2; } void CheckerBoard3DTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float2 dudv1 = texValues[--(*texValuesSize)]; const float2 dudv2 = texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = fmax(dudv1, dudv2); } #endif //------------------------------------------------------------------------------ // Mix texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_TEX_MIX) void MixTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float amt = clamp(texValues[--(*texValuesSize)], 0.f, 1.f);; const float value1 = texValues[--(*texValuesSize)]; const float value2 = texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = Lerp(amt, value1, value2); } void MixTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float3 amt = clamp(texValues[--(*texValuesSize)], 0.f, 1.f); const float3 value1 = texValues[--(*texValuesSize)]; const float3 value2 = texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = mix(value1, value2, amt); } void MixTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float2 dudv1 = texValues[--(*texValuesSize)]; const float2 dudv2 = texValues[--(*texValuesSize)]; const float2 dudv3 = texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = fmax(fmax(dudv1, dudv2), dudv3); } #endif //------------------------------------------------------------------------------ // FBM texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_FBM_TEX) void FBMTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float3 mapP = TextureMapping3D_Map(&texture->fbm.mapping, hitPoint); texValues[(*texValuesSize)++] = FBm(mapP, texture->fbm.omega, texture->fbm.octaves); } void FBMTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float3 mapP = TextureMapping3D_Map(&texture->fbm.mapping, hitPoint); texValues[(*texValuesSize)++] = FBm(mapP, texture->fbm.omega, texture->fbm.octaves); } void FBMTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = (float2)(DUDV_VALUE, DUDV_VALUE); } #endif //------------------------------------------------------------------------------ // Marble texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_MARBLE) // Evaluate marble spline at _t_ __constant float MarbleTexture_c[9][3] = { { .58f, .58f, .6f}, { .58f, .58f, .6f}, { .58f, .58f, .6f}, { .5f, .5f, .5f}, { .6f, .59f, .58f}, { .58f, .58f, .6f}, { .58f, .58f, .6f}, {.2f, .2f, .33f}, { .58f, .58f, .6f} }; float3 MarbleTexture_Evaluate(__global Texture *texture, __global HitPoint *hitPoint) { const float3 P = texture->marble.scale * TextureMapping3D_Map(&texture->marble.mapping, hitPoint); float marble = P.y + texture->marble.variation * FBm(P, texture->marble.omega, texture->marble.octaves); float t = .5f + .5f * sin(marble); #define NC sizeof(MarbleTexture_c) / sizeof(MarbleTexture_c[0]) #define NSEG (NC-3) const int first = Floor2Int(t * NSEG); t = (t * NSEG - first); #undef NC #undef NSEG #define ASSIGN_CF3(a) (float3)(a[0], a[1], a[2]) const float3 c0 = ASSIGN_CF3(MarbleTexture_c[first]); const float3 c1 = ASSIGN_CF3(MarbleTexture_c[first + 1]); const float3 c2 = ASSIGN_CF3(MarbleTexture_c[first + 2]); const float3 c3 = ASSIGN_CF3(MarbleTexture_c[first + 3]); #undef ASSIGN_CF3 // Bezier spline evaluated with de Castilejau's algorithm float3 s0 = mix(c0, c1, t); float3 s1 = mix(c1, c2, t); float3 s2 = mix(c2, c3, t); s0 = mix(s0, s1, t); s1 = mix(s1, s2, t); // Extra scale of 1.5 to increase variation among colors return 1.5f * mix(s0, s1, t); } void MarbleTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = Spectrum_Y(MarbleTexture_Evaluate(texture, hitPoint)); } void MarbleTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = MarbleTexture_Evaluate(texture, hitPoint); } void MarbleTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = (float2)(DUDV_VALUE, DUDV_VALUE); } #endif //------------------------------------------------------------------------------ // Dots texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_DOTS) void DotsTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float value1 = texValues[--(*texValuesSize)]; const float value2 = texValues[--(*texValuesSize)]; const float2 uv = TextureMapping2D_Map(&texture->dots.mapping, hitPoint); const int sCell = Floor2Int(uv.s0 + .5f); const int tCell = Floor2Int(uv.s1 + .5f); // Return _insideDot_ result if point is inside dot if (Noise(sCell + .5f, tCell + .5f, .5f) > 0.f) { const float radius = .35f; const float maxShift = 0.5f - radius; const float sCenter = sCell + maxShift * Noise(sCell + 1.5f, tCell + 2.8f, .5f); const float tCenter = tCell + maxShift * Noise(sCell + 4.5f, tCell + 9.8f, .5f); const float ds = uv.s0 - sCenter, dt = uv.s1 - tCenter; if (ds * ds + dt * dt < radius * radius) { texValues[(*texValuesSize)++] = value1; return; } } texValues[(*texValuesSize)++] = value2; } void DotsTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float3 value1 = texValues[--(*texValuesSize)]; const float3 value2 = texValues[--(*texValuesSize)]; const float2 uv = TextureMapping2D_Map(&texture->dots.mapping, hitPoint); const int sCell = Floor2Int(uv.s0 + .5f); const int tCell = Floor2Int(uv.s1 + .5f); // Return _insideDot_ result if point is inside dot if (Noise(sCell + .5f, tCell + .5f, .5f) > 0.f) { const float radius = .35f; const float maxShift = 0.5f - radius; const float sCenter = sCell + maxShift * Noise(sCell + 1.5f, tCell + 2.8f, .5f); const float tCenter = tCell + maxShift * Noise(sCell + 4.5f, tCell + 9.8f, .5f); const float ds = uv.s0 - sCenter, dt = uv.s1 - tCenter; if (ds * ds + dt * dt < radius * radius) { texValues[(*texValuesSize)++] = value1; return; } } texValues[(*texValuesSize)++] = value2; } void DotsTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float2 dudv1 = texValues[--(*texValuesSize)]; const float2 dudv2 = texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = fmax(dudv1, dudv2); } #endif //------------------------------------------------------------------------------ // Brick texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_BRICK) bool BrickTexture_RunningAlternate(__global Texture *texture, const float3 p, float3 *i, float3 *b, int nWhole) { const float run = texture->brick.run; const float mortarwidth = texture->brick.mortarwidth; const float mortarheight = texture->brick.mortarheight; const float mortardepth = texture->brick.mortardepth; const float sub = nWhole + 0.5f; const float rsub = ceil(sub); (*i).z = floor(p.z); (*b).x = (p.x + (*i).z * run) / sub; (*b).y = (p.y + (*i).z * run) / sub; (*i).x = floor((*b).x); (*i).y = floor((*b).y); (*b).x = ((*b).x - (*i).x) * sub; (*b).y = ((*b).y - (*i).y) * sub; (*b).z = (p.z - (*i).z) * sub; (*i).x += floor((*b).x) / rsub; (*i).y += floor((*b).y) / rsub; (*b).x -= floor((*b).x); (*b).y -= floor((*b).y); return (*b).z > mortarheight && (*b).y > mortardepth && (*b).x > mortarwidth; } bool BrickTexture_Basket(__global Texture *texture, const float3 p, float3 *i) { const float mortarwidth = texture->brick.mortarwidth; const float mortardepth = texture->brick.mortardepth; const float proportion = texture->brick.proportion; const float invproportion = texture->brick.invproportion; (*i).x = floor(p.x); (*i).y = floor(p.y); float bx = p.x - (*i).x; float by = p.y - (*i).y; (*i).x += (*i).y - 2.f * floor(0.5f * (*i).y); const bool split = ((*i).x - 2.f * floor(0.5f * (*i).x)) < 1.f; if (split) { bx = fmod(bx, invproportion); (*i).x = floor(proportion * p.x) * invproportion; } else { by = fmod(by, invproportion); (*i).y = floor(proportion * p.y) * invproportion; } return by > mortardepth && bx > mortarwidth; } bool BrickTexture_Herringbone(__global Texture *texture, const float3 p, float3 *i) { const float mortarwidth = texture->brick.mortarwidth; const float mortarheight = texture->brick.mortarheight; const float proportion = texture->brick.proportion; const float invproportion = texture->brick.invproportion; (*i).y = floor(proportion * p.y); const float px = p.x + (*i).y * invproportion; (*i).x = floor(px); float bx = 0.5f * px - floor(px * 0.5f); bx *= 2.f; float by = proportion * p.y - floor(proportion * p.y); by *= invproportion; if (bx > 1.f + invproportion) { bx = proportion * (bx - 1.f); (*i).y -= floor(bx - 1.f); bx -= floor(bx); bx *= invproportion; by = 1.f; } else if (bx > 1.f) { bx = proportion * (bx - 1.f); (*i).y -= floor(bx - 1.f); bx -= floor(bx); bx *= invproportion; } return by > mortarheight && bx > mortarwidth; } bool BrickTexture_Running(__global Texture *texture, const float3 p, float3 *i, float3 *b) { const float run = texture->brick.run; const float mortarwidth = texture->brick.mortarwidth; const float mortarheight = texture->brick.mortarheight; const float mortardepth = texture->brick.mortardepth; (*i).z = floor(p.z); (*b).x = p.x + (*i).z * run; (*b).y = p.y - (*i).z * run; (*i).x = floor((*b).x); (*i).y = floor((*b).y); (*b).z = p.z - (*i).z; (*b).x -= (*i).x; (*b).y -= (*i).y; return (*b).z > mortarheight && (*b).y > mortardepth && (*b).x > mortarwidth; } bool BrickTexture_English(__global Texture *texture, const float3 p, float3 *i, float3 *b) { const float run = texture->brick.run; const float mortarwidth = texture->brick.mortarwidth; const float mortarheight = texture->brick.mortarheight; const float mortardepth = texture->brick.mortardepth; (*i).z = floor(p.z); (*b).x = p.x + (*i).z * run; (*b).y = p.y - (*i).z * run; (*i).x = floor((*b).x); (*i).y = floor((*b).y); (*b).z = p.z - (*i).z; const float divider = floor(fmod(fabs((*i).z), 2.f)) + 1.f; (*b).x = (divider * (*b).x - floor(divider * (*b).x)) / divider; (*b).y = (divider * (*b).y - floor(divider * (*b).y)) / divider; return (*b).z > mortarheight && (*b).y > mortardepth && (*b).x > mortarwidth; } bool BrickTexture_Evaluate(__global Texture *texture, __global HitPoint *hitPoint) { #define BRICK_EPSILON 1e-3f const float3 P = TextureMapping3D_Map(&texture->brick.mapping, hitPoint); const float offs = BRICK_EPSILON + texture->brick.mortarsize; float3 bP = P + (float3)(offs, offs, offs); // Normalize coordinates according brick dimensions bP.x /= texture->brick.brickwidth; bP.y /= texture->brick.brickdepth; bP.z /= texture->brick.brickheight; bP += VLOAD3F(&texture->brick.offsetx); float3 brickIndex; float3 bevel; bool b; switch (texture->brick.bond) { case FLEMISH: b = BrickTexture_RunningAlternate(texture, bP, &brickIndex, &bevel, 1); break; case RUNNING: b = BrickTexture_Running(texture, bP, &brickIndex, &bevel); break; case ENGLISH: b = BrickTexture_English(texture, bP, &brickIndex, &bevel); break; case HERRINGBONE: b = BrickTexture_Herringbone(texture, bP, &brickIndex); break; case BASKET: b = BrickTexture_Basket(texture, bP, &brickIndex); break; case KETTING: b = BrickTexture_RunningAlternate(texture, bP, &brickIndex, &bevel, 2); break; default: b = true; break; } return b; #undef BRICK_EPSILON } void BrickTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float value1 = texValues[--(*texValuesSize)]; const float value2 = texValues[--(*texValuesSize)]; const float value3 = texValues[--(*texValuesSize)]; if (BrickTexture_Evaluate(texture, hitPoint)) texValues[(*texValuesSize)++] = value1 * value3; else texValues[(*texValuesSize)++] = value2; } void BrickTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float3 value1 = texValues[--(*texValuesSize)]; const float3 value2 = texValues[--(*texValuesSize)]; const float3 value3 = texValues[--(*texValuesSize)]; if (BrickTexture_Evaluate(texture, hitPoint)) texValues[(*texValuesSize)++] = value1 * value3; else texValues[(*texValuesSize)++] = value2; } void BrickTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float2 dudv1 = texValues[--(*texValuesSize)]; const float2 dudv2 = texValues[--(*texValuesSize)]; const float2 dudv3 = texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = fmax(fmax(dudv1, dudv2), dudv3); } #endif //------------------------------------------------------------------------------ // Add texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_TEX_ADD) void AddTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float value = texValues[--(*texValuesSize)] + texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = value; } void AddTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float3 value = texValues[--(*texValuesSize)] + texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = value; } void AddTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float2 dudv1 = texValues[--(*texValuesSize)]; const float2 dudv2 = texValues[--(*texValuesSize)]; texValues[(*texValuesSize)++] = fmax(dudv1, dudv2); } #endif //------------------------------------------------------------------------------ // Windy texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_WINDY) void WindyTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float3 mapP = TextureMapping3D_Map(&texture->windy.mapping, hitPoint); const float windStrength = FBm(.1f * mapP, .5f, 3); const float waveHeight = FBm(mapP, .5f, 6); texValues[(*texValuesSize)++] = fabs(windStrength) * waveHeight; } void WindyTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float3 mapP = TextureMapping3D_Map(&texture->windy.mapping, hitPoint); const float windStrength = FBm(.1f * mapP, .5f, 3); const float waveHeight = FBm(mapP, .5f, 6); texValues[(*texValuesSize)++] = fabs(windStrength) * waveHeight; } void WindyTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = (float2)(DUDV_VALUE, DUDV_VALUE); } #endif //------------------------------------------------------------------------------ // Wrinkled texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_WRINKLED) void WrinkledTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float3 mapP = TextureMapping3D_Map(&texture->wrinkled.mapping, hitPoint); texValues[(*texValuesSize)++] = Turbulence(mapP, texture->wrinkled.omega, texture->wrinkled.octaves); } void WrinkledTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float3 mapP = TextureMapping3D_Map(&texture->wrinkled.mapping, hitPoint); texValues[(*texValuesSize)++] = Turbulence(mapP, texture->wrinkled.omega, texture->wrinkled.octaves); } void WrinkledTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = (float2)(DUDV_VALUE, DUDV_VALUE); } #endif //------------------------------------------------------------------------------ // UV texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_TEX_UV) void UVTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float2 uv = TextureMapping2D_Map(&texture->uvTex.mapping, hitPoint); texValues[(*texValuesSize)++] = Spectrum_Y((float3)(uv.s0 - Floor2Int(uv.s0), uv.s1 - Floor2Int(uv.s1), 0.f)); } void UVTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float2 uv = TextureMapping2D_Map(&texture->uvTex.mapping, hitPoint); texValues[(*texValuesSize)++] = (float3)(uv.s0 - Floor2Int(uv.s0), uv.s1 - Floor2Int(uv.s1), 0.f); } void UVTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = (float2)(DUDV_VALUE, DUDV_VALUE); } #endif //------------------------------------------------------------------------------ // Band texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_TEX_BAND) void BandTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float a = clamp(texValues[--(*texValuesSize)], 0.f, 1.f); const uint last = texture->band.size - 1; if (a < texture->band.offsets[0]) texValues[(*texValuesSize)++] = Spectrum_Y(VLOAD3F(&texture->band.values[0].r)); else if (a >= texture->band.offsets[last]) texValues[(*texValuesSize)++] = Spectrum_Y(VLOAD3F(&texture->band.values[last].r)); else { uint p = 0; for (; p <= last; ++p) { if (a < texture->band.offsets[p]) break; } const float p1 = Spectrum_Y(VLOAD3F(&texture->band.values[p - 1].r)); const float p0 = Spectrum_Y(VLOAD3F(&texture->band.values[p].r)); const float o1 = texture->band.offsets[p - 1]; const float o0 = texture->band.offsets[p]; texValues[(*texValuesSize)++] = Lerp((a - o1) / (o0 - o1), p1, p0); } } void BandTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float a = clamp(Spectrum_Y(texValues[--(*texValuesSize)]), 0.f, 1.f); const uint last = texture->band.size - 1; if (a < texture->band.offsets[0]) texValues[(*texValuesSize)++] = VLOAD3F(&texture->band.values[0].r); else if (a >= texture->band.offsets[last]) texValues[(*texValuesSize)++] = VLOAD3F(&texture->band.values[last].r); else { uint p = 0; for (; p <= last; ++p) { if (a < texture->band.offsets[p]) break; } const float3 p1 = VLOAD3F(&texture->band.values[p - 1].r); const float3 p0 = VLOAD3F(&texture->band.values[p].r); const float o1 = texture->band.offsets[p - 1]; const float o0 = texture->band.offsets[p]; texValues[(*texValuesSize)++] = Lerp3((a - o1) / (o0 - o1), p1, p0); } } void BandTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { // Nothing to do: //const float2 dudv = texValues[--(*texValuesSize)]; //texValues[(*texValuesSize)++] = dudv; } #endif //------------------------------------------------------------------------------ // HitPointColor texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_TEX_HITPOINTCOLOR) void HitPointColorTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = Spectrum_Y(VLOAD3F(&hitPoint->color.r)); } void HitPointColorTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = VLOAD3F(&hitPoint->color.r); } void HitPointColorTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = 0.f; } #endif //------------------------------------------------------------------------------ // HitPointAlpha texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_TEX_HITPOINTALPHA) void HitPointAlphaTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = hitPoint->alpha; } void HitPointAlphaTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const float alpha = hitPoint->alpha; texValues[(*texValuesSize)++] = (float3)(alpha, alpha, alpha); } void HitPointAlphaTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = 0.f; } #endif //------------------------------------------------------------------------------ // HitPointGrey texture //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_TEX_HITPOINTGREY) void HitPointGreyTexture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const uint channel = texture->hitPointGrey.channel; switch (channel) { case 0: texValues[*texValuesSize] = hitPoint->color.r; break; case 1: texValues[*texValuesSize] = hitPoint->color.g; break; case 2: texValues[*texValuesSize] = hitPoint->color.b; break; default: texValues[*texValuesSize] = Spectrum_Y(VLOAD3F(&hitPoint->color.r)); break; } ++(*texValuesSize); } void HitPointGreyTexture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { const uint channel = texture->hitPointGrey.channel; float v; switch (channel) { case 0: v = hitPoint->color.r; break; case 1: v = hitPoint->color.g; break; case 2: v = hitPoint->color.b; break; default: v = Spectrum_Y(VLOAD3F(&hitPoint->color.r)); break; } texValues[(*texValuesSize)++] = (float3)(v, v, v); } void HitPointGreyTexture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize) { texValues[(*texValuesSize)++] = 0.f; } #endif //------------------------------------------------------------------------------ // Generic texture functions with support for recursive textures //------------------------------------------------------------------------------ uint Texture_AddSubTexture(__global Texture *texture, __global Texture *todoTex[TEXTURE_STACK_SIZE], uint *todoTexSize TEXTURES_PARAM_DECL) { switch (texture->type) { #if defined(PARAM_ENABLE_TEX_SCALE) case SCALE_TEX: todoTex[(*todoTexSize)++] = &texs[texture->scaleTex.tex1Index]; todoTex[(*todoTexSize)++] = &texs[texture->scaleTex.tex2Index]; return 2; #endif #if defined (PARAM_ENABLE_FRESNEL_APPROX_N) case FRESNEL_APPROX_N: todoTex[(*todoTexSize)++] = &texs[texture->fresnelApproxN.texIndex]; return 1; #endif #if defined (PARAM_ENABLE_FRESNEL_APPROX_K) case FRESNEL_APPROX_K: todoTex[(*todoTexSize)++] = &texs[texture->fresnelApproxK.texIndex]; return 1; #endif #if defined(PARAM_ENABLE_CHECKERBOARD2D) case CHECKERBOARD2D: todoTex[(*todoTexSize)++] = &texs[texture->checkerBoard2D.tex1Index]; todoTex[(*todoTexSize)++] = &texs[texture->checkerBoard2D.tex2Index]; return 2; #endif #if defined(PARAM_ENABLE_CHECKERBOARD3D) case CHECKERBOARD3D: todoTex[(*todoTexSize)++] = &texs[texture->checkerBoard3D.tex1Index]; todoTex[(*todoTexSize)++] = &texs[texture->checkerBoard3D.tex2Index]; return 2; #endif #if defined (PARAM_ENABLE_TEX_MIX) case MIX_TEX: todoTex[(*todoTexSize)++] = &texs[texture->mixTex.amountTexIndex]; todoTex[(*todoTexSize)++] = &texs[texture->mixTex.tex1Index]; todoTex[(*todoTexSize)++] = &texs[texture->mixTex.tex2Index]; return 3; #endif #if defined(PARAM_ENABLE_DOTS) case DOTS: todoTex[(*todoTexSize)++] = &texs[texture->dots.insideIndex]; todoTex[(*todoTexSize)++] = &texs[texture->dots.outsideIndex]; return 2; #endif #if defined(PARAM_ENABLE_BRICK) case BRICK: todoTex[(*todoTexSize)++] = &texs[texture->brick.tex1Index]; todoTex[(*todoTexSize)++] = &texs[texture->brick.tex2Index]; todoTex[(*todoTexSize)++] = &texs[texture->brick.tex3Index]; return 3; #endif #if defined(PARAM_ENABLE_TEX_ADD) case ADD_TEX: todoTex[(*todoTexSize)++] = &texs[texture->addTex.tex1Index]; todoTex[(*todoTexSize)++] = &texs[texture->addTex.tex2Index]; return 2; #endif #if defined (PARAM_ENABLE_TEX_BAND) case BAND_TEX: todoTex[(*todoTexSize)++] = &texs[texture->band.amountTexIndex]; return 1; #endif #if defined (PARAM_ENABLE_TEX_HITPOINTGREY) case HITPOINTGREY: #endif #if defined (PARAM_ENABLE_TEX_HITPOINTALPHA) case HITPOINTALPHA: #endif #if defined (PARAM_ENABLE_TEX_HITPOINTCOLOR) case HITPOINTCOLOR: #endif #if defined (PARAM_ENABLE_TEX_UV) case UV_TEX: #endif #if defined (PARAM_ENABLE_WRINKLED) case WRINKLED: #endif #if defined (PARAM_ENABLE_WINDY) case WINDY: #endif #if defined (PARAM_ENABLE_MARBLE) case MARBLE: #endif #if defined (PARAM_ENABLE_FBM_TEX) case FBM_TEX: #endif #if defined(PARAM_ENABLE_TEX_CONST_FLOAT) case CONST_FLOAT: #endif #if defined(PARAM_ENABLE_TEX_CONST_FLOAT3) case CONST_FLOAT3: #endif #if defined(PARAM_ENABLE_TEX_IMAGEMAP) case IMAGEMAP: #endif default: return 0; } } //------------------------------------------------------------------------------ // Float texture channel //------------------------------------------------------------------------------ void Texture_EvaluateFloat(__global Texture *texture, __global HitPoint *hitPoint, float texValues[TEXTURE_STACK_SIZE], uint *texValuesSize IMAGEMAPS_PARAM_DECL) { switch (texture->type) { #if defined(PARAM_ENABLE_TEX_CONST_FLOAT) case CONST_FLOAT: ConstFloatTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_CONST_FLOAT3) case CONST_FLOAT3: ConstFloat3Texture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_IMAGEMAP) case IMAGEMAP: ImageMapTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize IMAGEMAPS_PARAM); break; #endif #if defined(PARAM_ENABLE_TEX_SCALE) case SCALE_TEX: ScaleTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_FRESNEL_APPROX_N) case FRESNEL_APPROX_N: FresnelApproxNTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_FRESNEL_APPROX_K) case FRESNEL_APPROX_K: FresnelApproxKTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_CHECKERBOARD2D) case CHECKERBOARD2D: CheckerBoard2DTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_CHECKERBOARD3D) case CHECKERBOARD3D: CheckerBoard3DTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_TEX_MIX) case MIX_TEX: MixTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_FBM_TEX) case FBM_TEX: FBMTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_MARBLE) case MARBLE: MarbleTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_DOTS) case DOTS: DotsTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_BRICK) case BRICK: BrickTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_ADD) case ADD_TEX: AddTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_WINDY) case WINDY: WindyTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_WRINKLED) case WRINKLED: WrinkledTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_UV) case UV_TEX: UVTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_BAND) case BAND_TEX: BandTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_HITPOINTCOLOR) case HITPOINTCOLOR: HitPointColorTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_HITPOINTALPHA) case HITPOINTALPHA: HitPointAlphaTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_HITPOINTGREY) case HITPOINTGREY: HitPointGreyTexture_EvaluateFloat(texture, hitPoint, texValues, texValuesSize); break; #endif default: // Do nothing break; } } float Texture_GetFloatValue(__global Texture *texture, __global HitPoint *hitPoint TEXTURES_PARAM_DECL) { __global Texture *todoTex[TEXTURE_STACK_SIZE]; uint todoTexSize = 0; __global Texture *pendingTex[TEXTURE_STACK_SIZE]; uint targetTexCount[TEXTURE_STACK_SIZE]; uint pendingTexSize = 0; float texValues[TEXTURE_STACK_SIZE]; uint texValuesSize = 0; const uint subTexCount = Texture_AddSubTexture(texture, todoTex, &todoTexSize TEXTURES_PARAM); if (subTexCount == 0) { // A fast path for evaluating non recursive textures Texture_EvaluateFloat(texture, hitPoint, texValues, &texValuesSize IMAGEMAPS_PARAM); } else { // Normal complex path for evaluating non recursive textures pendingTex[pendingTexSize] = texture; targetTexCount[pendingTexSize++] = subTexCount; do { if ((pendingTexSize > 0) && (texValuesSize == targetTexCount[pendingTexSize - 1])) { // Pop the a texture to do __global Texture *tex = pendingTex[--pendingTexSize]; Texture_EvaluateFloat(tex, hitPoint, texValues, &texValuesSize IMAGEMAPS_PARAM); continue; } if (todoTexSize > 0) { // Pop the a texture to do __global Texture *tex = todoTex[--todoTexSize]; // Add this texture to the list of pending one const uint subTexCount = Texture_AddSubTexture(tex, todoTex, &todoTexSize TEXTURES_PARAM); pendingTex[pendingTexSize] = tex; targetTexCount[pendingTexSize++] = subTexCount + texValuesSize; } } while ((todoTexSize > 0) || (pendingTexSize > 0)); } return texValues[0]; } //------------------------------------------------------------------------------ // Color texture channel //------------------------------------------------------------------------------ void Texture_EvaluateSpectrum(__global Texture *texture, __global HitPoint *hitPoint, float3 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize IMAGEMAPS_PARAM_DECL) { switch (texture->type) { #if defined(PARAM_ENABLE_TEX_CONST_FLOAT) case CONST_FLOAT: ConstFloatTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_CONST_FLOAT3) case CONST_FLOAT3: ConstFloat3Texture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_IMAGEMAP) case IMAGEMAP: ImageMapTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize IMAGEMAPS_PARAM); break; #endif #if defined(PARAM_ENABLE_TEX_SCALE) case SCALE_TEX: ScaleTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_FRESNEL_APPROX_N) case FRESNEL_APPROX_N: FresnelApproxNTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_FRESNEL_APPROX_K) case FRESNEL_APPROX_K: FresnelApproxKTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_CHECKERBOARD2D) case CHECKERBOARD2D: CheckerBoard2DTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_CHECKERBOARD3D) case CHECKERBOARD3D: CheckerBoard3DTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_TEX_MIX) case MIX_TEX: MixTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_FBM_TEX) case FBM_TEX: FBMTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_MARBLE) case MARBLE: MarbleTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_DOTS) case DOTS: DotsTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_BRICK) case BRICK: BrickTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_ADD) case ADD_TEX: AddTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_WINDY) case WINDY: WindyTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_WRINKLED) case WRINKLED: WrinkledTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_UV) case UV_TEX: UVTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_BAND) case BAND_TEX: BandTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_HITPOINTCOLOR) case HITPOINTCOLOR: HitPointColorTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_HITPOINTALPHA) case HITPOINTALPHA: HitPointAlphaTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_HITPOINTGREY) case HITPOINTGREY: HitPointGreyTexture_EvaluateSpectrum(texture, hitPoint, texValues, texValuesSize); break; #endif default: // Do nothing break; } } float3 Texture_GetSpectrumValue(__global Texture *texture, __global HitPoint *hitPoint TEXTURES_PARAM_DECL) { __global Texture *todoTex[TEXTURE_STACK_SIZE]; uint todoTexSize = 0; __global Texture *pendingTex[TEXTURE_STACK_SIZE]; uint targetTexCount[TEXTURE_STACK_SIZE]; uint pendingTexSize = 0; float3 texValues[TEXTURE_STACK_SIZE]; uint texValuesSize = 0; const uint subTexCount = Texture_AddSubTexture(texture, todoTex, &todoTexSize TEXTURES_PARAM); if (subTexCount == 0) { // A fast path for evaluating non recursive textures Texture_EvaluateSpectrum(texture, hitPoint, texValues, &texValuesSize IMAGEMAPS_PARAM); } else { // Normal complex path for evaluating non recursive textures pendingTex[pendingTexSize] = texture; targetTexCount[pendingTexSize++] = subTexCount; do { if ((pendingTexSize > 0) && (texValuesSize == targetTexCount[pendingTexSize - 1])) { // Pop the a texture to do __global Texture *tex = pendingTex[--pendingTexSize]; Texture_EvaluateSpectrum(tex, hitPoint, texValues, &texValuesSize IMAGEMAPS_PARAM); continue; } if (todoTexSize > 0) { // Pop the a texture to do __global Texture *tex = todoTex[--todoTexSize]; // Add this texture to the list of pending one const uint subTexCount = Texture_AddSubTexture(tex, todoTex, &todoTexSize TEXTURES_PARAM); pendingTex[pendingTexSize] = tex; targetTexCount[pendingTexSize++] = subTexCount + texValuesSize; } } while ((todoTexSize > 0) || (pendingTexSize > 0)); } return texValues[0]; } //------------------------------------------------------------------------------ // DuDv texture information //------------------------------------------------------------------------------ void Texture_EvaluateDuDv(__global Texture *texture, __global HitPoint *hitPoint, float2 texValues[TEXTURE_STACK_SIZE], uint *texValuesSize IMAGEMAPS_PARAM_DECL) { switch (texture->type) { #if defined(PARAM_ENABLE_TEX_CONST_FLOAT) case CONST_FLOAT: ConstFloatTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_CONST_FLOAT3) case CONST_FLOAT3: ConstFloat3Texture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_IMAGEMAP) case IMAGEMAP: ImageMapTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_SCALE) case SCALE_TEX: ScaleTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_FRESNEL_APPROX_N) case FRESNEL_APPROX_N: FresnelApproxNTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_FRESNEL_APPROX_K) case FRESNEL_APPROX_K: FresnelApproxKTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_CHECKERBOARD2D) case CHECKERBOARD2D: CheckerBoard2DTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_CHECKERBOARD3D) case CHECKERBOARD3D: CheckerBoard3DTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_TEX_MIX) case MIX_TEX: MixTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_FBM_TEX) case FBM_TEX: FBMTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_MARBLE) case MARBLE: MarbleTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_DOTS) case DOTS: DotsTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined (PARAM_ENABLE_BRICK) case BRICK: BrickTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_ADD) case ADD_TEX: AddTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_WINDY) case WINDY: WindyTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_WRINKLED) case WRINKLED: WrinkledTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_UV) case UV_TEX: UVTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_BAND) case BAND_TEX: BandTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_HITPOINTCOLOR) case HITPOINTCOLOR: HitPointColorTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_HITPOINTALPHA) case HITPOINTALPHA: HitPointAlphaTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif #if defined(PARAM_ENABLE_TEX_HITPOINTGREY) case HITPOINTGREY: HitPointGreyTexture_EvaluateDuDv(texture, hitPoint, texValues, texValuesSize); break; #endif default: // Do nothing break; } } float2 Texture_GetDuDv(__global Texture *texture, __global HitPoint *hitPoint TEXTURES_PARAM_DECL) { __global Texture *todoTex[TEXTURE_STACK_SIZE]; uint todoTexSize = 0; __global Texture *pendingTex[TEXTURE_STACK_SIZE]; uint targetTexCount[TEXTURE_STACK_SIZE]; uint pendingTexSize = 0; float2 texValues[TEXTURE_STACK_SIZE]; uint texValuesSize = 0; const uint subTexCount = Texture_AddSubTexture(texture, todoTex, &todoTexSize TEXTURES_PARAM); if (subTexCount == 0) { // A fast path for evaluating non recursive textures Texture_EvaluateDuDv(texture, hitPoint, texValues, &texValuesSize IMAGEMAPS_PARAM); } else { // Normal complex path for evaluating non recursive textures pendingTex[pendingTexSize] = texture; targetTexCount[pendingTexSize++] = subTexCount; do { if ((pendingTexSize > 0) && (texValuesSize == targetTexCount[pendingTexSize - 1])) { // Pop the a texture to do __global Texture *tex = pendingTex[--pendingTexSize]; Texture_EvaluateDuDv(tex, hitPoint, texValues, &texValuesSize IMAGEMAPS_PARAM); continue; } if (todoTexSize > 0) { // Pop the a texture to do __global Texture *tex = todoTex[--todoTexSize]; // Add this texture to the list of pending one const uint subTexCount = Texture_AddSubTexture(tex, todoTex, &todoTexSize TEXTURES_PARAM); pendingTex[pendingTexSize] = tex; targetTexCount[pendingTexSize++] = subTexCount + texValuesSize; } } while ((todoTexSize > 0) || (pendingTexSize > 0)); } return texValues[0]; } #line 2 "material_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ //------------------------------------------------------------------------------ // Generic material related functions //------------------------------------------------------------------------------ float SchlickDistribution_SchlickZ(const float roughness, float cosNH) { const float d = 1.f + (roughness - 1.f) * cosNH * cosNH; return (roughness > 0.f) ? (roughness / (d * d)) : INFINITY; } float SchlickDistribution_SchlickA(const float3 H, const float anisotropy) { const float h = sqrt(H.x * H.x + H.y * H.y); if (h > 0.f) { const float w = (anisotropy > 0.f ? H.x : H.y) / h; const float p = 1.f - fabs(anisotropy); return sqrt(p / (p * p + w * w * (1.f - p * p))); } return 1.f; } float SchlickDistribution_D(const float roughness, const float3 wh, const float anisotropy) { const float cosTheta = fabs(wh.z); return SchlickDistribution_SchlickZ(roughness, cosTheta) * SchlickDistribution_SchlickA(wh, anisotropy) * M_1_PI_F; } float SchlickDistribution_SchlickG(const float roughness, const float costheta) { return costheta / (costheta * (1.f - roughness) + roughness); } float SchlickDistribution_G(const float roughness, const float3 fixedDir, const float3 sampledDir) { return SchlickDistribution_SchlickG(roughness, fabs(fixedDir.z)) * SchlickDistribution_SchlickG(roughness, fabs(sampledDir.z)); } float GetPhi(const float a, const float b) { return M_PI_F * .5f * sqrt(a * b / (1.f - a * (1.f - b))); } void SchlickDistribution_SampleH(const float roughness, const float anisotropy, const float u0, const float u1, float3 *wh, float *d, float *pdf) { float u1x4 = u1 * 4.f; // Values of roughness < .0001f seems to trigger some kind of exceptions with // AMD OpenCL on GPUs. The result is a nearly freeze of the PC. const float cos2Theta = (roughness < .0001f) ? 1.f : (u0 / (roughness * (1.f - u0) + u0)); const float cosTheta = sqrt(cos2Theta); const float sinTheta = sqrt(1.f - cos2Theta); const float p = 1.f - fabs(anisotropy); float phi; if (u1x4 < 1.f) { phi = GetPhi(u1x4 * u1x4, p * p); } else if (u1x4 < 2.f) { u1x4 = 2.f - u1x4; phi = M_PI_F - GetPhi(u1x4 * u1x4, p * p); } else if (u1x4 < 3.f) { u1x4 -= 2.f; phi = M_PI_F + GetPhi(u1x4 * u1x4, p * p); } else { u1x4 = 4.f - u1x4; phi = M_PI_F * 2.f - GetPhi(u1x4 * u1x4, p * p); } if (anisotropy > 0.f) phi += M_PI_F * .5f; *wh = (float3)(sinTheta * cos(phi), sinTheta * sin(phi), cosTheta); *d = SchlickDistribution_SchlickZ(roughness, cosTheta) * SchlickDistribution_SchlickA(*wh, anisotropy) * M_1_PI_F; *pdf = *d; } float SchlickDistribution_Pdf(const float roughness, const float3 wh, const float anisotropy) { return SchlickDistribution_D(roughness, wh, anisotropy); } float3 FresnelSlick_Evaluate(const float3 normalIncidence, const float cosi) { return normalIncidence + (WHITE - normalIncidence) * pow(1.f - cosi, 5.f); } float3 FrDiel2(const float cosi, const float3 cost, const float3 eta) { float3 Rparl = eta * cosi; Rparl = (cost - Rparl) / (cost + Rparl); float3 Rperp = eta * cost; Rperp = (cosi - Rperp) / (cosi + Rperp); return (Rparl * Rparl + Rperp * Rperp) * .5f; } float3 FrFull(const float cosi, const float3 cost, const float3 eta, const float3 k) { const float3 tmp = (eta * eta + k * k) * (cosi * cosi) + (cost * cost); const float3 Rparl2 = (tmp - (2.f * cosi * cost) * eta) / (tmp + (2.f * cosi * cost) * eta); const float3 tmp_f = (eta * eta + k * k) * (cost * cost) + (cosi * cosi); const float3 Rperp2 = (tmp_f - (2.f * cosi * cost) * eta) / (tmp_f + (2.f * cosi * cost) * eta); return (Rparl2 + Rperp2) * 0.5f; } float3 FresnelGeneral_Evaluate(const float3 eta, const float3 k, const float cosi) { float3 sint2 = fmax(0.f, 1.f - cosi * cosi); if (cosi > 0.f) sint2 /= eta * eta; else sint2 *= eta * eta; sint2 = Spectrum_Clamp(sint2); const float3 cost2 = 1.f - sint2; if (cosi > 0.f) { const float3 a = 2.f * k * k * sint2; return FrFull(cosi, Spectrum_Sqrt((cost2 + Spectrum_Sqrt(cost2 * cost2 + a * a)) / 2.f), eta, k); } else { const float3 a = 2.f * k * k * sint2; const float3 d2 = eta * eta + k * k; return FrFull(-cosi, Spectrum_Sqrt((cost2 + Spectrum_Sqrt(cost2 * cost2 + a * a)) / 2.f), eta / d2, -k / d2); } } float3 FresnelCauchy_Evaluate(const float eta, const float cosi) { // Compute indices of refraction for dielectric const bool entering = (cosi > 0.f); // Compute _sint_ using Snell's law const float eta2 = eta * eta; const float sint2 = (entering ? 1.f / eta2 : eta2) * fmax(0.f, 1.f - cosi * cosi); // Handle total internal reflection if (sint2 >= 1.f) return WHITE; else return FrDiel2(fabs(cosi), sqrt(fmax(0.f, 1.f - sint2)), entering ? eta : 1.f / eta); } //------------------------------------------------------------------------------ // Matte material //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_MAT_MATTE) float3 MatteMaterial_Evaluate(__global Material *material, __global HitPoint *hitPoint, const float3 lightDir, const float3 eyeDir, BSDFEvent *event, float *directPdfW TEXTURES_PARAM_DECL) { if (directPdfW) *directPdfW = fabs(lightDir.z * M_1_PI_F); *event = DIFFUSE | REFLECT; const float3 kd = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->matte.kdTexIndex], hitPoint TEXTURES_PARAM)); return M_1_PI_F * kd; } float3 MatteMaterial_Sample(__global Material *material, __global HitPoint *hitPoint, const float3 fixedDir, float3 *sampledDir, const float u0, const float u1, float *pdfW, float *cosSampledDir, BSDFEvent *event TEXTURES_PARAM_DECL) { if (fabs(fixedDir.z) < DEFAULT_COS_EPSILON_STATIC) return BLACK; *sampledDir = (signbit(fixedDir.z) ? -1.f : 1.f) * CosineSampleHemisphereWithPdf(u0, u1, pdfW); *cosSampledDir = fabs((*sampledDir).z); if (*cosSampledDir < DEFAULT_COS_EPSILON_STATIC) return BLACK; *event = DIFFUSE | REFLECT; const float3 kd = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->matte.kdTexIndex], hitPoint TEXTURES_PARAM)); return M_1_PI_F * kd; } #endif //------------------------------------------------------------------------------ // Mirror material //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_MAT_MIRROR) float3 MirrorMaterial_Sample(__global Material *material, __global HitPoint *hitPoint, const float3 fixedDir, float3 *sampledDir, const float u0, const float u1, float *pdfW, float *cosSampledDir, BSDFEvent *event TEXTURES_PARAM_DECL) { *event = SPECULAR | REFLECT; *sampledDir = (float3)(-fixedDir.x, -fixedDir.y, fixedDir.z); *pdfW = 1.f; *cosSampledDir = fabs((*sampledDir).z); const float3 kr = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->mirror.krTexIndex], hitPoint TEXTURES_PARAM)); // The cosSampledDir is used to compensate the other one used inside the integrator return kr / (*cosSampledDir); } #endif //------------------------------------------------------------------------------ // Glass material //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_MAT_GLASS) float3 GlassMaterial_Sample(__global Material *material, __global HitPoint *hitPoint, const float3 localFixedDir, float3 *localSampledDir, const float u0, const float u1, const float passThroughEvent, float *pdfW, float *absCosSampledDir, BSDFEvent *event TEXTURES_PARAM_DECL) { const float3 kt = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->glass.ktTexIndex], hitPoint TEXTURES_PARAM)); const float3 kr = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->glass.krTexIndex], hitPoint TEXTURES_PARAM)); const bool isKtBlack = Spectrum_IsBlack(kt); const bool isKrBlack = Spectrum_IsBlack(kr); if (isKtBlack && isKrBlack) return BLACK; const bool entering = (CosTheta(localFixedDir) > 0.f); const float nc = Texture_GetFloatValue(&texs[material->glass.ousideIorTexIndex], hitPoint TEXTURES_PARAM); const float nt = Texture_GetFloatValue(&texs[material->glass.iorTexIndex], hitPoint TEXTURES_PARAM); const float ntc = nt / nc; const float eta = entering ? (nc / nt) : ntc; const float costheta = CosTheta(localFixedDir); // Decide to transmit or reflect const float threshold = isKrBlack ? 1.f : (isKtBlack ? 0.f : .5f); float3 result; if (passThroughEvent < threshold) { // Transmit // Compute transmitted ray direction const float sini2 = SinTheta2(localFixedDir); const float eta2 = eta * eta; const float sint2 = eta2 * sini2; // Handle total internal reflection for transmission if (sint2 >= 1.f) return BLACK; const float cost = sqrt(fmax(0.f, 1.f - sint2)) * (entering ? -1.f : 1.f); *localSampledDir = (float3)(-eta * localFixedDir.x, -eta * localFixedDir.y, cost); *absCosSampledDir = fabs(CosTheta(*localSampledDir)); *event = SPECULAR | TRANSMIT; *pdfW = threshold; //if (!hitPoint.fromLight) result = (1.f - FresnelCauchy_Evaluate(ntc, cost)) * eta2; //else // result = (1.f - FresnelCauchy_Evaluate(ntc, costheta)); result *= kt; } else { // Reflect *localSampledDir = (float3)(-localFixedDir.x, -localFixedDir.y, localFixedDir.z); *absCosSampledDir = fabs(CosTheta(*localSampledDir)); *event = SPECULAR | REFLECT; *pdfW = 1.f - threshold; result = kr * FresnelCauchy_Evaluate(ntc, costheta); } // The absCosSampledDir is used to compensate the other one used inside the integrator return result / (*absCosSampledDir); } #endif //------------------------------------------------------------------------------ // Metal material //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_MAT_METAL) float3 GlossyReflection(const float3 fixedDir, const float exponent, const float u0, const float u1) { // Ray from outside going in ? const bool into = (fixedDir.z > 0.f); const float3 shadeN = (float3)(0.f, 0.f, into ? 1.f : -1.f); const float phi = 2.f * M_PI_F * u0; const float cosTheta = pow(1.f - u1, exponent); const float sinTheta = sqrt(fmax(0.f, 1.f - cosTheta * cosTheta)); const float x = cos(phi) * sinTheta; const float y = sin(phi) * sinTheta; const float z = cosTheta; const float3 dir = -fixedDir; const float dp = dot(shadeN, dir); const float3 w = dir - (2.f * dp) * shadeN; const float3 u = normalize(cross( (fabs(shadeN.x) > .1f) ? ((float3)(0.f, 1.f, 0.f)) : ((float3)(1.f, 0.f, 0.f)), w)); const float3 v = cross(w, u); return x * u + y * v + z * w; } float3 MetalMaterial_Sample(__global Material *material, __global HitPoint *hitPoint, const float3 fixedDir, float3 *sampledDir, const float u0, const float u1, float *pdfW, float *cosSampledDir, BSDFEvent *event TEXTURES_PARAM_DECL) { const float e = 1.f / (Texture_GetFloatValue(&texs[material->metal.expTexIndex], hitPoint TEXTURES_PARAM) + 1.f); *sampledDir = GlossyReflection(fixedDir, e, u0, u1); if ((*sampledDir).z * fixedDir.z > 0.f) { *event = SPECULAR | REFLECT; *pdfW = 1.f; *cosSampledDir = fabs((*sampledDir).z); const float3 kt = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->metal.krTexIndex], hitPoint TEXTURES_PARAM)); // The cosSampledDir is used to compensate the other one used inside the integrator return kt / (*cosSampledDir); } else return BLACK; } #endif //------------------------------------------------------------------------------ // ArchGlass material //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_MAT_ARCHGLASS) float3 ArchGlassMaterial_Sample(__global Material *material, __global HitPoint *hitPoint, const float3 localFixedDir, float3 *localSampledDir, const float u0, const float u1, const float passThroughEvent, float *pdfW, float *absCosSampledDir, BSDFEvent *event TEXTURES_PARAM_DECL) { const float3 kt = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->glass.ktTexIndex], hitPoint TEXTURES_PARAM)); const float3 kr = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->glass.krTexIndex], hitPoint TEXTURES_PARAM)); const bool isKtBlack = Spectrum_IsBlack(kt); const bool isKrBlack = Spectrum_IsBlack(kr); if (isKtBlack && isKrBlack) return BLACK; const bool entering = (CosTheta(localFixedDir) > 0.f); const float nc = Texture_GetFloatValue(&texs[material->glass.ousideIorTexIndex], hitPoint TEXTURES_PARAM); const float nt = Texture_GetFloatValue(&texs[material->glass.iorTexIndex], hitPoint TEXTURES_PARAM); const float ntc = nt / nc; const float eta = nc / nt; const float costheta = CosTheta(localFixedDir); // Decide to transmit or reflect const float threshold = isKrBlack ? 1.f : (isKtBlack ? 0.f : .5f); float3 result; if (passThroughEvent < threshold) { // Transmit // Compute transmitted ray direction const float sini2 = SinTheta2(localFixedDir); const float eta2 = eta * eta; const float sint2 = eta2 * sini2; // Handle total internal reflection for transmission if (sint2 >= 1.f) return BLACK; *localSampledDir = -localFixedDir; *absCosSampledDir = fabs(CosTheta(*localSampledDir)); *event = SPECULAR | TRANSMIT; *pdfW = threshold; //if (!hitPoint.fromLight) { if (entering) result = BLACK; else result = FresnelCauchy_Evaluate(ntc, -costheta); //} else { // if (entering) // result = FresnelCauchy_Evaluate(ntc, costheta); // else // result = BLACK; //} result *= 1.f + (1.f - result) * (1.f - result); result = 1.f - result; result *= kt; } else { // Reflect if (costheta <= 0.f) return BLACK; *localSampledDir = (float3)(-localFixedDir.x, -localFixedDir.y, localFixedDir.z); *absCosSampledDir = fabs(CosTheta(*localSampledDir)); *event = SPECULAR | REFLECT; *pdfW = 1.f - threshold; result = kr * FresnelCauchy_Evaluate(ntc, costheta); } // The absCosSampledDir is used to compensate the other one used inside the integrator return result / (*absCosSampledDir); } float3 ArchGlassMaterial_GetPassThroughTransparency(__global Material *material, __global HitPoint *hitPoint, const float3 localFixedDir, const float passThroughEvent TEXTURES_PARAM_DECL) { const float3 kt = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->glass.ktTexIndex], hitPoint TEXTURES_PARAM)); const float3 kr = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->glass.krTexIndex], hitPoint TEXTURES_PARAM)); const bool isKtBlack = Spectrum_IsBlack(kt); const bool isKrBlack = Spectrum_IsBlack(kr); if (isKtBlack && isKrBlack) return BLACK; const bool entering = (CosTheta(localFixedDir) > 0.f); const float nc = Texture_GetFloatValue(&texs[material->glass.ousideIorTexIndex], hitPoint TEXTURES_PARAM); const float nt = Texture_GetFloatValue(&texs[material->glass.iorTexIndex], hitPoint TEXTURES_PARAM); const float ntc = nt / nc; const float eta = nc / nt; const float costheta = CosTheta(localFixedDir); // Decide to transmit or reflect const float threshold = isKrBlack ? 1.f : (isKtBlack ? 0.f : .5f); if (passThroughEvent < threshold) { // Transmit // Compute transmitted ray direction const float sini2 = SinTheta2(localFixedDir); const float eta2 = eta * eta; const float sint2 = eta2 * sini2; // Handle total internal reflection for transmission if (sint2 >= 1.f) return BLACK; float3 result; //if (!hitPoint.fromLight) { if (entering) result = BLACK; else result = FresnelCauchy_Evaluate(ntc, -costheta); //} else { // if (entering) // result = FresnelCauchy_Evaluate(ntc, costheta); // else // result = Spectrum(); //} result *= 1.f + (1.f - result) * (1.f - result); result = 1.f - result; return kt * result; } else return BLACK; } #endif //------------------------------------------------------------------------------ // NULL material //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_MAT_NULL) float3 NullMaterial_Sample(__global Material *material, __global HitPoint *hitPoint, const float3 fixedDir, float3 *sampledDir, const float u0, const float u1, float *pdfW, float *cosSampledDir, BSDFEvent *event TEXTURES_PARAM_DECL) { *sampledDir = -fixedDir; *cosSampledDir = 1.f; *pdfW = 1.f; *event = SPECULAR | TRANSMIT; return WHITE; } #endif //------------------------------------------------------------------------------ // MatteTranslucent material //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_MAT_MATTETRANSLUCENT) float3 MatteTranslucentMaterial_Evaluate(__global Material *material, __global HitPoint *hitPoint, const float3 lightDir, const float3 eyeDir, BSDFEvent *event, float *directPdfW TEXTURES_PARAM_DECL) { const float cosSampledDir = dot(lightDir, eyeDir); const float3 r = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->matteTranslucent.krTexIndex], hitPoint TEXTURES_PARAM)); const float3 t = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->matteTranslucent.ktTexIndex], hitPoint TEXTURES_PARAM)) * // Energy conservation (1.f - r); if (directPdfW) *directPdfW = .5f * fabs(lightDir.z * M_1_PI_F); if (cosSampledDir > 0.f) { *event = DIFFUSE | REFLECT; return r * M_1_PI_F; } else { *event = DIFFUSE | TRANSMIT; return t * M_1_PI_F; } } float3 MatteTranslucentMaterial_Sample(__global Material *material, __global HitPoint *hitPoint, const float3 fixedDir, float3 *sampledDir, const float u0, const float u1, const float passThroughEvent, float *pdfW, float *cosSampledDir, BSDFEvent *event TEXTURES_PARAM_DECL) { if (fabs(fixedDir.z) < DEFAULT_COS_EPSILON_STATIC) return BLACK; *sampledDir = CosineSampleHemisphereWithPdf(u0, u1, pdfW); *cosSampledDir = fabs((*sampledDir).z); if (*cosSampledDir < DEFAULT_COS_EPSILON_STATIC) return BLACK; *pdfW *= .5f; const float3 r = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->matteTranslucent.krTexIndex], hitPoint TEXTURES_PARAM)); const float3 t = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->matteTranslucent.ktTexIndex], hitPoint TEXTURES_PARAM)) * // Energy conservation (1.f - r); if (passThroughEvent < .5f) { *sampledDir *= (signbit(fixedDir.z) ? -1.f : 1.f); *event = DIFFUSE | REFLECT; return r * M_1_PI_F; } else { *sampledDir *= -(signbit(fixedDir.z) ? -1.f : 1.f); *event = DIFFUSE | TRANSMIT; return t * M_1_PI_F; } } #endif //------------------------------------------------------------------------------ // Glossy2 material // // LuxRender Glossy2 material porting. //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_MAT_GLOSSY2) float SchlickBSDF_CoatingWeight(const float3 ks, const float3 fixedDir) { // No sampling on the back face if (fixedDir.z <= 0.f) return 0.f; // Approximate H by using reflection direction for wi const float u = fabs(fixedDir.z); const float3 S = FresnelSlick_Evaluate(ks, u); // Ensures coating is never sampled less than half the time // unless we are on the back face return .5f * (1.f + Spectrum_Filter(S)); } float3 SchlickBSDF_CoatingF(const float3 ks, const float roughness, const float anisotropy, const int multibounce, const float3 fixedDir, const float3 sampledDir) { // No sampling on the back face if (fixedDir.z <= 0.f) return BLACK; const float coso = fabs(fixedDir.z); const float cosi = fabs(sampledDir.z); const float3 wh = normalize(fixedDir + sampledDir); const float3 S = FresnelSlick_Evaluate(ks, fabs(dot(sampledDir, wh))); const float G = SchlickDistribution_G(roughness, fixedDir, sampledDir); // Multibounce - alternative with interreflection in the coating creases float factor = SchlickDistribution_D(roughness, wh, anisotropy) * G; //if (!fromLight) factor = factor / 4.f * coso + (multibounce ? cosi * clamp((1.f - G) / (4.f * coso * cosi), 0.f, 1.f) : 0.f); //else // factor = factor / (4.f * cosi) + // (multibounce ? coso * Clamp((1.f - G) / (4.f * cosi * coso), 0.f, 1.f) : 0.f); // The cosi is used to compensate the other one used inside the integrator factor /= cosi; return factor * S; } float3 SchlickBSDF_CoatingAbsorption(const float cosi, const float coso, const float3 alpha, const float depth) { if (depth > 0.f) { // 1/cosi+1/coso=(cosi+coso)/(cosi*coso) const float depthFactor = depth * (cosi + coso) / (cosi * coso); return Spectrum_Exp(alpha * -depthFactor); } else return WHITE; } float3 SchlickBSDF_CoatingSampleF(const float3 ks, const float roughness, const float anisotropy, const int multibounce, const float3 fixedDir, float3 *sampledDir, float u0, float u1, float *pdf) { // No sampling on the back face if (fixedDir.z <= 0.f) return BLACK; float3 wh; float d, specPdf; SchlickDistribution_SampleH(roughness, anisotropy, u0, u1, &wh, &d, &specPdf); const float cosWH = dot(fixedDir, wh); *sampledDir = 2.f * cosWH * wh - fixedDir; if (((*sampledDir).z < DEFAULT_COS_EPSILON_STATIC) || (fixedDir.z * (*sampledDir).z < 0.f)) return BLACK; const float coso = fabs(fixedDir.z); const float cosi = fabs((*sampledDir).z); *pdf = specPdf / (4.f * cosWH); if (*pdf <= 0.f) return BLACK; float3 S = FresnelSlick_Evaluate(ks, cosWH); const float G = SchlickDistribution_G(roughness, fixedDir, *sampledDir); //CoatingF(sw, *wi, wo, f_); S *= d * G / (4.f * coso) + (multibounce ? cosi * clamp((1.f - G) / (4.f * coso * cosi), 0.f, 1.f) : 0.f); // The cosi is used to compensate the other one used inside the integrator return S / cosi; } float SchlickBSDF_CoatingPdf(const float roughness, const float anisotropy, const float3 fixedDir, const float3 sampledDir) { // No sampling on the back face if (fixedDir.z <= 0.f) return 0.f; const float3 wh = normalize(fixedDir + sampledDir); return SchlickDistribution_Pdf(roughness, wh, anisotropy) / (4.f * fabs(dot(fixedDir, wh))); } float3 Glossy2Material_Evaluate(__global Material *material, __global HitPoint *hitPoint, const float3 lightDir, const float3 eyeDir, BSDFEvent *event, float *directPdfW TEXTURES_PARAM_DECL) { const float3 fixedDir = eyeDir; const float3 sampledDir = lightDir; const float3 baseF = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->glossy2.kdTexIndex], hitPoint TEXTURES_PARAM)) * M_1_PI_F; if (eyeDir.z <= 0.f) { // Back face: no coating if (directPdfW) *directPdfW = fabs(sampledDir.z * M_1_PI_F); *event = DIFFUSE | REFLECT; return baseF; } // Front face: coating+base *event = GLOSSY | REFLECT; float3 ks = Texture_GetSpectrumValue(&texs[material->glossy2.ksTexIndex], hitPoint TEXTURES_PARAM); #if defined(PARAM_ENABLE_MAT_GLOSSY2_INDEX) const float i = Texture_GetFloatValue(&texs[material->glossy2.indexTexIndex], hitPoint TEXTURES_PARAM); if (i > 0.f) { const float ti = (i - 1.f) / (i + 1.f); ks *= ti * ti; } #endif ks = Spectrum_Clamp(ks); const float u = clamp(Texture_GetFloatValue(&texs[material->glossy2.nuTexIndex], hitPoint TEXTURES_PARAM), 6e-3f, 1.f); #if defined(PARAM_ENABLE_MAT_GLOSSY2_ANISOTROPIC) const float v = clamp(Texture_GetFloatValue(&texs[material->glossy2.nvTexIndex], hitPoint TEXTURES_PARAM), 6e-3f, 1.f); const float u2 = u * u; const float v2 = v * v; const float anisotropy = (u2 < v2) ? (1.f - u2 / v2) : (v2 / u2 - 1.f); const float roughness = u * v; #else const float anisotropy = 0.f; const float roughness = u * u; #endif if (directPdfW) { const float wCoating = SchlickBSDF_CoatingWeight(ks, fixedDir); const float wBase = 1.f - wCoating; *directPdfW = wBase * fabs(sampledDir.z * M_1_PI_F) + wCoating * SchlickBSDF_CoatingPdf(roughness, anisotropy, fixedDir, sampledDir); } // Absorption const float cosi = fabs(sampledDir.z); const float coso = fabs(fixedDir.z); #if defined(PARAM_ENABLE_MAT_GLOSSY2_ANISOTROPIC) const float3 alpha = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->glossy2.kaTexIndex], hitPoint TEXTURES_PARAM)); const float d = Texture_GetFloatValue(&texs[material->glossy2.depthTexIndex], hitPoint TEXTURES_PARAM); const float3 absorption = SchlickBSDF_CoatingAbsorption(cosi, coso, alpha, d); #else const float3 absorption = WHITE; #endif // Coating fresnel factor const float3 H = normalize(fixedDir + sampledDir); const float3 S = FresnelSlick_Evaluate(ks, fabs(dot(sampledDir, H))); #if defined(PARAM_ENABLE_MAT_GLOSSY2_MULTIBOUNCE) const int multibounce = material->glossy2.multibounce; #else const int multibounce = 0; #endif const float3 coatingF = SchlickBSDF_CoatingF(ks, roughness, anisotropy, multibounce, fixedDir, sampledDir); // Blend in base layer Schlick style // assumes coating bxdf takes fresnel factor S into account return coatingF + absorption * (WHITE - S) * baseF; } float3 Glossy2Material_Sample(__global Material *material, __global HitPoint *hitPoint, const float3 fixedDir, float3 *sampledDir, const float u0, const float u1, const float passThroughEvent, float *pdfW, float *cosSampledDir, BSDFEvent *event TEXTURES_PARAM_DECL) { if (fabs(fixedDir.z) < DEFAULT_COS_EPSILON_STATIC) return BLACK; float3 ks = Texture_GetSpectrumValue(&texs[material->glossy2.ksTexIndex], hitPoint TEXTURES_PARAM); #if defined(PARAM_ENABLE_MAT_GLOSSY2_INDEX) const float i = Texture_GetFloatValue(&texs[material->glossy2.indexTexIndex], hitPoint TEXTURES_PARAM); if (i > 0.f) { const float ti = (i - 1.f) / (i + 1.f); ks *= ti * ti; } #endif ks = Spectrum_Clamp(ks); const float u = clamp(Texture_GetFloatValue(&texs[material->glossy2.nuTexIndex], hitPoint TEXTURES_PARAM), 6e-3f, 1.f); #if defined(PARAM_ENABLE_MAT_GLOSSY2_ANISOTROPIC) const float v = clamp(Texture_GetFloatValue(&texs[material->glossy2.nvTexIndex], hitPoint TEXTURES_PARAM), 6e-3f, 1.f); const float u2 = u * u; const float v2 = v * v; const float anisotropy = (u2 < v2) ? (1.f - u2 / v2) : (v2 / u2 - 1.f); const float roughness = u * v; #else const float anisotropy = 0.f; const float roughness = u * u; #endif // Coating is used only on the front face const float wCoating = (fixedDir.z <= 0.f) ? 0.f : SchlickBSDF_CoatingWeight(ks, fixedDir); const float wBase = 1.f - wCoating; const float3 baseF = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->glossy2.kdTexIndex], hitPoint TEXTURES_PARAM)) * M_1_PI_F; #if defined(PARAM_ENABLE_MAT_GLOSSY2_MULTIBOUNCE) const int multibounce = material->glossy2.multibounce; #else const int multibounce = 0; #endif float basePdf, coatingPdf; float3 coatingF; if (passThroughEvent < wBase) { // Sample base BSDF (Matte BSDF) *sampledDir = (signbit(fixedDir.z) ? -1.f : 1.f) * CosineSampleHemisphereWithPdf(u0, u1, &basePdf); *cosSampledDir = fabs((*sampledDir).z); if (*cosSampledDir < DEFAULT_COS_EPSILON_STATIC) return BLACK; // Evaluate coating BSDF (Schlick BSDF) coatingF = SchlickBSDF_CoatingF(ks, roughness, anisotropy, multibounce, fixedDir, *sampledDir); coatingPdf = SchlickBSDF_CoatingPdf(roughness, anisotropy, fixedDir, *sampledDir); *event = DIFFUSE | REFLECT; } else { // Sample coating BSDF (Schlick BSDF) coatingF = SchlickBSDF_CoatingSampleF(ks, roughness, anisotropy, multibounce, fixedDir, sampledDir, u0, u1, &coatingPdf); if (Spectrum_IsBlack(coatingF)) return BLACK; *cosSampledDir = fabs((*sampledDir).z); if (*cosSampledDir < DEFAULT_COS_EPSILON_STATIC) return BLACK; // Evaluate base BSDF (Matte BSDF) basePdf = fabs((*sampledDir).z * M_1_PI_F); *event = GLOSSY | REFLECT; } *pdfW = coatingPdf * wCoating + basePdf * wBase; if (fixedDir.z > 0.f) { // Front face reflection: coating+base // Absorption const float cosi = fabs((*sampledDir).z); const float coso = fabs(fixedDir.z); #if defined(PARAM_ENABLE_MAT_GLOSSY2_ANISOTROPIC) const float3 alpha = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->glossy2.kaTexIndex], hitPoint TEXTURES_PARAM)); const float d = Texture_GetFloatValue(&texs[material->glossy2.depthTexIndex], hitPoint TEXTURES_PARAM); const float3 absorption = SchlickBSDF_CoatingAbsorption(cosi, coso, alpha, d); #else const float3 absorption = WHITE; #endif // Coating fresnel factor const float3 H = normalize(fixedDir + *sampledDir); const float3 S = FresnelSlick_Evaluate(ks, fabs(dot(*sampledDir, H))); // Blend in base layer Schlick style // coatingF already takes fresnel factor S into account return coatingF + absorption * (WHITE - S) * baseF; } else { // Back face reflection: base return baseF; } } #endif //------------------------------------------------------------------------------ // Metal2 material // // LuxRender Metal2 material porting. //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_MAT_METAL2) float3 Metal2Material_Evaluate(__global Material *material, __global HitPoint *hitPoint, const float3 lightDir, const float3 eyeDir, BSDFEvent *event, float *directPdfW TEXTURES_PARAM_DECL) { const float3 fixedDir = eyeDir; const float3 sampledDir = lightDir; const float u = clamp(Texture_GetFloatValue(&texs[material->metal2.nuTexIndex], hitPoint TEXTURES_PARAM), 6e-3f, 1.f); #if defined(PARAM_ENABLE_MAT_METAL2_ANISOTROPIC) const float v = clamp(Texture_GetFloatValue(&texs[material->metal2.nvTexIndex], hitPoint TEXTURES_PARAM), 6e-3f, 1.f); const float u2 = u * u; const float v2 = v * v; const float anisotropy = (u2 < v2) ? (1.f - u2 / v2) : (v2 / u2 - 1.f); const float roughness = u * v; #else const float anisotropy = 0.f; const float roughness = u * u; #endif const float3 wh = normalize(fixedDir + sampledDir); const float cosWH = dot(fixedDir, wh); if (directPdfW) *directPdfW = SchlickDistribution_Pdf(roughness, wh, anisotropy) / (4.f * fabs(dot(fixedDir, wh))); const float3 nVal = Texture_GetSpectrumValue(&texs[material->metal2.nTexIndex], hitPoint TEXTURES_PARAM); const float3 kVal = Texture_GetSpectrumValue(&texs[material->metal2.kTexIndex], hitPoint TEXTURES_PARAM); const float3 F = FresnelGeneral_Evaluate(nVal, kVal, cosWH); const float G = SchlickDistribution_G(roughness, fixedDir, sampledDir); const float coso = fabs(fixedDir.z); const float cosi = fabs(sampledDir.z); float factor = SchlickDistribution_D(roughness, wh, anisotropy) * G; //if (!hitPoint.fromLight) factor /= 4.f * coso; //else // factor /= 4.f * cosi; *event = GLOSSY | REFLECT; // The cosi is used to compensate the other one used inside the integrator factor /= cosi; return factor * F; } float3 Metal2Material_Sample(__global Material *material, __global HitPoint *hitPoint, const float3 fixedDir, float3 *sampledDir, const float u0, const float u1, float *pdfW, float *cosSampledDir, BSDFEvent *event TEXTURES_PARAM_DECL) { if (fabs(fixedDir.z) < DEFAULT_COS_EPSILON_STATIC) return BLACK; const float u = clamp(Texture_GetFloatValue(&texs[material->metal2.nuTexIndex], hitPoint TEXTURES_PARAM), 6e-3f, 1.f); #if defined(PARAM_ENABLE_MAT_METAL2_ANISOTROPIC) const float v = clamp(Texture_GetFloatValue(&texs[material->metal2.nvTexIndex], hitPoint TEXTURES_PARAM), 6e-3f, 1.f); const float u2 = u * u; const float v2 = v * v; const float anisotropy = (u2 < v2) ? (1.f - u2 / v2) : (v2 / u2 - 1.f); const float roughness = u * v; #else const float anisotropy = 0.f; const float roughness = u * u; #endif float3 wh; float d, specPdf; SchlickDistribution_SampleH(roughness, anisotropy, u0, u1, &wh, &d, &specPdf); const float cosWH = dot(fixedDir, wh); *sampledDir = 2.f * cosWH * wh - fixedDir; const float coso = fabs(fixedDir.z); const float cosi = fabs((*sampledDir).z); *cosSampledDir = cosi; if ((*cosSampledDir < DEFAULT_COS_EPSILON_STATIC) || (fixedDir.z * (*sampledDir).z < 0.f)) return BLACK; *pdfW = specPdf / (4.f * cosWH); if (*pdfW <= 0.f) return BLACK; const float G = SchlickDistribution_G(roughness, fixedDir, *sampledDir); const float3 nVal = Texture_GetSpectrumValue(&texs[material->metal2.nTexIndex], hitPoint TEXTURES_PARAM); const float3 kVal = Texture_GetSpectrumValue(&texs[material->metal2.kTexIndex], hitPoint TEXTURES_PARAM); const float3 F = FresnelGeneral_Evaluate(nVal, kVal, cosWH); float factor = d * G; //if (!fromLight) factor /= 4.f * coso; //else // factor /= cosi; *event = GLOSSY | REFLECT; // The cosi is used to compensate the other one used inside the integrator factor /= cosi; return factor * F; } #endif //------------------------------------------------------------------------------ // RoughGlass material //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_MAT_ROUGHGLASS) float3 RoughGlassMaterial_Evaluate(__global Material *material, __global HitPoint *hitPoint, const float3 localLightDir, const float3 localEyeDir, BSDFEvent *event, float *directPdfW TEXTURES_PARAM_DECL) { const float3 kt = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->roughglass.ktTexIndex], hitPoint TEXTURES_PARAM)); const float3 kr = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->roughglass.krTexIndex], hitPoint TEXTURES_PARAM)); const bool isKtBlack = Spectrum_IsBlack(kt); const bool isKrBlack = Spectrum_IsBlack(kr); if (isKtBlack && isKrBlack) return BLACK; //const float3 localFixedDir = hitPoint.fromLight ? localLightDir : localEyeDir; //const float3 localSampledDir = hitPoint.fromLight ? localEyeDir : localLightDir; const float3 localFixedDir = localEyeDir; const float3 localSampledDir = localLightDir; const float nc = Texture_GetFloatValue(&texs[material->roughglass.ousideIorTexIndex], hitPoint TEXTURES_PARAM); const float nt = Texture_GetFloatValue(&texs[material->roughglass.iorTexIndex], hitPoint TEXTURES_PARAM); const float ntc = nt / nc; const float u = clamp(Texture_GetFloatValue(&texs[material->roughglass.nuTexIndex], hitPoint TEXTURES_PARAM), 6e-3f, 1.f); #if defined(PARAM_ENABLE_MAT_ROUGHGLASS_ANISOTROPIC) const float v = clamp(Texture_GetFloatValue(&texs[material->roughglass.nvTexIndex], hitPoint TEXTURES_PARAM), 6e-3f, 1.f); const float u2 = u * u; const float v2 = v * v; const float anisotropy = (u2 < v2) ? (1.f - u2 / v2) : (v2 / u2 - 1.f); const float roughness = u * v; #else const float anisotropy = 0.f; const float roughness = u * u; #endif const float threshold = isKrBlack ? 1.f : (isKtBlack ? 0.f : .5f); if (dot(localFixedDir, localSampledDir) < 0.f) { // Transmit const bool entering = (CosTheta(localFixedDir) > 0.f); const float eta = entering ? (nc / nt) : ntc; float3 wh = eta * localFixedDir + localSampledDir; if (wh.z < 0.f) wh = -wh; //const float lengthSquared = wh.LengthSquared(); const float lengthSquared = dot(wh, wh); if (!(lengthSquared > 0.f)) return BLACK; wh /= sqrt(lengthSquared); const float cosThetaI = fabs(CosTheta(localSampledDir)); const float cosThetaIH = fabs(dot(localSampledDir, wh)); const float cosThetaOH = dot(localFixedDir, wh); const float D = SchlickDistribution_D(roughness, wh, anisotropy); const float G = SchlickDistribution_G(roughness, localFixedDir, localSampledDir); const float specPdf = SchlickDistribution_Pdf(roughness, wh, anisotropy); const float3 F = FresnelCauchy_Evaluate(ntc, cosThetaOH); if (directPdfW) *directPdfW = threshold * specPdf * fabs(cosThetaOH) / lengthSquared; //if (reversePdfW) // *reversePdfW = threshold * specPdf * cosThetaIH * eta * eta / lengthSquared; float3 result = (fabs(cosThetaOH) * cosThetaIH * D * G / (cosThetaI * lengthSquared)) * kt * (1.f - F); // This is a porting of LuxRender code and there, the result is multiplied // by Dot(ns, wl). So I have to remove that term. result /= fabs(CosTheta(localLightDir)); return result; } else { // Reflect const float cosThetaO = fabs(CosTheta(localFixedDir)); const float cosThetaI = fabs(CosTheta(localSampledDir)); if (cosThetaO == 0.f || cosThetaI == 0.f) return BLACK; float3 wh = localFixedDir + localSampledDir; if (all(isequal(wh, BLACK))) return BLACK; wh = normalize(wh); if (wh.z < 0.f) wh = -wh; float cosThetaH = dot(localEyeDir, wh); const float D = SchlickDistribution_D(roughness, wh, anisotropy); const float G = SchlickDistribution_G(roughness, localFixedDir, localSampledDir); const float specPdf = SchlickDistribution_Pdf(roughness, wh, anisotropy); const float3 F = FresnelCauchy_Evaluate(ntc, cosThetaH); if (directPdfW) *directPdfW = (1.f - threshold) * specPdf / (4.f * fabs(dot(localFixedDir, wh))); //if (reversePdfW) // *reversePdfW = (1.f - threshold) * specPdf / (4.f * fabs(dot(localSampledDir, wh)); float3 result = (D * G / (4.f * cosThetaI)) * kr * F; // This is a porting of LuxRender code and there, the result is multiplied // by Dot(ns, wl). So I have to remove that term. result /= fabs(CosTheta(localLightDir)); return result; } } float3 RoughGlassMaterial_Sample(__global Material *material, __global HitPoint *hitPoint, const float3 localFixedDir, float3 *localSampledDir, const float u0, const float u1, const float passThroughEvent, float *pdfW, float *absCosSampledDir, BSDFEvent *event TEXTURES_PARAM_DECL) { if (fabs(localFixedDir.z) < DEFAULT_COS_EPSILON_STATIC) return BLACK; const float3 kt = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->roughglass.ktTexIndex], hitPoint TEXTURES_PARAM)); const float3 kr = Spectrum_Clamp(Texture_GetSpectrumValue(&texs[material->roughglass.krTexIndex], hitPoint TEXTURES_PARAM)); const bool isKtBlack = Spectrum_IsBlack(kt); const bool isKrBlack = Spectrum_IsBlack(kr); if (isKtBlack && isKrBlack) return BLACK; const float u = clamp(Texture_GetFloatValue(&texs[material->roughglass.nuTexIndex], hitPoint TEXTURES_PARAM), 6e-3f, 1.f); #if defined(PARAM_ENABLE_MAT_ROUGHGLASS_ANISOTROPIC) const float v = clamp(Texture_GetFloatValue(&texs[material->roughglass.nvTexIndex], hitPoint TEXTURES_PARAM), 6e-3f, 1.f); const float u2 = u * u; const float v2 = v * v; const float anisotropy = (u2 < v2) ? (1.f - u2 / v2) : (v2 / u2 - 1.f); const float roughness = u * v; #else const float anisotropy = 0.f; const float roughness = u * u; #endif float3 wh; float d, specPdf; SchlickDistribution_SampleH(roughness, anisotropy, u0, u1, &wh, &d, &specPdf); if (wh.z < 0.f) wh = -wh; const float cosThetaOH = dot(localFixedDir, wh); const float nc = Texture_GetFloatValue(&texs[material->roughglass.ousideIorTexIndex], hitPoint TEXTURES_PARAM); const float nt = Texture_GetFloatValue(&texs[material->roughglass.iorTexIndex], hitPoint TEXTURES_PARAM); const float ntc = nt / nc; const float coso = fabs(localFixedDir.z); // Decide to transmit or reflect const float threshold = isKrBlack ? 1.f : (isKtBlack ? 0.f : .5f); float3 result; if (passThroughEvent < threshold) { // Transmit const bool entering = (CosTheta(localFixedDir) > 0.f); const float eta = entering ? (nc / nt) : ntc; const float eta2 = eta * eta; const float sinThetaIH2 = eta2 * fmax(0.f, 1.f - cosThetaOH * cosThetaOH); if (sinThetaIH2 >= 1.f) return BLACK; float cosThetaIH = sqrt(1.f - sinThetaIH2); if (entering) cosThetaIH = -cosThetaIH; const float length = eta * cosThetaOH + cosThetaIH; *localSampledDir = length * wh - eta * localFixedDir; const float lengthSquared = length * length; *pdfW = specPdf * fabs(cosThetaIH) / lengthSquared; if (*pdfW <= 0.f) return BLACK; const float cosi = fabs((*localSampledDir).z); *absCosSampledDir = cosi; const float G = SchlickDistribution_G(roughness, localFixedDir, *localSampledDir); float factor = d * G * fabs(cosThetaOH) / specPdf; //if (!hitPoint.fromLight) { const float3 F = FresnelCauchy_Evaluate(ntc, cosThetaIH); result = (factor / coso) * kt * (1.f - F); //} else { // const Spectrum F = FresnelCauchy_Evaluate(ntc, cosThetaOH); // result = (factor / cosi) * kt * (Spectrum(1.f) - F); //} // This is a porting of LuxRender code and there, the result is multiplied // by Dot(ns, wi)/pdf if reverse==true and by Dot(ns. wo)/pdf if reverse==false. // So I have to remove that terms. //result *= *pdfW / ((!hitPoint.fromLight) ? cosi : coso); result *= *pdfW / cosi; *pdfW *= threshold; *event = GLOSSY | TRANSMIT; } else { // Reflect *pdfW = specPdf / (4.f * fabs(cosThetaOH)); if (*pdfW <= 0.f) return BLACK; *localSampledDir = 2.f * cosThetaOH * wh - localFixedDir; const float cosi = fabs((*localSampledDir).z); *absCosSampledDir = cosi; if ((*absCosSampledDir < DEFAULT_COS_EPSILON_STATIC) || (localFixedDir.z * (*localSampledDir).z < 0.f)) return BLACK; const float G = SchlickDistribution_G(roughness, localFixedDir, *localSampledDir); float factor = d * G * fabs(cosThetaOH) / specPdf; const float3 F = FresnelCauchy_Evaluate(ntc, cosThetaOH); //factor /= (!hitPoint.fromLight) ? coso : cosi; factor /= coso; result = factor * F * kr; // This is a porting of LuxRender code and there, the result is multiplied // by Dot(ns, wi)/pdf if reverse==true and by Dot(ns. wo)/pdf if reverse==false. // So I have to remove that terms. //result *= *pdfW / ((!hitPoint.fromLight) ? cosi : coso); result *= *pdfW / cosi; *pdfW *= (1.f - threshold); *event = GLOSSY | REFLECT; } return result; } #endif #line 2 "material_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ //------------------------------------------------------------------------------ // Generic material functions // // They include the support for all material but Mix // (because OpenCL doesn't support recursion) //------------------------------------------------------------------------------ bool Material_IsDeltaNoMix(__global Material *material) { switch (material->type) { // Non Specular materials #if defined (PARAM_ENABLE_MAT_ROUGHGLASS) case ROUGHGLASS: #endif #if defined (PARAM_ENABLE_MAT_METAL2) case METAL2: #endif #if defined (PARAM_ENABLE_MAT_GLOSSY2) case GLOSSY2: #endif #if defined (PARAM_ENABLE_MAT_MATTETRANSLUCENT) case MATTETRANSLUCENT: #endif #if defined (PARAM_ENABLE_MAT_MATTE) case MATTE: return false; #endif // Specular materials #if defined (PARAM_ENABLE_MAT_MIRROR) case MIRROR: #endif #if defined (PARAM_ENABLE_MAT_GLASS) case GLASS: #endif #if defined (PARAM_ENABLE_MAT_METAL) case METAL: #endif #if defined (PARAM_ENABLE_MAT_ARCHGLASS) case ARCHGLASS: #endif #if defined (PARAM_ENABLE_MAT_NULL) case NULLMAT: #endif default: return true; } } BSDFEvent Material_GetEventTypesNoMix(__global Material *mat) { switch (mat->type) { #if defined (PARAM_ENABLE_MAT_MATTE) case MATTE: return DIFFUSE | REFLECT; #endif #if defined (PARAM_ENABLE_MAT_MIRROR) case MIRROR: return SPECULAR | REFLECT; #endif #if defined (PARAM_ENABLE_MAT_GLASS) case GLASS: return SPECULAR | REFLECT | TRANSMIT; #endif #if defined (PARAM_ENABLE_MAT_METAL) case METAL: return SPECULAR | REFLECT; #endif #if defined (PARAM_ENABLE_MAT_ARCHGLASS) case ARCHGLASS: return SPECULAR | REFLECT | TRANSMIT; #endif #if defined (PARAM_ENABLE_MAT_NULL) case NULLMAT: return SPECULAR | TRANSMIT; #endif #if defined (PARAM_ENABLE_MAT_MATTETRANSLUCENT) case MATTETRANSLUCENT: return DIFFUSE | REFLECT | TRANSMIT; #endif #if defined (PARAM_ENABLE_MAT_GLOSSY2) case GLOSSY2: return DIFFUSE | GLOSSY | REFLECT; #endif #if defined (PARAM_ENABLE_MAT_METAL2) case METAL2: return GLOSSY | REFLECT; #endif #if defined (PARAM_ENABLE_MAT_ROUGHGLASS) case ROUGHGLASS: return GLOSSY | REFLECT | TRANSMIT; #endif default: return NONE; } } float3 Material_SampleNoMix(__global Material *material, __global HitPoint *hitPoint, const float3 fixedDir, float3 *sampledDir, const float u0, const float u1, #if defined(PARAM_HAS_PASSTHROUGH) const float passThroughEvent, #endif float *pdfW, float *cosSampledDir, BSDFEvent *event TEXTURES_PARAM_DECL) { switch (material->type) { #if defined (PARAM_ENABLE_MAT_MATTE) case MATTE: return MatteMaterial_Sample(material, hitPoint, fixedDir, sampledDir, u0, u1, pdfW, cosSampledDir, event TEXTURES_PARAM); #endif #if defined (PARAM_ENABLE_MAT_MIRROR) case MIRROR: return MirrorMaterial_Sample(material, hitPoint, fixedDir, sampledDir, u0, u1, pdfW, cosSampledDir, event TEXTURES_PARAM); #endif #if defined (PARAM_ENABLE_MAT_GLASS) case GLASS: return GlassMaterial_Sample(material, hitPoint, fixedDir, sampledDir, u0, u1, passThroughEvent, pdfW, cosSampledDir, event TEXTURES_PARAM); #endif #if defined (PARAM_ENABLE_MAT_METAL) case METAL: return MetalMaterial_Sample(material, hitPoint, fixedDir, sampledDir, u0, u1, pdfW, cosSampledDir, event TEXTURES_PARAM); #endif #if defined (PARAM_ENABLE_MAT_ARCHGLASS) case ARCHGLASS: return ArchGlassMaterial_Sample(material, hitPoint, fixedDir, sampledDir, u0, u1, passThroughEvent, pdfW, cosSampledDir, event TEXTURES_PARAM); #endif #if defined (PARAM_ENABLE_MAT_NULL) case NULLMAT: return NullMaterial_Sample(material, hitPoint, fixedDir, sampledDir, u0, u1, pdfW, cosSampledDir, event TEXTURES_PARAM); #endif #if defined (PARAM_ENABLE_MAT_MATTETRANSLUCENT) case MATTETRANSLUCENT: return MatteTranslucentMaterial_Sample(material, hitPoint, fixedDir, sampledDir, u0, u1, passThroughEvent, pdfW, cosSampledDir, event TEXTURES_PARAM); #endif #if defined (PARAM_ENABLE_MAT_GLOSSY2) case GLOSSY2: return Glossy2Material_Sample(material, hitPoint, fixedDir, sampledDir, u0, u1, passThroughEvent, pdfW, cosSampledDir, event TEXTURES_PARAM); #endif #if defined (PARAM_ENABLE_MAT_METAL2) case METAL2: return Metal2Material_Sample(material, hitPoint, fixedDir, sampledDir, u0, u1, pdfW, cosSampledDir, event TEXTURES_PARAM); #endif #if defined (PARAM_ENABLE_MAT_ROUGHGLASS) case ROUGHGLASS: return RoughGlassMaterial_Sample(material, hitPoint, fixedDir, sampledDir, u0, u1, passThroughEvent, pdfW, cosSampledDir, event TEXTURES_PARAM); #endif default: return BLACK; } } float3 Material_EvaluateNoMix(__global Material *material, __global HitPoint *hitPoint, const float3 lightDir, const float3 eyeDir, BSDFEvent *event, float *directPdfW TEXTURES_PARAM_DECL) { switch (material->type) { #if defined (PARAM_ENABLE_MAT_MATTE) case MATTE: return MatteMaterial_Evaluate(material, hitPoint, lightDir, eyeDir, event, directPdfW TEXTURES_PARAM); #endif #if defined (PARAM_ENABLE_MAT_MATTETRANSLUCENT) case MATTETRANSLUCENT: return MatteTranslucentMaterial_Evaluate(material, hitPoint, lightDir, eyeDir, event, directPdfW TEXTURES_PARAM); #endif #if defined (PARAM_ENABLE_MAT_GLOSSY2) case GLOSSY2: return Glossy2Material_Evaluate(material, hitPoint, lightDir, eyeDir, event, directPdfW TEXTURES_PARAM); #endif #if defined (PARAM_ENABLE_MAT_METAL2) case METAL2: return Metal2Material_Evaluate(material, hitPoint, lightDir, eyeDir, event, directPdfW TEXTURES_PARAM); #endif #if defined (PARAM_ENABLE_MAT_ROUGHGLASS) case ROUGHGLASS: return RoughGlassMaterial_Evaluate(material, hitPoint, lightDir, eyeDir, event, directPdfW TEXTURES_PARAM); #endif #if defined (PARAM_ENABLE_MAT_MIRROR) case MIRROR: #endif #if defined (PARAM_ENABLE_MAT_GLASS) case GLASS: #endif #if defined (PARAM_ENABLE_MAT_METAL) case METAL: #endif #if defined (PARAM_ENABLE_MAT_ARCHGLASS) case ARCHGLASS: #endif #if defined (PARAM_ENABLE_MAT_NULL) case NULLMAT: #endif default: return BLACK; } } float3 Material_GetEmittedRadianceNoMix(__global Material *material, __global HitPoint *hitPoint TEXTURES_PARAM_DECL) { const uint emitTexIndex = material->emitTexIndex; if (emitTexIndex == NULL_INDEX) return BLACK; return Texture_GetSpectrumValue(&texs[emitTexIndex], hitPoint TEXTURES_PARAM); } #if defined(PARAM_HAS_BUMPMAPS) float2 Material_GetBumpTexValueNoMix(__global Material *material, __global HitPoint *hitPoint TEXTURES_PARAM_DECL) { const uint bumpTexIndex = material->bumpTexIndex; if (bumpTexIndex == NULL_INDEX) return (float2)(0.f, 0.f); __global Texture *tex = &texs[bumpTexIndex]; const float2 dudv = Texture_GetDuDv(tex, hitPoint TEXTURES_PARAM); const float b0 = Texture_GetFloatValue(tex, hitPoint TEXTURES_PARAM); const float3 hitPointP = VLOAD3F(&hitPoint->p.x); const float2 hitPointUV = VLOAD2F(&hitPoint->uv.u); // Placing the the following line here as a workaround for an AMD OpenCL compiler bug VSTORE2F((float2)(hitPointUV.s0 + dudv.s0, hitPointUV.s1), &hitPoint->uv.u); float dbdu; if (dudv.s0 > 0.f) { // This is a simple trick. The correct code would require true differential information. VSTORE3F((float3)(hitPointP.x + dudv.s0, hitPointP.y, hitPointP.z), &hitPoint->p.x); //VSTORE2F((float2)(hitPointUV.s0 + dudv.s0, hitPointUV.s1), &hitPoint->uv.u); const float bu = Texture_GetFloatValue(tex, hitPoint TEXTURES_PARAM); dbdu = (bu - b0) / dudv.s0; } else dbdu = 0.f; // Placing the the following line here as a workaround for an AMD OpenCL compiler bug VSTORE2F((float2)(hitPointUV.s0, hitPointUV.s1 + dudv.s1), &hitPoint->uv.u); float dbdv; if (dudv.s1 > 0.f) { // This is a simple trick. The correct code would require true differential information. VSTORE3F((float3)(hitPointP.x, hitPointP.y + dudv.s1, hitPointP.z), &hitPoint->p.x); //VSTORE2F((float2)(hitPointUV.s0, hitPointUV.s1 + dudv.s1), &hitPoint->uv.u); const float bv = Texture_GetFloatValue(tex, hitPoint TEXTURES_PARAM); dbdv = (bv - b0) / dudv.s1; } else dbdv = 0.f; // Restore p and uv value VSTORE3F(hitPointP, &hitPoint->p.x); VSTORE2F(hitPointUV, &hitPoint->uv.u); return (float2)(dbdu, dbdv); } #endif #if defined(PARAM_HAS_NORMALMAPS) float3 Material_GetNormalTexValueNoMix(__global Material *material, __global HitPoint *hitPoint TEXTURES_PARAM_DECL) { const uint normalTexIndex = material->normalTexIndex; if (normalTexIndex == NULL_INDEX) return BLACK; return Texture_GetSpectrumValue(&texs[normalTexIndex], hitPoint TEXTURES_PARAM); } #endif float3 Material_GetPassThroughTransparencyNoMix(__global Material *material, __global HitPoint *hitPoint, const float3 fixedDir, const float passThroughEvent TEXTURES_PARAM_DECL) { switch (material->type) { #if defined (PARAM_ENABLE_MAT_ARCHGLASS) case ARCHGLASS: return ArchGlassMaterial_GetPassThroughTransparency(material, hitPoint, fixedDir, passThroughEvent TEXTURES_PARAM); #endif #if defined (PARAM_ENABLE_MAT_NULL) case NULLMAT: return WHITE; #endif default: return BLACK; } } //------------------------------------------------------------------------------ // Mix material // // This requires a quite complex implementation because OpenCL doesn't support // recursion. //------------------------------------------------------------------------------ #if defined (PARAM_ENABLE_MAT_MIX) #define MIX_STACK_SIZE 16 BSDFEvent MixMaterial_IsDelta(__global Material *material MATERIALS_PARAM_DECL) { __global Material *materialStack[MIX_STACK_SIZE]; materialStack[0] = material; int stackIndex = 0; while (stackIndex >= 0) { // Extract a material from the stack __global Material *m = materialStack[stackIndex--]; // Check if it is a Mix material too if (m->type == MIX) { // Add both material to the stack materialStack[++stackIndex] = &mats[m->mix.matAIndex]; materialStack[++stackIndex] = &mats[m->mix.matBIndex]; } else { // Normal GetEventTypes() evaluation if (!Material_IsDeltaNoMix(m)) return false; } } return true; } BSDFEvent MixMaterial_GetEventTypes(__global Material *material MATERIALS_PARAM_DECL) { BSDFEvent event = NONE; __global Material *materialStack[MIX_STACK_SIZE]; materialStack[0] = material; int stackIndex = 0; while (stackIndex >= 0) { // Extract a material from the stack __global Material *m = materialStack[stackIndex--]; // Check if it is a Mix material too if (m->type == MIX) { // Add both material to the stack materialStack[++stackIndex] = &mats[m->mix.matAIndex]; materialStack[++stackIndex] = &mats[m->mix.matBIndex]; } else { // Normal GetEventTypes() evaluation event |= Material_GetEventTypesNoMix(m); } } return event; } float3 MixMaterial_Evaluate(__global Material *material, __global HitPoint *hitPoint, const float3 lightDir, const float3 eyeDir, BSDFEvent *event, float *directPdfW MATERIALS_PARAM_DECL) { __global Material *materialStack[MIX_STACK_SIZE]; float totalWeightStack[MIX_STACK_SIZE]; // Push the root Mix material materialStack[0] = material; totalWeightStack[0] = 1.f; int stackIndex = 0; // Setup the results float3 result = BLACK; *event = NONE; if (directPdfW) *directPdfW = 0.f; while (stackIndex >= 0) { // Extract a material from the stack __global Material *m = materialStack[stackIndex]; float totalWeight = totalWeightStack[stackIndex--]; // Check if it is a Mix material too if (m->type == MIX) { // Add both material to the stack const float factor = Texture_GetFloatValue(&texs[m->mix.mixFactorTexIndex], hitPoint TEXTURES_PARAM); const float weight2 = clamp(factor, 0.f, 1.f); const float weight1 = 1.f - weight2; materialStack[++stackIndex] = &mats[m->mix.matAIndex]; totalWeightStack[stackIndex] = totalWeight * weight1; materialStack[++stackIndex] = &mats[m->mix.matBIndex]; totalWeightStack[stackIndex] = totalWeight * weight2; } else { // Normal Evaluate() evaluation if (totalWeight > 0.f) { BSDFEvent eventMat; float directPdfWMat; const float3 resultMat = Material_EvaluateNoMix(m, hitPoint, lightDir, eyeDir, &eventMat, &directPdfWMat TEXTURES_PARAM); if (!Spectrum_IsBlack(resultMat)) { *event |= eventMat; result += totalWeight * resultMat; if (directPdfW) *directPdfW += totalWeight * directPdfWMat; } } } } return result; } float3 MixMaterial_Sample(__global Material *material, __global HitPoint *hitPoint, const float3 fixedDir, float3 *sampledDir, const float u0, const float u1, const float passEvent, float *pdfW, float *cosSampledDir, BSDFEvent *event MATERIALS_PARAM_DECL) { __global Material *evaluationMatList[MIX_STACK_SIZE]; float parentWeightList[MIX_STACK_SIZE]; int evaluationListSize = 0; // Setup the results float3 result = BLACK; *pdfW = 0.f; // Look for a no Mix material to sample __global Material *currentMixMat = material; float passThroughEvent = passEvent; float parentWeight = 1.f; for (;;) { const float factor = Texture_GetFloatValue(&texs[currentMixMat->mix.mixFactorTexIndex], hitPoint TEXTURES_PARAM); const float weight2 = clamp(factor, 0.f, 1.f); const float weight1 = 1.f - weight2; const bool sampleMatA = (passThroughEvent < weight1); const float weightFirst = sampleMatA ? weight1 : weight2; const float weightSecond = sampleMatA ? weight2 : weight1; const float passThroughEventFirst = sampleMatA ? (passThroughEvent / weight1) : (passThroughEvent - weight1) / weight2; const uint matIndexFirst = sampleMatA ? currentMixMat->mix.matAIndex : currentMixMat->mix.matBIndex; const uint matIndexSecond = sampleMatA ? currentMixMat->mix.matBIndex : currentMixMat->mix.matAIndex; // Sample the first material, evaluate the second __global Material *matFirst = &mats[matIndexFirst]; __global Material *matSecond = &mats[matIndexSecond]; //---------------------------------------------------------------------- // Add the second material to the evaluation list //---------------------------------------------------------------------- if (weightSecond > 0.f) { evaluationMatList[evaluationListSize] = matSecond; parentWeightList[evaluationListSize++] = parentWeight * weightSecond; } //---------------------------------------------------------------------- // Sample the first material //---------------------------------------------------------------------- // Check if it is a Mix material too if (matFirst->type == MIX) { // Make the first material the current currentMixMat = matFirst; passThroughEvent = passThroughEventFirst; parentWeight *= weightFirst; } else { // Sample the first material float pdfWMatFirst; const float3 sampleResult = Material_SampleNoMix(matFirst, hitPoint, fixedDir, sampledDir, u0, u1, passThroughEventFirst, &pdfWMatFirst, cosSampledDir, event TEXTURES_PARAM); if (all(isequal(sampleResult, BLACK))) return BLACK; const float weight = parentWeight * weightFirst; *pdfW += weight * pdfWMatFirst; result += weight * sampleResult; // I can stop now break; } } while (evaluationListSize > 0) { // Extract the material to evaluate __global Material *evalMat = evaluationMatList[--evaluationListSize]; const float evalWeight = parentWeightList[evaluationListSize]; // Evaluate the material // Check if it is a Mix material too BSDFEvent eventMat; float pdfWMat; float3 eval; if (evalMat->type == MIX) { eval = MixMaterial_Evaluate(evalMat, hitPoint, *sampledDir, fixedDir, &eventMat, &pdfWMat MATERIALS_PARAM); } else { eval = Material_EvaluateNoMix(evalMat, hitPoint, *sampledDir, fixedDir, &eventMat, &pdfWMat TEXTURES_PARAM); } if (!Spectrum_IsBlack(eval)) { result += evalWeight * eval; *pdfW += evalWeight * pdfWMat; } } return result; } float3 MixMaterial_GetEmittedRadiance(__global Material *material, __global HitPoint *hitPoint MATERIALS_PARAM_DECL) { __global Material *materialStack[MIX_STACK_SIZE]; float totalWeightStack[MIX_STACK_SIZE]; // Push the root Mix material materialStack[0] = material; totalWeightStack[0] = 1.f; int stackIndex = 0; // Setup the results float3 result = BLACK; while (stackIndex >= 0) { // Extract a material from the stack __global Material *m = materialStack[stackIndex]; float totalWeight = totalWeightStack[stackIndex--]; if (m->type == MIX) { const float factor = Texture_GetFloatValue(&texs[m->mix.mixFactorTexIndex], hitPoint TEXTURES_PARAM); const float weight2 = clamp(factor, 0.f, 1.f); const float weight1 = 1.f - weight2; if (weight1 > 0.f) { materialStack[++stackIndex] = &mats[m->mix.matAIndex]; totalWeightStack[stackIndex] = totalWeight * weight1; } if (weight2 > 0.f) { materialStack[++stackIndex] = &mats[m->mix.matBIndex]; totalWeightStack[stackIndex] = totalWeight * weight2; } } else { const float3 emitRad = Material_GetEmittedRadianceNoMix(m, hitPoint TEXTURES_PARAM); if (!Spectrum_IsBlack(emitRad)) result += totalWeight * emitRad; } } return result; } #if defined(PARAM_HAS_BUMPMAPS) float2 MixMaterial_GetBumpTexValue(__global Material *material, __global HitPoint *hitPoint MATERIALS_PARAM_DECL) { __global Material *materialStack[MIX_STACK_SIZE]; float totalWeightStack[MIX_STACK_SIZE]; // Push the root Mix material materialStack[0] = material; totalWeightStack[0] = 1.f; int stackIndex = 0; // Setup the results float2 result = (float2)(0.f, 0.f); while (stackIndex >= 0) { // Extract a material from the stack __global Material *m = materialStack[stackIndex]; float totalWeight = totalWeightStack[stackIndex--]; if (m->type == MIX) { const float factor = Texture_GetFloatValue(&texs[m->mix.mixFactorTexIndex], hitPoint TEXTURES_PARAM); const float weight2 = clamp(factor, 0.f, 1.f); const float weight1 = 1.f - weight2; if (weight1 > 0.f) { materialStack[++stackIndex] = &mats[m->mix.matAIndex]; totalWeightStack[stackIndex] = totalWeight * weight1; } if (weight2 > 0.f) { materialStack[++stackIndex] = &mats[m->mix.matBIndex]; totalWeightStack[stackIndex] = totalWeight * weight2; } } else { const float2 value = Material_GetBumpTexValueNoMix(m, hitPoint TEXTURES_PARAM); if ((value.s0 != 0.f) || (value.s1 != 0.f)) result += totalWeight * value; } } return result; } #endif #if defined(PARAM_HAS_NORMALMAPS) float3 MixMaterial_GetNormalTexValue(__global Material *material, __global HitPoint *hitPoint MATERIALS_PARAM_DECL) { __global Material *materialStack[MIX_STACK_SIZE]; float totalWeightStack[MIX_STACK_SIZE]; // Push the root Mix material materialStack[0] = material; totalWeightStack[0] = 1.f; int stackIndex = 0; // Setup the results float3 result = BLACK; while (stackIndex >= 0) { // Extract a material from the stack __global Material *m = materialStack[stackIndex]; float totalWeight = totalWeightStack[stackIndex--]; if (m->type == MIX) { const float factor = Texture_GetFloatValue(&texs[m->mix.mixFactorTexIndex], hitPoint TEXTURES_PARAM); const float weight2 = clamp(factor, 0.f, 1.f); const float weight1 = 1.f - weight2; if (weight1 > 0.f) { materialStack[++stackIndex] = &mats[m->mix.matAIndex]; totalWeightStack[stackIndex] = totalWeight * weight1; } if (weight2 > 0.f) { materialStack[++stackIndex] = &mats[m->mix.matBIndex]; totalWeightStack[stackIndex] = totalWeight * weight2; } } else { const float3 value = Material_GetNormalTexValueNoMix(m, hitPoint TEXTURES_PARAM); if (!Spectrum_IsBlack(value)) result += totalWeight * value; } } return result; } #endif float3 MixMaterial_GetPassThroughTransparency(__global Material *material, __global HitPoint *hitPoint, const float3 fixedDir, const float passEvent MATERIALS_PARAM_DECL) { __global Material *currentMixMat = material; float passThroughEvent = passEvent; for (;;) { const float factor = Texture_GetFloatValue(&texs[currentMixMat->mix.mixFactorTexIndex], hitPoint TEXTURES_PARAM); const float weight2 = clamp(factor, 0.f, 1.f); const float weight1 = 1.f - weight2; const bool sampleMatA = (passThroughEvent < weight1); passThroughEvent = sampleMatA ? (passThroughEvent / weight1) : (passThroughEvent - weight1) / weight2; const uint matIndex = sampleMatA ? currentMixMat->mix.matAIndex : currentMixMat->mix.matBIndex; __global Material *mat = &mats[matIndex]; if (mat->type == MIX) { currentMixMat = mat; continue; } else return Material_GetPassThroughTransparencyNoMix(mat, hitPoint, fixedDir, passThroughEvent TEXTURES_PARAM); } } #endif //------------------------------------------------------------------------------ // Generic material functions with Mix support //------------------------------------------------------------------------------ BSDFEvent Material_GetEventTypes(__global Material *material MATERIALS_PARAM_DECL) { #if defined (PARAM_ENABLE_MAT_MIX) if (material->type == MIX) return MixMaterial_GetEventTypes(material MATERIALS_PARAM); else #endif return Material_GetEventTypesNoMix(material); } bool Material_IsDelta(__global Material *material MATERIALS_PARAM_DECL) { #if defined (PARAM_ENABLE_MAT_MIX) if (material->type == MIX) return MixMaterial_IsDelta(material MATERIALS_PARAM); else #endif return Material_IsDeltaNoMix(material); } float3 Material_Evaluate(__global Material *material, __global HitPoint *hitPoint, const float3 lightDir, const float3 eyeDir, BSDFEvent *event, float *directPdfW MATERIALS_PARAM_DECL) { #if defined (PARAM_ENABLE_MAT_MIX) if (material->type == MIX) return MixMaterial_Evaluate(material, hitPoint, lightDir, eyeDir, event, directPdfW MATERIALS_PARAM); else #endif return Material_EvaluateNoMix(material, hitPoint, lightDir, eyeDir, event, directPdfW TEXTURES_PARAM); } float3 Material_Sample(__global Material *material, __global HitPoint *hitPoint, const float3 fixedDir, float3 *sampledDir, const float u0, const float u1, #if defined(PARAM_HAS_PASSTHROUGH) const float passThroughEvent, #endif float *pdfW, float *cosSampledDir, BSDFEvent *event MATERIALS_PARAM_DECL) { #if defined (PARAM_ENABLE_MAT_MIX) if (material->type == MIX) return MixMaterial_Sample(material, hitPoint, fixedDir, sampledDir, u0, u1, passThroughEvent, pdfW, cosSampledDir, event MATERIALS_PARAM); else #endif return Material_SampleNoMix(material, hitPoint, fixedDir, sampledDir, u0, u1, #if defined(PARAM_HAS_PASSTHROUGH) passThroughEvent, #endif pdfW, cosSampledDir, event TEXTURES_PARAM); } float3 Material_GetEmittedRadiance(__global Material *material, __global HitPoint *hitPoint MATERIALS_PARAM_DECL) { #if defined (PARAM_ENABLE_MAT_MIX) if (material->type == MIX) return MixMaterial_GetEmittedRadiance(material, hitPoint MATERIALS_PARAM); else #endif return Material_GetEmittedRadianceNoMix(material, hitPoint TEXTURES_PARAM); } #if defined(PARAM_HAS_BUMPMAPS) float2 Material_GetBumpTexValue(__global Material *material, __global HitPoint *hitPoint MATERIALS_PARAM_DECL) { #if defined (PARAM_ENABLE_MAT_MIX) if (material->type == MIX) return MixMaterial_GetBumpTexValue(material, hitPoint MATERIALS_PARAM); else #endif return Material_GetBumpTexValueNoMix(material, hitPoint TEXTURES_PARAM); } #endif #if defined(PARAM_HAS_NORMALMAPS) float3 Material_GetNormalTexValue(__global Material *material, __global HitPoint *hitPoint MATERIALS_PARAM_DECL) { #if defined (PARAM_ENABLE_MAT_MIX) if (material->type == MIX) return MixMaterial_GetNormalTexValue(material, hitPoint MATERIALS_PARAM); else #endif return Material_GetNormalTexValueNoMix(material, hitPoint TEXTURES_PARAM); } #endif #if defined(PARAM_HAS_PASSTHROUGH) float3 Material_GetPassThroughTransparency(__global Material *material, __global HitPoint *hitPoint, const float3 fixedDir, const float passThroughEvent MATERIALS_PARAM_DECL) { #if defined (PARAM_ENABLE_MAT_MIX) if (material->type == MIX) return MixMaterial_GetPassThroughTransparency(material, hitPoint, fixedDir, passThroughEvent MATERIALS_PARAM); else #endif return Material_GetPassThroughTransparencyNoMix(material, hitPoint, fixedDir, passThroughEvent TEXTURES_PARAM); } #endif #line 2 "camera_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ #if defined(PARAM_CAMERA_ENABLE_OCULUSRIFT_BARREL) void Camera_OculusRiftBarrelPostprocess(const float x, const float y, float *barrelX, float *barrelY) { // Express the sample in coordinates relative to the eye center float ex, ey; if (x < .5f) { // A left eye sample ex = x * 4.f - 1.f; ey = y * 2.f - 1.f; } else { // A right eye sample ex = (x - .5f) * 4.f - 1.f; ey = y * 2.f - 1.f; } if ((ex == 0.f) && (ey == 0.f)) { *barrelX = 0.f; *barrelY = 0.f; return; } // Distance from the eye center const float distance = sqrt(ex * ex + ey * ey); // "Push" the sample away base on the distance from the center const float scale = 1.f / 1.4f; const float k0 = 1.f; const float k1 = .22f; const float k2 = .23f; const float k3 = 0.f; const float distance2 = distance * distance; const float distance4 = distance2 * distance2; const float distance6 = distance2 * distance4; const float fr = scale * (k0 + k1 * distance2 + k2 * distance4 + k3 * distance6); ex *= fr; ey *= fr; // Clamp the coordinates ex = clamp(ex, -1.f, 1.f); ey = clamp(ey, -1.f, 1.f); if (x < .5f) { *barrelX = (ex + 1.f) * .25f; *barrelY = (ey + 1.f) * .5f; } else { *barrelX = (ex + 1.f) * .25f + .5f; *barrelY = (ey + 1.f) * .5f; } } #endif void Camera_GenerateRay( __global Camera *camera, __global Ray *ray, const float scrSampleX, const float scrSampleY #if defined(PARAM_CAMERA_HAS_DOF) , const float dofSampleX, const float dofSampleY #endif ) { #if defined(PARAM_CAMERA_ENABLE_HORIZ_STEREO) // Left eye or right eye const uint transIndex = (scrSampleX < .5f) ? 0 : 1; #else const uint transIndex = 0; #endif float ssx, ssy; #if defined(PARAM_CAMERA_ENABLE_HORIZ_STEREO) && defined(PARAM_CAMERA_ENABLE_OCULUSRIFT_BARREL) Camera_OculusRiftBarrelPostprocess(scrSampleX, scrSampleY, &ssx, &ssy); #else ssx = scrSampleX; ssy = scrSampleY; #endif const float screenX = min(ssx * PARAM_IMAGE_WIDTH, (float)(PARAM_IMAGE_WIDTH - 1)); const float screenY = min((1.f - ssy) * PARAM_IMAGE_HEIGHT, (float)(PARAM_IMAGE_HEIGHT - 1)); float3 Pras = (float3)(screenX, screenY, 0.f); float3 rayOrig = Transform_ApplyPoint(&camera->rasterToCamera[transIndex], Pras); float3 rayDir = rayOrig; const float hither = camera->hither; #if defined(PARAM_CAMERA_HAS_DOF) // Sample point on lens float lensU, lensV; ConcentricSampleDisk(dofSampleX, dofSampleY, &lensU, &lensV); const float lensRadius = camera->lensRadius; lensU *= lensRadius; lensV *= lensRadius; // Compute point on plane of focus const float focalDistance = camera->focalDistance; const float dist = focalDistance - hither; const float ft = dist / rayDir.z; float3 Pfocus; Pfocus = rayOrig + rayDir * ft; // Update ray for effect of lens const float k = dist / focalDistance; rayOrig.x += lensU * k; rayOrig.y += lensV * k; rayDir = Pfocus - rayOrig; #endif rayDir = normalize(rayDir); const float maxt = (camera->yon - hither) / rayDir.z; // Transform ray in world coordinates rayOrig = Transform_ApplyPoint(&camera->cameraToWorld[transIndex], rayOrig); rayDir = Transform_ApplyVector(&camera->cameraToWorld[transIndex], rayDir); Ray_Init3(ray, rayOrig, rayDir, maxt); /*printf("(%f, %f, %f) (%f, %f, %f) [%f, %f]\n", ray->o.x, ray->o.y, ray->o.z, ray->d.x, ray->d.y, ray->d.z, ray->mint, ray->maxt);*/ } #line 2 "light_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ //------------------------------------------------------------------------------ // InfiniteLight //------------------------------------------------------------------------------ #if defined(PARAM_HAS_INFINITELIGHT) float3 InfiniteLight_GetRadiance( __global InfiniteLight *infiniteLight, const float3 dir IMAGEMAPS_PARAM_DECL) { __global ImageMap *imageMap = &imageMapDescs[infiniteLight->imageMapIndex]; __global float *pixels = ImageMap_GetPixelsAddress( imageMapBuff, imageMap->pageIndex, imageMap->pixelsIndex); const float3 localDir = normalize(Transform_InvApplyVector(&infiniteLight->light2World, -dir)); const float2 uv = (float2)( SphericalPhi(localDir) * (1.f / (2.f * M_PI_F)), SphericalTheta(localDir) * M_1_PI_F); // TextureMapping2D_Map() is expendaded here const float2 scale = VLOAD2F(&infiniteLight->mapping.uvMapping2D.uScale); const float2 delta = VLOAD2F(&infiniteLight->mapping.uvMapping2D.uDelta); const float2 mapUV = uv * scale + delta; return VLOAD3F(&infiniteLight->gain.r) * ImageMap_GetSpectrum( pixels, imageMap->width, imageMap->height, imageMap->channelCount, mapUV.s0, mapUV.s1); } #endif //------------------------------------------------------------------------------ // SktLight //------------------------------------------------------------------------------ #if defined(PARAM_HAS_SKYLIGHT) float SkyLight_PerezBase(__global float *lam, const float theta, const float gamma) { return (1.f + lam[1] * exp(lam[2] / cos(theta))) * (1.f + lam[3] * exp(lam[4] * gamma) + lam[5] * cos(gamma) * cos(gamma)); } float SkyLight_RiAngleBetween(const float thetav, const float phiv, const float theta, const float phi) { const float cospsi = sin(thetav) * sin(theta) * cos(phi - phiv) + cos(thetav) * cos(theta); if (cospsi >= 1.f) return 0.f; if (cospsi <= -1.f) return M_PI_F; return acos(cospsi); } float3 SkyLight_ChromaticityToSpectrum(float Y, float x, float y) { float X, Z; if (y != 0.f) X = (x / y) * Y; else X = 0.f; if (y != 0.f && Y != 0.f) Z = (1.f - x - y) / y * Y; else Z = 0.f; // Assuming sRGB (D65 illuminant) return (float3)(3.2410f * X - 1.5374f * Y - 0.4986f * Z, -0.9692f * X + 1.8760f * Y + 0.0416f * Z, 0.0556f * X - 0.2040f * Y + 1.0570f * Z); } float3 SkyLight_GetSkySpectralRadiance(__global SkyLight *skyLight, const float theta, const float phi) { // Add bottom half of hemisphere with horizon colour const float theta_fin = fmin(theta, (M_PI_F * .5f) - .001f); const float gamma = SkyLight_RiAngleBetween(theta, phi, skyLight->thetaS, skyLight->phiS); // Compute xyY values const float x = skyLight->zenith_x * SkyLight_PerezBase(skyLight->perez_x, theta_fin, gamma); const float y = skyLight->zenith_y * SkyLight_PerezBase(skyLight->perez_y, theta_fin, gamma); const float Y = skyLight->zenith_Y * SkyLight_PerezBase(skyLight->perez_Y, theta_fin, gamma); return SkyLight_ChromaticityToSpectrum(Y, x, y); } float3 SkyLight_GetRadiance(__global SkyLight *skyLight, const float3 dir) { const float3 localDir = normalize(Transform_InvApplyVector(&skyLight->light2World, -dir)); const float theta = SphericalTheta(localDir); const float phi = SphericalPhi(localDir); const float3 s = SkyLight_GetSkySpectralRadiance(skyLight, theta, phi); return VLOAD3F(&skyLight->gain.r) * s; } #endif //------------------------------------------------------------------------------ // SunLight //------------------------------------------------------------------------------ #if defined(PARAM_HAS_SUNLIGHT) float3 SunLight_Illuminate(__global SunLight *sunLight, const float u0, const float u1, float3 *dir, float *distance, float *directPdfW) { const float cosThetaMax = sunLight->cosThetaMax; const float3 sunDir = VLOAD3F(&sunLight->sunDir.x); *dir = UniformSampleCone(u0, u1, cosThetaMax, VLOAD3F(&sunLight->x.x), VLOAD3F(&sunLight->y.x), sunDir); // Check if the point can be inside the sun cone of light const float cosAtLight = dot(sunDir, *dir); if (cosAtLight <= cosThetaMax) return BLACK; *distance = INFINITY; *directPdfW = UniformConePdf(cosThetaMax); return VLOAD3F(&sunLight->sunColor.r); } float3 SunLight_GetRadiance(__global SunLight *sunLight, const float3 dir, float *directPdfA) { // Make the sun visible only if relsize has been changed (in order // to avoid fireflies). if (sunLight->relSize > 5.f) { const float cosThetaMax = sunLight->cosThetaMax; const float3 sunDir = VLOAD3F(&sunLight->sunDir.x); if ((cosThetaMax < 1.f) && (dot(-dir, sunDir) > cosThetaMax)) { if (directPdfA) *directPdfA = UniformConePdf(cosThetaMax); return VLOAD3F(&sunLight->sunColor.r); } } return BLACK; } #endif //------------------------------------------------------------------------------ // TriangleLight //------------------------------------------------------------------------------ float3 TriangleLight_Illuminate(__global TriangleLight *triLight, __global HitPoint *tmpHitPoint, const float3 p, const float u0, const float u1, const float passThroughEvent, float3 *dir, float *distance, float *directPdfW MATERIALS_PARAM_DECL) { const float3 p0 = VLOAD3F(&triLight->v0.x); const float3 p1 = VLOAD3F(&triLight->v1.x); const float3 p2 = VLOAD3F(&triLight->v2.x); float b0, b1, b2; float3 samplePoint = Triangle_Sample( p0, p1, p2, u0, u1, &b0, &b1, &b2); const float3 sampleN = Triangle_GetGeometryNormal(p0, p1, p2); // Light sources are supposed to be flat *dir = samplePoint - p; const float distanceSquared = dot(*dir, *dir);; *distance = sqrt(distanceSquared); *dir /= (*distance); const float cosAtLight = dot(sampleN, -(*dir)); if (cosAtLight < DEFAULT_COS_EPSILON_STATIC) return BLACK; *directPdfW = triLight->invArea * distanceSquared / cosAtLight; const float2 uv0 = VLOAD2F(&triLight->uv0.u); const float2 uv1 = VLOAD2F(&triLight->uv1.u); const float2 uv2 = VLOAD2F(&triLight->uv2.u); const float2 triUV = Triangle_InterpolateUV(uv0, uv1, uv2, b0, b1, b2); VSTORE3F(-sampleN, &tmpHitPoint->fixedDir.x); VSTORE3F(samplePoint, &tmpHitPoint->p.x); VSTORE2F(triUV, &tmpHitPoint->uv.u); VSTORE3F(sampleN, &tmpHitPoint->geometryN.x); VSTORE3F(sampleN, &tmpHitPoint->shadeN.x); #if defined(PARAM_HAS_PASSTHROUGH) tmpHitPoint->passThroughEvent = passThroughEvent; #endif return Material_GetEmittedRadiance(&mats[triLight->materialIndex], tmpHitPoint MATERIALS_PARAM); } float3 TriangleLight_GetRadiance(__global TriangleLight *triLight, __global HitPoint *hitPoint, float *directPdfA MATERIALS_PARAM_DECL) { const float3 dir = VLOAD3F(&hitPoint->fixedDir.x); const float3 hitPointNormal = VLOAD3F(&hitPoint->geometryN.x); const float cosOutLight = dot(hitPointNormal, dir); if (cosOutLight <= 0.f) return BLACK; if (directPdfA) *directPdfA = triLight->invArea; return Material_GetEmittedRadiance(&mats[triLight->materialIndex], hitPoint MATERIALS_PARAM); } #line 2 "filter_types.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ //------------------------------------------------------------------------------ // Pixel related functions //------------------------------------------------------------------------------ void PixelIndex2XY(const uint index, uint *x, uint *y) { *y = index / PARAM_IMAGE_WIDTH; *x = index - (*y) * PARAM_IMAGE_WIDTH; } uint XY2PixelIndex(const uint x, const uint y) { return x + y * PARAM_IMAGE_WIDTH; } uint XY2FrameBufferIndex(const int x, const int y) { return x + 1 + (y + 1) * (PARAM_IMAGE_WIDTH + 2); } bool IsValidPixelXY(const int x, const int y) { return (x >= 0) && (x < PARAM_IMAGE_WIDTH) && (y >= 0) && (y < PARAM_IMAGE_HEIGHT); } //------------------------------------------------------------------------------ // Image filtering related functions //------------------------------------------------------------------------------ #if (PARAM_IMAGE_FILTER_TYPE == 0) // Nothing #elif (PARAM_IMAGE_FILTER_TYPE == 1) // Box Filter float ImageFilter_Evaluate(const float x, const float y) { return 1.f; } #elif (PARAM_IMAGE_FILTER_TYPE == 2) float Gaussian(const float d, const float expv) { return max(0.f, exp(-PARAM_IMAGE_FILTER_GAUSSIAN_ALPHA * d * d) - expv); } // Gaussian Filter float ImageFilter_Evaluate(const float x, const float y) { return Gaussian(x, exp(-PARAM_IMAGE_FILTER_GAUSSIAN_ALPHA * PARAM_IMAGE_FILTER_WIDTH_X * PARAM_IMAGE_FILTER_WIDTH_X)) * Gaussian(y, exp(-PARAM_IMAGE_FILTER_GAUSSIAN_ALPHA * PARAM_IMAGE_FILTER_WIDTH_Y * PARAM_IMAGE_FILTER_WIDTH_Y)); } #elif (PARAM_IMAGE_FILTER_TYPE == 3) float Mitchell1D(float x) { const float B = PARAM_IMAGE_FILTER_MITCHELL_B; const float C = PARAM_IMAGE_FILTER_MITCHELL_C; if (x >= 1.f) return 0.f; x = fabs(2.f * x); if (x > 1.f) return (((-B / 6.f - C) * x + (B + 5.f * C)) * x + (-2.f * B - 8.f * C)) * x + (4.f / 3.f * B + 4.f * C); else return ((2.f - 1.5f * B - C) * x + (-3.f + 2.f * B + C)) * x * x + (1.f - B / 3.f); } // Mitchell Filter float ImageFilter_Evaluate(const float x, const float y) { const float distance = native_sqrt( x * x * (1.f / (PARAM_IMAGE_FILTER_WIDTH_X * PARAM_IMAGE_FILTER_WIDTH_X)) + y * y * (1.f / (PARAM_IMAGE_FILTER_WIDTH_Y * PARAM_IMAGE_FILTER_WIDTH_Y))); return Mitchell1D(distance); } #else Error: unknown image filter !!! #endif #if defined(PARAM_USE_PIXEL_ATOMICS) void AtomicAdd(__global float *val, const float delta) { union { float f; unsigned int i; } oldVal; union { float f; unsigned int i; } newVal; do { oldVal.f = *val; newVal.f = oldVal.f + delta; } while (atomic_cmpxchg((__global unsigned int *)val, oldVal.i, newVal.i) != oldVal.i); } #endif void Pixel_AddRadiance(__global Pixel *pixel, const float3 rad, const float weight) { /*if (isnan(rad->r) || isinf(rad->r) || isnan(rad->g) || isinf(rad->g) || isnan(rad->b) || isinf(rad->b) || isnan(weight) || isinf(weight)) printf(\"NaN/Inf. error: (%f, %f, %f) [%f]\\n\", rad->r, rad->g, rad->b, weight);*/ float4 s; s.xyz = rad; s.w = 1.f; s *= weight; #if defined(PARAM_USE_PIXEL_ATOMICS) AtomicAdd(&pixel->c.r, s.x); AtomicAdd(&pixel->c.g, s.y); AtomicAdd(&pixel->c.b, s.z); AtomicAdd(&pixel->count, s.w); #else float4 p = VLOAD4F(&(pixel->c.r)); p += s; VSTORE4F(p, &(pixel->c.r)); #endif } #if defined(PARAM_ENABLE_ALPHA_CHANNEL) void Pixel_AddAlpha(__global AlphaPixel *apixel, const float alpha, const float weight) { #if defined(PARAM_USE_PIXEL_ATOMICS) AtomicAdd(&apixel->alpha, weight * alpha); #else apixel->alpha += weight * alpha; #endif } #endif #if (PARAM_IMAGE_FILTER_TYPE == 1) || (PARAM_IMAGE_FILTER_TYPE == 2) || (PARAM_IMAGE_FILTER_TYPE == 3) void Pixel_AddFilteredRadiance(__global Pixel *pixel, const float3 rad, const float distX, const float distY, const float weight) { const float filterWeight = ImageFilter_Evaluate(distX, distY); Pixel_AddRadiance(pixel, rad, weight * filterWeight); } #if defined(PARAM_ENABLE_ALPHA_CHANNEL) void Pixel_AddFilteredAlpha(__global AlphaPixel *apixel, const float alpha, const float distX, const float distY, const float weight) { const float filterWeight = ImageFilter_Evaluate(distX, distY); Pixel_AddAlpha(apixel, alpha, weight * filterWeight); } #endif #endif #if (PARAM_IMAGE_FILTER_TYPE == 0) void SplatSample(__global Pixel *frameBuffer, const float scrX, const float scrY, const float3 radiance, #if defined(PARAM_ENABLE_ALPHA_CHANNEL) __global AlphaPixel *alphaFrameBuffer, const float alpha, #endif const float weight) { const uint x = min((uint)floor(PARAM_IMAGE_WIDTH * scrX + .5f), (uint)(PARAM_IMAGE_WIDTH - 1)); const uint y = min((uint)floor(PARAM_IMAGE_HEIGHT * scrY + .5f), (uint)(PARAM_IMAGE_HEIGHT - 1)); __global Pixel *pixel = &frameBuffer[XY2FrameBufferIndex(x, y)]; Pixel_AddRadiance(pixel, radiance, weight); #if defined(PARAM_ENABLE_ALPHA_CHANNEL) __global AlphaPixel *apixel = &alphaFrameBuffer[XY2FrameBufferIndex(x, y)]; Pixel_AddAlpha(apixel, alpha, weight); #endif } #elif (PARAM_IMAGE_FILTER_TYPE == 1) || (PARAM_IMAGE_FILTER_TYPE == 2) || (PARAM_IMAGE_FILTER_TYPE == 3) void SplatSample(__global Pixel *frameBuffer, const float scrX, const float scrY, const float3 radiance, #if defined(PARAM_ENABLE_ALPHA_CHANNEL) __global AlphaPixel *alphaFrameBuffer, const float alpha, #endif const float weight) { const float px = PARAM_IMAGE_WIDTH * scrX + .5f; const float py = PARAM_IMAGE_HEIGHT * scrY + .5f; const uint x = min((uint)floor(px), (uint)(PARAM_IMAGE_WIDTH - 1)); const uint y = min((uint)floor(py), (uint)(PARAM_IMAGE_HEIGHT - 1)); const float sx = px - (float)x; const float sy = py - (float)y; { __global Pixel *pixel = &frameBuffer[XY2FrameBufferIndex(x - 1, y - 1)]; Pixel_AddFilteredRadiance(pixel, radiance, sx + 1.f, sy + 1.f, weight); pixel = &frameBuffer[XY2FrameBufferIndex(x, y - 1)]; Pixel_AddFilteredRadiance(pixel, radiance, sx, sy + 1.f, weight); pixel = &frameBuffer[XY2FrameBufferIndex(x + 1, y - 1)]; Pixel_AddFilteredRadiance(pixel, radiance, sx - 1.f, sy + 1.f, weight); pixel = &frameBuffer[XY2FrameBufferIndex(x - 1, y)]; Pixel_AddFilteredRadiance(pixel, radiance, sx + 1.f, sy, weight); pixel = &frameBuffer[XY2FrameBufferIndex(x, y)]; Pixel_AddFilteredRadiance(pixel, radiance, sx, sy, weight); pixel = &frameBuffer[XY2FrameBufferIndex(x + 1, y)]; Pixel_AddFilteredRadiance(pixel, radiance, sx - 1.f, sy, weight); pixel = &frameBuffer[XY2FrameBufferIndex(x - 1, y + 1)]; Pixel_AddFilteredRadiance(pixel, radiance, sx + 1.f, sy - 1.f, weight); pixel = &frameBuffer[XY2FrameBufferIndex(x, y + 1)]; Pixel_AddFilteredRadiance(pixel, radiance, sx, sy - 1.f, weight); pixel = &frameBuffer[XY2FrameBufferIndex(x + 1, y + 1)]; Pixel_AddFilteredRadiance(pixel, radiance, sx - 1.f, sy - 1.f, weight); } #if defined(PARAM_ENABLE_ALPHA_CHANNEL) { __global AlphaPixel *apixel = &alphaFrameBuffer[XY2FrameBufferIndex(x - 1, y - 1)]; Pixel_AddFilteredAlpha(apixel, alpha, sx + 1.f, sy + 1.f, weight); apixel = &alphaFrameBuffer[XY2FrameBufferIndex(x, y - 1)]; Pixel_AddFilteredAlpha(apixel, alpha, sx, sy + 1.f, weight); apixel = &alphaFrameBuffer[XY2FrameBufferIndex(x + 1, y - 1)]; Pixel_AddFilteredAlpha(apixel, alpha, sx - 1.f, sy + 1.f, weight); apixel = &alphaFrameBuffer[XY2FrameBufferIndex(x - 1, y)]; Pixel_AddFilteredAlpha(apixel, alpha, sx + 1.f, sy, weight); apixel = &alphaFrameBuffer[XY2FrameBufferIndex(x, y)]; Pixel_AddFilteredAlpha(apixel, alpha, sx, sy, weight); apixel = &alphaFrameBuffer[XY2FrameBufferIndex(x + 1, y)]; Pixel_AddFilteredAlpha(apixel, alpha, sx - 1.f, sy, weight); apixel = &alphaFrameBuffer[XY2FrameBufferIndex(x - 1, y + 1)]; Pixel_AddFilteredAlpha(apixel, alpha, sx + 1.f, sy - 1.f, weight); apixel = &alphaFrameBuffer[XY2FrameBufferIndex(x, y + 1)]; Pixel_AddFilteredAlpha(apixel, alpha, sx, sy - 1.f, weight); apixel = &alphaFrameBuffer[XY2FrameBufferIndex(x + 1, y + 1)]; Pixel_AddFilteredAlpha(apixel, alpha, sx - 1.f, sy - 1.f, weight); } #endif } #else Error: unknown image filter !!! #endif #line 2 "sampler_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ //------------------------------------------------------------------------------ // Random Sampler Kernel //------------------------------------------------------------------------------ #if (PARAM_SAMPLER_TYPE == 0) #define Sampler_GetSamplePath(index) (Rnd_FloatValue(seed)) #define Sampler_GetSamplePathVertex(depth, index) (Rnd_FloatValue(seed)) __global float *Sampler_GetSampleData(__global Sample *sample, __global float *samplesData) { const size_t gid = get_global_id(0); return &samplesData[gid * TOTAL_U_SIZE]; } __global float *Sampler_GetSampleDataPathBase(__global Sample *sample, __global float *sampleData) { return sampleData; } __global float *Sampler_GetSampleDataPathVertex(__global Sample *sample, __global float *sampleDataPathBase, const uint depth) { return &sampleDataPathBase[IDX_BSDF_OFFSET + depth * VERTEX_SAMPLE_SIZE]; } void Sampler_Init(Seed *seed, __global Sample *sample, __global float *sampleData) { sampleData[IDX_SCREEN_X] = Rnd_FloatValue(seed); sampleData[IDX_SCREEN_Y] = Rnd_FloatValue(seed); VSTORE3F(BLACK, &sample->radiance.r); #if defined(PARAM_ENABLE_ALPHA_CHANNEL) sample->alpha = 1.f; #endif } void Sampler_NextSample( __global Sample *sample, __global float *sampleData, Seed *seed, __global Pixel *frameBuffer #if defined(PARAM_ENABLE_ALPHA_CHANNEL) , __global AlphaPixel *alphaFrameBuffer #endif ) { SplatSample(frameBuffer, sampleData[IDX_SCREEN_X], sampleData[IDX_SCREEN_Y], VLOAD3F(&sample->radiance.r), #if defined(PARAM_ENABLE_ALPHA_CHANNEL) alphaFrameBuffer, sample->alpha, #endif 1.f); // Move to the next assigned pixel sampleData[IDX_SCREEN_X] = Rnd_FloatValue(seed); sampleData[IDX_SCREEN_Y] = Rnd_FloatValue(seed); VSTORE3F(BLACK, &sample->radiance.r); #if defined(PARAM_ENABLE_ALPHA_CHANNEL) sample->alpha = 1.f; #endif } #endif //------------------------------------------------------------------------------ // Metropolis Sampler Kernel //------------------------------------------------------------------------------ #if (PARAM_SAMPLER_TYPE == 1) #define Sampler_GetSamplePath(index) (sampleDataPathBase[index]) #define Sampler_GetSamplePathVertex(depth, index) (sampleDataPathVertexBase[index]) __global float *Sampler_GetSampleData(__global Sample *sample, __global float *samplesData) { const size_t gid = get_global_id(0); return &samplesData[gid * (2 * TOTAL_U_SIZE)]; } __global float *Sampler_GetSampleDataPathBase(__global Sample *sample, __global float *sampleData) { return &sampleData[sample->proposed * TOTAL_U_SIZE]; } __global float *Sampler_GetSampleDataPathVertex(__global Sample *sample, __global float *sampleDataPathBase, const uint depth) { return &sampleDataPathBase[IDX_BSDF_OFFSET + depth * VERTEX_SAMPLE_SIZE]; } void LargeStep(Seed *seed, const uint largeStepCount, __global float *proposedU) { for (int i = 0; i < TOTAL_U_SIZE; ++i) proposedU[i] = Rnd_FloatValue(seed); } float Mutate(Seed *seed, const float x) { const float s1 = 1.f / 512.f; const float s2 = 1.f / 16.f; const float randomValue = Rnd_FloatValue(seed); const float dx = s1 / (s1 / s2 + fabs(2.f * randomValue - 1.f)) - s1 / (s1 / s2 + 1.f); float mutatedX = x; if (randomValue < 0.5f) { mutatedX += dx; mutatedX = (mutatedX < 1.f) ? mutatedX : (mutatedX - 1.f); } else { mutatedX -= dx; mutatedX = (mutatedX < 0.f) ? (mutatedX + 1.f) : mutatedX; } return mutatedX; } float MutateScaled(Seed *seed, const float x, const float range) { const float s1 = 32.f; const float randomValue = Rnd_FloatValue(seed); const float dx = range / (s1 / (1.f + s1) + (s1 * s1) / (1.f + s1) * fabs(2.f * randomValue - 1.f)) - range / s1; float mutatedX = x; if (randomValue < 0.5f) { mutatedX += dx; mutatedX = (mutatedX < 1.f) ? mutatedX : (mutatedX - 1.f); } else { mutatedX -= dx; mutatedX = (mutatedX < 0.f) ? (mutatedX + 1.f) : mutatedX; } return mutatedX; } void SmallStep(Seed *seed, __global float *currentU, __global float *proposedU) { proposedU[IDX_SCREEN_X] = MutateScaled(seed, currentU[IDX_SCREEN_X], PARAM_SAMPLER_METROPOLIS_IMAGE_MUTATION_RANGE); proposedU[IDX_SCREEN_Y] = MutateScaled(seed, currentU[IDX_SCREEN_Y], PARAM_SAMPLER_METROPOLIS_IMAGE_MUTATION_RANGE); for (int i = IDX_SCREEN_Y + 1; i < TOTAL_U_SIZE; ++i) proposedU[i] = Mutate(seed, currentU[i]); } void Sampler_Init(Seed *seed, __global Sample *sample, __global float *sampleData) { sample->totalI = 0.f; sample->largeMutationCount = 1.f; sample->current = NULL_INDEX; sample->proposed = 1; sample->smallMutationCount = 0; sample->consecutiveRejects = 0; sample->weight = 0.f; VSTORE3F(BLACK, &sample->currentRadiance.r); VSTORE3F(BLACK, &sample->radiance.r); #if defined(PARAM_ENABLE_ALPHA_CHANNEL) sample->currentAlpha = 1.f; sample->alpha = 1.f; #endif __global float *sampleDataPathBase = Sampler_GetSampleDataPathBase(sample, sampleData); LargeStep(seed, 0, sampleDataPathBase); VSTORE3F(BLACK, &sample->radiance.r); #if defined(PARAM_ENABLE_ALPHA_CHANNEL) sample->alpha = 1.f; #endif } void Sampler_NextSample( __global Sample *sample, __global float *sampleData, Seed *seed, __global Pixel *frameBuffer #if defined(PARAM_ENABLE_ALPHA_CHANNEL) , __global AlphaPixel *alphaFrameBuffer #endif ) { //-------------------------------------------------------------------------- // Accept/Reject the sample //-------------------------------------------------------------------------- uint current = sample->current; uint proposed = sample->proposed; const float3 radiance = VLOAD3F(&sample->radiance.r); if (current == NULL_INDEX) { // It is the very first sample, I have still to initialize the current // sample VSTORE3F(radiance, &sample->currentRadiance.r); #if defined(PARAM_ENABLE_ALPHA_CHANNEL) sample->currentAlpha = sample->alpha; #endif sample->totalI = Spectrum_Y(radiance); current = proposed; proposed ^= 1; } else { const float3 currentL = VLOAD3F(&sample->currentRadiance.r); #if defined(PARAM_ENABLE_ALPHA_CHANNEL) const float currentAlpha = sample->currentAlpha; #endif const float currentI = Spectrum_Y(currentL); const float3 proposedL = radiance; #if defined(PARAM_ENABLE_ALPHA_CHANNEL) const float proposedAlpha = sample->alpha; #endif float proposedI = Spectrum_Y(proposedL); proposedI = isinf(proposedI) ? 0.f : proposedI; float totalI = sample->totalI; uint largeMutationCount = sample->largeMutationCount; uint smallMutationCount = sample->smallMutationCount; if (smallMutationCount == 0) { // It is a large mutation totalI += Spectrum_Y(proposedL); largeMutationCount += 1; sample->totalI = totalI; sample->largeMutationCount = largeMutationCount; } const float meanI = (totalI > 0.f) ? (totalI / largeMutationCount) : 1.f; // Calculate accept probability from old and new image sample uint consecutiveRejects = sample->consecutiveRejects; float accProb; if ((currentI > 0.f) && (consecutiveRejects < PARAM_SAMPLER_METROPOLIS_MAX_CONSECUTIVE_REJECT)) accProb = min(1.f, proposedI / currentI); else accProb = 1.f; const float newWeight = accProb + ((smallMutationCount == 0) ? 1.f : 0.f); float weight = sample->weight; weight += 1.f - accProb; const float rndVal = Rnd_FloatValue(seed); /*if (get_global_id(0) == 0) printf(\"[%d] Current: (%f, %f, %f) [%f] Proposed: (%f, %f, %f) [%f] accProb: %f <%f>\\n\", consecutiveRejects, currentL.r, currentL.g, currentL.b, weight, proposedL.r, proposedL.g, proposedL.b, newWeight, accProb, rndVal);*/ float3 contrib; #if defined(PARAM_ENABLE_ALPHA_CHANNEL) float contribAlpha; #endif float norm; float scrX, scrY; if ((accProb == 1.f) || (rndVal < accProb)) { /*if (get_global_id(0) == 0) printf(\"\\t\\tACCEPTED !\\n\");*/ // Add accumulated contribution of previous reference sample norm = weight / (currentI / meanI + PARAM_SAMPLER_METROPOLIS_LARGE_STEP_RATE); contrib = currentL; #if defined(PARAM_ENABLE_ALPHA_CHANNEL) contribAlpha = currentAlpha; #endif scrX = sampleData[current * TOTAL_U_SIZE + IDX_SCREEN_X]; scrY = sampleData[current * TOTAL_U_SIZE + IDX_SCREEN_Y]; current ^= 1; proposed ^= 1; consecutiveRejects = 0; weight = newWeight; VSTORE3F(proposedL, &sample->currentRadiance.r); #if defined(PARAM_ENABLE_ALPHA_CHANNEL) sample->currentAlpha = proposedAlpha; #endif } else { /*if (get_global_id(0) == 0) printf(\"\\t\\tREJECTED !\\n\");*/ // Add contribution of new sample before rejecting it norm = newWeight / (proposedI / meanI + PARAM_SAMPLER_METROPOLIS_LARGE_STEP_RATE); contrib = proposedL; #if defined(PARAM_ENABLE_ALPHA_CHANNEL) contribAlpha = proposedAlpha; #endif scrX = sampleData[proposed * TOTAL_U_SIZE + IDX_SCREEN_X]; scrY = sampleData[proposed * TOTAL_U_SIZE + IDX_SCREEN_Y]; ++consecutiveRejects; } if (norm > 0.f) { /*if (get_global_id(0) == 0) printf(\"\\t\\tContrib: (%f, %f, %f) [%f] consecutiveRejects: %d\\n\", contrib.r, contrib.g, contrib.b, norm, consecutiveRejects);*/ SplatSample(frameBuffer, scrX, scrY, contrib, #if defined(PARAM_ENABLE_ALPHA_CHANNEL) alphaFrameBuffer, contribAlpha, #endif norm); } sample->weight = weight; sample->consecutiveRejects = consecutiveRejects; } sample->current = current; sample->proposed = proposed; //-------------------------------------------------------------------------- // Mutate the sample //-------------------------------------------------------------------------- __global float *proposedU = &sampleData[proposed * TOTAL_U_SIZE]; if (Rnd_FloatValue(seed) < PARAM_SAMPLER_METROPOLIS_LARGE_STEP_RATE) { LargeStep(seed, sample->largeMutationCount, proposedU); sample->smallMutationCount = 0; } else { __global float *currentU = &sampleData[current * TOTAL_U_SIZE]; SmallStep(seed, currentU, proposedU); sample->smallMutationCount += 1; } VSTORE3F(BLACK, &sample->radiance.r); #if defined(PARAM_ENABLE_ALPHA_CHANNEL) sample->alpha = 1.f; #endif } #endif //------------------------------------------------------------------------------ // Sobol Sampler Kernel //------------------------------------------------------------------------------ #if (PARAM_SAMPLER_TYPE == 2) uint SobolSampler_SobolDimension(const uint index, const uint dimension) { const uint offset = dimension * SOBOL_BITS; uint result = 0; uint i = index; for (uint j = 0; i; i >>= 1, j++) { if (i & 1) result ^= SOBOL_DIRECTIONS[offset + j]; } return result; } float SobolSampler_GetSample(__global Sample *sample, const uint index) { const uint pass = sample->pass; const uint result = SobolSampler_SobolDimension(pass, index); const float r = result * (1.f / 0xffffffffu); // Cranley-Patterson rotation to reduce visible regular patterns const float shift = (index & 1) ? sample->rng0 : sample->rng1; return r + shift - floor(r + shift); } #define Sampler_GetSamplePath(index) (SobolSampler_GetSample(sample, index)) #define Sampler_GetSamplePathVertex(depth, index) ((depth > PARAM_SAMPLER_SOBOL_MAXDEPTH) ? \ Rnd_FloatValue(seed) : \ SobolSampler_GetSample(sample, IDX_BSDF_OFFSET + (depth - 1) * VERTEX_SAMPLE_SIZE + index)) __global float *Sampler_GetSampleData(__global Sample *sample, __global float *samplesData) { const size_t gid = get_global_id(0); return &samplesData[gid * TOTAL_U_SIZE]; } __global float *Sampler_GetSampleDataPathBase(__global Sample *sample, __global float *sampleData) { return sampleData; } __global float *Sampler_GetSampleDataPathVertex(__global Sample *sample, __global float *sampleDataPathBase, const uint depth) { return &sampleDataPathBase[IDX_BSDF_OFFSET + depth * VERTEX_SAMPLE_SIZE]; } void Sampler_Init(Seed *seed, __global Sample *sample, __global float *sampleData) { VSTORE3F(BLACK, &sample->radiance.r); #if defined(PARAM_ENABLE_ALPHA_CHANNEL) sample->alpha = 1.f; #endif sample->rng0 = Rnd_FloatValue(seed); sample->rng1 = Rnd_FloatValue(seed); sample->pass = PARAM_SAMPLER_SOBOL_STARTOFFSET; const uint pixelIndex = get_global_id(0) + (PARAM_TASK_COUNT * PARAM_DEVICE_INDEX / PARAM_DEVICE_COUNT); sample->pixelIndex = pixelIndex; uint x, y; PixelIndex2XY(pixelIndex, &x, &y); sampleData[IDX_SCREEN_X] = (x + Sampler_GetSamplePath(IDX_SCREEN_X)) * (1.f / PARAM_IMAGE_WIDTH); sampleData[IDX_SCREEN_Y] = (y + Sampler_GetSamplePath(IDX_SCREEN_Y)) * (1.f / PARAM_IMAGE_HEIGHT); VSTORE3F(BLACK, &sample->radiance.r); #if defined(PARAM_ENABLE_ALPHA_CHANNEL) sample->alpha = 1.f; #endif } void Sampler_NextSample( __global Sample *sample, __global float *sampleData, Seed *seed, __global Pixel *frameBuffer #if defined(PARAM_ENABLE_ALPHA_CHANNEL) , __global AlphaPixel *alphaFrameBuffer #endif ) { SplatSample(frameBuffer, sampleData[IDX_SCREEN_X], sampleData[IDX_SCREEN_Y], VLOAD3F(&sample->radiance.r), #if defined(PARAM_ENABLE_ALPHA_CHANNEL) alphaFrameBuffer, sample->alpha, #endif 1.f); // Move to the next assigned pixel uint nextPixelIndex = sample->pixelIndex + PARAM_TASK_COUNT; if (nextPixelIndex > PARAM_IMAGE_WIDTH * PARAM_IMAGE_HEIGHT) { nextPixelIndex = get_global_id(0); sample->pass += 1; } sample->pixelIndex = nextPixelIndex; uint x, y; PixelIndex2XY(nextPixelIndex, &x, &y); sampleData[IDX_SCREEN_X] = (x + Sampler_GetSamplePath(IDX_SCREEN_X)) * (1.f / PARAM_IMAGE_WIDTH); sampleData[IDX_SCREEN_Y] = (y + Sampler_GetSamplePath(IDX_SCREEN_Y)) * (1.f / PARAM_IMAGE_HEIGHT); VSTORE3F(BLACK, &sample->radiance.r); #if defined(PARAM_ENABLE_ALPHA_CHANNEL) sample->alpha = 1.f; #endif } #endif #line 2 "bsdf_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ void BSDF_Init( __global BSDF *bsdf, //const bool fromL, __global Mesh *meshDescs, __global uint *meshMats, #if (PARAM_DL_LIGHT_COUNT > 0) __global uint *meshTriLightDefsOffset, #endif __global Point *vertices, #if defined(PARAM_HAS_NORMALS_BUFFER) __global Vector *vertNormals, #endif #if defined(PARAM_HAS_UVS_BUFFER) __global UV *vertUVs, #endif #if defined(PARAM_HAS_COLS_BUFFER) __global Spectrum *vertCols, #endif #if defined(PARAM_HAS_ALPHAS_BUFFER) __global float *vertAlphas, #endif __global Triangle *triangles, __global Ray *ray, __global RayHit *rayHit #if defined(PARAM_HAS_PASSTHROUGH) , const float u0 #endif #if defined(PARAM_HAS_BUMPMAPS) || defined(PARAM_HAS_NORMALMAPS) MATERIALS_PARAM_DECL #endif ) { //bsdf->fromLight = fromL; #if defined(PARAM_HAS_PASSTHROUGH) bsdf->hitPoint.passThroughEvent = u0; #endif const float3 rayOrig = VLOAD3F(&ray->o.x); const float3 rayDir = VLOAD3F(&ray->d.x); const float3 hitPointP = rayOrig + rayHit->t * rayDir; VSTORE3F(hitPointP, &bsdf->hitPoint.p.x); VSTORE3F(-rayDir, &bsdf->hitPoint.fixedDir.x); const uint meshIndex = rayHit->meshIndex; const uint triangleIndex = rayHit->triangleIndex; __global Mesh *meshDesc = &meshDescs[meshIndex]; __global Point *iVertices = &vertices[meshDesc->vertsOffset]; __global Triangle *iTriangles = &triangles[meshDesc->trisOffset]; // Get the material const uint matIndex = meshMats[meshIndex]; bsdf->materialIndex = matIndex; //-------------------------------------------------------------------------- // Get face normal //-------------------------------------------------------------------------- const float b1 = rayHit->b1; const float b2 = rayHit->b2; // Geometry normal expressed in local coordinates float3 geometryN = Mesh_GetGeometryNormal(iVertices, iTriangles, triangleIndex); // Transform to global coordinates geometryN = normalize(Transform_InvApplyNormal(&meshDesc->trans, geometryN)); // Store the geometry normal VSTORE3F(geometryN, &bsdf->hitPoint.geometryN.x); // The shading normal float3 shadeN; #if defined(PARAM_HAS_NORMALS_BUFFER) if (meshDesc->normalsOffset != NULL_INDEX) { __global Vector *iVertNormals = &vertNormals[meshDesc->normalsOffset]; // Shading normal expressed in local coordinates shadeN = Mesh_InterpolateNormal(iVertNormals, iTriangles, triangleIndex, b1, b2); // Transform to global coordinates shadeN = normalize(Transform_InvApplyNormal(&meshDesc->trans, shadeN)); } else #endif shadeN = geometryN; // Shade normal can not yet be stored because of normal and bump mapping //-------------------------------------------------------------------------- // Get UV coordinate //-------------------------------------------------------------------------- float2 hitPointUV; #if defined(PARAM_HAS_UVS_BUFFER) if (meshDesc->uvsOffset != NULL_INDEX) { __global UV *iVertUVs = &vertUVs[meshDesc->uvsOffset]; hitPointUV = Mesh_InterpolateUV(iVertUVs, iTriangles, triangleIndex, b1, b2); } else #endif hitPointUV = 0.f; VSTORE2F(hitPointUV, &bsdf->hitPoint.uv.u); //-------------------------------------------------------------------------- // Get color value //-------------------------------------------------------------------------- #if defined(PARAM_ENABLE_TEX_HITPOINTCOLOR) || defined(PARAM_ENABLE_TEX_HITPOINTGREY) float3 hitPointColor; #if defined(PARAM_HAS_COLS_BUFFER) if (meshDesc->colsOffset != NULL_INDEX) { __global Spectrum *iVertCols = &vertCols[meshDesc->colsOffset]; hitPointColor = Mesh_InterpolateColor(iVertCols, iTriangles, triangleIndex, b1, b2); } else #endif hitPointColor = WHITE; VSTORE3F(hitPointColor, &bsdf->hitPoint.color.r); #endif //-------------------------------------------------------------------------- // Get alpha value //-------------------------------------------------------------------------- #if defined(PARAM_ENABLE_TEX_HITPOINTALPHA) float hitPointAlpha; #if defined(PARAM_HAS_ALPHAS_BUFFER) if (meshDesc->colsOffset != NULL_INDEX) { __global float *iVertAlphas = &vertAlphas[meshDesc->alphasOffset]; hitPointAlpha = Mesh_InterpolateAlpha(iVertAlphas, iTriangles, triangleIndex, b1, b2); } else #endif hitPointAlpha = 1.f; bsdf->hitPoint.alpha = hitPointAlpha; #endif //-------------------------------------------------------------------------- #if (PARAM_DL_LIGHT_COUNT > 0) // Check if it is a light source bsdf->triangleLightSourceIndex = meshTriLightDefsOffset[meshIndex]; #endif #if defined(PARAM_HAS_BUMPMAPS) || defined(PARAM_HAS_NORMALMAPS) __global Material *mat = &mats[matIndex]; #if defined(PARAM_HAS_NORMALMAPS) //-------------------------------------------------------------------------- // Check if I have to apply normal mapping //-------------------------------------------------------------------------- const float3 normalColor = Material_GetNormalTexValue(mat, &bsdf->hitPoint MATERIALS_PARAM); if (!Spectrum_IsBlack(normalColor)) { const float3 xyz = 2.f * normalColor - 1.f; float3 v1, v2; CoordinateSystem(shadeN, &v1, &v2); shadeN = normalize((float3)( v1.x * xyz.x + v2.x * xyz.y + shadeN.x * xyz.z, v1.y * xyz.x + v2.y * xyz.y + shadeN.y * xyz.z, v1.z * xyz.x + v2.z * xyz.y + shadeN.z * xyz.z)); } #endif #if defined(PARAM_HAS_BUMPMAPS) //-------------------------------------------------------------------------- // Check if I have to apply bump mapping //-------------------------------------------------------------------------- const float2 bumpUV = Material_GetBumpTexValue(mat, &bsdf->hitPoint MATERIALS_PARAM); if ((bumpUV.s0 != 0.f) || (bumpUV.s1 != 0.f)) { float3 v1, v2; CoordinateSystem(shadeN, &v1, &v2); shadeN = normalize((float3)( v1.x * bumpUV.s0 + v2.x * bumpUV.s1 + shadeN.x, v1.y * bumpUV.s0 + v2.y * bumpUV.s1 + shadeN.y, v1.z * bumpUV.s0 + v2.z * bumpUV.s1 + shadeN.z)); } #endif #endif Frame_SetFromZ(&bsdf->frame, shadeN); VSTORE3F(shadeN, &bsdf->hitPoint.shadeN.x); } float3 BSDF_Evaluate(__global BSDF *bsdf, const float3 generatedDir, BSDFEvent *event, float *directPdfW MATERIALS_PARAM_DECL) { //const Vector &eyeDir = fromLight ? generatedDir : hitPoint.fixedDir; //const Vector &lightDir = fromLight ? hitPoint.fixedDir : generatedDir; const float3 eyeDir = VLOAD3F(&bsdf->hitPoint.fixedDir.x); const float3 lightDir = generatedDir; const float3 geometryN = VLOAD3F(&bsdf->hitPoint.geometryN.x); const float dotLightDirNG = dot(lightDir, geometryN); const float absDotLightDirNG = fabs(dotLightDirNG); const float dotEyeDirNG = dot(eyeDir, geometryN); const float absDotEyeDirNG = fabs(dotEyeDirNG); if ((absDotLightDirNG < DEFAULT_COS_EPSILON_STATIC) || (absDotEyeDirNG < DEFAULT_COS_EPSILON_STATIC)) return BLACK; __global Material *mat = &mats[bsdf->materialIndex]; const float sideTest = dotEyeDirNG * dotLightDirNG; const BSDFEvent matEvent = Material_GetEventTypes(mat MATERIALS_PARAM); if (((sideTest > 0.f) && !(matEvent & REFLECT)) || ((sideTest < 0.f) && !(matEvent & TRANSMIT))) return BLACK; __global Frame *frame = &bsdf->frame; const float3 localLightDir = Frame_ToLocal(frame, lightDir); const float3 localEyeDir = Frame_ToLocal(frame, eyeDir); const float3 result = Material_Evaluate(mat, &bsdf->hitPoint, localLightDir, localEyeDir, event, directPdfW MATERIALS_PARAM); // Adjoint BSDF // if (fromLight) { // const float absDotLightDirNS = AbsDot(lightDir, shadeN); // const float absDotEyeDirNS = AbsDot(eyeDir, shadeN); // return result * ((absDotLightDirNS * absDotEyeDirNG) / (absDotEyeDirNS * absDotLightDirNG)); // } else return result; } float3 BSDF_Sample(__global BSDF *bsdf, const float u0, const float u1, float3 *sampledDir, float *pdfW, float *cosSampledDir, BSDFEvent *event MATERIALS_PARAM_DECL) { const float3 fixedDir = VLOAD3F(&bsdf->hitPoint.fixedDir.x); const float3 localFixedDir = Frame_ToLocal(&bsdf->frame, fixedDir); float3 localSampledDir; const float3 result = Material_Sample(&mats[bsdf->materialIndex], &bsdf->hitPoint, localFixedDir, &localSampledDir, u0, u1, #if defined(PARAM_HAS_PASSTHROUGH) bsdf->hitPoint.passThroughEvent, #endif pdfW, cosSampledDir, event MATERIALS_PARAM); if (Spectrum_IsBlack(result)) return 0.f; *sampledDir = Frame_ToWorld(&bsdf->frame, localSampledDir); // Adjoint BSDF // if (fromLight) { // const float absDotFixedDirNS = fabsf(localFixedDir.z); // const float absDotSampledDirNS = fabsf(localSampledDir.z); // const float absDotFixedDirNG = AbsDot(fixedDir, geometryN); // const float absDotSampledDirNG = AbsDot(*sampledDir, geometryN); // return result * ((absDotFixedDirNS * absDotSampledDirNG) / (absDotSampledDirNS * absDotFixedDirNG)); // } else return result; } bool BSDF_IsDelta(__global BSDF *bsdf MATERIALS_PARAM_DECL) { return Material_IsDelta(&mats[bsdf->materialIndex] MATERIALS_PARAM); } #if (PARAM_DL_LIGHT_COUNT > 0) float3 BSDF_GetEmittedRadiance(__global BSDF *bsdf, __global TriangleLight *triLightDefs, float *directPdfA MATERIALS_PARAM_DECL) { const uint triangleLightSourceIndex = bsdf->triangleLightSourceIndex; if (triangleLightSourceIndex == NULL_INDEX) return BLACK; else return TriangleLight_GetRadiance(&triLightDefs[triangleLightSourceIndex], &bsdf->hitPoint, directPdfA MATERIALS_PARAM); } #endif #if defined(PARAM_HAS_PASSTHROUGH) float3 BSDF_GetPassThroughTransparency(__global BSDF *bsdf MATERIALS_PARAM_DECL) { const float3 localFixedDir = Frame_ToLocal(&bsdf->frame, VLOAD3F(&bsdf->hitPoint.fixedDir.x)); return Material_GetPassThroughTransparency(&mats[bsdf->materialIndex], &bsdf->hitPoint, localFixedDir, bsdf->hitPoint.passThroughEvent MATERIALS_PARAM); } #endif #line 2 "scene_funcs.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ float Scene_PickLightPdf() { #if defined(PARAM_HAS_SUNLIGHT) && (PARAM_DL_LIGHT_COUNT > 0) return 1.f / (1.f + PARAM_DL_LIGHT_COUNT); #elif defined(PARAM_HAS_SUNLIGHT) return 1.f; #elif (PARAM_DL_LIGHT_COUNT > 0) return 1.f / PARAM_DL_LIGHT_COUNT; #else return 0; #endif } #line 2 "datatypes.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ //------------------------------------------------------------------------------ // Some OpenCL specific definition //------------------------------------------------------------------------------ #if defined(SLG_OPENCL_KERNEL) #if defined(PARAM_USE_PIXEL_ATOMICS) #pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable #endif #if defined(PARAM_HAS_SUNLIGHT) & !defined(PARAM_DIRECT_LIGHT_SAMPLING) Error: PARAM_HAS_SUNLIGHT requires PARAM_DIRECT_LIGHT_SAMPLING ! #endif #ifndef TRUE #define TRUE 1 #endif #ifndef FALSE #define FALSE 0 #endif #endif //------------------------------------------------------------------------------ // GPUTask data types //------------------------------------------------------------------------------ typedef enum { RT_NEXT_VERTEX, GENERATE_DL_RAY, RT_DL, GENERATE_NEXT_VERTEX_RAY, SPLAT_SAMPLE } PathState; // This is defined only under OpenCL because of variable size structures #if defined(SLG_OPENCL_KERNEL) typedef struct { PathState state; unsigned int depth; Spectrum throughput; BSDF bsdf; // Variable size structure } PathStateBase; typedef struct { // Radiance to add to the result if light source is visible Spectrum lightRadiance; float lastPdfW; int lastSpecular; #if (PARAM_DL_LIGHT_COUNT > 0) // This is used by TriangleLight_Illuminate() to temporary store the // point on the light sources HitPoint tmpHitPoint; #endif } PathStateDirectLight; typedef struct { float passThroughEvent; // The passthrough sample used for the shadow ray BSDF passThroughBsdf; } PathStateDirectLightPassThrough; typedef struct { // The task seed Seed seed; // The set of Samples assigned to this task Sample sample; // The state used to keep track of the rendered path PathStateBase pathStateBase; #if defined(PARAM_DIRECT_LIGHT_SAMPLING) PathStateDirectLight directLightState; #if defined(PARAM_HAS_PASSTHROUGH) PathStateDirectLightPassThrough passThroughState; #endif #endif } GPUTask; #endif typedef struct { unsigned int sampleCount; } GPUTaskStats; #line 2 "patchocl_kernels.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ // List of symbols defined at compile time: // PARAM_TASK_COUNT // PARAM_IMAGE_WIDTH // PARAM_IMAGE_HEIGHT // PARAM_RAY_EPSILON_MIN // PARAM_RAY_EPSILON_MAX // PARAM_MAX_PATH_DEPTH // PARAM_RR_DEPTH // PARAM_RR_CAP // PARAM_HAS_IMAGEMAPS // PARAM_HAS_PASSTHROUGH // PARAM_USE_PIXEL_ATOMICS // PARAM_HAS_BUMPMAPS // PARAM_HAS_NORMALMAPS // PARAM_ACCEL_BVH or PARAM_ACCEL_MBVH or PARAM_ACCEL_QBVH or PARAM_ACCEL_MQBVH // PARAM_DEVICE_INDEX // PARAM_DEVICE_COUNT // To enable single material support // PARAM_ENABLE_MAT_MATTE // PARAM_ENABLE_MAT_MIRROR // PARAM_ENABLE_MAT_GLASS // PARAM_ENABLE_MAT_METAL // PARAM_ENABLE_MAT_ARCHGLASS // PARAM_ENABLE_MAT_MIX // PARAM_ENABLE_MAT_NULL // PARAM_ENABLE_MAT_MATTETRANSLUCENT // To enable single texture support // PARAM_ENABLE_TEX_CONST_FLOAT // PARAM_ENABLE_TEX_CONST_FLOAT3 // PARAM_ENABLE_TEX_CONST_FLOAT4 // PARAM_ENABLE_TEX_IMAGEMAP // PARAM_ENABLE_TEX_SCALE // (optional) // PARAM_DIRECT_LIGHT_SAMPLING // PARAM_DL_LIGHT_COUNT // (optional) // PARAM_CAMERA_HAS_DOF // (optional) // PARAM_HAS_INFINITELIGHT // (optional, requires PARAM_DIRECT_LIGHT_SAMPLING) // PARAM_HAS_SUNLIGHT // (optional) // PARAM_HAS_SKYLIGHT // (optional) // PARAM_IMAGE_FILTER_TYPE (0 = No filter, 1 = Box, 2 = Gaussian, 3 = Mitchell) // PARAM_IMAGE_FILTER_WIDTH_X // PARAM_IMAGE_FILTER_WIDTH_Y // (Box filter) // (Gaussian filter) // PARAM_IMAGE_FILTER_GAUSSIAN_ALPHA // (Mitchell filter) // PARAM_IMAGE_FILTER_MITCHELL_B // PARAM_IMAGE_FILTER_MITCHELL_C // (optional) // PARAM_SAMPLER_TYPE (0 = Inlined Random, 1 = Metropolis, 2 = Sobol) // (Metropolis) // PARAM_SAMPLER_METROPOLIS_LARGE_STEP_RATE // PARAM_SAMPLER_METROPOLIS_MAX_CONSECUTIVE_REJECT // PARAM_SAMPLER_METROPOLIS_IMAGE_MUTATION_RANGE // (Sobol) // PARAM_SAMPLER_SOBOL_STARTOFFSET // PARAM_SAMPLER_SOBOL_MAXDEPTH // (optional) // PARAM_ENABLE_ALPHA_CHANNEL // (optional) // PARAM_HAS_NORMALS_BUFFER // PARAM_HAS_UVS_BUFFER // PARAM_HAS_COLS_BUFFER // PARAM_HAS_ALPHAS_BUFFER //------------------------------------------------------------------------------ // Init Kernel //------------------------------------------------------------------------------ void GenerateCameraPath( __global GPUTask *task, __global float *sampleData, __global Camera *camera, __global Ray *ray, Seed *seed) { #if (PARAM_SAMPLER_TYPE == 0) const float scrSampleX = sampleData[IDX_SCREEN_X]; const float scrSampleY = sampleData[IDX_SCREEN_Y]; #if defined(PARAM_CAMERA_HAS_DOF) const float dofSampleX = Rnd_FloatValue(seed); const float dofSampleY = Rnd_FloatValue(seed); #endif #if defined(PARAM_HAS_PASSTHROUGH) const float eyePassthrough = Rnd_FloatValue(seed); #endif #endif #if (PARAM_SAMPLER_TYPE == 1) __global Sample *sample = &task->sample; __global float *sampleDataPathBase = Sampler_GetSampleDataPathBase(sample, sampleData); const float scrSampleX = Sampler_GetSamplePath(IDX_SCREEN_X); const float scrSampleY = Sampler_GetSamplePath(IDX_SCREEN_Y); #if defined(PARAM_CAMERA_HAS_DOF) const float dofSampleX = Sampler_GetSamplePath(IDX_DOF_X); const float dofSampleY = Sampler_GetSamplePath(IDX_DOF_Y); #endif #if defined(PARAM_HAS_PASSTHROUGH) const float eyePassthrough = Sampler_GetSamplePath(IDX_EYE_PASSTHROUGH); #endif #endif #if (PARAM_SAMPLER_TYPE == 2) __global Sample *sample = &task->sample; const float scrSampleX = sampleData[IDX_SCREEN_X]; const float scrSampleY = sampleData[IDX_SCREEN_Y]; #if defined(PARAM_CAMERA_HAS_DOF) const float dofSampleX = Sampler_GetSamplePath(IDX_DOF_X); const float dofSampleY = Sampler_GetSamplePath(IDX_DOF_Y); #endif #if defined(PARAM_HAS_PASSTHROUGH) const float eyePassthrough = Sampler_GetSamplePath(IDX_EYE_PASSTHROUGH); #endif #endif Camera_GenerateRay(camera, ray, scrSampleX, scrSampleY #if defined(PARAM_CAMERA_HAS_DOF) , dofSampleX, dofSampleY #endif ); // Initialize the path state task->pathStateBase.state = RT_NEXT_VERTEX; task->pathStateBase.depth = 1; VSTORE3F(WHITE, &task->pathStateBase.throughput.r); #if defined(PARAM_DIRECT_LIGHT_SAMPLING) task->directLightState.lastPdfW = 1.f; task->directLightState.lastSpecular = TRUE; #endif #if defined(PARAM_HAS_PASSTHROUGH) // This is a bit tricky. I store the passThroughEvent in the BSDF // before of the initialization because it can be use during the // tracing of next path vertex ray. task->pathStateBase.bsdf.hitPoint.passThroughEvent = eyePassthrough; #endif } __kernel __attribute__((work_group_size_hint(64, 1, 1))) void Init( uint seedBase, __global GPUTask *tasks, __global float *samplesData, __global GPUTaskStats *taskStats, __global Ray *rays, __global Camera *camera ) { const size_t gid = get_global_id(0); if (gid >= PARAM_TASK_COUNT) return; // Initialize the task __global GPUTask *task = &tasks[gid]; // Initialize random number generator Seed seed; Rnd_Init(seedBase + gid, &seed); // Initialize the sample and path __global Sample *sample = &task->sample; __global float *sampleData = Sampler_GetSampleData(sample, samplesData); Sampler_Init(&seed, sample, sampleData); GenerateCameraPath(task, sampleData, camera, &rays[gid], &seed); // Save the seed task->seed.s1 = seed.s1; task->seed.s2 = seed.s2; task->seed.s3 = seed.s3; __global GPUTaskStats *taskStat = &taskStats[gid]; taskStat->sampleCount = 0; } //------------------------------------------------------------------------------ // AdvancePaths Kernel //------------------------------------------------------------------------------ float RussianRouletteProb(const float3 color) { return clamp(Spectrum_Filter(color), PARAM_RR_CAP, 1.f); } __kernel __attribute__((work_group_size_hint(64, 1, 1))) void AdvancePaths( __global GPUTask *tasks, __global GPUTaskStats *taskStats, __global float *samplesData, __global Ray *rays, __global RayHit *rayHits, __global Pixel *frameBuffer, __global Material *mats, __global Texture *texs, __global uint *meshMats, __global Mesh *meshDescs, __global Point *vertices, #if defined(PARAM_HAS_NORMALS_BUFFER) __global Vector *vertNormals, #endif #if defined(PARAM_HAS_UVS_BUFFER) __global UV *vertUVs, #endif #if defined(PARAM_HAS_COLS_BUFFER) __global Spectrum *vertCols, #endif #if defined(PARAM_HAS_ALPHAS_BUFFER) __global float *vertAlphas, #endif __global Triangle *triangles, __global Camera *camera #if defined(PARAM_HAS_INFINITELIGHT) , __global InfiniteLight *infiniteLight #endif #if defined(PARAM_HAS_SUNLIGHT) , __global SunLight *sunLight #endif #if defined(PARAM_HAS_SKYLIGHT) , __global SkyLight *skyLight #endif #if (PARAM_DL_LIGHT_COUNT > 0) , __global TriangleLight *triLightDefs , __global uint *meshTriLightDefsOffset #endif #if defined(PARAM_IMAGEMAPS_PAGE_0) , __global ImageMap *imageMapDescs, __global float *imageMapBuff0 #endif #if defined(PARAM_IMAGEMAPS_PAGE_1) , __global float *imageMapBuff1 #endif #if defined(PARAM_IMAGEMAPS_PAGE_2) , __global float *imageMapBuff2 #endif #if defined(PARAM_IMAGEMAPS_PAGE_3) , __global float *imageMapBuff3 #endif #if defined(PARAM_IMAGEMAPS_PAGE_4) , __global float *imageMapBuff4 #endif #if defined(PARAM_IMAGEMAPS_PAGE_5) , __global float *imageMapBuff5 #endif #if defined(PARAM_IMAGEMAPS_PAGE_6) , __global float *imageMapBuff6 #endif #if defined(PARAM_IMAGEMAPS_PAGE_7) , __global float *imageMapBuff7 #endif #if defined(PARAM_ENABLE_ALPHA_CHANNEL) , __global AlphaPixel *alphaFrameBuffer #endif ) { const size_t gid = get_global_id(0); if (gid >= PARAM_TASK_COUNT) return; __global GPUTask *task = &tasks[gid]; // Read the path state PathState pathState = task->pathStateBase.state; const uint depth = task->pathStateBase.depth; __global BSDF *bsdf = &task->pathStateBase.bsdf; __global Sample *sample = &task->sample; __global float *sampleData = Sampler_GetSampleData(sample, samplesData); __global float *sampleDataPathBase = Sampler_GetSampleDataPathBase(sample, sampleData); #if (PARAM_SAMPLER_TYPE != 0) // Used by Sampler_GetSamplePathVertex() macro __global float *sampleDataPathVertexBase = Sampler_GetSampleDataPathVertex( sample, sampleDataPathBase, depth); #endif // Read the seed Seed seedValue; seedValue.s1 = task->seed.s1; seedValue.s2 = task->seed.s2; seedValue.s3 = task->seed.s3; // This trick is required by Sampler_GetSample() macro Seed *seed = &seedValue; #if defined(PARAM_HAS_IMAGEMAPS) // Initialize image maps page pointer table __global float *imageMapBuff[PARAM_IMAGEMAPS_COUNT]; #if defined(PARAM_IMAGEMAPS_PAGE_0) imageMapBuff[0] = imageMapBuff0; #endif #if defined(PARAM_IMAGEMAPS_PAGE_1) imageMapBuff[1] = imageMapBuff1; #endif #if defined(PARAM_IMAGEMAPS_PAGE_2) imageMapBuff[2] = imageMapBuff2; #endif #if defined(PARAM_IMAGEMAPS_PAGE_3) imageMapBuff[3] = imageMapBuff3; #endif #if defined(PARAM_IMAGEMAPS_PAGE_4) imageMapBuff[4] = imageMapBuff4; #endif #if defined(PARAM_IMAGEMAPS_PAGE_5) imageMapBuff[5] = imageMapBuff5; #endif #if defined(PARAM_IMAGEMAPS_PAGE_6) imageMapBuff[6] = imageMapBuff6; #endif #if defined(PARAM_IMAGEMAPS_PAGE_7) imageMapBuff[7] = imageMapBuff7; #endif #endif __global Ray *ray = &rays[gid]; __global RayHit *rayHit = &rayHits[gid]; const bool rayMiss = (rayHit->meshIndex == NULL_INDEX); //-------------------------------------------------------------------------- // Evaluation of the Path finite state machine. // // From: RT_NEXT_VERTEX // To: SPLAT_SAMPLE or GENERATE_DL_RAY //-------------------------------------------------------------------------- if (pathState == RT_NEXT_VERTEX) { if (!rayMiss) { // Something was hit BSDF_Init(bsdf, meshDescs, meshMats, #if (PARAM_DL_LIGHT_COUNT > 0) meshTriLightDefsOffset, #endif vertices, #if defined(PARAM_HAS_NORMALS_BUFFER) vertNormals, #endif #if defined(PARAM_HAS_UVS_BUFFER) vertUVs, #endif #if defined(PARAM_HAS_COLS_BUFFER) vertCols, #endif #if defined(PARAM_HAS_ALPHAS_BUFFER) vertAlphas, #endif triangles, ray, rayHit #if defined(PARAM_HAS_PASSTHROUGH) , task->pathStateBase.bsdf.hitPoint.passThroughEvent #endif #if defined(PARAM_HAS_BUMPMAPS) || defined(PARAM_HAS_NORMALMAPS) MATERIALS_PARAM #endif ); #if defined(PARAM_HAS_PASSTHROUGH) const float3 passThroughTrans = BSDF_GetPassThroughTransparency(bsdf MATERIALS_PARAM); if (!Spectrum_IsBlack(passThroughTrans)) { const float3 pathThroughput = VLOAD3F(&task->pathStateBase.throughput.r) * passThroughTrans; VSTORE3F(pathThroughput, &task->pathStateBase.throughput.r); // It is a pass through point, continue to trace the ray ray->mint = rayHit->t + MachineEpsilon_E(rayHit->t); // Keep the same path state } else #endif { #if (PARAM_DL_LIGHT_COUNT > 0) // Check if it is a light source (note: I can hit only triangle area light sources) if (bsdf->triangleLightSourceIndex != NULL_INDEX) { float directPdfA; const float3 emittedRadiance = BSDF_GetEmittedRadiance(bsdf, triLightDefs, &directPdfA MATERIALS_PARAM); if (!Spectrum_IsBlack(emittedRadiance)) { // Add emitted radiance float weight = 1.f; if (!task->directLightState.lastSpecular) { const float lightPickProb = Scene_PickLightPdf(); const float directPdfW = PdfAtoW(directPdfA, rayHit->t, fabs(dot(VLOAD3F(&bsdf->hitPoint.fixedDir.x), VLOAD3F(&bsdf->hitPoint.shadeN.x)))); // MIS between BSDF sampling and direct light sampling weight = PowerHeuristic(task->directLightState.lastPdfW, directPdfW * lightPickProb); } float3 radiance = VLOAD3F(&sample->radiance.r); const float3 pathThroughput = VLOAD3F(&task->pathStateBase.throughput.r); const float3 le = pathThroughput * weight * emittedRadiance; radiance += le; VSTORE3F(radiance, &sample->radiance.r); } } #endif #if defined(PARAM_HAS_SUNLIGHT) || (PARAM_DL_LIGHT_COUNT > 0) // Direct light sampling pathState = GENERATE_DL_RAY; #else // Sample next path vertex pathState = GENERATE_NEXT_VERTEX_RAY; #endif } } else { //------------------------------------------------------------------ // Nothing was hit, get environmental lights radiance //------------------------------------------------------------------ #if defined(PARAM_HAS_INFINITELIGHT) || defined(PARAM_HAS_SKYLIGHT) || defined(PARAM_HAS_SUNLIGHT) float3 radiance = VLOAD3F(&sample->radiance.r); const float3 pathThroughput = VLOAD3F(&task->pathStateBase.throughput.r); const float3 dir = -VLOAD3F(&ray->d.x); float3 lightRadiance = BLACK; #if defined(PARAM_HAS_INFINITELIGHT) lightRadiance += InfiniteLight_GetRadiance(infiniteLight, dir IMAGEMAPS_PARAM); #endif #if defined(PARAM_HAS_SKYLIGHT) lightRadiance += SkyLight_GetRadiance(skyLight, dir); #endif #if defined(PARAM_HAS_SUNLIGHT) float directPdfW; const float3 sunRadiance = SunLight_GetRadiance(sunLight, dir, &directPdfW); if (!Spectrum_IsBlack(sunRadiance)) { // MIS between BSDF sampling and direct light sampling const float weight = (task->directLightState.lastSpecular ? 1.f : PowerHeuristic(task->directLightState.lastPdfW, directPdfW)); lightRadiance += weight * sunRadiance; } #endif radiance += pathThroughput * lightRadiance; VSTORE3F(radiance, &sample->radiance.r); #endif #if defined(PARAM_ENABLE_ALPHA_CHANNEL) if (depth == 1) sample->alpha = 0.f; #endif pathState = SPLAT_SAMPLE; } } //-------------------------------------------------------------------------- // Evaluation of the Path finite state machine. // // From: RT_DL // To: GENERATE_NEXT_VERTEX_RAY //-------------------------------------------------------------------------- #if defined(PARAM_HAS_SUNLIGHT) || (PARAM_DL_LIGHT_COUNT > 0) if (pathState == RT_DL) { pathState = GENERATE_NEXT_VERTEX_RAY; if (rayMiss) { // Nothing was hit, the light source is visible float3 radiance = VLOAD3F(&sample->radiance.r); const float3 lightRadiance = VLOAD3F(&task->directLightState.lightRadiance.r); radiance += lightRadiance; VSTORE3F(radiance, &sample->radiance.r); } #if defined(PARAM_HAS_PASSTHROUGH) else { BSDF_Init(&task->passThroughState.passThroughBsdf, meshDescs, meshMats, #if (PARAM_DL_LIGHT_COUNT > 0) meshTriLightDefsOffset, #endif vertices, #if defined(PARAM_HAS_NORMALS_BUFFER) vertNormals, #endif #if defined(PARAM_HAS_UVS_BUFFER) vertUVs, #endif #if defined(PARAM_HAS_COLS_BUFFER) vertCols, #endif #if defined(PARAM_HAS_ALPHAS_BUFFER) vertAlphas, #endif triangles, ray, rayHit, task->passThroughState.passThroughEvent #if defined(PARAM_HAS_BUMPMAPS) || defined(PARAM_HAS_NORMALMAPS) MATERIALS_PARAM #endif ); const float3 passthroughTrans = BSDF_GetPassThroughTransparency(&task->passThroughState.passThroughBsdf MATERIALS_PARAM); if (!Spectrum_IsBlack(passthroughTrans)) { const float3 lightRadiance = VLOAD3F(&task->directLightState.lightRadiance.r) * passthroughTrans; VSTORE3F(lightRadiance, &task->directLightState.lightRadiance.r); // It is a pass through point, continue to trace the ray ray->mint = rayHit->t + MachineEpsilon_E(rayHit->t); pathState = RT_DL; } } #endif } #endif //-------------------------------------------------------------------------- // Evaluation of the Path finite state machine. // // From: GENERATE_DL_RAY // To: GENERATE_NEXT_VERTEX_RAY or RT_DL //-------------------------------------------------------------------------- #if defined(PARAM_HAS_SUNLIGHT) || (PARAM_DL_LIGHT_COUNT > 0) if (pathState == GENERATE_DL_RAY) { pathState = GENERATE_NEXT_VERTEX_RAY; if (!BSDF_IsDelta(bsdf MATERIALS_PARAM)) { float3 lightRayDir; float distance, directPdfW; float3 lightRadiance; #if defined(PARAM_HAS_SUNLIGHT) && (PARAM_DL_LIGHT_COUNT > 0) // Pick a light source to sample const float lu0 = Sampler_GetSamplePathVertex(depth, IDX_DIRECTLIGHT_X); const float lightPickPdf = Scene_PickLightPdf(); const uint lightIndex = min((uint)floor((PARAM_DL_LIGHT_COUNT + 1) * lu0), (uint)(PARAM_DL_LIGHT_COUNT)); if (lightIndex == PARAM_DL_LIGHT_COUNT) { lightRadiance = SunLight_Illuminate( sunLight, Sampler_GetSamplePathVertex(depth, IDX_DIRECTLIGHT_Y), Sampler_GetSamplePathVertex(depth, IDX_DIRECTLIGHT_Z), &lightRayDir, &distance, &directPdfW); } else { lightRadiance = TriangleLight_Illuminate( &triLightDefs[lightIndex], &task->directLightState.tmpHitPoint, VLOAD3F(&bsdf->hitPoint.p.x), Sampler_GetSamplePathVertex(depth, IDX_DIRECTLIGHT_Y), Sampler_GetSamplePathVertex(depth, IDX_DIRECTLIGHT_Z), Sampler_GetSamplePathVertex(depth, IDX_DIRECTLIGHT_W), &lightRayDir, &distance, &directPdfW MATERIALS_PARAM); } #elif (PARAM_DL_LIGHT_COUNT > 0) // Pick a light source to sample const float lu0 = Sampler_GetSamplePathVertex(depth, IDX_DIRECTLIGHT_X); const float lightPickPdf = Scene_PickLightPdf(); const uint lightIndex = min((uint)floor(PARAM_DL_LIGHT_COUNT * lu0), (uint)(PARAM_DL_LIGHT_COUNT - 1)); lightRadiance = TriangleLight_Illuminate( &triLightDefs[lightIndex], &task->directLightState.tmpHitPoint, VLOAD3F(&bsdf->hitPoint.p.x), Sampler_GetSamplePathVertex(depth, IDX_DIRECTLIGHT_Y), Sampler_GetSamplePathVertex(depth, IDX_DIRECTLIGHT_Z), Sampler_GetSamplePathVertex(depth, IDX_DIRECTLIGHT_W), &lightRayDir, &distance, &directPdfW MATERIALS_PARAM); #elif defined(PARAM_HAS_SUNLIGHT) // Pick a light source to sample const float lightPickPdf = 1.f; lightRadiance = SunLight_Illuminate( sunLight, Sampler_GetSamplePathVertex(depth, IDX_DIRECTLIGHT_Y), Sampler_GetSamplePathVertex(depth, IDX_DIRECTLIGHT_Z), &lightRayDir, &distance, &directPdfW); #endif // Setup the shadow ray if (!Spectrum_IsBlack(lightRadiance)) { BSDFEvent event; float bsdfPdfW; const float3 bsdfEval = BSDF_Evaluate(bsdf, lightRayDir, &event, &bsdfPdfW MATERIALS_PARAM); if (!Spectrum_IsBlack(bsdfEval)) { const float3 pathThroughput = VLOAD3F(&task->pathStateBase.throughput.r); const float cosThetaToLight = fabs(dot(lightRayDir, VLOAD3F(&bsdf->hitPoint.shadeN.x))); const float directLightSamplingPdfW = directPdfW * lightPickPdf; const float factor = cosThetaToLight / directLightSamplingPdfW; // Russian Roulette bsdfPdfW *= (depth >= PARAM_RR_DEPTH) ? RussianRouletteProb(bsdfEval) : 1.f; // MIS between direct light sampling and BSDF sampling const float weight = PowerHeuristic(directLightSamplingPdfW, bsdfPdfW); VSTORE3F((weight * factor) * pathThroughput * bsdfEval * lightRadiance, &task->directLightState.lightRadiance.r); #if defined(PARAM_HAS_PASSTHROUGH) task->passThroughState.passThroughEvent = Sampler_GetSamplePathVertex(depth, IDX_DIRECTLIGHT_A); #endif // Setup the shadow ray const float3 hitPoint = VLOAD3F(&bsdf->hitPoint.p.x); const float epsilon = fmax(MachineEpsilon_E_Float3(hitPoint), MachineEpsilon_E(distance)); Ray_Init4(ray, hitPoint, lightRayDir, epsilon, distance - epsilon); pathState = RT_DL; } } } } #endif //-------------------------------------------------------------------------- // Evaluation of the Path finite state machine. // // From: GENERATE_NEXT_VERTEX_RAY // To: SPLAT_SAMPLE or RT_NEXT_VERTEX //-------------------------------------------------------------------------- if (pathState == GENERATE_NEXT_VERTEX_RAY) { if (depth < PARAM_MAX_PATH_DEPTH) { // Sample the BSDF __global BSDF *bsdf = &task->pathStateBase.bsdf; float3 sampledDir; float lastPdfW; float cosSampledDir; BSDFEvent event; const float3 bsdfSample = BSDF_Sample(bsdf, Sampler_GetSamplePathVertex(depth, IDX_BSDF_X), Sampler_GetSamplePathVertex(depth, IDX_BSDF_Y), &sampledDir, &lastPdfW, &cosSampledDir, &event MATERIALS_PARAM); const bool lastSpecular = ((event & SPECULAR) != 0); // Russian Roulette const float rrProb = RussianRouletteProb(bsdfSample); const bool rrEnabled = (depth >= PARAM_RR_DEPTH) && !lastSpecular; const bool rrContinuePath = !rrEnabled || (Sampler_GetSamplePathVertex(depth, IDX_RR) < rrProb); const bool continuePath = !Spectrum_IsBlack(bsdfSample) && rrContinuePath; if (continuePath) { if (rrEnabled) lastPdfW *= rrProb; // Russian Roulette float3 throughput = VLOAD3F(&task->pathStateBase.throughput.r); throughput *= bsdfSample * (cosSampledDir / lastPdfW); VSTORE3F(throughput, &task->pathStateBase.throughput.r); Ray_Init2(ray, VLOAD3F(&bsdf->hitPoint.p.x), sampledDir); task->pathStateBase.depth = depth + 1; #if defined(PARAM_HAS_SUNLIGHT) || (PARAM_DL_LIGHT_COUNT > 0) task->directLightState.lastPdfW = lastPdfW; task->directLightState.lastSpecular = lastSpecular; #endif #if defined(PARAM_HAS_PASSTHROUGH) // This is a bit tricky. I store the passThroughEvent in the BSDF // before of the initialization because it can be use during the // tracing of next path vertex ray. // This sampleDataPathVertexBase is used inside Sampler_GetSamplePathVertex() macro __global float *sampleDataPathVertexBase = Sampler_GetSampleDataPathVertex( sample, sampleDataPathBase, depth + 1); task->pathStateBase.bsdf.hitPoint.passThroughEvent = Sampler_GetSamplePathVertex(depth + 1, IDX_PASSTHROUGH); #endif pathState = RT_NEXT_VERTEX; } else pathState = SPLAT_SAMPLE; } else pathState = SPLAT_SAMPLE; } //-------------------------------------------------------------------------- // Evaluation of the Path finite state machine. // // From: SPLAT_SAMPLE // To: RT_NEXT_VERTEX //-------------------------------------------------------------------------- if (pathState == SPLAT_SAMPLE) { Sampler_NextSample(sample, sampleData, seed, frameBuffer #if defined(PARAM_ENABLE_ALPHA_CHANNEL) , alphaFrameBuffer #endif ); taskStats[gid].sampleCount += 1; GenerateCameraPath(task, sampleData, camera, ray, seed); // task->pathStateBase.state is set to RT_NEXT_VERTEX inside Sampler_NextSample() => GenerateCameraPath() } else { // Save the state task->pathStateBase.state = pathState; } //-------------------------------------------------------------------------- // Save the seed task->seed.s1 = seed->s1; task->seed.s2 = seed->s2; task->seed.s3 = seed->s3; } #line 2 "rtpatchocl_kernels.cl" /*************************************************************************** * Copyright (C) 1998-2013 by authors (see AUTHORS.txt) * * * * This file is part of LuxRays. * * * * LuxRays 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 3 of the License, or * * (at your option) any later version. * * * * LuxRays 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. * * * * You should have received a copy of the GNU General Public License * * along with this program. If not, see . * * * * LuxRays website: http://www.luxrender.net * ***************************************************************************/ #define FRAMEBUFFER_WIDTH (PARAM_IMAGE_WIDTH + 2) #define FRAMEBUFFER_HEIGHT (PARAM_IMAGE_HEIGHT + 2) //------------------------------------------------------------------------------ // ClearFrameBuffer Kernel //------------------------------------------------------------------------------ __kernel __attribute__((work_group_size_hint(64, 1, 1))) void ClearFrameBuffer( __global Pixel *frameBuffer) { const size_t gid = get_global_id(0); if (gid >= FRAMEBUFFER_WIDTH * FRAMEBUFFER_HEIGHT) return; VSTORE4F(0.f, &frameBuffer[gid].c.r); } //------------------------------------------------------------------------------ // ClearScreenBuffer Kernel //------------------------------------------------------------------------------ __kernel __attribute__((work_group_size_hint(64, 1, 1))) void ClearScreenBuffer( __global Spectrum *screenBuffer) { const size_t gid = get_global_id(0); if (gid >= PARAM_IMAGE_WIDTH * PARAM_IMAGE_HEIGHT) return; VSTORE3F(0.f, &screenBuffer[gid].r); } //------------------------------------------------------------------------------ // NormalizeFrameBuffer Kernel //------------------------------------------------------------------------------ __kernel __attribute__((work_group_size_hint(64, 1, 1))) void NormalizeFrameBuffer( __global Pixel *frameBuffer) { const size_t gid = get_global_id(0); if (gid >= FRAMEBUFFER_WIDTH * FRAMEBUFFER_HEIGHT) return; float4 rgbc = VLOAD4F(&frameBuffer[gid].c.r); if (rgbc.w > 0.f) { const float k = 1.f / rgbc.w; rgbc.s0 *= k; rgbc.s1 *= k; rgbc.s2 *= k; VSTORE4F(rgbc, &frameBuffer[gid].c.r); } else VSTORE4F(0.f, &frameBuffer[gid].c.r); } //------------------------------------------------------------------------------ // MergeFrameBuffer Kernel //------------------------------------------------------------------------------ __kernel __attribute__((work_group_size_hint(64, 1, 1))) void MergeFrameBuffer( __global Pixel *srcFrameBuffer, __global Pixel *dstFrameBuffer) { const size_t gid = get_global_id(0); if (gid >= FRAMEBUFFER_WIDTH * FRAMEBUFFER_HEIGHT) return; float4 srcRGBC = VLOAD4F(&srcFrameBuffer[gid].c.r); if (srcRGBC.w > 0.f) { const float4 dstRGBC = VLOAD4F(&dstFrameBuffer[gid].c.r); VSTORE4F(srcRGBC + dstRGBC, &dstFrameBuffer[gid].c.r); } } //------------------------------------------------------------------------------ // Image filtering kernels //------------------------------------------------------------------------------ void ApplyBlurFilterXR1( __global Spectrum *src, __global Spectrum *dst, const float aF, const float bF, const float cF ) { // Do left edge float3 a; float3 b = VLOAD3F(&src[0].r); float3 c = VLOAD3F(&src[1].r); const float leftTotF = bF + cF; const float3 bLeftK = bF / leftTotF; const float3 cLeftK = cF / leftTotF; VSTORE3F(bLeftK * b + cLeftK * c, &dst[0].r); // Main loop const float totF = aF + bF + cF; const float3 aK = aF / totF; const float3 bK = bF / totF; const float3 cK = cF / totF; for (unsigned int x = 1; x < PARAM_IMAGE_WIDTH - 1; ++x) { a = b; b = c; c = VLOAD3F(&src[x + 1].r); VSTORE3F(aK * a + bK * b + cK * c, &dst[x].r); } // Do right edge const float rightTotF = aF + bF; const float3 aRightK = aF / rightTotF; const float3 bRightK = bF / rightTotF; a = b; b = c; VSTORE3F(aRightK * a + bRightK * b, &dst[PARAM_IMAGE_WIDTH - 1].r); } void ApplyBlurFilterYR1( __global Spectrum *src, __global Spectrum *dst, const float aF, const float bF, const float cF ) { // Do left edge float3 a; float3 b = VLOAD3F(&src[0].r); float3 c = VLOAD3F(&src[PARAM_IMAGE_WIDTH].r); const float leftTotF = bF + cF; const float3 bLeftK = bF / leftTotF; const float3 cLeftK = cF / leftTotF; VSTORE3F(bLeftK * b + cLeftK * c, &dst[0].r); // Main loop const float totF = aF + bF + cF; const float3 aK = aF / totF; const float3 bK = bF / totF; const float3 cK = cF / totF; for (unsigned int y = 1; y < PARAM_IMAGE_HEIGHT - 1; ++y) { const unsigned index = y * PARAM_IMAGE_WIDTH; a = b; b = c; c = VLOAD3F(&src[index + PARAM_IMAGE_WIDTH].r); VSTORE3F(aK * a + bK * b + cK * c, &dst[index].r); } // Do right edge const float rightTotF = aF + bF; const float3 aRightK = aF / rightTotF; const float3 bRightK = bF / rightTotF; a = b; b = c; VSTORE3F(aRightK * a + bRightK * b, &dst[(PARAM_IMAGE_HEIGHT - 1) * PARAM_IMAGE_WIDTH].r); } __kernel __attribute__((work_group_size_hint(64, 1, 1))) void ApplyGaussianBlurFilterXR1( __global Spectrum *src, __global Spectrum *dst, const float weight ) { const size_t gid = get_global_id(0); if (gid >= PARAM_IMAGE_HEIGHT) return; src += gid * PARAM_IMAGE_WIDTH; dst += gid * PARAM_IMAGE_WIDTH; const float aF = .15f; const float bF = 1.f; const float cF = .15f; ApplyBlurFilterXR1(src, dst, aF, bF, cF); } __kernel __attribute__((work_group_size_hint(64, 1, 1))) void ApplyGaussianBlurFilterYR1( __global Spectrum *src, __global Spectrum *dst, const float weight ) { const size_t gid = get_global_id(0); if (gid >= PARAM_IMAGE_WIDTH) return; src += gid; dst += gid; const float aF = weight; const float bF = 1.f; const float cF = weight; ApplyBlurFilterYR1(src, dst, aF, bF, cF); } //------------------------------------------------------------------------------ // Linear Tone Map Kernel //------------------------------------------------------------------------------ __kernel __attribute__((work_group_size_hint(64, 1, 1))) void ToneMapLinear( __global Pixel *src, __global Pixel *dst) { const int gid = get_global_id(0); if (gid >= FRAMEBUFFER_WIDTH * FRAMEBUFFER_HEIGHT) return; const float4 k = (float4)(PARAM_TONEMAP_LINEAR_SCALE, PARAM_TONEMAP_LINEAR_SCALE, PARAM_TONEMAP_LINEAR_SCALE, 1.f); const float4 sp = VLOAD4F(&src[gid].c.r); VSTORE4F(k * sp, &dst[gid].c.r); } //------------------------------------------------------------------------------ // UpdateScreenBuffer Kernel //------------------------------------------------------------------------------ float Radiance2PixelFloat(const float x) { return pow(clamp(x, 0.f, 1.f), 1.f / PARAM_GAMMA); } __kernel void UpdateScreenBuffer( __global Pixel *srcFrameBuffer, __global Spectrum *screenBuffer) { const int gid = get_global_id(0); if (gid >= FRAMEBUFFER_WIDTH * FRAMEBUFFER_HEIGHT) return; const int x = gid % FRAMEBUFFER_WIDTH - 2; const int y = gid / FRAMEBUFFER_WIDTH - 2; if ((x < 0) || (y < 0) || (x >= PARAM_IMAGE_WIDTH) || (y >= PARAM_IMAGE_HEIGHT)) return; float4 newRgbc = VLOAD4F(&srcFrameBuffer[gid].c.r); if (newRgbc.s3 > 0.f) { float3 newRgb = (float3)( newRgb.s0 = Radiance2PixelFloat(newRgbc.s0), newRgb.s1 = Radiance2PixelFloat(newRgbc.s1), newRgb.s2 = Radiance2PixelFloat(newRgbc.s2)); const float3 oldRgb = VLOAD3F(&screenBuffer[(x + y * PARAM_IMAGE_WIDTH)].r); // Blend old and new RGB value in for ghost effect VSTORE3F(mix(oldRgb, newRgb, .85f), &screenBuffer[(x + y * PARAM_IMAGE_WIDTH)].r); } } #endif