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

svn merge ^/trunk/blender -r42053:42069

Browse Source
This commit is contained in:
Campbell Barton 2011-11-22 15:33:44 +00:00
commit c6391e8d23
18 changed files with 141 additions and 107 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

35
intern/cycles/device/device_opencl.cpp
View File

@ -260,12 +260,9 @@ public:
return true;
}
bool build_kernel(const string& kernel_path)
string kernel_build_options()
{
string build_options = "";
build_options += "-I " + kernel_path + ""; /* todo: escape path */
build_options += " -cl-fast-relaxed-math ";
string build_options = " -cl-fast-relaxed-math ";
/* Full Shading only on NVIDIA cards at the moment */
char vendor[256];
@ -273,14 +270,19 @@ public:
clGetPlatformInfo(cpPlatform, CL_PLATFORM_NAME, sizeof(vendor), &vendor, NULL);
string name = vendor;
if (name == "NVIDIA CUDA") {
build_options += "-D __SVM__ ";
build_options += "-D __EMISSION__ ";
build_options += "-D __TEXTURES__ ";
build_options += "-D __HOLDOUT__ ";
build_options += "-D __MULTI_CLOSURE__ ";
}
if(name == "NVIDIA CUDA")
build_options += "-D__KERNEL_SHADING__ -D__MULTI_CLOSURE__ ";
return build_options;
}
bool build_kernel(const string& kernel_path)
{
string build_options = "";
build_options += "-I " + kernel_path + ""; /* todo: escape path, but it doesn't get parsed correct? */
build_options += kernel_build_options();
ciErr = clBuildProgram(cpProgram, 0, NULL, build_options.c_str(), NULL, NULL);
if(ciErr != CL_SUCCESS) {
@ -344,6 +346,9 @@ public:
md5.append((uint8_t*)name, strlen(name));
md5.append((uint8_t*)driver, strlen(driver));
string options = kernel_build_options();
md5.append((uint8_t*)options.c_str(), options.size());
return md5.get_hex();
}
@ -563,24 +568,20 @@ public:
cl_int set_kernel_arg_mem(cl_kernel kernel, int *narg, const char *name)
{
cl_mem ptr;
cl_int size, err = 0;
cl_int err = 0;
if(mem_map.find(name) != mem_map.end()) {
device_memory *mem = mem_map[name];
ptr = CL_MEM_PTR(mem->device_pointer);
size = mem->data_width;
}
else {
/* work around NULL not working, even though the spec says otherwise */
ptr = CL_MEM_PTR(null_mem);
size = 1;
}
err |= clSetKernelArg(kernel, (*narg)++, sizeof(ptr), (void*)&ptr);
opencl_assert(err);
err |= clSetKernelArg(kernel, (*narg)++, sizeof(size), (void*)&size);
opencl_assert(err);
return err;
}

12
intern/cycles/kernel/kernel.cl
View File

@ -33,8 +33,7 @@ __kernel void kernel_ocl_path_trace(
__global uint *rng_state,
#define KERNEL_TEX(type, ttype, name) \
__global type *name, \
int name##_width,
__global type *name,
#include "kernel_textures.h"
int sample,
@ -45,8 +44,7 @@ __kernel void kernel_ocl_path_trace(
kg->data = data;
#define KERNEL_TEX(type, ttype, name) \
kg->name = name; \
kg->name##_width = name##_width;
kg->name = name;
#include "kernel_textures.h"
int x = sx + get_global_id(0);
@ -62,8 +60,7 @@ __kernel void kernel_ocl_tonemap(
__global float4 *buffer,
#define KERNEL_TEX(type, ttype, name) \
__global type *name, \
int name##_width,
__global type *name,
#include "kernel_textures.h"
int sample, int resolution,
@ -74,8 +71,7 @@ __kernel void kernel_ocl_tonemap(
kg->data = data;
#define KERNEL_TEX(type, ttype, name) \
kg->name = name; \
kg->name##_width = name##_width;
kg->name = name;
#include "kernel_textures.h"
int x = sx + get_global_id(0);

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

@ -127,8 +127,8 @@ __device void camera_sample_orthographic(KernelGlobals *kg, float raster_x, floa
__device void camera_sample(KernelGlobals *kg, int x, int y, float filter_u, float filter_v, float lens_u, float lens_v, Ray *ray)
{
/* pixel filter */
float raster_x = x + kernel_tex_interp(__filter_table, filter_u);
float raster_y = y + kernel_tex_interp(__filter_table, filter_v);
float raster_x = x + kernel_tex_interp(__filter_table, filter_u, FILTER_TABLE_SIZE);
float raster_y = y + kernel_tex_interp(__filter_table, filter_v, FILTER_TABLE_SIZE);
/* motion blur */
//ray->time = lerp(time_t, kernel_data.cam.shutter_open, kernel_data.cam.shutter_close);

6
intern/cycles/kernel/kernel_compat_cpu.h
View File

@ -55,8 +55,10 @@ template<typename T> struct texture {
return ((__m128i*)data)[index];
}*/
float interp(float x)
float interp(float x, int size)
{
kernel_assert(size == width);
x = clamp(x, 0.0f, 1.0f)*width;
int index = min((int)x, width-1);
@ -151,7 +153,7 @@ typedef texture_image<uchar4> texture_image_uchar4;
#define kernel_tex_fetch(tex, index) (kg->tex.fetch(index))
#define kernel_tex_fetch_m128(tex, index) (kg->tex.fetch_m128(index))
#define kernel_tex_fetch_m128i(tex, index) (kg->tex.fetch_m128i(index))
#define kernel_tex_interp(tex, t) (kg->tex.interp(t))
#define kernel_tex_interp(tex, t, size) (kg->tex.interp(t, size))
#define kernel_tex_image_interp(tex, x, y) (kg->tex.interp(x, y))
#define kernel_data (kg->__data)

2
intern/cycles/kernel/kernel_compat_cuda.h
View File

@ -55,7 +55,7 @@ typedef texture<uchar4, 2, cudaReadModeNormalizedFloat> texture_image_uchar4;
/* Macros to handle different memory storage on different devices */
#define kernel_tex_fetch(t, index) tex1Dfetch(t, index)
#define kernel_tex_interp(t, x) tex1D(t, x)
#define kernel_tex_interp(t, x, size) tex1D(t, x)
#define kernel_tex_image_interp(t, x, y) tex2D(t, x, y)
#define kernel_data __data

2
intern/cycles/kernel/kernel_compat_opencl.h
View File

@ -100,7 +100,7 @@ __device float kernel_tex_interp_(__global float *data, int width, float x)
/* data lookup defines */
#define kernel_data (*kg->data)
#define kernel_tex_interp(t, x) kernel_tex_interp_(kg->t, kg->t##_width, x)
#define kernel_tex_interp(t, x, size) kernel_tex_interp_(kg->t, size, x)
#define kernel_tex_fetch(t, index) kg->t[index]
/* define NULL */

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

@ -77,8 +77,7 @@ typedef struct KernelGlobals {
__constant KernelData *data;
#define KERNEL_TEX(type, ttype, name) \
__global type *name; \
int name##_width;
__global type *name;
#include "kernel_textures.h"
} KernelGlobals;

2
intern/cycles/kernel/kernel_shader.h
View File

@ -226,7 +226,7 @@ __device void shader_setup_from_displace(KernelGlobals *kg, ShaderData *sd,
Ng = triangle_normal_MT(kg, prim, &shader);
/* force smooth shading for displacement */
sd->shader |= SHADER_SMOOTH_NORMAL;
shader |= SHADER_SMOOTH_NORMAL;
/* watch out: no instance transform currently */

48
intern/cycles/kernel/kernel_types.h
View File

@ -25,9 +25,30 @@
CCL_NAMESPACE_BEGIN
#define OBJECT_SIZE 16
#define LIGHT_SIZE 4
/* constants */
#define OBJECT_SIZE 16
#define LIGHT_SIZE 4
#define FILTER_TABLE_SIZE 256
/* device capabilities */
#ifdef __KERNEL_CPU__
#define __KERNEL_SHADING__
#define __KERNEL_ADV_SHADING__
#endif
#ifdef __KERNEL_CUDA__
#define __KERNEL_SHADING__
#if __CUDA_ARCH__ >= 200
#define __KERNEL_ADV_SHADING__
#endif
#endif
#ifdef __KERNEL_OPENCL__
//#define __KERNEL_SHADING__
//#define __KERNEL_ADV_SHADING__
#endif
/* kernel features */
#define __SOBOL__
#define __INSTANCING__
#define __DPDU__
@ -39,27 +60,20 @@ CCL_NAMESPACE_BEGIN
#define __CAMERA_CLIPPING__
#define __INTERSECTION_REFINE__
#ifndef __KERNEL_OPENCL__
#ifdef __KERNEL_SHADING__
#define __SVM__
#define __EMISSION__
#define __TEXTURES__
#define __HOLDOUT__
#endif
#ifdef __KERNEL_ADV_SHADING__
#define __MULTI_CLOSURE__
#define __TRANSPARENT_SHADOWS__
#endif
//#define __MULTI_LIGHT__
#endif
#ifdef __KERNEL_CPU__
#define __MULTI_CLOSURE__
#define __TRANSPARENT_SHADOWS__
//#define __OSL__
#endif
#ifdef __KERNEL_CUDA__
#if __CUDA_ARCH__ >= 200
#define __MULTI_CLOSURE__
#define __TRANSPARENT_SHADOWS__
#endif
#endif
//#define __SOBOL_FULL_SCREEN__
//#define __MODIFY_TP__
//#define __QBVH__

4
intern/cycles/render/filter.cpp
View File

@ -21,6 +21,8 @@
#include "filter.h"
#include "scene.h"
#include "kernel_types.h"
#include "util_algorithm.h"
#include "util_debug.h"
#include "util_math.h"
@ -51,7 +53,7 @@ static float filter_func_gaussian(float v, float width)
static vector<float> filter_table(FilterType type, float width)
{
const int filter_table_size = 256;
const int filter_table_size = FILTER_TABLE_SIZE;
vector<float> filter_table_cdf(filter_table_size+1);
vector<float> filter_table(filter_table_size+1);
float (*filter_func)(float, float) = NULL;

2
release/scripts/startup/bl_ui/properties_game.py
View File

@ -408,7 +408,7 @@ class SceneButtonsPanel():
class SCENE_PT_game_navmesh(SceneButtonsPanel, Panel):
bl_label = "Navigation mesh"
bl_default_closed = True
bl_options = {'DEFAULT_CLOSED'}
COMPAT_ENGINES = {'BLENDER_GAME'}
@classmethod

1
release/scripts/startup/bl_ui/space_view3d.py
View File

@ -262,6 +262,7 @@ class VIEW3D_MT_uv_map(Menu):
layout.separator()
layout.operator_context = 'EXEC_REGION_WIN'
layout.operator("uv.project_from_view")
layout.operator("uv.project_from_view", text="Project from View (Bounds)").scale_to_bounds = True

52
source/blender/editors/object/object_constraint.c
View File

@ -691,6 +691,7 @@ static int childof_set_inverse_exec (bContext *C, wmOperator *op)
Object *ob = ED_object_active_context(C);
bConstraint *con = edit_constraint_property_get(op, ob, CONSTRAINT_TYPE_CHILDOF);
bChildOfConstraint *data= (con) ? (bChildOfConstraint *)con->data : NULL;
bConstraint *lastcon = NULL;
bPoseChannel *pchan= NULL;
/* despite 3 layers of checks, we may still not be able to find a constraint */
@ -703,27 +704,45 @@ static int childof_set_inverse_exec (bContext *C, wmOperator *op)
/* nullify inverse matrix first */
unit_m4(data->invmat);
/* try to find a pose channel */
/* try to find a pose channel - assume that this is the constraint owner */
// TODO: get from context instead?
if (ob && ob->pose)
pchan= get_active_posechannel(ob);
/* calculate/set inverse matrix */
/* calculate/set inverse matrix:
* We just calculate all transform-stack eval up to but not including this constraint.
* This is because inverse should just inverse correct for just the constraint's influence
* when it gets applied; that is, at the time of application, we don't know anything about
* what follows.
*/
if (pchan) {
float pmat[4][4], cinf;
float imat[4][4], tmat[4][4];
float pmat[4][4];
/* make copy of pchan's original pose-mat (for use later) */
/* 1. calculate posemat where inverse doesn't exist yet (inverse was cleared above),
* to use as baseline ("pmat") to derive delta from. This extra calc saves users
* from having pressing "Clear Inverse" first
*/
where_is_pose(scene, ob);
copy_m4_m4(pmat, pchan->pose_mat);
/* disable constraint for pose to be solved without it */
cinf= con->enforce;
con->enforce= 0.0f;
/* 2. knock out constraints starting from this one */
lastcon = pchan->constraints.last;
pchan->constraints.last = con->prev;
/* solve pose without constraint */
if (con->prev) {
/* new end must not point to this one, else this chain cutting is useless */
con->prev->next = NULL;
}
else {
/* constraint was first */
pchan->constraints.first = NULL;
}
/* 3. solve pose without disabled constraints */
where_is_pose(scene, ob);
/* determine effect of constraint by removing the newly calculated
/* 4. determine effect of constraint by removing the newly calculated
* pchan->pose_mat from the original pchan->pose_mat, thus determining
* the effect of the constraint
*/
@ -731,8 +750,19 @@ static int childof_set_inverse_exec (bContext *C, wmOperator *op)
mul_m4_m4m4(tmat, imat, pmat);
invert_m4_m4(data->invmat, tmat);
/* recalculate pose with new inv-mat */
con->enforce= cinf;
/* 5. restore constraints */
pchan->constraints.last = lastcon;
if (con->prev) {
/* hook up prev to this one again */
con->prev->next = con;
}
else {
/* set as first again */
pchan->constraints.first = con;
}
/* 6. recalculate pose with new inv-mat applied */
where_is_pose(scene, ob);
}
else if (ob) {

23
source/blender/editors/sculpt_paint/paint_image.c
View File

@ -300,6 +300,7 @@ typedef struct ProjPaintState {
short do_occlude; /* Use raytraced occlusion? - ortherwise will paint right through to the back*/
short do_backfacecull; /* ignore faces with normals pointing away, skips a lot of raycasts if your normals are correctly flipped */
short do_mask_normal; /* mask out pixels based on their normals */
short do_new_shading_nodes; /* cache scene_use_new_shading_nodes value */
float normal_angle; /* what angle to mask at*/
float normal_angle_inner;
float normal_angle_range; /* difference between normal_angle and normal_angle_inner, for easy access */
@ -518,17 +519,16 @@ static Image *imapaint_face_image(const ImagePaintState *s, int face_index)
return ima;
}
static Image *project_paint_face_image(const ProjPaintState *ps, int face_index)
static Image *project_paint_face_image(const ProjPaintState *ps, MTFace *dm_mtface, int face_index)
{
Image *ima;
if(scene_use_new_shading_nodes(ps->scene)) {
if(ps->do_new_shading_nodes) { /* cached scene_use_new_shading_nodes result */
MFace *mf = ps->dm_mface+face_index;
ED_object_get_active_image(ps->ob, mf->mat_nr, &ima, NULL, NULL);
}
else {
MTFace *tf = ps->dm_mtface+face_index;
ima = tf->tpage;
ima = dm_mtface[face_index].tpage;
}
return ima;
@ -725,7 +725,7 @@ static int project_paint_PickColor(const ProjPaintState *ps, float pt[2], float
interp_v2_v2v2v2(uv, tf->uv[0], tf->uv[2], tf->uv[3], w);
}
ima = project_paint_face_image(ps, face_index);
ima = project_paint_face_image(ps, ps->dm_mtface, face_index);
ibuf = ima->ibufs.first; /* we must have got the imbuf before getting here */
if (!ibuf) return 0;
@ -1091,8 +1091,8 @@ static int check_seam(const ProjPaintState *ps, const int orig_face, const int o
/* Only need to check if 'i2_fidx' is valid because we know i1_fidx is the same vert on both faces */
if (i2_fidx != -1) {
Image *tpage = project_paint_face_image(ps, face_index);
Image *orig_tpage = project_paint_face_image(ps, orig_face);
Image *tpage = project_paint_face_image(ps, ps->dm_mtface, face_index);
Image *orig_tpage = project_paint_face_image(ps, ps->dm_mtface, orig_face);
/* This IS an adjacent face!, now lets check if the UVs are ok */
tf = ps->dm_mtface + face_index;
@ -1349,7 +1349,7 @@ static float project_paint_uvpixel_mask(
if (ps->do_layer_stencil) {
/* another UV layers image is masking this one's */
ImBuf *ibuf_other;
Image *other_tpage = project_paint_face_image(ps, face_index);
Image *other_tpage = project_paint_face_image(ps, ps->dm_mtface_stencil, face_index);
const MTFace *tf_other = ps->dm_mtface_stencil + face_index;
if (other_tpage && (ibuf_other = BKE_image_get_ibuf(other_tpage, NULL))) {
@ -1506,7 +1506,7 @@ static ProjPixel *project_paint_uvpixel_init(
if (ps->tool==PAINT_TOOL_CLONE) {
if (ps->dm_mtface_clone) {
ImBuf *ibuf_other;
Image *other_tpage = project_paint_face_image(ps, face_index);
Image *other_tpage = project_paint_face_image(ps, ps->dm_mtface_clone, face_index);
const MTFace *tf_other = ps->dm_mtface_clone + face_index;
if (other_tpage && (ibuf_other = BKE_image_get_ibuf(other_tpage, NULL))) {
@ -2746,7 +2746,7 @@ static void project_bucket_init(const ProjPaintState *ps, const int thread_index
face_index = GET_INT_FROM_POINTER(node->link);
/* Image context switching */
tpage = project_paint_face_image(ps, face_index);
tpage = project_paint_face_image(ps, ps->dm_mtface, face_index);
if (tpage_last != tpage) {
tpage_last = tpage;
@ -3252,7 +3252,7 @@ static void project_paint_begin(ProjPaintState *ps)
}
#endif
tpage = project_paint_face_image(ps, face_index);
tpage = project_paint_face_image(ps, ps->dm_mtface, face_index);
if (tpage && ((((Mesh *)ps->ob->data)->editflag & ME_EDIT_PAINT_MASK)==0 || mf->flag & ME_FACE_SEL)) {
@ -4718,6 +4718,7 @@ static void project_state_init(bContext *C, Object *ob, ProjPaintState *ps)
ps->do_backfacecull = (settings->imapaint.flag & IMAGEPAINT_PROJECT_BACKFACE) ? 0 : 1;
ps->do_occlude = (settings->imapaint.flag & IMAGEPAINT_PROJECT_XRAY) ? 0 : 1;
ps->do_mask_normal = (settings->imapaint.flag & IMAGEPAINT_PROJECT_FLAT) ? 0 : 1;
ps->do_new_shading_nodes = scene_use_new_shading_nodes(scene); /* only cache the value */
if (ps->tool == PAINT_TOOL_CLONE)
ps->do_layer_clone = (settings->imapaint.flag & IMAGEPAINT_PROJECT_LAYER_CLONE);

2
source/blender/editors/space_buttons/buttons_texture.c
View File

@ -257,8 +257,8 @@ void buttons_texture_context_compute(const bContext *C, SpaceButs *sbuts)
if(!scene_use_new_shading_nodes(scene)) {
if(ct) {
MEM_freeN(ct);
BLI_freelistN(&ct->users);
MEM_freeN(ct);
sbuts->texuser= NULL;
}

6
source/blender/editors/space_clip/tracking_ops.c
View File

@ -2251,7 +2251,7 @@ static int detect_features_exec(bContext *C, wmOperator *op)
MovieTrackingTrack *track= clip->tracking.tracks.first;
int placement= RNA_enum_get(op->ptr, "placement");
int margin= RNA_int_get(op->ptr, "margin");
int min_trackness= RNA_int_get(op->ptr, "min_trackness");
int min_trackability= RNA_int_get(op->ptr, "min_trackability");
int min_distance= RNA_int_get(op->ptr, "min_distance");
int place_outside_layer= 0;
bGPDlayer *layer= NULL;
@ -2270,7 +2270,7 @@ static int detect_features_exec(bContext *C, wmOperator *op)
track= track->next;
}
BKE_tracking_detect_fast(&clip->tracking, ibuf, sc->user.framenr, margin, min_trackness, min_distance, layer, place_outside_layer);
BKE_tracking_detect_fast(&clip->tracking, ibuf, sc->user.framenr, margin, min_trackability, min_distance, layer, place_outside_layer);
IMB_freeImBuf(ibuf);
@ -2303,7 +2303,7 @@ void CLIP_OT_detect_features(wmOperatorType *ot)
/* properties */
RNA_def_enum(ot->srna, "placement", placement_items, 0, "Placement", "Placement for detected features");
RNA_def_int(ot->srna, "margin", 16, 0, INT_MAX, "Margin", "Only corners further than margin pixels from the image edges are considered", 0, 300);
RNA_def_int(ot->srna, "min_trackness", 16, 0, INT_MAX, "Trackness", "Minimum score to add a corner", 0, 300);
RNA_def_int(ot->srna, "min_trackability", 16, 0, INT_MAX, "Trackability", "Minimum trackability score to add a corner", 0, 300);
RNA_def_int(ot->srna, "min_distance", 120, 0, INT_MAX, "Distance", "Minimal distance accepted between two corners", 0, 300);
}

38
source/blender/modifiers/intern/MOD_ocean.c
View File

@ -404,7 +404,8 @@ static DerivedMesh *doOcean(ModifierData *md, Object *ob,
if (omd->cached == TRUE) {
if (!omd->oceancache) init_cache_data(ob, omd);
BKE_simulate_ocean_cache(omd->oceancache, md->scene->r.cfra);
} else {
}
else {
simulate_ocean_modifier(omd);
}
@ -441,43 +442,26 @@ static DerivedMesh *doOcean(ModifierData *md, Object *ob,
mf = dm->getFaceArray(dm);
for (i = 0; i < num_faces; i++, mf++) {
for (j=0; j<4; j++) {
if (j == 3 && !mf->v4) continue;
switch(j) {
case 0:
u = ocean_co(omd, mv[mf->v1].co[0]);
v = ocean_co(omd, mv[mf->v1].co[1]);
break;
case 1:
u = ocean_co(omd, mv[mf->v2].co[0]);
v = ocean_co(omd, mv[mf->v2].co[1]);
break;
case 2:
u = ocean_co(omd, mv[mf->v3].co[0]);
v = ocean_co(omd, mv[mf->v3].co[1]);
break;
case 3:
u = ocean_co(omd, mv[mf->v4].co[0]);
v = ocean_co(omd, mv[mf->v4].co[1]);
break;
}
j= mf->v4 ? 3 : 2;
do {
const float *co= mv[*(&mf->v1 + j)].co;
u = ocean_co(omd, co[0]);
v = ocean_co(omd, co[1]);
if (omd->oceancache && omd->cached==TRUE) {
BKE_ocean_cache_eval_uv(omd->oceancache, &ocr, cfra, u, v);
foam = ocr.foam;
CLAMP(foam, 0.0f, 1.0f);
} else {
}
else {
BKE_ocean_eval_uv(omd->ocean, &ocr, u, v);
foam = BKE_ocean_jminus_to_foam(ocr.Jminus, omd->foam_coverage);
}
cf = (char)(foam*255);
cf = (char)(foam * 255);
mc[i*4 + j].r = mc[i*4 + j].g = mc[i*4 + j].b = cf;
mc[i*4 + j].a = 255;
}
} while (j--);
}
}

6
source/blender/python/generic/py_capi_utils.c
View File

@ -385,8 +385,12 @@ const char *PyC_UnicodeAsByte(PyObject *py_str, PyObject **coerce)
if (PyBytes_Check(py_str)) {
return PyBytes_AS_STRING(py_str);
}
else if ((*coerce= PyUnicode_EncodeFSDefault(py_str))) {
return PyBytes_AS_STRING(*coerce);
}
else {
return PyBytes_AS_STRING((*coerce= PyUnicode_EncodeFSDefault(py_str)));
/* leave error raised from EncodeFS */
return NULL;
}
}
}