<div dir="ltr">Hey, <div class="gmail_extra"> <div class="gmail_quote">2013/12/9 Stephan Bergmann <<a href="mailto:sbergman@redhat.com" target="_blank">sbergman@redhat.com</a>> <blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"> I was trying to find out why on my MacBook recent master "make check" reproducibly fails CppunitTest_sc_opencl with <blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"> warn:sc.opencl:13647:1:sc/source/core/opencl/formulagroupcl.cxx:2737: CL_BUILD_PROGRAM_FAILURE, status -2, log "Error getting function data from server" </blockquote> [...] <blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"> warn:sc.opencl:13647:1:sc/source/core/opencl/formulagroupcl.cxx:2737: CL_BUILD_PROGRAM_FAILURE, status -2, log "Error getting function data from server" </blockquote> [...] <blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"> warn:sc.opencl:13647:1:sc/source/core/opencl/formulagroupcl.cxx:2737: CL_BUILD_PROGRAM_FAILURE, status -2, log "Error getting function data from server" </blockquote> [...] <blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"> warn:sc.opencl:13647:1:sc/source/core/opencl/formulagroupcl.cxx:2737: CL_BUILD_PROGRAM_FAILURE, status -2, log "Error getting function data from server" </blockquote> [...] <blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"> ##Failure Location unknown## : Error Test name: ScOpenclTest::testFinacialFormula uncaught exception of unknown type ##Failure Location unknown## : Error Test name: ScOpenclTest::testFinancialCoupdaysncFormula uncaught exception of unknown type ##Failure Location unknown## : Error Test name: ScOpenclTest::testFinacialYIELDFormula uncaught exception of unknown type ##Failure Location unknown## : Error Test name: ScOpenclTest::testFinacialPriceFormula uncaught exception of unknown type Failures !!! Run: 174 Failure total: 4 Failures: 0 Errors: 4 </blockquote> while just "make CppunitTest_sc_opencl" succeeds. Turns out the machine has two GPUs, a built-in Intel HD Graphics 4000 and an additional NVIDIA GeForce GT 650M, and that it apparently executes the OpenCL code on the latter (where it fails) only if there's enough graphics power demand, like in the case of a full (non-headless on Mac) "make check". (Though it looks like it doesn't suffice to disable "System Preferences... - Energy Saver - Automatic graphics switching" to make it use the NVIDIA GPU for the OpenCL code even in the "make CppunitTest_sc_opencl" scenario.) </blockquote><div> </div><div>The selection code is based on a heuristic. We compile a small test program and rank the difference OpenCL devices/platforms based on the performance. So it might be normal that you see different results depending on the workload, ... </div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"> And using the NVIDIA GPU to compile some of the OpenCL programs apparently causes a CL_BUILD_PROGRAM_FAILURE from clBuildProgram. However, the "Error getting function data from server" program build log does not give much of a clue what's going wrong. (And inducing something like a syntax error into the OpenCL program code would result in a much more verbose log pointing out the compilation error, so it can't be something trivial like that.) The OpenCL program code in the first of the four failing cases is <blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"> #pragma OPENCL EXTENSION cl_khr_fp64: enable int isNan(double a) { return isnan(a); } double fsum_count(double a, double b, __private int *p) { bool t = isNan(a); (*p) += t?0:1; return t?b:a+b; } double fsum(double a, double b) { return isNan(a)?b:a+b; } double legalize(double a, double b) { return isNan(a)?b:a;} double fsub(double a, double b) { return a-b; } double fdiv(double a, double b) { return a/b; } double strequal(unsigned a, unsigned b) { return (a==b)?1.0:0; } double mcw_fmin(double a, double b) { return fmin(a, b); } double mcw_fmax(double a, double b) { return fmax(a, b); } bool IsLeapYear( int n ); double coupdaybs( int nSettle,int nMat,int nFreq,int nBase); double coupdays(int nSettle,int nMat,int nFreq,int nBase); double coupdaysnc( int nSettle,int nMat,int nFreq,int nBase); double coupnum( int nSettle,int nMat,int nFreq,int nBase); double getPrice_(int nSettle, int nMat, double fRate, double fYield, double fRedemp, int nFreq, int nBase ); double getYield_( int nNullDate, int nSettle, int nMat, double fCoup,double fPrice,double fRedemp, int nFreq, int nBase); int DateToDays( int nDay, int nMonth, int nYear ); int DaysInMonth( int nMonth, int nYear ); int GetDaysInYear( int nNullDate, int nDate, int nMode ); int GetDaysInYears( int nYear1, int nYear2 ); int GetNullDate(void); int getDaysInMonthRange( int nFrom, int nTo,int b30Days,int year); int getDaysInYearRange( int nFrom, int nTo,int b30Days ); int getDiff(int rFrom,int rTo,int fDay,int fMonth,int fYear,int fbLastDayMode,int fbLastDay,int fb30Days,int fbUSMode,int fnDay,int tDay,int tMonth,int tYear,int tbLastDayMode,int tbLastDay,int tb30Days,int tbUSMode,int tnDay); int lcl_Getcoupdaybs(int nNullDate,int nSettle, int nMat,int nFreq,int nBase); int lcl_Getcoupdays(int nNullDate,int nSettle, int nMat,int nFreq,int nBase); int lcl_Getcoupnum(int nNullDate,int nSettle, int nMat,int nFreq); void DaysToDate( int nDays, int *rDay, int* rMonth, int* rYear ); void ScaDate( int nNullDate, int nDate, int nBase,int *nOrigDay, int *nMonth,int *nYear,int *bLastDayMode,int *bLastDay,int *b30Days,int *bUSMode,int *nDay); void addMonths(int b30Days,int bLastDay,int *nDay,int nOrigDay,int *nMonth,int nMonthCount,int *year); bool IsLeapYear( int n ) { return ( (( ( n % 4 ) == 0 ) && ( ( n % 100 ) != 0)) || ( ( n % 400 ) == 0 ) ); } double coupdaybs( int nSettle,int nMat,int nFreq,int nBase) { int nNullDate=GetNullDate(); return lcl_Getcoupdaybs(nNullDate, nSettle, nMat,nFreq, nBase); } double coupdays(int nSettle,int nMat,int nFreq,int nBase) { int nNullDate=GetNullDate(); if( nBase == 1 ) return lcl_Getcoupdays(nNullDate, nSettle, nMat,nFreq, nBase); else return (double)GetDaysInYear(0,0,nBase)/nFreq; } double coupdaysnc( int nSettle,int nMat,int nFreq,int nBase) { int nNullDate=GetNullDate(); if((nBase != 0) && (nBase != 4)) { int aDate = nMat; int rDay=0,rMonth=0, rYear=0;int mDay=0,mMonth=0, mYear=0;int sDay=0,sMonth=0, sYear=0; int rbLastDayMode=0, rbLastDay=0,rb30Days=0,rbUSMode=0,rnDay=0; int sbLastDayMode=0, sbLastDay=0,sb30Days=0,sbUSMode=0,snDay=0; ScaDate( nNullDate,aDate,nBase,&rDay,&rMonth,&rYear,&rbLastDayMode,&rbLastDay,&rb30Days,&rbUSMode,&rnDay); ScaDate( nNullDate,nSettle,nBase,&sDay,&sMonth,&sYear,&sbLastDayMode,&sbLastDay,&sb30Days,&sbUSMode,&snDay); rYear= sYear; nSettle=nSettle+nNullDate; aDate=DateToDays( rDay,rMonth,rYear ); if( aDate > nSettle ) { rYear-= 1; aDate=DateToDays( rDay,rMonth,rYear ); } while(aDate <= nSettle ) { addMonths(rb30Days,rbLastDay,&rnDay,rDay,&rMonth,12/nFreq,&rYear); aDate=DateToDays( rDay,rMonth,rYear ); } return getDiff( nSettle, aDate, sDay, sMonth, sYear, sbLastDayMode, sbLastDay, sb30Days, sbUSMode, snDay, rDay, rMonth, rYear, rbLastDayMode, rbLastDay, rb30Days, rbUSMode, rnDay); } else return coupdays(nSettle,nMat,nFreq,nBase)- coupdaybs( nSettle,nMat,nFreq,nBase); } double coupnum( int nSettle,int nMat,int nFreq,int nBase) { int nNullDate=GetNullDate(); return lcl_Getcoupnum(nNullDate,nSettle,nMat,nFreq); } double getPrice_(int nSettle, int nMat, double fRate, double fYield, double fRedemp, int nFreq, int nBase ) { double fFreq = nFreq; double fE = coupdays( nSettle, nMat, nFreq, nBase ); double fDSC_E = coupdaysnc( nSettle, nMat, nFreq, nBase ) / fE; double fN = coupnum( nSettle, nMat, nFreq, nBase ); double fA = coupdaybs( nSettle, nMat, nFreq, nBase ); double fRet = fRedemp / ( pow( 1.0 + fYield / fFreq, fN - 1.0 + fDSC_E ) ); fRet -= 100.0 * fRate / fFreq * fA / fE; double fT1 = 100.0 * fRate / fFreq; double fT2 = 1.0 + fYield / fFreq; for( double fK = 0.0 ; fK < fN ; fK+=1.0 ) fRet += fT1 / pow( fT2, fK + fDSC_E ); return fRet; } double getYield_( int nNullDate, int nSettle, int nMat, double fCoup,double fPrice,double fRedemp, int nFreq, int nBase ) { double fRate = fCoup; double fPriceN = 0.0; double fYield1 = 0.0; double fYield2 = 1.0; double fPrice1 = getPrice_(nSettle, nMat, fRate, fYield1, fRedemp, nFreq, nBase ); double fPrice2 = getPrice_(nSettle, nMat, fRate, fYield2, fRedemp, nFreq, nBase ); double fYieldN = ( fYield2 - fYield1 ) * 0.5; for( unsigned int nIter = 0 ; nIter < 100 && fPriceN != fPrice ; nIter++ ) { fPriceN = getPrice_(nSettle, nMat, fRate, fYieldN, fRedemp, nFreq, nBase ); if( fPrice == fPrice1 ) return fYield1; else if( fPrice == fPrice2 ) return fYield2; else if( fPrice == fPriceN ) return fYieldN; else if( fPrice < fPrice2 ) { fYield2 *= 2.0; fPrice2 = getPrice_(nSettle, nMat, fRate, fYield2, fRedemp, nFreq, nBase ); fYieldN = ( fYield2 - fYield1 ) * 0.5; } else { if( fPrice < fPriceN ) { fYield1 = fYieldN; fPrice1 = fPriceN; } else { fYield2 = fYieldN; fPrice2 = fPriceN; } fYieldN = fYield2 - ( fYield2 - fYield1 ) * ( ( fPrice - fPrice2 ) / ( fPrice1 - fPrice2 ) ); } } return fYieldN; } int DateToDays( int nDay, int nMonth, int nYear ) { int nDays = ((int)nYear-1) * 365; nDays += ((nYear-1) / 4) - ((nYear-1) / 100) + ((nYear-1) / 400); for( int i = 1; i < nMonth; i++ ) nDays += DaysInMonth(i,nYear); nDays += nDay; return nDays; } int DaysInMonth( int nMonth, int nYear ) { int aDaysInMonth[12] = { 31, 28, 31, 30, 31, 30, 31, 31, 30, 31, 30, 31 }; if ( nMonth != 2 ) return aDaysInMonth[nMonth-1]; else { if ( IsLeapYear(nYear) ) return aDaysInMonth[nMonth-1] + 1; else return aDaysInMonth[nMonth-1]; } } int GetDaysInYear( int nNullDate, int nDate, int nMode ) { switch( nMode ) { case 0: case 2: case 4: return 360; case 1: { int nD=0, nM=0, nY=0; nDate += nNullDate; DaysToDate( nDate, &nD, &nM, &nY ); return IsLeapYear( nY )? 366 : 365; } case 3: return 365; } } int GetDaysInYears( int nYear1, int nYear2 ) { int nLeaps = 0; for( int n = nYear1 ; n <= nYear2 ; n++ ) { if( IsLeapYear( n ) ) nLeaps++; } int nSum = 1; nSum += nYear2; nSum -= nYear1; nSum *= 365; nSum += nLeaps; return nSum; } int GetNullDate(void) { return DateToDays(30,12,1899 ); } int getDaysInMonthRange( int nFrom, int nTo,int b30Days,int year) { if( nFrom > nTo ) return 0; int nRet = 0; if( b30Days ) nRet = (nTo - nFrom + 1) * 30; else { for( int nMonthIx = nFrom; nMonthIx <= nTo; ++nMonthIx ) nRet += b30Days ? 30 : DaysInMonth( nMonthIx, year ); } return nRet; } int getDaysInYearRange( int nFrom, int nTo,int b30Days ) { if( nFrom > nTo ) return 0; return b30Days ? ((nTo - nFrom + 1) * 360) : GetDaysInYears( nFrom, nTo); } int getDiff(int rFrom,int rTo,int fDay,int fMonth,int fYear,int fbLastDayMode,int fbLastDay,int fb30Days,int fbUSMode,int fnDay,int tDay,int tMonth,int tYear,int tbLastDayMode,int tbLastDay,int tb30Days,int tbUSMode,int tnDay) { if(rFrom>rTo) { int d=fDay;fDay=tDay;tDay=d; int m=fMonth;fMonth=tMonth;tMonth=m; int y=fYear;fYear=tYear;tYear=y; int a=fbLastDayMode;fbLastDayMode=tbLastDayMode;tbLastDayMode=a; int b=fbLastDay;fbLastDay=tbLastDay;tbLastDay=b; int c=fb30Days;fb30Days=tb30Days;tb30Days=c; int e=fbUSMode;fbUSMode=tbUSMode;tbUSMode=e; int f=fnDay;fnDay=tnDay;tnDay=f; } int nDiff=0; if( tb30Days ) { if( tbUSMode ) { if( ((fMonth == 2) || (fnDay < 30)) && (tDay == 31) ) tnDay = 31; else if( (tMonth == 2) && tbLastDay ) tnDay = DaysInMonth( 2, tYear ); } else { if( (fMonth == 2) && (fnDay == 30) ) fnDay = DaysInMonth( 2, fYear ); if( (tMonth == 2) && (tnDay == 30) ) tnDay = DaysInMonth( 2, tYear ); } } if( (fYear < tYear) || ((fYear == tYear) && (fMonth < tMonth)) ) { nDiff = (fb30Days? 30:DaysInMonth(fMonth,fYear) )- fnDay + 1; fDay = fnDay = 1; fbLastDay = 0; addMonths(fb30Days,fbLastDay,&fnDay,fDay,&fMonth,1,&fYear); if( fYear < tYear ) { nDiff += getDaysInMonthRange( fMonth, 12,fb30Days,fYear); addMonths(fb30Days,fbLastDay,&fnDay,fDay,&fMonth,13-fMonth,&fYear); nDiff += getDaysInYearRange( fYear, tYear - 1,fb30Days); fYear+=tYear - fYear; } nDiff += getDaysInMonthRange(fMonth, tMonth - 1,fb30Days ,fYear ); addMonths(fb30Days,fbLastDay,&fnDay,fDay,&fMonth,tMonth-fMonth,&fYear); } nDiff += tnDay - fnDay; return nDiff > 0 ? nDiff : 0; } int lcl_Getcoupdaybs(int nNullDate,int nSettle, int nMat,int nFreq,int nBase) { int aDate = nMat; int rDay=0,rMonth=0, rYear=0;int mDay=0,mMonth=0, mYear=0;int sDay=0,sMonth=0, sYear=0; int rbLastDayMode=0, rbLastDay=0,rb30Days=0,rbUSMode=0,rnDay=0; int sbLastDayMode=0, sbLastDay=0,sb30Days=0,sbUSMode=0,snDay=0; ScaDate( nNullDate,aDate,nBase,&rDay,&rMonth,&rYear,&rbLastDayMode,&rbLastDay,&rb30Days,&rbUSMode,&rnDay); ScaDate( nNullDate,nSettle,nBase,&sDay,&sMonth,&sYear,&sbLastDayMode,&sbLastDay,&sb30Days,&sbUSMode,&snDay); rYear= sYear; nSettle=nSettle+nNullDate; aDate=DateToDays( rDay,rMonth,rYear ); if( aDate < nSettle ) { rYear+= 1; aDate=DateToDays( rDay,rMonth,rYear ); } while(aDate > nSettle ) { addMonths(rb30Days,rbLastDay,&rnDay,rDay,&rMonth,-1*(12/nFreq),&rYear); aDate=DateToDays( rDay,rMonth,rYear ); } return getDiff( aDate, nSettle, rDay, rMonth, rYear, rbLastDayMode, rbLastDay, rb30Days, rbUSMode, rnDay, sDay, sMonth, sYear, sbLastDayMode,sbLastDay, sb30Days, sbUSMode, snDay); } int lcl_Getcoupdays(int nNullDate,int nSettle, int nMat,int nFreq,int nBase) { int aDate = nMat; int rDay=0,rMonth=0, rYear=0;int mDay=0,mMonth=0, mYear=0;int sDay=0,sMonth=0, sYear=0; int rbLastDayMode=0, rbLastDay=0,rb30Days=0,rbUSMode=0,rnDay=0; int sbLastDayMode=0, sbLastDay=0,sb30Days=0,sbUSMode=0,snDay=0; ScaDate( nNullDate,aDate,nBase,&rDay,&rMonth,&rYear,&rbLastDayMode,&rbLastDay,&rb30Days,&rbUSMode,&rnDay); ScaDate( nNullDate,nSettle,nBase,&sDay,&sMonth,&sYear,&sbLastDayMode,&sbLastDay,&sb30Days,&sbUSMode,&snDay); rYear= sYear; nSettle=nSettle+nNullDate; aDate=DateToDays( rDay,rMonth,rYear ); if( aDate < nSettle ) { rYear+= 1; aDate=DateToDays( rDay,rMonth,rYear ); } while(aDate > nSettle ) { addMonths(rb30Days,rbLastDay,&rnDay,rDay,&rMonth,-1*(12/nFreq),&rYear); aDate=DateToDays( rDay,rMonth,rYear ); } int aNextDate=aDate;int aDay=rDay,aMonth=rMonth, aYear=rYear; int abLastDayMode=rbLastDayMode, abLastDay=rbLastDay,ab30Days=rb30Days,abUSMode=rbUSMode,anDay=rnDay; addMonths(ab30Days,abLastDay,&anDay,aDay,&aMonth,12/nFreq,&aYear); return getDiff( aDate, aNextDate, rDay, rMonth, rYear, rbLastDayMode, rbLastDay, rb30Days, rbUSMode, rnDay, aDay, aMonth, aYear, abLastDayMode,abLastDay, ab30Days, abUSMode, anDay); } int lcl_Getcoupnum(int nNullDate,int nSettle, int nMat,int nFreq) { int aDate = nMat;int rDay=0,rMonth=0, rYear=0;int mDay=0,mMonth=0, mYear=0; int sDay=0,sMonth=0, sYear=0; DaysToDate(aDate+nNullDate,&rDay, &rMonth, &rYear ); DaysToDate(nMat+nNullDate,&mDay, &mMonth, &mYear ); DaysToDate(nSettle+nNullDate,&sDay, &sMonth, &sYear ); rYear= sYear; nSettle=nSettle+nNullDate; aDate=DateToDays( rDay,rMonth,rYear ); if( aDate < nSettle ) rYear+= 1; int d=DateToDays( rDay,rMonth,rYear ); int nMonthCount=-1*(12 / nFreq); while(d > nSettle ) { int nNewMonth = nMonthCount + rMonth; if( nNewMonth > 12 ) { --nNewMonth; rYear+=nNewMonth / 12; rMonth = nNewMonth % 12 + 1; } else if( nNewMonth < 1 ) { rYear+= nNewMonth / 12 - 1; rMonth = nNewMonth % 12 + 12; } else rMonth = nNewMonth; d=DateToDays( rDay,rMonth,rYear ); } int n=(mYear-rYear)*12+mMonth-rMonth; n=n*nFreq/12; return n; } void DaysToDate( int nDays, int *rDay, int* rMonth, int* rYear ) { int nTempDays; int i = 0; bool bCalc; do { nTempDays = nDays; *rYear = (int)((nTempDays / 365) - i); nTempDays -= ((int) *rYear -1) * 365; nTempDays -= (( *rYear -1) / 4) - (( *rYear -1) / 100) +((*rYear -1) / 400); bCalc = false; if ( nTempDays < 1 ) { i++; bCalc = true; } else { if ( nTempDays > 365 ) { if ( (nTempDays != 366) || !IsLeapYear( *rYear ) ) { i--; bCalc = true; } } } } while ( bCalc ); *rMonth = 1; while ( (int)nTempDays > DaysInMonth( *rMonth, *rYear ) ) { nTempDays -= DaysInMonth( *rMonth, *rYear ); *rMonth+=1; } *rDay = (int)nTempDays; } void ScaDate( int nNullDate, int nDate, int nBase,int *nOrigDay, int *nMonth,int *nYear,int *bLastDayMode,int *bLastDay,int *b30Days,int *bUSMode,int *nDay) { DaysToDate( nNullDate + nDate, nOrigDay, nMonth, nYear ); *bLastDayMode = (nBase != 5); *bLastDay = (*nOrigDay >= DaysInMonth( *nMonth, *nYear )); *b30Days = (nBase == 0) || (nBase == 4); *bUSMode = (nBase == 0); if( *b30Days) { *nDay = min( *nOrigDay, 30); if( *bLastDay || (*nDay >=DaysInMonth( *nMonth, *nYear )) ) *nDay = 30; } else { int nLastDay = DaysInMonth( *nMonth, *nYear ); *nDay = *bLastDay ? nLastDay : min( *nOrigDay, nLastDay ); } } void addMonths(int b30Days,int bLastDay,int *nDay,int nOrigDay,int *nMonth,int nMonthCount,int *year) { int nNewMonth = nMonthCount + *nMonth; if( nNewMonth > 12 ) { --nNewMonth; *year+=nNewMonth / 12 ; *nMonth = ( nNewMonth % 12 ) + 1; } else if( nNewMonth < 1 ) { *year+= nNewMonth / 12 - 1 ; *nMonth = nNewMonth % 12 + 12 ; } else *nMonth = nNewMonth ; if( b30Days ) { *nDay = min( nOrigDay, 30); if( bLastDay || (*nDay >= DaysInMonth( *nMonth, *year )) ) *nDay = 30; } else { int nLastDay = DaysInMonth( *nMonth, *year ); *nDay = bLastDay ? nLastDay : min( nOrigDay, nLastDay ); } } double tmp0_0_Yield(__global double *tmp0_0_0,__global double *tmp0_0_1,__global double *tmp0_0_2,__global double *tmp0_0_3,__global double *tmp0_0_4,__global double *tmp0_0_5,__global double *tmp0_0_6) { double tmp = 0; int gid0 = get_global_id(0); double tmp000; double tmp001; double tmp002; double tmp003; double tmp004; double tmp005; double tmp006; int buffer_tmp000_len = 11; int buffer_tmp001_len = 11; int buffer_tmp002_len = 11; int buffer_tmp003_len = 11; int buffer_tmp004_len = 11; int buffer_tmp005_len = 11; int buffer_tmp006_len = 11; if(gid0>=buffer_tmp000_len || isNan((gid0 < 11?tmp0_0_0[gid0]:NAN))) tmp000 = 0; else tmp000 = (gid0 < 11?tmp0_0_0[gid0]:NAN); if(gid0>=buffer_tmp001_len || isNan((gid0 < 11?tmp0_0_1[gid0]:NAN))) tmp001 = 0; else tmp001 = (gid0 < 11?tmp0_0_1[gid0]:NAN); if(gid0>=buffer_tmp002_len || isNan((gid0 < 11?tmp0_0_2[gid0]:NAN))) tmp002 = 0; else tmp002 = (gid0 < 11?tmp0_0_2[gid0]:NAN); if(gid0>=buffer_tmp003_len || isNan((gid0 < 11?tmp0_0_3[gid0]:NAN))) tmp003 = 0; else tmp003 = (gid0 < 11?tmp0_0_3[gid0]:NAN); if(gid0>=buffer_tmp004_len || isNan((gid0 < 11?tmp0_0_4[gid0]:NAN))) tmp004 = 0; else tmp004 = (gid0 < 11?tmp0_0_4[gid0]:NAN); if(gid0>=buffer_tmp005_len || isNan((gid0 < 11?tmp0_0_5[gid0]:NAN))) tmp005 = 0; else tmp005 = (gid0 < 11?tmp0_0_5[gid0]:NAN); if(gid0>=buffer_tmp006_len || isNan((gid0 < 11?tmp0_0_6[gid0]:NAN))) tmp006 = 0; else tmp006 = (gid0 < 11?tmp0_0_6[gid0]:NAN); tmp = getYield_(GetNullDate(),tmp000,tmp001,tmp002,tmp003,tmp004,tmp005,tmp006); return tmp; } double tmp0_nop(__global double *tmp0_0_0, __global double *tmp0_0_1, __global double *tmp0_0_2, __global double *tmp0_0_3, __global double *tmp0_0_4, __global double *tmp0_0_5, __global double *tmp0_0_6) { double tmp = 0; int gid0 = get_global_id(0); double tmpBottom; tmp = tmp0_0_Yield(tmp0_0_0, tmp0_0_1, tmp0_0_2, tmp0_0_3, tmp0_0_4, tmp0_0_5, tmp0_0_6); return tmp; } __kernel void DynamicKernel_nop_Yield(__global double *result, __global double *tmp0_0_0, __global double *tmp0_0_1, __global double *tmp0_0_2, __global double *tmp0_0_3, __global double *tmp0_0_4, __global double *tmp0_0_5, __global double *tmp0_0_6) { int gid0 = get_global_id(0); result[gid0] = tmp0_nop(tmp0_0_0,tmp0_0_1,tmp0_0_2,tmp0_0_3,tmp0_0_4,tmp0_0_5,tmp0_0_6); } </blockquote> and some experimentation revealed that replacing the call <blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"> return getDiff( nSettle, aDate, sDay, sMonth, sYear, sbLastDayMode, sbLastDay, sb30Days, sbUSMode, snDay, rDay, rMonth, rYear, rbLastDayMode, rbLastDay, rb30Days, rbUSMode, rnDay); </blockquote> towards the end of the definition of coupdaysnc with <blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"> return 0.0; </blockquote> prevents the build failure (though the test now trivially fails with a wrong result, of course). However, I have no idea what's wrong there. OpenCL forbids recursive functions, but this doesn't look like it is (indirectly) recursive. Maybe it is a problem of the NVIDIA GPU rather than our code? </blockquote><div> </div><div>It looks like another broken OpenCL compiler. So basically I'd need from you the output of clinfo if available or otherwise the generated file of <a href="http://dev-www.libreoffice.org/opencl/Linux/">http://dev-www.libreoffice.org/opencl/Linux/</a> which hopefully runs on Mac as well. I'll use these information to block broken compilers. However it will take me some time to implement that because "OpenCL device selection 2.0" no longer uses the code for disabling devices. </div><div>Regards, </div><div>Markus </div></div></div></div>