Re: [RFC PATCH net-next 13/13] drm/amdkfd: add IPsec full-packet offload

[sashiko-bot@kernel.org]

Thank you for your contribution! Sashiko AI review found 23 potential issue(s) to consider:
- [Critical] The anti-replay sliding window is maintained per RX queue instead of globally per SA, allowing attackers to bypass anti-replay protection.
- [Critical] An undersized packet causes an integer underflow in the GPU shader's ciphertext length calculation, leading to an out-of-bounds VRAM read that crashes the GPU (Remote Denial of Service).
- [High] The GPU shader does not include the high-order 32 bits of the Extended Sequence Number (`seq_hi`) in the AAD computation, causing all ESN packets to fail ICV verification.
- [High] The GPU shader is hardcoded to process exactly 254 ciphertext blocks, breaking Jumbo frame support and causing out-of-bounds VRAM reads on large packets.
- [High] Missing `dma_wmb()` in `knod_ipsec_write_sa_entry` causes a TOCTOU race condition between the CPU and the GPU shader, leading to GPU page faults.
- [High] If `xa_insert` fails during SA addition, the active GPU SA entry is not rolled back, leading to a GPU use-after-free.
- [High] Freeing SA VRAM allocations in `xdo_state_free` after an RCU grace period does not protect asynchronous GPU executions, causing a VRAM use-after-free.
- [High] The GPU verdict only returns the bare slot index without an SA version or SPI check. If a slot is freed and rapidly reallocated, an old decrypted packet will be incorrectly attributed to the new SA.
- [High] Premature `KNOD_IPSEC_PASS` assignment for bypassed/missed packets causes data corruption via early netmem recycling.
- [High] Unlocked access to `slot->x` in the dispatcher and NAPI hook races with `xfrm_state` teardown, leading to UAF.
- [High] Use-After-Free of `ipsec_priv` during device teardown due to missing RCU read lock in XFRM callbacks.
- [High] The in-kernel KAT selftest backs up, zeroes, and restores the live GPU SA table without locks, allowing concurrent traffic or configurations to permanently corrupt the table.
- [High] Unsafe allocation from `page_pool` in preemptible kthread context.
- [High] The SDMA engine writes up to 40 bytes past the end of the `PAGE_SIZE` delivery pool buffer for valid jumbo packets due to a flawed length check.
- [High] The CPU passes unmapped NIC DMA addresses from the page pool to the AMDGPU SDMA engine, which will cause a GPU VM or IOMMU fault.
- [High] Timeouts in `knod_ipsec_fence_wait` and `knod_ipsec_dispatch_wait` fail open, leading to processing of incomplete GPU/SDMA operations.
- [Medium] Use of `this_cpu_inc()` and `this_cpu_ptr()` in preemptible contexts causes smp_processor_id() warnings and stat corruption.
- [Medium] The RX shader checks the IP version but fails to verify that the protocol is ESP (IPPROTO_ESP / 50), allowing attackers to drop valid non-ESP traffic.
- [Medium] `knod_ipsec_xdo_state_update_stats` overwrites `x->curlft` instead of adding the hardware delta, destroying software-accumulated statistics.
- [Medium] Extremely large array allocation on the kernel stack in `knod_ipsec_kat_h_one`.
- [Medium] The self-test inadvertently unparks a dispatcher thread that may have been intentionally stopped by the control plane.
- [Medium] Missing RCU read lock around `xa_load` in `knod_ipsec_xdo_state_update_stats`.
- [Medium] Racy zeroing of remote per-CPU statistics.
--

--- Patch [13]: [RFC PATCH net-next 13/13] drm/amdkfd: add IPsec full-packet offload ---
commit a531874de77887797921e4f4aa4a3418e089408c
Author: Taehee Yoo <ap420073@gmail.com>
Subject: drm/amdkfd: add IPsec full-packet offload

This patch adds a KNOD feature for GPU-offloaded IPsec ESP, allowing
XFRM_DEV_OFFLOAD_PACKET SAs to decrypt AES-GCM payloads directly on the GPU.

> diff --git a/drivers/gpu/drm/amd/amdkfd/knod/ipsec_fused_gfx10.h b/drivers/gpu/drm/amd/amdkfd/knod/ipsec_fused_gfx10.h
> --- a/drivers/gpu/drm/amd/amdkfd/knod/ipsec_fused_gfx10.h
> +++ b/drivers/gpu/drm/amd/amdkfd/knod/ipsec_fused_gfx10.h
> @@ -216,8 +216,6 @@ static inline int kfd_ipsec_gen_fused_shader_gfx10(void *vbuf)
>  	_E(emit_gfx10_s_and_b32_p, I10(buf, n), P_S(28), P_I(0xF), P_S(28));
>  
>  	/* Check IPv4 (version==4) */
>  	_E(emit_gfx10_s_cmp_eq_u32, I10(buf, n), P_S(28), P_I(4));

[Severity: Medium]
Does the shader only check the IP version without verifying that the
protocol is ESP (IPPROTO_ESP)? Could valid non-ESP traffic that
accidentally matches an active SPI be dropped?

>  	br_ipv4 = _BR(emit_gfx10_s_cbranch_scc1, I10(buf, n), 0);
>  

[ ... ]

> @@ -493,8 +491,6 @@ static inline int kfd_ipsec_gen_fused_shader_gfx10(void *vbuf)
>  	_E(emit_gfx10_v_readfirstlane_b32, I10(buf, n), SR_CTEXT_LEN,
>  	   VR_SAVE_PKTLEN);
>  	_E(emit_gfx10_s_sub_u32_p, I10(buf, n), P_S(SR_CTEXT_LEN),
>  	   P_S(SR_CTEXT_LEN), P_S(28));

[Severity: Critical]
If bd->len is smaller than the combined headers and ICV overhead, could this
calculation underflow, producing a massive block count that leads to
out-of-bounds VRAM reads and a GPU crash?

>  	/* s[SR_NBLOCKS_GCM] = (ctext_len + 15) >> 4 */
>  	_E(emit_gfx10_s_add_u32, I10(buf, n), P_S(SR_NBLOCKS_GCM), P_I(15),

[ ... ]

> @@ -626,8 +622,6 @@ static inline int kfd_ipsec_gen_fused_shader_gfx10(void *vbuf)
>  	 * here then free for reuse by the ctext section.
>  	 */
>  	_E(emit_gfx10_v_sub_nc_u32, I10(buf, n), P_V(VR_TMP),
>  	   P_S(SR_TOTAL_GHASH_BLK), P_V(VR_TID));

[Severity: High]
With the workgroup size hardcoded to 256 threads, won't packets exceeding
254 ciphertext blocks (like MTU 9000 jumbo frames) fail decryption and
cause out-of-bounds VRAM reads when accessing the H-power table?

>  	_E(emit_gfx10_v_add_nc_u32, I10(buf, n), P_V(VR_TMP),
>  	   P_L(0xFFFFFFFF), P_V(VR_TMP));

[ ... ]

> @@ -649,8 +643,6 @@ static inline int kfd_ipsec_gen_fused_shader_gfx10(void *vbuf)
>  	 */
>  	_E(emit_gfx10_v_mov_b32_e32, I10(buf, n), P_V(VR_DATA0),
>  	   P_V(VR_SAVE_SPI));
>  	_E(emit_gfx10_v_mov_b32_e32, I10(buf, n), P_V(VR_DATA1),
>  	   P_V(VR_SAVE_SEQ));

[Severity: High]
If the control plane accepts SAs with ESN enabled, but the shader only uses
the 32-bit seq_lo when assembling the AAD, won't all valid ESN packets fail
the ICV verification?

>  	/* DATA2, DATA3 already 0 */
>  

> diff --git a/drivers/gpu/drm/amd/amdkfd/knod/knod_ipsec.h b/drivers/gpu/drm/amd/amdkfd/knod/knod_ipsec.h
> --- a/drivers/gpu/drm/amd/amdkfd/knod/knod_ipsec.h
> +++ b/drivers/gpu/drm/amd/amdkfd/knod/knod_ipsec.h
> @@ -266,4 +266,3 @@ struct knod_ipsec_sa_slot {
>  	 * queue - do not touch from control plane while SA is active.
>  	 */
>  	struct knod_ipsec_sa_window win[KNOD_SPSC_MAX];

[Severity: Critical]
Does maintaining the anti-replay sliding window per RX queue rather than
globally per SA allow an attacker to bypass replay protection by altering
outer headers to change the RSS hash?

>  };
>  

> diff --git a/drivers/gpu/drm/amd/amdkfd/knod/knod_ipsec.c b/drivers/gpu/drm/amd/amdkfd/knod/knod_ipsec.c
> --- a/drivers/gpu/drm/amd/amdkfd/knod/knod_ipsec.c
> +++ b/drivers/gpu/drm/amd/amdkfd/knod/knod_ipsec.c
> @@ -329,4 +329,3 @@ static void knod_ipsec_write_sa_entry(struct knod_ipsec_priv *priv,
>  	e = (struct knod_ipsec_sa_entry *)priv->sa_table->kaddr;
>  	e += slot_idx;
>  	memset(e, 0, sizeof(*e));

[Severity: High]
Could clearing the SA entry here without a dma_wmb() allow the GPU shader
to match a valid spi but read a zeroed key_gpu_addr or htable_gpu_addr,
causing a fatal GPU page fault?

>  
>  	if (!slot || !x) {

[ ... ]

> @@ -424,4 +423,3 @@ static int knod_ipsec_xdo_state_add(struct knod_dev *knodev,
>  
>  	if (x->props.flags & XFRM_STATE_ESN)
>  		flags |= 1u << 0;

[Severity: High]
Are ESN packets expected to be supported by the shader? If the control
plane accepts SAs with ESN enabled, but the shader lacks full seq_hi
support for the AAD, won't all valid ESN packets drop?

>  
>  	/* Store SPI in host byte order (little-endian on this platform).

[ ... ]

> @@ -493,4 +491,3 @@ static int knod_ipsec_xdo_state_add(struct knod_dev *knodev,
>  	err = xa_insert(&priv->spi_to_slot, spi, slot, GFP_KERNEL);
>  	if (err)
>  		goto err_replay;

[Severity: High]
If xa_insert() fails here, the error path frees the VRAM allocations but
leaves the SA entry active in the GPU sa_table. Could a concurrent GPU
dispatch dereference these freed VRAM allocations?

>  
>  	x->xso.offload_handle = (unsigned long)slot;
>  
>  	mutex_unlock(&priv->slot_lock);
>  
>  	this_cpu_inc(priv->stats->sa_add);

[Severity: Medium]
Can using this_cpu_inc() in process contexts trigger smp_processor_id()
warnings and corrupt statistics upon migration?

>  
>  	e_dbg = (struct knod_ipsec_sa_entry *)priv->sa_table->kaddr + slot_idx;

[ ... ]

> @@ -524,4 +521,3 @@ static void knod_ipsec_xdo_state_delete(struct knod_dev *knodev,
>  	spi = be32_to_cpu(x->id.spi);
>  
>  	mutex_lock(&priv->slot_lock);

[Severity: High]
Since XFRM callbacks like knod_ipsec_xdo_state_delete() can be invoked from
process context without rcu_read_lock(), could knod_ipsec_nod_exit() free
priv concurrently, causing a Use-After-Free?

>  	slot = knod_ipsec_lookup_slot_by_spi(priv, spi);
>  	if (!slot) {

[ ... ]

> @@ -561,4 +557,3 @@ static void knod_ipsec_xdo_state_free(struct knod_dev *knodev,
>  		knod_free_mem(priv->knod, slot->htable_mem);
>  	if (slot->key_mem)
>  		knod_free_mem(priv->knod, slot->key_mem);

[Severity: High]
Because the GPU dispatcher doesn't participate in the CPU RCU grace period,
is it possible for an in-flight GPU dispatch to access key_mem or
htable_mem while it is being freed here?

>  
>  	memset(slot, 0, sizeof(*slot));

[ ... ]

> @@ -597,4 +592,3 @@ static void knod_ipsec_xdo_state_update_stats(struct knod_dev *knodev,
>  	spi = be32_to_cpu(x->id.spi);
>  	slot = xa_load(&priv->spi_to_slot, spi);

[Severity: Medium]
Since this can be called from process context without holding the RCU read
lock, is it safe to call xa_load() here?

>  	if (!slot || !slot->active)
>  		return;
>  
>  	/* Read GPU-side per-SA counters (atomically updated by shader). */
>  	gs = (struct knod_ipsec_sa_gpu_stats *)
>  		((u8 *)priv->sa_table->kaddr + KNOD_IPSEC_STATS_REGION_OFF +
>  		 slot->slot_idx * KNOD_IPSEC_SA_STATS_SIZE);
>  
>  	x->curlft.packets = le64_to_cpu(READ_ONCE(gs->rx_packets));
>  	x->curlft.bytes   = le64_to_cpu(READ_ONCE(gs->rx_bytes));

[Severity: Medium]
Does assigning the hardware counter directly to x->curlft.packets overwrite
any software-accumulated statistics (e.g., from bypassed packets)? Should
this add the hardware delta instead?

>  }
>  

[ ... ]

> @@ -885,4 +878,3 @@ static void knod_ipsec_nod_exit(struct knod_dev *knodev)
>  	knodev->post_copy = NULL;
>  	WRITE_ONCE(ipsec_priv, NULL);
>  	synchronize_net();

[Severity: High]
Is it safe to assume all readers are finished here? As asked earlier, if
XFRM callbacks run without rcu_read_lock(), wouldn't this synchronize_net()
fail to wait for them, resulting in a Use-After-Free of priv below?

>  
>  	knod_ipsec_disp_destroy_all(priv);

[ ... ]

> @@ -1287,4 +1279,3 @@ static void knod_ipsec_fill_dispatch(struct knod *knod,
>  {
>  	p->workgroup_size_x = 256;
>  	p->grid_size_x = 256;

[Severity: High]
With the workgroup size hardcoded to 256 threads, won't packets exceeding
254 ciphertext blocks (like MTU 9000 jumbo frames) fail decryption?

>  	p->grid_size_y = max(work->nr_packets, 1);
>  	p->private_segment_size = 0;

[ ... ]

> @@ -1406,4 +1397,3 @@ static void knod_ipsec_fence_wait(struct knod_ipsec_work *w)
>  	while (!knod_ipsec_fence_passed(w)) {
>  		if (--timeout <= 0)
>  			break;

[Severity: High]
If the bounded spin loops in knod_ipsec_fence_wait() and
knod_ipsec_dispatch_wait() time out, does the dispatcher unconditionally
proceed and process incomplete GPU/SDMA operations?

>  		cpu_relax();
>  	}

[ ... ]

> @@ -1513,4 +1503,3 @@ static void knod_ipsec_finish_rx_deliver(struct knod_ipsec_dispatcher *disp,
>  			n_sdma++;
>  
>  			bd->act = KNOD_IPSEC_PASS;

[Severity: High]
For bypassed or missed packets, does assigning KNOD_IPSEC_PASS immediately
after submitting the asynchronous SDMA copy allow the NIC NAPI to recycle
and overwrite the source buffer before the copy completes?

>  			continue;
>  		}

[ ... ]

> @@ -1535,4 +1524,3 @@ static void knod_ipsec_finish_rx_deliver(struct knod_ipsec_dispatcher *disp,
>  		inner_len = bd->len;
>  		if (inner_len == 0 || inner_len > PAGE_SIZE) {

[Severity: High]
Since copies[0].dst is advanced by l3_hdr_len below, could an inner_len
close to PAGE_SIZE cause the SDMA engine to write up to 40 bytes past the
end of the PAGE_SIZE delivery pool buffer?

>  			if (stats_on)
>  				s->rx_drop_malformed++;

[ ... ]

> @@ -1544,4 +1532,3 @@ static void knod_ipsec_finish_rx_deliver(struct knod_ipsec_dispatcher *disp,
>  
>  		netmem = pool ? page_pool_dev_alloc_netmems(pool) : 0;

[Severity: High]
Because the dispatcher kthread runs with preemption enabled, can calling
page_pool_dev_alloc_netmems() corrupt the page_pool unprotected allocation
cache?

>  		if (!netmem) {
>  			if (stats_on)

[ ... ]

> @@ -1553,4 +1540,3 @@ static void knod_ipsec_finish_rx_deliver(struct knod_ipsec_dispatcher *disp,
>  		pkt_gaddr = le64_to_cpu(param->sub[i].pkt_addr);
>  		out_gaddr = le64_to_cpu(param->sub[i].out_addr);
>  		dst_base = page_pool_get_dma_addr_netmem(netmem);

[Severity: High]
Since these host pages are retrieved from the NIC page pool, are they
mapped into the AMDGPU virtual address space or IOMMU domain? Could passing
them to the SDMA engine trigger a GPU VM or IOMMU fault?

>  
>  		if (pkt_mode == XFRM_MODE_TRANSPORT) {

[ ... ]

> @@ -1785,4 +1771,3 @@ static bool knod_ipsec_post_copy(struct knod_dev *knodev, struct sk_buff *skb,
>  
>  	slot = &priv->slots[sa_slot];
>  	if (!slot->active) {

[Severity: High]
If an SA is rapidly deleted and its slot reallocated while its decrypted
packets are still pending in the delivery ring, could this accept those
packets under the new SA context since there is no SPI or version
validation?

>  		if (stats_on)
>  			s->rx_drop_no_sa++;

[ ... ]

> @@ -1800,4 +1785,3 @@ static bool knod_ipsec_post_copy(struct knod_dev *knodev, struct sk_buff *skb,
>  	seq = ((u64)desc->feat.seq_hi << 32) | (u64)desc->feat.seq_lo;
>  	if (!knod_ipsec_sa_window_check(&slot->win[queue_idx], seq)) {

[Severity: Critical]
Does maintaining the anti-replay sliding window per RX queue here rather
than globally per SA allow an attacker to bypass replay protection by
altering outer headers to change the RSS hash?

>  		if (stats_on)
>  			s->rx_drop_replay++;

[ ... ]

> @@ -1861,4 +1845,3 @@ static bool knod_ipsec_post_copy(struct knod_dev *knodev, struct sk_buff *skb,
>  			return false;
>  		}
>  		xfrm_state_hold(slot->x);

[Severity: High]
Is it safe to access and take a reference on slot->x here without locks? If
the SA is deleted concurrently, could this lead to a Use-After-Free when
accessing the freed xfrm_state?

>  		sp->xvec[sp->len++] = slot->x;
>  		sp->olen++;

[ ... ]

> @@ -1983,4 +1966,3 @@ static void knod_ipsec_dispatch_wait(struct knod_ipsec_dispatcher *disp,
>  	while (READ_ONCE(signal->value) > work->sigval) {
>  		if (--timeout <= 0) {

[Severity: High]
Does breaking early on this timeout cause the dispatcher to process
incomplete GPU/SDMA operations?

>  			pr_warn_ratelimited("knod_ipsec: GPU signal timeout nr=%d\n",
>  					    work->nr_packets);

[ ... ]

> @@ -2166,4 +2148,3 @@ static bool knod_ipsec_dispatcher_try_rx(struct knod_ipsec_dispatcher *disp,
>  				param->sub[si].bd_addr  =
>  					cpu_to_le64(bd_gaddr);
>  				param->sub[si].pkt_len  =
>  					cpu_to_le32(bd->len);

[Severity: Critical]
If bd->len is smaller than the combined headers and ICV overhead, could
this calculation underflow later in the shader, producing a massive block
count that leads to out-of-bounds VRAM reads and a GPU crash?

>  				param->sub[si].result_seq = 0;
>  

[ ... ]

> @@ -2178,4 +2159,3 @@ static bool knod_ipsec_dispatcher_try_rx(struct knod_ipsec_dispatcher *disp,
>  
>  	if (static_branch_unlikely(&ipsec_stats_enabled_key)) {
>  		struct knod_ipsec_stats *cs = this_cpu_ptr(priv->stats);

[Severity: Medium]
Can using this_cpu_ptr() here in preemptible kthread contexts trigger
smp_processor_id() warnings and corrupt statistics upon migration?

>  
>  		cs->rx_peek_total += total_n;

[ ... ]

> @@ -2328,4 +2308,3 @@ static bool knod_ipsec_finalise_inflight(struct knod_ipsec_dispatcher *disp,
>  	if (work->n_sdma_pending > 0) {
>  		if (stats_on) {
>  			struct knod_ipsec_stats *s = this_cpu_ptr(priv->stats);

[Severity: Medium]
Is it safe to use this_cpu_ptr() here when preemption isn't disabled?

>  			u64 build_ns = work->t_build_end - work->t_build_start;
>  			u64 gpu_ns = t_gpu_end > work->t_submit

[ ... ]

> @@ -2581,4 +2560,3 @@ static void knod_ipsec_dispatcher_drain(struct knod_ipsec_dispatcher *disp)
>  		if (w->state == KNOD_WORK_INFLIGHT) {
>  			knod_ipsec_dispatch_wait(disp, w);
>  			if (!knod_ipsec_finalise_inflight(disp, w)) {

[Severity: High]
If knod_ipsec_dispatch_wait() times out, does this unconditionally proceed
to finalise the work item with incomplete GPU results?

>  				/*
>  				 * SDMA started - spin-wait, we can't

[ ... ]

> @@ -2782,4 +2760,3 @@ static int knod_ipsec_kat_h_one(const u8 *key, int key_len)
>  {
>  	u8 h_table[KNOD_GCM_H_TABLE_SIZE];

[Severity: Medium]
Given that KNOD_GCM_H_TABLE_SIZE is 4096 bytes, could allocating this array
on the kernel stack cause a stack overflow?

>  	struct aes_enckey enckey;
>  	u8 zero[AES_BLOCK_SIZE] = {};

[ ... ]

> @@ -2868,4 +2845,3 @@ static int knod_ipsec_run_shader_kat_n(struct knod_ipsec_priv *priv, int nr)
>  	 * ALIGNed to 8 pages in the alloc path to dodge the KNOD VRAM
>  	 * 7-page alloc bug; the logical payload is still SA_BO_SIZE.
>  	 */
>  	memset(e, 0, KNOD_IPSEC_SA_BO_SIZE);

[Severity: High]
Does clearing the live GPU SA table here without holding priv->slot_lock
permanently overwrite any SA modifications made by concurrent
administrative operations?

>  	for (i = 0; i < nr; i++) {
>  		e[i].spi    = cpu_to_le32(SPI_KAT_BASE | (u32)i);

[ ... ]

> @@ -3012,4 +2988,3 @@ static int knod_ipsec_run_shader_kat_n(struct knod_ipsec_priv *priv, int nr)
>  	knod_ipsec_dispatch_and_wait(&priv->disp[0], dbg_work);
>  	if (priv->disp[0].kthread)
>  		kthread_unpark(priv->disp[0].kthread);

[Severity: Medium]
If the device was previously stopped via knod_ipsec_nod_stop(), does this
unconditionally unpark the thread and violate the stopped state of the
device?

>  	ret = 0;
>  	pr_info("knod_ipsec: shader-dispatch[%d] submitted, signal %lld -> (done)\n",

[ ... ]

> @@ -3338,4 +3313,3 @@ static int knod_ipsec_run_crypto_kat(struct knod_ipsec_priv *priv)
>  	}
>  	memcpy(sa_backup, priv->sa_table->kaddr, KNOD_IPSEC_SA_BO_SIZE);
>  	memset(priv->sa_table->kaddr, 0, KNOD_IPSEC_SA_BO_SIZE);

[Severity: High]
Can modifying the live SA table here without locks cause permanent
corruption if a user adds or deletes an SA during this selftest?

>  
>  	sa_entry = (struct knod_ipsec_sa_entry *)priv->sa_table->kaddr;

[ ... ]

> @@ -3934,4 +3908,3 @@ static ssize_t knod_ipsec_stats_reset_write(struct file *file,
>  		u64 sa_rekey = s->sa_rekey;
>  
>  		memset(s, 0, sizeof(*s));

[Severity: Medium]
Does using memset() on remote per-CPU statistics without synchronization
cause data races and corruption if those CPUs are concurrently updating
their local stats?

>  		s->sa_add = sa_add;
>  		s->sa_del = sa_del;

-- 
Sashiko AI review · https://sashiko.dev/#/patchset/20260719175857.4071636-1-ap420073@gmail.com?part=13