Extras/RigidBodyGpuPipeline/opencl/primitives/Adl/CL/AdlKernelUtilsCL.inl

   1 /*
   2 Copyright (c) 2012 Advanced Micro Devices, Inc.
   3
   4 This software is provided 'as-is', without any express or implied warranty.
   5 In no event will the authors be held liable for any damages arising from the use of this software.
   6 Permission is granted to anyone to use this software for any purpose,
   7 including commercial applications, and to alter it and redistribute it freely,
   8 subject to the following restrictions:
   9
  10 1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required.
  11 2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.
  12 3. This notice may not be removed or altered from any source distribution.
  13 */
  14 //Originally written by Takahiro Harada
  15
  16
  17
  18
  19
  20 namespace adl
  21 {
  22
  23 struct KernelCL : public Kernel
  24 {
  25         cl_kernel& getKernel() { return (cl_kernel&)m_kernel; }
  26 };
  27
  28 static const char* strip(const char* name, const char* pattern)
  29 {
  30           size_t const patlen = strlen(pattern);
  31         size_t patcnt = 0;
  32           const char * oriptr;
  33           const char * patloc;
  34                 // find how many times the pattern occurs in the original string
  35           for (oriptr = name; patloc = strstr(oriptr, pattern); oriptr = patloc + patlen)
  36           {
  37                 patcnt++;
  38           }
  39           return oriptr;
  40 }
  41
  42 static bool isFileUpToDate(const char* binaryFileName,const char* srcFileName)
  43
  44 {
  45         bool fileUpToDate = false;
  46
  47         bool binaryFileValid=false;
  48         FILETIME modtimeBinary;
  49
  50         int nameLength = (int)strlen(binaryFileName)+1;
  51 #ifdef UNICODE
  52         WCHAR* fName = new WCHAR[nameLength];
  53         MultiByteToWideChar(CP_ACP,0,binaryFileName,-1, fName, nameLength);
  54         HANDLE binaryFileHandle = CreateFile(fName,GENERIC_READ,0,0,OPEN_EXISTING,FILE_ATTRIBUTE_NORMAL,0);
  55         delete [] fName;
  56 #else
  57         HANDLE binaryFileHandle = CreateFile(binaryFileName,GENERIC_READ,0,0,OPEN_EXISTING,FILE_ATTRIBUTE_NORMAL,0);
  58 #endif
  59         if (binaryFileHandle ==INVALID_HANDLE_VALUE)
  60         {
  61                 DWORD errorCode;
  62                 errorCode = GetLastError();
  63                 switch (errorCode)
  64                 {
  65                 case ERROR_FILE_NOT_FOUND:
  66                         {
  67                                 debugPrintf("\nCached file not found %s\n", binaryFileName);
  68                                 break;
  69                         }
  70                 case ERROR_PATH_NOT_FOUND:
  71                         {
  72                                 debugPrintf("\nCached file path not found %s\n", binaryFileName);
  73                                 break;
  74                         }
  75                 default:
  76                         {
  77                                 debugPrintf("\nFailed reading cached file with errorCode = %d\n", errorCode);
  78                         }
  79                 }
  80         } else
  81         {
  82                 if (GetFileTime(binaryFileHandle, NULL, NULL, &modtimeBinary)==0)
  83                 {
  84                         DWORD errorCode;
  85                         errorCode = GetLastError();
  86                         debugPrintf("\nGetFileTime errorCode = %d\n", errorCode);
  87                 } else
  88                 {
  89                         binaryFileValid = true;
  90                 }
  91                 CloseHandle(binaryFileHandle);
  92         }
  93
  94         if (binaryFileValid)
  95         {
  96 #ifdef UNICODE
  97                 int nameLength = (int)strlen(srcFileName)+1;
  98                 WCHAR* fName = new WCHAR[nameLength];
  99                 MultiByteToWideChar(CP_ACP,0,srcFileName,-1, fName, nameLength);
 100                 HANDLE srcFileHandle = CreateFile(fName,GENERIC_READ,0,0,OPEN_EXISTING,FILE_ATTRIBUTE_NORMAL,0);
 101                 delete [] fName;
 102 #else
 103                 HANDLE srcFileHandle = CreateFile(srcFileName,GENERIC_READ,0,0,OPEN_EXISTING,FILE_ATTRIBUTE_NORMAL,0);
 104 #endif
 105                 if (srcFileHandle!=INVALID_HANDLE_VALUE)
 106                 {
 107                         FILETIME modtimeSrc;
 108                         if (GetFileTime(srcFileHandle, NULL, NULL, &modtimeSrc)==0)
 109                         {
 110                                 DWORD errorCode;
 111                                 errorCode = GetLastError();
 112                                 debugPrintf("\nGetFileTime errorCode = %d\n", errorCode);
 113                         }
 114                         if (  ( modtimeSrc.dwHighDateTime < modtimeBinary.dwHighDateTime)
 115                                 ||(( modtimeSrc.dwHighDateTime == modtimeBinary.dwHighDateTime)&&(modtimeSrc.dwLowDateTime <= modtimeBinary.dwLowDateTime)))
 116                         {
 117                                 fileUpToDate=true;
 118                         } else
 119                         {
 120                                 debugPrintf("\nCached binary file found (%s), but out-of-date\n",binaryFileName);
 121                         }
 122                         CloseHandle(srcFileHandle);
 123                 }
 124                 else
 125                 {
 126 #ifdef _DEBUG
 127                         DWORD errorCode;
 128                         errorCode = GetLastError();
 129                         switch (errorCode)
 130                         {
 131                         case ERROR_FILE_NOT_FOUND:
 132                                 {
 133                                         debugPrintf("\nSrc file not found %s\n", srcFileName);
 134                                         break;
 135                                 }
 136                         case ERROR_PATH_NOT_FOUND:
 137                                 {
 138                                         debugPrintf("\nSrc path not found %s\n", srcFileName);
 139                                         break;
 140                                 }
 141                         default:
 142                                 {
 143                                         debugPrintf("\nnSrc file reading errorCode = %d\n", errorCode);
 144                                 }
 145                         }
 146                         ADLASSERT(0);
 147 #else
 148                         //if we cannot find the source, assume it is OK in release builds
 149                         fileUpToDate = true;
 150 #endif
 151                 }
 152         }
 153
 154
 155         return fileUpToDate;
 156 }
 157
 158 template<>
 159 void KernelBuilder<TYPE_CL>::setFromFile( const Device* deviceData, const char* fileName, const char* option, bool addExtension,
 160         bool cacheKernel)
 161 {
 162         m_deviceData = deviceData;
 163
 164         char fileNameWithExtension[256];
 165
 166         if( addExtension )
 167                 sprintf_s( fileNameWithExtension, "%s.cl", fileName );
 168         else
 169                 sprintf_s( fileNameWithExtension, "%s", fileName );
 170
 171         class File
 172         {
 173                 public:
 174                         __inline
 175                         bool open(const char* fileNameWithExtension)
 176                         {
 177                                 size_t      size;
 178                                 char*       str;
 179
 180                                 // Open file stream
 181                                 std::fstream f(fileNameWithExtension, (std::fstream::in | std::fstream::binary));
 182
 183                                 // Check if we have opened file stream
 184                                 if (f.is_open()) {
 185                                         size_t  sizeFile;
 186                                         // Find the stream size
 187                                         f.seekg(0, std::fstream::end);
 188                                         size = sizeFile = (size_t)f.tellg();
 189                                         f.seekg(0, std::fstream::beg);
 190
 191                                         str = new char[size + 1];
 192                                         if (!str) {
 193                                                 f.close();
 194                                                 return  NULL;
 195                                         }
 196
 197                                         // Read file
 198                                         f.read(str, sizeFile);
 199                                         f.close();
 200                                         str[size] = '\0';
 201
 202                                         m_source  = str;
 203
 204                                         delete[] str;
 205
 206                                         return true;
 207                                 }
 208
 209                                 return false;
 210                         }
 211                         const std::string& getSource() const {return m_source;}
 212
 213                 private:
 214                         std::string m_source;
 215         };
 216
 217         cl_program& program = (cl_program&)m_ptr;
 218         cl_int status = 0;
 219
 220         bool cacheBinary = cacheKernel;
 221 #if defined(ADL_CL_FORCE_UNCACHE_KERNEL)
 222         cacheBinary = false;
 223 #endif
 224
 225         char binaryFileName[512];
 226         {
 227                 char deviceName[256];
 228                 deviceData->getDeviceName(deviceName);
 229                 char driverVersion[256];
 230                 const DeviceCL* dd = (const DeviceCL*) deviceData;
 231                 clGetDeviceInfo(dd->m_deviceIdx, CL_DRIVER_VERSION, 256, &driverVersion, NULL);
 232                 const char* strippedFileName = strip(fileName,"\\");
 233                 strippedFileName = strip(strippedFileName,"/");
 234
 235                 sprintf_s(binaryFileName,"cache/%s.%s.%s.bin",strippedFileName, deviceName,driverVersion );
 236         }
 237
 238         bool upToDate = isFileUpToDate(binaryFileName,fileNameWithExtension);
 239
 240         if( cacheBinary && upToDate)
 241         {
 242                 FILE* file = fopen(binaryFileName, "rb");
 243
 244                 if( file )
 245                 {
 246                         fseek( file, 0L, SEEK_END );
 247                         size_t binarySize = ftell( file );
 248
 249                         rewind( file );
 250                         char* binary = new char[binarySize];
 251                         fread( binary, sizeof(char), binarySize, file );
 252                         fclose( file );
 253
 254                         if (binarySize)
 255                         {
 256                                 const DeviceCL* dd = (const DeviceCL*) deviceData;
 257                                 program = clCreateProgramWithBinary( dd->m_context, 1, &dd->m_deviceIdx, &binarySize, (const unsigned char**)&binary, 0, &status );
 258                                 ADLASSERT( status == CL_SUCCESS );
 259                                 status = clBuildProgram( program, 1, &dd->m_deviceIdx, option, 0, 0 );
 260                                 ADLASSERT( status == CL_SUCCESS );
 261                         if( status != CL_SUCCESS )
 262                         {
 263                                 char *build_log;
 264                                 size_t ret_val_size;
 265                                 clGetProgramBuildInfo(program, dd->m_deviceIdx, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
 266                                 build_log = new char[ret_val_size+1];
 267                                 clGetProgramBuildInfo(program, dd->m_deviceIdx, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
 268
 269                                 build_log[ret_val_size] = '\0';
 270
 271                                 debugPrintf("%s\n", build_log);
 272
 273                                 delete build_log;
 274                                 ADLASSERT(0);
 275                                 }
 276
 277                         }
 278                 }
 279         }
 280         if( !m_ptr )
 281         {
 282                 File kernelFile;
 283                 ADLASSERT( kernelFile.open( fileNameWithExtension ) );
 284                 const char* source = kernelFile.getSource().c_str();
 285                 setFromSrc( m_deviceData, source, option );
 286
 287                 if( cacheBinary )
 288                 {       //      write to binary
 289                         size_t binarySize;
 290                         status = clGetProgramInfo( program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binarySize, 0 );
 291                         ADLASSERT( status == CL_SUCCESS );
 292
 293                         char* binary = new char[binarySize];
 294
 295                         status = clGetProgramInfo( program, CL_PROGRAM_BINARIES, sizeof(char*), &binary, 0 );
 296                         ADLASSERT( status == CL_SUCCESS );
 297
 298                         {
 299                                 FILE* file = fopen(binaryFileName, "wb");
 300                                 if (file)
 301                                 {
 302                                         fwrite( binary, sizeof(char), binarySize, file );
 303                                         fclose( file );
 304                                 }
 305                         }
 306
 307                         delete [] binary;
 308                 }
 309         }
 310 }
 311
 312
 313
 314 template<>
 315 void KernelBuilder<TYPE_CL>::setFromSrcCached( const Device* deviceData, const char* src, const char* fileName, const char* option )
 316 {
 317         m_deviceData = deviceData;
 318
 319         bool cacheBinary = true;
 320         cl_program& program = (cl_program&)m_ptr;
 321         cl_int status = 0;
 322
 323         char binaryFileName[512];
 324         {
 325                 char deviceName[256];
 326                 deviceData->getDeviceName(deviceName);
 327                 char driverVersion[256];
 328                 const DeviceCL* dd = (const DeviceCL*) deviceData;
 329                 clGetDeviceInfo(dd->m_deviceIdx, CL_DRIVER_VERSION, 256, &driverVersion, NULL);
 330
 331                 const char* strippedFileName = strip(fileName,"\\");
 332                 strippedFileName = strip(strippedFileName,"/");
 333
 334                 sprintf_s(binaryFileName,"cache/%s.%s.%s.bin",strippedFileName, deviceName,driverVersion );
 335         }
 336
 337
 338         char fileNameWithExtension[256];
 339         sprintf_s(fileNameWithExtension,"%s.cl",fileName, ".cl");
 340
 341         bool upToDate = isFileUpToDate(binaryFileName,fileNameWithExtension);
 342
 343
 344         if( cacheBinary )
 345         {
 346
 347                 bool fileUpToDate = isFileUpToDate(binaryFileName,fileNameWithExtension);
 348
 349                 if( fileUpToDate)
 350                 {
 351                         FILE* file = fopen(binaryFileName, "rb");
 352                         if (file)
 353                         {
 354                                 fseek( file, 0L, SEEK_END );
 355                                 size_t binarySize = ftell( file );
 356                                 rewind( file );
 357                                 char* binary = new char[binarySize];
 358                                 fread( binary, sizeof(char), binarySize, file );
 359                                 fclose( file );
 360
 361                                 const DeviceCL* dd = (const DeviceCL*) deviceData;
 362                                 program = clCreateProgramWithBinary( dd->m_context, 1, &dd->m_deviceIdx, &binarySize, (const unsigned char**)&binary, 0, &status );
 363                                 ADLASSERT( status == CL_SUCCESS );
 364                                 status = clBuildProgram( program, 1, &dd->m_deviceIdx, option, 0, 0 );
 365                                 ADLASSERT( status == CL_SUCCESS );
 366
 367                                 if( status != CL_SUCCESS )
 368                                 {
 369                                         char *build_log;
 370                                         size_t ret_val_size;
 371                                         clGetProgramBuildInfo(program, dd->m_deviceIdx, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
 372                                         build_log = new char[ret_val_size+1];
 373                                         clGetProgramBuildInfo(program, dd->m_deviceIdx, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
 374
 375                                         build_log[ret_val_size] = '\0';
 376
 377                                         debugPrintf("%s\n", build_log);
 378
 379                                         delete build_log;
 380                                         ADLASSERT(0);
 381                                 }
 382                                 delete[] binary;
 383                         }
 384                 }
 385         }
 386
 387
 388         if( !m_ptr )
 389         {
 390
 391                 setFromSrc( deviceData, src, option );
 392
 393                 if( cacheBinary )
 394                 {       //      write to binary
 395                         cl_uint numAssociatedDevices;
 396                         status = clGetProgramInfo( program, CL_PROGRAM_NUM_DEVICES, sizeof(cl_uint), &numAssociatedDevices, 0 );
 397                         ADLASSERT( status == CL_SUCCESS );
 398                         if (numAssociatedDevices==1)
 399                         {
 400
 401
 402                                 size_t binarySize;
 403                                 status = clGetProgramInfo( program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binarySize, 0 );
 404                                 ADLASSERT( status == CL_SUCCESS );
 405
 406                                 char* binary = new char[binarySize];
 407
 408                                 status = clGetProgramInfo( program, CL_PROGRAM_BINARIES, sizeof(char*), &binary, 0 );
 409                                 ADLASSERT( status == CL_SUCCESS );
 410
 411                                 {
 412                                         FILE* file = fopen(binaryFileName, "wb");
 413                                         if (file)
 414                                         {
 415                                                 fwrite( binary, sizeof(char), binarySize, file );
 416                                                 fclose( file );
 417                                         }
 418                                 }
 419
 420                                 delete [] binary;
 421                         }
 422                 }
 423         }
 424 }
 425
 426
 427 template<>
 428 void KernelBuilder<TYPE_CL>::setFromSrc( const Device* deviceData, const char* src, const char* option )
 429 {
 430         ADLASSERT( deviceData->m_type == TYPE_CL );
 431         m_deviceData = deviceData;
 432         const DeviceCL* dd = (const DeviceCL*) deviceData;
 433
 434         cl_program& program = (cl_program&)m_ptr;
 435         cl_int status = 0;
 436         size_t srcSize[] = {strlen( src )};
 437         program = clCreateProgramWithSource( dd->m_context, 1, &src, srcSize, &status );
 438         ADLASSERT( status == CL_SUCCESS );
 439         status = clBuildProgram( program, 1, &dd->m_deviceIdx, option, NULL, NULL );
 440         if( status != CL_SUCCESS )
 441         {
 442                 char *build_log;
 443                 size_t ret_val_size;
 444                 clGetProgramBuildInfo(program, dd->m_deviceIdx, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size);
 445                 build_log = new char[ret_val_size+1];
 446                 clGetProgramBuildInfo(program, dd->m_deviceIdx, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL);
 447
 448                 build_log[ret_val_size] = '\0';
 449
 450                 debugPrintf("%s\n", build_log);
 451                 printf("%s\n", build_log);
 452
 453                 ADLASSERT(0);
 454                 delete build_log;
 455
 456         }
 457 }
 458
 459 template<>
 460 KernelBuilder<TYPE_CL>::~KernelBuilder()
 461 {
 462         cl_program program = (cl_program)m_ptr;
 463         clReleaseProgram( program );
 464 }
 465
 466 template<>
 467 void KernelBuilder<TYPE_CL>::createKernel( const char* funcName, Kernel& kernelOut )
 468 {
 469         KernelCL* clKernel = (KernelCL*)&kernelOut;
 470
 471         cl_program program = (cl_program)m_ptr;
 472         cl_int status = 0;
 473         clKernel->getKernel() = clCreateKernel(program, funcName, &status );
 474         ADLASSERT( status == CL_SUCCESS );
 475
 476         kernelOut.m_type = TYPE_CL;
 477 }
 478
 479 template<>
 480 void KernelBuilder<TYPE_CL>::deleteKernel( Kernel& kernel )
 481 {
 482         KernelCL* clKernel = (KernelCL*)&kernel;
 483         clReleaseKernel( clKernel->getKernel() );
 484 }
 485
 486
 487
 488 class LauncherCL
 489 {
 490         public:
 491                 typedef Launcher::BufferInfo BufferInfo;
 492
 493                 __inline
 494                 static void setBuffers( Launcher* launcher, BufferInfo* buffInfo, int n );
 495                 template<typename T>
 496                 __inline
 497                 static void setConst( Launcher* launcher, Buffer<T>& constBuff, const T& consts );
 498                 __inline
 499                 static void launch2D( Launcher* launcher, int numThreadsX, int numThreadsY, int localSizeX, int localSizeY );
 500 };
 501
 502 void LauncherCL::setBuffers( Launcher* launcher, BufferInfo* buffInfo, int n )
 503 {
 504         KernelCL* clKernel = (KernelCL*)launcher->m_kernel;
 505         for(int i=0; i<n; i++)
 506         {
 507                 Buffer<int>* buff = (Buffer<int>*)buffInfo[i].m_buffer;
 508                 cl_int status = clSetKernelArg( clKernel->getKernel(), launcher->m_idx++, sizeof(cl_mem), &buff->m_ptr );
 509                 ADLASSERT( status == CL_SUCCESS );
 510         }
 511 }
 512
 513 template<typename T>
 514 void LauncherCL::setConst( Launcher* launcher, Buffer<T>& constBuff, const T& consts )
 515 {
 516         KernelCL* clKernel = (KernelCL*)launcher->m_kernel;
 517         int sz=sizeof(T);
 518         cl_int status = clSetKernelArg( clKernel->getKernel(), launcher->m_idx++, sz, &consts );
 519         ADLASSERT( status == CL_SUCCESS );
 520 }
 521
 522 void LauncherCL::launch2D( Launcher* launcher, int numThreadsX, int numThreadsY, int localSizeX, int localSizeY )
 523 {
 524         KernelCL* clKernel = (KernelCL*)launcher->m_kernel;
 525         const DeviceCL* ddcl = (const DeviceCL*)launcher->m_deviceData;
 526         size_t gRange[3] = {1,1,1};
 527         size_t lRange[3] = {1,1,1};
 528         lRange[0] = localSizeX;
 529         lRange[1] = localSizeY;
 530         gRange[0] = max((size_t)1, (numThreadsX/lRange[0])+(!(numThreadsX%lRange[0])?0:1));
 531         gRange[0] *= lRange[0];
 532         gRange[1] = max((size_t)1, (numThreadsY/lRange[1])+(!(numThreadsY%lRange[1])?0:1));
 533         gRange[1] *= lRange[1];
 534
 535         cl_int status = clEnqueueNDRangeKernel( ddcl->m_commandQueue,
 536                 clKernel->getKernel(), 2, NULL, gRange, lRange, 0,0,0 );
 537         ADLASSERT( status == CL_SUCCESS );
 538 }
 539
 540
 541 };