Module: Ignis::JIT::DriverAPIBindings

Extended by:
FFI::Library
Defined in:
lib/nvruby/jit/driver_api_bindings.rb

Overview

CUDA Driver API FFI bindings for module and kernel management Provides low-level access to load compiled code and execute kernels

Constant Summary collapse

CUDA_SUCCESS =

CUDA Driver result codes

0
CUDA_ERROR_INVALID_VALUE = 
1
CUDA_ERROR_OUT_OF_MEMORY = 
2
CUDA_ERROR_NOT_INITIALIZED = 
3
CUDA_ERROR_DEINITIALIZED = 
4
CUDA_ERROR_PROFILER_DISABLED = 
5
CUDA_ERROR_NO_DEVICE = 
100
CUDA_ERROR_INVALID_DEVICE = 
101
CUDA_ERROR_INVALID_IMAGE = 
200
CUDA_ERROR_INVALID_CONTEXT = 
201
CUDA_ERROR_CONTEXT_ALREADY_CURRENT = 
202
CUDA_ERROR_MAP_FAILED = 
205
CUDA_ERROR_UNMAP_FAILED = 
206
CUDA_ERROR_ARRAY_IS_MAPPED = 
207
CUDA_ERROR_ALREADY_MAPPED = 
208
CUDA_ERROR_NO_BINARY_FOR_GPU = 
209
CUDA_ERROR_ALREADY_ACQUIRED = 
210
CUDA_ERROR_NOT_MAPPED = 
211
CUDA_ERROR_NOT_MAPPED_AS_ARRAY = 
212
CUDA_ERROR_NOT_MAPPED_AS_POINTER = 
213
CUDA_ERROR_ECC_UNCORRECTABLE = 
214
CUDA_ERROR_UNSUPPORTED_LIMIT = 
215
CUDA_ERROR_CONTEXT_ALREADY_IN_USE = 
216
CUDA_ERROR_PEER_ACCESS_UNSUPPORTED = 
217
CUDA_ERROR_INVALID_PTX = 
218
CUDA_ERROR_INVALID_GRAPHICS_CONTEXT = 
219
CUDA_ERROR_INVALID_SOURCE = 
300
CUDA_ERROR_FILE_NOT_FOUND = 
301
CUDA_ERROR_SHARED_OBJECT_SYMBOL_NOT_FOUND = 
302
CUDA_ERROR_SHARED_OBJECT_INIT_FAILED = 
303
CUDA_ERROR_OPERATING_SYSTEM = 
304
CUDA_ERROR_INVALID_HANDLE = 
400
CUDA_ERROR_ILLEGAL_STATE = 
401
CUDA_ERROR_NOT_FOUND = 
500
CUDA_ERROR_NOT_READY = 
600
CUDA_ERROR_ILLEGAL_ADDRESS = 
700
CUDA_ERROR_LAUNCH_OUT_OF_RESOURCES = 
701
CUDA_ERROR_LAUNCH_TIMEOUT = 
702
CUDA_ERROR_LAUNCH_INCOMPATIBLE_TEXTURING = 
703
CUDA_ERROR_PEER_ACCESS_ALREADY_ENABLED = 
704
CUDA_ERROR_PEER_ACCESS_NOT_ENABLED = 
705
CUDA_ERROR_PRIMARY_CONTEXT_ACTIVE = 
708
CUDA_ERROR_CONTEXT_IS_DESTROYED = 
709
CUDA_ERROR_ASSERT = 
710
CUDA_ERROR_TOO_MANY_PEERS = 
711
CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED = 
712
CUDA_ERROR_HOST_MEMORY_NOT_REGISTERED = 
713
CUDA_ERROR_HARDWARE_STACK_ERROR = 
714
CUDA_ERROR_ILLEGAL_INSTRUCTION = 
715
CUDA_ERROR_MISALIGNED_ADDRESS = 
716
CUDA_ERROR_INVALID_ADDRESS_SPACE = 
717
CUDA_ERROR_INVALID_PC = 
718
CUDA_ERROR_LAUNCH_FAILED = 
719
CUDA_ERROR_UNKNOWN = 
999
CU_LAUNCH_PARAM_END =

cuLaunchKernel extra parameter constants

0
CU_LAUNCH_PARAM_BUFFER_POINTER = 
1
CU_LAUNCH_PARAM_BUFFER_SIZE = 
2

Class Method Summary collapse

Class Method Details

.check_result!(result, context) ⇒ void

This method returns an undefined value.

Check CUDA Driver result and raise on error

Parameters:

  • result (Integer) —

    CUDA result code

  • context (String) —

    Context for error message

Raises:



243
244
245
246
247
 
# File 'lib/nvruby/jit/driver_api_bindings.rb', line 243

def check_result!(result, context)
  return if result == CUDA_SUCCESS

  raise CudaDriverError.new(result, context: context)
end

.context_synchronize ⇒ void

This method returns an undefined value.

Synchronize current context

Raises:



231
232
233
234
235
236
 
# File 'lib/nvruby/jit/driver_api_bindings.rb', line 231

def context_synchronize
  ensure_loaded!

  result = cuCtxSynchronize()
  check_result!(result, "cuCtxSynchronize")
end

.ensure_loaded! ⇒ void

This method returns an undefined value.

Ensure CUDA Driver API is loaded and initialized

Raises:



86
87
88
89
90
91
92
93
94
95
96
 
# File 'lib/nvruby/jit/driver_api_bindings.rb', line 86

def ensure_loaded!
  @mutex.synchronize do
    return if @loaded

    load_cuda_driver!
    attach_driver_functions!
    initialize_driver!
    @loaded = true
    Ignis.logger.info("CUDA Driver API bindings initialized")
  end
end

.get_current_context ⇒ FFI::Pointer

Get current CUDA context

Returns:

  • (FFI::Pointer) —

    Context handle (CUcontext)

Raises:



182
183
184
185
186
187
188
189
190
 
# File 'lib/nvruby/jit/driver_api_bindings.rb', line 182

def get_current_context
  ensure_loaded!

  ctx_ptr = FFI::MemoryPointer.new(:pointer)
  result = cuCtxGetCurrent(ctx_ptr)
  check_result!(result, "cuCtxGetCurrent")

  ctx_ptr.read_pointer
end

.get_device_compute_capability(device_id) ⇒ Array<Integer>

Get device compute capability

Parameters:

  • device_id (Integer) —

    Device index

Returns:

  • (Array<Integer>) —
    major, minor

    compute capability

Raises:



207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
 
# File 'lib/nvruby/jit/driver_api_bindings.rb', line 207

def get_device_compute_capability(device_id)
  ensure_loaded!

  device_ptr = FFI::MemoryPointer.new(:int)
  result = cuDeviceGet(device_ptr, device_id)
  check_result!(result, "cuDeviceGet")

  device = device_ptr.read_int

  major_ptr = FFI::MemoryPointer.new(:int)
  minor_ptr = FFI::MemoryPointer.new(:int)

  result = cuDeviceGetAttribute(major_ptr, 75, device) # CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR
  check_result!(result, "cuDeviceGetAttribute(major)")

  result = cuDeviceGetAttribute(minor_ptr, 76, device) # CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR
  check_result!(result, "cuDeviceGetAttribute(minor)")

  [major_ptr.read_int, minor_ptr.read_int]
end

.get_module_function(cuda_module, function_name) ⇒ FFI::Pointer

Get a kernel function from a loaded module

Parameters:

  • cuda_module (FFI::Pointer) —

    Module handle

  • function_name (String) —

    Kernel function name

Returns:

  • (FFI::Pointer) —

    Function handle (CUfunction)

Raises:



126
127
128
129
130
131
132
133
134
135
136
 
# File 'lib/nvruby/jit/driver_api_bindings.rb', line 126

def get_module_function(cuda_module, function_name)
  ensure_loaded!

  func_ptr = FFI::MemoryPointer.new(:pointer)
  name_ptr = FFI::MemoryPointer.from_string(function_name)

  result = cuModuleGetFunction(func_ptr, cuda_module, name_ptr)
  check_result!(result, "cuModuleGetFunction(#{function_name})")

  func_ptr.read_pointer
end

.launch_kernel(function, grid_dim:, block_dim:, kernel_params:, shared_mem: 0, stream: nil) ⇒ void

This method returns an undefined value.

Launch a kernel function

Parameters:

  • function (FFI::Pointer) —

    Function handle (CUfunction)

  • grid_dim (Array<Integer>) —

    Grid dimensions [x, y, z]

  • block_dim (Array<Integer>) —

    Block dimensions [x, y, z]

  • kernel_params (Array<FFI::Pointer>) —

    Kernel parameter pointers

  • shared_mem (Integer) (defaults to: 0) —

    Dynamic shared memory in bytes

  • stream (FFI::Pointer, nil) (defaults to: nil) —

    CUDA stream (nil for default)

Raises:



147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
 
# File 'lib/nvruby/jit/driver_api_bindings.rb', line 147

def launch_kernel(function, grid_dim:, block_dim:, kernel_params:, shared_mem: 0, stream: nil)
  ensure_loaded!

  grid_x, grid_y, grid_z = normalize_dims(grid_dim)
  block_x, block_y, block_z = normalize_dims(block_dim)

  params_ptr = build_params_array(kernel_params)
  stream_ptr = stream || FFI::Pointer::NULL

  result = cuLaunchKernel(
    function,
    grid_x, grid_y, grid_z,
    block_x, block_y, block_z,
    shared_mem,
    stream_ptr,
    params_ptr,
    nil
  )

  check_result!(result, "cuLaunchKernel")
end

.load_module_data(data) ⇒ FFI::Pointer

Load a module from CUBIN/PTX data

Parameters:

  • data (String) —

    CUBIN or PTX binary data

Returns:

  • (FFI::Pointer) —

    Module handle (CUmodule)

Raises:



108
109
110
111
112
113
114
115
116
117
118
119
 
# File 'lib/nvruby/jit/driver_api_bindings.rb', line 108

def load_module_data(data)
  ensure_loaded!

  module_ptr = FFI::MemoryPointer.new(:pointer)
  data_ptr = FFI::MemoryPointer.new(:char, data.bytesize)
  data_ptr.put_bytes(0, data)

  result = cuModuleLoadData(module_ptr, data_ptr)
  check_result!(result, "cuModuleLoadData")

  module_ptr.read_pointer
end

.loaded? ⇒ Boolean

Check if Driver API is loaded

Returns:

  • (Boolean) 


100
101
102
 
# File 'lib/nvruby/jit/driver_api_bindings.rb', line 100

def loaded?
  @mutex.synchronize { @loaded }
end

.set_current_context(context) ⇒ void

This method returns an undefined value.

Set current CUDA context

Parameters:

  • context (FFI::Pointer) —

    Context handle

Raises:



196
197
198
199
200
201
 
# File 'lib/nvruby/jit/driver_api_bindings.rb', line 196

def set_current_context(context)
  ensure_loaded!

  result = cuCtxSetCurrent(context)
  check_result!(result, "cuCtxSetCurrent")
end

.unload_module(cuda_module) ⇒ void

This method returns an undefined value.

Unload a module

Parameters:

  • cuda_module (FFI::Pointer) —

    Module handle



172
173
174
175
176
177
 
# File 'lib/nvruby/jit/driver_api_bindings.rb', line 172

def unload_module(cuda_module)
  return if cuda_module.nil? || cuda_module.null?

  ensure_loaded!
  cuModuleUnload(cuda_module)
end