python - pyopenCL，openCL，无法在GPU上构建程序

Question

我有一段内核源代码，可以在我的 PC 上的 G970 上运行，但无法在我的 2015 年初配备 Iris 6100 1536MB 显卡的 MacBook pro 上编译。

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() 

此 get_devices() 显示我在“Apple”上有“Intel(R) Core(TM) i5-5287U CPU @ 2.90GHz”，地址为 0xffffffff>，在“Apple”上有“Intel(R) Iris(TM) Graphics 6100” ' 在 0x1024500。

内核将在 CPU 上正确运行。但是当我在 GPU 上构建程序时。它返回:

---------------------------------------------------------------------------
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)

有什么建议为什么这不会运行吗？我运行的是 2015 年初的 MacBook Pro、Sierra 10.12.5。print cl.version.VERSION 返回2015.2.4

这是内核代码:

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;  

}
"""

谢谢

家骏

Answer 1

在这种情况下，您应该尝试查询构建的错误。在类似的内核代码错误中您可以做的另一件事是您可以使用离线编译器。每个 OpenCL 实现者都有离线编译器。

您可以在这里找到英特尔的 OpenCL 离线编译器:https://software.intel.com/en-us/articles/programming-with-the-intel-sdk-for-opencl-applications-development-tools

AMD 有一个名为 CodeXL 的工具，您还可以在其中进行离线编译，看看您的内核代码是否可以编译。

这里是ARM OpenCL离线编译器:https://developer.arm.com/products/software-development-tools/graphics-development-tools/mali-offline-compiler/downloads

Intel 支持最高 OpenCL 2.1，而 ARM 支持最高 1.1。因此，您可以选择其中任何一个来编译您的内核代码，以轻松找出错误或错误。

您的内核中的问题是以下行:

float v[dim];

OpenCL C 规范不允许可变长度数组，离线编译器会给出以下错误:

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

您可以修复该行以克服错误，从现在开始，您可以检查您的内核是否可以使用离线编译器进行编译。

编辑:在规范中，有一个脚注解释了不支持可变长度数组。您可以在这里看到它:

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