| Age | Commit message (Collapse) | Author |
|---|
|
|
|
* Add test case for missing import attribution
Add a test case that imports a non-existent file, followed by a valid
file. Tests for absence of a bug where slang reports existent files as
non-existent if they're imported after a non-existent file.
* Skip processing imports after errors
Skip processing additional imports after the first error. This
behavior is already observed in Linkage::loadSourceModuleImpl, but
since that happenes after import processing already started, a false
diagnostic gets generated for a missing import.
By hoisting this check out before the import is processed, the
diagnostic message for a missing file is no longer erroneously
generated.
Fixes #6453
* Revert "Skip processing imports after errors"
This reverts commit 6b2fef09782414de4c5e017c4ecb5f2affa0c199.
* Remove early abort of import processing
Partial revert of commit 04f1bad
Reverts an early return in Linkage::loadSourceModuleImpl() whenever any
error diagnostic message has already been generated. This was causing
earlier errors to prevent subsequent imports from succeeding, and was
misattributing them to a missing file.
Fixes #6453
---------
Co-authored-by: Yong He <yonghe@outlook.com>
|
|
|
|
* Print CUDA version on CI servers
|
|
|
|
* Use aliased SPIRV-Headers::SPIRV-Headers to also work with an installed SPIRV-Headers
SPIRV-Headers standalone is only defined when using sources directly.
When consuming an installed SPIRV-Headers via find_package, the full SPIRV-Headers::SPIRV-Headers must be used.
The full syntax is supported by both source and installed builds.
* Fix SLANG_USE_SYSTEM_SPIRV_HEADERS
- Use find_package to bring in SPIRV-Headers cmake targets
- Set SPIRV-Headers_SOURCE_DIR as a workaround when including
spirv-tools
- Query cmake for SLANG_SPIRV_HEADERS_INCLUDE_DIR location, supporting
default, SLANG_OVERRIDE_SPIRV_HEADERS_PATH and find_package builds.
- Cleanup unnecessary SPIRV_HEADER_DIR (unconditionally overwritten in
spirv-tools)
|
|
|
|
Close #6176.
If the struct has a `no_diff` member, it should not be its Differential
type. We miss this check.
|
|
* enable building slang-rhi-tests
* re-enable running slang-rhi-tests
* format code
* fix typo
* update slang-rhi
* build slang-rhi-tests without glfw dependency
* skip fence tests
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
Co-authored-by: amey asgaonkar <160177341+aasgaonkar@users.noreply.github.com>
|
|
On Linux/MacOS, the test file lists were not in a sorted order and the order was non-deterministic.
It causes intermittent CI failures.
This commit is to make the list sorted and CI more reliable.
|
|
* Add new capdef for lss intrinsics
Fixes #7426
Raygen shaders need to be supported for only hitobject APIs. So we need
a special capability for that, instead of a common one.
* regenerate command line reference
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
Recent separate debugging support added two new functions which broke
backwards compatibility. This change restores the old API and moves the
new functions to an IComponentType2 interface which can be used if
separate debug files are needed.
|
|
* Fix issue of missing scope for 'Differential' type
When we synthesize the struct decl for Differential type, we should
add the ownedScope for this decl, because the scope is used in lots
of locations in the following synthesized processes, e.g. constructor
synthesize. And that could cause surprising behavior, e.g. the 'this'
expression could access the members of parent struct decl.
Fix the issue by adding the scope. The containerDecl will be the
Differential struct decl itself, parent scope will be the parent struct.
* Add a unit-test
|
|
Results of these tests had been marked ignored, because they failed on
VK with the GLSL backend. This change removes them from the
expected-failure.txt file and adds the correct command line option to
avoid using the GLSL target.
Addresses concern raised on #7282
|
|
* Add the missing generator expression for install build targts
Fixes #7351.
* Formatted
---------
Co-authored-by: Harsh Aggarwal (NVIDIA) <haaggarwal@nvidia.com>
|
|
* Diagnose on use of struct inheritance.
* fix test.
* Fix tests.
* fix.
---------
Co-authored-by: ArielG-NV <159081215+ArielG-NV@users.noreply.github.com>
|
|
This PR replaces enable/disable style C function calls with C++ RAII style code.
In debug build, when an assertion failed in between enable and disable functions, an exception is thrown and the disable function is not called. RAII style code is safer for an exception
|
|
fixes: #7417
Define was called `SLANG_ENABLE_IR_BREAK_ALLOC,`. This is incorrect, supposed to be `SLANG_ENABLE_IR_BREAK_ALLOC`.
|
|
Missing DebugLine in some basic blocks that include OpBranchConditional
causes invalid line number '0' presented in the line table of '.debug_line'
section. Emiting Debugline before IfElse fixes the issue.
Modified maybeEmitDebugLine() to handle the case without Stmt.
|
|
* Fix an issue in extension override.
* Fix typo in comment.
|
|
|
|
Apply argument buffer tier2 rule when using parameter block for Metal target.
Close #6803.
|
|
This allows checking capabilities in any stage, needed specifically for
the hlsl_2018 capability which is defined for sm_5_1 and above. Stage
specific capabilities such as cs_5_1 would not find this in any stage
other than compute, so we need to restrict the check to only desired
stages.
|
|
* Better handling for 16-byte boundary of d3d cbuffer
Fixes #6921
D3D cbuffers have slightly different packing rules that allow packing
vectors into a 16-byte slot at element alignments, except when
a field would cross a 16-byte boundary. In that case, we need to
realign the field to the next 16-byte boundary.
In particular, this impacts vec3s, which are not a power of two in
size and thus require slightly different alignment logic, compared to
std430 and std140. (Example: a float and float3 should fit together in
that order in a single slot.)
Adds test cases.
Adds documentation page for GLSL target
|
|
|
|
* Handle pointer types when getting type cast style
Closes https://github.com/shader-slang/slang/issues/6025
* Move vertex shader out parameters to return type for Metal
Closes https://github.com/shader-slang/slang/issues/6025
* More asserts
* Make struct instead of tuple
* More layout preservation
* Handle same function result
* more layout
* remove layout
* a
* more debug code
* more debug code
* a
* layout working
* refactored
* more tests
* more tests
* fuse loops
* remove unused comments
* Correct filecheck usage
* debug code
* correct name and order of filecheck vars
* simplify
* Address review comments
fix warning
* simplify handling of simple vertex shaders
|
|
* WiP: Add coopvec support for Optix
* format code
* fix minor issues
* Fix review comments
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
* Add release workflow targeting glibc 2.17
* format code (#7366)
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
* Update comment
---------
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
Most of what this change does is straightforward: take all the places in the code that used to operate directly on `ContainerDecl::members` and related fields, and instead have them call into a smaller set of accessor methods defined on `ContainerDecl`.
The primary motivation for making this change is that in order to implement on-demand loading of members from serialized AST modules, we need a way to identify and intercept the "demand" for those members.
On-demand loading benefits from having all accesses to the members of a `ContainerDecl` be as narrow as possible.
If a part of the code only need a member at a specific index, it should say so.
If it only needs access to members with a specific name, or a given subclass of `Decl`, then it should say so.
A secondary motivation for this change is that there have recently been several changes that added complexity and special cases by introducing code that operated on (and *mutated*) the member list of a container decl in ways that the existing code had never done before.
Any code that mutates the member list of a `ContainerDecl` needs to be sure to not disrupt the invariants that the lookup acceleration structures currently rely on.
One of the recent changes added a declaration-to-index map to the set of acceleration structures (with different validation/invalidation behavior than the others...) while other recent changes would remove or insert declarations in ways that could change the indices of other declarations in the same container.
It is not clear if any of these pieces of code were aware of the others, and the invariants that might be expected or broken along the way.
This change bottlenecks the vast majority of accesses to the members of a `ContainerDecl` through the following operations:
* Getting a `List` of all of the direct member declarations of a container
* Get the number of direct member declarations, and accessing them by index.
* Looking up the list of direct member declarations with a given name.
* Adding a new direct member declaration to the end of the list.
Some other operations are layered on top of those (e.g., getting a list of all the direct member declarations of a given C++ class).
These layered operations are still centralized on the `ContainerDecl`, with the intention that we *can* change them to be non-layered implementations if we ever need to for performance (e.g., by building a lookup structure for finding member declarations by their type).
The exceptional cases of access/mutation on the direct members of a `ContainerDecl` have also been encapsulated, but rather than expose what would risk appearing like general-purpose accessors (e.g., `removeDecl(d)`, `setDecl(index)`, etc.), these operations have been explicitly named after the specific use case that they serve in the codebase today, to discourage others from using them for more kinds of operations we'd rather not support.
These operations have also been given parameter signatures that match their use cases, to make it so that even somebody determined to abuse them would have to invent suitable arguments out of thin air.
In the case of the declaration-to-index mapping, this change eliminates that acceleration structure, in favor or slightly more complicated (and possibly inefficient, yes) code at the use site.
Over time, it would be good to closely scrutinize each of the use cases that requires more complicated interaction with the members of a `ContainerDecl`, to see whether any of them can be reframed in terms of the more basic operations, or if there is some clean abstraction we can introduce to make operations that mutate the member list feel like... hacky.
|
|
Added error checking to reject interface types as the right-hand side
of is and as operators. Enhanced semantic analysis with new diagnostic
30301 and comprehensive test coverage.
|
|
* Remove unused compile definitions
No need for full external path; slang is already linking with the
SPIRV-Headers cmake target which adds the proper include flags
|
|
* Add command line option for separate debug info
Add command line arg -separate-debug-info which, if provided, produces
both a .spv and a .dbg.spv file. The .dbg.spv file contains full debug
info and the .spv file has all debug info stripped out.
Also add a DebugBuildIdentifier instruction to store a unique hash in
both the output files, so they can be more easily matched together.
A matching API is provided to allow using the Slang API to retrieve a
base and debug SPIRV as well as the debug build identifier string.
|
|
* Address issues with GLSL style global in/out vars (#6669)
Asserts and segfaults were observed trying to compile a simple
vertex shader like:
````
in int2 inPos;
[shader("vertex")]
main(uniform int2 test1, int2 test2, out float4 pos: SV_Position)
void main()
{
// Bogus use of all input vars to prevent optimizing out.
pos = float4(inPos.x, test1.x, test2.y, 0);
}
````
Further investigation found that while replacing "uniform int2 test1"
with "int2 test1" allowed for successful compilation, the resulting
output shader would have overlapping location qualifiers. For example,
compiling the above with "int2 test1" to glsl might give:
````
...
layout(location = 0) in ivec2 test1_0;
layout(location = 1) in ivec2 test2_0;
layout(location = 0) in ivec2 translatedGlobalParams_inPos_0;
...
````
This was because Slang does not actually support mixing GLSL style global
in/out vars and entry point params. However, this is never checked for
or noted in documentation. Slang source also assumes input shaders do not
mix these and these assumptions ultimately led to the observed asserts
and seg faults when using uniform entry point params.
This change makes updates to throw an error when the compiler detects that
it is trying to translate global in/out variables into entry point params
when an entry point already contains parameters, allowing for compilation
to fail gracefully.
Certain tests have been updated to avoid mixing GLSL style global in/out
vars and entry point params. This was mostly for tests that were using
functions like WaveGetLaneIndex which use global in vars for certain
platforms (see __builtinWaveLaneIndex).
* Address issues with GLSL style global in/out vars - updates 1 (#6669)
Update addresses review feedback to support mixing GLSL-flavored global
in/out vars and entrypoint parameters when either all global in/out vars
or all entry point params have a system value binding semantic.
* Address issues with GLSL style global in/out vars - updates 2 (#6669)
This update attempts to actually allow mixing GLSL style global in
vars and entry point vars.
Change attempts to recalculate offsets when adding the global input
vars into the recreated entry point params layout.
Additional updates were made to:
-resolve further issues uncovered with entry point uniform params.
-Address improper use of SV_DispatchThreadID in wave-get-lane-index.slang
for metal. "thread_position_in_grid" is not supported for signed integer
scalars or vectors.
-Fix a spirv casting conflict due to the implementation of
gl_PrimitiveID.get conflicting with PrimitiveIndex().
-Add a call to remove a global var in replaceUsesOfGlobalVar(). The global
var is already replaced in this function and keeping it around can prevent
it from being cleaned up by DCE if it still has decorations.
* format code
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
* Disable Link-Time-Optimization by default
LTO was requested for the release package a while ago.
When we added it, LTO was enabled by default although it was needed only for
the release packages.
Later we found that the Release build cannot be incrementally recompiled
when LTO is enabled. It sometimes works fine, but it required full recompilation
when it doesn't work. We added a new CMake option, `SLANG_ENABLE_RELEASE_LTO`,
to disable it for developers. But many Slang developers don't know the
option exists.
I was going to update the document, CONTRIBUTING.md, but I thought it
will be better to change the default behavior.
* Fix a compiler warning treated as an error on linux
A padding variable was uninitialized, which is fine, but the compiler
was complaining about it.
* Fix other gcc error for uninitialized variable
* Fix more compile warning treated as error
* Fix compiler warning from gcc 11
It appears that this is a valid warning that the `delete this` is done
on an offset 8 when the class uses multiple inheritance.
The compiler warning is following:
```
In file included from /home/runner/work/slang/slang/source/core/slang-memory-file-system.h:5,
from /home/runner/work/slang/slang/tools/slang-unit-test/unit-test-module-ptr.cpp:3:
In destructor ‘virtual Slang::ComBaseObject::~ComBaseObject()’,
inlined from ‘uint32_t Slang::ComBaseObject::_releaseImpl()’ at /home/runner/work/slang/slang/source/core/slang-com-object.h:49:16,
inlined from ‘virtual uint32_t Slang::MemoryFileSystem::release()’ at /home/runner/work/slang/slang/source/core/slang-memory-file-system.h:34:5:
/home/runner/work/slang/slang/source/core/slang-com-object.h:33:31: error: ‘void operator delete(void*, std::size_t)’ called on pointer ‘<unknown>’ with nonzero offset 8 [-Werror=free-nonheap-object]
33 | virtual ~ComBaseObject() {}
| ^
In destructor ‘virtual Slang::ComBaseObject::~ComBaseObject()’,
inlined from ‘uint32_t Slang::ComBaseObject::_releaseImpl()’ at /home/runner/work/slang/slang/source/core/slang-com-object.h:49:16,
inlined from ‘virtual uint32_t Slang::MemoryFileSystem::release()’ at /home/runner/work/slang/slang/source/core/slang-memory-file-system.h:34:5,
inlined from ‘Slang::ComPtr<T>::~ComPtr() [with T = ISlangMutableFileSystem]’ at /home/runner/work/slang/slang/include/slang-com-ptr.h:113:34,
inlined from ‘void _modulePtr_impl(UnitTestContext*)’ at /home/runner/work/slang/slang/tools/slang-unit-test/unit-test-module-ptr.cpp:92:1:
/home/runner/work/slang/slang/source/core/slang-com-object.h:33:31: error: ‘void operator delete(void*, std::size_t)’ called on pointer ‘<unknown>’ with nonzero offset 8 [-Werror=free-nonheap-object]
33 | virtual ~ComBaseObject() {}
| ^
/home/runner/work/slang/slang/tools/slang-unit-test/unit-test-module-ptr.cpp: In function ‘void _modulePtr_impl(UnitTestContext*)’:
/home/runner/work/slang/slang/tools/slang-unit-test/unit-test-module-ptr.cpp:36:69: note: returned from ‘void* operator new(std::size_t)’
36 | ComPtr<ISlangMutableFileSystem>(new Slang::MemoryFileSystem());
| ^
```
The problem is on the fact that `ComBaseObject` is not the first in the
multiple inheritance:
```
class MemoryFileSystem : public ISlangMutableFileSystem, public ComBaseObject
{
public:
// ISlangUnknown
SLANG_COM_BASE_IUNKNOWN_ALL
```
It should be:
```
class MemoryFileSystem : public ComBaseObject, public ISlangMutableFileSystem
```
The chain of ComObject release is little complicated and it is easy to
make a mistake. Here is summary with details,
1. `release()` is declared as a pure-virtual in ISlangUnknown, which is
one of the base classes of `ISlangMutableFileSystem`.
```
struct ISlangUnknown
{
virtual SLANG_NO_THROW uint32_t SLANG_MCALL release() = 0;
```
2. `release()` is implemented with the macro
`SLANG_COM_BASE_IUNKNOWN_RELEASE`.
```
SLANG_NO_THROW uint32_t SLANG_MCALL release() SLANG_OVERRIDE \
{ \
return _releaseImpl(); \
}
inline uint32_t ComBaseObject::_releaseImpl()
{
// Check there is a ref count to avoid underflow
SLANG_ASSERT(m_refCount != 0);
const uint32_t count = --m_refCount;
if (count == 0)
{
delete this;
}
return count;
}
```
3. The instance of `MemoryFileSystem` is handled by ComPtr. And
`ComPtr::~ComPtr()` calls the `release()`.
```
ComPtr<ISlangMutableFileSystem> memoryFileSystem =
ComPtr<ISlangMutableFileSystem>(new Slang::MemoryFileSystem());
SLANG_FORCE_INLINE ~ComPtr()
{
if (m_ptr)
((Ptr)m_ptr)->release();
}
```
4. When `delete this` is called, because ComBaseObject is not the first
in the multiple inheritance, `this` is 8 byte off from the actual
instance address.
A fix for this is to change the order of the inheritance and make
ComBaseObject to be the first in the order.
|
|
it (#7357)
* Disable coopvec tests until we have the driver support
* Remove conflicting agility sdk binaries
|
|
* update slang-rhi
* adapt to new slang-rhi API
* enable slang-rhi agility sdk
* fix handling empty list
* disable failing slang-rhi tests
* format code
* fix slang-rhi-tests ci step
* skip running slang-rhi-tests
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
|
|
Close #7315.
We have couple mis-definition in capability.
sm_50 shouldn't require cuda compute_9_0, drop it to compute_6_0
unpack should only require compute_6_0
subgroup_ballot will require sm_60
Co-authored-by: Yong He <yonghe@outlook.com>
|
|
slang-test is ran with a simple test, tests/render/check-backend-support-on-ci.slang.
And the github CI,yaml will check if certain keywords are found.
This is to prevent the regression on the CI server where a set of backend was supported at some point and it stopped supporting it siliently.
|
|
Update spirv-tools to take the DebugBuildIdentifier changes. These are
needed for upcoming separate debug spirv output.
Additionally update spirv-headers to fd966619.
|
|
WGSL doesn't support isnan and isinf, because it assumes that it always uses fast-math and fast-math doesnt' handle NaN as defined in IEEE standard.
The initial implementation used a clever workaround but it stopped working from some point.
This PR implemented isnan and isinf with a bitwise operation, which can be expensive.
But that seems to be an only option at the moment.
|
|
* Fix#6993 - Emit Diagnostic Warning and Fix SIGSEGV
* Update external/slang-rhi submodule
* Add checks for valid stage names for paq in SemanticsVisitor check
* format code
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
Co-authored-by: Ellie Hermaszewska <ellieh@nvidia.com>
|
|
* Fix missing debug info in the included slang file
Issue:
https://github.com/shader-slang/slang/issues/7271
Debug info including DebugFunction, DebugLocation, and DebugValue
are missing in IR for "#included" Slang shader file.
The included shader file was not added to TranslationUnit's source
file list, therefore mapSourceFileToDebugSourceInst.add() was not
called for the source in generateIRForTranslationUnit(), and later
mapSourceFileToDebugSourceInst.tryGetValue() could not get value
for the source to add DebugLocationDecoration, which led to missing
DebugFunction, DebugLocation and other debug info for the included
file in IR.
Adding the include file in TranslationUnit's source file list fixes
the issue.
* Add source file using PreprocessorHandler
Call _addSourceFile from FrontEndPreprocessorHandler::handleFileDependency.
* Just use FrontEndPreprocessorHandler
* Make _addSourceFile public
* format code
* Distingush the included source file
* Add m_includedFileSet to avoid adding dup file
HashSet<SourceFile*> m_includedFileSet;
---------
Co-authored-by: ArielG-NV <159081215+ArielG-NV@users.noreply.github.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
|
|
This reverts commit 2d98381501140bdf519f2a28c417b205187525c7.
|
|
* Add legalization for 0-sized arrays.
* Allow 0-sized arrays in the front-end.
* More tests.
* Add `Conditional<T, hasValue>` type to core module.
* Update toc.
* Fix wording.
* Update test.
|
|
* Make interface types non c-style.
* Make Optional<T> work with autodiff and existential types.
* Fix.
* patch behind slang 2026.
* Fix warnings.
* cleanup.
* Fix tests.
* Fix.
* Fix com interface lowering.
* Add comment to test.
* regenerate command line reference
* Add test for passing `none` to autodiff function.
* Fix recording of `getDynamicObjectRTTIBytes`.
* Fix nested Optional types.
---------
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
|
|
* Break down RecordReply to individual tests to avoid timeout
In Debug build, RecordReplay unit-test was timing out.
It was running six tests all in one unit-test, but this commit breaks it
down to individual test so that each unit test can be done within the
timeout limit.
This issue has seen only in Debug build but it has been unnoticed
because even when the test failed with test-server, it was still passing
on its retry because the time-out applies only when using test-server.
* Reduce the retry from 2 times to 1 time
* Remove RecordReplay from expected failure
|
|
|
|
|
