| Age | Commit message (Collapse) | Author |
|
* Backend for Multiple Entry Points
Introduces the basic backend on the compiler for zero or more entry
points. Entry points have been extended to lists for several functions,
with loopFunctions have been extended to take in entry points and
indices as appropriate, to allow for multiple entry points once the
frontend is expanded. Several functions are currently being assumed to
have a single entry point for simplicity and provide a work in progress
commit.
* Progress on debugging fixes
* Tests passing
* Refactored emitEntryPoints
* Updated lists to be by constant reference
* Fixes to formatting
* Refactoring updates for the compiler
* Fix for compilation errors
* Reformatting
* More reformatting
* Moved struct around to help with compilation
Co-authored-by: Tim Foley <tfoleyNV@users.noreply.github.com>
|
|
|
|
* Try to fix problem with C++ extractor concating tokens producing an erroneous result.
* Improve naming/comments around C++ extractor fix.
* Another small improvement around space concating when outputing token list.
* Handle some more special cases for consecutive tokens for C++ extractor concat of tokens.
* WIP AST serialization.
* Comment out so compile works.
* More work on AST serialization.
* WIP AST serialize.
* WIP AST Serialization - handling more types.
* WIP: Compiles but not all types are converted, as not all List element types are handled.
* Compiles with array types.
* Finish off AST serialization of remaining types.
* Remove ComputedLayoutModifier and TupleVarModifier.
* Add fields to ASTSerialClass type.
* Construct AST type layout.
* AST Serialization working for writing to ASTSerialWriter.
* Removed call to ASTSerialization::selfTest in session creation.
* Fixes for gcc.
* Diagnostics handling - better handling of dashify.
* Improve comment around DiagnosticLookup.
* Updated VS project.
* Write out as a Stream, taking into account alignment.
* First pass at serializing in AST.
* Added support for deserializing arrays.
* Small bug fixes.
* Fix problem calculating layout.
Split out loading on entries.
* Fix typo in AST conversion.
* Add some flags to control AST dumping.
* Fix bug from a typo.
* Special case handling of Name* in AST serialization.
* Special case handling of Token lexemes, make Names on read.
* Documentation on AST serialization.
* ASTSerialTestUtil - put AST testing functions.
Fix typo that broke compilation.
* Fix typo.
|
|
|
|
|
|
|
|
|
|
|
|
|
|
* Try to fix problem with C++ extractor concating tokens producing an erroneous result.
* Improve naming/comments around C++ extractor fix.
* Another small improvement around space concating when outputing token list.
* Handle some more special cases for consecutive tokens for C++ extractor concat of tokens.
* WIP AST serialization.
* Comment out so compile works.
* More work on AST serialization.
* WIP AST serialize.
* WIP AST Serialization - handling more types.
* WIP: Compiles but not all types are converted, as not all List element types are handled.
* Compiles with array types.
* Finish off AST serialization of remaining types.
* Remove ComputedLayoutModifier and TupleVarModifier.
* Add fields to ASTSerialClass type.
* Construct AST type layout.
* AST Serialization working for writing to ASTSerialWriter.
* Removed call to ASTSerialization::selfTest in session creation.
* Fixes for gcc.
* Diagnostics handling - better handling of dashify.
* Improve comment around DiagnosticLookup.
* Updated VS project.
* Write out as a Stream, taking into account alignment.
* First pass at serializing in AST.
* Added support for deserializing arrays.
* Small bug fixes.
* Fix problem calculating layout.
Split out loading on entries.
* Fix typo in AST conversion.
|
|
|
|
|
|
|
|
-Lower interfaces into actual `IRInterfaceType` insts.
-Lower `DeclRef<AssocTypeDecl>` into `IRAssociatedType`
-Generate proper IRType for generic functions.
-Add a test case exercising dynamic dispatching a generic static function through an associated type.
-Bug fixes for the test case.
|
|
* Try to fix problem with C++ extractor concating tokens producing an erroneous result.
* Improve naming/comments around C++ extractor fix.
* Another small improvement around space concating when outputing token list.
* Handle some more special cases for consecutive tokens for C++ extractor concat of tokens.
* WIP AST serialization.
* Comment out so compile works.
* More work on AST serialization.
* WIP AST serialize.
* WIP AST Serialization - handling more types.
* WIP: Compiles but not all types are converted, as not all List element types are handled.
* Compiles with array types.
* Finish off AST serialization of remaining types.
* Remove ComputedLayoutModifier and TupleVarModifier.
* Add fields to ASTSerialClass type.
* Construct AST type layout.
* AST Serialization working for writing to ASTSerialWriter.
* Removed call to ASTSerialization::selfTest in session creation.
* Fixes for gcc.
* Diagnostics handling - better handling of dashify.
* Improve comment around DiagnosticLookup.
* Updated VS project.
|
|
|
|
struct-inheritance-and-interfaces
|
|
|
|
The main new feature that works here is that a derived `struct` type can satisfy one or more interface requirements using methods it inherited from a base `struct` type:
```hlsl
interface ICounter { [mutating] void increment(); }
struct CounterBase { int val; [mutating] void increment() { val++; } }
struct ResetableCounter : CounterBase, ICounter
{
[mutating] void reset() { val = 0; }
}
```
Here the derived `ResetableCounter` type is satisfying the `increment()` requirement from `ICounter` using the inherited `CounterBase` method instead of one defined on `ResetableCounter`.
The crux of the problem here was that after lowering to HLSL/GLSL, the above code looks something like:
```hlsl
struct CounterBase { int val; };
void CounterBase_increment(in out CounterBase this) { this.val++; }
struct ResetableCounter { CounterBase base; }
void ResetableCounter_reset(in out ResetableCounter this) { this.base.val = 0; }
```
The central problem is that `CounterBase_increment` here is not type-compatible what we expect to find in the witness table for `ResetableCounter : ICounter`: the `this` parameter has the wrong type!
The basic solution strategy here is to intercept the search for a witness to sastify an interface requirement in `findWitnessForInterfaceRequirement` (those witnesses get collected into a witness table). The revised logic first looks for an exact match, which will only consider members introduced for the type itself, and not those introduced by base types.
If an exact match for a method requirement is not found, the semantic checker then tries to *synthesize* a witness for the requirement, which more or less amounts to generating a function like:
```hlsl
[mutating] void ResetableCounter::synthesized_increment()
{
this.increment();
}
```
The body of that synthesized method will type-check just fine in this case (because it desugars into `this.base.increment()`, more or less), and thus the synthesized method declaration can be used as the actual witness that drives downstream code generation.
Details:
* I added some options to lookup to allow us to explicitly skip member lookup through base interfaces; this should make sure that we don't accidentally satisfy an interface requirement using a member of the same or another interface (since such members are conceptually `abstract`).
* As it originally stood, the semantic checker was allowing `CounterBase.increment()` to satisfy the `increment()` requirement of `ResetableCounter` directly, with the result that we got invalid HLSL/GLSL code as output. In order to avoid this and other bad cases, I made sure that the "exact match" case of requirement satisfaction ignores members that included any "breadcrumbs" in the lookup result item (since the breadcrumbs would all indicate transformations that needed to be applied to `this` to find the right member).
* If we eventually have targets where `this` is passed by pointer/reference in all cases, then all of this work is not needed for the common case of single inheritance, and the base-type method should be usable as a witness directly. I don't see any easy way to handle that special case without producing target-dependent code in the front-end. It might be that we need an IR pass that can detect functions that are trivial "forwarding" functions and replace them with the function they forward to.
* This change includes a test case that should have come along with the original PR that started adding struct inheritance
Caveats:
* The comments in this change talk about things like allowing a method with a default parameter to satisfy a requirement without that parameter. That scenario won't actually work at present because we still have an enormous hack in our logic for checking methods against requirements: we don't actually consider their signatures! I couldn't fold a fix for that issue into this change because there are subtle corner cases around associated types that we need to handle correctly (which were part of the reason why the checking is as hacked as it is)
* This change does *not* try to test or address the case where we want to have a `Derived` type conform to `ISomething` because it inherits from `Base` and `Base : ISomething`. That case has its own details that need to be worked out, but ideally can follow a similar implementation strategy when it comes to re-using methods from `Base` to satisfy requirement on `Derived`.
|
|
|
|
* Associate a downstream compiler for prelude lookup even if output is source.
* Remove LanguageStyle and just use SourceLanguage instread.
* Added set/getPrelude.
Made prelude work on source language.
* Fix typo in method name replacement.
get/SetPrelude get/setLanguagePrelude
* Fix issue because of method name change.
* Remove getPreludeDownstreamCompilerForTarget
|
|
* * Fix output in slang repro command line
* Profile uses lowerCamel method names (had mix of upper and lower)
* Rename slang-serialize-state/SerializeStateUtil to slang-repro and ReproUtil.
|
|
* * Remove UniformState and UniformEntryPointParams types
* Put all output C++ source in an anonymous namespace
* If SLANG_PRELUDE_NAMESPACE is set, make what it defines available in generated file.
* Fix signature issue in performance-profile.slang
* Context -> KernelContext to avoid ambiguity.
* Fix issues around dynamic dispatch and anonymous namespace.
* Fix typo.
|
|
(#1395)
Co-authored-by: Tim Foley <tfoleyNV@users.noreply.github.com>
|
|
|
|
This was an oversight in the stdlib, and the `!=` definition follows the `==` in a straightforward fashion.
|
|
* Add IR pass to lower generics into ordinary functions.
* Fix project files
* Emit dynamic C++ code for simple generics and witness tables.
Fixes #1386.
* Remove -dump-ir flag.
* Fixups.
|
|
* Releases compile request if there is an error.
* Arrange so that caller can clean up CompileRequest so don't have to capture all paths.
|
|
|
|
Fixes #890.
Co-authored-by: Tim Foley <tfoleyNV@users.noreply.github.com>
|
|
* Remove implicit conversions to `void`
Fixes #1372
The standard library code had accidentally introduced implicit-conversion `__init` operations on the `void` type that accepted each of the other basic types, so that a function written like:
```hlsl
void bad() { return 1; }
```
would translate to:
```hlsl
void bad() { return (void)1; }
```
The dual problesm are that the input code should have produced a diagnostic of some kind, and the output code doesn't appear to compile correctly through fxc.
This change introduces several fixes aimed at this issue:
* First, the problem in the stdlib code is plugged: we don't introduce implicit conversion operations *to* or *from* `void` (we'd only been banning it in one direction before)
* Next, an explicit `__init` was added to `void` that accepts *any* type so that existing HLSL code that might do `(void) someExpression` to ignore a result will continue to work. This is a compatibility feature, and it might be argued that we should at least warn when it is used. Note that this function is expected to never appear in output HLSL/GLSL because its result will never be used, and it is marked `[__readNone]` allowing calls to it to be eliminated as dead code.
* During IR lowering, we now take care to only emit the `IRReturnVal` instruction type if there is a non-`void` value being returned, and use `IRReturnVoid` for both the case where no expression was used in the `return` statement *and* the case where an expression of type `void` is returned.
* A test case was added to confirm that returning `1` from a `void` function isn't allowed, while returning `(void) 1` *is*.
The net result of these changes is that we now produce an error for the bad input code, we allow explicit casts to `void` as a compatibility feature, and we are more robust about treating `void` as if it is an ordinary type in the front-end.
* fixup: missing file
|
|
|
|
IRWitnessTable values (#1387)
* Generate IRType for interfaces, and use them as the type of IRWitnessTable values.
This results the following IR for the included test case:
```
[export("_S3tu010IInterface7Computep1pii")]
let %1 : _ = key
[export("_ST3tu010IInterface")]
[nameHint("IInterface")]
interface %IInterface : _(%1);
[export("_S3tu04Impl7Computep1pii")]
[nameHint("Impl.Compute")]
func %Implx5FCompute : Func(Int, Int)
{
block %2(
[nameHint("inVal")]
param %inVal : Int):
let %3 : Int = mul(%inVal, %inVal)
return_val(%3)
}
[export("_SW3tu04Impl3tu010IInterface")]
witness_table %4 : %IInterface
{
witness_table_entry(%1,%Implx5FCompute)
}
```
* Fixes per code review comments.
Moved interface type reference in IRWitnessTable from their type to operand[0].
* Fix typo in comment.
|
|
|
|
* Diagnose circularly-defined constants
Work on #1374
This change diagnoses cases like the following:
```hlsl
static const int kCircular = kCircular;
static const int kInfinite = kInfinite + 1;
static const int kHere = kThere;
static const int kThere = kHere;
```
By diagnosing these as errors in the front-end we protect against infinite recursion leading to stack overflow crashes.
The basic approach is to have front-end constant folding track variables that are in use when folding a sub-expression, and then diagnosing an error if the same variable is encountered again while it is in use. In order to make sure the error occurs whether or not the constant is referenced, we invoke constant folding on all `static const` integer variables.
Limitations:
* This only works for integers, since that is all front-end constant folding applies to. A future change can/should catch circularity in constants at the IR level (and handle more types).
* This only works for constants. Circular references in the definition of a global variable are harder to diagnose, but at least shouldn't result in compiler crashes.
* This doesn't work across modules, or through generic specialization: anything that requires global knowledge won't be checked
* fixup: missing files
* fixup: review feedback
|
|
|
|
|
|
Fixes #1377
|
|
|
|
* Add a ASTBuilder to a Module
Only construct on valid ASTBuilder (was being called on nullptr on occassion)
* Add nodes to ASTBuilder.
* Compiles with RefPtr removed from AST node types.
* Initialize all AST node pointer variables in headers to nullptr;
* Initialize AST node variables as nullptr.
Make ASTBuilder keep a ref on node types.
Make SyntaxParseCallback returns a NodeBase
* Don't release canonicalType on dtor (managed by ASTBuilder).
* Give ASTBuilders a name and id, to help in debugging.
For now destroy the session TypeCache, to stop it holding things released when the compile request destroys ASTBuilders.
* Moved the TypeCheckingCache over to Linkage from Session.
* NodeBase no longer derived from RefObject.
* Only add/dtor nodes that need destruction.
First pass compile on linux.
|
|
|
|
|
|
|
|
* Fixes for active mask synthesis + tests
There are two fixes here:
* The code generation that follows active mask synthesis was requiring CUDA SM architecture version 7.0 for one of the introduced instructions, but not all of them. This change centralizes the handling of upgrading the required CUDA SM architecture version, and makes sure that the instructions introduced by active mask synthesis request version 7.0.
* The tests for active mask synthesis were not flagged as requiring the `cuda_sm_7_0` feature when invoking `render-test-tool`, which meant they run but produce unexpected results when invoked on a GPU without the required semantics for functions like `__ballot_sync()`. This change adds the missing `-render-feature cuda_sm_7_0` to those tests.
* fixup: mark more tests that rely on implicit active mask
|
|
Only construct on valid ASTBuilder (was being called on nullptr on occassion)
|
|
|
|
|
|
* First steps toward inheritance for struct types
This change adds the ability for a `struct` type to declare a base type that is another `struct`:
```hlsl
struct Base
{
int baseMember;
}
struct Derived : Base
{
int derivedMember;
}
```
The semantics of the feature are that code like the above desugars into code like:
```hlsl
struct Base
{
int baseMember;
}
struct Derived
{
Base _base;
int derivedMember;
}
```
At points where a member from the base type is being projected out, or the value is being implicitly cast to the base type, the compiler transforms the code to reference the implicitly-generated `_base` member. That means code like this:
```hlsl
void f(Base b);
...
Derived d = ...;
int x = d.baseMember;
f(d);
```
gets transformed into a form like this:
```hlsl
void f(Base b);
...
Derived d = ...;
int x = d._base.baseMember;
f(d._base);
```
Note that as a result of this choice, the behavior when passing a `Derived` value to a function that expects a `Base` (including to inherited member functions) is that of "object shearing" from the C++ world: the called function can only "see" the `Base` part of the argument, and any operations performed on it will behave as if the value was indeed a `Base`. There is no polymorphism going on because Slang doesn't currently have `virtual` methods.
In an attempt to work toward inheritance being a robust feature, this change adds a bunch of more detailed logic for checking the bases of various declarations:
* An `interface` declaration is only allowed to inherit from other `interface`s
* An `extension` declaration can only introduce inheritance from `interface`s
* A `struct` declaration can only inherit from at most one other `struct`, and that `struct` must be the first entry in the list of bases
This change also adds a mechanism to control whether a `struct` or `interface` in one module can inherit from a `struct` or `interface` declared in another module:
* If the base declaration is marked `[open]`, then the inheritance is allowed
* If the base declaration is marked `[sealed]`, then the inheritance is allowed
* If it is not marked otherwise, a `struct` is implicitly `[sealed]`
* If it is not marked otherwise, an `interface` is implicitly `[open]`
These seem like reasonable defaults. In order to safeguard the standard library a bit, the interfaces for builtin types have been marked `[sealed]` to make sure that a user cannot declare a `struct` and then mark it as a `BuiltinFloatingPointType`. This step should bring us a bit closer to being able to document and expose these interfaces for built-in types so that users can write code that is generic over them.
There are some big caveats with this work, such that it really only represents a stepping-stone toward a usable inheritance feature. The most important caveats are:
* If a `Derived` type tries to conform to an interface, such that one or more interface requirements are satisfied with members inherited from the `Base` type, that is likely to cause a crash or incorrect code generation.
* If a `Derived` type tries to inherit from a `Base` type that conforms to one or more interfaces, the witness table generated for the conformance of `Derived` to that interface is likely to lead to a crash or incorrect code generation.
It is clear that solving both of those issues will be necessary before we can really promote `struct` inheritance as a feature for users to try out.
* fixup: trying to appease clang error
* fixups: review feedback
|
|
* Make getSup work with more general non-virtual 'virtual' mechanism.
* WIP: Non virtual AST types.
* Project change.
* Type doesn't implement equalsImpl
* Fix macro invocation
Make Overridden functions public to make simply accessible by base types.
* Use SLANG_UNEXPECTED.
* GetScalarType -> getScalarType
Use SLANG_UNEXPECTED instead on ASSERT in NamedExpressionType and TypeType
|
|
|