2
* Copyright 2011, Blender Foundation.
4
* This program is free software; you can redistribute it and/or
5
* modify it under the terms of the GNU General Public License
6
* as published by the Free Software Foundation; either version 2
7
* of the License, or (at your option) any later version.
9
* This program is distributed in the hope that it will be useful,
10
* but WITHOUT ANY WARRANTY; without even the implied warranty of
11
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
12
* GNU General Public License for more details.
14
* You should have received a copy of the GNU General Public License
15
* along with this program; if not, write to the Free Software Foundation,
16
* Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
2
* Copyright 2011-2013 Blender Foundation
4
* Licensed under the Apache License, Version 2.0 (the "License");
5
* you may not use this file except in compliance with the License.
6
* You may obtain a copy of the License at
8
* http://www.apache.org/licenses/LICENSE-2.0
10
* Unless required by applicable law or agreed to in writing, software
11
* distributed under the License is distributed on an "AS IS" BASIS,
12
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13
* See the License for the specific language governing permissions and
14
* limitations under the License
19
17
/* CUDA kernel entry points */
35
33
kernel_path_trace(NULL, buffer, rng_state, sample, x, y, offset, stride);
38
extern "C" __global__ void kernel_cuda_tonemap(uchar4 *rgba, float *buffer, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
40
int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
41
int y = sy + blockDim.y*blockIdx.y + threadIdx.y;
43
if(x < sx + sw && y < sy + sh)
44
kernel_film_tonemap(NULL, rgba, buffer, sample, x, y, offset, stride);
36
#ifdef __BRANCHED_PATH__
37
extern "C" __global__ void kernel_cuda_branched_path_trace(float *buffer, uint *rng_state, int sample, int sx, int sy, int sw, int sh, int offset, int stride)
39
int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
40
int y = sy + blockDim.y*blockIdx.y + threadIdx.y;
42
if(x < sx + sw && y < sy + sh)
43
kernel_branched_path_trace(NULL, buffer, rng_state, sample, x, y, offset, stride);
47
extern "C" __global__ void kernel_cuda_convert_to_byte(uchar4 *rgba, float *buffer, float sample_scale, int sx, int sy, int sw, int sh, int offset, int stride)
49
int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
50
int y = sy + blockDim.y*blockIdx.y + threadIdx.y;
52
if(x < sx + sw && y < sy + sh)
53
kernel_film_convert_to_byte(NULL, rgba, buffer, sample_scale, x, y, offset, stride);
56
extern "C" __global__ void kernel_cuda_convert_to_half_float(uchar4 *rgba, float *buffer, float sample_scale, int sx, int sy, int sw, int sh, int offset, int stride)
58
int x = sx + blockDim.x*blockIdx.x + threadIdx.x;
59
int y = sy + blockDim.y*blockIdx.y + threadIdx.y;
61
if(x < sx + sw && y < sy + sh)
62
kernel_film_convert_to_half_float(NULL, rgba, buffer, sample_scale, x, y, offset, stride);
47
65
extern "C" __global__ void kernel_cuda_shader(uint4 *input, float4 *output, int type, int sx)
added
removed
intern/cycles/kernel/kernel.cu
