Class: Barracuda::Program

Inherits:

Object

• Object
• Barracuda::Program

Defined in:

ext/barracuda.c

Instance Method Summary

• #compile(source) ⇒ Object
• #initialize(*args) ⇒ Object constructor
• #method_missing(*args) ⇒ Object

Constructor Details

#initialize(*args) ⇒ Object

# File 'ext/barracuda.c', line 452

static VALUE
program_initialize(int argc, VALUE *argv, VALUE self)
{
    VALUE source;
    
    rb_scan_args(argc, argv, "01", &source);
    if (source != Qnil) {
        program_compile(self, source);
    }
    
    return self;
}

Dynamic Method Handling

This class handles dynamic methods through the method_missing method

#method_missing(*args) ⇒ Object

# File 'ext/barracuda.c', line 508

static VALUE
program_method_missing(int argc, VALUE *argv, VALUE self)
{
    int i;
    size_t local = 0, global = 0;
    cl_kernel kernel;
    cl_command_queue commands;
    GET_PROGRAM();
    
    StringValue(argv[0]);
    kernel = clCreateKernel(program->program, RSTRING_PTR(argv[0]), &err);
    if (!kernel || err != CL_SUCCESS) {
        rb_raise(rb_eNoMethodError, "no kernel method '%s'", RSTRING_PTR(argv[0]));
    }
    
    commands = clCreateCommandQueue(context, device_id, 0, &err);
    if (!commands) {
        clReleaseKernel(kernel);
        rb_raise(rb_eOpenCLError, "could not execute kernel method '%s'", RSTRING_PTR(argv[0]));
    }

    for (i = 1; i < argc; i++) {
        VALUE item = argv[i];
        err = !CL_SUCCESS;
        
        if (i == argc - 1 && TYPE(item) == T_HASH) {
            VALUE worker_size = rb_hash_aref(item, ID2SYM(id_times));
            if (RTEST(worker_size) && TYPE(worker_size) == T_FIXNUM) {
                global = FIX2UINT(worker_size);
            }
            else {
                CLEAN();
                rb_raise(rb_eArgError, "opts hash must be {:times => INT_VALUE}, got %s",
                    RSTRING_PTR(rb_inspect(item)));
            }
            break;
        }
        
        if (TYPE(item) == T_ARRAY) {
            /* create buffer from arg */
            VALUE buf = buffer_s_allocate(rb_cBuffer);
            item = buffer_initialize(1, &item, buf);
        }

        if (CLASS_OF(item) == rb_cOutputBuffer) {
            struct buffer *buffer;
            Data_Get_Struct(item, struct buffer, buffer);
            err = clSetKernelArg(kernel, i - 1, sizeof(cl_mem), &buffer->data);
            if (buffer->num_items > global) {
                global = buffer->num_items;
            }
        }
        else if (CLASS_OF(item) == rb_cBuffer) {
            struct buffer *buffer;
            Data_Get_Struct(item, struct buffer, buffer);

            buffer_write(item);
            clEnqueueWriteBuffer(commands, buffer->data, CL_TRUE, 0, 
                buffer->num_items * buffer->member_size, buffer->cachebuf, 0, NULL, NULL);
            err = clSetKernelArg(kernel, i - 1, sizeof(cl_mem), &buffer->data);
            if (buffer->num_items > global) {
                global = buffer->num_items;
            }
        }
        else {
            unsigned long data_ptr[16]; // a buffer of data
            size_t data_size_t;
            VALUE data_type, data_size;
            
            if (CLASS_OF(item) == rb_cType) {
                data_type = rb_funcall(type_object(item), id_data_type, 0);
            }
            else {
                data_type = rb_funcall(item, id_data_type, 0);
            }
            data_size = rb_hash_aref(rb_hTypes, data_type);
            if (NIL_P(data_size)) {
                CLEAN();
                rb_raise(rb_eRuntimeError, "invalid data type for %s", 
                    RSTRING_PTR(rb_inspect(item)));
            }
            
            data_size_t = FIX2UINT(data_size);
            type_to_native(item, SYM2ID(data_type), (void *)data_ptr);
            err = clSetKernelArg(kernel, i - 1, data_size_t, data_ptr);
        }

        if (err != CL_SUCCESS) {
            CLEAN();
            rb_raise(rb_eArgError, "invalid kernel method parameter: %s", RSTRING_PTR(rb_inspect(item)));
        }
    }
    
    err = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &local, NULL);
    ERROR("failed to retrieve kernel work group info");
    
    { /* global work size must be power of 2, greater than 3 and not smaller than local */
        size_t size = 4;
        while (size < global) size *= 2;
        global = size;
        if (global < local) global = local;
    }

    clEnqueueNDRangeKernel(commands, kernel, 1, NULL, &global, &local, 0, NULL, NULL);
    if (err) { CLEAN(); rb_raise(rb_eOpenCLError, "failed to execute kernel method"); }
    
    clFinish(commands);
    
    for (i = 1; i < argc; i++) {
        VALUE item = argv[i];
        if (CLASS_OF(item) == rb_cOutputBuffer) {
            struct buffer *buffer;
            Data_Get_Struct(item, struct buffer, buffer);
            err = clEnqueueReadBuffer(commands, buffer->data, CL_TRUE, 0, 
                buffer->num_items * buffer->member_size, buffer->cachebuf, 0, NULL, NULL);
            ERROR("failed to read output buffer");
            buffer_read(item);
        }
    }

    CLEAN();
    return Qnil;
}

Instance Method Details

#compile(source) ⇒ Object

# File 'ext/barracuda.c', line 465

static VALUE
program_compile(VALUE self, VALUE source)
{
    const char *c_source;
    GET_PROGRAM();
    StringValue(source);
    
    if (program->program) {
        clReleaseProgram(program->program);
        program->program = 0;
    }
    
    c_source = StringValueCStr(source);
    program->program = clCreateProgramWithSource(context, 1, &c_source, NULL, &err);
    if (!program->program) {
        program->program = 0;
        rb_raise(rb_eOpenCLError, "failed to create compute program");
    }

    err = clBuildProgram(program->program, 0, NULL, NULL, NULL, NULL);
    if (err != CL_SUCCESS) {
        size_t len;
        char buffer[2048];

        clGetProgramBuildInfo(program->program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
        clReleaseProgram(program->program);
        program->program = 0;
        rb_raise(rb_eProgramSyntaxError, "%s", buffer);
    }
    
    return Qtrue;
}