From 35e25937ad05e409340e7cd356c3ce1a45a5a3f9 Mon Sep 17 00:00:00 2001 From: Ludovic Pouzenc Date: Sun, 3 Mar 2013 20:27:37 +0000 Subject: Avancee sur le C++. Boring parts disparait presque, au profit d'une classe OpenCLMeshKit. Il ne manque que execKernel() a code dans cette classe. Dans gpudataviz.cpp, il faut changer le code OpenGL pour afficher le maillage et pas un isodecaheron de test. git-svn-id: file:///var/svn/2013-gpudataviz/trunk@18 371a6b4a-a258-45f8-9dcc-bdd82ce0ac9d --- src/opencl_mesh_kit.cpp | 225 ++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 225 insertions(+) create mode 100644 src/opencl_mesh_kit.cpp (limited to 'src/opencl_mesh_kit.cpp') 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 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()" << 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::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::iterator ii = kernels.begin(); ii != kernels.end(); ++ii ) { + clReleaseKernel((*ii).second); + } + kernels.clear(); +} + +cl_int OpenCLMeshKit::resetVBO() { + cl_int res; + std::map user_kernels=kernels; + + std::list 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() { } + -- cgit v1.2.3