CVE-2026-106281
Overview
Changed Functions
| Function | Change | Notes |
|---|---|---|
ifsrc/tint/lang/glsl/writer/printer/printer.cc |
modified | |
ifsrc/tint/lang/hlsl/writer/printer/printer.cc |
modified | |
ifsrc/tint/lang/msl/writer/printer/printer.cc |
modified |
Files Changed
src/tint/lang/hlsl/writer/printer/printer.ccsrc/tint/lang/msl/writer/printer/printer.ccsrc/tint/lang/glsl/writer/printer/printer.cc
Patch
From d0acacc83ce879a5076ce3c10c08e16b77e2e821 Mon Sep 17 00:00:00 2001 From: dan sinclair <[email protected]> Date: Tue, 01 Sep 2026 12:13:41 -0700 Subject: [PATCH] Fix issue with emit_continue block maintenance. The `emit_continuing_` callback in the printers is currently maintained through a TINT_SCOPED_ASSIGNMENT. This will store the old value away and put a new one in place. This works up until the point when you assign a function pointer while executing the outer function pointer. The act of copying away does a move/destroy of the current function, which is being executed. If this shows up as an issue depends on how the compiler structures the functions and other surrounding code. This has only been reproducible on certain printers under ASAN with a GCC build. Fixed: 554992296 Change-Id: If96d213ac9605871e5e84ec673fcd1df82c2642f Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/338575 Reviewed-by: Kai Ninomiya <[email protected]> Reviewed-by: James Price <[email protected]> Commit-Queue: dan sinclair <[email protected]> --- diff --git a/src/tint/lang/glsl/writer/printer/printer.cc b/src/tint/lang/glsl/writer/printer/printer.cc index c7ffb7d..5497bf6 100644 --- a/src/tint/lang/glsl/writer/printer/printer.cc +++ b/src/tint/lang/glsl/writer/printer/printer.cc @@ -27,8 +27,10 @@ #include "src/tint/lang/glsl/writer/printer/printer.h" +#include <memory> #include <string> #include <utility> +#include <vector> #include "src/tint/lang/core/constant/splat.h" #include "src/tint/lang/core/enums.h" @@ -228,7 +230,7 @@ struct_to_padding_struct_ids_; /// Block to emit for a continuing - std::function<void()> emit_continuing_; + std::vector<std::unique_ptr<std::function<void()>>> emit_continuing_; /// @returns `true` if @p ident should be renamed bool ShouldRename(std::string_view ident) { @@ -506,8 +508,9 @@ void EmitDiscard() { Line() << "discard;"; } void EmitContinue(const core::ir::Continue* c) { - if (emit_continuing_) { - emit_continuing_(); + if (!emit_continuing_.empty()) { + auto fn = emit_continuing_.back().get(); + (*fn)(); } if (c->Block() != c->Loop()->Body()) { Line() << "continue;"; @@ -527,18 +530,17 @@ // } // } - auto emit_continuing = [&] { - Line() << "{"; - { - const ScopedIndent si(current_buffer_); - EmitBlock(l->Continuing()); - } - Line() << "}"; - }; - TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing); - Line() << "{"; { + emit_continuing_.push_back(std::make_unique<std::function<void()>>([&] { + Line() << "{"; + { + const ScopedIndent si(current_buffer_); + EmitBlock(l->Continuing()); + } + Line() << "}"; + })); + ScopedIndent init(current_buffer_); EmitBlock(l->Initializer()); @@ -548,6 +550,8 @@ EmitBlock(l->Body()); } Line() << "}"; + + emit_continuing_.pop_back(); } Line() << "}"; } diff --git a/src/tint/lang/hlsl/writer/printer/printer.cc b/src/tint/lang/hlsl/writer/printer/printer.cc index beb8fae..e14f6a7 100644 --- a/src/tint/lang/hlsl/writer/printer/printer.cc +++ b/src/tint/lang/hlsl/writer/printer/printer.cc @@ -211,7 +211,7 @@ const core::ir::Block* current_block_ = nullptr; /// Block to emit for a continuing - std::function<void()> emit_continuing_; + std::vector<std::unique_ptr<std::function<void()>>> emit_continuing_; /// `true` if the linalg.h header has been included. bool linalg_included = false; @@ -494,8 +494,9 @@ } void EmitContinue(const core::ir::Continue* c) { - if (emit_continuing_) { - emit_continuing_(); + if (!emit_continuing_.empty()) { + auto fn = emit_continuing_.back().get(); + (*fn)(); } if (c->Block() != c->Loop()->Body()) { Line() << "continue;"; @@ -558,20 +559,19 @@ Line() << "while(true) {"; } - auto emit_continuing = [&] { - if (uses_for_loop_update) { - return; - } - Line() << "{"; - { - const ScopedIndent si(current_buffer_); - EmitBlock(l->Continuing()); - } - Line() << "}"; - }; - TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing); - { + emit_continuing_.push_back(std::make_unique<std::function<void()>>([&] { + if (uses_for_loop_update) { + return; + } + Line() << "{"; + { + const ScopedIndent si(current_buffer_); + EmitBlock(l->Continuing()); + } + Line() << "}"; + })); + const ScopedIndent si(current_buffer_); if (has_loop_condition) { EmitBlock(l->Body(), [&analysis](const core::ir::Instruction* inst) { @@ -580,6 +580,8 @@ } else { EmitBlock(l->Body()); } + + emit_continuing_.pop_back(); } Line() << "}"; } diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc index abc29ee..5452d56 100644 --- a/src/tint/lang/msl/writer/printer/printer.cc +++ b/src/tint/lang/msl/writer/printer/printer.cc @@ -30,8 +30,10 @@ #include <atomic> #include <cstddef> #include <cstdint> +#include <memory> #include <string> #include <utility> +#include <vector> #include "src/tint/lang/core/constant/splat.h" #include "src/tint/lang/core/constant/string.h" @@ -202,7 +204,7 @@ std::string array_template_name_; /// Block to emit for a continuing - std::function<void()> emit_continuing_; + std::vector<std::unique_ptr<std::function<void()>>> emit_continuing_; /// @returns the name of the templated `tint_array` helper type, generating it if needed const std::string& ArrayTemplateName() { @@ -730,8 +732,9 @@ } void EmitContinue(const core::ir::Continue* c) { - if (emit_continuing_) { - emit_continuing_(); + if (!emit_continuing_.empty()) { + auto fn = emit_continuing_.back().get(); + (*fn)(); } if (c->Block() != c->Loop()->Body()) { Line() << "continue;"; @@ -749,18 +752,17 @@ // } // } - auto emit_continuing = [&] { - Line() << "{"; - { - const ScopedIndent si(current_buffer_); - EmitBlock(l->Continuing()); - } - Line() << "}"; - }; - TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing); - Line() << "{"; { + emit_continuing_.push_back(std::make_unique<std::function<void()>>([&] { + Line() << "{"; + { + const ScopedIndent si(current_buffer_); + EmitBlock(l->Continuing()); + } + Line() << "}"; + })); + ScopedIndent init(current_buffer_); EmitBlock(l->Initializer()); @@ -770,6 +772,8 @@ EmitBlock(l->Body()); } Line() << "}"; + + emit_continuing_.pop_back(); } Line() << "}"; }
Original Bug Report
tint hlsl heap-use-after-free
Not sure if this is security-related or not. But considering the UAF, I’d like to report this as security bug. Feel free to downgrade this.
The following shader:
@compute @workgroup_size(1)
fn main() {
loop {
continuing {
loop { break; }
break if true;
}
}
}
Resulted in this output:
==181384==ERROR: AddressSanitizer: heap-use-after-free on address 0x503000025ff8 at pc 0x55f77e79ae1d bp 0x7ffe89951d80 sp 0x7ffe89951d70
READ of size 8 at 0x503000025ff8 thread T0
#0 0x55f77e79ae1c in operator() src/tint/lang/hlsl/writer/printer/printer.cc:570
#1 0x55f77e7e6694 in __invoke_impl<void, tint::hlsl::writer::(anonymous namespace)::Printer::EmitLoop(const tint::core::ir::Loop*)::<lambda()>&> /usr/include/c++/13/bits/invoke.h:61
#2 0x55f77e7e62fe in __invoke_r<void, tint::hlsl::writer::(anonymous namespace)::Printer::EmitLoop(const tint::core::ir::Loop*)::<lambda()>&> /usr/include/c++/13/bits/invoke.h:111
#3 0x55f77e7e4216 in _M_invoke /usr/include/c++/13/bits/std_function.h:290
#4 0x55f77e4b0a42 in std::function<void ()>::operator()() const /usr/include/c++/13/bits/std_function.h:591
#5 0x55f77e79a002 in EmitContinue src/tint/lang/hlsl/writer/printer/printer.cc:498
#6 0x55f77e7d6ed8 in operator() src/tint/lang/hlsl/writer/printer/printer.cc:354
#7 0x55f77e7d818e in operator()<tint::hlsl::writer::(anonymous namespace)::Printer::EmitBlock<tint::hlsl::writer::(anonymous namespace)::Printer::EmitBlock(const tint::core::ir::Block*)::<lambda(const tint::core::ir::Instruction*)>&>(const tint::core::ir::Block*, tint::hlsl::writer::(anonymous namespace)::Printer::EmitBlock(const tint::core::ir::Block*)::<lambda(const tint::core::ir::Instruction*)>&)::<lambda(const tint::core::ir::Continue*)> > src/tint/utils/rtti/switch.h:347
#8 0x55f77e7dae38 in Switch<> src/tint/utils/rtti/switch.h:358
#9 0x55f77e7dc30b in EmitBlock<tint::hlsl::writer::(anonymous namespace)::Printer::EmitBlock(const tint::core::ir::Block*)::<lambda(const tint::core::ir::Instruction*)>&> src/tint/lang/hlsl/writer/printer/printer.cc:346
#10 0x55f77e796f3c in EmitBlock src/tint/lang/hlsl/writer/printer/printer.cc:335
#11 0x55f77e79bf11 in EmitLoop src/tint/lang/hlsl/writer/printer/printer.cc:581
--- freed by ---
freed by thread T0 here:
#0 0x7f2522fa65e8 in operator delete(void*, unsigned long) ../../../../src/libsanitizer/asan/asan_new_delete.cpp:164
#1 0x55f77e7e6792 in _M_destroy /usr/include/c++/13/bits/std_function.h:175
#2 0x55f77e7e652e in _M_manager /usr/include/c++/13/bits/std_function.h:203
#3 0x55f77e7e437d in _M_manager /usr/include/c++/13/bits/std_function.h:282
#4 0x55f77da23e4c in std::_Function_base::~_Function_base() /usr/include/c++/13/bits/std_function.h:244
#5 0x55f77e4a0729 in std::function<void ()>::~function() /usr/include/c++/13/bits/std_function.h:334
#6 0x55f77e5a3273 in std::function<void ()>::operator=(std::function<void ()> const&) /usr/include/c++/13/bits/std_function.h:471
--- previously allocated ---
previously allocated by thread T0 here:
#0 0x7f2522fa5548 in operator new(unsigned long) ../../../../src/libsanitizer/asan/asan_new_delete.cpp:95
#1 0x55f77e7e67f0 in _M_create<const tint::hlsl::writer::(anonymous namespace)::Printer::EmitLoop(const tint::core::ir::Loop*)::<lambda()>&> /usr/include/c++/13/bits/std_function.h:161
#2 0x55f77e7e66f1 in _M_init_functor<const tint::hlsl::writer::(anonymous namespace)::Printer::EmitLoop(const tint::core::ir::Loop*)::<lambda()>&> /usr/include/c++/13/bits/std_function.h:215
#3 0x55f77e7e6520 in _M_manager /usr/include/c++/13/bits/std_function.h:198
SUMMARY: AddressSanitizer: heap-use-after-free src/tint/lang/hlsl/writer/printer/printer.cc:570 in operator()
To reproduce:
tint ./min.wgsl --format hlsl
Only the HLSL backend is affected — --format spirv|msl|glsl|wgsl all
complete cleanly on the same shader. The crash needs a loop nested
inside another loop’s continuing block; moving the inner loop out of
continuing into the loop body makes it go away. Nothing else in the
original 22-line reproducer was required: no storage buffer, no
if/else, no loop variable.
Analysis:
Printer::EmitLoop points the emit_continuing_ member at a lambda
that lives on its own stack frame:
auto emit_continuing = [&] { ... EmitBlock(l->Continuing()); ... };
TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
With a nested loop the assignment nests too. The outer EmitLoop installs
its lambda and starts emitting; inside the continuing block it reaches
the inner loop, whose own EmitLoop installs and then restores
emit_continuing_ on scope exit. That restore is a
std::function::operator=, which destroys the callable the member was
holding — the ASan report shows the free coming from
std::function::operator= at std_function.h:471 and the allocation from
_M_init_functor for the EmitLoop lambda.
EmitContinue (printer.cc:498) then invokes emit_continuing_, reading
the freed closure at printer.cc:570.
Commit:
c1fcba648220f2080851f5d9ded33b8c53e9df14
Build configuration:
cmake -GNinja -DCMAKE_BUILD_TYPE=Debug -DDAWN_FETCH_DEPENDENCIES=ON \
-DDAWN_ALWAYS_ASSERT=ON -DTINT_ENABLE_IR_VALIDATION_ASSERTS=ON \
-DDAWN_ENABLE_ASAN=ON -DDAWN_ENABLE_UBSAN=ON \
-DTINT_BUILD_SPV_WRITER=ON -DTINT_BUILD_MSL_WRITER=ON \
-DTINT_BUILD_HLSL_WRITER=ON -DTINT_BUILD_GLSL_WRITER=ON
Operating System:
Ubuntu 24.04, Docker fusion-fuzz-tint:latest
This bug was found by fusion-fuzz