1
0
Fork 0

inicializacny commit

climage
Dušan Poizl 13 years ago
commit 5449500dd1
  1. 17
      cellspace.cpp
  2. 34
      cellspace.h
  3. 183
      clutility.cpp
  4. 12
      clutility.h
  5. 18
      head.cl
  6. 94
      kernel.cl
  7. 62
      machine.cpp
  8. 24
      machine.h
  9. 8
      main.cpp
  10. 51
      model.cpp
  11. 23
      model.h

@ -0,0 +1,17 @@
#include "cellspace.h"
CellSpace::CellSpace(cl_context context, int width, int height, int num_neighbors, int len_queue)
{
cl_int err_code;
space = clCreateBuffer(context, CL_MEM_READ_WRITE, 1000*sizeof(float), NULL, &err_code);
if(err_code != CL_SUCCESS)
{
cerr << "Nepodarilo sa vytvorit buffer " << err_code << endl;
throw err_code;
}
}
CellSpace::~CellSpace()
{
clReleaseMemObject(space);
}

@ -0,0 +1,34 @@
#ifndef _CELLSPACE_H_NOU_
#define _CELLSPACE_H_NOU_
#include <CL/cl.h>
#include <iostream>
using namespace std;
/*typedef struct
{
float time;
float state;
}event;
typedef struct
{
float state;
event neighbors[8][4];
int cir_buf[8][2];
}cell;*/
class CellSpace
{
friend class Model;
private:
cl_mem space;
int w,h,num_neighbors,length_queue;
public:
CellSpace(cl_context context, int width, int height, int num_neighbors, int len_queue);
~CellSpace();
void renderImage(unsigned char *data, int width, int height);
};
#endif // _CELLSPACE_H_NOU_

@ -0,0 +1,183 @@
#include "clutility.h"
void getDeviceInfo(cl_device_id device)
{
cl_uint uint;
cl_ulong ulong;
size_t size;
cl_bool b;
cl_device_fp_config fp_conf;
cl_device_mem_cache_type mem_cache_type;
cl_device_type type;
char info[10000];
clGetDeviceInfo(device, CL_DEVICE_VENDOR_ID, sizeof(cl_uint), &uint, NULL);
cout << "Device vendor ID: " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_TYPE, sizeof(cl_device_type), &type, NULL);
cout << "Device type: ";
if(type & CL_DEVICE_TYPE_DEFAULT)cout << "DEFAULT ";
if(type & CL_DEVICE_TYPE_CPU)cout << "CPU ";
if(type & CL_DEVICE_TYPE_GPU)cout << "GPU ";
if(type & CL_DEVICE_TYPE_ACCELERATOR)cout << "ACCELERATOR ";
cout << endl;
clGetDeviceInfo(device, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), &uint, NULL);
cout << "Max compute units: " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS, sizeof(cl_uint), &uint, NULL);
cout << "Max work item dimensions: " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), &size, NULL);
cout << "Max work group size: " << size << endl;
clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(size_t), &size, NULL);
cout << "Max work item sizes: " << size << endl;
clGetDeviceInfo(device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR, sizeof(cl_uint), &uint, NULL);
cout << "Prefered vector width char: " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT, sizeof(cl_uint), &uint, NULL);
cout << "Prefered vector width short: " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT, sizeof(cl_uint), &uint, NULL);
cout << "Prefered vector width int: " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG, sizeof(cl_uint), &uint, NULL);
cout << "Prefered vector width long: " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT, sizeof(cl_uint), &uint, NULL);
cout << "Prefered vector width float: " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE, sizeof(cl_uint), &uint, NULL);
cout << "Prefered vector width double: " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_MAX_CLOCK_FREQUENCY, sizeof(cl_uint), &uint, NULL);
cout << "Max clock frequency: " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_ADDRESS_BITS, sizeof(cl_uint), &uint, NULL);
cout << "Address bits: " << uint << endl;
/*clGetDeviceInfo(device, CL_DEVICE_MAX_READ_IMAGE_ARGS, sizeof(), &, NULL);
cout << ": " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_MAX_WRITE_IMAGE_ARGS, sizeof(), &, NULL);
cout << ": " << uint << endl;*/
clGetDeviceInfo(device, CL_DEVICE_MAX_MEM_ALLOC_SIZE, sizeof(cl_ulong), &ulong, NULL);
cout << "Max mem alloc size: " << ulong/1024/1024 << " MiB" << endl;
/*clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof(), &, NULL);
cout << ": " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof(), &, NULL);
cout << ": " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_IMAGE3D_MAX_WIDTH, sizeof(), &, NULL);
cout << ": " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_IMAGE3D_MAX_HEIGHT, sizeof(), &, NULL);
cout << ": " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_IMAGE3D_MAX_DEPTH, sizeof(), &, NULL);
cout << ": " << uint << endl;*/
clGetDeviceInfo(device, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &b, NULL);
cout << "Image support: " << b << endl;
clGetDeviceInfo(device, CL_DEVICE_MAX_PARAMETER_SIZE, sizeof(size_t), &size, NULL);
cout << "Max parameter size: " << size << endl;
/*clGetDeviceInfo(device, CL_DEVICE_MAX_SAMPLERS, sizeof(), &, NULL);
cout << ": " << uint << endl;*/
clGetDeviceInfo(device, CL_DEVICE_MEM_BASE_ADDR_ALIGN, sizeof(cl_uint), &uint, NULL);
cout << "Mem base address align: " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE, sizeof(cl_uint), &uint, NULL);
cout << "Min data type align size: " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_SINGLE_FP_CONFIG, sizeof(cl_device_fp_config), &fp_conf, NULL);
cout << "Single FP config: ";
if(fp_conf & CL_FP_DENORM)cout << "DENORM ";
if(fp_conf & CL_FP_INF_NAN)cout << "INF_NAN ";
if(fp_conf & CL_FP_ROUND_TO_NEAREST)cout << "ROUND_TO_NEAREST ";
if(fp_conf & CL_FP_ROUND_TO_ZERO)cout << "ROUND_TO_ZERO ";
if(fp_conf & CL_FP_ROUND_TO_INF)cout << "ROUND_TO_INF ";
if(fp_conf & CL_FP_FMA)cout << "FMA ";
cout << endl;
clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_CACHE_TYPE, sizeof(cl_device_mem_cache_type), &mem_cache_type, NULL);
cout << "Global mem cache type: ";
if(mem_cache_type == CL_NONE)cout << "NONE ";
if(mem_cache_type == CL_READ_ONLY_CACHE)cout << "READ_ONLY_CACHE ";
if(mem_cache_type == CL_READ_WRITE_CACHE)cout << "READ_WRITE_CACHE ";
cout << endl;
clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE, sizeof(cl_uint), &uint, NULL);
cout << "Global mem cacheline size: " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_CACHE_SIZE, sizeof(cl_ulong), &ulong, NULL);
cout << "Global mem cache size: " << ulong/1024 << " kiB" << endl;
clGetDeviceInfo(device, CL_DEVICE_GLOBAL_MEM_SIZE, sizeof(cl_ulong), &ulong, NULL);
cout << "Global mem size: " << ulong/1024/1024 << " MiB" << endl;
clGetDeviceInfo(device, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE, sizeof(cl_ulong), &ulong, NULL);
cout << "Max constant buffer size: " << ulong/1024 << " kiB" << endl;
clGetDeviceInfo(device, CL_DEVICE_MAX_CONSTANT_ARGS, sizeof(cl_ulong), &ulong, NULL);
cout << "Max constant args: " << ulong << endl;
clGetDeviceInfo(device, CL_DEVICE_LOCAL_MEM_TYPE, sizeof(cl_device_local_mem_type), &mem_cache_type, NULL);
cout << "Local mem type: ";
if(mem_cache_type == CL_LOCAL)cout << "LOCAL ";
if(mem_cache_type == CL_GLOBAL)cout << "GLOBAL ";
cout << endl;
clGetDeviceInfo(device, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(cl_ulong), &ulong, NULL);
cout << "Local mem size: " << ulong/1024 << " kiB" << endl;
/*clGetDeviceInfo(device, CL_DEVICE_ERROR_CORRECTION_SUPPORT, sizeof(), &, NULL);
cout << ": " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_PROFILING_TIMER_RESOLUTION, sizeof(), &, NULL);
cout << ": " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_ENDIAN_LITTLE, sizeof(), &, NULL);
cout << ": " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_AVAILABLE, sizeof(), &, NULL);
cout << ": " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_COMPILER_AVAILABLE, sizeof(), &, NULL);
cout << ": " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_EXECUTION_CAPABILITIES, sizeof(), &, NULL);
cout << ": " << uint << endl;
clGetDeviceInfo(device, CL_DEVICE_QUEUE_PROPERTIES, sizeof(), &, NULL);
cout << ": " << uint << endl;*/
clGetDeviceInfo(device, CL_DEVICE_NAME, 10000, &info, NULL);
cout << "Name: " << info << endl;
clGetDeviceInfo(device, CL_DEVICE_VENDOR, 10000, &info, NULL);
cout << "Vendor: " << info << endl;
clGetDeviceInfo(device, CL_DRIVER_VERSION, 10000, &info, NULL);
cout << "Driver version: " << info << endl;
clGetDeviceInfo(device, CL_DEVICE_PROFILE, 10000, &info, NULL);
cout << "Profile: " << info << endl;
clGetDeviceInfo(device, CL_DEVICE_VERSION, 10000, &info, NULL);
cout << "Version: " << info << endl;
clGetDeviceInfo(device, CL_DEVICE_EXTENSIONS, 10000, &info, NULL);
cout << "Extensions: " << info << endl;
}
void getError(cl_int err_code)
{
switch(err_code)
{
case CL_DEVICE_NOT_FOUND: cerr << "DEVICE_NOT_FOUND" << endl; break;
case CL_DEVICE_NOT_AVAILABLE: cerr << "DEVICE_NOT_AVAILABLE" << endl; break;
case CL_COMPILER_NOT_AVAILABLE: cerr << "COMPILER_NOT_AVAILABLE" << endl; break;
case CL_MEM_OBJECT_ALLOCATION_FAILURE: cerr << "MEM_OBJECT_ALLOCATION_FAILURE" << endl; break;
case CL_OUT_OF_RESOURCES: cerr << "OUT_OF_RESOURCES" << endl; break;
case CL_OUT_OF_HOST_MEMORY: cerr << "OUT_OF_HOST_MEMORY" << endl; break;
case CL_PROFILING_INFO_NOT_AVAILABLE: cerr << "PROFILING_INFO_NOT_AVAILABLE" << endl; break;
case CL_MEM_COPY_OVERLAP: cerr << "MEM_COPY_OVERLAP" << endl; break;
case CL_IMAGE_FORMAT_MISMATCH: cerr << "IMAGE_FORMAT_MISMATCH" << endl; break;
case CL_IMAGE_FORMAT_NOT_SUPPORTED: cerr << "IMAGE_FORMAT_NOT_SUPPORTED" << endl; break;
case CL_BUILD_PROGRAM_FAILURE: cerr << "BUILD_PROGRAM_FAILURE" << endl; break;
case CL_MAP_FAILURE: cerr << "MAP_FAILURE" << endl; break;
case CL_INVALID_VALUE: cerr << "INVALID_VALUE" << endl; break;
case CL_INVALID_DEVICE_TYPE: cerr << "INVALID_DEVICE_TYPE" << endl; break;
case CL_INVALID_PLATFORM: cerr << "INVALID_PLATFORM" << endl; break;
case CL_INVALID_DEVICE: cerr << "INVALID_DEVICE" << endl; break;
case CL_INVALID_CONTEXT: cerr << "INVALID_CONTEXT" << endl; break;
case CL_INVALID_QUEUE_PROPERTIES: cerr << "INVALID_QUEUE_PROPERTIES" << endl; break;
case CL_INVALID_COMMAND_QUEUE: cerr << "INVALID_COMMAND_QUEUE" << endl; break;
case CL_INVALID_HOST_PTR: cerr << "INVALID_HOST_PTR" << endl; break;
case CL_INVALID_MEM_OBJECT: cerr << "INVALID_MEM_OBJECT" << endl; break;
case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR: cerr << "INVALID_IMAGE_FORMAT_DESCRIPTOR" << endl; break;
case CL_INVALID_IMAGE_SIZE: cerr << "INVALID_IMAGE_SIZE" << endl; break;
case CL_INVALID_SAMPLER: cerr << "INVALID_SAMPLER" << endl; break;
case CL_INVALID_BINARY: cerr << "INVALID_BINARY" << endl; break;
case CL_INVALID_BUILD_OPTIONS: cerr << "INVALID_BUILD_OPTIONS" << endl; break;
case CL_INVALID_PROGRAM: cerr << "INVALID_PROGRAM" << endl; break;
case CL_INVALID_PROGRAM_EXECUTABLE: cerr << "INVALID_PROGRAM_EXECUTABLE" << endl; break;
case CL_INVALID_KERNEL_NAME: cerr << "INVALID_KERNEL_NAME" << endl; break;
case CL_INVALID_KERNEL_DEFINITION: cerr << "INVALID_KERNEL_DEFINITION" << endl; break;
case CL_INVALID_KERNEL: cerr << "INVALID_KERNEL" << endl; break;
case CL_INVALID_ARG_INDEX: cerr << "INVALID_ARG_INDEX" << endl; break;
case CL_INVALID_ARG_VALUE: cerr << "INVALID_ARG_VALUE" << endl; break;
case CL_INVALID_ARG_SIZE: cerr << "INVALID_ARG_SIZE" << endl; break;
case CL_INVALID_KERNEL_ARGS: cerr << "INVALID_KERNEL_ARGS" << endl; break;
case CL_INVALID_WORK_DIMENSION: cerr << "INVALID_WORK_DIMENSION" << endl; break;
case CL_INVALID_WORK_GROUP_SIZE: cerr << "INVALID_WORK_GROUP_SIZE" << endl; break;
case CL_INVALID_WORK_ITEM_SIZE: cerr << "INVALID_WORK_ITEM_SIZE" << endl; break;
case CL_INVALID_GLOBAL_OFFSET: cerr << "INVALID_GLOBAL_OFFSET" << endl; break;
case CL_INVALID_EVENT_WAIT_LIST: cerr << "INVALID_EVENT_WAIT_LIST" << endl; break;
case CL_INVALID_EVENT: cerr << "INVALID_EVENT" << endl; break;
case CL_INVALID_OPERATION: cerr << "INVALID_OPERATION" << endl; break;
case CL_INVALID_GL_OBJECT: cerr << "INVALID_GL_OBJECT" << endl; break;
case CL_INVALID_BUFFER_SIZE: cerr << "INVALID_BUFFER_SIZE" << endl; break;
case CL_INVALID_MIP_LEVEL: cerr << "INVALID_MIP_LEVEL" << endl; break;
default: break;
}
}

@ -0,0 +1,12 @@
#ifndef _CLUTILITY_H_NOU_
#define _CLUTILITY_H_NOU_
#include <iostream>
#include <CL/cl.h>
using namespace std;
void getDeviceInfo(cl_device_id device);
void getError(cl_int err_code);
#endif // _CLUTILITY_H_NOU_

@ -0,0 +1,18 @@
typedef struct
{
int a[4];
float b[4];
} struktura;
typedef struct
{
float time;
float state;
}event;
typedef struct
{
float state;
event neighbors[8][4];
int cir_buf[8][2];
}cell;

@ -0,0 +1,94 @@
#define NUM_NEIGHBORS 8
typedef struct
{
float time;
float state;
}event;
typedef struct
{
event events[4];
int num;
}queue;
typedef struct
{
float state;
queue queues[NUM_NEIGHBORS];
}cell;
void neighbors_fill(float *neighbors)
{
}
float transition()
{
/*rule : 1 100 { (0,0) = 1 and (truecount = 3 or truecount = 4) }
rule : 1 100 { (0,0) = 0 and truecount = 2 }
rule : 0 100 { t }
if()*/
return 0.0f;
}
int wrap(int x, int y)
{
int r = x%y;
if(r<0)return y+r;
else return r;
}
__kernel void size()
{
int r = remainder(-22, 5);
printf("hello %d\n", r);
}
__kernel void take_step(__global cell *cellspace)
{
size_t x = get_global_id(0);
size_t y = get_global_id(0);
float susedia[9];
neighbors_fill(susedia);
printf("%f %f %f\n", susedia[0], susedia[1], susedia[2]);
}
__kernel void init(__global cell *cell_space, __constant int *neighbors, __global float *init_state)
{
size_t x = get_global_id(0);
size_t y = get_global_id(1);
size_t width = get_global_size(0);
size_t height = get_global_size(1);
cell_space[x+y*height].state = init_state[x+y*height];
for(int i=0;i<NUM_NEIGHBORS;i++)
{
size_t pos = wrap(x+neighbors[i*2], width)+wrap(y+neighbors[i*2+1], height)*width;
cell_space[pos].queues[i].events[0].state = init_state[x*width+y];
cell_space[pos].queues[i].events[0].time = 0.0f;
cell_space[pos].queues[i].num = 1;
}
}
/* Usporiada udalosti v jednotlivych frontach podla casu */
__kernel void sort_queue(__global cell *cell_space)
{
size_t x = get_global_id(0);
size_t y = get_global_id(1);
size_t width = get_global_size(0);
__global cell *c = cell_space+x+y*width;
int min;
event ev;
for(int i=0;i<NUM_NEIGHBORS;i++)//prejdem vsetkych susedov
{
for(int o=0;o<c->queues[i].num;o++)
{
min = o;
for(int p=o+1;p<c->queues[i].num;p++)if(c->queues[i].events[p].time < c->queues[i].events[min].time)min=p;//najdem udalost s najmensim casom
ev = c->queues[i].events[min];//vymenim
c->queues[i].events[min] = c->queues[i].events[o];
c->queues[i].events[o] = ev;
}
}
}

@ -0,0 +1,62 @@
#include "machine.h"
Machine::Machine()
{
cl_int err_code;
context = clCreateContextFromType(NULL, CL_DEVICE_TYPE_ALL, NULL, NULL, &err_code);
if(err_code != CL_SUCCESS)
{
cerr << "Nepodarilo sa vytvorit kontext ";getError(err_code);
throw err_code;
}
clGetContextInfo(context, CL_CONTEXT_DEVICES, sizeof(cl_device_id), &device, NULL);
getDeviceInfo(device);
model = new Model(context);
try
{
model->loadModel();
}
catch(cl_int err)
{
if(err == CL_BUILD_PROGRAM_FAILURE)
{
char build_log[10000];
clGetProgramBuildInfo(model->program, device, CL_PROGRAM_BUILD_LOG, 10000, build_log, NULL);
cerr << "Build log: " << endl << build_log << endl;
}
else getError(err_code);
return;
}
model->setKernel("size");
cout << "kernel vytvoreny" << endl;
queue = clCreateCommandQueue(context, device, NULL, &err_code);
if(err_code != CL_SUCCESS)
{
cerr << "Nepodarilo sa vytvorit prikazovu frontu ";getError(err_code);
throw err_code;
}
size_t global_s[] = {1};
err_code = clEnqueueNDRangeKernel(queue, model->kernel, 1, NULL, global_s, NULL, NULL, NULL, NULL);
if(err_code != CL_SUCCESS)
{
cerr << "Nepodarilo sa spustit kernel ";getError(err_code);
throw err_code;
}
clFinish(queue);
cout << "koniec" << endl;
cin.get();
}
Machine::~Machine()
{
delete model;
clReleaseCommandQueue(queue);
clReleaseContext(context);
}
Model* Machine::getModel()
{
return model;
}

@ -0,0 +1,24 @@
#ifndef _MACHINE_H_NOU_
#define _MACHINE_H_NOU_
#include <iostream>
#include <CL/cl.h>
#include "clutility.h"
#include "model.h"
using namespace std;
class Machine
{
private:
cl_context context;
cl_device_id device;
cl_command_queue queue;
Model *model;
public:
Machine();
~Machine();
Model* getModel();
};
#endif

@ -0,0 +1,8 @@
#include <iostream>
#include "machine.h"
int main(int argc, char* argv[])
{
Machine automata;
return 0;
}

@ -0,0 +1,51 @@
#include "model.h"
Model::Model(cl_context context)
{
this->parent_context = context;
cell_space = new CellSpace(context, 512, 512, 8, 4);
program = NULL;
kernel = NULL;
}
Model::~Model()
{
delete cell_space;
clReleaseProgram(program);
clReleaseKernel(kernel);
}
void Model::loadModel()
{
cl_int err_code;
const char *kernel_source = "#include \"kernel.cl\"\n";
if(kernel)clReleaseKernel(kernel);
if(program)clReleaseProgram(program);
program = clCreateProgramWithSource(parent_context, 1, &kernel_source, NULL, &err_code);
if(err_code != CL_SUCCESS)
{
cerr << "Nepodarilo sa vytvorit program " << err_code << endl;
throw err_code;
}
err_code = clBuildProgram(program, NULL, NULL, "-I.", NULL, NULL);
if(err_code != CL_SUCCESS)
{
cerr << "Nepodarilo sa skompilovat program " << err_code << endl;
throw err_code;
}
}
void Model::setKernel(const char *name)
{
if(kernel)clReleaseKernel(kernel);
cl_int err_code;
kernel = clCreateKernel(program, name, &err_code);
if(err_code != CL_SUCCESS)
{
cerr << "Nepodarilo sa vytvorit kernel " << err_code << endl;
throw err_code;
}
}

@ -0,0 +1,23 @@
#ifndef _MODEL_H_NOU_
#define _MODEL_H_NOU_
#include "cellspace.h"
class Model
{
friend class Machine;
private:
CellSpace *cell_space;
cl_context parent_context;
cl_kernel kernel;
cl_program program;
public:
Model(cl_context context);
~Model();
void loadModel();
void setKernel(const char *name);
protected:
};
#endif // _MODEL_H_NOU_
Loading…
Cancel
Save