lormutryas
BANNED | Редактировать | Профиль | Сообщение | Цитировать | Сообщить модератору
пытаюсь обуздать свою видюху
программа из opencl programming guide rev 2.7
#include <CL/cl.h>
#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include "Timer.h"
#define NDEVS 2
const char* kernel_source=
"#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics:enable\n"
"#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics:enable\n"
"#pragma OPENCL EXTENSION cl_amd_printf : enable\n"
"__kernel void minp(__global uint4 *src,\n"
"__global uint *gmin,\n"
"__local uint *lmin,\n"
"__global uint *dbg,\n"
"int nitems,\n"
"uint dev)\n"
"{\n"
"printf(\"ok\\n\");\n"
"uint count=(nitems/4)/get_global_size(0);\n"
"uint idx=(dev==0)?get_global_id(0)*count:get_global_id(0);\n"
"uint stride=(dev==0)?1:get_global_size(0);\n"
"uint pmin=(uint)-1;\n"
"for(int n=0;n<count;n++,idx+=stride)\n"
"{\n"
"pmin=min(pmin,src[idx].x);\n"
"pmin=min(pmin,src[idx].y);\n"
"pmin=min(pmin,src[idx].z);\n"
"pmin=min(pmin,src[idx].w);\n"
"}\n"
"if(get_local_id(0)==0)\n"
"lmin[0]=(uint)-1;\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"(void) atom_min(lmin,pmin);\n"
"barrier(CLK_LOCAL_MEM_FENCE);\n"
"if(get_local_id(0)==0)\n"
"gmin[get_group_id(0)]=lmin[0];\n"
"if(get_global_id(0)==0)\n"
"{\n"
"dbg[0]=get_num_groups(0);\n"
"dbg[1]=get_global_size(0);\n"
"dbg[2]=count;\n"
"dbg[3]=stride;\n"
"}\n"
"}\n"
"__kernel void reduce(__global uint4 *src,\n"
"__global uint *gmin)\n"
"{\n"
"(void) atom_min(gmin,gmin[get_global_id(0)]);\n"
"}\n";
int main(int argc,char ** argv)
{
cl_platform_id platform;
int dev,nw;
cl_device_type devs[NDEVS]={CL_DEVICE_TYPE_CPU,CL_DEVICE_TYPE_GPU};
cl_uint *src_ptr;
unsigned int num_src_items = 4096*4096;
time_t ltime;
time(<ime);
src_ptr = (cl_uint *)malloc(num_src_items*sizeof(cl_uint));
cl_uint a = (cl_uint) ltime,
b = (cl_uint) ltime;
cl_uint min = (cl_uint) - 1;
for(int i=0;i<num_src_items;i++)
{
src_ptr[i]=(cl_uint)(b=(a*(b&65535))+(b>>16));
min=src_ptr[i]<min?src_ptr[i]:min;
}
clGetPlatformIDs(1,&platform,NULL);
for(dev=1;dev<NDEVS;dev++)
{
cl_device_id device;
cl_context context;
cl_command_queue queue;
cl_program program;
cl_kernel minp;
cl_kernel reduce;
cl_mem src_buf;
cl_mem dst_buf;
cl_mem dbg_buf;
cl_uint *dst_ptr,*dbg_ptr;
printf("\n%s",dev==0 ? "CPU\n\0":"GPU\n\0");
clGetDeviceIDs(platform,devs[dev],1,&device,NULL);
cl_uint compute_units;
size_t global_work_size;
size_t local_work_size;
size_t num_groups;
clGetDeviceInfo(device,CL_DEVICE_MAX_COMPUTE_UNITS,sizeof(cl_uint),&compute_units,NULL);
printf("%d\n",compute_units);
if(devs[dev]==CL_DEVICE_TYPE_CPU)
{
global_work_size=compute_units*1;
local_work_size=1;
}
else
{
cl_uint ws=64;
global_work_size=compute_units*ws*7;
while((num_src_items/4)%global_work_size!=0) global_work_size+=ws;
local_work_size=ws;
};
num_groups=global_work_size/local_work_size;
context=clCreateContext(NULL,1,&device,NULL,NULL,NULL);
queue=clCreateCommandQueue(context,device,0,NULL);
if(queue==NULL)
{
printf("Compute device setup failed\n");
return(-1);
}
program=clCreateProgramWithSource(context,1,&kernel_source,NULL,NULL);
cl_uint ret=clBuildProgram(program,1,&device,"-save-temps",NULL,NULL);
if(ret!=CL_SUCCESS)
{
printf("clBuildProgram failed: %d\n",ret);
char buf[0x10000];
clGetProgramBuildInfo(program,device,CL_PROGRAM_BUILD_LOG,0x10000,buf,NULL);
printf("\n%s\n",buf);
return(-1);
}
minp=clCreateKernel(program,"minp",NULL);
reduce=clCreateKernel(program,"reduce",NULL);
src_buf=clCreateBuffer(context,CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,num_src_items*sizeof(cl_uint),src_ptr,NULL);
dst_buf=clCreateBuffer(context,CL_MEM_READ_WRITE,num_groups*sizeof(cl_uint),NULL,NULL);
dbg_buf=clCreateBuffer(context,CL_MEM_WRITE_ONLY,global_work_size*sizeof(cl_uint),NULL,NULL);
clSetKernelArg(minp,0,sizeof(void*),(void*)src_buf);
clSetKernelArg(minp,1,sizeof(void*),(void*)dst_buf);
clSetKernelArg(minp,2,1*sizeof(cl_uint),(void*)NULL);
clSetKernelArg(minp,3,sizeof(void*),(void*)dbg_buf);
clSetKernelArg(minp,4,sizeof(num_src_items),(void*)&num_src_items);
clSetKernelArg(minp,5,sizeof(dev),(void*)&dev);
clSetKernelArg(reduce,0,sizeof(void*),(void*)&src_buf);
clSetKernelArg(reduce,1,sizeof(void*),(void*)&dst_buf);
CPerfCounter t;
t.Reset();
t.Start();
#define NLOOPS 500
cl_event ev;
int nloops=NLOOPS;
while(nloops--)
{
clEnqueueNDRangeKernel(queue,minp,1,NULL,&global_work_size,&local_work_size,0,NULL,&ev);
clEnqueueNDRangeKernel(queue,reduce,1,NULL,&num_groups,NULL,1,&ev,NULL);
};
clFinish(queue);
t.Stop();
printf("B/W %.2f GB/sec, ",((float)num_src_items*sizeof(cl_uint)*NLOOPS)/t.GetElapsedTime()/1e9);
dst_ptr=(cl_uint*)clEnqueueMapBuffer(queue,dst_buf,CL_TRUE,CL_MAP_READ,0,num_groups*sizeof(cl_uint),0,NULL,NULL,NULL);
dbg_ptr=(cl_uint*)clEnqueueMapBuffer(queue,dbg_buf,CL_TRUE,CL_MAP_READ,0,global_work_size*sizeof(cl_uint),0,NULL,NULL,NULL);
printf("%d groups,%d threads,count %d,stride %d\n",dbg_ptr[0],dbg_ptr[1],dbg_ptr[2],dbg_ptr[3]);
if(dst_ptr[0]==min)
printf("result correct\n");
else
printf("result incorrect\n");
}
printf("\n");
return 0;
}
падает в enqueuendrangekernel для cpu и gpu
в segmentation fault
вопрос - где я накосячил?
система debian 7.6 amd64, gpu 7870, драйвера 14.4
год тому назад у меня получилось завести эту простыню, а сейчас никак, видно ошибка где-то в вызовах в тексте самого мануала, а где точно не знаю, прошу кого-то пройтись по исходнику и задетектить этого жука, исходник скопировал из редактора один в один, число исполнительных юнитов определяет верно, для цпу 1, для гпу 20, значит работает сс нужным девайсом, но что ему надо, уже сломал чашку, готовя кофе..?!! |
Всего записей: 113 | Зарегистр. 26-07-2014 | Отправлено: 00:49 18-09-2014 | Исправлено: lormutryas, 07:02 18-09-2014 |
|
