OpenClカーネルのループが大域メモリ浮動小数点配列をローリングする

Question

基本的な並列プログラミングの概念がわかりません。以下のカーネルは、私が抱えている問題を再現する簡単な/工夫された例です。値を計算し、それを「ブロック」内のすべてのアイテムに割り当てるために、ポイント内のすべての値を使用しようとします。私はこれらの配列のサイズの制限を押してみたいです。 "ブロック"配列を大量に（1億以上の）浮動小数点数にすることができますが、clEnqueueNDRangeKernelの直後にclFinishを呼び出した後、 "points"が〜100,000浮動小数点数でいっぱいになると "無効なコマンドキュー"エラーが発生します。なぜあなたが私に理解を助けることができますか？

__kernel void openClTesting (__global float *blocks, __global float *points, int pointsCount) 
    { 
     int globalId = get_global_id(0); 
     int count = 0; 
     for (int i = 0; i < pointsCount; i++) 
     { 
      count++; 
     } 
     blocks[globalId] = count; 
    }; 

いくつかのデバイス情報：

CL_DEVICE_LOCAL_MEM_SIZE = 49,152 
CL_DEVICE_GLOBAL_MEM_SIZE = 2,147,483,648 
CL_DEVICE_MAX_MEM_ALLOC_SIZE = 536,870,912 

ホスト・コード：私は気付か

#include "stdafx.h" 
#include "CL\opencl.h" 
#include <iostream> 
#include <fstream> 
#include <string> 
#include <stddef.h> 
#include <stdlib.h> 
#include <stdio.h> 

#define NUM_POINTS 100000 
#define NUM_BLOCKS 100000000 

struct openClData 
{ 
cl_device_id deviceId = NULL; 
cl_uint numDevices; 
cl_uint numPlatforms; 
cl_int ret; 
cl_platform_id *platforms = NULL; 
cl_context context; 
cl_command_queue commandQueue; 
cl_program program; 
cl_kernel kernel; 
char* kernelCode; 
cl_uint kernelCodeSize; 
size_t globalItemSize; 
size_t localItemSize = 1; 
}; 


char* getKernelCode(); 
void printErrorLog(openClData oclData); 
void printRet(openClData oclData, int line); 
int countFileChars(const char *fileName); 

int _tmain(int argc, _TCHAR* argv[]) 
{ 
openClData oclData; 
oclData.globalItemSize = NUM_POINTS; 
oclData.kernelCode = getKernelCode(); 
std::cout << oclData.kernelCode << std::endl; 
oclData.kernelCodeSize = strlen(oclData.kernelCode); 

int numPoints = NUM_POINTS; 
int numBlocks = NUM_BLOCKS; 
cl_long localMemSize = 0, globalMemSize = 0, maxAllocMemSize = 0; 
float *blocks = new float[numBlocks]{0}; 
float *points = new float[numPoints]{0}; 

//prepare platform, device, context and command queue 
oclData.ret = clGetPlatformIDs(0, NULL, &oclData.numPlatforms); 
printRet(oclData, __LINE__); 
oclData.platforms = (cl_platform_id *)malloc(oclData.numPlatforms * sizeof(cl_platform_id)); 
oclData.ret = clGetPlatformIDs(oclData.numPlatforms, oclData.platforms, NULL); 
printRet(oclData, __LINE__); 
oclData.ret = clGetDeviceIDs(oclData.platforms[0], CL_DEVICE_TYPE_GPU, 1, &oclData.deviceId, &oclData.numDevices); 
printRet(oclData, __LINE__); 
oclData.context = clCreateContext(NULL, 1, &oclData.deviceId, NULL, NULL, &oclData.ret); 
printRet(oclData, __LINE__); 
oclData.commandQueue = clCreateCommandQueue(oclData.context, oclData.deviceId, 0, &oclData.ret); 
printRet(oclData, __LINE__); 
//prepare cl_mem objects 
cl_mem memObjBlocks = clCreateBuffer(oclData.context, CL_MEM_READ_WRITE, sizeof(float) * numBlocks, NULL, &oclData.ret); 
printRet(oclData, __LINE__); 
cl_mem memObjPoints = clCreateBuffer(oclData.context, CL_MEM_READ_WRITE, sizeof(float) * numPoints, NULL, &oclData.ret); 
printRet(oclData, __LINE__); 
oclData.ret = clEnqueueWriteBuffer(oclData.commandQueue, memObjBlocks, CL_TRUE, 0, sizeof(float) * numBlocks, blocks, 0, NULL, NULL); 
printRet(oclData, __LINE__); 
oclData.ret = clEnqueueWriteBuffer(oclData.commandQueue, memObjPoints, CL_TRUE, 0, sizeof(float) * numPoints, points, 0, NULL, NULL); 
printRet(oclData, __LINE__); 
//prepare program 
oclData.program = clCreateProgramWithSource(oclData.context, 1, (const char**)&oclData.kernelCode, (const size_t *)&oclData.kernelCodeSize, &oclData.ret); 
printRet(oclData, __LINE__); 
oclData.ret = clBuildProgram(oclData.program, 1, &oclData.deviceId, NULL, NULL, NULL); 
printRet(oclData, __LINE__); 
if (oclData.ret == CL_BUILD_PROGRAM_FAILURE) printErrorLog(oclData); 
oclData.kernel = clCreateKernel(oclData.program, "openClTesting", &oclData.ret); 
printRet(oclData, __LINE__); 
//set arguments 
oclData.ret = clSetKernelArg(oclData.kernel, 0, sizeof(cl_mem), &memObjBlocks); 
printRet(oclData, __LINE__); 
oclData.ret = clSetKernelArg(oclData.kernel, 1, sizeof(cl_mem), &memObjPoints); 
printRet(oclData, __LINE__); 
oclData.ret = clSetKernelArg(oclData.kernel, 2, sizeof(int), &numPoints); 
printRet(oclData, __LINE__); 
//run 
oclData.ret = clEnqueueNDRangeKernel(oclData.commandQueue, oclData.kernel, 1, NULL, &oclData.globalItemSize, &oclData.localItemSize, 0, NULL, NULL); 
printRet(oclData, __LINE__); 
oclData.ret = clFinish(oclData.commandQueue); 
printRet(oclData, __LINE__); 
oclData.ret = clEnqueueReadBuffer(oclData.commandQueue, memObjBlocks, CL_TRUE, 0, sizeof(float) * numBlocks, blocks, 0, NULL, NULL); 
printRet(oclData, __LINE__); 
oclData.ret = clFinish(oclData.commandQueue); 
printRet(oclData, __LINE__); 
//print some device info 
oclData.ret = clGetDeviceInfo(oclData.deviceId, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), &localMemSize, 0); 
std::cout << "CL_DEVICE_LOCAL_MEM_SIZE = " << localMemSize << '\n'; 
oclData.ret = clGetDeviceInfo(oclData.deviceId, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(cl_long), &globalMemSize, 0); 
std::cout << "CL_DEVICE_GLOBAL_MEM_SIZE = " << globalMemSize << '\n'; 
oclData.ret = clGetDeviceInfo(oclData.deviceId, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_long), &maxAllocMemSize, 0); 
std::cout << "CL_DEVICE_MAX_MEM_ALLOC_SIZE = " << maxAllocMemSize << '\n'; 

//clean up 
oclData.ret = clFlush(oclData.commandQueue); 
printRet(oclData, __LINE__); 
oclData.ret = clFinish(oclData.commandQueue); 
printRet(oclData, __LINE__); 
oclData.ret = clReleaseKernel(oclData.kernel); 
printRet(oclData, __LINE__); 
oclData.ret = clReleaseProgram(oclData.program); 
printRet(oclData, __LINE__); 
oclData.ret = clReleaseMemObject(memObjBlocks); 
printRet(oclData, __LINE__); 
oclData.ret = clReleaseMemObject(memObjPoints); 
printRet(oclData, __LINE__); 
oclData.ret = clReleaseCommandQueue(oclData.commandQueue); 
printRet(oclData, __LINE__); 
oclData.ret = clReleaseContext(oclData.context); 
printRet(oclData, __LINE__); 
for (size_t i = 0; i < 10; i++) 
{ 
    std::cout << blocks[i] << std::endl; 
} 
delete blocks; 
delete points; 
return 0; 
} 

char* getKernelCode() 
{ 
char* kernelCode = 
    "__kernel void openClTesting (__global float *blocks, __global float *points, int pointsCount)" 
    "{" 
    " int globalId = get_global_id(0);" 
    " int count = 0;" 
    " for (int i = 0; i < pointsCount; i++)" 
    " {" 
    "  count++;" 
    " }" 
    "blocks[globalId] = count;" 
    "}"; 
return kernelCode; 
} 

void printErrorLog(openClData oclData) 
{ 
size_t log_size; 
clGetProgramBuildInfo(oclData.program, oclData.deviceId, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size); 
char *log = (char *)malloc(log_size); 
clGetProgramBuildInfo(oclData.program, oclData.deviceId, CL_PROGRAM_BUILD_LOG, log_size, log, NULL); 
std::cout << log; 
free(log); 
} 

void printRet(openClData oclData, int line) 
{ 
std::cout << line << ", " << oclData.ret << std::endl; 
} 

int countFileChars(const char *fileName) 
{ 
std::ifstream ifs(fileName); 
ifs.seekg(0, std::ios_base::end); 
size_t count = ifs.tellg(); 
ifs.seekg(0, std::ios_base::beg); 
return count; 
} 

Answer 1

いくつかのこと：

• あなたはNUM_POINTS作業項目を起動しているが、それぞれの結果をに書き込みます- NUM_BLOCKSのアイテムがあります。したがって、NUM_POINTSがNUM_BLOCKSより大きい場合、これは未定義の動作です。また、NUM_BLOCKSを変更すると何もしない理由についても説明しています（上記の制限の外）.NUM_BLOCKSの値はメモリ割り当てとは別に効果がありません。 （見つかったメモリの割り当て制限は、実装のCL_DEVICE_MAX_MEM_ALLOC_SIZEの値とほぼ一致します）。
• ここでは、カーネルのタイムアウト条件が発生している可能性があります。 1回の作業項目で100000回のループ反復がかなり多くなります。 OpenCLの実装によっては、実行に時間がかかりすぎるとカーネルが強制終了される可能性があります。使用可能なスレッド並列性をより有効に活用し、ループするのではなく作業項目間でワークをより水平に分割することを検討してください。短期間で実行される多くの作業項目は、通常、長時間実行されるものよりも優れています。一般的なノートで

Answer 2

それはあなたがお使いのコンピューティングデバイスを並列に実行することができ、ワークグループの数に並列処理削減する単一の作業項目から構成するすべてのOpenCLのワークグループを強制するので、localItemSize = 1;は、避けるべきですこれは実行可能な作業項目の数よりはるかに少なくなります。あなたはNUM_POINTS作業グループを作成しているので、これはまた、あなたのエラーの原因かもしれ

clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalSize, NULL, 0, NULL, NULL); 

しかし：あなたは、単にそれ自身で妥当な値をOpenCLの実装図を持っている代わりに、ローカル項目のサイズにNULLを渡すことができますデバイス上のキューのサイズはメモリが制限されています（CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE）。