Overview

High
Severity
—
CVSS
No
Exploited ITW
Fixed
Fix Status
ImpactUse after free in Tint
DescriptionUse after free in Tint
ComponentTint
Sub-componentDawn>Tint
Bug ClassUAF
Tracker554992296
Fix commitd0acacc83ce8 (dawn) +52/-42
CISA KEVNot listed
Credited0599jiangyc
Disclosed2026-10-06

Changed Functions

FunctionChangeNotes
if
src/tint/lang/glsl/writer/printer/printer.cc
modified
if
src/tint/lang/hlsl/writer/printer/printer.cc
modified
if
src/tint/lang/msl/writer/printer/printer.cc
modified

Files Changed

  • src/tint/lang/hlsl/writer/printer/printer.cc
  • src/tint/lang/msl/writer/printer/printer.cc
  • src/tint/lang/glsl/writer/printer/printer.cc
From d0acacc83ce879a5076ce3c10c08e16b77e2e821 Mon Sep 17 00:00:00 2001
From: dan sinclair <dsinclair@chromium.org>
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 <kainino@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
---
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() << "}";
     }
Loading diff…

Original Bug Report

reported by 05...@gmail.com

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

View on issue tracker