diff options
Diffstat (limited to 'src/opencl_mesh_kit.cpp')
-rw-r--r-- | src/opencl_mesh_kit.cpp | 225 |
1 files changed, 225 insertions, 0 deletions
diff --git a/src/opencl_mesh_kit.cpp b/src/opencl_mesh_kit.cpp new file mode 100644 index 0000000..a3eb6dc --- /dev/null +++ b/src/opencl_mesh_kit.cpp @@ -0,0 +1,225 @@ +#include "opencl_mesh_kit.hpp" + +// TODO : print streamsdk::getOpenCLErrorCodeStr(res) +#define CL_RETURN_VAL_IF_FAIL(val, expr) do { \ + cl_int res=(expr); \ + if ( res != CL_SUCCESS ) { \ + std::cerr << "file " << __FILE__ << ": line " << __LINE__ << " (" << __PRETTY_FUNCTION__ \ + << "): '" << "expr" << "' failed (return code : " << res << ")" << std::endl; \ + return val; \ + } \ +} while(0) + +cl_int OpenCLMeshKit::initCL(intptr_t gl_display, intptr_t gl_context, intptr_t gl_vbo, size_t meshWidth, size_t meshHeight, size_t groupSize) { + cl_uint id, numPlatforms; + cl_int res; + char pbuf[100]; + cl_platform_id *platforms, platform; + bool usableDeviceFound=false; + + this->meshWidth = meshWidth; + this->meshHeight = meshHeight; + this->groupSize = groupSize; + + // Get platform count + CL_RETURN_VAL_IF_FAIL(10, + clGetPlatformIDs(0, NULL, &numPlatforms) + ); + + std::cout << "Detected " << numPlatforms << " platform(s)" << std::endl; + if ( ! ( numPlatforms > 0 ) ) return 2; + + // Allocate room for all platform IDs + platforms = new cl_platform_id[numPlatforms]; + + // Get platform IDs + CL_RETURN_VAL_IF_FAIL(11, + clGetPlatformIDs(numPlatforms, platforms, &numPlatforms) + ); + + // Enumerate platforms and grab informations + for(id=0;id<numPlatforms;id++) { + platform=platforms[id]; + + CL_RETURN_VAL_IF_FAIL(12, + clGetPlatformInfo(platform, CL_PLATFORM_VENDOR, sizeof(pbuf), pbuf, NULL) + ); + std::cout << "Platform " << id << " : " << pbuf << std::endl; + + // Dynamically get the function pointer for clGetGLConetextInfoKHR + clGetGLContextInfoKHR_fn clGetGLContextInfoKHR_proc = (clGetGLContextInfoKHR_fn) clGetExtensionFunctionAddressForPlatform(platform, "clGetGLContextInfoKHR"); + if (!clGetGLContextInfoKHR_proc) { + std::cerr << "clGetExtensionFunctionAddressForPlatform(platform, clGetGLContextInfoKHR) failed" << std::endl; + continue; + } + + // Try to get the device corresponding to the GL context/display on this platform + cl_context_properties cpsGL[] = { + CL_CONTEXT_PLATFORM, (cl_context_properties)platform, + CL_GLX_DISPLAY_KHR, gl_display, + CL_GL_CONTEXT_KHR, gl_context, + 0 + }; + + std::cout << "cl_context_properties cpsGL :" << std::endl; + std::cout << "\tCL_CONTEXT_PLATFORM :" << (void *)cpsGL[1] << std::endl; + std::cout << "\tCL_GLX_DISPLAY_KHR :" << (void *)cpsGL[3] << std::endl; + std::cout << "\tCL_GL_CONTEXT_KHR :" << (void *)cpsGL[5] << std::endl; + + size_t deviceSize=0; + // get deviceSize (should be 1*sizeof(cl_device_id) with CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR) + res=clGetGLContextInfoKHR_proc(cpsGL,CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR,0,NULL,&deviceSize); + if ( res!=CL_SUCCESS || deviceSize!=1*sizeof(cl_device_id)) { + std::cerr << "clGetGLContextInfoKHR_proc(cpsGL,CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR,0,...) failed" << std::endl; + std::cerr << " (return code : " << res << ")" << std::endl; + continue; + } + + cl_dev=0; + res=clGetGLContextInfoKHR_proc(cpsGL,CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR,deviceSize,&cl_dev,NULL); + if ( res!=CL_SUCCESS || cl_dev==0 ) { + std::cerr << "clGetGLContextInfoKHR_proc(cpsGL,CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR," << deviceSize << ",...) failed" << std::endl; + std::cerr << " (return code : " << res << ")" << std::endl; + continue; + } + + std::cout << "cl_device :" << (void *)cl_dev << std::endl; + + cl_ctx = clCreateContext(cpsGL,1,&cl_dev,0,0,&res); + if ( res!=CL_SUCCESS ) { + std::cerr << "clCreateContext() failed" << std::endl; + std::cerr << " (return code : " << res << ")" << std::endl; + continue; + } + + cl_cq = clCreateCommandQueue(cl_ctx,cl_dev,0,&res); + if ( res!=CL_SUCCESS ) { + std::cerr << "clCreateCommandQueue() failed" << std::endl; + std::cerr << " (return code : " << res << ")" << std::endl; + continue; + } + + usableDeviceFound=true; + break; + } + + if (! usableDeviceFound) { + std::cerr << "No OpenCL device has been successfully initialized" << std::endl; + return 13; + } + + cl_vbo = clCreateFromGLBuffer(cl_ctx, CL_MEM_WRITE_ONLY, gl_vbo, &res); + if ( res!=CL_SUCCESS ) { + std::cerr << "clCreateFromGLBuffer(..., gl_vbo, &res) failed" << std::endl; + return 14; + } + + std::cout << "OpenCL initialization done." << std::endl; + return 0; + +} + +cl_int OpenCLMeshKit::compileKernels(std::list<std::string> names, const char source[], size_t sourceLen) { + cl_int res=0; + + const char *p_source=source; + + cl_program program = clCreateProgramWithSource(cl_ctx, 1, &p_source, &sourceLen,&res); + if ( res!=CL_SUCCESS ) { + std::cerr << "Failed to clCreateProgramWithSource(<source of zero_z kernel>)" << std::endl; + return 21; + } + + res = clBuildProgram(program, 1, &cl_dev, "", NULL, NULL); + if ( res!=CL_SUCCESS ) { + std::cerr << "Failed to clBuildProgram()" << std::endl; + return 22; + } + + for (std::list<std::string>::iterator ii = names.begin(); ii != names.end(); ++ii) { + std::string kName = (*ii); + char *kNameZTS = new char[kName.length()+1]; + std::strcpy(kNameZTS, kName.c_str()); + + cl_kernel kernel = clCreateKernel(program,kNameZTS,&res); + delete [] kNameZTS; + + if ( res!=CL_SUCCESS ) { + std::cerr << "Failed to clCreateKernel(program,\"" << kName << "\",&res);" << std::endl; + return 23; + } + kernels[kName]=kernel; + } + + return 0; +} + +cl_int OpenCLMeshKit::execKernel(std::string kernelName) { + /*TODO + +cl_int execKernel(cl_context cl_ctx, cl_command_queue commandQueue, cl_kernel kernel, size_t mesh_width, size_t mesh_height, size_t group_size, cl_mem cl_vbo, float time) { + + cl_int res; + cl_event eventND[1]; + + // Set local and global work group sizes + size_t globalWorkSize[2], localWorkSize[2]; + globalWorkSize[0]=mesh_width; + globalWorkSize[1]=mesh_height; + localWorkSize[0]=group_size; + localWorkSize[1]=1; + + res=clEnqueueAcquireGLObjects(commandQueue, 1, &cl_vbo, 0, 0, NULL); + + res=clSetKernelArg(kernel, 0, sizeof(cl_mem), &cl_vbo); // float4 *pos + res=clSetKernelArg(kernel, 1, sizeof(cl_uint), (void *)&mesh_width); + res=clSetKernelArg(kernel, 2, sizeof(cl_uint), (void *)&mesh_height); + res=clSetKernelArg(kernel, 3, sizeof(float), &time); + + // Execute kernel on given device + res=clEnqueueNDRangeKernel(commandQueue, kernel, 2, NULL, globalWorkSize, localWorkSize, 0, NULL, eventND); + if ( res != CL_SUCCESS ) { + fputs("Failed to clEnqueueNDRangeKernel()\n", stderr); + return 1; + } + + res=clFlush(commandQueue); + res=clWaitForEvents(1,eventND); //XXX: SimpleGL utilise une attente active, pourquoi ? + res=clReleaseEvent(eventND[0]); + res=clEnqueueReleaseGLObjects(commandQueue, 1, &cl_vbo, 0, 0, 0); + res=clFinish(commandQueue); + + return CL_SUCCESS; +} +*/ + return 0; +} + +void OpenCLMeshKit::releaseKernels() { + for (std::map<std::string,cl_kernel>::iterator ii = kernels.begin(); ii != kernels.end(); ++ii ) { + clReleaseKernel((*ii).second); + } + kernels.clear(); +} + +cl_int OpenCLMeshKit::resetVBO() { + cl_int res; + std::map<std::string, cl_kernel> user_kernels=kernels; + + std::list<std::string> n; n.push_back("zero_z"); + res = compileKernels(n, kernel_src_zero_z, sizeof(kernel_src_zero_z)); + + if(res==0) res = execKernel("zero_z"); + + releaseKernels(); + kernels=user_kernels; + + return res; +} + +void OpenCLMeshKit::setGroupSize(size_t groupSize) { + this->groupSize=groupSize; +} + +OpenCLMeshKit::~OpenCLMeshKit() { } + |