JNIEXPORT jlong JNICALL Java_org_lwjgl_opencl_CL10GL_nclCreateFromGLTexture3D(JNIEnv *__env, jclass clazz, jlong contextAddress, jlong flags, jint texture_target, jint miplevel, jint texture, jlong errcode_retAddress, jlong __functionAddress) {
cl_context context = (cl_context)(intptr_t)contextAddress;
cl_int *errcode_ret = (cl_int *)(intptr_t)errcode_retAddress;
clCreateFromGLTexture3DPROC clCreateFromGLTexture3D = (clCreateFromGLTexture3DPROC)(intptr_t)__functionAddress;
UNUSED_PARAMS(__env, clazz)
return (jlong)(intptr_t)clCreateFromGLTexture3D(context, flags, texture_target, miplevel, texture, errcode_ret);
}
ClState::Buffer ClState::createFromGLTexture(
const Context& c,
BufferFlag flag,
uint32 gl_target,
uint32 gl_id){
ensure(gl_id != 0);
cl_int err= 0;
Buffer b;
// clCreateFromGLTexture*D is deprecated, but OpenCL 1.1 is what we use
if (gl_target == GL_TEXTURE_2D){
b.id= clCreateFromGLTexture2D( c.id,
flag,
gl_target,
0, // miplevel
gl_id,
&err);
}
else if (gl_target == GL_TEXTURE_3D){
b.id= clCreateFromGLTexture3D( c.id,
flag,
gl_target,
0, // miplevel
gl_id,
&err);
}
else {
// OpenCL 1.2 has clCreateFromGLTexture which would allow tex arrays
throw global::Exception(
"ClState::createFromGLTexture(..): unsupported tex target");
}
errorCheck("ClState::createFromGLTexture(..): clCreateFromGLTexture failed: ", err);
return b;
}
sge::opencl::memory_object::image::volume::volume(
context::object &_context,
memory_object::flags_field const &_mem_flags,
renderer::texture::volume &_renderer_texture)
:
impl_(),
image_format_(),
size_(
fcppt::math::dim::structure_cast<
sge::opencl::dim3,
fcppt::cast::size_fun
>(
_renderer_texture.size()
)
)
{
cl_int error_code;
impl_ =
clCreateFromGLTexture3D(
_context.impl(),
sge::opencl::impl::memory_object::to_opencl_mem_flags(
_mem_flags),
dynamic_cast<renderer::opengl::texture::base &>(
_renderer_texture).type().get(),
// mip level
0,
dynamic_cast<renderer::opengl::texture::base &>(
_renderer_texture).id().get(),
&error_code);
opencl::impl::handle_error(
error_code,
FCPPT_TEXT("clCreateFromGLTexture3D()"));
error_code =
clGetImageInfo(
impl_,
CL_IMAGE_FORMAT,
sizeof(cl_image_format),
&image_format_,
nullptr);
opencl::impl::handle_error(
error_code,
FCPPT_TEXT("clGetImageInfo(image format)"));
}
void vglClAlloc(VglImage* img)
{
glFlush();
glFinish();
cl_int err_cl;
if (img->oclPtr == NULL)
{
if (img->ndim == 2)
{
img->oclPtr = clCreateFromGLTexture2D(cl.context, CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, img->tex, &err_cl);
vglClCheckError(err_cl, (char*) "clCreateFromGLTexture");
}
else if(img->ndim == 3)
{
img->oclPtr = clCreateFromGLTexture3D(cl.context, CL_MEM_READ_WRITE, GL_TEXTURE_3D, 0, img->tex, &err_cl);
vglClCheckError(err_cl, (char*) "clCreateFromGLTexture");
}
}
}
/*!
Creates a 3D OpenCL image object from the specified \a mipmapLevel,
OpenGL \a texture object, and \a access mode.
The \a type must be \c{GL_TEXTURE_3D}. The \a texture does not need
to be bound to an OpenGL texture target.
This function will only work if supportsObjectSharing() is true.
\sa createTexture2D()
*/
QCLImage3D QCLContextGL::createTexture3D
(GLenum type, GLuint texture, GLint mipmapLevel, 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 = clCreateFromGLTexture3D
(contextId(), flags, type, mipmapLevel, texture, &error);
reportError("QCLContextGL::createGLTexture3D:", error);
if (mem)
return QCLImage3D(this, mem);
else
return QCLImage3D();
#else
Q_UNUSED(type);
Q_UNUSED(texture);
Q_UNUSED(mipmapLevel);
Q_UNUSED(access);
reportError("QCLContextGL::createGLTexture3D:", CL_INVALID_VALUE);
return QCLImage3D();
#endif
}
/// 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;
}
}
JNIEXPORT jlong JNICALL Java_org_lwjgl_opencl_CL10GL_nclCreateFromGLTexture3D(JNIEnv *env, jclass clazz, jlong context, jlong flags, jint target, jint miplevel, jint texture, jlong errcode_ret, jlong function_pointer) {
cl_int *errcode_ret_address = (cl_int *)(intptr_t)errcode_ret;
clCreateFromGLTexture3DPROC clCreateFromGLTexture3D = (clCreateFromGLTexture3DPROC)((intptr_t)function_pointer);
cl_mem __result = clCreateFromGLTexture3D((cl_context)(intptr_t)context, flags, target, miplevel, texture, errcode_ret_address);
return (intptr_t)__result;
}
void OpenCLExecuter::ocl_filter_shared(void)
{
cl_int err; // debugging variables
size_t szParmDataBytes; // Byte size of context information
cl_mem src_buffer; // OpenCL device source buffer
cl_mem dst_buffer; // OpenCL device source buffer
cl_sampler sampler; // OpenCL sampler
cl_kernel ckKernel; // OpenCL kernel
int iNumElements = volobj->texwidth*volobj->texheight*volobj->texdepth; // Length of float arrays to process
// set Local work size dimensions
// size_t local_threads[3] ={256,256,64};
// set Global work size dimensions
// size_t global_threads[3] ={roundup((int) volobj->texwidth/local_threads[0], 0)*local_threads[0], roundup((int) volobj->texheight/local_threads[1], 0)*local_threads[1], roundup((int) volobj->texdepth/local_threads[2], 0)*local_threads[2]};
// set Global work size dimensions
size_t global_threads[3] ={volobj->texwidth, volobj->texheight, volobj->texdepth};
// allocate the source buffer memory object
src_buffer = clCreateFromGLTexture3D (ocl_wrapper->context, CL_MEM_READ_WRITE, GL_TEXTURE_3D, 0, volobj->TEXTURE3D_RED, &err);
printf("OPENCL: clCreateFromGLTexture3D: %s\n", ocl_wrapper->get_error(err));
// allocate the destination buffer memory object
dst_buffer = clCreateBuffer (ocl_wrapper->context, CL_MEM_READ_WRITE, sizeof(unsigned char) * iNumElements, NULL, &err);
printf("OPENCL: clCreateBuffer: %s\n", ocl_wrapper->get_error(err));
// create a sampler object
sampler = clCreateSampler(ocl_wrapper->context, CL_FALSE, CL_ADDRESS_CLAMP, CL_FILTER_NEAREST, &err);
printf("OPENCL: clCreateSampler: %s\n", ocl_wrapper->get_error(err));
// Create the kernel
ckKernel = clCreateKernel (cpProgram, "myFunc", &err);
printf("OPENCL: clCreateKernel: %s\n", ocl_wrapper->get_error(err));
// Set the Argument values
err = clSetKernelArg (ckKernel, 0, sizeof(cl_mem), (void*)&src_buffer);
printf("OPENCL: clSetKernelArg: %s\n", ocl_wrapper->get_error(err));
err = clSetKernelArg (ckKernel, 1, sizeof(cl_mem), (void*)&dst_buffer);
printf("OPENCL: clSetKernelArg: %s\n", ocl_wrapper->get_error(err));
err = clSetKernelArg (ckKernel, 2, sizeof(sampler), (void*)&sampler);
printf("OPENCL: clSetKernelArg: %s\n", ocl_wrapper->get_error(err));
size_t local;
err = clGetKernelWorkGroupInfo(ckKernel, ocl_wrapper->devices[ocl_wrapper->deviceUsed], CL_KERNEL_LOCAL_MEM_SIZE , sizeof(local), &local, NULL);
printf("OPENCL: clGetKernelWorkGroupInfo (kernel memory): %s\n", ocl_wrapper->get_error(err));
printf("OPENCL: Kernel local memory use: %d Bytes\n", (int)local);
// grab input data from OpenGL, compute, copy the results back to OpenGL
// Runs asynchronous to host, up until blocking clFinish at the end
glFinish();
glFlush();
// grab the OpenGL texture object for read/writing from OpenCL
err = clEnqueueAcquireGLObjects(ocl_wrapper->commandQue, 1, &src_buffer, 0,NULL,NULL);
printf("OPENCL: clEnqueueAcquireGLObjects: %s\n", ocl_wrapper->get_error(err));
// Execute a kernel
err = clEnqueueNDRangeKernel (ocl_wrapper->commandQue, ckKernel, 3, NULL, global_threads, NULL, 0, NULL, NULL);
printf("OPENCL: clEnqueueNDRangeKernel: %s\n", ocl_wrapper->get_error(err));
/*
// Blocking read of results from GPU to Host
int size = volobj->texwidth*volobj->texheight*volobj->texdepth;
unsigned char* result = new unsigned char[size];
err = clEnqueueReadBuffer (ocl_wrapper->commandQue, dst_buffer, CL_TRUE, 0, sizeof(unsigned char) * iNumElements, result, 0, NULL, NULL);
printf("OPENCL: clEnqueueReadBuffer: %s\n", ocl_wrapper->get_error(err));
for(int i=0; i<size; i++) volobj->texture3d[3*i+0] = result[i];
delete[] result;
*/
// copy OpenCL buffer to OpenGl texture
size_t corigin[3] = {0,0,0};
size_t cdimensions[3] = {(unsigned int)volobj->texwidth, (unsigned int)volobj->texheight, (unsigned int)volobj->texdepth};
err = clEnqueueCopyBufferToImage(ocl_wrapper->commandQue , dst_buffer, src_buffer, 0, corigin, cdimensions, 0, NULL, NULL);
printf("OPENCL: clEnqueueCopyBufferToImage: %s\n", ocl_wrapper->get_error(err));
//make sure we block until we are done.
//err = clFinish(ocl_wrapper->commandQue);
//printf("OPENCL: clFinish: %s\n", ocl_wrapper->get_error(err));
//release opengl objects now
err = clEnqueueReleaseGLObjects(ocl_wrapper->commandQue, 1, &src_buffer, 0,0,0);
printf("OPENCL: clEnqueueAcquireGLObjects: %s\n", ocl_wrapper->get_error(err));
// Cleanup allocated objects
printf("OPENCL: Releasing kernel memory\n");
if(ckKernel)clReleaseKernel(ckKernel);
//need to release any other OpenCL memory objects here
if(src_buffer)clReleaseMemObject(src_buffer);
if(dst_buffer)clReleaseMemObject(dst_buffer);
}
