内核函数变量累加崩溃的问题

  • 5 replies
  • 3629 views
内核函数变量累加崩溃的问题
« 于: 三月 19, 2018, 09:56:52 am »
 本帖最后由 Loke 于 2018-3-19 10:44 编辑

内核函数如下:
__kernel void helloworld(__global float *pIn, __global float *pOut)  
{  
       int iNum = get_global_id(0);  

        float dd = 0;
        for(int i = 0;i < 65535;i++)
        {
           dd += 0.01;
        }

    pOut[iNum] = pIn[iNum] + dd;

}  


主机端:

int main()  
{  
   //获取平台、设备、上下文、编译内核等初始化操作
    ............


    // 创建输入内存对象  

        memInutBuffer = clCreateBuffer(  
                Context,  
                CL_MEM_WRITE_ONLY |CL_MEM_ALLOC_HOST_PTR,
                (65535 ) * sizeof(float),         // 输入内存空间大小  
                NULL,  
                NULL);  

        float* mapPtr = (float*)clEnqueueMapBuffer(
                CommandQueue,
                memInutBuffer,
                CL_TRUE,
                CL_MAP_WRITE,
                0,
                (65535 ) * sizeof(float),
                0,
                NULL,
                NULL,
                NULL
                );

          for(int i = 0;i<65535;i++)
          {
                  mapPtr = 1;
         }

  clEnqueueUnmapMemObject(
                  CommandQueue,
                  memInutBuffer,
                  mapPtr,
                  0,
                  NULL,
                  NULL);


    // 创建输出内存对象  
    memOutputBuffer = clCreateBuffer(  
        Context,  
        CL_MEM_READ_ONLY |  CL_MEM_ALLOC_HOST_PTR,            
        (65535) * sizeof(float),    
        NULL,  
        NULL);  

    if ((NULL == memInutBuffer) || (NULL == memOutputBuffer))  
    {  
        cout << "Error creating memory objects" << endl;  
        return 0;  
    }  

    //--------------------------8. 创建内核对象-------------------------------------  
     .........

    //----------------------------9. 设置内核参数----------------------------------  
    iStatus = clSetKernelArg(Kernel,  
        0,      // 参数索引  
        sizeof(cl_mem),  
        (void *)&memInutBuffer);  

    iStatus |= clSetKernelArg(Kernel, 1, sizeof(cl_mem), (void *)&memOutputBuffer);  


    // --------------------------10.运行内核---------------------------------  
    uiGlobal_Work_Size[0] = 65535;  
    size_t  localSize[] = {1};

    // 利用命令队列使将再设备上执行的内核排队  
    iStatus = clEnqueueNDRangeKernel(  
        CommandQueue,  
        Kernel,  
        1,  
        NULL,  
        uiGlobal_Work_Size,  // 确定内核在设备上的多个处理单元间的分布  
        localSize,                // 确定内核在设备上的多个处理单元间的分布  
        0,  
        NULL,  
        NULL);  


    if (CL_SUCCESS != iStatus)  
    {  
        cout << iStatus << endl;  
        return 0;  
    }  

    // ----------------------------11. 将输出读取到主机内存  

        float * p =NULL;
        p = (float *)clEnqueueMapBuffer(
                        CommandQueue,
                        memOutputBuffer,
                        CL_TRUE,
                        CL_MAP_READ,
                        0,
                        (65535) * sizeof(float),
                        0,
                        NULL,
                        NULL,
                        NULL
                        );
       

        for(int i = 0;i<4;i++)
        cout << p<< endl;  
       
       clEnqueueUnmapMemObject(
                        CommandQueue,
                        memOutputBuffer,
                        p,
                        0,
                        NULL,
                        NULL);

        iStatus = clReleaseMemObject(memOutputBuffer);  


    // -------------------------------13. 释放资源--------------------------------  
        ......
    return 0;  
}  


问题是:     当我读出内存数据,打印p,出现内存读取位置错误。而将内核函数的  dd += 0.01; 改为dd = 0.01 ;或不将累加的dd赋值,如pOut[iNum] = pIn[iNum] ,可以得到预期结果。或者将内核函数的累加次数减少,如将i<65535改为 i < 40000,也能得到正确结果。 将全局工作项 uiGlobal_Work_Size[0] = 65535;  改为很小的如32,也能得到预期结果。opencl新手,请教大家这是什么原因导致的,如果一定要累加的值,应如何解决。谢谢!

(无标题)
« 回复 #1 于: 三月 19, 2018, 10:12:02 am »
GPU 是 NVIDIA GTX 1060 3GB

(无标题)
« 回复 #2 于: 三月 21, 2018, 10:38:09 am »
这个kernel代码非常简单,局部变量不多,而且主要就是一个for循环。因此我基本怀疑是编译器对循环展开优化导致的这一系列的问题。因此我建议楼主把OpenCL C的编译选项设置为-O0看看结果何如。

(无标题)
« 回复 #3 于: 三月 21, 2018, 05:27:16 pm »
这个kernel代码非常简单,局部变量不多,而且主要就是一个for循环。因此我基本怀疑是编译器对循环展开优化 ...

感谢您的建议,我试试看结果如何。另外,在原来基础上,当我把内存映射p = (float *)clEnqueueMapBuffer(CommandQueue,memOutputBuffer,CL_TRUE,CL_MAP_READ,0,(65535)*sizeof(float),0,NULL,NULL,NULL);改为读clEnqueueReadBuffer(CommandQueue,memOutputBuffer, CL_TRUE, 0,  (65535) * sizeof(float),p,  0,  NULL,  NULL)后,加上累加的dd得到了预期结果。不知道是不是映射的原因,导致加上累加值后指针错误。

(无标题)
« 回复 #4 于: 三月 21, 2018, 09:41:51 pm »
感谢您的建议,我试试看结果如何。另外,在原来基础上,当我把内存映射p = (float *)clEnqueueMapBuffer( ...

哦……这个也有可能。Map操作其实主要针对核心GPU的,可以将主机端要访问的存储器对象映射到GPU专用的VRAM空间区域,这样可以实现零拷贝。不过对于独立GPU来说就没啥效果了,所以直接用Read/Write更干脆利落,呼呼。不过也不排除是NV的驱动可能有些问题。

(无标题)
« 回复 #5 于: 三月 23, 2018, 11:13:04 am »
哦……这个也有可能。Map操作其实主要针对核心GPU的,可以将主机端要访问的存储器对象映射到GPU专用的VRA ...

嗯嗯,学习了。又做了试验,将Write改回原来的Map后,无论编译优化选项改不改,莫名其妙地又得到了预期结果。可能确实NV的驱动有问题,或者其他还没排除的问题。总之,用Read/Write的话,能保证得到预期结果。