svn merge ^/trunk/blender -r48674:48681

This commit is contained in:
Campbell Barton 2012-07-06 14:25:49 +00:00
commit ec233cd747
10 changed files with 218 additions and 51 deletions

@ -798,7 +798,7 @@ if env['OURPLATFORM'] == 'win64-mingw':
dllsources.append('${LCGDIR}/thumbhandler/lib/BlendThumb64.dll')
dllsources.append('${LCGDIR}/binaries/libgcc_s_sjlj-1.dll')
dllsources.append('${LCGDIR}/binaries/libwinpthread-1.dll')
dllsources.append('${LCGDIR}/binaries/libstdc++-6.dll)')
dllsources.append('${LCGDIR}/binaries/libstdc++-6.dll')
dllsources.append('#source/icons/blender.exe.manifest')
windlls = env.Install(dir=env['BF_INSTALLDIR'], source = dllsources)

@ -57,7 +57,7 @@ void COM_execute(RenderData *rd, bNodeTree *editingtree, int rendering)
/* set progress bar to 0% and status to init compositing*/
editingtree->progress(editingtree->prh, 0.0);
bool twopass = (editingtree->flag&NTREE_TWO_PASS) > 0 || rendering;
bool twopass = (editingtree->flag&NTREE_TWO_PASS) > 0 && !rendering;
/* initialize execution system */
if (twopass) {
ExecutionSystem *system = new ExecutionSystem(rd, editingtree, rendering, twopass);

@ -51,6 +51,68 @@ __kernel void bokehBlurKernel(__read_only image2d_t boundingBox, __read_only ima
write_imagef(output, coords, color);
}
//KERNEL --- DEFOCUS /VARIABLESIZEBOKEHBLUR ---
__kernel void defocusKernel(__read_only image2d_t inputImage, __read_only image2d_t bokehImage,
__read_only image2d_t inputDepth, __read_only image2d_t inputSize,
__write_only image2d_t output, int2 offsetInput, int2 offsetOutput,
int step, int maxBlur, float threshold, int2 dimension, int2 offset)
{
float4 color = {1.0f, 0.0f, 0.0f, 1.0f};
int2 coords = {get_global_id(0), get_global_id(1)};
coords += offset;
const int2 realCoordinate = coords + offsetOutput;
float4 readColor;
float4 bokeh;
float tempSize;
float tempDepth;
float4 multiplier_accum = {1.0f, 1.0f, 1.0f, 1.0f};
float4 color_accum;
int minx = max(realCoordinate.s0 - maxBlur, 0);
int miny = max(realCoordinate.s1 - maxBlur, 0);
int maxx = min(realCoordinate.s0 + maxBlur, dimension.s0);
int maxy = min(realCoordinate.s1 + maxBlur, dimension.s1);
{
int2 inputCoordinate = realCoordinate - offsetInput;
float size = read_imagef(inputSize, SAMPLER_NEAREST, inputCoordinate).s0;
float depth = read_imagef(inputDepth, SAMPLER_NEAREST, inputCoordinate).s0 + threshold;
color_accum = read_imagef(inputImage, SAMPLER_NEAREST, inputCoordinate);
for (int ny = miny; ny < maxy; ny += step) {
for (int nx = minx; nx < maxx; nx += step) {
if (nx >= 0 && nx < dimension.s0 && ny >= 0 && ny < dimension.s1) {
inputCoordinate.s0 = nx - offsetInput.s0;
inputCoordinate.s1 = ny - offsetInput.s1;
tempDepth = read_imagef(inputDepth, SAMPLER_NEAREST, inputCoordinate).s0;
if (tempDepth < depth) {
tempSize = read_imagef(inputSize, SAMPLER_NEAREST, inputCoordinate).s0;
if ((size > threshold && tempSize > threshold) || tempSize <= threshold) {
float dx = nx - realCoordinate.s0;
float dy = ny - realCoordinate.s1;
if (dx != 0 || dy != 0) {
if (tempSize >= fabs(dx) && tempSize >= fabs(dy)) {
float2 uv = { 256.0f + dx * 256.0f / tempSize, 256.0f + dy * 256.0f / tempSize};
bokeh = read_imagef(bokehImage, SAMPLER_NEAREST, uv);
readColor = read_imagef(inputImage, SAMPLER_NEAREST, inputCoordinate);
color_accum += bokeh*readColor;
multiplier_accum += bokeh;
}
}
}
}
}
}
}
}
color = color_accum * (1.0f / multiplier_accum);
write_imagef(output, coords, color);
}
// KERNEL --- DILATE ---
__kernel void dilateKernel(__read_only image2d_t inputImage, __write_only image2d_t output,
int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension,

@ -16,7 +16,7 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope
" coords += offset;\n" \
" float tempBoundingBox;\n" \
" float4 color = {0.0f,0.0f,0.0f,0.0f};\n" \
" float4 multiplier = {0.0f,0.0f,0.0f,0.0f};\n" \
" float4 multiplyer = {0.0f,0.0f,0.0f,0.0f};\n" \
" float4 bokeh;\n" \
" const float radius2 = radius*2.0f;\n" \
" const int2 realCoordinate = coords + offsetOutput;\n" \
@ -40,10 +40,10 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope
" uv.x = ((realCoordinate.x-nx)/radius2)*bokehImageDim.x+bokehImageCenter.x;\n" \
" bokeh = read_imagef(bokehImage, SAMPLER_NEAREST, uv);\n" \
" color += bokeh * read_imagef(inputImage, SAMPLER_NEAREST, inputXy);\n" \
" multiplier += bokeh;\n" \
" multiplyer += bokeh;\n" \
" }\n" \
" }\n" \
" color /= multiplier;\n" \
" color /= multiplyer;\n" \
"\n" \
" } else {\n" \
" int2 imageCoordinates = realCoordinate - offsetInput;\n" \
@ -53,6 +53,68 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope
" write_imagef(output, coords, color);\n" \
"}\n" \
"\n" \
"//KERNEL --- DEFOCUS /VARIABLESIZEBOKEHBLUR ---\n" \
"__kernel void defocusKernel(__read_only image2d_t inputImage, __read_only image2d_t bokehImage,\n" \
" __read_only image2d_t inputDepth, __read_only image2d_t inputSize,\n" \
" __write_only image2d_t output, int2 offsetInput, int2 offsetOutput,\n" \
" int step, int maxBlur, float threshold, int2 dimension, int2 offset)\n" \
"{\n" \
" float4 color = {1.0f, 0.0f, 0.0f, 1.0f};\n" \
" int2 coords = {get_global_id(0), get_global_id(1)};\n" \
" coords += offset;\n" \
" const int2 realCoordinate = coords + offsetOutput;\n" \
"\n" \
" float4 readColor;\n" \
" float4 bokeh;\n" \
" float tempSize;\n" \
" float tempDepth;\n" \
" float4 multiplier_accum = {1.0f, 1.0f, 1.0f, 1.0f};\n" \
" float4 color_accum;\n" \
"\n" \
" int minx = max(realCoordinate.s0 - maxBlur, 0);\n" \
" int miny = max(realCoordinate.s1 - maxBlur, 0);\n" \
" int maxx = min(realCoordinate.s0 + maxBlur, dimension.s0);\n" \
" int maxy = min(realCoordinate.s1 + maxBlur, dimension.s1);\n" \
"\n" \
" {\n" \
" int2 inputCoordinate = realCoordinate - offsetInput;\n" \
" float size = read_imagef(inputSize, SAMPLER_NEAREST, inputCoordinate).s0;\n" \
" float depth = read_imagef(inputDepth, SAMPLER_NEAREST, inputCoordinate).s0 + threshold;\n" \
" color_accum = read_imagef(inputImage, SAMPLER_NEAREST, inputCoordinate);\n" \
"\n" \
" for (int ny = miny; ny < maxy; ny += step) {\n" \
" for (int nx = minx; nx < maxx; nx += step) {\n" \
" if (nx >= 0 && nx < dimension.s0 && ny >= 0 && ny < dimension.s1) {\n" \
" inputCoordinate.s0 = nx - offsetInput.s0;\n" \
" inputCoordinate.s1 = ny - offsetInput.s1;\n" \
" tempDepth = read_imagef(inputDepth, SAMPLER_NEAREST, inputCoordinate).s0;\n" \
" if (tempDepth < depth) {\n" \
" tempSize = read_imagef(inputSize, SAMPLER_NEAREST, inputCoordinate).s0;\n" \
"\n" \
" if ((size > threshold && tempSize > threshold) || tempSize <= threshold) {\n" \
" float dx = nx - realCoordinate.s0;\n" \
" float dy = ny - realCoordinate.s1;\n" \
" if (dx != 0 || dy != 0) {\n" \
" if (tempSize >= fabs(dx) && tempSize >= fabs(dy)) {\n" \
" float2 uv = { 256.0f + dx * 256.0f / tempSize, 256.0f + dy * 256.0f / tempSize};\n" \
" bokeh = read_imagef(bokehImage, SAMPLER_NEAREST, uv);\n" \
" readColor = read_imagef(inputImage, SAMPLER_NEAREST, inputCoordinate);\n" \
" color_accum += bokeh*readColor;\n" \
" multiplier_accum += bokeh;\n" \
" }\n" \
" }\n" \
" }\n" \
" }\n" \
" }\n" \
" }\n" \
" }\n" \
" }\n" \
"\n" \
" color = color_accum * (1.0f / multiplier_accum);\n" \
" write_imagef(output, coords, color);\n" \
"}\n" \
"\n" \
"\n" \
"// KERNEL --- DILATE ---\n" \
"__kernel void dilateKernel(__read_only image2d_t inputImage, __write_only image2d_t output,\n" \
" int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension,\n" \
@ -70,9 +132,9 @@ const char * clkernelstoh_COM_OpenCLKernels_cl = "/// This file contains all ope
" int2 inputXy;\n" \
"\n" \
" for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) {\n" \
" const float deltaY = (realCoordinate.y - ny);\n" \
" for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx ++, inputXy.x++) {\n" \
" const float deltaX = (realCoordinate.x - nx);\n" \
" const float deltaY = (realCoordinate.y - ny);\n" \
" const float measuredDistance = deltaX*deltaX+deltaY*deltaY;\n" \
" if (measuredDistance <= distanceSquared) {\n" \
" value = max(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0);\n" \

@ -22,6 +22,7 @@
#include "COM_VariableSizeBokehBlurOperation.h"
#include "BLI_math.h"
#include "COM_OpenCLDevice.h"
extern "C" {
#include "RE_pipeline.h"
@ -38,6 +39,7 @@ VariableSizeBokehBlurOperation::VariableSizeBokehBlurOperation() : NodeOperation
#endif
this->addOutputSocket(COM_DT_COLOR);
this->setComplex(true);
this->setOpenCL(true);
this->m_inputProgram = NULL;
this->m_inputBokehProgram = NULL;
@ -128,6 +130,33 @@ void VariableSizeBokehBlurOperation::executePixel(float *color, int x, int y, Me
}
static cl_kernel defocusKernel = 0;
void VariableSizeBokehBlurOperation::executeOpenCL(OpenCLDevice* device,
MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer,
MemoryBuffer **inputMemoryBuffers, list<cl_mem> *clMemToCleanUp,
list<cl_kernel> *clKernelsToCleanUp)
{
if (!defocusKernel) {
defocusKernel = device->COM_clCreateKernel("defocusKernel", NULL);
}
cl_int step = this->getStep();
cl_int maxBlur = this->m_maxBlur;
cl_float threshold = this->m_threshold;
device->COM_clAttachMemoryBufferToKernelParameter(defocusKernel, 0, -1, clMemToCleanUp, inputMemoryBuffers, this->m_inputProgram);
device->COM_clAttachMemoryBufferToKernelParameter(defocusKernel, 1, -1, clMemToCleanUp, inputMemoryBuffers, this->m_inputBokehProgram);
device->COM_clAttachMemoryBufferToKernelParameter(defocusKernel, 2, 5, clMemToCleanUp, inputMemoryBuffers, this->m_inputDepthProgram);
device->COM_clAttachMemoryBufferToKernelParameter(defocusKernel, 3, -1, clMemToCleanUp, inputMemoryBuffers, this->m_inputSizeProgram);
device->COM_clAttachOutputMemoryBufferToKernelParameter(defocusKernel, 4, clOutputBuffer);
device->COM_clAttachMemoryBufferOffsetToKernelParameter(defocusKernel, 6, outputMemoryBuffer);
clSetKernelArg(defocusKernel, 7, sizeof(cl_int), &step);
clSetKernelArg(defocusKernel, 8, sizeof(cl_int), &maxBlur);
clSetKernelArg(defocusKernel, 9, sizeof(cl_float), &threshold);
device->COM_clAttachSizeToKernelParameter(defocusKernel, 10, this);
device->COM_clEnqueueRange(defocusKernel, outputMemoryBuffer, 11, this);
}
void VariableSizeBokehBlurOperation::deinitExecution()
{
this->m_inputProgram = NULL;

@ -62,7 +62,7 @@ public:
void setThreshold(float threshold) { this->m_threshold = threshold; }
void executeOpenCL(OpenCLDevice* device, MemoryBuffer *outputMemoryBuffer, cl_mem clOutputBuffer, MemoryBuffer **inputMemoryBuffers, list<cl_mem> *clMemToCleanUp, list<cl_kernel> *clKernelsToCleanUp);
};
#ifdef COM_DEFOCUS_SEARCH

@ -2996,13 +2996,15 @@ void draw_nodespace_back_pix(const bContext *C, ARegion *ar, SpaceNode *snode)
if (snode->flag & (SNODE_SHOW_R | SNODE_SHOW_G | SNODE_SHOW_B)) {
int ofs;
#ifdef __BIG_ENDIAN__
if (snode->flag & SNODE_SHOW_R) ofs = 2;
else if (snode->flag & SNODE_SHOW_G) ofs = 1;
else ofs = 0;
#else
if (snode->flag & SNODE_SHOW_R) ofs = 1;
else if (snode->flag & SNODE_SHOW_G) ofs = 2;
else ofs = 3;
if (ENDIAN_ORDER == B_ENDIAN) {
ofs = 3 - ofs;
}
#endif
glPixelZoom(snode->zoom, snode->zoom);
/* swap bytes, so alpha is most significant one, then just draw it as luminance int */
@ -3014,12 +3016,14 @@ void draw_nodespace_back_pix(const bContext *C, ARegion *ar, SpaceNode *snode)
else if (snode->flag & SNODE_SHOW_ALPHA) {
glPixelZoom(snode->zoom, snode->zoom);
/* swap bytes, so alpha is most significant one, then just draw it as luminance int */
if (ENDIAN_ORDER == B_ENDIAN)
glPixelStorei(GL_UNPACK_SWAP_BYTES, 1);
#ifdef __BIG_ENDIAN__
glPixelStorei(GL_UNPACK_SWAP_BYTES, 1);
#endif
glaDrawPixelsSafe(x, y, ibuf->x, ibuf->y, ibuf->x, GL_LUMINANCE, GL_UNSIGNED_INT, display_buffer);
#ifdef __BIG_ENDIAN__
glPixelStorei(GL_UNPACK_SWAP_BYTES, 0);
#endif
glPixelZoom(1.0f, 1.0f);
}
else if (snode->flag & SNODE_USE_ALPHA) {

@ -2858,9 +2858,10 @@ static void rna_def_space_node(BlenderRNA *brna)
{SNODE_USE_ALPHA, "COLOR_ALPHA", ICON_IMAGE_RGB_ALPHA, "Color and Alpha",
"Draw image with RGB colors and alpha transparency"},
{SNODE_SHOW_ALPHA, "ALPHA", ICON_IMAGE_ALPHA, "Alpha", "Draw alpha transparency channel"},
{SNODE_SHOW_R, "RED", 0, "Red", ""},
{SNODE_SHOW_G, "GREEN", 0, "Green", ""},
{SNODE_SHOW_B, "BLUE", 0, "Blue", ""},
/* XXX, we could use better icons here */
{SNODE_SHOW_R, "RED", ICON_COLOR, "Red", ""},
{SNODE_SHOW_G, "GREEN", ICON_COLOR, "Green", ""},
{SNODE_SHOW_B, "BLUE", ICON_COLOR, "Blue", ""},
{0, NULL, 0, NULL, NULL}
};

@ -114,15 +114,14 @@ static DerivedMesh *applyModifier(ModifierData *md, Object *ob,
ParticleSimulationData sim;
ParticleSystem *psys = NULL;
ParticleData *pa = NULL, *pars = NULL;
MFace *mface, *orig_mface;
MPoly *mpoly, *orig_mpoly;
MLoop *mloop, *orig_mloop;
MVert *mvert, *orig_mvert;
int i, totvert, totpart = 0, totface, maxvert, maxface, first_particle = 0;
int i, totvert, totpoly, totloop, maxvert, maxpoly, maxloop, totpart = 0, first_particle = 0;
short track = ob->trackflag % 3, trackneg, axis = pimd->axis;
float max_co = 0.0, min_co = 0.0, temp_co[3], cross[3];
float *size = NULL;
DM_ensure_tessface(dm); /* BMESH - UNTIL MODIFIER IS UPDATED FOR MPoly */
trackneg = ((ob->trackflag > 2) ? 1 : 0);
if (pimd->ob == ob) {
@ -175,15 +174,16 @@ static DerivedMesh *applyModifier(ModifierData *md, Object *ob,
pars = psys->particles;
totvert = dm->getNumVerts(dm);
totface = dm->getNumTessFaces(dm);
totpoly = dm->getNumPolys(dm);
totloop = dm->getNumLoops(dm);
maxvert = totvert * totpart;
maxface = totface * totpart;
maxpoly = totpoly * totpart;
maxloop = totloop * totpart;
psys->lattice = psys_get_lattice(&sim);
if (psys->flag & (PSYS_HAIR_DONE | PSYS_KEYED) || psys->pointcache->flag & PTCACHE_BAKED) {
float min_r[3], max_r[3];
INIT_MINMAX(min_r, max_r);
dm->getMinMax(dm, min_r, max_r);
@ -191,7 +191,7 @@ static DerivedMesh *applyModifier(ModifierData *md, Object *ob,
max_co = max_r[track];
}
result = CDDM_from_template(dm, maxvert, dm->getNumEdges(dm) * totpart, maxface, 0, 0);
result = CDDM_from_template(dm, maxvert, dm->getNumEdges(dm) * totpart, 0, maxloop, maxpoly);
mvert = result->getVertArray(result);
orig_mvert = dm->getVertArray(dm);
@ -263,29 +263,31 @@ static DerivedMesh *applyModifier(ModifierData *md, Object *ob,
add_v3_v3(mv->co, state.co);
}
mface = result->getTessFaceArray(result);
orig_mface = dm->getTessFaceArray(dm);
mpoly = result->getPolyArray(result);
orig_mpoly = dm->getPolyArray(dm);
mloop = result->getLoopArray(result);
orig_mloop = dm->getLoopArray(dm);
for (i = 0; i < maxface; i++) {
MFace *inMF;
MFace *mf = mface + i;
for (i = 0; i < maxpoly; i++) {
MPoly *inMP = orig_mpoly + i % totpoly;
MPoly *mp = mpoly + i;
if (pimd->flag & eParticleInstanceFlag_Parents) {
if (i / totface >= psys->totpart) {
if (i / totpoly >= psys->totpart) {
if (psys->part->childtype == PART_CHILD_PARTICLES) {
pa = psys->particles + (psys->child + i / totface - psys->totpart)->parent;
pa = psys->particles + (psys->child + i / totpoly - psys->totpart)->parent;
}
else {
pa = NULL;
}
}
else {
pa = pars + i / totface;
pa = pars + i / totpoly;
}
}
else {
if (psys->part->childtype == PART_CHILD_PARTICLES) {
pa = psys->particles + (psys->child + i / totface)->parent;
pa = psys->particles + (psys->child + i / totpoly)->parent;
}
else {
pa = NULL;
@ -298,19 +300,23 @@ static DerivedMesh *applyModifier(ModifierData *md, Object *ob,
if (pa->alive == PARS_DEAD && (pimd->flag & eParticleInstanceFlag_Dead) == 0) continue;
}
inMF = orig_mface + i % totface;
DM_copy_poly_data(dm, result, i % totface, i, 1);
*mf = *inMF;
DM_copy_poly_data(dm, result, i % totpoly, i, 1);
*mp = *inMP;
mp->loopstart += (i / totpoly) * totloop;
mf->v1 += (i / totface) * totvert;
mf->v2 += (i / totface) * totvert;
mf->v3 += (i / totface) * totvert;
if (mf->v4) {
mf->v4 += (i / totface) * totvert;
{
MLoop *inML = orig_mloop + inMP->loopstart;
MLoop *ml = mloop + mp->loopstart;
int j = mp->totloop;
DM_copy_loop_data(dm, result, inMP->loopstart, mp->loopstart, j);
for (; j; j--, ml++, inML++) {
ml->v = inML->v + ((i / totpoly) * totvert);
}
}
}
CDDM_calc_edges_tessface(result);
CDDM_calc_edges(result);
if (psys->lattice) {
end_latt_deform(psys->lattice);
@ -320,7 +326,6 @@ static DerivedMesh *applyModifier(ModifierData *md, Object *ob,
if (size)
MEM_freeN(size);
CDDM_tessfaces_to_faces(result); /*builds ngon faces from tess (mface) faces*/
CDDM_calc_normals(result);
return result;

@ -176,12 +176,16 @@ RenderResult *RE_engine_begin_result(RenderEngine *engine, int x, int y, int w,
result = render_result_new(re, &disprect, 0, RR_USE_MEM, layername);
/* todo: make this thread safe */
BLI_addtail(&engine->fullresult, result);
result->tilerect.xmin += re->disprect.xmin;
result->tilerect.xmax += re->disprect.xmin;
result->tilerect.ymin += re->disprect.ymin;
result->tilerect.ymax += re->disprect.ymin;
/* can be NULL if we CLAMP the width or height to 0 */
if (result) {
BLI_addtail(&engine->fullresult, result);
result->tilerect.xmin += re->disprect.xmin;
result->tilerect.xmax += re->disprect.xmin;
result->tilerect.ymin += re->disprect.ymin;
result->tilerect.ymax += re->disprect.ymin;
}
return result;
}