module BOAST::CUDARuntime

@private

Private Instance Methods

copy_array_param_from_ruby(par, param, ruby_param ) click to toggle source
# File lib/BOAST/Runtime/CUDARuntime.rb, line 53
    def copy_array_param_from_ruby(par, param, ruby_param )
      rb_ptr = Variable::new("_boast_rb_ptr", CustomType, :type_name => "VALUE")
      (rb_ptr === ruby_param).pr
      get_output.print <<EOF
  if ( IsNArray(_boast_rb_ptr) ) {
    struct NARRAY *_boast_n_ary;
    size_t _boast_array_size;
    Data_Get_Struct(_boast_rb_ptr, struct NARRAY, _boast_n_ary);
    _boast_array_size = _boast_n_ary->total * na_sizeof[_boast_n_ary->type];
    cudaMalloc( (void **) &#{par}, _boast_array_size);
    cudaMemcpy(#{par}, (void *) _boast_n_ary->ptr, _boast_array_size, cudaMemcpyHostToDevice);
  } else {
    rb_raise(rb_eArgError, "Wrong type of argument for %s, expecting array!", "#{param}");
  }
EOF
    end
copy_array_param_to_ruby(par, param, ruby_param) click to toggle source
# File lib/BOAST/Runtime/CUDARuntime.rb, line 145
    def copy_array_param_to_ruby(par, param, ruby_param)
      rb_ptr = Variable::new("_boast_rb_ptr", CustomType, :type_name => "VALUE")
      (rb_ptr === ruby_param).pr
      get_output.print <<EOF
  if ( IsNArray(_boast_rb_ptr) ) {
EOF
      if param.direction == :out or param.direction == :inout then
        get_output.print <<EOF
    struct NARRAY *_boast_n_ary;
    size_t _boast_array_size;
    Data_Get_Struct(_boast_rb_ptr, struct NARRAY, _boast_n_ary);
    _boast_array_size = _boast_n_ary->total * na_sizeof[_boast_n_ary->type];
    cudaMemcpy((void *) _boast_n_ary->ptr, #{par}, _boast_array_size, cudaMemcpyDeviceToHost);
EOF
      end
      get_output.print <<EOF
    cudaFree( (void *) #{par});
  } else {
    rb_raise(rb_eArgError, "Wrong type of argument for %s, expecting array!", "#{param}");
  }
EOF
    end
create_procedure_call() click to toggle source
# File lib/BOAST/Runtime/CUDARuntime.rb, line 138
def create_procedure_call
  get_output.print "  #{TimerProbe::RESULT} = "
  get_output.print " #{method_name}_wrapper( "
  get_output.print create_procedure_call_parameters.join(", ")
  get_output.puts " );"
end
create_procedure_call_parameters() click to toggle source
# File lib/BOAST/Runtime/CUDARuntime.rb, line 134
def create_procedure_call_parameters
  return create_procedure_call_parameters_old + ["_boast_block_number", "_boast_block_size", "_boast_params._boast_repeat"]
end
create_procedure_call_parameters_old()
create_wrapper() click to toggle source
# File lib/BOAST/Runtime/CUDARuntime.rb, line 131
def create_wrapper
end
Also aliased as: create_wrapper_old
create_wrapper_old()
Alias for: create_wrapper
fill_decl_module_params() click to toggle source
# File lib/BOAST/Runtime/CUDARuntime.rb, line 70
    def fill_decl_module_params
      fill_decl_module_params_old
      get_output.print <<EOF
  size_t _boast_block_size[3] = {1,1,1};
  size_t _boast_block_number[3] = {1,1,1};
  int64_t _boast_duration;
EOF
    end
Also aliased as: fill_decl_module_params_old
fill_decl_module_params_old()
fill_library_header() click to toggle source
# File lib/BOAST/Runtime/CUDARuntime.rb, line 23
def fill_library_header
  fill_library_header_old
  get_output.puts "#include <cuda.h>"
end
Also aliased as: fill_library_header_old
fill_library_header_old()
Alias for: fill_library_header
fill_library_source() click to toggle source
# File lib/BOAST/Runtime/CUDARuntime.rb, line 28
    def fill_library_source
      fill_library_source_old
      get_output.write <<EOF
extern "C" {
  #{@procedure.send(:boast_header_s,CUDA)}{
    int _boast_i;
    dim3 dimBlock(_boast_block_size[0], _boast_block_size[1], _boast_block_size[2]);
    dim3 dimGrid(_boast_block_number[0], _boast_block_number[1], _boast_block_number[2]);
    cudaEvent_t __start, __stop;
    float __time;
    cudaEventCreate(&__start);
    cudaEventCreate(&__stop);
    cudaEventRecord(__start, 0);
    for( _boast_i = 0; _boast_i < _boast_repeat; _boast_i ++) {
      #{@procedure.name}<<<dimGrid,dimBlock>>>(#{@procedure.parameters.join(", ")});
    }
    cudaEventRecord(__stop, 0);
    cudaEventSynchronize(__stop);
    cudaEventElapsedTime(&__time, __start, __stop);
    return (unsigned long long int)((double)__time*(double)1e6);
  }
}
EOF
    end
Also aliased as: fill_library_source_old
fill_library_source_old()
Alias for: fill_library_source
fill_module_header() click to toggle source
# File lib/BOAST/Runtime/CUDARuntime.rb, line 18
def fill_module_header
  fill_module_header_old
  get_output.puts "#include <cuda_runtime.h>"
end
Also aliased as: fill_module_header_old
fill_module_header_old()
Alias for: fill_module_header
get_params_value() click to toggle source
# File lib/BOAST/Runtime/CUDARuntime.rb, line 79
    def get_params_value
      get_params_value_old
      get_output.print <<EOF
  if( _boast_rb_opts != Qnil ) {
    VALUE _boast_rb_array_data = Qnil;
    int _boast_i;
    _boast_rb_ptr = rb_hash_aref(_boast_rb_opts, ID2SYM(rb_intern("block_size")));
    if( _boast_rb_ptr != Qnil ) {
      if (TYPE(_boast_rb_ptr) != T_ARRAY)
        rb_raise(rb_eArgError, "Cuda option block_size should be an array");
      for(_boast_i=0; _boast_i<3; _boast_i++) {
        _boast_rb_array_data = rb_ary_entry(_boast_rb_ptr, _boast_i);
        if( _boast_rb_array_data != Qnil )
          _boast_block_size[_boast_i] = (size_t) NUM2LONG( _boast_rb_array_data );
      }
    } else {
      _boast_rb_ptr = rb_hash_aref(_boast_rb_opts, ID2SYM(rb_intern("local_work_size")));
      if( _boast_rb_ptr != Qnil ) {
        if (TYPE(_boast_rb_ptr) != T_ARRAY)
          rb_raise(rb_eArgError, "Cuda option local_work_size should be an array");
        for(_boast_i=0; _boast_i<3; _boast_i++) {
          _boast_rb_array_data = rb_ary_entry(_boast_rb_ptr, _boast_i);
          if( _boast_rb_array_data != Qnil )
            _boast_block_size[_boast_i] = (size_t) NUM2LONG( _boast_rb_array_data );
        }
      }
    }
    _boast_rb_ptr = rb_hash_aref(_boast_rb_opts, ID2SYM(rb_intern("block_number")));
    if( _boast_rb_ptr != Qnil ) {
      if (TYPE(_boast_rb_ptr) != T_ARRAY)
        rb_raise(rb_eArgError, "Cuda option block_number should be an array");
      for(_boast_i=0; _boast_i<3; _boast_i++) {
        _boast_rb_array_data = rb_ary_entry(_boast_rb_ptr, _boast_i);
        if( _boast_rb_array_data != Qnil )
          _boast_block_number[_boast_i] = (size_t) NUM2LONG( _boast_rb_array_data );
      }
    } else {
      _boast_rb_ptr = rb_hash_aref(_boast_rb_opts, ID2SYM(rb_intern("global_work_size")));
      if( _boast_rb_ptr != Qnil ) {
        if (TYPE(_boast_rb_ptr) != T_ARRAY)
          rb_raise(rb_eArgError, "Cuda option global_work_size should be an array");
        for(_boast_i=0; _boast_i<3; _boast_i++) {
          _boast_rb_array_data = rb_ary_entry(_boast_rb_ptr, _boast_i);
          if( _boast_rb_array_data != Qnil )
            _boast_block_number[_boast_i] = (size_t) NUM2LONG( _boast_rb_array_data ) / _boast_block_size[_boast_i];
        }
      }
    }
  }
EOF
    end
Also aliased as: get_params_value_old
get_params_value_old()
Alias for: get_params_value
store_results() click to toggle source
# File lib/BOAST/Runtime/CUDARuntime.rb, line 168
def store_results
  store_results_old
  get_output.print "  rb_hash_aset(_boast_stats,ID2SYM(rb_intern(\"duration\")),rb_float_new((double)_boast_duration*(double)1e-9));\n"
end
Also aliased as: store_results_old
store_results_old()
Alias for: store_results