Module: Cumo::CUDA::NVRTC

Defined in:
ext/cumo/cuda/nvrtc.c

Class Method Summary collapse

Class Method Details

.nvrtcAddNameExpression(prog, name_expression) ⇒ Object

Names a template instantiation or other C++ entity whose mangled name the program is to report after it is compiled.

Parameters:

  • prog (Integer)
  • name_expression (String) —

    such as "kernel"

Raises:



259
260
261
262
263
264
265
266
267
268
 
# File 'ext/cumo/cuda/nvrtc.c', line 259

static VALUE
rb_nvrtcAddNameExpression(VALUE self, VALUE prog, VALUE name_expression)
{
    nvrtcProgram _prog;
    StringValueCStr(name_expression);
    _prog = (nvrtcProgram)cumo_cuda_handle_get(&programs, prog, "nvrtcProgram");
    check_status(nvrtcAddNameExpression(_prog, RSTRING_PTR(name_expression)));
    RB_GC_GUARD(name_expression);
    return Qnil;
}

.nvrtcCompileProgram(prog, options) ⇒ Object 



203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
 
# File 'ext/cumo/cuda/nvrtc.c', line 203

static VALUE
rb_nvrtcCompileProgram(VALUE self, VALUE prog, VALUE options)
{
    nvrtcResult status;
    nvrtcProgram _prog = (nvrtcProgram)cumo_cuda_handle_get(&programs, prog, "nvrtcProgram");
    int _numOptions = ary_len(options, "options");
    const char** _options = NULL;
    VALUE strings;

    strings = rb_ary_new_capa(_numOptions);
    push_strings(strings, options, _numOptions);

    if (_numOptions > 0) {
        _options = alloc_cstrs(_numOptions);
        fill_cstrs(_options, strings, 0, _numOptions);
    }

    {
        struct nvrtcCompileProgramParam param = {_prog, _numOptions, _options};
        status = (nvrtcResult)rb_thread_call_without_gvl(nvrtcCompileProgram_without_gvl_cb, &param, NULL, NULL);
    }

    free(_options);
    RB_GC_GUARD(strings);
    check_status(status);
    return Qnil;
}

.nvrtcCreateProgram(src, name, headers, includeNames) ⇒ Object 



112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
 
# File 'ext/cumo/cuda/nvrtc.c', line 112

static VALUE
rb_nvrtcCreateProgram(
        VALUE self,
        VALUE src,
        VALUE name,
        VALUE headers,
        VALUE includeNames)
{
    nvrtcResult status;
    nvrtcProgram _prog;
    const char* _src = StringValueCStr(src);
    const char* _name = StringValueCStr(name);
    int _numHeaders = ary_len(headers, "headers");
    const char** _headers = NULL;
    const char** _includeNames = NULL;
    VALUE strings;

    // nvrtcCreateProgram reads numHeaders entries out of includeNames as well,
    // so a shorter array would be read past its end.
    if (ary_len(includeNames, "include_names") != _numHeaders) {
        rb_raise(rb_eArgError,
                 "headers and include_names must have the same length (%d vs %ld)",
                 _numHeaders, RARRAY_LEN(includeNames));
    }

    strings = rb_ary_new_capa((long)_numHeaders * 2);
    push_strings(strings, headers, _numHeaders);
    push_strings(strings, includeNames, _numHeaders);

    if (_numHeaders > 0) {
        // Both vectors come out of one allocation, so there is one pointer to
        // free rather than two.
        _headers = alloc_cstrs((long)_numHeaders * 2);
        _includeNames = _headers + _numHeaders;
        fill_cstrs(_headers, strings, 0, _numHeaders);
        fill_cstrs(_includeNames, strings, _numHeaders, _numHeaders);
    }

    {
        struct nvrtcCreateProgramParam param = {&_prog, _src, _name, _numHeaders, _headers, _includeNames};
        status = (nvrtcResult)rb_thread_call_without_gvl(nvrtcCreateProgram_without_gvl_cb, &param, NULL, NULL);
    }

    free(_headers);
    RB_GC_GUARD(strings);
    check_status(status);
    cumo_cuda_handle_set_add(&programs, (size_t)_prog);
    return SIZET2NUM((size_t)_prog);
}

.nvrtcDestroyProgram(prog) ⇒ Object 



175
176
177
178
179
180
181
182
183
184
185
186
 
# File 'ext/cumo/cuda/nvrtc.c', line 175

static VALUE
rb_nvrtcDestroyProgram(VALUE self, VALUE prog)
{
    nvrtcResult status;
    nvrtcProgram _prog = (nvrtcProgram)cumo_cuda_handle_take(&programs, prog, "nvrtcProgram");

    struct nvrtcDestroyProgramParam param = {&_prog};
    status = (nvrtcResult)rb_thread_call_without_gvl(nvrtcDestroyProgram_without_gvl_cb, &param, NULL, NULL);

    check_status(status);
    return Qnil;
}

.nvrtcGetLoweredName(prog, name_expression) ⇒ String

Returns the mangled name of a name expression given before the program was compiled, which is the name cuModuleGetFunction takes.

Parameters:

  • prog (Integer)
  • name_expression (String)

Returns:

  • (String)

Raises:



279
280
281
282
283
284
285
286
287
288
289
 
# File 'ext/cumo/cuda/nvrtc.c', line 279

static VALUE
rb_nvrtcGetLoweredName(VALUE self, VALUE prog, VALUE name_expression)
{
    nvrtcProgram _prog;
    const char *lowered = NULL;
    StringValueCStr(name_expression);
    _prog = (nvrtcProgram)cumo_cuda_handle_get(&programs, prog, "nvrtcProgram");
    check_status(nvrtcGetLoweredName(_prog, RSTRING_PTR(name_expression), &lowered));
    RB_GC_GUARD(name_expression);
    return rb_str_new_cstr(lowered);
}

.nvrtcGetProgramLog(prog) ⇒ Object 



291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
 
# File 'ext/cumo/cuda/nvrtc.c', line 291

static VALUE
rb_nvrtcGetProgramLog(VALUE self, VALUE prog)
{
    nvrtcResult status;
    nvrtcProgram _prog = (nvrtcProgram)cumo_cuda_handle_get(&programs, prog, "nvrtcProgram");
    size_t _logSizeRet;
    char *_log;
    VALUE log;

    status = nvrtcGetProgramLogSize(_prog, &_logSizeRet);
    check_status(status);

    log = rb_str_new(NULL, _logSizeRet);
    _log = RSTRING_PTR(log);
    status = nvrtcGetProgramLog(_prog, _log);
    check_status(status);

    return log;
}

.nvrtcGetPTX(prog) ⇒ Object 



231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
 
# File 'ext/cumo/cuda/nvrtc.c', line 231

static VALUE
rb_nvrtcGetPTX(VALUE self, VALUE prog)
{
    nvrtcResult status;
    nvrtcProgram _prog = (nvrtcProgram)cumo_cuda_handle_get(&programs, prog, "nvrtcProgram");
    size_t _ptxSizeRet;
    char *_ptx;
    VALUE ptx;

    status = nvrtcGetPTXSize(_prog, &_ptxSizeRet);
    check_status(status);

    ptx = rb_str_new(NULL, _ptxSizeRet);
    _ptx = RSTRING_PTR(ptx);
    status = nvrtcGetPTX(_prog, _ptx);
    check_status(status);

    return ptx;
}

.nvrtcVersion ⇒ Object 



24
25
26
27
28
29
30
31
32
33
34
35
36
37
 
# File 'ext/cumo/cuda/nvrtc.c', line 24

static VALUE
rb_nvrtcVersion(VALUE self)
{
    int _major, _minor;
    nvrtcResult status;
    VALUE major, minor;

    status = nvrtcVersion(&_major, &_minor);

    check_status(status);
    major = INT2NUM(_major);
    minor = INT2NUM(_minor);
    return rb_ary_new3(2, major, minor);
}