Включение стандартных заголовков C в код CUDA NVRTC

Question

Я пишу ядро ​​CUDA, которое компилируется во время выполнения с использованием NVRTC (CUDA версия 9.2 с NVRTC версия 7.5), которому требуется заголовок stdint.h, чтобы иметь типы int32_t и т. Д.

Если я напишу исходный код ядра без include, он будет работать правильно.Например, ядро ​​

extern "C" __global__ void f() { ... }

Компилируется в код PTX, где f определяется как .visible .entry f.

Но если исходный код ядра

#include <stdint.h>
extern "C" __global__ void f() { ... }

, он сообщаетA function without execution space annotations (__host__/__device__/__global__) is considered a host function, and host functions are not allowed in JIT mode. (также без extern "C").

Передача -default-device делает код PTX .visible .func f, поэтому функция не может быть вызвана с хоста.

Есть ли способвключить заголовки в исходный код и все еще иметь функцию ввода __global__?Или, альтернативно, способ узнать, какое соглашение о целочисленном размере используется компилятором NVRTC, чтобы типы int32_t и т. Д. Можно было определять вручную?

Редактировать: Пример программыэто показывает проблему:

#include <cstdlib>
#include <string>
#include <vector>
#include <memory>
#include <cassert>
#include <iostream>

#include <cuda.h>
#include <cuda_runtime.h>
#include <nvrtc.h>

[[noreturn]] void fail(const std::string& msg, int code) {
    std::cerr << "error: " << msg << " (" << code << ')' << std::endl;
    std::exit(EXIT_FAILURE);
}


std::unique_ptr<char[]> compile_to_ptx(const char* program_source) {
    nvrtcResult rv;

    // create nvrtc program
    nvrtcProgram prog;
    rv = nvrtcCreateProgram(
        &prog,
        program_source,
        "program.cu",
        0,
        nullptr,
        nullptr
    );
    if(rv != NVRTC_SUCCESS) fail("nvrtcCreateProgram", rv);

    // compile nvrtc program
    std::vector<const char*> options = {
        "--gpu-architecture=compute_30"
    };
    //options.push_back("-default-device");
    rv = nvrtcCompileProgram(prog, options.size(), options.data());
    if(rv != NVRTC_SUCCESS) {
        std::size_t log_size;
        rv = nvrtcGetProgramLogSize(prog, &log_size);
        if(rv != NVRTC_SUCCESS) fail("nvrtcGetProgramLogSize", rv);

        auto log = std::make_unique<char[]>(log_size);
        rv = nvrtcGetProgramLog(prog, log.get());
        if(rv != NVRTC_SUCCESS) fail("nvrtcGetProgramLog", rv);
        assert(log[log_size - 1] == '\0');

        std::cerr << "Compile error; log:\n" << log.get() << std::endl;

        fail("nvrtcCompileProgram", rv);
    }

    // get ptx code
    std::size_t ptx_size;
    rv = nvrtcGetPTXSize(prog, &ptx_size);
    if(rv != NVRTC_SUCCESS) fail("nvrtcGetPTXSize", rv);

    auto ptx = std::make_unique<char[]>(ptx_size);
    rv = nvrtcGetPTX(prog, ptx.get());
    if(rv != NVRTC_SUCCESS) fail("nvrtcGetPTX", rv);
    assert(ptx[ptx_size - 1] == '\0');

    nvrtcDestroyProgram(&prog);

    return ptx;
}

const char program_source[] = R"%%%(
//#include <stdint.h>
extern "C" __global__ void f(int* in, int* out) {
    out[threadIdx.x] = in[threadIdx.x];
}
)%%%";

int main() {
    CUresult rv;

    // initialize CUDA
    rv = cuInit(0);
    if(rv != CUDA_SUCCESS) fail("cuInit", rv);

    // compile program to ptx
    auto ptx = compile_to_ptx(program_source);
    std::cout << "PTX code:\n" << ptx.get() << std::endl;
}

Когда //#include <stdint.h> в исходном коде ядра не комментируется, он больше не компилируется.Когда //options.push_back("-default-device"); не комментируется, он компилируется, но не помечает функцию f как .entry.

CMakeLists.txt для его компиляции (требуется API драйвера CUDA + NVRTC)

cmake_minimum_required(VERSION 3.4)
project(cudabug CXX)

find_package(CUDA REQUIRED)

set(CMAKE_CXX_STANDARD 14)
set(CMAKE_CXX_STANDARD_REQUIRED 14)

add_executable(cudabug cudabug.cc)
include_directories(SYSTEM ${CUDA_INCLUDE_DIRS})
link_directories(${CUDA_LIBRARY_DIRS})
target_link_libraries(cudabug PUBLIC ${CUDA_LIBRARIES} nvrtc cuda)

Answer 1

[Предисловие: это очень хакерский ответ, специфичный для цепочки инструментов GNU (хотя я подозреваю, что проблема в этом вопросе также специфична для цепочки инструментов GNU)].

Может показаться, чтопроблема здесь в стандартном заголовке GNU features.h, который вытягивается в stdint.h и который в итоге определяет множество функций-заглушек, которые имеют пространство компиляции по умолчанию __host__ и приводят к взрыву nvrtc.Также кажется, что опция -default-device приведет к разрешенному набору функций компилятора glibC, который приводит к сбою всего компилятора nvrtc.

Вы можете победить это (очень хакерским способом), предварительно определив набор функций длястандартная библиотека, которая исключает все функции хоста.Изменение кода ядра JIT на

const char program_source[] = R"%%%(
#define __ASSEMBLER__
#define __extension__
#include <stdint.h>
extern "C" __global__ void f(int32_t* in, int32_t* out) {
    out[threadIdx.x] = in[threadIdx.x];
}
)%%%";

дало мне следующее:

$ nvcc -std=c++14 -ccbin=g++-7 jit_header.cu -o jitheader -lnvrtc -lcuda
$ ./jitheader 
PTX code:
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-24330188
// Cuda compilation tools, release 9.2, V9.2.148
// Based on LLVM 3.4svn
//

.version 6.2
.target sm_30
.address_size 64

    // .globl   f

.visible .entry f(
    .param .u64 f_param_0,
    .param .u64 f_param_1
)
{
    .reg .b32   %r<3>;
    .reg .b64   %rd<8>;


    ld.param.u64    %rd1, [f_param_0];
    ld.param.u64    %rd2, [f_param_1];
    cvta.to.global.u64  %rd3, %rd2;
    cvta.to.global.u64  %rd4, %rd1;
    mov.u32     %r1, %tid.x;
    mul.wide.u32    %rd5, %r1, 4;
    add.s64     %rd6, %rd4, %rd5;
    ld.global.u32   %r2, [%rd6];
    add.s64     %rd7, %rd3, %rd5;
    st.global.u32   [%rd7], %r2;
    ret;
}

Большое предостережение: это сработало в системе glibC, на которой я его пробовал.Вероятно, он не будет работать с другими наборами инструментов или реализациями libC (если, действительно, у них есть эта проблема).