Class: SGC::CU::CUFunction

Inherits:
Object
  • Object
show all
Defined in:
lib/cuda/driver/function.rb

Instance Method Summary collapse

Instance Method Details

#attribute(attrib) ⇒ Integer

Returns The particular attribute attrib of this kernel function.

Examples:

Get function attribute.

func.attribute(:MAX_THREADS_PER_BLOCK)    #=> 512
func.attribute(:SHARED_SIZE_BYTES)        #=> 44
func.attribute(:NUM_REGS)                 #=> 3

Parameters:

Returns:

  • (Integer) —

    The particular attribute attrib of this kernel function.



206
207
208
209
210
211
 
# File 'lib/cuda/driver/function.rb', line 206

def attribute(attrib)
    p = FFI::MemoryPointer.new(:int)
    status = API::cuFuncGetAttribute(p, attrib, self.to_api)
    Pvt::handle_error(status, "Failed to query function attribute: attribute = #{attrib}.")
    p.read_int
end

#block_shape=(xdim) ⇒ Object #block_shape=(xdim, ydim) ⇒ Object #block_shape=(xdim, ydim, zdim) ⇒ Object

Deprecated.

Use #launch_kernel.

Set the block dimensions to use for next launch.

Parameters:

  • xdim (Integer) —

    The size of the x dimension.

  • ydim (Integer) —

    The size of the y dimension. Defaults to 1.

  • zdim (Integer) —

    The size of the z dimension. Defaults to 1.



138
139
140
141
142
143
144
 
# File 'lib/cuda/driver/function.rb', line 138

def block_shape=(*args)
    xdim, ydim, zdim = args.flatten
    ydim = 1 if ydim.nil?
    zdim = 1 if zdim.nil?
    status = API::cuFuncSetBlockShape(self.to_api, xdim, ydim, zdim)
    Pvt::handle_error(status, "Failed to set function block shape: (x,y,z) = (#{xdim},#{ydim},#{zdim}).")
end

#cache_config=(conf) ⇒ Object

Set the preferred cache configuration (CUFunctionCache) to use for next launch.

Parameters:



216
217
218
219
 
# File 'lib/cuda/driver/function.rb', line 216

def cache_config=(conf)
    status = API::cuFuncSetCacheConfig(self.to_api, conf)
    Pvt::handle_error(status, "Failed to set function cache config: config = #{conf}.")
end

#launch ⇒ CUFunction

Deprecated.

Use #launch_kernel.

Launch this kernel function with 1x1x1 grid of blocks to execute on the current CUDA device.

Returns:



161
162
163
164
165
 
# File 'lib/cuda/driver/function.rb', line 161

def launch
    status = API::cuLaunch(self.to_api)
    Pvt::handle_error(status, "Failed to launch kernel function on 1x1x1 grid of blocks.")
    self
end

#launch_grid(xdim) ⇒ CUFunction #launch_grid(xdim, ydim) ⇒ CUFunction

Deprecated.

Use #launch_kernel.

Launch this kernel function with grid dimensions (xdim, ydim) to execute on the current CUDA device.

Parameters:

  • xdim (Integer) —

    The x dimensional size of the grid to launch.

  • ydim (Integer) (defaults to: 1) —

    The y dimensional size of the grid to launch. Defaults to 1.

Returns:



176
177
178
179
180
 
# File 'lib/cuda/driver/function.rb', line 176

def launch_grid(xdim, ydim = 1)
    status = API::cuLaunchGrid(self.to_api, xdim, ydim)
    Pvt::handle_error(status, "Failed to launch kernel function on #{xdim}x#{ydim} grid of blocks.")
    self
end

#launch_grid_async(xdim, stream) ⇒ Object #launch_grid_async(xdim, ydim, stream) ⇒ Object

Deprecated.

Use #launch_kernel.

Launch this kernel function with grid dimensions (xdim, ydim) on stream asynchronously to execute on the current CUDA device. Setting stream to anything other than an instance of CUStream will execute on the default stream 0.

Parameters:

  • xdim (Integer) —

    The x dimensional size



191
192
193
194
195
196
 
# File 'lib/cuda/driver/function.rb', line 191

def launch_grid_async(xdim, ydim = 1, stream)
    s = Pvt::parse_stream(stream)
    status = API::cuLaunchGridAsync(self.to_api, xdim, ydim, s)
    Pvt::handle_error(status, "Failed to launch kernel function asynchronously on #{xdim}x#{ydim} grid of blocks.")
    self
end

#launch_kernel(grid_xdim, grid_ydim, grid_zdim, block_xdim, block_ydim, block_zdim, shared_mem_size, stream, params) ⇒ CUFunction

TODO:

Add support for other C data types for the kernel function parameters.

Launch this kernel function with full configuration parameters and function parameters to execute on the current CUDA device.

Parameters:

  • grid_xdim (Integer) —

    The x dimensional size of the grid to launch.

  • grid_ydim (Integer) —

    The y dimensional size of the grid to launch.

  • grid_zdim (Integer) —

    The z dimensional size of the grid to launch.

  • block_xdim (Integer) —

    The x dimensional size of a block in the grid.

  • block_ydim (Integer) —

    The y dimensional size of a block in the grid.

  • block_zdim (Integer) —

    The z dimensional size of a block in the grid.

  • shared_mem_size (Integer) —

    Number of bytes of dynamic shared memory for each thread block.

  • stream (Integer, CUStream) —

    The stream to launch this kernel function on. Setting stream to anything other than an instance of CUStream will execute on the default stream 0.

  • params (Array<Fixnum, Float, CUDevicePtr>) —

    The list of parameters to pass in for the kernel function launch.

    • A Fixnum is mapped to a C type int.

    • A Float is mapped to a C type float.

Returns:



238
239
240
241
242
243
244
245
246
247
 
# File 'lib/cuda/driver/function.rb', line 238

def launch_kernel(grid_xdim, grid_ydim, grid_zdim, block_xdim, block_ydim, block_zdim, shared_mem_size, stream, params)
    p = parse_params(params)
    s = Pvt::parse_stream(stream)
    status = API::cuLaunchKernel(self.to_api, grid_xdim, grid_ydim, grid_zdim, block_xdim, block_ydim, block_zdim, shared_mem_size, s, p, nil)
    Pvt::handle_error(status, "Failed to launch kernel function.\n" +
        "* #{grid_xdim} x #{grid_ydim} x #{grid_zdim} grid\n" +
        "* #{block_xdim} x #{block_ydim} x #{block_zdim} blocks\n" +
        "* shared memory size = #{shared_mem_size}")
    self
end

#param=(*args) ⇒ Object

Deprecated.

Use #launch_kernel.

Set the argument list of subsequent function call to arg1, arg2, *other_args.

Parameters:

  • *args —

    The list of arguments to pass to the kernel function.



38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
 
# File 'lib/cuda/driver/function.rb', line 38

def param=(*args)
    offset = 0
    args.flatten.each do |x|
        case x
            when Fixnum
                p = FFI::MemoryPointer.new(:int)
                p.write_int(x)
                size = 4
            when Float
                p = FFI::MemoryPointer.new(:float)
                p.write_float(x)
                size = 4
            when CUDevicePtr
                p = FFI::MemoryPointer.new(:CUDevicePtr)
                API::write_cudeviceptr(p, x.to_api)
                size = p.size
            else
                raise TypeError, "Invalid type of argument #{x.to_s}."
        end
        offset = align_up(offset, size)
        status = API::cuParamSetv(self.to_api, offset, p, size)
        Pvt::handle_error(status, "Failed to set function parameters: offset = #{offset}, value = #{x}.")
        offset += size
    end

    status = API::cuParamSetSize(self.to_api, offset)
    Pvt::handle_error(status, "Failed to set function parameter size: size = #{offset}.")
end

#param_set_size(nbytes) ⇒ CUFunction

Deprecated.

Use #launch_kernel.

Set the function parameter size to nbytes.

Parameters:

  • nbytes (Integer) —

    The parameter size in bytes.

Returns:



115
116
117
118
119
 
# File 'lib/cuda/driver/function.rb', line 115

def param_set_size(nbytes)
    status = API::cuParamSetSize(self.to_api, nbytes)
    Pvt::handle_error(status, "Failed to set function parameter size: size = #{nbytes}.")
    self
end

#param_set_texref(texunit, texref) ⇒ Object

Deprecated.

Raises:

  • (NotImplementedError) 


124
125
126
 
# File 'lib/cuda/driver/function.rb', line 124

def param_set_texref(texunit, texref)
    raise NotImplementedError
end

#param_setf(offset, value) ⇒ CUFunction

Deprecated.

Use #launch_kernel.

Set a float parameter to the function’s argument list at offset with value.

Parameters:

  • offset (Integer) —

    Number of bytes to offset.

  • value (Float) —

    The floating-point value to set as the function parameter.

Returns:



74
75
76
77
78
 
# File 'lib/cuda/driver/function.rb', line 74

def param_setf(offset, value)
    status = API::cuParamSetf(self.to_api, offset, value)
    Pvt::handle_error(status, "Failed to set function float parameter: offset = #{offset}, value = #{value}.")
    self
end

#param_seti(offset, value) ⇒ CUFunction

Deprecated.

Use #launch_kernel.

Set an integer parameter to the function’s argument list at offset with value.

Parameters:

  • offset (Integer) —

    Number of bytes to offset.

  • value (Integer) —

    The integer value to set as the function parameter.

Returns:



87
88
89
90
91
 
# File 'lib/cuda/driver/function.rb', line 87

def param_seti(offset, value)
    status = API::cuParamSeti(self.to_api, offset, value)
    Pvt::handle_error(status, "Failed to set function integer parameter: offset = #{offset}, value = #{value}")
    self
end

#param_setv(offset, ptr, nbytes) ⇒ CUFunction

Deprecated.

Use #launch_kernel.

Set an arbitrary data to the function’s argument list at offset with ptr pointed nbytes data.

Parameters:

  • offset (Integer) —

    Number of bytes to offset.

  • ptr (CUDevicePtr) —

    A device pointer pointing to an arbitrary data to be used as the function parameter.

  • nbytes (Integer) —

    The size of the arbitrary data in bytes.

Returns:



101
102
103
104
105
106
107
 
# File 'lib/cuda/driver/function.rb', line 101

def param_setv(offset, ptr, nbytes)
    p = FFI::MemoryPointer.new(:pointer)
    API::write_size_t(p, ptr.to_api.to_i) # Workaround broken p.write_pointer() on 64bit pointer.
    status = API::cuParamSetv(self.to_api, offset, p, nbytes)
    Pvt::handle_error(status, "Failed to set function arbitrary parameter: offset = #{offset}, size = #{nbytes}.")
    self
end

#shared_size=(nbytes) ⇒ Object

Deprecated.

Use #launch_kernel.

Set the dynamic shared-memory size to use for next launch.

Parameters:

  • nbytes (Integer) —

    Number of bytes.



151
152
153
154
 
# File 'lib/cuda/driver/function.rb', line 151

def shared_size=(nbytes)
    status = API::cuFuncSetSharedSize(self.to_api, nbytes)
    Pvt::handle_error(status, "Failed to set function shared memory size: #{nbytes}.")
end