kannonier8/blender
1
0
Fork 0
forked from bartvdbraak/blender
Code Pull Requests Activity

Cycles: Make it possible to access KernelGlobals from split data initialization function

Browse Source
This commit is contained in:
Sergey Sharybin 2017-03-08 11:02:54 +01:00
parent ef7c36f5ed
commit 712f7c3640
10 changed files with 125 additions and 97 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
intern/cycles/device/device_cuda.cpp
View File

@ -44,7 +44,7 @@
#include "util_types.h"
#include "util_time.h"
#include "split/kernel_split_data.h"
#include "split/kernel_split_data_types.h"
CCL_NAMESPACE_BEGIN

2
intern/cycles/device/device_split_kernel.cpp
View File

@ -17,7 +17,7 @@
#include "device_split_kernel.h"
#include "kernel_types.h"
#include "kernel_split_data.h"
#include "kernel_split_data_types.h"
#include "util_time.h"

2
intern/cycles/device/opencl/opencl_split.cpp
View File

@ -21,7 +21,7 @@
#include "buffers.h"
#include "kernel_types.h"
#include "kernel_split_data.h"
#include "kernel_split_data_types.h"
#include "device_split_kernel.h"

1
intern/cycles/kernel/CMakeLists.txt
View File

@ -209,6 +209,7 @@ set(SRC_SPLIT_HEADERS
split/kernel_shadow_blocked.h
split/kernel_split_common.h
split/kernel_split_data.h
split/kernel_split_data_types.h
)
# CUDA module

4
intern/cycles/kernel/kernel_globals.h
View File

@ -16,6 +16,9 @@
/* Constant Globals */
#ifndef __KERNEL_GLOBALS_H__
#define __KERNEL_GLOBALS_H__
CCL_NAMESPACE_BEGIN
/* On the CPU, we pass along the struct KernelGlobals to nearly everywhere in
@ -153,3 +156,4 @@ ccl_device float lookup_table_read_2D(KernelGlobals *kg, float x, float y, int o
CCL_NAMESPACE_END
#endif /* __KERNEL_GLOBALS_H__ */

2
intern/cycles/kernel/osl/osl_closures.cpp
View File

@ -42,7 +42,7 @@
#include "kernel_types.h"
#include "kernel_compat_cpu.h"
#include "split/kernel_split_data.h"
#include "split/kernel_split_data_types.h"
#include "kernel_globals.h"
#include "kernel_montecarlo.h"
#include "kernel_random.h"

2
intern/cycles/kernel/osl/osl_services.cpp
View File

@ -39,7 +39,7 @@
#include "util_string.h"
#include "kernel_compat_cpu.h"
#include "split/kernel_split_data.h"
#include "split/kernel_split_data_types.h"
#include "kernel_globals.h"
#include "kernel_random.h"
#include "kernel_projection.h"

2
intern/cycles/kernel/osl/osl_shader.cpp
View File

@ -19,7 +19,7 @@
#include "kernel_compat_cpu.h"
#include "kernel_montecarlo.h"
#include "kernel_types.h"
#include "split/kernel_split_data.h"
#include "split/kernel_split_data_types.h"
#include "kernel_globals.h"
#include "geom/geom_object.h"

96
intern/cycles/kernel/split/kernel_split_data.h
View File

@ -17,85 +17,12 @@
#ifndef __KERNEL_SPLIT_DATA_H__
#define __KERNEL_SPLIT_DATA_H__
#include "kernel_split_data_types.h"
#include "kernel_globals.h"
CCL_NAMESPACE_BEGIN
/* parameters used by the split kernels, we use a single struct to avoid passing these to each kernel */
typedef struct SplitParams {
int x;
int y;
int w;
int h;
int offset;
int stride;
ccl_global uint *rng_state;
int start_sample;
int end_sample;
ccl_global unsigned int *work_pools;
unsigned int num_samples;
ccl_global int *queue_index;
int queue_size;
ccl_global char *use_queues_flag;
ccl_global float *buffer;
} SplitParams;
/* Global memory variables [porting]; These memory is used for
* co-operation between different kernels; Data written by one
* kernel will be available to another kernel via this global
* memory.
*/
/* SPLIT_DATA_ENTRY(type, name, num) */
#if defined(WITH_CYCLES_DEBUG) || defined(__KERNEL_DEBUG__)
/* DebugData memory */
# define SPLIT_DATA_DEBUG_ENTRIES \
SPLIT_DATA_ENTRY(DebugData, debug_data, 1)
#else
# define SPLIT_DATA_DEBUG_ENTRIES
#endif
#define SPLIT_DATA_ENTRIES \
SPLIT_DATA_ENTRY(ccl_global RNG, rng, 1) \
SPLIT_DATA_ENTRY(ccl_global float3, throughput, 1) \
SPLIT_DATA_ENTRY(ccl_global float, L_transparent, 1) \
SPLIT_DATA_ENTRY(PathRadiance, path_radiance, 1) \
SPLIT_DATA_ENTRY(ccl_global Ray, ray, 1) \
SPLIT_DATA_ENTRY(ccl_global PathState, path_state, 1) \
SPLIT_DATA_ENTRY(Intersection, isect, 1) \
SPLIT_DATA_ENTRY(ccl_global float3, ao_alpha, 1) \
SPLIT_DATA_ENTRY(ccl_global float3, ao_bsdf, 1) \
SPLIT_DATA_ENTRY(ccl_global Ray, ao_light_ray, 1) \
SPLIT_DATA_ENTRY(ccl_global BsdfEval, bsdf_eval, 1) \
SPLIT_DATA_ENTRY(ccl_global int, is_lamp, 1) \
SPLIT_DATA_ENTRY(ccl_global Ray, light_ray, 1) \
SPLIT_DATA_ENTRY(Intersection, isect_shadow, 2) \
SPLIT_DATA_ENTRY(ccl_global int, queue_data, (NUM_QUEUES*2)) /* TODO(mai): this is too large? */ \
SPLIT_DATA_ENTRY(ccl_global uint, work_array, 1) \
SPLIT_DATA_ENTRY(ShaderData, sd, 1) \
SPLIT_DATA_ENTRY(ShaderData, sd_DL_shadow, 2) \
SPLIT_DATA_DEBUG_ENTRIES \
/* struct that holds pointers to data in the shared state buffer */
typedef struct SplitData {
#define SPLIT_DATA_ENTRY(type, name, num) type *name;
SPLIT_DATA_ENTRIES
#undef SPLIT_DATA_ENTRY
/* this is actually in a separate buffer from the rest of the split state data (so it can be read back from
* the host easily) but is still used the same as the other data so we have it here in this struct as well
*/
ccl_global char *ray_state;
} SplitData;
/* TODO: find a way to get access to kg here */
ccl_device_inline size_t split_data_buffer_size(ccl_global void *kg, size_t num_elements)
ccl_device_inline size_t split_data_buffer_size(KernelGlobals *kg, size_t num_elements)
{
(void)kg; /* Unused on CPU. */
@ -107,7 +34,7 @@ ccl_device_inline size_t split_data_buffer_size(ccl_global void *kg, size_t num_
return size;
}
ccl_device_inline void split_data_init(ccl_global void *kg,
ccl_device_inline void split_data_init(KernelGlobals *kg,
ccl_global SplitData *split_data,
size_t num_elements,
ccl_global void *data,
@ -125,19 +52,6 @@ ccl_device_inline void split_data_init(ccl_global void *kg,
split_data->ray_state = ray_state;
}
#ifndef __KERNEL_CUDA__
# define kernel_split_state (kg->split_data)
# define kernel_split_params (kg->split_param_data)
#else
__device__ SplitData __split_data;
# define kernel_split_state (__split_data)
__device__ SplitParams __split_param_data;
# define kernel_split_params (__split_param_data)
#endif /* __KERNEL_CUDA__ */
CCL_NAMESPACE_END
#endif /* __KERNEL_SPLIT_DATA_H__ */

109
intern/cycles/kernel/split/kernel_split_data_types.h Normal file
View File

@ -0,0 +1,109 @@
/*
* Copyright 2011-2016 Blender Foundation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/
#ifndef __KERNEL_SPLIT_DATA_TYPES_H__
#define __KERNEL_SPLIT_DATA_TYPES_H__
CCL_NAMESPACE_BEGIN
/* parameters used by the split kernels, we use a single struct to avoid passing these to each kernel */
typedef struct SplitParams {
int x;
int y;
int w;
int h;
int offset;
int stride;
ccl_global uint *rng_state;
int start_sample;
int end_sample;
ccl_global unsigned int *work_pools;
unsigned int num_samples;
ccl_global int *queue_index;
int queue_size;
ccl_global char *use_queues_flag;
ccl_global float *buffer;
} SplitParams;
/* Global memory variables [porting]; These memory is used for
* co-operation between different kernels; Data written by one
* kernel will be available to another kernel via this global
* memory.
*/
/* SPLIT_DATA_ENTRY(type, name, num) */
#if defined(WITH_CYCLES_DEBUG) || defined(__KERNEL_DEBUG__)
/* DebugData memory */
# define SPLIT_DATA_DEBUG_ENTRIES \
SPLIT_DATA_ENTRY(DebugData, debug_data, 1)
#else
# define SPLIT_DATA_DEBUG_ENTRIES
#endif
#define SPLIT_DATA_ENTRIES \
SPLIT_DATA_ENTRY(ccl_global RNG, rng, 1) \
SPLIT_DATA_ENTRY(ccl_global float3, throughput, 1) \
SPLIT_DATA_ENTRY(ccl_global float, L_transparent, 1) \
SPLIT_DATA_ENTRY(PathRadiance, path_radiance, 1) \
SPLIT_DATA_ENTRY(ccl_global Ray, ray, 1) \
SPLIT_DATA_ENTRY(ccl_global PathState, path_state, 1) \
SPLIT_DATA_ENTRY(Intersection, isect, 1) \
SPLIT_DATA_ENTRY(ccl_global float3, ao_alpha, 1) \
SPLIT_DATA_ENTRY(ccl_global float3, ao_bsdf, 1) \
SPLIT_DATA_ENTRY(ccl_global Ray, ao_light_ray, 1) \
SPLIT_DATA_ENTRY(ccl_global BsdfEval, bsdf_eval, 1) \
SPLIT_DATA_ENTRY(ccl_global int, is_lamp, 1) \
SPLIT_DATA_ENTRY(ccl_global Ray, light_ray, 1) \
SPLIT_DATA_ENTRY(Intersection, isect_shadow, 2) \
SPLIT_DATA_ENTRY(ccl_global int, queue_data, (NUM_QUEUES*2)) /* TODO(mai): this is too large? */ \
SPLIT_DATA_ENTRY(ccl_global uint, work_array, 1) \
SPLIT_DATA_ENTRY(ShaderData, sd, 1) \
SPLIT_DATA_ENTRY(ShaderData, sd_DL_shadow, 2) \
SPLIT_DATA_DEBUG_ENTRIES \
/* struct that holds pointers to data in the shared state buffer */
typedef struct SplitData {
#define SPLIT_DATA_ENTRY(type, name, num) type *name;
SPLIT_DATA_ENTRIES
#undef SPLIT_DATA_ENTRY
/* this is actually in a separate buffer from the rest of the split state data (so it can be read back from
* the host easily) but is still used the same as the other data so we have it here in this struct as well
*/
ccl_global char *ray_state;
} SplitData;
#ifndef __KERNEL_CUDA__
# define kernel_split_state (kg->split_data)
# define kernel_split_params (kg->split_param_data)
#else
__device__ SplitData __split_data;
# define kernel_split_state (__split_data)
__device__ SplitParams __split_param_data;
# define kernel_split_params (__split_param_data)
#endif /* __KERNEL_CUDA__ */
CCL_NAMESPACE_END
#endif /* __KERNEL_SPLIT_DATA_TYPES_H__ */