In the previous article in this series, I introduced the notion of making a series of Fuses based on the lessons from Vivo & Lowe’s Book of Shaders. I provided the template I usually start from and a brief rundown of how the parts work. Today, we’ll look at the simplest of kernels. As pointed out at BoS, the customary first program people write when learning a new language is one that prints the message “Hello World!” to the standard output. Graphics programming is, by its nature, not well suited to printing messages of this kind, so this “Hello World” program instead fills the output image with a solid color. The OpenCL program is a little longer than the GLSL one given at BoS, but not much. Most of the added complexity comes from the interface between the Fuse and OCL.
BoS’ function looks like this:
void main() {
gl_FragColor = vec4(0.340,0.000,0.340,1.000);
}
We’ll have to jump through a few more hoops than that. Since we’re not interfacing directly with a live image buffer, we need to pass an image into the function and include a way for OCL to tell Fusion exactly which pixel in that image it’s processing. Here’s the OCL code:
clsource = [[
__kernel
void hello(FuWriteImage_t img, int2 imgsize)
{
int2 ipos = (int2)(get_global_id(1), get_global_id(0));
const float4 col = (float4)(0.3, 0.4, 0.5, 1.0);
FuWriteImagef(img, ipos, imgsize, col);
}
]]
Remember from the previous article that we feed the OCL code into the OCLManager as a string. Therefore, every bit of OpenCL language code needs to appear between the [[ ]].
Next, we have the __kernel qualifier, which has special meaning to the OCLManager.
In the function declaration: void hello(FuWriteImage_t img, int2 imgsize), the “void” keyword indicates that the function does not return a value other than “true” to the function that called it. The output pixel color is written directly to the image buffer instead of going back out through a returned value.
Unlike Lua, OpenCL requires variable types to be explicitly declared. In Lua, you can store anything in any variable, regardless of its type. You can replace a float value with a string, and Lua doesn’t care. In OpenCL, data types are fixed—a variable that has been declared to hold an integer vector (int2) can never receive a string or a single float value. FuWriteImage_t img declares a variable called img intended to hold a writeable image buffer. int2 imgsize declares a two-dimensional integer matrix (technically a vector, although not used as such) called imgsize which will be used to pass the width and height of the image into the function.
In the function body, we have another int2 declaration:
int2 ipos = (int2)(get_global_id(1), get_global_id(0));
get_global_id is a built-in variable that uniquely identifies the work item in the current thread. In this case, Fusion will assign a global_id based on the coordinates of the pixel under consideration. For some reason the coordinates are reversed, so global_id(0) is y, and global_id(1) is x. Once again, the variable is declared as an int2, and the data that goes into it must also be identified as int2. Apparently OpenCL won’t infer the data type from context—you have to tell it exactly what each piece of data is as you use it. For scalar (one-dimensional) values, this usually means putting a f or i after the value, like so: float myvar = 3f;
Notice also the semi-colon at the end of the line. Unlike Lua, OpenCL statements must always terminate in a semi-colon.
float4 col = (float4)(0.3, 0.4, 0.5, 1.0);
This is the line that actually does something. A variable of type float4 called “col” is declared and assigned the vector [0.3, 0.4, 0.5, 1.0]. In RGB, this is a pale blue solid. Hopefully it should be fairly clear how that relates to the Shader Language example given at the beginning.
FuWriteImagef(img, ipos, imgsize, col);
The FuWriteImagef() function is called, which takes four arguments: the image to be written, which was passed to the program in its own arguments; the coordinates of the pixel that is to be written, given by the work item id; the size of the image, again passed in as an argument from the Fuse; and the value to be assigned to the pixel, which we assigned in the previous line. The function sends that information back out to the Fuse.
Input and Output
Now all that remains is the issue of how to pass data into and out of the OpenCL kernel. For that, we go to the Process() function and add some lines after we’ve initialized the working and output images.
local success = false
if prog then
local imgcl = prog:CreateImage(img, "write")
if imgcl then
local kernel = prog:CreateKernel("hello")
if kernel then
prog:SetArg(kernel, 0, imgcl)
prog:SetArgInt(kernel, 1, img.Width, img.Height)
success = prog:RunKernel(kernel)
end
if success then
success = prog:Download(imgcl, out)
end
end
end
First, we initialize a flag that we can use to prevent writing an image if the OCL execution fails for some reason. Then we test to make sure the OpenCL code compiled properly.
Remember that prog is a handle that contains the OpenCL program. CreateImage is a method available in prog that converts a Fusion image into an image buffer that OpenCL can use. The image is made writable with the “write” argument. That could also be “read” for a read-only image or “readwrite” in case we need OpenCL to both read and write to the buffer. This buffer is stored in imgcl. We then immediately test to be sure that imgcl was created successfully.
The CreateKernel() method is used to prepare the OpenCL program to be sent out to the processors as a work item. We give it the name of the kernel we want to use, then test to be certain that it was created.
Finally, we get to the point where we can start feeding OpenCL the information it needs to do its job. For that, we use the SetArg() and SetArgInt() methods. The SetArg() method takes several arguments. The first is always the kernel that is meant to receive the data, followed by an index indicating which of the kernel’s arguments this data is for. Recall from the OCL function declaration that hello() expects an image and a 2d integer matrix as input. The order is important because the image argument has index of 0, and the int2 has an index of 1. Every following argument to SetArg() is the expected values. In this case, it only gets one thing: imgcl.
SetArgInt() works exactly the same way, except that it only sends out integer values. This is necessary because Lua’s data typing is so relaxed and OpenCL’s is so restrictive. This time, since the kernel expects a matrix, we send the width and height values one at a time as separate arguments to the SetArgInt() method.
Once the kernel has received its arguments, it is executed with RunKernel(). The void function type means that once it is done it will return 1 (or true). That return code is stored in the success variable, allowing us to test whether the kernel completed operation.
If it did, then the Download() method is called to convert the imgcl image buffer back into a Fusion image and store it in out.
Finally, we use OutImage:Set(req,out) to assign out to the actual output of the Fuse. And here is the utterly unexciting result:
Next time we’ll look at methods for applying temporal and spatial variation to the image. In addition, we’ll set up a control panel with controls to choose the color and select which algorithm to use.
<< Previous Article — Fuse Template and Structure
Next Article — Position and Time >>
