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 <[email protected]> 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/[email protected]?part=13
