2 Copyright (c) 2012 Advanced Micro Devices, Inc.
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:
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.
14 //Originally written by Takahiro Harada
23 struct KernelCL : public Kernel
25 cl_kernel& getKernel() { return (cl_kernel&)m_kernel; }
28 static const char* strip(const char* name, const char* pattern)
30 size_t const patlen = strlen(pattern);
34 // find how many times the pattern occurs in the original string
35 for (oriptr = name; patloc = strstr(oriptr, pattern); oriptr = patloc + patlen)
42 static bool isFileUpToDate(const char* binaryFileName,const char* srcFileName)
45 bool fileUpToDate = false;
47 bool binaryFileValid=false;
48 FILETIME modtimeBinary;
50 int nameLength = (int)strlen(binaryFileName)+1;
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);
57 HANDLE binaryFileHandle = CreateFile(binaryFileName,GENERIC_READ,0,0,OPEN_EXISTING,FILE_ATTRIBUTE_NORMAL,0);
59 if (binaryFileHandle ==INVALID_HANDLE_VALUE)
62 errorCode = GetLastError();
65 case ERROR_FILE_NOT_FOUND:
67 debugPrintf("\nCached file not found %s\n", binaryFileName);
70 case ERROR_PATH_NOT_FOUND:
72 debugPrintf("\nCached file path not found %s\n", binaryFileName);
77 debugPrintf("\nFailed reading cached file with errorCode = %d\n", errorCode);
82 if (GetFileTime(binaryFileHandle, NULL, NULL, &modtimeBinary)==0)
85 errorCode = GetLastError();
86 debugPrintf("\nGetFileTime errorCode = %d\n", errorCode);
89 binaryFileValid = true;
91 CloseHandle(binaryFileHandle);
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);
103 HANDLE srcFileHandle = CreateFile(srcFileName,GENERIC_READ,0,0,OPEN_EXISTING,FILE_ATTRIBUTE_NORMAL,0);
105 if (srcFileHandle!=INVALID_HANDLE_VALUE)
108 if (GetFileTime(srcFileHandle, NULL, NULL, &modtimeSrc)==0)
111 errorCode = GetLastError();
112 debugPrintf("\nGetFileTime errorCode = %d\n", errorCode);
114 if ( ( modtimeSrc.dwHighDateTime < modtimeBinary.dwHighDateTime)
115 ||(( modtimeSrc.dwHighDateTime == modtimeBinary.dwHighDateTime)&&(modtimeSrc.dwLowDateTime <= modtimeBinary.dwLowDateTime)))
120 debugPrintf("\nCached binary file found (%s), but out-of-date\n",binaryFileName);
122 CloseHandle(srcFileHandle);
128 errorCode = GetLastError();
131 case ERROR_FILE_NOT_FOUND:
133 debugPrintf("\nSrc file not found %s\n", srcFileName);
136 case ERROR_PATH_NOT_FOUND:
138 debugPrintf("\nSrc path not found %s\n", srcFileName);
143 debugPrintf("\nnSrc file reading errorCode = %d\n", errorCode);
148 //if we cannot find the source, assume it is OK in release builds
159 void KernelBuilder<TYPE_CL>::setFromFile( const Device* deviceData, const char* fileName, const char* option, bool addExtension,
162 m_deviceData = deviceData;
164 char fileNameWithExtension[256];
167 sprintf_s( fileNameWithExtension, "%s.cl", fileName );
169 sprintf_s( fileNameWithExtension, "%s", fileName );
175 bool open(const char* fileNameWithExtension)
181 std::fstream f(fileNameWithExtension, (std::fstream::in | std::fstream::binary));
183 // Check if we have opened file stream
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);
191 str = new char[size + 1];
198 f.read(str, sizeFile);
211 const std::string& getSource() const {return m_source;}
214 std::string m_source;
217 cl_program& program = (cl_program&)m_ptr;
220 bool cacheBinary = cacheKernel;
221 #if defined(ADL_CL_FORCE_UNCACHE_KERNEL)
225 char binaryFileName[512];
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,"/");
235 sprintf_s(binaryFileName,"cache/%s.%s.%s.bin",strippedFileName, deviceName,driverVersion );
238 bool upToDate = isFileUpToDate(binaryFileName,fileNameWithExtension);
240 if( cacheBinary && upToDate)
242 FILE* file = fopen(binaryFileName, "rb");
246 fseek( file, 0L, SEEK_END );
247 size_t binarySize = ftell( file );
250 char* binary = new char[binarySize];
251 fread( binary, sizeof(char), binarySize, file );
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 )
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);
269 build_log[ret_val_size] = '\0';
271 debugPrintf("%s\n", build_log);
283 ADLASSERT( kernelFile.open( fileNameWithExtension ) );
284 const char* source = kernelFile.getSource().c_str();
285 setFromSrc( m_deviceData, source, option );
290 status = clGetProgramInfo( program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binarySize, 0 );
291 ADLASSERT( status == CL_SUCCESS );
293 char* binary = new char[binarySize];
295 status = clGetProgramInfo( program, CL_PROGRAM_BINARIES, sizeof(char*), &binary, 0 );
296 ADLASSERT( status == CL_SUCCESS );
299 FILE* file = fopen(binaryFileName, "wb");
302 fwrite( binary, sizeof(char), binarySize, file );
315 void KernelBuilder<TYPE_CL>::setFromSrcCached( const Device* deviceData, const char* src, const char* fileName, const char* option )
317 m_deviceData = deviceData;
319 bool cacheBinary = true;
320 cl_program& program = (cl_program&)m_ptr;
323 char binaryFileName[512];
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);
331 const char* strippedFileName = strip(fileName,"\\");
332 strippedFileName = strip(strippedFileName,"/");
334 sprintf_s(binaryFileName,"cache/%s.%s.%s.bin",strippedFileName, deviceName,driverVersion );
338 char fileNameWithExtension[256];
339 sprintf_s(fileNameWithExtension,"%s.cl",fileName, ".cl");
341 bool upToDate = isFileUpToDate(binaryFileName,fileNameWithExtension);
347 bool fileUpToDate = isFileUpToDate(binaryFileName,fileNameWithExtension);
351 FILE* file = fopen(binaryFileName, "rb");
354 fseek( file, 0L, SEEK_END );
355 size_t binarySize = ftell( file );
357 char* binary = new char[binarySize];
358 fread( binary, sizeof(char), binarySize, file );
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 );
367 if( status != CL_SUCCESS )
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);
375 build_log[ret_val_size] = '\0';
377 debugPrintf("%s\n", build_log);
391 setFromSrc( deviceData, src, option );
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)
403 status = clGetProgramInfo( program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &binarySize, 0 );
404 ADLASSERT( status == CL_SUCCESS );
406 char* binary = new char[binarySize];
408 status = clGetProgramInfo( program, CL_PROGRAM_BINARIES, sizeof(char*), &binary, 0 );
409 ADLASSERT( status == CL_SUCCESS );
412 FILE* file = fopen(binaryFileName, "wb");
415 fwrite( binary, sizeof(char), binarySize, file );
428 void KernelBuilder<TYPE_CL>::setFromSrc( const Device* deviceData, const char* src, const char* option )
430 ADLASSERT( deviceData->m_type == TYPE_CL );
431 m_deviceData = deviceData;
432 const DeviceCL* dd = (const DeviceCL*) deviceData;
434 cl_program& program = (cl_program&)m_ptr;
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 )
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);
448 build_log[ret_val_size] = '\0';
450 debugPrintf("%s\n", build_log);
451 printf("%s\n", build_log);
460 KernelBuilder<TYPE_CL>::~KernelBuilder()
462 cl_program program = (cl_program)m_ptr;
463 clReleaseProgram( program );
467 void KernelBuilder<TYPE_CL>::createKernel( const char* funcName, Kernel& kernelOut )
469 KernelCL* clKernel = (KernelCL*)&kernelOut;
471 cl_program program = (cl_program)m_ptr;
473 clKernel->getKernel() = clCreateKernel(program, funcName, &status );
474 ADLASSERT( status == CL_SUCCESS );
476 kernelOut.m_type = TYPE_CL;
480 void KernelBuilder<TYPE_CL>::deleteKernel( Kernel& kernel )
482 KernelCL* clKernel = (KernelCL*)&kernel;
483 clReleaseKernel( clKernel->getKernel() );
491 typedef Launcher::BufferInfo BufferInfo;
494 static void setBuffers( Launcher* launcher, BufferInfo* buffInfo, int n );
497 static void setConst( Launcher* launcher, Buffer<T>& constBuff, const T& consts );
499 static void launch2D( Launcher* launcher, int numThreadsX, int numThreadsY, int localSizeX, int localSizeY );
502 void LauncherCL::setBuffers( Launcher* launcher, BufferInfo* buffInfo, int n )
504 KernelCL* clKernel = (KernelCL*)launcher->m_kernel;
505 for(int i=0; i<n; i++)
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 );
514 void LauncherCL::setConst( Launcher* launcher, Buffer<T>& constBuff, const T& consts )
516 KernelCL* clKernel = (KernelCL*)launcher->m_kernel;
518 cl_int status = clSetKernelArg( clKernel->getKernel(), launcher->m_idx++, sz, &consts );
519 ADLASSERT( status == CL_SUCCESS );
522 void LauncherCL::launch2D( Launcher* launcher, int numThreadsX, int numThreadsY, int localSizeX, int localSizeY )
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];
535 cl_int status = clEnqueueNDRangeKernel( ddcl->m_commandQueue,
536 clKernel->getKernel(), 2, NULL, gRange, lRange, 0,0,0 );
537 ADLASSERT( status == CL_SUCCESS );