Merge branch 'master' into blender2.8

CMakeLists.txt

@@ -415,12 +415,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)
@@ -531,6 +533,9 @@ if(WIN32)
 
 	set(WINDOWS_CODESIGN_PFX_PASSWORD CACHE STRING  "password for pfx file used for codesigning.")
 	mark_as_advanced(WINDOWS_CODESIGN_PFX_PASSWORD)
+
+	option(WINDOWS_USE_VISUAL_STUDIO_FOLDERS "Organize the visual studio project according to source folders." ON)
+	mark_as_advanced(WINDOWS_USE_VISUAL_STUDIO_FOLDERS)
 endif()
 
 # avoid using again

build_files/cmake/macros.cmake

@@ -242,6 +242,13 @@ function(blender_add_lib__impl
 	# listed is helpful for IDE's (QtCreator/MSVC)
 	blender_source_group("${sources}")
 
+	#if enabled, set the FOLDER property for visual studio projects 
+	if(WINDOWS_USE_VISUAL_STUDIO_FOLDERS)
+		get_filename_component(FolderDir ${CMAKE_CURRENT_SOURCE_DIR} DIRECTORY)
+		string(REPLACE ${CMAKE_SOURCE_DIR} "" FolderDir ${FolderDir})
+		set_target_properties(${name} PROPERTIES FOLDER ${FolderDir})
+	endif()
+
 	list_assert_duplicates("${sources}")
 	list_assert_duplicates("${includes}")
 	# Not for system includes because they can resolve to the same path

build_files/cmake/platform/platform_win32.cmake

@@ -31,6 +31,12 @@ endif()
 
 # Libraries configuration for Windows when compiling with MSVC.
 
+set_property(GLOBAL PROPERTY USE_FOLDERS ${WINDOWS_USE_VISUAL_STUDIO_FOLDERS})
+
+if(NOT WITH_PYTHON_MODULE)
+	set_property(DIRECTORY PROPERTY VS_STARTUP_PROJECT blender)
+endif()
+
 macro(warn_hardcoded_paths package_name
 	)
 	if(WITH_WINDOWS_FIND_MODULES)

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
 }

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 ";

intern/cycles/CMakeLists.txt

@@ -9,6 +9,7 @@ endif()
 # External Libraries
 
 include(cmake/external_libs.cmake)
+include(cmake/macros.cmake)
 
 # Build Flags
 # todo: this code could be refactored a bit to avoid duplication
@@ -242,6 +243,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 +273,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()
 

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()

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: nvrtcGetProgramLogSize not resolved\n");
+		return false;
+	}
+
+	if(!nvrtcGetProgramLog) {
+		fprintf(stderr, "Error: nvrtcGetProgramLog not resolved\n");
+		return false;
+	}
+
+	if(!nvrtcGetPTXSize) {
+		fprintf(stderr, "Error: nvrtcGetPTXSize not resolved\n");
+		return false;
+	}
+
+	if(!nvrtcGetPTX) {
+		fprintf(stderr, "Error: nvrtcGetPTX not resolved\n");
+		return false;
+	}
+
+	return true;
+}
+
+bool parse_parameters(int argc, const char **argv, CompilationSettings &settings)
+{
+	OIIO::ArgParse ap;
+	ap.options("Usage: cycles_cubin_cc [options]",
+		"-target %d", &settings.target_arch, "target shader model",
+		"-m %d", &settings.bits, "Cuda architecture bits",
+		"-i %s", &settings.input_file, "Input source filename",
+		"-o %s", &settings.output_file, "Output cubin filename",
+		"-I %L", &settings.includes, "Add additional includepath",
+		"-D %L", &settings.defines, "Add additional defines",
+		"-v", &settings.verbose, "Use verbose logging",
+		"--use_fast_math", &settings.fast_math, "Use fast math",
+		"-cuda-toolkit-dir %s", &settings.cuda_toolkit_dir, "path to the cuda toolkit binary directory",
+		NULL);
+
+	if(ap.parse(argc, argv) < 0) {
+		fprintf(stderr, "%s\n", ap.geterror().c_str());
+		ap.usage();
+		return false;
+	}
+
+	if(!settings.output_file.size()) {
+		fprintf(stderr, "Error: Output file not set(-o), required\n\n");
+		return false;
+	}
+
+	if(!settings.input_file.size()) {
+		fprintf(stderr, "Error: Input file not set(-i, required\n\n");
+		return false;
+	}
+
+	if(!settings.target_arch) {
+		fprintf(stderr, "Error: target shader model not set (-target), required\n\n");
+		return false;
+	}
+
+	return true;
+}
+
+int main(int argc, const char **argv)
+{
+	CompilationSettings settings;
+
+	if(!parse_parameters(argc, argv, settings)) {
+		fprintf(stderr, "Error: invalid parameters, exiting\n");
+		exit(EXIT_FAILURE);
+	}
+
+	if(!init(settings)) {
+		fprintf(stderr, "Error: initialization error, exiting\n");
+		exit(EXIT_FAILURE);
+	}
+
+	if(!compile_cuda(settings)) {
+		fprintf(stderr, "Error: compilation error, exiting\n");
+		exit(EXIT_FAILURE);
+	}
+
+	if(!link_ptxas(settings)) {
+		exit(EXIT_FAILURE);
+	}
+
+	return 0;
+}

intern/cycles/bvh/CMakeLists.txt

@@ -34,4 +34,4 @@ set(SRC_HEADERS
 include_directories(${INC})
 include_directories(SYSTEM ${INC_SYS})
 
-add_library(cycles_bvh ${SRC} ${SRC_HEADERS})
+cycles_add_library(cycles_bvh ${SRC} ${SRC_HEADERS})

intern/cycles/cmake/macros.cmake

@@ -0,0 +1,12 @@
+function(cycles_set_solution_folder target)
+	if(WINDOWS_USE_VISUAL_STUDIO_FOLDERS)
+		get_filename_component(folderdir ${CMAKE_CURRENT_SOURCE_DIR} DIRECTORY)
+		string(REPLACE ${CMAKE_SOURCE_DIR} "" folderdir ${folderdir})
+		set_target_properties(${target} PROPERTIES FOLDER ${folderdir})
+	endif()
+endfunction()
+
+macro(cycles_add_library target)
+	add_library(${target} ${ARGN})
+	cycles_set_solution_folder(${target})
+endmacro()

intern/cycles/device/CMakeLists.txt

@@ -77,4 +77,4 @@ endif()
 include_directories(${INC})
 include_directories(SYSTEM ${INC_SYS})
 
-add_library(cycles_device ${SRC} ${SRC_OPENCL} ${SRC_HEADERS})
+cycles_add_library(cycles_device ${SRC} ${SRC_OPENCL} ${SRC_HEADERS})

intern/cycles/device/device_cuda.cpp

@@ -348,7 +348,6 @@ public:
 	        const DeviceRequestedFeatures& requested_features,
 	        bool filter=false, bool split=false)
 	{
-		const int cuda_version = cuewCompilerVersion();
 		const int machine = system_cpu_bits();
 		const string source_path = path_get("source");
 		const string include_path = source_path;
@@ -356,10 +355,8 @@ public:
 		                              "--ptxas-options=\"-v\" "
 		                              "--use_fast_math "
 		                              "-DNVCC "
-		                              "-D__KERNEL_CUDA_VERSION__=%d "
 		                               "-I\"%s\"",
 		                              machine,
-		                              cuda_version,
 		                              include_path.c_str());
 		if(!filter && use_adaptive_compilation()) {
 			cflags += " " + requested_features.get_build_options();

intern/cycles/graph/CMakeLists.txt

@@ -19,5 +19,5 @@ set(SRC_HEADERS
 include_directories(${INC})
 include_directories(SYSTEM ${INC_SYS})
 
-add_library(cycles_graph ${SRC} ${SRC_HEADERS})
+cycles_add_library(cycles_graph ${SRC} ${SRC_HEADERS})
 

intern/cycles/kernel/CMakeLists.txt

@@ -356,55 +356,67 @@ if(WITH_CYCLES_CUDA_BINARIES)
 	set(cuda_cubins)
 
 	macro(CYCLES_CUDA_KERNEL_ADD arch name flags sources experimental)
+		set(cuda_cubin ${name}_${arch}.cubin)
+		set(cuda_kernel_src "/kernels/cuda/${name}.cu")
+
+		set(cuda_flags
+			-D CCL_NAMESPACE_BEGIN=
+			-D CCL_NAMESPACE_END=
+			-D NVCC
+			-m ${CUDA_BITS}
+			-I ${CMAKE_CURRENT_SOURCE_DIR}/..
+			-I ${CMAKE_CURRENT_SOURCE_DIR}/kernels/cuda
+			--use_fast_math
+			-o ${CMAKE_CURRENT_BINARY_DIR}/${cuda_cubin})
+
 		if(${experimental})
-			set(flags ${flags} -D__KERNEL_EXPERIMENTAL__)
+			set(cuda_flags ${cuda_flags} -D __KERNEL_EXPERIMENTAL__)
 			set(name ${name}_experimental)
 		endif()
 
-		set(cuda_cubin ${name}_${arch}.cubin)
-
 		if(WITH_CYCLES_DEBUG)
-			set(cuda_debug_flags "-D__KERNEL_DEBUG__")
-		else()
-			set(cuda_debug_flags "")
+			set(cuda_flags ${cuda_flags} -D __KERNEL_DEBUG__)
 		endif()
 
-		set(cuda_nvcc_command ${CUDA_NVCC_EXECUTABLE})
-		set(cuda_nvcc_version ${CUDA_VERSION})
+		if(WITH_CYCLES_CUBIN_COMPILER)
+			string(SUBSTRING ${arch} 3 -1 CUDA_ARCH)
 
-		set(cuda_version_flags "-D__KERNEL_CUDA_VERSION__=${cuda_nvcc_version}")
-		set(cuda_math_flags "--use_fast_math")
-
-		set(cuda_kernel_src "/kernels/cuda/${name}.cu")
-
-		add_custom_command(
-			OUTPUT ${cuda_cubin}
-			COMMAND ${cuda_nvcc_command}
-					-arch=${arch}
-					${CUDA_NVCC_FLAGS}
-					-m${CUDA_BITS}
-					--cubin ${CMAKE_CURRENT_SOURCE_DIR}${cuda_kernel_src}
-					-o ${CMAKE_CURRENT_BINARY_DIR}/${cuda_cubin}
-					--ptxas-options="-v"
-					${cuda_arch_flags}
-					${cuda_version_flags}
-					${cuda_math_flags}
-					${flags}
-					${cuda_debug_flags}
-					-I${CMAKE_CURRENT_SOURCE_DIR}/..
-					-DCCL_NAMESPACE_BEGIN=
-					-DCCL_NAMESPACE_END=
-					-DNVCC
-			DEPENDS ${sources})
+			# Needed to find libnvrtc-builtins.so. Can't do it from inside
+			# cycles_cubin_cc since the env variable is read before main()
+			if(APPLE)
+				set(CUBIN_CC_ENV ${CMAKE_COMMAND}
+					-E env DYLD_LIBRARY_PATH="${CUDA_TOOLKIT_ROOT_DIR}/lib")
+			elseif(UNIX)
+				set(CUBIN_CC_ENV ${CMAKE_COMMAND}
+					-E env LD_LIBRARY_PATH="${CUDA_TOOLKIT_ROOT_DIR}/lib64")
+			endif()
 
+			add_custom_command(
+				OUTPUT ${cuda_cubin}
+				COMMAND ${CUBIN_CC_ENV}
+						"$<TARGET_FILE:cycles_cubin_cc>"
+						-target ${CUDA_ARCH}
+						-i ${CMAKE_CURRENT_SOURCE_DIR}${cuda_kernel_src}
+						${cuda_flags}
+						-v
+						-cuda-toolkit-dir "${CUDA_TOOLKIT_ROOT_DIR}"
+				DEPENDS ${sources} cycles_cubin_cc)
+		else()
+			add_custom_command(
+				OUTPUT ${cuda_cubin}
+				COMMAND ${CUDA_NVCC_EXECUTABLE}
+						-arch=${arch}
+						${CUDA_NVCC_FLAGS}
+						--cubin
+						${CMAKE_CURRENT_SOURCE_DIR}${cuda_kernel_src}
+						--ptxas-options="-v"
+						${cuda_flags}
+				DEPENDS ${sources})
+		endif()
 		delayed_install("${CMAKE_CURRENT_BINARY_DIR}" "${cuda_cubin}" ${CYCLES_INSTALL_PATH}/lib)
 		list(APPEND cuda_cubins ${cuda_cubin})
 
-		unset(cuda_extra_flags)
 		unset(cuda_debug_flags)
-
-		unset(cuda_nvcc_command)
-		unset(cuda_nvcc_version)
 	endmacro()
 
 	foreach(arch ${CYCLES_CUDA_BINARIES_ARCH})
@@ -412,17 +424,18 @@ if(WITH_CYCLES_CUDA_BINARIES)
 			message(STATUS "CUDA binaries for ${arch} disabled, not supported by CUDA 9.")
 		else()
 			# Compile regular kernel
-			CYCLES_CUDA_KERNEL_ADD(${arch} kernel "" "${cuda_sources}" FALSE)
 			CYCLES_CUDA_KERNEL_ADD(${arch} filter "" "${cuda_filter_sources}" FALSE)
+			CYCLES_CUDA_KERNEL_ADD(${arch} kernel "" "${cuda_sources}" FALSE)
 
 			if(WITH_CYCLES_CUDA_SPLIT_KERNEL_BINARIES)
 				# Compile split kernel
-				CYCLES_CUDA_KERNEL_ADD(${arch} kernel_split "-D__SPLIT__" ${cuda_sources} FALSE)
+				CYCLES_CUDA_KERNEL_ADD(${arch} kernel_split "-D __SPLIT__" ${cuda_sources} FALSE)
 			endif()
 		endif()
 	endforeach()
 
 	add_custom_target(cycles_kernel_cuda ALL DEPENDS ${cuda_cubins})
+	cycles_set_solution_folder(cycles_kernel_cuda)
 endif()
 
 # OSL module
@@ -465,7 +478,7 @@ if(CXX_HAS_AVX2)
 	set_source_files_properties(kernels/cpu/filter_avx2.cpp PROPERTIES COMPILE_FLAGS "${CYCLES_AVX2_KERNEL_FLAGS}")
 endif()
 
-add_library(cycles_kernel
+cycles_add_library(cycles_kernel
 	${SRC_CPU_KERNELS}
 	${SRC_CUDA_KERNELS}
 	${SRC_OPENCL_KERNELS}

intern/cycles/kernel/kernel_compat_cuda.h

@@ -30,10 +30,22 @@
 #  define __NODES_FEATURES__ NODE_FEATURE_ALL
 #endif
 
-#include <cuda.h>
-#include <cuda_fp16.h>
-#include <float.h>
-#include <stdint.h>
+/* Manual definitions so we can compile without CUDA toolkit. */
+
+typedef unsigned int uint32_t;
+typedef unsigned long long uint64_t;
+typedef unsigned short half;
+typedef unsigned long long CUtexObject;
+
+#define FLT_MAX 1.175494350822287507969e-38f
+#define FLT_MIN 340282346638528859811704183484516925440.0f
+
+__device__ half __float2half(const float f)
+{
+       half val;
+       asm("{  cvt.rn.f16.f32 %0, %1;}\n" : "=h"(val) : "f"(f));
+       return val;
+}
 
 /* Qualifier wrappers for different names on different devices */
 

intern/cycles/kernel/osl/CMakeLists.txt

@@ -30,5 +30,5 @@ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${RTTI_DISABLE_FLAGS}")
 include_directories(${INC})
 include_directories(SYSTEM ${INC_SYS})
 
-add_library(cycles_kernel_osl ${SRC} ${HEADER_SRC})
+cycles_add_library(cycles_kernel_osl ${SRC} ${HEADER_SRC})
 

intern/cycles/kernel/shaders/CMakeLists.txt

@@ -104,6 +104,7 @@ set(SRC_OSO
 # TODO, add a module to compile OSL
 foreach(_file ${SRC_OSL})
 	set(_OSL_FILE ${CMAKE_CURRENT_SOURCE_DIR}/${_file})
+	set_source_files_properties(${_file} PROPERTIES HEADER_FILE_ONLY TRUE)
 	string(REPLACE ".osl" ".oso" _OSO_FILE ${_OSL_FILE})
 	string(REPLACE ${CMAKE_SOURCE_DIR} ${CMAKE_BINARY_DIR} _OSO_FILE ${_OSO_FILE})
 	add_custom_command(
@@ -118,7 +119,8 @@ foreach(_file ${SRC_OSL})
 	unset(_OSO_FILE)
 endforeach()
 
-add_custom_target(cycles_osl_shaders ALL DEPENDS ${SRC_OSO} ${SRC_OSL_HEADERS} ${OSL_COMPILER})
+add_custom_target(cycles_osl_shaders ALL DEPENDS ${SRC_OSO} ${SRC_OSL_HEADERS} ${OSL_COMPILER} SOURCES ${SRC_OSL})
+cycles_set_solution_folder(cycles_osl_shaders)
 
 # CMAKE_CURRENT_SOURCE_DIR is already included in OSO paths
 delayed_install("" "${SRC_OSO}" ${CYCLES_INSTALL_PATH}/shader)

intern/cycles/render/CMakeLists.txt

@@ -71,4 +71,4 @@ include_directories(SYSTEM ${INC_SYS})
 
 add_definitions(${GL_DEFINITIONS})
 
-add_library(cycles_render ${SRC} ${SRC_HEADERS})
+cycles_add_library(cycles_render ${SRC} ${SRC_HEADERS})

intern/cycles/subd/CMakeLists.txt

@@ -24,4 +24,4 @@ set(SRC_HEADERS
 include_directories(${INC})
 include_directories(SYSTEM ${INC_SYS})
 
-add_library(cycles_subd ${SRC} ${SRC_HEADERS})
+cycles_add_library(cycles_subd ${SRC} ${SRC_HEADERS})

intern/cycles/util/CMakeLists.txt

@@ -127,4 +127,4 @@ include_directories(SYSTEM ${INC_SYS})
 
 add_definitions(${GL_DEFINITIONS})
 
-add_library(cycles_util ${SRC} ${SRC_HEADERS})
+cycles_add_library(cycles_util ${SRC} ${SRC_HEADERS})