OpenCL 學習step by step (5) 使用二維NDRange workgroup

      在本教程中，我們使用二維NDRange來設定workgroup，這樣在opencl中，workitme的組織形式是二維的，Kernel中 的代碼也要做相應的改變，我們先看一下clEnqueueNDRangeKernel函數的變化。首先我們指定了workgroup size為localx*localy,通常這個值為64的倍數，但最好不要超過256。

//執行kernel,Range用2維，work itmes size為width*height,

cl_event ev;

size_t globalThreads[] = {width, height};

size_t localx, localy;

if(width/8 > 4)

    localx = 16;

else if(width < 8)

    localx = width;

else localx = 8;

if(height/8 > 4)

    localy = 16;

else if (height < 8)

    localy = height;

else localy = 8;

size_t localThreads[] = {localx, localy}; // localx*localy應該是64的倍數

printf("global_work_size =(%d,%d), local_work_size=(%d, %d)\n",width,height,localx,localy);

clTimer.Reset();

clTimer.Start();

clEnqueueNDRangeKernel( queue,

    kernel,

    2,

    NULL,

    globalThreads,

    localThreads, 0, NULL, &ev);

注意：在上面代碼中，定義global threads以及local threads數量，都是通過二維數組的方式進行的。

    新的Kernel代碼如下：

      我們在kernel中增加了#pragma OPENCL EXTENSION cl_amd_printf : enable，以便在kernel中通過printf函數進行debug，這是AMD的一個擴充。printf還可以直接列印出float4這樣的向量，比如printf(“%v4f”, vec)。

      另外，在main.cpp中增加一行代碼：

//告訴driver dump il和isa檔案

_putenv("GPU_DUMP_DEVICE_KERNEL=3");

      我們可以在程式目錄dump出il和isa形式的kernel檔案，對于熟悉isa彙編的人，這是一個很好的調試performance的方法。

     在最新的app sdk 2.7及以後的sdk中，在kernel中使用printf的時候，這個程式會hang在哪兒，以前沒這種情況。

程式執行界面。

完整的代碼請參考：

工程檔案gclTutorial4

代碼下載下傳：

<a href="http://files.cnblogs.com/mikewolf2002/gclTutorial.zip">http://files.cnblogs.com/mikewolf2002/gclTutorial.zip</a>

<a href="http://files.cnblogs.com/mikewolf2002/gclTutorial4.zip"></a>