CudaLCP.inl

Go to the documentation of this file.
00001 /******************************************************************************
00002 *       SOFA, Simulation Open-Framework Architecture, version 1.0 beta 3      *
00003 *                (c) 2006-2008 MGH, INRIA, USTL, UJF, CNRS                    *
00004 *                                                                             *
00005 * This library is free software; you can redistribute it and/or modify it     *
00006 * under the terms of the GNU Lesser General Public License as published by    *
00007 * the Free Software Foundation; either version 2.1 of the License, or (at     *
00008 * your option) any later version.                                             *
00009 *                                                                             *
00010 * This library is distributed in the hope that it will be useful, but WITHOUT *
00011 * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or       *
00012 * FITNESS FOR A PARTICULAR PURPOSE. See the GNU Lesser General Public License *
00013 * for more details.                                                           *
00014 *                                                                             *
00015 * You should have received a copy of the GNU Lesser General Public License    *
00016 * along with this library; if not, write to the Free Software Foundation,     *
00017 * Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301 USA.          *
00018 *******************************************************************************
00019 *                               SOFA :: Modules                               *
00020 *                                                                             *
00021 * Authors: The SOFA Team and external contributors (see Authors.txt)          *
00022 *                                                                             *
00023 * Contact information: contact@sofa-framework.org                             *
00024 ******************************************************************************/
00025 int CudaLCP_MultVector_ResSize(unsigned int dim) {
00026     return (dim+BSIZE-1)/BSIZE;
00027 }
00028 
00029 int CudaNLCP_MultVector_ResSize(unsigned int dim) {
00030     return (dim+MBSIZE-1)/MBSIZE;
00031 }
00032 
00034 
00035 void CudaLCP_MultVectorf(int dim,int index, const void * m,const void * f,void * r) {
00036     dim3 threads(BSIZE,1);
00037     dim3 grid((dim+BSIZE-1)/BSIZE,1);
00038 
00039     CudaLCP_MultVector_kernel<float><<< grid, threads, threads.x*sizeof(float) >>>(dim, index, (const float*)m, (const float*)f, (float*)r, BSIZE/2);
00040 }
00041 void CudaLCP_MultVectord(int dim,int index, const void * m,const void * f,void * r) {
00042 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00043     myprintf("CUDA ERROR: double precision not supported.\n");
00044 #else
00045     dim3 threads(BSIZE,1);
00046     dim3 grid((dim+BSIZE-1)/BSIZE,1);
00047 
00048     CudaLCP_MultVector_kernel<double><<< grid, threads, threads.x*sizeof(double) >>>(dim, index, (const double*)m, (const double*)f, (double*)r, BSIZE/2);
00049 #endif
00050 }
00051 
00052 void CudaLCP_ComputeErrorf(int compteur2,int sizeTmp, const void * tmp, const void * M,const void * q,void * f,void * error) {
00053     dim3 threads(sizeTmp,1);
00054     dim3 grid(1,1);
00055     int offset;
00056     if (sizeTmp==1) offset = 0;
00057     else {
00058         offset = 1;
00059         while (offset*2 < sizeTmp)  offset *= 2;
00060     }
00061 
00062     CudaLCP_ComputeError_kernel<float><<< grid, threads, threads.x*sizeof(float) >>>(compteur2,(const float*)tmp, (const float*)M,(const float*)q, (float*)f,(float*)error,offset);
00063 }
00064 void CudaLCP_ComputeErrord(int compteur2,int sizeTmp, const void * tmp, const void * M,const void * q,void * f,void * error) {
00065 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00066     myprintf("CUDA ERROR: double precision not supported.\n");
00067 #else
00068     dim3 threads(sizeTmp,1);
00069     dim3 grid(1,1);
00070     int offset;
00071     if (sizeTmp==1) offset = 0;
00072     else {
00073         offset = 1;
00074         while (offset*2 < sizeTmp)  offset *= 2;
00075     }
00076 
00077     CudaLCP_ComputeError_kernel<double><<< grid, threads, threads.x*sizeof(double) >>>(compteur2,(const double*)tmp, (const double*)M,(const double*)q, (double*)f,(double*)error,offset);
00078 #endif
00079 }
00080 
00082 
00083 /*
00084 for (int i=0;i<dim;i++) {
00085     cuda_res[i] = cuda_q[i];
00086     for (int j=0;j<dim;j++) {
00087         if (j>i) cuda_res[i] += cuda_M[i][j] * cuda_f[j];
00088 }
00089 }
00090 */
00091 void CudaLCP_MultIndepf(int dim,const void * m,int pM,const void * f,void * tmp,int pTmp) {
00092     dim3 threads(BSIZE,1);  
00093     dim3 grid((dim+BSIZE-1)/BSIZE,dim);
00094 
00095     CudaLCP_MultIndep_kernel<float><<< grid, threads,0>>>(dim, (const float*)m,pM,(const float*)f,(float*)tmp,pTmp,BSIZE/2);
00096 }
00097 void CudaLCP_MultIndepd(int dim,const void * m,int pM,const void * f,void * tmp,int pTmp) {
00098 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00099     myprintf("CUDA ERROR: double precision not supported.\n");
00100 #else
00101     dim3 threads(BSIZE,1);  
00102     dim3 grid((dim+BSIZE-1)/BSIZE,dim);
00103 
00104     CudaLCP_MultIndep_kernel<double><<< grid, threads,0>>>(dim, (const double*)m,pM,(const double*)f,(double*)tmp,pTmp,BSIZE/2);
00105 #endif
00106 }
00107 
00108 /*
00109 for (int j=0;j<dim;j++) {
00110     cuda_res[j] = cuda_q[i];
00111     for (int i=0;i<tmpsize;i++) {
00112         cuda_res[j] += cuda_tmp[i][j];
00113 }
00114 }
00115 */
00116 void CudaLCP_AddIndepf(int dim,int tmpsize,const void * q,const void * tmp,int pTmp,void * res) {
00117     dim3 threads(tmpsize,1);    
00118     dim3 grid(dim,1);
00119 
00120     CudaLCP_AddIndep_kernel<float><<< grid, threads,threads.x*sizeof(float)>>>(dim,tmpsize,(const float*)q,(const float*)tmp,pTmp,(float*)res,tmpsize);
00121 }
00122 void CudaLCP_AddIndepd(int dim,int tmpsize,const void * q,const void * tmp,int pTmp,void * res) {
00123 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00124     myprintf("CUDA ERROR: double precision not supported.\n");
00125 #else
00126     dim3 threads(tmpsize,1);    
00127     dim3 grid(dim,1);
00128 
00129     CudaLCP_AddIndep_kernel<double><<< grid, threads,threads.x*sizeof(double)>>>(dim,tmpsize,(const double*)q,(const double*)tmp,pTmp,(double*)res,tmpsize);
00130 #endif
00131 }
00132    
00133 void CudaLCP_AddIndepAndUpdatef(int dim,int tmpsize,const void * m,const void * q,const void * tmp,int pTmp,void * f,void * res,void * err) {
00134     dim3 threads(tmpsize,1);    
00135     dim3 grid(dim,1);
00136     
00137     CudaLCP_AddIndepAndUpdate_kernel<float><<< grid, threads,threads.x*sizeof(float)>>>(dim,tmpsize,(const float*)m,(const float*)q,(const float*)tmp,pTmp,(float*)f,(float*)res,(float*)err,tmpsize);
00138 }
00139 void CudaLCP_AddIndepAndUpdated(int dim,int tmpsize,const void * m,const void * q,const void * tmp,int pTmp,void * f,void * res,void * err) {
00140 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00141     myprintf("CUDA ERROR: double precision not supported.\n");
00142 #else
00143     dim3 threads(tmpsize,1);    
00144     dim3 grid(dim,1); 
00145     
00146     CudaLCP_AddIndepAndUpdate_kernel<double><<< grid, threads,threads.x*sizeof(double)>>>(dim,tmpsize,(const double*)m,(const double*)q,(const double*)tmp,pTmp,(double*)f,(double*)res,(double*)err,tmpsize);
00147 #endif
00148 } 
00149 
00150 /*
00151 float f_1 = cuda_f[compteur2];
00152 if (cuda_res[compteur2]<0) cuda_f[compteur2] = -cuda_res[compteur2]/cuda_M[compteur2][compteur2];
00153 else cuda_f[compteur2]=0.0;
00154 cuda_err[0] += fabs(cuda_M[compteur2][compteur2] * (cuda_f[compteur2] - f_1));
00155 cuda_err[1] = f_1;
00156 
00157 for (int compteur3=0;compteur3<dim;compteur3++) {
00158     if (compteur3!=compteur2) cuda_res[compteur3] += cuda_M[compteur3][compteur2] * (cuda_f[compteur2] - cuda_err[1]);
00159 }   
00160 */      
00161 void CudaLCP_ComputeNextIter_V2f(int dim,int compteur2,const void * m,int mP,const void * q,void * f,void * err,void * res) {
00162     dim3 threads(BSIZE,1);
00163     dim3 grid((dim+BSIZE-1)/BSIZE,1);
00164     
00165     CudaLCP_ComputeNextIter_kernel_V2<float><<< grid, threads,0 >>>(dim,compteur2,(const float*) m,mP,(const float *) q,(float*)f, (float*) err,(float *)res);
00166 }    
00167 void CudaLCP_ComputeNextIter_V2d(int dim,int compteur2,const void * m,int mP,const void * q,void * f,void * err,void * res) {
00168 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00169     myprintf("CUDA ERROR: double precision not supported.\n");
00170 #else
00171     dim3 threads(BSIZE,1);
00172     dim3 grid((dim+BSIZE-1)/BSIZE,1);
00173     
00174     CudaLCP_ComputeNextIter_kernel_V2<double><<< grid, threads,0 >>>(dim,compteur2,(const double*) m,mP,(const double *) q,(double*)f, (double*) err,(double *)res);
00175 #endif
00176 }   
00177 
00179 
00180 void CudaLCP_ComputeNextIter_V3_OneKernelf(int dim,const void * m,int mP,const void * q,void * f,void * err,void * res) {
00181     dim3 threads(dim,1);
00182     dim3 grid(1,1);
00183     
00184     CudaLCP_ComputeNextIter_V3_OneKernel_kernel<float><<< grid, threads,0 >>>(dim,(const float*) m,mP,(const float *) q,(float*)f, (float*) err,(float *)res);
00185 }   
00186 void CudaLCP_ComputeNextIter_V3_OneKerneld(int dim,const void * m,int mP,const void * q,void * f,void * err,void * res) {
00187 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00188     myprintf("CUDA ERROR: double precision not supported.\n");
00189 #else
00190     dim3 threads(dim,1);
00191     dim3 grid(1,1);
00192     
00193     CudaLCP_ComputeNextIter_V3_OneKernel_kernel<double><<< grid, threads,0 >>>(dim,(const double*) m,mP,(const double *) q,(double*)f, (double*) err,(double *)res);
00194 #endif
00195 }   
00196 
00198 
00199 /*
00200 for (int block=0; block<d; block++) {   
00201     int compteur2 = debutblock+block;   
00202     
00203     float f_1 = cuda_f[compteur2];
00204     
00205     float f_2;
00206     if (cuda_res[compteur2]<0) f_2 = -cuda_res[compteur2]/cuda_M[compteur2][compteur2];
00207     else f_2=0.0;
00208                     
00209     cuda_f[compteur2] = f_2;
00210     cuda_res[compteur2] = cuda_q[compteur2];    
00211     cuda_err[0] += fabs(cuda_M[compteur2][compteur2] * (f_2 - f_1));
00212                     
00213     for (int compteur3=debutblock;compteur3<debutblock+d;compteur3++) {
00214         if (compteur3!=compteur2) cuda_res[compteur3] += cuda_M[compteur2][compteur3] * f_2;
00215 }   
00216 }
00217 */
00218 void CudaLCP_ComputeNextIter_V4_DepKernelf(int d,int debutblock,const void * m,int mP,const void * q,void * f,void * err,void * res) {
00219     dim3 threads(BSIZE,1);
00220     dim3 grid(1,1);
00221         
00222     CudaLCP_ComputeNextIter_V4_DepKernel_kernel<float><<< grid, threads,0>>>(d,debutblock,(const float*) m,mP,(const float *) q,(float*)f, (float*) err,(float *)res);
00223 }
00224 void CudaLCP_ComputeNextIter_V4_DepKerneld(int d,int debutblock,const void * m,int mP,const void * q,void * f,void * err,void * res) {
00225 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00226     myprintf("CUDA ERROR: double precision not supported.\n");
00227 #else
00228     dim3 threads(BSIZE,1);
00229     dim3 grid(1,1);
00230         
00231     CudaLCP_ComputeNextIter_V4_DepKernel_kernel<double><<< grid, threads,0>>>(d,debutblock,(const double*) m,mP,(const double *) q,(double*)f, (double*) err,(double *)res);
00232 #endif
00233 }
00234 
00235 /*
00236 for (int compteur4=0;compteur4<dim-d;compteur4++) {
00237     int ligne = compteur4;
00238     if (compteur4>=debutblock) ligne+=BSIZE;
00239     
00240     for (int k=debutblock;k<debutblock+d;k++) {
00241         cuda_res[ligne]+= cuda_M[k][ligne] * cuda_f[k];
00242 }
00243 }   
00244 */
00245 void CudaLCP_ComputeNextIter_V4_InDepKernelf(int dim,int d,int debutBlock,const void * m,int mP,const void * f,void * res) {
00246     dim3 threads(BSIZE,1);
00247     dim3 grid(1,dim-d);
00248     
00249     CudaLCP_ComputeNextIter_V4_InDepKernel_kernel<float><<< grid, threads,0>>>(dim,debutBlock,(const float*) m,mP,(const float*)f, (float *)res,BSIZE/2);
00250 }
00251 void CudaLCP_ComputeNextIter_V4_InDepKerneld(int dim,int d,int debutBlock,const void * m,int mP,const void * f,void * res) {
00252 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00253     myprintf("CUDA ERROR: double precision not supported.\n");
00254 #else
00255     dim3 threads(BSIZE,1);
00256     dim3 grid(1,dim-d);
00257     
00258     CudaLCP_ComputeNextIter_V4_InDepKernel_kernel<double><<< grid, threads,0>>>(dim,debutBlock,(const double*) m,mP,(const double*)f, (double *)res,BSIZE/2);
00259 #endif
00260 }
00261 
00263 
00264 /*
00265 for (int compteur2=0; compteur2<BSIZE;compteur2++) {    //calcule le bloc en haut a gauche pour lancer le calcul
00266     float f_2;
00267     if (cuda_res[compteur2]<0) f_2 = -cuda_res[compteur2]/cuda_M[compteur2][compteur2];
00268     else f_2=0.0;
00269                         
00270     cuda_f[compteur2] = f_2;
00271     cuda_res[compteur2] = cuda_q[compteur2];
00272                         
00273     for (int compteur3=0;compteur3<BSIZE;compteur3++) {
00274         if (compteur3!=compteur2) cuda_res[compteur3] += cuda_M[compteur3][compteur2] * f_2;
00275 }
00276 }
00277 */  
00278 void CudaLCP_ComputeNextIter_V5_FirstKernelf(const void * m,int mP,const void * q,void * f,void * res) {
00279     dim3 threads(BSIZE,1);
00280     dim3 grid(1,1);
00281     
00282     CudaLCP_ComputeNextIter_V5_FirstKernel_kernel<float><<< grid, threads,0 >>>((const float*) m,mP,(const float *) q,(float*)f,(float *)res);
00283     
00284 }
00285 void CudaLCP_ComputeNextIter_V5_FirstKerneld(const void * m,int mP,const void * q,void * f,void * res) {
00286 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00287     myprintf("CUDA ERROR: double precision not supported.\n");
00288 #else
00289     dim3 threads(BSIZE,1);
00290     dim3 grid(1,1);
00291     
00292     CudaLCP_ComputeNextIter_V5_FirstKernel_kernel<double><<< grid, threads,0 >>>((const double*) m,mP,(const double *) q,(double*)f,(double *)res);
00293 #endif  
00294 }
00295 
00296 /*
00297 for (int j=0;j<d;j++) {
00298     int compteur2=ligne+j;//ligne du bloc en cours
00299     
00300     float acc=0.0;
00301     
00302     for (int index1=0;index1<BSIZE;index1++) {//calcul du 1er bloc independant
00303         int dc = index1+colone;
00304         if (dc<dim) acc += cuda_M[dc][compteur2] * cuda_f[dc];
00305 }   
00306                     
00307     float r_calc = cuda_res[compteur2]+acc;
00308     float f_1 = cuda_f[compteur2];
00309     
00310     float f_2;
00311     if (r_calc<0) f_2 = -r_calc/cuda_M[compteur2][compteur2];
00312     else f_2=0.0;
00313     
00314     cuda_f[compteur2] = f_2;
00315     cuda_res[compteur2] = cuda_q[compteur2];
00316     cuda_err[0] += fabs(cuda_M[compteur2][compteur2] * (f_2 - f_1));
00317     
00318     for (int index1=0;index1<BSIZE;index1++) {//calcul du 1er bloc independant
00319         int compteur3 = ligne+index1;
00320         if ((compteur3<dim) && (compteur3!=compteur2)) cuda_res[compteur3] += cuda_M[compteur2][compteur3] * f_2;
00321 }   
00322 }   
00323                 
00324 for (int j=0;j<nbth;j++) {
00325     int dl=j;   
00326     if (ligne-BSIZE<0) dl+=BSIZE;
00327     else if (j>=ligne-BSIZE) dl+=2*BSIZE;
00328 
00329     for (int i=0;i<BSIZE;i++) {
00330         int dc = colone+i;
00331         if (dc<dim) cuda_res[dl]+= cuda_M[dc][dl] * cuda_f[dc];
00332 }
00333 }   
00334 */
00335 void CudaLCP_ComputeNextIter_V5_SecondKernelf(int dim,int nbth,int d,int ligne,int colone,const void * m,int mP,const void * q,void * f,void * err,void * res) {
00336     dim3 threads(BSIZE,1);
00337     dim3 grid(1,nbth+1);
00338     
00339     CudaLCP_ComputeNextIter_V5_SecondKernel_kernel<float><<< grid, threads,0 /*,threads.x*sizeof(float)*/ >>>(dim,d,ligne,colone,(const float*) m,mP,(const float *) q,(float*)f,(float*)err,(float *)res,BSIZE/2);
00340 }
00341 void CudaLCP_ComputeNextIter_V5_SecondKerneld(int dim,int nbth,int d,int ligne,int colone,const void * m,int mP,const void * q,void * f,void * err,void * res) {
00342 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00343     myprintf("CUDA ERROR: double precision not supported.\n");
00344 #else
00345     dim3 threads(BSIZE,1);
00346     dim3 grid(1,nbth+1);
00347     
00348     CudaLCP_ComputeNextIter_V5_SecondKernel_kernel<double><<< grid, threads,0 /*,threads.x*sizeof(double)*/ >>>(dim,d,ligne,colone,(const double*) m,mP,(const double *) q,(double*)f,(double*)err,(double *)res,BSIZE/2);
00349 #endif
00350 }
00351 
00353 
00354 void CudaLCP_FullKernel_V6f(int dim,int itMax,float tol,const void * m,int mP,const void * q,void * f,void * err,void * share) {
00355     dim3 threads(BSIZE_C,1);
00356     dim3 grid(1,NB_MULTIPROC);
00357 
00358     CudaLCP_FullKernel_V6_kernel<<< grid, threads,threads.x*threads.y*sizeof(float)>>>(dim,dim*itMax,tol,(const float *) m,mP,(const float *) q,(float *) f,(float *) err,(int *) share);
00359 
00360 }
00361 void CudaLCP_FullKernel_V6d(int dim,int itMax,float tol,const void * m,int mP,const void * q,void * f,void * err,void * share) {
00362 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00363     myprintf("CUDA ERROR: double precision not supported.\n");
00364 #else       
00365     dim3 threads(BSIZE_C,1);
00366     dim3 grid(1,NB_MULTIPROC);
00367 
00368     CudaLCP_FullKernel_V6_kernel<<< grid, threads,threads.x*threads.y*sizeof(double)>>>(dim,dim*itMax,(double)tol,(const double *) m,mP,(const double *) q,(double *) f,(double *) err,(int *) share);
00369 #endif
00370 }
00371 
00373 
00374 void CudaLCP_FullKernel_V7f(int dim,int itMax,float tol,const void * m,int mP,const void * q,void * f,void * err,void * share) {
00375     dim3 threads(BSIZE_C,BSIZE_L);
00376     dim3 grid(1,NB_MULTIPROC);
00377     int dim_n = (dim+BSIZE_L-1)/BSIZE_L * BSIZE_L;
00378 
00379     CudaLCP_FullKernel_V7_kernel<<< grid, threads,0>>>(dim,dim_n,dim_n*itMax,tol,(const float *) m,mP,(const float *) q,(float *) f,(float *) err,(int *) share);
00380 }
00381 void CudaLCP_FullKernel_V7d(int dim,int itMax,float tol,const void * m,int mP,const void * q,void * f,void * err,void * share) {
00382 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00383     myprintf("CUDA ERROR: double precision not supported.\n");
00384 #else       
00385     dim3 threads(BSIZE_C,BSIZE_L);
00386     dim3 grid(1,NB_MULTIPROC);
00387     int dim_n = (dim+BSIZE_L-1)/BSIZE_L * BSIZE_L;
00388 
00389     CudaLCP_FullKernel_V7_kernel<<< grid, threads,0>>>(dim,dim_n,dim_n*itMax,(double)tol,(const double *) m,mP,(const double *) q,(double *) f,(double *) err,(int *) share);
00390 
00391 #endif
00392 }           
00393 
00395 
00396 void CudaLCP_FullKernel_V8f(int dim,int itMax,float tol,const void * m,int mP,const void * q,void * f,void * err,void * share) {
00397     dim3 threads(BSIZE_C,BSIZE_C);
00398     dim3 grid(1,NB_MULTIPROC);
00399     int dim_n = (dim+BSIZE_C-1)/BSIZE_C * BSIZE_C;
00400 
00401     CudaLCP_FullKernel_V8_kernel<<< grid, threads,0>>>(dim,dim_n,dim_n*itMax,tol,(const float *) m,mP,(const float *) q,(float *) f,(float *) err,(int *) share);
00402 }
00403 void CudaLCP_FullKernel_V8d(int dim,int itMax,float tol,const void * m,int mP,const void * q,void * f,void * err,void * share) {
00404 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00405     myprintf("CUDA ERROR: double precision not supported.\n");
00406 #else       
00407     dim3 threads(BSIZE_C,BSIZE_C);
00408     dim3 grid(1,NB_MULTIPROC);
00409     int dim_n = (dim+BSIZE_C-1)/BSIZE_C * BSIZE_C;
00410 
00411     CudaLCP_FullKernel_V8_kernel<<< grid, threads,0>>>(dim,dim_n,dim_n*itMax,(double)tol,(const double *) m,mP,(const double *) q,(double *) f,(double *) err,(int *) share);
00412 
00413 #endif
00414 }           
00415 
00417 
00418 /*
00419     for (int i=0;i<numContacts;i++) {
00420         int index0 = 3*i;
00421         int index1 = index0+1;
00422         int index2 = index0+2;
00423         
00424         cuda_res[index0] = cuda_q[index0];
00425         cuda_res[index1] = cuda_q[index1];
00426         cuda_res[index2] = cuda_q[index2];
00427         
00428         for (int j=index2+1;j<dim;j++) {
00429             cuda_res[index0] += cuda_M[j][index0] * cuda_f[j];
00430             cuda_res[index1] += cuda_M[j][index1] * cuda_f[j];
00431             cuda_res[index2] += cuda_M[j][index2] * cuda_f[j];
00432 }
00433 }   
00434 */
00435 void CudaNLCP_MultIndepf(int dim,const void * m,int pM,const void * f,void * tmp,int pTmp) {
00436     dim3 threads(MBSIZE,1); 
00437     dim3 grid((dim+MBSIZE-1)/MBSIZE,dim);
00438 
00439     CudaNLCP_MultIndep_kernel<float><<< grid, threads,0>>>(dim, (const float*)m,pM,(const float*)f,(float*)tmp,pTmp);
00440 }
00441 void CudaNLCP_MultIndepd(int dim,const void * m,int pM,const void * f,void * tmp,int pTmp) {
00442 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00443     myprintf("CUDA ERROR: double precision not supported.\n");
00444 #else       
00445     dim3 threads(MBSIZE,1); 
00446     dim3 grid((dim+MBSIZE-1)/MBSIZE,dim);
00447 
00448     CudaNLCP_MultIndep_kernel<double><<< grid, threads,0>>>(dim, (const double*)m,pM,(const double*)f,(double*)tmp,pTmp);
00449 #endif
00450 }
00451 
00452 void CudaNLCP_AddIndepf(int dim,int tmpsize,const void * q,const void * tmp,int pTmp,void * res) {
00453     dim3 threads(tmpsize,1);    
00454     dim3 grid(dim,1);
00455 
00456     CudaNLCP_AddIndep_kernel<float><<< grid, threads,threads.x*sizeof(float)>>>(dim,tmpsize,(const float*)q,(const float*)tmp,pTmp,(float*)res);
00457 }
00458 void CudaNLCP_AddIndepd(int dim,int tmpsize,const void * q,const void * tmp,int pTmp,void * res) {
00459 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00460     myprintf("CUDA ERROR: double precision not supported.\n");
00461 #else
00462     dim3 threads(tmpsize,1);    
00463     dim3 grid(dim,1);
00464 
00465     CudaNLCP_AddIndep_kernel<double><<< grid, threads,threads.x*sizeof(double)>>>(dim,tmpsize,(const double*)q,(const double*)tmp,pTmp,(double*)res);
00466 #endif
00467 }
00468 
00469 void CudaNLCP_ComputeNextIter_V1_InDepKernelf(int dim,int d,int debutBlock,const void * m,int mP,const void * f,void * res) {
00470     dim3 threads(MBSIZE,1);
00471     dim3 grid(1,dim-d); 
00472     
00473     CudaNLCP_ComputeNextIter_V1_InDepKernel_kernel<float><<< grid, threads,0>>>(dim,debutBlock,(const float*) m,mP,(const float*)f, (float *)res);
00474 }
00475 void CudaNLCP_ComputeNextIter_V1_InDepKerneld(int dim,int d,int debutBlock,const void * m,int mP,const void * f,void * res) {
00476 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00477     myprintf("CUDA ERROR: double precision not supported.\n");
00478 #else       
00479     dim3 threads(MBSIZE,1);
00480     dim3 grid(1,dim-d);
00481     
00482     CudaNLCP_ComputeNextIter_V1_InDepKernel_kernel<double><<< grid, threads,0>>>(dim,debutBlock,(const double*) m,mP,(const double*)f, (double *)res);
00483 #endif
00484 }
00485 
00486 void CudaNLCP_ComputeNextIter_V1_DepKernelf(int d,int debutblock,float mu,const void * m,int mP,const void * q,void * f,void * err,void * res) {
00487     dim3 threads(MBSIZE,1);
00488     dim3 grid(1,1);
00489         
00490     CudaNLCP_ComputeNextIter_V1_DepKernel_kernel<float><<< grid, threads,0>>>(d,debutblock,mu,(const float*) m,mP,(const float *) q,(float*)f, (float*) err,(float *)res);
00491 }
00492 void CudaNLCP_ComputeNextIter_V1_DepKerneld(int d,int debutblock,float mu,const void * m,int mP,const void * q,void * f,void * err,void * res) {
00493 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00494     myprintf("CUDA ERROR: double precision not supported.\n");
00495 #else       
00496     dim3 threads(MBSIZE,1);
00497     dim3 grid(1,1);
00498         
00499     CudaNLCP_ComputeNextIter_V1_DepKernel_kernel<double><<< grid, threads,0>>>(d,debutblock,(double)mu,(const double*) m,mP,(const double *) q,(double*)f, (double*) err,(double *)res);
00500 #endif
00501 }
00502 
00503 void CudaNLCP_FullKernel_V2f(int dim,int itMax,float tol,float mu,const void * m,int mP,const void * q,void * f,void * err,void * share) {
00504     dim3 threads(MBSIZE_C,MBSIZE_C);
00505     dim3 grid(1,NB_MULTIPROC);
00506     int dim_n = (dim+MBSIZE_C-1)/MBSIZE_C * MBSIZE_C;
00507 
00508     CudaNLCP_FullKernel_V2_kernel<float><<< grid, threads,0>>>(dim,dim_n,dim_n*itMax,tol,mu,(const float *) m,mP,(const float *) q,(float *) f,(float *) err,(int *) share);
00509 }
00510 void CudaNLCP_FullKernel_V2d(int dim,int itMax,float tol,float mu,const void * m,int mP,const void * q,void * f,void * err,void * share) {
00511 #if !defined(__CUDA_ARCH__) ||  __CUDA_ARCH__ < 130
00512     myprintf("CUDA ERROR: double precision not supported.\n");
00513 #else       
00514     dim3 threads(MBSIZE_C,MBSIZE_C);
00515     dim3 grid(1,NB_MULTIPROC);
00516     int dim_n = (dim+MBSIZE_C-1)/MBSIZE_C * MBSIZE_C;
00517 
00518     CudaNLCP_FullKernel_V2_kernel<double><<< grid, threads,0>>>(dim,dim_n,dim_n*itMax,tol,mu,(const double *) m,mP,(const double *) q,(double *) f,(double *) err,(int *) share);
00519 
00520 #endif
00521 }