Skip to main content

Browse the archive

https://github.com/shader-slang/slang
06 April 2024, 09:30:39 UTC
sort by:
Revision Author Date Message Commit Date
dfc9100 jsmall-nvidia 02 July 2020, 20:32:16 UTC Bug fix in C++ extractor (#1429) * Fix bug from change in diagnostics. * Catch exceptions and display a message on problem with C++ extractor. 02 July 2020, 20:32:16 UTC
5fc0185 Tim Foley 02 July 2020, 18:45:59 UTC Attempt to silence some warnings (#1428) * Attempt to silence some warnings This is an attempt to change code in `slang-ast-serialize.cpp` so that it doesn't trigger a warning(-as-error) in one of our build configurations. The original code is fine in terms of expressing its intent, so the right answer might actually be to silence the warning. * fixup: make sure to actually initialize 02 July 2020, 18:45:59 UTC
54675a3 jsmall-nvidia 02 July 2020, 13:09:52 UTC Only call m_api functions if m_api has been validly set on dtor of VulkanDeviceQueue. (#1426) 02 July 2020, 13:09:52 UTC
6cbb88f jsmall-nvidia 01 July 2020, 21:44:46 UTC Disable dynamic dispatch tests on CUDA - as fails with exception about unhandled op. (#1425) 01 July 2020, 21:44:46 UTC
8c33e7b jsmall-nvidia 01 July 2020, 20:11:56 UTC Ignore tests that don't have all the rendering APIs they require available. (#1419) 01 July 2020, 20:11:56 UTC
5c15329 jsmall-nvidia 01 July 2020, 18:20:42 UTC Fix bug in slang-dxc-support where it didn't get the source path correctly (#1420) * Fix handling of UniformState from #1396 * * Fix bug in slang-dxc-support where it didn't get the source path correctly * Make entryPointIndices const List<Int>& 01 July 2020, 18:20:42 UTC
69a0595 jsmall-nvidia 01 July 2020, 16:45:02 UTC Fix handling of UniformState from #1396 (#1417) 01 July 2020, 16:45:02 UTC
8ced9d2 Tim Foley 30 June 2020, 19:25:27 UTC Clean up unused code for IR object ownership (#1416) There was a small but non-trivial amount of code across `IRModule`, the `ObjectScopeManager`, and `StringRepresentationCache` that had to do with managing the lifetimes of `RefObject`s that might be referenced by IR instructions (and thus need to be kept alive for the lifetime of the IR module). We have long since migrated to a model where IR instruction do not include owned references to `RefObject`s, so these facilities weren't actually needed. This streamlines `IRModule`'s declaration, and trims code that we aren't actually using. One note for the future is that the `StringRepresentationCache` no longer does what its name implies (it is not a cache of `StringRepresentation`s), so we should consider giving it a more narrowly scoped name. I didn't include that in this change because I wanted to keep the diffs narrow and easy to review. A follow-on renaming change should be trivial if/when we can agree on what the type should be called at this point. Alternatively, we could simply bake the functionality of `StringRepresentationCache` into he IR deserialiation logic itself, since that is the only code using it. 30 June 2020, 19:25:27 UTC
dc44b08 Tim Foley 30 June 2020, 17:01:09 UTC Initial work on property declarations (#1410) * Initial work on property declarations Introduction ============ The main feature added here is support for `property` declarations, which provide a nicer experience for working with getter/setter pairs. If existing code had something like this: ```hlsl struct Sphere { float4 centerAndRadius; // xyz: center, w: radius float3 getCenter() { return centerAndRadius.xyz; } void setCenter(float3 newValue) { centerAndRadius.xyz = newValue; } // similarly for radius... } void someFunc(in out Sphere s) { float3 c = s.getCenter(); s.setCenter(c + offset); } ``` It can be expressed instead using a `property` declaration for `center`: ```hlsl struct Sphere { float4 centerAndRadius; // xyz: center, w: radius property center : float3 { get { return centerAndRadius.xyz; } set(newValue) { centerAndRadius.xyz = newValue; } } // similarly for radius... } void someFunc(in out Sphere s) { float3 c = s.center; s.center = c + offset; } ``` The benefits at the declaration site aren't that signficiant (e.g., in the example above we actually have slightly more lines of code), but the improvement in code clarity for users is significant. Having `property` declarations should also make it easier to migrate from a simple field to a property with more complex logic without having to first abstract the use-site code using a getter and setter. An important future benefit of `property` syntax will be if we allow `interface`s to include `property` requirements, and then also allow those requirements to be satisfied by ordinary fields in concrete types. Subscripts ---------- The Slang compiler already has limited (stdlib-use-only) support for `__subscript` declarations, which are conceptually similar to `operator[]` from the C++ world, but are expressed in a way that is more in line with `subscript` declarations in Swift. A `SubscriptDecl` in the AST contains zero or more `AccessorDecl`s, which correspond to the `get` and `set` clauses inside the original declaration (there is also a case for a `__ref` accessor, to handle the case where access needs to return a single address/reference that can be atomically mutated). A major goal of the implementation here is to re-use as much of the infrastructure as possible for `__subscript` declarations when implementing `property` declarations. Nonmutating Setters ------------------- One additional thing added in this change is the ability to mark a `set` accessor on either a subscript or a property as `[nonmutating]`, and indeed all of the existing `set` accessors declared in the stdlib have been marked this way. The need for this modifier is a bit subtle. If we think about a typical subscript or property: ```hlsl struct MyThing { int f; property p : int { get { return f; } set(newValue) { f = newValue; } } } ``` it is clear we want the `set` accessor to translate to output HLSL as something like: ``` void MyThing_p_set(inout MyThing this, int newValue) { this.f = newValue; } ``` Note how the implicit `this` parameter is `inout` even though we didn't mark anything as `[mutating]`. This is the obvious thing a user would expect us to generate given a property declaration. Now consider a case like the following: ```hlsl struct MyThing { RWStructuredBuffer<int> storage; property p : int { get { return storage[0]; } set(newValue) { storage[0] = newValue; } } } ``` This new declaration doesn't require (or want) an `inout` `this` parameter at all: ``` void MyThing_p_set(MyThing this, int newValue) { this.storage[0] = newValue; } ``` In fact, given the limitations in the current Slang compiler around functions that return resource types (or use them for `inout` parameters), we can only support a `set` operation like this if we can ensure that the `this` parameter is considered to be `in` instead of `inout`. This is exactly the behavior we allow users to opt into with a `[nonmutating] set` declaration. All of the subscript operations in the stdlib today have `set` accessors that don't actually change the value of `this` that they act on (e.g., storing into a `RWStructuredBuffer` using its `operator[]` doesn't change the value of the `RWStructuredBuffer` variable -- just its contents). We'd gotten away without this detail so far just because `set` accessors were only being declared in the stdlib and they were all implicitly `[nonmutating]` anyway, so it never surfaced as an issue that the code we generated assumed a setter wouldn't change `this`. Implementation ============== Parser and AST -------------- Adding a new AST node for `PropertyDecl` and the relevant parsing logic was mostly straightforward. The biggest change was allowing a `set` declaration to introduce an explicit name for the parameter that represents the new value to be set. This change also adds a `[nonmutating]` attribute as a dual to `[mutating]`, for reasons I will get to later. Semantic Checking ----------------- The `getTypeForDeclRef` logic was updated to allow references to `property` declarations. Some of the semantic checking work for subscripts was pulled out into re-usable subroutines to allow it to be shared by `__subscript` and `property` declarations. The checking of accessor declarations, which sets their result type based on the type of the outer `__subscript` was changed to also handle an outer `property`. Some special-case logic was added for checking of `set` declarations to make sure that their parameter is given the expected type. Some logic around deciding whether or not `this` is mutable had to be updated to correctly note that `this` should be mutable by default in a `set` accessor, with an explicit `[nonmutating]` modifier required to opt out of this default. (This is the inverse of how a typical method or `get` accessor works). IR Lowering ----------- The good news is that after IR lowering, access to properties turns into ordinary function calls (equivalent to what hand-written getters and setters would produce), so that subsequent compiler steps (including all the target-specific emit logic) doesn't have to care about the new feature. The bad news is that adding `property` declarations has revealed a few holes in how IR lowering was handling `__subscript` declarations and their accessors, so that it didn't trivially work for the new case as-is. The IR lowering pass already has the `LoweredValInfo` type that abstractly represents a value that resulted from lowering some AST code to the IR. One of the cases of `LoweredValInfo` was `BoundSubscript` that represented an expression of the form `baseVal[someIndex]` where the AST-level expression referenced a `__subscript` declaration. The key feature of `BoundSubscript` is that it avoided deciding whether to invoke the getter, the setter, or both "too early" and instead tried to only invoke the expected/required operations on-demand. This change generalizes `BoundSubscript` to handle `property` references as well, so it changes to `BoundStorage`. Making the type handle user-defined property declarations required fixing a bunch of issues: * When building up argument lists in the IR, we need to know whether an argument corresponds to an `in` or an `out`/`inout` parameter, to decide whether to pass the value directly or a pointer to the value. Some of the logic in the lowering pass had been playing fast and loose with this, so this change tries to make sure that whenever we care computing a list of `IRInst*` that represent the arguments to a call we have the information about the corresponding parameter. * Similarly, when emitting a call to an accessor in the IR, the information about the expected type of the callee was missing/unavailable, and the code was incorrectly building up the expected type of the callee based on the types of the arguments at the call site. The logic has been changed so that we can extract the expected signature of an accessor (how it will be translated to the IR) using the same logic that is used to produce the actual `IRFunc` for the accessor (so hopefully both will always agree). * Dealing with `in` vs. `inout` differences around parameters means also dealing with the "fixup" code that is used to assign from the temporary used to pass an `inout` argument back into the actual l-value expression that was used. That logic has all been hoisted out of the expression visitor(s) and into the global scope. Future Work =========== The entire approach to handling l-values in the IR lowering pass is broken, and it is in need a of a complete rewrite based on new first-principles design goals. While something like `LoweredValInfo` is decent for abstracting over the easy cases of r-values, addresses, and a few complicated l-value cases like swizzling, it just doesn't scale to highly abstract l-values like we get from `__subcript` and `property` declarations, nor other corner cases of l-values that we need to handle (e.g., passing an `int` to an `inout float` parameter is allowed in HLSL, and performs conversions in both directions!). It Should be Easy (TM) to extend the logic that tries to synthesize an interface conformance witness method when there isn't an exact match to also support synthesizing a property declaration (plus its accessors) to witness a required property when the type has a field of the same name/type. * fixup: pedantic template parsing error (thanks, clang!) * fixup: cleanups and review feedback * Removed some `#ifdef`'d out code from merge change * Added proper diagnostics for accessor parameter constraints, which led to some fixes/refactorings * Added a test case for the accessor-related diagnostics 30 June 2020, 17:01:09 UTC
47b43f8 Dietrich Geisler 29 June 2020, 21:42:12 UTC Backend for Multiple Entry Points (#1411) * 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> 29 June 2020, 21:42:12 UTC
3e8bdb6 Yong He 26 June 2020, 18:59:33 UTC Merge pull request #1408 from csyonghe/dyndispatch2 Dynamic dispatch for generic interface requirements and `associatedtype` 26 June 2020, 18:59:33 UTC
4e44398 Yong He 26 June 2020, 17:20:01 UTC Merge remote-tracking branch 'official/master' into dyndispatch2 26 June 2020, 17:20:01 UTC
d084f63 jsmall-nvidia 26 June 2020, 16:40:31 UTC AST serialize improvements (#1412) * 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. 26 June 2020, 16:40:31 UTC
4cf7119 Yong He 26 June 2020, 02:56:39 UTC Add a TODO comment for generic interface requirement key 26 June 2020, 02:56:39 UTC
dd88ba1 Yong He 26 June 2020, 02:03:51 UTC Fixes 26 June 2020, 02:03:51 UTC
5b57195 Yong He 26 June 2020, 01:26:12 UTC Fixes. 26 June 2020, 01:26:12 UTC
09c64ac Yong He 25 June 2020, 23:29:39 UTC Merge remote-tracking branch 'official/master' into dyndispatch2 25 June 2020, 23:29:39 UTC
218a39b Yong He 25 June 2020, 23:29:07 UTC remove ThisPointerDecoration, generate IRInterfaceType in one pass 25 June 2020, 23:29:07 UTC
509e36b Yong He 25 June 2020, 21:01:33 UTC Remove interfaceType operand from lookup_witness_method inst 25 June 2020, 21:01:33 UTC
892acc4 jsmall-nvidia 25 June 2020, 20:41:14 UTC AST Serialize Reading (#1409) * 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. 25 June 2020, 20:41:14 UTC
a1fed5e Yong He 25 June 2020, 20:19:45 UTC Partial fixes to code review comments 25 June 2020, 20:23:28 UTC
ffa9a35 Yong He 25 June 2020, 01:09:40 UTC Fix `lowerFuncType` and small bug fixes. 25 June 2020, 03:25:49 UTC
161c525 Yong He 24 June 2020, 21:22:52 UTC Fixes. 25 June 2020, 01:10:15 UTC
0ca75fe Yong He 24 June 2020, 20:16:11 UTC Dynamic dispatch for generic interface requirements. -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. 25 June 2020, 01:10:15 UTC
3fe4f53 Dietrich Geisler 24 June 2020, 21:22:58 UTC Heterogeneous example (#1399) * Introduced heterogeneous example. Example includes C++ source and header files, and does not currently make use of the associated slang file when building. The intent of this commit is to introduce the example as a baseline for later updates as the heterogeneous model is expanded. * Changing namespace * Renamed and rewrote README * Updated example to account for compiler updates * Updated path Co-authored-by: Tim Foley <tfoleyNV@users.noreply.github.com> 24 June 2020, 21:22:58 UTC
ae41db8 jsmall-nvidia 24 June 2020, 17:56:06 UTC AST Serialization writing (#1407) * 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. 24 June 2020, 17:56:06 UTC
b595dd0 Yong He 19 June 2020, 20:16:22 UTC Merge pull request #1403 from tfoleyNV/struct-inheritance-and-interfaces Work on struct inheritance and interfaces 19 June 2020, 20:16:22 UTC
1988011 Tim Foley 19 June 2020, 18:50:15 UTC fixup: review feedback 19 June 2020, 18:50:15 UTC
11e377a Tim Foley 19 June 2020, 18:48:28 UTC Merge remote-tracking branch 'origin/master' into struct-inheritance-and-interfaces 19 June 2020, 18:48:28 UTC
110d15b Yong He 19 June 2020, 18:19:51 UTC Dynamic dispatch for static member functions of associatedtypes. (#1404) 19 June 2020, 18:19:51 UTC
fc4342b Tim Foley 19 June 2020, 16:25:44 UTC fixup: actually make the test case test something 19 June 2020, 16:25:44 UTC
5fbb9ff Yong He 19 June 2020, 06:15:39 UTC Merge pull request #1401 from jsmall-nvidia/feature/prelude-fix Prelude fix/disable memaccess warning on gcc 19 June 2020, 06:15:39 UTC
0eddf45 Tim Foley 17 June 2020, 21:55:46 UTC Work on 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`. 18 June 2020, 23:00:40 UTC
515d8eb Yong He 18 June 2020, 22:06:14 UTC Merge branch 'master' into feature/prelude-fix 18 June 2020, 22:06:14 UTC
aa6aca4 Yong He 18 June 2020, 22:05:58 UTC Merge pull request #1400 from csyonghe/dyndispatch Dynamic dispatch non-static functions. 18 June 2020, 22:05:58 UTC
82ba914 Tim Foley 18 June 2020, 20:40:08 UTC Merge branch 'master' into dyndispatch 18 June 2020, 20:40:08 UTC
5952e3b jsmall-nvidia 18 June 2020, 20:39:06 UTC Prelude is associated with SourceLanguage (#1398) * 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 18 June 2020, 20:39:06 UTC
e2d2102 jsmall-nvidia 18 June 2020, 18:46:14 UTC Try using cmath or math.h depending on compiler to avoid issues around isinf etc. 18 June 2020, 18:46:14 UTC
dfbe3cf jsmall-nvidia 18 June 2020, 18:17:57 UTC Fix and improvements around repro (#1397) * * 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. 18 June 2020, 18:17:57 UTC
48da3ed jsmall-nvidia 18 June 2020, 16:44:51 UTC #include <cmath> Use SLANG_PRELUDE_STD macro to prefix functions that may need to be specified in std:: namespace. 18 June 2020, 16:44:51 UTC
5a86cd4 jsmall-nvidia 18 June 2020, 15:38:30 UTC Improvements around C++ code generation (#1396) * * 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. 18 June 2020, 15:38:30 UTC
f9b5f18 jsmall-nvidia 18 June 2020, 15:36:58 UTC * Fix warnings from prelude * Make compilation work on gcc by disabling -Wclass-mem-access 18 June 2020, 15:36:58 UTC
31ae346 jsmall-nvidia 18 June 2020, 12:10:47 UTC Associate a downstream compiler for prelude lookup even if output is source. (#1395) Co-authored-by: Tim Foley <tfoleyNV@users.noreply.github.com> 18 June 2020, 12:10:47 UTC
8c6e02b Yong He 18 June 2020, 05:41:56 UTC Dynamic dipatch non-static functions. 18 June 2020, 06:47:11 UTC
d1a8cd2 Tim Foley 17 June 2020, 23:30:18 UTC Add != operator for enum types (#1394) This was an oversight in the stdlib, and the `!=` definition follows the `==` in a straightforward fashion. 17 June 2020, 23:30:18 UTC
cd7f01b Yong He 17 June 2020, 20:08:27 UTC Generate dynamic C++ code for the minimal test case. (#1391) * 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. 17 June 2020, 20:08:27 UTC
ca503d4 jsmall-nvidia 17 June 2020, 16:15:29 UTC Hotfix/slangc unreleased compile request (#1393) * Releases compile request if there is an error. * Arrange so that caller can clean up CompileRequest so don't have to capture all paths. 17 June 2020, 16:15:29 UTC
40370ac Yong He 16 June 2020, 19:15:26 UTC Merge pull request #1392 from tfoleyNV/premake-bat Add a batch file for invoke premake 16 June 2020, 19:15:26 UTC
aa925d3 Tim Foley 16 June 2020, 16:54:06 UTC Add a batch file for invoke premake This change adds `./premake.bat` to the repository, which users in Windows (64-bit) can use to conveniently invoke the copy of premake that is pulled via the `slang-binaries` submodule. It should be possible to pass whatever options you passed to `premake5.exe` through to `premake.bat`. E.g., if you invoke: ``` .\premake.bat vs2015 ``` then you should get the desired results for the project/solution files we want to check in. 16 June 2020, 16:54:06 UTC
8ec293c Yong He 16 June 2020, 01:10:02 UTC Merge pull request #1390 from csyonghe/glsl-loop Emit [[dont_unroll]] GLSL attribute for [loop] attribute. 16 June 2020, 01:10:02 UTC
926e4bb Tim Foley 15 June 2020, 20:56:11 UTC Merge branch 'master' into glsl-loop 15 June 2020, 20:56:11 UTC
3461ed4 Yong He 15 June 2020, 20:55:56 UTC Specialize function calls involving array arguments. (#1389) Fixes #890. Co-authored-by: Tim Foley <tfoleyNV@users.noreply.github.com> 15 June 2020, 20:55:56 UTC
d84cfb7 Tim Foley 15 June 2020, 19:05:04 UTC Remove implicit conversions to `void` (#1388) * 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 15 June 2020, 19:05:04 UTC
7e7425d Yong He 15 June 2020, 16:05:49 UTC Merge branch 'master' into glsl-loop 15 June 2020, 16:05:49 UTC
90444f8 Yong He 15 June 2020, 16:04:53 UTC Generate IRType for interfaces, and reference them as `operand[0]` in 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. 15 June 2020, 16:04:53 UTC
04a81ab Yong He 13 June 2020, 07:24:21 UTC Emit [[dont_unroll]] attribute in GLSL 13 June 2020, 07:25:12 UTC
36a06f1 Tim Foley 12 June 2020, 20:30:32 UTC Diagnose circularly-defined constants (#1384) * 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 12 June 2020, 20:30:32 UTC
2359921 Yong He 12 June 2020, 00:13:27 UTC Merge pull request #1383 from csyonghe/dyndispatch Add compiler flag to disable specialization pass. 12 June 2020, 00:13:27 UTC
8452129 Yong He 11 June 2020, 18:10:40 UTC Merge branch 'master' into dyndispatch 11 June 2020, 18:10:40 UTC
1c77c44 jsmall-nvidia 11 June 2020, 18:06:27 UTC Fix problem with C++ extractor ernoneous concating of type tokens (#1382) * 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. 11 June 2020, 18:06:27 UTC
8de0a2e Yong He 10 June 2020, 21:57:30 UTC Add compiler flag to disable specialization pass. 10 June 2020, 21:57:30 UTC
98459ba Yong He 09 June 2020, 17:35:26 UTC Merge pull request #1381 from csyonghe/master Generate .tar.gz file in linux release 09 June 2020, 17:35:26 UTC
00e0e25 Yong He 08 June 2020, 23:16:12 UTC Generate .tar.gz file in linux release 08 June 2020, 23:17:34 UTC
78696a6 jsmall-nvidia 08 June 2020, 19:28:48 UTC Small fixes/improvements based on review. (#1379) 08 June 2020, 19:28:48 UTC
b3fbb92 Yong He 08 June 2020, 16:28:03 UTC Merge pull request #1378 from csyonghe/fix Filter lookup results from interfaces in `visitMemberExpr`. 08 June 2020, 16:28:03 UTC
956ede9 Yong He 06 June 2020, 02:38:43 UTC Filter lookup results from interfaces in `visitMemberExpr`. Fixes #1377 06 June 2020, 02:43:30 UTC
7d4432b Yong He 06 June 2020, 02:34:55 UTC Merge pull request #1375 from csyonghe/findtypebynamefix Fix FindTypeByName reflection API not finding stdlib types. 06 June 2020, 02:34:55 UTC
52026c7 Yong He 06 June 2020, 01:34:24 UTC Merge branch 'master' into findtypebynamefix 06 June 2020, 01:34:24 UTC
43c1467 jsmall-nvidia 05 June 2020, 22:20:09 UTC ASTNodes use MemoryArena (#1376) * 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. 05 June 2020, 22:20:09 UTC
92fc3aa Yong He 05 June 2020, 20:01:06 UTC Merge branch 'master' into findtypebynamefix 05 June 2020, 20:01:06 UTC
e3e1cf2 Yong He 05 June 2020, 20:00:32 UTC Merge pull request #1371 from csyonghe/loop_attrib Emit [loop] attribute to output HLSL. 05 June 2020, 20:00:32 UTC
389be08 Yong He 05 June 2020, 19:57:57 UTC Fix FindTypeByName reflection API not finding stdlib types. 05 June 2020, 19:57:57 UTC
00db821 Tim Foley 05 June 2020, 17:00:29 UTC Merge branch 'master' into loop_attrib 05 June 2020, 17:00:29 UTC
3bb7807 Tim Foley 05 June 2020, 17:00:15 UTC Fixes for active mask synthesis + tests (#1370) * 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 05 June 2020, 17:00:15 UTC
ecac0c7 jsmall-nvidia 04 June 2020, 21:15:38 UTC Add a ASTBuilder to a Module (#1369) Only construct on valid ASTBuilder (was being called on nullptr on occassion) 04 June 2020, 21:15:38 UTC
899824e Yong He 04 June 2020, 21:07:30 UTC Add missing loop-attribs.slang.hlsl for the test case 04 June 2020, 21:07:30 UTC
2c30d36 Yong He 04 June 2020, 21:06:10 UTC Remove aborting in emitLoopControlDecoration default case. 04 June 2020, 21:06:10 UTC
5e52b33 Yong He 04 June 2020, 21:04:54 UTC Change loop-attrib test case to CROSS_COMPILE per review comments 04 June 2020, 21:04:54 UTC
4e2f277 Yong He 04 June 2020, 19:42:17 UTC Emit [loop] attribute to output HLSL. 04 June 2020, 19:42:17 UTC
f3d637b Tim Foley 04 June 2020, 18:53:13 UTC First steps toward inheritance for struct types (#1366) * 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 04 June 2020, 18:53:13 UTC
1b8731c jsmall-nvidia 03 June 2020, 21:22:48 UTC Devirtualize AST types (#1368) * 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 03 June 2020, 21:22:48 UTC
4e3e7f2 jsmall-nvidia 03 June 2020, 13:41:31 UTC Disable CUDA active mask tests as failing on CI. (#1367) 03 June 2020, 13:41:31 UTC
d386e27 jsmall-nvidia 02 June 2020, 19:26:51 UTC Added spGetBuildTagString. (#1365) 02 June 2020, 19:26:51 UTC
f87b632 jsmall-nvidia 02 June 2020, 18:05:35 UTC Make stdlib path just be the filename. (#1364) * Made bad-operaor-call available on all targets. Fix the line filename to not inclue path, to avoid paths being absolute and therefores value be host environment dependent (causing tests to fail). * Disable on linux because theres still a problem on gcc x86 where the file path is different. * Fix to some typos in bad-operator-call.slang * Fix diagnostic for bad-operator-call.slang 02 June 2020, 18:05:35 UTC
926a0c5 Dietrich Geisler 02 June 2020, 16:12:35 UTC Working matrix swizzle (#1354) * Working matrix swizzle. Supports one and zero indexing and multiple elements. Performs semantic checking of the swizzle. Matrix swizzles are transformed into a vector of indexing operations during lowering to the IR. This change does not handle matrix swizzle as lvalues. * Renaming * Added missing semicolon * Initialize variable for gcc * Added the expect file for diagnostics * Matrix swizzle updated per PR feedback * Stylistic fix * Formatting fixes * Fix compiling with AST change. Change indentation. Co-authored-by: jsmall-nvidia <jsmall@nvidia.com> 02 June 2020, 16:12:35 UTC
8acb704 jsmall-nvidia 29 May 2020, 19:56:28 UTC Bug fix problem with ray tracing from fragment shader (#1362) * Added GLSL_460 if ray tracing is used on fragment shader. Moved GLSL specific setup init function. * Split out _requireRayTracing method. 29 May 2020, 19:56:28 UTC
9773495 jsmall-nvidia 29 May 2020, 18:26:48 UTC NodeBase types constructed with astNodeType member set (#1363) * Maked Substituions derived from NodeBase * * Add astNodeTYpe field to NodeBase * Make Substitutions derived from NodeBase * Make all construction through ASTBuilder * Make getClassInfo non virtual (just uses the astNodeType) 29 May 2020, 18:26:48 UTC
45e414f jsmall-nvidia 29 May 2020, 15:01:24 UTC Update target compat doc to include Wave intrinsics for CUDA (#1361) 29 May 2020, 15:01:24 UTC
f3d7042 jsmall-nvidia 29 May 2020, 12:36:10 UTC Feature/ast syntax standard (#1360) * Small improvements to documentation and code around DiagnosticSink * Made methods/functions in slang-syntax.h be lowerCamel Removed some commented out source (was placed elsewhere in code) * Making AST related methods and function lowerCamel. Made IsLeftValue -> isLeftValue. 29 May 2020, 12:36:10 UTC
95597d7 jsmall-nvidia 28 May 2020, 19:32:19 UTC Small improvements to documentation and code around DiagnosticSink (#1359) 28 May 2020, 19:32:19 UTC
c2d3134 jsmall-nvidia 28 May 2020, 18:01:51 UTC WIP: ASTBuilder (#1358) * Compiles. * Small tidy up around session/ASTBuilder. * Tests are now passing. * Fix Visual Studio project. * Fix using new X to use builder when protectedness of Ctor is not enough. Substitute->substitute * Add some missing ast nodes created outside of ASTBuilder. * Compile time check that ASTBuilder is making an AST type. * Moced findClasInfo and findSyntaxClass (essentially the same thing) to SharedASTBuilder from Session. 28 May 2020, 18:01:51 UTC
e5d0f33 Tim Foley 26 May 2020, 22:11:22 UTC Synthesize "active mask" for CUDA (#1352) * Synthesize "active mask" for CUDA The Big Picture =============== The most important change here is to `hlsl.meta.slang`, where the declaration of `WaveGetActiveMask()` is changed so that instead of mapping to `__activemask()` on CUDA (which is semantically incorrect) it maps to a dedicated IR instruction. The other `WaveActive*()` intrinsics that make use of the implicit "active mask" concept had already been changed in #1336 so that they explicitly translate to call the equivalent `WaveMask*()` intrinsic with the result of `WaveGetActiveMask()`. As a result, all of the `WaveActive*()` functions are now no different from a user-defined function that uses `WaveGetActiveMask()`. The bulk of the work in this change goes into an IR pass to replace the new instruction for getting the active mask gets replaced with appropriately computed values before we generate output CUDA code. That work is in `slang-ir-synthesize-active-mask.{h,cpp}`. Utilities ========= There are a few pieces of code that were helpful in writing the main pass but that can be explained separately: * IR instructions were added corresponding to the Slang `WaveMaskBallot()` and `WaveMaskMatch()` functions, which map to the CUDA `__ballot_sync()` and `__match_any_sync()` operations, respectively. These are only implemented for the CUDA target because they are only being generated as part of our CUDA-only pass. * The `IRDominatorTree` type was updated to make it a bit more robust in the presence of unreachable blocks in the CFG. It is possible that the same ends could be achieved more efficiently by folding the corner cases into the main logic, but I went ahead and made things very explicit for now. * I added an `IREdge` utility type to better encapsulate the way that certain code operating on the predecessors/successors of an `IRBlock` were using an `IRUse*` to represent a control-flow edge. The `IREdge` type makes the logic of those operations more explicit. A future change should proably change it so that `IRBlock::getPredecessors()` and `getSuccessors()` are instead `getIncomingEdges()` and `getOutgoingEdges()` and work as iterators over `IREdge` values, given the way that the predecessor and successor lists today can contain duplicates. * Using the above `IREdge` type, the logic for detecting and break critical edges was broken down into something that is a bit more clear (I hope), and that also factors out the breaking of an edge (by inserting a block along it) into a reusable subroutine. The Main Pass ============= The implementation of the new pass is in `slang-ir-synthesize-active-mask.cpp`, and that file attempts to include enough comments to make the logic clear. A brief summary for the benefit of the commit history: * The first order of business is to identify functions that need to have the active mask value piped into them, and to add an additional parameter to them so that the active mask is passed down explicitly. Call sites are adjusted to pass down the active mask which can then result in new functions being identified as needing the active mask. * The next challenge is for a function that uses the active mask, to compute the active mask value to use in each basic block. The entry block can easily use the active mask value that was passed in, while other blocks need more work. * When doing a conditional branch, we can compute the new mask for the block we branch to as a function of the existing mask and the branch condition. E.g., the value `WaveMaskBallot(existingMask, condition)` can be used as the mask for the "then" block of an `if` statement. * When control flow paths need to "reconverge" at a point after a structured control-flow statement, we need to insert logic to synchronize and re-build the mask that will execute after the statement, while also excluding any lanes/threads that exited the statement in other ways (e.g., an early `return` from the function). The explanation here is fairly hand-wavy, but the actual pass uses much more crisp definitions, so the code itself should be inspected if you care about the details. Tests ===== The tests for the new feature are all under `tests/hlsl-intrinsic/active-mask/`. Most of them stress a single control-flow construct (`if`, `switch`, or loop) and write out the value of `WaveGetActiveMask()` at various points in the code. In practice, our definition of the active mask doesn't always agree with what D3D/Vulkan implementations seem to produce in practice, and as a result a certain amount of effort has gone into adding tweaks to the tests that force them to produce the expected output on existing graphics APIs. These tweaks usually amount to introducing conditional branches that aren't actually conditional in practice (the branch condition is always `true` or always `false` at runtime), in order to trick some simplistic analysis approaches that downstream compilers seem to employ. One test case currently fails on our CUDA target (`switch-trivial-fallthrough.slang`) and has been disabled. This is an expected failure, because making it produce the expected value requires a bit of detailed/careful coding that would add a lot of additional complexity to this change. It seemed better to leave that as future work. Future Work =========== * As discussed under "Tests" above, the handling of simple `switch` statements in the current pass is incomplete. * There's an entire can of worms to be dealt with around the handling of fall-through for `switch`. * The current work also doesn't handle `discard` statements, which is unimportant right now (CUDA doesn't have fragment shaders), but might matter if we decide to synthesize masks for other targets. Similar work would probably be needed if we ever have `throw` or other non-local control flow that crosses function boundaries. * An important optimization opportunity is being left on the floor in this change. When block that comes "after" a structured control-flow region (which is encoded explicitly in Slang IR and SPIR-V) post-dominates the entry block of the region, then we know that the active mask when exiting the region must be the same as the mask when entering the region, and there is no need to insert explicit code to cause "re-convergence." This should be addressed in a follow-on change once we add code to Slang for computing a post-dominator tree from a function CFG. * Related to the above, the decision-making around whether a basic block "needs" the active mask is perhaps too conservative, since it decides that any block that precedes one needing the active mask also needs it. This isn't true in cases where the active mask for a merge block can be inferred by post-dominance (as described above), so that the blocks that branch to it don't need to compute an active mask at all. * If/when we extend the CPU target to support these operations (along with SIMD code generation, I assume), we will also need to synthesize an active mask on that platform, but the approach taken here (which pretty much relies on support for CUDA "cooperative groups") wouldn't seem to apply in the SIMD case. * Similarly, the approach taken to computing the active mask here requires a new enough CUDA SM architecture version to support explicit cooperative groups. If we want to run on older CUDA-supporting architectures, we will need a new and potentially very different strategy. * Because the new pass here changes the signature of functions that require the active mask (and not those that don't), it creates possible problems for generating code that uses dynamic dispatch (via function pointers). In principle, we need to know at a call site whether or not the callee uses the active mask. There are multiple possible solutions to this problem, and they'd need to be worked through before we can make the implicit active mask and dynamic dispatch be mutually compatible. * Related to changing function signatures: no effort is made in this pass to clean up the IR type of the functions it modifies, so there could technically be mismatches between the IR type of a function and its actual signature. If/when this causes problems for downstream passes we probably need to do some cleanup. * fixup: backslash-escaped lines I did some "ASCII art" sorts of diagrams to explain cases in the CFG, and some of those diagrams used backslash (`\`) characters as the last character on the line, causing them to count as escaped newlines for C/C++. The gcc compiler apparently balked at those lines, since they made some of the single-line comments into multi-line comments. I solved the problem by adding a terminating column of `|` characters at the end of each line that was part of an ASCII art diagram. * fixup: typos Co-authored-by: jsmall-nvidia <jsmall@nvidia.com> 26 May 2020, 22:11:22 UTC
b136904 jsmall-nvidia 26 May 2020, 17:53:10 UTC Improvements around hashing (#1355) * Fields from upper to lower case in slang-ast-decl.h * Lower camel field names in slang-ast-stmt.h * Fix fields in slang-ast-expr.h * slang-ast-type.h make fields lowerCamel. * slang-ast-base.h members functions lowerCamel. * Method names in slang-ast-type.h to lowerCamel. * GetCanonicalType -> getCanonicalType * Substitute -> substitute * Equals -> equals ToString -> toString * ParentDecl -> parentDecl Members -> members * * Make hash code types explicit * Use HashCode as return type of GetHashCode * Added conversion from double to int64_t * Split Stable from other hash functions * toHash32/64 to convert a HashCode to the other styles. GetHashCode32/64 -> getHashCode32/64 GetStableHashCode32/64 -> getStableHashCode32/64 * Other Get/Stable/HashCode32/64 fixes * GetHashCode -> getHashCode * Equals -> equals * CreateCanonicalType -> createCanonicalType * Catches of polymorphic types should be through references otherwise slicing can occur. * Fixes for newer verison of gcc. Fix hashing problem on gcc for Dictionary. * Another fix for GetHashPos * Fix signed issue around GetHashPos 26 May 2020, 17:53:10 UTC
ee2ec68 Yong He 25 May 2020, 21:53:15 UTC Merge pull request #1356 from csyonghe/master Remove non-ascii characters from source file. 25 May 2020, 21:53:15 UTC
fd28dcf Yong He 25 May 2020, 20:07:26 UTC Remove non-ascii characters from source file. 25 May 2020, 20:07:26 UTC
076a4c0 jsmall-nvidia 22 May 2020, 18:21:37 UTC Tidy up around AST nodes (#1353) * Fields from upper to lower case in slang-ast-decl.h * Lower camel field names in slang-ast-stmt.h * Fix fields in slang-ast-expr.h * slang-ast-type.h make fields lowerCamel. * slang-ast-base.h members functions lowerCamel. * Method names in slang-ast-type.h to lowerCamel. * GetCanonicalType -> getCanonicalType * Substitute -> substitute * Equals -> equals ToString -> toString * ParentDecl -> parentDecl Members -> members 22 May 2020, 18:21:37 UTC
daf53bb jsmall-nvidia 21 May 2020, 19:36:02 UTC Non virtual accept implementation on AST types (#1351) * First pass impl of making accept on AST node types non virtual. * A single switch for ITypeVistor on Val type. * Use ORIGIN to choose ITypeVisitor dispatch. * Don't use ORIGIN - we don't need special handling for ITypeVisitor on Val derived types. 21 May 2020, 19:36:02 UTC
d90ae36 jsmall-nvidia 21 May 2020, 18:06:18 UTC AST dump improvements (#1350) * Add support for parsing array types to C++ extractor. * C++ extractor looks for 'balanced tokens'. Use for extracting array suffixes. * First pass at field dumping. * Update project for field dumping. * WIP AST Dumper. * More AST dump compiling. * Fix bug in StringSlicePool where it doesn't use the copy of the UnownedStringSlice in the map. * Add support for SLANG_RELFECTED and SLANG_UNREFLECTED More AST dump support. * Support for hierarchical dumping/flat dumping. Use SourceWriter to dump. * Add -dump-ast command line option. * Add fixes to VS project to incude AST dump. * Fix compilation on gcc. * Add fix for type ambiguity issue on x86 VS. * Fixes from merge of reducing Token size. * Fix comment about using SourceWriter. * Improvement AST dumping around * Pointers (write as hex) * Scope * Turn off some unneeded fields in AST hierarchy * Only output the initial module in full. 21 May 2020, 18:06:18 UTC
96a00c8 jsmall-nvidia 20 May 2020, 14:56:49 UTC AST dumping via C++ Extractor reflection (#1348) * Add support for parsing array types to C++ extractor. * C++ extractor looks for 'balanced tokens'. Use for extracting array suffixes. * First pass at field dumping. * Update project for field dumping. * WIP AST Dumper. * More AST dump compiling. * Fix bug in StringSlicePool where it doesn't use the copy of the UnownedStringSlice in the map. * Add support for SLANG_RELFECTED and SLANG_UNREFLECTED More AST dump support. * Support for hierarchical dumping/flat dumping. Use SourceWriter to dump. * Add -dump-ast command line option. * Add fixes to VS project to incude AST dump. * Fix compilation on gcc. * Add fix for type ambiguity issue on x86 VS. * Fixes from merge of reducing Token size. * Fix comment about using SourceWriter. 20 May 2020, 14:56:49 UTC
c54c957 jsmall-nvidia 19 May 2020, 19:37:40 UTC Reduce the size of Token (#1349) * Token size on 64 bits is 24 bytes (from 40). On 32 bits is 16 bytes from 24. * Added hasContent method to Token. Some other small improvements around Token. 19 May 2020, 19:37:40 UTC
back to top