pyopenCL, openCL, Impossible de construire le programme sur GPU

Question

J'ai une source de noyau qui fonctionne sur le G970 sur mon PC mais ne compilera pas sur mon début 2015 MacBook pro avec le graphisme Iris 6100 1536MB.

platform = cl.get_platforms()[0] 
device = platform.get_devices()[1] # Get the GPU ID 
ctx  = cl.Context([device])  # Tell CL to use GPU 
queue = cl.CommandQueue(ctx)  # Create a command queue for the target device. 
# program = cl.Program(ctx, kernelsource).build() 
print platform.get_devices() 

Ce get_devices() show je 'Intel (R) Core (TM) i5-5287U CPU @ 2.90GHz', « Intel (R) Iris (TM) sur 'Apple' à 0xffffffff> Graphics 6100. 'sur' Apple 'à 0x1024500.

Le noyau fonctionnera correctement sur la CPU. Mais quand je construis le programme sur GPU. Il renvoie:

--------------------------------------------------------------------------- 
RuntimeError        Traceback (most recent call last) 
<ipython-input-44-e2b6e1b931de> in <module>() 
     3 ctx  = cl.Context([device])  # Tell CL to use GPU 
     4 queue = cl.CommandQueue(ctx)  # Create a command queue for the target device. 
----> 5 program = cl.Program(ctx, kernelsource).build() 
     6 
     7 

/usr/local/lib/python2.7/site-packages/pyopencl-2015.2.4-py2.7-macosx-10.11-x86_64.egg/pyopencl/__init__.pyc in build(self, options, devices, cache_dir) 
    393       self._context, self._source, options, devices, 
    394       cache_dir=cache_dir), 
--> 395      options=options, source=self._source) 
    396 
    397    del self._context 

/usr/local/lib/python2.7/site-packages/pyopencl-2015.2.4-py2.7-macosx-10.11-x86_64.egg/pyopencl/__init__.pyc in _build_and_catch_errors(self, build_func, options, source) 
    428   # Python 3.2 outputs the whole list of currently active exceptions 
    429   # This serves to remove one (redundant) level from that nesting. 
--> 430   raise err 
    431 
    432  # }}} 

RuntimeError: clbuildprogram failed: BUILD_PROGRAM_FAILURE - 

Build on <pyopencl.Device 'Intel(R) Iris(TM) Graphics 6100' on 'Apple' at 0x1024500>: 

Cannot select: 0x7f94b30a5110: i64,ch = dynamic_stackalloc 0x7f94b152a290, 0x7f94b30a4f10, 0x7f94b3092c10 [ORD=7] [ID=54] 
    0x7f94b30a4f10: i64 = and 0x7f94b30a4c10, 0x7f94b3092b10 [ORD=7] [ID=52] 
    0x7f94b30a4c10: i64 = add 0x7f94b30a6610, 0x7f94b3092a10 [ORD=7] [ID=49] 
     0x7f94b30a6610: i64 = shl 0x7f94b3092d10, 0x7f94b3092e10 [ID=46] 
     0x7f94b3092d10: i64 = bitcast 0x7f94b30a4810 [ID=41] 
      0x7f94b30a4810: v2i32 = IGILISD::MOVSWZ 0x7f94b3092710, 0x7f94b30a2810, 0x7f94b30a2810, 0x7f94b30a2810 [ID=32] 
      0x7f94b3092710: i32,ch = CopyFromReg 0x7f94b152a290, 0x7f94b3092610 [ORD=5] [ID=22] 
       0x7f94b3092610: i32 = Register %vreg60 [ORD=5] [ID=1] 
      0x7f94b30a2810: i32 = Constant<0> [ORD=1] [ID=7] 
      0x7f94b30a2810: i32 = Constant<0> [ORD=1] [ID=7] 
      0x7f94b30a2810: i32 = Constant<0> [ORD=1] [ID=7] 
     0x7f94b3092e10: i64 = bitcast 0x7f94b30a3f10 [ID=38] 
      0x7f94b30a3f10: v2i32 = IGILISD::MOVSWZ 0x7f94b30a4510, 0x7f94b30a2810, 0x7f94b30a2810, 0x7f94b30a2810 [ID=29] 
      0x7f94b30a4510: i32 = Constant<2> [ID=19] 
      0x7f94b30a2810: i32 = Constant<0> [ORD=1] [ID=7] 
      0x7f94b30a2810: i32 = Constant<0> [ORD=1] [ID=7] 
      0x7f94b30a2810: i32 = Constant<0> [ORD=1] [ID=7] 
     0x7f94b3092a10: i64 = bitcast 0x7f94b30a4b10 [ID=40] 
     0x7f94b30a4b10: v2i32 = IGILISD::MOVSWZ 0x7f94b30a4e10, 0x7f94b30a2810, 0x7f94b30a2810, 0x7f94b30a2810 [ID=31] 
      0x7f94b30a4e10: i32 = Constant<7> [ID=21] 
      0x7f94b30a2810: i32 = Constant<0> [ORD=1] [ID=7] 
      0x7f94b30a2810: i32 = Constant<0> [ORD=1] [ID=7] 
      0x7f94b30a2810: i32 = Constant<0> [ORD=1] [ID=7] 
    0x7f94b3092b10: i64 = bitcast 0x7f94b3092910 [ID=39] 
     0x7f94b3092910: v2i32 = IGILISD::MOVSWZ 0x7f94b30a5010, 0x7f94b30a4210, 0x7f94b30a2810, 0x7f94b30a2810 [ID=30] 
     0x7f94b30a5010: i32 = Constant<-8> [ID=20] 
     0x7f94b30a4210: i32 = Constant<-1> [ORD=3] [ID=10] 
     0x7f94b30a2810: i32 = Constant<0> [ORD=1] [ID=7] 
     0x7f94b30a2810: i32 = Constant<0> [ORD=1] [ID=7] 
    0x7f94b3092c10: i64 = bitcast 0x7f94b3092810 [ID=35] 
    0x7f94b3092810: v2i32 = IGILISD::MOVSWZ 0x7f94b30a2810, 0x7f94b30a2810, 0x7f94b30a2810, 0x7f94b30a2810 [ID=27] 
     0x7f94b30a2810: i32 = Constant<0> [ORD=1] [ID=7] 
     0x7f94b30a2810: i32 = Constant<0> [ORD=1] [ID=7] 
     0x7f94b30a2810: i32 = Constant<0> [ORD=1] [ID=7] 
     0x7f94b30a2810: i32 = Constant<0> [ORD=1] [ID=7] 
In function: trajectories 
(options: -I /usr/local/lib/python2.7/site-packages/pyopencl-2015.2.4-py2.7-macosx-10.11-x86_64.egg/pyopencl/cl) 
(source saved as /var/folders/p2/jd7m10gs5k1_q6hx5kvktkcc0000gn/T/tmpWQmCKr.cl) 

Une suggestion pourquoi cela ne fonctionnera pas? Je cours début 2015 MacBook Pro, Sierra 10.12.5. imprimer cl.version.VERSION retour 2015.2.4

Voici le code du noyau:

kernelsource = """ 
__kernel void trajectories(
    // TODO: adjust argtypes above if this is changed 
    const int N, 
    const int dim, 
    __constant float* data, 
    const int nrParticles, 
    __global float* pos, 
    __global float* vel, 
    const int nrSteps, 
    __global float* trj, 
    __global float* sigarr, 
    const float sigma, 
    const float mass, 
    const float alpha, // alpha is resistance in reverse. 
    const float dt 
){ 
    int i,k,step; 
    float h, sigsum, hexp; 
    int pidx = get_global_id(0); // global ID used as particle index 
    int ofs = pidx * nrSteps * dim; 
    int accofs = ofs + (nrSteps-1) * dim; // use last trj point to tmp store acc vector 
    float v[dim]; 
    float sigma2 = sigma*sigma; 
    float m = mass/sigma2; 
    float dt_over_m = dt /m; 
    for(step=0; step<nrSteps; step++){ 
     for(k=0; k<dim; k++) 
     { 
      trj[accofs+k]=0; 
     } 
     for(i=0; i<N; i++) 
     { 

      h=0; // to store ||data[i]-x||**2 
      for(k=0; k<dim; k++) 
      { 
       v[k] = pos[pidx*dim+k] - data[i*dim + k]; 
       h += v[k]*v[k];  //h == force1p_sum 
      }; 
      hexp = exp(-h/sigma2)/sigma2; 

      for(k=0; k<dim; k++) 
      { 
       trj[accofs+k] += -(hexp) * v[k]; 
      };   
     }; 
     sigsum = 0; 
     for(k=0; k<dim; k++) 
     { 
      vel[pidx*dim+k]  = alpha * vel[pidx*dim+k] + dt_over_m * trj[accofs+k];  // vel = alpha*vel + acc*dt 
      pos[pidx*dim+k] += dt * vel[pidx*dim+k];      // pos = pos + vel*dt 
      sigsum    += vel[pidx*dim+k] * vel[pidx*dim+k]; // v^2 for kinetic energy 
      trj[ofs+step*dim+k] = pos[pidx*dim+k];    // write to result vector 

     }; 
     sigarr[pidx*nrSteps+step] = sigsum;     // sig = | vel | 
    } 
    for(step=0; step<nrSteps-2; step++) 
    { 
     sigarr[pidx*nrSteps+step] = sigarr[pidx*nrSteps+step+2] - sigarr[pidx*nrSteps+step+1]; 
    }; 
    sigarr[pidx*nrSteps+nrSteps-1] = sigarr[pidx*nrSteps+nrSteps-2] = 0; 

} 
""" 

Merci

Jiajun

Answer 1

Vous devriez essayer d'interroger l'erreur de la construction dans ce cas. Une autre chose que vous pouvez faire dans les erreurs de code noyau similaires est que vous pouvez utiliser des compilateurs hors ligne. Chaque implémenteur OpenCL a un compilateur hors ligne.

Vous pouvez trouver OpenCL compilateur en ligne d'Intel ici: https://software.intel.com/en-us/articles/programming-with-the-intel-sdk-for-opencl-applications-development-tools

AMD a un outil appelé CodeXL, dans lequel vous pouvez également faire la compilation en mode hors connexion pour voir si votre code du noyau compile.

Voici le compilateur ARM OpenCL en ligne: https://developer.arm.com/products/software-development-tools/graphics-development-tools/mali-offline-compiler/downloads

soutien Intel est à OpenCL 2.1 tandis que ARM prend en charge jusqu'à 1.1. Ainsi, vous pouvez choisir n'importe lequel d'entre eux pour compiler votre code noyau afin de trouver facilement des bogues ou des erreurs.

Le problème dans votre noyau est la ligne suivante:

float v[dim]; 

spécification OpenCL C ne permet pas de tableaux de longueur variable et le compilateur offline donne l'erreur suivante:

ERROR: <source>:22:12: error: variable length arrays are not supported in OpenCL 

Vous pouvez résoudre ce problème ligne pour surmonter l'erreur et à partir de maintenant, vous pouvez vérifier si votre noyau peut être compilé avec le compilateur hors ligne.

EDIT: Dans la spécification, il existe une note de bas de page expliquant que les tableaux de longueur variable ne sont pas pris en charge. Vous pouvez le voir ici:

https://www.khronos.org/registry/OpenCL/specs/opencl-2.0-openclc.pdf#page=31