Forge
opencl/stream.cpp
/*******************************************************
* Copyright (c) 2015-2019, ArrayFire
* All rights reserved.
*
* This file is distributed under 3-clause BSD license.
* The complete license agreement can be obtained at:
* http://arrayfire.com/licenses/BSD-3-Clause
********************************************************/
#include <forge.h>
#include "cl_helpers.h"
#include <cmath>
#include <ctime>
#include <vector>
#include <sstream>
#include <iostream>
#include <iterator>
#include <algorithm>
using namespace cl;
using namespace std;
const unsigned DIMX = 640;
const unsigned DIMY = 480;
const float MINIMUM = 1.0f;
const float MAXIMUM = 20.f;
const float STEP = 2.0f;
const int NELEMS = (MAXIMUM-MINIMUM+1)/STEP;
#define USE_FORGE_OPENCL_COPY_HELPERS
#include <ComputeCopy.h>
static const std::string streamKernel =
R"EOK(
constant float AF_BLUE[4] = {0.0588f , 0.1137f , 0.2745f , 1.0f};
constant float AF_ORANGE[4] = {0.8588f , 0.6137f , 0.0745f , 1.0f};
kernel
void genColorsKernel(global float* colors, int NELEMS)
{
const size_t nelems = NELEMS * NELEMS * NELEMS;
int i = get_global_id(0);
if (i<nelems) {
if (i%2==0) {
colors[3*i+0] = AF_ORANGE[0];
colors[3*i+1] = AF_ORANGE[1];
colors[3*i+2] = AF_ORANGE[2];
} else {
colors[3*i+0] = AF_BLUE[0];
colors[3*i+1] = AF_BLUE[1];
colors[3*i+2] = AF_BLUE[2];
}
}
}
kernel
void pointGenKernel(global float* points, global float* dirs, int nBBS0,
int NELEMS, float MINIMUM, float STEP)
{
int k = get_group_id(0) / nBBS0;
int i = get_local_size(0) * (get_group_id(0)-k*nBBS0) + get_local_id(0);
int j = get_global_id(1);
if (i<NELEMS && j<NELEMS && k<NELEMS) {
float x = MINIMUM + i*STEP;
float y = MINIMUM + j*STEP;
float z = MINIMUM + k*STEP;
int id = i + j * NELEMS + k * NELEMS*NELEMS;
points[3*id+0] = x;
points[3*id+1] = y;
points[3*id+2] = z;
dirs[3*id+0] = x-10.f;
dirs[3*id+1] = y-10.f;
dirs[3*id+2] = z-10.f;
}
}
)EOK";
inline int divup(int a, int b)
{
return (a+b-1)/b;
}
void generatePoints(Buffer& points, Buffer &dirs, Buffer& colors,
CommandQueue& queue, Device& device)
{
static bool compileFlag = true;
static cl::Program prog;
static cl::Kernel pointGenKernel;
static cl::Kernel colorsKernel;
if (compileFlag) {
try {
prog = cl::Program(queue.getInfo<CL_QUEUE_CONTEXT>(), streamKernel, false);
std::vector<cl::Device> devs;
devs.push_back(device);
prog.build(devs);
pointGenKernel = cl::Kernel(prog, "pointGenKernel");
colorsKernel = cl::Kernel(prog, "genColorsKernel");
} catch (cl::Error err) {
std::cout << "Compile Errors: " << std::endl;
std::cout << err.what() << err.err() << std::endl;
std::cout << prog.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device) << std::endl;
exit(255);
}
std::cout<< "Kernels compiled successfully" << std::endl;
compileFlag = false;
}
static const NDRange local(8, 8);
int blk_x = divup(NELEMS, local[0]);
int blk_y = divup(NELEMS, local[1]);
NDRange global(NELEMS * local[0] * blk_x, local[1] * blk_y);
pointGenKernel.setArg(0, points);
pointGenKernel.setArg(1, dirs);
pointGenKernel.setArg(2, blk_x);
pointGenKernel.setArg(3, NELEMS);
pointGenKernel.setArg(4, MINIMUM);
pointGenKernel.setArg(5, STEP);
queue.enqueueNDRangeKernel(pointGenKernel, cl::NullRange, global, local);
const int numElems = NELEMS*NELEMS*NELEMS;
static const NDRange thrds(64, 1);
NDRange glob(thrds[0] * divup(numElems, thrds[0]), thrds[1]);
colorsKernel.setArg(0, colors);
colorsKernel.setArg(1, NELEMS);
queue.enqueueNDRangeKernel(colorsKernel, cl::NullRange, glob, thrds);
}
int main(void)
{
try {
/*
* First Forge call should be a window creation call
* so that necessary OpenGL context is created for any
* other forge::* object to be created successfully
*/
forge::Window wnd(DIMX, DIMY, "3D Vector Field Demo");
wnd.makeCurrent();
chart.setAxesLimits(MINIMUM-1.0f, MAXIMUM,
MINIMUM-1.0f, MAXIMUM,
MINIMUM-1.0f, MAXIMUM);
chart.setAxesTitles("x-axis", "y-axis", "z-axis");
int numElems = NELEMS*NELEMS*NELEMS;
forge::VectorField field = chart.vectorField(numElems, forge::f32);
field.setColor(0.f, 1.f, 0.f, 1.f);
/*
* Helper function to create a CLGL interop context.
* This function checks for if the extension is available
* and creates the context on the appropriate device.
* Note: context and queue are defined in cl_helpers.h
*/
context = createCLGLContext(wnd);
Device device = context.getInfo<CL_CONTEXT_DEVICES>()[0];
queue = CommandQueue(context, device);
cl::Buffer points(context, CL_MEM_READ_WRITE, sizeof(float)*3*numElems);
cl::Buffer colors(context, CL_MEM_READ_WRITE, sizeof(float)*3*numElems);
cl::Buffer dirs(context, CL_MEM_READ_WRITE, sizeof(float)*3*numElems);
GfxHandle* handles[3];
createGLBuffer(&handles[0], field.vertices(), FORGE_VERTEX_BUFFER);
createGLBuffer(&handles[1], field.colors(), FORGE_VERTEX_BUFFER);
createGLBuffer(&handles[2], field.directions(), FORGE_VERTEX_BUFFER);
generatePoints(points, dirs, colors, queue, device);
copyToGLBuffer(handles[0], (ComputeResourceHandle)points(), field.verticesSize());
copyToGLBuffer(handles[1], (ComputeResourceHandle)colors(), field.colorsSize());
copyToGLBuffer(handles[2], (ComputeResourceHandle)dirs(), field.directionsSize());
do {
wnd.draw(chart);
} while(!wnd.close());
releaseGLBuffer(handles[0]);
releaseGLBuffer(handles[1]);
releaseGLBuffer(handles[2]);
} catch (forge::Error err) {
std::cout << err.what() << "(" << err.err() << ")" << std::endl;
} catch (cl::Error err) {
std::cout << err.what() << "(" << err.err() << ")" << std::endl;
}
return 0;
}