sandey/blender
1
0
Fork 0
forked from bartvdbraak/blender
Code Pull Requests Activity

Cycles:

Browse Source 
* OpenCL now only uses GPU/Accelerator devices, it's only confusing if CPU
  device is used, easy to enable in the code for debugging.
* OpenCL kernel binaries are now cached for faster startup after the first
  time compiling.
* CUDA kernels can now be compiled and cached at runtime if the CUDA toolkit
  is installed. This means that even if the build does not have CUDA enabled,
  it's still possible to use it as long as you install the toolkit.
This commit is contained in:
Brecht Van Lommel 2011-09-09 12:04:39 +00:00
parent 9b31cba74e
commit cfbd6cf154
11 changed files with 317 additions and 46 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
intern/cycles/CMakeLists.txt
View File

@ -46,7 +46,7 @@ if(WITH_CYCLES_MULTI)
endif()
if(WITH_CYCLES_CUDA)
add_definitions(-DWITH_CUDA)
add_definitions(-DWITH_CUDA_BINARIES)
endif()
if(WITH_CYCLES_OSL)
@ -58,6 +58,7 @@ if(WITH_CYCLES_PARTIO)
endif()
add_definitions(-DWITH_OPENCL)
add_definitions(-DWITH_CUDA)
include_directories(
${BOOST_INCLUDE_DIR}

6
intern/cycles/blender/addon/engine.py
View File

@ -21,7 +21,11 @@ import bpy
def init():
import libcycles_blender as lib
import os.path
lib.init(os.path.dirname(__file__))
path = os.path.dirname(__file__)
user_path = os.path.dirname(os.path.abspath(bpy.utils.user_resource('CONFIG', '')))
lib.init(path, user_path)
def create(engine, data, scene, region = 0, v3d = 0, rv3d = 0):
import libcycles_blender as lib

6
intern/cycles/blender/blender_python.cpp
View File

@ -28,12 +28,12 @@ CCL_NAMESPACE_BEGIN
static PyObject *init_func(PyObject *self, PyObject *args)
{
const char *path;
const char *path, *user_path;
if(!PyArg_ParseTuple(args, "s", &path))
if(!PyArg_ParseTuple(args, "ss", &path, &user_path))
return NULL;
path_init(path);
path_init(path, user_path);
Py_INCREF(Py_None);
return Py_None;

103
intern/cycles/device/device_cuda.cpp
View File

@ -28,7 +28,9 @@
#include "util_map.h"
#include "util_opengl.h"
#include "util_path.h"
#include "util_system.h"
#include "util_types.h"
#include "util_time.h"
CCL_NAMESPACE_BEGIN
@ -125,6 +127,15 @@ public:
} \
}
bool cuda_error(CUresult result)
{
if(result == CUDA_SUCCESS)
return false;
fprintf(stderr, "CUDA error: %s\n", cuda_error_string(result));
return true;
}
void cuda_push_context()
{
cuda_assert(cuCtxSetCurrent(cuContext))
@ -140,17 +151,26 @@ public:
background = background_;
cuDevId = 0;
cuDevice = 0;
cuContext = 0;
/* intialize */
cuda_assert(cuInit(0))
if(cuda_error(cuInit(0)))
return;
/* setup device and context */
cuda_assert(cuDeviceGet(&cuDevice, cuDevId))
if(cuda_error(cuDeviceGet(&cuDevice, cuDevId)))
return;
CUresult result;
if(background)
cuda_assert(cuCtxCreate(&cuContext, 0, cuDevice))
result = cuCtxCreate(&cuContext, 0, cuDevice);
else
cuda_assert(cuGLCtxCreate(&cuContext, 0, cuDevice))
result = cuGLCtxCreate(&cuContext, 0, cuDevice);
if(cuda_error(result))
return;
cuda_pop_context();
}
@ -173,21 +193,80 @@ public:
return string("CUDA ") + deviceName;
}
string compile_kernel()
{
/* compute cubin name */
int major, minor;
cuDeviceComputeCapability(&major, &minor, cuDevId);
/* attempt to use kernel provided with blender */
string cubin = path_get(string_printf("lib/kernel_sm_%d%d.cubin", major, minor));
if(path_exists(cubin))
return cubin;
/* not found, try to use locally compiled kernel */
string kernel_path = path_get("kernel");
string md5 = path_files_md5_hash(kernel_path);
cubin = string_printf("cycles_kernel_sm%d%d_%s.cubin", major, minor, md5.c_str());;
cubin = path_user_get(path_join("cache", cubin));
/* if exists already, use it */
if(path_exists(cubin))
return cubin;
/* if not, find CUDA compiler */
string nvcc = cuCompilerPath();
if(nvcc == "") {
fprintf(stderr, "CUDA nvcc compiler not found. Install CUDA toolkit in default location.\n");
return "";
}
/* compile */
string kernel = path_join(kernel_path, "kernel.cu");
string include = kernel_path;
const int machine = system_cpu_bits();
const int maxreg = 24;
double starttime = time_dt();
printf("Compiling CUDA kernel ...\n");
string command = string_printf("%s -arch=sm_%d%d -m%d --cubin \"%s\" --use_fast_math "
"-o \"%s\" --ptxas-options=\"-v\" --maxrregcount=%d --opencc-options -OPT:Olimit=0 -I\"%s\" -DNVCC",
nvcc.c_str(), major, minor, machine, kernel.c_str(), cubin.c_str(), maxreg, include.c_str());
system(command.c_str());
/* verify if compilation succeeded */
if(!path_exists(cubin)) {
fprintf(stderr, "CUDA kernel compilation failed.\n");
return "";
}
printf("Kernel compilation finished in %.2lfs.\n", time_dt() - starttime);
return cubin;
}
bool load_kernels()
{
CUresult result;
int major, minor;
/* check if cuda init succeeded */
if(cuContext == 0)
return false;
cuda_push_context();
/* get kernel */
string cubin = compile_kernel();
if(cubin == "")
return false;
/* open module */
cuDeviceComputeCapability(&major, &minor, cuDevId);
string cubin = path_get(string_printf("lib/kernel_sm_%d%d.cubin", major, minor));
cuda_push_context();
result = cuModuleLoad(&cuModule, cubin.c_str());
if(result != CUDA_SUCCESS)
fprintf(stderr, "Failed loading CUDA kernel %s (%s).\n", cubin.c_str(), cuda_error_string(result));
CUresult result = cuModuleLoad(&cuModule, cubin.c_str());
if(cuda_error(result))
fprintf(stderr, "Failed loading CUDA kernel %s.\n", cubin.c_str());
cuda_pop_context();

149
intern/cycles/device/device_opencl.cpp
View File

@ -27,6 +27,7 @@
#include "util_map.h"
#include "util_math.h"
#include "util_md5.h"
#include "util_opencl.h"
#include "util_opengl.h"
#include "util_path.h"
@ -118,7 +119,7 @@ public:
void opencl_assert(cl_int err)
{
if(err != CL_SUCCESS) {
printf("error (%d): %s\n", err, opencl_error_string(err));
fprintf(stderr, "OpenCL error (%d): %s\n", err, opencl_error_string(err));
#ifndef NDEBUG
abort();
#endif
@ -157,7 +158,7 @@ public:
cpPlatform = platform_ids[0]; /* todo: pick specified platform && device */
ciErr = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_ALL, 1, &cdDevice, NULL);
ciErr = clGetDeviceIDs(cpPlatform, CL_DEVICE_TYPE_GPU|CL_DEVICE_TYPE_ACCELERATOR, 1, &cdDevice, NULL);
if(opencl_error(ciErr))
return;
@ -208,38 +209,67 @@ public:
return true;
}
bool load_kernels()
bool load_binary(const string& kernel_path, const string& clbin)
{
/* verify if device was initialized */
if(!device_initialized)
/* read binary into memory */
vector<uint8_t> binary;
if(!path_read_binary(clbin, binary)) {
fprintf(stderr, "OpenCL failed to read cached binary %s.\n", clbin.c_str());
return false;
}
/* create program */
cl_int status;
size_t size = binary.size();
const uint8_t *bytes = &binary[0];
cpProgram = clCreateProgramWithBinary(cxContext, 1, &cdDevice,
&size, &bytes, &status, &ciErr);
if(opencl_error(status) || opencl_error(ciErr)) {
fprintf(stderr, "OpenCL failed create program from cached binary %s.\n", clbin.c_str());
return false;
}
if(!build_kernel(kernel_path))
return false;
/* verify we have right opencl version */
if(!opencl_version_check())
return true;
}
bool save_binary(const string& clbin)
{
size_t size = 0;
clGetProgramInfo(cpProgram, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &size, NULL);
if(!size)
return false;
/* we compile kernels consisting of many files. unfortunately opencl
kernel caches do not seem to recognize changes in included files.
so we force recompile on changes by adding the md5 hash of all files */
string kernel_path = path_get("kernel");
string kernel_md5 = path_files_md5_hash(kernel_path);
vector<uint8_t> binary(size);
uint8_t *bytes = &binary[0];
string source = "#include \"kernel.cl\" // " + kernel_md5 + "\n";
size_t source_len = source.size();
const char *source_str = source.c_str();
clGetProgramInfo(cpProgram, CL_PROGRAM_BINARIES, sizeof(uint8_t*), &bytes, NULL);
if(!path_write_binary(clbin, binary)) {
fprintf(stderr, "OpenCL failed to write cached binary %s.\n", clbin.c_str());
return false;
}
return true;
}
bool build_kernel(const string& kernel_path)
{
string build_options = "";
build_options += "-I " + kernel_path + ""; /* todo: escape path */
build_options += " -cl-fast-relaxed-math -cl-strict-aliasing";
cpProgram = clCreateProgramWithSource(cxContext, 1, &source_str, &source_len, &ciErr);
if(opencl_error(ciErr))
return false;
ciErr = clBuildProgram(cpProgram, 0, NULL, build_options.c_str(), NULL, NULL);
if(ciErr != CL_SUCCESS) {
/* show build errors */
char *build_log;
size_t ret_val_size;
@ -256,6 +286,87 @@ public:
return false;
}
return true;
}
bool compile_kernel(const string& kernel_path, const string& kernel_md5)
{
/* we compile kernels consisting of many files. unfortunately opencl
kernel caches do not seem to recognize changes in included files.
so we force recompile on changes by adding the md5 hash of all files */
string source = "#include \"kernel.cl\" // " + kernel_md5 + "\n";
size_t source_len = source.size();
const char *source_str = source.c_str();
cpProgram = clCreateProgramWithSource(cxContext, 1, &source_str, &source_len, &ciErr);
if(opencl_error(ciErr))
return false;
double starttime = time_dt();
printf("Compiling OpenCL kernel ...\n");
if(!build_kernel(kernel_path))
return false;
printf("Kernel compilation finished in %.2lfs.\n", time_dt() - starttime);
return true;
}
string device_md5_hash()
{
MD5Hash md5;
char version[256], driver[256], name[256], vendor[256];
clGetPlatformInfo(cpPlatform, CL_PLATFORM_VENDOR, sizeof(vendor), &vendor, NULL);
clGetDeviceInfo(cdDevice, CL_DEVICE_VERSION, sizeof(version), &version, NULL);
clGetDeviceInfo(cdDevice, CL_DEVICE_NAME, sizeof(name), &name, NULL);
clGetDeviceInfo(cdDevice, CL_DRIVER_VERSION, sizeof(driver), &driver, NULL);
md5.append((uint8_t*)vendor, strlen(vendor));
md5.append((uint8_t*)version, strlen(version));
md5.append((uint8_t*)name, strlen(name));
md5.append((uint8_t*)driver, strlen(driver));
return md5.get_hex();
}
bool load_kernels()
{
/* verify if device was initialized */
if(!device_initialized) {
fprintf(stderr, "OpenCL: failed to initialize device.\n");
return false;
}
/* verify we have right opencl version */
if(!opencl_version_check())
return false;
/* md5 hash to detect changes */
string kernel_path = path_get("kernel");
string kernel_md5 = path_files_md5_hash(kernel_path);
string device_md5 = device_md5_hash();
/* try to use cache binary */
string clbin = string_printf("cycles_kernel_%s_%s.clbin", device_md5.c_str(), kernel_md5.c_str());;
clbin = path_user_get(path_join("cache", clbin));
if(path_exists(clbin)) {
/* if exists already, try use it */
if(!load_binary(kernel_path, clbin))
return false;
}
else {
/* compile kernel */
if(!compile_kernel(kernel_path, kernel_md5))
return false;
/* save binary for reuse */
save_binary(clbin);
}
/* find kernels */
ckPathTraceKernel = clCreateKernel(cpProgram, "kernel_ocl_path_trace", &ciErr);
if(opencl_error(ciErr))

1
intern/cycles/kernel/CMakeLists.txt
View File

@ -132,6 +132,7 @@ endif()
#delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${kernel_preprocessed}" ${CYCLES_INSTALL_PATH}/kernel)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernel.cl" ${CYCLES_INSTALL_PATH}/kernel)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "kernel.cu" ${CYCLES_INSTALL_PATH}/kernel)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${headers}" ${CYCLES_INSTALL_PATH}/kernel)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${svm_headers}" ${CYCLES_INSTALL_PATH}/kernel/svm)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${util_headers}" ${CYCLES_INSTALL_PATH}/kernel)

6
intern/cycles/kernel/kernel_compat_cuda.h
View File

@ -21,14 +21,14 @@
#define __KERNEL_GPU__
#define __KERNEL_CUDA__
#define CCL_NAMESPACE_BEGIN
#define CCL_NAMESPACE_END
#include <cuda.h>
#include <float.h>
#include "util_types.h"
CCL_NAMESPACE_BEGIN
/* Qualifier wrappers for different names on different devices */
#define __device __device__ __inline__
@ -60,7 +60,5 @@ typedef texture<uchar4, 2, cudaReadModeNormalizedFloat> texture_image_uchar4;
#define kernel_data __data
CCL_NAMESPACE_END
#endif /* __KERNEL_COMPAT_CUDA_H__ */

2
intern/cycles/render/session.cpp
View File

@ -384,7 +384,7 @@ void Session::run_cpu()
void Session::run()
{
/* load kernels */
progress.set_status("Loading render kernels (may take a few minutes)");
progress.set_status("Loading render kernels (may take a few minutes the first time)");
if(!device->load_kernels()) {
progress.set_status("Failed loading render kernel, see console for errors");

23
intern/cycles/util/util_cuda.cpp
View File

@ -16,6 +16,8 @@
* Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
*/
#include <stdlib.h>
#include "util_cuda.h"
#include "util_debug.h"
#include "util_dynlib.h"
@ -371,6 +373,11 @@ bool cuLibraryInit()
/* cuda 4.0 */
CUDA_LIBRARY_FIND(cuCtxSetCurrent);
#ifndef WITH_CUDA_BINARIES
if(cuCompilerPath() == "")
return false;
#endif
/* success */
result = true;
@ -379,13 +386,23 @@ bool cuLibraryInit()
string cuCompilerPath()
{
/* todo: better nvcc detection */
#ifdef _WIN32
string nvcc = "C:/CUDA/bin/nvcc.exe";
const char *defaultpath = "C:/CUDA/bin";
const char *executable = "nvcc.exe";
#else
string nvcc = "/usr/local/cuda/bin/nvcc";
const char *defaultpath = "/usr/local/cuda/bin";
const char *executable = "nvcc";
#endif
const char *binpath = getenv("CUDA_BIN_PATH");
string nvcc;
if(binpath)
nvcc = path_join(binpath, executable);
else
nvcc = path_join(defaultpath, executable);
return (path_exists(nvcc))? nvcc: "";
}

57
intern/cycles/util/util_path.cpp
View File

@ -24,6 +24,8 @@
#include <OpenImageIO/sysutil.h>
OIIO_NAMESPACE_USING
#include <stdio.h>
#define BOOST_FILESYSTEM_VERSION 2
#include <boost/filesystem.hpp>
@ -32,10 +34,12 @@ OIIO_NAMESPACE_USING
CCL_NAMESPACE_BEGIN
static string cached_path = "";
static string cached_user_path = "";
void path_init(const string& path)
void path_init(const string& path, const string& user_path)
{
cached_path = path;
cached_user_path = user_path;
}
string path_get(const string& sub)
@ -46,6 +50,14 @@ string path_get(const string& sub)
return path_join(cached_path, sub);
}
string path_user_get(const string& sub)
{
if(cached_user_path == "")
cached_user_path = path_dirname(Sysutil::this_program_path());
return path_join(cached_user_path, sub);
}
string path_filename(const string& path)
{
return boost::filesystem::path(path).filename();
@ -97,5 +109,48 @@ string path_files_md5_hash(const string& dir)
return hash.get_hex();
}
bool path_write_binary(const string& path, const vector<uint8_t>& binary)
{
/* write binary file from memory */
boost::filesystem::create_directories(path_dirname(path));
FILE *f = fopen(path.c_str(), "wb");
if(!f)
return false;
if(binary.size() > 0)
fwrite(&binary[0], sizeof(uint8_t), binary.size(), f);
fclose(f);
return true;
}
bool path_read_binary(const string& path, vector<uint8_t>& binary)
{
binary.resize(boost::filesystem::file_size(path));
/* read binary file into memory */
FILE *f = fopen(path.c_str(), "rb");
if(!f)
return false;
if(binary.size() == 0) {
fclose(f);
return false;
}
if(fread(&binary[0], sizeof(uint8_t), binary.size(), f) != binary.size()) {
fclose(f);
return false;
}
fclose(f);
return true;
}
CCL_NAMESPACE_END

7
intern/cycles/util/util_path.h
View File

@ -25,11 +25,13 @@
* then makes all paths relative to that. */
#include "util_string.h"
#include "util_vector.h"
CCL_NAMESPACE_BEGIN
void path_init(const string& path = "");
void path_init(const string& path = "", const string& user_path = "");
string path_get(const string& sub = "");
string path_user_get(const string& sub = "");
string path_filename(const string& path);
string path_dirname(const string& path);
@ -39,6 +41,9 @@ string path_escape(const string& path);
bool path_exists(const string& path);
string path_files_md5_hash(const string& dir);
bool path_write_binary(const string& path, const vector<uint8_t>& binary);
bool path_read_binary(const string& path, vector<uint8_t>& binary);
CCL_NAMESPACE_END
#endif