diff options
author | Campbell Barton <ideasman42@gmail.com> | 2012-08-09 23:59:36 +0400 |
---|---|---|
committer | Campbell Barton <ideasman42@gmail.com> | 2012-08-09 23:59:36 +0400 |
commit | 7a760b4804ab4574e01d7f6d75fc52f93aa2a3b3 (patch) | |
tree | ae5d20f881beb00420dd31bd79d9eba033b0ee48 | |
parent | 9f30c7147cdb726e503ff37ba585f5e35b090d8c (diff) |
generate COM_OpenCLKernels.cl.h automatically at build time, this allows editing COM_OpenCLKernels.cl and rebuilding and means we dont have to have both files in svn.
updates made to cmake and scons.
-rw-r--r-- | SConstruct | 44 | ||||
-rw-r--r-- | build_files/cmake/data_to_c.cmake | 25 | ||||
-rwxr-xr-x | release/datafiles/clkernelstoh.py | 70 | ||||
-rw-r--r-- | source/blender/compositor/CMakeLists.txt | 18 | ||||
-rw-r--r-- | source/blender/compositor/SConscript | 3 | ||||
-rw-r--r-- | source/blender/compositor/intern/COM_WorkScheduler.cpp | 3 | ||||
-rw-r--r-- | source/blender/compositor/operations/COM_OpenCLKernels.cl.h | 250 |
7 files changed, 92 insertions, 321 deletions
diff --git a/SConstruct b/SConstruct index 681e3964a13..faf1835e57d 100644 --- a/SConstruct +++ b/SConstruct @@ -445,6 +445,50 @@ if not os.path.isdir ( B.root_build_dir): # if not os.path.isdir(B.doc_build_dir) and env['WITH_BF_DOCS']: # os.makedirs ( B.doc_build_dir ) +################################### +# Ensure all data files are valid # +################################### +if not os.path.isdir ( B.root_build_dir + 'data_headers'): + os.makedirs ( B.root_build_dir + 'data_headers' ) +# use for includes +env['DATA_HEADERS'] = "#" + env['BF_BUILDDIR'] + "/data_headers" +def ensure_data(FILE_FROM, FILE_TO, VAR_NAME): + if os.sep == "\\": + FILE_FROM = FILE_FROM.replace("/", "\\") + FILE_TO = FILE_TO.replace("/", "\\") + + # first check if we need to bother. + if os.path.exists(FILE_TO): + if os.path.getmtime(FILE_FROM) < os.path.getmtime(FILE_TO): + return + + print(B.bc.HEADER + "Generating: " + B.bc.ENDC + "%r" % os.path.basename(FILE_TO)) + fpin = open(FILE_FROM, "rb") + fpin.seek(0, os.SEEK_END) + size = fpin.tell() + fpin.seek(0) + + fpout = open(FILE_TO, "w") + fpout.write("int %s_size = %d;\n" % (VAR_NAME, size)) + fpout.write("char %s[] = {\n" % VAR_NAME) + + while size > 0: + size -= 1 + if size % 32 == 31: + fpout.write("\n") + + fpout.write("%3d," % ord(fpin.read(1))) + fpout.write("\n 0};\n\n") + + fpin.close() + fpout.close() + +ensure_data("source/blender/compositor/operations/COM_OpenCLKernels.cl", + B.root_build_dir + "data_headers/COM_OpenCLKernels.cl.h", + "clkernelstoh_COM_OpenCLKernels_cl") + +##### END DATAFILES ########## + Help(opts.GenerateHelpText(env)) # default is new quieter output, but if you need to see the diff --git a/build_files/cmake/data_to_c.cmake b/build_files/cmake/data_to_c.cmake new file mode 100644 index 00000000000..b8b18269dc8 --- /dev/null +++ b/build_files/cmake/data_to_c.cmake @@ -0,0 +1,25 @@ +# cmake script, to be called on its own with 3 defined args +# +# - FILE_FROM +# - FILE_TO +# - VAR_NAME + +# not highly optimal, may replace with generated C program like makesdna +file(READ ${FILE_FROM} file_from_string HEX) +string(LENGTH ${file_from_string} _max_index) +math(EXPR size_on_disk ${_max_index}/2) + +file(REMOVE ${FILE_TO}) + +file(APPEND ${FILE_TO} "int ${VAR_NAME}_size = ${size_on_disk};\n") +file(APPEND ${FILE_TO} "char ${VAR_NAME}[] = {") + +set(_index 0) + +while(NOT _index EQUAL _max_index) + string(SUBSTRING "${file_from_string}" ${_index} 2 _pair) + file(APPEND ${FILE_TO} "0x${_pair},") + math(EXPR _index ${_index}+2) +endwhile() +# null terminator not essential but good if we want plane strings encoded +file(APPEND ${FILE_TO} "0x00};\n") diff --git a/release/datafiles/clkernelstoh.py b/release/datafiles/clkernelstoh.py deleted file mode 100755 index 9c24c9e2d03..00000000000 --- a/release/datafiles/clkernelstoh.py +++ /dev/null @@ -1,70 +0,0 @@ -#!/usr/bin/python -# -*- coding: utf-8 -*- - -# ***** BEGIN GPL LICENSE BLOCK ***** -# -# This program 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 2 -# of the License, or (at your option) any later version. -# -# This program 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, write to the Free Software Foundation, -# Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA. -# -# The Original Code is Copyright (C) 2012 Blender Foundation. -# All rights reserved. -# -# Contributor(s): Jeroen Bakker -# -# ***** END GPL LICENCE BLOCK ***** - -# <pep8 compliant> - -import sys -import os - -if len(sys.argv) < 2: - sys.stdout.write("Usage: clkernelstoh <cl_file>\n") - sys.exit(1) - -filename = sys.argv[1] - -try: - fpin = open(filename, "r") -except: - sys.stdout.write("Unable to open input %s\n" % sys.argv[1]) - sys.exit(1) - -if filename[0:2] == "." + os.sep: - filename = filename[2:] - -cname = filename + ".h" -sys.stdout.write("Making H file <%s>\n" % cname) - -filename = filename.split("/")[-1].split("\\")[-1] -filename = filename.replace(".", "_") - -try: - fpout = open(cname, "w") -except: - sys.stdout.write("Unable to open output %s\n" % cname) - sys.exit(1) - -fpout.write("/* clkernelstoh output of file <%s> */\n\n" % filename) -fpout.write("const char * clkernelstoh_%s = " % filename) - -lines = fpin.readlines() -for line in lines: - fpout.write("\"") - fpout.write(line.rstrip()) - fpout.write("\\n\" \\\n") -fpout.write("\"\\0\";\n") - -fpin.close() -fpout.close() diff --git a/source/blender/compositor/CMakeLists.txt b/source/blender/compositor/CMakeLists.txt index 4dc111cebd2..550441be9b7 100644 --- a/source/blender/compositor/CMakeLists.txt +++ b/source/blender/compositor/CMakeLists.txt @@ -50,6 +50,21 @@ set(INC_SYS ) +# --- data file --- +# ... may make this a macro +list(APPEND INC + ${CMAKE_CURRENT_BINARY_DIR}/operations +) +add_custom_command( + OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/operations/COM_OpenCLKernels.cl.h + COMMAND ${CMAKE_COMMAND} + -DFILE_FROM=${CMAKE_CURRENT_SOURCE_DIR}/operations/COM_OpenCLKernels.cl + -DFILE_TO=${CMAKE_CURRENT_BINARY_DIR}/operations/COM_OpenCLKernels.cl.h + -DVAR_NAME=clkernelstoh_COM_OpenCLKernels_cl + -P ${CMAKE_SOURCE_DIR}/build_files/cmake/data_to_c.cmake + DEPENDS operations/COM_OpenCLKernels.cl) +# --- end data file -- + set(SRC COM_compositor.h COM_defines.h @@ -638,6 +653,9 @@ set(SRC operations/COM_MaskOperation.cpp operations/COM_MaskOperation.h + + # generated file + ${CMAKE_CURRENT_BINARY_DIR}/operations/COM_OpenCLKernels.cl.h ) blender_add_lib(bf_compositor "${SRC}" "${INC}" "${INC_SYS}") diff --git a/source/blender/compositor/SConscript b/source/blender/compositor/SConscript index d7f18cfa436..9f947ca7327 100644 --- a/source/blender/compositor/SConscript +++ b/source/blender/compositor/SConscript @@ -11,4 +11,7 @@ incs += '../opencl ../nodes ../nodes/intern ../nodes/composite ' if env['OURPLATFORM'] in ('win32-vc', 'win32-mingw', 'linuxcross', 'win64-vc'): incs += ' ' + env['BF_PTHREADS_INC'] +# data files +incs += ' ' + env['DATA_HEADERS'] + env.BlenderLib ( 'bf_composite', sources, Split(incs), defines=defs, libtype=['core'], priority = [164] ) diff --git a/source/blender/compositor/intern/COM_WorkScheduler.cpp b/source/blender/compositor/intern/COM_WorkScheduler.cpp index 849c4affd9c..8c42b05edca 100644 --- a/source/blender/compositor/intern/COM_WorkScheduler.cpp +++ b/source/blender/compositor/intern/COM_WorkScheduler.cpp @@ -288,7 +288,8 @@ void WorkScheduler::initialize() g_context = clCreateContext(NULL, numberOfDevices, cldevices, clContextError, NULL, &error); if (error != CL_SUCCESS) { printf("CLERROR[%d]: %s\n", error, clewErrorString(error)); } - g_program = clCreateProgramWithSource(g_context, 1, &clkernelstoh_COM_OpenCLKernels_cl, 0, &error); + const char *cl_str[2] = {clkernelstoh_COM_OpenCLKernels_cl, NULL}; + g_program = clCreateProgramWithSource(g_context, 1, cl_str, 0, &error); error = clBuildProgram(g_program, numberOfDevices, cldevices, 0, 0, 0); if (error != CL_SUCCESS) { cl_int error2; diff --git a/source/blender/compositor/operations/COM_OpenCLKernels.cl.h b/source/blender/compositor/operations/COM_OpenCLKernels.cl.h deleted file mode 100644 index e72ed5bf755..00000000000 --- a/source/blender/compositor/operations/COM_OpenCLKernels.cl.h +++ /dev/null @@ -1,250 +0,0 @@ -/* clkernelstoh output of file <COM_OpenCLKernels_cl> */ - -const char * clkernelstoh_COM_OpenCLKernels_cl = "/*\n" \ -" * Copyright 2011, Blender Foundation.\n" \ -" *\n" \ -" * This program is free software; you can redistribute it and/or\n" \ -" * modify it under the terms of the GNU General Public License\n" \ -" * as published by the Free Software Foundation; either version 2\n" \ -" * of the License, or (at your option) any later version.\n" \ -" *\n" \ -" * This program is distributed in the hope that it will be useful,\n" \ -" * but WITHOUT ANY WARRANTY; without even the implied warranty of\n" \ -" * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the\n" \ -" * GNU General Public License for more details.\n" \ -" *\n" \ -" * You should have received a copy of the GNU General Public License\n" \ -" * along with this program; if not, write to the Free Software Foundation,\n" \ -" * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.\n" \ -" *\n" \ -" * Contributor:\n" \ -" * Jeroen Bakker\n" \ -" * Monique Dewanchand\n" \ -" */\n" \ -"\n" \ -"/// This file contains all opencl kernels for node-operation implementations\n" \ -"\n" \ -"// Global SAMPLERS\n" \ -"const sampler_t SAMPLER_NEAREST = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;\n" \ -"const sampler_t SAMPLER_NEAREST_CLAMP = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;\n" \ -"\n" \ -"__constant const int2 zero = {0,0};\n" \ -"\n" \ -"// KERNEL --- BOKEH BLUR ---\n" \ -"__kernel void bokehBlurKernel(__read_only image2d_t boundingBox, __read_only image2d_t inputImage,\n" \ -" __read_only image2d_t bokehImage, __write_only image2d_t output,\n" \ -" int2 offsetInput, int2 offsetOutput, int radius, int step, int2 dimension, int2 offset)\n" \ -"{\n" \ -" int2 coords = {get_global_id(0), get_global_id(1)};\n" \ -" coords += offset;\n" \ -" float tempBoundingBox;\n" \ -" float4 color = {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" \ -"\n" \ -" tempBoundingBox = read_imagef(boundingBox, SAMPLER_NEAREST, coords).s0;\n" \ -"\n" \ -" if (tempBoundingBox > 0.0f && radius > 0 ) {\n" \ -" const int2 bokehImageDim = get_image_dim(bokehImage);\n" \ -" const int2 bokehImageCenter = bokehImageDim/2;\n" \ -" const int2 minXY = max(realCoordinate - radius, zero);\n" \ -" const int2 maxXY = min(realCoordinate + radius, dimension);\n" \ -" int nx, ny;\n" \ -"\n" \ -" float2 uv;\n" \ -" int2 inputXy;\n" \ -"\n" \ -" for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny +=step, inputXy.y+=step) {\n" \ -" uv.y = ((realCoordinate.y-ny)/radius2)*bokehImageDim.y+bokehImageCenter.y;\n" \ -"\n" \ -" for (nx = minXY.x, inputXy.x = nx - offsetInput.x; nx < maxXY.x ; nx +=step, inputXy.x+=step) {\n" \ -" 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" \ -" multiplyer += bokeh;\n" \ -" }\n" \ -" }\n" \ -" color /= multiplyer;\n" \ -"\n" \ -" } else {\n" \ -" int2 imageCoordinates = realCoordinate - offsetInput;\n" \ -" color = read_imagef(inputImage, SAMPLER_NEAREST, imageCoordinates);\n" \ -" }\n" \ -"\n" \ -" 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 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 tempColor;\n" \ -" float4 bokeh;\n" \ -" float size;\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_center = read_imagef(inputSize, SAMPLER_NEAREST, inputCoordinate).s0;\n" \ -" color_accum = read_imagef(inputImage, SAMPLER_NEAREST, inputCoordinate);\n" \ -" readColor = color_accum;\n" \ -"\n" \ -" if (size_center > threshold) {\n" \ -" for (int ny = miny; ny < maxy; ny += step) {\n" \ -" inputCoordinate.s1 = ny - offsetInput.s1;\n" \ -" float dy = ny - realCoordinate.s1;\n" \ -" for (int nx = minx; nx < maxx; nx += step) {\n" \ -" float dx = nx - realCoordinate.s0;\n" \ -" if (dx != 0 || dy != 0) {\n" \ -" inputCoordinate.s0 = nx - offsetInput.s0;\n" \ -" size = read_imagef(inputSize, SAMPLER_NEAREST, inputCoordinate).s0;\n" \ -" if (size > threshold) {\n" \ -" if (size >= fabs(dx) && size >= fabs(dy)) {\n" \ -" float2 uv = {256.0f + dx * 255.0f / size,\n" \ -" 256.0f + dy * 255.0f / size};\n" \ -" bokeh = read_imagef(bokehImage, SAMPLER_NEAREST, uv);\n" \ -" tempColor = read_imagef(inputImage, SAMPLER_NEAREST, inputCoordinate);\n" \ -" color_accum += bokeh * tempColor;\n" \ -" multiplier_accum += bokeh;\n" \ -" }\n" \ -" }\n" \ -" }\n" \ -" }\n" \ -" }\n" \ -" }\n" \ -"\n" \ -" color = color_accum * (1.0f / multiplier_accum);\n" \ -"\n" \ -" /* blend in out values over the threshold, otherwise we get sharp, ugly transitions */\n" \ -" if ((size_center > threshold) &&\n" \ -" (size_center < threshold * 2.0f))\n" \ -" {\n" \ -" /* factor from 0-1 */\n" \ -" float fac = (size_center - threshold) / threshold;\n" \ -" color = (readColor * (1.0f - fac)) + (color * fac);\n" \ -" }\n" \ -"\n" \ -" write_imagef(output, coords, color);\n" \ -" }\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" \ -" int2 offset)\n" \ -"{\n" \ -" int2 coords = {get_global_id(0), get_global_id(1)};\n" \ -" coords += offset;\n" \ -" const int2 realCoordinate = coords + offsetOutput;\n" \ -"\n" \ -" const int2 minXY = max(realCoordinate - scope, zero);\n" \ -" const int2 maxXY = min(realCoordinate + scope, dimension);\n" \ -"\n" \ -" float value = 0.0f;\n" \ -" int nx, ny;\n" \ -" 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 measuredDistance = deltaX * deltaX + deltaY * deltaY;\n" \ -" if (measuredDistance <= distanceSquared) {\n" \ -" value = max(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0);\n" \ -" }\n" \ -" }\n" \ -" }\n" \ -"\n" \ -" float4 color = {value,0.0f,0.0f,0.0f};\n" \ -" write_imagef(output, coords, color);\n" \ -"}\n" \ -"\n" \ -"// KERNEL --- DILATE ---\n" \ -"__kernel void erodeKernel(__read_only image2d_t inputImage, __write_only image2d_t output,\n" \ -" int2 offsetInput, int2 offsetOutput, int scope, int distanceSquared, int2 dimension,\n" \ -" int2 offset)\n" \ -"{\n" \ -" int2 coords = {get_global_id(0), get_global_id(1)};\n" \ -" coords += offset;\n" \ -" const int2 realCoordinate = coords + offsetOutput;\n" \ -"\n" \ -" const int2 minXY = max(realCoordinate - scope, zero);\n" \ -" const int2 maxXY = min(realCoordinate + scope, dimension);\n" \ -"\n" \ -" float value = 1.0f;\n" \ -" int nx, ny;\n" \ -" int2 inputXy;\n" \ -"\n" \ -" for (ny = minXY.y, inputXy.y = ny - offsetInput.y ; ny < maxXY.y ; ny ++, inputXy.y++) {\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 = min(value, read_imagef(inputImage, SAMPLER_NEAREST, inputXy).s0);\n" \ -" }\n" \ -" }\n" \ -" }\n" \ -"\n" \ -" float4 color = {value,0.0f,0.0f,0.0f};\n" \ -" write_imagef(output, coords, color);\n" \ -"}\n" \ -"\n" \ -"// KERNEL --- DIRECTIONAL BLUR ---\n" \ -"__kernel void directionalBlurKernel(__read_only image2d_t inputImage, __write_only image2d_t output,\n" \ -" int2 offsetOutput, int iterations, float scale, float rotation, float2 translate,\n" \ -" float2 center, int2 offset)\n" \ -"{\n" \ -" int2 coords = {get_global_id(0), get_global_id(1)};\n" \ -" coords += offset;\n" \ -" const int2 realCoordinate = coords + offsetOutput;\n" \ -"\n" \ -" float4 col;\n" \ -" float2 ltxy = translate;\n" \ -" float lsc = scale;\n" \ -" float lrot = rotation;\n" \ -"\n" \ -" col = read_imagef(inputImage, SAMPLER_NEAREST, realCoordinate);\n" \ -"\n" \ -" /* blur the image */\n" \ -" for (int i = 0; i < iterations; ++i) {\n" \ -" const float cs = cos(lrot), ss = sin(lrot);\n" \ -" const float isc = 1.0f / (1.0f + lsc);\n" \ -"\n" \ -" const float v = isc * (realCoordinate.s1 - center.s1) + ltxy.s1;\n" \ -" const float u = isc * (realCoordinate.s0 - center.s0) + ltxy.s0;\n" \ -" float2 uv = {\n" \ -" cs * u + ss * v + center.s0,\n" \ -" cs * v - ss * u + center.s1\n" \ -" };\n" \ -"\n" \ -" col += read_imagef(inputImage, SAMPLER_NEAREST_CLAMP, uv);\n" \ -"\n" \ -" /* double transformations */\n" \ -" ltxy += translate;\n" \ -" lrot += rotation;\n" \ -" lsc += scale;\n" \ -" }\n" \ -"\n" \ -" col *= (1.0f/(iterations+1));\n" \ -"\n" \ -" write_imagef(output, coords, col);\n" \ -"}\n" \ -"\0"; |