slang.git/source/slang/slang-ir-explicit-global-context.cpp, branch master

Update function type after inserting KernelContext parameter (#8551)

2025-09-29T06:33:45+00:00

Without this, there are functions with missing parameters in their type in the IR after running the `introduceExplicitGlobalContext` pass: ``` [layout(%15)] [export("_SV4test12outputBuffer")] [nameHint("outputBuffer")] let %outputBuffer : _ = key [noSideEffect] [export("_S4test7dostuffp1pi_ff")] [nameHint("dostuff")] func %dostuff : Func(Float, Float) { block %34( [nameHint("f")] param %f : Float, [nameHint("kernelContext")] param %kernelContext : Ptr(%KernelContext, 0 : UInt64, 1 : UInt64)): let %35 : Float = mul(%f, %f) let %36 : Ptr(ConstantBuffer(%GlobalParams, DefaultLayout), 0 : UInt64, 1 : UInt64) = get_field_addr(%kernelContext, %globalParams) let %37 : ConstantBuffer(%GlobalParams, DefaultLayout) = load(%36) let %38 : Ptr(RWStructuredBuffer(Float, DefaultLayout, %20)) = get_field_addr(%37, %outputBuffer) let %39 : RWStructuredBuffer(Float, DefaultLayout, %20) = load(%38) let %40 : Ptr(Float) = rwstructuredBufferGetElementPtr(%39, 1 : Int) let %41 : Float = load(%40) let %42 : Float = mul(%35, %41) return_val(%42) } ``` Not sure why this doesn't seem to negatively affect existing targets, but it sure is an issue for the LLVM target I'm working on. I could've left this fix for that PR, but I want to check now if this causes any issues with the existing targets using the CI. This also happens with the entry point functions, where the function type is not updated after adding `ComputeThreadVaryingInput`. This had no effect in the C++ target because `convertEntryPointPtrParamsToRawPtrs(irModule);` is called right after and fixes it.

[CBP] Pointer frontend changes + groupshared pointer support (#7848)

2025-08-29T22:52:34+00:00

Resolves #7628 Resolves: #8197 Primary Goals: 1. Add `Access` to pointer 2. AddressSpace::GroupShared support for pointers (SPIR-V) 3. Add `__getAddress()` to replace `&` * `&` is not updated to `require(cpu)` since slangpy uses `&`. This means we must: (1) merge PR; (2) replace `&` with `__getAddress()`; (3) add `require(cpu)` to `&` Changes: * Added to `Ptr` the `Access` generic argument & logic (for `Access::Read`). * Moved the generic argument `AddressSpace` from `Ptr` to the end of the type. * Added pointer casting support between any `Ptr` as long as the `AddressSpace` is the same * Disallow globallycoherent T* and coherent T* * Disallow const T*, T const*, and const T* * Fixed .natvis display of `ConstantValue` `ValOperandNode` * Support generic resolution of type-casted integers * Added `VariablePointer` emitting for spirv + other minor logic needed for groupshared pointers Breaking Changes: * Anyone using the `AddressSpace` of `Ptr` will now have to account for the `Access` argument * we disallow various syntax paired with `Ptr` and `T*` --------- Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>

Address issues with GLSL style global in/out vars (#6669) (#6998)

2025-06-06T20:13:43+00:00

* 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>

Remove unnecessary parameters from Metal entry point signature (#6131)

2025-01-22T16:57:53+00:00

* fix metal entry point global params * address review comments, cleanup and test * remove dead code * undo accidental change * address review comments and cleanup * minor fix and cleanup --------- Co-authored-by: Yong He

Support specialization constant on WGSL and Metal. (#5780)

2024-12-06T08:55:35+00:00

Move switch statement bodies to their own lines (#5493)

2024-11-05T17:47:26+00:00

* Move switch statement bodies to their own lines * format --------- Co-authored-by: Yong He

format

2024-10-29T06:49:26+00:00

* format * Minor test fixes * enable checking cpp format in ci

Remove using SpvStorageClass values casted into AddressSpace values (#4861)

2024-08-19T22:06:34+00:00

* Remove using SpvStorageClass values casted into AddressSpace values Also removes support for specific storage classes in __target_intrinsic snippets * remove SLANG_RETURN_NEVER macro * squash warnings * Make nonexhaustive switch statement error on gcc * Add SLANG_EXHAUSTIVE_SWITCH_BEGIN/END macros --------- Co-authored-by: Yong He

Move SPIRV global variables into a context variable (#4741)

2024-07-31T03:04:08+00:00

Overhaul IR lowering of pointer types. (#4710)

2024-07-25T22:00:14+00:00

* Overhaul IR lowering of pointer types. * Propagate address space in IRBuilder. * Fixup. * Fix. * Fix. * Change how Ptr type is printed to text. * Fix.