class Yadriggy::C::OclCodeGen
OpenCL-code generator
Constants
- HelperSource
Public Instance Methods
build_cmd()
click to toggle source
Calls superclass method
Yadriggy::C::CodeGen#build_cmd
# File lib/yadriggy/c/opencl.rb, line 174 def build_cmd super + Config::OpenCLoptions end
expand_functions(func_names, func_types)
click to toggle source
# File lib/yadriggy/c/opencl.rb, line 209 def expand_functions(func_names, func_types) voidFunc = MethodType.new([], Void) return func_names + ['ocl_init', 'ocl_finish'], func_types + [MethodType.new([Integer], Void), MethodType.new([], Void)] end
headers()
click to toggle source
Calls superclass method
Yadriggy::C::CodeGen#headers
# File lib/yadriggy/c/opencl.rb, line 178 def headers() super Config::OpenCLHeaders.each {|h| @printer << h << :nl } @printer << :nl end
preamble()
click to toggle source
Calls superclass method
Yadriggy::C::CodeGen#preamble
# File lib/yadriggy/c/opencl.rb, line 197 def preamble super @printer << 'int ocl_init(int);' << :nl @printer << 'void ocl_finish();' << :nl << :nl print_kernel_source @printer << HelperSource print_ocl_init print_ocl_finish print_callers end
variable_declarations()
click to toggle source
Calls superclass method
Yadriggy::C::CodeGen#variable_declarations
# File lib/yadriggy/c/opencl.rb, line 184 def variable_declarations() super @gvariables.each do |obj, name| if obj.is_a?(CType::OclArray) @printer << 'static cl_mem ' << name << ';' << :nl end end @typechecker.blocks.each do |obj, name_and_vars| @printer << 'static cl_kernel ' << name_and_vars[0] << ';' << :nl end @printer << :nl end
Private Instance Methods
c_type(type)
click to toggle source
Calls superclass method
Yadriggy::C::CodeGen#c_type
# File lib/yadriggy/c/opencl.rb, line 290 def c_type(type) if @printer.is_a?(KernelPrinter) && (type == RubyClass::Integer || type == Integer) return 'int' else super end end
print_callers()
click to toggle source
# File lib/yadriggy/c/opencl.rb, line 355 def print_callers @typechecker.blocks.each do |blk, name_and_vars| name = name_and_vars[0] @printer << 'static void ' << name << '_call(size_t p0' name_and_vars[1].each_with_index do |name_type, i| type = name_type[1] tname = if type <= RubyClass[CType::OclArray] 'cl_mem' else c_type(type) end @printer << ', ' << tname << " p#{i + 1}" end @printer << ') {' @printer.down @printer << 'size_t global;' << :nl @printer << 'int err = 0;' << :nl i = name_and_vars[1].size name_and_vars[2].each do |obj| i += 1 @printer << "cl_mem p#{i} = #{@gvariables[obj]};" << :nl end i.times do |j| @printer << "err |= clSetKernelArg(#{name}, #{j}, sizeof(p#{j + 1}), &p#{j + 1});" << :nl end @printer << 'ocl_err_check(err, "clSetKernelArg");' << :nl @printer << 'global = p0;' << :nl @printer << 'ocl_err_check(clEnqueueNDRangeKernel(commands, ' @printer << name << ', 1, NULL, &global, NULL, 0, NULL, NULL), "clEnqueueNDRangeKernel");' << :nl @printer << 'clFinish(commands);' << :nl @printer.up @printer << '}' << :nl end end
print_create_buffer_code()
click to toggle source
generate the code for creating buffers.
# File lib/yadriggy/c/opencl.rb, line 343 def print_create_buffer_code @gvariables.each do |obj, name| if obj.is_a?(CType::OclArray) @printer << " #{name} = clCreateBuffer(context, CL_MEM_READ_WRITE,"\ " sizeof(float) * #{obj.size}, NULL, NULL);" << :nl @printer << " if (!#{name}) {" << :nl @printer << ' fprintf(stderr, "error: clCreateBuffer\n");' @printer << ' return 1; }' << :nl end end end
print_create_kernel_code()
click to toggle source
generate the code for creating kernels.
# File lib/yadriggy/c/opencl.rb, line 330 def print_create_kernel_code @printer << ' int err;' << :nl @typechecker.blocks.each do |blk, name_vars| @printer << ' ' << name_vars[0] @printer << ' = clCreateKernel(program, "' << name_vars[0] << '", &err);' << :nl @printer << ' if (err != CL_SUCCESS) {' << :nl @printer << ' fprintf(stderr, "error: clCreateKernel\n");' << :nl @printer << ' return 1; }' << :nl end @printer << :nl end
print_kernel_source()
click to toggle source
generate kenrel source.
# File lib/yadriggy/c/opencl.rb, line 249 def print_kernel_source @printer << 'static const char* kernelSource = ' << :nl << '"' @printer = KernelPrinter.new(@printer) @typechecker.blocks.each do |blk, name_vars| func_name = name_vars[0] @printer << '__kernel void ' << func_name << '(' all_free_vars = name_vars[1].to_a all_free_vars += name_vars[2].map do |obj| [@gvariables[obj], obj.class] end first = true all_free_vars.each do |name, type| if first then first = false else @printer << ', ' end print_type_in_kernel(type) @printer << name end @printer << ') {' @printer.down local_var_declarations(blk) @printer << 'int ' << blk.params[0].name @printer << ' = get_global_id(0);' << :nl check(blk.body) @printer << ';' unless blk.body.is_a?(Exprs) @printer.up @printer << '}' << :nl end @printer = @printer.printer @printer << ' ";' << :nl << :nl end
print_ocl_finish()
click to toggle source
# File lib/yadriggy/c/opencl.rb, line 228 def print_ocl_finish @printer << 'void ocl_finish() {' << :nl @printer << ' if (!ocl_initialized) return;' << :nl @printer << ' ocl_initialized = 0;' << :nl @gvariables.each do |obj, name| if obj.is_a?(CType::OclArray) @printer << " clReleaseMemObject(#{name});" << :nl end end @typechecker.blocks.each do |obj, name_vars| @printer << ' clReleaseKernel(' << name_vars[0] << ');' << :nl end @printer << ' clReleaseProgram(program);' << :nl @printer << ' clReleaseCommandQueue(commands);' << :nl @printer << ' clReleaseContext(context);' << :nl @printer << '}' << :nl << :nl end
print_ocl_init()
click to toggle source
# File lib/yadriggy/c/opencl.rb, line 218 def print_ocl_init @printer << 'int ocl_init(int is_gpu) {' << :nl @printer << ' if (ocl_initialized) return 0;' << :nl @printer << ' ocl_initialized = 1;' << :nl @printer << ' if (ocl_init0(is_gpu)) return 1;' << :nl << :nl print_create_kernel_code print_create_buffer_code @printer << :nl << ' return 0;' << :nl << '}' << :nl << :nl end
print_type_in_kernel(type)
click to toggle source
@param [Type|Class] type
# File lib/yadriggy/c/opencl.rb, line 282 def print_type_in_kernel(type) if type == CType::OclArray || type == RubyClass[CType::OclArray] @printer << '__global float* ' else @printer << c_type(type) << ' ' end end