Method: BOAST::CKernel#fill_module

Defined in:

lib/BOAST/CKernel.rb

#fill_module(module_file, module_name) ⇒ Object

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

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 "("
  params = []
  if(@lang == BOAST::FORTRAN) then
    @procedure.parameters.each { |param|
      if param.dimension then
        params.push( param.name )
      else
        params.push( "&"+param.name )
      end
    }
  else
    @procedure.parameters.each { |param|
      if param.dimension then
        params.push( param.name )
      elsif param.direction == :out or param.direction == :inout then
        params.push( "&"+param.name )
      else
        params.push( param.name )
      end
    }
  end
  if @lang == BOAST::CUDA then
    params.push( "block_number", "block_size" )
  end
  module_file.print params.join(", ")
  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