// Window resize handler callback
//*****************************************************************************
void Reshape(int w, int h)
{
w = MAX(w,1);
h = MAX(h,1);
iGraphicsWinWidth = w;
iGraphicsWinHeight = h;
glBindTexture(GL_TEXTURE_2D, tex_screen);
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, w, h, 0, GL_RGBA, GL_UNSIGNED_BYTE, NULL);
image_width = w;
image_height = h;
num_texels = image_width * image_height;
num_values = num_texels * 4;
size_tex_data = sizeof(GLubyte) * num_values;
if( cl_pbos[0] != 0 ) {
// update sizes of pixel buffer objects
glBindBuffer(GL_ARRAY_BUFFER, pbo_source);
glBufferData(GL_ARRAY_BUFFER, size_tex_data, NULL, GL_DYNAMIC_DRAW);
glBindBuffer(GL_ARRAY_BUFFER, pbo_dest);
glBufferData(GL_ARRAY_BUFFER, size_tex_data, NULL, GL_DYNAMIC_DRAW);
glBindBuffer(GL_ARRAY_BUFFER,0);
// release current mem objects
clReleaseMemObject(cl_pbos[0]);
clReleaseMemObject(cl_pbos[1]);
// create new objects for the current sizes
if( bGLinterop ) {
cl_pbos[0] = clCreateFromGLBuffer(cxGPUContext, CL_MEM_READ_ONLY, pbo_source, &ciErrNum);
cl_pbos[1] = clCreateFromGLBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, pbo_dest, &ciErrNum);
} else {
cl_pbos[0] = clCreateBuffer(cxGPUContext, CL_MEM_READ_ONLY, 4 * image_width * image_height, NULL, &ciErrNum);
cl_pbos[1] = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, 4 * image_width * image_height, NULL, &ciErrNum);
}
// update kernel arguments
clSetKernelArg(ckKernel, 0, sizeof(cl_mem), (void *) &(cl_pbos[0]));
clSetKernelArg(ckKernel, 1, sizeof(cl_mem), (void *) &(cl_pbos[1]));
clSetKernelArg(ckKernel, 2, sizeof(cl_int), &image_width);
clSetKernelArg(ckKernel, 3, sizeof(cl_int), &image_height);
}
glutPostRedisplay();
}
void BounceKernel::Load(cl_GLuint pos_buffer, size_t num_objects)
{
m_num_objects = num_objects;
std::string program_code = ReadFile(ProgramFileName, m_success);
if (!m_success)
{
m_log += "[FAIL] Unable to read program file: " + ProgramFileName + "\n";
return;
}
KernelBaseCL::Load(program_code, KernelName);
if (!m_success)
return;
cl_int cl_err;
cl_mem buf = clCreateFromGLBuffer(m_system.Context(), CL_MEM_READ_WRITE, pos_buffer, &cl_err);
if (cl_err != CL_SUCCESS)
{
m_success = false;
m_log += "[FAIL] Unable to create position buffer: EC = " + std::to_string(cl_err) + "\n";
return;
}
m_position = buf;
m_success = true;
}
sge::opencl::memory_object::buffer::buffer(
context::object &_context,
sge::renderer::vertex::buffer &_vb,
memory_object::renderer_buffer_lock_mode const _lock_mode)
:
impl_(nullptr),
byte_size_(
static_cast<byte_size::value_type>(
_vb.linear_size() * _vb.format().get().stride().get()))
{
cl_int error_code;
impl_ =
clCreateFromGLBuffer(
_context.impl(),
sge::opencl::impl::memory_object::renderer_buffer_lock_mode_to_cl_mem_flags(
_lock_mode),
dynamic_cast<sge::renderer::opengl::buffer::base &>(
_vb).id().get(),
&error_code);
opencl::impl::handle_error(
error_code,
FCPPT_TEXT("clCreateFromGLBuffer"));
}
bool GLBuffer::bufferData(const GLvoid *data, GLsizeiptr size, AccessType at, GLenum usage)
{
assert(m_ctx != nullptr);
OGLF->glGetError(); // clear any previous errors
OGLF->glBindBuffer(GL_ARRAY_BUFFER, m_vbo);
OGLF->glBufferData(GL_ARRAY_BUFFER, size, data, usage);
{
GLenum err = OGLF->glGetError();
if (err != GL_NO_ERROR)
{
ERRORM("Failed to buffer data: " << ogl::errorToStr(err));
return false;
}
}
OGLF->glBindBuffer(GL_ARRAY_BUFFER, 0);
cl_int err = CL_SUCCESS;
cl_mem mem = clCreateFromGLBuffer(m_ctx, at, m_vbo, &err);
if (err != CL_SUCCESS)
{
ERRORM("Failed to create OpenCL buffer: " << ocl::errorToStr(err));
return false;
}
clReleaseMemObject(m_mem);
m_mem = mem;
return true;
}
void buffers()
{
focl[0] = clCreateBuffer(mycontext,CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,
SIZE*sizeof(cl_float8),&f[0][0][0],&err);
focl[1] = clCreateBuffer(mycontext,CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,
SIZE*sizeof(cl_float8),&f[1][0][0],&err);
dist_ocl = clCreateBuffer(mycontext,CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
SIZE*sizeof(int),&dist[0],&err);
omega_ocl = clCreateBuffer(mycontext,CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
DIRECTIONS*DIRECTIONS*sizeof(cl_float8),&omega[0][0],&err);
eye_ocl = clCreateBuffer(mycontext,CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
3*sizeof(cl_float4),&eye[0][0],&err);
rbuffer_ocl = clCreateFromGLBuffer(mycontext,CL_MEM_READ_WRITE,OGL_VBO,&err);
nbuffer_ocl = clCreateBuffer(mycontext,CL_MEM_READ_WRITE,WIDTH*LENGTH*
sizeof(cl_float4),NULL,&err);
clSetKernelArg(mykrn_update,2,sizeof(cl_mem),(void *)&dist_ocl);
clSetKernelArg(mykrn_update,3,sizeof(cl_mem),(void *)&omega_ocl);
clSetKernelArg(mykrn_heights,0,sizeof(cl_mem), (void *)&rbuffer_ocl);
clSetKernelArg(mykrn_normals,0,sizeof(cl_mem), (void *)&rbuffer_ocl);
clSetKernelArg(mykrn_normals,1,sizeof(cl_mem), (void *)&nbuffer_ocl);
clSetKernelArg(mykrn_colors,0,sizeof(cl_mem), (void *)&rbuffer_ocl);
clSetKernelArg(mykrn_colors,1,sizeof(cl_float4), &lightdir);
clSetKernelArg(mykrn_colors,2,sizeof(cl_mem), (void *)&eye_ocl);
clSetKernelArg(mykrn_colors,3,sizeof(cl_mem), (void *)&nbuffer_ocl);
}
bool
OsdCLGLVertexBuffer::allocate(cl_context clContext) {
assert(clContext);
// create GL buffer first
int size = _numElements * _numVertices * sizeof(float);
GLint prev = 0;
glGetIntegerv(GL_ARRAY_BUFFER_BINDING, &prev);
glGenBuffers(1, &_vbo);
glBindBuffer(GL_ARRAY_BUFFER, _vbo);
glBufferData(GL_ARRAY_BUFFER, size, 0, GL_STREAM_DRAW);
glBindBuffer(GL_ARRAY_BUFFER, prev);
if (glGetError() != GL_NO_ERROR) return false;
// register vbo as cl memory
cl_int err;
_clMemory = clCreateFromGLBuffer(clContext,
CL_MEM_READ_WRITE, _vbo, &err);
if (err != CL_SUCCESS) return false;
return true;
}
//--------------------------------------------------------------
bool RayTracingKernel::setup(cl_context _context, cl_command_queue _commandQ)
{
if (!ofxClKernel::setup(_context, _commandQ)) return false;
if (pbo)
{
// delete old buffer
clReleaseMemObject(pbo_cl);
glDeleteBuffersARB(1, &pbo);
}
// create pixel buffer object for display
glGenBuffersARB(1, &pbo);
glBindBufferARB(GL_PIXEL_UNPACK_BUFFER_ARB, pbo);
glBufferDataARB(GL_PIXEL_UNPACK_BUFFER_ARB, width * height * sizeof(GLubyte) * 4, 0, GL_STREAM_DRAW_ARB);
glBindBufferARB(GL_PIXEL_UNPACK_BUFFER_ARB, 0);
#ifdef GL_INTEROP
// create OpenCL buffer from GL PBO
pbo_cl = clCreateFromGLBuffer(context, CL_MEM_WRITE_ONLY, pbo, &clErr);
if (!checkOpenClError(clErr, "clCreateFromGLBuffer")) return false;
#else
pbo_cl = clCreateBuffer(context, CL_MEM_WRITE_ONLY, width * height * sizeof(GLubyte) * 4, NULL, &clErr);
if (!checkOpenClError(clErr, "clCreateBuffer")) return false;
#endif
clErr = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &pbo_cl);
if (!checkOpenClError(clErr, "clSetKernelArg 0")) return false;
}
void configure_shared_data() {
int err;
/* Create vertex array objects */
glGenVertexArrays(1, &vao);
glBindVertexArray(vao);
/* Create vertex buffers */
glGenBuffers(1, &vbo);
/* VBO for coordinates of first square */
glBindBuffer(GL_ARRAY_BUFFER, vbo);
glBufferData(GL_ARRAY_BUFFER, 4 * NUM_VERTICES*sizeof(GLfloat),
NULL, GL_DYNAMIC_DRAW);
glVertexAttribPointer(0, 4, GL_FLOAT, GL_FALSE, 0, 0);
glEnableVertexAttribArray(0);
/* Create memory objects from the VBOs */
vertex_buffer = clCreateFromGLBuffer(context, CL_MEM_WRITE_ONLY,
vbo, &err);
if(err < 0) {
perror("Couldn't create a buffer object from the VBO");
exit(1);
}
/* Set kernel arguments */
err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &vertex_buffer);
err |= clSetKernelArg(kernel, 1, sizeof(float), &tick);
if(err < 0) {
printf("Couldn't set a kernel argument");
exit(1);
};
}
EXTERN_C_ENTER
JNIEXPORT jlong JNICALL Java_org_lwjgl_opencl_CL10GL_nclCreateFromGLBuffer(JNIEnv *__env, jclass clazz, jlong contextAddress, jlong flags, jint bufobj, jlong errcode_retAddress, jlong __functionAddress) {
cl_context context = (cl_context)(intptr_t)contextAddress;
cl_int *errcode_ret = (cl_int *)(intptr_t)errcode_retAddress;
clCreateFromGLBufferPROC clCreateFromGLBuffer = (clCreateFromGLBufferPROC)(intptr_t)__functionAddress;
UNUSED_PARAMS(__env, clazz)
return (jlong)(intptr_t)clCreateFromGLBuffer(context, flags, bufobj, errcode_ret);
}
//--------------------------------------------------------------------------------------
// Name: Initialize()
// Desc: Initialize the OpenCL implementation of the cloth simulation
//--------------------------------------------------------------------------------------
BOOL CClothSimCL::Initialize( cl_context context, cl_device_id device )
{
// Call the base class first
if( !CClothSim::Initialize() )
return FALSE;
// Generate VBO for previous position
glGenBuffers( 1, &m_hPrevPositionVBO );
glBindBuffer( GL_ARRAY_BUFFER, m_hPrevPositionVBO );
glBufferData( GL_ARRAY_BUFFER, m_uiNumVerts * 4 * sizeof(float), NULL, GL_DYNAMIC_DRAW );
glBufferSubData( GL_ARRAY_BUFFER, 0, m_uiNumVerts * 4 * sizeof(float), m_pVerts );
glBindBuffer( GL_ARRAY_BUFFER, 0 );
if( !InitKernels( context, device ) )
return FALSE;
// Create OpenCL memory objects for the VBOs
UINT32 vbos[NUM_VBOS] =
{
m_hPositionVBO,
m_hPrevPositionVBO,
m_hNormalVBO,
m_hTangentVBO,
m_hBitangentVBO,
m_hTextureVBO,
m_hBackNormalVBO,
m_hBackTangentVBO,
m_hBackBitangentVBO
};
cl_int errNum = 0;
for( INT32 i = 0; i < NUM_VBOS; i++ )
{
m_vboMem[i] = clCreateFromGLBuffer( context, CL_MEM_READ_WRITE, vbos[i], &errNum );
if( errNum != CL_SUCCESS )
{
FrmLogMessage( "Error creating OpenCL memory object from GL VBO." );
return FALSE;
}
}
// Initialize the base distances memory object
if( !InitConstraintsBaseDists( context ) )
return FALSE;
m_vertsCopyMem = clCreateBuffer( context, CL_MEM_READ_WRITE, sizeof(cl_float4) * m_uiNumVerts, NULL, &errNum);
if( errNum != CL_SUCCESS )
{
FrmLogMessage( "Error creating OpenCL memory object." );
return FALSE;
}
return TRUE;
}
ClState::Buffer ClState::createFromGLBuffer(const Context& c, BufferFlag flag, uint32 gl_id){
ensure(gl_id != 0);
cl_int err= 0;
Buffer b;
b.id= clCreateFromGLBuffer( c.id,
flag,
gl_id,
&err);
errorCheck("ClState::createFromGLBuffer(..): clCreateFromGLBuffer failed: ", err);
return (b);
}
void InitCL()
{
int i;
cl_platform_id myplatform;
cl_device_id *mydevice;
cl_int err;
char* oclsource;
size_t program_length;
unsigned int gpudevcount;
err = RGUGetPlatformID(&myplatform);
err = clGetDeviceIDs(myplatform,CL_DEVICE_TYPE_GPU,0,NULL,&gpudevcount);
mydevice = new cl_device_id[gpudevcount];
err = clGetDeviceIDs(myplatform,CL_DEVICE_TYPE_GPU,gpudevcount,mydevice,NULL);
// You need all these to get full interoperability with OpenGL:
cl_context_properties props[] = {
CL_GL_CONTEXT_KHR, (cl_context_properties)glXGetCurrentContext(),
CL_GLX_DISPLAY_KHR, (cl_context_properties)glXGetCurrentDisplay(),
CL_CONTEXT_PLATFORM, (cl_context_properties)myplatform,
0};
mycontext = clCreateContext(props,1,&mydevice[0],NULL,NULL,&err);
mycommandqueue = clCreateCommandQueue(mycontext,mydevice[0],0,&err);
oclsource = RGULoadProgSource("particle_system.cl", "", &program_length);
myprogram = clCreateProgramWithSource(mycontext,1,(const char **)&oclsource,
&program_length, &err);
if(err==CL_SUCCESS) fprintf(stderr,"create ok\n");
else fprintf(stderr,"create err %d\n",err);
clBuildProgram(myprogram, 0, NULL, NULL, NULL, NULL);
mykernel = clCreateKernel(myprogram, "VVerlet", &err);
if(err==CL_SUCCESS) fprintf(stderr,"build ok\n");
else fprintf(stderr,"build err %d\n",err);
glBindBuffer(GL_ARRAY_BUFFER, OGL_VBO);
glBufferData(GL_ARRAY_BUFFER, DATA_SIZE, &host_position[0][0], GL_DYNAMIC_DRAW);
oclvbo = clCreateFromGLBuffer(mycontext,CL_MEM_WRITE_ONLY,OGL_VBO,&err);
dev_velocity = clCreateBuffer(mycontext,CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,
DATA_SIZE,&host_velocity[0][0],&err);
dev_rseed = clCreateBuffer(mycontext,CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,
NUMBER_OF_PARTICLES*sizeof(float),&host_rseed[0],&err);
oclcbo = clCreateBuffer(mycontext,CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,
DATA_SIZE,&host_color[0],&err);
clSetKernelArg(mykernel,0,sizeof(cl_mem),(void *)&oclvbo);
clSetKernelArg(mykernel,1,sizeof(cl_mem),(void *)&dev_velocity);
clSetKernelArg(mykernel,2,sizeof(cl_mem),(void *)&dev_rseed);
clSetKernelArg(mykernel,3,sizeof(cl_mem),(void *)&oclcbo);
}
/// Tells LibOI that the image source is located in OpenGL device memory at the location
/// specified. You must also indicate whether the OpenGL location is a
/// OPENGL_FRAMEBUFFER | OPENGL_TEXTUREBUFFER
/// All subsequent CopyImageToBuffer commands will read from this location.
void CLibOI::SetImageSource(GLuint gl_device_memory, LibOIEnums::ImageTypes type)
{
mImageType = type;
int status = CL_SUCCESS;
unsigned int CLVersion = mOCL->GetOpenCLVersion();
switch(type)
{
case LibOIEnums::OPENGL_FRAMEBUFFER:
mImage_gl = clCreateFromGLBuffer(mOCL->GetContext(), CL_MEM_READ_ONLY, gl_device_memory, &status);
CHECK_OPENCL_ERROR(status, "clCreateFromGLBuffer failed.");
break;
case LibOIEnums::OPENGL_TEXTUREBUFFER:
// OpenCL's support for texture buffers changed between v1.0 and v1.2,
// but we might be linking with an older library, hence we use both
// runtime version detection and compile-time detection
if(CLVersion == 100 || CLVersion == 110 || CLVersion == 000)
{
mImage_gl = clCreateFromGLTexture2D(mOCL->GetContext(), CL_MEM_READ_ONLY, GL_TEXTURE_2D, 0, gl_device_memory, &status);
}
#if MAX_OPENCL_VERSION >= 120
//#if MAX_OPENCL_VERSION_MAJOR >= 1 && MAX_OPENCL_VERSION_MAJOR >= 2
else if(CLVersion == 120)
{
mImage_gl = clCreateFromGLTexture(mOCL->GetContext(), CL_MEM_READ_ONLY, GL_TEXTURE_2D, 0, gl_device_memory, &status);
}
#endif // MAX_OPENCL_VERSION >= 1.2
else // mImage_gl == NULL
{
throw runtime_error("No function for accessing OpenGL memory is defined for this OpenCL context!");
}
CHECK_OPENCL_ERROR(status, "clCreateFromGLTexture failed.");
break;
case LibOIEnums::OPENGL_RENDERBUFFER:
// TODO: note that the clCreateFromGLTexture2D was depreciated in the OpenCL 1.2 specifications.
mImage_gl = clCreateFromGLRenderbuffer(mOCL->GetContext(), CL_MEM_READ_ONLY, gl_device_memory, &status);
CHECK_OPENCL_ERROR(status, "clCreateFromGLRenderbuffer failed.");
break;
default:
// We don't know what type of image this is!
assert(false);
break;
}
}
bool CreateMemObjects()
{
cl_int errNum;
cl_vbo_mem = clCreateFromGLBuffer(context, CL_MEM_READ_WRITE, vbo, &errNum );
if(errNum != CL_SUCCESS)
{
std::cerr<< "Failed creating memory from GL buffer." << std::endl;
return false;
}
return true;
}
static buffer from_gl_buffer(const context &context,
GLuint bufobj,
cl_mem_flags flags = read_write)
{
cl_int error = 0;
cl_mem mem = clCreateFromGLBuffer(context, flags, bufobj, &error);
if(!mem){
BOOST_THROW_EXCEPTION(runtime_exception(error));
}
buffer buf(mem);
clReleaseMemObject(mem);
return buf;
}
bool
CLGLVertexBuffer::allocate(cl_context clContext) {
assert(clContext);
// create GL buffer first
int size = _numElements * _numVertices * sizeof(float);
glGenBuffers(1, &_vbo);
#if defined(GL_EXT_direct_state_access)
if (glNamedBufferDataEXT) {
glNamedBufferDataEXT(_vbo, size, 0, GL_DYNAMIC_DRAW);
} else {
#else
{
#endif
GLint prev = 0;
glGetIntegerv(GL_ARRAY_BUFFER_BINDING, &prev);
glBindBuffer(GL_ARRAY_BUFFER, _vbo);
glBufferData(GL_ARRAY_BUFFER, size, 0, GL_DYNAMIC_DRAW);
glBindBuffer(GL_ARRAY_BUFFER, prev);
}
// register vbo as cl memory
cl_int err;
_clMemory = clCreateFromGLBuffer(clContext,
CL_MEM_READ_WRITE, _vbo, &err);
if (err != CL_SUCCESS) return false;
return true;
}
void
CLGLVertexBuffer::map(cl_command_queue queue) {
if (_clMapped) return; // XXX: what if another queue is given?
_clQueue = queue;
clEnqueueAcquireGLObjects(queue, 1, &_clMemory, 0, 0, 0);
_clMapped = true;
}
void
CLGLVertexBuffer::unmap() {
if (! _clMapped) return;
clEnqueueReleaseGLObjects(_clQueue, 1, &_clMemory, 0, 0, 0);
_clMapped = false;
}
} // end namespace Osd
// Initialize GL
//*****************************************************************************
void initPixelBuffer()
{
ciErrNum = CL_SUCCESS;
if (pbo) {
// delete old buffer
clReleaseMemObject(pbo_cl);
glDeleteBuffersARB(1, &pbo);
}
// create pixel buffer object for display
glGenBuffersARB(1, &pbo);
glBindBufferARB(GL_PIXEL_UNPACK_BUFFER_ARB, pbo);
glBufferDataARB(GL_PIXEL_UNPACK_BUFFER_ARB, width * height * sizeof(GLubyte) * 4, 0, GL_STREAM_DRAW_ARB);
glBindBufferARB(GL_PIXEL_UNPACK_BUFFER_ARB, 0);
if( g_glInterop )
{
// create OpenCL buffer from GL PBO
pbo_cl = clCreateFromGLBuffer(cxGPUContext,CL_MEM_WRITE_ONLY, pbo, &ciErrNum);
//oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
printf("Create from GL buffer error %i\n",ciErrNum);
}
else
{
pbo_cl = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY, width * height * sizeof(GLubyte) * 4, NULL, &ciErrNum);
}
// calculate new grid size
//gridSize[0] = floor(width/LOCAL_SIZE_X);
//gridSize[1] = floor(height/LOCAL_SIZE_Y);
gridSize[0] = width;
gridSize[1] = height;
int xBlocks = (size_t)ceil((float)width / (float)LOCAL_SIZE_X);
int yBlocks = (size_t)ceil((float)height / (float)LOCAL_SIZE_Y);
// Calculate total number of threads (this is done to guarantee that total number of threads is multiple of local work size, required by OpenCL)
gridSize[0] = xBlocks * LOCAL_SIZE_X;
gridSize[1] = yBlocks * LOCAL_SIZE_Y;
ciErrNum |= clSetKernelArg(ckKernel, 0, sizeof(cl_mem), (void *) &pbo_cl);
ciErrNum |= clSetKernelArg(ckKernel, 1, sizeof(unsigned int), &width);
ciErrNum |= clSetKernelArg(ckKernel, 2, sizeof(unsigned int), &height);
printf("Error first kernel args is %i\n",ciErrNum);
//oclCheckErrorEX(ciErrNum, CL_SUCCESS, pCleanup);
}
void OpenCLBuffer::initFromGLObject(GLuint glBufferObject,
cl_mem_flags memFlags)
{
ofLog(OF_LOG_VERBOSE, "OpenCLBuffer::initFromGLObject");
init();
cl_int err;
clMemObject= clCreateFromGLBuffer(pOpenCL->getContext(), memFlags, glBufferObject, &err);
assert(err != CL_INVALID_CONTEXT);
assert(err != CL_INVALID_VALUE);
assert(err != CL_INVALID_GL_OBJECT);
assert(err != CL_OUT_OF_HOST_MEMORY);
assert(err == CL_SUCCESS);
assert(clMemObject);
hasCorrespondingGLObject = true;
}
static VALUE
rcl_mem_create_from_gl_buffer(VALUE self, VALUE context,
VALUE flags, VALUE bufobj)
{
EXPECT_RCL_TYPE(context, Context);
EXPECT_FIXNUM(flags);
EXPECT_FIXNUM(bufobj);
cl_context cxt = ContextPtr(context);
cl_mem_flags mf = FIX2INT(flags);
cl_GLuint glbuf = FIX2UINT(bufobj);
cl_int res;
cl_mem mem = clCreateFromGLBuffer(cxt, mf, glbuf, &res);
CHECK_AND_RAISE(res);
return RMemory(mem);
}
Particles::Particles(size_t number)
: _number_particles(number) {
cl_int err;
glGenVertexArrays(1, &this->_vao);
this->allocate_buffer(number);
// init openCL things
this->_init_cl_context();
this->_init_particles_sphere_program = cl::Program::Program(
program_from_kernel(this->_cl_context(), "src/kernels/init_particles_sphere.cl"));
if (this->_init_particles_sphere_program() == 0)
{
dprintf(2, "Error while trying to load program...\n");
throw std::exception();
}
this->_init_particles_sphere_kernel = cl::Kernel::Kernel(
this->_init_particles_sphere_program, "init_particles_sphere", &err);
if (err != CL_SUCCESS)
{
dprintf(2, "Error while trying to load kernel...\n");
throw std::exception();
}
this->_positions_cl = clCreateFromGLBuffer(this->_cl_context(),
CL_MEM_READ_WRITE, this->_positions_vbo, &err);
if (err != CL_SUCCESS)
{
dprintf(2, "Cannot create buffer (status: %d) from opengl\n", err);
throw std::exception();
}
// openGL
this->_gl_program = program_from_shaders("src/shaders/particles.vert",
"src/shaders/particles.frag");
if (this->_gl_program == 0)
{
dprintf(2, "Error while trying to load shaders...\n");
throw std::exception();
}
}
/*!
Creates an OpenCL memory buffer from the OpenGL buffer object
\a bufobj, with the specified \a access mode.
This function will only work if supportsObjectSharing() is true.
*/
QCLBuffer QCLContextGL::createGLBuffer(GLuint bufobj, QCLMemoryObject::Access access)
{
#ifndef QT_NO_CL_OPENGL
cl_int error = CL_INVALID_CONTEXT;
cl_mem_flags flags = cl_mem_flags(access);
cl_mem mem = clCreateFromGLBuffer
(contextId(), flags, bufobj, &error);
reportError("QCLContextGL::createGLBuffer:", error);
if (mem)
return QCLBuffer(this, mem);
else
return QCLBuffer();
#else
Q_UNUSED(bufobj);
Q_UNUSED(access);
reportError("QCLContextGL::createGLBuffer:", CL_INVALID_VALUE);
return QCLBuffer();
#endif
}
// Handles a screen window resize, re-allocating affected buffers etc.
int resizeInterop(struct Interop* interop, int width, int height)
{
// Bind the OpenGL buffers
size_t bufferSize = (size_t)width * (size_t)height * 4 * sizeof(float);
glViewport(0, 0, width, height);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, interop->glBuffer);
glBufferData(GL_PIXEL_UNPACK_BUFFER, (GLsizeiptr)bufferSize, NULL, GL_DYNAMIC_COPY);
if (PrintErr((int)glGetError())) return -1;
freeMem(interop, 0);
cl_int openclError = 0;
cl_mem openclBuffer = clCreateFromGLBuffer(interop->context, CL_MEM_READ_WRITE,
interop->glBuffer, &openclError);
if (PrintErr(openclError))
return -1;
if (PrintErr(clEnqueueAcquireGLObjects(interop->command_queue, 1,
&openclBuffer, 0, NULL, NULL)))
return -1;
if (PrintErr(addMem(interop, 0, openclBuffer, 0)))
return -1;
struct cl_mem_list* memList = interop->clMems;
while (memList)
{
if (memList->key != 0 && memList->memorySize == (size_t)0)
{
clReleaseMemObject(memList->memory);
memList->memory = clCreateBuffer(interop->context, CL_MEM_READ_WRITE, bufferSize,
NULL, &openclError);
if (PrintErr(openclError && "clCreateBuffer()"))
return -1;
}
memList = memList->next;
}
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
glBindTexture(GL_TEXTURE_2D, interop->glTexture);
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, width, height, 0, GL_RGBA, GL_FLOAT, NULL);
if (PrintErr((int)glGetError())) return -1;
return 0;
}
///
// Create memory objects used as the arguments to kernels in OpenCL
// The memory objects are created from existing OpenGL buffers and textures
//
bool CreateMemObjects(cl_context context, GLuint texture, GLuint vbo, cl_mem *p_cl_vbo_mem, cl_mem *p_cl_tex_mem)
{
cl_int errNum;
*p_cl_vbo_mem = clCreateFromGLBuffer(context, CL_MEM_READ_WRITE, vbo, &errNum);
if (errNum != CL_SUCCESS)
{
std::cerr << "Failed creating memory from GL buffer." << std::endl;
return false;
}
*p_cl_tex_mem = clCreateFromGLTexture2D(context, CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, texture, &errNum);//GL_TEXTURE_RECTANGLE
if (errNum != CL_SUCCESS)
{
std::cerr << "Failed creating memory from GL texture." << std::endl;
return false;
}
return true;
}
static int ocl_register_buf(QSP_ARG_DECL Data_Obj *dp)
{
if( opengl_prohibited )
error1("ocl_register_buf: Need to specify GL window BEFORE initializing OpenCL!?");
#ifdef HAVE_OPENGL
cl_mem img;
cl_int status;
// Texture2D deprecated on Apple
//fprintf(stderr,"obj %s has texture id %d\n",OBJ_NAME(dp),OBJ_TEX_ID(dp));
//fprintf(stderr,"obj %s has platform device %s\n",OBJ_NAME(dp),PFDEV_NAME(OBJ_PFDEV(dp)));
//advise("ocl_register_buf calling clCreateFromGLBuffer");
//longlist(QSP_ARG dp);
// Used to call clCreateFromGLTexture, but this works:
img = clCreateFromGLBuffer(
OCLDEV_CTX( OBJ_PFDEV(dp) ), // OCL context
CL_MEM_READ_WRITE, // flags
OBJ_TEX_ID(dp), // from glBufferData?
&status);
if( status != CL_SUCCESS ){
report_ocl_error(status, "clCreateFromGLTexture");
return -1;
} else {
SET_OBJ_DATA_PTR(dp,img);
}
// dp is a special buffer object...
//cl_mem memobj;
//cl_mem = clCreate
return 0;
#else // ! HAVE_OPENGL
warn("ocl_register_buf: Sorry, no OpenGL support in this build!?");
return -1;
#endif // ! HAVE_OPENGL
}
void BufferObject::create(
const kvs::cl::Context& context,
const kvs::glew::VertexBufferObject& vbo )
{
if ( m_is_created )
{
kvsMessageError( "Buffer is already created." );
return;
}
cl_int result = CL_SUCCESS;
const cl_mem_flags flags = static_cast<cl_mem_flags>( m_access_type );
m_memory = clCreateFromGLBuffer(
context.context(), flags, vbo.id(), &result );
if ( result != CL_SUCCESS )
{
kvsMessageError( "OpenCL; %s.", kvs::cl::ErrorString( result ) );
return;
}
m_is_created = true;
m_size = 0;
}
/// Tells LibOI that the image source is located in OpenGL device memory at the location
/// specified. You must also indicate whether the OpenGL location is a
/// OPENGL_FRAMEBUFFER | OPENGL_TEXTUREBUFFER
/// All subsequent CopyImageToBuffer commands will read from this location.
void CLibOI::SetImageSource(GLuint gl_device_memory, LibOIEnums::ImageTypes type)
{
mImageType = type;
int status = CL_SUCCESS;
switch(type)
{
case LibOIEnums::OPENGL_FRAMEBUFFER:
mImage_gl = clCreateFromGLBuffer(mOCL->GetContext(), CL_MEM_READ_ONLY, gl_device_memory, &status);
CHECK_OPENCL_ERROR(status, "clCreateFromGLBuffer failed.");
break;
case LibOIEnums::OPENGL_TEXTUREBUFFER:
#if defined(DETECTED_OPENCL_1_0) || defined(DETECTED_OPENCL_1_1) || defined(DETECTED_OPENCL_UNKNOWN_VERSION)
mImage_gl = clCreateFromGLTexture3D(mOCL->GetContext(), CL_MEM_READ_ONLY, GL_TEXTURE_3D, 0, gl_device_memory, &status);
#else
mImage_gl = clCreateFromGLTexture(mOCL->GetContext(), CL_MEM_READ_ONLY, GL_TEXTURE_2D_ARRAY, 0, gl_device_memory, &status);
#endif // defined(DETECTED_OPENCL_1_0) || defined(DETECTED_OPENCL_1_1)
CHECK_OPENCL_ERROR(status, "clCreateFromGLTexture failed.");
break;
case LibOIEnums::OPENGL_RENDERBUFFER:
// TODO: note that the clCreateFromGLTexture2D was depreciated in the OpenCL 1.2 specifications.
mImage_gl = clCreateFromGLRenderbuffer(mOCL->GetContext(), CL_MEM_READ_ONLY, gl_device_memory, &status);
CHECK_OPENCL_ERROR(status, "clCreateFromGLRenderbuffer failed.");
break;
default:
// We don't know what type of image this is!
assert(false);
break;
}
}
void GLWidget::initializeCL()
{
qDebug() << "Initializing OpenCL";
if (!setupOpenCLGL(clContext, clQueue, clDevice)) {
qDebug() << "OpenCL initialization error";
return;
}
cl_int error;
loadKernel(clContext, &clKernel, clDevice, "../src/vboproc.cl", "vboproc");
// Creo OpenCL buffer a partir del OpenGL buffer
qDebug() << "Creando OpenCL buffer.";
clvbo = clCreateFromGLBuffer(clContext, CL_MEM_READ_WRITE, particlesVBO->bufferId(), &error);
if (checkError(error, "clCreateFromGLBuffer")) {
qDebug() << "OpenCL initialization error";
return;
}
// Setean los parametros del kernel, y luego se encola su ejecucion
qDebug() << "Seteo los parametros del kernel.";
error = clSetKernelArg(clKernel, 0, sizeof(cl_mem), (void*)&clvbo);
error |= clSetKernelArg(clKernel, 1, sizeof(cl_int), (void*)&vertexNumber);
const float cubeLims[]= {cubeLimits.x(), cubeLimits.y(), cubeLimits.z()};
error |= clSetKernelArg(clKernel, 2, sizeof(cl_float3), (void*)&cubeLims);
error |= clSetKernelArg(clKernel, 3, sizeof(cl_float), (void*)×tep);
if(checkError(error, "clSetKernelArg")) {
qDebug() << "OpenCL initialization error";
return;
}
// Una vez creado el kernel, decremento la referencia al programa creado
qDebug() << "OpenCL initialized successfully";
}
void configure_shared_data() {
int err;
/* Create and configure pixel buffer */
glGenBuffers(1, &pbo);
glBindBuffer(GL_ARRAY_BUFFER, pbo);
glBufferData(GL_ARRAY_BUFFER, width*height*sizeof(char),
NULL, GL_STATIC_DRAW);
glBindBuffer(GL_ARRAY_BUFFER, 0);
out_buffer = clCreateFromGLBuffer(context, CL_MEM_WRITE_ONLY,
pbo, &err);
if(err < 0) {
perror("Couldn't create a buffer object from the PBO");
exit(1);
}
err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &out_buffer);
if(err < 0) {
printf("Couldn't set a kernel argument");
exit(1);
};
}
JNIEXPORT jlong JNICALL Java_org_lwjgl_opencl_CL10GL_nclCreateFromGLBuffer(JNIEnv *env, jclass clazz, jlong context, jlong flags, jint bufobj, jlong errcode_ret, jlong function_pointer) {
cl_int *errcode_ret_address = (cl_int *)(intptr_t)errcode_ret;
clCreateFromGLBufferPROC clCreateFromGLBuffer = (clCreateFromGLBufferPROC)((intptr_t)function_pointer);
cl_mem __result = clCreateFromGLBuffer((cl_context)(intptr_t)context, flags, bufobj, errcode_ret_address);
return (intptr_t)__result;
}
game::WaterLogicComponent::WaterLogicComponent(game::ActorWPtr actorWPtr) {
auto &&model = std::dynamic_pointer_cast<WaterModelComponent>(actorWPtr.lock()->getComponent(ComponentType::MODEL_COMPONENT));
auto &openCLSystem = gamesystem::OpenCLSystem::getInstance();
if (!openCLSystem.TryLoadKernel("Kernel/test.cl", "hash_particles", _hash_particles_kernel)) {
LOG(ERROR) << "Failed to load hash_particles kernel";
event::EventManager::getInstance().TriggerEvent(std::make_shared<event::OnWindowClose>());
}
if (!openCLSystem.TryLoadKernel("Kernel/test.cl", "sort_post_pass", _sort_post_pass_kernel)) {
LOG(ERROR) << "Failed to load sort_post_pass kernel";
event::EventManager::getInstance().TriggerEvent(std::make_shared<event::OnWindowClose>());
}
if (!openCLSystem.TryLoadKernel("Kernel/test.cl", "indexx", _index_kernel)) {
LOG(ERROR) << "Failed to load indexx kernel";
event::EventManager::getInstance().TriggerEvent(std::make_shared<event::OnWindowClose>());
}
if (!openCLSystem.TryLoadKernel("Kernel/test.cl", "index_post_pass", _index_post_pass_kernel)) {
LOG(ERROR) << "Failed to load index_post_pass kernel";
event::EventManager::getInstance().TriggerEvent(std::make_shared<event::OnWindowClose>());
}
if (!openCLSystem.TryLoadKernel("Kernel/test.cl", "neighbour_map", _neighbour_map_kernel)) {
LOG(ERROR) << "Failed to load neighbour_map kernel";
event::EventManager::getInstance().TriggerEvent(std::make_shared<event::OnWindowClose>());
}
if (!openCLSystem.TryLoadKernel("Kernel/test.cl", "find_voxel_neighbours", _find_voxel_neighbours_kernel)) {
LOG(ERROR) << "Failed to load find_voxel_neighbours kernel";
event::EventManager::getInstance().TriggerEvent(std::make_shared<event::OnWindowClose>());
}
if (!openCLSystem.TryLoadKernel("Kernel/test.cl", "compute_density_pressure", _compute_density_pressure_kernel)) {
LOG(ERROR) << "Failed to load compute_density_pressure kernel";
event::EventManager::getInstance().TriggerEvent(std::make_shared<event::OnWindowClose>());
}
if (!openCLSystem.TryLoadKernel("Kernel/test.cl", "compute_acceleration", _compute_acceleration_kernel)) {
LOG(ERROR) << "Failed to load compute_acceleration kernel";
event::EventManager::getInstance().TriggerEvent(std::make_shared<event::OnWindowClose>());
}
if (!openCLSystem.TryLoadKernel("Kernel/test.cl", "integrate", _integrate_kernel)) {
LOG(ERROR) << "Failed to load integrate kernel";
event::EventManager::getInstance().TriggerEvent(std::make_shared<event::OnWindowClose>());
}
auto context = openCLSystem.getContext();
cl_int errNum = 0;
_particle_count = model->particle_count();
auto& config = gamesystem::ConfigSystem::getInstance();
cl_float4 lbf{0,0,0,0};
cl_float4 rtb{2,2,2,0};
cl_float h = ::atof(config["h"].get().c_str());
cl_float ro0 = ::atof(config["ro0"].get().c_str());
cl_float m = 0.02;
cl_float4 g{0,-10,0,0};
cl_float mi{8};
cl_float k{2};
cl_float speed_loss{1};
unsigned int voxelsX = (unsigned int) (fabs((float)((rtb.s[0] - lbf.s[0])/(2*h))) + 0.5);
unsigned int voxelsY = (unsigned int) (fabs((float)((rtb.s[1] - lbf.s[1])/(2*h))) + 0.5);
unsigned int voxelsZ = (unsigned int) (fabs((float)((rtb.s[2] - lbf.s[2])/(2*h))) + 0.5);
unsigned voxelCount = voxelsX * voxelsY * voxelsZ;
clppProgram::setBasePath("Kernel");
cont.clContext = openCLSystem.getContext();
cont.clDevice = openCLSystem.getDevice();
cont.clPlatform = openCLSystem.getPlatform();
cont.clQueue = openCLSystem.getCommandQueue();
_positions.resize(_particle_count);
_position_cl = clCreateFromGLBuffer(context, CL_MEM_READ_WRITE, model->position_vbo, &errNum);
_velocities.resize(_particle_count);
_velocity_cl = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_float4) * _velocities.size(), &_velocities[0], &errNum);
_accelerations.resize(_particle_count);
_acceleration_cl = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_float4) * _accelerations.size(), &_accelerations[0], &errNum);
_sorted_positions.resize(_particle_count);
_sorted_position_cl = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_float4) * _sorted_positions.size(), &_sorted_positions[0], &errNum);
_sorted_velocities.resize(_particle_count);
_sorted_velocity_cl = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_float4) * _sorted_velocities.size(), &_sorted_velocities[0], &errNum);
_grid_voxel_indices.resize(voxelCount + 1);
_grid_voxel_indices.back() = _particle_count;
_grid_voxel_index_cl = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_int) * _grid_voxel_indices.size(), &_grid_voxel_indices[0], &errNum);
_neighbour_map.resize(neighbour_count*_particle_count);
_neighbour_map_cl = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_int) * _neighbour_map.size(), &_neighbour_map[0], &errNum);
_voxel_positions.resize(_particle_count);
_voxel_positions_cl = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_int2) * _voxel_positions.size(), &_voxel_positions[0], &errNum);
_voxel_neighbours.resize(64*voxelCount);
_voxel_neighbours_cl = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_int) * _voxel_neighbours.size(), &_voxel_neighbours[0], &errNum);
_density_pressure.resize(_particle_count);
_density_pressure_cl = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(cl_float2) * _density_pressure.size(), &_density_pressure[0], &errNum);
sort = clpp::createBestSortKV(&cont, _voxel_positions.size(), false);
sort->pushCLDatas(_voxel_positions_cl, _voxel_positions.size());
errNum = clSetKernelArg(_hash_particles_kernel, 0, sizeof(cl_mem), &_position_cl);
errNum = clSetKernelArg(_hash_particles_kernel, 1, sizeof(cl_mem), &_voxel_positions_cl);
errNum = clSetKernelArg(_hash_particles_kernel, 2, sizeof(cl_float4), &lbf);
errNum = clSetKernelArg(_hash_particles_kernel, 3, sizeof(cl_float4), &rtb);
errNum = clSetKernelArg(_hash_particles_kernel, 4, sizeof(cl_float), &h);
errNum = clSetKernelArg(_sort_post_pass_kernel, 0, sizeof(cl_mem), &_position_cl);
errNum = clSetKernelArg(_sort_post_pass_kernel, 1, sizeof(cl_mem), &_velocity_cl);
errNum = clSetKernelArg(_sort_post_pass_kernel, 2, sizeof(cl_mem), &_sorted_position_cl);
errNum = clSetKernelArg(_sort_post_pass_kernel, 3, sizeof(cl_mem), &_sorted_velocity_cl);
errNum = clSetKernelArg(_sort_post_pass_kernel, 4, sizeof(cl_mem), &_voxel_positions_cl);
errNum = clSetKernelArg(_index_kernel, 0, sizeof(cl_mem), &_grid_voxel_index_cl);
errNum = clSetKernelArg(_index_kernel, 1, sizeof(cl_mem), &_voxel_positions_cl);
errNum = clSetKernelArg(_index_kernel, 2, sizeof(cl_uint), &_particle_count);
errNum = clSetKernelArg(_index_post_pass_kernel, 0, sizeof(cl_mem), &_grid_voxel_index_cl);
/*
__global int* neighbourMap,
__global int2* voxelParticle,
__global int* gridVoxelIndex,
__global float4* sortedPositions,
__global int* randomInts, //should be particle_count elements
__global int* voxelNeighbourMap,
int neighboursToFind,
float4 lbf,
float4 rtb,
float h
*/
errNum = clSetKernelArg(_neighbour_map_kernel, 0, sizeof(cl_mem), &_neighbour_map_cl);
errNum = clSetKernelArg(_neighbour_map_kernel, 1, sizeof(cl_mem), &_voxel_positions_cl);
errNum = clSetKernelArg(_neighbour_map_kernel, 2, sizeof(cl_mem), &_grid_voxel_index_cl);
errNum = clSetKernelArg(_neighbour_map_kernel, 3, sizeof(cl_mem), &_sorted_position_cl);
errNum = clSetKernelArg(_neighbour_map_kernel, 4, sizeof(cl_mem), &_voxel_neighbours_cl);
errNum = clSetKernelArg(_neighbour_map_kernel, 5, sizeof(cl_int), &neighbour_count);
errNum = clSetKernelArg(_neighbour_map_kernel, 6, sizeof(cl_float4), &lbf);
errNum = clSetKernelArg(_neighbour_map_kernel, 7, sizeof(cl_float4), &rtb);
errNum = clSetKernelArg(_neighbour_map_kernel, 8, sizeof(cl_float), &h);
errNum = clSetKernelArg(_find_voxel_neighbours_kernel, 0, sizeof(cl_mem), &_voxel_neighbours_cl);
errNum = clSetKernelArg(_find_voxel_neighbours_kernel, 1, sizeof(cl_float4), &lbf);
errNum = clSetKernelArg(_find_voxel_neighbours_kernel, 2, sizeof(cl_float4), &rtb);
errNum = clSetKernelArg(_find_voxel_neighbours_kernel, 3, sizeof(cl_float), &h);
auto commandQueue = openCLSystem.getCommandQueue();
size_t voxelWorkSize = _grid_voxel_indices.size() - 1;
errNum = clEnqueueNDRangeKernel(commandQueue, _find_voxel_neighbours_kernel, 1, NULL, &voxelWorkSize, NULL, 0,0,0 );
clFinish(commandQueue);
/* errNum = clEnqueueReadBuffer(commandQueue, _voxel_neighbours_cl, CL_TRUE, 0, _voxel_neighbours.size() * sizeof(cl_int), &_voxel_neighbours[0], 0, 0, NULL);
for(int i=0;i<_voxel_neighbours.size()/64;i++) {
cout<<i<<endl;
for (int j = 0; j < 8; j++) {
cout<<_voxel_neighbours[i*64+j*8] << " " << _voxel_neighbours[i*64+j*8+1] << " " << _voxel_neighbours[i*64+j*8+2] << " " << _voxel_neighbours[i*64+j*8+3] << " "
<<_voxel_neighbours[i*64+j*8+4] << " " << _voxel_neighbours[i*64+j*8+5] << " " << _voxel_neighbours[i*64+j*8+6] << " " << _voxel_neighbours[i*64+j*8+7] << std::endl;
}
}*/
errNum = clSetKernelArg(_compute_density_pressure_kernel, 0, sizeof(cl_mem), &_sorted_position_cl);
errNum = clSetKernelArg(_compute_density_pressure_kernel, 1, sizeof(cl_mem), &_neighbour_map_cl);
errNum = clSetKernelArg(_compute_density_pressure_kernel, 2, sizeof(cl_mem), &_density_pressure_cl);
errNum = clSetKernelArg(_compute_density_pressure_kernel, 3, sizeof(cl_int), &neighbour_count);
errNum = clSetKernelArg(_compute_density_pressure_kernel, 4, sizeof(cl_float), &m);
errNum = clSetKernelArg(_compute_density_pressure_kernel, 5, sizeof(cl_float), &h);
errNum = clSetKernelArg(_compute_density_pressure_kernel, 6, sizeof(cl_float), &k);
errNum = clSetKernelArg(_compute_density_pressure_kernel, 7, sizeof(cl_float), &ro0);
errNum = clSetKernelArg(_compute_acceleration_kernel, 0, sizeof(cl_mem), &_sorted_position_cl);
errNum = clSetKernelArg(_compute_acceleration_kernel, 1, sizeof(cl_mem), &_sorted_velocity_cl);
errNum = clSetKernelArg(_compute_acceleration_kernel, 2, sizeof(cl_mem), &_acceleration_cl);
errNum = clSetKernelArg(_compute_acceleration_kernel, 3, sizeof(cl_mem), &_density_pressure_cl);
errNum = clSetKernelArg(_compute_acceleration_kernel, 4, sizeof(cl_mem), &_neighbour_map_cl);
errNum = clSetKernelArg(_compute_acceleration_kernel, 5, sizeof(cl_float4), &g);
errNum = clSetKernelArg(_compute_acceleration_kernel, 6, sizeof(cl_float), &m);
errNum = clSetKernelArg(_compute_acceleration_kernel, 7, sizeof(cl_float), &h);
errNum = clSetKernelArg(_compute_acceleration_kernel, 8, sizeof(cl_float), &mi);
errNum = clSetKernelArg(_compute_acceleration_kernel, 9, sizeof(cl_int), &neighbour_count);
errNum = clSetKernelArg(_integrate_kernel, 0, sizeof(cl_mem), &_position_cl);
errNum = clSetKernelArg(_integrate_kernel, 1, sizeof(cl_mem), &_velocity_cl);
errNum = clSetKernelArg(_integrate_kernel, 2, sizeof(cl_mem), &_acceleration_cl);
errNum = clSetKernelArg(_integrate_kernel, 3, sizeof(cl_mem), &_voxel_positions_cl);
errNum = clSetKernelArg(_integrate_kernel, 4, sizeof(cl_float4), &lbf);
errNum = clSetKernelArg(_integrate_kernel, 5, sizeof(cl_float4), &rtb);
errNum = clSetKernelArg(_integrate_kernel, 6, sizeof(cl_float), &speed_loss);
}
