skepu_opencl_helpers.h

Go to the documentation of this file.

#ifndef SKEPU_CUDA_HELPER_H
#define SKEPU_CUDA_HELPER_H

#include <stdio.h>
#include <string.h>
#include <stdlib.h>

#include "device_cl.h"


namespace skepu
{


#ifndef MIN
#define MIN(a,b) ((a < b) ? a : b)
#endif

#ifndef MAX
#define MAX(a,b) ((a > b) ? a : b)
#endif


// Give a little more for Windows : the console window often disapears before we can read the message
#ifdef _WIN32
# if 1//ndef UNICODE
#  ifdef _DEBUG // Do this only in debug mode...
inline void VSPrintf(FILE *file, LPCSTR fmt, ...)
{
   size_t fmt2_sz   = 2048;
   char *fmt2       = (char*)malloc(fmt2_sz);
   va_list  vlist;
   va_start(vlist, fmt);
   while((_vsnprintf(fmt2, fmt2_sz, fmt, vlist)) < 0) // means there wasn't anough room
   {
      fmt2_sz *= 2;
      if(fmt2) free(fmt2);
      fmt2 = (char*)malloc(fmt2_sz);
   }
   OutputDebugStringA(fmt2);
   fprintf(file, fmt2);
   free(fmt2);
}
#   define FPRINTF(a) VSPrintf a
#  else //debug
#   define FPRINTF(a) fprintf a
// For other than Win32
#  endif //debug
# else //unicode
// Unicode case... let's give-up for now and keep basic printf
#   define FPRINTF(a) fprintf a
# endif //unicode
#else //win32
#   define FPRINTF(a) fprintf a
#endif //win32


template <typename T>
void copyDeviceToHost(T *hostPtr, cl_mem devPtr, size_t numElements, Device_CL* device, size_t offset)
{
   if(devPtr != NULL && hostPtr != NULL)
   {
      DEBUG_TEXT_LEVEL2("** DEVICE_TO_HOST OpenCL: "<< numElements <<"!!!\n")

      cl_int err;

      size_t sizeVec;

      sizeVec = numElements*sizeof(T);

      err = clEnqueueReadBuffer(device->getQueue(), devPtr, CL_TRUE, offset, sizeVec, (void*)hostPtr, 0, NULL, NULL);

      if(err != CL_SUCCESS)
      {
         FPRINTF((stderr, "Error copying data from device\n"));
      }
   }
}



template <typename T>
void copyHostToDevice(T *hostPtr, cl_mem devPtr, size_t numElements, Device_CL* device, size_t offset)
{
   if(hostPtr != NULL && devPtr != NULL)
   {
      DEBUG_TEXT_LEVEL2("** HOST_TO_DEVICE OpenCL: "<< numElements <<"!!!\n")

      cl_int err;

      size_t sizeVec;

      sizeVec = numElements*sizeof(T);

      err = clEnqueueWriteBuffer(device->getQueue(), devPtr, CL_TRUE, offset, sizeVec, (void*)hostPtr, 0, NULL, NULL);

      if(err != CL_SUCCESS)
      {
         FPRINTF((stderr, "Error copying data to device\n"));
      }
   }
}


template <typename T>
inline cl_mem allocateOpenCLMemory(size_t size, Device_CL* device)
{
   DEBUG_TEXT_LEVEL2("** ALLOC OpenCL: "<< size <<"!!!\n")

   cl_int err;
   cl_mem devicePointer;

   size_t sizeVec = size*sizeof(T);

   devicePointer = clCreateBuffer(device->getContext(), CL_MEM_READ_WRITE, sizeVec, NULL, &err);
   if(err != CL_SUCCESS)
   {
      FPRINTF((stderr, "Error allocating memory on device\n"));
   }

   return devicePointer;
}


template <typename T>
inline void freeOpenCLMemory(cl_mem d_pointer)
{
   DEBUG_TEXT_LEVEL2("** DE-ALLOC OpenCL !!!\n")

//  if(d_pointer!=NULL)
   {
      if(clReleaseMemObject(d_pointer) != CL_SUCCESS)
         FPRINTF((stderr, "Error releasing memory on device\n"));
   }
}



template <typename T>
void ExecuteReduceOnADevice(size_t  n, const size_t  &numThreads, const size_t  &numBlocks, _cl_mem*& in_p, _cl_mem*& out_p, cl_kernel &kernel, Device_CL *device)
{
   cl_int err;

   size_t globalWorkSize[1];
   size_t localWorkSize[1];

   size_t sharedMemSize = (numThreads <= 32) ? 2 * numThreads * sizeof(T) : numThreads * sizeof(T);

   // Sets the kernel arguments for first reduction
   clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&in_p);
   clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&out_p);
   clSetKernelArg(kernel, 2, sizeof(size_t), (void*)&n);
   clSetKernelArg(kernel, 3, sharedMemSize, NULL);

   globalWorkSize[0] = numBlocks * numThreads;
   localWorkSize[0] = numThreads;

   // First reduce all elements blockwise so that each block produces one element.
   err = clEnqueueNDRangeKernel(device->getQueue(), kernel, 1, NULL, globalWorkSize, localWorkSize, 0, NULL, NULL);
   if(err != CL_SUCCESS)
   {
      std::cerr<<"Error launching kernel RowWise!! 1st\n";
   }

   // Sets the kernel arguments for second reduction
   n = numBlocks;
   clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&out_p);
   clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&out_p);
   clSetKernelArg(kernel, 2, sizeof(size_t), (void*)&n);
   clSetKernelArg(kernel, 3, sharedMemSize, NULL);

   globalWorkSize[0] = 1 * numThreads;
   localWorkSize[0] = numThreads;

   // Reduces the elements from the previous reduction in a single block to produce the scalar result.
   err = clEnqueueNDRangeKernel(device->getQueue(), kernel, 1, NULL, globalWorkSize, localWorkSize, 0, NULL, NULL);
   if(err != CL_SUCCESS)
   {
      std::cerr<<"Error launching kernel RowWise!! 2nd\n";
   }
}




void replaceTextInString(std::string& text, std::string find, std::string replace)
{
   std::string::size_type pos=0;
   while((pos = text.find(find, pos)) != std::string::npos)
   {
      text.erase(pos, find.length());
      text.insert(pos, replace);
      pos+=replace.length();
   }
}

void printCLError(cl_int Err, std::string s = std::string())
{
  std::string msg;
  if (Err != CL_SUCCESS)
  {
    switch(Err)
    {
      case CL_DEVICE_NOT_FOUND:
        msg = "Device not found";
      break;
      case CL_DEVICE_NOT_AVAILABLE:
        msg = "Device not available";
      break;
      case CL_COMPILER_NOT_AVAILABLE:
        msg = "Compiler not available";
      break;
      case CL_MEM_OBJECT_ALLOCATION_FAILURE:
        msg = "Memory object allocation failure";
      break;
      case CL_OUT_OF_RESOURCES:
        msg = "Out of resources";
      break;
      case CL_OUT_OF_HOST_MEMORY:
        msg = "Out of host memory";
      break;
      case CL_PROFILING_INFO_NOT_AVAILABLE:
        msg = "Profiling info not available";
      break;
      case CL_MEM_COPY_OVERLAP:
        msg = "Memory copy overlap";
      break;
      case CL_IMAGE_FORMAT_MISMATCH:
        msg = "Image format mismatch";
      break;
      case CL_IMAGE_FORMAT_NOT_SUPPORTED:
        msg = "Image format not supported";
      break;
      case CL_BUILD_PROGRAM_FAILURE:
        msg = "Build program failure";
      break;
      case CL_MAP_FAILURE:
        msg = "Map failure";
      break;
      case CL_MISALIGNED_SUB_BUFFER_OFFSET:
        msg = "Misaligned sub buffer offset";
      break;
      case CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST:
        msg = "Exec status error for events in wait list";
      break;
      case CL_INVALID_VALUE:
        msg = "Invalid value";
      break;
      case CL_INVALID_DEVICE_TYPE:
        msg = "Invalid device type";
      break;
      case CL_INVALID_PLATFORM:
        msg = "Invalid platform";
      break;
      case CL_INVALID_DEVICE:
        msg = "Invalid device";
      break;
      case CL_INVALID_CONTEXT:
        msg = "Invalid context";
      break;
      case CL_INVALID_QUEUE_PROPERTIES:
        msg = "Invalid queue properties";
      break;
      case CL_INVALID_COMMAND_QUEUE:
        msg = "Invalid command queue";
      break;
      case CL_INVALID_HOST_PTR:
        msg = "Invalid host pointer";
      break;
      case CL_INVALID_MEM_OBJECT:
        msg = "Invalid memory object";
      break;
      case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:
        msg = "Invalid image format descriptor";
      break;
      case CL_INVALID_IMAGE_SIZE:
        msg = "Invalid image size";
      break;
      case CL_INVALID_SAMPLER:
        msg = "Invalid sampler";
      break;
      case CL_INVALID_BINARY:
        msg = "Invalid binary";
      break;
      case CL_INVALID_BUILD_OPTIONS:
        msg = "Invalid build options";
      break;
      case CL_INVALID_PROGRAM:
        msg = "Invalid program";
      break;
      case CL_INVALID_PROGRAM_EXECUTABLE:
        msg = "Invalid program executable";
      break;
      case CL_INVALID_KERNEL_NAME:
        msg = "Invalid kernel name";
      break;
      case CL_INVALID_KERNEL_DEFINITION:
        msg = "Invalid kernel definition";
      break;
      case CL_INVALID_KERNEL:
        msg = "Invalid kernel";
      break;
      case CL_INVALID_ARG_INDEX:
        msg = "Invalid argument index";
      break;
      case CL_INVALID_ARG_VALUE:
        msg = "Invalid argument value";
      break;
      case CL_INVALID_ARG_SIZE:
        msg = "Invalid argument size";
      break;
      case CL_INVALID_KERNEL_ARGS:
        msg = "Invalid kernel arguments";
      break;
      case CL_INVALID_WORK_DIMENSION:
        msg = "Invalid work dimension";
      break;
      case CL_INVALID_WORK_GROUP_SIZE:
        msg = "Invalid work group size";
      break;
      case CL_INVALID_WORK_ITEM_SIZE:
        msg = "Invalid work item size";
      break;
      case CL_INVALID_GLOBAL_OFFSET:
        msg = "Invalid global offset";
      break;
      case CL_INVALID_EVENT_WAIT_LIST:
        msg = "Invalid event wait list";
      break;
      case CL_INVALID_EVENT:
        msg = "Invalid event";
      break;
      case CL_INVALID_OPERATION:
        msg = "Invalid operation";
      break;
      case CL_INVALID_GL_OBJECT:
        msg = "Invalid GL object";
      break;
      case CL_INVALID_BUFFER_SIZE:
        msg = "Invalid buffer size";
      break;
      case CL_INVALID_MIP_LEVEL:
        msg = "Invalid MIP level";
      break;
      case CL_INVALID_GLOBAL_WORK_SIZE:
        msg = "Invalid global work size";
      break;
      case CL_INVALID_PROPERTY:
        msg = "Invalid property";
      break;
      default:
        msg = "Unknown error";
      break;
    }
    SKEPU_ERROR(s<<" OpenCL error code "<<Err<<" "<<msg<<" \n");
  }
}


}

#endif