blender
/
blender
129
398
Fork 572
Code Issues 6.2k Pull Requests 748 Projects 19 Wiki Activity

Cycles: disable NanoVDB for AMD OpenCL 

It is causing issue with AMD OpenCL drivers, due to a potential driver bug.

Ref T84461
This commit is contained in:
Brecht Van Lommel 2021-03-29 22:58:19 +02:00
parent 661e6e0966
commit 91c44fe885
Notes: blender-bot 2023-04-19 22:54:54 +02:00 
Referenced by issue #87724, Segmentation Fault when rendering with Cycles on 2.92
Referenced by issue #87471, Volume objects not rendered on Cycles with GPU and AMD OpenCL
Referenced by issue #86610, Blender Crashing when rendering with Opencl GPU on Ubuntu.
Referenced by issue #86248, Crash using GPU Compute in Cycles on AMD in 2.92 and 2.93 but not 2.91
Referenced by issue #86182, Blender crashes with cycles and GPU rendering with AMD cards
Referenced by issue #85698, Blender crashes when rendering with OpenCL
Referenced by issue #85099, Cycles GPU render crashes in 2.92, 2.93 on Ubuntu with rx580
Referenced by issue #84579, OpenCL error: CL_INVALID_KERNEL_ARGS in clEnqueueNDRangeKernel()
Referenced by issue #84461, Cycles: Split kernel error with OpenCL
17 changed files with 143 additions and 98 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
intern/cycles/blender/blender_image.cpp
View File

@ -29,7 +29,7 @@ BlenderImageLoader::BlenderImageLoader(BL::Image b_image, int frame)
{
}
bool BlenderImageLoader::load_metadata(ImageMetaData &metadata)
bool BlenderImageLoader::load_metadata(const ImageDeviceFeatures &, ImageMetaData &metadata)
{
metadata.width = b_image.size()[0];
metadata.height = b_image.size()[1];
@ -171,7 +171,7 @@ BlenderPointDensityLoader::BlenderPointDensityLoader(BL::Depsgraph b_depsgraph,
{
}
bool BlenderPointDensityLoader::load_metadata(ImageMetaData &metadata)
bool BlenderPointDensityLoader::load_metadata(const ImageDeviceFeatures &, ImageMetaData &metadata)
{
metadata.channels = 4;
metadata.width = b_node.resolution();

4
intern/cycles/blender/blender_image.h
View File

@ -27,7 +27,7 @@ class BlenderImageLoader : public ImageLoader {
public:
BlenderImageLoader(BL::Image b_image, int frame);
bool load_metadata(ImageMetaData &metadata) override;
bool load_metadata(const ImageDeviceFeatures &features, ImageMetaData &metadata) override;
bool load_pixels(const ImageMetaData &metadata,
void *pixels,
const size_t pixels_size,
@ -44,7 +44,7 @@ class BlenderPointDensityLoader : public ImageLoader {
public:
BlenderPointDensityLoader(BL::Depsgraph depsgraph, BL::ShaderNodeTexPointDensity b_node);
bool load_metadata(ImageMetaData &metadata) override;
bool load_metadata(const ImageDeviceFeatures &features, ImageMetaData &metadata) override;
bool load_pixels(const ImageMetaData &metadata,
void *pixels,
const size_t pixels_size,

2
intern/cycles/blender/blender_volume.cpp
View File

@ -41,7 +41,7 @@ class BlenderSmokeLoader : public ImageLoader {
mesh_texture_space(b_mesh, texspace_loc, texspace_size);
}
bool load_metadata(ImageMetaData &metadata) override
bool load_metadata(const ImageDeviceFeatures &, ImageMetaData &metadata) override
{
if (!b_domain) {
return false;

2
intern/cycles/device/device.cpp
View File

@ -619,6 +619,7 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
info.num = 0;
info.has_half_images = true;
info.has_nanovdb = true;
info.has_volume_decoupled = true;
info.has_branched_path = true;
info.has_adaptive_stop_per_sample = true;
@ -665,6 +666,7 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
/* Accumulate device info. */
info.has_half_images &= device.has_half_images;
info.has_nanovdb &= device.has_nanovdb;
info.has_volume_decoupled &= device.has_volume_decoupled;
info.has_branched_path &= device.has_branched_path;
info.has_adaptive_stop_per_sample &= device.has_adaptive_stop_per_sample;

2
intern/cycles/device/device.h
View File

@ -78,6 +78,7 @@ class DeviceInfo {
int num;
bool display_device; /* GPU is used as a display device. */
bool has_half_images; /* Support half-float textures. */
bool has_nanovdb; /* Support NanoVDB volumes. */
bool has_volume_decoupled; /* Decoupled volume shading. */
bool has_branched_path; /* Supports branched path tracing. */
bool has_adaptive_stop_per_sample; /* Per-sample adaptive sampling stopping. */
@ -99,6 +100,7 @@ class DeviceInfo {
cpu_threads = 0;
display_device = false;
has_half_images = false;
has_nanovdb = false;
has_volume_decoupled = false;
has_branched_path = true;
has_adaptive_stop_per_sample = false;

1
intern/cycles/device/device_cpu.cpp
View File

@ -1654,6 +1654,7 @@ void device_cpu_info(vector<DeviceInfo> &devices)
info.has_adaptive_stop_per_sample = true;
info.has_osl = true;
info.has_half_images = true;
info.has_nanovdb = true;
info.has_profiling = true;
info.denoisers = DENOISER_NLM;
if (openimagedenoise_supported()) {

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

@ -128,6 +128,7 @@ void device_cuda_info(vector<DeviceInfo> &devices)
info.num = num;
info.has_half_images = (major >= 3);
info.has_nanovdb = true;
info.has_volume_decoupled = false;
info.has_adaptive_stop_per_sample = false;
info.denoisers = DENOISER_NLM;

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

@ -126,6 +126,9 @@ void device_opencl_info(vector<DeviceInfo> &devices)
/* Check OpenCL extensions */
info.has_half_images = platform_device.device_extensions.find("cl_khr_fp16") != string::npos;
/* Disabled for now due to apparent AMD driver bug. */
info.has_nanovdb = platform_name != "AMD Accelerated Parallel Processing";
devices.push_back(info);
num_devices++;
}

4
intern/cycles/device/opencl/device_opencl_impl.cpp
View File

@ -2036,7 +2036,9 @@ string OpenCLDevice::kernel_build_options(const string *debug_src)
# endif
# ifdef WITH_NANOVDB
build_options += "-DWITH_NANOVDB ";
if (info.has_nanovdb) {
build_options += "-DWITH_NANOVDB ";
}
# endif
return build_options;

18
intern/cycles/render/image.cpp
View File

@ -303,7 +303,8 @@ ImageManager::ImageManager(const DeviceInfo &info)
animation_frame = 0;
/* Set image limits */
has_half_images = info.has_half_images;
features.has_half_float = info.has_half_images;
features.has_nanovdb = info.has_nanovdb;
}
ImageManager::~ImageManager()
@ -347,7 +348,7 @@ void ImageManager::load_image_metadata(Image *img)
metadata = ImageMetaData();
metadata.colorspace = img->params.colorspace;
if (img->loader->load_metadata(metadata)) {
if (img->loader->load_metadata(features, metadata)) {
assert(metadata.type != IMAGE_DATA_NUM_TYPES);
}
else {
@ -356,15 +357,10 @@ void ImageManager::load_image_metadata(Image *img)
metadata.detect_colorspace();
/* No half textures on OpenCL, use full float instead. */
if (!has_half_images) {
if (metadata.type == IMAGE_DATA_TYPE_HALF4) {
metadata.type = IMAGE_DATA_TYPE_FLOAT4;
}
else if (metadata.type == IMAGE_DATA_TYPE_HALF) {
metadata.type = IMAGE_DATA_TYPE_FLOAT;
}
}
assert(features.has_half_float ||
(metadata.type != IMAGE_DATA_TYPE_HALF4 && metadata.type != IMAGE_DATA_TYPE_HALF));
assert(features.has_nanovdb || (metadata.type != IMAGE_DATA_TYPE_NANOVDB_FLOAT ||
metadata.type != IMAGE_DATA_TYPE_NANOVDB_FLOAT3));
img->need_metadata = false;
}

12
intern/cycles/render/image.h
View File

@ -97,6 +97,13 @@ class ImageMetaData {
void detect_colorspace();
};
/* Information about supported features that Image loaders can use. */
class ImageDeviceFeatures {
public:
bool has_half_float;
bool has_nanovdb;
};
/* Image loader base class, that can be subclassed to load image data
* from custom sources (file, memory, procedurally generated, etc). */
class ImageLoader {
@ -105,7 +112,7 @@ class ImageLoader {
virtual ~ImageLoader(){};
/* Load metadata without actual image yet, should be fast. */
virtual bool load_metadata(ImageMetaData &metadata) = 0;
virtual bool load_metadata(const ImageDeviceFeatures &features, ImageMetaData &metadata) = 0;
/* Load actual image contents. */
virtual bool load_pixels(const ImageMetaData &metadata,
@ -212,7 +219,8 @@ class ImageManager {
private:
bool need_update_;
bool has_half_images;
ImageDeviceFeatures features;
thread_mutex device_mutex;
thread_mutex images_mutex;

4
intern/cycles/render/image_oiio.cpp
View File

@ -30,7 +30,7 @@ OIIOImageLoader::~OIIOImageLoader()
{
}
bool OIIOImageLoader::load_metadata(ImageMetaData &metadata)
bool OIIOImageLoader::load_metadata(const ImageDeviceFeatures &features, ImageMetaData &metadata)
{
/* Perform preliminary checks, with meaningful logging. */
if (!path_exists(filepath.string())) {
@ -76,7 +76,7 @@ bool OIIOImageLoader::load_metadata(ImageMetaData &metadata)
}
/* check if it's half float */
if (spec.format == TypeDesc::HALF) {
if (spec.format == TypeDesc::HALF && features.has_half_float) {
is_half = true;
}

2
intern/cycles/render/image_oiio.h
View File

@ -26,7 +26,7 @@ class OIIOImageLoader : public ImageLoader {
OIIOImageLoader(const string &filepath);
~OIIOImageLoader();
bool load_metadata(ImageMetaData &metadata) override;
bool load_metadata(const ImageDeviceFeatures &features, ImageMetaData &metadata) override;
bool load_pixels(const ImageMetaData &metadata,
void *pixels,

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

@ -40,7 +40,7 @@ SkyLoader::SkyLoader(float sun_elevation,
SkyLoader::~SkyLoader(){};
bool SkyLoader::load_metadata(ImageMetaData &metadata)
bool SkyLoader::load_metadata(const ImageDeviceFeatures &, ImageMetaData &metadata)
{
metadata.width = 512;
metadata.height = 128;

2
intern/cycles/render/image_sky.h
View File

@ -34,7 +34,7 @@ class SkyLoader : public ImageLoader {
float ozone_density);
~SkyLoader();
bool load_metadata(ImageMetaData &metadata) override;
bool load_metadata(const ImageDeviceFeatures &features, ImageMetaData &metadata) override;
bool load_pixels(const ImageMetaData &metadata,
void *pixels,

175
intern/cycles/render/image_vdb.cpp
View File

@ -34,7 +34,7 @@ VDBImageLoader::~VDBImageLoader()
{
}
bool VDBImageLoader::load_metadata(ImageMetaData &metadata)
bool VDBImageLoader::load_metadata(const ImageDeviceFeatures &features, ImageMetaData &metadata)
{
#ifdef WITH_OPENVDB
if (!grid) {
@ -56,55 +56,71 @@ bool VDBImageLoader::load_metadata(ImageMetaData &metadata)
if (grid->isType<openvdb::FloatGrid>()) {
metadata.channels = 1;
# ifdef WITH_NANOVDB
nanogrid = nanovdb::openToNanoVDB(*openvdb::gridConstPtrCast<openvdb::FloatGrid>(grid));
if (features.has_nanovdb) {
nanogrid = nanovdb::openToNanoVDB(*openvdb::gridConstPtrCast<openvdb::FloatGrid>(grid));
}
# endif
}
else if (grid->isType<openvdb::Vec3fGrid>()) {
metadata.channels = 3;
# ifdef WITH_NANOVDB
nanogrid = nanovdb::openToNanoVDB(*openvdb::gridConstPtrCast<openvdb::Vec3fGrid>(grid));
if (features.has_nanovdb) {
nanogrid = nanovdb::openToNanoVDB(*openvdb::gridConstPtrCast<openvdb::Vec3fGrid>(grid));
}
# endif
}
else if (grid->isType<openvdb::BoolGrid>()) {
metadata.channels = 1;
# ifdef WITH_NANOVDB
nanogrid = nanovdb::openToNanoVDB(
openvdb::FloatGrid(*openvdb::gridConstPtrCast<openvdb::BoolGrid>(grid)));
if (features.has_nanovdb) {
nanogrid = nanovdb::openToNanoVDB(
openvdb::FloatGrid(*openvdb::gridConstPtrCast<openvdb::BoolGrid>(grid)));
}
# endif
}
else if (grid->isType<openvdb::DoubleGrid>()) {
metadata.channels = 1;
# ifdef WITH_NANOVDB
nanogrid = nanovdb::openToNanoVDB(
openvdb::FloatGrid(*openvdb::gridConstPtrCast<openvdb::DoubleGrid>(grid)));
if (features.has_nanovdb) {
nanogrid = nanovdb::openToNanoVDB(
openvdb::FloatGrid(*openvdb::gridConstPtrCast<openvdb::DoubleGrid>(grid)));
}
# endif
}
else if (grid->isType<openvdb::Int32Grid>()) {
metadata.channels = 1;
# ifdef WITH_NANOVDB
nanogrid = nanovdb::openToNanoVDB(
openvdb::FloatGrid(*openvdb::gridConstPtrCast<openvdb::Int32Grid>(grid)));
if (features.has_nanovdb) {
nanogrid = nanovdb::openToNanoVDB(
openvdb::FloatGrid(*openvdb::gridConstPtrCast<openvdb::Int32Grid>(grid)));
}
# endif
}
else if (grid->isType<openvdb::Int64Grid>()) {
metadata.channels = 1;
# ifdef WITH_NANOVDB
nanogrid = nanovdb::openToNanoVDB(
openvdb::FloatGrid(*openvdb::gridConstPtrCast<openvdb::Int64Grid>(grid)));
if (features.has_nanovdb) {
nanogrid = nanovdb::openToNanoVDB(
openvdb::FloatGrid(*openvdb::gridConstPtrCast<openvdb::Int64Grid>(grid)));
}
# endif
}
else if (grid->isType<openvdb::Vec3IGrid>()) {
metadata.channels = 3;
# ifdef WITH_NANOVDB
nanogrid = nanovdb::openToNanoVDB(
openvdb::Vec3fGrid(*openvdb::gridConstPtrCast<openvdb::Vec3IGrid>(grid)));
if (features.has_nanovdb) {
nanogrid = nanovdb::openToNanoVDB(
openvdb::Vec3fGrid(*openvdb::gridConstPtrCast<openvdb::Vec3IGrid>(grid)));
}
# endif
}
else if (grid->isType<openvdb::Vec3dGrid>()) {
metadata.channels = 3;
# ifdef WITH_NANOVDB
nanogrid = nanovdb::openToNanoVDB(
openvdb::Vec3fGrid(*openvdb::gridConstPtrCast<openvdb::Vec3dGrid>(grid)));
if (features.has_nanovdb) {
nanogrid = nanovdb::openToNanoVDB(
openvdb::Vec3fGrid(*openvdb::gridConstPtrCast<openvdb::Vec3dGrid>(grid)));
}
# endif
}
else if (grid->isType<openvdb::MaskGrid>()) {
@ -118,21 +134,25 @@ bool VDBImageLoader::load_metadata(ImageMetaData &metadata)
}
# ifdef WITH_NANOVDB
metadata.byte_size = nanogrid.size();
if (metadata.channels == 1) {
metadata.type = IMAGE_DATA_TYPE_NANOVDB_FLOAT;
}
else {
metadata.type = IMAGE_DATA_TYPE_NANOVDB_FLOAT3;
}
# else
if (metadata.channels == 1) {
metadata.type = IMAGE_DATA_TYPE_FLOAT;
}
else {
metadata.type = IMAGE_DATA_TYPE_FLOAT4;
if (nanogrid) {
metadata.byte_size = nanogrid.size();
if (metadata.channels == 1) {
metadata.type = IMAGE_DATA_TYPE_NANOVDB_FLOAT;
}
else {
metadata.type = IMAGE_DATA_TYPE_NANOVDB_FLOAT3;
}
}
else
# endif
{
if (metadata.channels == 1) {
metadata.type = IMAGE_DATA_TYPE_FLOAT;
}
else {
metadata.type = IMAGE_DATA_TYPE_FLOAT4;
}
}
/* Set transform from object space to voxel index. */
openvdb::math::Mat4f grid_matrix = grid->transform().baseMap()->getAffineMap()->getMat4();
@ -143,13 +163,18 @@ bool VDBImageLoader::load_metadata(ImageMetaData &metadata)
}
}
Transform texture_to_index;
# ifdef WITH_NANOVDB
Transform texture_to_index = transform_identity();
# else
openvdb::Coord min = bbox.min();
Transform texture_to_index = transform_translate(min.x(), min.y(), min.z()) *
transform_scale(dim.x(), dim.y(), dim.z());
if (nanogrid) {
texture_to_index = transform_identity();
}
else
# endif
{
openvdb::Coord min = bbox.min();
texture_to_index = transform_translate(min.x(), min.y(), min.z()) *
transform_scale(dim.x(), dim.y(), dim.z());
}
metadata.transform_3d = transform_inverse(index_to_object * texture_to_index);
metadata.use_transform_3d = true;
@ -165,48 +190,52 @@ bool VDBImageLoader::load_pixels(const ImageMetaData &, void *pixels, const size
{
#ifdef WITH_OPENVDB
# ifdef WITH_NANOVDB
memcpy(pixels, nanogrid.data(), nanogrid.size());
# else
if (grid->isType<openvdb::FloatGrid>()) {
openvdb::tools::Dense<float, openvdb::tools::LayoutXYZ> dense(bbox, (float *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::FloatGrid>(grid), dense);
}
else if (grid->isType<openvdb::Vec3fGrid>()) {
openvdb::tools::Dense<openvdb::Vec3f, openvdb::tools::LayoutXYZ> dense(
bbox, (openvdb::Vec3f *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::Vec3fGrid>(grid), dense);
}
else if (grid->isType<openvdb::BoolGrid>()) {
openvdb::tools::Dense<float, openvdb::tools::LayoutXYZ> dense(bbox, (float *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::BoolGrid>(grid), dense);
}
else if (grid->isType<openvdb::DoubleGrid>()) {
openvdb::tools::Dense<float, openvdb::tools::LayoutXYZ> dense(bbox, (float *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::DoubleGrid>(grid), dense);
}
else if (grid->isType<openvdb::Int32Grid>()) {
openvdb::tools::Dense<float, openvdb::tools::LayoutXYZ> dense(bbox, (float *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::Int32Grid>(grid), dense);
}
else if (grid->isType<openvdb::Int64Grid>()) {
openvdb::tools::Dense<float, openvdb::tools::LayoutXYZ> dense(bbox, (float *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::Int64Grid>(grid), dense);
}
else if (grid->isType<openvdb::Vec3IGrid>()) {
openvdb::tools::Dense<openvdb::Vec3f, openvdb::tools::LayoutXYZ> dense(
bbox, (openvdb::Vec3f *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::Vec3IGrid>(grid), dense);
}
else if (grid->isType<openvdb::Vec3dGrid>()) {
openvdb::tools::Dense<openvdb::Vec3f, openvdb::tools::LayoutXYZ> dense(
bbox, (openvdb::Vec3f *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::Vec3dGrid>(grid), dense);
}
else if (grid->isType<openvdb::MaskGrid>()) {
openvdb::tools::Dense<float, openvdb::tools::LayoutXYZ> dense(bbox, (float *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::MaskGrid>(grid), dense);
if (nanogrid) {
memcpy(pixels, nanogrid.data(), nanogrid.size());
}
else
# endif
{
if (grid->isType<openvdb::FloatGrid>()) {
openvdb::tools::Dense<float, openvdb::tools::LayoutXYZ> dense(bbox, (float *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::FloatGrid>(grid), dense);
}
else if (grid->isType<openvdb::Vec3fGrid>()) {
openvdb::tools::Dense<openvdb::Vec3f, openvdb::tools::LayoutXYZ> dense(
bbox, (openvdb::Vec3f *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::Vec3fGrid>(grid), dense);
}
else if (grid->isType<openvdb::BoolGrid>()) {
openvdb::tools::Dense<float, openvdb::tools::LayoutXYZ> dense(bbox, (float *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::BoolGrid>(grid), dense);
}
else if (grid->isType<openvdb::DoubleGrid>()) {
openvdb::tools::Dense<float, openvdb::tools::LayoutXYZ> dense(bbox, (float *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::DoubleGrid>(grid), dense);
}
else if (grid->isType<openvdb::Int32Grid>()) {
openvdb::tools::Dense<float, openvdb::tools::LayoutXYZ> dense(bbox, (float *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::Int32Grid>(grid), dense);
}
else if (grid->isType<openvdb::Int64Grid>()) {
openvdb::tools::Dense<float, openvdb::tools::LayoutXYZ> dense(bbox, (float *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::Int64Grid>(grid), dense);
}
else if (grid->isType<openvdb::Vec3IGrid>()) {
openvdb::tools::Dense<openvdb::Vec3f, openvdb::tools::LayoutXYZ> dense(
bbox, (openvdb::Vec3f *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::Vec3IGrid>(grid), dense);
}
else if (grid->isType<openvdb::Vec3dGrid>()) {
openvdb::tools::Dense<openvdb::Vec3f, openvdb::tools::LayoutXYZ> dense(
bbox, (openvdb::Vec3f *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::Vec3dGrid>(grid), dense);
}
else if (grid->isType<openvdb::MaskGrid>()) {
openvdb::tools::Dense<float, openvdb::tools::LayoutXYZ> dense(bbox, (float *)pixels);
openvdb::tools::copyToDense(*openvdb::gridConstPtrCast<openvdb::MaskGrid>(grid), dense);
}
}
return true;
#else
(void)pixels;

3
intern/cycles/render/image_vdb.h
View File

@ -33,7 +33,8 @@ class VDBImageLoader : public ImageLoader {
VDBImageLoader(const string &grid_name);
~VDBImageLoader();
virtual bool load_metadata(ImageMetaData &metadata) override;
virtual bool load_metadata(const ImageDeviceFeatures &features,
ImageMetaData &metadata) override;
virtual bool load_pixels(const ImageMetaData &metadata,
void *pixels,