Compare commits

...

520 Commits

Author SHA1 Message Date
f240ac1673 Improve formating 2021-11-24 21:41:25 +01:00
a590474f4c Fix compilation & dependency depth 2021-11-24 20:09:36 +01:00
b2462b6f5c Make shaders sources from draw included in the dependency library. 2021-11-24 19:46:00 +01:00
21ee89c52f Fix compilation issues on MSVC and a bug in builder 2021-11-24 18:56:51 +01:00
2c95da88aa GPUShaderDependency: Initial Commit
This is a prototype to support `#include` directive inside glsl sources.
The sources are aggregated into a list before being translated to byte_array.
This list is used to generate a mapping between the filename and the
associated byte_array. For each byte_array, we search for any `#include`
and store the file to merge.

At runtime, for one input filename we concatenate all byte_arrays that
are needed following the include order and avoiding double include.

This is meant to evolve into a fully supported `#include` system.
2021-11-24 17:58:06 +01:00
7358a5aba2 GPUShaderDescriptor: Initial Commit
This is a first draft of what the Shader Descriptor system could be.

A shader descriptor provides a way to define shader structure, resources
and interfaces. This makes for a quick way to provide backend agnostic
binding informations while also making shader variations easy to declare.
2021-11-24 17:52:39 +01:00
9159295c3c Clang Tidy: ignore some passes that changed or were added in version 13
I get hundreds of clang-tidy errors without ignoring those passes right now.
To not forget about the passes, I added them to T78535.
2021-11-23 19:45:05 +01:00
62a04f7aa6 Cleanup: clang tidy
The parameter name was inconsistent with the declaration.
2021-11-23 19:38:22 +01:00
38a3819171 Merge branch 'blender-v3.0-release' 2021-11-23 19:04:44 +01:00
ea93e5df6c Asset: Merge asset library/list refresh operators
In rBdcdbaf89bd11, I introduced a new operator
(`file.asset_library_refresh()`) to handle Asset Browser refreshing more
separate from File Browser refreshing. However, there already was
`asset.asset_list_refresh()`, which at this point only works for asset
view templates, but was intended to cover the Asset Browser case in
future too. This would happen once the Asset Browser uses the asset list
design of the asset view template.

So rather than having two operators for refreshing asset library data,
have one that just handles both cases, until they converge into one.
This avoids changes to the Python API in future (deprecating/changing
operators).

Differential Revision: https://developer.blender.org/D13239
2021-11-23 18:57:25 +01:00
60befc8f02 Clean-up: Fix BLI_rect.h collision with windows.h
windows.h `#defines rct1` as a number which is
problematic if we include `BLI_rect.h` after
`windows.h` .

by renaming `rct1/2` to `rct_a/b` we side step
the collision and straighten up the naming with
the functions directly above it.
2021-11-23 10:51:09 -07:00
62b50c612f Cleanup: Else after return, other simplifications
`std::stringstream` already returns a `std::string`, and there is no
particular reason to use short here instead of int.
2021-11-23 12:49:45 -05:00
c09e8a3590 Merge branch 'blender-v3.0-release' 2021-11-23 18:35:56 +01:00
c0a2b21744 Merge branch 'blender-v3.0-release' 2021-11-23 18:04:28 +01:00
b40e930ac7 Merge branch 'blender-v3.0-release' 2021-11-23 17:52:30 +01:00
e4986f92f3 Geometry Nodes: Node execution time overlay
Adds a new overlay called "Timings" to the Geometry Node editor.
This shows the node execution time in milliseconds above the node.
For group nodes and frames, the total time for all nodes inside
(recursively) is shown. Group output node shows the node tree total.
The code is prepared for easily adding new rows of information
to the box above the node in the future.

Differential Revision: https://developer.blender.org/D13256
2021-11-23 17:37:31 +01:00
fab39440e9 Cleanup: Simplify geometry node function names
With this commit, we no longer use the prefixes for every node type
function like `geo_node_translate_instances_`. They just added more
places to change when adding a new node, for no real benefit.

Differential Revision: https://developer.blender.org/D13337
2021-11-23 10:55:51 -05:00
a9eb4e6f59 Merge branch 'blender-v3.0-release' 2021-11-23 16:41:53 +01:00
db450c9320 Merge branch 'blender-v3.0-release' 2021-11-23 16:38:30 +01:00
71c80bd939 Merge branch 'blender-v3.0-release' 2021-11-23 16:32:13 +01:00
b716a771b4 Merge branch 'blender-v3.0-release' 2021-11-23 15:46:28 +01:00
fca8eb0185 Cleanup: Suppress clang-tidy warning. 2021-11-23 15:41:28 +01:00
2c2b79191f Merge branch 'blender-v3.0-release' 2021-11-23 15:41:09 +01:00
4b13dcaf02 Merge branch 'blender-v3.0-release' 2021-11-23 15:35:09 +01:00
8897e0aa8f Fix add-on Preferences using the .blend file icon, not the Blender logo
Intention of the icon is to mark add-ons that are official/bundled.
Doesn't make much sense to use the .blend file icon for that. It's
arguable if the Blender logo should be used for this, but the file icon
is definitely the wrong choice.
2021-11-23 15:30:28 +01:00
1df8abff25 Geometry Nodes: add namespace for every file
This puts all static functions in geometry node files into a new
namespace. This allows using unity build which can improve
compile times significantly (P2578).

* The name space name is derived from the file name. That makes
  it possible to write some tooling that checks the names later on.
  The file name extension (`cc`) is added to the namespace name as
  well. This also possibly simplifies tooling but also makes it more
  obvious that this namespace is specific to a file.
* In the register function of every node, I added a namespace alias
  `namespace file_ns = blender::nodes::node_geo_*_cc;`. This avoids
  some duplication of the file name and may also simplify tooling,
  because this line is easy to detect. The name `file_ns` stands for "file
  namespace" and also indicates that this namespace corresponds to
  the current file. In the beginning I used `node_ns` but `file_ns` is more
  generic which may make it more suitable when we want to use unity
  builds outside of the nodes modules in the future.
* Some node files contain code that is actually shared between
  different nodes. For now I left that code in the `blender::nodes`
  namespace and moved it to the top of the file (couldn't move it to
  the bottom in all cases, so I just moved it to the top everywhere).
  As a separate cleanup step, this shared code should actually be
  moved to a separate file.

Differential Revision: https://developer.blender.org/D13330
2021-11-23 14:56:01 +01:00
47276b8470 Geometry Nodes: reduce overhead when processing single values
Currently the geometry nodes evaluator always stores a field for every
type that supports it, even if it is just a single value. This results in a lot
of overhead when there are many sockets that just contain a single
value, which is often the case.

This introduces a new `ValueOrField<T>` type that is used by the geometry
nodes evaluator. Now a field will only be created when it is actually
necessary. See D13307 for more details. In extrem cases this can speed
up the evaluation 2-3x (those cases are probably never hit in practice
though, but it's good to get rid of unnecessary overhead nevertheless).

Differential Revision: https://developer.blender.org/D13307
2021-11-23 14:49:26 +01:00
0bedd5d14f Merge branch 'blender-v3.0-release' 2021-11-23 14:39:55 +01:00
d7b7cbb047 Merge branch 'blender-v3.0-release' 2021-11-23 14:36:57 +01:00
89b927a720 Cleanup: Silence compilation warning.
For now made DRW_notify_view_update_offscreen static.
2021-11-23 14:30:14 +01:00
fecdf9d44b Merge branch 'blender-v3.0-release'
Conflicts:
	source/blender/editors/transform/transform_generics.c
2021-11-23 10:17:24 -03:00
a6d1a2d3fc Merge branch 'blender-v3.0-release' 2021-11-23 14:08:53 +01:00
cf299bee80 Merge remote-tracking branch 'origin/blender-v3.0-release' 2021-11-23 13:02:25 +01:00
f392ce50c4 Merge remote-tracking branch 'origin/blender-v3.0-release' 2021-11-23 12:39:26 +01:00
605cdc4346 BKE LibLink/Append: Add mechanism for external code to loop over link/append context items.
Will be required for python's `bpy.data.libraries.load()` refactor.
2021-11-23 12:18:37 +01:00
0452a04f1a BKE link/append: Add optional blendfile handle to libraries.
This enables calling code to deal with the blendfile handle themselves,
BKE_blendfile_link then just borrows, uses this handle and does not
release it.

Needed e.g. for python's libcontext system to use new
BKE_blendfile_link_append code.

Part of T91414: Unify link/append between WM operators and BPY context
manager API, and cleanup usages of `BKE_library_make_local`.
2021-11-23 12:17:16 +01:00
6c8f73b220 Cleanup: remove unnecessary "extern" 2021-11-23 12:05:50 +01:00
b9b98448a6 Cleanup: add missing pragma once 2021-11-23 12:00:11 +01:00
fbb4a7eb43 BKE link/append: Add optional blendfile handle to libraries.
This enables calling code to deal with the blendfile handle themselves,
BKE_blendfile_link then just borrows, uses this handle and does not
release it.

Needed e.g. for python's libcontext system to use new
BKE_blendfile_link_append code.

Part of T91414: Unify link/append between WM operators and BPY context
manager API, and cleanup usages of `BKE_library_make_local`.
2021-11-23 10:38:51 +01:00
f657356062 Merge branch 'blender-v3.0-release' 2021-11-23 09:44:04 +01:00
28870a8f89 Cleanup: Use new CollectionRef::empty() method
Use the new CollectionRef::empty() method in all locations where
appropriate.

Differential Revision: https://developer.blender.org/D13276
2021-11-22 23:47:26 -08:00
59754ef0b2 Cleanup: silent clang-tidy warning NULL vs nullptr. 2021-11-23 08:39:26 +01:00
Peter Fog
34370d9fdf VSE: Add drag and drop handler for preview area
For some users, dropping assets into preview area may be more practical
due to space constraints or it may be just more intuitive.

Reviewed By: ISS

Differential Revision: https://developer.blender.org/D13311
2021-11-23 05:40:08 +01:00
Peter Fog
b42494cf6b VSE: Support drag and drop for datablocks
For using the Outliner and/or the Asset Browser as scene independent
tools to organize a/v source material is necessary for the users to be
able to drag and drop data blocks into the VSE. This was also an
unfulfilled design target for the Outliner Gsoc project.

Datablocks won't be used directly. Path to file will be passed to strip
add operator instead.

Reviewed By: ISS

Differential Revision: https://developer.blender.org/D13304
2021-11-23 05:39:10 +01:00
d1a4e043bd Merge branch 'blender-v3.0-release' 2021-11-23 00:57:15 +01:00
8600d4491f Fix: Const warning in editmesh_knife.c
Fixes a warning caused by freeing a const pointer.
This commit removes the const modifier.

Differential Revision: https://developer.blender.org/D13321
2021-11-22 23:40:21 +01:00
456d5e14b8 Merge branch 'blender-v3.0-release' 2021-11-22 21:26:39 +01:00
48c2b4012f Merge branch 'blender-v3.0-release' 2021-11-22 21:06:10 +01:00
29681f186e Fix T93283: Cycles render error with CUDA CPU + GPU after recent optimization
BVH2 triangle intersection was broken on the GPU since packed floats can't
be loaded directly into SSE. The better long term solution for performance
would be to build a BVH2 for GPU and Embree for CPU, similar to what we do
for OptiX.
2021-11-22 21:02:46 +01:00
Takahiro Shizuki
73b1ad1920 IME: Fix Multi-Window Duplicated First Character
Fix problem with duplicated initial character when initiating or
switching to new windows. This is done by updating our copies of state
and modes from the new window when it receives WM_IME_SETCONTEXT
message. This problem and fix are only for the Windows platform.
2021-11-22 10:44:34 -08:00
bfff9ca5f1 Fix confusing new Cycles UI terminology
* Rename "Auto Tiles" to "Use Tiling", it's not really automatic and
  confusing with the old auto tile size add-on.
* Rename "Adaptive" scrambling distance to "Automatic", to avoid confusion
  with adaptive sampling.
2021-11-22 10:44:34 -08:00
cc6bcb53b2 Geometry Nodes: Rename legacy node files
This will be useful to solve some issues with unity builds,
which compiles different files together to improve build times.
2021-11-22 13:16:54 -05:00
5ad4ca4e02 CurveEval: Add total_length() and total_control_point_size() methods
Add the following methods to the CurveEval class:
total_length() : returns the total length of the curve without needing to
                    allocate a new array
total_control_point_size() : returns the total number of control points without
                    needing to allocate a new array
2021-11-22 12:07:59 -06:00
0129178376 UI: Use a map for block name lookups
Use a map to speed up search for UI block names.
Time to redraw the node editor was decreased from
around 75-120ms to 40-70ms in a tree with many
Geometry Nodes.

Differential Revision: https://developer.blender.org/D13225
2021-11-22 18:23:54 +01:00
55c82d8380 Fix T84493 issue with selection after boolean.
According to Blender selection rules, selections should be flushed
to containing elements. Added an EDMB_select_flush() after edit
mode booleans or intersects are done. Hopefully this doesn't break
any scripts that might have been depending on the old (broken) behavior.
2021-11-22 11:47:46 -05:00
1706bf7780 Merge branch 'blender-v3.0-release' 2021-11-22 17:32:23 +01:00
059da44fbc BKE Link/Append: Use BLO's LibraryLink_Params.
This allows to reduce signature of several functions, and make it eaiser
to integrate more higher-level usages later on.

This should be a non-behavioral-change commit.

Part of T91414: Unify link/append between WM operators and BPY context
manager API, and cleanup usages of `BKE_library_make_local`.
2021-11-22 17:14:21 +01:00
d7cf7d4048 Link/Append: Move main linking code (from WM) to use new instantiation code in BKE.
This removes the last main usage of BLO's instantiation code.

Also required some limited refactoring of BKE_blendfile_link_append's
instantiation to make it more modular, and usable by both linking and
appending code paths.

NOTE: This should be a non-behavioral change commit.

Part of T91414: Unify link/append between WM operators and BPY context
manager API, and cleanup usages of `BKE_library_make_local`.
2021-11-22 16:55:36 +01:00
fe274d91a1 Link/Append: Move most of core link/append code from WM to new BKE_blendflie_link_append module.
This will allow to expose all those advanced features of the WM
operators to other parts of the code, like the python library context
manager, copy/paste code, etc.

This is expected to be a strictly no-behavioral-change commit.

Part of T91414: Unify link/append between WM operators and BPY context
manager API, and cleanup usages of `BKE_library_make_local`.

Maniphest Tasks: T91414

Differential Revision: https://developer.blender.org/D13222
2021-11-22 16:52:17 +01:00
9c2a4d158c Merge remote-tracking branch 'origin/blender-v3.0-release' 2021-11-22 16:47:50 +01:00
Jeroen Bakker
0624235900 Moviecache: Fix potential memory corruption.
`IMB_moviecache` is implemented as a singleton. When destructing the
singleton via `IMB_moviecache_destruct` it will not be created anymore
resulting inusage of unallocated memory and potentional memory
corruption.

When running blender this doesn't happen, but when creating images in
test cases the moviecache should be able to be recreated after it is
destroyed.

Reviewed By: sergey

Differential Revision: https://developer.blender.org/D13287
2021-11-22 15:52:30 +01:00
be876b8db6 Painting: Performance curve masks.
This patch separates the static-part from the dynamic-part when
generate brush masks. This makes the generation of brush masks 2-5 times
faster depending on the size of the brush.

More improvements can be done, this was just low hanging fruit.
2021-11-22 15:36:03 +01:00
178947dec3 Merge branch 'blender-v3.0-release' 2021-11-22 09:34:51 -05:00
77ddc6e350 Fix T93250: Crash in spreadsheet tree view after recent commit
It seems that update_from_old assumed there would be an old
tree view available in the old block. This works for the asset browser
because the tree is always drawn, but for the spreadsheet that depends
on having an active object, which isn't necessarily always true.

Differential Revision: https://developer.blender.org/D13301
2021-11-22 08:47:02 -05:00
7e8f9213e9 Merge branch 'blender-v3.0-release' 2021-11-22 13:49:37 +01:00
9b6f3d2d0a Cleanup: Removed unused code in curve mask creation.
Generating curve mask for 2d texture painting had some hard-coded
parameters that eventually weren't used in the algorithm (hardness and
rotation of the brush). This patch removes these parameters.
2021-11-22 11:22:58 +01:00
9bbb5f5a6a Painting: migrated curve mask generation to CPP.
Curve mask generation is done during 2d texture painting. There are some
performance issues in this part of the code. Before addressing those we
move the code to CPP.
2021-11-22 10:46:33 +01:00
29f6ec56e6 Cleanup: Make parameter const (BKE_brush_curve_strength_clamped). 2021-11-22 10:46:33 +01:00
31864a40ba Cleanup: use simple data member instead of callback
This really doesn't have to be a callback currently, since it is always
the same `CPPType` for a socket type.
2021-11-22 10:18:08 +01:00
c850189adf Cleanup: make naming more consistent 2021-11-22 09:48:36 +01:00
db20837c3a Merge branch 'blender-v3.0-release' 2021-11-22 09:32:59 +01:00
fb470c256a Fix T93256: Instances to points node broken after recent commit
When 97533eede4 added the instance domain, it didn't change
the domain that instance attributes are read from in this node.
2021-11-21 11:37:35 -05:00
873f6148ad Functions: remove test for dynamic name
This was broken in rB6ee2abde82ef121cd6e927995053ac33afdbb438.
2021-11-21 13:08:23 +01:00
940e6525c7 Functions: fix compile error in tests 2021-11-21 13:06:05 +01:00
15011e0b70 Functions: use static string for parameter names
The idea behind this change is the same as in
rB6ee2abde82ef121cd6e927995053ac33afdbb438.

A `MultiFunction::debug_parameter_name` method could be
added separately when necessary.
2021-11-21 12:57:34 +01:00
6ee2abde82 Functions: use static names for multi-functions
Previously, the function names were stored in `std::string` and were often
created dynamically (especially when the function just output a constant).
This resulted in a lot of overhead.

Now the function name is just a `const char *` that should be statically
allocated. This is good enough for the majority of cases. If a multi-function
needs a more dynamic name, it can override the `MultiFunction::debug_name`
method.

In my test file with >400,000 simple math nodes, the execution time improves from
3s to 1s.
2021-11-21 12:48:07 +01:00
d455eadcd8 Cleanup: Move menu search template to C++
This allows the use of better/more readable data structures,
and will also make some refactors to the search button easier.
The build completed on the buildbot for all platforms.
2021-11-20 16:30:53 -05:00
b3ee9f44cf Merge branch 'blender-v3.0-release' 2021-11-20 18:00:46 +01:00
e949ac0bfc Cleanup: unset 'FILE_ENTRY_PREVIEW_LOADING' at the end
Although this function only runs on the main thread, it seems safer to
clear the flag only after setting the result.
2021-11-20 10:34:29 -03:00
59ffe1c5b1 Merge branch 'blender-v3.0-release' 2021-11-20 14:05:26 +01:00
3d447b6335 Merge branch 'blender-v3.0-release' 2021-11-20 12:47:31 +01:00
3baaab15fc Cleanup: Else after return 2021-11-19 23:48:51 -05:00
411261fb32 Merge branch 'blender-v3.0-release' 2021-11-20 01:45:06 -03:00
15ecd47b96 Geometry Nodes: Instance attributes in Transfer/Capture nodes
Updates the Transfer Attributes and Capture Attributes nodes
to support attributes from instances.

Differential Revision: https://developer.blender.org/D13292
2021-11-19 23:40:08 +01:00
01df48a983 Refactor: Port spreadsheet data set to UI tree view
This patch removes a bunch of specific code for drawing the spreadsheet
data set region, which was an overly specific solution for a generic UI.
Nowadays, the UI tree view API used for asset browser catalogs is a much
better way to implement this behavior.

To make this possible, the tree view API is extended in a few ways.
Collapsibility can now be turned off, and whether an item should
be active is moved to a separate virtual function.

The only visual change is that the items are now drawn in a box,
just like the asset catalog.

Differential Revision: https://developer.blender.org/D13198
2021-11-19 17:36:11 -05:00
a0780ad625 Fix: Exception in tree view code in new case
This isn't a problem in 3.0 or master, but I'm porting the spreadsheet
data set region to a tree view and ran into this. This line needs to
whether the function is empty before calling it.

Differential Revision: https://developer.blender.org/D13197
2021-11-19 17:28:08 -05:00
cfbc9df60e Cleanup: Use canonical variable name
All `ARegion` variables should be called `region` unless there is
a good reason not to, since that is the convention.
2021-11-19 16:10:37 -05:00
217d0a1524 Cleanup: Use vector instead of linked list 2021-11-19 15:57:31 -05:00
51a7961e09 Cleanup: Simplify node editor link dragging storage
Now that `node_intern.hh` is a C++ header, we can use C++ types
there. This patch replaces the linked list of dragged links with a
vector. Also, the list of drag operator custom data, `nldrag`, doesn't
seem to need to be a list at all, so I just made it a unique pointer.

Differential Revision: https://developer.blender.org/D13252
2021-11-19 15:36:32 -05:00
c3fed4d463 Merge branch 'blender-v3.0-release' 2021-11-19 16:57:33 -03:00
50ad0e15fe Merge branch 'blender-v3.0-release' 2021-11-19 19:24:55 +01:00
fabd088067 Update splash for Blender 3.x development series
CC-BY Blender Studio https://studio.blender.org

Update the splash artwork for the daily builds (`master`) to
celebrate the beginning of a new major series in Blender.

The badges of the Development Fund membership levels have been removed for
a simpler design, and the font matches the one used elsewhere in blender.org
2021-11-19 18:39:14 +01:00
1222c45544 Fix: Use the instances domain in instance transform nodes
Instance attributes exist on the instance domain after rB97533eede44421,
so these fields should be evaluated on that domain.
2021-11-19 12:26:48 -05:00
ba8dd0f24f Spreadsheet: Support instances component viewer node columns
After rB97533eede444217b, instances have their own attribute domain,
but the spreadsheet code worked under the assumption that the component
used the point domain. Old files have to re-select the instances data
source to make it properly active
2021-11-19 12:21:54 -05:00
06a2e2b28c Merge branch 'blender-v3.0-release' 2021-11-19 18:05:17 +01:00
ef687bd7c2 Merge branch 'blender-v3.0-release' 2021-11-19 12:00:35 -05:00
97533eede4 Geometry Nodes: Support custom instance attributes
Adds an attribute provider for instance attributes.
A new domain `ATTR_DOMAIN_INSTANCE` is implemented.
Instance attributes are not yet realized correctly.

Differential Revision: D13149
2021-11-19 17:53:48 +01:00
9e3a913b35 Merge remote-tracking branch 'origin/blender-v3.0-release' 2021-11-19 16:39:28 +01:00
ec71054a9b Merge branch 'blender-v3.0-release'
Conflicts:
	source/blender/blenkernel/BKE_blender_version.h
	source/blender/blenloader/intern/versioning_300.c
2021-11-19 16:10:28 +01:00
330290d2a8 Cleanup: typos in comments 2021-11-19 15:55:47 +01:00
d7aaa145c6 Merge branch 'blender-v3.0-release' 2021-11-19 15:30:54 +01:00
0852805ed7 Merge branch 'blender-v3.0-release' 2021-11-19 15:58:37 +02:00
06691d1b21 Tests: disable Cycles volume test when WITH_MOD_FLUID is off 2021-11-19 13:21:48 +01:00
1b94c53aa6 Cleanup: fix typos in comments and docs
Contributed by luzpaz.

Differential Revision: https://developer.blender.org/D10447
2021-11-19 13:02:16 +01:00
7d5ef64bfb Cleanup: fix typos in comments and docs
Contributed by luzpaz.

Differential Revision: https://developer.blender.org/D13264
2021-11-19 12:46:49 +01:00
eb071c9ff4 Merge branch 'blender-v3.0-release' 2021-11-19 10:16:30 +01:00
48e64a5fb5 Merge branch 'blender-v3.0-release' 2021-11-19 10:09:29 +01:00
992634427e Nodes: add bf_nodes_geometry library
Separating geometry nodes into a new library will make it
easier to improve compile times with features like unity
builds and precompiled headers.

Differential Revision: https://developer.blender.org/D13261
2021-11-19 09:09:14 +01:00
b8dc845e57 Merge branch 'blender-v3.0-release' 2021-11-19 06:22:47 +01:00
5ed3a5d023 Cleanup: Add an empty() method to RNA's CollectionRef class
The existing RNA CollectionRef class only offers a length() operation
which is sometimes used for checking if the collection is empty. This is
inefficient for certain collection types which do not have a native
length member; the entire list is iterated to find the count.

This patch creates an explicit empty() method to be used in such cases
for better semantics. Additionally, many collection types will behave
more efficiently when using the new method instead of checking length.

Making use of the new method will follow separately.

Differential Revision: https://developer.blender.org/D12314
2021-11-18 14:32:56 -08:00
3f288e9414 Node Editor: Change minimum zoom
Now with Geometry Nodes in Blender,
trees can become much bigger than before.
This changes the minimum zoom value in the node editor.

Differential Revision: https://developer.blender.org/D13254
2021-11-18 21:24:55 +01:00
b2d37c35d0 Node Editor: Skip socket drawing on low zoom
Socket drawing can be heavy with many nodes.
This patch skips drawing them on scale < 0.2
when they are barely visible anyway.

Differential Revision: https://developer.blender.org/D13255
2021-11-18 21:21:10 +01:00
167ee8f2c7 Merge branch 'blender-v3.0-release' 2021-11-18 19:37:48 +01:00
f190f2d267 Revert fixes 2021-11-18 14:19:59 -03:00
0c33411bdd Revert "Revert "Revert "Revert "Allow navigating while transforming""""
This reverts commit 717a971035.
2021-11-18 14:15:08 -03:00
ea42c1a22e Revert "Revert "Revert "Revert "Adjust snap source drawing when adding multiple snap points""""
This reverts commit b8bf40ed4b.
2021-11-18 14:14:57 -03:00
f61a73093b Revert "Revert "Revert "Revert "Transform: interactive mode for editing a 'Snap Source'""""
This reverts commit 701f2dfd5b.
2021-11-18 14:14:51 -03:00
ada6742601 Merge remote-tracking branch 'origin/blender-v3.0-release' 2021-11-18 17:58:26 +01:00
701f2dfd5b Revert "Revert "Revert "Transform: interactive mode for editing a 'Snap Source'"""
This reverts commit 25fa6c74b9.
2021-11-18 13:55:24 -03:00
b8bf40ed4b Revert "Revert "Revert "Adjust snap source drawing when adding multiple snap points"""
This reverts commit c7f9a782aa.
2021-11-18 13:55:21 -03:00
717a971035 Revert "Revert "Revert "Allow navigating while transforming"""
This reverts commit 2a9cfdac7e.
2021-11-18 13:55:17 -03:00
Germano Cavalcante
2a9cfdac7e Revert "Revert "Allow navigating while transforming""
This reverts commit 5e6fdaa07f.
2021-11-18 13:52:39 -03:00
c7f9a782aa Revert "Revert "Adjust snap source drawing when adding multiple snap points""
This reverts commit 77df32548b.
2021-11-18 13:52:30 -03:00
Germano Cavalcante
25fa6c74b9 Revert "Revert "Transform: interactive mode for editing a 'Snap Source'""
This reverts commit 805181bffa.
2021-11-18 13:52:18 -03:00
3531021d1b Cleanup: Simplify declarations in C++ header
Using `struct` everywhere is unnecessary in C++, and the typedefs
are also unnecessary.
2021-11-18 11:46:44 -05:00
Germano Cavalcante
805181bffa Revert "Transform: interactive mode for editing a 'Snap Source'"
This reverts commit f19bd637e2.
2021-11-18 13:42:45 -03:00
77df32548b Revert "Adjust snap source drawing when adding multiple snap points"
This reverts commit cb3ba68ec4.
2021-11-18 13:42:36 -03:00
Germano Cavalcante
5e6fdaa07f Revert "Allow navigating while transforming"
This reverts commit 1d1855e95f.
2021-11-18 13:42:31 -03:00
Germano Cavalcante
1d1855e95f Allow navigating while transforming
This feature has been desired for some time:
- https://rightclickselect.com/p/ui/Tqbbbc/allow-navigating-while-transforming (See comments);
- D1583;
- T37427;

In short, blocking navigation during transform limits the user to move the object only to visible areas within the screen and hinders the allocation of objects within closed meshes.

The node editor is also impaired because some nodes are far between them and the connectors are too small.

The only disadvantage of this patch (as I see it) is the conflict with the existing key map:
MIDDLEMOUSE:
- enable axis constrain in 3D view;

WHEELDOWNMOUSE, WHEELUPMOUSE, PAGEUPKEY, PAGEDOWNKEY:
- change the threshold of the proportional edit;

So the patch solution was to change these keymaps:
- MIDDLEMOUSE to Alt+MIDDLEMOUSE;
- WHEELDOWNMOUSE, WHEELUPMOUSE, PAGEUPKEY, PAGEDOWNKEY to Alt+(corresponding key);

When you use this new keymap for the first time in the proportional edit, it may seem strange due to the custom of using it (both in View2D and View3D).
But quickly the user gets used to it.

Alternatively we can add an option to the user preferences ([] Allow navigating while transforming). (I'm not much fan of this option).

The patch was done on branch2.8. But maybe it's a good idea to apply it to 2.79

Differential Revision: https://developer.blender.org/D2624
2021-11-18 13:14:18 -03:00
cb3ba68ec4 Adjust snap source drawing when adding multiple snap points 2021-11-18 13:14:18 -03:00
Germano Cavalcante
f19bd637e2 Transform: interactive mode for editing a 'Snap Source'
This patch implements part of what was stated in {T66484}, with respect to `Base Point`.

## Introduction

The snapping feature of the transform tools has a variety of applications:
- Organization of nodes.
- Positioning of frames in precise time units.
- Retopology with snap to face
- Creation of armatures with bone positioning through the snap to volume
- Precise positioning of 3D or 2D objects in the surrounding geometry (CAD modeling)

The goal of this document is to make it more powerful for precision modeling and still supporting the old use cases without extra complexity.
The main topic addressed here is the introduction of a **interactive mode for setting a snap source** (See terminology).

## Terminology

* **Snap Source**: 3d coordinate * we want to snap from. (Currently defined by the `Snap With` options: `Closest`, `Center`, `Median` and `Active`).
* **Snap Target**: 3d coordinate*  we want to snap to. (Vertices, Edges, Faces, Grid...)

## Interactive Mode for Editing a Snap Source

Currently the fixed snap point can only be obtained through the `Snap With` options. So it's a little tricky for the user to define a snap source point having so much geometry on an object.
Because of this, the user needs to resort to impractical solutions to get a point in the geometry.
See example of an impractical use:
{F11714181, layout=left, width=960, alt="The user used the cursor (which can be snapped) to choose the snap origin point."}
The user used the cursor (which can be snapped) to choose the snap source point.

While it is possible to work around this current limitation, it is important to reduce the number of steps and allow the user to set a snap source point through an optional interactive mode during a transformation.

The proposed solution is to be able to move the current snap source point through a modal modifier activated with a key (eg. B).
The snap source point can thus "snap" to the elements in the scene (vertex, mid-edge, Lamp, …) during this mode.
{F9122814, layout=left, width=960, alt="Base Point Snap, example of transform operation via the shortcut (not the tool). After pressing g and the snap base change shortcut (e.g., shift + ctrl) the user set the base point. The base point is then visible until the end of the operation. The z axis constrains the final position."}

## Implementation Details

- The feature will only be available in 3D View.
- The feature will only be available for `Move`, `Rotate` and `Scale` transform modes.
- The snap source editing will be enabled with a single click on the modifier key (B).
- Having a snap point indicated, the new snap origin point will be confirmed with the same buttons that confirms the transformation (but the transformation will not be concluded).
- The snap source editing can be canceled with the same key that activated it (B).
- If the transformation is done with "release_confirm" (common for gizmos), the new feature cannot be enabled.
- During the transformation, when enabling the feature, if the snap option is turned off in the scene, the snap will be forced on throughout the rest of the transformation (unless interactive mode is canceled).
- During a transformation, if no snap target is set for an element in the scene (Vertex, Grid...), the snap targets to geometry Vertex, Edge, Face, Center of Edge and Perpendicular of Edge will be set automatically.
- Snap cannot be turned off during the snap source editing.
- Constraint or similar modification features will not be available during the snap source editing.
- Text input will not be available during the snap source editing.
- When adding multiple snap points (A) the new prone snap source point will be indicated with an "X" drawing.
{F11817267}

Maniphest Tasks: T66484

Differential Revision: https://developer.blender.org/D9415
2021-11-18 13:14:18 -03:00
67b4eecac9 Cleanup: Use const arguments 2021-11-18 11:03:18 -05:00
Nikhil Shringarpurey
dd31b8bd50 UI: Use full word "Start" instead of "Sta"
Differential Revision: https://developer.blender.org/D13098
2021-11-18 10:45:10 -05:00
5816eb4c56 Merge remote-tracking branch 'origin/blender-v3.0-release' 2021-11-18 16:36:11 +01:00
beb9e332ca Merge remote-tracking branch 'origin/blender-v3.0-release' 2021-11-18 16:21:17 +01:00
d1f944c186 Cycles: declare constants at program scope on Metal
MSL requires that constant address space literals be declared at program
scope. This patch moves the `blackbody_table_r/g/b` and `cie_colour_match`
constants into separate files so they can be declared at the appropriate scope.

Ref T92212

Differential Revision: https://developer.blender.org/D13241
2021-11-18 14:38:05 +01:00
d19e35873f Cycles: several small fixes and additions for MSL
This patch contains many small leftover fixes and additions that are
required for Metal-enablement:

- Address space fixes and a few other small compile fixes
- Addition of missing functionality to the Metal adapter headers
- Addition of various scattered `__KERNEL_METAL__` blocks (e.g. for
  atomic support & maths functions)

Ref T92212

Differential Revision: https://developer.blender.org/D13263
2021-11-18 14:38:02 +01:00
c0d52db783 Merge branch 'blender-v3.0-release' 2021-11-18 14:33:43 +01:00
2b63a76041 Merge branch 'blender-v3.0-release' 2021-11-18 11:53:46 +01:00
William Leeson
e1a3b697ec Merge branch 'blender-v3.0-release' to pick up D13262. 2021-11-18 09:41:11 +01:00
Félix
9cf3d841a8 VSE: Add meta.separate() Python API function
This function can be used to "dissolve" meta strip anywhere in strip
hierarchy. This has same effect as `meta_separate` operator.

Reviewed By: ISS

Differential Revision: https://developer.blender.org/T91005
2021-11-18 03:55:48 +01:00
f1f7a8b018 Merge branch 'blender-v3.0-release' 2021-11-18 03:03:13 +01:00
032ab0270d Merge branch 'blender-v3.0-release' 2021-11-18 02:23:17 +01:00
ceec400975 Merge branch 'blender-v3.0-release' 2021-11-18 01:40:37 +01:00
d6b5251572 Merge branch 'blender-v3.0-release' 2021-11-18 01:13:22 +01:00
fa7a6d67a8 Fix Cycles CUDA/HIP compiler error after recent changes 2021-11-17 19:56:18 +01:00
Sebastian Herholz
d9bc8f189c Cycles: add build option to enable a debugging feature for MIS
This patch adds a CMake option "WITH_CYCLES_DEBUG" which builds cycles with
a feature that allows debugging/selecting the direct-light sampling strategy.
The same option may later be used to add other debugging features that could
affect performance in release builds.

The three options are:
* Forward path tracing (e.g., via BSDF or phase function)
* Next-event estimation
* Multiple importance sampling combination of the previous two methods

Such a feature is useful for debugging light different sampling, evaluation,
and pdf methods (e.g., for light sources and BSDFs).

Differential Revision: https://developer.blender.org/D13152
2021-11-17 18:03:56 +01:00
063ad8635e Cycles: reduce triangle memory usage with packed_float3
Depends on D13243

Differential Revision: https://developer.blender.org/D13244
2021-11-17 17:29:41 +01:00
9937d5379c Cycles: add packed_float3 type for storage
Introduce a packed_float3 type for smaller storage that is exactly 3
floats, instead of 4. For computation float3 is still used since it can
use SIMD instructions.

Ref T92212

Differential Revision: https://developer.blender.org/D13243
2021-11-17 17:29:41 +01:00
89d5714d8f Build: match GCC and Clang float conversion warnings in Cycles 2021-11-17 17:29:41 +01:00
ea7efa5569 Merge branch 'blender-v3.0-release' 2021-11-17 15:41:32 +01:00
c3422c48ad Cleanup: remove dummy multi function 2021-11-17 12:33:18 +01:00
e5f05bc7a6 Cleanup: Painting - reduce reallocation of same memory.
Curve mask is freed/allocated every time, but could still reuse the
previous allocated buffer when the diameter of the brush doesn't change.
2021-11-17 12:30:42 +01:00
f5dde382af Cleanup: use same function for updating internal links for all nodes
Previously, node types had a callback that creates internal links. Pretty
much all nodes used the same callback though. The exceptions are the
reroute node (which probably shouldn't be mutable anyway) and some
input/output nodes that are not mutable.

Removing the callback helps with D13246, because it makes it easier
to reason about which internal links are created and when they change.
In the future, the internal links should be part of the node declaration.
2021-11-17 11:52:54 +01:00
83a4d51997 Cleanup: Remove unused show_samples() device code in Cycles. 2021-11-17 11:16:48 +01:00
c2ab47e729 Cleanup: change node socket availability in a single place
This cleans up part of the code that still set the flag manually. Also, this
change helps with D13246 because it makes it easier to tag the node
tree as changed when the availability of a socket changed.
2021-11-17 11:11:28 +01:00
473be239c3 LineArt: Improve certain edge cases in occlusion
This patch includes:
View vector fix for ortho back face.
Point on segment logic correction.
Better handling of boundary cases.

See review page for detailed description.

Reviewed By: Sebastian Parborg (zeddb)

Differential Revision: https://developer.blender.org/D13143
2021-11-17 14:30:08 +08:00
e3c974b7e4 Merge branch 'blender-v3.0-release' 2021-11-17 05:43:34 +01:00
ecad33f214 UI: Use ampersand instad of 'and' in labels
- When and is used in labels use ampsand
- When used in description use 'and'
2021-11-16 21:38:03 -05:00
61bffa565e Fix T90412: Inconsistency in mask strip color
This seems to be oversight in 271231f58e where strip color was
defined only for light theme.
2021-11-17 03:27:13 +01:00
f72dc00569 Merge branch 'blender-v3.0-release' 2021-11-17 03:12:19 +01:00
4c988eb3e1 Fix error with makefiles compilation
Use 'template' keyword to treat 'is' as a dependent template name
2021-11-16 20:46:33 -03:00
luzpaz
dea26253a0 cleanup: fix typos in comments and docs
Followup to https://developer.blender.org/D10288

Reviewed By: Blendify

Differential Revision: https://developer.blender.org/D10346
2021-11-16 18:45:10 -05:00
8290edefad Cleanup: Use bool instead of int 2021-11-16 17:20:31 -06:00
2f39b45e8c Merge branch 'blender-v3.0-release' 2021-11-16 17:50:24 -05:00
93f26d652e Merge branch 'blender-v3.0-release' 2021-11-16 23:21:13 +01:00
fbf4fe6963 Fix missing Cycles implicit float/double conversion error with Clang
Since these are errors with GCC and Visual Studio.
2021-11-16 23:18:25 +01:00
59da22c309 Merge branch 'blender-v3.0-release' 2021-11-16 22:40:08 +01:00
cf83719761 Geometry Nodes: Small improvements to object info node
This commit contains a few mostly-related changes to this node:
 -  Add a warning when retrieving the geometry from the modifier object.
 - Only create the output geometry when it is necessary.
 - Decompose transform matrices in a more friendly way.
 - Use default return callbacks like other newer nodes.

Differential Revision: https://developer.blender.org/D13232
2021-11-16 15:09:00 -06:00
c9fb08e075 Merge branch 'blender-v3.0-release' 2021-11-16 14:55:13 -06:00
Germano Cavalcante
9d7422b817 File Browser: Improve usage of threads in the creation of thumbnails
Due to asynchronous process, the preview for a given image may be
generated several times.

This regenerates many thumbs unnecessarily.

The solution is to add the `FILE_ENTRY_PREVIEW_LOADING` flag for file
entries that are still in the thread queue.

So this flag is checked not to redraw the thumb when it is still being
created on a different thread.

Differential Revision: https://developer.blender.org/D11150
2021-11-16 14:10:24 -03:00
917218269e Merge branch 'blender-v3.0-release' 2021-11-16 13:59:28 -03:00
ba6427adfa Merge branch 'blender-v3.0-release' 2021-11-16 17:25:48 +01:00
bee7a56687 Cleanup: document that MEM_dupallocN is NULL-safe
Add comment explaining `MEM_dupallocN` is NULL-safe, in that it returns
NULL when it receives a NULL pointer. This is currently true for both
implementations of the function (`MEM_lockfree_dupallocN` and
`MEM_guarded_dupallocN`), and will be expected of other implementations
as well.

No functional changes.
2021-11-16 17:11:45 +01:00
Michael Jones
64003fa4b0 Cycles: Adapt volumetric lambda functions to work on MSL
This patch adapts the existing volumetric read/write lambda functions for Metal. Lambda expressions are not supported on MSL, so two new macros `VOLUME_READ_LAMBDA` and `VOLUME_WRITE_LAMBDA` have been defined with a default implementation which, on Metal, is overridden to use inline function objects.

This patch also removes the last remaining mention of the now-unused `ccl_addr_space`.

Ref T92212

Reviewed By: leesonw

Maniphest Tasks: T92212

Differential Revision: https://developer.blender.org/D13234
2021-11-16 13:42:23 +00:00
85ac9b8584 Merge branch 'blender-v3.0-release' 2021-11-16 14:39:51 +01:00
12a986c9b5 Merge remote-tracking branch 'origin/blender-v3.0-release' 2021-11-16 13:10:35 +01:00
7da714f387 Merge branch 'blender-v3.0-release' 2021-11-16 10:57:26 +01:00
d4c868da9f Geometry Nodes: refactor virtual array system
Goals of this refactor:
* Simplify creating virtual arrays.
* Simplify passing virtual arrays around.
* Simplify converting between typed and generic virtual arrays.
* Reduce memory allocations.

As a quick reminder, a virtual arrays is a data structure that behaves like an
array (i.e. it can be accessed using an index). However, it may not actually
be stored as array internally. The two most important implementations
of virtual arrays are those that correspond to an actual plain array and those
that have the same value for every index. However, many more
implementations exist for various reasons (interfacing with legacy attributes,
unified iterator over all points in multiple splines, ...).

With this refactor the core types (`VArray`, `GVArray`, `VMutableArray` and
`GVMutableArray`) can be used like "normal values". They typically live
on the stack. Before, they were usually inside a `std::unique_ptr`. This makes
passing them around much easier. Creation of new virtual arrays is also
much simpler now due to some constructors. Memory allocations are
reduced by making use of small object optimization inside the core types.

Previously, `VArray` was a class with virtual methods that had to be overridden
to change the behavior of a the virtual array. Now,`VArray` has a fixed size
and has no virtual methods. Instead it contains a `VArrayImpl` that is
similar to the old `VArray`. `VArrayImpl` should rarely ever be used directly,
unless a new virtual array implementation is added.

To support the small object optimization for many `VArrayImpl` classes,
a new `blender::Any` type is added. It is similar to `std::any` with two
additional features. It has an adjustable inline buffer size and alignment.
The inline buffer size of `std::any` can't be relied on and is usually too
small for our use case here. Furthermore, `blender::Any` can store
additional user-defined type information without increasing the
stack size.

Differential Revision: https://developer.blender.org/D12986
2021-11-16 10:16:30 +01:00
6d35972b06 Merge branch 'blender-v3.0-release' 2021-11-16 09:58:47 +01:00
57ed435def Cleanup: Use C++ matrix identity constructor 2021-11-15 23:24:16 -06:00
7e42ae7c1a Cleanup: Typo in comments 2021-11-15 22:08:28 +01:00
165cacc6f0 VSE: Use alpha over as default blend mode
With transform tools, it is expected to see backgroud image when overlay
is transformed.

Alpha over caused performance to be not optimal when used with opaque
media. This should be addressed with D12914 at least partially. There
may be some corner cases not addressed.

Differential Revision: https://developer.blender.org/D12952
2021-11-15 21:07:25 +01:00
62da6ffe08 VSE: Use early out for aplha over blending
When scaling down image, users expect to see background, which doesn't
currently happen in VSE. This is because strips use cross blend mode by
default, because alpha over is much slower. Reason is, because any area
of image can be transparent, and therefore it can't have early out
implemented in a way that cross blend mode can.

Flag images rendered by codecs that don't support transparency as fully
opaque and implement a form of early out for alpha over blend mode.

When rendering image stack, 2-input effects are ignored on the
"way down". Alpha over needs rendered overlay image to decide whether it
will use only overlay or background too. Therefore overlay can be
rendered safely before it is used. Image flags can be checked and it can
be freed if needed. Freeing doesn't cause any performance degradation,
because image is always stored in cache.

This feature does not improve blend mode performance. In summary, it
only allowes for having alpha over blend mode on background images
without suffering from lower performance.

Reviewed By: sergey

Differential Revision: https://developer.blender.org/D12914
2021-11-15 21:03:43 +01:00
46f5f60c13 Merge branch 'blender-v3.0-release' 2021-11-15 20:33:34 +01:00
10a6a540af Cleanup: remove unnecessary functions
Those functions were more useful when `FieldInferencingInterface`
was still declared further down in `node.cc`.
2021-11-15 18:04:03 +01:00
8976b72843 Merge branch 'blender-v3.0-release' 2021-11-15 09:10:19 -06:00
c3472cb11c Fix T93074: Gpencil cutter not using flat caps in middle cuts
When cut an stroke using the option Flat Caps, the falt was not done if the cut was done in the middle of the stroke.

Now the flat is applied to the segments created and also some cleanup of the code done.
2021-11-15 12:17:11 +01:00
c2c65cc4bf Merge branch 'blender-v3.0-release' 2021-11-15 02:39:33 +01:00
7e82c840b7 Fix text editor auto-close with quotes
Back-spacing a quote from the beginning of a line
would delete the quote in-front instead of doing nothing.
2021-11-14 11:26:06 +11:00
2549384baa Cleanup: minor tweaks to auto-close
Spelling and failure to reuse variable missed in review.
2021-11-14 11:11:20 +11:00
73047c69ea BLF: Use Floats for Font Point Sizes
Allow the use of floating-point values for font point sizes, which
allows greater precision and flexibility for text output.

See D8960 for more information, details, and justification.

Differential Revision: https://developer.blender.org/D8960

Reviewed by Campbell Barton
2021-11-13 09:39:18 -08:00
e8a8bb67fc Cleanup: Correct order of guard and lock in moviecache_valfree
Fix own mistake in rB7061d1e39fe

In my attempt to quickly address T92838, along with the original bug, I
made a nonsensical choice to use the limiter lock to guard the check
against the cache item itself. While harmless, it is not necessary and
semantically wrong / potentially confusing to future readers of the code.

Differential Revision: https://developer.blender.org/D13122
2021-11-12 20:47:26 -08:00
ab9ec193c3 Fix splash screen showing on startup with files loaded by scripts
Suppressing the splash was only done when passing in an argument from
the command line.

Remove G.file_loaded, as it is misleading, only set once on startup,
replace with G.relbase_valid which is used everywhere else to check
if the file path should be used.
2021-11-13 14:05:27 +11:00
Matheus Santos
c4ea5cb1a3 Text Editor: Auto close relevant characters
Support the ability to close relevant characters like '(', '[' and '{'.

It will also delete the pair character if they're empty.

Ref D13119

Reviewed By: campbellbarton
2021-11-13 13:56:31 +11:00
1143bf281a Cleanup: spelling in comments, comment block formatting 2021-11-13 13:07:13 +11:00
acc800d24d Cleanup: clang-format 2021-11-13 12:47:18 +11:00
dc378bf1a4 Merge branch 'blender-v3.0-release' 2021-11-12 21:57:44 -03:00
60b8eb30bb Merge branch 'blender-v3.0-release' 2021-11-12 21:36:35 -03:00
738f4fbc5e Revert "Fix T92636: Vector math node link disconnects when loading old file"
This reverts commit 6b4ca78108.

A simpler fix was used for 3.0, but rBd845ba481c6d2ef already contained
a more complete solution to the problem of inconsistent socket ids.
2021-11-12 17:14:27 -06:00
ebb4aba325 Merge branch 'blender-v3.0-release' 2021-11-12 17:12:39 -06:00
ec432ae998 Merge branch 'blender-v3.0-release' 2021-11-12 14:06:26 -06:00
55c69373e8 Cleanup: split 'initSnappingMode' into more specific functions
This helps to reuse small regions of a function's code elsewhere.

The logic had to be reorganized, theoretically it should behave the same way.
2021-11-12 16:30:01 -03:00
8b13cf5667 Cleanup: move 'imm_drawcircball' to 'gpu_immediate_util.c' 2021-11-12 16:30:01 -03:00
5941c39fbf Cleanup: fix some comments in the transform code 2021-11-12 16:30:01 -03:00
30f9034182 Cleanup: use 't->tsnap.mode' in transform code
This also prevents different snap modes from being used at the same time in the code.
2021-11-12 16:30:01 -03:00
e5a7dd8ab6 Cleanup: unify snap modes to geometry in a single flag
This combination was being repeated in some places.
2021-11-12 16:30:01 -03:00
5b787c24fb Merge branch 'blender-v3.0-release' 2021-11-12 13:26:26 -06:00
1b55b911f2 Merge branch 'blender-v3.0-release' 2021-11-12 20:04:05 +01:00
d845ba481c Fix T91826: Inconsistent node socket name identifier separator
Previously both `.` and `_` were used as separators when finding
a unique name for a socket. This removes the use of `.`, since `_`
was more common. It also does versioning for all of a file's node
trees to make sure that they all use the `_` convention.

Differential Revision: https://developer.blender.org/D13181
2021-11-12 12:22:43 -06:00
cbca71a7cf Cleanup: Move remaning node editor files to C++
Differential Revision: https://developer.blender.org/D13200
2021-11-12 12:12:27 -06:00
809ae823b7 Merge branch 'blender-v3.0-release' 2021-11-12 19:00:23 +01:00
8a8bf99717 Merge remote-tracking branch 'origin/blender-v3.0-release' 2021-11-12 18:23:44 +01:00
Yevgeny Makarov
a89529d8db UI: Do not shade alpha when blending colors
UI_GetThemeColorBlendShade4fv incorrectly changing alpha by the amount
of the shading offset.

See D9944 for more details.

Differential Revision: https://developer.blender.org/D9944

Reviewed by Hans Goudey
2021-11-12 08:55:20 -08:00
9f5290e3bc Fix T93007: Cycles not updating for animated Object properties like color 2021-11-12 08:55:20 -08:00
c671b5eee4 Fix Cycles ray visibility panel missing for volume objects 2021-11-12 08:55:20 -08:00
ddf66cd060 Merge remote-tracking branch 'origin/blender-v3.0-release' 2021-11-12 14:36:55 +01:00
c7a88cf91a Merge branch 'blender-v3.0-release' 2021-11-13 00:31:03 +11:00
a470e3c9d1 Merge branch 'blender-v3.0-release' 2021-11-13 00:20:07 +11:00
3fe735d371 Geometry Nodes: Add Outer Points Selection to Star
Adds a boolean field output containing a selection of the
points of the star that are controlled by the outer radius
of the star.

Differential Revision: https://developer.blender.org/D13097
2021-11-12 06:44:41 -06:00
a47359ff36 Merge branch 'blender-v3.0-release' 2021-11-12 11:52:49 +01:00
26502f3d89 Merge remote-tracking branch 'origin/blender-v3.0-release' 2021-11-12 10:52:12 +01:00
86ca206db8 Cleanup/document BKE_blender_copybuffer.
* Rename the 'copy' functions to make it clear they belong to the same
  'group' and are to be used together.
* Fix `flag` parameter of `BKE_copybuffer_paste` being a short instead
  of an int.
* Improve documentation.
2021-11-12 10:20:49 +01:00
William Leeson
32c7687859 Fix T92601: Disable profiling when the profiler is deemed not active.
Adds a method to profiler that can be used to check if it is active.
This is used to determine if stop_profiling and start_profiling
should be called.

| patch | Juans Scene UI 256 samples | Juans Scene bg 256 samples | junkshop UI | junkshop bg |
| No patch | 6:16.59 | 4:05.37 | 2:08.48 | 1:59.7 |
| D13187   | 4:12.15 | 3:57.36 | 2:07.25 | 1:58.16 |
| D13185   | 4.11.18 |3:54.74 | 2:07.44 | 1:58.03 |
| D13190   | 4:12.39 | 3:55.42 | 2:07.62 | 1:58.68 |

UI - means rendered from within Blender
bg - means rendered from the command line using ##blender -b scene.blend -f 1##

Reviewed By: sergey, brecht

Maniphest Tasks: T92601

Differential Revision: https://developer.blender.org/D13190
2021-11-12 10:01:48 +01:00
a87253942d Cleanup: Remove GHOST_isUpsideDownContext.
GHOST API only has a header definition. No implementation or usage.
2021-11-12 09:38:25 +01:00
ae74ad191c Merge branch 'blender-v3.0-release' 2021-11-12 18:35:22 +11:00
aa1c44a113 Merge branch 'blender-v3.0-release' 2021-11-12 18:35:19 +11:00
5c0d4753cf Merge branch 'blender-v3.0-release' 2021-11-12 18:35:16 +11:00
02333544d1 Merge branch 'blender-v3.0-release' 2021-11-12 18:35:13 +11:00
1e1c870001 Cleanup: Improve comment 2021-11-11 21:57:40 -06:00
0533f2851e Geometry Nodes: change selection output order in Cylinder node
This new order is a bit more intuitive.
2021-11-11 19:53:02 +01:00
d6e682a7b0 Merge branch 'blender-v3.0-release' 2021-11-11 19:50:24 +01:00
e61da8e4fb Merge branch 'blender-v3.0-release' 2021-11-11 11:47:43 -06:00
50f32025ac Merge branch 'blender-v3.0-release' 2021-11-11 18:27:31 +01:00
c9c7658926 Geometry Nodes: Add Offset to Handle Position Node
Adds a vector offset field to the "Curve Handle Position Node".
This vector is added to the incoming position (which is the
implicit handle position if not connected) which will set the
position of the handle. Default is (0,0,0)

Differential Revision: https://developer.blender.org/D13035
2021-11-11 10:58:23 -06:00
3ca41b7312 Merge branch 'blender-v3.0-release' 2021-11-11 10:16:35 -06:00
25e7365d0d Cleanup CUDA / HIP comments
Remove outdated CUDA comments for bindless textures and cleanup some HIP comments that still mentioned CUDA.

Differential Revision: https://developer.blender.org/D13189
2021-11-11 16:37:29 +01:00
52c617802f Merge branch 'blender-v3.0-release' 2021-11-11 09:27:42 -06:00
ce395c84a3 Merge branch 'blender-v3.0-release' 2021-11-11 15:29:35 +01:00
06a74e7816 LibLink/Append tests: Add basic testing of bpy.data.libraries.load code. 2021-11-11 14:54:26 +01:00
9f31b9b7d3 Merge branch 'blender-v3.0-release' 2021-11-11 14:33:28 +01:00
4a98faf9f1 Merge branch 'blender-v3.0-release' 2021-11-11 21:34:05 +11:00
8c240f50b2 Merge branch 'blender-v3.0-release' 2021-11-11 21:34:02 +11:00
Andrii
c63e735f6b Cycles: Add sample offset option
This patch exposes the sampling offset option to Blender. It is located in the "Sampling > Advanced" panel.
For example, this can be useful to parallelize rendering and distribute different chunks of samples for each computer to render.

---

I also had to add this option to `RenderWork` and `RenderScheduler` classes so that the sample count in the status string can be calculated correctly.

Reviewed By: leesonw

Differential Revision: https://developer.blender.org/D13086
2021-11-11 09:39:25 +01:00
b8d53b703a Merge branch 'blender-v3.0-release' 2021-11-11 15:00:23 +11:00
9787b46f09 Merge branch 'blender-v3.0-release' 2021-11-11 15:00:20 +11:00
03f0be35d6 Merge branch 'blender-v3.0-release' 2021-11-11 15:00:17 +11:00
e1bd4bbb66 UI: Introduce View pie in more editors
#### Motivation

The View pie menu is a convenient way to access operators such as `Frame Selected` and `Frame All` which are usually mapped to `PERIOD` or `HOME` keys on the right side of most keyboard, making it hard hard to reach with the left hand.

The motivation for this patch comes from working with a 75% keyboard (no numpad). Most laptops face a similar problem.

#### Implementation

The View pie menu has been added to the following editors and sub-modes where applicable:

* Node Editor
* Video Sequencer
* Dopesheet
* Graph
* NLA
* Image
* Clip
* Outliner

More options could definitely be added to this menu for convenience, as long as it maintains the common options in the same place (Frame Selected on the left, Frame All on the right).

For positioning I went with the following layout:
{F11791186, size=full}

I've added `Zoom 1:1`to the Image Editor and the VSE Preview since there is no way to reset the zoom on keyboards without numpad (unless Emulate Numpad is turned on).

The Outliner uses `Show Active` and `Show Hierarchy` which are the closest ones to the equivalent in other editors. Should `Show Active` be renamed to `Frame Selected`?

The shortcut assigned is the same as the 3D Viewport (`ACCENT_GRAVE`).

#### Screenshots

Node Editor
{F11778387, size=full}

Dopesheet
{F11778400, size=full}

Graph
{F11778403, size=full}

Image Editor (Paint and View)
{F11791113, size=full}

Image Editor (Mask)
{F11791114, size=full}

UV Editor
{F11791119, size=full}

Clip Editor (Tracking)
{F11791137, size=full}

Clip Editor (Mask)
{F11791140, size=full}

Clip Editor (Graph)
{F11791151, size=full}
View operators are not yet implemented in Clip Editor Dopesheet mode (left a note about this in the menu poll).

Reviewed By: #user_interface, campbellbarton

Differential Revision: https://developer.blender.org/D13169
2021-11-11 01:18:50 +01:00
f1a8644121 Cleanup: Move interface_region_search.c to C++
This will be helpful for solving a bug with search during animation
playback, T89313. I tested this on all platforms on the buildbot.
2021-11-10 15:49:49 -06:00
1ec7075ff2 Merge branch 'blender-v3.0-release' 2021-11-10 15:44:13 -06:00
9ca8bf0b29 Merge branch 'blender-v3.0-release' 2021-11-10 22:28:03 +01:00
3fa86f4b28 Merge branch 'blender-v3.0-release' 2021-11-10 20:19:09 +01:00
abf62d06d1 Merge branch 'blender-v3.0-release' 2021-11-10 11:11:58 -06:00
cc17ed26ce Merge branch 'blender-v3.0-release' 2021-11-10 10:52:34 -06:00
e9b7e5e0b9 Merge branch 'blender-v3.0-release' 2021-11-10 10:47:09 -06:00
f565620435 Fix T92985: CUDA errors with Cycles film convert kernels
rB3a4c8f406a3a3bf0627477c6183a594fa707a6e2 changed the macros that create the film
convert kernel entry points, but in the process accidentally changed the parameter definition
to one of those (which caused CUDA launch and misaligned address errors) and changed the
implementation as well. This restores the correct implementation from before.

In addition, the `ccl_gpu_kernel_threads` macro did not work as intended and caused the
generated launch bounds to end up with an incorrect input for the second parameter (it was
set to "thread_num_registers", rather than the result of the block number calculation). I'm
not entirely sure why, as the macro definition looked sound to me. Decided to simply go with
two separate macros instead, to simplify and solve this.

Also changed how state is captured with the `ccl_gpu_kernel_lambda` macro slightly, to avoid
a compiler warning (expression has no effect) that otherwise occurred.

Maniphest Tasks: T92985

Differential Revision: https://developer.blender.org/D13175
2021-11-10 15:49:50 +01:00
a6e4cb092e Merge branch 'blender-v3.0-release' 2021-11-10 13:53:44 +01:00
53468c2b13 Merge branch 'blender-v3.0-release' 2021-11-10 13:43:08 +01:00
03e22da665 Merge branch 'blender-v3.0-release' 2021-11-10 12:47:01 +01:00
ff0c42acfc Merge branch 'blender-v3.0-release' 2021-11-10 21:22:01 +11:00
a49d6a5350 Merge branch 'blender-v3.0-release' 2021-11-10 21:21:58 +11:00
aa440923c8 Fix: Hide selection value in resample curve node, fix order
Selection sockets are meant to come right after the geometry,
this was missed in review of rBa7672caeb255e3.
Also, the selection value was not hidden.
2021-11-09 17:23:11 -06:00
1f6010e609 Cleanup: Move info_stats.c to C++ 2021-11-09 15:57:58 -06:00
Michael Jones
3a4c8f406a Cycles: Adapt shared kernel/device/gpu layer for MSL
This patch adapts the shared kernel entrypoints so that they can be compiled as MSL (Metal Shading Language). Where possible, the adaptations avoid changes in common code.

In MSL, kernel function inputs are explicitly bound to resources. In the case of argument buffers, we declare a struct containing the kernel arguments, accessible via device pointer. This differs from CUDA and HIP where kernel function arguments are declared as traditional C-style function parameters. This patch adapts the entrypoints declared in kernel.h so that they can be translated via a new `ccl_gpu_kernel_signature` macro into the required parameter struct + kernel entrypoint pairing for MSL.

MSL buffer attribution must be applied to function parameters or non-static class data members. To allow universal access to the integrator state, kernel data, and texture fetch adapters, we wrap all of the shared kernel code in a `MetalKernelContext` class. This is achieved by bracketing the appropriate kernel headers with "context_begin.h" and "context_end.h" on Metal. When calling deeper into the kernel code, we must reference the context class (e.g. `context.integrator_init_from_camera`). This extra prefixing is performed by a set of defines in "context_end.h". These will require explicit maintenance if entrypoints change. We invite discussion on more maintainable ways to enforce correctness.

Lambda expressions are not supported on MSL, so a new `ccl_gpu_kernel_lambda` macro generates an inline function object and optionally capturing any required state. This yields the same behaviour. This approach is applied to all parallel_... implementations which are templated by operation. The lambda expressions in the film_convert... kernels don't adapt cleanly to use function objects. However, these entrypoints can be macro-generated more concisely to avoid lambda expressions entirely, instead relying on constant folding to handle the pixel/channel conversions.

A separate implementation of `gpu_parallel_active_index_array` is provided for Metal to workaround some subtle differences in SIMD width, and also to encapsulate some required thread parameters which must be declared as explicit entrypoint function parameters.

Ref T92212

Reviewed By: brecht

Maniphest Tasks: T92212

Differential Revision: https://developer.blender.org/D13109
2021-11-09 21:43:10 +00:00
4648c4990c Merge branch 'blender-v3.0-release' 2021-11-09 13:08:36 -06:00
a356e4fb3f Merge branch 'blender-v3.0-release' 2021-11-09 17:53:15 +01:00
7383f95443 Merge branch 'blender-v3.0-release' 2021-11-09 11:43:18 -05:00
ed0df0f3c6 Merge branch 'blender-v3.0-release' 2021-11-09 10:15:20 -06:00
e5fb5c9d7b Merge branch 'blender-v3.0-release' 2021-11-09 17:07:41 +01:00
Demeter Dzadik
e452c43fd6 Let Unlink Action operator have an undo step
I noticed while rigging a character and editing actions that the Unlink Action operator had no undo step. Doesn't feel intentional, so this patch adds the necessary flags.

Reviewed By: mont29

Differential Revision: https://developer.blender.org/D12346
2021-11-09 16:58:39 +01:00
0bdf9d10a4 Merge branch 'blender-v3.0-release' 2021-11-09 16:46:58 +01:00
fb0ae66ee5 Merge branch 'blender-v3.0-release' 2021-11-09 16:10:19 +01:00
9dc3f454d9 Merge branch 'blender-v3.0-release' 2021-11-09 11:27:25 -03:00
Demeter Dzadik
c092cc35b3 Expose BLI_string_flip_side_name as bpy.utils.flip_name
Expose a new function in `bpy.utils.flip_name(name, strip_number=False)
that allows flipping bone names, eg "Bone.L" -> "Bone.R".

Useful for add-ons to avoid re-implementing Blender's name flipping.

Ref D12322
2021-11-10 01:19:04 +11:00
accdd4c1bc Merge branch 'blender-v3.0-release' 2021-11-09 15:08:37 +01:00
625349a6bd Cleanup: spelling, C style comments 2021-11-10 00:56:17 +11:00
65bbac6692 Cleanup: clang-format 2021-11-10 00:55:38 +11:00
faeb2cc900 Merge branch 'blender-v3.0-release' 2021-11-09 14:49:47 +01:00
0bcf014bcf Merge branch 'blender-v3.0-release' 2021-11-10 00:38:51 +11:00
41b0820ddd Merge branch 'blender-v3.0-release' 2021-11-09 13:31:33 +01:00
6c24cafecc Fix T92876: Cycles incorrect volume emission + absorption handling 2021-11-09 13:13:56 +01:00
cb487b6507 Asset Catalogs: add test for proper shortening of simple names
Catalog simple names are supposed to fit into the DNA field `char
AssetMetaData::catalog_simple_name[64]`, and thus should be shortened
appropriately. This was already happening, but is now also covered by a
test.

No functional changes.
2021-11-09 13:13:30 +01:00
09f1be53d8 Fix T92950: spreadsheet shows 0 instances when there are instances
Fix found by @erik85.
2021-11-09 13:10:13 +01:00
Demeter Dzadik
de8e13036b Armature Make/Clear Parent: Grey out options that don't do anything
In armature edit mode, the Make/Clear Parent operators don't do anything
in various cases, but only one of these cases was previously indicated,
and it was indicated by hiding the option completely instead of graying
it out.

Clear Parent (Alt+P) problems fixed:
- "Clear Parent" option always showed up, even when none of the selected
  bones had a parent.
- "Disconnect Bone" option always showed up, even when use_connected on
  all selected bones was already false.

Make Parent (Ctrl+P) problems fixed:
- "Keep Offset" option didn't show up when all selected bones' parent
  was already the active bone. This was correct, and this patch tries to
  make all behaviours consistent with this.
- "Connected" option always showed up, even when all selected bones'
  parent was already the active bone, and they all had use_connect set
  to True.

With this patch all options show up all the time, but in cases where
they would do nothing, they will be grayed out.

Reviewed By: sybren

Differential Revision: https://developer.blender.org/D6100
2021-11-09 12:16:30 +01:00
Cody Winchester
4e2478940e Alembic: Allow exporting of animated vertex colors
Allow exporting of animated vertex colors to Alembic.

The changes are made to be in line with the way the UV Maps are written.
Each vertex color gets a OC4fGeomParam created and mapped into the
CDStreamConfig to avoid recreating the Param on each frame.

The time sample index is also stored in the config now and set onto the
UV and Vertex Color params each frame. Without this the exports would
get inconsistent timing results where animated UV maps and Vertex Colors
were not playing back at the original speed.

Reviewed By: sybren

Maniphest Tasks: T88074

Differential Revision: https://developer.blender.org/D11278
2021-11-09 10:54:13 +01:00
6b0a6c2ca9 Merge branch 'blender-v3.0-release' 2021-11-09 10:34:07 +01:00
ad679ee747 Merge branch 'blender-v3.0-release' 2021-11-09 10:07:44 +01:00
a7540f4b36 Merge branch 'blender-v3.0-release' 2021-11-09 17:11:35 +11:00
2772a033c9 Merge branch 'blender-v3.0-release' 2021-11-09 15:49:41 +11:00
8772a6fb9b Merge branch 'blender-v3.0-release' 2021-11-09 15:49:38 +11:00
Jeducious
d5d97e4169 Fix T92704: Redrawing while saving crashes outside the main thread
If the blend file is saved from a script in another thread,
like the render thread for example, Blender will crash on the call that
redraws the UI.

Ref D13140
2021-11-09 15:28:00 +11:00
e0dae0f98f Merge branch 'blender-v3.0-release' 2021-11-08 16:57:19 +01:00
fe2ed4a229 Merge branch 'blender-v3.0-release' 2021-11-08 18:44:39 +03:00
33beec1cec Cleanup: remove redundant arg when forcing zero initialization 2021-11-08 12:41:30 -03:00
e1c4e5df22 GPencil: New option to export PDF full scene
This new mode export all frames of the scene.

Reviewed By: pepeland

Differential Revision: https://developer.blender.org/D13055
2021-11-08 16:03:30 +01:00
c865577643 GPUTest: Add support to test on Windows.
On windows the OpenGL context wasn't activated when created, on Linux it
is. This patch will activate the context in gpu/draw test cases.
2021-11-08 15:49:51 +01:00
Jarrett Johnson
495e60c0da Basic engine shaders test
This patch adds shader compilation tests for the basic engine in `shaders_test.cc`

Addresses T92701

Reviewed By: jbakker

Differential Revision: https://developer.blender.org/D13066
2021-11-08 15:43:50 +01:00
bb6547cb5f Merge branch 'blender-v3.0-release' 2021-11-09 00:25:49 +11:00
c55d0ebea5 Merge branch 'blender-v3.0-release' 2021-11-09 00:25:46 +11:00
Bastien Montagne
d6e2210935 Fix inconsistent creation of NodeTreeTypeUndefined type of node tree
Currently, when creating a new node tree ID, its `typeinfo` is set to
`NodeTreeTypeUndefined`, but its `type` enum value is left to `0`,
aka `NTREE_SHADER`.

This patch adds a new `NTREE_UNDIFINED` value, and use it for
`NodeTreeTypeUndefined` types of node trees.

NOTE: While it is not clear whether that actually fixes issues currently,
quite a bit of code still relies on the value of `type`, so think it
makes sense to sanitize this.

NOTE: Would have been ideal to reserve `0` value to undefined type,
but at this point this is not possible anymore, so chose to use `-2` instead.

Reviewed By: JacquesLucke

Differential Revision: https://developer.blender.org/D13123
2021-11-08 12:24:47 +01:00
fc373af8f5 Nodes: store socket declaration reference in socket
Previously, to get the declaration of a socket, one had to go
through `node->declaration`. Now this indirection is not necessary
anymore. This makes it easier to add more per-socket information
into the declaration and accessing it in various places.

Currently, this system is used by socket descriptions and node warnings
for unsupported geometry component types.
2021-11-08 12:24:01 +01:00
09cef0fc00 Fix ASAN issue in image editor.
Adding virtual constructors to base classes.
2021-11-08 11:37:42 +01:00
ee4966d146 Merge branch 'blender-v3.0-release' 2021-11-08 17:19:08 +11:00
c3f5fca8a2 Cleanup: avoid error prone struct declarations in C++
Reference struct members by name instead relying on their order.
This also simplifies moving back to named members when all compilers
we use support them.
2021-11-08 17:00:36 +11:00
de581a2302 Fix reading the 3rd value of 2D cursors when transforming
Out of bounds read and potential out-of-bounds write when transforming
the 2D cursor for image editor and sequencer.

While this didn't cause user visible bugs in my tests,
it's error prone and should be avoided.

Use TransData2D for 2D cursors.
2021-11-08 17:00:36 +11:00
fb4b737518 CMake: add missing headers to CMake lists 2021-11-08 17:00:36 +11:00
27b37517f8 Cleanup: use static sets 2021-11-08 17:00:36 +11:00
42df2a7b57 Cleanup: Grammar in comments 2021-11-07 23:20:29 -06:00
0f80602632 Cleanup: remove references to non-existent 'mtexpoly' 2021-11-08 14:14:16 +11:00
b24a03e635 Cleanup: remove duplicate doc-strings
Internal struct ObTfmBack had out of sync doc-strings
for members duplicated from Object.

Remove the doc-strings as there is this is just temporary storage.
2021-11-08 14:14:15 +11:00
ed24b7d9a2 Cleanup: spelling in comments 2021-11-08 14:14:14 +11:00
3b726cfee8 Cleanup: compiler warnings 2021-11-08 14:14:13 +11:00
0654c41b0c Cleanup: use doxygen for BLF glyph
- Use doxy formatted functions.
- Use doxy sections.
2021-11-08 13:37:51 +11:00
4f387e66ac Merge branch 'blender-v3.0-release' 2021-11-07 20:45:19 -03:00
892da668dc Cleanup: Clang tidy 2021-11-07 00:39:20 -05:00
eddf5ad581 BLF: Refactor blf_glyph.c
Cleanup and Simplification of blf_glyph.c

See D13095 for details.

Differential Revision: https://developer.blender.org/D13095

Reviewed by Campbell Barton
2021-11-06 11:29:48 -07:00
6002914f14 Cleanup: Convert mesh.c to C++
This commit renames mesh.c to mesh.cc and makes
it compile in C++. Can be useful in the future to be able
to use C++ functionality in existing and new functions.

Differential Revision: https://developer.blender.org/D13134
2021-11-06 19:16:37 +01:00
f315a46982 Nodes: add preview image storage to node group
This is part of T92811.

Differential Revision: https://developer.blender.org/D13105
2021-11-06 16:43:26 +01:00
81baeec59b Cleanup: remove window_manager & editor includes from BLF
Remove the need to include the window manager & editor functions
in low level font rendering code.

- The default font size is now set when changed in the preferences.
- Flushing cache is set as a callback.
2021-11-06 16:28:02 +11:00
a804a11db1 Merge branch 'blender-v3.0-release' 2021-11-05 17:22:06 -05:00
48841c479f Merge branch 'blender-v3.0-release' 2021-11-05 18:00:01 -04:00
9e611c5616 Merge branch 'blender-v3.0-release' 2021-11-05 16:33:08 -05:00
f0bc7f3261 Merge branch 'blender-v3.0-release' 2021-11-05 20:40:02 +01:00
81bee0e75a UI: Fix minor theme mismatch
Pie menu got wrong item highlight and options settings were outdated.
2021-11-05 19:20:47 +01:00
3211c80a31 Fix T92815: Incorrect handling of evaluated meshes from curves
Evaluated meshes from curves are presented to render engines as
separate instance objects now, just like evaluated meshes from other
object types like point clouds and volumes. For that reason, cycles
should not consider curve objects as geometry (previously it did,
meaning it retrieved a second mesh from the curve object as well
as the temporary evaluated mesh geometry).

Further, avoid adding a curve object's evaluated mesh as data_eval,
since that is special behavior for meshes that is arbitrary. Adding an
evaluated mesh there but not an evalauted pointcloud is arbitrary,
for example. Retrieve the evaluated mesh in from the geometry set
in BKE_object_get_evaluated_mesh now, to support that change.

This gets us closer to a place where all of an object's evaluated data
is stored in geometry_set_eval, and we just have helper functions
to access specific geometry components.

Differential Revision: https://developer.blender.org/D13118
2021-11-05 19:20:47 +01:00
8d2a0d9b4c UI: Apply recent theme fixes for Preferences saved in 3.1 builds
Followup to e65230f0c0.

Pablo and I decided it's fine to reset themes again when saved with the
recent 3.1 builds.

This needed to be done a bit careful, since a normal version patch
resetting the theme would've reset the theme for anybody opening
preferences of a 3.0 build (even the final release build) in a 3.1
build. So make sure the theme is at least from a 3.1 build (but not
newer then this commit of course).
2021-11-05 19:00:10 +01:00
cc49c479a7 Cleanup: Use reference for non-optional C++ parameter
A reference makes clear that NULL is not an expected value. So it's the
prefered way of passing a `const` input parameter (at least if it may
not be cheap to copy).
2021-11-05 18:27:55 +01:00
aaf86bad87 Merge branch 'blender-v3.0-release' 2021-11-05 11:19:29 -05:00
212dcd6075 Merge branch 'blender-v3.0-release' 2021-11-05 10:57:04 -05:00
bbd8d33453 Merge branch 'blender-v3.0-release' 2021-11-05 16:23:56 +01:00
f415b41a94 Merge branch 'blender-v3.0-release' 2021-11-05 09:59:26 -05:00
625b2f59f0 Fix GCC warnings after own recent commit
Caused by 4e09fd76bc.
2021-11-05 15:32:11 +01:00
2986924301 Cleanup: Remove misleading comment
Python isn't doing any conversion here. We just do a regular lookup of
the given enum identifier in the RNA enum definition.
2021-11-05 15:32:11 +01:00
29efd26e71 Cleanup: Remove redundant scope qualifiers. 2021-11-05 15:18:58 +01:00
885c79915f Merge branch 'blender-v3.0-release' 2021-11-05 15:08:09 +01:00
a0f50c1890 Merge branch 'blender-v3.0-release' 2021-11-05 15:06:10 +01:00
4e09fd76bc Cleanup (UI): Add/use type for operator context enum
Adds a `wmOperatorCallContext` typedef for the existing `WM_OP_XXX`
operator context enum. This adds type safety, allows the compiler to
produce better warnings and helps understanding what a variable is for.

Differential Revision: https://developer.blender.org/D13113

Reviewed by: Campbell Barton
2021-11-05 14:57:26 +01:00
35198606d5 Cleanup: Remove wrong comments in versioning
This comment is from the block at the end of the versioning functions,
where we have an unversioned block to collect versioning code that
doesn't require immediate version bumping. The comment was probably just
copied over with the code when bumping the version eventually.
2021-11-05 14:53:57 +01:00
32c90d2d7c Cleanup: Split image engine into ImageEngine, SpaceAccessor and DrawingMode.
Image engine is used to draw an image into a space. The current
structure wasn't clear and couldn't be easilly extended. This refactor
spliced the image draw engine into 3 main components.

- Space accessors: contains an interface to communicate with space data
  (Image editor, UV Editor, Node Editor) in a common way. This reduced
  the branching in the code base.
- DrawingMode: contains an interface to the used tactic to draw an image
  inside the space framebuffer. Currently only one mode is implemented;
  in the future there could be a separate drawing mode for huge images.
- ImageEngine: the core that connects the draw manager with the space
  data and drawing mode.
2021-11-05 13:25:00 +01:00
d7f4fdf845 GPencil: Fix dash modifier missing vertex color.
The original code did not copy vertex color to the generated stroke,
now fixed.
2021-11-05 20:02:35 +08:00
016a575002 Merge branch 'blender-v3.0-release' 2021-11-05 16:05:00 +11:00
3a4dade2f8 Merge branch 'blender-v3.0-release' 2021-11-05 16:04:57 +11:00
b9968b83ad Merge branch 'blender-v3.0-release' 2021-11-05 16:04:52 +11:00
40d090cc72 Merge branch 'blender-v3.0-release' 2021-11-05 16:04:43 +11:00
5cc21b095a Merge branch 'blender-v3.0-release' 2021-11-05 16:04:40 +11:00
7061d1e39f Fix T92740: Missing lock around the image CacheLimiter
A recent change exposed this long-standing race. Simply protect the
MEM_CacheLimiter with its lock now. Additionally, guard against
unmanaging an already destroyed cache handle.

Ref T92740, T92838
2021-11-04 20:58:32 -07:00
1b6238edba Merge branch 'blender-v3.0-release' 2021-11-05 13:45:32 +11:00
5b7a14c019 Cleanup: use doxy sections in graph editor module 2021-11-05 13:38:58 +11:00
9d2e325694 Cleanup: spelling/typos 2021-11-05 13:38:58 +11:00
Christoph Lendenfeld
3364a5bea6 Cleanup: move code in graph_slider_ops
Future operators can use the same code,
so it is moved up to disassociate it from decimate

No functional changes

Reviewed by: Sybren A. Stüvel
Differential Revision: https://developer.blender.org/D12489
Ref: D12489
2021-11-04 21:35:28 +00:00
Christoph Lendenfeld
6986b43b3d Cleanup: renames in graph_slider_ops
This patch renames:
* tDecimateGraphOp to tGraphSliderOp
* dgo to gso (to match with the struct rename)
* decimate_reset_bezts to reset_bezts to indicate it can be used by other functions

No functional changes

Reviewed by: Sybren A. Stüvel
Differential Revision: https://developer.blender.org/D12490
Ref: D12490
2021-11-04 21:20:26 +00:00
df3e30398f Merge branch 'blender-v3.0-release' 2021-11-04 16:56:32 -03:00
556c71a84a GPencil: New option to Merge All layers in active one
This new option allows to combine all layers in the active one. Also the merge down option has been improved.

Reviewed By: mendio, pablo vazquez (UI)

Differential Revision: https://developer.blender.org/D13054
2021-11-04 19:40:55 +01:00
1c6d3d614a Minor tweaks to new append code.
* Name generated 'append' collection, instead of getting a generic
  meaningless name.
* Do not check if a collections's objects are already instantiated, when
  we already know that we want to instantiate that collection.
2021-11-04 19:00:22 +01:00
101fa4a425 Merge branch 'blender-v3.0-release' 2021-11-04 18:32:37 +01:00
2373ce7fcf Merge branch 'blender-v3.0-release' 2021-11-04 17:27:51 +01:00
a72b26527d Merge branch 'blender-v3.0-release' 2021-11-04 16:50:53 +01:00
2becb3e9af Merge branch 'blender-v3.0-release' 2021-11-04 10:10:52 -05:00
37b862fa6c Cleanup: Remove operator context override for drop-box operators
Drop-boxes should act on the context determined through the exact cursor
location. There should be no need to override that, basically by the
nature of how drop-boxes work.
So Campbell and I agreed on removing this.

If we wanted to support it, we'd have to restore the operator context
when drawing drop-boxes, see
https://developer.blender.org/T92501#1247581.
2021-11-04 16:02:54 +01:00
db43d19c16 Merge branch 'blender-v3.0-release' 2021-11-04 15:32:50 +01:00
d3328fabc9 Attempt to fix failing lib link test on windows.
According to https://docs.python.org/3/library/os.html#os.rename,
`os.rename` has os-specific behavior, and will fail in case you attempt
to rename to an existing file on windows.

So using `os.replace` instead, which should be os-agnostic.

NOTE: Fact that temp test directory is not cleared after tests are
sucessfully ran does not sound great...
2021-11-04 15:06:58 +01:00
8eff1eca52 Merge remote-tracking branch 'origin/blender-v3.0-release' 2021-11-04 14:50:48 +01:00
c641107c95 Fix T92800: (UI) Radial control values get wrong color assigned
The global theme state didn't get updated or unset properly when drawing
overlays. Now paint cursors use the theme settings of the space they are
in, while global overlays use the global fallback, which is the main 3D
View region.
2021-11-04 11:52:11 +01:00
978ef093db Cleanup: fix compiler warnings
Those were introduced in rBccead2ed9c6121c42a516712da38a2faec877e2f.
2021-11-04 11:29:16 +01:00
2d6d8fc7ca Attempt fix for new lib reload/relocate tests on windows.
Try splitting them into their own class.
2021-11-04 11:25:30 +01:00
82b20b6975 Merge branch 'blender-v3.0-release' 2021-11-04 10:51:40 +01:00
feaf5b95e0 Merge branch 'blender-v3.0-release' 2021-11-04 08:35:46 +01:00
36a6528723 Merge branch 'blender-v3.0-release' 2021-11-04 18:10:07 +11:00
d7f9f083d4 Merge branch 'blender-v3.0-release' 2021-11-04 14:51:58 +11:00
7c188d8241 Merge branch 'blender-v3.0-release' 2021-11-04 14:51:55 +11:00
9cd5b3c9b6 Merge branch 'blender-v3.0-release' 2021-11-04 14:51:52 +11:00
682f1548be Merge branch 'blender-v3.0-release' 2021-11-04 14:51:50 +11:00
c0fdaf700a Fix the GPencil stroke not sticking to other strokes
Issue seen when setting `Stroke Placement` of type `Stroke`.

Regression introduced in {rBaa0ac0035a0d3601672a0c732e3f8f932a36fc04}.
2021-11-03 18:51:25 -03:00
0c6b815855 Geometry Nodes: Add Length Output to Curve Parameter Node
Adds a length output to the curve parameter node which returns the
length of a spline at each point, or the length of the curve at
each spline depending on the domain.

Differential Revision: https://developer.blender.org/D12882
2021-11-03 15:05:46 -05:00
d6ed9c2b40 Merge branch 'blender-v3.0-release' 2021-11-03 14:41:50 -05:00
431524aebc Geometry Nodes: Selection outputs for Cone and Cylinder
This adds Top, Bottom and Side selections to the Primitive
Mesh nodes Cone and Cylinder.
2021-11-03 20:20:15 +01:00
ccead2ed9c Spreadsheet: Display geometry volume component grids
This shows a geometry's volume grids in the spreadsheet.
Three columns are displayed:
 - Name: The text name of each grid
 - Data type: Float, Vector, etc.
 - Class: Fog volume, Level Set, or unkown

In the future, values of the voxels themselves could be displayed,
but that is a much more complex problem, with important performance
implications, etc.

Differential Revision: https://developer.blender.org/D13049
2021-11-03 13:45:51 -05:00
Nikhil Shringarpurey
4e5537d841 Geometry Nodes: Add tooltips to primitive node inputs
Building on the work in rBef45399f3be0, this commits adds
tooltips to the inputs for the default primitives nodes.

Differential Revision: https://developer.blender.org/D12640
2021-11-03 13:25:44 -05:00
7aaedc09c7 Merge branch 'blender-v3.0-release' 2021-11-03 18:16:16 +01:00
e10caf6fe3 Merge remote-tracking branch 'origin/blender-v3.0-release' 2021-11-03 18:07:36 +01:00
a7672caeb2 Geometry Nodes: Add Selection Input to Resample Curves Node
Add a boolean input to the resample curve node that indicates which
splines should be resampled and which should be unchanged.

Differential Revision: https://developer.blender.org/D13064
2021-11-03 11:55:24 -05:00
de2988ea1b Cleanup: Remove effect-less const
Using `const` on an enum type returned by value doesn't have an effect.
2021-11-03 17:28:52 +01:00
debf4b70db Cleanup: Avoid redundant template parameter in BLI serializing API
The `ContainerValue` template can obtain the type of the contained value
via the given `Container` type, simply using `Container::value_type`.
Use this as the default way to determine the value type which simplifies
using the template. If necessary the value type can be passed explicitly
still.
2021-11-03 17:28:20 +01:00
Germano Cavalcante
aa0ac0035a GPencil and Annotation: Use cached depth to perform depth testing operations
Operations such as erasing with occlusion and drawing on the surface
require reading the depth buffer.

However, this is being done with minimal efficiency.

Currently, to read the depth corresponding to each point of the new stroke,
a ReadPixel is called to send a message to the GPU and read the depth of
the corresponding pixel in the VRAM.

The communication between GPU and CPU is known to be a slow operation so
it is good to be avoided.

Therefore, save the entire depth buffer in a cache to be read directly
from the RAM.

(Also the `ED_view3d_autodist_depth` and `ED_view3d_autodist_depth_seg` have
been removed since they are no longer used).

Reviewed By: antoniov, fclem

Differential Revision: https://developer.blender.org/D10894
2021-11-03 12:10:37 -03:00
8b516d8712 Include node name for socket animation channel UI
The channel names were often indistingushable in animation editors.
Now include the node _name_ (unfortunately, getting the _label_ could
result in bad performance in some circustances -- see previous version
of D13085).
Similar to what rB77744b581d08 did for some VSE strip properties.

ref. T91917

Maniphest Tasks: T91917

Differential Revision: https://developer.blender.org/D13085
2021-11-03 15:03:40 +01:00
2a88343213 Merge branch 'blender-v3.0-release' 2021-11-03 14:24:33 +01:00
04d35f9315 Merge branch 'blender-v3.0-release' 2021-11-03 10:23:55 -03:00
3532da44ee Merge branch 'blender-v3.0-release' 2021-11-03 22:26:39 +11:00
5095e4fc22 UI: Display disabled-hint when dragging material outside object mode
Display a "disabled hint" (text explaining why something isn't possible)
when dragging a material over the 3D View, while being in edit mode or
so (anything that isn't object mode).
2021-11-03 11:57:05 +01:00
bdf6665e3a Merge branch 'blender-v3.0-release' 2021-11-03 10:54:53 +01:00
e1f1b0841d Merge branch 'blender-v3.0-release' 2021-11-03 15:56:11 +11:00
cea7ee7582 Merge branch 'blender-v3.0-release' 2021-11-03 11:49:58 +11:00
ac0eefe26f Merge branch 'blender-v3.0-release' 2021-11-02 18:22:50 -05:00
c4b73847d3 BLF: Remove Thread Locking For Font Metrics
This patch removes the need to lock the thread just to get to some
generic (not glyph-specific) font metrics.

See D12976 for more details.

Differential Revision: https://developer.blender.org/D12976

Reviewed by Campbell Barton
2021-11-02 14:28:25 -07:00
8c58838f6a Merge branch 'blender-v3.0-release' 2021-11-02 15:23:59 -05:00
5cd1210b52 Merge branch 'blender-v3.0-release' 2021-11-02 13:44:17 -05:00
3dcd042bfc Merge branch 'blender-v3.0-release' 2021-11-02 18:38:59 +01:00
53fdde3f64 Merge branch 'blender-v3.0-release' 2021-11-02 18:14:59 +01:00
0a254109b8 Merge branch 'blender-v3.0-release' 2021-11-02 17:52:03 +01:00
242ad4bd0f Merge branch 'blender-v3.0-release' 2021-11-02 17:49:52 +01:00
27621490c2 Merge branch 'blender-v3.0-release' 2021-11-02 18:33:25 +02:00
12bf4adbe3 Merge branch 'blender-v3.0-release' 2021-11-02 17:05:03 +01:00
2c23256288 Merge branch 'blender-v3.0-release' 2021-11-02 16:03:17 +00:00
c01b3c534b Merge branch 'blender-v3.0-release' 2021-11-02 15:42:09 +01:00
1b2342b4d3 Tests: Add basic unittest for library reload and relocate operators. 2021-11-02 15:33:30 +01:00
a7e92843f7 Fix: Build error on windows.
External symbols in C files need
to be marked as such otherwise
the linker will not find them.
2021-11-02 07:43:44 -06:00
e64d4d0200 Merge branch 'blender-v3.0-release' 2021-11-02 07:59:35 -05:00
0daf429591 Merge branch 'blender-v3.0-release' 2021-11-02 23:38:01 +11:00
b7dc667eb2 Merge branch 'blender-v3.0-release' 2021-11-02 12:31:05 +01:00
7a4ee2fd4f Merge branch 'blender-v3.0-release' 2021-11-02 22:06:35 +11:00
52f4a908f7 Removed compilation warning nullptr check in image engine. 2021-11-02 12:00:07 +01:00
ffd3dd6376 Merge branch 'blender-v3.0-release' 2021-11-02 11:17:53 +01:00
a2f5a10129 Merge branch 'blender-v3.0-release' 2021-11-02 11:09:09 +01:00
223f2b27d1 Merge branch 'blender-v3.0-release' 2021-11-02 09:57:16 +01:00
21e168069d Merge branch 'blender-v3.0-release' 2021-11-02 19:52:01 +11:00
8ca6e51ade Cleanup: clang-tidy 2021-11-02 19:50:53 +11:00
47d12268e3 Cleanup: Change image engine to CPP.
Added namespace blender::draw::image_engine. Code is still C-a-like.
Only changed the obvious code style to CPP (structs, nullptr, casts).
2021-11-02 09:16:33 +01:00
9cc05fe9c4 Merge branch 'blender-v3.0-release' 2021-11-02 16:08:13 +11:00
af9e0409f1 Merge branch 'blender-v3.0-release' 2021-11-02 16:08:10 +11:00
7b436ead6b Merge branch 'blender-v3.0-release' 2021-11-02 16:08:07 +11:00
5363437555 Merge branch 'blender-v3.0-release' 2021-11-02 00:46:32 -03:00
75f5edcaf3 Merge branch 'blender-v3.0-release' 2021-11-01 14:48:20 -05:00
6ddbcaa096 Merge branch 'blender-v3.0-release' 2021-11-01 14:42:49 -05:00
e045249a28 Merge branch 'blender-v3.0-release' 2021-11-01 19:54:58 +01:00
85176c86f0 Fix T92722: Error when saving new render preset
The preset menu name was not renamed in 4ddad5a7ee.
2021-11-01 19:25:11 +01:00
765c2cc6c7 Merge branch 'blender-v3.0-release' 2021-11-01 13:17:54 -05:00
339fd8027f Merge branch 'blender-v3.0-release' 2021-11-01 17:28:40 +01:00
adc540cf7c Fix T92655: spreadsheet_duplicate Split Exception
Check SpaceSpreadsheet's runtime is not null when trying to duplicate
the data when doing an area split.

See D13047 for further details.

Differential Revision: https://developer.blender.org/D13047

Reviewed by Jacques Lucke
2021-11-01 08:52:00 -07:00
1704a394d8 Fix T92689: Assert loading file with a sound.
Dummy mistake in rBc8c53ceecc30 (boolean inversion).
2021-11-01 16:44:12 +01:00
d56d3fc6b1 Merge branch 'blender-v3.0-release' 2021-11-01 16:43:26 +01:00
6acba759e0 Merge branch 'blender-v3.0-release' 2021-11-01 11:58:11 -03:00
8fbbd69946 Fix T92629: Crash on mesh separate after rB43bc494892c3.
rB43bc494892c3 switched `BKE_libblock_relink_to_newid` to use new ID
remapping and libquery code.

However, that new code does protect by default against remapping an
objects's data pointer when that object is in Edit mode, since this is
not a behavior that generic BKE code can handle (due to required editing
data for most obdata types when in edit mode).

So specific code that does create new IDs and need remapping in Edit
mode has to pass specific exception flags to remaping code.

This commit adds those remapping flags to `BKE_libblock_relink_to_newid`
and add said exception flag to the remapping call from
`ED_object_add_duplicate` when the object is in edit mode.
2021-11-01 15:46:28 +01:00
e85e126e3f IDRemap: Add option to force remapping obdata in edit mode.
In theory we should never allow remapping of Objects' obdata ID pointer
when the object is in Edit mode. But there are some cases were this is
needed, so adding yet another exception option to remapping flags.

Preliminary change to fix T92629.
2021-11-01 15:46:28 +01:00
64de6ad4fe Fix use-after-free in image code 2021-11-01 15:36:09 +01:00
2fb725ea30 Cleanup: Unused argument
Fixes strict compiler warnings.
2021-11-01 15:14:49 +01:00
9de4f64197 Fix compile error on Windows
Caused by 7150f919d3. This undid part of 79a88b5e91. Added a comment
for why this include is needed, to avoid this error from happening
again.
2021-11-01 14:44:18 +01:00
6897c2141e Merge branch 'blender-v3.0-release' 2021-11-01 13:39:24 +01:00
7150f919d3 Cleanup: Remove unused headers in asset files
Also move system includes first, like we have it elsewhere in Blender.
2021-11-01 13:21:07 +01:00
0eb63328e8 Merge branch 'blender-v3.0-release' 2021-11-01 13:02:48 +01:00
d07e3bde20 Fix Cycles tests after recent logging changes
The constant folding tests rely on logging sync.
2021-11-01 12:47:03 +01:00
9111ea78ac Localize image mutex lock into runtime field of Image datablock
Allows to avoid a global lock being held while reading files from disk,
solving performance issues when Cycles needs to read a lot of packed
images.

Simple test file F11597666

Differential Revision: https://developer.blender.org/D13032
2021-11-01 12:47:03 +01:00
b6dd5be213 Merge branch 'blender-v3.0-release' 2021-11-01 22:25:27 +11:00
b5eada7f69 Merge branch 'blender-v3.0-release' 2021-11-01 22:25:24 +11:00
3f0991266f Merge branch 'blender-v3.0-release' 2021-11-01 12:15:09 +01:00
a96b2f39b8 Geometry Nodes: improve check if object has geometry set instances
The improves playback speed in my instance heavy scene from ~3.7 fps to ~3.9 fps.
2021-11-01 12:00:41 +01:00
Jarrett Johnson
7dd84f05aa Pointcloud selection support
This patch adds support for selecting pointclouds.

Since pointclouds were not properly drawn to the selection buffer (as diagonsed by output from `glReadPixels` and Renderdoc), they were not able to be selectable by depth picking or occlusion queries. In `basic_engine`, objects were rendered with a shader which draws to a depth buffer but only assumes a single position vertex attribute. Pointclouds, though, require at least another vertex attribute `pos_inst` which provides the instance offsets. Thus, this patch adds another shader variant for pointclouds which supports these two attributes and renders the points appropriately.

{F11652666}

Addresses T92415

Reviewed By: fclem

Differential Revision: https://developer.blender.org/D13059
2021-11-01 11:41:13 +01:00
fe44001215 Revert "Pointcloud selection support"
This reverts commit a50f8b3fd8.
2021-11-01 11:40:34 +01:00
Jarrett Johnson
a50f8b3fd8 Pointcloud selection support
This patch adds support for selecting pointclouds.

Since pointclouds were not properly drawn to the selection buffer (as diagonsed by output from `glReadPixels` and Renderdoc), they were not able to be selectable by depth picking or occlusion queries. In `basic_engine`, objects were rendered with a shader which draws to a depth buffer but only assumes a single position vertex attribute. Pointclouds, though, require at least another vertex attribute `pos_inst` which provides the instance offsets. Thus, this patch adds another shader variant for pointclouds which supports these two attributes and renders the points appropriately.

{F11652666}

Addresses T92415

Reviewed By: fclem

Differential Revision: https://developer.blender.org/D13059
2021-11-01 11:25:23 +01:00
8379eefafb Cycles: Enable debug symbols for Clang
Debug symbols were disabled for Clang at some point due to link issues.
This is no longer the case for any reasonably modern version of Clang.
So this patch removes the check in question.

Differential Revision: https://developer.blender.org/D13045

Reviewed By: brecht
2021-11-01 11:38:03 +02:00
5327413b37 Cleanup: Remove Cycles device checks for half float.
All supported devices support half float now, so we can remove the check.

Differential Revision: https://developer.blender.org/D13021
2021-11-01 10:18:30 +01:00
289f013e64 Merge branch 'blender-v3.0-release' 2021-11-01 17:48:50 +11:00
49a0453799 Merge branch 'blender-v3.0-release' 2021-11-01 17:22:44 +11:00
de4793e0e6 Merge branch 'blender-v3.0-release' 2021-11-01 17:16:53 +11:00
55ce05e0bb Merge branch 'blender-v3.0-release' 2021-11-01 17:16:50 +11:00
346a812d7e Fix scale cage gizmo in pose-mode
The active objects matrix was ignored when calculating the cage.
2021-11-01 17:16:45 +11:00
1e749d0602 Cleanup: spelling, use C comments 2021-11-01 14:00:58 +11:00
b99d6e1bed Fix errors in BKE_appdir_font_folder_default
- Missing NULL check for the HOME environment variable.
- The user preference path was written to even when the
  path didn't exist.
2021-11-01 13:43:35 +11:00
e2937ff24f Merge branch 'blender-v3.0-release' 2021-11-01 13:15:10 +11:00
c312c71969 Fix: Build error on all platforms
Types were used before being declared.
2021-10-31 11:52:24 -06:00
Christoph Lendenfeld
b2e9f35c5e Cleanup: Extract function to store bezt arrays
The code to store an original bezt array previously lived in
`graphkeys_decimate_invoke`.
Since future graph slider operators will need this function as well,
it has been extracted.

No functional changes.

Reviewed by: Sybren A. Stüvel
Differential Revision: https://developer.blender.org/D12487
Ref: D12487
2021-10-31 11:28:10 +00:00
Christoph Lendenfeld
1b6daa871d Cleanup: Extract keyframe filter to constant
An int flag is used to filter animation channels for
operators to work on. The flag was duplicated multiple times.
This patch removes the duplication by creating a constant

Reviewed by: Sybren A. Stüvel
Differential Revision: https://developer.blender.org/D12486
Ref: D12486
2021-10-31 11:19:40 +00:00
Christoph Lendenfeld
4e502bb6d2 Merge branch 'blender-v3.0-release' 2021-10-31 11:10:11 +00:00
ae9052a33e Cleanup: Simplify logic for adding grid in points to volume node
Instead of creating a separate grid first and then merging the points
to volume grid, use the recently added `BKE_volume_grid_add_vdb`
helper function for this purpose.
2021-10-30 17:26:18 -05:00
9cfffe8468 UI: Open File Browser with Thumbnails for Fonts
When browsing to open a font file, open File Browser in Thumbnail View
and sorting by name, instead of using the last-used states.

See D13040 for more details.

Differential Revision: https://developer.blender.org/D13040

Reviewed by Julian Eisel
2021-10-30 14:02:34 -07:00
03a962d8ca Merge branch 'blender-v3.0-release' 2021-10-30 16:45:04 +02:00
02a9377da0 UI: Default Fonts Folder for Mac and Linux
Initial defaults for userdef->fontdir for Mac and Linux.

See D12802 for more details.

Differential Revision: https://developer.blender.org/D12802

Reviewed by Campbell Barton
2021-10-29 17:15:22 -07:00
dcdbaf89bd Asset Browser: Correct name & tooltip for asset list refresh operator
The name and tooltip were talking about file-lists, which exposes the
fact that the Asset Browser uses the File Browser code in the UI, which
we shouldn't do. This can confuse users.
Instead have a dedicated operator for the Asset Browser with a proper
name and tooltip.
2021-10-29 16:55:20 +02:00
0c3da7f370 Merge branch 'blender-v3.0-release' 2021-10-29 09:42:24 -05:00
837b72fa83 Merge branch 'blender-v3.0-release' 2021-10-29 11:25:32 -03:00
d18d87d3e7 Fix T92576: Crash switching from Asset Browser to File Brower
The asset catalog filtering data needs to be cleared when with the other
asset library data of the file list. This is done when changing between
asset and file browser (and in other cases).
2021-10-29 16:00:43 +02:00
e1db6dc11b Fix crash on undo after recent lib_query refactor.
Forgot that scene uses part of its ID looping code for specific undo
handling.

Caused by rBe3b2f0fd6ff9.
2021-10-29 12:54:32 +02:00
c8c53ceecc LibQuery: Remove last 'bool returns' from public API.
Those were used in a very few places to detect whether iteration should
be stopped or not, but one can use `BKE_lib_query_foreachid_iter_stop`
now for that.

Also fix early break handling in embedded IDs processing.

Fix T90922: Fix return policy inconsistency in `scene_foreach_id`.
2021-10-29 11:32:26 +02:00
259731909c LibQuery: Fix potential memleak in recursive case.
In case `library_foreach_ID_link` would return early in recursive
process, it would not properly free its utils data.

Also add proper iteration break in case some sub-calls requested it.

Finally, make this function return a boolean to know whether iteration
should be stopped or not (will be used in future commit to fix this
handling in embedded IDs case).

Part of T90922: Fix return policy inconsistency in `scene_foreach_id`.
2021-10-29 11:32:26 +02:00
e3b2f0fd6f LibQuery: Add macro to help break looping when requested.
The new `BKE_LIB_FOREACHID_PROCESS_FUNCTION_CALL` execute the given
statement and then check status of `LibraryForeachIDData` data, and
return in case stop of iteration is requested.

This is very similar to the other `BKE_LIB_FOREACHID_PROCESS_` existing
macros, and allows us to properly break iteration when a sub-function
has requested it.

Part of T90922: Fix return policy inconsistency in `scene_foreach_id`.
2021-10-29 11:32:26 +02:00
51c1c1cd93 Fix potential early-return in WM foreach_id process.
Add a function to check if iteration over ID usages should stop (using
internal `IDWALK_STOP` status flag).

Use it in `BKE_LIB_FOREACHID_PROCESS_` macros, and in
`window_manager_foreach_id` to handle properly the active workspace case
(previous code could skip the call to `BKE_workspace_active_set` in case
iteration over ID usages was stopped by callback on that specific ID
usage).

Part of T90922: Fix return policy inconsistency in `scene_foreach_id`.
2021-10-29 11:32:26 +02:00
eae59645de Cleanup: Add some comments to some sub-function of foreach_id process. 2021-10-29 11:32:26 +02:00
c112418e95 Merge branch 'blender-v3.0-release' 2021-10-29 11:05:20 +02:00
43bc494892 IDManagement: Remove deprecated BKE_libblock_relink_to_newid usages.
Move all usages to new `BKE_libblock_relink_to_newid_new`, and rename
that one to `BKE_libblock_relink_to_newid`.

Fix T91413.
2021-10-29 10:45:48 +02:00
fb688c8d5c Merge branch 'blender-v3.0-release' 2021-10-29 10:24:37 +02:00
261bb766fb Merge branch 'blender-v3.0-release' 2021-10-29 10:11:26 +02:00
1688cb27cd Merge branch 'blender-v3.0-release' 2021-10-29 09:49:48 +02:00
657923cf93 Merge branch 'blender-v3.0-release' 2021-10-29 15:01:50 +11:00
57f7650dc7 Merge branch 'blender-v3.0-release' 2021-10-29 14:01:19 +11:00
99fbf1716f Merge branch 'blender-v3.0-release' 2021-10-29 13:54:43 +11:00
38fc19d643 Cleanup: rename blf_utf8_next_fast to blf_glyph_from_utf8_and_step
Calling this 'fast' no longer made sense as the slower code-path
has been removed.
2021-10-29 13:49:31 +11:00
1e2589bfa5 Cleanup: remove redundant BLI_UTF8_ERR check 2021-10-29 13:15:39 +11:00
0e71162e68 Cleanup: resolve cast warnings 2021-10-29 13:14:23 +11:00
70947ebc65 BLF Refactor: blf_utf8_next_fast
Simplification of BLF glyph loading

See D13026 for details.

Differential Revision: https://developer.blender.org/D13026

Reviewed by Campbell Barton
2021-10-28 18:49:21 -07:00
59534dbee2 BLF Refactor: blf_kerning_step_fast
Simplification of BLF Kerning

See D13015 for more details.

Differential Revision: https://developer.blender.org/D13015

Reviewed by Campbell Barton
2021-10-28 18:20:34 -07:00
0283a22456 Merge branch 'blender-v3.0-release' 2021-10-29 10:53:42 +11:00
8eb10de739 Merge branch 'blender-v3.0-release' 2021-10-29 10:53:39 +11:00
adf82fe943 Merge branch 'blender-v3.0-release' 2021-10-28 22:42:04 +02:00
fc36772b06 Tests: minor updates to benchmark script for running on buildbot
* graph command accepts folder of json files as input
* reset command clears log files
2021-10-28 22:41:40 +02:00
731926e70e Merge branch 'blender-v3.0-release' 2021-10-28 14:25:14 -05:00
a6af0e570d Merge branch 'blender-v3.0-release' 2021-10-28 21:20:25 +02:00
ec9357a94e Merge branch 'blender-v3.0-release' 2021-10-28 16:02:06 +02:00
60b278a3bb Merge branch 'blender-v3.0-release' 2021-10-28 15:52:19 +02:00
5f1107ffaf Merge branch 'blender-v3.0-release' 2021-10-28 12:01:01 +02:00
2f8ed53d6f Merge remote-tracking branch 'origin/blender-v3.0-release' 2021-10-28 11:23:25 +02:00
c1cfb475b3 LineArt: Fix(unreported) Material mask panel logic
The logic should be: show material mask panel if in_front is on,
it was inverted unintentionally.
2021-10-28 17:17:48 +08:00
4b57d5a9a0 LineArt: Fix(unreported) depsgraph camera error
This fixes unintentional line art error when custom camera doesn't
exist, now not adding custom camera relation in this case.
2021-10-28 17:13:22 +08:00
4adde62f60 Python doc generator: add missing selected_ids context key
Add the context key I introduced in rB03c0581c6ed to the Python API docs
generator.

No functional changes to Blender.
2021-10-28 10:32:13 +02:00
2501d00268 Python doc generator: raise explanatory error when context key is missing
When a new key is added to the context, it also needs to be added to the
`sphinx_doc_gen.py` file for generating the Python API documentation.
When this isn't done, the script would raise a generic `KeyError`. Now
it explains what needs to be updated to solve the problem.

No functional changes to Blender.
2021-10-28 10:32:13 +02:00
45439dfe4c Preferences: remove special case for copying previous settings
This was only needed for skipping version numbers when the numbering
scheme changed for 3.0.
2021-10-28 17:59:51 +11:00
5568599015 Merge branch 'blender-v3.0-release' 2021-10-28 17:49:15 +11:00
a7879dea7c Merge branch 'blender-v3.0-release' 2021-10-28 17:49:12 +11:00
43c603c2ff Merge branch 'blender-v3.0-release' 2021-10-28 17:45:28 +11:00
4979537a65 Merge branch 'blender-v3.0-release' 2021-10-28 16:56:11 +11:00
99b6127b73 Merge branch 'blender-v3.0-release' 2021-10-28 16:56:08 +11:00
c647bd899f Merge branch 'blender-v3.0-release' 2021-10-28 14:08:57 +11:00
a2f0f98271 Merge branch 'blender-v3.0-release' 2021-10-27 19:07:40 -03:00
7d2c759054 Merge branch 'blender-v3.0-release' 2021-10-27 23:22:37 +02:00
8f02de3de7 Cleanup: remove redundant variable
`free_tooltip` is no longer needed.
2021-10-27 18:11:47 -03:00
e4a5fd4298 Fix broken Python API doc generation after addition of selected_ids 2021-10-27 22:16:36 +02:00
39c11c03d0 Asset Browser: Increase size of search button a bit
Before this, the search button was quite small really, not much text would fit
into it. Increase the size a bit, but not too much to still make the layout
work in smaller area sizes.
2021-10-27 18:56:41 +02:00
3e32a68f38 UI: Refactor how dragging onto text buttons works, fixing issues
There was a bunch of special handling to support dropping data-blocks onto
string or search-menu buttons, to change the value of these. This refactor
makes that case use the normal drop-box design, where an operator is executed
on drop that gets input properties set by the drop-box. This should also make
it easier to add support for dragging assets into these buttons.

In addition this fixes an issue: Two tooltips were shown when dragging assets
over text buttons. None should be shown, because this isn't supported.
2021-10-27 18:51:44 +02:00
bfd2921d38 Revert "Blender 3.0 bcon3 (beta)"
This reverts commit f7a3450e63.
2021-10-27 18:44:28 +02:00
c346bb1990 Revert "3.0 splashscreen"
This reverts commit 78c1c71988.
2021-10-27 18:43:27 +02:00
17efd14682 Merge branch 'blender-v3.0-release' 2021-10-27 18:43:18 +02:00
dab3591588 Blender 3.1 bcon1 - alpha
Bump the version number for the new release cycle.
2021-10-27 18:40:49 +02:00
762 changed files with 19511 additions and 12300 deletions

View File

@@ -269,5 +269,9 @@ StatementMacros:
- PyObject_HEAD
- PyObject_VAR_HEAD
StatementMacros:
- GPU_STAGE_INTERFACE_CREATE
- GPU_SHADER_DESCRIPTOR
MacroBlockBegin: "^BSDF_CLOSURE_CLASS_BEGIN$"
MacroBlockEnd: "^BSDF_CLOSURE_CLASS_END$"

View File

@@ -12,6 +12,8 @@ Checks: >
-readability-avoid-const-params-in-decls,
-readability-simplify-boolean-expr,
-readability-make-member-function-const,
-readability-suspicious-call-argument,
-readability-redundant-member-init,
-readability-misleading-indentation,
@@ -25,6 +27,8 @@ Checks: >
-bugprone-branch-clone,
-bugprone-macro-parentheses,
-bugprone-reserved-identifier,
-bugprone-easily-swappable-parameters,
-bugprone-implicit-widening-of-multiplication-result,
-bugprone-sizeof-expression,
-bugprone-integer-division,
@@ -40,7 +44,8 @@ Checks: >
-modernize-pass-by-value,
# Cannot be enabled yet, because using raw string literals in tests breaks
# the windows compiler currently.
-modernize-raw-string-literal
-modernize-raw-string-literal,
-modernize-return-braced-init-list
CheckOptions:
- key: modernize-use-default-member-init.UseAssignment

View File

@@ -411,6 +411,7 @@ option(WITH_CYCLES "Enable Cycles Render Engine" ON)
option(WITH_CYCLES_OSL "Build Cycles with OpenShadingLanguage support" ON)
option(WITH_CYCLES_EMBREE "Build Cycles with Embree support" ON)
option(WITH_CYCLES_LOGGING "Build Cycles with logging support" ON)
option(WITH_CYCLES_DEBUG "Build Cycles with options useful for debugging (e.g., MIS)" OFF)
option(WITH_CYCLES_STANDALONE "Build Cycles standalone application" OFF)
option(WITH_CYCLES_STANDALONE_GUI "Build Cycles standalone with GUI" OFF)
@@ -1069,7 +1070,7 @@ if(MSVC)
add_definitions(-D__LITTLE_ENDIAN__)
# OSX-Note: as we do cross-compiling with specific set architecture,
# endianess-detection and auto-setting is counterproductive
# endianness-detection and auto-setting is counterproductive
# so we just set endianness according CMAKE_OSX_ARCHITECTURES
elseif(CMAKE_OSX_ARCHITECTURES MATCHES i386 OR CMAKE_OSX_ARCHITECTURES MATCHES x86_64 OR CMAKE_OSX_ARCHITECTURES MATCHES arm64)
@@ -1759,7 +1760,7 @@ endif()
set(CMAKE_CXX_STANDARD 17)
# If C++17 is not available, downgrading to an earlier standard is NOT OK.
set(CMAKE_CXX_STANDARD_REQUIRED ON)
# Do not enable compiler specific language extentions.
# Do not enable compiler specific language extensions.
set(CMAKE_CXX_EXTENSIONS OFF)
# Make MSVC properly report the value of the __cplusplus preprocessor macro

View File

@@ -51,7 +51,7 @@ Other Convenience Targets
* config: Run cmake configuration tool to set build options.
* deps: Build library dependencies (intended only for platform maintainers).
The existance of locally build dependancies overrides the pre-built dependencies from subversion.
The existance of locally build dependencies overrides the pre-built dependencies from subversion.
These must be manually removed from '../lib/' to go back to using the pre-compiled libraries.
Project Files

View File

@@ -17,7 +17,7 @@
# ***** END GPL LICENSE BLOCK *****
########################################################################
# Copy all generated files to the proper strucure as blender prefers
# Copy all generated files to the proper structure as blender prefers
########################################################################
if(NOT DEFINED HARVEST_TARGET)

View File

@@ -39,7 +39,7 @@ endif()
set(DOWNLOAD_DIR "${CMAKE_CURRENT_BINARY_DIR}/downloads" CACHE STRING "Path for downloaded files")
# This path must be hard-coded like this, so that the GNUmakefile knows where it is and can pass it to make_source_archive.py:
set(PACKAGE_DIR "${CMAKE_CURRENT_BINARY_DIR}/packages")
option(PACKAGE_USE_UPSTREAM_SOURCES "Use soures upstream to download the package sources, when OFF the blender mirror will be used" ON)
option(PACKAGE_USE_UPSTREAM_SOURCES "Use sources upstream to download the package sources, when OFF the blender mirror will be used" ON)
file(TO_CMAKE_PATH ${DOWNLOAD_DIR} DOWNLOAD_DIR)
file(TO_CMAKE_PATH ${PACKAGE_DIR} PACKAGE_DIR)

View File

@@ -24,7 +24,7 @@ if(MSVC)
add_custom_command(
OUTPUT ${PYTARGET}/bin/python${PYTHON_POSTFIX}.exe
COMMAND echo packaging python
COMMAND echo this should ouput at ${PYTARGET}/bin/python${PYTHON_POSTFIX}.exe
COMMAND echo this should output at ${PYTARGET}/bin/python${PYTHON_POSTFIX}.exe
COMMAND ${CMAKE_COMMAND} -E make_directory ${PYTARGET}/libs
COMMAND ${CMAKE_COMMAND} -E copy ${PYSRC}/libs/python${PYTHON_SHORT_VERSION_NO_DOTS}.lib ${PYTARGET}/libs/python${PYTHON_SHORT_VERSION_NO_DOTS}.lib
COMMAND ${CMAKE_COMMAND} -E copy ${PYSRC}/python.exe ${PYTARGET}/bin/python.exe
@@ -43,7 +43,7 @@ if(MSVC)
add_custom_command(
OUTPUT ${PYTARGET}/bin/python${PYTHON_POSTFIX}.exe
COMMAND echo packaging python
COMMAND echo this should ouput at ${PYTARGET}/bin/python${PYTHON_POSTFIX}.exe
COMMAND echo this should output at ${PYTARGET}/bin/python${PYTHON_POSTFIX}.exe
COMMAND ${CMAKE_COMMAND} -E make_directory ${PYTARGET}/libs
COMMAND ${CMAKE_COMMAND} -E copy ${PYSRC}/libs/python${PYTHON_SHORT_VERSION_NO_DOTS}${PYTHON_POSTFIX}.lib ${PYTARGET}/libs/python${PYTHON_SHORT_VERSION_NO_DOTS}${PYTHON_POSTFIX}.lib
COMMAND ${CMAKE_COMMAND} -E copy ${PYSRC}/python${PYTHON_POSTFIX}.exe ${PYTARGET}/bin/python${PYTHON_POSTFIX}.exe

View File

@@ -1826,7 +1826,7 @@ compile_OCIO() {
# Force linking against static libs
#rm -f $_inst/lib/*.so*
# Additional depencencies
# Additional dependencies
#cp ext/dist/lib/libtinyxml.a $_inst/lib
#cp ext/dist/lib/libyaml-cpp.a $_inst/lib

View File

@@ -180,7 +180,7 @@ def create_nb_project_main():
f.write(' </logicalFolder>\n')
f.write(' </logicalFolder>\n')
# default, but this dir is infact not in blender dir so we can ignore it
# default, but this dir is in fact not in blender dir so we can ignore it
# f.write(' <sourceFolderFilter>^(nbproject)$</sourceFolderFilter>\n')
f.write(r' <sourceFolderFilter>^(nbproject|__pycache__|.*\.py|.*\.html|.*\.blend)$</sourceFolderFilter>\n')

View File

@@ -529,7 +529,7 @@ function(SETUP_LIBDIRS)
# NOTE: For all new libraries, use absolute library paths.
# This should eventually be phased out.
# APPLE plaform uses full paths for linking libraries, and avoids link_directories.
# APPLE platform uses full paths for linking libraries, and avoids link_directories.
if(NOT MSVC AND NOT APPLE)
link_directories(${JPEG_LIBPATH} ${PNG_LIBPATH} ${ZLIB_LIBPATH} ${FREETYPE_LIBPATH})

View File

@@ -27,7 +27,7 @@ if(WITH_WINDOWS_BUNDLE_CRT)
# Install the CRT to the blender.crt Sub folder.
install(FILES ${CMAKE_INSTALL_SYSTEM_RUNTIME_LIBS} DESTINATION ./blender.crt COMPONENT Libraries)
# Generating the manifest is a relativly expensive operation since
# Generating the manifest is a relatively expensive operation since
# it is collecting an sha1 hash for every file required. so only do
# this work when the libs have either changed or the manifest does
# not exist yet.

View File

@@ -5,38 +5,38 @@
update-code:
git:
submodules:
- branch: blender-v3.0-release
- branch: master
commit_id: HEAD
path: release/scripts/addons
- branch: blender-v3.0-release
- branch: master
commit_id: HEAD
path: release/scripts/addons_contrib
- branch: blender-v3.0-release
- branch: master
commit_id: HEAD
path: release/datafiles/locale
- branch: blender-v3.0-release
- branch: master
commit_id: HEAD
path: source/tools
svn:
libraries:
darwin-arm64:
branch: tags/blender-3.0-release
branch: trunk
commit_id: HEAD
path: lib/darwin_arm64
darwin-x86_64:
branch: tags/blender-3.0-release
branch: trunk
commit_id: HEAD
path: lib/darwin
linux-x86_64:
branch: tags/blender-3.0-release
branch: trunk
commit_id: HEAD
path: lib/linux_centos7_x86_64
windows-amd64:
branch: tags/blender-3.0-release
branch: trunk
commit_id: HEAD
path: lib/win64_vc15
tests:
branch: tags/blender-3.0-release
branch: trunk
commit_id: HEAD
path: lib/tests
benchmarks:

View File

@@ -38,7 +38,7 @@ PROJECT_NAME = Blender
# could be handy for archiving the generated documentation or if some version
# control system is used.
PROJECT_NUMBER = V3.0
PROJECT_NUMBER = V3.1
# Using the PROJECT_BRIEF tag one can provide an optional one line description
# for a project that appears at the top of each page and should give viewer a

View File

@@ -11,7 +11,7 @@ import queue
execution_queue = queue.Queue()
# This function can savely be called in another thread.
# This function can safely be called in another thread.
# The function will be executed when the timer runs the next time.
def run_in_main_thread(function):
execution_queue.put(function)

View File

@@ -728,7 +728,7 @@ Abusing RNA property callbacks
------------------------------
Python-defined RNA properties can have custom callbacks. Trying to perform complex operations
from there, like calling an operator, may work, but is not officialy recommended nor supported.
from there, like calling an operator, may work, but is not officially recommended nor supported.
Main reason is that those callback should be very fast, but additionally, it may for example
create issues with undo/redo system (most operators store an history step, and editing an RNA

View File

@@ -1224,7 +1224,10 @@ def pycontext2sphinx(basepath):
while char_array[i] is not None:
member = ctypes.string_at(char_array[i]).decode(encoding="ascii")
fw(".. data:: %s\n\n" % member)
member_type, is_seq = context_type_map[member]
try:
member_type, is_seq = context_type_map[member]
except KeyError:
raise SystemExit("Error: context key %r not found in context_type_map; update %s" % (member, __file__)) from None
fw(" :type: %s :class:`bpy.types.%s`\n\n" % ("sequence of " if is_seq else "", member_type))
unique.add(member)
i += 1
@@ -2251,7 +2254,7 @@ def main():
# First monkey patch to load in fake members.
setup_monkey_patch()
# Perform changes to Blender it's self.
# Perform changes to Blender itself.
setup_data = setup_blender()
# eventually, create the dirs

View File

@@ -226,6 +226,9 @@ add_definitions(
-DCCL_NAMESPACE_END=}
)
if(WITH_CYCLES_DEBUG)
add_definitions(-DWITH_CYCLES_DEBUG)
endif()
if(WITH_CYCLES_STANDALONE_GUI)
add_definitions(-DWITH_CYCLES_STANDALONE_GUI)
endif()
@@ -334,7 +337,7 @@ else()
endif()
# Warnings
if(CMAKE_COMPILER_IS_GNUCXX)
if(CMAKE_COMPILER_IS_GNUCXX OR CMAKE_C_COMPILER_ID MATCHES "Clang")
ADD_CHECK_CXX_COMPILER_FLAG(CMAKE_CXX_FLAGS _has_cxxflag_float_conversion "-Werror=float-conversion")
ADD_CHECK_CXX_COMPILER_FLAG(CMAKE_CXX_FLAGS _has_cxxflag_double_promotion "-Werror=double-promotion")
ADD_CHECK_CXX_COMPILER_FLAG(CMAKE_CXX_FLAGS _has_no_error_unused_macros "-Wno-error=unused-macros")

View File

@@ -138,11 +138,6 @@ endif()
blender_add_lib(bf_intern_cycles "${SRC}" "${INC}" "${INC_SYS}" "${LIB}")
# avoid link failure with clang 3.4 debug
if(CMAKE_C_COMPILER_ID MATCHES "Clang" AND NOT ${CMAKE_C_COMPILER_VERSION} VERSION_LESS '3.4')
string(APPEND CMAKE_CXX_FLAGS_DEBUG " -gline-tables-only")
endif()
add_dependencies(bf_intern_cycles bf_rna)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${ADDON_FILES}" ${CYCLES_INSTALL_PATH})

View File

@@ -218,6 +218,12 @@ enum_denoising_prefilter = (
('ACCURATE', "Accurate", "Prefilter noisy guiding passes before denoising color. Improves quality when guiding passes are noisy using extra processing time", 3),
)
enum_direct_light_sampling_type = (
('MULTIPLE_IMPORTANCE_SAMPLING', "Multiple Importance Sampling", "Multiple importance sampling is used to combine direct light contributions from next-event estimation and forward path tracing", 0),
('FORWARD_PATH_TRACING', "Forward Path Tracing", "Direct light contributions are only sampled using forward path tracing", 1),
('NEXT_EVENT_ESTIMATION', "Next-Event Estimation", "Direct light contributions are only sampled using next-event estimation", 2),
)
def update_render_passes(self, context):
scene = context.scene
view_layer = context.view_layer
@@ -325,6 +331,13 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
default=1024,
)
sample_offset: IntProperty(
name="Sample Offset",
description="Number of samples to skip when starting render",
min=0, max=(1 << 24),
default=0,
)
time_limit: FloatProperty(
name="Time Limit",
description="Limit the render time (excluding synchronization time)."
@@ -415,6 +428,13 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
default=0,
)
direct_light_sampling_type: EnumProperty(
name="Direct Light Sampling Type",
description="The type of strategy used for sampling direct light contributions",
items=enum_direct_light_sampling_type,
default='MULTIPLE_IMPORTANCE_SAMPLING',
)
min_light_bounces: IntProperty(
name="Min Light Bounces",
description="Minimum number of light bounces. Setting this higher reduces noise in the first bounces, "

View File

@@ -290,6 +290,9 @@ class CYCLES_RENDER_PT_sampling_advanced(CyclesButtonsPanel, Panel):
col.active = not (cscene.use_adaptive_sampling and cscene.use_preview_adaptive_sampling)
col.prop(cscene, "sampling_pattern", text="Pattern")
col = layout.column(align=True)
col.prop(cscene, "sample_offset")
layout.separator()
heading = layout.column(align=True, heading="Scrambling Distance")

View File

@@ -199,7 +199,7 @@ static bool ObtainCacheParticleUV(Hair *hair,
b_mesh->uv_layers.begin(l);
float2 uv = zero_float2();
if (b_mesh->uv_layers.length())
if (!b_mesh->uv_layers.empty())
b_psys.uv_on_emitter(psmd, *b_pa, pa_no, uv_num, &uv.x);
CData->curve_uv.push_back_slow(uv);
@@ -261,7 +261,7 @@ static bool ObtainCacheParticleVcol(Hair *hair,
b_mesh->vertex_colors.begin(l);
float4 vcol = make_float4(0.0f, 0.0f, 0.0f, 1.0f);
if (b_mesh->vertex_colors.length())
if (!b_mesh->vertex_colors.empty())
b_psys.mcol_on_emitter(psmd, *b_pa, pa_no, vcol_num, &vcol.x);
CData->curve_vcol.push_back_slow(vcol);

View File

@@ -334,7 +334,7 @@ bool BlenderDisplayDriver::update_begin(const Params &params,
/* Update PBO dimensions if needed.
*
* NOTE: Allocate the PBO for the the size which will fit the final render resolution (as in,
* NOTE: Allocate the PBO for the size which will fit the final render resolution (as in,
* at a resolution divider 1. This was we don't need to recreate graphics interoperability
* objects which are costly and which are tied to the specific underlying buffer size.
* The downside of this approach is that when graphics interoperability is not used we are

View File

@@ -555,7 +555,7 @@ static void attr_create_vertex_color(Scene *scene, Mesh *mesh, BL::Mesh &b_mesh,
/* Create uv map attributes. */
static void attr_create_uv_map(Scene *scene, Mesh *mesh, BL::Mesh &b_mesh)
{
if (b_mesh.uv_layers.length() != 0) {
if (!b_mesh.uv_layers.empty()) {
for (BL::MeshUVLoopLayer &l : b_mesh.uv_layers) {
const bool active_render = l.active_render();
AttributeStandard uv_std = (active_render) ? ATTR_STD_UV : ATTR_STD_NONE;
@@ -619,7 +619,7 @@ static void attr_create_uv_map(Scene *scene, Mesh *mesh, BL::Mesh &b_mesh)
static void attr_create_subd_uv_map(Scene *scene, Mesh *mesh, BL::Mesh &b_mesh, bool subdivide_uvs)
{
if (b_mesh.uv_layers.length() != 0) {
if (!b_mesh.uv_layers.empty()) {
BL::Mesh::uv_layers_iterator l;
int i = 0;
@@ -951,7 +951,7 @@ static void create_mesh(Scene *scene,
N = attr_N->data_float3();
/* create generated coordinates from undeformed coordinates */
const bool need_default_tangent = (subdivision == false) && (b_mesh.uv_layers.length() == 0) &&
const bool need_default_tangent = (subdivision == false) && (b_mesh.uv_layers.empty()) &&
(mesh->need_attribute(scene, ATTR_STD_UV_TANGENT));
if (mesh->need_attribute(scene, ATTR_STD_GENERATED) || need_default_tangent) {
Attribute *attr = attributes.add(ATTR_STD_GENERATED);

View File

@@ -392,6 +392,12 @@ void BlenderSync::sync_integrator(BL::ViewLayer &b_view_layer, bool background)
integrator->set_ao_bounces(0);
}
#ifdef WITH_CYCLES_DEBUG
DirectLightSamplingType direct_light_sampling_type = (DirectLightSamplingType)get_enum(
cscene, "direct_light_sampling_type", DIRECT_LIGHT_SAMPLING_NUM, DIRECT_LIGHT_SAMPLING_MIS);
integrator->set_direct_light_sampling_type(direct_light_sampling_type);
#endif
const DenoiseParams denoise_params = get_denoise_params(b_scene, b_view_layer, background);
integrator->set_use_denoise(denoise_params.use);
@@ -835,18 +841,25 @@ SessionParams BlenderSync::get_session_params(BL::RenderEngine &b_engine,
/* samples */
int samples = get_int(cscene, "samples");
int preview_samples = get_int(cscene, "preview_samples");
int sample_offset = get_int(cscene, "sample_offset");
if (background) {
params.samples = samples;
params.sample_offset = sample_offset;
}
else {
params.samples = preview_samples;
if (params.samples == 0)
if (params.samples == 0) {
params.samples = INT_MAX;
}
params.sample_offset = 0;
}
/* Clamp sample offset. */
params.sample_offset = clamp(params.sample_offset, 0, Integrator::MAX_SAMPLES);
/* Clamp samples. */
params.samples = min(params.samples, Integrator::MAX_SAMPLES);
params.samples = clamp(params.samples, 0, Integrator::MAX_SAMPLES - params.sample_offset);
/* Viewport Performance */
params.pixel_size = b_engine.get_preview_pixel_size(b_scene);
@@ -865,7 +878,7 @@ SessionParams BlenderSync::get_session_params(BL::RenderEngine &b_engine,
/* Time limit. */
if (background) {
params.time_limit = get_float(cscene, "time_limit");
params.time_limit = (double)get_float(cscene, "time_limit");
}
else {
/* For the viewport it kind of makes more sense to think in terms of the noise floor, which is

View File

@@ -303,7 +303,7 @@ static inline string image_user_file_path(BL::ImageUser &iuser,
string filepath_str = string(filepath);
if (load_tiled && ima.source() == BL::Image::source_TILED) {
string udim;
if (ima.tiles.length() > 0) {
if (!ima.tiles.empty()) {
udim = to_string(ima.tiles[0].number());
}
string_replace(filepath_str, udim, "<UDIM>");
@@ -647,7 +647,7 @@ static inline Mesh::SubdivisionType object_subdivision_type(BL::Object &b_ob,
{
PointerRNA cobj = RNA_pointer_get(&b_ob.ptr, "cycles");
if (cobj.data && b_ob.modifiers.length() > 0 && experimental) {
if (cobj.data && !b_ob.modifiers.empty() && experimental) {
BL::Modifier mod = b_ob.modifiers[b_ob.modifiers.length() - 1];
bool enabled = preview ? mod.show_viewport() : mod.show_render();

View File

@@ -303,7 +303,7 @@ static void rtc_error_func(void *, enum RTCError, const char *str)
VLOG(1) << str;
}
static double progress_start_time = 0.0f;
static double progress_start_time = 0.0;
static bool rtc_progress_func(void *user_ptr, const double n)
{

View File

@@ -153,7 +153,7 @@ void BVHNode::update_time()
namespace {
struct DumpTraversalContext {
/* Descriptor of wile where writing is happening. */
/* Descriptor of while where writing is happening. */
FILE *stream;
/* Unique identifier of the node current. */
int id;

View File

@@ -178,7 +178,7 @@ class InnerNode : public BVHNode {
reset_unused_children();
}
/* NOTE: This function is only used during binary BVH builder, and it
/* NOTE: This function is only used during binary BVH builder, and it's
* supposed to be configured to have 2 children which will be filled-in in a
* bit. But this is important to have children reset to NULL. */
explicit InnerNode(const BoundBox &bounds) : BVHNode(bounds), num_children_(0)

View File

@@ -88,7 +88,7 @@ endmacro()
function(cycles_link_directories)
if(APPLE)
# APPLE plaform uses full paths for linking libraries, and avoids link_directories.
# APPLE platform uses full paths for linking libraries, and avoids link_directories.
return()
endif()

View File

@@ -38,7 +38,6 @@ void device_cpu_info(vector<DeviceInfo> &devices)
info.id = "CPU";
info.num = 0;
info.has_osl = true;
info.has_half_images = true;
info.has_nanovdb = true;
info.has_profiling = true;
if (openimagedenoise_supported()) {

View File

@@ -93,11 +93,6 @@ CPUDevice::~CPUDevice()
texture_info.free();
}
bool CPUDevice::show_samples() const
{
return (info.cpu_threads == 1);
}
BVHLayoutMask CPUDevice::get_bvh_layout_mask() const
{
BVHLayoutMask bvh_layout_mask = BVH_LAYOUT_BVH2;

View File

@@ -60,8 +60,6 @@ class CPUDevice : public Device {
CPUDevice(const DeviceInfo &info_, Stats &stats_, Profiler &profiler_);
~CPUDevice();
virtual bool show_samples() const override;
virtual BVHLayoutMask get_bvh_layout_mask() const override;
/* Returns true if the texture info was copied to the device (meaning, some more

View File

@@ -144,7 +144,6 @@ void device_cuda_info(vector<DeviceInfo> &devices)
info.description = string(name);
info.num = num;
info.has_half_images = (major >= 3);
info.has_nanovdb = true;
info.denoisers = 0;

View File

@@ -46,12 +46,6 @@ bool CUDADevice::have_precompiled_kernels()
return path_exists(cubins_path);
}
bool CUDADevice::show_samples() const
{
/* The CUDADevice only processes one tile at a time, so showing samples is fine. */
return true;
}
BVHLayoutMask CUDADevice::get_bvh_layout_mask() const
{
return BVH_LAYOUT_BVH2;
@@ -242,6 +236,10 @@ string CUDADevice::compile_kernel_get_common_cflags(const uint kernel_features)
cflags += " -DWITH_NANOVDB";
# endif
# ifdef WITH_CYCLES_DEBUG
cflags += " -DWITH_CYCLES_DEBUG";
# endif
return cflags;
}
@@ -932,7 +930,6 @@ void CUDADevice::tex_alloc(device_texture &mem)
{
CUDAContextScope scope(this);
/* General variables for both architectures */
string bind_name = mem.name;
size_t dsize = datatype_size(mem.data_type);
size_t size = mem.memory_size();
@@ -1095,7 +1092,6 @@ void CUDADevice::tex_alloc(device_texture &mem)
if (mem.info.data_type != IMAGE_DATA_TYPE_NANOVDB_FLOAT &&
mem.info.data_type != IMAGE_DATA_TYPE_NANOVDB_FLOAT3) {
/* Kepler+, bindless textures. */
CUDA_RESOURCE_DESC resDesc;
memset(&resDesc, 0, sizeof(resDesc));

View File

@@ -76,8 +76,6 @@ class CUDADevice : public Device {
static bool have_precompiled_kernels();
virtual bool show_samples() const override;
virtual BVHLayoutMask get_bvh_layout_mask() const override;
void set_error(const string &error) override;

View File

@@ -286,7 +286,6 @@ DeviceInfo Device::get_multi_device(const vector<DeviceInfo> &subdevices,
info.description = "Multi Device";
info.num = 0;
info.has_half_images = true;
info.has_nanovdb = true;
info.has_osl = true;
info.has_profiling = true;
@@ -333,7 +332,6 @@ 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_osl &= device.has_osl;
info.has_profiling &= device.has_profiling;

View File

@@ -73,7 +73,6 @@ class DeviceInfo {
int num;
bool display_device; /* GPU is used as a display device. */
bool has_nanovdb; /* Support NanoVDB volumes. */
bool has_half_images; /* Support half-float textures. */
bool has_osl; /* Support Open Shading Language. */
bool has_profiling; /* Supports runtime collection of profiling info. */
bool has_peer_memory; /* GPU has P2P access to memory of another GPU. */
@@ -90,7 +89,6 @@ class DeviceInfo {
num = 0;
cpu_threads = 0;
display_device = false;
has_half_images = false;
has_nanovdb = false;
has_osl = false;
has_profiling = false;
@@ -151,10 +149,6 @@ class Device {
fprintf(stderr, "%s\n", error.c_str());
fflush(stderr);
}
virtual bool show_samples() const
{
return false;
}
virtual BVHLayoutMask get_bvh_layout_mask() const = 0;
/* statistics */

View File

@@ -148,7 +148,6 @@ void device_hip_info(vector<DeviceInfo> &devices)
info.description = string(name);
info.num = num;
info.has_half_images = true;
info.has_nanovdb = true;
info.denoisers = 0;

View File

@@ -47,12 +47,6 @@ bool HIPDevice::have_precompiled_kernels()
return path_exists(fatbins_path);
}
bool HIPDevice::show_samples() const
{
/* The HIPDevice only processes one tile at a time, so showing samples is fine. */
return true;
}
BVHLayoutMask HIPDevice::get_bvh_layout_mask() const
{
return BVH_LAYOUT_BVH2;
@@ -243,7 +237,7 @@ string HIPDevice::compile_kernel(const uint kernel_features, const char *name, c
hipGetDeviceProperties(&props, hipDevId);
/* gcnArchName can contain tokens after the arch name with features, ie.
"gfx1010:sramecc-:xnack-" so we tokenize it to get the first part. */
* `gfx1010:sramecc-:xnack-` so we tokenize it to get the first part. */
char *arch = strtok(props.gcnArchName, ":");
if (arch == NULL) {
arch = props.gcnArchName;
@@ -374,10 +368,9 @@ string HIPDevice::compile_kernel(const uint kernel_features, const char *name, c
bool HIPDevice::load_kernels(const uint kernel_features)
{
/* TODO(sergey): Support kernels re-load for CUDA devices adaptive compile.
/* TODO(sergey): Support kernels re-load for HIP devices adaptive compile.
*
* Currently re-loading kernel will invalidate memory pointers,
* causing problems in cuCtxSynchronize.
* Currently re-loading kernels will invalidate memory pointers.
*/
if (hipModule) {
if (use_adaptive_compilation()) {
@@ -899,7 +892,6 @@ void HIPDevice::tex_alloc(device_texture &mem)
{
HIPContextScope scope(this);
/* General variables for both architectures */
string bind_name = mem.name;
size_t dsize = datatype_size(mem.data_type);
size_t size = mem.memory_size();

View File

@@ -75,8 +75,6 @@ class HIPDevice : public Device {
static bool have_precompiled_kernels();
virtual bool show_samples() const override;
virtual BVHLayoutMask get_bvh_layout_mask() const override;
void set_error(const string &error) override;
@@ -93,9 +91,7 @@ class HIPDevice : public Device {
virtual string compile_kernel_get_common_cflags(const uint kernel_features);
string compile_kernel(const uint kernel_features,
const char *name,
const char *base = "hip");
string compile_kernel(const uint kernel_features, const char *name, const char *base = "hip");
virtual bool load_kernels(const uint kernel_features) override;
void reserve_local_memory(const uint kernel_features);

View File

@@ -48,7 +48,7 @@ class HIPDeviceGraphicsInterop : public DeviceGraphicsInterop {
HIPDeviceQueue *queue_ = nullptr;
HIPDevice *device_ = nullptr;
/* OpenGL PBO which is currently registered as the destination for the CUDA buffer. */
/* OpenGL PBO which is currently registered as the destination for the HIP buffer. */
uint opengl_pbo_id_ = 0;
/* Buffer area in pixels of the corresponding PBO. */
int64_t buffer_area_ = 0;

View File

@@ -23,7 +23,7 @@ CCL_NAMESPACE_BEGIN
device_memory::device_memory(Device *device, const char *name, MemoryType type)
: data_type(device_type_traits<uchar>::data_type),
data_elements(device_type_traits<uchar>::num_elements_cpu),
data_elements(device_type_traits<uchar>::num_elements),
data_size(0),
device_size(0),
data_width(0),

View File

@@ -81,155 +81,140 @@ static constexpr size_t datatype_size(DataType datatype)
template<typename T> struct device_type_traits {
static const DataType data_type = TYPE_UNKNOWN;
static const size_t num_elements_cpu = sizeof(T);
static const size_t num_elements_gpu = sizeof(T);
static const size_t num_elements = sizeof(T);
};
template<> struct device_type_traits<uchar> {
static const DataType data_type = TYPE_UCHAR;
static const size_t num_elements_cpu = 1;
static const size_t num_elements_gpu = 1;
static_assert(sizeof(uchar) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 1;
static_assert(sizeof(uchar) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<uchar2> {
static const DataType data_type = TYPE_UCHAR;
static const size_t num_elements_cpu = 2;
static const size_t num_elements_gpu = 2;
static_assert(sizeof(uchar2) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 2;
static_assert(sizeof(uchar2) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<uchar3> {
static const DataType data_type = TYPE_UCHAR;
static const size_t num_elements_cpu = 3;
static const size_t num_elements_gpu = 3;
static_assert(sizeof(uchar3) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 3;
static_assert(sizeof(uchar3) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<uchar4> {
static const DataType data_type = TYPE_UCHAR;
static const size_t num_elements_cpu = 4;
static const size_t num_elements_gpu = 4;
static_assert(sizeof(uchar4) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 4;
static_assert(sizeof(uchar4) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<uint> {
static const DataType data_type = TYPE_UINT;
static const size_t num_elements_cpu = 1;
static const size_t num_elements_gpu = 1;
static_assert(sizeof(uint) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 1;
static_assert(sizeof(uint) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<uint2> {
static const DataType data_type = TYPE_UINT;
static const size_t num_elements_cpu = 2;
static const size_t num_elements_gpu = 2;
static_assert(sizeof(uint2) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 2;
static_assert(sizeof(uint2) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<uint3> {
static const DataType data_type = TYPE_UINT;
static const size_t num_elements_cpu = 3;
static const size_t num_elements_gpu = 3;
static_assert(sizeof(uint3) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 3;
static_assert(sizeof(uint3) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<uint4> {
static const DataType data_type = TYPE_UINT;
static const size_t num_elements_cpu = 4;
static const size_t num_elements_gpu = 4;
static_assert(sizeof(uint4) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 4;
static_assert(sizeof(uint4) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<int> {
static const DataType data_type = TYPE_INT;
static const size_t num_elements_cpu = 1;
static const size_t num_elements_gpu = 1;
static_assert(sizeof(int) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 1;
static_assert(sizeof(int) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<int2> {
static const DataType data_type = TYPE_INT;
static const size_t num_elements_cpu = 2;
static const size_t num_elements_gpu = 2;
static_assert(sizeof(int2) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 2;
static_assert(sizeof(int2) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<int3> {
static const DataType data_type = TYPE_INT;
static const size_t num_elements_cpu = 4;
static const size_t num_elements_gpu = 3;
static_assert(sizeof(int3) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 4;
static_assert(sizeof(int3) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<int4> {
static const DataType data_type = TYPE_INT;
static const size_t num_elements_cpu = 4;
static const size_t num_elements_gpu = 4;
static_assert(sizeof(int4) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 4;
static_assert(sizeof(int4) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<float> {
static const DataType data_type = TYPE_FLOAT;
static const size_t num_elements_cpu = 1;
static const size_t num_elements_gpu = 1;
static_assert(sizeof(float) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 1;
static_assert(sizeof(float) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<float2> {
static const DataType data_type = TYPE_FLOAT;
static const size_t num_elements_cpu = 2;
static const size_t num_elements_gpu = 2;
static_assert(sizeof(float2) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 2;
static_assert(sizeof(float2) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<float3> {
/* float3 has different size depending on the device, can't use it for interchanging
* memory between CPU and GPU.
*
* Leave body empty to trigger a compile error if used. */
};
template<> struct device_type_traits<packed_float3> {
static const DataType data_type = TYPE_FLOAT;
static const size_t num_elements_cpu = 4;
static const size_t num_elements_gpu = 3;
static_assert(sizeof(float3) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 3;
static_assert(sizeof(packed_float3) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<float4> {
static const DataType data_type = TYPE_FLOAT;
static const size_t num_elements_cpu = 4;
static const size_t num_elements_gpu = 4;
static_assert(sizeof(float4) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 4;
static_assert(sizeof(float4) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<half> {
static const DataType data_type = TYPE_HALF;
static const size_t num_elements_cpu = 1;
static const size_t num_elements_gpu = 1;
static_assert(sizeof(half) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 1;
static_assert(sizeof(half) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<ushort4> {
static const DataType data_type = TYPE_UINT16;
static const size_t num_elements_cpu = 4;
static const size_t num_elements_gpu = 4;
static_assert(sizeof(ushort4) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 4;
static_assert(sizeof(ushort4) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<uint16_t> {
static const DataType data_type = TYPE_UINT16;
static const size_t num_elements_cpu = 1;
static const size_t num_elements_gpu = 1;
static_assert(sizeof(uint16_t) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 1;
static_assert(sizeof(uint16_t) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<half4> {
static const DataType data_type = TYPE_HALF;
static const size_t num_elements_cpu = 4;
static const size_t num_elements_gpu = 4;
static_assert(sizeof(half4) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 4;
static_assert(sizeof(half4) == num_elements * datatype_size(data_type));
};
template<> struct device_type_traits<uint64_t> {
static const DataType data_type = TYPE_UINT64;
static const size_t num_elements_cpu = 1;
static const size_t num_elements_gpu = 1;
static_assert(sizeof(uint64_t) == num_elements_cpu * datatype_size(data_type));
static const size_t num_elements = 1;
static_assert(sizeof(uint64_t) == num_elements * datatype_size(data_type));
};
/* Device Memory
@@ -325,9 +310,7 @@ template<typename T> class device_only_memory : public device_memory {
: device_memory(device, name, allow_host_memory_fallback ? MEM_READ_WRITE : MEM_DEVICE_ONLY)
{
data_type = device_type_traits<T>::data_type;
data_elements = max(device_is_cpu() ? device_type_traits<T>::num_elements_cpu :
device_type_traits<T>::num_elements_gpu,
1);
data_elements = max(device_type_traits<T>::num_elements, 1);
}
device_only_memory(device_only_memory &&other) noexcept : device_memory(std::move(other))
@@ -383,15 +366,11 @@ template<typename T> class device_only_memory : public device_memory {
template<typename T> class device_vector : public device_memory {
public:
/* Can only use this for types that have the same size on CPU and GPU. */
static_assert(device_type_traits<T>::num_elements_cpu ==
device_type_traits<T>::num_elements_gpu);
device_vector(Device *device, const char *name, MemoryType type)
: device_memory(device, name, type)
{
data_type = device_type_traits<T>::data_type;
data_elements = device_type_traits<T>::num_elements_cpu;
data_elements = device_type_traits<T>::num_elements;
modified = true;
need_realloc_ = true;

View File

@@ -109,14 +109,6 @@ class MultiDevice : public Device {
return error_msg;
}
virtual bool show_samples() const override
{
if (devices.size() > 1) {
return false;
}
return devices.front().device->show_samples();
}
virtual BVHLayoutMask get_bvh_layout_mask() const override
{
BVHLayoutMask bvh_layout_mask = BVH_LAYOUT_ALL;

View File

@@ -3,7 +3,7 @@ This program uses code from various sources, the default license is Apache 2.0
for all code, with the following exceptions.
Modified BSD License
* Code adapated from Open Shading Language
* Code adapted from Open Shading Language
* Sobol direction vectors
* Matrix inversion code from OpenEXR
* MD5 Hash code

View File

@@ -380,7 +380,10 @@ void PathTrace::path_trace(RenderWork &render_work)
PathTraceWork *path_trace_work = path_trace_works_[i].get();
PathTraceWork::RenderStatistics statistics;
path_trace_work->render_samples(statistics, render_work.path_trace.start_sample, num_samples);
path_trace_work->render_samples(statistics,
render_work.path_trace.start_sample,
num_samples,
render_work.path_trace.sample_offset);
const double work_time = time_dt() - work_start_time;
work_balance_infos_[i].time_spent += work_time;
@@ -850,7 +853,8 @@ void PathTrace::progress_update_if_needed(const RenderWork &render_work)
const uint64_t num_samples_added = uint64_t(tile_size.x) * tile_size.y *
render_work.path_trace.num_samples;
const int current_sample = render_work.path_trace.start_sample +
render_work.path_trace.num_samples;
render_work.path_trace.num_samples -
render_work.path_trace.sample_offset;
progress_->add_samples(num_samples_added, current_sample);
}

View File

@@ -76,7 +76,7 @@ class PathTraceDisplay {
/* Copy buffer of rendered pixels of a given size into a given position of the texture.
*
* This function does not acquire a lock. The reason for this is is to allow use of this function
* This function does not acquire a lock. The reason for this is to allow use of this function
* for partial updates from different devices. In this case the caller will acquire the lock
* once, update all the slices and release
* the lock once. This will ensure that draw() will never use partially updated texture. */

View File

@@ -75,7 +75,10 @@ class PathTraceWork {
/* Render given number of samples as a synchronous blocking call.
* The samples are added to the render buffer associated with this work. */
virtual void render_samples(RenderStatistics &statistics, int start_sample, int samples_num) = 0;
virtual void render_samples(RenderStatistics &statistics,
int start_sample,
int samples_num,
int sample_offset) = 0;
/* Copy render result from this work to the corresponding place of the GPU display.
*

View File

@@ -71,7 +71,8 @@ void PathTraceWorkCPU::init_execution()
void PathTraceWorkCPU::render_samples(RenderStatistics &statistics,
int start_sample,
int samples_num)
int samples_num,
int sample_offset)
{
const int64_t image_width = effective_buffer_params_.width;
const int64_t image_height = effective_buffer_params_.height;
@@ -99,6 +100,7 @@ void PathTraceWorkCPU::render_samples(RenderStatistics &statistics,
work_tile.w = 1;
work_tile.h = 1;
work_tile.start_sample = start_sample;
work_tile.sample_offset = sample_offset;
work_tile.num_samples = 1;
work_tile.offset = effective_buffer_params_.offset;
work_tile.stride = effective_buffer_params_.stride;

View File

@@ -48,7 +48,8 @@ class PathTraceWorkCPU : public PathTraceWork {
virtual void render_samples(RenderStatistics &statistics,
int start_sample,
int samples_num) override;
int samples_num,
int sample_offset) override;
virtual void copy_to_display(PathTraceDisplay *display,
PassMode pass_mode,

View File

@@ -250,7 +250,8 @@ void PathTraceWorkGPU::init_execution()
void PathTraceWorkGPU::render_samples(RenderStatistics &statistics,
int start_sample,
int samples_num)
int samples_num,
int sample_offset)
{
/* Limit number of states for the tile and rely on a greedy scheduling of tiles. This allows to
* add more work (because tiles are smaller, so there is higher chance that more paths will
@@ -261,6 +262,7 @@ void PathTraceWorkGPU::render_samples(RenderStatistics &statistics,
work_tile_scheduler_.reset(effective_buffer_params_,
start_sample,
samples_num,
sample_offset,
device_scene_->data.integrator.scrambling_distance);
enqueue_reset();

View File

@@ -46,7 +46,8 @@ class PathTraceWorkGPU : public PathTraceWork {
virtual void render_samples(RenderStatistics &statistics,
int start_sample,
int samples_num) override;
int samples_num,
int sample_offset) override;
virtual void copy_to_display(PathTraceDisplay *display,
PassMode pass_mode,

View File

@@ -88,6 +88,16 @@ int RenderScheduler::get_num_samples() const
return num_samples_;
}
void RenderScheduler::set_sample_offset(int sample_offset)
{
sample_offset_ = sample_offset;
}
int RenderScheduler::get_sample_offset() const
{
return sample_offset_;
}
void RenderScheduler::set_time_limit(double time_limit)
{
time_limit_ = time_limit;
@@ -110,13 +120,15 @@ int RenderScheduler::get_num_rendered_samples() const
return state_.num_rendered_samples;
}
void RenderScheduler::reset(const BufferParams &buffer_params, int num_samples)
void RenderScheduler::reset(const BufferParams &buffer_params, int num_samples, int sample_offset)
{
buffer_params_ = buffer_params;
update_start_resolution_divider();
set_num_samples(num_samples);
set_start_sample(sample_offset);
set_sample_offset(sample_offset);
/* In background mode never do lower resolution render preview, as it is not really supported
* by the software. */
@@ -171,7 +183,7 @@ void RenderScheduler::reset(const BufferParams &buffer_params, int num_samples)
void RenderScheduler::reset_for_next_tile()
{
reset(buffer_params_, num_samples_);
reset(buffer_params_, num_samples_, sample_offset_);
}
bool RenderScheduler::render_work_reschedule_on_converge(RenderWork &render_work)
@@ -317,6 +329,7 @@ RenderWork RenderScheduler::get_render_work()
render_work.path_trace.start_sample = get_start_sample_to_path_trace();
render_work.path_trace.num_samples = get_num_samples_to_path_trace();
render_work.path_trace.sample_offset = get_sample_offset();
render_work.init_render_buffers = (render_work.path_trace.start_sample == get_start_sample());
@@ -835,7 +848,7 @@ int RenderScheduler::get_num_samples_to_path_trace() const
* When time limit is not used the number of samples per render iteration is either increasing
* or stays the same, so there is no need to clamp number of samples calculated for occupancy.
*/
if (time_limit_ && state_.start_render_time) {
if (time_limit_ != 0.0 && state_.start_render_time != 0.0) {
const double remaining_render_time = max(
0.0, time_limit_ - (time_dt() - state_.start_render_time));
const double time_per_sample_average = path_trace_time_.get_average();

View File

@@ -39,6 +39,7 @@ class RenderWork {
struct {
int start_sample = 0;
int num_samples = 0;
int sample_offset = 0;
} path_trace;
struct {
@@ -125,6 +126,9 @@ class RenderScheduler {
void set_num_samples(int num_samples);
int get_num_samples() const;
void set_sample_offset(int sample_offset);
int get_sample_offset() const;
/* Time limit for the path tracing tasks, in minutes.
* Zero disables the limit. */
void set_time_limit(double time_limit);
@@ -150,7 +154,7 @@ class RenderScheduler {
/* Reset scheduler, indicating that rendering will happen from scratch.
* Resets current rendered state, as well as scheduling information. */
void reset(const BufferParams &buffer_params, int num_samples);
void reset(const BufferParams &buffer_params, int num_samples, int sample_offset);
/* Reset scheduler upon switching to a next tile.
* Will keep the same number of samples and full-frame render parameters, but will reset progress
@@ -419,6 +423,8 @@ class RenderScheduler {
int start_sample_ = 0;
int num_samples_ = 0;
int sample_offset_ = 0;
/* Limit in seconds for how long path tracing is allowed to happen.
* Zero means no limit is applied. */
double time_limit_ = 0.0;

View File

@@ -36,6 +36,7 @@ void WorkTileScheduler::set_max_num_path_states(int max_num_path_states)
void WorkTileScheduler::reset(const BufferParams &buffer_params,
int sample_start,
int samples_num,
int sample_offset,
float scrambling_distance)
{
/* Image buffer parameters. */
@@ -51,6 +52,7 @@ void WorkTileScheduler::reset(const BufferParams &buffer_params,
/* Samples parameters. */
sample_start_ = sample_start;
samples_num_ = samples_num;
sample_offset_ = sample_offset;
/* Initialize new scheduling. */
reset_scheduler_state();
@@ -111,6 +113,7 @@ bool WorkTileScheduler::get_work(KernelWorkTile *work_tile_, const int max_work_
work_tile.h = tile_size_.height;
work_tile.start_sample = sample_start_ + start_sample;
work_tile.num_samples = min(tile_size_.num_samples, samples_num_ - start_sample);
work_tile.sample_offset = sample_offset_;
work_tile.offset = offset_;
work_tile.stride = stride_;

View File

@@ -41,6 +41,7 @@ class WorkTileScheduler {
void reset(const BufferParams &buffer_params,
int sample_start,
int samples_num,
int sample_offset,
float scrambling_distance);
/* Get work for a device.
@@ -79,6 +80,7 @@ class WorkTileScheduler {
* (splitting into a smaller work tiles). */
int sample_start_ = 0;
int samples_num_ = 0;
int sample_offset_ = 0;
/* Tile size which be scheduled for rendering. */
TileSize tile_size_;

View File

@@ -39,6 +39,10 @@ set(SRC_KERNEL_DEVICE_HIP
device/hip/kernel.cpp
)
set(SRC_KERNEL_DEVICE_METAL
device/metal/kernel.metal
)
set(SRC_KERNEL_DEVICE_OPTIX
device/optix/kernel.cu
device/optix/kernel_shader_raytrace.cu
@@ -79,6 +83,13 @@ set(SRC_KERNEL_DEVICE_OPTIX_HEADERS
device/optix/globals.h
)
set(SRC_KERNEL_DEVICE_METAL_HEADERS
device/metal/compat.h
device/metal/context_begin.h
device/metal/context_end.h
device/metal/globals.h
)
set(SRC_KERNEL_CLOSURE_HEADERS
closure/alloc.h
closure/bsdf.h
@@ -262,6 +273,7 @@ set(SRC_KERNEL_UTIL_HEADERS
)
set(SRC_KERNEL_TYPES_HEADERS
tables.h
textures.h
types.h
)
@@ -399,12 +411,8 @@ if(WITH_CYCLES_CUDA_BINARIES)
-I ${CMAKE_CURRENT_SOURCE_DIR}/..
-I ${CMAKE_CURRENT_SOURCE_DIR}/device/cuda
--use_fast_math
-o ${CMAKE_CURRENT_BINARY_DIR}/${cuda_file})
if(${experimental})
set(cuda_flags ${cuda_flags} -D __KERNEL_EXPERIMENTAL__)
set(name ${name}_experimental)
endif()
-o ${CMAKE_CURRENT_BINARY_DIR}/${cuda_file}
-Wno-deprecated-gpu-targets)
if(WITH_NANOVDB)
set(cuda_flags ${cuda_flags}
@@ -412,6 +420,10 @@ if(WITH_CYCLES_CUDA_BINARIES)
-I "${NANOVDB_INCLUDE_DIR}")
endif()
if(WITH_CYCLES_DEBUG)
set(cuda_flags ${cuda_flags} -D WITH_CYCLES_DEBUG)
endif()
if(WITH_CYCLES_CUBIN_COMPILER)
string(SUBSTRING ${arch} 3 -1 CUDA_ARCH)
@@ -560,11 +572,6 @@ if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
-ffast-math
-o ${CMAKE_CURRENT_BINARY_DIR}/${hip_file})
if(${experimental})
set(hip_flags ${hip_flags} -D __KERNEL_EXPERIMENTAL__)
set(name ${name}_experimental)
endif()
if(WITH_NANOVDB)
set(hip_flags ${hip_flags}
-D WITH_NANOVDB
@@ -572,7 +579,7 @@ if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
endif()
if(WITH_CYCLES_DEBUG)
set(hip_flags ${hip_flags} -D __KERNEL_DEBUG__)
set(hip_flags ${hip_flags} -D WITH_CYCLES_DEBUG)
endif()
add_custom_command(
@@ -613,6 +620,10 @@ if(WITH_CYCLES_DEVICE_OPTIX AND WITH_CYCLES_CUDA_BINARIES)
-I "${NANOVDB_INCLUDE_DIR}")
endif()
if(WITH_CYCLES_DEBUG)
set(cuda_flags ${cuda_flags} -D WITH_CYCLES_DEBUG)
endif()
if(WITH_CYCLES_CUBIN_COMPILER)
# Needed to find libnvrtc-builtins.so. Can't do it from inside
# cycles_cubin_cc since the env variable is read before main()
@@ -701,7 +712,7 @@ if(WITH_COMPILER_ASAN)
string(APPEND CMAKE_CXX_FLAGS_RELWITHDEBINFO " -fno-sanitize=all")
string(APPEND CMAKE_CXX_FLAGS_DEBUG " -fno-sanitize=vptr")
elseif(CMAKE_C_COMPILER_ID MATCHES "Clang")
# With OSL, Cycles disables rtti in some modules, wich then breaks at linking
# With OSL, Cycles disables rtti in some modules, which then breaks at linking
# when trying to use vptr sanitizer (included into 'undefined' general option).
string(APPEND CMAKE_CXX_FLAGS_RELWITHDEBINFO " -fno-sanitize=vptr")
string(APPEND CMAKE_CXX_FLAGS_DEBUG " -fno-sanitize=vptr")
@@ -729,12 +740,14 @@ cycles_add_library(cycles_kernel "${LIB}"
${SRC_KERNEL_DEVICE_CUDA}
${SRC_KERNEL_DEVICE_HIP}
${SRC_KERNEL_DEVICE_OPTIX}
${SRC_KERNEL_DEVICE_METAL}
${SRC_KERNEL_HEADERS}
${SRC_KERNEL_DEVICE_CPU_HEADERS}
${SRC_KERNEL_DEVICE_GPU_HEADERS}
${SRC_KERNEL_DEVICE_CUDA_HEADERS}
${SRC_KERNEL_DEVICE_HIP_HEADERS}
${SRC_KERNEL_DEVICE_OPTIX_HEADERS}
${SRC_KERNEL_DEVICE_METAL_HEADERS}
)
source_group("bake" FILES ${SRC_KERNEL_BAKE_HEADERS})
@@ -746,6 +759,7 @@ source_group("device\\cuda" FILES ${SRC_KERNEL_DEVICE_CUDA} ${SRC_KERNEL_DEVICE_
source_group("device\\gpu" FILES ${SRC_KERNEL_DEVICE_GPU_HEADERS})
source_group("device\\hip" FILES ${SRC_KERNEL_DEVICE_HIP} ${SRC_KERNEL_DEVICE_HIP_HEADERS})
source_group("device\\optix" FILES ${SRC_KERNEL_DEVICE_OPTIX} ${SRC_KERNEL_DEVICE_OPTIX_HEADERS})
source_group("device\\metal" FILES ${SRC_KERNEL_DEVICE_METAL} ${SRC_KERNEL_DEVICE_METAL_HEADERS})
source_group("film" FILES ${SRC_KERNEL_FILM_HEADERS})
source_group("geom" FILES ${SRC_KERNEL_GEOM_HEADERS})
source_group("integrator" FILES ${SRC_KERNEL_INTEGRATOR_HEADERS})
@@ -778,6 +792,8 @@ delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_HIP}" ${CYCLES_
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_HIP_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/hip)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_OPTIX}" ${CYCLES_INSTALL_PATH}/source/kernel/device/optix)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_OPTIX_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/optix)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_METAL}" ${CYCLES_INSTALL_PATH}/source/kernel/device/metal)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_DEVICE_METAL_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/metal)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_FILM_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/film)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_GEOM_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/geom)
delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_KERNEL_INTEGRATOR_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/integrator)

View File

@@ -97,7 +97,7 @@ ccl_device_inline void sort_intersections_and_normals(ccl_private Intersection *
swapped = false;
for (int j = 0; j < num_hits - 1; ++j) {
if (hits[j].t > hits[j + 1].t) {
struct Intersection tmp_hit = hits[j];
Intersection tmp_hit = hits[j];
float3 tmp_Ng = Ng[j];
hits[j] = hits[j + 1];
Ng[j] = Ng[j + 1];

View File

@@ -18,6 +18,7 @@
#pragma once
#include "kernel/tables.h"
#include "kernel/types.h"
#include "kernel/util/profiling.h"

View File

@@ -52,8 +52,9 @@ typedef unsigned long long uint64_t;
#endif
#define ccl_device_noinline __device__ __noinline__
#define ccl_device_noinline_cpu ccl_device
#define ccl_device_inline_method ccl_device
#define ccl_global
#define ccl_static_constant __constant__
#define ccl_inline_constant __constant__
#define ccl_device_constant __constant__ __device__
#define ccl_constant const
#define ccl_gpu_shared __shared__
@@ -75,6 +76,7 @@ typedef unsigned long long uint64_t;
#define ccl_gpu_block_idx_x (blockIdx.x)
#define ccl_gpu_grid_dim_x (gridDim.x)
#define ccl_gpu_warp_size (warpSize)
#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))
#define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x)
#define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x)
@@ -84,7 +86,6 @@ typedef unsigned long long uint64_t;
#define ccl_gpu_syncthreads() __syncthreads()
#define ccl_gpu_ballot(predicate) __ballot_sync(0xFFFFFFFF, predicate)
#define ccl_gpu_shfl_down_sync(mask, var, detla) __shfl_down_sync(mask, var, detla)
#define ccl_gpu_popc(x) __popc(x)
/* GPU texture objects */

View File

@@ -92,12 +92,29 @@
/* Compute number of threads per block and minimum blocks per multiprocessor
* given the maximum number of registers per thread. */
#define ccl_gpu_kernel(block_num_threads, thread_num_registers) \
extern "C" __global__ void __launch_bounds__(block_num_threads, \
GPU_MULTIPRESSOR_MAX_REGISTERS / \
(block_num_threads * thread_num_registers))
#define ccl_gpu_kernel_threads(block_num_threads) \
extern "C" __global__ void __launch_bounds__(block_num_threads)
#define ccl_gpu_kernel_signature(name, ...) kernel_gpu_##name(__VA_ARGS__)
#define ccl_gpu_kernel_call(x) x
/* Define a function object where "func" is the lambda body, and additional parameters are used to
* specify captured state */
#define ccl_gpu_kernel_lambda(func, ...) \
struct KernelLambda { \
__VA_ARGS__; \
__device__ int operator()(const int state) \
{ \
return (func); \
} \
} ccl_gpu_kernel_lambda_pass
/* sanity checks */
#if GPU_KERNEL_BLOCK_NUM_THREADS > GPU_BLOCK_MAX_THREADS

View File

@@ -65,7 +65,9 @@ ccl_device float cubic_h1(float a)
/* Fast bicubic texture lookup using 4 bilinear lookups, adapted from CUDA samples. */
template<typename T>
ccl_device_noinline T kernel_tex_image_interp_bicubic(const TextureInfo &info, float x, float y)
ccl_device_noinline T kernel_tex_image_interp_bicubic(ccl_global const TextureInfo &info,
float x,
float y)
{
ccl_gpu_tex_object tex = (ccl_gpu_tex_object)info.data;
@@ -94,7 +96,7 @@ ccl_device_noinline T kernel_tex_image_interp_bicubic(const TextureInfo &info, f
/* Fast tricubic texture lookup using 8 trilinear lookups. */
template<typename T>
ccl_device_noinline T
kernel_tex_image_interp_tricubic(const TextureInfo &info, float x, float y, float z)
kernel_tex_image_interp_tricubic(ccl_global const TextureInfo &info, float x, float y, float z)
{
ccl_gpu_tex_object tex = (ccl_gpu_tex_object)info.data;
@@ -169,7 +171,7 @@ ccl_device T kernel_tex_image_interp_tricubic_nanovdb(S &s, float x, float y, fl
template<typename T>
ccl_device_noinline T kernel_tex_image_interp_nanovdb(
const TextureInfo &info, float x, float y, float z, uint interpolation)
ccl_global const TextureInfo &info, float x, float y, float z, uint interpolation)
{
using namespace nanovdb;
@@ -191,7 +193,7 @@ ccl_device_noinline T kernel_tex_image_interp_nanovdb(
ccl_device float4 kernel_tex_image_interp(KernelGlobals kg, int id, float x, float y)
{
const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
ccl_global const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
/* float4, byte4, ushort4 and half4 */
const int texture_type = info.data_type;
@@ -226,7 +228,7 @@ ccl_device float4 kernel_tex_image_interp_3d(KernelGlobals kg,
float3 P,
InterpolationType interp)
{
const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
ccl_global const TextureInfo &info = kernel_tex_fetch(__texture_info, id);
if (info.use_transform_3d) {
P = transform_point(&info.transform_3d, P);

File diff suppressed because it is too large Load Diff

View File

@@ -31,10 +31,43 @@ CCL_NAMESPACE_BEGIN
# define GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE 512
#endif
#ifdef __KERNEL_METAL__
struct ActiveIndexContext {
ActiveIndexContext(int _thread_index,
int _global_index,
int _threadgroup_size,
int _simdgroup_size,
int _simd_lane_index,
int _simd_group_index,
int _num_simd_groups,
threadgroup int *_simdgroup_offset)
: thread_index(_thread_index),
global_index(_global_index),
blocksize(_threadgroup_size),
ccl_gpu_warp_size(_simdgroup_size),
thread_warp(_simd_lane_index),
warp_index(_simd_group_index),
num_warps(_num_simd_groups),
warp_offset(_simdgroup_offset)
{
}
const int thread_index, global_index, blocksize, ccl_gpu_warp_size, thread_warp, warp_index,
num_warps;
threadgroup int *warp_offset;
template<uint blocksizeDummy, typename IsActiveOp>
void active_index_array(const uint num_states,
ccl_global int *indices,
ccl_global int *num_indices,
IsActiveOp is_active_op)
{
const uint state_index = global_index;
#else
template<uint blocksize, typename IsActiveOp>
__device__ void gpu_parallel_active_index_array(const uint num_states,
int *indices,
int *num_indices,
ccl_global int *indices,
ccl_global int *num_indices,
IsActiveOp is_active_op)
{
extern ccl_gpu_shared int warp_offset[];
@@ -45,43 +78,62 @@ __device__ void gpu_parallel_active_index_array(const uint num_states,
const uint warp_index = thread_index / ccl_gpu_warp_size;
const uint num_warps = blocksize / ccl_gpu_warp_size;
/* Test if state corresponding to this thread is active. */
const uint state_index = ccl_gpu_block_idx_x * blocksize + thread_index;
const uint is_active = (state_index < num_states) ? is_active_op(state_index) : 0;
#endif
/* For each thread within a warp compute how many other active states precede it. */
const uint thread_mask = 0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp);
const uint thread_offset = ccl_gpu_popc(ccl_gpu_ballot(is_active) & thread_mask);
/* Test if state corresponding to this thread is active. */
const uint is_active = (state_index < num_states) ? is_active_op(state_index) : 0;
/* Last thread in warp stores number of active states for each warp. */
if (thread_warp == ccl_gpu_warp_size - 1) {
warp_offset[warp_index] = thread_offset + is_active;
}
/* For each thread within a warp compute how many other active states precede it. */
const uint thread_offset = popcount(ccl_gpu_ballot(is_active) &
ccl_gpu_thread_mask(thread_warp));
ccl_gpu_syncthreads();
/* Last thread in block converts per-warp sizes to offsets, increments global size of
* index array and gets offset to write to. */
if (thread_index == blocksize - 1) {
/* TODO: parallelize this. */
int offset = 0;
for (int i = 0; i < num_warps; i++) {
int num_active = warp_offset[i];
warp_offset[i] = offset;
offset += num_active;
/* Last thread in warp stores number of active states for each warp. */
if (thread_warp == ccl_gpu_warp_size - 1) {
warp_offset[warp_index] = thread_offset + is_active;
}
const uint block_num_active = warp_offset[warp_index] + thread_offset + is_active;
warp_offset[num_warps] = atomic_fetch_and_add_uint32(num_indices, block_num_active);
ccl_gpu_syncthreads();
/* Last thread in block converts per-warp sizes to offsets, increments global size of
* index array and gets offset to write to. */
if (thread_index == blocksize - 1) {
/* TODO: parallelize this. */
int offset = 0;
for (int i = 0; i < num_warps; i++) {
int num_active = warp_offset[i];
warp_offset[i] = offset;
offset += num_active;
}
const uint block_num_active = warp_offset[warp_index] + thread_offset + is_active;
warp_offset[num_warps] = atomic_fetch_and_add_uint32(num_indices, block_num_active);
}
ccl_gpu_syncthreads();
/* Write to index array. */
if (is_active) {
const uint block_offset = warp_offset[num_warps];
indices[block_offset + warp_offset[warp_index] + thread_offset] = state_index;
}
}
ccl_gpu_syncthreads();
#ifdef __KERNEL_METAL__
}; /* end class ActiveIndexContext */
/* Write to index array. */
if (is_active) {
const uint block_offset = warp_offset[num_warps];
indices[block_offset + warp_offset[warp_index] + thread_offset] = state_index;
}
}
/* inject the required thread params into a struct, and redirect to its templated member function
*/
# define gpu_parallel_active_index_array \
ActiveIndexContext(metal_local_id, \
metal_global_id, \
metal_local_size, \
simdgroup_size, \
simd_lane_index, \
simd_group_index, \
num_simd_groups, \
simdgroup_offset) \
.active_index_array
#endif
CCL_NAMESPACE_END

View File

@@ -33,10 +33,12 @@ CCL_NAMESPACE_BEGIN
# define GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE 512
#endif
template<uint blocksize>
__device__ void gpu_parallel_prefix_sum(int *counter, int *prefix_sum, const int num_values)
__device__ void gpu_parallel_prefix_sum(const int global_id,
ccl_global int *counter,
ccl_global int *prefix_sum,
const int num_values)
{
if (!(ccl_gpu_block_idx_x == 0 && ccl_gpu_thread_idx_x == 0)) {
if (global_id != 0) {
return;
}

View File

@@ -33,16 +33,16 @@ CCL_NAMESPACE_BEGIN
#endif
#define GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY (~0)
template<uint blocksize, typename GetKeyOp>
__device__ void gpu_parallel_sorted_index_array(const uint num_states,
template<typename GetKeyOp>
__device__ void gpu_parallel_sorted_index_array(const uint state_index,
const uint num_states,
const int num_states_limit,
int *indices,
int *num_indices,
int *key_counter,
int *key_prefix_sum,
ccl_global int *indices,
ccl_global int *num_indices,
ccl_global int *key_counter,
ccl_global int *key_prefix_sum,
GetKeyOp get_key_op)
{
const uint state_index = ccl_gpu_block_idx_x * blocksize + ccl_gpu_thread_idx_x;
const int key = (state_index < num_states) ? get_key_op(state_index) :
GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY;

View File

@@ -45,8 +45,9 @@ typedef unsigned long long uint64_t;
#define ccl_device_forceinline __device__ __forceinline__
#define ccl_device_noinline __device__ __noinline__
#define ccl_device_noinline_cpu ccl_device
#define ccl_device_inline_method ccl_device
#define ccl_global
#define ccl_static_constant __constant__
#define ccl_inline_constant __constant__
#define ccl_device_constant __constant__ __device__
#define ccl_constant const
#define ccl_gpu_shared __shared__
@@ -74,6 +75,7 @@ typedef unsigned long long uint64_t;
#define ccl_gpu_block_idx_x (blockIdx.x)
#define ccl_gpu_grid_dim_x (gridDim.x)
#define ccl_gpu_warp_size (warpSize)
#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))
#define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x)
#define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x)
@@ -83,7 +85,6 @@ typedef unsigned long long uint64_t;
#define ccl_gpu_syncthreads() __syncthreads()
#define ccl_gpu_ballot(predicate) __ballot(predicate)
#define ccl_gpu_shfl_down_sync(mask, var, detla) __shfl_down(var, detla)
#define ccl_gpu_popc(x) __popc(x)
/* GPU texture objects */
typedef hipTextureObject_t ccl_gpu_tex_object;

View File

@@ -35,12 +35,29 @@
/* Compute number of threads per block and minimum blocks per multiprocessor
* given the maximum number of registers per thread. */
#define ccl_gpu_kernel(block_num_threads, thread_num_registers) \
extern "C" __global__ void __launch_bounds__(block_num_threads, \
GPU_MULTIPRESSOR_MAX_REGISTERS / \
(block_num_threads * thread_num_registers))
#define ccl_gpu_kernel_threads(block_num_threads) \
extern "C" __global__ void __launch_bounds__(block_num_threads)
#define ccl_gpu_kernel_signature(name, ...) kernel_gpu_##name(__VA_ARGS__)
#define ccl_gpu_kernel_call(x) x
/* Define a function object where "func" is the lambda body, and additional parameters are used to
* specify captured state */
#define ccl_gpu_kernel_lambda(func, ...) \
struct KernelLambda { \
__VA_ARGS__; \
__device__ int operator()(const int state) \
{ \
return (func); \
} \
} ccl_gpu_kernel_lambda_pass
/* sanity checks */
#if GPU_KERNEL_BLOCK_NUM_THREADS > GPU_BLOCK_MAX_THREADS

View File

@@ -34,6 +34,7 @@ using namespace metal;
#pragma clang diagnostic ignored "-Wunused-variable"
#pragma clang diagnostic ignored "-Wsign-compare"
#pragma clang diagnostic ignored "-Wuninitialized"
/* Qualifiers */
@@ -42,8 +43,9 @@ using namespace metal;
#define ccl_device_forceinline ccl_device
#define ccl_device_noinline ccl_device __attribute__((noinline))
#define ccl_device_noinline_cpu ccl_device
#define ccl_device_inline_method ccl_device
#define ccl_global device
#define ccl_static_constant static constant constexpr
#define ccl_inline_constant static constant constexpr
#define ccl_device_constant constant
#define ccl_constant const device
#define ccl_gpu_shared threadgroup
@@ -58,6 +60,122 @@ using namespace metal;
#define kernel_assert(cond)
#define ccl_gpu_global_id_x() metal_global_id
#define ccl_gpu_warp_size simdgroup_size
#define ccl_gpu_thread_idx_x simd_group_index
#define ccl_gpu_thread_mask(thread_warp) uint64_t((1ull << thread_warp) - 1)
#define ccl_gpu_ballot(predicate) ((uint64_t)((simd_vote::vote_t)simd_ballot(predicate)))
#define ccl_gpu_syncthreads() threadgroup_barrier(mem_flags::mem_threadgroup);
// clang-format off
/* kernel.h adapters */
#define ccl_gpu_kernel(block_num_threads, thread_num_registers)
#define ccl_gpu_kernel_threads(block_num_threads)
/* Convert a comma-separated list into a semicolon-separated list
* (so that we can generate a struct based on kernel entry-point parameters). */
#define FN0()
#define FN1(p1) p1;
#define FN2(p1, p2) p1; p2;
#define FN3(p1, p2, p3) p1; p2; p3;
#define FN4(p1, p2, p3, p4) p1; p2; p3; p4;
#define FN5(p1, p2, p3, p4, p5) p1; p2; p3; p4; p5;
#define FN6(p1, p2, p3, p4, p5, p6) p1; p2; p3; p4; p5; p6;
#define FN7(p1, p2, p3, p4, p5, p6, p7) p1; p2; p3; p4; p5; p6; p7;
#define FN8(p1, p2, p3, p4, p5, p6, p7, p8) p1; p2; p3; p4; p5; p6; p7; p8;
#define FN9(p1, p2, p3, p4, p5, p6, p7, p8, p9) p1; p2; p3; p4; p5; p6; p7; p8; p9;
#define FN10(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10;
#define FN11(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11;
#define FN12(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12;
#define FN13(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13;
#define FN14(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14;
#define FN15(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15;
#define FN16(p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16) p1; p2; p3; p4; p5; p6; p7; p8; p9; p10; p11; p12; p13; p14; p15; p16;
#define GET_LAST_ARG(p0, p1, p2, p3, p4, p5, p6, p7, p8, p9, p10, p11, p12, p13, p14, p15, p16, ...) p16
#define PARAMS_MAKER(...) GET_LAST_ARG(__VA_ARGS__, FN16, FN15, FN14, FN13, FN12, FN11, FN10, FN9, FN8, FN7, FN6, FN5, FN4, FN3, FN2, FN1, FN0)
/* Generate a struct containing the entry-point parameters and a "run"
* method which can access them implicitly via this-> */
#define ccl_gpu_kernel_signature(name, ...) \
struct kernel_gpu_##name \
{ \
PARAMS_MAKER(__VA_ARGS__)(__VA_ARGS__) \
void run(thread MetalKernelContext& context, \
threadgroup int *simdgroup_offset, \
const uint metal_global_id, \
const ushort metal_local_id, \
const ushort metal_local_size, \
uint simdgroup_size, \
uint simd_lane_index, \
uint simd_group_index, \
uint num_simd_groups) ccl_global const; \
}; \
kernel void kernel_metal_##name(device const kernel_gpu_##name *params_struct, \
constant KernelParamsMetal &ccl_restrict _launch_params_metal, \
constant MetalAncillaries *_metal_ancillaries, \
threadgroup int *simdgroup_offset[[ threadgroup(0) ]], \
const uint metal_global_id [[thread_position_in_grid]], \
const ushort metal_local_id [[thread_position_in_threadgroup]], \
const ushort metal_local_size [[threads_per_threadgroup]], \
uint simdgroup_size [[threads_per_simdgroup]], \
uint simd_lane_index [[thread_index_in_simdgroup]], \
uint simd_group_index [[simdgroup_index_in_threadgroup]], \
uint num_simd_groups [[simdgroups_per_threadgroup]]) { \
MetalKernelContext context(_launch_params_metal, _metal_ancillaries); \
params_struct->run(context, simdgroup_offset, metal_global_id, metal_local_id, metal_local_size, simdgroup_size, simd_lane_index, simd_group_index, num_simd_groups); \
} \
void kernel_gpu_##name::run(thread MetalKernelContext& context, \
threadgroup int *simdgroup_offset, \
const uint metal_global_id, \
const ushort metal_local_id, \
const ushort metal_local_size, \
uint simdgroup_size, \
uint simd_lane_index, \
uint simd_group_index, \
uint num_simd_groups) ccl_global const
#define ccl_gpu_kernel_call(x) context.x
/* define a function object where "func" is the lambda body, and additional parameters are used to specify captured state */
#define ccl_gpu_kernel_lambda(func, ...) \
struct KernelLambda \
{ \
KernelLambda(ccl_private MetalKernelContext &_context) : context(_context) {} \
ccl_private MetalKernelContext &context; \
__VA_ARGS__; \
int operator()(const int state) const { return (func); } \
} ccl_gpu_kernel_lambda_pass(context)
// clang-format on
/* volumetric lambda functions - use function objects for lambda-like functionality */
#define VOLUME_READ_LAMBDA(function_call) \
struct FnObjectRead { \
KernelGlobals kg; \
ccl_private MetalKernelContext *context; \
int state; \
\
VolumeStack operator()(const int i) const \
{ \
return context->function_call; \
} \
} volume_read_lambda_pass{kg, this, state};
#define VOLUME_WRITE_LAMBDA(function_call) \
struct FnObjectWrite { \
KernelGlobals kg; \
ccl_private MetalKernelContext *context; \
int state; \
\
void operator()(const int i, VolumeStack entry) const \
{ \
context->function_call; \
} \
} volume_write_lambda_pass{kg, this, state};
/* make_type definitions with Metal style element initializers */
#ifdef make_float2
# undef make_float2
@@ -112,6 +230,7 @@ using namespace metal;
#define sinhf(x) sinh(float(x))
#define coshf(x) cosh(float(x))
#define tanhf(x) tanh(float(x))
#define saturatef(x) saturate(float(x))
/* Use native functions with possibly lower precision for performance,
* no issues found so far. */
@@ -124,3 +243,43 @@ using namespace metal;
#define logf(x) trigmode::log(float(x))
#define NULL 0
#define __device__
/* texture bindings and sampler setup */
struct Texture2DParamsMetal {
texture2d<float, access::sample> tex;
};
struct Texture3DParamsMetal {
texture3d<float, access::sample> tex;
};
struct MetalAncillaries {
device Texture2DParamsMetal *textures_2d;
device Texture3DParamsMetal *textures_3d;
};
#include "util/half.h"
#include "util/types.h"
enum SamplerType {
SamplerFilterNearest_AddressRepeat,
SamplerFilterNearest_AddressClampEdge,
SamplerFilterNearest_AddressClampZero,
SamplerFilterLinear_AddressRepeat,
SamplerFilterLinear_AddressClampEdge,
SamplerFilterLinear_AddressClampZero,
SamplerCount
};
constant constexpr array<sampler, SamplerCount> metal_samplers = {
sampler(address::repeat, filter::nearest),
sampler(address::clamp_to_edge, filter::nearest),
sampler(address::clamp_to_zero, filter::nearest),
sampler(address::repeat, filter::linear),
sampler(address::clamp_to_edge, filter::linear),
sampler(address::clamp_to_zero, filter::linear),
};

View File

@@ -0,0 +1,79 @@
/*
* Copyright 2021 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.
*/
// clang-format off
/* Open the Metal kernel context class
* Necessary to access resource bindings */
class MetalKernelContext {
public:
constant KernelParamsMetal &launch_params_metal;
constant MetalAncillaries *metal_ancillaries;
MetalKernelContext(constant KernelParamsMetal &_launch_params_metal, constant MetalAncillaries * _metal_ancillaries)
: launch_params_metal(_launch_params_metal), metal_ancillaries(_metal_ancillaries)
{}
/* texture fetch adapter functions */
typedef uint64_t ccl_gpu_tex_object;
template<typename T>
inline __attribute__((__always_inline__))
T ccl_gpu_tex_object_read_2D(ccl_gpu_tex_object tex, float x, float y) const {
kernel_assert(0);
return 0;
}
template<typename T>
inline __attribute__((__always_inline__))
T ccl_gpu_tex_object_read_3D(ccl_gpu_tex_object tex, float x, float y, float z) const {
kernel_assert(0);
return 0;
}
// texture2d
template<>
inline __attribute__((__always_inline__))
float4 ccl_gpu_tex_object_read_2D(ccl_gpu_tex_object tex, float x, float y) const {
const uint tid(tex);
const uint sid(tex >> 32);
return metal_ancillaries->textures_2d[tid].tex.sample(metal_samplers[sid], float2(x, y));
}
template<>
inline __attribute__((__always_inline__))
float ccl_gpu_tex_object_read_2D(ccl_gpu_tex_object tex, float x, float y) const {
const uint tid(tex);
const uint sid(tex >> 32);
return metal_ancillaries->textures_2d[tid].tex.sample(metal_samplers[sid], float2(x, y)).x;
}
// texture3d
template<>
inline __attribute__((__always_inline__))
float4 ccl_gpu_tex_object_read_3D(ccl_gpu_tex_object tex, float x, float y, float z) const {
const uint tid(tex);
const uint sid(tex >> 32);
return metal_ancillaries->textures_3d[tid].tex.sample(metal_samplers[sid], float3(x, y, z));
}
template<>
inline __attribute__((__always_inline__))
float ccl_gpu_tex_object_read_3D(ccl_gpu_tex_object tex, float x, float y, float z) const {
const uint tid(tex);
const uint sid(tex >> 32);
return metal_ancillaries->textures_3d[tid].tex.sample(metal_samplers[sid], float3(x, y, z)).x;
}
# include "kernel/device/gpu/image.h"
// clang-format on

View File

@@ -0,0 +1,23 @@
/*
* Copyright 2021 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.
*/
}
; /* end of MetalKernelContext class definition */
/* Silently redirect into the MetalKernelContext instance */
/* NOTE: These macros will need maintaining as entry-points change. */
#undef kernel_integrator_state
#define kernel_integrator_state context.launch_params_metal.__integrator_state

View File

@@ -0,0 +1,51 @@
/*
* Copyright 2021 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.
*/
/* Constant Globals */
#include "kernel/types.h"
#include "kernel/util/profiling.h"
#include "kernel/integrator/state.h"
CCL_NAMESPACE_BEGIN
typedef struct KernelParamsMetal {
#define KERNEL_TEX(type, name) ccl_global const type *name;
#include "kernel/textures.h"
#undef KERNEL_TEX
const IntegratorStateGPU __integrator_state;
const KernelData data;
} KernelParamsMetal;
typedef struct KernelGlobalsGPU {
int unused[1];
} KernelGlobalsGPU;
typedef ccl_global const KernelGlobalsGPU *ccl_restrict KernelGlobals;
#define kernel_data launch_params_metal.data
#define kernel_integrator_state launch_params_metal.__integrator_state
/* data lookup defines */
#define kernel_tex_fetch(tex, index) launch_params_metal.tex[index]
#define kernel_tex_array(tex) launch_params_metal.tex
CCL_NAMESPACE_END

View File

@@ -0,0 +1,25 @@
/*
* Copyright 2021 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.
*/
/* Metal kernel entry points */
// clang-format off
#include "kernel/device/metal/compat.h"
#include "kernel/device/metal/globals.h"
#include "kernel/device/gpu/kernel.h"
// clang-format on

View File

@@ -49,10 +49,11 @@ typedef unsigned long long uint64_t;
__device__ __forceinline__ // Function calls are bad for OptiX performance, so inline everything
#define ccl_device_inline ccl_device
#define ccl_device_forceinline ccl_device
#define ccl_device_inline_method ccl_device
#define ccl_device_noinline __device__ __noinline__
#define ccl_device_noinline_cpu ccl_device
#define ccl_global
#define ccl_static_constant __constant__
#define ccl_inline_constant __constant__
#define ccl_device_constant __constant__ __device__
#define ccl_constant const
#define ccl_gpu_shared __shared__
@@ -76,6 +77,7 @@ typedef unsigned long long uint64_t;
#define ccl_gpu_block_idx_x (blockIdx.x)
#define ccl_gpu_grid_dim_x (gridDim.x)
#define ccl_gpu_warp_size (warpSize)
#define ccl_gpu_thread_mask(thread_warp) uint(0xFFFFFFFF >> (ccl_gpu_warp_size - thread_warp))
#define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x)
#define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x)
@@ -85,7 +87,6 @@ typedef unsigned long long uint64_t;
#define ccl_gpu_syncthreads() __syncthreads()
#define ccl_gpu_ballot(predicate) __ballot_sync(0xFFFFFFFF, predicate)
#define ccl_gpu_shfl_down_sync(mask, var, detla) __shfl_down_sync(mask, var, detla)
#define ccl_gpu_popc(x) __popc(x)
/* GPU texture objects */

View File

@@ -21,6 +21,8 @@
#include "kernel/device/gpu/image.h" /* Texture lookup uses normal CUDA intrinsics. */
#include "kernel/tables.h"
#include "kernel/integrator/state.h"
#include "kernel/integrator/state_flow.h"
#include "kernel/integrator/state_util.h"
@@ -44,7 +46,7 @@ template<typename T> ccl_device_forceinline T *get_payload_ptr_2()
ccl_device_forceinline int get_object_id()
{
#ifdef __OBJECT_MOTION__
/* Always get the the instance ID from the TLAS
/* Always get the instance ID from the TLAS
* There might be a motion transform node between TLAS and BLAS which does not have one. */
return optixGetInstanceIdFromHandle(optixGetTransformListHandle(0));
#else
@@ -159,9 +161,9 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
/* Record geometric normal. */
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
const float3 tri_a = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 0));
const float3 tri_b = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 1));
const float3 tri_c = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 2));
const float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0);
const float3 tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1);
const float3 tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
/* Continue tracing (without this the trace call would return after the first hit). */

View File

@@ -151,7 +151,8 @@ ccl_device_forceinline ccl_global float *kernel_accum_pixel_render_buffer(
ccl_device_inline int kernel_accum_sample(KernelGlobals kg,
ConstIntegratorState state,
ccl_global float *ccl_restrict render_buffer,
int sample)
int sample,
int sample_offset)
{
if (kernel_data.film.pass_sample_count == PASS_UNUSED) {
return sample;
@@ -159,7 +160,9 @@ ccl_device_inline int kernel_accum_sample(KernelGlobals kg,
ccl_global float *buffer = kernel_accum_pixel_render_buffer(kg, state, render_buffer);
return atomic_fetch_and_add_uint32((uint *)(buffer) + kernel_data.film.pass_sample_count, 1);
return atomic_fetch_and_add_uint32(
(ccl_global uint *)(buffer) + kernel_data.film.pass_sample_count, 1) +
sample_offset;
}
ccl_device void kernel_accum_adaptive_buffer(KernelGlobals kg,
@@ -550,7 +553,7 @@ ccl_device_inline void kernel_accum_background(KernelGlobals kg,
const bool is_transparent_background_ray,
ccl_global float *ccl_restrict render_buffer)
{
float3 contribution = INTEGRATOR_STATE(state, path, throughput) * L;
float3 contribution = float3(INTEGRATOR_STATE(state, path, throughput)) * L;
kernel_accum_clamp(kg, &contribution, INTEGRATOR_STATE(state, path, bounce) - 1);
ccl_global float *buffer = kernel_accum_pixel_render_buffer(kg, state, render_buffer);

View File

@@ -27,7 +27,12 @@ CCL_NAMESPACE_BEGIN
* Lookup of attributes is different between OSL and SVM, as OSL is ustring
* based while for SVM we use integer ids. */
ccl_device_inline uint subd_triangle_patch(KernelGlobals kg, ccl_private const ShaderData *sd);
/* Patch index for triangle, -1 if not subdivision triangle */
ccl_device_inline uint subd_triangle_patch(KernelGlobals kg, ccl_private const ShaderData *sd)
{
return (sd->prim != PRIM_NONE) ? kernel_tex_fetch(__tri_patch, sd->prim) : ~0;
}
ccl_device_inline uint attribute_primitive_type(KernelGlobals kg, ccl_private const ShaderData *sd)
{
@@ -106,9 +111,9 @@ ccl_device Transform primitive_attribute_matrix(KernelGlobals kg,
{
Transform tfm;
tfm.x = kernel_tex_fetch(__attributes_float3, desc.offset + 0);
tfm.y = kernel_tex_fetch(__attributes_float3, desc.offset + 1);
tfm.z = kernel_tex_fetch(__attributes_float3, desc.offset + 2);
tfm.x = kernel_tex_fetch(__attributes_float4, desc.offset + 0);
tfm.y = kernel_tex_fetch(__attributes_float4, desc.offset + 1);
tfm.z = kernel_tex_fetch(__attributes_float4, desc.offset + 2);
return tfm;
}

View File

@@ -126,8 +126,8 @@ ccl_device float3 curve_attribute_float3(KernelGlobals kg,
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
int k1 = k0 + 1;
float3 f0 = float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset + k0));
float3 f1 = float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset + k1));
float3 f0 = kernel_tex_fetch(__attributes_float3, desc.offset + k0);
float3 f1 = kernel_tex_fetch(__attributes_float3, desc.offset + k1);
# ifdef __RAY_DIFFERENTIALS__
if (dx)
@@ -149,7 +149,7 @@ ccl_device float3 curve_attribute_float3(KernelGlobals kg,
if (desc.element & (ATTR_ELEMENT_CURVE | ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
const int offset = (desc.element == ATTR_ELEMENT_CURVE) ? desc.offset + sd->prim :
desc.offset;
return float4_to_float3(kernel_tex_fetch(__attributes_float3, offset));
return kernel_tex_fetch(__attributes_float3, offset);
}
else {
return make_float3(0.0f, 0.0f, 0.0f);
@@ -168,8 +168,8 @@ ccl_device float4 curve_attribute_float4(KernelGlobals kg,
int k0 = curve.first_key + PRIMITIVE_UNPACK_SEGMENT(sd->type);
int k1 = k0 + 1;
float4 f0 = kernel_tex_fetch(__attributes_float3, desc.offset + k0);
float4 f1 = kernel_tex_fetch(__attributes_float3, desc.offset + k1);
float4 f0 = kernel_tex_fetch(__attributes_float4, desc.offset + k0);
float4 f1 = kernel_tex_fetch(__attributes_float4, desc.offset + k1);
# ifdef __RAY_DIFFERENTIALS__
if (dx)
@@ -191,7 +191,7 @@ ccl_device float4 curve_attribute_float4(KernelGlobals kg,
if (desc.element & (ATTR_ELEMENT_CURVE | ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
const int offset = (desc.element == ATTR_ELEMENT_CURVE) ? desc.offset + sd->prim :
desc.offset;
return kernel_tex_fetch(__attributes_float3, offset);
return kernel_tex_fetch(__attributes_float4, offset);
}
else {
return make_float4(0.0f, 0.0f, 0.0f, 0.0f);

View File

@@ -48,8 +48,8 @@ ccl_device_inline void motion_curve_keys_for_step_linear(KernelGlobals kg,
offset += step * numkeys;
keys[0] = kernel_tex_fetch(__attributes_float3, offset + k0);
keys[1] = kernel_tex_fetch(__attributes_float3, offset + k1);
keys[0] = kernel_tex_fetch(__attributes_float4, offset + k0);
keys[1] = kernel_tex_fetch(__attributes_float4, offset + k1);
}
}
@@ -106,10 +106,10 @@ ccl_device_inline void motion_curve_keys_for_step(KernelGlobals kg,
offset += step * numkeys;
keys[0] = kernel_tex_fetch(__attributes_float3, offset + k0);
keys[1] = kernel_tex_fetch(__attributes_float3, offset + k1);
keys[2] = kernel_tex_fetch(__attributes_float3, offset + k2);
keys[3] = kernel_tex_fetch(__attributes_float3, offset + k3);
keys[0] = kernel_tex_fetch(__attributes_float4, offset + k0);
keys[1] = kernel_tex_fetch(__attributes_float4, offset + k1);
keys[2] = kernel_tex_fetch(__attributes_float4, offset + k2);
keys[3] = kernel_tex_fetch(__attributes_float4, offset + k3);
}
}

View File

@@ -43,9 +43,9 @@ ccl_device_inline void motion_triangle_verts_for_step(KernelGlobals kg,
{
if (step == numsteps) {
/* center step: regular vertex location */
verts[0] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
verts[1] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
verts[2] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
verts[0] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 0);
verts[1] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 1);
verts[2] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 2);
}
else {
/* center step not store in this array */
@@ -54,9 +54,9 @@ ccl_device_inline void motion_triangle_verts_for_step(KernelGlobals kg,
offset += step * numverts;
verts[0] = float4_to_float3(kernel_tex_fetch(__attributes_float3, offset + tri_vindex.x));
verts[1] = float4_to_float3(kernel_tex_fetch(__attributes_float3, offset + tri_vindex.y));
verts[2] = float4_to_float3(kernel_tex_fetch(__attributes_float3, offset + tri_vindex.z));
verts[0] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.x);
verts[1] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.y);
verts[2] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.z);
}
}
@@ -70,9 +70,9 @@ ccl_device_inline void motion_triangle_normals_for_step(KernelGlobals kg,
{
if (step == numsteps) {
/* center step: regular vertex location */
normals[0] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.x));
normals[1] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.y));
normals[2] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.z));
normals[0] = kernel_tex_fetch(__tri_vnormal, tri_vindex.x);
normals[1] = kernel_tex_fetch(__tri_vnormal, tri_vindex.y);
normals[2] = kernel_tex_fetch(__tri_vnormal, tri_vindex.z);
}
else {
/* center step is not stored in this array */
@@ -81,9 +81,9 @@ ccl_device_inline void motion_triangle_normals_for_step(KernelGlobals kg,
offset += step * numverts;
normals[0] = float4_to_float3(kernel_tex_fetch(__attributes_float3, offset + tri_vindex.x));
normals[1] = float4_to_float3(kernel_tex_fetch(__attributes_float3, offset + tri_vindex.y));
normals[2] = float4_to_float3(kernel_tex_fetch(__attributes_float3, offset + tri_vindex.z));
normals[0] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.x);
normals[1] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.y);
normals[2] = kernel_tex_fetch(__attributes_float3, offset + tri_vindex.z);
}
}

View File

@@ -163,19 +163,7 @@ ccl_device_inline bool motion_triangle_intersect(KernelGlobals kg,
motion_triangle_vertices(kg, fobject, prim, time, verts);
/* Ray-triangle intersection, unoptimized. */
float t, u, v;
if (ray_triangle_intersect(P,
dir,
tmax,
#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
(ssef *)verts,
#else
verts[0],
verts[1],
verts[2],
#endif
&u,
&v,
&t)) {
if (ray_triangle_intersect(P, dir, tmax, verts[0], verts[1], verts[2], &u, &v, &t)) {
#ifdef __VISIBILITY_FLAG__
/* Visibility flag test. we do it here under the assumption
* that most triangles are culled by node flags.
@@ -229,19 +217,7 @@ ccl_device_inline bool motion_triangle_intersect_local(KernelGlobals kg,
motion_triangle_vertices(kg, local_object, prim, time, verts);
/* Ray-triangle intersection, unoptimized. */
float t, u, v;
if (!ray_triangle_intersect(P,
dir,
tmax,
# if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
(ssef *)verts,
# else
verts[0],
verts[1],
verts[2],
# endif
&u,
&v,
&t)) {
if (!ray_triangle_intersect(P, dir, tmax, verts[0], verts[1], verts[2], &u, &v, &t)) {
return false;
}

View File

@@ -380,7 +380,7 @@ ccl_device float3 patch_eval_float3(KernelGlobals kg,
*dv = make_float3(0.0f, 0.0f, 0.0f);
for (int i = 0; i < num_control; i++) {
float3 v = float4_to_float3(kernel_tex_fetch(__attributes_float3, offset + indices[i]));
float3 v = kernel_tex_fetch(__attributes_float3, offset + indices[i]);
val += v * weights[i];
if (du)
@@ -417,7 +417,7 @@ ccl_device float4 patch_eval_float4(KernelGlobals kg,
*dv = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
for (int i = 0; i < num_control; i++) {
float4 v = kernel_tex_fetch(__attributes_float3, offset + indices[i]);
float4 v = kernel_tex_fetch(__attributes_float4, offset + indices[i]);
val += v * weights[i];
if (du)

View File

@@ -284,18 +284,33 @@ ccl_device_inline float4 primitive_motion_vector(KernelGlobals kg,
int numverts, numkeys;
object_motion_info(kg, sd->object, NULL, &numverts, &numkeys);
/* lookup attributes */
motion_pre = primitive_surface_attribute_float3(kg, sd, desc, NULL, NULL);
desc.offset += (sd->type & PRIMITIVE_ALL_TRIANGLE) ? numverts : numkeys;
motion_post = primitive_surface_attribute_float3(kg, sd, desc, NULL, NULL);
#ifdef __HAIR__
if (is_curve_primitive && (sd->object_flag & SD_OBJECT_HAS_VERTEX_MOTION) == 0) {
object_position_transform(kg, sd, &motion_pre);
object_position_transform(kg, sd, &motion_post);
if (is_curve_primitive) {
motion_pre = float4_to_float3(curve_attribute_float4(kg, sd, desc, NULL, NULL));
desc.offset += numkeys;
motion_post = float4_to_float3(curve_attribute_float4(kg, sd, desc, NULL, NULL));
/* Curve */
if ((sd->object_flag & SD_OBJECT_HAS_VERTEX_MOTION) == 0) {
object_position_transform(kg, sd, &motion_pre);
object_position_transform(kg, sd, &motion_post);
}
}
else
#endif
if (sd->type & PRIMITIVE_ALL_TRIANGLE) {
/* Triangle */
if (subd_triangle_patch(kg, sd) == ~0) {
motion_pre = triangle_attribute_float3(kg, sd, desc, NULL, NULL);
desc.offset += numverts;
motion_post = triangle_attribute_float3(kg, sd, desc, NULL, NULL);
}
else {
motion_pre = subd_triangle_attribute_float3(kg, sd, desc, NULL, NULL);
desc.offset += numverts;
motion_post = subd_triangle_attribute_float3(kg, sd, desc, NULL, NULL);
}
}
}
/* object motion. note that depending on the mesh having motion vectors, this

View File

@@ -20,13 +20,6 @@
CCL_NAMESPACE_BEGIN
/* Patch index for triangle, -1 if not subdivision triangle */
ccl_device_inline uint subd_triangle_patch(KernelGlobals kg, ccl_private const ShaderData *sd)
{
return (sd->prim != PRIM_NONE) ? kernel_tex_fetch(__tri_patch, sd->prim) : ~0;
}
/* UV coords of triangle within patch */
ccl_device_inline void subd_triangle_patch_uv(KernelGlobals kg,
@@ -443,8 +436,8 @@ ccl_device_noinline float3 subd_triangle_attribute_float3(KernelGlobals kg,
if (dy)
*dy = make_float3(0.0f, 0.0f, 0.0f);
return float4_to_float3(
kernel_tex_fetch(__attributes_float3, desc.offset + subd_triangle_patch_face(kg, patch)));
return kernel_tex_fetch(__attributes_float3,
desc.offset + subd_triangle_patch_face(kg, patch));
}
else if (desc.element == ATTR_ELEMENT_VERTEX || desc.element == ATTR_ELEMENT_VERTEX_MOTION) {
float2 uv[3];
@@ -452,10 +445,10 @@ ccl_device_noinline float3 subd_triangle_attribute_float3(KernelGlobals kg,
uint4 v = subd_triangle_patch_indices(kg, patch);
float3 f0 = float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset + v.x));
float3 f1 = float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset + v.y));
float3 f2 = float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset + v.z));
float3 f3 = float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset + v.w));
float3 f0 = kernel_tex_fetch(__attributes_float3, desc.offset + v.x);
float3 f1 = kernel_tex_fetch(__attributes_float3, desc.offset + v.y);
float3 f2 = kernel_tex_fetch(__attributes_float3, desc.offset + v.z);
float3 f3 = kernel_tex_fetch(__attributes_float3, desc.offset + v.w);
if (subd_triangle_patch_num_corners(kg, patch) != 4) {
f1 = (f1 + f0) * 0.5f;
@@ -484,10 +477,10 @@ ccl_device_noinline float3 subd_triangle_attribute_float3(KernelGlobals kg,
float3 f0, f1, f2, f3;
f0 = float4_to_float3(kernel_tex_fetch(__attributes_float3, corners[0] + desc.offset));
f1 = float4_to_float3(kernel_tex_fetch(__attributes_float3, corners[1] + desc.offset));
f2 = float4_to_float3(kernel_tex_fetch(__attributes_float3, corners[2] + desc.offset));
f3 = float4_to_float3(kernel_tex_fetch(__attributes_float3, corners[3] + desc.offset));
f0 = kernel_tex_fetch(__attributes_float3, corners[0] + desc.offset);
f1 = kernel_tex_fetch(__attributes_float3, corners[1] + desc.offset);
f2 = kernel_tex_fetch(__attributes_float3, corners[2] + desc.offset);
f3 = kernel_tex_fetch(__attributes_float3, corners[3] + desc.offset);
if (subd_triangle_patch_num_corners(kg, patch) != 4) {
f1 = (f1 + f0) * 0.5f;
@@ -513,7 +506,7 @@ ccl_device_noinline float3 subd_triangle_attribute_float3(KernelGlobals kg,
if (dy)
*dy = make_float3(0.0f, 0.0f, 0.0f);
return float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset));
return kernel_tex_fetch(__attributes_float3, desc.offset);
}
else {
if (dx)
@@ -590,7 +583,7 @@ ccl_device_noinline float4 subd_triangle_attribute_float4(KernelGlobals kg,
if (dy)
*dy = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
return kernel_tex_fetch(__attributes_float3,
return kernel_tex_fetch(__attributes_float4,
desc.offset + subd_triangle_patch_face(kg, patch));
}
else if (desc.element == ATTR_ELEMENT_VERTEX || desc.element == ATTR_ELEMENT_VERTEX_MOTION) {
@@ -599,10 +592,10 @@ ccl_device_noinline float4 subd_triangle_attribute_float4(KernelGlobals kg,
uint4 v = subd_triangle_patch_indices(kg, patch);
float4 f0 = kernel_tex_fetch(__attributes_float3, desc.offset + v.x);
float4 f1 = kernel_tex_fetch(__attributes_float3, desc.offset + v.y);
float4 f2 = kernel_tex_fetch(__attributes_float3, desc.offset + v.z);
float4 f3 = kernel_tex_fetch(__attributes_float3, desc.offset + v.w);
float4 f0 = kernel_tex_fetch(__attributes_float4, desc.offset + v.x);
float4 f1 = kernel_tex_fetch(__attributes_float4, desc.offset + v.y);
float4 f2 = kernel_tex_fetch(__attributes_float4, desc.offset + v.z);
float4 f3 = kernel_tex_fetch(__attributes_float4, desc.offset + v.w);
if (subd_triangle_patch_num_corners(kg, patch) != 4) {
f1 = (f1 + f0) * 0.5f;
@@ -642,10 +635,10 @@ ccl_device_noinline float4 subd_triangle_attribute_float4(KernelGlobals kg,
color_uchar4_to_float4(kernel_tex_fetch(__attributes_uchar4, corners[3] + desc.offset)));
}
else {
f0 = kernel_tex_fetch(__attributes_float3, corners[0] + desc.offset);
f1 = kernel_tex_fetch(__attributes_float3, corners[1] + desc.offset);
f2 = kernel_tex_fetch(__attributes_float3, corners[2] + desc.offset);
f3 = kernel_tex_fetch(__attributes_float3, corners[3] + desc.offset);
f0 = kernel_tex_fetch(__attributes_float4, corners[0] + desc.offset);
f1 = kernel_tex_fetch(__attributes_float4, corners[1] + desc.offset);
f2 = kernel_tex_fetch(__attributes_float4, corners[2] + desc.offset);
f3 = kernel_tex_fetch(__attributes_float4, corners[3] + desc.offset);
}
if (subd_triangle_patch_num_corners(kg, patch) != 4) {
@@ -672,7 +665,7 @@ ccl_device_noinline float4 subd_triangle_attribute_float4(KernelGlobals kg,
if (dy)
*dy = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
return kernel_tex_fetch(__attributes_float3, desc.offset);
return kernel_tex_fetch(__attributes_float4, desc.offset);
}
else {
if (dx)

View File

@@ -29,9 +29,9 @@ ccl_device_inline float3 triangle_normal(KernelGlobals kg, ccl_private ShaderDat
{
/* load triangle vertices */
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, sd->prim);
const float3 v0 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
const float3 v1 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
const float3 v2 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
const float3 v0 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 0);
const float3 v1 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 1);
const float3 v2 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 2);
/* return normal */
if (sd->object_flag & SD_OBJECT_NEGATIVE_SCALE_APPLIED) {
@@ -54,9 +54,9 @@ ccl_device_inline void triangle_point_normal(KernelGlobals kg,
{
/* load triangle vertices */
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
float3 v0 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
float3 v1 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
float3 v2 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
float3 v0 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 0);
float3 v1 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 1);
float3 v2 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 2);
/* compute point */
float t = 1.0f - u - v;
*P = (u * v0 + v * v1 + t * v2);
@@ -78,9 +78,9 @@ ccl_device_inline void triangle_point_normal(KernelGlobals kg,
ccl_device_inline void triangle_vertices(KernelGlobals kg, int prim, float3 P[3])
{
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
P[0] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
P[1] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
P[2] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
P[0] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 0);
P[1] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 1);
P[2] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 2);
}
/* Triangle vertex locations and vertex normals */
@@ -91,12 +91,12 @@ ccl_device_inline void triangle_vertices_and_normals(KernelGlobals kg,
float3 N[3])
{
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
P[0] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
P[1] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
P[2] = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
N[0] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.x));
N[1] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.y));
N[2] = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.z));
P[0] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 0);
P[1] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 1);
P[2] = kernel_tex_fetch(__tri_verts, tri_vindex.w + 2);
N[0] = kernel_tex_fetch(__tri_vnormal, tri_vindex.x);
N[1] = kernel_tex_fetch(__tri_vnormal, tri_vindex.y);
N[2] = kernel_tex_fetch(__tri_vnormal, tri_vindex.z);
}
/* Interpolate smooth vertex normal from vertices */
@@ -106,9 +106,9 @@ triangle_smooth_normal(KernelGlobals kg, float3 Ng, int prim, float u, float v)
{
/* load triangle vertices */
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
float3 n0 = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.x));
float3 n1 = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.y));
float3 n2 = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.z));
float3 n0 = kernel_tex_fetch(__tri_vnormal, tri_vindex.x);
float3 n1 = kernel_tex_fetch(__tri_vnormal, tri_vindex.y);
float3 n2 = kernel_tex_fetch(__tri_vnormal, tri_vindex.z);
float3 N = safe_normalize((1.0f - u - v) * n2 + u * n0 + v * n1);
@@ -120,9 +120,9 @@ ccl_device_inline float3 triangle_smooth_normal_unnormalized(
{
/* load triangle vertices */
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
float3 n0 = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.x));
float3 n1 = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.y));
float3 n2 = float4_to_float3(kernel_tex_fetch(__tri_vnormal, tri_vindex.z));
float3 n0 = kernel_tex_fetch(__tri_vnormal, tri_vindex.x);
float3 n1 = kernel_tex_fetch(__tri_vnormal, tri_vindex.y);
float3 n2 = kernel_tex_fetch(__tri_vnormal, tri_vindex.z);
/* ensure that the normals are in object space */
if (sd->object_flag & SD_OBJECT_TRANSFORM_APPLIED) {
@@ -145,9 +145,9 @@ ccl_device_inline void triangle_dPdudv(KernelGlobals kg,
{
/* fetch triangle vertex coordinates */
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, prim);
const float3 p0 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 0));
const float3 p1 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 1));
const float3 p2 = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex.w + 2));
const float3 p0 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 0);
const float3 p1 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 1);
const float3 p2 = kernel_tex_fetch(__tri_verts, tri_vindex.w + 2);
/* compute derivatives of P w.r.t. uv */
*dPdu = (p0 - p2);
@@ -267,15 +267,15 @@ ccl_device float3 triangle_attribute_float3(KernelGlobals kg,
if (desc.element & (ATTR_ELEMENT_VERTEX | ATTR_ELEMENT_VERTEX_MOTION)) {
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, sd->prim);
f0 = float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset + tri_vindex.x));
f1 = float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset + tri_vindex.y));
f2 = float4_to_float3(kernel_tex_fetch(__attributes_float3, desc.offset + tri_vindex.z));
f0 = kernel_tex_fetch(__attributes_float3, desc.offset + tri_vindex.x);
f1 = kernel_tex_fetch(__attributes_float3, desc.offset + tri_vindex.y);
f2 = kernel_tex_fetch(__attributes_float3, desc.offset + tri_vindex.z);
}
else {
const int tri = desc.offset + sd->prim * 3;
f0 = float4_to_float3(kernel_tex_fetch(__attributes_float3, tri + 0));
f1 = float4_to_float3(kernel_tex_fetch(__attributes_float3, tri + 1));
f2 = float4_to_float3(kernel_tex_fetch(__attributes_float3, tri + 2));
f0 = kernel_tex_fetch(__attributes_float3, tri + 0);
f1 = kernel_tex_fetch(__attributes_float3, tri + 1);
f2 = kernel_tex_fetch(__attributes_float3, tri + 2);
}
#ifdef __RAY_DIFFERENTIALS__
@@ -298,7 +298,7 @@ ccl_device float3 triangle_attribute_float3(KernelGlobals kg,
if (desc.element & (ATTR_ELEMENT_FACE | ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
const int offset = (desc.element == ATTR_ELEMENT_FACE) ? desc.offset + sd->prim :
desc.offset;
return float4_to_float3(kernel_tex_fetch(__attributes_float3, offset));
return kernel_tex_fetch(__attributes_float3, offset);
}
else {
return make_float3(0.0f, 0.0f, 0.0f);
@@ -318,16 +318,16 @@ ccl_device float4 triangle_attribute_float4(KernelGlobals kg,
if (desc.element & (ATTR_ELEMENT_VERTEX | ATTR_ELEMENT_VERTEX_MOTION)) {
const uint4 tri_vindex = kernel_tex_fetch(__tri_vindex, sd->prim);
f0 = kernel_tex_fetch(__attributes_float3, desc.offset + tri_vindex.x);
f1 = kernel_tex_fetch(__attributes_float3, desc.offset + tri_vindex.y);
f2 = kernel_tex_fetch(__attributes_float3, desc.offset + tri_vindex.z);
f0 = kernel_tex_fetch(__attributes_float4, desc.offset + tri_vindex.x);
f1 = kernel_tex_fetch(__attributes_float4, desc.offset + tri_vindex.y);
f2 = kernel_tex_fetch(__attributes_float4, desc.offset + tri_vindex.z);
}
else {
const int tri = desc.offset + sd->prim * 3;
if (desc.element == ATTR_ELEMENT_CORNER) {
f0 = kernel_tex_fetch(__attributes_float3, tri + 0);
f1 = kernel_tex_fetch(__attributes_float3, tri + 1);
f2 = kernel_tex_fetch(__attributes_float3, tri + 2);
f0 = kernel_tex_fetch(__attributes_float4, tri + 0);
f1 = kernel_tex_fetch(__attributes_float4, tri + 1);
f2 = kernel_tex_fetch(__attributes_float4, tri + 2);
}
else {
f0 = color_srgb_to_linear_v4(
@@ -359,7 +359,7 @@ ccl_device float4 triangle_attribute_float4(KernelGlobals kg,
if (desc.element & (ATTR_ELEMENT_FACE | ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
const int offset = (desc.element == ATTR_ELEMENT_FACE) ? desc.offset + sd->prim :
desc.offset;
return kernel_tex_fetch(__attributes_float3, offset);
return kernel_tex_fetch(__attributes_float4, offset);
}
else {
return make_float4(0.0f, 0.0f, 0.0f, 0.0f);

View File

@@ -37,27 +37,11 @@ ccl_device_inline bool triangle_intersect(KernelGlobals kg,
{
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
const ssef *ssef_verts = (ssef *)&kg->__tri_verts.data[tri_vindex];
#else
const float4 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
const float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
#endif
float t, u, v;
if (ray_triangle_intersect(P,
dir,
tmax,
#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
ssef_verts,
#else
float4_to_float3(tri_a),
float4_to_float3(tri_b),
float4_to_float3(tri_c),
#endif
&u,
&v,
&t)) {
if (ray_triangle_intersect(P, dir, tmax, tri_a, tri_b, tri_c, &u, &v, &t)) {
#ifdef __VISIBILITY_FLAG__
/* Visibility flag test. we do it here under the assumption
* that most triangles are culled by node flags.
@@ -106,27 +90,11 @@ ccl_device_inline bool triangle_intersect_local(KernelGlobals kg,
const int prim = kernel_tex_fetch(__prim_index, prim_addr);
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, prim).w;
# if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
const ssef *ssef_verts = (ssef *)&kg->__tri_verts.data[tri_vindex];
# else
const float3 tri_a = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 0)),
tri_b = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 1)),
tri_c = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 2));
# endif
const float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
float t, u, v;
if (!ray_triangle_intersect(P,
dir,
tmax,
# if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
ssef_verts,
# else
tri_a,
tri_b,
tri_c,
# endif
&u,
&v,
&t)) {
if (!ray_triangle_intersect(P, dir, tmax, tri_a, tri_b, tri_c, &u, &v, &t)) {
return false;
}
@@ -178,11 +146,6 @@ ccl_device_inline bool triangle_intersect_local(KernelGlobals kg,
isect->t = t;
/* Record geometric normal. */
# if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
const float3 tri_a = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 0)),
tri_b = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 1)),
tri_c = float4_to_float3(kernel_tex_fetch(__tri_verts, tri_vindex + 2));
# endif
local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
return false;
@@ -223,9 +186,9 @@ ccl_device_inline float3 triangle_refine(KernelGlobals kg,
P = P + D * t;
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, isect_prim).w;
const float4 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
const packed_float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
float3 edge1 = make_float3(tri_a.x - tri_c.x, tri_a.y - tri_c.y, tri_a.z - tri_c.z);
float3 edge2 = make_float3(tri_b.x - tri_c.x, tri_b.y - tri_c.y, tri_b.z - tri_c.z);
float3 tvec = make_float3(P.x - tri_c.x, P.y - tri_c.y, P.z - tri_c.z);
@@ -280,9 +243,9 @@ ccl_device_inline float3 triangle_refine_local(KernelGlobals kg,
# ifdef __INTERSECTION_REFINE__
const uint tri_vindex = kernel_tex_fetch(__tri_vindex, isect_prim).w;
const float4 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
const packed_float3 tri_a = kernel_tex_fetch(__tri_verts, tri_vindex + 0),
tri_b = kernel_tex_fetch(__tri_verts, tri_vindex + 1),
tri_c = kernel_tex_fetch(__tri_verts, tri_vindex + 2);
float3 edge1 = make_float3(tri_a.x - tri_c.x, tri_a.y - tri_c.y, tri_a.z - tri_c.z);
float3 edge2 = make_float3(tri_b.x - tri_c.x, tri_b.y - tri_c.y, tri_b.z - tri_c.z);
float3 tvec = make_float3(P.x - tri_c.x, P.y - tri_c.y, P.z - tri_c.z);

View File

@@ -75,7 +75,7 @@ ccl_device float4 volume_attribute_float4(KernelGlobals kg,
const AttributeDescriptor desc)
{
if (desc.element & (ATTR_ELEMENT_OBJECT | ATTR_ELEMENT_MESH)) {
return kernel_tex_fetch(__attributes_float3, desc.offset);
return kernel_tex_fetch(__attributes_float4, desc.offset);
}
else if (desc.element == ATTR_ELEMENT_VOXEL) {
/* todo: optimize this so we don't have to transform both here and in

View File

@@ -65,7 +65,8 @@ ccl_device bool integrator_init_from_bake(KernelGlobals kg,
}
/* Always count the sample, even if the camera sample will reject the ray. */
const int sample = kernel_accum_sample(kg, state, render_buffer, scheduled_sample);
const int sample = kernel_accum_sample(
kg, state, render_buffer, scheduled_sample, tile->sample_offset);
/* Setup render buffers. */
const int index = INTEGRATOR_STATE(state, path, render_pixel_index);

View File

@@ -89,7 +89,8 @@ ccl_device bool integrator_init_from_camera(KernelGlobals kg,
* This logic allows to both count actual number of samples per pixel, and to add samples to this
* pixel after it was converged and samples were added somewhere else (in which case the
* `scheduled_sample` will be different from actual number of samples in this pixel). */
const int sample = kernel_accum_sample(kg, state, render_buffer, scheduled_sample);
const int sample = kernel_accum_sample(
kg, state, render_buffer, scheduled_sample, tile->sample_offset);
/* Initialize random number seed for path. */
const uint rng_hash = path_rng_hash_init(kg, sample, x, y);

View File

@@ -20,7 +20,6 @@
#include "kernel/integrator/shader_eval.h"
#include "kernel/light/light.h"
#include "kernel/light/sample.h"
#include "kernel/sample/mis.h"
CCL_NAMESPACE_BEGIN
@@ -81,8 +80,7 @@ ccl_device float3 integrator_eval_background_shader(KernelGlobals kg,
/* multiple importance sampling, get background light pdf for ray
* direction, and compute weight with respect to BSDF pdf */
const float pdf = background_light_pdf(kg, ray_P - ray_D * mis_ray_t, ray_D);
const float mis_weight = power_heuristic(mis_ray_pdf, pdf);
const float mis_weight = light_sample_mis_weight_forward(kg, mis_ray_pdf, pdf);
L *= mis_weight;
}
# endif
@@ -169,7 +167,7 @@ ccl_device_inline void integrate_distant_lights(KernelGlobals kg,
/* multiple importance sampling, get regular light pdf,
* and compute weight with respect to BSDF pdf */
const float mis_ray_pdf = INTEGRATOR_STATE(state, path, mis_ray_pdf);
const float mis_weight = power_heuristic(mis_ray_pdf, ls.pdf);
const float mis_weight = light_sample_mis_weight_forward(kg, mis_ray_pdf, ls.pdf);
light_eval *= mis_weight;
}

View File

@@ -84,7 +84,7 @@ ccl_device_inline void integrate_light(KernelGlobals kg,
/* multiple importance sampling, get regular light pdf,
* and compute weight with respect to BSDF pdf */
const float mis_ray_pdf = INTEGRATOR_STATE(state, path, mis_ray_pdf);
const float mis_weight = power_heuristic(mis_ray_pdf, ls.pdf);
const float mis_weight = light_sample_mis_weight_forward(kg, mis_ray_pdf, ls.pdf);
light_eval *= mis_weight;
}

View File

@@ -95,8 +95,8 @@ ccl_device_inline void integrate_transparent_volume_shadow(KernelGlobals kg,
shader_setup_from_volume(kg, shadow_sd, &ray);
const float step_size = volume_stack_step_size(
kg, [=](const int i) { return integrator_state_read_shadow_volume_stack(state, i); });
VOLUME_READ_LAMBDA(integrator_state_read_shadow_volume_stack(state, i));
const float step_size = volume_stack_step_size(kg, volume_read_lambda_pass);
volume_shadow_heterogeneous(kg, state, &ray, shadow_sd, throughput, step_size);
}

View File

@@ -27,8 +27,6 @@
#include "kernel/light/light.h"
#include "kernel/light/sample.h"
#include "kernel/sample/mis.h"
CCL_NAMESPACE_BEGIN
ccl_device_forceinline void integrate_surface_shader_setup(KernelGlobals kg,
@@ -95,8 +93,7 @@ ccl_device_forceinline void integrate_surface_emission(KernelGlobals kg,
/* Multiple importance sampling, get triangle light pdf,
* and compute weight with respect to BSDF pdf. */
float pdf = triangle_light_pdf(kg, sd, t);
float mis_weight = power_heuristic(bsdf_pdf, pdf);
float mis_weight = light_sample_mis_weight_forward(kg, bsdf_pdf, pdf);
L *= mis_weight;
}
@@ -155,7 +152,7 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
bsdf_eval_mul3(&bsdf_eval, light_eval / ls.pdf);
if (ls.shader & SHADER_USE_MIS) {
const float mis_weight = power_heuristic(ls.pdf, bsdf_pdf);
const float mis_weight = light_sample_mis_weight_nee(kg, ls.pdf, bsdf_pdf);
bsdf_eval_mul(&bsdf_eval, mis_weight);
}
@@ -195,12 +192,13 @@ ccl_device_forceinline void integrate_surface_direct_light(KernelGlobals kg,
const float3 throughput = INTEGRATOR_STATE(state, path, throughput) * bsdf_eval_sum(&bsdf_eval);
if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
const float3 pass_diffuse_weight = (bounce == 0) ?
bsdf_eval_pass_diffuse_weight(&bsdf_eval) :
INTEGRATOR_STATE(state, path, pass_diffuse_weight);
const float3 pass_glossy_weight = (bounce == 0) ?
bsdf_eval_pass_glossy_weight(&bsdf_eval) :
INTEGRATOR_STATE(state, path, pass_glossy_weight);
const packed_float3 pass_diffuse_weight =
(bounce == 0) ? packed_float3(bsdf_eval_pass_diffuse_weight(&bsdf_eval)) :
INTEGRATOR_STATE(state, path, pass_diffuse_weight);
const packed_float3 pass_glossy_weight = (bounce == 0) ?
packed_float3(
bsdf_eval_pass_glossy_weight(&bsdf_eval)) :
INTEGRATOR_STATE(state, path, pass_glossy_weight);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_diffuse_weight) = pass_diffuse_weight;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = pass_glossy_weight;
}

View File

@@ -27,8 +27,6 @@
#include "kernel/light/light.h"
#include "kernel/light/sample.h"
#include "kernel/sample/mis.h"
CCL_NAMESPACE_BEGIN
#ifdef __VOLUME__
@@ -78,9 +76,8 @@ ccl_device_inline bool shadow_volume_shader_sample(KernelGlobals kg,
ccl_private ShaderData *ccl_restrict sd,
ccl_private float3 *ccl_restrict extinction)
{
shader_eval_volume<true>(kg, state, sd, PATH_RAY_SHADOW, [=](const int i) {
return integrator_state_read_shadow_volume_stack(state, i);
});
VOLUME_READ_LAMBDA(integrator_state_read_shadow_volume_stack(state, i))
shader_eval_volume<true>(kg, state, sd, PATH_RAY_SHADOW, volume_read_lambda_pass);
if (!(sd->flag & SD_EXTINCTION)) {
return false;
@@ -98,9 +95,8 @@ ccl_device_inline bool volume_shader_sample(KernelGlobals kg,
ccl_private VolumeShaderCoefficients *coeff)
{
const uint32_t path_flag = INTEGRATOR_STATE(state, path, flag);
shader_eval_volume<false>(kg, state, sd, path_flag, [=](const int i) {
return integrator_state_read_volume_stack(state, i);
});
VOLUME_READ_LAMBDA(integrator_state_read_volume_stack(state, i))
shader_eval_volume<false>(kg, state, sd, path_flag, volume_read_lambda_pass);
if (!(sd->flag & (SD_EXTINCTION | SD_SCATTER | SD_EMISSION))) {
return false;
@@ -772,7 +768,7 @@ ccl_device_forceinline void integrate_volume_direct_light(
const float phase_pdf = shader_volume_phase_eval(kg, sd, phases, ls->D, &phase_eval);
if (ls->shader & SHADER_USE_MIS) {
float mis_weight = power_heuristic(ls->pdf, phase_pdf);
float mis_weight = light_sample_mis_weight_nee(kg, ls->pdf, phase_pdf);
bsdf_eval_mul(&phase_eval, mis_weight);
}
@@ -805,9 +801,10 @@ ccl_device_forceinline void integrate_volume_direct_light(
const float3 throughput_phase = throughput * bsdf_eval_sum(&phase_eval);
if (kernel_data.kernel_features & KERNEL_FEATURE_LIGHT_PASSES) {
const float3 pass_diffuse_weight = (bounce == 0) ?
one_float3() :
INTEGRATOR_STATE(state, path, pass_diffuse_weight);
const packed_float3 pass_diffuse_weight = (bounce == 0) ?
packed_float3(one_float3()) :
INTEGRATOR_STATE(
state, path, pass_diffuse_weight);
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_diffuse_weight) = pass_diffuse_weight;
INTEGRATOR_STATE_WRITE(shadow_state, shadow_path, pass_glossy_weight) = zero_float3();
}
@@ -932,8 +929,8 @@ ccl_device VolumeIntegrateEvent volume_integrate(KernelGlobals kg,
VOLUME_SAMPLE_DISTANCE;
/* Step through volume. */
const float step_size = volume_stack_step_size(
kg, [=](const int i) { return integrator_state_read_volume_stack(state, i); });
VOLUME_READ_LAMBDA(integrator_state_read_volume_stack(state, i))
const float step_size = volume_stack_step_size(kg, volume_read_lambda_pass);
/* TODO: expensive to zero closures? */
VolumeIntegrateResult result = {};

View File

@@ -40,15 +40,15 @@ KERNEL_STRUCT_MEMBER(shadow_path, uint16_t, queued_kernel, KERNEL_FEATURE_PATH_T
/* enum PathRayFlag */
KERNEL_STRUCT_MEMBER(shadow_path, uint32_t, flag, KERNEL_FEATURE_PATH_TRACING)
/* Throughput. */
KERNEL_STRUCT_MEMBER(shadow_path, float3, throughput, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_path, packed_float3, throughput, KERNEL_FEATURE_PATH_TRACING)
/* Throughput for shadow pass. */
KERNEL_STRUCT_MEMBER(shadow_path,
float3,
packed_float3,
unshadowed_throughput,
KERNEL_FEATURE_SHADOW_PASS | KERNEL_FEATURE_AO_ADDITIVE)
/* Ratio of throughput to distinguish diffuse / glossy / transmission render passes. */
KERNEL_STRUCT_MEMBER(shadow_path, float3, pass_diffuse_weight, KERNEL_FEATURE_LIGHT_PASSES)
KERNEL_STRUCT_MEMBER(shadow_path, float3, pass_glossy_weight, KERNEL_FEATURE_LIGHT_PASSES)
KERNEL_STRUCT_MEMBER(shadow_path, packed_float3, pass_diffuse_weight, KERNEL_FEATURE_LIGHT_PASSES)
KERNEL_STRUCT_MEMBER(shadow_path, packed_float3, pass_glossy_weight, KERNEL_FEATURE_LIGHT_PASSES)
/* Number of intersections found by ray-tracing. */
KERNEL_STRUCT_MEMBER(shadow_path, uint16_t, num_hits, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_END(shadow_path)
@@ -56,8 +56,8 @@ KERNEL_STRUCT_END(shadow_path)
/********************************** Shadow Ray *******************************/
KERNEL_STRUCT_BEGIN(shadow_ray)
KERNEL_STRUCT_MEMBER(shadow_ray, float3, P, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float3, D, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, packed_float3, P, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, packed_float3, D, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float, t, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float, time, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(shadow_ray, float, dP, KERNEL_FEATURE_PATH_TRACING)

View File

@@ -59,12 +59,12 @@ KERNEL_STRUCT_MEMBER(path, float, min_ray_pdf, KERNEL_FEATURE_PATH_TRACING)
/* Continuation probability for path termination. */
KERNEL_STRUCT_MEMBER(path, float, continuation_probability, KERNEL_FEATURE_PATH_TRACING)
/* Throughput. */
KERNEL_STRUCT_MEMBER(path, float3, throughput, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(path, packed_float3, throughput, KERNEL_FEATURE_PATH_TRACING)
/* Ratio of throughput to distinguish diffuse / glossy / transmission render passes. */
KERNEL_STRUCT_MEMBER(path, float3, pass_diffuse_weight, KERNEL_FEATURE_LIGHT_PASSES)
KERNEL_STRUCT_MEMBER(path, float3, pass_glossy_weight, KERNEL_FEATURE_LIGHT_PASSES)
KERNEL_STRUCT_MEMBER(path, packed_float3, pass_diffuse_weight, KERNEL_FEATURE_LIGHT_PASSES)
KERNEL_STRUCT_MEMBER(path, packed_float3, pass_glossy_weight, KERNEL_FEATURE_LIGHT_PASSES)
/* Denoising. */
KERNEL_STRUCT_MEMBER(path, float3, denoising_feature_throughput, KERNEL_FEATURE_DENOISING)
KERNEL_STRUCT_MEMBER(path, packed_float3, denoising_feature_throughput, KERNEL_FEATURE_DENOISING)
/* Shader sorting. */
/* TODO: compress as uint16? or leave out entirely and recompute key in sorting code? */
KERNEL_STRUCT_MEMBER(path, uint32_t, shader_sort_key, KERNEL_FEATURE_PATH_TRACING)
@@ -73,8 +73,8 @@ KERNEL_STRUCT_END(path)
/************************************** Ray ***********************************/
KERNEL_STRUCT_BEGIN(ray)
KERNEL_STRUCT_MEMBER(ray, float3, P, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(ray, float3, D, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(ray, packed_float3, P, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(ray, packed_float3, D, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(ray, float, t, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(ray, float, time, KERNEL_FEATURE_PATH_TRACING)
KERNEL_STRUCT_MEMBER(ray, float, dP, KERNEL_FEATURE_PATH_TRACING)
@@ -96,10 +96,10 @@ KERNEL_STRUCT_END(isect)
/*************** Subsurface closure state for subsurface kernel ***************/
KERNEL_STRUCT_BEGIN(subsurface)
KERNEL_STRUCT_MEMBER(subsurface, float3, albedo, KERNEL_FEATURE_SUBSURFACE)
KERNEL_STRUCT_MEMBER(subsurface, float3, radius, KERNEL_FEATURE_SUBSURFACE)
KERNEL_STRUCT_MEMBER(subsurface, packed_float3, albedo, KERNEL_FEATURE_SUBSURFACE)
KERNEL_STRUCT_MEMBER(subsurface, packed_float3, radius, KERNEL_FEATURE_SUBSURFACE)
KERNEL_STRUCT_MEMBER(subsurface, float, anisotropy, KERNEL_FEATURE_SUBSURFACE)
KERNEL_STRUCT_MEMBER(subsurface, float3, Ng, KERNEL_FEATURE_SUBSURFACE)
KERNEL_STRUCT_MEMBER(subsurface, packed_float3, Ng, KERNEL_FEATURE_SUBSURFACE)
KERNEL_STRUCT_END(subsurface)
/********************************** Volume Stack ******************************/

View File

@@ -18,6 +18,14 @@
CCL_NAMESPACE_BEGIN
/* Volumetric read/write lambda functions - default implementations */
#ifndef VOLUME_READ_LAMBDA
# define VOLUME_READ_LAMBDA(function_call) \
auto volume_read_lambda_pass = [=](const int i) { return function_call; };
# define VOLUME_WRITE_LAMBDA(function_call) \
auto volume_write_lambda_pass = [=](const int i, VolumeStack entry) { function_call; };
#endif
/* Volume Stack
*
* This is an array of object/shared ID's that the current segment of the path
@@ -88,26 +96,18 @@ ccl_device void volume_stack_enter_exit(KernelGlobals kg,
IntegratorState state,
ccl_private const ShaderData *sd)
{
volume_stack_enter_exit(
kg,
sd,
[=](const int i) { return integrator_state_read_volume_stack(state, i); },
[=](const int i, const VolumeStack entry) {
integrator_state_write_volume_stack(state, i, entry);
});
VOLUME_READ_LAMBDA(integrator_state_read_volume_stack(state, i))
VOLUME_WRITE_LAMBDA(integrator_state_write_volume_stack(state, i, entry))
volume_stack_enter_exit(kg, sd, volume_read_lambda_pass, volume_write_lambda_pass);
}
ccl_device void shadow_volume_stack_enter_exit(KernelGlobals kg,
IntegratorShadowState state,
ccl_private const ShaderData *sd)
{
volume_stack_enter_exit(
kg,
sd,
[=](const int i) { return integrator_state_read_shadow_volume_stack(state, i); },
[=](const int i, const VolumeStack entry) {
integrator_state_write_shadow_volume_stack(state, i, entry);
});
VOLUME_READ_LAMBDA(integrator_state_read_shadow_volume_stack(state, i))
VOLUME_WRITE_LAMBDA(integrator_state_write_shadow_volume_stack(state, i, entry))
volume_stack_enter_exit(kg, sd, volume_read_lambda_pass, volume_write_lambda_pass);
}
/* Clean stack after the last bounce.

View File

@@ -676,19 +676,7 @@ ccl_device_forceinline void triangle_light_sample(KernelGlobals kg,
ls->D = z * B + safe_sqrtf(1.0f - z * z) * safe_normalize(C_ - dot(C_, B) * B);
/* calculate intersection with the planar triangle */
if (!ray_triangle_intersect(P,
ls->D,
FLT_MAX,
#if defined(__KERNEL_SSE2__) && defined(__KERNEL_SSE__)
(ssef *)V,
#else
V[0],
V[1],
V[2],
#endif
&ls->u,
&ls->v,
&ls->t)) {
if (!ray_triangle_intersect(P, ls->D, FLT_MAX, V[0], V[1], V[2], &ls->u, &ls->v, &ls->t)) {
ls->pdf = 0.0f;
return;
}

View File

@@ -22,6 +22,7 @@
#include "kernel/light/light.h"
#include "kernel/sample/mapping.h"
#include "kernel/sample/mis.h"
CCL_NAMESPACE_BEGIN
@@ -268,4 +269,36 @@ ccl_device_inline void light_sample_to_volume_shadow_ray(
shadow_ray_setup(sd, ls, P, ray);
}
ccl_device_inline float light_sample_mis_weight_forward(KernelGlobals kg,
const float forward_pdf,
const float nee_pdf)
{
#ifdef WITH_CYCLES_DEBUG
if (kernel_data.integrator.direct_light_sampling_type == DIRECT_LIGHT_SAMPLING_FORWARD) {
return 1.0f;
}
else if (kernel_data.integrator.direct_light_sampling_type == DIRECT_LIGHT_SAMPLING_NEE) {
return 0.0f;
}
else
#endif
return power_heuristic(forward_pdf, nee_pdf);
}
ccl_device_inline float light_sample_mis_weight_nee(KernelGlobals kg,
const float nee_pdf,
const float forward_pdf)
{
#ifdef WITH_CYCLES_DEBUG
if (kernel_data.integrator.direct_light_sampling_type == DIRECT_LIGHT_SAMPLING_FORWARD) {
return 0.0f;
}
else if (kernel_data.integrator.direct_light_sampling_type == DIRECT_LIGHT_SAMPLING_NEE) {
return 1.0f;
}
else
#endif
return power_heuristic(nee_pdf, forward_pdf);
}
CCL_NAMESPACE_END

View File

@@ -55,7 +55,7 @@ if(APPLE)
# Disable allocation warning on macOS prior to 10.14: the OSLRenderServices
# contains member which is 64 bytes aligned (cache inside of OIIO's
# unordered_map_concurrent). This is not something what the SDK supportsm, but
# since we take care of allocations ourselves is is OK to ignore the
# since we take care of allocations ourselves is OK to ignore the
# diagnostic message.
string(APPEND CMAKE_CXX_FLAGS " -faligned-allocation")
endif()

Some files were not shown because too many files have changed in this diff Show More