ホストプログラム
[cpp]
#include<iostream>
#include<cstdio>
#include<cstdlib>
#include<CL/cl.h>
#include<string>
#define MAX_SOURCE_SIZE (0x100000)
#define MEM_SIZE (1000)
using namespace std;
int main()
{
//char string[MEM_SIZE];
FILE *fp;
char fileName[] = "./calc.cl";
char *source_str;
size_t source_size;
fp = fopen(fileName, "r");
if(!fp)
{
cout << "file failed" <<endl;
exit(1);
}
source_str = (char *)malloc(MAX_SOURCE_SIZE);
source_size = fread(source_str,1,MAX_SOURCE_SIZE,fp);
fclose(fp);
cl_int status;
cl_platform_id platforms[2];
cl_uint num_platforms;
status = clGetPlatformIDs(2, platforms, &num_platforms);
//cout << status << endl;
cl_device_id devices;
cl_uint num_devices;
cl_context context;
status = clGetDeviceIDs(platforms[1], CL_DEVICE_TYPE_GPU, 4, &devices, &num_devices);
//cout << status << endl;
context = clCreateContext(NULL, 4, &devices, NULL, NULL, &status);
//cout << status << endl;
cl_command_queue queue;
queue = clCreateCommandQueue(context, devices, 0, &status);
//cout <<"queue:"<<status << endl;
cl_program program;
program = clCreateProgramWithSource(context, 1, (const char**)&source_str, &source_size, &status);
//cout << status << endl;
status = clBuildProgram(program, num_devices, &devices, NULL, NULL, NULL);
//cout << status << endl;
/*////////////////////////
size_t logsize;
status = clGetProgramBuildInfo(program,devices,CL_PROGRAM_BUILD_LOG,0,NULL,&logsize);
if(status == CL_SUCCESS)
{
char *logbuffer;
logbuffer = new char[logsize + 1];
if(logbuffer == NULL)
{
cout << "memory allocation failed."<<endl;
//return;
}
status = clGetProgramBuildInfo(program,devices,CL_PROGRAM_BUILD_LOG,logsize,logbuffer,NULL);
if(status == CL_SUCCESS)
{
logbuffer[logsize]=’\0’;
cout << "build log" << endl;
cout << logbuffer << endl;
}
delete [] logbuffer;
}
*////////////////////////
cl_kernel kernel;
kernel = clCreateKernel(program, "calc", &status);
//cout << status << endl;
float Outx[MEM_SIZE];
float Outy[MEM_SIZE];
bool OutCheck[MEM_SIZE][MEM_SIZE];
Outx[0] = -2;
Outy[0] = -2;
for(int i= 0; i < MEM_SIZE-1 ;i++)
{
Outx[i+1] = Outx[i] + (4.0 / MEM_SIZE);
Outy[i+1] = Outy[i] + (4.0 / MEM_SIZE);
for(int j = 0; j < MEM_SIZE ;j++)
{
OutCheck[i][j] = 0;
}
}
cl_mem memoutx, memouty, memoutcheck;
memoutx = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(float), NULL, &status);
//cout <<status<<endl;
memouty = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * sizeof(float), NULL, &status);
//cout <<status<<endl;
memoutcheck = clCreateBuffer(context, CL_MEM_READ_WRITE, MEM_SIZE * MEM_SIZE * sizeof(bool), NULL, &status);
//cout <<status<<endl;
status = clEnqueueWriteBuffer( queue, memoutx, CL_TRUE, 0, MEM_SIZE * sizeof( float ), Outx, 0, NULL, NULL );
//cout << status << endl;
status = clEnqueueWriteBuffer( queue, memouty, CL_TRUE, 0, MEM_SIZE * sizeof( float ), Outy, 0, NULL, NULL );
//cout << status << endl;
status = clEnqueueWriteBuffer( queue, memoutcheck, CL_TRUE, 0, MEM_SIZE * MEM_SIZE * sizeof( bool ), OutCheck, 0, NULL, NULL );
//cout << status << endl;
status = clSetKernelArg(kernel, 0,sizeof(cl_mem),(void *)&memoutx);
status = clSetKernelArg(kernel, 1,sizeof(cl_mem),(void *)&memouty);
status = clSetKernelArg(kernel, 2,sizeof(cl_mem),(void *)&memoutcheck);
//cout <<"カーネルに渡す引数セット："<<status<<endl;
//status = clEnqueueTask(queue, kernel, 0, NULL, NULL);
size_t global_work_size[] = {MEM_SIZE * MEM_SIZE};
size_t local_work_size[] = {MEM_SIZE};
status = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
if( status < 0){
cout <<"実行："<<status<<endl;
cout << CL_DEVICE_MAX_WORK_GROUP_SIZE << endl;
exit(0);
}
//status = clEnqueueReadBuffer(queue, memoutx, CL_TRUE, 0, MEM_SIZE * sizeof(float), Out, 0, NULL, NULL);
//status = clEnqueueReadBuffer(queue, memouty, CL_TRUE, 0, MEM_SIZE * sizeof(float), Out, 0, NULL, NULL);
status = clEnqueueReadBuffer(queue, memoutcheck, CL_TRUE, 0, MEM_SIZE * MEM_SIZE * sizeof(bool), OutCheck, 0, NULL, NULL);
//cout <<"結果取り出し："<<status<<endl;
for(int i = 0;i < MEM_SIZE; i++)
{
for(int j = 0;j < MEM_SIZE; j++)
{
if(OutCheck[i][j] == 0){
cout << Outx[i]<< " " << Outy[j] << endl;
}
}
}
clReleaseMemObject(memoutx);
clReleaseMemObject(memouty);
clReleaseMemObject(memoutcheck);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseCommandQueue(queue);
clReleaseContext(context);
}
[/cpp]
カーネルプログラム
[cpp]
#define MEM_SIZE (1000)
__kernel void
calc(__global float *outx,
__global float *outy,
__global bool outcheck[MEM_SIZE][MEM_SIZE] )
{
#define MaxCalcNum 100
int width = 2, height = 2, j = 0;
int out[MEM_SIZE*MEM_SIZE];
#pragma OPENCL EXTENSION cl_khr_fp64:enable
//int index_x = get_global_id(0);
//int index_y = get_global_id(1);
int index_x = get_group_id(0);
int index_y = get_local_id(0);
if(index_x > MEM_SIZE)index_x = 1;
if(index_y > MEM_SIZE)index_y = 1;
double prevX = 0.0f;
double prevY = 0.0f;
for(int i = 0; i < MaxCalcNum ; i++)
{
double newX = (prevX * prevX) – (prevY * prevY) + outx[index_x];
double newY = 2.0f * prevX * prevY + outy[index_y];
if ((newX * newX) + (newY * newY) > 4.0f)
{
outcheck[index_x][index_y] = -1;
break;
}
else
{
prevX = newX;
prevY = newY;
}
}
}
[/cpp]
画像出力
$ g++ -o main_cp main_cp.cpp -lOpenCL
$ ./main_cp > mandel.txt
$ gnuplot
> set terminal png
> set out "mandel.png"
> plot "mandel.txt"