Class: SGC::CU::CUFunction

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

Instance Method Summary (collapse)

Instance Method Details

- (Integer) attribute(attrib)

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


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

- (Object) block_shape=(xdim) - (Object) block_shape=(xdim, ydim) - (Object) block_shape=(xdim, ydim, zdim)

Deprecated.

Set the block dimensions to use for next launch.



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

- (Object) cache_config=(conf)

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



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

- (CUFunction) launch

Deprecated.

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



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

- (CUFunction) launch_grid(xdim) - (CUFunction) launch_grid(xdim, ydim)

Deprecated.

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



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

- (Object) launch_grid_async(xdim, stream) - (Object) launch_grid_async(xdim, ydim, stream)

Deprecated.

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.



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

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

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.



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

- (Object) param=(*args)

Deprecated.

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



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

- (CUFunction) param_set_size(nbytes)

Deprecated.

Set the function parameter size to nbytes.



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

- (Object) param_set_texref(texunit, texref)

Deprecated.

Raises:

  • (NotImplementedError)


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

def param_set_texref(texunit, texref)
    raise NotImplementedError
end

- (CUFunction) param_setf(offset, value)

Deprecated.

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



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

- (CUFunction) param_seti(offset, value)

Deprecated.

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



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

- (CUFunction) param_setv(offset, ptr, nbytes)

Deprecated.

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



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

- (Object) shared_size=(nbytes)

Deprecated.

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



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