Module: Ignis::JIT::NVRTCBindings

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

Overview

NVRTC (NVIDIA Runtime Compilation) library FFI bindings Provides runtime compilation of CUDA C++ source code to PTX/CUBIN

Constant Summary collapse

NVRTC_SUCCESS =

NVRTC Result codes

0
NVRTC_ERROR_OUT_OF_MEMORY = 
1
NVRTC_ERROR_PROGRAM_CREATION_FAILURE = 
2
NVRTC_ERROR_INVALID_INPUT = 
3
NVRTC_ERROR_INVALID_PROGRAM = 
4
NVRTC_ERROR_INVALID_OPTION = 
5
NVRTC_ERROR_COMPILATION = 
6
NVRTC_ERROR_BUILTIN_OPERATION_FAILURE = 
7
NVRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION = 
8
NVRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION = 
9
NVRTC_ERROR_NAME_EXPRESSION_NOT_VALID = 
10
NVRTC_ERROR_INTERNAL_ERROR = 
11

Class Method Summary collapse

Class Method Details

.check_result!(result, context) ⇒ void

This method returns an undefined value.

Check NVRTC result and raise on error

Parameters:

  • result (Integer) —

    NVRTC result code

  • context (String) —

    Context for error message

Raises:



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

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

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

.compile_program(program, options: []) ⇒ void

This method returns an undefined value.

Compile an NVRTC program

Parameters:

  • program (FFI::Pointer) —

    Program handle

  • options (Array<String>) (defaults to: []) —

    Compilation options

Raises:

  • (NVRTCError) —

    If compilation fails (includes error log)



111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
 
# File 'lib/nvruby/jit/nvrtc_bindings.rb', line 111

def compile_program(program, options: [])
  ensure_loaded!

  if options.any?
    options_array = FFI::MemoryPointer.new(:pointer, options.size)
    options.each_with_index do |opt, i|
      options_array.put_pointer(i * FFI::Pointer.size, FFI::MemoryPointer.from_string(opt))
    end
  else
    options_array = nil
  end

  result = nvrtcCompileProgram(program, options.size, options_array)

  if result != NVRTC_SUCCESS
    log = get_program_log(program)
    raise NVRTCError.new(result, compilation_log: log)
  end

  nil
end

.create_program(source, name: "kernel.cu", headers: [], header_names: []) ⇒ FFI::Pointer

Create an NVRTC program from source code

Parameters:

  • source (String) —

    CUDA C++ source code

  • name (String) (defaults to: "kernel.cu") —

    Program name (for error messages)

  • headers (Array<String>) (defaults to: []) —

    Header contents (optional)

  • header_names (Array<String>) (defaults to: []) —

    Header names (optional)

Returns:

  • (FFI::Pointer) —

    Program handle

Raises:



75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
 
# File 'lib/nvruby/jit/nvrtc_bindings.rb', line 75

def create_program(source, name: "kernel.cu", headers: [], header_names: [])
  ensure_loaded!

  prog_ptr = FFI::MemoryPointer.new(:pointer)
  source_ptr = FFI::MemoryPointer.from_string(source)
  name_ptr = FFI::MemoryPointer.from_string(name)

  num_headers = headers.size

  if num_headers.positive?
    headers_ptr = FFI::MemoryPointer.new(:pointer, num_headers)
    header_names_ptr = FFI::MemoryPointer.new(:pointer, num_headers)

    headers.each_with_index do |header, i|
      headers_ptr.put_pointer(i * FFI::Pointer.size, FFI::MemoryPointer.from_string(header))
    end

    header_names.each_with_index do |header_name, i|
      header_names_ptr.put_pointer(i * FFI::Pointer.size, FFI::MemoryPointer.from_string(header_name))
    end
  else
    headers_ptr = nil
    header_names_ptr = nil
  end

  result = nvrtcCreateProgram(prog_ptr, source_ptr, name_ptr, num_headers, headers_ptr, header_names_ptr)
  check_result!(result, "nvrtcCreateProgram")

  prog_ptr.read_pointer
end

.destroy_program(program) ⇒ void

This method returns an undefined value.

Destroy an NVRTC program

Parameters:

  • program (FFI::Pointer) —

    Program handle



216
217
218
219
220
221
222
223
224
 
# File 'lib/nvruby/jit/nvrtc_bindings.rb', line 216

def destroy_program(program)
  return if program.nil? || program.null?

  ensure_loaded!

  prog_ptr = FFI::MemoryPointer.new(:pointer)
  prog_ptr.write_pointer(program)
  nvrtcDestroyProgram(prog_ptr)
end

.ensure_loaded! ⇒ void

This method returns an undefined value.

Ensure NVRTC library is loaded

Raises:



36
37
38
39
40
41
42
43
44
45
 
# File 'lib/nvruby/jit/nvrtc_bindings.rb', line 36

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

    Ignis::CUDA::LibraryLoader.load_library(:nvrtc)
    attach_nvrtc_functions!
    @loaded = true
    Ignis.logger.info("NVRTC bindings initialized")
  end
end

.get_cubin(program) ⇒ String

Get compiled CUBIN binary

Parameters:

  • program (FFI::Pointer) —

    Program handle

Returns:

  • (String) —

    CUBIN binary data

Raises:



171
172
173
174
175
176
177
178
179
180
181
 
# File 'lib/nvruby/jit/nvrtc_bindings.rb', line 171

def get_cubin(program)
  ensure_loaded!

  cubin_size = get_cubin_size(program)
  cubin_ptr = FFI::MemoryPointer.new(:char, cubin_size)

  result = nvrtcGetCUBIN(program, cubin_ptr)
  check_result!(result, "nvrtcGetCUBIN")

  cubin_ptr.read_bytes(cubin_size)
end

.get_cubin_size(program) ⇒ Integer

Get compiled CUBIN size

Parameters:

  • program (FFI::Pointer) —

    Program handle

Returns:

  • (Integer) —

    CUBIN size in bytes

Raises:



157
158
159
160
161
162
163
164
165
 
# File 'lib/nvruby/jit/nvrtc_bindings.rb', line 157

def get_cubin_size(program)
  ensure_loaded!

  size_ptr = FFI::MemoryPointer.new(:size_t)
  result = nvrtcGetCUBINSize(program, size_ptr)
  check_result!(result, "nvrtcGetCUBINSize")

  size_ptr.read(:size_t)
end

.get_program_log(program) ⇒ String

Get the program compilation log

Parameters:

  • program (FFI::Pointer) —

    Program handle

Returns:

  • (String) —

    Compilation log



136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
 
# File 'lib/nvruby/jit/nvrtc_bindings.rb', line 136

def get_program_log(program)
  ensure_loaded!

  size_ptr = FFI::MemoryPointer.new(:size_t)
  result = nvrtcGetProgramLogSize(program, size_ptr)
  return "" unless result == NVRTC_SUCCESS

  log_size = size_ptr.read(:size_t)
  return "" if log_size.zero?

  log_ptr = FFI::MemoryPointer.new(:char, log_size)
  result = nvrtcGetProgramLog(program, log_ptr)
  return "" unless result == NVRTC_SUCCESS

  log_ptr.read_string(log_size - 1)
end

.get_ptx(program) ⇒ String

Get compiled PTX code

Parameters:

  • program (FFI::Pointer) —

    Program handle

Returns:

  • (String) —

    PTX code

Raises:



201
202
203
204
205
206
207
208
209
210
211
 
# File 'lib/nvruby/jit/nvrtc_bindings.rb', line 201

def get_ptx(program)
  ensure_loaded!

  ptx_size = get_ptx_size(program)
  ptx_ptr = FFI::MemoryPointer.new(:char, ptx_size)

  result = nvrtcGetPTX(program, ptx_ptr)
  check_result!(result, "nvrtcGetPTX")

  ptx_ptr.read_string(ptx_size - 1)
end

.get_ptx_size(program) ⇒ Integer

Get compiled PTX size

Parameters:

  • program (FFI::Pointer) —

    Program handle

Returns:

  • (Integer) —

    PTX size in bytes

Raises:



187
188
189
190
191
192
193
194
195
 
# File 'lib/nvruby/jit/nvrtc_bindings.rb', line 187

def get_ptx_size(program)
  ensure_loaded!

  size_ptr = FFI::MemoryPointer.new(:size_t)
  result = nvrtcGetPTXSize(program, size_ptr)
  check_result!(result, "nvrtcGetPTXSize")

  size_ptr.read(:size_t)
end

.loaded? ⇒ Boolean

Check if NVRTC is loaded

Returns:

  • (Boolean) 


49
50
51
 
# File 'lib/nvruby/jit/nvrtc_bindings.rb', line 49

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

.version ⇒ String

Get NVRTC version

Returns:

  • (String) —

    Version string (e.g., “12.6”)

Raises:



56
57
58
59
60
61
62
63
64
65
66
 
# File 'lib/nvruby/jit/nvrtc_bindings.rb', line 56

def version
  ensure_loaded!

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

  result = nvrtcVersion(major_ptr, minor_ptr)
  check_result!(result, "nvrtcVersion")

  "#{major_ptr.read_int}.#{minor_ptr.read_int}"
end