Skip to content

Instantly share code, notes, and snippets.

@Robadob

Robadob/jitify.hpp Secret

Last active March 2, 2020 16:07
Show Gist options
  • Save Robadob/172c44e73b3bacd45ae33c3dec48a6db to your computer and use it in GitHub Desktop.
Save Robadob/172c44e73b3bacd45ae33c3dec48a6db to your computer and use it in GitHub Desktop.
Download ZIP
CUDA_ERROR_NOT_SUPPORTED from cuModuleLoadDataEx()
Raw
jitify.hpp
/*
* Copyright (c) 2017-2019, NVIDIA CORPORATION. All rights reserved.
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions
* are met:
* * Redistributions of source code must retain the above copyright
* notice, this list of conditions and the following disclaimer.
* * Redistributions in binary form must reproduce the above copyright
* notice, this list of conditions and the following disclaimer in the
* documentation and/or other materials provided with the distribution.
* * Neither the name of NVIDIA CORPORATION nor the names of its
* contributors may be used to endorse or promote products derived
* from this software without specific prior written permission.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY
* EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR
* PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR
* CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL,
* EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO,
* PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR
* PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY
* OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
*/
/*
-----------
Jitify 0.9
-----------
A C++ library for easy integration of CUDA runtime compilation into
existing codes.
--------------
How to compile
--------------
Compiler dependencies: <jitify.hpp>, -std=c++11
Linker dependencies: dl cuda nvrtc
--------------------------------------
Embedding source files into executable
--------------------------------------
g++ ... -ldl -rdynamic
-Wl,-b,binary,my_kernel.cu,include/my_header.cuh,-b,default nvcc ... -ldl
-Xcompiler "-rdynamic
-Wl\,-b\,binary\,my_kernel.cu\,include/my_header.cuh\,-b\,default"
JITIFY_INCLUDE_EMBEDDED_FILE(my_kernel_cu);
JITIFY_INCLUDE_EMBEDDED_FILE(include_my_header_cuh);
----
TODO
----
Extract valid compile options and pass the rest to cuModuleLoadDataEx
See if can have stringified headers automatically looked-up
by having stringify add them to a (static) global map.
The global map can be updated by creating a static class instance
whose constructor performs the registration.
Can then remove all headers from JitCache constructor in example code
See other TODOs in code
*/
/*! \file jitify.hpp
* \brief The Jitify library header
*/
/*! \mainpage Jitify - A C++ library that simplifies the use of NVRTC
* \p Use class jitify::JitCache to manage and launch JIT-compiled CUDA
* kernels.
*
* \p Use namespace jitify::reflection to reflect types and values into
* code-strings.
*
* \p Use JITIFY_INCLUDE_EMBEDDED_FILE() to declare files that have been
* embedded into the executable using the GCC linker.
*
* \p Use jitify::parallel_for and JITIFY_LAMBDA() to generate and launch
* simple kernels.
*/
#pragma once
#ifndef JITIFY_THREAD_SAFE
#define JITIFY_THREAD_SAFE 1
#endif
// WAR for MSVC not correctly defining __cplusplus (before MSVC 2017)
#ifdef _MSVC_LANG
#pragma push_macro("__cplusplus")
#undef __cplusplus
#define __cplusplus _MSVC_LANG
#endif
#if defined(_WIN32) || defined(_WIN64)
// WAR for strtok_r being called strtok_s on Windows
#pragma push_macro("strtok_r")
#undef strtok_r
#define strtok_r strtok_s
#endif
#include <dlfcn.h>
#include <stdint.h>
#include <algorithm>
#include <cstring> // For strtok_r etc.
#include <deque>
#include <fstream>
#include <iomanip>
#include <iostream>
#include <map>
#include <memory>
#include <sstream>
#include <stdexcept>
#include <string>
#include <typeinfo>
#include <unordered_map>
#include <unordered_set>
#include <vector>
#if JITIFY_THREAD_SAFE
#include <mutex>
#endif
#include <cuda.h>
#include <cuda_runtime_api.h> // For dim3, cudaStream_t
#if CUDA_VERSION >= 8000
#define NVRTC_GET_TYPE_NAME 1
#endif
#include <nvrtc.h>
#ifndef JITIFY_PRINT_LOG
#define JITIFY_PRINT_LOG 1
#endif
#if JITIFY_PRINT_ALL
#define JITIFY_PRINT_INSTANTIATION 1
#define JITIFY_PRINT_SOURCE 1
#define JITIFY_PRINT_LOG 1
#define JITIFY_PRINT_PTX 1
#define JITIFY_PRINT_LINKER_LOG 1
#define JITIFY_PRINT_LAUNCH 1
#define JITIFY_PRINT_HEADER_PATHS 1
#endif
#define JITIFY_FORCE_UNDEFINED_SYMBOL(x) void* x##_forced = (void*)&x
/*! Include a source file that has been embedded into the executable using the
* GCC linker.
* \param name The name of the source file (<b>not</b> as a string), which must
* be sanitized by replacing non-alpha-numeric characters with underscores.
* E.g., \code{.cpp}JITIFY_INCLUDE_EMBEDDED_FILE(my_header_h)\endcode will
* include the embedded file "my_header.h".
* \note Files declared with this macro can be referenced using
* their original (unsanitized) filenames when creating a \p
* jitify::Program instance.
*/
#define JITIFY_INCLUDE_EMBEDDED_FILE(name) \
extern "C" uint8_t _jitify_binary_##name##_start[] asm("_binary_" #name \
"_start"); \
extern "C" uint8_t _jitify_binary_##name##_end[] asm("_binary_" #name \
"_end"); \
JITIFY_FORCE_UNDEFINED_SYMBOL(_jitify_binary_##name##_start); \
JITIFY_FORCE_UNDEFINED_SYMBOL(_jitify_binary_##name##_end)
/*! Jitify library namespace
*/
namespace jitify {
/*! Source-file load callback.
*
* \param filename The name of the requested source file.
* \param tmp_stream A temporary stream that can be used to hold source code.
* \return A pointer to an input stream containing the source code, or NULL
* to defer loading of the file to Jitify's file-loading mechanisms.
*/
typedef std::istream* (*file_callback_type)(std::string filename,
std::iostream& tmp_stream);
// Exclude from Doxygen
//! \cond
class JitCache;
// Simple cache using LRU discard policy
template <typename KeyType, typename ValueType>
class ObjectCache {
public:
typedef KeyType key_type;
typedef ValueType value_type;
private:
typedef std::map<key_type, value_type> object_map;
typedef std::deque<key_type> key_rank;
typedef typename key_rank::iterator rank_iterator;
object_map _objects;
key_rank _ranked_keys;
size_t _capacity;
inline void discard_old(size_t n = 0) {
if (n > _capacity) {
throw std::runtime_error("Insufficient capacity in cache");
}
while (_objects.size() > _capacity - n) {
key_type discard_key = _ranked_keys.back();
_ranked_keys.pop_back();
_objects.erase(discard_key);
}
}
public:
inline ObjectCache(size_t capacity = 8) : _capacity(capacity) {}
inline void resize(size_t capacity) {
_capacity = capacity;
this->discard_old();
}
inline bool contains(const key_type& k) const { return _objects.count(k); }
inline void touch(const key_type& k) {
if (!this->contains(k)) {
throw std::runtime_error("Key not found in cache");
}
rank_iterator rank = std::find(_ranked_keys.begin(), _ranked_keys.end(), k);
if (rank != _ranked_keys.begin()) {
// Move key to front of ranks
_ranked_keys.erase(rank);
_ranked_keys.push_front(k);
}
}
inline value_type& get(const key_type& k) {
if (!this->contains(k)) {
throw std::runtime_error("Key not found in cache");
}
this->touch(k);
return _objects[k];
}
inline value_type& insert(const key_type& k,
const value_type& v = value_type()) {
this->discard_old(1);
_ranked_keys.push_front(k);
return _objects.insert(std::make_pair(k, v)).first->second;
}
template <typename... Args>
inline value_type& emplace(const key_type& k, Args&&... args) {
this->discard_old(1);
// Note: Use of piecewise_construct allows non-movable non-copyable types
auto iter = _objects
.emplace(std::piecewise_construct, std::forward_as_tuple(k),
std::forward_as_tuple(args...))
.first;
_ranked_keys.push_front(iter->first);
return iter->second;
}
};
namespace detail {
// Convenience wrapper for std::vector that provides handy constructors
template <typename T>
class vector : public std::vector<T> {
typedef std::vector<T> super_type;
public:
vector() : super_type() {}
vector(size_t n) : super_type(n) {} // Note: Not explicit, allows =0
vector(std::vector<T> const& vals) : super_type(vals) {}
template <int N>
vector(T const (&vals)[N]) : super_type(vals, vals + N) {}
#if defined __cplusplus && __cplusplus >= 201103L
vector(std::vector<T>&& vals) : super_type(vals) {}
vector(std::initializer_list<T> vals) : super_type(vals) {}
#endif
};
// Helper functions for parsing/manipulating source code
inline std::string replace_characters(std::string str,
std::string const& oldchars,
char newchar) {
size_t i = str.find_first_of(oldchars);
while (i != std::string::npos) {
str[i] = newchar;
i = str.find_first_of(oldchars, i + 1);
}
return str;
}
inline std::string sanitize_filename(std::string name) {
return replace_characters(name, "/\\.-: ?%*|\"<>", '_');
}
class EmbeddedData {
void* _app;
EmbeddedData(EmbeddedData const&);
EmbeddedData& operator=(EmbeddedData const&);
public:
EmbeddedData() {
_app = dlopen(NULL, RTLD_LAZY);
if (!_app) {
throw std::runtime_error(std::string("dlopen failed: ") + dlerror());
}
dlerror(); // Clear any existing error
}
~EmbeddedData() {
if (_app) {
dlclose(_app);
}
}
const uint8_t* operator[](std::string key) const {
key = sanitize_filename(key);
key = "_binary_" + key;
uint8_t const* data = (uint8_t const*)dlsym(_app, key.c_str());
if (!data) {
throw std::runtime_error(std::string("dlsym failed: ") + dlerror());
}
return data;
}
const uint8_t* begin(std::string key) const {
return (*this)[key + "_start"];
}
const uint8_t* end(std::string key) const { return (*this)[key + "_end"]; }
};
inline bool is_tokenchar(char c) {
return (c >= 'a' && c <= 'z') || (c >= 'A' && c <= 'Z') ||
(c >= '0' && c <= '9') || c == '_';
}
inline std::string replace_token(std::string src, std::string token,
std::string replacement) {
size_t i = src.find(token);
while (i != std::string::npos) {
if (i == 0 || i == src.size() - token.size() ||
(!is_tokenchar(src[i - 1]) && !is_tokenchar(src[i + token.size()]))) {
src.replace(i, token.size(), replacement);
i += replacement.size();
}
else {
i += token.size();
}
i = src.find(token, i);
}
return src;
}
inline std::string path_base(std::string p) {
// "/usr/local/myfile.dat" -> "/usr/local"
// "foo/bar" -> "foo"
// "foo/bar/" -> "foo/bar"
#if defined _WIN32 || defined _WIN64
char sep = '\\';
#else
char sep = '/';
#endif
size_t i = p.find_last_of(sep);
if (i != std::string::npos) {
return p.substr(0, i);
}
else {
return "";
}
}
inline std::string path_join(std::string p1, std::string p2) {
#ifdef _WIN32
char sep = '\\';
#else
char sep = '/';
#endif
if (p1.size() && p2.size() && p2[0] == sep) {
throw std::invalid_argument("Cannot join to absolute path");
}
if (p1.size() && p1[p1.size() - 1] != sep) {
p1 += sep;
}
return p1 + p2;
}
// Elides "/." and "/.." tokens from path.
inline std::string path_simplify(const std::string& path) {
std::vector<std::string> dirs;
std::string cur_dir;
bool after_slash = false;
for (int i = 0; i < (int)path.size(); ++i) {
if (path[i] == '/') {
if (after_slash) continue; // Ignore repeat slashes
after_slash = true;
if (cur_dir == ".." && !dirs.empty() && dirs.back() != "..") {
if (dirs.size() == 1 && dirs.front().empty()) {
throw std::runtime_error(
"Invalid path: back-traversals exceed depth of absolute path");
}
dirs.pop_back();
}
else if (cur_dir != ".") { // Ignore /./
dirs.push_back(cur_dir);
}
cur_dir.clear();
}
else {
after_slash = false;
cur_dir.push_back(path[i]);
}
}
if (!after_slash) {
dirs.push_back(cur_dir);
}
std::stringstream ss;
for (int i = 0; i < (int)dirs.size() - 1; ++i) {
ss << dirs[i] << "/";
}
if (!dirs.empty()) ss << dirs.back();
if (after_slash) ss << "/";
return ss.str();
}
inline unsigned long long hash_larson64(const char* s,
unsigned long long seed = 0) {
unsigned long long hash = seed;
while (*s) {
hash = hash * 101 + *s++;
}
return hash;
}
inline uint64_t hash_combine(uint64_t a, uint64_t b) {
// Note: The magic number comes from the golden ratio
return a ^ (0x9E3779B97F4A7C17ull + b + (b >> 2) + (a << 6));
}
inline bool extract_include_info_from_compile_error(std::string log,
std::string& name,
std::string& parent,
int& line_num) {
static const std::string pattern = "cannot open source file \"";
size_t beg = log.find(pattern);
if (beg == std::string::npos) {
return false;
}
beg += pattern.size();
size_t end = log.find("\"", beg);
name = log.substr(beg, end - beg);
size_t line_beg = log.rfind("\n", beg);
if (line_beg == std::string::npos) {
line_beg = 0;
}
else {
line_beg += 1;
}
size_t split = log.find("(", line_beg);
parent = log.substr(line_beg, split - line_beg);
line_num = atoi(
log.substr(split + 1, log.find(")", split + 1) - (split + 1)).c_str());
return true;
}
inline bool is_include_directive_with_quotes(const std::string& source,
int line_num) {
// TODO: Check each find() for failure.
size_t beg = 0;
for (int i = 1; i < line_num; ++i) {
beg = source.find("\n", beg) + 1;
}
beg = source.find("include", beg) + 7;
beg = source.find_first_of("\"<", beg);
return source[beg] == '"';
}
inline std::string comment_out_code_line(int line_num, std::string source) {
size_t beg = 0;
for (int i = 1; i < line_num; ++i) {
beg = source.find("\n", beg) + 1;
}
return (source.substr(0, beg) + "//" + source.substr(beg));
}
inline void print_with_line_numbers(std::string const& source) {
int linenum = 1;
std::stringstream source_ss(source);
for (std::string line; std::getline(source_ss, line); ++linenum) {
std::cout << std::setfill(' ') << std::setw(3) << linenum << " " << line
<< std::endl;
}
}
inline void print_compile_log(std::string program_name,
std::string const& log) {
std::cout << "---------------------------------------------------"
<< std::endl;
std::cout << "--- JIT compile log for " << program_name << " ---"
<< std::endl;
std::cout << "---------------------------------------------------"
<< std::endl;
std::cout << log << std::endl;
std::cout << "---------------------------------------------------"
<< std::endl;
}
inline std::vector<std::string> split_string(std::string str,
long maxsplit = -1,
std::string delims = " \t") {
std::vector<std::string> results;
if (maxsplit == 0) {
results.push_back(str);
return results;
}
// Note: +1 to include NULL-terminator
std::vector<char> v_str(str.c_str(), str.c_str() + (str.size() + 1));
char* c_str = v_str.data();
char* saveptr = c_str;
char* token = nullptr;
for (long i = 0; i != maxsplit; ++i) {
token = ::strtok_r(c_str, delims.c_str(), &saveptr);
c_str = 0;
if (!token) {
return results;
}
results.push_back(token);
}
// Check if there's a final piece
token += ::strlen(token) + 1;
if (token - v_str.data() < (ptrdiff_t)str.size()) {
// Find the start of the final piece
token += ::strspn(token, delims.c_str());
if (*token) {
results.push_back(token);
}
}
return results;
}
static const std::map<std::string, std::string>& get_jitsafe_headers_map();
inline bool load_source(
std::string filename, std::map<std::string, std::string>& sources,
std::string current_dir = "",
std::vector<std::string> include_paths = std::vector<std::string>(),
file_callback_type file_callback = 0,
std::map<std::string, std::string>* fullpaths = nullptr,
bool search_current_dir = true) {
std::istream* source_stream = 0;
std::stringstream string_stream;
std::ifstream file_stream;
// First detect direct source-code string ("my_program\nprogram_code...")
size_t newline_pos = filename.find("\n");
if (newline_pos != std::string::npos) {
std::string source = filename.substr(newline_pos + 1);
filename = filename.substr(0, newline_pos);
string_stream << source;
source_stream = &string_stream;
}
if (sources.count(filename)) {
// Already got this one
return true;
}
if (!source_stream) {
std::string fullpath = path_join(current_dir, filename);
// Try loading from callback
if (!file_callback ||
!(source_stream = file_callback(fullpath, string_stream))) {
// Try loading as embedded file
EmbeddedData embedded;
std::string source;
try {
source.assign(embedded.begin(fullpath), embedded.end(fullpath));
string_stream << source;
source_stream = &string_stream;
}
catch (std::runtime_error const&) {
// Try loading from filesystem
bool found_file = false;
if (search_current_dir) {
file_stream.open(fullpath.c_str());
if (file_stream) {
source_stream = &file_stream;
found_file = true;
}
}
// Search include directories
if (!found_file) {
for (int i = 0; i < (int)include_paths.size(); ++i) {
fullpath = path_join(include_paths[i], filename);
file_stream.open(fullpath.c_str());
if (file_stream) {
source_stream = &file_stream;
found_file = true;
break;
}
}
if (!found_file) {
// Try loading from builtin headers
fullpath = path_join("__jitify_builtin", filename);
auto it = get_jitsafe_headers_map().find(filename);
if (it != get_jitsafe_headers_map().end()) {
string_stream << it->second;
source_stream = &string_stream;
}
else {
return false;
}
}
}
}
}
if (fullpaths) {
// Record the full file path corresponding to this include name.
(*fullpaths)[filename] = path_simplify(fullpath);
}
}
sources[filename] = std::string();
std::string& source = sources[filename];
std::string line;
size_t linenum = 0;
unsigned long long hash = 0;
bool pragma_once = false;
bool remove_next_blank_line = false;
while (std::getline(*source_stream, line)) {
++linenum;
// HACK WAR for static variables not allowed on the device (unless
// __shared__)
// TODO: This breaks static member variables
// line = replace_token(line, "static const", "/*static*/ const");
// TODO: Need to watch out for /* */ comments too
std::string cleanline =
line.substr(0, line.find("//")); // Strip line comments
// if( cleanline.back() == "\r" ) { // Remove Windows line ending
// cleanline = cleanline.substr(0, cleanline.size()-1);
//}
// TODO: Should trim whitespace before checking .empty()
if (cleanline.empty() && remove_next_blank_line) {
remove_next_blank_line = false;
continue;
}
// Maintain a file hash for use in #pragma once WAR
hash = hash_larson64(line.c_str(), hash);
if (cleanline.find("#pragma once") != std::string::npos) {
pragma_once = true;
// Note: This is an attempt to recover the original line numbering,
// which otherwise gets off-by-one due to the include guard.
remove_next_blank_line = true;
// line = "//" + line; // Comment out the #pragma once line
continue;
}
// HACK WAR for Thrust using "#define FOO #pragma bar"
size_t pragma_beg = cleanline.find("#pragma ");
if (pragma_beg != std::string::npos) {
std::string line_after_pragma = line.substr(pragma_beg);
std::vector<std::string> pragma_split =
split_string(line_after_pragma, 2);
line =
(line.substr(0, pragma_beg) + "_Pragma(\"" + pragma_split[1] + "\")");
if (pragma_split.size() == 3) {
line += " " + pragma_split[2];
}
}
source += line + "\n";
}
// HACK TESTING (WAR for cub)
// source = "#define cudaDeviceSynchronize() cudaSuccess\n" + source;
////source = "cudaError_t cudaDeviceSynchronize() { return cudaSuccess; }\n" +
/// source;
// WAR for #pragma once causing problems when there are multiple inclusions
// of the same header from different paths.
if (pragma_once) {
std::stringstream ss;
ss << std::uppercase << std::hex << std::setw(8) << std::setfill('0')
<< hash;
std::string include_guard_name = "_JITIFY_INCLUDE_GUARD_" + ss.str() + "\n";
std::string include_guard_header;
include_guard_header += "#ifndef " + include_guard_name;
include_guard_header += "#define " + include_guard_name;
std::string include_guard_footer;
include_guard_footer += "#endif // " + include_guard_name;
source = include_guard_header + source + "\n" + include_guard_footer;
}
// return filename;
return true;
}
} // namespace detail
//! \endcond
/*! Jitify reflection utilities namespace
*/
namespace reflection {
// Provides type and value reflection via a function 'reflect':
// reflect<Type>() -> "Type"
// reflect(value) -> "(T)value"
// reflect<VAL>() -> "VAL"
// reflect<Type,VAL> -> "VAL"
// reflect_template<float,NonType<int,7>,char>() -> "<float,7,char>"
// reflect_template({"float", "7", "char"}) -> "<float,7,char>"
/*! A wrapper class for non-type template parameters.
*/
template <typename T, T VALUE_>
struct NonType {
#if defined __cplusplus && __cplusplus >= 201103L
constexpr
#endif
static T VALUE = VALUE_;
};
// Forward declaration
template <typename T>
inline std::string reflect(T const& value);
//! \cond
namespace detail {
template <typename T>
inline std::string value_string(const T& x) {
std::stringstream ss;
ss << x;
return ss.str();
}
// WAR for non-printable characters
template <>
inline std::string value_string<char>(const char& x) {
std::stringstream ss;
ss << (int)x;
return ss.str();
}
template <>
inline std::string value_string<signed char>(const signed char& x) {
std::stringstream ss;
ss << (int)x;
return ss.str();
}
template <>
inline std::string value_string<unsigned char>(const unsigned char& x) {
std::stringstream ss;
ss << (int)x;
return ss.str();
}
template <>
inline std::string value_string<wchar_t>(const wchar_t& x) {
std::stringstream ss;
ss << (long)x;
return ss.str();
}
// Specialisation for bool true/false literals
template <>
inline std::string value_string<bool>(const bool& x) {
return x ? "true" : "false";
}
//#if CUDA_VERSION < 8000
#ifdef _MSC_VER // MSVC compiler
inline std::string demangle(const char* verbose_name) {
// Strips annotations from the verbose name returned by typeid(X).name()
std::string result = verbose_name;
result = jitify::detail::replace_token(result, "__ptr64", "");
result = jitify::detail::replace_token(result, "__cdecl", "");
result = jitify::detail::replace_token(result, "class", "");
result = jitify::detail::replace_token(result, "struct", "");
return result;
}
#else // not MSVC
#include <cxxabi.h>
inline std::string demangle(const char* mangled_name) {
size_t bufsize = 1024;
auto buf = std::unique_ptr<char, decltype(free)*>(
reinterpret_cast<char*>(malloc(bufsize)), free);
std::string demangled_name;
int status;
char* demangled_ptr =
abi::__cxa_demangle(mangled_name, buf.get(), &bufsize, &status);
if (status == 0) {
demangled_name = demangled_ptr; // all worked as expected
}
else if (status == -2) {
demangled_name = mangled_name; // we interpret this as plain C name
}
else if (status == -1) {
throw std::runtime_error(
std::string("memory allocation failure in __cxa_demangle"));
}
else if (status == -3) {
throw std::runtime_error(std::string("invalid argument to __cxa_demangle"));
}
return demangled_name;
}
#endif // not MSVC
//#endif // CUDA_VERSION < 8000
template <typename T>
struct type_reflection {
inline static std::string name() {
//#if CUDA_VERSION < 8000
// WAR for typeid discarding cv qualifiers on value-types
// We use a pointer type to maintain cv qualifiers, then strip out the '*'
std::string no_cv_name = demangle(typeid(T).name());
std::string ptr_name = demangle(typeid(T*).name());
// Find the right '*' by diffing the type name and ptr name
// Note that the '*' will also be prefixed with the cv qualifiers
size_t diff_begin =
std::mismatch(no_cv_name.begin(), no_cv_name.end(), ptr_name.begin())
.first -
no_cv_name.begin();
size_t star_begin = ptr_name.find("*", diff_begin);
if (star_begin == std::string::npos) {
throw std::runtime_error("Type reflection failed: " + ptr_name);
}
std::string name =
ptr_name.substr(0, star_begin) + ptr_name.substr(star_begin + 1);
return name;
//#else
// std::string ret;
// nvrtcResult status = nvrtcGetTypeName<T>(&ret);
// if( status != NVRTC_SUCCESS ) {
// throw std::runtime_error(std::string("nvrtcGetTypeName
// failed:
//")+ nvrtcGetErrorString(status));
// }
// return ret;
//#endif
}
};
template <typename T, T VALUE>
struct type_reflection<NonType<T, VALUE> > {
inline static std::string name() {
return jitify::reflection::reflect(VALUE);
}
};
} // namespace detail
//! \endcond
/*! Create an Instance object that contains a const reference to the
* value. We use this to wrap abstract objects from which we want to extract
* their type at runtime (e.g., derived type). This is used to facilitate
* templating on derived type when all we know at compile time is abstract
* type.
*/
template <typename T>
struct Instance {
const T& value;
Instance(const T& value) : value(value) {}
};
/*! Create an Instance object from which we can extract the value's run-time
* type.
* \param value The const value to be captured.
*/
template <typename T>
inline Instance<T const> instance_of(T const& value) {
return Instance<T const>(value);
}
/*! A wrapper used for representing types as values.
*/
template <typename T>
struct Type {};
// Type reflection
// E.g., reflect<float>() -> "float"
// Note: This strips trailing const and volatile qualifiers
/*! Generate a code-string for a type.
* \code{.cpp}reflect<float>() --> "float"\endcode
*/
template <typename T>
inline std::string reflect() {
return detail::type_reflection<T>::name();
}
// Value reflection
// E.g., reflect(3.14f) -> "(float)3.14"
/*! Generate a code-string for a value.
* \code{.cpp}reflect(3.14f) --> "(float)3.14"\endcode
*/
template <typename T>
inline std::string reflect(T const& value) {
return "(" + reflect<T>() + ")" + detail::value_string(value);
}
// Non-type template arg reflection (implicit conversion to int64_t)
// E.g., reflect<7>() -> "(int64_t)7"
/*! Generate a code-string for an integer non-type template argument.
* \code{.cpp}reflect<7>() --> "(int64_t)7"\endcode
*/
template <long long N>
inline std::string reflect() {
return reflect<NonType<int, N> >();
}
// Non-type template arg reflection (explicit type)
// E.g., reflect<int,7>() -> "(int)7"
/*! Generate a code-string for a generic non-type template argument.
* \code{.cpp} reflect<int,7>() --> "(int)7" \endcode
*/
template <typename T, T N>
inline std::string reflect() {
return reflect<NonType<T, N> >();
}
// Type reflection via value
// E.g., reflect(Type<float>()) -> "float"
/*! Generate a code-string for a type wrapped as a Type instance.
* \code{.cpp}reflect(Type<float>()) --> "float"\endcode
*/
template <typename T>
inline std::string reflect(jitify::reflection::Type<T>) {
return reflect<T>();
}
/*! Generate a code-string for a type wrapped as an Instance instance.
* \code{.cpp}reflect(Instance<float>(3.1f)) --> "float"\endcode
* or more simply when passed to a instance_of helper
* \code{.cpp}reflect(instance_of(3.1f)) --> "float"\endcodei
* This is specifically for the case where we want to extract the run-time
* type, e.g., derived type, of an object pointer.
*/
template <typename T>
inline std::string reflect(jitify::reflection::Instance<T>& value) {
return detail::demangle(typeid(value.value).name());
}
// Type from value
// E.g., type_of(3.14f) -> Type<float>()
/*! Create a Type object representing a value's type.
* \param value The value whose type is to be captured.
*/
template <typename T>
inline Type<T> type_of(T& value) {
return Type<T>();
}
/*! Create a Type object representing a value's type.
* \param value The const value whose type is to be captured.
*/
template <typename T>
inline Type<T const> type_of(T const& value) {
return Type<T const>();
}
#if __cplusplus >= 201103L
// Multiple value reflections one call, returning list of strings
template <typename... Args>
inline std::vector<std::string> reflect_all(Args... args) {
return{ reflect(args)... };
}
#endif // __cplusplus >= 201103L
inline std::string reflect_list(jitify::detail::vector<std::string> const& args,
std::string opener = "",
std::string closer = "") {
std::stringstream ss;
ss << opener;
for (int i = 0; i < (int)args.size(); ++i) {
if (i > 0) ss << ",";
ss << args[i];
}
ss << closer;
return ss.str();
}
// Template instantiation reflection
// inline std::string reflect_template(std::vector<std::string> const& args) {
inline std::string reflect_template(
jitify::detail::vector<std::string> const& args) {
// Note: The space in " >" is a WAR to avoid '>>' appearing
return reflect_list(args, "<", " >");
}
#if __cplusplus >= 201103L
// TODO: See if can make this evaluate completely at compile-time
template <typename... Ts>
inline std::string reflect_template() {
return reflect_template({ reflect<Ts>()... });
// return reflect_template<sizeof...(Ts)>({reflect<Ts>()...});
}
#endif
} // namespace reflection
//! \cond
namespace detail {
class CUDAKernel {
std::vector<std::string> _link_files;
std::vector<std::string> _link_paths;
CUlinkState _link_state;
CUmodule _module;
CUfunction _kernel;
std::string _func_name;
std::string _ptx;
std::map<std::string, std::string> _constant_map;
std::vector<CUjit_option> _opts;
std::vector<void*> _optvals;
#ifdef JITIFY_PRINT_LINKER_LOG
static const unsigned int _log_size = 8192;
char _info_log[_log_size];
#endif
inline void cuda_safe_call(CUresult res) const {
if (res != CUDA_SUCCESS) {
const char* msg;
cuGetErrorName(res, &msg);
throw std::runtime_error(msg);
}
}
inline void create_module(std::vector<std::string> link_files,
std::vector<std::string> link_paths) {
#ifndef JITIFY_PRINT_LINKER_LOG
// WAR since linker log does not seem to be constructed using a single call
// to cuModuleLoadDataEx.
if (link_files.empty()) {
assert(cudaDeviceSynchronize() == CUDA_SUCCESS);
cuda_safe_call(cuCtxSynchronize());
// CUDA_ERROR_NOT_SUPPORTED thrown by following method
cuda_safe_call(cuModuleLoadDataEx(&_module, _ptx.c_str(), _opts.size(),
_opts.data(), _optvals.data()));
}
else
#endif
{
cuda_safe_call(cuLinkCreate(_opts.size(), _opts.data(), _optvals.data(),
&_link_state));
cuda_safe_call(cuLinkAddData(_link_state, CU_JIT_INPUT_PTX,
(void*)_ptx.c_str(), _ptx.size(),
"jitified_source.ptx", 0, 0, 0));
for (int i = 0; i < (int)link_files.size(); ++i) {
std::string link_file = link_files[i];
#if defined _WIN32 || defined _WIN64
link_file = link_file + ".lib";
#else
link_file = "lib" + link_file + ".a";
#endif
CUresult result = cuLinkAddFile(_link_state, CU_JIT_INPUT_LIBRARY,
link_file.c_str(), 0, 0, 0);
int path_num = 0;
while (result == CUDA_ERROR_FILE_NOT_FOUND &&
path_num < (int)link_paths.size()) {
std::string filename = path_join(link_paths[path_num++], link_file);
result = cuLinkAddFile(_link_state, CU_JIT_INPUT_LIBRARY,
filename.c_str(), 0, 0, 0);
}
#if JITIFY_PRINT_LOG
if (result == CUDA_ERROR_FILE_NOT_FOUND) {
std::cout << "Error: Device library not found: " << link_file
<< std::endl;
}
#endif
cuda_safe_call(result);
}
size_t cubin_size;
void* cubin;
cuda_safe_call(cuLinkComplete(_link_state, &cubin, &cubin_size));
cuda_safe_call(cuModuleLoadData(&_module, cubin));
}
#ifdef JITIFY_PRINT_LINKER_LOG
std::cout << "---------------------------------------" << std::endl;
std::cout << "--- Linker for "
<< reflection::detail::demangle(_func_name.c_str()) << " ---"
<< std::endl;
std::cout << "---------------------------------------" << std::endl;
std::cout << _info_log << std::endl;
std::cout << "---------------------------------------" << std::endl;
#endif
cuda_safe_call(cuModuleGetFunction(&_kernel, _module, _func_name.c_str()));
}
inline void destroy_module() {
if (_link_state) {
cuda_safe_call(cuLinkDestroy(_link_state));
}
_link_state = 0;
if (_module) {
cuModuleUnload(_module);
}
_module = 0;
}
// create a map of constants in the ptx file mapping demangled to mangled name
inline void create_constant_map() {
size_t pos = 0;
while (pos < _ptx.size()) {
pos = _ptx.find(".const .align", pos);
if (pos == std::string::npos) break;
size_t end = _ptx.find(";", pos);
std::string line = _ptx.substr(pos, end - pos);
size_t constant_start = line.find_last_of(" ") + 1;
size_t constant_end = line.find_last_of("[");
std::string entry =
line.substr(constant_start, constant_end - constant_start);
#ifdef _MSC_VER // interpret anything that doesn't begine ? as unmangled name
std::string key = (entry[0] != '?')
? entry.c_str()
: reflection::detail::demangle(entry.c_str());
#else // interpret anything that doesn't begine _Z as unmangled name
std::string key = (entry[0] != '_' && entry[1] != 'Z')
? entry.c_str()
: reflection::detail::demangle(entry.c_str());
#endif
_constant_map[key] = entry;
pos = end;
}
}
inline void set_linker_log() {
#ifdef JITIFY_PRINT_LINKER_LOG
_opts.push_back(CU_JIT_INFO_LOG_BUFFER);
_optvals.push_back((void*)_info_log);
_opts.push_back(CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES);
_optvals.push_back((void*)(long)_log_size);
_opts.push_back(CU_JIT_LOG_VERBOSE);
_optvals.push_back((void*)1);
#endif
}
public:
inline CUDAKernel() : _link_state(0), _module(0), _kernel(0) {}
inline CUDAKernel(const CUDAKernel& other) = delete;
inline CUDAKernel& operator=(const CUDAKernel& other) = delete;
inline CUDAKernel(CUDAKernel&& other) = delete;
inline CUDAKernel& operator=(CUDAKernel&& other) = delete;
inline CUDAKernel(const char* func_name, const char* ptx,
std::vector<std::string> link_files,
std::vector<std::string> link_paths, unsigned int nopts = 0,
CUjit_option* opts = 0, void** optvals = 0)
: _link_files(link_files),
_link_paths(link_paths),
_link_state(0),
_module(0),
_kernel(0),
_func_name(func_name),
_ptx(ptx),
_opts(opts, opts + nopts),
_optvals(optvals, optvals + nopts) {
this->set_linker_log();
this->create_module(link_files, link_paths);
this->create_constant_map();
}
inline CUDAKernel& set(const char* func_name, const char* ptx,
std::vector<std::string> link_files,
std::vector<std::string> link_paths,
unsigned int nopts = 0, CUjit_option* opts = 0,
void** optvals = 0) {
this->destroy_module();
_func_name = func_name;
_ptx = ptx;
_link_files = link_files;
_link_paths = link_paths;
_opts.assign(opts, opts + nopts);
_optvals.assign(optvals, optvals + nopts);
this->set_linker_log();
this->create_module(link_files, link_paths);
this->create_constant_map();
return *this;
}
inline ~CUDAKernel() { this->destroy_module(); }
inline operator CUfunction() const { return _kernel; }
inline CUresult launch(dim3 grid, dim3 block, unsigned int smem,
CUstream stream, std::vector<void*> arg_ptrs) const {
return cuLaunchKernel(_kernel, grid.x, grid.y, grid.z, block.x, block.y,
block.z, smem, stream, arg_ptrs.data(), NULL);
}
inline CUdeviceptr get_constant_ptr(const char* name) const {
CUdeviceptr const_ptr = 0;
auto constant = _constant_map.find(name);
if (constant != _constant_map.end()) {
cuda_safe_call(
cuModuleGetGlobal(&const_ptr, 0, _module, constant->second.c_str()));
}
else {
throw std::runtime_error(std::string("failed to look up constant ") +
name);
}
return const_ptr;
}
const std::string& function_name() const { return _func_name; }
const std::string& ptx() const { return _ptx; }
const std::vector<std::string>& link_files() const { return _link_files; }
const std::vector<std::string>& link_paths() const { return _link_paths; }
};
static const char* jitsafe_header_preinclude_h = R"(
//// WAR for Thrust (which appears to have forgotten to include this in result_of_adaptable_function.h
//#include <type_traits>
//// WAR for Thrust (which appear to have forgotten to include this in error_code.h)
//#include <string>
// WAR for Thrust (which only supports gnuc, clang or msvc)
#define __GNUC__ 4
// WAR for generics/shfl.h
#define THRUST_STATIC_ASSERT(x)
// WAR for CUB
#ifdef __host__
#undef __host__
#endif
#define __host__
// WAR to allow exceptions to be parsed
#define try
#define catch(...)
)";
static const char* jitsafe_header_float_h =
"#pragma once\n"
"\n"
"inline __host__ __device__ float jitify_int_as_float(int i) "
"{ union FloatInt { float f; int i; } fi; fi.i = i; return fi.f; }\n"
"inline __host__ __device__ double jitify_longlong_as_double(long long i) "
"{ union DoubleLongLong { double f; long long i; } fi; fi.i = i; return "
"fi.f; }\n"
"#define FLT_RADIX 2\n"
"#define FLT_MANT_DIG 24\n"
"#define DBL_MANT_DIG 53\n"
"#define FLT_DIG 6\n"
"#define DBL_DIG 15\n"
"#define FLT_MIN_EXP -125\n"
"#define DBL_MIN_EXP -1021\n"
"#define FLT_MIN_10_EXP -37\n"
"#define DBL_MIN_10_EXP -307\n"
"#define FLT_MAX_EXP 128\n"
"#define DBL_MAX_EXP 1024\n"
"#define FLT_MAX_10_EXP 38\n"
"#define DBL_MAX_10_EXP 308\n"
"#define FLT_MAX jitify_int_as_float(2139095039)\n"
"#define DBL_MAX jitify_longlong_as_double(9218868437227405311)\n"
"#define FLT_EPSILON jitify_int_as_float(872415232)\n"
"#define DBL_EPSILON jitify_longlong_as_double(4372995238176751616)\n"
"#define FLT_MIN jitify_int_as_float(8388608)\n"
"#define DBL_MIN jitify_longlong_as_double(4503599627370496)\n"
"#define FLT_ROUNDS 1\n"
"#if defined __cplusplus && __cplusplus >= 201103L\n"
"#define FLT_EVAL_METHOD 0\n"
"#define DECIMAL_DIG 21\n"
"#endif\n";
static const char* jitsafe_header_limits_h =
"#pragma once\n"
"\n"
"#if defined _WIN32 || defined _WIN64\n"
" #define __WORDSIZE 32\n"
"#else\n"
" #if defined __x86_64__ && !defined __ILP32__\n"
" #define __WORDSIZE 64\n"
" #else\n"
" #define __WORDSIZE 32\n"
" #endif\n"
"#endif\n"
"#define MB_LEN_MAX 16\n"
"#define CHAR_BIT 8\n"
"#define SCHAR_MIN (-128)\n"
"#define SCHAR_MAX 127\n"
"#define UCHAR_MAX 255\n"
"#ifdef __CHAR_UNSIGNED__\n"
" #define CHAR_MIN 0\n"
" #define CHAR_MAX UCHAR_MAX\n"
"#else\n"
" #define CHAR_MIN SCHAR_MIN\n"
" #define CHAR_MAX SCHAR_MAX\n"
"#endif\n"
"#define SHRT_MIN (-32768)\n"
"#define SHRT_MAX 32767\n"
"#define USHRT_MAX 65535\n"
"#define INT_MIN (-INT_MAX - 1)\n"
"#define INT_MAX 2147483647\n"
"#define UINT_MAX 4294967295U\n"
"#if __WORDSIZE == 64\n"
" # define LONG_MAX 9223372036854775807L\n"
"#else\n"
" # define LONG_MAX 2147483647L\n"
"#endif\n"
"#define LONG_MIN (-LONG_MAX - 1L)\n"
"#if __WORDSIZE == 64\n"
" #define ULONG_MAX 18446744073709551615UL\n"
"#else\n"
" #define ULONG_MAX 4294967295UL\n"
"#endif\n"
"#define LLONG_MAX 9223372036854775807LL\n"
"#define LLONG_MIN (-LLONG_MAX - 1LL)\n"
"#define ULLONG_MAX 18446744073709551615ULL\n";
static const char* jitsafe_header_iterator =
"#pragma once\n"
"\n"
"namespace __jitify_iterator_ns {\n"
"struct output_iterator_tag {};\n"
"struct input_iterator_tag {};\n"
"struct forward_iterator_tag {};\n"
"struct bidirectional_iterator_tag {};\n"
"struct random_access_iterator_tag {};\n"
"template<class Iterator>\n"
"struct iterator_traits {\n"
" typedef typename Iterator::iterator_category iterator_category;\n"
" typedef typename Iterator::value_type value_type;\n"
" typedef typename Iterator::difference_type difference_type;\n"
" typedef typename Iterator::pointer pointer;\n"
" typedef typename Iterator::reference reference;\n"
"};\n"
"template<class T>\n"
"struct iterator_traits<T*> {\n"
" typedef random_access_iterator_tag iterator_category;\n"
" typedef T value_type;\n"
" typedef ptrdiff_t difference_type;\n"
" typedef T* pointer;\n"
" typedef T& reference;\n"
"};\n"
"template<class T>\n"
"struct iterator_traits<T const*> {\n"
" typedef random_access_iterator_tag iterator_category;\n"
" typedef T value_type;\n"
" typedef ptrdiff_t difference_type;\n"
" typedef T const* pointer;\n"
" typedef T const& reference;\n"
"};\n"
"} // namespace __jitify_iterator_ns\n"
"namespace std { using namespace __jitify_iterator_ns; }\n"
"using namespace __jitify_iterator_ns;\n";
// TODO: This is incomplete; need floating point limits
static const char* jitsafe_header_limits =
"#pragma once\n"
"#include <climits>\n"
"\n"
"namespace std {\n"
"// TODO: Floating-point limits\n"
"namespace __jitify_detail {\n"
"template<class T, T Min, T Max, int Digits=-1>\n"
"struct IntegerLimits {\n"
" static inline __host__ __device__ T min() { return Min; }\n"
" static inline __host__ __device__ T max() { return Max; }\n"
"#if __cplusplus >= 201103L\n"
" static constexpr inline __host__ __device__ T lowest() noexcept {\n"
" return Min;\n"
" }\n"
"#endif // __cplusplus >= 201103L\n"
" enum {\n"
" is_specialized = true,\n"
" digits = (Digits == -1) ? (int)(sizeof(T)*8 - (Min != 0)) "
": Digits,\n"
" digits10 = (digits * 30103) / 100000,\n"
" is_signed = ((T)(-1)<0),\n"
" is_integer = true,\n"
" is_exact = true,\n"
" radix = 2,\n"
" is_bounded = true,\n"
" is_modulo = false\n"
" };\n"
"};\n"
"} // namespace detail\n"
"template<typename T> struct numeric_limits {\n"
" enum { is_specialized = false };\n"
"};\n"
"template<> struct numeric_limits<bool> : public "
"__jitify_detail::IntegerLimits<bool, false, true,1> {};\n"
"template<> struct numeric_limits<char> : public "
"__jitify_detail::IntegerLimits<char, CHAR_MIN, CHAR_MAX> "
"{};\n"
"template<> struct numeric_limits<signed char> : public "
"__jitify_detail::IntegerLimits<signed char, SCHAR_MIN,SCHAR_MAX> "
"{};\n"
"template<> struct numeric_limits<unsigned char> : public "
"__jitify_detail::IntegerLimits<unsigned char, 0, UCHAR_MAX> "
"{};\n"
"template<> struct numeric_limits<wchar_t> : public "
"__jitify_detail::IntegerLimits<wchar_t, INT_MIN, INT_MAX> {};\n"
"template<> struct numeric_limits<short> : public "
"__jitify_detail::IntegerLimits<short, SHRT_MIN, SHRT_MAX> "
"{};\n"
"template<> struct numeric_limits<unsigned short> : public "
"__jitify_detail::IntegerLimits<unsigned short, 0, USHRT_MAX> "
"{};\n"
"template<> struct numeric_limits<int> : public "
"__jitify_detail::IntegerLimits<int, INT_MIN, INT_MAX> {};\n"
"template<> struct numeric_limits<unsigned int> : public "
"__jitify_detail::IntegerLimits<unsigned int, 0, UINT_MAX> "
"{};\n"
"template<> struct numeric_limits<long> : public "
"__jitify_detail::IntegerLimits<long, LONG_MIN, LONG_MAX> "
"{};\n"
"template<> struct numeric_limits<unsigned long> : public "
"__jitify_detail::IntegerLimits<unsigned long, 0, ULONG_MAX> "
"{};\n"
"template<> struct numeric_limits<long long> : public "
"__jitify_detail::IntegerLimits<long long, LLONG_MIN,LLONG_MAX> "
"{};\n"
"template<> struct numeric_limits<unsigned long long> : public "
"__jitify_detail::IntegerLimits<unsigned long long,0, ULLONG_MAX> "
"{};\n"
"//template<typename T> struct numeric_limits { static const bool "
"is_signed = ((T)(-1)<0); };\n"
"} // namespace std\n";
// TODO: This is highly incomplete
static const char* jitsafe_header_type_traits = R"(
#pragma once
#if __cplusplus >= 201103L
namespace __jitify_type_traits_ns {
template<bool B, class T = void> struct enable_if {};
template<class T> struct enable_if<true, T> { typedef T type; };
#if __cplusplus >= 201402L
template< bool B, class T = void > using enable_if_t = typename enable_if<B,T>::type;
#endif
struct true_type {
enum { value = true };
operator bool() const { return true; }
};
struct false_type {
enum { value = false };
operator bool() const { return false; }
};
template<typename T> struct is_floating_point : false_type {};
template<> struct is_floating_point<float> : true_type {};
template<> struct is_floating_point<double> : true_type {};
template<> struct is_floating_point<long double> : true_type {};
template<class T> struct is_integral : false_type {};
template<> struct is_integral<bool> : true_type {};
template<> struct is_integral<char> : true_type {};
template<> struct is_integral<signed char> : true_type {};
template<> struct is_integral<unsigned char> : true_type {};
template<> struct is_integral<short> : true_type {};
template<> struct is_integral<unsigned short> : true_type {};
template<> struct is_integral<int> : true_type {};
template<> struct is_integral<unsigned int> : true_type {};
template<> struct is_integral<long> : true_type {};
template<> struct is_integral<unsigned long> : true_type {};
template<> struct is_integral<long long> : true_type {};
template<> struct is_integral<unsigned long long> : true_type {};
template<typename T> struct is_signed : false_type {};
template<> struct is_signed<float> : true_type {};
template<> struct is_signed<double> : true_type {};
template<> struct is_signed<long double> : true_type {};
template<> struct is_signed<signed char> : true_type {};
template<> struct is_signed<short> : true_type {};
template<> struct is_signed<int> : true_type {};
template<> struct is_signed<long> : true_type {};
template<> struct is_signed<long long> : true_type {};
template<typename T> struct is_unsigned : false_type {};
template<> struct is_unsigned<unsigned char> : true_type {};
template<> struct is_unsigned<unsigned short> : true_type {};
template<> struct is_unsigned<unsigned int> : true_type {};
template<> struct is_unsigned<unsigned long> : true_type {};
template<> struct is_unsigned<unsigned long long> : true_type {};
template<typename T, typename U> struct is_same : false_type {};
template<typename T> struct is_same<T,T> : true_type {};
template<class T> struct is_array : false_type {};
template<class T> struct is_array<T[]> : true_type {};
template<class T, size_t N> struct is_array<T[N]> : true_type {};
//partial implementation only of is_function
template<class> struct is_function : false_type { };
template<class Ret, class... Args> struct is_function<Ret(Args...)> : true_type {}; //regular
template<class Ret, class... Args> struct is_function<Ret(Args......)> : true_type {}; // variadic
template<class> struct result_of;
template<class F, typename... Args>
struct result_of<F(Args...)> {
// TODO: This is a hack; a proper implem is quite complicated.
typedef typename F::result_type type;
};
template <class T> struct remove_reference { typedef T type; };
template <class T> struct remove_reference<T&> { typedef T type; };
template <class T> struct remove_reference<T&&> { typedef T type; };
#if __cplusplus >= 201402L
template< class T > using remove_reference_t = typename remove_reference<T>::type;
#endif
template<class T> struct remove_extent { typedef T type; };
template<class T> struct remove_extent<T[]> { typedef T type; };
template<class T, size_t N> struct remove_extent<T[N]> { typedef T type; };
#if __cplusplus >= 201402L
template< class T > using remove_extent_t = typename remove_extent<T>::type;
#endif
template< class T > struct remove_const { typedef T type; };
template< class T > struct remove_const<const T> { typedef T type; };
template< class T > struct remove_volatile { typedef T type; };
template< class T > struct remove_volatile<volatile T> { typedef T type; };
template< class T > struct remove_cv { typedef typename remove_volatile<typename remove_const<T>::type>::type type; };
#if __cplusplus >= 201402L
template< class T > using remove_cv_t = typename remove_cv<T>::type;
template< class T > using remove_const_t = typename remove_const<T>::type;
template< class T > using remove_volatile_t = typename remove_volatile<T>::type;
#endif
template<bool B, class T, class F> struct conditional { typedef T type; };
template<class T, class F> struct conditional<false, T, F> { typedef F type; };
#if __cplusplus >= 201402L
template< bool B, class T, class F > using conditional_t = typename conditional<B,T,F>::type;
#endif
namespace __jitify_detail {
template< class T, bool is_function_type = false > struct add_pointer { using type = typename remove_reference<T>::type*; };
template< class T > struct add_pointer<T, true> { using type = T; };
template< class T, class... Args > struct add_pointer<T(Args...), true> { using type = T(*)(Args...); };
template< class T, class... Args > struct add_pointer<T(Args..., ...), true> { using type = T(*)(Args..., ...); };
}
template< class T > struct add_pointer : __jitify_detail::add_pointer<T, is_function<T>::value> {};
#if __cplusplus >= 201402L
template< class T > using add_pointer_t = typename add_pointer<T>::type;
#endif
template< class T > struct decay {
private:
typedef typename remove_reference<T>::type U;
public:
typedef typename conditional<is_array<U>::value, typename remove_extent<U>::type*,
typename conditional<is_function<U>::value,typename add_pointer<U>::type,typename remove_cv<U>::type
>::type>::type type;
};
#if __cplusplus >= 201402L
template< class T > using decay_t = typename decay<T>::type;
#endif
} // namespace __jtiify_type_traits_ns
namespace std { using namespace __jitify_type_traits_ns; }
using namespace __jitify_type_traits_ns;
#endif // c++11
)";
// TODO: INT_FAST8_MAX et al. and a few other misc constants
static const char* jitsafe_header_stdint_h =
"#pragma once\n"
"#include <climits>\n"
"namespace __jitify_stdint_ns {\n"
"typedef signed char int8_t;\n"
"typedef signed short int16_t;\n"
"typedef signed int int32_t;\n"
"typedef signed long long int64_t;\n"
"typedef signed char int_fast8_t;\n"
"typedef signed short int_fast16_t;\n"
"typedef signed int int_fast32_t;\n"
"typedef signed long long int_fast64_t;\n"
"typedef signed char int_least8_t;\n"
"typedef signed short int_least16_t;\n"
"typedef signed int int_least32_t;\n"
"typedef signed long long int_least64_t;\n"
"typedef signed long long intmax_t;\n"
"typedef signed long intptr_t; //optional\n"
"typedef unsigned char uint8_t;\n"
"typedef unsigned short uint16_t;\n"
"typedef unsigned int uint32_t;\n"
"typedef unsigned long long uint64_t;\n"
"typedef unsigned char uint_fast8_t;\n"
"typedef unsigned short uint_fast16_t;\n"
"typedef unsigned int uint_fast32_t;\n"
"typedef unsigned long long uint_fast64_t;\n"
"typedef unsigned char uint_least8_t;\n"
"typedef unsigned short uint_least16_t;\n"
"typedef unsigned int uint_least32_t;\n"
"typedef unsigned long long uint_least64_t;\n"
"typedef unsigned long long uintmax_t;\n"
"typedef unsigned long uintptr_t; //optional\n"
"#define INT8_MIN SCHAR_MIN\n"
"#define INT16_MIN SHRT_MIN\n"
"#define INT32_MIN INT_MIN\n"
"#define INT64_MIN LLONG_MIN\n"
"#define INT8_MAX SCHAR_MAX\n"
"#define INT16_MAX SHRT_MAX\n"
"#define INT32_MAX INT_MAX\n"
"#define INT64_MAX LLONG_MAX\n"
"#define UINT8_MAX UCHAR_MAX\n"
"#define UINT16_MAX USHRT_MAX\n"
"#define UINT32_MAX UINT_MAX\n"
"#define UINT64_MAX ULLONG_MAX\n"
"#define INTPTR_MIN LONG_MIN\n"
"#define INTMAX_MIN LLONG_MIN\n"
"#define INTPTR_MAX LONG_MAX\n"
"#define INTMAX_MAX LLONG_MAX\n"
"#define UINTPTR_MAX ULONG_MAX\n"
"#define UINTMAX_MAX ULLONG_MAX\n"
"#define PTRDIFF_MIN INTPTR_MIN\n"
"#define PTRDIFF_MAX INTPTR_MAX\n"
"#define SIZE_MAX UINT64_MAX\n"
"} // namespace __jitify_stdint_ns\n"
"namespace std { using namespace __jitify_stdint_ns; }\n"
"using namespace __jitify_stdint_ns;\n";
// TODO: offsetof
static const char* jitsafe_header_stddef_h =
"#pragma once\n"
"#include <climits>\n"
"namespace __jitify_stddef_ns {\n"
//"enum { NULL = 0 };\n"
"typedef unsigned long size_t;\n"
"typedef signed long ptrdiff_t;\n"
"#if __cplusplus >= 201103L\n"
"typedef decltype(nullptr) nullptr_t;\n"
"#if defined(_MSC_VER)\n"
" typedef double max_align_t;\n"
"#elif defined(__APPLE__)\n"
" typedef long double max_align_t;\n"
"#else\n"
" // Define max_align_t to match the GCC definition.\n"
" typedef struct {\n"
" long long __jitify_max_align_nonce1\n"
" __attribute__((__aligned__(__alignof__(long long))));\n"
" long double __jitify_max_align_nonce2\n"
" __attribute__((__aligned__(__alignof__(long double))));\n"
" } max_align_t;\n"
"#endif\n"
"#endif // __cplusplus >= 201103L\n"
"#if __cplusplus >= 201703L\n"
"enum class byte : unsigned char {};\n"
"#endif // __cplusplus >= 201703L\n"
"} // namespace __jitify_stddef_ns\n"
"namespace std { using namespace __jitify_stddef_ns; }\n"
"using namespace __jitify_stddef_ns;\n";
static const char* jitsafe_header_stdlib_h =
"#pragma once\n"
"#include <stddef.h>\n";
static const char* jitsafe_header_stdio_h =
"#pragma once\n"
"#include <stddef.h>\n"
"#define FILE int\n"
"int fflush ( FILE * stream );\n"
"int fprintf ( FILE * stream, const char * format, ... );\n";
static const char* jitsafe_header_string_h =
"#pragma once\n"
"char* strcpy ( char * destination, const char * source );\n"
"int strcmp ( const char * str1, const char * str2 );\n"
"char* strerror( int errnum );\n";
static const char* jitsafe_header_cstring =
"#pragma once\n"
"\n"
"namespace __jitify_cstring_ns {\n"
"char* strcpy ( char * destination, const char * source );\n"
"int strcmp ( const char * str1, const char * str2 );\n"
"char* strerror( int errnum );\n"
"} // namespace __jitify_cstring_ns\n"
"namespace std { using namespace __jitify_cstring_ns; }\n"
"using namespace __jitify_cstring_ns;\n";
// HACK TESTING (WAR for cub)
static const char* jitsafe_header_iostream =
"#pragma once\n"
"#include <ostream>\n"
"#include <istream>\n";
// HACK TESTING (WAR for Thrust)
static const char* jitsafe_header_ostream =
"#pragma once\n"
"\n"
"namespace __jitify_ostream_ns {\n"
"template<class CharT,class Traits=void>\n" // = std::char_traits<CharT>
// >\n"
"struct basic_ostream {\n"
"};\n"
"typedef basic_ostream<char> ostream;\n"
"ostream& endl(ostream& os);\n"
"ostream& operator<<( ostream&, ostream& (*f)( ostream& ) );\n"
"template< class CharT, class Traits > basic_ostream<CharT, Traits>& endl( "
"basic_ostream<CharT, Traits>& os );\n"
"template< class CharT, class Traits > basic_ostream<CharT, Traits>& "
"operator<<( basic_ostream<CharT,Traits>& os, const char* c );\n"
"#if __cplusplus >= 201103L\n"
"template< class CharT, class Traits, class T > basic_ostream<CharT, "
"Traits>& operator<<( basic_ostream<CharT,Traits>&& os, const T& value );\n"
"#endif // __cplusplus >= 201103L\n"
"} // namespace __jitify_ostream_ns\n"
"namespace std { using namespace __jitify_ostream_ns; }\n"
"using namespace __jitify_ostream_ns;\n";
static const char* jitsafe_header_istream =
"#pragma once\n"
"\n"
"namespace __jitify_istream_ns {\n"
"template<class CharT,class Traits=void>\n" // = std::char_traits<CharT>
// >\n"
"struct basic_istream {\n"
"};\n"
"typedef basic_istream<char> istream;\n"
"} // namespace __jitify_istream_ns\n"
"namespace std { using namespace __jitify_istream_ns; }\n"
"using namespace __jitify_istream_ns;\n";
static const char* jitsafe_header_sstream =
"#pragma once\n"
"#include <ostream>\n"
"#include <istream>\n";
static const char* jitsafe_header_utility =
"#pragma once\n"
"namespace __jitify_utility_ns {\n"
"template<class T1, class T2>\n"
"struct pair {\n"
" T1 first;\n"
" T2 second;\n"
" inline pair() {}\n"
" inline pair(T1 const& first_, T2 const& second_)\n"
" : first(first_), second(second_) {}\n"
" // TODO: Standard includes many more constructors...\n"
" // TODO: Comparison operators\n"
"};\n"
"template<class T1, class T2>\n"
"pair<T1,T2> make_pair(T1 const& first, T2 const& second) {\n"
" return pair<T1,T2>(first, second);\n"
"}\n"
"} // namespace __jitify_utility_ns\n"
"namespace std { using namespace __jitify_utility_ns; }\n"
"using namespace __jitify_utility_ns;\n";
// TODO: incomplete
static const char* jitsafe_header_vector =
"#pragma once\n"
"namespace __jitify_vector_ns {\n"
"template<class T, class Allocator=void>\n" // = std::allocator> \n"
"struct vector {\n"
"};\n"
"} // namespace __jitify_vector_ns\n"
"namespace std { using namespace __jitify_vector_ns; }\n"
"using namespace __jitify_vector_ns;\n";
// TODO: incomplete
static const char* jitsafe_header_string =
"#pragma once\n"
"namespace __jitify_string_ns {\n"
"template<class CharT,class Traits=void,class Allocator=void>\n"
"struct basic_string {\n"
"basic_string();\n"
"basic_string( const CharT* s );\n" //, const Allocator& alloc =
// Allocator() );\n"
"const CharT* c_str() const;\n"
"bool empty() const;\n"
"void operator+=(const char *);\n"
"void operator+=(const basic_string &);\n"
"};\n"
"typedef basic_string<char> string;\n"
"} // namespace __jitify_string_ns\n"
"namespace std { using namespace __jitify_string_ns; }\n"
"using namespace __jitify_string_ns;\n";
// TODO: incomplete
static const char* jitsafe_header_stdexcept =
"#pragma once\n"
"namespace __jitify_stdexcept_ns {\n"
"struct runtime_error {\n"
"explicit runtime_error( const std::string& what_arg );"
"explicit runtime_error( const char* what_arg );"
"virtual const char* what() const;\n"
"};\n"
"} // namespace __jitify_stdexcept_ns\n"
"namespace std { using namespace __jitify_stdexcept_ns; }\n"
"using namespace __jitify_stdexcept_ns;\n";
// TODO: incomplete
static const char* jitsafe_header_complex =
"#pragma once\n"
"namespace __jitify_complex_ns {\n"
"template<typename T>\n"
"class complex {\n"
" T _real;\n"
" T _imag;\n"
"public:\n"
" complex() : _real(0), _imag(0) {}\n"
" complex(T const& real, T const& imag)\n"
" : _real(real), _imag(imag) {}\n"
" complex(T const& real)\n"
" : _real(real), _imag(static_cast<T>(0)) {}\n"
" T const& real() const { return _real; }\n"
" T& real() { return _real; }\n"
" void real(const T &r) { _real = r; }\n"
" T const& imag() const { return _imag; }\n"
" T& imag() { return _imag; }\n"
" void imag(const T &i) { _imag = i; }\n"
" complex<T>& operator+=(const complex<T> z)\n"
" { _real += z.real(); _imag += z.imag(); return *this; }\n"
"};\n"
"template<typename T>\n"
"complex<T> operator*(const complex<T>& lhs, const complex<T>& rhs)\n"
" { return complex<T>(lhs.real()*rhs.real()-lhs.imag()*rhs.imag(),\n"
" lhs.real()*rhs.imag()+lhs.imag()*rhs.real()); }\n"
"template<typename T>\n"
"complex<T> operator*(const complex<T>& lhs, const T & rhs)\n"
" { return complexs<T>(lhs.real()*rhs,lhs.imag()*rhs); }\n"
"template<typename T>\n"
"complex<T> operator*(const T& lhs, const complex<T>& rhs)\n"
" { return complexs<T>(rhs.real()*lhs,rhs.imag()*lhs); }\n"
"} // namespace __jitify_complex_ns\n"
"namespace std { using namespace __jitify_complex_ns; }\n"
"using namespace __jitify_complex_ns;\n";
// TODO: This is incomplete (missing binary and integer funcs, macros,
// constants, types)
static const char* jitsafe_header_math =
"#pragma once\n"
"namespace __jitify_math_ns {\n"
"#if __cplusplus >= 201103L\n"
"#define DEFINE_MATH_UNARY_FUNC_WRAPPER(f) \\\n"
" inline double f(double x) { return ::f(x); } \\\n"
" inline float f##f(float x) { return ::f(x); } \\\n"
" /*inline long double f##l(long double x) { return ::f(x); }*/ \\\n"
" inline float f(float x) { return ::f(x); } \\\n"
" /*inline long double f(long double x) { return ::f(x); }*/\n"
"#else\n"
"#define DEFINE_MATH_UNARY_FUNC_WRAPPER(f) \\\n"
" inline double f(double x) { return ::f(x); } \\\n"
" inline float f##f(float x) { return ::f(x); } \\\n"
" /*inline long double f##l(long double x) { return ::f(x); }*/\n"
"#endif\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(cos)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(sin)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(tan)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(acos)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(asin)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(atan)\n"
"template<typename T> inline T atan2(T y, T x) { return ::atan2(y, x); }\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(cosh)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(sinh)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(tanh)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(exp)\n"
"template<typename T> inline T frexp(T x, int* exp) { return ::frexp(x, "
"exp); }\n"
"template<typename T> inline T ldexp(T x, int exp) { return ::ldexp(x, "
"exp); }\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(log)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(log10)\n"
"template<typename T> inline T modf(T x, T* intpart) { return ::modf(x, "
"intpart); }\n"
"template<typename T> inline T pow(T x, T y) { return ::pow(x, y); }\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(sqrt)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(ceil)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(floor)\n"
"template<typename T> inline T fmod(T n, T d) { return ::fmod(n, d); }\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(fabs)\n"
"template<typename T> inline T abs(T x) { return ::abs(x); }\n"
"#if __cplusplus >= 201103L\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(acosh)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(asinh)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(atanh)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(exp2)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(expm1)\n"
"template<typename T> inline int ilogb(T x) { return ::ilogb(x); }\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(log1p)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(log2)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(logb)\n"
"template<typename T> inline T scalbn (T x, int n) { return ::scalbn(x, "
"n); }\n"
"template<typename T> inline T scalbln(T x, long n) { return ::scalbn(x, "
"n); }\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(cbrt)\n"
"template<typename T> inline T hypot(T x, T y) { return ::hypot(x, y); }\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(erf)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(erfc)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(tgamma)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(lgamma)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(trunc)\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(round)\n"
"template<typename T> inline long lround(T x) { return ::lround(x); }\n"
"template<typename T> inline long long llround(T x) { return ::llround(x); "
"}\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(rint)\n"
"template<typename T> inline long lrint(T x) { return ::lrint(x); }\n"
"template<typename T> inline long long llrint(T x) { return ::llrint(x); "
"}\n"
"DEFINE_MATH_UNARY_FUNC_WRAPPER(nearbyint)\n"
// TODO: remainder, remquo, copysign, nan, nextafter, nexttoward, fdim,
// fmax, fmin, fma
"#endif\n"
"#undef DEFINE_MATH_UNARY_FUNC_WRAPPER\n"
"} // namespace __jitify_math_ns\n"
"namespace std { using namespace __jitify_math_ns; }\n"
"#define M_PI 3.14159265358979323846\n"
// Note: Global namespace already includes CUDA math funcs
"//using namespace __jitify_math_ns;\n";
static const char* jitsafe_header_memory_h = R"(
#pragma once
#include <string.h>
)";
// TODO: incomplete
static const char* jitsafe_header_mutex = R"(
#pragma once
#if __cplusplus >= 201103L
namespace __jitify_mutex_ns {
class mutex {
public:
void lock();
bool try_lock();
void unlock();
};
} // namespace __jitify_mutex_ns
namespace std { using namespace __jitify_mutex_ns; }
using namespace __jitify_mutex_ns;
#endif
)";
static const char* jitsafe_header_algorithm = R"(
#pragma once
#if __cplusplus >= 201103L
namespace __jitify_algorithm_ns {
#if __cplusplus == 201103L
#define JITIFY_CXX14_CONSTEXPR
#else
#define JITIFY_CXX14_CONSTEXPR constexpr
#endif
template<class T> JITIFY_CXX14_CONSTEXPR const T& max(const T& a, const T& b)
{
return (b > a) ? b : a;
}
template<class T> JITIFY_CXX14_CONSTEXPR const T& min(const T& a, const T& b)
{
return (b < a) ? b : a;
}
#endif
} // namespace __jitify_algorithm_ns
namespace std { using namespace __jitify_algorithm_ns; }
using namespace __jitify_algorithm_ns;
#endif
)";
static const char* jitsafe_header_time_h = R"(
#pragma once
#define NULL 0
#define CLOCKS_PER_SEC 1000000
namespace __jitify_time_ns {
typedef unsigned long size_t;
typedef long clock_t;
typedef long time_t;
struct tm {
int tm_sec;
int tm_min;
int tm_hour;
int tm_mday;
int tm_mon;
int tm_year;
int tm_wday;
int tm_yday;
int tm_isdst;
};
#if __cplusplus >= 201703L
struct timespec {
time_t tv_sec;
long tv_nsec;
};
#endif
} // namespace __jitify_time_ns
namespace std { using namespace __jitify_time_ns; }
using namespace __jitify_time_ns;
)";
// WAR: These need to be pre-included as a workaround for NVRTC implicitly using
// /usr/include as an include path. The other built-in headers will be included
// lazily as needed.
static const char* preinclude_jitsafe_header_names[] = {
"jitify_preinclude.h",
"limits.h",
"math.h",
"memory.h",
"stdint.h",
"stdlib.h",
"stdio.h",
"string.h",
"time.h",
};
template <class T, size_t N>
size_t array_size(T(&)[N]) {
return N;
}
const int preinclude_jitsafe_headers_count =
array_size(preinclude_jitsafe_header_names);
static const std::map<std::string, std::string>& get_jitsafe_headers_map() {
static const std::map<std::string, std::string> jitsafe_headers_map = {
{ "jitify_preinclude.h", jitsafe_header_preinclude_h },
{ "float.h", jitsafe_header_float_h },
{ "cfloat", jitsafe_header_float_h },
{ "limits.h", jitsafe_header_limits_h },
{ "climits", jitsafe_header_limits_h },
{ "stdint.h", jitsafe_header_stdint_h },
{ "cstdint", jitsafe_header_stdint_h },
{ "stddef.h", jitsafe_header_stddef_h },
{ "cstddef", jitsafe_header_stddef_h },
{ "stdlib.h", jitsafe_header_stdlib_h },
{ "cstdlib", jitsafe_header_stdlib_h },
{ "stdio.h", jitsafe_header_stdio_h },
{ "cstdio", jitsafe_header_stdio_h },
{ "string.h", jitsafe_header_string_h },
{ "cstring", jitsafe_header_cstring },
{ "iterator", jitsafe_header_iterator },
{ "limits", jitsafe_header_limits },
{ "type_traits", jitsafe_header_type_traits },
{ "utility", jitsafe_header_utility },
{ "math.h", jitsafe_header_math },
{ "cmath", jitsafe_header_math },
{ "memory.h", jitsafe_header_memory_h },
{ "complex", jitsafe_header_complex },
{ "iostream", jitsafe_header_iostream },
{ "ostream", jitsafe_header_ostream },
{ "istream", jitsafe_header_istream },
{ "sstream", jitsafe_header_sstream },
{ "vector", jitsafe_header_vector },
{ "string", jitsafe_header_string },
{ "stdexcept", jitsafe_header_stdexcept },
{ "mutex", jitsafe_header_mutex },
{ "algorithm", jitsafe_header_algorithm },
{ "time.h", jitsafe_header_time_h },
{ "ctime", jitsafe_header_time_h },
};
return jitsafe_headers_map;
}
inline void add_options_from_env(std::vector<std::string>& options) {
// Add options from environment variable
const char* env_options = std::getenv("JITIFY_OPTIONS");
if (env_options) {
std::stringstream ss;
ss << env_options;
std::string opt;
while (!(ss >> opt).fail()) {
options.push_back(opt);
}
}
// Add options from JITIFY_OPTIONS macro
#ifdef JITIFY_OPTIONS
#define JITIFY_TOSTRING_IMPL(x) #x
#define JITIFY_TOSTRING(x) JITIFY_TOSTRING_IMPL(x)
std::stringstream ss;
ss << JITIFY_TOSTRING(JITIFY_OPTIONS);
std::string opt;
while (!(ss >> opt).fail()) {
options.push_back(opt);
}
#undef JITIFY_TOSTRING
#undef JITIFY_TOSTRING_IMPL
#endif
}
inline void detect_and_add_cuda_arch(std::vector<std::string>& options) {
for (int i = 0; i < (int)options.size(); ++i) {
// Note that this will also match the middle of "--gpu-architecture".
if (options[i].find("-arch") != std::string::npos) {
// Arch already specified in options
return;
}
}
// Use the compute capability of the current device
// TODO: Check these API calls for errors
cudaError_t status;
int device;
status = cudaGetDevice(&device);
if (status != cudaSuccess) {
throw std::runtime_error(
std::string(
"Failed to detect GPU architecture: cudaGetDevice failed: ") +
cudaGetErrorString(status));
}
int cc_major;
cudaDeviceGetAttribute(&cc_major, cudaDevAttrComputeCapabilityMajor, device);
int cc_minor;
cudaDeviceGetAttribute(&cc_minor, cudaDevAttrComputeCapabilityMinor, device);
int cc = cc_major * 10 + cc_minor;
// Note: We must limit the architecture to the max supported by the current
// version of NVRTC, otherwise newer hardware will cause errors
// on older versions of CUDA.
// TODO: It would be better to detect this somehow, rather than hard-coding it
// Tegra chips do not have forwards compatibility so we need to special case
// them.
bool is_tegra = ((cc_major == 3 && cc_minor == 2) || // Logan
(cc_major == 5 && cc_minor == 3) || // Erista
(cc_major == 6 && cc_minor == 2) || // Parker
(cc_major == 7 && cc_minor == 2)); // Xavier
if (!is_tegra) {
// ensure that future CUDA versions just work (even if suboptimal)
const int cuda_major = std::min(10, CUDA_VERSION / 1000);
// clang-format off
switch (cuda_major) {
case 10: cc = std::min(cc, 75); break; // Turing
case 9: cc = std::min(cc, 70); break; // Volta
case 8: cc = std::min(cc, 61); break; // Pascal
case 7: cc = std::min(cc, 52); break; // Maxwell
default:
throw std::runtime_error("Unexpected CUDA major version " +
std::to_string(cuda_major));
}
// clang-format on
}
std::stringstream ss;
ss << cc;
options.push_back("-arch=compute_" + ss.str());
}
inline void detect_and_add_cxx11_flag(std::vector<std::string>& options) {
// Reverse loop so we can erase on the fly.
for (int i = (int)options.size() - 1; i >= 0; --i) {
if (options[i].find("-std=c++98") != std::string::npos) {
// NVRTC doesn't support specifying c++98 explicitly, so we remove it.
options.erase(options.begin() + i);
return;
}
else if (options[i].find("-std") != std::string::npos) {
// Some other standard was explicitly specified, don't change anything.
return;
}
}
// Jitify must be compiled with C++11 support, so we default to enabling it
// for the JIT-compiled code too.
options.push_back("-std=c++11");
}
inline void split_compiler_and_linker_options(
std::vector<std::string> options,
std::vector<std::string>* compiler_options,
std::vector<std::string>* linker_files,
std::vector<std::string>* linker_paths) {
for (int i = 0; i < (int)options.size(); ++i) {
std::string opt = options[i];
std::string flag = opt.substr(0, 2);
std::string value = opt.substr(2);
if (flag == "-l") {
linker_files->push_back(value);
}
else if (flag == "-L") {
linker_paths->push_back(value);
}
else {
compiler_options->push_back(opt);
}
}
}
inline bool pop_remove_unused_globals_flag(std::vector<std::string>* options) {
auto it = std::remove_if(
options->begin(), options->end(), [](const std::string& opt) {
return opt.find("-remove-unused-globals") != std::string::npos;
});
if (it != options->end()) {
options->resize(it - options->begin());
return true;
}
return false;
}
inline std::string ptx_parse_decl_name(const std::string& line) {
size_t name_end = line.find_first_of("[;");
if (name_end == std::string::npos) {
throw std::runtime_error(
"Failed to parse .global declaration in PTX: expected a semicolon");
}
size_t name_start_minus1 = line.find_last_of(" \t", name_end);
if (name_start_minus1 == std::string::npos) {
throw std::runtime_error(
"Failed to parse .global declaration in PTX: expected whitespace");
}
size_t name_start = name_start_minus1 + 1;
std::string name = line.substr(name_start, name_end - name_start);
return name;
}
inline void ptx_remove_unused_globals(std::string* ptx) {
std::istringstream iss(*ptx);
std::vector<std::string> lines;
std::unordered_map<size_t, std::string> line_num_to_global_name;
std::unordered_set<std::string> name_set;
for (std::string line; std::getline(iss, line);) {
size_t line_num = lines.size();
lines.push_back(line);
auto terms = split_string(line);
if (terms.size() <= 1) continue; // Ignore lines with no arguments
if (terms[0].substr(0, 2) == "//") continue; // Ignore comment lines
if (terms[0].substr(0, 7) == ".global") {
line_num_to_global_name.emplace(line_num, ptx_parse_decl_name(line));
continue;
}
if (terms[0][0] == '.') continue; // Ignore .version, .reg, .param etc.
// Note: The first term will always be an instruction name; starting at 1
// also allows unchecked inspection of the previous term.
for (int i = 1; i < (int)terms.size(); ++i) {
if (terms[i].substr(0, 2) == "//") break; // Ignore comments
// Note: The characters '.' and '%' are not treated as delimiters.
const char* token_delims = " \t()[]{},;+-*/~&|^?:=!<>\"'\\";
for (auto token : split_string(terms[i], -1, token_delims)) {
if ( // Ignore non-names
!(std::isalpha(token[0]) || token[0] == '_') ||
token.find('.') != std::string::npos ||
// Ignore variable/parameter declarations
terms[i - 1][0] == '.' ||
// Ignore branch instructions
(token == "bra" && terms[i - 1][0] == '@') ||
// Ignore branch labels
(token.substr(0, 2) == "BB" &&
terms[i - 1].substr(0, 3) == "bra")) {
continue;
}
name_set.insert(token);
}
}
}
std::ostringstream oss;
for (size_t line_num = 0; line_num < lines.size(); ++line_num) {
auto it = line_num_to_global_name.find(line_num);
if (it != line_num_to_global_name.end()) {
const std::string& name = it->second;
if (!name_set.count(name)) {
continue; // Remove unused .global declaration.
}
}
oss << lines[line_num] << '\n';
}
*ptx = oss.str();
}
inline nvrtcResult compile_kernel(std::string program_name,
std::map<std::string, std::string> sources,
std::vector<std::string> options,
std::string instantiation = "",
std::string* log = 0, std::string* ptx = 0,
std::string* mangled_instantiation = 0) {
std::string program_source = sources[program_name];
// Build arrays of header names and sources
std::vector<const char*> header_names_c;
std::vector<const char*> header_sources_c;
int num_headers = sources.size() - 1;
header_names_c.reserve(num_headers);
header_sources_c.reserve(num_headers);
typedef std::map<std::string, std::string> source_map;
for (source_map::const_iterator iter = sources.begin(); iter != sources.end();
++iter) {
std::string const& name = iter->first;
std::string const& code = iter->second;
if (name == program_name) {
continue;
}
header_names_c.push_back(name.c_str());
header_sources_c.push_back(code.c_str());
}
// TODO: This WAR is expected to be unnecessary as of CUDA > 10.2.
bool should_remove_unused_globals =
detail::pop_remove_unused_globals_flag(&options);
std::vector<const char*> options_c(options.size() + 2);
options_c[0] = "--device-as-default-execution-space";
options_c[1] = "--pre-include=jitify_preinclude.h";
for (int i = 0; i < (int)options.size(); ++i) {
options_c[i + 2] = options[i].c_str();
}
#if CUDA_VERSION < 8000
std::string inst_dummy;
if (!instantiation.empty()) {
// WAR for no nvrtcAddNameExpression before CUDA 8.0
// Force template instantiation by adding dummy reference to kernel
inst_dummy = "__jitify_instantiation";
program_source +=
"\nvoid* " + inst_dummy + " = (void*)" + instantiation + ";\n";
}
#endif
#define CHECK_NVRTC(call) \
do { \
nvrtcResult ret = call; \
if (ret != NVRTC_SUCCESS) { \
return ret; \
} \
} while (0)
nvrtcProgram nvrtc_program;
CHECK_NVRTC(nvrtcCreateProgram(
&nvrtc_program, program_source.c_str(), program_name.c_str(), num_headers,
header_sources_c.data(), header_names_c.data()));
#if CUDA_VERSION >= 8000
if (!instantiation.empty()) {
CHECK_NVRTC(nvrtcAddNameExpression(nvrtc_program, instantiation.c_str()));
}
#endif
nvrtcResult ret =
nvrtcCompileProgram(nvrtc_program, options_c.size(), options_c.data());
if (log) {
size_t logsize;
CHECK_NVRTC(nvrtcGetProgramLogSize(nvrtc_program, &logsize));
std::vector<char> vlog(logsize, 0);
CHECK_NVRTC(nvrtcGetProgramLog(nvrtc_program, vlog.data()));
log->assign(vlog.data(), logsize);
if (ret != NVRTC_SUCCESS) {
return ret;
}
}
if (ptx) {
size_t ptxsize;
CHECK_NVRTC(nvrtcGetPTXSize(nvrtc_program, &ptxsize));
std::vector<char> vptx(ptxsize);
CHECK_NVRTC(nvrtcGetPTX(nvrtc_program, vptx.data()));
ptx->assign(vptx.data(), ptxsize);
if (should_remove_unused_globals) {
detail::ptx_remove_unused_globals(ptx);
}
}
if (!instantiation.empty() && mangled_instantiation) {
#if CUDA_VERSION >= 8000
const char* mangled_instantiation_cstr;
// Note: The returned string pointer becomes invalid after
// nvrtcDestroyProgram has been called, so we save it.
CHECK_NVRTC(nvrtcGetLoweredName(nvrtc_program, instantiation.c_str(),
&mangled_instantiation_cstr));
*mangled_instantiation = mangled_instantiation_cstr;
#else
// Extract mangled kernel template instantiation from PTX
inst_dummy += " = "; // Note: This must match how the PTX is generated
int mi_beg = ptx->find(inst_dummy) + inst_dummy.size();
int mi_end = ptx->find(";", mi_beg);
*mangled_instantiation = ptx->substr(mi_beg, mi_end - mi_beg);
#endif
}
CHECK_NVRTC(nvrtcDestroyProgram(&nvrtc_program));
#undef CHECK_NVRTC
return NVRTC_SUCCESS;
}
inline void load_program(std::string const& cuda_source,
std::vector<std::string> const& headers,
file_callback_type file_callback,
std::vector<std::string>* include_paths,
std::map<std::string, std::string>* program_sources,
std::vector<std::string>* program_options,
std::string* program_name) {
// Extract include paths from compile options
std::vector<std::string>::iterator iter = program_options->begin();
while (iter != program_options->end()) {
std::string const& opt = *iter;
if (opt.substr(0, 2) == "-I") {
include_paths->push_back(opt.substr(2));
iter = program_options->erase(iter);
}
else {
++iter;
}
}
// Load program source
if (!detail::load_source(cuda_source, *program_sources, "", *include_paths,
file_callback)) {
throw std::runtime_error("Source not found: " + cuda_source);
}
*program_name = program_sources->begin()->first;
// Maps header include names to their full file paths.
std::map<std::string, std::string> header_fullpaths;
// Load header sources
for (std::string const& header : headers) {
if (!detail::load_source(header, *program_sources, "", *include_paths,
file_callback, &header_fullpaths)) {
// **TODO: Deal with source not found
throw std::runtime_error("Source not found: " + header);
}
}
#if JITIFY_PRINT_SOURCE
std::string& program_source = (*program_sources)[*program_name];
std::cout << "---------------------------------------" << std::endl;
std::cout << "--- Source of " << *program_name << " ---" << std::endl;
std::cout << "---------------------------------------" << std::endl;
detail::print_with_line_numbers(program_source);
std::cout << "---------------------------------------" << std::endl;
#endif
std::vector<std::string> compiler_options, linker_files, linker_paths;
detail::split_compiler_and_linker_options(*program_options, &compiler_options,
&linker_files, &linker_paths);
// If no arch is specified at this point we use whatever the current
// context is. This ensures we pick up the correct internal headers
// for arch-dependent compilation, e.g., some intrinsics are only
// present for specific architectures.
detail::detect_and_add_cuda_arch(compiler_options);
detail::detect_and_add_cxx11_flag(compiler_options);
// Iteratively try to compile the sources, and use the resulting errors to
// identify missing headers.
std::string log;
nvrtcResult ret;
while ((ret = detail::compile_kernel(*program_name, *program_sources,
compiler_options, "", &log)) ==
NVRTC_ERROR_COMPILATION) {
std::string include_name;
std::string include_parent;
int line_num = 0;
if (!detail::extract_include_info_from_compile_error(
log, include_name, include_parent, line_num)) {
#if JITIFY_PRINT_LOG
detail::print_compile_log(*program_name, log);
#endif
// There was a non include-related compilation error
// TODO: How to handle error?
throw std::runtime_error("Runtime compilation failed");
}
bool is_included_with_quotes = false;
if (program_sources->count(include_parent)) {
const std::string& parent_source = (*program_sources)[include_parent];
is_included_with_quotes =
is_include_directive_with_quotes(parent_source, line_num);
}
// Try to load the new header
// Note: This fullpath lookup is needed because the compiler error
// messages have the include name of the header instead of its full path.
std::string include_parent_fullpath = header_fullpaths[include_parent];
std::string include_path = detail::path_base(include_parent_fullpath);
if (detail::load_source(include_name, *program_sources, include_path,
*include_paths, file_callback, &header_fullpaths,
is_included_with_quotes)) {
#if JITIFY_PRINT_HEADER_PATHS
std::cout << "Found #include " << include_name << " from "
<< include_parent << ":" << line_num << " ["
<< include_parent_fullpath << "]"
<< " at:\n " << header_fullpaths[include_name] << std::endl;
#endif
}
else { // Failed to find header file.
// Comment-out the include line and print a warning
if (!program_sources->count(include_parent)) {
// ***TODO: Unless there's another mechanism (e.g., potentially
// the parent path vs. filename problem), getting
// here means include_parent was found automatically
// in a system include path.
// We need a WAR to zap it from *its parent*.
typedef std::map<std::string, std::string> source_map;
for (source_map::const_iterator it = program_sources->begin();
it != program_sources->end(); ++it) {
std::cout << " " << it->first << std::endl;
}
throw std::out_of_range(include_parent +
" not in loaded sources!"
" This may be due to a header being loaded by"
" NVRTC without Jitify's knowledge.");
}
std::string& parent_source = (*program_sources)[include_parent];
parent_source = detail::comment_out_code_line(line_num, parent_source);
#if JITIFY_PRINT_LOG
std::cout << include_parent << "(" << line_num
<< "): warning: " << include_name << ": [jitify] File not found"
<< std::endl;
#endif
}
}
if (ret != NVRTC_SUCCESS) {
#if JITIFY_PRINT_LOG
if (ret == NVRTC_ERROR_INVALID_OPTION) {
std::cout << "Compiler options: ";
for (int i = 0; i < (int)compiler_options.size(); ++i) {
std::cout << compiler_options[i] << " ";
}
std::cout << std::endl;
}
#endif
throw std::runtime_error(std::string("NVRTC error: ") +
nvrtcGetErrorString(ret));
}
}
inline void instantiate_kernel(
std::string const& program_name,
std::map<std::string, std::string> const& program_sources,
std::string const& instantiation, std::vector<std::string> const& options,
std::string* log, std::string* ptx, std::string* mangled_instantiation,
std::vector<std::string>* linker_files,
std::vector<std::string>* linker_paths) {
std::vector<std::string> compiler_options;
detail::split_compiler_and_linker_options(options, &compiler_options,
linker_files, linker_paths);
nvrtcResult ret =
detail::compile_kernel(program_name, program_sources, compiler_options,
instantiation, log, ptx, mangled_instantiation);
#if JITIFY_PRINT_LOG
if (log->size() > 1) {
detail::print_compile_log(program_name, *log);
}
#endif
if (ret != NVRTC_SUCCESS) {
throw std::runtime_error(std::string("NVRTC error: ") +
nvrtcGetErrorString(ret));
}
#if JITIFY_PRINT_PTX
std::cout << "---------------------------------------" << std::endl;
std::cout << *mangled_instantiation << std::endl;
std::cout << "---------------------------------------" << std::endl;
std::cout << "--- PTX for " << mangled_instantiation << " in " << program_name
<< " ---" << std::endl;
std::cout << "---------------------------------------" << std::endl;
std::cout << *ptx << std::endl;
std::cout << "---------------------------------------" << std::endl;
#endif
}
inline void get_1d_max_occupancy(CUfunction func,
CUoccupancyB2DSize smem_callback, size_t* smem,
int max_block_size, unsigned int flags,
int* grid, int* block) {
if (!func) {
throw std::runtime_error(
"Kernel pointer is NULL; you may need to define JITIFY_THREAD_SAFE "
"1");
}
CUresult res = cuOccupancyMaxPotentialBlockSizeWithFlags(
grid, block, func, smem_callback, *smem, max_block_size, flags);
if (res != CUDA_SUCCESS) {
const char* msg;
cuGetErrorName(res, &msg);
throw std::runtime_error(msg);
}
if (smem_callback) {
*smem = smem_callback(*block);
}
}
} // namespace detail
//! \endcond
class KernelInstantiation;
class Kernel;
class Program;
class JitCache;
struct ProgramConfig {
std::vector<std::string> options;
std::vector<std::string> include_paths;
std::string name;
typedef std::map<std::string, std::string> source_map;
source_map sources;
};
class JitCache_impl {
friend class Program_impl;
friend class KernelInstantiation_impl;
friend class KernelLauncher_impl;
typedef uint64_t key_type;
jitify::ObjectCache<key_type, detail::CUDAKernel> _kernel_cache;
jitify::ObjectCache<key_type, ProgramConfig> _program_config_cache;
std::vector<std::string> _options;
#if JITIFY_THREAD_SAFE
std::mutex _kernel_cache_mutex;
std::mutex _program_cache_mutex;
#endif
public:
inline JitCache_impl(size_t cache_size)
: _kernel_cache(cache_size), _program_config_cache(cache_size) {
detail::add_options_from_env(_options);
// Bootstrap the cuda context to avoid errors
cudaFree(0);
}
};
class Program_impl {
// A friendly class
friend class Kernel_impl;
friend class KernelLauncher_impl;
friend class KernelInstantiation_impl;
// TODO: This can become invalid if JitCache is destroyed before the
// Program object is. However, this can't happen if JitCache
// instances are static.
JitCache_impl& _cache;
uint64_t _hash;
ProgramConfig* _config;
void load_sources(std::string source, std::vector<std::string> headers,
std::vector<std::string> options,
file_callback_type file_callback);
public:
inline Program_impl(JitCache_impl& cache, std::string source,
jitify::detail::vector<std::string> headers = 0,
jitify::detail::vector<std::string> options = 0,
file_callback_type file_callback = 0);
#if __cplusplus >= 201103L
inline Program_impl(Program_impl const&) = default;
inline Program_impl(Program_impl&&) = default;
#endif
inline std::vector<std::string> const& options() const {
return _config->options;
}
inline std::string const& name() const { return _config->name; }
inline ProgramConfig::source_map const& sources() const {
return _config->sources;
}
inline std::vector<std::string> const& include_paths() const {
return _config->include_paths;
}
};
class Kernel_impl {
friend class KernelLauncher_impl;
friend class KernelInstantiation_impl;
Program_impl _program;
std::string _name;
std::vector<std::string> _options;
uint64_t _hash;
public:
inline Kernel_impl(Program_impl const& program, std::string name,
jitify::detail::vector<std::string> options = 0);
#if __cplusplus >= 201103L
inline Kernel_impl(Kernel_impl const&) = default;
inline Kernel_impl(Kernel_impl&&) = default;
#endif
};
class KernelInstantiation_impl {
friend class KernelLauncher_impl;
Kernel_impl _kernel;
uint64_t _hash;
std::string _template_inst;
std::vector<std::string> _options;
detail::CUDAKernel* _cuda_kernel;
inline void print() const;
void build_kernel();
public:
inline KernelInstantiation_impl(
Kernel_impl const& kernel, std::vector<std::string> const& template_args);
#if __cplusplus >= 201103L
inline KernelInstantiation_impl(KernelInstantiation_impl const&) = default;
inline KernelInstantiation_impl(KernelInstantiation_impl&&) = default;
#endif
detail::CUDAKernel const& cuda_kernel() const { return *_cuda_kernel; }
};
class KernelLauncher_impl {
KernelInstantiation_impl _kernel_inst;
dim3 _grid;
dim3 _block;
size_t _smem;
cudaStream_t _stream;
public:
inline KernelLauncher_impl(KernelInstantiation_impl const& kernel_inst,
dim3 grid, dim3 block, size_t smem = 0,
cudaStream_t stream = 0)
: _kernel_inst(kernel_inst),
_grid(grid),
_block(block),
_smem(smem),
_stream(stream) {}
#if __cplusplus >= 201103L
inline KernelLauncher_impl(KernelLauncher_impl const&) = default;
inline KernelLauncher_impl(KernelLauncher_impl&&) = default;
#endif
inline CUresult launch(
jitify::detail::vector<void*> arg_ptrs,
jitify::detail::vector<std::string> arg_types = 0) const;
};
/*! An object representing a configured and instantiated kernel ready
* for launching.
*/
class KernelLauncher {
std::unique_ptr<KernelLauncher_impl const> _impl;
public:
inline KernelLauncher(KernelInstantiation const& kernel_inst, dim3 grid,
dim3 block, size_t smem = 0, cudaStream_t stream = 0);
// Note: It's important that there is no implicit conversion required
// for arg_ptrs, because otherwise the parameter pack version
// below gets called instead (probably resulting in a segfault).
/*! Launch the kernel.
*
* \param arg_ptrs A vector of pointers to each function argument for the
* kernel.
* \param arg_types A vector of function argument types represented
* as code-strings. This parameter is optional and is only used to print
* out the function signature.
*/
inline CUresult launch(
std::vector<void*> arg_ptrs = std::vector<void*>(),
jitify::detail::vector<std::string> arg_types = 0) const {
return _impl->launch(arg_ptrs, arg_types);
}
#if __cplusplus >= 201103L
// Regular function call syntax
/*! Launch the kernel.
*
* \see launch
*/
template <typename... ArgTypes>
inline CUresult operator()(ArgTypes... args) const {
return this->launch(args...);
}
/*! Launch the kernel.
*
* \param args Function arguments for the kernel.
*/
template <typename... ArgTypes>
inline CUresult launch(ArgTypes... args) const {
return this->launch(std::vector<void*>({ (void*)&args... }),
{ reflection::reflect<ArgTypes>()... });
}
#endif
};
/*! An object representing a kernel instantiation made up of a Kernel and
* template arguments.
*/
class KernelInstantiation {
friend class KernelLauncher;
std::unique_ptr<KernelInstantiation_impl const> _impl;
public:
inline KernelInstantiation(Kernel const& kernel,
std::vector<std::string> const& template_args);
/*! Implicit conversion to the underlying CUfunction object.
*
* \note This allows use of CUDA APIs like
* cuOccupancyMaxActiveBlocksPerMultiprocessor.
*/
inline operator CUfunction() const { return _impl->cuda_kernel(); }
/*! Configure the kernel launch.
*
* \see configure
*/
inline KernelLauncher operator()(dim3 grid, dim3 block, size_t smem = 0,
cudaStream_t stream = 0) const {
return this->configure(grid, block, smem, stream);
}
/*! Configure the kernel launch.
*
* \param grid The thread grid dimensions for the launch.
* \param block The thread block dimensions for the launch.
* \param smem The amount of shared memory to dynamically allocate, in
* bytes.
* \param stream The CUDA stream to launch the kernel in.
*/
inline KernelLauncher configure(dim3 grid, dim3 block, size_t smem = 0,
cudaStream_t stream = 0) const {
return KernelLauncher(*this, grid, block, smem, stream);
}
/*! Configure the kernel launch with a 1-dimensional block and grid chosen
* automatically to maximise occupancy.
*
* \param max_block_size The upper limit on the block size, or 0 for no
* limit.
* \param smem The amount of shared memory to dynamically allocate, in bytes.
* \param smem_callback A function returning smem for a given block size (overrides \p smem).
* \param stream The CUDA stream to launch the kernel in.
* \param flags The flags to pass to cuOccupancyMaxPotentialBlockSizeWithFlags.
*/
inline KernelLauncher configure_1d_max_occupancy(
int max_block_size = 0, size_t smem = 0,
CUoccupancyB2DSize smem_callback = 0, cudaStream_t stream = 0,
unsigned int flags = 0) const {
int grid;
int block;
CUfunction func = _impl->cuda_kernel();
detail::get_1d_max_occupancy(func, smem_callback, &smem, max_block_size,
flags, &grid, &block);
return this->configure(grid, block, smem, stream);
}
inline CUdeviceptr get_constant_ptr(const char* name) const {
return _impl->cuda_kernel().get_constant_ptr(name);
}
const std::string& mangled_name() const {
return _impl->cuda_kernel().function_name();
}
const std::string& ptx() const { return _impl->cuda_kernel().ptx(); }
const std::vector<std::string>& link_files() const {
return _impl->cuda_kernel().link_files();
}
const std::vector<std::string>& link_paths() const {
return _impl->cuda_kernel().link_paths();
}
};
/*! An object representing a kernel made up of a Program, a name and options.
*/
class Kernel {
friend class KernelInstantiation;
std::unique_ptr<Kernel_impl const> _impl;
public:
Kernel(Program const& program, std::string name,
jitify::detail::vector<std::string> options = 0);
/*! Instantiate the kernel.
*
* \param template_args A vector of template arguments represented as
* code-strings. These can be generated using
* \code{.cpp}jitify::reflection::reflect<type>()\endcode or
* \code{.cpp}jitify::reflection::reflect(value)\endcode
*
* \note Template type deduction is not possible, so all types must be
* explicitly specified.
*/
// inline KernelInstantiation instantiate(std::vector<std::string> const&
// template_args) const {
inline KernelInstantiation instantiate(
std::vector<std::string> const& template_args =
std::vector<std::string>()) const {
return KernelInstantiation(*this, template_args);
}
#if __cplusplus >= 201103L
// Regular template instantiation syntax (note limited flexibility)
/*! Instantiate the kernel.
*
* \note The template arguments specified on this function are
* used to instantiate the kernel. Non-type template arguments must
* be wrapped with
* \code{.cpp}jitify::reflection::NonType<type,value>\endcode
*
* \note Template type deduction is not possible, so all types must be
* explicitly specified.
*/
template <typename... TemplateArgs>
inline KernelInstantiation instantiate() const {
return this->instantiate(
std::vector<std::string>({ reflection::reflect<TemplateArgs>()... }));
}
// Template-like instantiation syntax
// E.g., instantiate(myvar,Type<MyType>())(grid,block)
/*! Instantiate the kernel.
*
* \param targs The template arguments for the kernel, represented as
* values. Types must be wrapped with
* \code{.cpp}jitify::reflection::Type<type>()\endcode or
* \code{.cpp}jitify::reflection::type_of(value)\endcode
*
* \note Template type deduction is not possible, so all types must be
* explicitly specified.
*/
template <typename... TemplateArgs>
inline KernelInstantiation instantiate(TemplateArgs... targs) const {
return this->instantiate(
std::vector<std::string>({ reflection::reflect(targs)... }));
}
#endif
};
/*! An object representing a program made up of source code, headers
* and options.
*/
class Program {
friend class Kernel;
std::unique_ptr<Program_impl const> _impl;
public:
Program(JitCache& cache, std::string source,
jitify::detail::vector<std::string> headers = 0,
jitify::detail::vector<std::string> options = 0,
file_callback_type file_callback = 0);
/*! Select a kernel.
*
* \param name The name of the kernel (unmangled and without
* template arguments).
* \param options A vector of options to be passed to the NVRTC
* compiler when compiling this kernel.
*/
inline Kernel kernel(std::string name,
jitify::detail::vector<std::string> options = 0) const {
return Kernel(*this, name, options);
}
/*! Select a kernel.
*
* \see kernel
*/
inline Kernel operator()(
std::string name, jitify::detail::vector<std::string> options = 0) const {
return this->kernel(name, options);
}
};
/*! An object that manages a cache of JIT-compiled CUDA kernels.
*
*/
class JitCache {
friend class Program;
std::unique_ptr<JitCache_impl> _impl;
public:
/*! JitCache constructor.
* \param cache_size The number of kernels to hold in the cache
* before overwriting the least-recently-used ones.
*/
enum { DEFAULT_CACHE_SIZE = 128 };
JitCache(size_t cache_size = DEFAULT_CACHE_SIZE)
: _impl(new JitCache_impl(cache_size)) {}
/*! Create a program.
*
* \param source A string containing either the source filename or
* the source itself; in the latter case, the first line must be
* the name of the program.
* \param headers A vector of strings representing the source of
* each header file required by the program. Each entry can be
* either the header filename or the header source itself; in
* the latter case, the first line must be the name of the header
* (i.e., the name by which the header is #included).
* \param options A vector of options to be passed to the
* NVRTC compiler. Include paths specified with \p -I
* are added to the search paths used by Jitify. The environment
* variable JITIFY_OPTIONS can also be used to define additional
* options.
* \param file_callback A pointer to a callback function that is
* invoked whenever a source file needs to be loaded. Inside this
* function, the user can either load/specify the source themselves
* or defer to Jitify's file-loading mechanisms.
* \note Program or header source files referenced by filename are
* looked-up using the following mechanisms (in this order):
* \note 1) By calling file_callback.
* \note 2) By looking for the file embedded in the executable via the GCC
* linker.
* \note 3) By looking for the file in the filesystem.
*
* \note Jitify recursively scans all source files for \p #include
* directives and automatically adds them to the set of headers needed
* by the program.
* If a \p #include directive references a header that cannot be found,
* the directive is automatically removed from the source code to prevent
* immediate compilation failure. This may result in compilation errors
* if the header was required by the program.
*
* \note Jitify automatically includes NVRTC-safe versions of some
* standard library headers.
*/
inline Program program(std::string source,
jitify::detail::vector<std::string> headers = 0,
jitify::detail::vector<std::string> options = 0,
file_callback_type file_callback = 0) {
return Program(*this, source, headers, options, file_callback);
}
};
inline Program::Program(JitCache& cache, std::string source,
jitify::detail::vector<std::string> headers,
jitify::detail::vector<std::string> options,
file_callback_type file_callback)
: _impl(new Program_impl(*cache._impl, source, headers, options,
file_callback)) {}
inline Kernel::Kernel(Program const& program, std::string name,
jitify::detail::vector<std::string> options)
: _impl(new Kernel_impl(*program._impl, name, options)) {}
inline KernelInstantiation::KernelInstantiation(
Kernel const& kernel, std::vector<std::string> const& template_args)
: _impl(new KernelInstantiation_impl(*kernel._impl, template_args)) {}
inline KernelLauncher::KernelLauncher(KernelInstantiation const& kernel_inst,
dim3 grid, dim3 block, size_t smem,
cudaStream_t stream)
: _impl(new KernelLauncher_impl(*kernel_inst._impl, grid, block, smem,
stream)) {}
inline std::ostream& operator<<(std::ostream& stream, dim3 d) {
if (d.y == 1 && d.z == 1) {
stream << d.x;
}
else {
stream << "(" << d.x << "," << d.y << "," << d.z << ")";
}
return stream;
}
inline CUresult KernelLauncher_impl::launch(
jitify::detail::vector<void*> arg_ptrs,
jitify::detail::vector<std::string> arg_types) const {
#if JITIFY_PRINT_LAUNCH
Kernel_impl const& kernel = _kernel_inst._kernel;
std::string arg_types_string =
(arg_types.empty() ? "..." : reflection::reflect_list(arg_types));
std::cout << "Launching " << kernel._name << _kernel_inst._template_inst
<< "<<<" << _grid << "," << _block << "," << _smem << "," << _stream
<< ">>>"
<< "(" << arg_types_string << ")" << std::endl;
#endif
if (!_kernel_inst._cuda_kernel) {
throw std::runtime_error(
"Kernel pointer is NULL; you may need to define JITIFY_THREAD_SAFE 1");
}
return _kernel_inst._cuda_kernel->launch(_grid, _block, _smem, _stream,
arg_ptrs);
}
inline KernelInstantiation_impl::KernelInstantiation_impl(
Kernel_impl const& kernel, std::vector<std::string> const& template_args)
: _kernel(kernel), _options(kernel._options) {
_template_inst =
(template_args.empty() ? ""
: reflection::reflect_template(template_args));
using detail::hash_combine;
using detail::hash_larson64;
_hash = _kernel._hash;
_hash = hash_combine(_hash, hash_larson64(_template_inst.c_str()));
JitCache_impl& cache = _kernel._program._cache;
uint64_t cache_key = _hash;
#if JITIFY_THREAD_SAFE
std::lock_guard<std::mutex> lock(cache._kernel_cache_mutex);
#endif
if (cache._kernel_cache.contains(cache_key)) {
#if JITIFY_PRINT_INSTANTIATION
std::cout << "Found ";
this->print();
#endif
_cuda_kernel = &cache._kernel_cache.get(cache_key);
}
else {
#if JITIFY_PRINT_INSTANTIATION
std::cout << "Building ";
this->print();
#endif
_cuda_kernel = &cache._kernel_cache.emplace(cache_key);
this->build_kernel();
}
}
inline void KernelInstantiation_impl::print() const {
std::string options_string = reflection::reflect_list(_options);
std::cout << _kernel._name << _template_inst << " [" << options_string << "]"
<< std::endl;
}
inline void KernelInstantiation_impl::build_kernel() {
Program_impl const& program = _kernel._program;
std::string instantiation = _kernel._name + _template_inst;
std::string log, ptx, mangled_instantiation;
std::vector<std::string> linker_files, linker_paths;
detail::instantiate_kernel(program.name(), program.sources(), instantiation,
_options, &log, &ptx, &mangled_instantiation,
&linker_files, &linker_paths);
_cuda_kernel->set(mangled_instantiation.c_str(), ptx.c_str(), linker_files,
linker_paths);
}
Kernel_impl::Kernel_impl(Program_impl const& program, std::string name,
jitify::detail::vector<std::string> options)
: _program(program), _name(name), _options(options) {
// Merge options from parent
_options.insert(_options.end(), _program.options().begin(),
_program.options().end());
detail::detect_and_add_cuda_arch(_options);
detail::detect_and_add_cxx11_flag(_options);
std::string options_string = reflection::reflect_list(_options);
using detail::hash_combine;
using detail::hash_larson64;
_hash = _program._hash;
_hash = hash_combine(_hash, hash_larson64(_name.c_str()));
_hash = hash_combine(_hash, hash_larson64(options_string.c_str()));
}
Program_impl::Program_impl(JitCache_impl& cache, std::string source,
jitify::detail::vector<std::string> headers,
jitify::detail::vector<std::string> options,
file_callback_type file_callback)
: _cache(cache) {
// Compute hash of source, headers and options
std::string options_string = reflection::reflect_list(options);
using detail::hash_combine;
using detail::hash_larson64;
_hash = hash_combine(hash_larson64(source.c_str()),
hash_larson64(options_string.c_str()));
for (size_t i = 0; i < headers.size(); ++i) {
_hash = hash_combine(_hash, hash_larson64(headers[i].c_str()));
}
_hash = hash_combine(_hash, (uint64_t)file_callback);
// Add pre-include built-in JIT-safe headers
for (int i = 0; i < detail::preinclude_jitsafe_headers_count; ++i) {
const char* hdr_name = detail::preinclude_jitsafe_header_names[i];
const std::string& hdr_source =
detail::get_jitsafe_headers_map().at(hdr_name);
headers.push_back(std::string(hdr_name) + "\n" + hdr_source);
}
// Merge options from parent
options.insert(options.end(), _cache._options.begin(), _cache._options.end());
// Load sources
#if JITIFY_THREAD_SAFE
std::lock_guard<std::mutex> lock(cache._program_cache_mutex);
#endif
if (!cache._program_config_cache.contains(_hash)) {
_config = &cache._program_config_cache.insert(_hash);
this->load_sources(source, headers, options, file_callback);
}
else {
_config = &cache._program_config_cache.get(_hash);
}
}
inline void Program_impl::load_sources(std::string source,
std::vector<std::string> headers,
std::vector<std::string> options,
file_callback_type file_callback) {
_config->options = options;
detail::load_program(source, headers, file_callback, &_config->include_paths,
&_config->sources, &_config->options, &_config->name);
}
#if __cplusplus >= 201103L
enum Location { HOST, DEVICE };
/*! Specifies location and parameters for execution of an algorithm.
* \param stream The CUDA stream on which to execute.
* \param headers A vector of headers to include in the code.
* \param options Options to pass to the NVRTC compiler.
* \param file_callback See jitify::Program.
* \param block_size The size of the CUDA thread block with which to
* execute.
* \param cache_size The number of kernels to store in the cache
* before overwriting the least-recently-used ones.
*/
struct ExecutionPolicy {
/*! Location (HOST or DEVICE) on which to execute.*/
Location location;
/*! List of headers to include when compiling the algorithm.*/
std::vector<std::string> headers;
/*! List of compiler options.*/
std::vector<std::string> options;
/*! Optional callback for loading source files.*/
file_callback_type file_callback;
/*! CUDA stream on which to execute.*/
cudaStream_t stream;
/*! CUDA device on which to execute.*/
int device;
/*! CUDA block size with which to execute.*/
int block_size;
/*! The number of instantiations to store in the cache before overwriting
* the least-recently-used ones.*/
size_t cache_size;
ExecutionPolicy(Location location_ = DEVICE,
jitify::detail::vector<std::string> headers_ = 0,
jitify::detail::vector<std::string> options_ = 0,
file_callback_type file_callback_ = 0,
cudaStream_t stream_ = 0, int device_ = 0,
int block_size_ = 256,
size_t cache_size_ = JitCache::DEFAULT_CACHE_SIZE)
: location(location_),
headers(headers_),
options(options_),
file_callback(file_callback_),
stream(stream_),
device(device_),
block_size(block_size_),
cache_size(cache_size_) {}
};
template <class Func>
class Lambda;
/*! An object that captures a set of variables for use in a parallel_for
* expression. See JITIFY_CAPTURE().
*/
class Capture {
public:
std::vector<std::string> _arg_decls;
std::vector<void*> _arg_ptrs;
public:
template <typename... Args>
inline Capture(std::vector<std::string> arg_names, Args const&... args)
: _arg_ptrs{ (void*)&args... } {
std::vector<std::string> arg_types = { reflection::reflect<Args>()... };
_arg_decls.resize(arg_names.size());
for (int i = 0; i < (int)arg_names.size(); ++i) {
_arg_decls[i] = arg_types[i] + " " + arg_names[i];
}
}
};
/*! An object that captures the instantiated Lambda function for use
in a parallel_for expression and the function string for NVRTC
compilation
*/
template <class Func>
class Lambda {
public:
Capture _capture;
std::string _func_string;
Func _func;
public:
inline Lambda(Capture const& capture, std::string func_string, Func func)
: _capture(capture), _func_string(func_string), _func(func) {}
};
template <typename T>
inline Lambda<T> make_Lambda(Capture const& capture, std::string func,
T lambda) {
return Lambda<T>(capture, func, lambda);
}
#define JITIFY_CAPTURE(...) \
jitify::Capture(jitify::detail::split_string(#__VA_ARGS__, -1, ","), \
__VA_ARGS__)
#define JITIFY_MAKE_LAMBDA(capture, x, ...) \
jitify::make_Lambda(capture, std::string(#__VA_ARGS__), \
[x](int i) { __VA_ARGS__; })
#define JITIFY_ARGS(...) __VA_ARGS__
#define JITIFY_LAMBDA_(x, ...) \
JITIFY_MAKE_LAMBDA(JITIFY_CAPTURE(x), JITIFY_ARGS(x), __VA_ARGS__)
// macro sequence to strip surrounding brackets
#define JITIFY_STRIP_PARENS(X) X
#define JITIFY_PASS_PARAMETERS(X) JITIFY_STRIP_PARENS(JITIFY_ARGS X)
/*! Creates a Lambda object with captured variables and a function
* definition.
* \param capture A bracket-enclosed list of variables to capture.
* \param ... The function definition.
*
* \code{.cpp}
* float* capture_me;
* int capture_me_too;
* auto my_lambda = JITIFY_LAMBDA( (capture_me, capture_me_too),
* capture_me[i] = i*capture_me_too );
* \endcode
*/
#define JITIFY_LAMBDA(capture, ...) \
JITIFY_LAMBDA_(JITIFY_ARGS(JITIFY_PASS_PARAMETERS(capture)), \
JITIFY_ARGS(__VA_ARGS__))
// TODO: Try to implement for_each that accepts iterators instead of indices
// Add compile guard for NOCUDA compilation
/*! Call a function for a range of indices
*
* \param policy Determines the location and device parameters for
* execution of the parallel_for.
* \param begin The starting index.
* \param end The ending index.
* \param lambda A Lambda object created using the JITIFY_LAMBDA() macro.
*
* \code{.cpp}
* char const* in;
* float* out;
* parallel_for(0, 100, JITIFY_LAMBDA( (in, out), {char x = in[i]; out[i] =
* x*x; } ); \endcode
*/
template <typename IndexType, class Func>
CUresult parallel_for(ExecutionPolicy policy, IndexType begin, IndexType end,
Lambda<Func> const& lambda) {
using namespace jitify;
if (policy.location == HOST) {
#ifdef _OPENMP
#pragma omp parallel for
#endif
for (IndexType i = begin; i < end; i++) {
lambda._func(i);
}
return CUDA_SUCCESS; // FIXME - replace with non-CUDA enum type?
}
thread_local static JitCache kernel_cache(policy.cache_size);
std::vector<std::string> arg_decls;
arg_decls.push_back("I begin, I end");
arg_decls.insert(arg_decls.end(), lambda._capture._arg_decls.begin(),
lambda._capture._arg_decls.end());
std::stringstream source_ss;
source_ss << "parallel_for_program\n";
for (auto const& header : policy.headers) {
std::string header_name = header.substr(0, header.find("\n"));
source_ss << "#include <" << header_name << ">\n";
}
source_ss << "template<typename I>\n"
"__global__\n"
"void parallel_for_kernel("
<< reflection::reflect_list(arg_decls)
<< ") {\n"
" I i0 = threadIdx.x + blockDim.x*blockIdx.x;\n"
" for( I i=i0+begin; i<end; i+=blockDim.x*gridDim.x ) {\n"
" "
<< "\t" << lambda._func_string << ";\n"
<< " }\n"
"}\n";
Program program = kernel_cache.program(source_ss.str(), policy.headers,
policy.options, policy.file_callback);
std::vector<void*> arg_ptrs;
arg_ptrs.push_back(&begin);
arg_ptrs.push_back(&end);
arg_ptrs.insert(arg_ptrs.end(), lambda._capture._arg_ptrs.begin(),
lambda._capture._arg_ptrs.end());
size_t n = end - begin;
dim3 block(policy.block_size);
dim3 grid(std::min((n - 1) / block.x + 1, size_t(65535)));
cudaSetDevice(policy.device);
return program.kernel("parallel_for_kernel")
.instantiate<IndexType>()
.configure(grid, block, 0, policy.stream)
.launch(arg_ptrs);
}
#endif // __cplusplus >= 201103L
namespace experimental {
using jitify::file_callback_type;
namespace serialization {
namespace detail {
// This should be incremented whenever the serialization format changes in any
// incompatible way.
static constexpr const size_t kSerializationVersion = 1;
inline void serialize(std::ostream& stream, size_t u) {
uint64_t u64 = u;
stream.write(reinterpret_cast<char*>(&u64), sizeof(u64));
}
inline bool deserialize(std::istream& stream, size_t* size) {
uint64_t u64;
stream.read(reinterpret_cast<char*>(&u64), sizeof(u64));
*size = u64;
return stream.good();
}
inline void serialize(std::ostream& stream, std::string const& s) {
serialize(stream, s.size());
stream.write(s.data(), s.size());
}
inline bool deserialize(std::istream& stream, std::string* s) {
size_t size;
if (!deserialize(stream, &size)) return false;
s->resize(size);
if (s->size()) {
stream.read(&(*s)[0], s->size());
}
return stream.good();
}
inline void serialize(std::ostream& stream, std::vector<std::string> const& v) {
serialize(stream, v.size());
for (auto const& s : v) {
serialize(stream, s);
}
}
inline bool deserialize(std::istream& stream, std::vector<std::string>* v) {
size_t size;
if (!deserialize(stream, &size)) return false;
v->resize(size);
for (auto& s : *v) {
if (!deserialize(stream, &s)) return false;
}
return true;
}
inline void serialize(std::ostream& stream,
std::map<std::string, std::string> const& m) {
serialize(stream, m.size());
for (auto const& kv : m) {
serialize(stream, kv.first);
serialize(stream, kv.second);
}
}
inline bool deserialize(std::istream& stream,
std::map<std::string, std::string>* m) {
size_t size;
if (!deserialize(stream, &size)) return false;
for (size_t i = 0; i < size; ++i) {
std::string key;
if (!deserialize(stream, &key)) return false;
if (!deserialize(stream, &(*m)[key])) return false;
}
return true;
}
template <typename T, typename... Rest>
inline void serialize(std::ostream& stream, T const& value, Rest... rest) {
serialize(stream, value);
serialize(stream, rest...);
}
template <typename T, typename... Rest>
inline bool deserialize(std::istream& stream, T* value, Rest... rest) {
if (!deserialize(stream, value)) return false;
return deserialize(stream, rest...);
}
inline void serialize_magic_number(std::ostream& stream) {
stream.write("JTFY", 4);
serialize(stream, kSerializationVersion);
}
inline bool deserialize_magic_number(std::istream& stream) {
char magic_number[4] = { 0, 0, 0, 0 };
stream.read(&magic_number[0], 4);
if (!(magic_number[0] == 'J' && magic_number[1] == 'T' &&
magic_number[2] == 'F' && magic_number[3] == 'Y')) {
return false;
}
size_t serialization_version;
if (!deserialize(stream, &serialization_version)) return false;
return serialization_version == kSerializationVersion;
}
} // namespace detail
template <typename... Values>
inline std::string serialize(Values const&... values) {
std::ostringstream ss(std::stringstream::out | std::stringstream::binary);
detail::serialize_magic_number(ss);
detail::serialize(ss, values...);
return ss.str();
}
template <typename... Values>
inline bool deserialize(std::string const& serialized, Values*... values) {
std::istringstream ss(serialized,
std::stringstream::in | std::stringstream::binary);
if (!detail::deserialize_magic_number(ss)) return false;
return detail::deserialize(ss, values...);
}
} // namespace serialization
class Program;
class Kernel;
class KernelInstantiation;
class KernelLauncher;
/*! An object representing a program made up of source code, headers
* and options.
*/
class Program {
private:
friend class KernelInstantiation;
std::string _name;
std::vector<std::string> _options;
std::map<std::string, std::string> _sources;
// Private constructor used by deserialize()
Program() {}
public:
/*! Create a program.
*
* \param source A string containing either the source filename or
* the source itself; in the latter case, the first line must be
* the name of the program.
* \param headers A vector of strings representing the source of
* each header file required by the program. Each entry can be
* either the header filename or the header source itself; in
* the latter case, the first line must be the name of the header
* (i.e., the name by which the header is #included).
* \param options A vector of options to be passed to the
* NVRTC compiler. Include paths specified with \p -I
* are added to the search paths used by Jitify. The environment
* variable JITIFY_OPTIONS can also be used to define additional
* options.
* \param file_callback A pointer to a callback function that is
* invoked whenever a source file needs to be loaded. Inside this
* function, the user can either load/specify the source themselves
* or defer to Jitify's file-loading mechanisms.
* \note Program or header source files referenced by filename are
* looked-up using the following mechanisms (in this order):
* \note 1) By calling file_callback.
* \note 2) By looking for the file embedded in the executable via the GCC
* linker.
* \note 3) By looking for the file in the filesystem.
*
* \note Jitify recursively scans all source files for \p #include
* directives and automatically adds them to the set of headers needed
* by the program.
* If a \p #include directive references a header that cannot be found,
* the directive is automatically removed from the source code to prevent
* immediate compilation failure. This may result in compilation errors
* if the header was required by the program.
*
* \note Jitify automatically includes NVRTC-safe versions of some
* standard library headers.
*/
Program(std::string const& cuda_source,
std::vector<std::string> const& given_headers = {},
std::vector<std::string> const& given_options = {},
file_callback_type file_callback = nullptr) {
// Add pre-include built-in JIT-safe headers
std::vector<std::string> headers = given_headers;
for (int i = 0; i < detail::preinclude_jitsafe_headers_count; ++i) {
const char* hdr_name = detail::preinclude_jitsafe_header_names[i];
const std::string& hdr_source =
detail::get_jitsafe_headers_map().at(hdr_name);
headers.push_back(std::string(hdr_name) + "\n" + hdr_source);
}
_options = given_options;
detail::add_options_from_env(_options);
std::vector<std::string> include_paths;
detail::load_program(cuda_source, headers, file_callback, &include_paths,
&_sources, &_options, &_name);
}
/*! Restore a serialized program.
*
* \param serialized_program The serialized program to restore.
*
* \see serialize
*/
static Program deserialize(std::string const& serialized_program) {
Program program;
if (!serialization::deserialize(serialized_program, &program._name,
&program._options, &program._sources)) {
throw std::runtime_error("Failed to deserialize program");
}
return program;
}
/*! Save the program.
*
* \see deserialize
*/
std::string serialize() const {
// Note: Must update kSerializationVersion if this is changed.
return serialization::serialize(_name, _options, _sources);
};
/*! Select a kernel.
*
* \param name The name of the kernel (unmangled and without
* template arguments).
* \param options A vector of options to be passed to the NVRTC
* compiler when compiling this kernel.
*/
Kernel kernel(std::string const& name,
std::vector<std::string> const& options = {}) const;
};
class Kernel {
friend class KernelInstantiation;
Program const* _program;
std::string _name;
std::vector<std::string> _options;
public:
Kernel(Program const* program, std::string const& name,
std::vector<std::string> const& options = {})
: _program(program), _name(name), _options(options) {}
/*! Instantiate the kernel.
*
* \param template_args A vector of template arguments represented as
* code-strings. These can be generated using
* \code{.cpp}jitify::reflection::reflect<type>()\endcode or
* \code{.cpp}jitify::reflection::reflect(value)\endcode
*
* \note Template type deduction is not possible, so all types must be
* explicitly specified.
*/
KernelInstantiation instantiate(
std::vector<std::string> const& template_args =
std::vector<std::string>()) const;
KernelInstantiation instantiate2(unsigned int nopts = 0,
CUjit_option* opts = 0, void** optvals = 0) const;
// Regular template instantiation syntax (note limited flexibility)
/*! Instantiate the kernel.
*
* \note The template arguments specified on this function are
* used to instantiate the kernel. Non-type template arguments must
* be wrapped with
* \code{.cpp}jitify::reflection::NonType<type,value>\endcode
*
* \note Template type deduction is not possible, so all types must be
* explicitly specified.
*/
template <typename... TemplateArgs>
KernelInstantiation instantiate() const;
// Template-like instantiation syntax
// E.g., instantiate(myvar,Type<MyType>())(grid,block)
/*! Instantiate the kernel.
*
* \param targs The template arguments for the kernel, represented as
* values. Types must be wrapped with
* \code{.cpp}jitify::reflection::Type<type>()\endcode or
* \code{.cpp}jitify::reflection::type_of(value)\endcode
*
* \note Template type deduction is not possible, so all types must be
* explicitly specified.
*/
template <typename... TemplateArgs>
KernelInstantiation instantiate(TemplateArgs... targs) const;
};
class KernelInstantiation {
friend class KernelLauncher;
std::unique_ptr<detail::CUDAKernel> _cuda_kernel;
// Private constructor used by deserialize()
KernelInstantiation(std::string const& func_name, std::string const& ptx,
std::vector<std::string> const& link_files,
std::vector<std::string> const& link_paths)
: _cuda_kernel(new detail::CUDAKernel(func_name.c_str(), ptx.c_str(),
link_files, link_paths)) {}
public:
KernelInstantiation(Kernel const& kernel,
std::vector<std::string> const& template_args, unsigned int nopts = 0,
CUjit_option* opts = 0, void** optvals = 0) {
Program const* program = kernel._program;
std::string template_inst =
(template_args.empty() ? ""
: reflection::reflect_template(template_args));
std::string instantiation = kernel._name + template_inst;
std::vector<std::string> options;
options.insert(options.begin(), program->_options.begin(),
program->_options.end());
options.insert(options.begin(), kernel._options.begin(),
kernel._options.end());
detail::detect_and_add_cuda_arch(options);
detail::detect_and_add_cxx11_flag(options);
std::string log, ptx, mangled_instantiation;
std::vector<std::string> linker_files, linker_paths;
detail::instantiate_kernel(program->_name, program->_sources, instantiation,
options, &log, &ptx, &mangled_instantiation,
&linker_files, &linker_paths);
_cuda_kernel.reset(new detail::CUDAKernel(mangled_instantiation.c_str(),
ptx.c_str(), linker_files,
linker_paths, nopts, opts, optvals));
}
/*! Implicit conversion to the underlying CUfunction object.
*
* \note This allows use of CUDA APIs like
* cuOccupancyMaxActiveBlocksPerMultiprocessor.
*/
operator CUfunction() const { return *_cuda_kernel; }
/*! Restore a serialized kernel instantiation.
*
* \param serialized_kernel_inst The serialized kernel instantiation to
* restore.
*
* \see serialize
*/
static KernelInstantiation deserialize(
std::string const& serialized_kernel_inst) {
std::string func_name, ptx;
std::vector<std::string> link_files, link_paths;
if (!serialization::deserialize(serialized_kernel_inst, &func_name, &ptx,
&link_files, &link_paths)) {
throw std::runtime_error("Failed to deserialize kernel instantiation");
}
return KernelInstantiation(func_name, ptx, link_files, link_paths);
}
/*! Save the program.
*
* \see deserialize
*/
std::string serialize() const {
// Note: Must update kSerializationVersion if this is changed.
return serialization::serialize(
_cuda_kernel->function_name(), _cuda_kernel->ptx(),
_cuda_kernel->link_files(), _cuda_kernel->link_paths());
}
/*! Configure the kernel launch.
*
* \param grid The thread grid dimensions for the launch.
* \param block The thread block dimensions for the launch.
* \param smem The amount of shared memory to dynamically allocate, in
* bytes.
* \param stream The CUDA stream to launch the kernel in.
*/
KernelLauncher configure(dim3 grid, dim3 block, size_t smem = 0,
cudaStream_t stream = 0) const;
/*! Configure the kernel launch with a 1-dimensional block and grid chosen
* automatically to maximise occupancy.
*
* \param max_block_size The upper limit on the block size, or 0 for no
* limit.
* \param smem The amount of shared memory to dynamically allocate, in bytes.
* \param smem_callback A function returning smem for a given block size
* (overrides \p smem).
* \param stream The CUDA stream to launch the kernel in.
* \param flags The flags to pass to
* cuOccupancyMaxPotentialBlockSizeWithFlags.
*/
KernelLauncher configure_1d_max_occupancy(
int max_block_size = 0, size_t smem = 0,
CUoccupancyB2DSize smem_callback = 0, cudaStream_t stream = 0,
unsigned int flags = 0) const;
CUdeviceptr get_constant_ptr(const char* name) const {
return _cuda_kernel->get_constant_ptr(name);
}
const std::string& mangled_name() const {
return _cuda_kernel->function_name();
}
const std::string& ptx() const { return _cuda_kernel->ptx(); }
const std::vector<std::string>& link_files() const {
return _cuda_kernel->link_files();
}
const std::vector<std::string>& link_paths() const {
return _cuda_kernel->link_paths();
}
};
class KernelLauncher {
KernelInstantiation const* _kernel_inst;
dim3 _grid;
dim3 _block;
size_t _smem;
cudaStream_t _stream;
public:
KernelLauncher(KernelInstantiation const* kernel_inst, dim3 grid, dim3 block,
size_t smem = 0, cudaStream_t stream = 0)
: _kernel_inst(kernel_inst),
_grid(grid),
_block(block),
_smem(smem),
_stream(stream) {}
// Note: It's important that there is no implicit conversion required
// for arg_ptrs, because otherwise the parameter pack version
// below gets called instead (probably resulting in a segfault).
/*! Launch the kernel.
*
* \param arg_ptrs A vector of pointers to each function argument for the
* kernel.
* \param arg_types A vector of function argument types represented
* as code-strings. This parameter is optional and is only used to print
* out the function signature.
*/
CUresult launch(std::vector<void*> arg_ptrs = {},
std::vector<std::string> arg_types = {}) const {
#if JITIFY_PRINT_LAUNCH
std::string arg_types_string =
(arg_types.empty() ? "..." : reflection::reflect_list(arg_types));
std::cout << "Launching " << _kernel_inst->_cuda_kernel->function_name()
<< "<<<" << _grid << "," << _block << "," << _smem << ","
<< _stream << ">>>"
<< "(" << arg_types_string << ")" << std::endl;
#endif
return _kernel_inst->_cuda_kernel->launch(_grid, _block, _smem, _stream,
arg_ptrs);
}
/*! Launch the kernel.
*
* \param args Function arguments for the kernel.
*/
template <typename... ArgTypes>
CUresult launch(ArgTypes... args) const {
return this->launch(std::vector<void*>({ (void*)&args... }),
{ reflection::reflect<ArgTypes>()... });
}
};
inline Kernel Program::kernel(std::string const& name,
std::vector<std::string> const& options) const {
return Kernel(this, name, options);
}
inline KernelInstantiation Kernel::instantiate2(unsigned int nopts,
CUjit_option* opts, void** optvals) const {
return KernelInstantiation(*this, {}, nopts, opts, optvals);
}
inline KernelInstantiation Kernel::instantiate(
std::vector<std::string> const& template_args) const {
return KernelInstantiation(*this, template_args);
}
template <typename... TemplateArgs>
inline KernelInstantiation Kernel::instantiate() const {
return this->instantiate(
std::vector<std::string>({ reflection::reflect<TemplateArgs>()... }));
}
template <typename... TemplateArgs>
inline KernelInstantiation Kernel::instantiate(TemplateArgs... targs) const {
return this->instantiate(
std::vector<std::string>({ reflection::reflect(targs)... }));
}
inline KernelLauncher KernelInstantiation::configure(
dim3 grid, dim3 block, size_t smem, cudaStream_t stream) const {
return KernelLauncher(this, grid, block, smem, stream);
}
inline KernelLauncher KernelInstantiation::configure_1d_max_occupancy(
int max_block_size, size_t smem, CUoccupancyB2DSize smem_callback,
cudaStream_t stream, unsigned int flags) const {
int grid;
int block;
CUfunction func = *_cuda_kernel;
detail::get_1d_max_occupancy(func, smem_callback, &smem, max_block_size,
flags, &grid, &block);
return this->configure(grid, block, smem, stream);
}
} // namespace experimental
} // namespace jitify
#if defined(_WIN32) || defined(_WIN64)
#pragma pop_macro("strtok_r")
#endif
#ifdef _MSVC_LANG
#pragma pop_macro("__cplusplus")
#endif
Raw
kernel.cu
/**
* Build this with -rdc=true
* You will need to link against : nvrtc, cuda, cudart_static, dl
*/
#include <string>
#include <vector>
#include <memory>
#include <cassert>
#include <cuda_runtime.h>
#include <nvrtc.h>
#include <cctype> // Required to build Jitify header under VS2015
#include "jitify/jitify.hpp"
__device__ int TEST_DEVICE_SYMBOL;
__constant__ int TEST_CONSTANT_SYMBOL;
int main() {
const char* rtc_kernel = R"###(my_program
// Extern as these symbols are defined elsewhere
extern __device__ int TEST_DEVICE_SYMBOL;
extern __constant__ int TEST_CONSTANT_SYMBOL;
__global__ void simple_test(const int * const TEST_DEVICE_POINTER) {
printf("TEST_DEVICE_POINTER: %s\n", (*TEST_DEVICE_POINTER == 12) ? "Success" : "Failure");
printf("TEST_DEVICE_SYMBOL: %s\n", (TEST_DEVICE_SYMBOL == 12) ? "Success" : "Failure");
printf("TEST_CONSTANT_SYMBOL: %s\n", (TEST_CONSTANT_SYMBOL == 12) ? "Success" : "Failure");
})###";
// Set TEST_SYMBOL
const int t_TEST_SYMBOL = 12;
assert(cudaMemcpyToSymbol(TEST_DEVICE_SYMBOL, &t_TEST_SYMBOL, sizeof(int)) == CUDA_SUCCESS);
assert(cudaMemcpyToSymbol(TEST_CONSTANT_SYMBOL, &t_TEST_SYMBOL, sizeof(int)) == CUDA_SUCCESS);
int *d_TEST_DEVICE_POINTER = nullptr;
assert(cudaMalloc(&d_TEST_DEVICE_POINTER, sizeof(int)) == CUDA_SUCCESS);
assert(cudaMemcpy(d_TEST_DEVICE_POINTER, &t_TEST_SYMBOL, sizeof(int), cudaMemcpyHostToDevice) == CUDA_SUCCESS);
// Build RTC kernel
std::vector<std::string> options;
// cuda path
std::string include_cuda;
include_cuda = "-I" + std::string(std::getenv("CUDA_PATH")) + "\\include";
options.push_back(include_cuda);
// rdc
std::string rdc;
rdc = "-rdc=true";
options.push_back(rdc);
#define __EXPERIMENTAL
#ifndef __EXPERIMENTAL
static jitify::JitCache jit_cache;
auto program = std::make_shared<jitify::Program>(jit_cache, rtc_kernel, 0, options);
#else
auto program = std::make_shared<jitify::experimental::Program>(rtc_kernel, std::vector<std::string>(), options);
#endif
// Launch RTC kernel
dim3 grid(1);
dim3 block(1);
auto kernel = program->kernel("simple_test");
const unsigned int nopts = 3;
CUjit_option opts[3] = { CU_JIT_GLOBAL_SYMBOL_COUNT, CU_JIT_GLOBAL_SYMBOL_NAMES, CU_JIT_GLOBAL_SYMBOL_ADDRESSES };
void* optvals[3];
unsigned int optvals_0 = 2;
optvals[0] = &optvals_0;
char* optvals_1[2] = {"TEST_DEVICE_SYMBOL", "TEST_CONSTANT_SYMBOL"};
optvals[1] = &optvals_1;
void *optvals_2[2];
assert(cudaGetSymbolAddress(&optvals_2[0], TEST_DEVICE_SYMBOL) == CUDA_SUCCESS);
assert(cudaGetSymbolAddress(&optvals_2[1], TEST_CONSTANT_SYMBOL) == CUDA_SUCCESS);
optvals[2] = &optvals_2;
auto kernelInstance = kernel.instantiate2(nopts, reinterpret_cast<CUjit_option*>(&opts), reinterpret_cast<void**>(&optvals));
auto kernelLauncher = kernelInstance.configure(grid, block);
assert(kernelLauncher.launch({ reinterpret_cast<void*>(&d_TEST_DEVICE_POINTER) }) == CUresult::CUDA_SUCCESS);
// Cleanup
assert(cudaDeviceSynchronize() == CUDA_SUCCESS);
assert(cudaFree(d_TEST_DEVICE_POINTER) == CUDA_SUCCESS);
#ifdef _WIN32
// Prevent window closing as probably launched via visual studio debugger
printf("Hit enter to exit.\n");
getchar();
#endif
return EXIT_SUCCESS;
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment