[Bf-blender-cvs] [a5052770b85] master: cycles: Add an nvrtc based cubin cli compiler.
Ray Molenkamp
noreply at git.blender.org
Sat Feb 3 18:59:18 CET 2018
Commit: a5052770b85fefe00511886429e6fc1f5056e1e8
Author: Ray Molenkamp
Date: Sat Feb 3 10:59:09 2018 -0700
Branches: master
https://developer.blender.org/rBa5052770b85fefe00511886429e6fc1f5056e1e8
cycles: Add an nvrtc based cubin cli compiler.
nvcc is very picky regarding compiler versions, severely limiting the compiler we can use, this commit adds a nvrtc based compiler that'll allow us to build the cubins even if the host compiler is unsupported. for details see D2913.
Differential Revision: http://developer.blender.org/D2913
===================================================================
M CMakeLists.txt
M extern/cuew/include/cuew.h
M extern/cuew/src/cuew.c
M intern/cycles/CMakeLists.txt
M intern/cycles/app/CMakeLists.txt
A intern/cycles/app/cycles_cubin_cc.cpp
M intern/cycles/device/device_cuda.cpp
M intern/cycles/kernel/CMakeLists.txt
M intern/cycles/kernel/kernel_compat_cuda.h
===================================================================
diff --git a/CMakeLists.txt b/CMakeLists.txt
index 2152f6e51ae..213b5e96f06 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -409,12 +409,14 @@ option(WITH_CYCLES_STANDALONE_GUI "Build Cycles standalone with GUI" OFF)
option(WITH_CYCLES_OSL "Build Cycles with OSL support" ${_init_CYCLES_OSL})
option(WITH_CYCLES_OPENSUBDIV "Build Cycles with OpenSubdiv support" ${_init_CYCLES_OPENSUBDIV})
option(WITH_CYCLES_CUDA_BINARIES "Build Cycles CUDA binaries" OFF)
+option(WITH_CYCLES_CUBIN_COMPILER "Build cubins with nvrtc based compiler instead of nvcc" OFF)
set(CYCLES_CUDA_BINARIES_ARCH sm_20 sm_21 sm_30 sm_35 sm_37 sm_50 sm_52 sm_60 sm_61 CACHE STRING "CUDA architectures to build binaries for")
mark_as_advanced(CYCLES_CUDA_BINARIES_ARCH)
unset(PLATFORM_DEFAULT)
option(WITH_CYCLES_LOGGING "Build Cycles with logging support" ON)
option(WITH_CYCLES_DEBUG "Build Cycles with extra debug capabilities" OFF)
option(WITH_CYCLES_NATIVE_ONLY "Build Cycles with native kernel only (which fits current CPU, use for development only)" OFF)
+mark_as_advanced(WITH_CYCLES_CUBIN_COMPILER)
mark_as_advanced(WITH_CYCLES_LOGGING)
mark_as_advanced(WITH_CYCLES_DEBUG)
mark_as_advanced(WITH_CYCLES_NATIVE_ONLY)
diff --git a/extern/cuew/include/cuew.h b/extern/cuew/include/cuew.h
index 0eace96bc3f..f5009d4f2c7 100644
--- a/extern/cuew/include/cuew.h
+++ b/extern/cuew/include/cuew.h
@@ -1323,6 +1323,7 @@ int cuewInit(void);
const char *cuewErrorString(CUresult result);
const char *cuewCompilerPath(void);
int cuewCompilerVersion(void);
+int cuewNvrtcVersion(void);
#ifdef __cplusplus
}
diff --git a/extern/cuew/src/cuew.c b/extern/cuew/src/cuew.c
index 962059bfcce..b68dc597049 100644
--- a/extern/cuew/src/cuew.c
+++ b/extern/cuew/src/cuew.c
@@ -329,7 +329,7 @@ int cuewInit(void) {
#ifdef _WIN32
/* Expected in c:/windows/system or similar, no path needed. */
const char *cuda_paths[] = {"nvcuda.dll", NULL};
- const char *nvrtc_paths[] = {"nvrtc.dll", NULL};
+ const char *nvrtc_paths[] = {"nvrtc64_80.dll", "nvrtc64_90.dll", "nvrtc64_91.dll", NULL};
#elif defined(__APPLE__)
/* Default installation path. */
const char *cuda_paths[] = {"/usr/local/cuda/lib/libcuda.dylib", NULL};
@@ -766,6 +766,15 @@ const char *cuewCompilerPath(void) {
return NULL;
}
+int cuewNvrtcVersion(void) {
+ int major, minor;
+ if (nvrtcVersion) {
+ nvrtcVersion(&major, &minor);
+ return 10 * major + minor;
+ }
+ return 0;
+}
+
int cuewCompilerVersion(void) {
const char *path = cuewCompilerPath();
const char *marker = "Cuda compilation tools, release ";
diff --git a/intern/cycles/CMakeLists.txt b/intern/cycles/CMakeLists.txt
index 543fb5b2cf8..3da1170ec77 100644
--- a/intern/cycles/CMakeLists.txt
+++ b/intern/cycles/CMakeLists.txt
@@ -242,6 +242,24 @@ if(CMAKE_COMPILER_IS_GNUCXX)
unset(_has_no_error_unused_macros)
endif()
+if(WITH_CYCLES_CUDA_BINARIES AND (NOT WITH_CYCLES_CUBIN_COMPILER))
+ if(MSVC)
+ set(MAX_MSVC 1800)
+ if(${CUDA_VERSION} EQUAL "8.0")
+ set(MAX_MSVC 1900)
+ elseif(${CUDA_VERSION} EQUAL "9.0")
+ set(MAX_MSVC 1910)
+ elseif(${CUDA_VERSION} EQUAL "9.1")
+ set(MAX_MSVC 1911)
+ endif()
+ if (NOT MSVC_VERSION LESS ${MAX_MSVC})
+ message(STATUS "nvcc not supported for this compiler version, using cycles_cubin_cc instead.")
+ set(WITH_CYCLES_CUBIN_COMPILER ON)
+ endif()
+ unset(MAX_MSVC)
+ endif()
+endif()
+
# Subdirectories
@@ -254,7 +272,7 @@ if(WITH_CYCLES_NETWORK)
add_definitions(-DWITH_NETWORK)
endif()
-if(WITH_CYCLES_STANDALONE OR WITH_CYCLES_NETWORK)
+if(WITH_CYCLES_STANDALONE OR WITH_CYCLES_NETWORK OR WITH_CYCLES_CUBIN_COMPILER)
add_subdirectory(app)
endif()
diff --git a/intern/cycles/app/CMakeLists.txt b/intern/cycles/app/CMakeLists.txt
index 08a3931ef46..9ebeceb1659 100644
--- a/intern/cycles/app/CMakeLists.txt
+++ b/intern/cycles/app/CMakeLists.txt
@@ -120,3 +120,27 @@ if(WITH_CYCLES_NETWORK)
endif()
unset(SRC)
endif()
+
+if(WITH_CYCLES_CUBIN_COMPILER)
+ # 32 bit windows is special, nvrtc is not supported on x86, so even
+ # though we are building 32 bit blender a 64 bit cubin_cc will have
+ # to be build to compile the cubins.
+ if(MSVC AND NOT CMAKE_CL_64)
+ Message("cycles_cubin_cc not supported on x86")
+ else()
+ set(SRC cycles_cubin_cc.cpp)
+ set(INC ../../../extern/cuew/include)
+ add_executable(cycles_cubin_cc ${SRC})
+ include_directories(${INC})
+ target_link_libraries(cycles_cubin_cc
+ extern_cuew
+ ${OPENIMAGEIO_LIBRARIES}
+ ${PLATFORM_LINKLIBS}
+ )
+ if(NOT CYCLES_STANDALONE_REPOSITORY)
+ target_link_libraries(cycles_cubin_cc bf_intern_guardedalloc)
+ endif()
+ unset(SRC)
+ unset(INC)
+ endif()
+endif()
diff --git a/intern/cycles/app/cycles_cubin_cc.cpp b/intern/cycles/app/cycles_cubin_cc.cpp
new file mode 100644
index 00000000000..c1f3974be6d
--- /dev/null
+++ b/intern/cycles/app/cycles_cubin_cc.cpp
@@ -0,0 +1,284 @@
+/*
+ * Copyright 2017 Blender Foundation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <stdio.h>
+#include <stdint.h>
+
+#include <string>
+#include <vector>
+
+#include <OpenImageIO/argparse.h>
+#include <OpenImageIO/filesystem.h>
+
+#include "cuew.h"
+
+#ifdef _MSC_VER
+# include <Windows.h>
+#endif
+
+using std::string;
+using std::vector;
+
+class CompilationSettings
+{
+public:
+ CompilationSettings()
+ : target_arch(0),
+ bits(64),
+ verbose(false),
+ fast_math(false)
+ {}
+
+ string cuda_toolkit_dir;
+ string input_file;
+ string output_file;
+ string ptx_file;
+ vector<string> defines;
+ vector<string> includes;
+ int target_arch;
+ int bits;
+ bool verbose;
+ bool fast_math;
+};
+
+bool compile_cuda(CompilationSettings &settings)
+{
+ const char* headers[] = {"stdlib.h" , "float.h", "math.h", "stdio.h"};
+ const char* header_content[] = {"\n", "\n", "\n", "\n"};
+
+ printf("Building %s\n", settings.input_file.c_str());
+
+ string code;
+ if(!OIIO::Filesystem::read_text_file(settings.input_file, code)) {
+ fprintf(stderr, "Error: unable to read %s\n", settings.input_file.c_str());
+ return false;
+ }
+
+ vector<string> options;
+ for(size_t i = 0; i < settings.includes.size(); i++) {
+ options.push_back("-I" + settings.includes[i]);
+ }
+
+ for(size_t i = 0; i < settings.defines.size(); i++) {
+ options.push_back("-D" + settings.defines[i]);
+ }
+
+ options.push_back("-arch=compute_" + std::to_string(settings.target_arch));
+ options.push_back("--device-as-default-execution-space");
+ if(settings.fast_math)
+ options.push_back("--use_fast_math");
+
+ nvrtcProgram prog;
+ nvrtcResult result = nvrtcCreateProgram(&prog,
+ code.c_str(), // buffer
+ NULL, // name
+ sizeof(headers) / sizeof(void*), // numHeaders
+ header_content, // headers
+ headers); // includeNames
+
+ if(result != NVRTC_SUCCESS) {
+ fprintf(stderr, "Error: nvrtcCreateProgram failed (%x)\n\n", result);
+ return false;
+ }
+
+ /* Tranfer options to a classic C array. */
+ vector<const char*> opts(options.size());
+ for(size_t i = 0; i < options.size(); i++) {
+ opts[i] = options[i].c_str();
+ }
+
+ result = nvrtcCompileProgram(prog, options.size(), &opts[0]);
+
+ if(result != NVRTC_SUCCESS) {
+ fprintf(stderr, "Error: nvrtcCompileProgram failed (%x)\n\n", result);
+
+ size_t log_size;
+ nvrtcGetProgramLogSize(prog, &log_size);
+
+ vector<char> log(log_size);
+ nvrtcGetProgramLog(prog, &log[0]);
+ fprintf(stderr, "%s\n", &log[0]);
+
+ return false;
+ }
+
+ /* Retrieve the ptx code. */
+ size_t ptx_size;
+ result = nvrtcGetPTXSize(prog, &ptx_size);
+ if(result != NVRTC_SUCCESS) {
+ fprintf(stderr, "Error: nvrtcGetPTXSize failed (%x)\n\n", result);
+ return false;
+ }
+
+ vector<char> ptx_code(ptx_size);
+ result = nvrtcGetPTX(prog, &ptx_code[0]);
+ if(result != NVRTC_SUCCESS) {
+ fprintf(stderr, "Error: nvrtcGetPTX failed (%x)\n\n", result);
+ return false;
+ }
+
+ /* Write a file in the temp folder with the ptx code. */
+ settings.ptx_file = OIIO::Filesystem::temp_directory_path() + "/" + OIIO::Filesystem::unique_path();
+ FILE * f= fopen(settings.ptx_file.c_str(), "wb");
+ fwrite(&ptx_code[0], 1, ptx_size, f);
+ fclose(f);
+
+ return true;
+}
+
+bool link_ptxas(CompilationSettings &settings)
+{
+ string cudapath = "";
+ if(settings.cuda_toolkit_dir.size())
+ cudapath = settings.cuda_toolkit_dir + "/bin/";
+
+ string ptx = "\"" +cudapath + "ptxas\" " + settings.ptx_file +
+ " -o " + settings.output_file +
+ " --gpu-name sm_" + std::to_string(settings.target_arch) +
+ " -m" + std::to_string(settings.bits);
+
+ if(settings.verbose)
+ ptx += " --verbose";
+
+ int pxresult = system(ptx.c_str());
+ if(pxresult) {
+ fprintf(stderr, "Error: ptxas failed (%x)\n\n", pxresult);
+ return false;
+ }
+
+ if(!OIIO::Filesystem::remove(settings.ptx_file)) {
+ fprintf(stderr, "Error: removing %s\n\n", settings.ptx_file.c_str());
+ }
+
+ return true;
+}
+
+bool init(CompilationSettings &settings)
+{
+#ifdef _MSC_VER
+ if(settings.cuda_toolkit_dir.size()) {
+ SetDllDirectory((settings.cuda_toolkit_dir + "/bin").c_str());
+ }
+#endif
+
+ int cuewresult = cuewInit();
+ if(cuewresult != CUEW_SUCCESS) {
+ fprintf(stderr, "Error: cuew init fialed (0x%x)\n\n", cuewresult);
+ return false;
+ }
+
+ if(cuewNvrtcVersion() < 80) {
+ fprintf(stderr, "Error: only cuda 8 and higher is supported, %d\n\n", cuewCompilerVersion());
+ return false;
+ }
+
+ if(!nvrtcCreateProgram) {
+ fprintf(stderr, "Error: nvrtcCreateProgram not resolved\n");
+ return false;
+ }
+
+ if(!nvrtcCompileProgram) {
+ fprintf(stderr, "Error: nvrtcCompileProgram not resolved\n");
+ return false;
+ }
+
+ if(!nvrtcGetProgramLogSize) {
+ fprintf(stderr, "Error: nvrtcGetPr
@@ Diff output truncated at 10240 characters. @@
More information about the Bf-blender-cvs
mailing list