From 309dd11c84fbe1109a046701d6212e23be995618 Mon Sep 17 00:00:00 2001 From: maddyscientist Date: Fri, 9 Jun 2023 17:26:46 -0700 Subject: [PATCH 1/6] WAR when nvrtc fails to correctly deal with __has_include --- jitify.hpp | 41 +++++++++++++++++++++++++++++++++++++++++ 1 file changed, 41 insertions(+) diff --git a/jitify.hpp b/jitify.hpp index 21692e4..847d093 100644 --- a/jitify.hpp +++ b/jitify.hpp @@ -687,6 +687,47 @@ inline bool load_source( comment; } + // WAR where nvrtc can fail to correctly return if an include is present + if (cleanline.find("#if __has_include") != std::string::npos) { + // check for angle bracket include + size_t start = cleanline.find("<") + 1; + size_t count = cleanline.find(">") - start; + std::string has_include_name = cleanline.substr(start, count); + +#if JITIFY_PRINT_HEADER_PATHS + std::cout << "Found #if __has_include(<" << has_include_name << ">)" << " from " + << filename << ":" << linenum << std::endl; +#endif + // Try loading from filesystem + bool found_file = false; + std::string has_include_fullpath = path_join(current_dir, has_include_name); + if (search_current_dir) { + file_stream.open(has_include_fullpath.c_str()); + if (file_stream) found_file = true; + } + // Search include directories + if (!found_file) { + for (int i = 0; i < (int)include_paths.size(); ++i) { + has_include_fullpath = path_join(include_paths[i], has_include_name); + file_stream.open(has_include_fullpath.c_str()); + if (file_stream) { + found_file = true; + break; + } + } + if (!found_file) { + // Try loading from builtin headers + has_include_fullpath = path_join("__jitify_builtin", has_include_name); + auto it = get_jitsafe_headers_map().find(has_include_name); + if (it != get_jitsafe_headers_map().end()) { + found_file = true; + } + } + } + + line = found_file ? "#if 1" : "#if 0"; + } + source += line + "\n"; } // HACK TESTING (WAR for cub) From e38de0e06a633ed3c815dd91ec7f1edae3fceb24 Mon Sep 17 00:00:00 2001 From: maddyscientist Date: Fri, 9 Jun 2023 17:27:35 -0700 Subject: [PATCH 2/6] Remove now unnecessary includes builtin_numeric_cuda_std_limits_program since __has_include issue is now WARed --- jitify_test.cu | 2 -- 1 file changed, 2 deletions(-) diff --git a/jitify_test.cu b/jitify_test.cu index 4a79806..1fa8ada 100644 --- a/jitify_test.cu +++ b/jitify_test.cu @@ -961,8 +961,6 @@ static const char* const builtin_numeric_cuda_std_limits_program_source = "builtin_numeric_cuda_std_limits_program\n" "#include \n" "#include \n" - "#include \n" // test fails without this explicit include - "#include \n" "struct MyType {};\n" "namespace cuda {\n" "namespace std {\n" From 6f9e8acf4157bff4c283a2b108ee5fc0375ee8db Mon Sep 17 00:00:00 2001 From: maddyscientist Date: Mon, 12 Jun 2023 22:50:54 -0700 Subject: [PATCH 3/6] __has_include parsing is now more robust and handles quotation includes --- jitify.hpp | 95 ++++++++++++++++++++++++++++++++++++------------------ 1 file changed, 64 insertions(+), 31 deletions(-) diff --git a/jitify.hpp b/jitify.hpp index 847d093..4c74eac 100644 --- a/jitify.hpp +++ b/jitify.hpp @@ -688,44 +688,77 @@ inline bool load_source( } // WAR where nvrtc can fail to correctly return if an include is present - if (cleanline.find("#if __has_include") != std::string::npos) { - // check for angle bracket include - size_t start = cleanline.find("<") + 1; - size_t count = cleanline.find(">") - start; - std::string has_include_name = cleanline.substr(start, count); + size_t has_include_start = cleanline.find("__has_include"); + if (has_include_start != std::string::npos) { + // find subsequent opening and closing braces + size_t open = cleanline.find("(", has_include_start); + size_t close = cleanline.find(")", open); + if (!(open == std::string::npos || close == std::string::npos)) { + std::string has_name = cleanline.substr(open + 1, close - open - 1); + + size_t header_start = 0; + size_t header_count = 0; + bool quote_include = false; + if (has_name.find("<") != std::string::npos) { + header_start = has_name.find("<") + 1; + header_count = has_name.find(">") - header_start; + } else if (has_name.find("\"") != std::string::npos) { + quote_include = true; + header_start = has_name.find("\"") + 1; + header_count = has_name.find("\"", header_start) - header_start; + if (has_name.find("\"", header_start) == std::string::npos) + throw std::runtime_error("Malformed __has_include statement (" + + filename + ":" + std::to_string(linenum) + + ")"); + } + + if (header_count != 0) { + std::string has_include_name = + has_name.substr(header_start, header_count); #if JITIFY_PRINT_HEADER_PATHS - std::cout << "Found #if __has_include(<" << has_include_name << ">)" << " from " - << filename << ":" << linenum << std::endl; + std::cout << "Found #if __has_include(" << has_name << ")" + << " from " << filename << ":" << linenum << std::endl; #endif - // Try loading from filesystem - bool found_file = false; - std::string has_include_fullpath = path_join(current_dir, has_include_name); - if (search_current_dir) { - file_stream.open(has_include_fullpath.c_str()); - if (file_stream) found_file = true; - } - // Search include directories - if (!found_file) { - for (int i = 0; i < (int)include_paths.size(); ++i) { - has_include_fullpath = path_join(include_paths[i], has_include_name); - file_stream.open(has_include_fullpath.c_str()); - if (file_stream) { - found_file = true; - break; + // Try loading from filesystem + bool found_file = false; + std::string has_include_fullpath = + path_join(current_dir, has_include_name); + if (quote_include) { + file_stream.open(has_include_fullpath.c_str()); + if (file_stream) found_file = true; } - } - if (!found_file) { - // Try loading from builtin headers - has_include_fullpath = path_join("__jitify_builtin", has_include_name); - auto it = get_jitsafe_headers_map().find(has_include_name); - if (it != get_jitsafe_headers_map().end()) { - found_file = true; + // Search include directories + if (!found_file) { + for (int i = 0; i < (int)include_paths.size(); ++i) { + has_include_fullpath = + path_join(include_paths[i], has_include_name); + file_stream.open(has_include_fullpath.c_str()); + if (file_stream) { + found_file = true; + break; + } + } + if (!found_file) { + // Try loading from builtin headers + has_include_fullpath = + path_join("__jitify_builtin", has_include_name); + auto it = get_jitsafe_headers_map().find(has_include_name); + if (it != get_jitsafe_headers_map().end()) { + found_file = true; + } + } + } + + if (found_file) { + line = cleanline.substr(0, has_include_start) + "(1)" + + cleanline.substr(close + 1); + } else { + line = cleanline.substr(0, has_include_start) + "(0)" + + cleanline.substr(close + 1); } } } - - line = found_file ? "#if 1" : "#if 0"; } source += line + "\n"; From f42f49f8f970fccaa24adce393fe3fe1cbdbb97a Mon Sep 17 00:00:00 2001 From: maddyscientist Date: Mon, 12 Jun 2023 22:52:36 -0700 Subject: [PATCH 4/6] Added unit test for checking __has_include processing --- jitify_test.cu | 56 ++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 56 insertions(+) diff --git a/jitify_test.cu b/jitify_test.cu index 1fa8ada..477e7e6 100644 --- a/jitify_test.cu +++ b/jitify_test.cu @@ -1149,3 +1149,59 @@ TEST(JitifyTest, AssertHeader) { .configure(grid, block) .launch())); } + +static const char* const has_include_source = R"( + #if __has_include() + #else + #error __has_include failed + #endif + + #if 1 && __has_include() + #else + #error __has_include failed + #endif + + #if __has_include() && 0 + #error __has_include failed + #else + #endif + + #if __has_include("") + #else + #error __has_include failed + #endif + + #if __has_include("limits") + #else + #error __has_include failed + #endif + + #if __has_include("example_headers/my_header1.cuh") + #else + #error __has_include failed + #endif + + // check we don't touch these + #if defined(__has_include) + #else + #error __has_include failed + #endif + + #if !defined(__has_include) + #error __has_include failed + #endif + + __global__ void has_include_kernel() { } +)"; + +TEST(JitifyTest, HasInclude) { + // Checks that cassert works as expected + jitify::JitCache kernel_cache; + auto program = kernel_cache.program(has_include_source); + dim3 grid(1); + dim3 block(1); + CHECK_CUDA((program.kernel("has_include_kernel") + .instantiate<>() + .configure(grid, block) + .launch())); +} From 80f3920148adaa9b74a23286371c78302e48d85e Mon Sep 17 00:00:00 2001 From: maddyscientist Date: Tue, 5 Sep 2023 13:45:19 -0700 Subject: [PATCH 5/6] Restore AssertHeader test as being the final test --- jitify_test.cu | 28 ++++++++++++++-------------- 1 file changed, 14 insertions(+), 14 deletions(-) diff --git a/jitify_test.cu b/jitify_test.cu index 477e7e6..46d0ac7 100644 --- a/jitify_test.cu +++ b/jitify_test.cu @@ -1136,20 +1136,6 @@ TEST(JitifyTest, EnvVarOptions) { setenv("JITIFY_OPTIONS", "", true); } -// NOTE: This MUST be the last test in the file, due to sticky CUDA error. -TEST(JitifyTest, AssertHeader) { - // Checks that cassert works as expected - jitify::JitCache kernel_cache; - auto program = - kernel_cache.program(assert_program_source, {}, {"-I" CUDA_INC_DIR}); - dim3 grid(1); - dim3 block(1); - CHECK_CUDA((program.kernel("my_assert_kernel") - .instantiate<>() - .configure(grid, block) - .launch())); -} - static const char* const has_include_source = R"( #if __has_include() #else @@ -1205,3 +1191,17 @@ TEST(JitifyTest, HasInclude) { .configure(grid, block) .launch())); } + +// NOTE: This MUST be the last test in the file, due to sticky CUDA error. +TEST(JitifyTest, AssertHeader) { + // Checks that cassert works as expected + jitify::JitCache kernel_cache; + auto program = + kernel_cache.program(assert_program_source, {}, {"-I" CUDA_INC_DIR}); + dim3 grid(1); + dim3 block(1); + CHECK_CUDA((program.kernel("my_assert_kernel") + .instantiate<>() + .configure(grid, block) + .launch())); +} From 39ecb4c58871659ce65a6389a9614edc91c6fb56 Mon Sep 17 00:00:00 2001 From: maddyscientist Date: Tue, 5 Sep 2023 13:49:25 -0700 Subject: [PATCH 6/6] Some cleanup --- jitify.hpp | 37 +++++++++++++++---------------------- 1 file changed, 15 insertions(+), 22 deletions(-) diff --git a/jitify.hpp b/jitify.hpp index 4c74eac..68d18d7 100644 --- a/jitify.hpp +++ b/jitify.hpp @@ -696,21 +696,19 @@ inline bool load_source( if (!(open == std::string::npos || close == std::string::npos)) { std::string has_name = cleanline.substr(open + 1, close - open - 1); - size_t header_start = 0; - size_t header_count = 0; - bool quote_include = false; - if (has_name.find("<") != std::string::npos) { - header_start = has_name.find("<") + 1; - header_count = has_name.find(">") - header_start; - } else if (has_name.find("\"") != std::string::npos) { - quote_include = true; - header_start = has_name.find("\"") + 1; - header_count = has_name.find("\"", header_start) - header_start; - if (has_name.find("\"", header_start) == std::string::npos) - throw std::runtime_error("Malformed __has_include statement (" + - filename + ":" + std::to_string(linenum) + - ")"); - } + if (has_name.find("<") == std::string::npos && has_name.find("\"") == std::string::npos) + throw std::runtime_error("Malformed __has_include statement (" + + filename + ":" + std::to_string(linenum) + + ")"); + // are we using quote includes or angle brackets? + // (test using angle brackets, since quotes are valid around angle brackets) + bool quote_include = has_name.find("<") == std::string::npos; + size_t header_start = (quote_include ? has_name.find("\"") : has_name.find("<")) + 1; + size_t header_count = has_name.find(quote_include ? "\"" : ">", header_start) - header_start; + if (has_name.find(quote_include ? "\"" : ">", header_start) == std::string::npos) + throw std::runtime_error("Malformed __has_include statement (" + + filename + ":" + std::to_string(linenum) + + ")"); if (header_count != 0) { std::string has_include_name = @@ -750,13 +748,8 @@ inline bool load_source( } } - if (found_file) { - line = cleanline.substr(0, has_include_start) + "(1)" + - cleanline.substr(close + 1); - } else { - line = cleanline.substr(0, has_include_start) + "(0)" + - cleanline.substr(close + 1); - } + line = cleanline.substr(0, has_include_start) + (found_file ? "(1)" : "(0)") + + cleanline.substr(close + 1); } } }