include header in nvrtc

Hello,

I am currently trying to include a header file in NVRTC, but I can’t get it to work. It currently looks like this:

nvrtcCreateProgram( &program,     
    kernel_as_string.c_str(),
    "cuda_program",           
    1,                         
    (const char**)("/usr/lib/llvm-4.0/lib/clang/4.0.0/include/"),                   
    (const char**)("stddef.h")              
);

But after I am launching the Kernel the process gets finished with exit code 139:
Process finished with exit code 139 (interrupted by signal 11: SIGSEGV)

Does someone know what I am doing wrong?

A segfault can normally be isolated to a single line of host code.

nvrtcCreateProgram doesn’t “launch” a kernel.

However you have said

“but after I am launching the Kernel”…

so it’s not clear what is happening, or where the seg fault is actually occurring.

it may not have anything to do with your include header.

Which line of code is causing the seg fault?

You are completely right. I thought the segmentation fault appeared when I executed the kernel but indeed it is appearing at the line with nvrtcCreateProgram.

these lines won’t work:

(const char**)("/usr/lib/llvm-4.0/lib/clang/4.0.0/include/"),                   
    (const char**)("stddef.h")

a character string like that is a const char *, simply casting it to const char ** is going to cause trouble. The API is designed to accept an array of strings for each parameter. You need to actually provide a proper pointer to an array of strings. Here’s an example:

https://github.com/NVIDIA/jitify/blob/master/jitify.hpp#L1625

In fact, you may just want to study and use that jitify header library.

I’d also suggest reading this:

https://stackoverflow.com/questions/40087364/how-do-you-include-standard-cuda-libraries-to-link-with-nvrtc-code

to understand what you are really expected to provide for those parameters. Again, maybe just use jitify instead.

Thanks. Just casting char* to char** was the problem. I fixed segmentation fault now.

But when I am including a header file the following error occurs: “stddef.h(1): error: expected a declaration”. My code looks like this now:

// create program
nvrtcProgram program;
int num_headers = 1;
std::vector<const char*> header_names;
std::vector<const char*> header_sources;
header_names.push_back("stddef.h");
header_sources.push_back("/usr/lib/llvm-4.0/lib/clang/4.0.0/include/");

NVRTC_SAFE_CALL
(
    nvrtcCreateProgram( &program,                   // prog
        _kernel_as_string.c_str(),  // buffer
        "cuda_program",             // name
        num_headers,                          // numHeaders
        &header_sources[0],                    // headers
        &header_names[0]               // includeNames
    );
);

Is it correct that I am providing the path to the header file in header_sources and the name of the header file in header_names?

Using jitify is not a possibility in my case unfortunately.

my previous comment here was incorrect. Please refer to my comment later in this thread:

https://devtalk.nvidia.com/default/topic/1028233/cuda-programming-and-performance/include-header-in-nvrtc/post/5230929/#5230929

I understand that there is potentially more complexity to it and that I have to include other files, which are included in stddef.h as well. But that error message makes no sense to me, since the first line of the file is a comment.

Here are the first three lines of the file:

/*===---- stddef.h - Basic type definitions --------------------------------===
 *
 * Copyright (c) 2008 Eli Friedman

Thats why I thought that I may have done something wrong.

It’s possible you may have done something wrong. I can’t tell from just a couple lines of code, and no program to look at.

I would start by building a really simple test case, and use a really simple include file test, that doesn’t have any other dependencies, and get that working first. I’d be willing to bet stddef.h includes other include files, and that may be troublesome. But that doesn’t seem to be the issue at hand.

since multi-line comment is started in first line of stddef.h, the real cause of problem may be chars after the comment

I built a simple test case, but I am still getting an error message when I am including a header file.

My header file looks like this and is stored in /home/martin/:

#ifndef OCAL_TEST_HEADER_H
#define OCAL_TEST_HEADER_H

typedef float real;
#endif //OCAL_TEST_HEADER_H

My kernel looks like this:

#include "test_header.h"
extern "C" __global__ void simple (float* input, float* output){
    output[0] = input[0];
}

And this is how I create the program with nvrtc:

// create program
nvrtcProgram program;
int num_headers = 1;
std::vector<const char*> header_names;
std::vector<const char*> header_sources;
header_names.push_back("test_header.h");
header_sources.push_back("/home/martin/");

NVRTC_SAFE_CALL
(
    nvrtcCreateProgram( &program,     
        kernel_as_string.c_str(),  
        "cuda_program",             
        num_headers,                         
        &header_sources[0],                   
        &header_names[0]              
    );
);

But I still get the error message. This is the complete log of nvrtc:

test_header.h(1): warning: last line of file ends without a newline
test_header.h(1): error: expected a declaration
At end of source: warning: parsing restarts here after previous syntax error
1 error detected in the compilation of “cuda_program”.

When I am leaving out the line with the include everything works just fine.

I built a simple test case, but I am still getting an error message when I am including a header file.

My header file looks like this and is stored in /home/martin/:

#ifndef OCAL_TEST_HEADER_H
#define OCAL_TEST_HEADER_H

typedef float real;
#endif //OCAL_TEST_HEADER_H

My kernel looks like this:

#include "test_header.h"
extern "C" __global__ void simple (float* input, float* output){
    output[0] = input[0];
}

And this is how I create the program with nvrtc:

// create program
nvrtcProgram program;
int num_headers = 1;
std::vector<const char*> header_names;
std::vector<const char*> header_sources;
header_names.push_back("test_header.h");
header_sources.push_back("/home/martin/");

NVRTC_SAFE_CALL
(
    nvrtcCreateProgram( &program,     
        kernel_as_string.c_str(),  
        "cuda_program",             
        num_headers,                         
        &header_sources[0],                   
        &header_names[0]              
    );
);

But I still get the error message. This is the complete log of nvrtc:

test_header.h(1): warning: last line of file ends without a newline
test_header.h(1): error: expected a declaration
At end of source: warning: parsing restarts here after previous syntax error
1 error detected in the compilation of “cuda_program”.

When I am leaving out the line with the include everything works just fine.

well, add empty line at the end of test_header.h

If I add a new line at the end of the file, the same warning occurs.

But the warning is not the problem; it is only a warning. The error in the log is the problem (test_header.h(1): error: expected a declaration).

A description of how to use these parameters is contained here:

http://docs.nvidia.com/cuda/nvrtc/index.html#basic-usage

You are still using them incorrectly.

The nvrtc mechanism gives 2 methods by which an include file can be incorporated into your program:

  1. By passing the entire contents of the include file to nvrtcCreateProgram
  2. By passing just the path of the include file to nvrtcCompileProgram

The second method is probably easier. The first method is provided in case you cannot be certain that a particular include file will be available or findable at the place you expect in the filesystem, when the program is being compiled at runtime.

The first method is not really outlined anywhere (although it is offered via jitify). An example of the second method is available in the CUDA sample codes, in the file /usr/local/cuda/samples/common/inc/nvrtc_helper.h Study the function compileFileToPTX in that helper file, and note how the cooperative_groups.h header file is passed to the nvrtcCompileProgram function.

If you want to use the first method, don’t pass the path to the file in the header_sources, pass the complete contents of the include file in header_sources. This is the proximal reason for the error you are getting. The compiler is attempting to interpret /home/martin/ as the C/C++ code contents of your header file.