Skip to content
Snippets Groups Projects
Commit 7a760b48 authored by Campbell Barton's avatar Campbell Barton
Browse files

generate COM_OpenCLKernels.cl.h automatically at build time, this allows... 

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.
parent 9f30c714
Branches
Tags
No related merge requests found
Show whitespace changes
Inline Side-by-side
Showing
with 92 additions and 321 deletions
SConstruct
+ 44
− 0
View file @ 7a760b48
......@@ -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
......
build_files/cmake/data_to_c.cmake 0 → 100644
+ 25
− 0
View file @ 7a760b48
# 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")
release/datafiles/clkernelstoh.py deleted 100755 → 0
+ 0
− 70
View file @ 9f30c714
#!/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()
source/blender/compositor/CMakeLists.txt
+ 18
− 0
View file @ 7a760b48
......@@ -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}")
source/blender/compositor/SConscript
+ 3
− 0
View file @ 7a760b48
......@@ -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] )
source/blender/compositor/intern/COM_WorkScheduler.cpp
+ 2
− 1
View file @ 7a760b48
......@@ -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;
......
source/blender/compositor/operations/COM_OpenCLKernels.cl.h deleted 100644 → 0
+ 0
− 250
View file @ 9f30c714
/* 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";
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Please to comment