Compare commits
32
Commits
| Author | SHA1 | Date | |
|---|---|---|---|
|
|
f4db549061 | ||
|
|
4debd96adc | ||
|
|
30a4a95b7e | ||
|
|
a38aa58c0e | ||
|
|
e126a20df7 | ||
|
|
52cf2ee418 | ||
|
|
83a1e7e74b | ||
|
|
6069b7cfc0 | ||
|
|
41cc6d78db | ||
|
|
8561a817ea | ||
|
|
cfaf643d0e | ||
|
|
c505c4ca5f | ||
|
|
318f0e488e | ||
|
|
b22c9b7c84 | ||
|
|
bd491d025e | ||
|
|
6140c72aa6 | ||
|
|
7fcdc3f71c | ||
|
|
22feb208c2 | ||
|
|
ab65bbad4b | ||
|
|
64c66949e7 | ||
|
|
2fe18079a3 | ||
|
|
e0a66c7bc6 | ||
|
|
e4e5088ed6 | ||
|
|
aea281dcad | ||
|
|
b5aba16147 | ||
|
|
4ba1678561 | ||
|
|
8417f8eb04 | ||
|
|
55029e388f | ||
|
|
89c033bda9 | ||
|
|
0b57ae6b03 | ||
|
|
6588caafa6 | ||
|
|
b11fb97626 |
@@ -12,13 +12,18 @@ on:
|
||||
# - cron: "0 0 * * 0"
|
||||
push:
|
||||
branches:
|
||||
# - master
|
||||
- main
|
||||
- '*analysis*'
|
||||
# Allow manual kicking off of the workflow from github.com
|
||||
workflow_dispatch:
|
||||
|
||||
permissions: read-all
|
||||
|
||||
# Allow subsequent pushes to the same PR or REF to cancel any previous jobs.
|
||||
concurrency:
|
||||
group: ${{ github.workflow }}-${{ github.event.pull_request.number || github.ref }}
|
||||
cancel-in-progress: true
|
||||
|
||||
|
||||
jobs:
|
||||
|
||||
@@ -30,16 +35,18 @@ jobs:
|
||||
fail-fast: false
|
||||
matrix:
|
||||
include:
|
||||
- desc: sonary gcc9/C++14 llvm13 py3.9 boost1.76 exr3.1 oiio2.3 avx2
|
||||
- desc: sonar gcc9/C++14 llvm13 py3.9 boost1.76 exr3.1 oiio2.3 avx2
|
||||
nametag: static-analysis-sonar
|
||||
os: ubuntu-latest
|
||||
os: ubuntu-20.04-8c-32g-300h
|
||||
# os: ubuntu-latest
|
||||
vfxyear: 2022
|
||||
vfxsuffix: -clang13
|
||||
cxx_std: 14
|
||||
openimageio_ver: master
|
||||
python_ver: 3.9
|
||||
pybind11_ver: v2.9.0
|
||||
simd: avx2,f16c
|
||||
batched: b8_AVX2
|
||||
batched: b8_AVX2,b16_AVX512
|
||||
coverage: 1
|
||||
# skip_tests: 1
|
||||
sonar: 1
|
||||
@@ -53,7 +60,7 @@ jobs:
|
||||
TESTRENDER_AA=1
|
||||
OSL_TESTSUITE_SKIP_DIFF=1
|
||||
CTEST_TEST_TIMEOUT=1200
|
||||
OSL_CMAKE_FLAGS="-DOSL_TEST_BIG_TIMEOUT=1200"
|
||||
OSL_CMAKE_FLAGS="-DOSL_TEST_BIG_TIMEOUT=1200 -DUSE_QT=0"
|
||||
CTEST_EXCLUSIONS="broken|noise-reg.regress|noise-gabor-reg.regress"
|
||||
|
||||
runs-on: ${{ matrix.os }}
|
||||
@@ -65,10 +72,12 @@ jobs:
|
||||
CMAKE_CXX_STANDARD: ${{matrix.cxx_std}}
|
||||
FMT_VERSION: ${{matrix.fmt_ver}}
|
||||
OPENEXR_VERSION: ${{matrix.openexr_ver}}
|
||||
OPENIMAGEIO_VERSION: ${{matrix.openimageio_ver}}
|
||||
PYBIND11_VERSION: ${{matrix.pybind11_ver}}
|
||||
PYTHON_VERSION: ${{matrix.python_ver}}
|
||||
USE_BATCHED: ${{matrix.batched}}
|
||||
USE_SIMD: ${{matrix.simd}}
|
||||
PARALLEL: 24
|
||||
# DEBUG_CI: 1
|
||||
steps:
|
||||
# We would like to use harden-runner, but it flags too many false
|
||||
@@ -95,6 +104,13 @@ jobs:
|
||||
run: |
|
||||
${{matrix.setenvs}}
|
||||
src/build-scripts/ci-startup.bash
|
||||
- name: Remove existing OIIO
|
||||
if: matrix.openimageio_ver != ''
|
||||
run: |
|
||||
sudo rm -rf /usr/local/include/OpenImageIO
|
||||
sudo rm -rf /usr/local/lib*/cmake/OpenImageIO
|
||||
sudo rm -rf /usr/local/lib*/libOpenImageIO*
|
||||
sudo rm -rf /usr/local/lib*/python3.9/site-packages/OpenImageIO*
|
||||
- name: Dependencies
|
||||
run: |
|
||||
${{matrix.depcmds}}
|
||||
|
||||
+89
-55
@@ -25,6 +25,11 @@ on:
|
||||
|
||||
permissions: read-all
|
||||
|
||||
# Allow subsequent pushes to the same PR or REF to cancel any previous jobs.
|
||||
concurrency:
|
||||
group: ${{ github.workflow }}-${{ github.event.pull_request.number || github.ref }}
|
||||
cancel-in-progress: true
|
||||
|
||||
|
||||
jobs:
|
||||
|
||||
@@ -40,7 +45,7 @@ jobs:
|
||||
vfxyear: 2019
|
||||
vfxsuffix: -clang9
|
||||
cxx_std: 14
|
||||
openimageio_ver: v2.2.17.0
|
||||
openimageio_ver: v2.3.17.0
|
||||
python_ver: 2.7
|
||||
pybind11_ver: v2.4.2
|
||||
- desc: clang9/C++14 llvm9 oiio-release boost1.66 avx2 exr2.3 py2.7
|
||||
@@ -72,7 +77,7 @@ jobs:
|
||||
os: ubuntu-latest
|
||||
vfxyear: 2020
|
||||
cxx_std: 14
|
||||
openimageio_ver: v2.2.17.0
|
||||
openimageio_ver: v2.3.17.0
|
||||
python_ver: 3.7
|
||||
pybind11_ver: v2.5.0
|
||||
simd: sse4.2
|
||||
@@ -83,7 +88,7 @@ jobs:
|
||||
vfxyear: 2021
|
||||
vfxsuffix: -clang11
|
||||
cxx_std: 17
|
||||
openimageio_ver: dev-2.2
|
||||
openimageio_ver: dev-2.3
|
||||
python_ver: 3.7
|
||||
pybind11_ver: v2.7.0
|
||||
simd: avx2,f16c
|
||||
@@ -149,12 +154,72 @@ jobs:
|
||||
vfxyear: 2019
|
||||
vfxsuffix: -clang9
|
||||
cxx_std: 14
|
||||
openimageio_ver: v2.2.17.0
|
||||
openimageio_ver: v2.3.17.0
|
||||
python_ver: 2.7
|
||||
pybind11_ver: v2.6.2
|
||||
simd: 0
|
||||
setenvs: export PUGIXML_VERSION=v1.8
|
||||
|
||||
# Address and leak sanitizers
|
||||
- desc: sanitizers
|
||||
nametag: sanitizer
|
||||
os: ubuntu-latest
|
||||
vfxyear: 2022
|
||||
vfxsuffix: -clang13
|
||||
cc_compiler: clang
|
||||
cxx_compiler: clang++
|
||||
cxx_std: 17
|
||||
openimageio_ver: master
|
||||
python_ver: 3.9
|
||||
pybind11_ver: v2.9.0
|
||||
simd: avx2,f16c
|
||||
batched: b8_AVX2
|
||||
setenvs: export OSL_CMAKE_FLAGS="-DSANITIZE=address -DOSL_TEST_BIG_TIMEOUT=1200"
|
||||
CMAKE_BUILD_TYPE=Debug
|
||||
CTEST_TEST_TIMEOUT=1200
|
||||
CTEST_EXCLUSIONS="broken|noise-reg.regress|noise-gabor-reg.regress|example-.*deformer|spline-reg.regress.batched|osl-imageio"
|
||||
OSL_OPTIONS=llvm_debugging_symbols=1
|
||||
TESTRENDER_AA=1
|
||||
OSL_TESTSUITE_SKIP_DIFF=1
|
||||
# Sanitizer notes:
|
||||
# - There are a handful of tests we disable entirely with
|
||||
# CTEST_EXCLUSIONS because they are irrelevant or just cannot be
|
||||
# made to be clean (and that's ok).
|
||||
# - Tests with sanitizers run slow, so we increase timeouts, and
|
||||
# also set AA to 1 for render tests to speed them up.
|
||||
# OSL_TESTSUITE_SKIP_DIFF prevents making the AA=1 image result
|
||||
# fail, but that's ok because every other CI test is checking
|
||||
# the results; here we have no reason to believe results will
|
||||
# differ from other tests, we just need to run with sanitizers
|
||||
# turned on.
|
||||
# - In testing.cmake, we rig texture tests to disable the leak
|
||||
# sanitizer, because we can't eliminate some completely benign
|
||||
# leaks from the texture system, and because they are called
|
||||
# from JITed code, we don't get a clean stack trace of things
|
||||
# that happen in OIIO and therefore we just can't enumerate the
|
||||
# allowed exceptions. So we trust that OIIO itself tests the
|
||||
# texture system adequately and is able to ignore the benign
|
||||
# "leaks."
|
||||
|
||||
# Test formatting. This test entry doesn't do a full build, it
|
||||
# just runs clang-format on everything, and passes if nothing is
|
||||
# misformatted. Upon failure, the build artifact will be the full
|
||||
# source code with the formatting fixed (diffs will also appear in
|
||||
# the console output).
|
||||
- desc: "clang-format"
|
||||
nametag: clang-format
|
||||
os: ubuntu-latest
|
||||
vfxyear: 2022
|
||||
vfxsuffix: -clang14
|
||||
cxx_std: 17
|
||||
openimageio_ver: release
|
||||
python_ver: 3.9
|
||||
simd: avx2,f16c
|
||||
batched: b8_AVX2
|
||||
skip_tests: 1
|
||||
setenvs: export BUILDTARGET=clang-format
|
||||
OPENIMAGEIO_CMAKE_FLAGS=-DUSE_PYTHON=0
|
||||
|
||||
runs-on: ${{matrix.os}}
|
||||
container:
|
||||
image: aswftesting/ci-osl:${{matrix.vfxyear}}${{matrix.vfxsuffix}}
|
||||
@@ -366,20 +431,25 @@ jobs:
|
||||
- desc: MacOS-11
|
||||
os: macos-11
|
||||
nametag: macos11-py39
|
||||
cxx_std: 17
|
||||
python_ver: 3.9
|
||||
aclang: 13
|
||||
openimageio_ver: release
|
||||
- desc: MacOS-12
|
||||
os: macos-12
|
||||
nametag: macos11-p310
|
||||
cc_compiler: /usr/local/opt/llvm@14/bin/clang
|
||||
cxx_compiler: /usr/local/opt/llvm@14/bin/clang++
|
||||
cxx_std: 17
|
||||
python_ver: "3.10"
|
||||
aclang: 13
|
||||
setenvs: export LLVMBREWVER="@14" DO_BREW_UPDATE=1
|
||||
- desc: MacOS-12
|
||||
os: macos-12
|
||||
nametag: macos11-p310
|
||||
cc_compiler: /usr/local/opt/llvm@15/bin/clang
|
||||
cxx_compiler: /usr/local/opt/llvm@15/bin/clang++
|
||||
cxx_std: 17
|
||||
python_ver: "3.10"
|
||||
aclang: 13
|
||||
setenvs: export LLVMBREWVER="@15" DO_BREW_UPDATE=1
|
||||
runs-on: ${{matrix.os}}
|
||||
env:
|
||||
CC: clang
|
||||
CXX: clang++
|
||||
CXX: ${{matrix.cxx_compiler}}
|
||||
CC: ${{matrix.cc_compiler}}
|
||||
CMAKE_CXX_STANDARD: ${{matrix.cxx_std}}
|
||||
PYTHON_VERSION: ${{matrix.python_ver}}
|
||||
OPENIMAGEIO_VERSION: ${{matrix.openimageio_ver}}
|
||||
@@ -393,10 +463,12 @@ jobs:
|
||||
uses: actions/cache@c3f1317a9e7b1ef106c153ac8c0f00fed3ddbc0d # v3.0.4
|
||||
with:
|
||||
path: /Users/runner/.ccache
|
||||
key: ${{github.job}}-${{matrix.os}}-${{steps.ccache_cache_keys.outputs.date}}
|
||||
restore-keys: ${{github.job}}-${{matrix.os}}-
|
||||
key: ${{github.job}}-${{matrix.nametag}}-${{steps.ccache_cache_keys.outputs.date}}
|
||||
restore-keys: ${{github.job}}-${{matrix.nametag}}-
|
||||
- name: Build setup
|
||||
run: src/build-scripts/ci-startup.bash
|
||||
run: |
|
||||
${{matrix.setenvs}}
|
||||
src/build-scripts/ci-startup.bash
|
||||
- name: Dependencies
|
||||
run: |
|
||||
source src/build-scripts/install_homebrew_deps.bash
|
||||
@@ -416,46 +488,8 @@ jobs:
|
||||
- uses: actions/upload-artifact@3cea5372237819ed00197afe530f5a7ea3e805c8 # v3.1.0
|
||||
if: failure()
|
||||
with:
|
||||
name: osl-${{github.job}}
|
||||
name: osl-${{github.job}}-${{matrix.nametag}}
|
||||
path: |
|
||||
build/*.cmake
|
||||
build/CMake*
|
||||
build/testsuite/*/*.*
|
||||
|
||||
|
||||
clang-format:
|
||||
# Test formatting. This test entry doesn't do a full build, it just runs
|
||||
# clang-format on everything, and passes if nothing is misformatted.
|
||||
# Upon failure, the build artifact will be the full source code with the
|
||||
# formatting fixed (diffs will also appear in the console output).
|
||||
name: "clang-format verification"
|
||||
runs-on: ubuntu-20.04
|
||||
container:
|
||||
image: aswf/ci-osl:2022-clang14
|
||||
env:
|
||||
BUILDTARGET: clang-format
|
||||
CMAKE_CXX_STANDARD: 17
|
||||
PYTHON_VERSION: 3.9
|
||||
steps:
|
||||
# We would like to use harden-runner, but it flags too many false
|
||||
# positives, every time we download a dependency. We should use it only
|
||||
# on CI runs where we are producing artifacts that users might rely on.
|
||||
# - name: Harden Runner
|
||||
# uses: step-security/harden-runner@248ae51c2e8cc9622ecf50685c8bf7150c6e8813 # v1.4.3
|
||||
# with:
|
||||
# egress-policy: audit # TODO: change to 'egress-policy: block' after couple of runs
|
||||
- uses: actions/checkout@2541b1294d2704b0964813337f33b291d3f8596b # v3.0.2
|
||||
- name: Build setup
|
||||
run: |
|
||||
src/build-scripts/ci-startup.bash
|
||||
- name: Dependencies
|
||||
run: |
|
||||
src/build-scripts/gh-installdeps.bash
|
||||
- name: Build
|
||||
run: |
|
||||
src/build-scripts/ci-build.bash
|
||||
- uses: actions/upload-artifact@3cea5372237819ed00197afe530f5a7ea3e805c8 # v3.1.0
|
||||
if: failure()
|
||||
with:
|
||||
name: osl-${{github.job}}
|
||||
path: src/*/*.{cpp,h}
|
||||
|
||||
@@ -9,4 +9,5 @@ build
|
||||
*.pyc
|
||||
*.log
|
||||
.vscode
|
||||
.envrc
|
||||
|
||||
|
||||
+22
-7
@@ -6,19 +6,23 @@ Release 1.13 -- ?? 2023 ?? (compared to 1.12)
|
||||
Dependency and standards requirements changes:
|
||||
* OptiX 6.0 support has been removed. For GPU rendering with OptiX, a minimum
|
||||
of OptiX 7.0 is required.
|
||||
* The minimum version OpenImageIO has been raised from 2.2 to 2.3. #1591
|
||||
(1.13.0.3)
|
||||
|
||||
OSL Language and oslc compiler:
|
||||
* Bug fix: shader params whose default values involve "init ops" could be
|
||||
incorrectly initialized if their values came from instance parameters set to
|
||||
zero. #1578 (1.13.0.1)
|
||||
zero. #1578 (1.13.0.2)
|
||||
|
||||
OSL Standard library:
|
||||
|
||||
API changes, new options, new ShadingSystem features (for renderer writers):
|
||||
* A new `shadingStateUniform` pointer field was added to ShaderGlobals for
|
||||
planned future expansion and improved GPU support. #1585 (1.13.1.0)
|
||||
|
||||
SIMD batched shading mode:
|
||||
* Fix some array overruns (asymptomatic, but still potentially buggy) in
|
||||
common_ancestor_between. #1577 (1.13.0.1)
|
||||
common_ancestor_between. #1577 (1.13.0.2)
|
||||
|
||||
OptiX rendering:
|
||||
* OptiX 6.0 support has been removed. For GPU rendering with OptiX, a minimum
|
||||
@@ -34,26 +38,35 @@ OptiX rendering:
|
||||
Performance improvements:
|
||||
|
||||
Bug fixes and other improvements (internals):
|
||||
* Fix memory leak in oslc ASTvariable_declaration. #1576 (1.13.0.1)
|
||||
* Fix memory leak in oslc ASTvariable_declaration. #1576 (1.13.0.2)
|
||||
* Fix memory leaks in LPE code in error conditions. #1593 #1594 (1.13.0.3)
|
||||
|
||||
Internals/developer concerns:
|
||||
* platform.h:
|
||||
- New `OSL_ALLOCA` macro to allocate on the stack. #1589 (1.13.0.3)
|
||||
|
||||
Build & test system improvements:
|
||||
* CMake build system and scripts:
|
||||
- The version number is now a cache variable, and so can be overridden at
|
||||
build time using `-DOSL_VERSION=...`. Use with extreme caution! #1579
|
||||
(1.13.0.1)
|
||||
(1.13.0.2)
|
||||
* Dependency version support:
|
||||
- Fix some issues that came up with the new icx 2022.0 compiler. #1601
|
||||
(1.13.1.0)
|
||||
* Testing and Continuous integration (CI) systems:
|
||||
- Dynamic analysis using address and leak sanitizers. #1581 (1.13.0.3)
|
||||
- `testshade --help` now prints TextureSystem options and all hardware
|
||||
information. #1584 (1.13.0.3)
|
||||
- Improvements to SonarCloud and coverage analysis. #1596 (1.13.0.3)
|
||||
* Platform support:
|
||||
|
||||
Documentation:
|
||||
* `doc/RELEASING.md` documents our release process and versioning policies.
|
||||
#1572 (1.13.0.1)
|
||||
#1572 (1.13.0.2)
|
||||
|
||||
|
||||
Release 1.12 -- beta 1 Sep 2022 (compared to 1.11)
|
||||
--------------------------------------------------
|
||||
Release 1.12 -- 1 Oct 2022 (compared to 1.11)
|
||||
----------------------------------------------
|
||||
Big Deal Changes:
|
||||
* **Batch shading**: A fully operational "batch shading" support when using
|
||||
CPUs supporting the Intel AVX2 or AVX512 instruction set architectures,
|
||||
@@ -295,6 +308,7 @@ Build & test system improvements:
|
||||
#1271 (1.12.0.1)
|
||||
- Build against LLVM 11 #1274 (1.12.0.1), LLVM 12 #1351 (1.12.1.0)
|
||||
#1412 (1.12.3), LLVM 13 #1420 (1.12.3), LLVM 14 #1492 (1.12.4.6)
|
||||
LLVM 15 #1592 (1.12.6.2)
|
||||
- Fix build break against recent OIIO master change where m_mutex field
|
||||
was removed from ImageInput. #1281 (1.12.0.1/1.11.9)
|
||||
- Work to ensure that OIIO will build correctly against the upcoming
|
||||
@@ -346,6 +360,7 @@ Build & test system improvements:
|
||||
#1501 (1.12.5.0)
|
||||
- Overhaul of GHA ci.yml to use strategy matrix. #1474 (1.12.4.4)
|
||||
- Test against clang 14 and fix new warnings. #1498 (1.12.5.0)
|
||||
- Test against clang 15 and fix new warnings. #1586 (1.12.6.2)
|
||||
- A scorecards workflow guards against a variety of security issues of
|
||||
the CI system itself. #1529 (1.12.6.0)
|
||||
- Enable SonarCloud static analysis. #1551 #1556 #1563 (1.12.6.0)
|
||||
|
||||
+1
-1
@@ -4,7 +4,7 @@
|
||||
|
||||
cmake_minimum_required (VERSION 3.12)
|
||||
|
||||
set (OSL_VERSION "1.13.0.2" CACHE STRING "Version")
|
||||
set (OSL_VERSION "1.13.1.0" CACHE STRING "Version")
|
||||
project (OSL VERSION ${OSL_VERSION}
|
||||
LANGUAGES CXX C
|
||||
HOMEPAGE_URL "https://github.com/AcademySoftwareFoundation/OpenShadingLanguage")
|
||||
|
||||
+13
-9
@@ -173,9 +173,16 @@ it really only applies to the particular file in which it appears.
|
||||
|
||||
#### Formatting
|
||||
|
||||
NEVER alter somebody else's code to reformat just because you found
|
||||
something that violates the rules. Let the group/author/leader know, and
|
||||
resist the temptation to change it yourself.
|
||||
We use clang-format to enforce formatting rules for the vast majority of the
|
||||
code base. It is automatically run during CI and is a CI failure if the code
|
||||
does not conform. There are, however, parts of the code base that are not
|
||||
clang-formatted, either because we have not yet put them under the
|
||||
clang-format regime, or because individual modules were deemed to be less
|
||||
clear with clang-format than with hand formatting.
|
||||
|
||||
For non-clang-format regions of code, NEVER alter somebody else's code to
|
||||
reformat just because you found something that violates the rules. Let the
|
||||
group/author/leader know, and resist the temptation to change it yourself.
|
||||
|
||||
Each line of text in your code should be at most 80 characters long.
|
||||
Exceptions are allowed for those rare cases where letting a line be longer
|
||||
@@ -217,13 +224,10 @@ Function calls should have a space between the function name and the opening
|
||||
parenthesis, NO space inside the parenthesis, except for a single blank
|
||||
space between each argument. For example:
|
||||
|
||||
x = foo (a, b); // Yes, this is always ok
|
||||
x = foo(a, b); // Yes
|
||||
|
||||
x = foo ( a, b ); // No
|
||||
x = foo (a,b); // No
|
||||
x = foo(a, b); // No
|
||||
x = foo(a); // Occasionally, this just looks better, when the function name is short,
|
||||
// and there's just one very short argument. What can I say, I do it too.
|
||||
|
||||
Function declarations: function names should begin at column 0 for a full
|
||||
function definition. (It's ok to violate this rule for very short inline
|
||||
@@ -233,7 +237,7 @@ functions within class definitions.)
|
||||
Here is a short code fragment that shows some of these rules in action:
|
||||
|
||||
static int
|
||||
function (int a, int b)
|
||||
function(int a, int b)
|
||||
{
|
||||
int x = a + b;
|
||||
if (a == 0 || b == 0) {
|
||||
@@ -244,7 +248,7 @@ Here is a short code fragment that shows some of these rules in action:
|
||||
}
|
||||
for (int i = 0; i < 3; ++i) {
|
||||
x += a * i;
|
||||
x *= foo (i); // function call
|
||||
x *= foo(i); // function call
|
||||
}
|
||||
return x;
|
||||
}
|
||||
|
||||
+14
-19
@@ -14,14 +14,14 @@ NEW or CHANGED dependencies since the last major release are **bold**.
|
||||
|
||||
* Build system: [CMake](https://cmake.org/) 3.12 or newer (tested through 3.24)
|
||||
|
||||
* A suitable **C++14 or C++17** compiler to build OSL itself, which may be any of:
|
||||
- **GCC 6.1 or newer** (tested through gcc 12.1)
|
||||
- Clang 3.4 or newer (tested through clang 14)
|
||||
- Microsoft Visual Studio **2017 or newer**
|
||||
- Intel C++ compiler icc **version 17 or newer** or LLVM-based icx compiler
|
||||
**version 2022 or newer**.
|
||||
* A suitable C++14 or C++17 compiler to build OSL itself, which may be any of:
|
||||
- GCC 6.1 or newer (tested through gcc 12.1)
|
||||
- Clang 3.4 or newer (tested through clang 15)
|
||||
- Microsoft Visual Studio 2017 or newer
|
||||
- Intel C++ compiler icc version 17 or newer or LLVM-based icx compiler
|
||||
version 2022 or newer.
|
||||
|
||||
* **[OpenImageIO](http://openimageio.org) 2.2 or newer** (tested through 2.4)
|
||||
* **[OpenImageIO](http://openimageio.org) 2.3 or newer** (tested through 2.4)
|
||||
|
||||
OSL uses OIIO both for its texture mapping functionality as well as
|
||||
numerous utility classes. If you are integrating OSL into an existing
|
||||
@@ -40,8 +40,8 @@ NEW or CHANGED dependencies since the last major release are **bold**.
|
||||
$OpenImageIO_ROOT/lib to be in your LD_LIBRARY_PATH (or
|
||||
DYLD_LIBRARY_PATH on OS X).
|
||||
|
||||
* **[LLVM](http://www.llvm.org) 9, 10, 11, 12, 13, or 14**, including clang
|
||||
libraries.
|
||||
* [LLVM](http://www.llvm.org) 9, 10, 11, 12, 13, 14, or 15, including
|
||||
clang libraries.
|
||||
|
||||
* (optional) For GPU rendering on NVIDIA GPUs:
|
||||
* [OptiX](https://developer.nvidia.com/rtx/ray-tracing/optix) 7.0 or higher.
|
||||
@@ -97,19 +97,14 @@ Here are the steps to check out, build, and test the OSL distribution:
|
||||
full errors, not warnings.
|
||||
|
||||
4. After compilation, you'll end up with a full OSL distribution in
|
||||
dist/ARCH, where ARCH is the architecture you are building on, one of
|
||||
"linux", "linux64", "macosx", "windows", or "windows64".
|
||||
dist/
|
||||
|
||||
Note: The default is to make an optimized "release" build. If
|
||||
instead type 'make debug' at the top level, you will end up with
|
||||
a debug build (no optimization, full symbols) in "dist/ARCH.debug".
|
||||
|
||||
5. Add the "dist/ARCH/bin" to your $PATH, and "dist/ARCH/lib" to your
|
||||
$LD_LIBRAY_PATH (or $DYLD_LIBRARY_PATH on OS X), or copy the contents
|
||||
5. Add the "dist/bin" to your $PATH, and "dist/lib" to your
|
||||
$LD_LIBRAY_PATH (or $DYLD_LIBRARY_PATH on MacOS), or copy the contents
|
||||
of those files to appropriate directories. Public include files
|
||||
(those needed when building applications that incorporate OSL)
|
||||
can be found in "dist/ARCH/include", and documentation can be found
|
||||
in "dist/ARCH/share/doc".
|
||||
can be found in "dist/include", and documentation can be found
|
||||
in "dist/share/doc".
|
||||
|
||||
6. After building (and setting your library path), you can run the
|
||||
test suite with:
|
||||
|
||||
@@ -181,14 +181,14 @@ else ifeq (${platform}, macosx)
|
||||
MY_CMAKE_FLAGS += \
|
||||
-DCMAKE_C_COMPILER=gcc-9 -DCMAKE_CXX_COMPILER=g++-9
|
||||
USE_LIBCPLUSPLUS := 0
|
||||
else ifeq (${COMPILER},clang11)
|
||||
else ifeq (${COMPILER},clang14)
|
||||
MY_CMAKE_FLAGS += \
|
||||
-DCMAKE_C_COMPILER=/usr/local/opt/llvm@11/bin/clang \
|
||||
-DCMAKE_CXX_COMPILER=/usr/local/opt/llvm@11/bin/clang++
|
||||
else ifeq (${COMPILER},clang12)
|
||||
-DCMAKE_C_COMPILER=/usr/local/opt/llvm@14/bin/clang \
|
||||
-DCMAKE_CXX_COMPILER=/usr/local/opt/llvm@14/bin/clang++
|
||||
else ifeq (${COMPILER},clang15)
|
||||
MY_CMAKE_FLAGS += \
|
||||
-DCMAKE_C_COMPILER=/usr/local/opt/llvm@12/bin/clang \
|
||||
-DCMAKE_CXX_COMPILER=/usr/local/opt/llvm@12/bin/clang++
|
||||
-DCMAKE_C_COMPILER=/usr/local/opt/llvm@15/bin/clang \
|
||||
-DCMAKE_CXX_COMPILER=/usr/local/opt/llvm@15/bin/clang++
|
||||
else ifeq (${COMPILER},clang)
|
||||
MY_CMAKE_FLAGS += -DCMAKE_C_COMPILER=clang -DCMAKE_CXX_COMPILER=clang++
|
||||
else ifeq (${COMPILER},gcc)
|
||||
|
||||
@@ -18,6 +18,7 @@ sonar.links.issue=https://github.com/AcademySoftwareFoundation/OpenShadingLangua
|
||||
# Source properties
|
||||
sonar.sources=src
|
||||
sonar.sourceEncoding=UTF-8
|
||||
sonar.exclusions=src/include/OSL/Imath/*.h,src/build-scripts/*
|
||||
|
||||
# C/C++ analyzer properties
|
||||
sonar.cfamily.build-wrapper-output=build/bw_output
|
||||
|
||||
@@ -32,9 +32,10 @@ cmake .. -G "$CMAKE_GENERATOR" \
|
||||
mkdir cmake-save || /bin/true
|
||||
cp -r CMake* *.cmake cmake-save
|
||||
|
||||
: ${BUILDTARGET:=install}
|
||||
if [[ "$BUILDTARGET" != "none" ]] ; then
|
||||
echo "Parallel build " ${CMAKE_BUILD_PARALLEL_LEVEL}
|
||||
time ${OSL_CMAKE_BUILD_WRAPPER} cmake --build . --target ${BUILDTARGET:=install} --config ${CMAKE_BUILD_TYPE}
|
||||
echo "Parallel build ${CMAKE_BUILD_PARALLEL_LEVEL} of target ${BUILDTARGET}"
|
||||
time ${OSL_CMAKE_BUILD_WRAPPER} cmake --build . --target ${BUILDTARGET} --config ${CMAKE_BUILD_TYPE}
|
||||
fi
|
||||
popd
|
||||
|
||||
|
||||
@@ -23,7 +23,7 @@ export DYLD_LIBRARY_PATH=$OSL_ROOT/lib:$DYLD_LIBRARY_PATH
|
||||
export LD_LIBRARY_PATH=$OSL_ROOT/lib:$LD_LIBRARY_PATH
|
||||
export OIIO_LIBRARY_PATH=$OSL_ROOT/lib:${OIIO_LIBRARY_PATH}
|
||||
export LSAN_OPTIONS=suppressions=$PWD/src/build-scripts/nosanitize.txt
|
||||
export ASAN_OPTIONS=print_suppressions=0
|
||||
export ASAN_OPTIONS=print_suppressions=0:detect_odr_violation=1
|
||||
|
||||
export USE_PYTHON=${USE_PYTHON:=1}
|
||||
export PYTHON_VERSION=${PYTHON_VERSION:="2.7"}
|
||||
@@ -61,6 +61,7 @@ echo "uname -a: " `uname -a`
|
||||
echo "uname -m: " `uname -m`
|
||||
echo "uname -s: " `uname -s`
|
||||
echo "uname -n: " `uname -n`
|
||||
echo "nproc " `nproc`
|
||||
pwd
|
||||
ls
|
||||
env | sort
|
||||
|
||||
@@ -42,8 +42,12 @@ if [[ "$ASWF_ORG" != "" ]] ; then
|
||||
|
||||
if [[ "$CXX" == "icpc" || "$CC" == "icc" || "$USE_ICC" != "" || "$CXX" == "icpx" || "$CC" == "icx" || "$USE_ICX" != "" ]] ; then
|
||||
sudo cp src/build-scripts/oneAPI.repo /etc/yum.repos.d
|
||||
sudo yum install -y intel-oneapi-compiler-dpcpp-cpp-and-cpp-classic
|
||||
set +e; source /opt/intel/oneapi/setvars.sh; set -e
|
||||
sudo yum install -y intel-oneapi-compiler-dpcpp-cpp-and-cpp-classic-2022.1.0-3768
|
||||
# Because multiple (possibly newer) versions of oneAPI may be installed,
|
||||
# use a config file to specify compiler and tbb versions
|
||||
# NOTE: oneAPI components have independent version numbering.
|
||||
set +e; source /opt/intel/oneapi/setvars.sh --config oneapi_2022.1.0.cfg; set -e
|
||||
|
||||
if [[ "$CXX" == "icpc" || "$CC" == "icc" || "$USE_ICC" != "" ]] ; then
|
||||
echo "Verifying installation of Intel(r) C++ Compiler:"
|
||||
icpc --version
|
||||
@@ -106,8 +110,11 @@ else
|
||||
sudo apt-key add GPG-PUB-KEY-INTEL-SW-PRODUCTS-2023.PUB
|
||||
echo "deb https://apt.repos.intel.com/oneapi all main" | sudo tee /etc/apt/sources.list.d/oneAPI.list
|
||||
time sudo apt-get update
|
||||
time sudo apt-get install -y intel-oneapi-compiler-dpcpp-cpp-and-cpp-classic
|
||||
set +e; source /opt/intel/oneapi/setvars.sh; set -e
|
||||
time sudo apt-get install -y intel-oneapi-compiler-dpcpp-cpp-and-cpp-classic=2022.1.0-3768
|
||||
# Because multiple (possibly newer) versions of oneAPI may be installed,
|
||||
# use a config file to specify compiler and tbb versions
|
||||
# NOTE: oneAPI components have independent version numbering.
|
||||
set +e; source /opt/intel/oneapi/setvars.sh --config oneapi_2022.1.0.cfg; set -e
|
||||
fi
|
||||
|
||||
source src/build-scripts/build_llvm.bash
|
||||
|
||||
@@ -20,7 +20,9 @@ if [[ `which brew` == "" ]] ; then
|
||||
fi
|
||||
|
||||
|
||||
#brew update >/dev/null
|
||||
if [[ "$DO_BREW_UPDATE" != "" ]] ; then
|
||||
brew update >/dev/null
|
||||
fi
|
||||
echo ""
|
||||
echo "Before my brew installs:"
|
||||
brew list --versions
|
||||
|
||||
@@ -1,5 +0,0 @@
|
||||
leak:Imf_2_2
|
||||
leak:Imf_2_3
|
||||
leak:Imf_2_4
|
||||
leak:libIlmImf
|
||||
leak:oslgram.y
|
||||
|
||||
@@ -0,0 +1,3 @@
|
||||
default=exclude
|
||||
compiler=2022.1.0
|
||||
tbb=2021.6.0
|
||||
@@ -392,22 +392,9 @@ if (SANITIZE AND (CMAKE_COMPILER_IS_GNUCC OR CMAKE_COMPILER_IS_CLANG))
|
||||
foreach (feature ${SANITIZE_FEATURE_LIST})
|
||||
message (STATUS " sanitize feature: ${feature}")
|
||||
add_compile_options (-fsanitize=${feature})
|
||||
set (CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} -fsanitize=${feature}")
|
||||
set (CMAKE_MODULE_LINKER_FLAGS "${CMAKE_MODULE_LINKER_FLAGS} -fsanitize=${feature}")
|
||||
set (CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -fsanitize=${feature}")
|
||||
add_link_options (-fsanitize=${feature})
|
||||
endforeach()
|
||||
add_compile_options (-g -fno-omit-frame-pointer)
|
||||
if (${CMAKE_SYSTEM_NAME} STREQUAL "Linux")
|
||||
set (SANITIZE_ON_LINUX 1)
|
||||
endif ()
|
||||
if (CMAKE_COMPILER_IS_GNUCC AND ${CMAKE_SYSTEM_NAME} STREQUAL "Linux")
|
||||
add_compile_options ("-fuse-ld=gold")
|
||||
set (CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} -fuse-ld=gold")
|
||||
set (CMAKE_MODULE_LINKER_FLAGS "${CMAKE_MODULE_LINKER_FLAGS} -fuse-ld=gold")
|
||||
set (CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -fuse-ld=gold")
|
||||
set (SANITIZE_LIBRARIES "asan;pthread")
|
||||
# set (SANITIZE_LIBRARIES "asan" "ubsan")
|
||||
endif()
|
||||
if (CMAKE_COMPILER_IS_GNUCC)
|
||||
# turn on glibcxx extra annotations to find vector writes past end
|
||||
add_definitions ("-D_GLIBCXX_SANITIZE_VECTOR=1")
|
||||
|
||||
@@ -105,7 +105,7 @@ endif ()
|
||||
|
||||
# OpenImageIO
|
||||
checked_find_package (OpenImageIO REQUIRED
|
||||
VERSION_MIN 2.2.6
|
||||
VERSION_MIN 2.3.17
|
||||
DEFINITIONS -DOIIO_HIDE_FORMAT=1)
|
||||
|
||||
checked_find_package (pugixml REQUIRED
|
||||
@@ -136,6 +136,14 @@ if (APPLE AND LLVM_VERSION VERSION_EQUAL 10.0.1 AND EXISTS "/usr/local/Cellar/ll
|
||||
" brew upgrade llvm \n"
|
||||
"${ColorReset}\n")
|
||||
endif ()
|
||||
if (LLVM_VERSION VERSION_GREATER_EQUAL 15.0
|
||||
AND ( (CMAKE_COMPILER_IS_APPLECLANG AND APPLECLANG_VERSION_STRING VERSION_LESS 15.0)
|
||||
OR (CMAKE_COMPILER_IS_CLANG AND CLANG_VERSION_STRING VERSION_LESS 15.0)))
|
||||
message (WARNING
|
||||
"${ColorYellow}"
|
||||
"If you are using LLVM 15 or higher, you should also use clang version "
|
||||
"15 or higher, or you may get build errors.${ColorReset}\n")
|
||||
endif ()
|
||||
|
||||
checked_find_package (partio)
|
||||
|
||||
|
||||
@@ -82,6 +82,11 @@ function ( EMBED_LLVM_BITCODE_IN_CPP src_list suffix output_name list_to_append_
|
||||
list (TRANSFORM OpenImageIO_INCLUDES PREPEND -I
|
||||
OUTPUT_VARIABLE ALL_OpenImageIO_INCLUDES)
|
||||
|
||||
if (${LLVM_VERSION} VERSION_GREATER_EQUAL 15.0)
|
||||
# Until we fully support opaque pointers, we need to disable
|
||||
# them when using LLVM 15.
|
||||
list (APPEND LLVM_COMPILE_FLAGS -Xclang -no-opaque-pointers)
|
||||
endif ()
|
||||
|
||||
# Command to turn the .cpp file into LLVM assembly language .s, into
|
||||
# LLVM bitcode .bc, then back into a C++ file with the bc embedded!
|
||||
|
||||
@@ -77,7 +77,7 @@ macro (setup_python_module)
|
||||
PRIVATE
|
||||
pybind11::pybind11
|
||||
Python::Python
|
||||
${lib_LIBS} ${SANITIZE_LIBRARIES})
|
||||
${lib_LIBS})
|
||||
if (APPLE)
|
||||
set_target_properties (${target_name} PROPERTIES LINK_FLAGS "-undefined dynamic_lookup")
|
||||
endif ()
|
||||
|
||||
+15
-4
@@ -35,6 +35,7 @@ macro (add_one_testsuite testname testsrcdir)
|
||||
cmake_parse_arguments (_tst "" "" "ENV;COMMAND" ${ARGN})
|
||||
set (testsuite "${CMAKE_SOURCE_DIR}/testsuite")
|
||||
set (testdir "${CMAKE_BINARY_DIR}/testsuite/${testname}")
|
||||
file (MAKE_DIRECTORY "${testdir}")
|
||||
if (NOT _tst_COMMAND)
|
||||
set (_tst_COMMAND ${Python_EXECUTABLE} "${testsuite}/runtest.py" ${testdir})
|
||||
if (MSVC_IDE)
|
||||
@@ -50,7 +51,17 @@ macro (add_one_testsuite testname testsrcdir)
|
||||
OSL_TESTSUITE_SRC=${testsrcdir}
|
||||
OSL_TESTSUITE_CUR=${testdir}
|
||||
)
|
||||
file (MAKE_DIRECTORY "${testdir}")
|
||||
|
||||
# We need to suppress the leak sanitizer for "texture" tests, because of a
|
||||
# minor leak inside OIIO (140 bytes! -- inconsequential) but which we
|
||||
# can't use the suppressions list becuase the leaks come from calls made
|
||||
# inside JITed code for which there is not a good stack trace. Any real
|
||||
# leaks from OIIO or its dependencies we assume will be caught by lsan on
|
||||
# the OIIO side, where there is no JIT to complicate things.
|
||||
if (${testname} MATCHES "^texture" OR ${testname} MATCHES "gettextureinfo")
|
||||
list (APPEND _tst_ENV ASAN_OPTIONS=detect_leaks=0:$ENV{ASAN_OPTIONS})
|
||||
endif ()
|
||||
|
||||
add_test ( NAME ${testname} COMMAND ${_tst_COMMAND} )
|
||||
# message ("Test -- env ${_tst_ENV} cmd ${_tst_COMMAND}")
|
||||
set_tests_properties (${testname} PROPERTIES ENVIRONMENT "${_tst_ENV}" )
|
||||
@@ -357,10 +368,10 @@ macro (osl_add_all_tests)
|
||||
endif ()
|
||||
|
||||
# Add tests that require the Python bindings if we built them.
|
||||
# We also exclude these tests if this is a sanitizer build on Linux,
|
||||
# because the Python interpreter itself won't be linked with the right asan
|
||||
# We also exclude these tests if this is a sanitizer build, because the
|
||||
# Python interpreter itself won't be linked with the right asan
|
||||
# libraries to run correctly.
|
||||
if (USE_PYTHON AND NOT SANITIZE_ON_LINUX)
|
||||
if (USE_PYTHON AND NOT SANITIZE)
|
||||
TESTSUITE ( python-oslquery )
|
||||
endif ()
|
||||
|
||||
|
||||
@@ -4611,7 +4611,7 @@ Example:
|
||||
\index{material closure functions}
|
||||
\index{functions!closures}
|
||||
|
||||
For \closurecolor functions, the return ``value'' is symbolic and is may
|
||||
For \closurecolor functions, the return ``value'' is symbolic and may
|
||||
be passed to an output variable or assigned to \Ci, to be evaluated at a
|
||||
later time convenient to the renderer in order to compute the exitant
|
||||
radiance in the direction {\cf -I}. But the shader itself cannot
|
||||
@@ -4882,7 +4882,7 @@ parameters control the contribution of each reflection/transmission lobe.
|
||||
|
||||
The closure may be vertically layered over a base BSDF for the surface beneath
|
||||
the dielectric layer. This is done using the layer() closure. By chaining
|
||||
multiple dielectric_bsdf closures you can describe a surface with multiple
|
||||
multiple {\cf dielectric_bsdf} closures you can describe a surface with multiple
|
||||
specular lobes. If transmission is enabled (transmission_tint $>$ 0.0) the
|
||||
closure may be layered over a VDF closure describing the surface interior to
|
||||
handle absorption and scattering inside the medium.
|
||||
@@ -5131,7 +5131,7 @@ is supported and controlled by the anisotropy input.
|
||||
\apiitem{anisotropy}
|
||||
\vspace{12pt}
|
||||
Scattering anisotropy [-1,1]. Negative values give backwards scattering,
|
||||
positive values give forward scatterixng, and 0.0 gives uniform scattering.
|
||||
positive values give forward scattering, and 0.0 gives uniform scattering.
|
||||
\apiend
|
||||
|
||||
% \vspace{-16pt}
|
||||
@@ -5149,7 +5149,7 @@ is supported and controlled by the anisotropy input.
|
||||
|
||||
Constructs a VDF for light passing through a dielectric homogeneous medium,
|
||||
such as glass or liquids. The parameters {\cf transmission_depth} and
|
||||
{\cf transmission_color} control the extinction coefficient of the medium in and
|
||||
{\cf transmission_color} control the extinction coefficient of the medium in an
|
||||
artist-friendly way. A priority can be set to determine the ordering of
|
||||
overlapping media.
|
||||
|
||||
@@ -5223,7 +5223,7 @@ $\mathrm{W}\cdot\mathrm{sr}^{-1}\cdot\mathrm{m}^{-2}$). This means that a surfac
|
||||
directly seen by the camera will directly reproduce the closure weight in the
|
||||
final pixel, regardless of being a surface or a volume.
|
||||
|
||||
For surface emission, If you divide {\cf emission} by
|
||||
For an emissive surface, if you divide the return value of {\cf uniform_edf()} by
|
||||
{\cf surfacearea() * M_PI}, then you can easily specify the total emissive
|
||||
power of the light (e.g., $\mathrm{W}$), regardless of its physical size.
|
||||
\apiend
|
||||
@@ -5344,7 +5344,7 @@ Nayar, ``Generalization of Lambert's Reflectance Model,'' Proceedings of
|
||||
SIGGRAPH 1994, pp.239-246 (July, 1994).
|
||||
|
||||
\begin{comment}
|
||||
Like all \closurecolors, the return ``value'' is symbolic and is may be
|
||||
Like all \closurecolors, the return ``value'' is symbolic and may be
|
||||
evaluated at a later time convenient to the renderer in order to compute
|
||||
the exitant radiance in the direction {\cf -I}. But aside from the
|
||||
fact that the shader cannot examine the numeric values of the
|
||||
@@ -5532,7 +5532,7 @@ regardless of its original direction.
|
||||
\indexapi{henyey_greenstein}
|
||||
Returns a \colorclosure that represents the directional volumetric
|
||||
scattering by small suspended particles. The {\cf g} parameter is the
|
||||
anisotropy factor, ranging from $-1$ to $1$), with positive values
|
||||
anisotropy factor, in the range $(-1, 1)$, with positive values
|
||||
indicating predominantly forward-scattering, negative values indicating
|
||||
predominantly back-scattering, and value of $g=0$ resulting in isotropic
|
||||
scattering.
|
||||
@@ -5563,7 +5563,7 @@ $\mathrm{W}\cdot\mathrm{sr}^{-1}\cdot\mathrm{m}^{-2}$). This means that a surfac
|
||||
directly seen by the camera will directly reproduce the closure weight in the
|
||||
final pixel, regardless of being a surface or a volume.
|
||||
|
||||
For surface emission, If you divide {\cf emission()} by
|
||||
For an emissive surface, if you divide the return value of {\cf emission()} by
|
||||
{\cf surfacearea() * M_PI}, then you can easily specify the total emissive
|
||||
power of the light (e.g., $\mathrm{W}$), regardless of its physical size.
|
||||
\apiend
|
||||
@@ -6834,7 +6834,7 @@ full control in the shader.
|
||||
\item[Shader.] A small self-contained program written in \langname, used
|
||||
to extend the functionality of a renderer with custom behavior of
|
||||
materials and lights. A particular shader may have multiple
|
||||
\emph{shader instances} within a scene, each of which has ts unique
|
||||
\emph{shader instances} within a scene, each of which has its unique
|
||||
\emph{instance parameters}, transformation, etc.
|
||||
|
||||
\item[Shader function.] A function written in \langname that may be
|
||||
|
||||
@@ -154,15 +154,6 @@ public:
|
||||
void debug_pop_inlined_function();
|
||||
void debug_set_location(OIIO::ustring sourcefile, int sourceline);
|
||||
|
||||
/// Create a new function (that will later be populated with
|
||||
/// instructions) with up to 4 args.
|
||||
OSL_DEPRECATED("Use make_function flavor that takes a cspan (1.12)")
|
||||
llvm::Function* make_function(const std::string& name, bool fastcall,
|
||||
llvm::Type* rettype, llvm::Type* arg1 = NULL,
|
||||
llvm::Type* arg2 = NULL,
|
||||
llvm::Type* arg3 = NULL,
|
||||
llvm::Type* arg4 = NULL);
|
||||
|
||||
/// Create a new function (that will later be populated with
|
||||
/// instructions) with a cspan of args.
|
||||
llvm::Function* make_function(const std::string& name, bool fastcall,
|
||||
@@ -289,12 +280,7 @@ public:
|
||||
Linkage default_linkage = Linkage::Internal,
|
||||
std::string* out_err = nullptr);
|
||||
|
||||
// OLD, might deprecate later
|
||||
/// Change symbols in the module that are marked as having external
|
||||
/// linkage to an alternate linkage that allows them to be discarded if
|
||||
/// not used within the module. Only do this for functions that start
|
||||
/// with prefix, and that DON'T match anything in the two exceptions
|
||||
/// lists.
|
||||
OSL_DEPRECATED("prune_and_internalize_module is better (1.13)")
|
||||
void internalize_module_functions(
|
||||
const std::string& prefix, const std::vector<std::string>& exceptions,
|
||||
const std::vector<std::string>& moreexceptions);
|
||||
|
||||
@@ -4,7 +4,6 @@
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <immintrin.h>
|
||||
#include <type_traits>
|
||||
|
||||
#include <OSL/oslconfig.h>
|
||||
@@ -23,6 +22,8 @@ using std::countr_zero;
|
||||
|
||||
#elif OSL_INTEL_CLASSIC_COMPILER_VERSION
|
||||
|
||||
#include <immintrin.h>
|
||||
|
||||
OSL_FORCEINLINE int popcount(uint32_t x) noexcept { return _mm_popcnt_u32(x);}
|
||||
OSL_FORCEINLINE int popcount(uint64_t x) noexcept { return _mm_popcnt_u64(x); }
|
||||
OSL_FORCEINLINE int countr_zero(uint32_t x) noexcept { return _bit_scan_forward(x); }
|
||||
|
||||
@@ -102,7 +102,7 @@ using OIIO::TextureOpt;
|
||||
// And some other things we borrow from OIIO...
|
||||
using OIIO::ErrorHandler;
|
||||
using OIIO::ustring;
|
||||
using OIIO::ustringHash;
|
||||
using OIIO::ustringhash;
|
||||
using OIIO::string_view;
|
||||
using OIIO::span;
|
||||
using OIIO::cspan;
|
||||
@@ -225,3 +225,30 @@ using ConstIndex = std::integral_constant<int, N>;
|
||||
|
||||
|
||||
OSL_NAMESPACE_EXIT
|
||||
|
||||
|
||||
namespace std { // not necessary in C++17, then we can just say std::hash
|
||||
#ifndef OIIO_USTRING_HAS_STDHASH
|
||||
// Not a new enough OIIO to define std::hash<ustring>, so we'll do it
|
||||
template<> struct hash<OSL::ustring> {
|
||||
std::size_t operator()(OSL::ustring u) const noexcept { return u.hash(); }
|
||||
};
|
||||
|
||||
template<> struct hash<OSL::ustringhash> {
|
||||
OSL_HOSTDEVICE constexpr std::size_t
|
||||
operator()(OSL::ustringhash u) const noexcept
|
||||
{
|
||||
return u.hash();
|
||||
}
|
||||
};
|
||||
#endif
|
||||
|
||||
// Not necessary once minimum OIIO defines operator< for ustringhash
|
||||
template<> struct less<OSL::ustringhash> {
|
||||
OSL_HOSTDEVICE constexpr bool operator()(OSL::ustringhash u,
|
||||
OSL::ustringhash v) const noexcept
|
||||
{
|
||||
return u.hash() < v.hash();
|
||||
}
|
||||
};
|
||||
} // namespace std
|
||||
|
||||
@@ -231,8 +231,7 @@ public:
|
||||
/// for debugging. (0)
|
||||
/// int llvm_dumpasm Print the CPU assembly code from the JIT (0)
|
||||
/// string llvm_prune_ir_strategy Strategy for pruning unnecessary
|
||||
/// IR (choices: "prune" [default],
|
||||
/// "internalize", or "none").
|
||||
/// IR (choices: "prune" [default], or "none").
|
||||
/// int max_local_mem_KB Error if shader group needs more than this
|
||||
/// much local storage to execute (1024K)
|
||||
/// string debug_groupname Name of shader group -- debug only this one
|
||||
@@ -673,6 +672,15 @@ public:
|
||||
ShaderGlobals& globals, void* userdata_base_ptr,
|
||||
void* output_base_ptr, bool run = true);
|
||||
|
||||
/// Future execute signature that will be range based. Shader globals will be
|
||||
/// obtained from renderer services.
|
||||
#if 0 // TODO in future PR
|
||||
bool execute(ShadingContext& ctx, ShaderGroup& group, OpaqueShadingStateUniformPtr ssu,
|
||||
int shadeindex_begin_at, int shadeindex_end_before
|
||||
void* userdata_base_ptr,
|
||||
void* output_base_ptr, bool run = true);
|
||||
#endif
|
||||
|
||||
// DEPRECATED(2.0): no shadeindex or base pointers
|
||||
bool execute(ShadingContext& ctx, ShaderGroup& group,
|
||||
ShaderGlobals& globals, bool run = true)
|
||||
|
||||
@@ -482,11 +482,19 @@
|
||||
# define OSL_DASSERT_MSG(x, ...) ((void)sizeof(x)) /*NOLINT*/
|
||||
#endif
|
||||
|
||||
#if defined(_MSC_VER) && !defined(__INTEL_COMPILER)
|
||||
#define OSL_STACK_ARRAY(TYPE, NAME, DYNAMIC_COUNT) TYPE * NAME = reinterpret_cast<TYPE *>(alloca(sizeof(TYPE)*(DYNAMIC_COUNT)))
|
||||
|
||||
/// OSL_ALLOCA is used to allocate smallish amount of memory on the stack,
|
||||
/// equivalent of C99 type var_name[size].
|
||||
///
|
||||
/// NOTE: in a debug build, this will assert for allocations >= 1MB, which is
|
||||
/// much too big. Hopefully this will keep us from abusing alloca and having
|
||||
/// stack overflows. The rule of thumb is that it's ok to use alloca for small
|
||||
/// things of bounded size, but not for anything that could be arbitrarily
|
||||
/// big, which should instead use heap allocation.
|
||||
#if defined(__GNUC__)
|
||||
# define OSL_ALLOCA(type, size) (assert(size < (1<<20)), (size) != 0 ? (reinterpret_cast<type*>(__builtin_alloca((size) * sizeof(type)))) : nullptr)
|
||||
#else
|
||||
// Utilize variable length array compiler extension when available
|
||||
#define OSL_STACK_ARRAY(TYPE, NAME, DYNAMIC_COUNT) TYPE NAME[DYNAMIC_COUNT]
|
||||
# define OSL_ALLOCA(type, size) (assert(size < (1<<20)), (size) != 0 ? (reinterpret_cast<type*>(alloca((size) * sizeof(type)))) : nullptr)
|
||||
#endif
|
||||
|
||||
OSL_NAMESPACE_ENTER
|
||||
|
||||
@@ -17,6 +17,9 @@ class RendererServices;
|
||||
/// coordinate transformation.
|
||||
typedef const void* TransformationPtr;
|
||||
|
||||
/// Type for an opaque pointer to whatever the Shading System uses to represent
|
||||
/// Uniform State.
|
||||
typedef const void* OpaqueShadingStateUniformPtr;
|
||||
|
||||
|
||||
/// The ShaderGlobals structure represents the state describing a particular
|
||||
@@ -83,6 +86,7 @@ struct ShaderGlobals {
|
||||
/// Back-pointer to the ShadingContext (set and used by OSL itself --
|
||||
/// renderers shouldn't mess with this at all).
|
||||
ShadingContext* context;
|
||||
OpaqueShadingStateUniformPtr shadingStateUniform;
|
||||
|
||||
/// Pointer to the RendererServices object. This is how OSL finds its
|
||||
/// way back to the renderer for callbacks.
|
||||
|
||||
@@ -2739,19 +2739,24 @@ foreach_unique(Wide<DataT, WidthT> wdata, Mask<WidthT> data_mask, FunctorT f)
|
||||
do {
|
||||
ActiveLane lead_lane(remaining_mask.first_on());
|
||||
DataT lead_data = wdata[lead_lane];
|
||||
Mask<WidthT> matching_lanes(false);
|
||||
// Use int instead of Mask<> to allow reduction clause in openmp simd declaration
|
||||
int matching_lanes_bits { 0 };
|
||||
OSL_FORCEINLINE_BLOCK
|
||||
{
|
||||
OSL_OMP_PRAGMA(omp simd simdlen(WidthT))
|
||||
OSL_OMP_PRAGMA(omp simd simdlen(WidthT)
|
||||
reduction(|
|
||||
: matching_lanes_bits))
|
||||
for (int lane = 0; lane < WidthT; ++lane) {
|
||||
// NOTE: the comparison ignores the remaining_mask
|
||||
bool lane_matches = (lead_data == wdata[lane]);
|
||||
// NOTE: using bitwise & to avoid branches
|
||||
if (lane_matches & remaining_mask[lane]) {
|
||||
matching_lanes.set_on(lane);
|
||||
// inline of Mask::set_on(lane)
|
||||
matching_lanes_bits |= 1 << lane;
|
||||
}
|
||||
}
|
||||
}
|
||||
Mask<WidthT> matching_lanes { matching_lanes_bits };
|
||||
|
||||
f(lead_data, matching_lanes);
|
||||
remaining_mask &= ~matching_lanes;
|
||||
|
||||
@@ -490,7 +490,7 @@ private:
|
||||
int m_main_method_start; ///< Instruction where 'main' starts
|
||||
bool m_declaring_shader_formals; ///< Are we declaring shader formals?
|
||||
std::set<std::pair<ustring, int>> m_nowarn_lines; ///< Lines for 'nowarn'
|
||||
typedef std::unordered_map<ustring, std::string, ustringHash> FCMap;
|
||||
typedef std::unordered_map<ustring, std::string> FCMap;
|
||||
FCMap m_filecontents_map; ///< Contents of source files we've seen
|
||||
ustring m_last_sourcefile; ///< Last filename for retrieve_source
|
||||
std::string* m_last_filecontents = nullptr; //< Last file contents
|
||||
|
||||
@@ -28,6 +28,10 @@
|
||||
#pragma clang diagnostic ignored "-Wparentheses-equality"
|
||||
#endif
|
||||
|
||||
#if (OSL_CLANG_VERSION >= 150000) || (OSL_INTEL_CLANG_VERSION >= 140000)
|
||||
#pragma clang diagnostic ignored "-Wunused-but-set-variable"
|
||||
#endif
|
||||
|
||||
#if defined(__GNUC__) || defined(__clang__)
|
||||
#pragma GCC visibility push(hidden)
|
||||
#endif
|
||||
|
||||
@@ -186,7 +186,7 @@ private:
|
||||
///
|
||||
class SymbolTable {
|
||||
public:
|
||||
typedef std::unordered_map<ustring, Symbol*, ustringHash> ScopeTable;
|
||||
typedef std::unordered_map<ustring, Symbol*> ScopeTable;
|
||||
typedef std::vector<ScopeTable> ScopeTableStack;
|
||||
typedef SymbolPtrVec::iterator iterator;
|
||||
typedef SymbolPtrVec::const_iterator const_iterator;
|
||||
|
||||
@@ -370,7 +370,13 @@ foreach(batched_target ${BATCHED_TARGET_LIST})
|
||||
endif ()
|
||||
|
||||
if (${TARGET_OPT_FMA} STREQUAL "noFMA")
|
||||
list (APPEND TARGET_CXX_OPTS "-mno-fma")
|
||||
# Do NOT use "-mno-fma" as it may lower ISA to one without any FMA
|
||||
# instructions (AVX512->AVX2). Our intent is to not generate
|
||||
# FMA's but still use the available instructions and register
|
||||
# width. However math library calls may use still end up using
|
||||
# FMA's, but lets try turning of fused math ops because many OSL
|
||||
# math functions are from OIIO/IMath headers generate non FMA
|
||||
list (APPEND TARGET_CXX_OPTS "-ffp-contract=off")
|
||||
endif ()
|
||||
|
||||
# large SIMD function loops will exceed llvm's -inline-threshold default of 225.
|
||||
@@ -388,7 +394,8 @@ foreach(batched_target ${BATCHED_TARGET_LIST})
|
||||
list (APPEND TARGET_CXX_OPTS "-fno-unroll-loops")
|
||||
|
||||
if (CMAKE_COMPILER_IS_INTELCLANG)
|
||||
list (APPEND TARGET_CXX_OPTS "-qopenmp-simd")
|
||||
# utilize Intel compiler's optimized openmp simd generation vs. open source "-qopenmp-simd"
|
||||
list (APPEND TARGET_CXX_OPTS "-fiopenmp-simd")
|
||||
else ()
|
||||
list (APPEND TARGET_CXX_OPTS "-fopenmp")
|
||||
endif ()
|
||||
@@ -434,7 +441,13 @@ foreach(batched_target ${BATCHED_TARGET_LIST})
|
||||
endif ()
|
||||
|
||||
if (${TARGET_OPT_FMA} STREQUAL "noFMA")
|
||||
list (APPEND TARGET_CXX_OPTS "-mno-fma")
|
||||
# Do NOT use "-mno-fma" as it may lower ISA to one without any FMA
|
||||
# instructions (AVX512->AVX2). Our intent is to not generate
|
||||
# FMA's but still use the available instructions and register
|
||||
# width. However math library calls may use still end up using
|
||||
# FMA's, but lets try turning of fused math ops because many OSL
|
||||
# math functions are from OIIO/IMath headers generate non FMA
|
||||
list (APPEND TARGET_CXX_OPTS "-ffp-contract=off")
|
||||
endif ()
|
||||
|
||||
list (APPEND TARGET_CXX_OPTS "-fopenmp-simd")
|
||||
|
||||
@@ -23,12 +23,12 @@ OSL_NAMESPACE_ENTER
|
||||
typedef std::set<int>
|
||||
IntSet; // probably faster to test for equality, unions and so
|
||||
|
||||
typedef std::unordered_set<ustring, ustringHash> SymbolSet;
|
||||
typedef std::unordered_set<ustring> SymbolSet;
|
||||
// This is for the transition table used in DfAutomata::State
|
||||
typedef std::unordered_map<ustring, int, ustringHash> SymbolToInt;
|
||||
typedef std::unordered_map<ustring, int> SymbolToInt;
|
||||
// And this is for the transition table in NdfAutomata which
|
||||
// has several movements for each symbol
|
||||
typedef std::unordered_map<ustring, IntSet, ustringHash> SymbolToIntList;
|
||||
typedef std::unordered_map<ustring, IntSet> SymbolToIntList;
|
||||
typedef std::unordered_map<int, int> HashIntInt;
|
||||
|
||||
// For the rules in the deterministic states, we don't need a real set
|
||||
|
||||
@@ -236,6 +236,7 @@ static ustring fields[] = { ustring("P"),
|
||||
ustring("tracedata"),
|
||||
ustring("objdata"),
|
||||
ustring("shadingcontext"),
|
||||
ustring("shadingStateUniform"),
|
||||
ustring("renderer"),
|
||||
ustring("object2common"),
|
||||
ustring("shader2common"),
|
||||
|
||||
@@ -240,7 +240,7 @@ private:
|
||||
static constexpr OpcodeTest no_test = nullptr;
|
||||
|
||||
// Map opcode names to a Dependency info
|
||||
std::unordered_map<ustring, Dependency, ustringHash> m_lookup;
|
||||
std::unordered_map<ustring, Dependency> m_lookup;
|
||||
|
||||
/// Return the ptr to the symbol that is the argnum-th argument to the
|
||||
/// given op in the instance.
|
||||
@@ -643,8 +643,8 @@ public:
|
||||
auto write_iter = begin_at(write_pos);
|
||||
|
||||
const int read_count = read_iter.depth() + 1;
|
||||
OSL_STACK_ARRAY(Position, read_path, read_count);
|
||||
int read_depth = 0;
|
||||
Position* read_path = OSL_ALLOCA(Position, read_count);
|
||||
int read_depth = 0;
|
||||
{
|
||||
// Loop structure is to allow
|
||||
// writing the end position into the read path
|
||||
@@ -658,8 +658,8 @@ public:
|
||||
}
|
||||
|
||||
const int writeCount = write_iter.depth() + 1;
|
||||
OSL_STACK_ARRAY(Position, write_path, writeCount);
|
||||
int write_depth = 0;
|
||||
Position* write_path = OSL_ALLOCA(Position, writeCount);
|
||||
int write_depth = 0;
|
||||
{
|
||||
// Loop structure is to allow
|
||||
// writing the end position into the write path
|
||||
|
||||
@@ -2691,21 +2691,9 @@ BatchedBackendLLVM::run()
|
||||
// entry points, as well as for all the external functions that are
|
||||
// just declarations (not definitions) in the module (which we have
|
||||
// conveniently stashed in external_function_names).
|
||||
#if 0
|
||||
std::vector<std::string> entry_function_names;
|
||||
entry_function_names.push_back (ll.func_name(init_func));
|
||||
for (int layer = 0; layer < nlayers; ++layer) {
|
||||
// set_inst (layer);
|
||||
llvm::Function* f = funcs[layer];
|
||||
if (f && group().is_entry_layer(layer))
|
||||
entry_function_names.push_back (ll.func_name(f));
|
||||
}
|
||||
ll.internalize_module_functions ("osl_", external_function_names, entry_function_names);
|
||||
#else
|
||||
std::unordered_set<llvm::Function*> external_functions;
|
||||
external_functions.insert(init_func);
|
||||
for (int layer = 0; layer < nlayers; ++layer) {
|
||||
// set_inst (layer);
|
||||
llvm::Function* f = funcs[layer];
|
||||
// If we plan to call bitcode_string of a layer's function after optimization
|
||||
// it may not exist after optimization unless we treat it as external.
|
||||
@@ -2714,7 +2702,6 @@ BatchedBackendLLVM::run()
|
||||
}
|
||||
}
|
||||
ll.prune_and_internalize_module(external_functions);
|
||||
#endif
|
||||
|
||||
// Debug code to dump the pre-optimized bitcode to a file
|
||||
if (llvm_debug() >= 2 || shadingsys().llvm_output_bitcode()) {
|
||||
|
||||
@@ -117,7 +117,6 @@ BatchedRendererServices<WidthT>::resolve_udim_uniform(
|
||||
if (!texture_handle)
|
||||
texture_handle = texturesys()->get_texture_handle(filename,
|
||||
texture_thread_info);
|
||||
#if OPENIMAGEIO_VERSION >= 20307
|
||||
if (texturesys()->is_udim(texture_handle)) {
|
||||
// Newer versions of the TextureSystem interface are able to determine the
|
||||
// specific UDIM tile we're using.
|
||||
@@ -131,9 +130,9 @@ BatchedRendererServices<WidthT>::resolve_udim_uniform(
|
||||
udim_handle = texture_handle;
|
||||
}
|
||||
return udim_handle;
|
||||
} else
|
||||
#endif
|
||||
} else {
|
||||
return texture_handle;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -151,7 +150,6 @@ BatchedRendererServices<WidthT>::resolve_udim(
|
||||
if (!texture_handle)
|
||||
texture_handle = texturesys()->get_texture_handle(filename,
|
||||
texture_thread_info);
|
||||
#if OPENIMAGEIO_VERSION >= 20307
|
||||
if (texturesys()->is_udim(texture_handle)) {
|
||||
// Newer versions of the TextureSystem interface are able to determine the
|
||||
// specific UDIM tile we're using.
|
||||
@@ -167,9 +165,9 @@ BatchedRendererServices<WidthT>::resolve_udim(
|
||||
}
|
||||
wresult[l] = udim_handle;
|
||||
});
|
||||
} else
|
||||
#endif
|
||||
} else {
|
||||
assign_all(wresult, texture_handle);
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
|
||||
@@ -2410,7 +2410,7 @@ DECLFOLDER(constfold_gettextureinfo)
|
||||
ustring filename = Filename.get_string();
|
||||
ustring dataname = Dataname.get_string();
|
||||
TypeDesc t = Data.typespec().simpletype();
|
||||
void* mydata = OIIO_ALLOCA(char, t.size());
|
||||
void* mydata = OSL_ALLOCA(char, t.size());
|
||||
// FIXME(ptex) -- exclude folding of ptex, since these things
|
||||
// can vary per face.
|
||||
ustring errormessage;
|
||||
@@ -2776,7 +2776,7 @@ DECLFOLDER(constfold_pointcloud_get)
|
||||
return 0;
|
||||
|
||||
// Must transfer to size_t array
|
||||
size_t* indices = OIIO_ALLOCA(size_t, count);
|
||||
size_t* indices = OSL_ALLOCA(size_t, count);
|
||||
for (int i = 0; i < count; ++i)
|
||||
indices[i] = Indices.get_int(i);
|
||||
|
||||
|
||||
@@ -128,9 +128,10 @@ ShadingContext::execute_init(ShaderGroup& sgroup, int shadeindex,
|
||||
RunLLVMGroupFunc run_func = sgroup.llvm_compiled_init();
|
||||
if (!run_func)
|
||||
return false;
|
||||
ssg.context = this;
|
||||
ssg.renderer = renderer();
|
||||
ssg.Ci = NULL;
|
||||
ssg.context = this;
|
||||
ssg.shadingStateUniform = &(shadingsys().m_shading_state_uniform);
|
||||
ssg.renderer = renderer();
|
||||
ssg.Ci = NULL;
|
||||
run_func(&ssg, m_heap.get(), userdata_base_ptr, output_base_ptr,
|
||||
shadeindex);
|
||||
}
|
||||
|
||||
@@ -115,7 +115,7 @@ private:
|
||||
};
|
||||
|
||||
typedef std::unordered_map<Query, QueryResult, QueryHash> QueryMap;
|
||||
typedef std::unordered_map<ustring, int, ustringHash> DocMap;
|
||||
typedef std::unordered_map<ustring, int> DocMap;
|
||||
|
||||
ShadingContext* m_context; // back-pointer to shading context
|
||||
|
||||
|
||||
@@ -466,11 +466,9 @@ ShaderInstance::evaluate_writes_globals_and_userdata_params()
|
||||
// userdata_params as accurately as we can based on what we know from
|
||||
// the symbol overrides. This is very important to get instance merging
|
||||
// working correctly.
|
||||
int p = 0;
|
||||
for (auto&& s : m_instoverrides) {
|
||||
if (!s.lockgeom())
|
||||
userdata_params(true);
|
||||
++p;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
@@ -216,6 +216,7 @@ BackendLLVM::llvm_type_sg()
|
||||
sg_types.push_back(vp); // opaque tracedata*
|
||||
sg_types.push_back(vp); // opaque objdata*
|
||||
sg_types.push_back(vp); // ShadingContext*
|
||||
sg_types.push_back(vp); // OpaqueShadingStateUniformPtr
|
||||
sg_types.push_back(vp); // RendererServices*
|
||||
sg_types.push_back(vp); // object2common
|
||||
sg_types.push_back(vp); // shader2common
|
||||
@@ -1501,21 +1502,6 @@ BackendLLVM::run()
|
||||
if (shadingsys().llvm_prune_ir_strategy() == "none") {
|
||||
// Do nothing! This is only useful for testing how much we reduce
|
||||
// optimization and JIT time with the other strategies.
|
||||
} else if (shadingsys().llvm_prune_ir_strategy() == "internalize") {
|
||||
// Mark as 'internal' all functions that start with "osl_" but are
|
||||
// but are not among the known entry points. LLVM can then quickly
|
||||
// realize that any 'internal' functions not called can be
|
||||
// discarded. This was our original stab at this strategy.
|
||||
std::vector<std::string> entry_function_names;
|
||||
entry_function_names.push_back(ll.func_name(init_func));
|
||||
for (int layer = 0; layer < nlayers; ++layer) {
|
||||
// set_inst (layer);
|
||||
llvm::Function* f = funcs[layer];
|
||||
if (f && group().is_entry_layer(layer))
|
||||
entry_function_names.push_back(ll.func_name(f));
|
||||
}
|
||||
ll.internalize_module_functions("osl_", external_function_names,
|
||||
entry_function_names);
|
||||
} else /* if (shadingsys().llvm_prune_ir_strategy() == "prune") */ {
|
||||
// New (2020) and default behavior, from Alex Wells:
|
||||
// Full prune of uncalled functions, and marking as 'internal' to
|
||||
|
||||
@@ -399,6 +399,11 @@ LLVM_Util::LLVM_Util(const PerThreadInfo& per_thread_info, int debuglevel,
|
||||
OIIO::spin_lock lock(llvm_global_mutex);
|
||||
if (!m_thread->llvm_context) {
|
||||
m_thread->llvm_context = new llvm::LLVMContext();
|
||||
#if OSL_LLVM_VERSION >= 150
|
||||
m_thread->llvm_context->setOpaquePointers(false);
|
||||
// FIXME: For now, keep using typed pointers. We're going to have
|
||||
// to fix this and switch to opaque pointers by llvm 16.
|
||||
#endif
|
||||
//static SetCommandLineOptionsForLLVM sSetCommandLineOptionsForLLVM;
|
||||
}
|
||||
|
||||
@@ -1584,29 +1589,26 @@ LLVM_Util::make_jit_execengine(std::string* err, TargetISA requestedISA,
|
||||
|
||||
|
||||
|
||||
namespace /*anonymous*/ {
|
||||
#ifdef OSL_DEV
|
||||
// The return value of llvm::StructLayout::getAlignment()
|
||||
// changed from an int to llvm::Align, hide with accessor function
|
||||
#if OSL_LLVM_VERSION < 100
|
||||
uint64_t
|
||||
inline uint64_t
|
||||
get_alignment(const llvm::StructLayout* layout)
|
||||
{
|
||||
# if OSL_LLVM_VERSION < 100
|
||||
return layout->getAlignment();
|
||||
}
|
||||
#else
|
||||
uint64_t
|
||||
get_alignment(const llvm::StructLayout* layout)
|
||||
{
|
||||
# else
|
||||
return layout->getAlignment().value();
|
||||
# endif
|
||||
}
|
||||
#endif
|
||||
} // namespace
|
||||
|
||||
|
||||
|
||||
void
|
||||
LLVM_Util::dump_struct_data_layout(llvm::Type* Ty)
|
||||
LLVM_Util::dump_struct_data_layout(llvm::Type* Ty OSL_MAYBE_UNUSED)
|
||||
{
|
||||
#ifdef OSL_DEV
|
||||
OSL_ASSERT(Ty);
|
||||
OSL_ASSERT(Ty->isStructTy());
|
||||
|
||||
@@ -1630,6 +1632,7 @@ LLVM_Util::dump_struct_data_layout(llvm::Type* Ty)
|
||||
}
|
||||
std::cout << std::endl;
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
@@ -1884,7 +1887,9 @@ LLVM_Util::setup_optimization_passes(int optlevel, bool target_host)
|
||||
// we might need to recreate that meta data in OSL's loop code to enable these passes
|
||||
mpm.add(llvm::createLoopRotatePass());
|
||||
mpm.add(llvm::createLICMPass());
|
||||
# if OSL_LLVM_VERSION < 150
|
||||
mpm.add(llvm::createLoopUnswitchPass(false));
|
||||
# endif
|
||||
// mpm.add(llvm::createInstructionCombiningPass());
|
||||
mpm.add(llvm::createIndVarSimplifyPass());
|
||||
// Don't think we emitted any idioms that should be converted to a loop
|
||||
@@ -1971,7 +1976,9 @@ LLVM_Util::setup_optimization_passes(int optlevel, bool target_host)
|
||||
mpm.add(llvm::createDeadCodeEliminationPass());
|
||||
mpm.add(llvm::createCFGSimplificationPass());
|
||||
|
||||
#if OSL_LLVM_VERSION < 150
|
||||
mpm.add(llvm::createArgumentPromotionPass());
|
||||
#endif
|
||||
mpm.add(llvm::createAggressiveDCEPass());
|
||||
mpm.add(llvm::createInstructionCombiningPass());
|
||||
mpm.add(llvm::createJumpThreadingPass());
|
||||
@@ -1992,7 +1999,9 @@ LLVM_Util::setup_optimization_passes(int optlevel, bool target_host)
|
||||
mpm.add(llvm::createCFGSimplificationPass());
|
||||
|
||||
mpm.add(llvm::createFunctionInliningPass());
|
||||
#if OSL_LLVM_VERSION < 150
|
||||
mpm.add(llvm::createArgumentPromotionPass());
|
||||
#endif
|
||||
mpm.add(llvm::createSROAPass());
|
||||
|
||||
mpm.add(llvm::createInstructionCombiningPass());
|
||||
@@ -2000,7 +2009,9 @@ LLVM_Util::setup_optimization_passes(int optlevel, bool target_host)
|
||||
mpm.add(llvm::createReassociatePass());
|
||||
mpm.add(llvm::createLoopRotatePass());
|
||||
mpm.add(llvm::createLICMPass());
|
||||
#if OSL_LLVM_VERSION < 150
|
||||
mpm.add(llvm::createLoopUnswitchPass(false));
|
||||
#endif
|
||||
mpm.add(llvm::createInstructionCombiningPass());
|
||||
mpm.add(llvm::createIndVarSimplifyPass());
|
||||
mpm.add(llvm::createLoopIdiomPass());
|
||||
@@ -2426,89 +2437,12 @@ LLVM_Util::validate_global_mappings(
|
||||
|
||||
|
||||
|
||||
// Soon to be deprecated
|
||||
// DEPRECATED(1.13)
|
||||
void
|
||||
LLVM_Util::internalize_module_functions(
|
||||
const std::string& prefix, const std::vector<std::string>& exceptions,
|
||||
const std::vector<std::string>& moreexceptions)
|
||||
{
|
||||
for (llvm::Function& func : module()->getFunctionList()) {
|
||||
llvm::Function* sym = &func;
|
||||
std::string symname = sym->getName().str();
|
||||
if (prefix.size() && !OIIO::Strutil::starts_with(symname, prefix))
|
||||
continue;
|
||||
bool needed = false;
|
||||
for (size_t i = 0, e = exceptions.size(); i < e; ++i)
|
||||
if (sym->getName() == exceptions[i]) {
|
||||
needed = true;
|
||||
// std::cout << " necessary LLVM module function "
|
||||
// << sym->getName().str() << "\n";
|
||||
break;
|
||||
}
|
||||
for (size_t i = 0, e = moreexceptions.size(); i < e; ++i)
|
||||
if (sym->getName() == moreexceptions[i]) {
|
||||
needed = true;
|
||||
// std::cout << " necessary LLVM module function "
|
||||
// << sym->getName().str() << "\n";
|
||||
break;
|
||||
}
|
||||
if (!needed) {
|
||||
llvm::GlobalValue::LinkageTypes linkage = sym->getLinkage();
|
||||
// std::cout << " unnecessary LLVM module function "
|
||||
// << sym->getName().str() << " linkage " << int(linkage) << "\n";
|
||||
if (linkage == llvm::GlobalValue::ExternalLinkage)
|
||||
sym->setLinkage(llvm::GlobalValue::LinkOnceODRLinkage);
|
||||
// ExternalLinkage means it's potentially externally callable,
|
||||
// and so will definitely have code generated.
|
||||
// LinkOnceODRLinkage keeps one copy so it can be inlined or
|
||||
// called internally to the module, but allows it to be
|
||||
// discarded otherwise.
|
||||
}
|
||||
}
|
||||
#if 0
|
||||
// I don't think we need to worry about linkage of global symbols, but
|
||||
// here is an example of how to iterate over the globals anyway.
|
||||
for (llvm::Module::global_iterator iter = module()->global_begin(); iter != module()->global_end(); iter++) {
|
||||
llvm::GlobalValue *sym = llvm::dyn_cast<llvm::GlobalValue>(iter);
|
||||
if (!sym)
|
||||
continue;
|
||||
std::string symname = sym->getName();
|
||||
if (prefix.size() && ! OIIO::Strutil::starts_with(symname, prefix))
|
||||
continue;
|
||||
bool needed = false;
|
||||
for (size_t i = 0, e = exceptions.size(); i < e; ++i)
|
||||
if (sym->getName() == exceptions[i]) {
|
||||
needed = true;
|
||||
break;
|
||||
}
|
||||
if (! needed) {
|
||||
llvm::GlobalValue::LinkageTypes linkage = sym->getLinkage();
|
||||
// std::cout << " unnecessary LLVM global " << sym->getName().str()
|
||||
// << " linkage " << int(linkage) << "\n";
|
||||
if (linkage == llvm::GlobalValue::ExternalLinkage)
|
||||
f->setLinkage (llvm::GlobalValue::LinkOnceODRLinkage);
|
||||
}
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
|
||||
llvm::Function*
|
||||
LLVM_Util::make_function(const std::string& name, bool fastcall,
|
||||
llvm::Type* rettype, llvm::Type* arg1,
|
||||
llvm::Type* arg2, llvm::Type* arg3, llvm::Type* arg4)
|
||||
{
|
||||
std::vector<llvm::Type*> argtypes;
|
||||
if (arg1)
|
||||
argtypes.emplace_back(arg1);
|
||||
if (arg2)
|
||||
argtypes.emplace_back(arg2);
|
||||
if (arg3)
|
||||
argtypes.emplace_back(arg3);
|
||||
if (arg4)
|
||||
argtypes.emplace_back(arg4);
|
||||
return make_function(name, fastcall, rettype, argtypes, false);
|
||||
}
|
||||
|
||||
|
||||
@@ -3663,9 +3597,14 @@ LLVM_Util::call_function(llvm::Value* func, cspan<llvm::Value*> args)
|
||||
#endif
|
||||
//llvm_gen_debug_printf (std::string("start ") + std::string(name));
|
||||
#if OSL_LLVM_VERSION >= 110
|
||||
OSL_PRAGMA_WARNING_PUSH
|
||||
OSL_GCC_PRAGMA(GCC diagnostic ignored "-Wdeprecated-declarations")
|
||||
// FIXME: This will need to be revisited for future LLVM 16+ that fully
|
||||
// drops typed pointers.
|
||||
llvm::Value* r = builder().CreateCall(
|
||||
llvm::cast<llvm::FunctionType>(func->getType()->getPointerElementType()),
|
||||
func, llvm::ArrayRef<llvm::Value*>(args.data(), args.size()));
|
||||
OSL_PRAGMA_WARNING_POP
|
||||
#else
|
||||
llvm::Value* r
|
||||
= builder().CreateCall(func, llvm::ArrayRef<llvm::Value*>(args.data(),
|
||||
@@ -3802,7 +3741,12 @@ LLVM_Util::op_load(llvm::Type* type, llvm::Value* ptr,
|
||||
llvm::Value*
|
||||
LLVM_Util::op_load(llvm::Value* ptr, const std::string& llname)
|
||||
{
|
||||
OSL_PRAGMA_WARNING_PUSH
|
||||
OSL_GCC_PRAGMA(GCC diagnostic ignored "-Wdeprecated-declarations")
|
||||
// FIXME: This will need to be revisited for future LLVM 16+ that fully
|
||||
// drops typed pointers.
|
||||
return op_load(ptr->getType()->getPointerElementType(), ptr, llname);
|
||||
OSL_PRAGMA_WARNING_POP
|
||||
}
|
||||
|
||||
|
||||
@@ -5320,8 +5264,13 @@ LLVM_Util::GEP(llvm::Type* type, llvm::Value* ptr, llvm::Value* elem,
|
||||
llvm::Value*
|
||||
LLVM_Util::GEP(llvm::Value* ptr, llvm::Value* elem, const std::string& llname)
|
||||
{
|
||||
OSL_PRAGMA_WARNING_PUSH
|
||||
OSL_GCC_PRAGMA(GCC diagnostic ignored "-Wdeprecated-declarations")
|
||||
// FIXME: This will need to be revisited for future LLVM 16+ that fully
|
||||
// drops typed pointers.
|
||||
return GEP(ptr->getType()->getScalarType()->getPointerElementType(), ptr,
|
||||
elem, llname);
|
||||
OSL_PRAGMA_WARNING_POP
|
||||
}
|
||||
|
||||
|
||||
@@ -5338,8 +5287,13 @@ LLVM_Util::GEP(llvm::Type* type, llvm::Value* ptr, int elem,
|
||||
llvm::Value*
|
||||
LLVM_Util::GEP(llvm::Value* ptr, int elem, const std::string& llname)
|
||||
{
|
||||
OSL_PRAGMA_WARNING_PUSH
|
||||
OSL_GCC_PRAGMA(GCC diagnostic ignored "-Wdeprecated-declarations")
|
||||
// FIXME: This will need to be revisited for future LLVM 16+ that fully
|
||||
// drops typed pointers.
|
||||
return GEP(ptr->getType()->getScalarType()->getPointerElementType(), ptr,
|
||||
elem, llname);
|
||||
OSL_PRAGMA_WARNING_POP
|
||||
}
|
||||
|
||||
|
||||
@@ -5357,8 +5311,13 @@ llvm::Value*
|
||||
LLVM_Util::GEP(llvm::Value* ptr, int elem1, int elem2,
|
||||
const std::string& llname)
|
||||
{
|
||||
OSL_PRAGMA_WARNING_PUSH
|
||||
OSL_GCC_PRAGMA(GCC diagnostic ignored "-Wdeprecated-declarations")
|
||||
// FIXME: This will need to be revisited for future LLVM 16+ that fully
|
||||
// drops typed pointers.
|
||||
return GEP(ptr->getType()->getScalarType()->getPointerElementType(), ptr,
|
||||
elem1, elem2, llname);
|
||||
OSL_PRAGMA_WARNING_POP
|
||||
}
|
||||
|
||||
|
||||
@@ -5851,6 +5810,7 @@ bool
|
||||
LLVM_Util::ptx_compile_group(llvm::Module* lib_module, const std::string& name,
|
||||
std::string& out)
|
||||
{
|
||||
#ifdef OSL_USE_OPTIX
|
||||
std::string target_triple = module()->getTargetTriple();
|
||||
|
||||
OSL_ASSERT(
|
||||
@@ -5894,9 +5854,9 @@ LLVM_Util::ptx_compile_group(llvm::Module* lib_module, const std::string& name,
|
||||
options.AllowFPOpFusion = llvm::FPOpFusion::Fast;
|
||||
options.NoZerosInBSS = 0;
|
||||
options.GuaranteedTailCallOpt = 0;
|
||||
#if OSL_LLVM_VERSION < 120
|
||||
# if OSL_LLVM_VERSION < 120
|
||||
options.StackAlignmentOverride = 0;
|
||||
#endif
|
||||
# endif
|
||||
options.UseInitArray = 0;
|
||||
|
||||
llvm::TargetMachine* target_machine = llvm_target->createTargetMachine(
|
||||
@@ -5928,15 +5888,15 @@ LLVM_Util::ptx_compile_group(llvm::Module* lib_module, const std::string& name,
|
||||
llvm::raw_svector_ostream assembly_stream(assembly);
|
||||
|
||||
// TODO: Make sure rounding modes, etc., are set correctly
|
||||
#if OSL_LLVM_VERSION >= 100
|
||||
# if OSL_LLVM_VERSION >= 100
|
||||
target_machine->addPassesToEmitFile(mod_pm, assembly_stream,
|
||||
nullptr, // FIXME: Correct?
|
||||
llvm::CGFT_AssemblyFile);
|
||||
#else
|
||||
# else
|
||||
target_machine->addPassesToEmitFile(mod_pm, assembly_stream,
|
||||
nullptr, // FIXME: Correct?
|
||||
llvm::TargetMachine::CGFT_AssemblyFile);
|
||||
#endif
|
||||
# endif
|
||||
|
||||
// Run the optimization passes on the functions
|
||||
fn_pm.doInitialization();
|
||||
@@ -5962,6 +5922,9 @@ LLVM_Util::ptx_compile_group(llvm::Module* lib_module, const std::string& name,
|
||||
delete linked_module;
|
||||
|
||||
return true;
|
||||
#else
|
||||
return false;
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
|
||||
@@ -72,7 +72,7 @@ private:
|
||||
int m_oso_major, m_oso_minor; ///< oso file format version
|
||||
int m_sym_default_index; ///< Next sym default value to fill in
|
||||
bool m_errors; ///< Did we hit any errors?
|
||||
typedef std::unordered_map<ustring, int, ustringHash> UstringIntMap;
|
||||
typedef std::unordered_map<ustring, int> UstringIntMap;
|
||||
UstringIntMap m_symmap; ///< map sym name to index
|
||||
};
|
||||
|
||||
|
||||
+61
-61
@@ -47,14 +47,14 @@ Parser::Parser(const std::vector<ustring>* user_events,
|
||||
m_minus_stop.insert(Labels::STOP);
|
||||
|
||||
if (user_events)
|
||||
for (size_t i = 0; i < user_events->size(); ++i) {
|
||||
m_label_position[(*user_events)[i]] = 0;
|
||||
m_basic_labels.insert((*user_events)[i]);
|
||||
for (auto&& e : *user_events) {
|
||||
m_label_position[e] = 0;
|
||||
m_basic_labels.insert(e);
|
||||
}
|
||||
if (user_scatterings)
|
||||
for (size_t i = 0; i < user_scatterings->size(); ++i) {
|
||||
m_label_position[(*user_scatterings)[i]] = 1;
|
||||
m_basic_labels.insert((*user_scatterings)[i]);
|
||||
for (auto&& s : *user_scatterings) {
|
||||
m_label_position[s] = 1;
|
||||
m_basic_labels.insert(s);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -65,9 +65,8 @@ Parser::buildStop(LPexp* etype, LPexp* scatter, const std::list<LPexp*>& custom)
|
||||
lpexp::Cat* cat = new lpexp::Cat();
|
||||
cat->append(etype);
|
||||
cat->append(scatter);
|
||||
for (std::list<LPexp*>::const_iterator i = custom.begin();
|
||||
i != custom.end(); ++i)
|
||||
cat->append(*i);
|
||||
for (const auto& p : custom)
|
||||
cat->append(p);
|
||||
|
||||
if (custom.size() < 5)
|
||||
cat->append(new lpexp::Repeat(new lpexp::Wildexp(m_basic_labels)));
|
||||
@@ -98,9 +97,8 @@ Parser::parseSymbol()
|
||||
if (sym != ".") {
|
||||
SymbolToInt::const_iterator i = m_label_position.find(sym);
|
||||
if (i == m_label_position.end()) {
|
||||
m_error = std::string("Unrecognized basic label: ")
|
||||
+ sym.c_str();
|
||||
return NULL;
|
||||
m_error = fmtformat("Unrecognized basic label: {}", sym);
|
||||
return nullptr;
|
||||
}
|
||||
int pos = i->second;
|
||||
basics[pos] = new lpexp::Symbol(sym);
|
||||
@@ -161,16 +159,17 @@ Parser::parseCat()
|
||||
if (head() == '|') {
|
||||
if (!explist.size()) {
|
||||
m_error = "No left expression to or with |";
|
||||
for (size_t i = 0; i < explist.size(); ++i)
|
||||
delete explist[i];
|
||||
return NULL;
|
||||
for (auto& exp : explist)
|
||||
delete exp;
|
||||
return nullptr;
|
||||
}
|
||||
next();
|
||||
LPexp* e = _parse();
|
||||
if (error()) {
|
||||
for (size_t i = 0; i < explist.size(); ++i)
|
||||
delete explist[i];
|
||||
return NULL;
|
||||
delete e;
|
||||
for (auto& exp : explist)
|
||||
delete exp;
|
||||
return nullptr;
|
||||
}
|
||||
if (explist.back()->getType() == lpexp::OR)
|
||||
((lpexp::Orlist*)explist.back())->append(e);
|
||||
@@ -183,9 +182,10 @@ Parser::parseCat()
|
||||
} else {
|
||||
LPexp* e = _parse();
|
||||
if (error()) {
|
||||
for (size_t i = 0; i < explist.size(); ++i)
|
||||
delete explist[i];
|
||||
return NULL;
|
||||
delete e;
|
||||
for (auto& exp : explist)
|
||||
delete exp;
|
||||
return nullptr;
|
||||
}
|
||||
explist.push_back(e);
|
||||
}
|
||||
@@ -194,13 +194,13 @@ Parser::parseCat()
|
||||
next();
|
||||
else if (endchar != 0) {
|
||||
m_error = "Reached end of line looking for )";
|
||||
for (size_t i = 0; i < explist.size(); ++i)
|
||||
delete explist[i];
|
||||
return NULL;
|
||||
for (auto& exp : explist)
|
||||
delete exp;
|
||||
return nullptr;
|
||||
}
|
||||
lpexp::Cat* cat = new lpexp::Cat();
|
||||
for (size_t i = 0; i < explist.size(); ++i)
|
||||
cat->append(explist[i]);
|
||||
for (auto& exp : explist)
|
||||
cat->append(exp);
|
||||
return cat;
|
||||
}
|
||||
|
||||
@@ -212,29 +212,29 @@ Parser::parseGroup()
|
||||
OSL_DASSERT(head() == '<');
|
||||
if (m_ingroup) {
|
||||
m_error = "No groups allowed inside of groups";
|
||||
return NULL;
|
||||
return nullptr;
|
||||
}
|
||||
int basicpos = 0;
|
||||
LPexp* basics[2] = { NULL, NULL };
|
||||
LPexp* basics[2] = { nullptr, nullptr };
|
||||
std::list<LPexp*> custom;
|
||||
#define THROWAWAY() \
|
||||
do { \
|
||||
for (int i = 0; i < 2; ++i) \
|
||||
if (basics[i]) \
|
||||
delete basics[i]; \
|
||||
for (std::list<LPexp*>::iterator i = custom.begin(); \
|
||||
i != custom.end(); ++i) \
|
||||
delete *i; \
|
||||
m_ingroup = false; \
|
||||
return NULL; \
|
||||
#define THROWAWAY() \
|
||||
do { \
|
||||
delete basics[0]; \
|
||||
delete basics[1]; \
|
||||
for (auto& p : custom) \
|
||||
delete p; \
|
||||
m_ingroup = false; \
|
||||
return nullptr; \
|
||||
} while (0)
|
||||
|
||||
m_ingroup = true;
|
||||
next();
|
||||
while (hasInput() && head() != '>') {
|
||||
LPexp* e = _parse();
|
||||
if (error())
|
||||
if (error()) {
|
||||
delete e;
|
||||
THROWAWAY();
|
||||
}
|
||||
if (basicpos < 2)
|
||||
basics[basicpos++] = e;
|
||||
else
|
||||
@@ -266,7 +266,7 @@ Parser::parseNegor()
|
||||
bool iscustom = false;
|
||||
ustring sym = parseRawSymbol(iscustom);
|
||||
if (error())
|
||||
return NULL;
|
||||
return nullptr;
|
||||
symlist.insert(sym);
|
||||
if (iscustom) {
|
||||
if (symlist.size() > 2 && pos != -1)
|
||||
@@ -278,7 +278,7 @@ Parser::parseNegor()
|
||||
SymbolToInt::const_iterator found = m_label_position.find(sym);
|
||||
if (found == m_label_position.end()) {
|
||||
m_error = "Unrecognized basic label";
|
||||
return NULL;
|
||||
return nullptr;
|
||||
}
|
||||
if (symlist.size() > 2 && found->second != pos)
|
||||
std::cerr
|
||||
@@ -290,11 +290,11 @@ Parser::parseNegor()
|
||||
if (!hasInput()) {
|
||||
m_error
|
||||
= "Reached end of line looking for ] to end an negative or list'";
|
||||
return NULL;
|
||||
return nullptr;
|
||||
}
|
||||
if (symlist.size() < 2) {
|
||||
m_error = "Empty or list [^] not allowed";
|
||||
return NULL;
|
||||
return nullptr;
|
||||
}
|
||||
next();
|
||||
lpexp::Wildexp* wildcard = new lpexp::Wildexp(symlist);
|
||||
@@ -307,12 +307,9 @@ Parser::parseNegor()
|
||||
return buildStop(new lpexp::Wildexp(m_minus_stop),
|
||||
new lpexp::Wildexp(m_minus_stop), custom);
|
||||
} else {
|
||||
LPexp* basics[2] = { NULL, NULL };
|
||||
basics[pos] = wildcard;
|
||||
for (int i = 0; i < 2; ++i)
|
||||
if (!basics[i])
|
||||
basics[i] = new lpexp::Wildexp(m_minus_stop);
|
||||
return buildStop(basics[0], basics[1], custom);
|
||||
LPexp* w2 = new lpexp::Wildexp(m_minus_stop);
|
||||
return buildStop(pos == 0 ? wildcard : w2, pos == 0 ? w2 : wildcard,
|
||||
custom);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -331,15 +328,16 @@ Parser::parseOrlist()
|
||||
while (hasInput() && head() != ']') {
|
||||
LPexp* e = _parse();
|
||||
if (error()) {
|
||||
delete e;
|
||||
delete orlist;
|
||||
return NULL;
|
||||
return nullptr;
|
||||
}
|
||||
orlist->append(e);
|
||||
}
|
||||
if (!hasInput()) {
|
||||
m_error = "Reached end of line looking for ] to end an or list";
|
||||
delete orlist;
|
||||
return NULL;
|
||||
return nullptr;
|
||||
}
|
||||
next();
|
||||
return orlist;
|
||||
@@ -374,11 +372,11 @@ Parser::parseRange()
|
||||
}
|
||||
next();
|
||||
if (secondnum.size())
|
||||
return std::pair<int, int>(atoi(firstnum.c_str()),
|
||||
atoi(secondnum.c_str()));
|
||||
return std::pair<int, int>(OIIO::Strutil::stoi(firstnum),
|
||||
OIIO::Strutil::stoi(secondnum));
|
||||
else
|
||||
return std::pair<int, int>(atoi(firstnum.c_str()),
|
||||
atoi(firstnum.c_str()));
|
||||
return std::pair<int, int>(OIIO::Strutil::stoi(firstnum),
|
||||
OIIO::Strutil::stoi(firstnum));
|
||||
}
|
||||
|
||||
|
||||
@@ -389,7 +387,7 @@ Parser::parseModifier(LPexp* e)
|
||||
if (hasInput()) {
|
||||
if (m_ingroup && (head() == '*' || head() == '{' || head() == '+')) {
|
||||
m_error = std::string("Repetitions not allowed inside '<...>'");
|
||||
return NULL;
|
||||
return nullptr;
|
||||
}
|
||||
if (head() == '*') {
|
||||
next();
|
||||
@@ -397,7 +395,7 @@ Parser::parseModifier(LPexp* e)
|
||||
} else if (head() == '{') {
|
||||
std::pair<int, int> range = parseRange();
|
||||
if (error())
|
||||
return NULL;
|
||||
return nullptr;
|
||||
if (range.second < 0) {
|
||||
lpexp::Cat* cat = new lpexp::Cat();
|
||||
cat->append(new lpexp::NRepeat(e, range.first, range.first));
|
||||
@@ -422,7 +420,7 @@ Parser::parseModifier(LPexp* e)
|
||||
LPexp*
|
||||
Parser::_parse()
|
||||
{
|
||||
LPexp* e;
|
||||
LPexp* e = nullptr;
|
||||
if (head() == '(')
|
||||
e = parseCat();
|
||||
else if (head() == '[')
|
||||
@@ -431,8 +429,10 @@ Parser::_parse()
|
||||
e = parseGroup();
|
||||
else
|
||||
e = parseSymbol();
|
||||
if (error())
|
||||
return NULL;
|
||||
if (error()) {
|
||||
delete e;
|
||||
return nullptr;
|
||||
}
|
||||
return parseModifier(e);
|
||||
}
|
||||
|
||||
@@ -448,7 +448,7 @@ Parser::parse(const char* text)
|
||||
if (hasInput())
|
||||
return parseCat();
|
||||
else
|
||||
return NULL;
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
|
||||
|
||||
@@ -265,9 +265,11 @@ ColorSystem::error(StringParam src, StringParam dst, Context context) const
|
||||
}
|
||||
|
||||
static inline const ColorSystem&
|
||||
op_color_colorsystem(void* sg)
|
||||
op_color_colorsystem(void* sg_)
|
||||
{
|
||||
return ((ShaderGlobals*)sg)->context->shadingsys().colorsystem();
|
||||
ShaderGlobals* sg = (ShaderGlobals*)sg_;
|
||||
ShadingStateUniform* ssu = (ShadingStateUniform*)sg->shadingStateUniform;
|
||||
return ssu->m_colorsystem;
|
||||
}
|
||||
|
||||
static inline OSL::ShadingContext*
|
||||
|
||||
+15
-12
@@ -138,10 +138,10 @@ osl_transformn_dvmdv(void* result, void* M_, void* v_)
|
||||
OSL_SHADEOP int
|
||||
osl_get_matrix(void* sg_, void* r, const char* from)
|
||||
{
|
||||
ShaderGlobals* sg = (ShaderGlobals*)sg_;
|
||||
ShadingContext* ctx = (ShadingContext*)sg->context;
|
||||
ShaderGlobals* sg = (ShaderGlobals*)sg_;
|
||||
ShadingStateUniform* ssu = (ShadingStateUniform*)sg->shadingStateUniform;
|
||||
if (HDSTR(from) == STRING_PARAMS(common)
|
||||
|| HDSTR(from) == ctx->shadingsys().commonspace_synonym()) {
|
||||
|| HDSTR(from) == ssu->m_commonspace_synonym) {
|
||||
MAT(r).makeIdentity();
|
||||
return true;
|
||||
}
|
||||
@@ -157,7 +157,7 @@ osl_get_matrix(void* sg_, void* r, const char* from)
|
||||
if (!ok) {
|
||||
MAT(r).makeIdentity();
|
||||
ShadingContext* ctx = (ShadingContext*)((ShaderGlobals*)sg)->context;
|
||||
if (ctx->shadingsys().unknown_coordsys_error())
|
||||
if (ssu->m_unknown_coordsys_error)
|
||||
ctx->errorfmt("Unknown transformation \"{}\"", from);
|
||||
}
|
||||
return ok;
|
||||
@@ -168,10 +168,10 @@ osl_get_matrix(void* sg_, void* r, const char* from)
|
||||
OSL_SHADEOP int
|
||||
osl_get_inverse_matrix(void* sg_, void* r, const char* to)
|
||||
{
|
||||
ShaderGlobals* sg = (ShaderGlobals*)sg_;
|
||||
ShadingContext* ctx = (ShadingContext*)sg->context;
|
||||
ShaderGlobals* sg = (ShaderGlobals*)sg_;
|
||||
ShadingStateUniform* ssu = (ShadingStateUniform*)sg->shadingStateUniform;
|
||||
if (HDSTR(to) == STRING_PARAMS(common)
|
||||
|| HDSTR(to) == ctx->shadingsys().commonspace_synonym()) {
|
||||
|| HDSTR(to) == ssu->m_commonspace_synonym) {
|
||||
MAT(r).makeIdentity();
|
||||
return true;
|
||||
}
|
||||
@@ -189,7 +189,7 @@ osl_get_inverse_matrix(void* sg_, void* r, const char* to)
|
||||
if (!ok) {
|
||||
MAT(r).makeIdentity();
|
||||
ShadingContext* ctx = (ShadingContext*)((ShaderGlobals*)sg)->context;
|
||||
if (ctx->shadingsys().unknown_coordsys_error())
|
||||
if (ssu->m_unknown_coordsys_error)
|
||||
ctx->errorfmt("Unknown transformation \"{}\"", to);
|
||||
}
|
||||
return ok;
|
||||
@@ -207,18 +207,21 @@ osl_get_inverse_matrix(void* sg_, void* r, const char* to);
|
||||
|
||||
|
||||
OSL_SHADEOP OSL_HOSTDEVICE int
|
||||
osl_prepend_matrix_from(void* sg, void* r, const char* from)
|
||||
osl_prepend_matrix_from(void* sg_, void* r, const char* from)
|
||||
{
|
||||
ShaderGlobals* sg = (ShaderGlobals*)sg_;
|
||||
Matrix44 m;
|
||||
bool ok = osl_get_matrix((ShaderGlobals*)sg, &m, from);
|
||||
bool ok = osl_get_matrix(sg, &m, from);
|
||||
if (ok)
|
||||
MAT(r) = m * MAT(r);
|
||||
#ifndef __CUDACC__
|
||||
// TODO: How do we manage this in OptiX?
|
||||
else {
|
||||
ShadingContext* ctx = (ShadingContext*)((ShaderGlobals*)sg)->context;
|
||||
if (ctx->shadingsys().unknown_coordsys_error())
|
||||
ShadingStateUniform* ssu = (ShadingStateUniform*)sg->shadingStateUniform;
|
||||
if (ssu->m_unknown_coordsys_error) {
|
||||
ShadingContext* ctx = (ShadingContext*)(sg)->context;
|
||||
ctx->errorfmt("Unknown transformation \"{}\"", from);
|
||||
}
|
||||
}
|
||||
#endif
|
||||
return ok;
|
||||
|
||||
@@ -33,6 +33,7 @@
|
||||
#include <OpenImageIO/ustring.h>
|
||||
|
||||
#include "osl_pvt.h"
|
||||
|
||||
#include <OSL/dual.h>
|
||||
#include <OSL/dual_vec.h>
|
||||
#include <OSL/genclosure.h>
|
||||
@@ -42,6 +43,8 @@
|
||||
#include <OSL/oslexec.h>
|
||||
#include <OSL/rendererservices.h>
|
||||
#include <OSL/shaderglobals.h>
|
||||
|
||||
#include "shading_state_uniform.h"
|
||||
#include "constantpool.h"
|
||||
#include "opcolor.h"
|
||||
|
||||
@@ -57,7 +60,6 @@ using OIIO::ParamValueList;
|
||||
using OIIO::RefCnt;
|
||||
using OIIO::spin_lock;
|
||||
using OIIO::spin_mutex;
|
||||
using OIIO::ustringHash;
|
||||
namespace Strutil = OIIO::Strutil;
|
||||
|
||||
|
||||
@@ -615,7 +617,10 @@ public:
|
||||
bool lockgeom_default() const { return m_lockgeom_default; }
|
||||
bool strict_messages() const { return m_strict_messages; }
|
||||
bool range_checking() const { return m_range_checking; }
|
||||
bool unknown_coordsys_error() const { return m_unknown_coordsys_error; }
|
||||
bool unknown_coordsys_error() const
|
||||
{
|
||||
return m_shading_state_uniform.m_unknown_coordsys_error;
|
||||
}
|
||||
bool connection_error() const { return m_connection_error; }
|
||||
bool relaxed_param_typecheck() const { return m_relaxed_param_typecheck; }
|
||||
int optimize() const { return m_optimize; }
|
||||
@@ -640,7 +645,10 @@ public:
|
||||
bool no_pointcloud() const { return m_no_pointcloud; }
|
||||
bool force_derivs() const { return m_force_derivs; }
|
||||
bool allow_shader_replacement() const { return m_allow_shader_replacement; }
|
||||
ustring commonspace_synonym() const { return m_commonspace_synonym; }
|
||||
ustring commonspace_synonym() const
|
||||
{
|
||||
return m_shading_state_uniform.m_commonspace_synonym;
|
||||
}
|
||||
|
||||
bool llvm_jit_fma() const { return m_llvm_jit_fma; }
|
||||
ustring llvm_jit_target() const { return m_llvm_jit_target; }
|
||||
@@ -693,8 +701,7 @@ public:
|
||||
void optimize_all_groups(int nthreads = 0, int mythread = 0,
|
||||
int totalthreads = 1, bool do_jit = true);
|
||||
|
||||
typedef std::unordered_map<ustring, OpDescriptor, ustringHash>
|
||||
OpDescriptorMap;
|
||||
typedef std::unordered_map<ustring, OpDescriptor> OpDescriptorMap;
|
||||
|
||||
/// Look up OpDescriptor for the named op, return NULL for unknown op.
|
||||
///
|
||||
@@ -719,7 +726,7 @@ public:
|
||||
|
||||
void count_noise(int number = 1) { m_stat_noise_calls += number; }
|
||||
|
||||
ColorSystem& colorsystem() { return m_colorsystem; }
|
||||
ColorSystem& colorsystem() { return m_shading_state_uniform.m_colorsystem; }
|
||||
|
||||
std::shared_ptr<OIIO::ColorConfig> colorconfig();
|
||||
|
||||
@@ -833,10 +840,9 @@ private:
|
||||
bool m_strict_messages; ///< Strict checking of message passing usage?
|
||||
bool m_error_repeats; ///< Allow repeats of identical err/warn?
|
||||
bool m_range_checking; ///< Range check arrays & components?
|
||||
bool m_unknown_coordsys_error; ///< Error to use unknown xform name?
|
||||
bool m_connection_error; ///< Error for ConnectShaders to fail?
|
||||
bool m_greedyjit; ///< JIT as much as we can?
|
||||
bool m_countlayerexecs; ///< Count number of layer execs?
|
||||
bool m_connection_error; ///< Error for ConnectShaders to fail?
|
||||
bool m_greedyjit; ///< JIT as much as we can?
|
||||
bool m_countlayerexecs; ///< Count number of layer execs?
|
||||
bool m_relaxed_param_typecheck; ///< Allow parameters to be set from isomorphic types (same data layout)
|
||||
int m_max_warnings_per_thread; ///< How many warnings to display per thread before giving up?
|
||||
int m_profile; ///< Level of profiling of shader execution
|
||||
@@ -885,7 +891,6 @@ private:
|
||||
std::string m_library_searchpath; ///< Library search path
|
||||
std::vector<std::string>
|
||||
m_library_searchpath_dirs; ///< All library searchpath dirs
|
||||
ustring m_commonspace_synonym; ///< Synonym for "common" space
|
||||
std::vector<ustring> m_raytypes; ///< Names of ray types
|
||||
std::vector<ustring> m_renderer_outputs; ///< Names of renderer outputs
|
||||
std::vector<SymLocationDesc> m_symlocs;
|
||||
@@ -901,8 +906,10 @@ private:
|
||||
int m_gpu_opt_error; ///< Error on inability to optimize
|
||||
///< away things that can't GPU.
|
||||
|
||||
ustring m_colorspace; ///< What RGB colors mean
|
||||
ColorSystem m_colorsystem; ///< Data for current colorspace
|
||||
ustring m_colorspace; ///< What RGB colors mean
|
||||
|
||||
ShadingStateUniform m_shading_state_uniform;
|
||||
|
||||
std::shared_ptr<OIIO::ColorConfig>
|
||||
m_colorconfig; ///< OIIO/OCIO color configuration
|
||||
|
||||
@@ -2352,8 +2359,7 @@ private:
|
||||
nullptr, &OIIO::aligned_free
|
||||
};
|
||||
size_t m_heapsize = 0;
|
||||
using RegexMap
|
||||
= std::unordered_map<ustring, std::unique_ptr<std::regex>, ustringHash>;
|
||||
using RegexMap = std::unordered_map<ustring, std::unique_ptr<std::regex>>;
|
||||
RegexMap m_regex_map; ///< Compiled regex's
|
||||
MessageList m_messages; ///< Message blackboard
|
||||
#if OSL_USE_BATCHED
|
||||
|
||||
@@ -27,6 +27,10 @@
|
||||
#pragma clang diagnostic ignored "-Wparentheses-equality"
|
||||
#endif
|
||||
|
||||
#if OSL_CLANG_VERSION >= 150000
|
||||
#pragma clang diagnostic ignored "-Wunused-but-set-variable"
|
||||
#endif
|
||||
|
||||
#if defined(__GNUC__) || defined(__clang__)
|
||||
#pragma GCC visibility push(hidden)
|
||||
#endif
|
||||
|
||||
@@ -12,8 +12,7 @@ OSL_NAMESPACE_ENTER
|
||||
namespace pvt {
|
||||
|
||||
#ifdef USE_PARTIO
|
||||
typedef std::unordered_map<ustring, std::unique_ptr<PointCloud>, ustringHash>
|
||||
PointCloudMap;
|
||||
typedef std::unordered_map<ustring, std::unique_ptr<PointCloud>> PointCloudMap;
|
||||
static PointCloudMap pointclouds;
|
||||
static OIIO::spin_mutex pointcloudmap_mutex;
|
||||
|
||||
@@ -255,7 +254,7 @@ RendererServices::pointcloud_get(ShaderGlobals* sg, ustring filename,
|
||||
// Actual data query
|
||||
if (partio_type == TypeString) {
|
||||
// strings are special cases because they are stored as int index
|
||||
int* strindices = OIIO_ALLOCA(int, count);
|
||||
int* strindices = OSL_ALLOCA(int, count);
|
||||
const_cast<Partio::ParticlesData*>(cloud)->data(
|
||||
*attr, count, (const Partio::ParticleIndex*)indices,
|
||||
/*sorted=*/false, (void*)strindices);
|
||||
@@ -385,7 +384,7 @@ osl_pointcloud_search(ShaderGlobals* sg, const char* filename, void* center,
|
||||
if (sizeof(int) == sizeof(size_t) && out_indices)
|
||||
indices = (size_t*)out_indices;
|
||||
else
|
||||
indices = OIIO_ALLOCA(size_t, max_points);
|
||||
indices = OSL_ALLOCA(size_t, max_points);
|
||||
|
||||
int count
|
||||
= sg->renderer->pointcloud_search(sg, USTR(filename), *((Vec3*)center),
|
||||
@@ -425,7 +424,7 @@ osl_pointcloud_get(ShaderGlobals* sg, const char* filename, void* in_indices,
|
||||
if (shadingsys.no_pointcloud()) // Debug mode to skip pointcloud expense
|
||||
return 0;
|
||||
|
||||
size_t* indices = OIIO_ALLOCA(size_t, count);
|
||||
size_t* indices = OSL_ALLOCA(size_t, count);
|
||||
for (int i = 0; i < count; ++i)
|
||||
indices[i] = ((int*)in_indices)[i];
|
||||
|
||||
|
||||
@@ -22,8 +22,8 @@ public:
|
||||
~PointCloud();
|
||||
static PointCloud* get(ustring filename, bool write = false);
|
||||
|
||||
typedef std::unordered_map<
|
||||
ustring, std::unique_ptr<Partio::ParticleAttribute>, ustringHash>
|
||||
typedef std::unordered_map<ustring,
|
||||
std::unique_ptr<Partio::ParticleAttribute>>
|
||||
AttributeMap;
|
||||
|
||||
const Partio::ParticlesData* read_access() const
|
||||
|
||||
@@ -114,14 +114,11 @@ RendererServices::good(TextureHandle* texture_handle)
|
||||
bool
|
||||
RendererServices::is_udim(TextureHandle* texture_handle)
|
||||
{
|
||||
#if OPENIMAGEIO_VERSION >= 20307
|
||||
return texturesys()->is_udim(texture_handle);
|
||||
#else
|
||||
return false;
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
|
||||
bool
|
||||
RendererServices::texture(ustring filename, TextureHandle* texture_handle,
|
||||
TexturePerthread* texture_thread_info,
|
||||
|
||||
@@ -481,7 +481,7 @@ private:
|
||||
ShaderGlobals m_shaderglobals; ///< Dummy ShaderGlobals
|
||||
|
||||
// Keep track of some things for the whole shader group:
|
||||
typedef std::unordered_map<ustring, ustring, ustringHash> ustringmap_t;
|
||||
typedef std::unordered_map<ustring, ustring> ustringmap_t;
|
||||
std::vector<ustringmap_t> m_params_holding_globals;
|
||||
///< Which params of each layer really just hold globals
|
||||
|
||||
|
||||
@@ -66,10 +66,10 @@ shade_image(ShadingSystem& shadingsys, ShaderGroup& group,
|
||||
|
||||
// Gather some information about the outputs once, rather than for
|
||||
// each pixel.
|
||||
const ShaderSymbol** output_sym = OIIO_ALLOCA(const ShaderSymbol*,
|
||||
outputs.size());
|
||||
TypeDesc* output_type = OIIO_ALLOCA(TypeDesc, outputs.size());
|
||||
int* output_nchans = OIIO_ALLOCA(int, outputs.size());
|
||||
const ShaderSymbol** output_sym = OSL_ALLOCA(const ShaderSymbol*,
|
||||
outputs.size());
|
||||
TypeDesc* output_type = OSL_ALLOCA(TypeDesc, outputs.size());
|
||||
int* output_nchans = OSL_ALLOCA(int, outputs.size());
|
||||
for (int i = 0; i < int(outputs.size()); ++i) {
|
||||
output_sym[i] = shadingsys.find_symbol(group, outputs[i]);
|
||||
output_type[i] = shadingsys.symbol_typedesc(output_sym[i]);
|
||||
|
||||
@@ -0,0 +1,26 @@
|
||||
// Copyright Contributors to the Open Shading Language project.
|
||||
// SPDX-License-Identifier: BSD-3-Clause
|
||||
// https://github.com/AcademySoftwareFoundation/OpenShadingLanguage
|
||||
|
||||
#include "opcolor.h"
|
||||
|
||||
OSL_NAMESPACE_ENTER
|
||||
|
||||
namespace pvt {
|
||||
|
||||
// Intent is for a device to export or serialize the ShadingStateUniform from the
|
||||
// ShadingSystem into a device side buffer that can be passed into
|
||||
// kernel launches.
|
||||
|
||||
// Requires data to be self-contained and flat in memory to be device friendly.
|
||||
|
||||
|
||||
struct ShadingStateUniform {
|
||||
ColorSystem m_colorsystem; ///< Data for current colorspace
|
||||
ustring m_commonspace_synonym; ///< Synonym for "common" space
|
||||
bool m_unknown_coordsys_error; ///< Error to use unknown xform name?
|
||||
};
|
||||
|
||||
} // namespace pvt
|
||||
|
||||
OSL_NAMESPACE_EXIT
|
||||
@@ -1002,7 +1002,6 @@ ShadingSystemImpl::ShadingSystemImpl(RendererServices* renderer,
|
||||
, m_strict_messages(true)
|
||||
, m_error_repeats(false)
|
||||
, m_range_checking(true)
|
||||
, m_unknown_coordsys_error(true)
|
||||
, m_connection_error(true)
|
||||
, m_greedyjit(false)
|
||||
, m_countlayerexecs(false)
|
||||
@@ -1049,7 +1048,6 @@ ShadingSystemImpl::ShadingSystemImpl(RendererServices* renderer,
|
||||
, m_llvm_profiling_events(0)
|
||||
, m_llvm_output_bitcode(0)
|
||||
, m_llvm_dumpasm(0)
|
||||
, m_commonspace_synonym("world")
|
||||
, m_max_local_mem_KB(2048)
|
||||
, m_compile_report(false)
|
||||
, m_buffer_printf(true)
|
||||
@@ -1071,6 +1069,9 @@ ShadingSystemImpl::ShadingSystemImpl(RendererServices* renderer,
|
||||
, m_stat_inst_merge_time(0)
|
||||
, m_stat_max_llvm_local_mem(0)
|
||||
{
|
||||
m_shading_state_uniform.m_commonspace_synonym = ustring("world");
|
||||
m_shading_state_uniform.m_unknown_coordsys_error = true;
|
||||
|
||||
m_stat_shaders_loaded = 0;
|
||||
m_stat_shaders_requested = 0;
|
||||
m_stat_groups = 0;
|
||||
@@ -1551,7 +1552,8 @@ ShadingSystemImpl::attribute(string_view name, TypeDesc type, const void* val)
|
||||
ATTR_SET_STRING("llvm_prune_ir_strategy", m_llvm_prune_ir_strategy);
|
||||
ATTR_SET("strict_messages", int, m_strict_messages);
|
||||
ATTR_SET("range_checking", int, m_range_checking);
|
||||
ATTR_SET("unknown_coordsys_error", int, m_unknown_coordsys_error);
|
||||
ATTR_SET("unknown_coordsys_error", int,
|
||||
m_shading_state_uniform.m_unknown_coordsys_error);
|
||||
ATTR_SET("connection_error", int, m_connection_error);
|
||||
ATTR_SET("greedyjit", int, m_greedyjit);
|
||||
ATTR_SET("relaxed_param_typecheck", int, m_relaxed_param_typecheck);
|
||||
@@ -1567,7 +1569,8 @@ ShadingSystemImpl::attribute(string_view name, TypeDesc type, const void* val)
|
||||
ATTR_SET("exec_repeat", int, m_exec_repeat);
|
||||
ATTR_SET("opt_warnings", int, m_opt_warnings);
|
||||
ATTR_SET("gpu_opt_error", int, m_gpu_opt_error);
|
||||
ATTR_SET_STRING("commonspace", m_commonspace_synonym);
|
||||
ATTR_SET_STRING("commonspace",
|
||||
m_shading_state_uniform.m_commonspace_synonym);
|
||||
ATTR_SET_STRING("debug_groupname", m_debug_groupname);
|
||||
ATTR_SET_STRING("debug_layername", m_debug_layername);
|
||||
ATTR_SET_STRING("opt_layername", m_opt_layername);
|
||||
@@ -1718,13 +1721,15 @@ ShadingSystemImpl::getattribute(string_view name, TypeDesc type, void* val)
|
||||
ATTR_DECODE("strict_messages", int, m_strict_messages);
|
||||
ATTR_DECODE("error_repeats", int, m_error_repeats);
|
||||
ATTR_DECODE("range_checking", int, m_range_checking);
|
||||
ATTR_DECODE("unknown_coordsys_error", int, m_unknown_coordsys_error);
|
||||
ATTR_DECODE("unknown_coordsys_error", int,
|
||||
m_shading_state_uniform.m_unknown_coordsys_error);
|
||||
ATTR_DECODE("connection_error", int, m_connection_error);
|
||||
ATTR_DECODE("greedyjit", int, m_greedyjit);
|
||||
ATTR_DECODE("countlayerexecs", int, m_countlayerexecs);
|
||||
ATTR_DECODE("relaxed_param_typecheck", int, m_relaxed_param_typecheck);
|
||||
ATTR_DECODE("max_warnings_per_thread", int, m_max_warnings_per_thread);
|
||||
ATTR_DECODE_STRING("commonspace", m_commonspace_synonym);
|
||||
ATTR_DECODE_STRING("commonspace",
|
||||
m_shading_state_uniform.m_commonspace_synonym);
|
||||
ATTR_DECODE_STRING("colorspace", m_colorspace);
|
||||
ATTR_DECODE_STRING("debug_groupname", m_debug_groupname);
|
||||
ATTR_DECODE_STRING("debug_layername", m_debug_layername);
|
||||
|
||||
@@ -60,7 +60,7 @@ invert_wide_matrix(Masked<Matrix44> wresult, Wide<const Matrix44> wmatrix)
|
||||
bool is_affine = true;
|
||||
if (wresult.mask()[lane]) {
|
||||
is_affine = test_if_affine(m);
|
||||
if (is_affine) {
|
||||
if (OSL_UNLIKELY(is_affine)) {
|
||||
Matrix44 r = OSL::affineInverse(m);
|
||||
wresult[ActiveLane(lane)] = r;
|
||||
}
|
||||
@@ -338,7 +338,7 @@ __OSL_MASKED_OP3(div, Wm, Wm, Wm)(void* wr_, void* wa_, void* wb_,
|
||||
bool is_affine = true;
|
||||
if (wresult.mask()[lane]) {
|
||||
is_affine = test_if_affine(b);
|
||||
if (is_affine) {
|
||||
if (OSL_UNLIKELY(is_affine)) {
|
||||
wresult[ActiveLane(lane)]
|
||||
= multiplyMatrixByMatrix(a, OSL::affineInverse(b));
|
||||
}
|
||||
@@ -434,7 +434,7 @@ __OSL_MASKED_OP3(div, Wm, Wf, Wm)(void* wr_, void* wa_, void* wb_,
|
||||
bool is_affine = true;
|
||||
if (wresult.mask()[lane]) {
|
||||
is_affine = test_if_affine(b);
|
||||
if (is_affine) {
|
||||
if (OSL_UNLIKELY(is_affine)) {
|
||||
Matrix44 r = a * OSL::affineInverse(b);
|
||||
|
||||
wresult[ActiveLane(lane)] = r;
|
||||
@@ -633,44 +633,56 @@ impl_get_varying_from_matrix_batched(BatchedShaderGlobals* bsg,
|
||||
// Deal with a varying 'from' space
|
||||
ustring commonspace_synonym = ctx->shadingsys().commonspace_synonym();
|
||||
|
||||
Mask commonSpaceMask(false);
|
||||
Mask shaderSpaceMask(false);
|
||||
Mask objectSpaceMask(false);
|
||||
Mask namedSpaceMask(false);
|
||||
// Use int instead of Mask<> to allow reduction clause in openmp simd declaration
|
||||
int common_space_bits { 0 };
|
||||
int shader_space_bits { 0 };
|
||||
int object_space_bits { 0 };
|
||||
int named_space_bits { 0 };
|
||||
|
||||
OSL_FORCEINLINE_BLOCK
|
||||
{
|
||||
OSL_OMP_PRAGMA(omp simd simdlen(__OSL_WIDTH))
|
||||
OSL_OMP_PRAGMA(omp simd simdlen(__OSL_WIDTH)
|
||||
reduction(|
|
||||
: common_space_bits, shader_space_bits,
|
||||
object_space_bits, named_space_bits))
|
||||
for (int lane = 0; lane < __OSL_WIDTH; ++lane) {
|
||||
ustring from = wFrom[lane];
|
||||
if (wMfrom.mask()[lane]) {
|
||||
if (from == Strings::common || from == commonspace_synonym) {
|
||||
commonSpaceMask.set_on(lane);
|
||||
// inline of Mask::set_on(lane)
|
||||
common_space_bits |= 1 << lane;
|
||||
} else if (from == Strings::shader) {
|
||||
shaderSpaceMask.set_on(lane);
|
||||
// inline of Mask::set_on(lane)
|
||||
shader_space_bits |= 1 << lane;
|
||||
} else if (from == Strings::object) {
|
||||
objectSpaceMask.set_on(lane);
|
||||
// inline of Mask::set_on(lane)
|
||||
object_space_bits |= 1 << lane;
|
||||
} else {
|
||||
namedSpaceMask.set_on(lane);
|
||||
// inline of Mask::set_on(lane)
|
||||
named_space_bits |= 1 << lane;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
Mask common_space_mask(common_space_bits);
|
||||
Mask shader_space_mask(shader_space_bits);
|
||||
Mask object_space_mask(object_space_bits);
|
||||
Mask named_space_mask(named_space_bits);
|
||||
|
||||
if (commonSpaceMask.any_on()) {
|
||||
Masked<Matrix44> mfrom(wMfrom.data(), commonSpaceMask);
|
||||
if (common_space_mask.any_on()) {
|
||||
Masked<Matrix44> mfrom(wMfrom.data(), common_space_mask);
|
||||
makeIdentity(mfrom);
|
||||
}
|
||||
const auto& sgbv = bsg->varying;
|
||||
if (shaderSpaceMask.any_on()) {
|
||||
Masked<Matrix44> mfrom(wMfrom.data(), shaderSpaceMask);
|
||||
if (shader_space_mask.any_on()) {
|
||||
Masked<Matrix44> mfrom(wMfrom.data(), shader_space_mask);
|
||||
ctx->batched<__OSL_WIDTH>().renderer()->get_matrix(bsg, mfrom,
|
||||
sgbv.shader2common,
|
||||
sgbv.time);
|
||||
// NOTE: matching scalar version of code which ignores the renderservices return value
|
||||
}
|
||||
if (objectSpaceMask.any_on()) {
|
||||
Masked<Matrix44> mfrom(wMfrom.data(), objectSpaceMask);
|
||||
if (object_space_mask.any_on()) {
|
||||
Masked<Matrix44> mfrom(wMfrom.data(), object_space_mask);
|
||||
ctx->batched<__OSL_WIDTH>().renderer()->get_matrix(bsg, mfrom,
|
||||
sgbv.object2common,
|
||||
sgbv.time);
|
||||
@@ -678,13 +690,13 @@ impl_get_varying_from_matrix_batched(BatchedShaderGlobals* bsg,
|
||||
}
|
||||
// Only named lookups can fail, so we can just subtract those lanes
|
||||
Mask succeeded(wMfrom.mask());
|
||||
if (namedSpaceMask.any_on()) {
|
||||
Masked<Matrix44> mfrom(wMfrom.data(), namedSpaceMask);
|
||||
if (named_space_mask.any_on()) {
|
||||
Masked<Matrix44> mfrom(wMfrom.data(), named_space_mask);
|
||||
|
||||
Mask success = ctx->batched<__OSL_WIDTH>().renderer()->get_matrix(
|
||||
bsg, mfrom, wFrom, sgbv.time);
|
||||
|
||||
Mask failedLanes = success.invert() & namedSpaceMask;
|
||||
Mask failedLanes = success.invert() & named_space_mask;
|
||||
if (failedLanes.any_on()) {
|
||||
Masked<Matrix44> mto_failed(wMfrom.data(), failedLanes);
|
||||
makeIdentity(mto_failed);
|
||||
@@ -756,56 +768,68 @@ impl_get_varying_to_matrix_masked(BatchedShaderGlobals* bsg,
|
||||
// Deal with a varying 'to' space
|
||||
ustring commonspace_synonym = ctx->shadingsys().commonspace_synonym();
|
||||
|
||||
Mask commonSpaceMask(false);
|
||||
Mask shaderSpaceMask(false);
|
||||
Mask objectSpaceMask(false);
|
||||
Mask namedSpaceMask(false);
|
||||
// Use int instead of Mask<> to allow reduction clause in openmp simd declaration
|
||||
int common_space_bits { 0 };
|
||||
int shader_space_bits { 0 };
|
||||
int object_space_bits { 0 };
|
||||
int named_space_bits { 0 };
|
||||
|
||||
OSL_FORCEINLINE_BLOCK
|
||||
{
|
||||
OSL_OMP_PRAGMA(omp simd simdlen(__OSL_WIDTH))
|
||||
OSL_OMP_PRAGMA(omp simd simdlen(__OSL_WIDTH)
|
||||
reduction(|
|
||||
: common_space_bits, shader_space_bits,
|
||||
object_space_bits, named_space_bits))
|
||||
for (int lane = 0; lane < __OSL_WIDTH; ++lane) {
|
||||
ustring to = wTo[lane];
|
||||
if (wMto.mask()[lane]) {
|
||||
if (to == Strings::common || to == commonspace_synonym) {
|
||||
commonSpaceMask.set_on(lane);
|
||||
// inline of Mask::set_on(lane)
|
||||
common_space_bits |= 1 << lane;
|
||||
} else if (to == Strings::shader) {
|
||||
shaderSpaceMask.set_on(lane);
|
||||
// inline of Mask::set_on(lane)
|
||||
shader_space_bits |= 1 << lane;
|
||||
} else if (to == Strings::object) {
|
||||
objectSpaceMask.set_on(lane);
|
||||
// inline of Mask::set_on(lane)
|
||||
object_space_bits |= 1 << lane;
|
||||
} else {
|
||||
namedSpaceMask.set_on(lane);
|
||||
// inline of Mask::set_on(lane)
|
||||
named_space_bits |= 1 << lane;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
Mask common_space_mask(common_space_bits);
|
||||
Mask shader_space_mask(shader_space_bits);
|
||||
Mask object_space_mask(object_space_bits);
|
||||
Mask named_space_mask(named_space_bits);
|
||||
|
||||
if (commonSpaceMask.any_on()) {
|
||||
Masked<Matrix44> mto(wMto.data(), commonSpaceMask);
|
||||
if (common_space_mask.any_on()) {
|
||||
Masked<Matrix44> mto(wMto.data(), common_space_mask);
|
||||
makeIdentity(mto);
|
||||
}
|
||||
const auto& sgbv = bsg->varying;
|
||||
if (shaderSpaceMask.any_on()) {
|
||||
Masked<Matrix44> mto(wMto.data(), shaderSpaceMask);
|
||||
if (shader_space_mask.any_on()) {
|
||||
Masked<Matrix44> mto(wMto.data(), shader_space_mask);
|
||||
dispatch_get_inverse_matrix(ctx->batched<__OSL_WIDTH>().renderer(), bsg,
|
||||
mto, sgbv.shader2common, sgbv.time);
|
||||
// NOTE: matching scalar version of code which ignores the renderservices return value
|
||||
}
|
||||
if (objectSpaceMask.any_on()) {
|
||||
Masked<Matrix44> mto(wMto.data(), objectSpaceMask);
|
||||
if (object_space_mask.any_on()) {
|
||||
Masked<Matrix44> mto(wMto.data(), object_space_mask);
|
||||
dispatch_get_inverse_matrix(ctx->batched<__OSL_WIDTH>().renderer(), bsg,
|
||||
mto, sgbv.object2common, sgbv.time);
|
||||
// NOTE: matching scalar version of code which ignores the renderservices return value
|
||||
}
|
||||
// Only named lookups can fail, so we can just subtract those lanes
|
||||
Mask succeeded(wMto.mask());
|
||||
if (namedSpaceMask.any_on()) {
|
||||
Masked<Matrix44> mto(wMto.data(), namedSpaceMask);
|
||||
if (named_space_mask.any_on()) {
|
||||
Masked<Matrix44> mto(wMto.data(), named_space_mask);
|
||||
|
||||
Mask success = dispatch_get_inverse_matrix(
|
||||
ctx->batched<__OSL_WIDTH>().renderer(), bsg, mto, wTo, sgbv.time);
|
||||
|
||||
Mask failedLanes = success.invert() & namedSpaceMask;
|
||||
Mask failedLanes = success.invert() & named_space_mask;
|
||||
if (failedLanes.any_on()) {
|
||||
Masked<Matrix44> mto(wMto.data(), failedLanes);
|
||||
makeIdentity(mto);
|
||||
|
||||
@@ -351,15 +351,20 @@ __OSL_MASKED_OP(getmessage)(void* bsg_, void* result, char* source_,
|
||||
auto found_lanes = mask & m->valid_mask;
|
||||
auto missing_lanes = mask & ~m->valid_mask & ~m->get_before_set_mask;
|
||||
|
||||
Mask lanes_set_by_deeper_layer(false);
|
||||
OSL_OMP_PRAGMA(omp simd simdlen(__OSL_WIDTH))
|
||||
// Use int instead of Mask<> to allow reduction clause in openmp simd declaration
|
||||
int lanes_set_by_deeper_layer_bits { 0 };
|
||||
OSL_OMP_PRAGMA(omp simd simdlen(__OSL_WIDTH)
|
||||
reduction(|
|
||||
: lanes_set_by_deeper_layer_bits))
|
||||
for (int lane = 0; lane < __OSL_WIDTH; ++lane) {
|
||||
int msg_layerid = msg_wlayeridx[lane];
|
||||
// NOTE: using bitwise & to avoid branches
|
||||
if (found_lanes[lane] & (msg_layerid > layeridx)) {
|
||||
lanes_set_by_deeper_layer.set_on(lane);
|
||||
// inline of Mask::set_on(lane)
|
||||
lanes_set_by_deeper_layer_bits |= 1 << lane;
|
||||
}
|
||||
}
|
||||
Mask lanes_set_by_deeper_layer(lanes_set_by_deeper_layer_bits);
|
||||
|
||||
lanes_set_by_deeper_layer.foreach ([=](ActiveLane lane) -> void {
|
||||
int msg_layerid = msg_wlayeridx[lane];
|
||||
|
||||
@@ -94,7 +94,7 @@ default_pointcloud_search(BatchedShaderGlobals* bsg, ustring filename,
|
||||
// indices array
|
||||
// TODO: evaluate if sg->context->alloc_scratch should be used instead?
|
||||
Partio::ParticleIndex* indices
|
||||
= (Partio::ParticleIndex*)OIIO_ALLOCA(size_t, max_points);
|
||||
= (Partio::ParticleIndex*)OSL_ALLOCA(size_t, max_points);
|
||||
|
||||
Wide<const OSL::Vec3> wcenter(wcenter_);
|
||||
|
||||
@@ -102,8 +102,8 @@ default_pointcloud_search(BatchedShaderGlobals* bsg, ustring filename,
|
||||
// and our batched representation is
|
||||
// structure of arrays (wide) so we need a scalar temporary
|
||||
// distances array
|
||||
float* dist2 = OIIO_ALLOCA(float, max_points);
|
||||
SortedPointRecord* sorted = OIIO_ALLOCA(SortedPointRecord, max_points);
|
||||
float* dist2 = OSL_ALLOCA(float, max_points);
|
||||
SortedPointRecord* sorted = OSL_ALLOCA(SortedPointRecord, max_points);
|
||||
auto windices = results.windices();
|
||||
auto wnum_points = results.wnum_points();
|
||||
results.mask().foreach ([=](ActiveLane lane) -> void {
|
||||
@@ -118,7 +118,7 @@ default_pointcloud_search(BatchedShaderGlobals* bsg, ustring filename,
|
||||
// indices at the same time.
|
||||
if (sort && count > 1) {
|
||||
//SortedPointRecord *sorted = (SortedPointRecord *) sg->context->alloc_scratch (count * sizeof(SortedPointRecord), sizeof(SortedPointRecord));
|
||||
//SortedPointRecord *sorted = OIIO_ALLOCA(SortedPointRecord, count);
|
||||
//SortedPointRecord *sorted = OSL_ALLOCA(SortedPointRecord, count);
|
||||
for (int i = 0; i < count; ++i)
|
||||
sorted[i] = SortedPointRecord(dist2[i], indices[i]);
|
||||
std::sort(sorted, sorted + count, SortedPointCompare());
|
||||
@@ -154,7 +154,7 @@ default_pointcloud_search(BatchedShaderGlobals* bsg, ustring filename,
|
||||
// We are going to need the positions if we need to compute
|
||||
// distance derivs
|
||||
//OSL::Vec3 *positions = (OSL::Vec3 *) sg->context->alloc_scratch (sizeof(OSL::Vec3) * count, sizeof(float));
|
||||
OSL::Vec3* positions = OIIO_ALLOCA(OSL::Vec3, count);
|
||||
OSL::Vec3* positions = OSL_ALLOCA(OSL::Vec3, count);
|
||||
// FIXME(Partio): this function really should be marked as const because it is just a wrapper of a private const method
|
||||
const_cast<Partio::ParticlesData*>(cloud)->data(
|
||||
*pos_attr, count, indices, true, (void*)positions);
|
||||
@@ -255,14 +255,14 @@ default_pointcloud_get(BatchedShaderGlobals* bsg, ustring filename,
|
||||
if (partio_type == OIIO::TypeString) {
|
||||
// strings are special cases because they are stored as int index
|
||||
// Ensure alloca's happen outside loops
|
||||
strindices = OIIO_ALLOCA(int, attr_type.numelements());
|
||||
strindices = OSL_ALLOCA(int, attr_type.numelements());
|
||||
} else {
|
||||
aos_buffer = OIIO_ALLOCA(char, attr_type.size());
|
||||
aos_buffer = OSL_ALLOCA(char, attr_type.size());
|
||||
}
|
||||
}
|
||||
|
||||
Partio::ParticleIndex* indices
|
||||
= (Partio::ParticleIndex*)OIIO_ALLOCA(size_t, windices.length());
|
||||
= (Partio::ParticleIndex*)OSL_ALLOCA(size_t, windices.length());
|
||||
|
||||
wout_data.mask().foreach ([=, &success](ActiveLane lane) -> void {
|
||||
int count = wnum_points[lane];
|
||||
|
||||
@@ -115,7 +115,7 @@ private:
|
||||
typedef bool (OSLToyRenderer::*AttrGetter)(ShaderGlobals* sg, bool derivs,
|
||||
ustring object, TypeDesc type,
|
||||
ustring name, void* val);
|
||||
typedef std::unordered_map<ustring, AttrGetter, ustringHash> AttrGetterMap;
|
||||
typedef std::unordered_map<ustring, AttrGetter> AttrGetterMap;
|
||||
AttrGetterMap m_attr_getters;
|
||||
|
||||
// Attribute getters
|
||||
|
||||
@@ -399,8 +399,8 @@ osl_texture(void* sg_, const char* name, void* handle, void* opt_, float s,
|
||||
{
|
||||
if (!handle)
|
||||
return 0;
|
||||
cudaTextureObject_t texID = cudaTextureObject_t(handle);
|
||||
float4 fromTexture = osl_tex2DLookup(handle, s, t);
|
||||
// cudaTextureObject_t texID = cudaTextureObject_t(handle);
|
||||
float4 fromTexture = osl_tex2DLookup(handle, s, t);
|
||||
// see note above
|
||||
// float4 fromTexture = tex2D<float4>(texID, s, t);
|
||||
*((float3*)result) = make_float3(fromTexture.x, fromTexture.y,
|
||||
|
||||
@@ -98,6 +98,7 @@ struct ShaderGlobals {
|
||||
float3 dPdtime;
|
||||
float3 Ps, dPsdx, dPsdy;
|
||||
void* renderstate;
|
||||
void* shadingStateUniform;
|
||||
void* tracedata;
|
||||
void* objdata;
|
||||
void* context;
|
||||
|
||||
@@ -93,9 +93,8 @@ private:
|
||||
OptixProgramGroup* pg);
|
||||
|
||||
std::string m_materials_ptx;
|
||||
std::unordered_map<OIIO::ustring, optix::TextureSampler, OIIO::ustringHash>
|
||||
m_samplers;
|
||||
std::unordered_map<OIIO::ustring, uint64_t, OIIO::ustringHash> m_globals_map;
|
||||
std::unordered_map<OIIO::ustring, optix::TextureSampler> m_samplers;
|
||||
std::unordered_map<OIIO::ustring, uint64_t> m_globals_map;
|
||||
};
|
||||
|
||||
|
||||
|
||||
@@ -132,7 +132,7 @@ private:
|
||||
typedef bool (SimpleRaytracer::*AttrGetter)(ShaderGlobals* sg, bool derivs,
|
||||
ustring object, TypeDesc type,
|
||||
ustring name, void* val);
|
||||
typedef std::unordered_map<ustring, AttrGetter, ustringHash> AttrGetterMap;
|
||||
typedef std::unordered_map<ustring, AttrGetter> AttrGetterMap;
|
||||
AttrGetterMap m_attr_getters;
|
||||
|
||||
// Attribute getters
|
||||
|
||||
@@ -42,6 +42,7 @@ static bool saveptx = false;
|
||||
static bool warmup = false;
|
||||
static bool profile = false;
|
||||
static bool O0 = false, O1 = false, O2 = false;
|
||||
static int llvm_opt = 1; // LLVM optimization level
|
||||
static bool debugnan = false;
|
||||
static bool debug_uninit = false;
|
||||
static bool userdata_isconnected = false;
|
||||
@@ -76,6 +77,15 @@ set_shadingsys_options()
|
||||
if (const char* opt_env = getenv("TESTSHADE_OPT")) // overrides opt
|
||||
opt = atoi(opt_env);
|
||||
shadingsys->attribute("optimize", opt);
|
||||
|
||||
// The cost of more optimization passes usually pays for itself by
|
||||
// reducing the number of instructions JIT ultimately has to lower to
|
||||
// the target ISA.
|
||||
if (const char* llvm_opt_env = getenv(
|
||||
"TESTSHADE_LLVM_OPT")) // overrides llvm_opt
|
||||
llvm_opt = atoi(llvm_opt_env);
|
||||
shadingsys->attribute("llvm_optimize", llvm_opt);
|
||||
|
||||
shadingsys->attribute("profile", int(profile));
|
||||
shadingsys->attribute("lockgeom", 1);
|
||||
shadingsys->attribute("debug_nan", debugnan);
|
||||
@@ -149,6 +159,8 @@ getargs(int argc, const char* argv[])
|
||||
.help("Do a little runtime shader optimization");
|
||||
ap.arg("-O2", &O2)
|
||||
.help("Do lots of runtime shader optimization");
|
||||
ap.arg("--llvm_opt %d:LEVEL", &llvm_opt)
|
||||
.help("LLVM JIT optimization level");
|
||||
ap.arg("--debugnan", &debugnan)
|
||||
.help("Turn on 'debugnan' mode");
|
||||
ap.arg("--path SEARCHPATH", &shaderpath)
|
||||
|
||||
@@ -62,33 +62,42 @@ target_link_libraries (testshade
|
||||
PRIVATE
|
||||
oslexec oslquery)
|
||||
|
||||
|
||||
install (TARGETS testshade RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} )
|
||||
|
||||
# The 'libtestshade' library
|
||||
add_library ( "libtestshade" ${testshade_srcs} )
|
||||
|
||||
set_target_properties (libtestshade
|
||||
PROPERTIES
|
||||
VERSION ${PROJECT_VERSION_MAJOR}.${PROJECT_VERSION_MINOR}.${PROJECT_VERSION_PATCH}
|
||||
SOVERSION ${SOVERSION}
|
||||
OUTPUT_NAME libtestshade${OSL_LIBNAME_SUFFIX}
|
||||
)
|
||||
|
||||
target_link_libraries (libtestshade
|
||||
PRIVATE
|
||||
oslexec oslquery)
|
||||
set_target_properties (libtestshade PROPERTIES PREFIX "")
|
||||
|
||||
install_targets ( libtestshade )
|
||||
|
||||
# The 'testshade_dso' executable
|
||||
add_executable ( testshade_dso testshade_dso.cpp )
|
||||
target_link_libraries (testshade_dso
|
||||
PRIVATE
|
||||
OpenImageIO::OpenImageIO
|
||||
${Boost_LIBRARIES} ${CMAKE_DL_LIBS} )
|
||||
install (TARGETS testshade_dso RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} )
|
||||
|
||||
osl_optix_target(testshade)
|
||||
osl_optix_target(libtestshade)
|
||||
|
||||
|
||||
# As explained in PR #39, problems were reported with calling OSL from a
|
||||
# Houdini plugin. So this section sets up a version of testshade as a DSO, to
|
||||
# be sure we can link it into an executable properly.
|
||||
#
|
||||
# But the fact that we use testshade.cpp twice -- once for regular testshade,
|
||||
# once for the DSO version -- is throwing off the code coverage test, so at
|
||||
# least for now, disable the dso version entirely when CODECOV is on.
|
||||
if (NOT CODECOV)
|
||||
# The 'libtestshade' library
|
||||
add_library ( "libtestshade" ${testshade_srcs} )
|
||||
|
||||
set_target_properties (libtestshade
|
||||
PROPERTIES
|
||||
VERSION ${PROJECT_VERSION_MAJOR}.${PROJECT_VERSION_MINOR}.${PROJECT_VERSION_PATCH}
|
||||
SOVERSION ${SOVERSION}
|
||||
OUTPUT_NAME libtestshade${OSL_LIBNAME_SUFFIX}
|
||||
)
|
||||
|
||||
target_link_libraries (libtestshade
|
||||
PRIVATE
|
||||
oslexec oslquery)
|
||||
set_target_properties (libtestshade PROPERTIES PREFIX "")
|
||||
|
||||
install_targets ( libtestshade )
|
||||
|
||||
# The 'testshade_dso' executable
|
||||
add_executable ( testshade_dso testshade_dso.cpp )
|
||||
target_link_libraries (testshade_dso
|
||||
PRIVATE
|
||||
OpenImageIO::OpenImageIO
|
||||
${Boost_LIBRARIES} ${CMAKE_DL_LIBS} )
|
||||
install (TARGETS testshade_dso RUNTIME DESTINATION ${CMAKE_INSTALL_BINDIR} )
|
||||
osl_optix_target(libtestshade)
|
||||
endif ()
|
||||
|
||||
@@ -99,8 +99,8 @@ public:
|
||||
|
||||
private:
|
||||
SimpleRenderer& m_sr;
|
||||
std::unordered_set<ustring, ustringHash> m_uniform_objects;
|
||||
std::unordered_set<ustring, ustringHash> m_uniform_attributes;
|
||||
std::unordered_set<ustring> m_uniform_objects;
|
||||
std::unordered_set<ustring> m_uniform_attributes;
|
||||
|
||||
// Attribute and userdata retrieval -- for fast dispatch, use a hash
|
||||
// table to map attribute names to functions that retrieve them. We
|
||||
@@ -110,15 +110,13 @@ private:
|
||||
typedef bool (BatchedSimpleRenderer::*VaryingAttrGetter)(ustring object,
|
||||
ustring name,
|
||||
MaskedData val);
|
||||
typedef std::unordered_map<ustring, VaryingAttrGetter, ustringHash>
|
||||
VaryingAttrGetterMap;
|
||||
typedef std::unordered_map<ustring, VaryingAttrGetter> VaryingAttrGetterMap;
|
||||
VaryingAttrGetterMap m_varying_attr_getters;
|
||||
|
||||
typedef bool (BatchedSimpleRenderer::*UniformAttrGetter)(ustring object,
|
||||
ustring name,
|
||||
RefData val);
|
||||
typedef std::unordered_map<ustring, UniformAttrGetter, ustringHash>
|
||||
UniformAttrGetterMap;
|
||||
typedef std::unordered_map<ustring, UniformAttrGetter> UniformAttrGetterMap;
|
||||
UniformAttrGetterMap m_uniform_attr_getters;
|
||||
|
||||
// Attribute getters
|
||||
|
||||
@@ -90,9 +90,8 @@ private:
|
||||
const unsigned long OSL_PRINTF_BUFFER_SIZE = 8 * 1024 * 1024;
|
||||
|
||||
std::string m_materials_ptx;
|
||||
std::unordered_map<OIIO::ustring, optix::TextureSampler, OIIO::ustringHash>
|
||||
m_samplers;
|
||||
std::unordered_map<OIIO::ustring, uint64_t, OIIO::ustringHash> m_globals_map;
|
||||
std::unordered_map<OIIO::ustring, optix::TextureSampler> m_samplers;
|
||||
std::unordered_map<OIIO::ustring, uint64_t> m_globals_map;
|
||||
|
||||
OSL::Matrix44 m_shader2common; // "shader" space to "common" space matrix
|
||||
OSL::Matrix44 m_object2common; // "object" space to "common" space matrix
|
||||
|
||||
@@ -183,7 +183,7 @@ protected:
|
||||
typedef bool (SimpleRenderer::*AttrGetter)(ShaderGlobals* sg, bool derivs,
|
||||
ustring object, TypeDesc type,
|
||||
ustring name, void* val);
|
||||
typedef std::unordered_map<ustring, AttrGetter, ustringHash> AttrGetterMap;
|
||||
typedef std::unordered_map<ustring, AttrGetter> AttrGetterMap;
|
||||
AttrGetterMap m_attr_getters;
|
||||
|
||||
// Attribute getters
|
||||
|
||||
@@ -604,6 +604,7 @@ print_info()
|
||||
TextureSystem* texturesys = TextureSystem::create();
|
||||
shadingsys = new ShadingSystem(rend, texturesys, &errhandler);
|
||||
rend->init_shadingsys(shadingsys);
|
||||
set_shadingsys_options();
|
||||
|
||||
std::cout << "\n" << shadingsys->getstats(5) << "\n";
|
||||
|
||||
@@ -630,6 +631,8 @@ getargs(int argc, const char* argv[])
|
||||
ap.arg("filename")
|
||||
.hidden()
|
||||
.action([&](cspan<const char*> argv){ stash_shader_arg(argv); });
|
||||
ap.arg("--help")
|
||||
.help("Print help message");
|
||||
ap.arg("-v", &verbose)
|
||||
.help("Verbose messages");
|
||||
ap.arg("-t %d:NTHREADS", &num_threads)
|
||||
@@ -1179,7 +1182,7 @@ save_outputs(SimpleRenderer* rend, ShadingSystem* shadingsys,
|
||||
} else if (t.basetype == TypeDesc::INT) {
|
||||
// We are outputting an integer variable, so we need to
|
||||
// convert it to floating point.
|
||||
float* pixel = OIIO_ALLOCA(float, nchans);
|
||||
float* pixel = OSL_ALLOCA(float, nchans);
|
||||
OIIO::convert_types(TypeDesc::BASETYPE(t.basetype), data,
|
||||
TypeDesc::FLOAT, pixel, nchans);
|
||||
outputimg->setpixel(x, y, &pixel[0]);
|
||||
@@ -1344,8 +1347,8 @@ batched_save_outputs(SimpleRenderer* rend, ShadingSystem* shadingsys,
|
||||
Wide<const int[], WidthT> batchResults(
|
||||
shadingsys->symbol_address(*ctx, out_symbol), nchans);
|
||||
// TODO: Try not to do alloca's inside a loop
|
||||
int* intPixel = OIIO_ALLOCA(int, nchans);
|
||||
float* floatPixel = OIIO_ALLOCA(float, nchans);
|
||||
int* intPixel = OSL_ALLOCA(int, nchans);
|
||||
float* floatPixel = OSL_ALLOCA(float, nchans);
|
||||
for (int batchIndex = 0; batchIndex < batchSize; ++batchIndex) {
|
||||
int x = bx[batchIndex];
|
||||
int y = by[batchIndex];
|
||||
|
||||
Binary file not shown.
@@ -27,23 +27,6 @@ static ustring u_NDC("NDC"), u_raster("raster");
|
||||
static ustring u_perspective("perspective");
|
||||
static ustring u_s("s"), u_t("t");
|
||||
|
||||
CudaGridRenderer::CudaGridRenderer()
|
||||
{
|
||||
// Set up the string table. This allocates a block of CUDA device memory to
|
||||
// hold all of the static strings used by the OSL shaders. The strings can
|
||||
// be accessed via OptiX variables that hold pointers to the table entries.
|
||||
_string_table.init();
|
||||
|
||||
// Register all of our string table entries
|
||||
for (auto&& gvar : _string_table.contents()) {
|
||||
register_global(gvar.first.c_str(), gvar.second);
|
||||
}
|
||||
}
|
||||
|
||||
CudaGridRenderer::~CudaGridRenderer()
|
||||
{
|
||||
_string_table.freetable();
|
||||
}
|
||||
|
||||
uint64_t
|
||||
CudaGridRenderer::register_global(const std::string& str, uint64_t value)
|
||||
|
||||
@@ -9,19 +9,15 @@
|
||||
#include <OSL/oslexec.h>
|
||||
#include <OSL/rendererservices.h>
|
||||
|
||||
#include "cuda_string_table.h"
|
||||
|
||||
using GlobalsMap
|
||||
= std::unordered_map<OIIO::ustring, uint64_t, OIIO::ustringHash>;
|
||||
using TextureSamplerMap
|
||||
= std::unordered_map<OIIO::ustring, cudaTextureObject_t, OIIO::ustringHash>;
|
||||
using GlobalsMap = std::unordered_map<OIIO::ustring, uint64_t>;
|
||||
using TextureSamplerMap = std::unordered_map<OIIO::ustring, cudaTextureObject_t>;
|
||||
|
||||
// Just use 4x4 matrix for transformations
|
||||
typedef OSL::Matrix44 Transformation;
|
||||
typedef std::map<OIIO::ustring, std::shared_ptr<Transformation>> TransformMap;
|
||||
|
||||
class CudaGridRenderer final : public OSL::RendererServices {
|
||||
CudaStringTable _string_table;
|
||||
TextureSamplerMap _samplers;
|
||||
GlobalsMap _globals_map;
|
||||
|
||||
@@ -36,8 +32,8 @@ class CudaGridRenderer final : public OSL::RendererServices {
|
||||
int _xres, _yres;
|
||||
|
||||
public:
|
||||
CudaGridRenderer();
|
||||
virtual ~CudaGridRenderer();
|
||||
CudaGridRenderer() {}
|
||||
virtual ~CudaGridRenderer() {}
|
||||
|
||||
uint64_t register_global(const std::string& str, uint64_t value);
|
||||
bool fetch_global(const std::string& str, uint64_t* value);
|
||||
|
||||
@@ -1,187 +0,0 @@
|
||||
// Copyright Contributors to the Open Shading Language project.
|
||||
// SPDX-License-Identifier: BSD-3-Clause
|
||||
// https://github.com/AcademySoftwareFoundation/OpenShadingLanguage
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <map>
|
||||
#include <unordered_map>
|
||||
#include <utility>
|
||||
|
||||
#include <cuda_runtime_api.h>
|
||||
|
||||
#include <OSL/oslexec.h>
|
||||
#include <OpenImageIO/ustring.h>
|
||||
|
||||
typedef std::unordered_map<OIIO::ustring, int64_t, OIIO::ustringHash>
|
||||
StringTableMap;
|
||||
|
||||
// The CudaStringTable manages a block of CUDA device memory designated
|
||||
// to hold all of the string constants that a shader might access during
|
||||
// execution.
|
||||
//
|
||||
// Any string that needs to be visible on the device needs to be added using the
|
||||
// addString function.
|
||||
//
|
||||
// Note: this is basically just a copy of the one from testrender
|
||||
class CudaStringTable {
|
||||
public:
|
||||
CudaStringTable() : m_ptr(nullptr), m_size(1 << 16), m_offset(0) {}
|
||||
|
||||
~CudaStringTable() { freetable(); }
|
||||
|
||||
// Allocate CUDA device memory for the raw string table and add the
|
||||
// "standard" strings declared in strdecls.h.
|
||||
void init()
|
||||
{
|
||||
OSL_ASSERT(!m_ptr && "StringTable should only be initialized once");
|
||||
|
||||
cudaMalloc(reinterpret_cast<void**>(&m_ptr), (m_size));
|
||||
|
||||
// Add the statically-declared strings to the table, and create OptiX
|
||||
// variables for them in the OSL::DeviceStrings namespace.
|
||||
//
|
||||
// The names of the variables created here must match the extern
|
||||
// variables declared in OSL/device_string.h for OptiX's variable
|
||||
// scoping mechanisms to work.
|
||||
|
||||
#define STRDECL(str, var_name) \
|
||||
addString(OIIO::ustring(str), \
|
||||
OIIO::ustring(OSL_NAMESPACE_STRING \
|
||||
"::DeviceStrings::" #var_name));
|
||||
#include <OSL/strdecls.h>
|
||||
#undef STRDECL
|
||||
}
|
||||
|
||||
const StringTableMap& contents() const { return m_addr_table; }
|
||||
|
||||
// Add a string to the table (if it hasn't already been added), and return
|
||||
// its address in device memory. Also, create an OptiX variable to hold the
|
||||
// address of each string variable, through which the strings will be
|
||||
// accessed during execution.
|
||||
//
|
||||
// Creating variables for the strings -- rather than using the string
|
||||
// addresses directly -- is necessary for keeping the generated PTX stable
|
||||
// from run to run, regardless of the order in which strings are added to
|
||||
// the table. This helps make the PTX more cacheable.
|
||||
//
|
||||
// This function should not be called when a kernel that could potentially
|
||||
// access the table is running, since it has the potential to invalidate
|
||||
// pointers if reallocation is triggered.
|
||||
uint64_t addString(OIIO::ustring str, OIIO::ustring var_name)
|
||||
{
|
||||
OSL_ASSERT(m_ptr && "StringTable has not been initialized");
|
||||
|
||||
// The strings are laid out in the table as a struct:
|
||||
//
|
||||
// struct TableRep {
|
||||
// size_t len;
|
||||
// size_t hash;
|
||||
// char str[len+1];
|
||||
// };
|
||||
|
||||
// Compute the size of the entry before adding it to the table
|
||||
size_t size = sizeof(size_t) + sizeof(size_t) + str.size() + 1;
|
||||
if (((m_offset + size) >= m_size)) {
|
||||
reallocTable();
|
||||
}
|
||||
|
||||
// It should be hard to trigger this assert, unless the table size is
|
||||
// very small and the string is very large.
|
||||
OSL_ASSERT(m_offset + size <= m_size
|
||||
&& "String table allocation error");
|
||||
|
||||
int offset = getOffset(str.string());
|
||||
if (offset < 0) {
|
||||
// Place the hash and length of the string before the characters
|
||||
size_t hash = str.hash();
|
||||
cudaMemcpy(m_ptr + m_offset, (void*)&hash, sizeof(size_t),
|
||||
cudaMemcpyHostToDevice);
|
||||
m_offset += sizeof(size_t);
|
||||
|
||||
size_t len = str.length();
|
||||
cudaMemcpy(m_ptr + m_offset, (void*)&len, sizeof(size_t),
|
||||
cudaMemcpyHostToDevice);
|
||||
m_offset += sizeof(size_t);
|
||||
|
||||
offset = m_offset;
|
||||
m_offset_map[str] = offset;
|
||||
m_name_map[str] = var_name;
|
||||
|
||||
// Copy the raw characters to the table
|
||||
cudaMemcpy(m_ptr + m_offset, str.c_str(), str.size() + 1,
|
||||
cudaMemcpyHostToDevice);
|
||||
m_offset += str.size() + 1;
|
||||
|
||||
// Align the offset for the next entry to 8-byte boundaries
|
||||
m_offset = (m_offset + 0x7u) & ~0x7u;
|
||||
} else if (!var_name.empty()) {
|
||||
// update what str points to
|
||||
m_name_map[str] = var_name;
|
||||
}
|
||||
|
||||
uint64_t addr = reinterpret_cast<uint64_t>(m_ptr + offset);
|
||||
|
||||
// Optionally create an OptiX variable for the string. It's not
|
||||
// necessary to create a variable for strings that do not appear by name
|
||||
// in compiled code (in either the OSL library functions or in the
|
||||
// renderer).
|
||||
if (!var_name.empty())
|
||||
m_addr_table[var_name] = addr;
|
||||
return addr;
|
||||
}
|
||||
|
||||
void freetable()
|
||||
{
|
||||
if (m_ptr)
|
||||
cudaFree(m_ptr);
|
||||
m_ptr = nullptr;
|
||||
}
|
||||
|
||||
private:
|
||||
// If a string has already been added to the table, return its offset in the
|
||||
// char array; otherwise, return -1.
|
||||
int getOffset(const std::string& str) const
|
||||
{
|
||||
auto it = m_offset_map.find(OIIO::ustring(str));
|
||||
return (it != m_offset_map.end()) ? it->second : -1;
|
||||
}
|
||||
|
||||
// Free the previous allocation, allocate a new block of GPU memory of twice
|
||||
// the size, copy the string contents into the new allocation, and update
|
||||
// the OptiX variables that hold the string addresses.
|
||||
void reallocTable()
|
||||
{
|
||||
m_size *= 2;
|
||||
cudaFree(m_ptr);
|
||||
cudaMalloc(reinterpret_cast<void**>(&m_ptr), (m_size));
|
||||
|
||||
// The offsets need to be recomputed
|
||||
m_offset = 0;
|
||||
m_offset_map.clear();
|
||||
|
||||
// Add the string collection to the newly-allocated memory
|
||||
for (auto& entry : m_name_map) {
|
||||
addString(entry.first, entry.second);
|
||||
}
|
||||
}
|
||||
|
||||
// A byte array containing the concatenation of all strings added to the
|
||||
// table, allocated in CUDA device memory. The hash value and length of each
|
||||
// string are stored in the 16 bytes preceding the raw characters.
|
||||
char* m_ptr;
|
||||
|
||||
// The size of the table in bytes.
|
||||
size_t m_size;
|
||||
|
||||
// The offset in the char array at which the next string will be added.
|
||||
int m_offset;
|
||||
|
||||
// The memory offsets associated with each canonical string.
|
||||
std::map<OIIO::ustring, int> m_offset_map;
|
||||
|
||||
// The variable names associated with each canonical string.
|
||||
std::map<OIIO::ustring, OIIO::ustring> m_name_map;
|
||||
|
||||
StringTableMap m_addr_table;
|
||||
};
|
||||
@@ -326,8 +326,8 @@ osl_texture(void* sg_, const char* name, void* handle, void* opt_, float s,
|
||||
{
|
||||
if (!handle)
|
||||
return 0;
|
||||
cudaTextureObject_t texID = cudaTextureObject_t(handle);
|
||||
float4 fromTexture = osl_tex2DLookup(handle, s, t);
|
||||
// cudaTextureObject_t texID = cudaTextureObject_t(handle);
|
||||
float4 fromTexture = osl_tex2DLookup(handle, s, t);
|
||||
// see note above
|
||||
// float4 fromTexture = tex2D<float4>(texID, s, t);
|
||||
*((float3*)result) = make_float3(fromTexture.x, fromTexture.y,
|
||||
|
||||
@@ -5,4 +5,6 @@
|
||||
# https://github.com/AcademySoftwareFoundation/OpenShadingLanguage
|
||||
|
||||
outputs = [ "out.exr" ]
|
||||
command = testrender("-r 256 256 -aa 4 bumptest.xml out.exr")
|
||||
command = testrender("-r 256 256 -aa 4 --llvm_opt 13 bumptest.xml out.exr")
|
||||
|
||||
# Note: we pick this test arbitrarily as the one to verify llvm_opt=13 works
|
||||
|
||||
@@ -7,4 +7,6 @@
|
||||
failthresh = 0.01
|
||||
failpercent = 1
|
||||
outputs = [ "out.exr" ]
|
||||
command = testrender("-r 256 256 -aa 4 cornell.xml out.exr")
|
||||
command = testrender("-r 256 256 -aa 4 --llvm_opt 12 cornell.xml out.exr")
|
||||
|
||||
# Note: we pick this test arbitrarily as the one to verify llvm_opt=12 works
|
||||
|
||||
Reference in New Issue
Block a user