Class: BOAST::CKernel

Inherits:

Object

• Object
• BOAST::CKernel

Includes:

Rake::DSL

Defined in:

lib/BOAST/CKernel.rb

Instance Attribute Summary

• #binary ⇒ Object

  Returns the value of attribute binary.

• #code ⇒ Object

  Returns the value of attribute code.

• #kernels ⇒ Object

  Returns the value of attribute kernels.

• #lang ⇒ Object

  Returns the value of attribute lang.

• #procedure ⇒ Object

  Returns the value of attribute procedure.

Instance Method Summary

• #build(options = {}) ⇒ Object
• #build_opencl(options) ⇒ Object
• #fill_code(source_file) ⇒ Object
• #fill_module(module_file, module_name) ⇒ Object
• #initialize(options = {}) ⇒ CKernel constructor

  A new instance of CKernel.

• #method_missing(meth, *args, &block) ⇒ Object
• #print ⇒ Object
• #setup_compiler(options = {}) ⇒ Object
• #to_s ⇒ Object
• #to_str ⇒ Object

Constructor Details

#initialize(options = {}) ⇒ CKernel

Returns a new instance of CKernel.

# File 'lib/BOAST/CKernel.rb', line 94

def initialize(options={})
  if options[:code] then
    @code = options[:code]
  elsif BOAST::get_chain_code
    @code = BOAST::get_output
    @code.seek(0,SEEK_END)
  else
    @code = StringIO::new
  end
  BOAST::set_output( @code )
  if options[:kernels] then
    @kernels = options[:kernels]
  else
    @kernels  = []
  end
  if options[:lang] then
    @lang = options[:lang]
  else
    @lang = BOAST::get_lang
  end
end

Dynamic Method Handling

This class handles dynamic methods through the method_missing method

#method_missing(meth, *args, &block) ⇒ Object

# File 'lib/BOAST/CKernel.rb', line 645

def method_missing(meth, *args, &block)
 if meth.to_s == "run" then
   self.build
   self.run(*args,&block)
 else
   super
 end
end

Instance Attribute Details

#binary ⇒ Object

Returns the value of attribute binary.

# File 'lib/BOAST/CKernel.rb', line 91

def binary
  @binary
end

#code ⇒ Object

Returns the value of attribute code.

# File 'lib/BOAST/CKernel.rb', line 88

def code
  @code
end

#kernels ⇒ Object

Returns the value of attribute kernels.

# File 'lib/BOAST/CKernel.rb', line 92

def kernels
  @kernels
end

#lang ⇒ Object

Returns the value of attribute lang.

# File 'lib/BOAST/CKernel.rb', line 90

def lang
  @lang
end

#procedure ⇒ Object

Returns the value of attribute procedure.

# File 'lib/BOAST/CKernel.rb', line 89

def procedure
  @procedure
end

Instance Method Details

#build(options = {}) ⇒ Object

# File 'lib/BOAST/CKernel.rb', line 341

def build(options = {})
  compiler_options = BOAST::get_compiler_options
  compiler_options.update(options)
  return build_opencl(comiler_options) if @lang == BOAST::CL
  ldflags = self.setup_compiler(compiler_options)
  extension = ".c" if @lang == BOAST::C
  extension = ".cu" if @lang == BOAST::CUDA
  extension = ".f90" if @lang == BOAST::FORTRAN
#temporary
  c_compiler = compiler_options[:CC]
  c_compiler = "cc" if not c_compiler
  linker = compiler_options[:LD]
  linker = c_compiler if not linker
#end temporary
  if options[:openmp] then
    openmp_ld_flags = BOAST::get_openmp_flags[linker]
      if not openmp_ld_flags then
        keys = BOAST::get_openmp_flags.keys
        keys.each { |k|
          openmp_ld_flags = BOAST::get_openmp_flags[k] if linker.match(k)
        }
      end
      raise "unkwown openmp flags for: #{linker}" if not openmp_ld_flags
      ldflags += " #{openmp_ld_flags}"
  end
  source_file = Tempfile::new([@procedure.name,extension])
  path = source_file.path
  target = path.chomp(File::extname(path))+".o"
  fill_code(source_file)
  source_file.close

  previous_lang = BOAST::get_lang
  previous_output = BOAST::get_output
  BOAST::set_lang(BOAST::C)
  module_file_name = File::split(path.chomp(File::extname(path)))[0] + "/Mod_" + File::split(path.chomp(File::extname(path)))[1].gsub("-","_") + ".c"
  module_name = File::split(module_file_name.chomp(File::extname(module_file_name)))[1]
  module_file = File::open(module_file_name,"w+")
  BOAST::set_output(module_file)
  fill_module(module_file, module_name)
  module_file.rewind
#     puts module_file.read
  module_file.close
  BOAST::set_lang(previous_lang)
  BOAST::set_output(previous_output)
  module_target = module_file_name.chomp(File::extname(module_file_name))+".o"
  module_final = module_file_name.chomp(File::extname(module_file_name))+".so"
  kernel_files = []
  @kernels.each { |kernel|
    kernel_file = Tempfile::new([kernel.procedure.name,".o"])
    kernel.binary.rewind
    kernel_file.write( kernel.binary.read )
    kernel_file.close
    kernel_files.push(kernel_file)
  }
  file module_final => [module_target, target] do
    #puts "#{linker} -shared -o #{module_final} #{module_target} #{target} #{kernel_files.join(" ")} -Wl,-Bsymbolic-functions -Wl,-z,relro -rdynamic -Wl,-export-dynamic #{ldflags}"
    sh "#{linker} -shared -o #{module_final} #{module_target} #{target} #{(kernel_files.collect {|f| f.path}).join(" ")} -Wl,-Bsymbolic-functions -Wl,-z,relro -rdynamic -Wl,-export-dynamic #{ldflags}"
  end
  Rake::Task[module_final].invoke
  require(module_final)
  eval "self.extend(#{module_name})"
  f = File::open(target,"rb")
  @binary = StringIO::new
  @binary.write( f.read )
  f.close
  File.unlink(target)
  File.unlink(module_target)
  File.unlink(module_file_name)
  File.unlink(module_final)
  kernel_files.each { |f|
    f.unlink
  }
  return self
end

#build_opencl(options) ⇒ Object

# File 'lib/BOAST/CKernel.rb', line 244

def build_opencl(options)
  require 'opencl_ruby_ffi'
  platform = nil
  platforms = OpenCL::get_platforms
  if options[:platform_vendor] then
    platforms.each{ |p|
      platform = p if p.vendor.match(options[:platform_vendor])
    }
  else
    platform = platforms.first
  end
  device = nil
  type = options[:device_type] ? options[:device_type] : OpenCL::Device::Type::ALL
  devices = platform.devices(type)
  if options[:device_name] then
    devices.each{ |d|
      device = d if d.name.match(options[:device_name])
    }
  else
    device = devices.first
  end
  @context = OpenCL::create_context([device])
  program = @context.create_program_with_source([@code.string])
  opts = options[:CLFLAGS]
  program.build(:options => options[:CLFLAGS])
  if options[:verbose] then
    program.build_log.each {|dev,log|
      STDERR.puts "#{device.name}: #{log}"
    }
  end
  @queue = @context.create_command_queue(device, :properties => OpenCL::CommandQueue::PROFILING_ENABLE)
  @kernel = program.create_kernel(@procedure.name)
  run_method = "def self.run(*args)\n  raise \"Wrong number of arguments \\\#{args.length} for \#{@procedure.parameters.length}\" if args.length > \#{@procedure.parameters.length+1} or args.length < \#{@procedure.parameters.length}\n  params = []\n  opts = {}\n  opts = args.pop if args.length == \#{@procedure.parameters.length+1}\n  @procedure.parameters.each_index { |i|\nif @procedure.parameters[i].dimension then\n  if @procedure.parameters[i].direction == :in and @procedure.parameters[i].direction == :out then\n    params[i] = @context.create_buffer( args[i].size * args[i].element_size )\n    @queue.enqueue_write_buffer( params[i], args[i], :blocking => true )\n  elsif @procedure.parameters[i].direction == :in then\n    params[i] = @context.create_buffer( args[i].size * args[i].element_size, :flags => OpenCL::Mem::Flags::READ_ONLY )\n    @queue.enqueue_write_buffer( params[i], args[i], :blocking => true )\n  elsif @procedure.parameters[i].direction == :out then\n    params[i] = @context.create_buffer( args[i].size * args[i].element_size, :flags => OpenCL::Mem::Flags::WRITE_ONLY )\n  else\n    params[i] = @context.create_buffer( args[i].size * args[i].element_size )\n  end\nelse\n  if @procedure.parameters[i].type.is_a?(Real) then\n    params[i] = OpenCL::Half::new(args[i]) if @procedure.parameters[i].type.size == 2\n    params[i] = OpenCL::Float::new(args[i]) if @procedure.parameters[i].type.size == 4\n    params[i] = OpenCL::Double::new(args[i]) if @procedure.parameters[i].type.size == 8\n  elsif @procedure.parameters[i].type.is_a?(Int) then\n    if @procedure.parameters[i].type.signed\n      params[i] = OpenCL::Char::new(args[i]) if @procedure.parameters[i].type.size == 1\n      params[i] = OpenCL::Short::new(args[i]) if @procedure.parameters[i].type.size == 2\n      params[i] = OpenCL::Int::new(args[i]) if @procedure.parameters[i].type.size == 4\n      params[i] = OpenCL::Long::new(args[i]) if @procedure.parameters[i].type.size == 8\n    else\n      params[i] = OpenCL::UChar::new(args[i]) if @procedure.parameters[i].type.size == 1\n      params[i] = OpenCL::UShort::new(args[i]) if @procedure.parameters[i].type.size == 2\n      params[i] = OpenCL::UInt::new(args[i]) if @procedure.parameters[i].type.size == 4\n      params[i] = OpenCL::ULong::new(args[i]) if @procedure.parameters[i].type.size == 8\n    end\n  else\n    params[i] = args[i]\n  end\nend\n  }\n  params.each_index{ |i|\n@kernel.set_arg(i, params[i])\n  }\n  event = @queue.enqueue_NDrange_kernel(@kernel, opts[:global_work_size], :local_work_size => opts[:local_work_size])\n  @procedure.parameters.each_index { |i|\nif @procedure.parameters[i].dimension then\n  if @procedure.parameters[i].direction == :in and @procedure.parameters[i].direction == :out then\n    @queue.enqueue_read_buffer( params[i], args[i], :blocking => true )\n  elsif @procedure.parameters[i].direction == :out then\n    @queue.enqueue_read_buffer( params[i], args[i], :blocking => true )\n  end\nend\n  }\n  result = {}\n  result[:start] = event.profiling_command_start\n  result[:end] = event.profiling_command_end\n  result[:duration] = (result[:end] - result[:start])/1000000000.0\n  return result\nend\n"
eval run_method
return self
end

#fill_code(source_file) ⇒ Object

# File 'lib/BOAST/CKernel.rb', line 416

def fill_code(source_file)
  @code.rewind
  source_file.puts "#include <inttypes.h>" if @lang == BOAST::C or @lang == BOAST::CUDA
  source_file.puts "#include <cuda.h>" if @lang == BOAST::CUDA
  source_file.write @code.read
  if @lang == BOAST::CUDA then
    source_file.write "extern \"C\" {\n  \#{@procedure.header(BOAST::CUDA,false)}{\ndim3 dimBlock(block_size[0], block_size[1], block_size[2]);\ndim3 dimGrid(block_number[0], block_number[1], block_number[2]);\ncudaEvent_t __start, __stop;\nfloat __time;\ncudaEventCreate(&__start);\ncudaEventCreate(&__stop);\ncudaEventRecord(__start, 0);\n\#{@procedure.name}<<<dimGrid,dimBlock>>>(\#{@procedure.parameters.join(\", \")});\ncudaEventRecord(__stop, 0);\ncudaEventSynchronize(__stop);\ncudaEventElapsedTime(&__time, __start, __stop);\nreturn (unsigned long long int)((double)__time*(double)1e6);\n  }\n}\n"
  end
  @code.rewind
end

#fill_module(module_file, module_name) ⇒ Object

# File 'lib/BOAST/CKernel.rb', line 444

def fill_module(module_file, module_name)
  module_file.write "#include \"ruby.h\"\n#include <inttypes.h>\n#include <time.h>\n#ifdef HAVE_NARRAY_H\n#include \"narray.h\"\n#endif\n"
  if( @lang == BOAST::CUDA ) then
    module_file.print "#include <cuda_runtime.h>\n"
  end
  module_file.print @procedure.header(@lang)
  module_file.write "VALUE \#{module_name} = Qnil;\nvoid Init_\#{module_name}();\nVALUE method_run(int argc, VALUE *argv, VALUE self);\nvoid Init_\#{module_name}() {\n  \#{module_name} = rb_define_module(\"\#{module_name}\");\n  rb_define_method(\#{module_name}, \"run\", method_run, -1);\n}\nVALUE method_run(int argc, VALUE *argv, VALUE self) {\n"
  if( @lang == BOAST::CUDA ) then
    module_file.write "  if( argc < \#{@procedure.parameters.length} || argc > \#{@procedure.parameters.length + 1} )\nrb_raise(rb_eArgError, \"wrong number of arguments for \#{@procedure.name} (%d for \#{@procedure.parameters.length})\", argc);\n  VALUE rb_opts;\n  VALUE rb_ptr;\n  size_t block_size[3] = {1,1,1};\n  size_t block_number[3] = {1,1,1};\n"
  else
    module_file.write "  if( argc != \#{@procedure.parameters.length} )\nrb_raise(rb_eArgError, \"wrong number of arguments for \#{@procedure.name} (%d for \#{@procedure.parameters.length})\", argc);\n  VALUE rb_ptr;\n"
  end
  argc = @procedure.parameters.length
  argv = Variable::new("argv",Real,{:dimension => [ Dimension::new(0,argc-1) ] })
  rb_ptr = Variable::new("rb_ptr",Int)
  @procedure.parameters.each { |param| 
    param_copy = param.copy
    param_copy.constant = nil
    param_copy.direction = nil
    param_copy.decl
  }
  @procedure.parameters.each_index do |i|
    param = @procedure.parameters[i]
    if not param.dimension then
      case param.type
        when Int 
          (param === FuncCall::new("NUM2INT", argv[i])).print if param.type.size == 4
          (param === FuncCall::new("NUM2LONG", argv[i])).print if param.type.size == 8
        when Real
          (param === FuncCall::new("NUM2DBL", argv[i])).print
      end
    else
      (rb_ptr === argv[i]).print
      if @lang == BOAST::CUDA then
        module_file.print "  if ( IsNArray(rb_ptr) ) {\nstruct NARRAY *n_ary;\nsize_t array_size;\nData_Get_Struct(rb_ptr, struct NARRAY, n_ary);\narray_size = n_ary->total * na_sizeof[n_ary->type];\ncudaMalloc( (void **) &\#{param.name}, array_size);\n"
        if param.direction == :in then
        module_file.print "cudaMemcpy(\#{param.name}, (void *) n_ary->ptr, array_size, cudaMemcpyHostToDevice);\n"
        end
        module_file.print "  } else\nrb_raise(rb_eArgError, \"wrong type of argument %d\", \#{i});\n  \n"
      else
        module_file.print "  if (TYPE(rb_ptr) == T_STRING) {\n\#{param.name} = (void *) RSTRING_PTR(rb_ptr);\n  } else if ( IsNArray(rb_ptr) ) {\nstruct NARRAY *n_ary;\nData_Get_Struct(rb_ptr, struct NARRAY, n_ary);\n\#{param.name} = (void *) n_ary->ptr;\n  } else\nrb_raise(rb_eArgError, \"wrong type of argument %d\", \#{i});\n"
      end
    end
  end
  if @lang == BOAST::CUDA then
    module_file.write "  if( argc == \#{@procedure.parameters.length + 1} ) {\nrb_opts = argv[argc -1];\nif ( rb_opts != Qnil ) {\n  VALUE rb_array_data = Qnil;\n  int i;\n  if (TYPE(rb_opts) != T_HASH)\n    rb_raise(rb_eArgError, \"Cuda options should be passed as a hash\");\n  rb_ptr = rb_hash_aref(rb_opts, ID2SYM(rb_intern(\"block_size\")));\n  if( rb_ptr != Qnil ) {\n    if (TYPE(rb_ptr) != T_ARRAY)\n      rb_raise(rb_eArgError, \"Cuda option block_size should be an array\");\n    for(i=0; i<3; i++) {\n      rb_array_data = rb_ary_entry(rb_ptr, i);\n      if( rb_array_data != Qnil )\n        block_size[i] = (size_t) NUM2LONG( rb_array_data );\n    }\n  }\n  rb_ptr = rb_hash_aref(rb_opts, ID2SYM(rb_intern(\"block_number\")));\n  if( rb_ptr != Qnil ) {\n    if (TYPE(rb_ptr) != T_ARRAY)\n      rb_raise(rb_eArgError, \"Cuda option block_number should be an array\");\n    for(i=0; i<3; i++) {\n      rb_array_data = rb_ary_entry(rb_ptr, i);\n      if( rb_array_data != Qnil )\n        block_number[i] = (size_t) NUM2LONG( rb_array_data );\n    }\n  }\n}\n  }\n"
  end
  module_file.print "  #{@procedure.properties[:return].type.decl} ret;\n" if @procedure.properties[:return]
  module_file.print "  VALUE stats = rb_hash_new();\n"
  module_file.print "  struct timespec start, stop;\n"
  module_file.print "  unsigned long long int duration;\n"
  module_file.print "  clock_gettime(CLOCK_REALTIME, &start);\n"
  if @lang == BOAST::CUDA then
    module_file.print "  duration = "
  elsif @procedure.properties[:return] then
    module_file.print "  ret = "
  end
  module_file.print "  #{@procedure.name}"
  module_file.print "_" if @lang == BOAST::FORTRAN
  module_file.print "_wrapper" if @lang == BOAST::CUDA
  module_file.print "("
  if(@lang == BOAST::FORTRAN) then
    params = []
    @procedure.parameters.each { |param|
      if param.dimension then
        params.push( param.name )
      else
        params.push( "&"+param.name )
      end
    }
    module_file.print params.join(", ")
  else
    module_file.print @procedure.parameters.join(", ") 
  end
  if @lang == BOAST::CUDA then
    module_file.print ", " if @procedure.parameters.length > 0
    module_file.print "block_number, block_size"
  end
  module_file.print "  );\n"
  module_file.print "  clock_gettime(CLOCK_REALTIME, &stop);\n"

  if @lang == BOAST::CUDA then
    @procedure.parameters.each_index do |i|
      param = @procedure.parameters[i]
      if param.dimension then
        (rb_ptr === argv[i]).print
        module_file.print "  if ( IsNArray(rb_ptr) ) {\n"
        if param.direction == :out then
        module_file.print "struct NARRAY *n_ary;\nsize_t array_size;\nData_Get_Struct(rb_ptr, struct NARRAY, n_ary);\narray_size = n_ary->total * na_sizeof[n_ary->type];\ncudaMemcpy(\#{param.name}, (void *) n_ary->ptr, array_size, cudaMemcpyDeviceToHost);\n"
        end
        module_file.print "cudaFree( (void *) \#{param.name});\n  } else\nrb_raise(rb_eArgError, \"wrong type of argument %d\", \#{i});\n  \n"
      end
    end
  end
  if @lang != BOAST::CUDA then
    module_file.print "  duration = (unsigned long long int)stop.tv_sec * (unsigned long long int)1000000000 + stop.tv_nsec;\n"
    module_file.print "  duration -= (unsigned long long int)start.tv_sec * (unsigned long long int)1000000000 + start.tv_nsec;\n"
  end
  module_file.print "  rb_hash_aset(stats,ID2SYM(rb_intern(\"duration\")),rb_float_new((double)duration*(double)1e-9));\n"
  if @procedure.properties[:return] then
    type_ret = @procedure.properties[:return].type
    module_file.print "  rb_hash_aset(stats,ID2SYM(rb_intern(\"return\")),rb_int_new((long long)ret));\n" if type_ret.kind_of?(Int) and type_ret.signed
    module_file.print "  rb_hash_aset(stats,ID2SYM(rb_intern(\"return\")),rb_int_new((unsigned long long)ret));\n" if type_ret.kind_of?(Int) and not type_ret.signed
    module_file.print "  rb_hash_aset(stats,ID2SYM(rb_intern(\"return\")),rb_float_new((double)ret));\n" if type_ret.kind_of?(Real)
  end
  module_file.print "  return stats;\n"
  module_file.print  "}"
end

#print ⇒ Object

# File 'lib/BOAST/CKernel.rb', line 116

def print
  @code.rewind
  puts @code.read
end

#setup_compiler(options = {}) ⇒ Object

# File 'lib/BOAST/CKernel.rb', line 131

def setup_compiler(options = {})
  Rake::Task::clear
  verbose = options[:verbose]
  verbose = BOAST::get_verbose if not verbose
  Rake::verbose(verbose)
  Rake::FileUtilsExt.verbose_flag=verbose
  f_compiler = options[:FC]
  c_compiler = options[:CC]
  cxx_compiler = options[:CXX]
  cuda_compiler = options[:NVCC]
  f_flags = options[:FCFLAGS]
  f_flags += " -fPIC"
  f_flags += " -fno-second-underscore" if f_compiler == 'g95'
  ld_flags = options[:LDFLAGS]
  cuda_flags = options[:NVCCFLAGS]
  cuda_flags += " --compiler-options '-fPIC'"


  includes = "-I#{RbConfig::CONFIG["archdir"]}"
  includes += " -I#{RbConfig::CONFIG["rubyhdrdir"]} -I#{RbConfig::CONFIG["rubyhdrdir"]}/#{RbConfig::CONFIG["arch"]}"
  includes += " -I#{RbConfig::CONFIG["rubyarchhdrdir"]}" if RbConfig::CONFIG["rubyarchhdrdir"]
  ld_flags += " -L#{RbConfig::CONFIG["libdir"]} #{RbConfig::CONFIG["LIBRUBYARG"]} -lrt"
  ld_flags += " -lcudart" if @lang == BOAST::CUDA
  narray_path = nil
  begin
    spec = Gem::Specification::find_by_name('narray')
    narray_path = spec.full_gem_path
  rescue Gem::LoadError => e
  rescue NoMethodError => e
    spec = Gem::available?('narray')
    if spec then
      require 'narray' 
      narray_path = Gem.loaded_specs['narray'].full_gem_path
    end
  end
  includes += " -I#{narray_path}" if narray_path
  cflags = options[:CFLAGS]
  cxxflags = options[:CXXFLAGS]
  cflags += " -fPIC #{includes}"
  cxxflags += " -fPIC #{includes}"
  cflags += " -DHAVE_NARRAY_H" if narray_path
  fcflags = f_flags
  cudaflags = cuda_flags

  if options[:openmp] then
    case @lang
    when BOAST::C
      openmp_c_flags = BOAST::get_openmp_flags[c_compiler]
      if not openmp_c_flags then
        keys = BOAST::get_openmp_flags.keys
        keys.each { |k|
          openmp_c_flags = BOAST::get_openmp_flags[k] if c_compiler.match(k)
        }
      end
      raise "unkwown openmp flags for: #{c_compiler}" if not openmp_c_flags
      cflags += " #{openmp_c_flags}"
      openmp_cxx_flags = BOAST::get_openmp_flags[cxx_compiler]
      if not openmp_cxx_flags then
        keys = BOAST::get_openmp_flags.keys
        keys.each { |k|
          openmp_cxx_flags = BOAST::get_openmp_flags[k] if cxx_compiler.match(k)
        }
      end
      raise "unkwown openmp flags for: #{cxx_compiler}" if not openmp_cxx_flags
      cxxflags += " #{openmp_cxx_flags}"
    when BOAST::FORTRAN
      openmp_f_flags = BOAST::get_openmp_flags[f_compiler]
      if not openmp_f_flags then
        keys = BOAST::get_openmp_flags.keys
        keys.each { |k|
          openmp_f_flags = BOAST::get_openmp_flags[k] if f_compiler.match(k)
        }
      end
      raise "unkwown openmp flags for: #{f_compiler}" if not openmp_f_flags
      fcflags += " #{openmp_f_flags}"
    end
  end

  runner = lambda { |t, call_string|
    if verbose then
      sh call_string
    else
      status, stdout, stderr = systemu call_string
      if not status.success? then
        puts stderr
        fail "#{t.source}: compilation failed"
      end
      status.success?
    end
  }

  rule '.o' => '.c' do |t|
    c_call_string = "#{c_compiler} #{cflags} -c -o #{t.name} #{t.source}"
    runner.call(t, c_call_string)
  end

  rule '.o' => '.f90' do |t|
    f_call_string = "#{f_compiler} #{fcflags} -c -o #{t.name} #{t.source}"
    runner.call(t, f_call_string)
  end

  rule '.o' => '.cpp' do |t|
    cxx_call_string = "#{cxx_compiler} #{cxxflags} -c -o #{t.name} #{t.source}"
    runner.call(t, cxx_call_string)
  end

  rule '.o' => '.cu' do |t|
    cuda_call_string = "#{cuda_compiler} #{cudaflags} -c -o #{t.name} #{t.source}"
    runner.call(t, cuda_call_string)
  end
  return ld_flags
end

#to_s ⇒ Object

# File 'lib/BOAST/CKernel.rb', line 126

def to_s
  @code.rewind
  return code.read
end

#to_str ⇒ Object

# File 'lib/BOAST/CKernel.rb', line 121

def to_str
  @code.rewind
  return code.read
end