# HG changeset patch # User John Tsiombikas # Date 1279910923 -3600 # Node ID deaf85acf6aff4694ba2b5f4fb5cbf882ab525a0 # Parent 575383f3a239f99fdc42d1add6606d0b7c99b937 interactive spheres diff -r 575383f3a239 -r deaf85acf6af Makefile --- a/Makefile Fri Jul 23 01:22:03 2010 +0100 +++ b/Makefile Fri Jul 23 19:48:43 2010 +0100 @@ -4,7 +4,15 @@ CXX = g++ CXXFLAGS = -pedantic -Wall -g -LDFLAGS = -framework OpenGL -framework GLUT -framework OpenCL +LDFLAGS = $(libgl) $(libcl) + +ifeq ($(shell uname -s), Darwin) + libgl = -framework OpenGL -framework GLUT + libcl = -framework OpenCL +else + libgl = -lGL -lglut + libcl = -lOpenCL +endif $(bin): $(obj) $(CXX) -o $@ $(obj) $(LDFLAGS) diff -r 575383f3a239 -r deaf85acf6af rt.cl --- a/rt.cl Fri Jul 23 01:22:03 2010 +0100 +++ b/rt.cl Fri Jul 23 19:48:43 2010 +0100 @@ -21,20 +21,18 @@ struct SurfPoint { float t; - float3 pos, norm; + float4 pos, norm; global const struct Sphere *obj; }; -struct Matrix4x4 { - float m[16]; -}; - #define EPSILON 1e-6 float4 shade(struct Ray ray, struct SurfPoint sp, global const struct Light *lights, int num_lights); bool intersect(struct Ray ray, global const struct Sphere *sph, struct SurfPoint *sp); -float3 reflect(float3 v, float3 n); +float4 reflect(float4 v, float4 n); +float4 transform(float4 v, global const float *xform); +struct Ray transform_ray(global const struct Ray *ray, global const float *xform); kernel void render(global float4 *fb, @@ -42,13 +40,13 @@ global const struct Sphere *sphlist, global const struct Light *lights, global const struct Ray *primrays, - global const struct Matrix4x4 xform) + global const float *xform) { int idx = get_global_id(0); - struct Ray ray = primrays[idx]; + struct Ray ray = transform_ray(primrays + idx, xform); + struct SurfPoint sp, sp0; - sp0.t = FLT_MAX; sp0.obj = 0; @@ -68,31 +66,31 @@ float4 shade(struct Ray ray, struct SurfPoint sp, global const struct Light *lights, int num_lights) { - float3 dcol = (float3)(0, 0, 0); - float3 scol = (float3)(0, 0, 0); + float4 dcol = (float4)(0, 0, 0, 0); + float4 scol = (float4)(0, 0, 0, 0); for(int i=0; ispow); - dcol += sp.obj->kd.xyz * diff * lights[i].color.xyz; - scol += sp.obj->ks.xyz * spec * lights[i].color.xyz; + dcol += sp.obj->kd * diff * lights[i].color; + scol += sp.obj->ks * spec * lights[i].color; } - return (float4)(dcol + scol, 1.0f); + return dcol + scol; } bool intersect(struct Ray ray, global const struct Sphere *sph, struct SurfPoint *sp) { - float3 dir = ray.dir.xyz; - float3 orig = ray.origin.xyz; - float3 spos = sph->pos.xyz; + float4 dir = ray.dir; + float4 orig = ray.origin; + float4 spos = sph->pos; float a = dot(dir, dir); float b = 2.0 * dir.x * (orig.x - spos.x) + @@ -123,7 +121,33 @@ return true; } -float3 reflect(float3 v, float3 n) +float4 reflect(float4 v, float4 n) { return 2.0f * dot(v, n) * n - v; } + +float4 transform(float4 v, global const float *xform) +{ + float4 res; + res.x = v.x * xform[0] + v.y * xform[4] + v.z * xform[8] + xform[12]; + res.y = v.x * xform[1] + v.y * xform[5] + v.z * xform[9] + xform[13]; + res.z = v.x * xform[2] + v.y * xform[6] + v.z * xform[10] + xform[14]; + res.w = 1.0; + return res; +} + +struct Ray transform_ray(global const struct Ray *ray, global const float *xform) +{ + struct Ray res; + float rot[16]; + + for(int i=0; i<16; i++) { + rot[i] = xform[i]; + } + rot[3] = rot[7] = rot[11] = rot[12] = rot[13] = rot[14] = 0.0f; + rot[15] = 1.0f; + + res.origin = transform(ray->origin, xform); + res.dir = transform(ray->dir, xform); + return res; +} diff -r 575383f3a239 -r deaf85acf6af src/clray.cc --- a/src/clray.cc Fri Jul 23 01:22:03 2010 +0100 +++ b/src/clray.cc Fri Jul 23 19:48:43 2010 +0100 @@ -21,6 +21,9 @@ static int xsz, ysz; static bool need_update = true; +static float cam_theta, cam_phi = 25.0; +static float cam_dist = 10.0; + int main(int argc, char **argv) { glutInitWindowSize(800, 600); @@ -63,7 +66,21 @@ void disp() { + glMatrixMode(GL_MODELVIEW); + glLoadIdentity(); + if(need_update) { + float mat[16]; + + glPushMatrix(); + glRotatef(cam_theta, 0, 1, 0); + glRotatef(cam_phi, 1, 0, 0); + glTranslatef(0, 0, -cam_dist); + + glGetFloatv(GL_MODELVIEW_MATRIX, mat); + set_xform(mat); + glPopMatrix(); + render(); glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, xsz, ysz, GL_RGBA, GL_FLOAT, fb); need_update = false; @@ -115,12 +132,47 @@ } } +static bool bnstate[32]; +static int prev_x, prev_y; + void mouse(int bn, int state, int x, int y) { + if(state == GLUT_DOWN) { + prev_x = x; + prev_y = y; + bnstate[bn] = true; + } else { + bnstate[bn] = false; + } } +#define ROT_SCALE 0.5 +#define PAN_SCALE 0.1 + void motion(int x, int y) { + int dx = x - prev_x; + int dy = y - prev_y; + prev_x = x; + prev_y = y; + + if(bnstate[0]) { + cam_theta += dx * ROT_SCALE; + cam_phi += dy * ROT_SCALE; + + if(cam_phi < -89) cam_phi = 89; + if(cam_phi > 89) cam_phi = 89; + + need_update = true; + glutPostRedisplay(); + } + if(bnstate[2]) { + cam_dist += dy * PAN_SCALE; + if(cam_dist < 0) cam_dist = 0; + + need_update = true; + glutPostRedisplay(); + } } bool write_ppm(const char *fname, float *fb, int xsz, int ysz) diff -r 575383f3a239 -r deaf85acf6af src/mesh.cc --- a/src/mesh.cc Fri Jul 23 01:22:03 2010 +0100 +++ b/src/mesh.cc Fri Jul 23 19:48:43 2010 +0100 @@ -2,6 +2,7 @@ #include #include #include +#include #include #include #include @@ -9,7 +10,7 @@ using namespace std; -#define COMMANDS \ +#define COMMANDS \ CMD(V), \ CMD(VN), \ CMD(VT), \ @@ -136,7 +137,7 @@ char cur_name[16]; obj_file obj; - + sprintf(cur_name, "default%02d.obj", seq++); obj.cur_obj = cur_name; @@ -144,7 +145,7 @@ for(;;) { Vector3 vec; obj_face face; - + char line[BUF_SZ]; fgets(line, sizeof line, fp); if(feof(fp)) { @@ -246,7 +247,7 @@ mat.kt = 1.0 - vmtl[i].alpha; mat.kr = 0.0; // TODO mat.spow = vmtl[i].shininess; - + matnames[vmtl[i].name] = i; } } @@ -356,7 +357,7 @@ if(added_tc) { obj->vt.pop_back(); } - + return mesh; } @@ -445,7 +446,7 @@ static int get_cmd(char *str) { char *s = str; - while((*s = toupper(*s))) s++; + while((*s = toupper(*s))) s++; for(int i=0; cmd_names[i]; i++) { if(strcmp(str, cmd_names[i]) == 0) { diff -r 575383f3a239 -r deaf85acf6af src/ocl.cc --- a/src/ocl.cc Fri Jul 23 01:22:03 2010 +0100 +++ b/src/ocl.cc Fri Jul 23 19:48:43 2010 +0100 @@ -1,10 +1,14 @@ +#define OCL_CC_ + #include #include #include +#include #include #include #include #include "ocl.h" +#include "ocl_errstr.h" class InitCL { @@ -31,6 +35,7 @@ static int devcmp(struct device_info *a, struct device_info *b); static const char *devtypestr(cl_device_type type); static void print_memsize(FILE *out, unsigned long memsz); +static const char *clstrerror(int err); static InitCL initcl; @@ -72,7 +77,7 @@ if(!(mem = clCreateBuffer(ctx, rdwr | CL_MEM_USE_HOST_PTR, sz, buf, &err))) { - fprintf(stderr, "failed to create memory buffer (%d)\n", err); + fprintf(stderr, "failed to create memory buffer: %s\n", clstrerror(err)); return 0; } @@ -98,7 +103,7 @@ int err; mbuf->ptr = clEnqueueMapBuffer(cmdq, mbuf->mem, 1, rdwr, 0, mbuf->size, 0, 0, 0, &err); if(!mbuf->ptr) { - fprintf(stderr, "failed to map buffer (%d)\n", err); + fprintf(stderr, "failed to map buffer: %s\n", clstrerror(err)); return 0; } return mbuf->ptr; @@ -116,7 +121,7 @@ int err; if((err = clEnqueueWriteBuffer(cmdq, mbuf->mem, 1, 0, sz, src, 0, 0, 0)) != 0) { - fprintf(stderr, "failed to write buffer (%d)\n", err); + fprintf(stderr, "failed to write buffer: %s\n", clstrerror(err)); return false; } return true; @@ -128,7 +133,7 @@ int err; if((err = clEnqueueReadBuffer(cmdq, mbuf->mem, 1, 0, sz, dest, 0, 0, 0)) != 0) { - fprintf(stderr, "failed to read buffer (%d)\n", err); + fprintf(stderr, "failed to read buffer: %s\n", clstrerror(err)); return false; } return true; @@ -250,7 +255,7 @@ char *errlog = (char*)alloca(sz + 1); clGetProgramBuildInfo(prog, devinf.id, CL_PROGRAM_BUILD_LOG, sz, errlog, 0); - fprintf(stderr, "failed to build program: (%d)\n%s\n", err, errlog); + fprintf(stderr, "failed to build program: %s\n%s\n", clstrerror(err), errlog); clReleaseProgram(prog); prog = 0; @@ -275,14 +280,14 @@ switch(args[i].type) { case ARGTYPE_INT: if((err = clSetKernelArg(kernel, i, sizeof(int), &args[i].v.ival)) != 0) { - fprintf(stderr, "failed to bind kernel argument: %d (%d)\n", (int)i, err); + fprintf(stderr, "failed to bind kernel argument %d: %s\n", (int)i, clstrerror(err)); goto fail; } break; case ARGTYPE_FLOAT: if((err = clSetKernelArg(kernel, i, sizeof(float), &args[i].v.fval)) != 0) { - fprintf(stderr, "failed to bind kernel argument: %d (%d)\n", (int)i, err); + fprintf(stderr, "failed to bind kernel argument %d: %s\n", (int)i, clstrerror(err)); goto fail; } break; @@ -292,7 +297,7 @@ CLMemBuffer *mbuf = args[i].v.mbuf; if((err = clSetKernelArg(kernel, i, sizeof mbuf->mem, &mbuf->mem)) != 0) { - fprintf(stderr, "failed to bind kernel argument: %d (%d)\n", (int)i, err); + fprintf(stderr, "failed to bind kernel argument %d: %s\n", (int)i, clstrerror(err)); goto fail; } } @@ -338,7 +343,7 @@ int err; if((err = clEnqueueNDRangeKernel(cmdq, kernel, dim, 0, global_size, 0, 0, 0, 0)) != 0) { - fprintf(stderr, "error executing kernel (%d)\n", err); + fprintf(stderr, "error executing kernel: %s\n", clstrerror(err)); return false; } return true; @@ -346,13 +351,36 @@ static int select_device(struct device_info *dev_inf, int (*devcmp)(struct device_info*, struct device_info*)) { - unsigned int i, j, num_dev, sel; + unsigned int i, j, num_dev, num_plat, sel, ret; cl_device_id dev[32]; + cl_platform_id plat[32]; dev_inf->work_item_sizes = 0; + if((ret = clGetPlatformIDs(32, plat, &num_plat)) != 0) { + fprintf(stderr, "clGetPlatformIDs failed: %s\n", clstrerror(ret)); + return -1; + } + if(!num_plat) { + fprintf(stderr, "OpenCL not available!\n"); + return -1; + } - clGetDeviceIDs(0, CL_DEVICE_TYPE_ALL, 32, dev, &num_dev); + for(i=0; i 0) { + return ""; + } + if(err <= -(int)(sizeof ocl_errstr / sizeof *ocl_errstr)) { + return ""; + } + return ocl_errstr[-err]; +} diff -r 575383f3a239 -r deaf85acf6af src/ocl.h --- a/src/ocl.h Fri Jul 23 01:22:03 2010 +0100 +++ b/src/ocl.h Fri Jul 23 19:48:43 2010 +0100 @@ -4,7 +4,7 @@ #include #include #ifndef __APPLE__ -#include +#include #else #include #endif diff -r 575383f3a239 -r deaf85acf6af src/ocl_errstr.h --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/src/ocl_errstr.h Fri Jul 23 19:48:43 2010 +0100 @@ -0,0 +1,64 @@ +#ifndef OCL_CC_ +#error "don't include ocl_errstr.h" +#endif + +static const char *ocl_errstr[] = { + "CL_SUCCESS", + "CL_DEVICE_NOT_FOUND", + "CL_DEVICE_NOT_AVAILABLE", + "CL_COMPILER_NOT_AVAILABLE", + "CL_MEM_OBJECT_ALLOCATION_FAILURE", + "CL_OUT_OF_RESOURCES", + "CL_OUT_OF_HOST_MEMORY", + "CL_PROFILING_INFO_NOT_AVAILABLE", + "CL_MEM_COPY_OVERLAP", + "CL_IMAGE_FORMAT_MISMATCH", + "CL_IMAGE_FORMAT_NOT_SUPPORTED", + "CL_BUILD_PROGRAM_FAILURE", + "CL_MAP_FAILURE", /* 12 */ + + "", "", /* 13, 14 */ + "", "", /* 15, 16 */ + "", "", /* 17, 18 */ + "", "", /* 19, 20 */ + "", "", /* 21, 22 */ + "", "", /* 23, 24 */ + "", "", /* 25, 26 */ + "", "", /* 27, 28 */ + "", /* 29 */ + + "CL_INVALID_VALUE", /* 30 */ + "CL_INVALID_DEVICE_TYPE", + "CL_INVALID_PLATFORM", + "CL_INVALID_DEVICE", + "CL_INVALID_CONTEXT", + "CL_INVALID_QUEUE_PROPERTIES", + "CL_INVALID_COMMAND_QUEUE", + "CL_INVALID_HOST_PTR", + "CL_INVALID_MEM_OBJECT", + "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR", + "CL_INVALID_IMAGE_SIZE", + "CL_INVALID_SAMPLER", + "CL_INVALID_BINARY", + "CL_INVALID_BUILD_OPTIONS", + "CL_INVALID_PROGRAM", + "CL_INVALID_PROGRAM_EXECUTABLE", + "CL_INVALID_KERNEL_NAME", + "CL_INVALID_KERNEL_DEFINITION", + "CL_INVALID_KERNEL", + "CL_INVALID_ARG_INDEX", + "CL_INVALID_ARG_VALUE", + "CL_INVALID_ARG_SIZE", + "CL_INVALID_KERNEL_ARGS", + "CL_INVALID_WORK_DIMENSION", + "CL_INVALID_WORK_GROUP_SIZE", + "CL_INVALID_WORK_ITEM_SIZE", + "CL_INVALID_GLOBAL_OFFSET", + "CL_INVALID_EVENT_WAIT_LIST", + "CL_INVALID_EVENT", + "CL_INVALID_OPERATION", + "CL_INVALID_GL_OBJECT", + "CL_INVALID_BUFFER_SIZE", + "CL_INVALID_MIP_LEVEL", + "CL_INVALID_GLOBAL_WORK_SIZE" +}; diff -r 575383f3a239 -r deaf85acf6af src/rt.cc --- a/src/rt.cc Fri Jul 23 01:22:03 2010 +0100 +++ b/src/rt.cc Fri Jul 23 19:48:43 2010 +0100 @@ -1,4 +1,5 @@ #include +#include #include #include #include "ocl.h" @@ -10,25 +11,21 @@ } __attribute__((packed)); struct Sphere { - cl_float4 pos; - cl_float4 kd, ks; - cl_float radius; - cl_float spow; - cl_float kr, kt; + float pos[4]; + float kd[4], ks[4]; + float radius; + float spow; + float kr, kt; } __attribute__((packed)); struct Ray { - cl_float4 origin, dir; + float origin[4], dir[4]; } __attribute__((packed)); struct Light { - cl_float4 pos, color; + float pos[4], color[4]; } __attribute__((packed)); -struct Matrix4x4 { - cl_float m[16]; -}; - static Ray get_primary_ray(int x, int y, int w, int h, float vfov_deg); static Ray *prim_rays; @@ -36,16 +33,16 @@ static int global_size; static Sphere sphlist[] = { - {{0, 0, 8, 1}, {0.7, 0.2, 0.15, 1}, {1, 1, 1, 1}, 1.0, 60, 0, 0}, - {{-0.2, 0.4, 5, 1}, {0.2, 0.9, 0.3, 1}, {1, 1, 1, 1}, 0.25, 40, 0, 0} + {{0, 0, 0, 1}, {0.7, 0.2, 0.15, 1}, {1, 1, 1, 1}, 1.0, 60, 0, 0}, + {{-0.2, 0.4, -3, 1}, {0.2, 0.9, 0.3, 1}, {1, 1, 1, 1}, 0.25, 40, 0, 0} }; static Light lightlist[] = { {{-10, 10, -20, 1}, {1, 1, 1, 1}} }; -static Matrix4x4 xform = { - {1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1} +static float xform[16] = { + 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1 }; static RendInfo rinf; @@ -72,7 +69,7 @@ /* setup opencl */ prog = new CLProgram("render"); if(!prog->load("rt.cl")) { - return 1; + return false; } /* setup argument buffers */ @@ -108,6 +105,16 @@ return true; } +void set_xform(float *matrix) +{ + CLMemBuffer *mbuf = prog->get_arg_buffer(5); + assert(mbuf); + + assert(map_mem_buffer(mbuf, MAP_WR) == xform); + memcpy(xform, matrix, sizeof xform); + unmap_mem_buffer(mbuf); +} + static Ray get_primary_ray(int x, int y, int w, int h, float vfov_deg) { float vfov = M_PI * vfov_deg / 180.0; diff -r 575383f3a239 -r deaf85acf6af src/rt.h --- a/src/rt.h Fri Jul 23 01:22:03 2010 +0100 +++ b/src/rt.h Fri Jul 23 19:48:43 2010 +0100 @@ -4,5 +4,6 @@ bool init_renderer(int xsz, int ysz, float *fb); void destroy_renderer(); bool render(); +void set_xform(float *matrix); #endif /* RT_H_ */