@@ -246,13 +246,8 @@ void Buffer::CopyTo( Buffer* buffer )
246
246
void Buffer::Clear ()
247
247
{
248
248
uint value = 0 ;
249
- #if 0
250
- memset( hostBuffer, 0, size );
251
- CopyToDevice();
252
- #else
253
249
cl_int error;
254
250
CHECKCL ( error = clEnqueueFillBuffer ( Kernel::GetQueue (), deviceBuffer, &value, 4 , 0 , size, 0 , 0 , 0 ) );
255
- #endif
256
251
}
257
252
258
253
// Kernel constructor
@@ -272,45 +267,7 @@ Kernel::Kernel( char* file, char* entryPoint )
272
267
if (isAmpere) csText = " #define ISAMPERE\n " + csText, vendorLines++;
273
268
if (isTuring) csText = " #define ISTURING\n " + csText, vendorLines++;
274
269
if (isPascal) csText = " #define ISPASCAL\n " + csText, vendorLines++;
275
- // expand #include directives: cl compiler doesn't support these natively
276
- // warning: this simple system does not handle nested includes.
277
- struct Include { int start, end; string file; } includes[64 ];
278
- int Ninc = 0 ;
279
- #if 0 // needed for NVIDIA OpenCL 1.0; this no longer is necessary
280
- while (1)
281
- {
282
- // see if any #includes remain
283
- size_t pos = csText.find( "#include" );
284
- if (pos == string::npos) break;
285
- // start of expanded source construction
286
- string tmp;
287
- if (pos > 0)
288
- tmp = csText.substr( 0, pos - 1 ) + "\n",
289
- includes[Ninc].start = LineCount( tmp ); // record first line of #include content
290
- else
291
- includes[Ninc].start = 0;
292
- // parse filename of include file
293
- pos = csText.find( "\"", pos + 1 );
294
- if (pos == string::npos) FatalError( "Expected \" after #include in shader." );
295
- size_t end = csText.find( "\"", pos + 1 );
296
- if (end == string::npos) FatalError( "Expected second \" after #include in shader." );
297
- string file = csText.substr( pos + 1, end - pos - 1 );
298
- // load include file content
299
- string incText = TextFileRead( file.c_str() );
300
- includes[Ninc].end = includes[Ninc].start + LineCount( incText );
301
- includes[Ninc++].file = file;
302
- if (incText.size() == 0) FatalError( "#include file not found:\n%s", file.c_str() );
303
- // cleanup include file content: we get some crap first sometimes, but why?
304
- int firstValidChar = 0;
305
- while (incText[firstValidChar] < 0) firstValidChar++;
306
- // add include file content and remainder of source to expanded source string
307
- tmp += incText.substr( firstValidChar, string::npos );
308
- tmp += csText.substr( end + 1, string::npos ) + "\n";
309
- // repeat until no #includes left
310
- csText = tmp;
311
- }
312
- #endif
313
- // attempt to compile the loaded and expanded source text
270
+ // attempt to compile the loaded source text
314
271
const char * source = csText.c_str ();
315
272
size_t size = strlen ( source );
316
273
cl_int error;
@@ -355,58 +312,9 @@ Kernel::Kernel( char* file, char* entryPoint )
355
312
FILE* f = fopen ( " errorlog.txt" , " wb" );
356
313
fwrite ( log, 1 , size, f );
357
314
fclose ( f );
358
- // find and display the first error. Note: platform specific sadly; code below is for NVIDIA
359
- char * errorString = strstr ( log, " : error:" );
360
- if (errorString)
361
- {
362
- int errorPos = (int )(errorString - log);
363
- while (errorPos > 0 ) if (log[errorPos - 1 ] == ' \n ' ) break ; else errorPos--;
364
- // translate file and line number of error and report
365
- log[errorPos + 2048 ] = 0 ;
366
- int lineNr = 0 , linePos = 0 ;
367
- char * lns = strstr ( log + errorPos, " >:" ), * eol;
368
- if (!lns) FatalError ( log + errorPos ); else
369
- {
370
- lns += 2 ;
371
- while (*lns >= ' 0' && *lns <= ' 9' ) lineNr = lineNr * 10 + (*lns++ - ' 0' );
372
- lns++; // proceed to line number
373
- while (*lns >= ' 0' && *lns <= ' 9' ) linePos = linePos * 10 + (*lns++ - ' 0' );
374
- lns += 9 ; // proceed to error message
375
- eol = lns;
376
- while (*eol != ' \n ' && *eol > 0 ) eol++;
377
- *eol = 0 ;
378
- lineNr--; // we count from 0 instead of 1
379
- // adjust file and linenr based on include file data
380
- string errorFile = file;
381
- bool errorInInclude = false ;
382
- for (int i = Ninc - 1 ; i >= 0 ; i--)
383
- {
384
- if (lineNr > includes[i].end )
385
- {
386
- for (int j = 0 ; j <= i; j++) lineNr -= includes[j].end - includes[j].start ;
387
- break ;
388
- }
389
- else if (lineNr > includes[i].start )
390
- {
391
- errorFile = includes[i].file ;
392
- lineNr -= includes[i].start ;
393
- errorInInclude = true ;
394
- break ;
395
- }
396
- }
397
- if (!errorInInclude) lineNr -= vendorLines;
398
- // present error message
399
- char t[1024 ];
400
- sprintf ( t, " file %s, line %i, pos %i:\n %s" , errorFile.c_str (), lineNr + 1 , linePos, lns );
401
- FatalError ( t, " Build error" );
402
- }
403
- }
404
- else
405
- {
406
- // error string has unknown format; just dump it to a window
407
- log[2048 ] = 0 ; // truncate very long logs
408
- FatalError ( log, " Build error" );
409
- }
315
+ // find and display the first errormat; just dump it to a window
316
+ log[2048 ] = 0 ; // truncate very long logs
317
+ FatalError ( log, " Build error" );
410
318
}
411
319
kernel = clCreateKernel ( program, entryPoint, &error );
412
320
if (kernel == 0 ) FatalError ( " clCreateKernel failed: entry point not found." );
@@ -674,4 +582,4 @@ void Kernel::Run2D( const int2 count, const int2 lsize, cl_event* eventToWaitFor
674
582
{
675
583
CHECKCL ( error = clEnqueueNDRangeKernel ( queue, kernel, 2 , 0 , workSize, localSize, eventToWaitFor ? 1 : 0 , eventToWaitFor, eventToSet ) );
676
584
}
677
- }
585
+ }
0 commit comments