[RFC PATCH net-next 07/13] drm/amdkfd: add knod provider core

[Taehee Yoo <ap420073@gmail.com>]

Add the knod accelerator provider, built into amdgpu.  It allocates and
drives GPU AQL/SDMA queues, manages GPU memory for the RX data path,
and registers a knod accelerator that binds to a NIC via the knod core.
Feature workers (BPF, IPsec) plug in on top through the accel ops.

Signed-off-by: Taehee Yoo <ap420073@gmail.com>
(cherry picked from commit acaff825a0e87188906fd59f8a377c04d0b0ed2e)
---
 drivers/gpu/drm/amd/amdgpu/Makefile        |    1 +
 drivers/gpu/drm/amd/amdkfd/Kconfig         |   11 +
 drivers/gpu/drm/amd/amdkfd/Makefile        |    4 +
 drivers/gpu/drm/amd/amdkfd/kfd_knod.c      | 2202 ++++++++++++++++++++
 drivers/gpu/drm/amd/amdkfd/knod/kfd_knod.h |  270 +++
 5 files changed, 2488 insertions(+)
 create mode 100644 drivers/gpu/drm/amd/amdkfd/kfd_knod.c
 create mode 100644 drivers/gpu/drm/amd/amdkfd/knod/kfd_knod.h

diff --git a/drivers/gpu/drm/amd/amdgpu/Makefile b/drivers/gpu/drm/amd/amdgpu/Makefile
index ba80542ead9d..7a9826df9f37 100644
--- a/drivers/gpu/drm/amd/amdgpu/Makefile
+++ b/drivers/gpu/drm/amd/amdgpu/Makefile
@@ -42,6 +42,7 @@ ccflags-y := -I$(FULL_AMD_PATH)/include/asic_reg \
 	-I$(FULL_AMD_DISPLAY_PATH)/dc \
 	-I$(FULL_AMD_DISPLAY_PATH)/amdgpu_dm \
 	-I$(FULL_AMD_PATH)/amdkfd \
+	-I$(FULL_AMD_PATH)/amdkfd/knod \
 	-I$(FULL_AMD_PATH)/ras/ras_mgr
 
 # Locally disable W=1 warnings enabled in drm subsystem Makefile
diff --git a/drivers/gpu/drm/amd/amdkfd/Kconfig b/drivers/gpu/drm/amd/amdkfd/Kconfig
index a5d7467c2f34..d93f1af749ff 100644
--- a/drivers/gpu/drm/amd/amdkfd/Kconfig
+++ b/drivers/gpu/drm/amd/amdkfd/Kconfig
@@ -38,3 +38,14 @@ config HSA_AMD_P2P
 	  GPUs with large memory BARs that expose the entire VRAM in PCIe bus
 	  address space within the physical address limits of the GPUs.
 
+config HSA_AMD_KNOD
+	bool "KNOD GPU network offload core"
+	depends on HSA_AMD
+	default y
+	help
+	  Core framework for KNOD, GPU-accelerated zero-copy network packet
+	  processing on AMD GPUs.  It is built into amdgpu and cannot be a
+	  module because it relies on non-exported amdgpu/amdkfd internals.
+
+	  Say N to drop the KNOD core along with the BPF and IPsec offloads
+	  layered on top of it.  If unsure, say Y.
diff --git a/drivers/gpu/drm/amd/amdkfd/Makefile b/drivers/gpu/drm/amd/amdkfd/Makefile
index 85fc67d521e5..1834faa54863 100644
--- a/drivers/gpu/drm/amd/amdkfd/Makefile
+++ b/drivers/gpu/drm/amd/amdkfd/Makefile
@@ -71,3 +71,7 @@ ifneq ($(CONFIG_HSA_AMD_SVM),)
 AMDKFD_FILES += $(AMDKFD_PATH)/kfd_svm.o \
 		$(AMDKFD_PATH)/kfd_migrate.o
 endif
+
+ifneq ($(CONFIG_HSA_AMD_KNOD),)
+AMDKFD_FILES += $(AMDKFD_PATH)/kfd_knod.o
+endif
diff --git a/drivers/gpu/drm/amd/amdkfd/kfd_knod.c b/drivers/gpu/drm/amd/amdkfd/kfd_knod.c
new file mode 100644
index 000000000000..03fff054dcb9
--- /dev/null
+++ b/drivers/gpu/drm/amd/amdkfd/kfd_knod.c
@@ -0,0 +1,2202 @@
+// SPDX-License-Identifier: GPL-2.0 OR MIT
+/*
+ * Copyright (c) 2021 Taehee Yoo <ap420073@gmail.com>
+ * Copyright (c) 2021 Hoyeon Lee <hoyeon.rhee@gmail.com>
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and associated documentation files (the "Software"),
+ * to deal in the Software without restriction, including without limitation
+ * the rights to use, copy, modify, merge, publish, distribute, sublicense,
+ * and/or sell copies of the Software, and to permit persons to whom the
+ * Software is furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in
+ * all copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL
+ * THE COPYRIGHT HOLDER(S) OR AUTHOR(S) BE LIABLE FOR ANY CLAIM, DAMAGES OR
+ * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE,
+ * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
+ * OTHER DEALINGS IN THE SOFTWARE.
+ *
+ */
+
+#include <linux/types.h>
+#include <linux/mutex.h>
+#include <linux/slab.h>
+#include <linux/sched.h>
+#include <linux/workqueue.h>
+#include <linux/file.h>
+#include <linux/dma-buf.h>
+#include <crypto/algapi.h>
+#include <crypto/internal/simd.h>
+#include "kfd_priv.h"
+#include "kfd_hsa.h"
+#include "kfd_knod.h"
+#include "kfd_topology.h"
+#include "kfd_device_queue_manager.h"
+#include "kfd_events.h"
+#include <crypto/skcipher.h>
+#include <crypto/internal/skcipher.h>
+#include <linux/pid.h>
+#include <linux/debugfs.h>
+#include <linux/sched/signal.h>
+#include "../amdgpu/amdgpu_amdkfd.h"
+#include "../amdgpu/amdgpu_gfx.h"
+#include "../amdgpu/./navi10_sdma_pkt_open.h"
+#include "amdgpu_dpm.h"
+#include "kgd_pp_interface.h"
+#include <linux/kthread.h>
+#include <linux/delay.h>
+#include <drm/ttm/ttm_tt.h>
+#include <linux/seq_file.h>
+#include "knod_bpf.h"
+#include <net/page_pool/helpers.h>
+#include <linux/netdevice.h>
+
+struct umh_data {
+	pid_t pid;
+};
+
+LIST_HEAD(ctx_list);
+
+/*
+ * AQL/SDMA queue pair count requested by the highest-demand accel
+ * consumer (currently knod_ipsec's parallel-dispatcher machinery).
+ * Each accel module sets this via knod_request_queue_cnt() during its
+ * module_init BEFORE the NOD attach happens, so knod_attach() creates
+ * a context with enough kaql[]/sdma[] pairs for the worst-case
+ * consumer.
+ *
+ * The value is a high-water mark across all accel types - whichever
+ * consumer asks for the most queue pairs wins; a consumer that needs
+ * fewer just leaves the extra pairs unused (a minor, harmless GPU
+ * resource waste).
+ *
+ * Default is 1 when nothing has called the setter - mirrors the
+ * historical single-queue behaviour.
+ */
+static int knod_requested_queue_cnt = 1;
+
+static struct knod_accel *accels[KNOD_MAX_AQL];
+static int nr_accels;
+static struct knod_accel_ops accel_ops;
+
+/*
+ * Thin wrappers around the KFD core entry points.  knod drives KFD from
+ * kernel context using kernel-internal types, so it calls these directly
+ * instead of going through the kfd_ioctl_* handlers (which marshal user
+ * data and stay static to KFD).
+ */
+static int knod_create_queue(struct kfd_process *p,
+			     struct queue_properties *qp, u32 gpu_id,
+			     u32 *queue_id, u64 *doorbell_offset)
+{
+	return kfd_create_queue(p, qp, gpu_id, queue_id, doorbell_offset);
+}
+
+static int knod_destroy_queue(struct kfd_process *p, u32 queue_id)
+{
+	int ret;
+
+	mutex_lock(&p->mutex);
+	ret = pqm_destroy_queue(&p->pqm, queue_id);
+	mutex_unlock(&p->mutex);
+	return ret;
+}
+
+static int knod_create_event(struct kfd_process *p, u32 event_type,
+			     bool auto_reset, u32 node_id,
+			     struct knod_event *event)
+{
+	u32 event_trigger_data;
+	u64 event_page_offset = 0;
+
+	return kfd_event_create(NULL, p, event_type, auto_reset, node_id,
+				&event->id, &event_trigger_data,
+				&event_page_offset, &event->slot);
+}
+
+static int knod_destroy_event(struct kfd_process *p, u32 event_id)
+{
+	return kfd_event_destroy(p, event_id);
+}
+
+struct knod_mem *__knod_alloc_mem(struct knod *knod, size_t size,
+				  int flags)
+{
+	struct kfd_process_device *pdd = knod->process->pdds[0];
+	struct kfd_node *kdev = pdd->dev;
+	struct knod_mem *mem;
+	int err;
+
+	mem = kzalloc_obj(struct knod_mem, GFP_KERNEL);
+	if (!mem)
+		return ERR_PTR(-ENOMEM);
+
+	size = ALIGN(size, PAGE_SIZE);
+	mem->flags = flags;
+	mem->size = size;
+	mem->gaddr = gen_pool_alloc(knod->pool, size);
+	if (!mem->gaddr) {
+		kfree(mem);
+		return ERR_PTR(-ENOMEM);
+	}
+
+	err = amdgpu_amdkfd_gpuvm_alloc_memory_of_gpu(kdev->adev, mem->gaddr,
+						      size, pdd->drm_priv,
+						      &mem->mem, NULL, flags,
+						      false);
+	if (err) {
+		knod_err(" failed to alloc mem\n");
+		gen_pool_free(knod->pool, mem->gaddr, mem->size);
+		kfree(mem);
+		return ERR_PTR(-ENOMEM);
+	}
+
+	err = amdgpu_amdkfd_gpuvm_sync_memory(kdev->adev, mem->mem, true);
+	if (err) {
+		pr_debug("Sync memory failed, wait interrupted by user signal\n");
+		amdgpu_amdkfd_gpuvm_free_memory_of_gpu(kdev->adev, mem->mem,
+						       pdd->drm_priv, NULL);
+		gen_pool_free(knod->pool, mem->gaddr, mem->size);
+		kfree(mem);
+		return ERR_PTR(-ENOMEM);
+	}
+
+	list_add_tail(&mem->list, &knod->active_list);
+
+	return mem;
+}
+
+int __knod_export_dma_buf(struct knod *knod, struct knod_mem *mem)
+{
+	struct dma_buf *dmabuf;
+	int err;
+
+	err = amdgpu_amdkfd_gpuvm_export_dmabuf(mem->mem, &dmabuf);
+	if (err) {
+		pr_debug("export dmabuf failed\n");
+		return -ENOMEM;
+	}
+
+	dma_buf_put(dmabuf);
+
+	return 0;
+}
+
+struct knod_mem *knod_alloc_mem(struct knod *knod, size_t size, int flags)
+{
+	struct knod_mem *mem;
+	int err;
+
+	mem = kzalloc_obj(struct knod_mem, GFP_KERNEL);
+	if (!mem)
+		return ERR_PTR(-ENOMEM);
+
+	size = ALIGN(size, PAGE_SIZE);
+	mem->flags = flags;
+	mem->size = size;
+	mem->gaddr = gen_pool_alloc(knod->pool, size);
+	if (!mem->gaddr) {
+		kfree(mem);
+		return ERR_PTR(-ENOMEM);
+	}
+
+	err = kfd_process_alloc_gpuvm(knod->process->pdds[0],
+				      mem->gaddr,
+				      mem->size,
+				      flags,
+				      &mem->mem,
+				      &mem->kaddr);
+
+	if (err) {
+		knod_err(" err = %d\n", err);
+		gen_pool_free(knod->pool, mem->gaddr, mem->size);
+		kfree(mem);
+		return ERR_PTR(-ENOMEM);
+	}
+
+	list_add_tail(&mem->list, &knod->active_list);
+
+	return mem;
+}
+EXPORT_SYMBOL(knod_alloc_mem);
+
+int __knod_map_kaddr(struct knod *knod, struct knod_mem *mem)
+{
+	int err = 0;
+
+	if (mem->flags & KFD_IOC_ALLOC_MEM_FLAGS_GTT) {
+		err = amdgpu_amdkfd_gpuvm_map_gtt_bo_to_kernel(mem->mem,
+							       &mem->kaddr,
+							       NULL);
+		if (err) {
+			pr_debug("Map GTT BO to kernel failed\n");
+			err = -ENOMEM;
+		}
+	} else if (mem->flags & KFD_IOC_ALLOC_MEM_FLAGS_VRAM) {
+		err = amdgpu_amdkfd_gpuvm_map_vram_bo_to_kernel(mem->mem,
+								&mem->kaddr,
+								NULL);
+		if (err) {
+			pr_debug("Map VRAM BO to kernel failed\n");
+			err = -ENOMEM;
+		}
+	}
+
+	return err;
+}
+
+int __knod_map_mem(struct knod *knod, struct knod_mem *mem)
+{
+	struct kfd_process_device *pdd = knod->process->pdds[0];
+	struct kfd_node *kdev = pdd->dev;
+	int err;
+
+	err = amdgpu_amdkfd_gpuvm_map_memory_to_gpu(kdev->adev, mem->mem,
+						    pdd->drm_priv);
+	if (err) {
+		knod_err(" failed to map gpu\n");
+		return 1;
+	}
+
+	err = amdgpu_amdkfd_gpuvm_sync_memory(kdev->adev, mem->mem, true);
+	if (err) {
+		pr_debug("Sync memory failed, wait interrupted by user signal\n");
+		return 1;
+	}
+
+	return 0;
+}
+EXPORT_SYMBOL(__knod_map_mem);
+
+void knod_free_mem(struct knod *knod, struct knod_mem *mem)
+{
+	if (mem) {
+		list_del_init(&mem->list);
+		kfd_process_free_gpuvm(mem->mem, knod->process->pdds[0],
+				       &mem->kaddr);
+		gen_pool_free(knod->pool, mem->gaddr, mem->size);
+		kfree(mem);
+	}
+}
+EXPORT_SYMBOL(knod_free_mem);
+
+void knod_sdma_copy(struct knod *knod, u64 dst_offset,
+		    u64 src_offset, int idx, int size)
+{
+	struct knod_sdma *sdma = &knod->sdma[idx];
+	u32 ring_mask = (sdma->sdma->size / 4) - 1;
+	u64 *wptr = (u64 *)sdma->queue->kaddr + 1;
+	u32 *ptr = sdma->sdma->kaddr;
+
+	ptr[sdma->idx++ & ring_mask] = SDMA_PKT_HEADER_OP(SDMA_OP_COPY) |
+		SDMA_PKT_HEADER_SUB_OP(SDMA_SUBOP_COPY_LINEAR) |
+		SDMA_PKT_COPY_LINEAR_HEADER_TMZ((0));
+	ptr[sdma->idx++ & ring_mask] = size - 1;
+	ptr[sdma->idx++ & ring_mask] = 0; /* src/dst endian swap */
+	ptr[sdma->idx++ & ring_mask] = lower_32_bits(src_offset);
+	ptr[sdma->idx++ & ring_mask] = upper_32_bits(src_offset);
+	ptr[sdma->idx++ & ring_mask] = lower_32_bits(dst_offset);
+	ptr[sdma->idx++ & ring_mask] = upper_32_bits(dst_offset);
+
+	*wptr += 7 * 4;
+}
+
+void knod_sdma_fence(struct knod *knod, u64 fence_addr, u32 fence_val,
+		     int idx)
+{
+	struct knod_sdma *sdma = &knod->sdma[idx];
+	u32 ring_mask = (sdma->sdma->size / 4) - 1;
+	u64 *wptr = (u64 *)sdma->queue->kaddr + 1;
+	u32 *ptr = sdma->sdma->kaddr;
+	u32 hdr = SDMA_PKT_HEADER_OP(SDMA_OP_FENCE);
+
+	if (knod->isa_version >= 10)
+		hdr |= SDMA_PKT_FENCE_HEADER_MTYPE(3);
+
+	ptr[sdma->idx++ & ring_mask] = hdr;
+	ptr[sdma->idx++ & ring_mask] = lower_32_bits(fence_addr);
+	ptr[sdma->idx++ & ring_mask] = upper_32_bits(fence_addr);
+	ptr[sdma->idx++ & ring_mask] = fence_val;
+
+	*wptr += 4 * 4;
+}
+
+void knod_sdma_trap(struct knod *knod, int idx)
+{
+	struct knod_sdma *sdma = &knod->sdma[idx];
+	u32 ring_mask = (sdma->sdma->size / 4) - 1;
+	u64 *wptr = (u64 *)sdma->queue->kaddr + 1;
+	u32 *ptr = sdma->sdma->kaddr;
+	u32 ctx = knod->sdma_event[idx].id & 0x0fffffff;
+	u64 slot_addr = knod->mailbox->gaddr +
+		knod->sdma_event[idx].slot * 8;
+	u32 hdr;
+
+	/* Write 0 to the signal page slot so lookup_signaled_event
+	 * sees it as non-UNSIGNALED (0xFFFFFFFFFFFFFFFF).
+	 */
+	hdr = SDMA_PKT_HEADER_OP(SDMA_OP_FENCE);
+	if (knod->isa_version >= 10)
+		hdr |= SDMA_PKT_FENCE_HEADER_MTYPE(3);
+
+	ptr[sdma->idx++ & ring_mask] = hdr;
+	ptr[sdma->idx++ & ring_mask] = lower_32_bits(slot_addr);
+	ptr[sdma->idx++ & ring_mask] = upper_32_bits(slot_addr);
+	ptr[sdma->idx++ & ring_mask] = 0;
+
+	ptr[sdma->idx++ & ring_mask] = SDMA_PKT_HEADER_OP(SDMA_OP_TRAP);
+	ptr[sdma->idx++ & ring_mask] = ctx;
+
+	*wptr += 6 * 4;
+}
+
+void knod_sdma_doorbell(struct knod *knod, int idx)
+{
+	struct knod_sdma *sdma = &knod->sdma[idx];
+	u64 *wptr = (u64 *)sdma->queue->kaddr + 1;
+
+	/* ensure the SDMA ring writes are visible before ringing doorbell */
+	wmb();
+	writeq(*wptr, sdma->doorbell);
+}
+
+u32 knod_sdma_submit(struct knod *knod, int idx,
+		     const struct knod_sdma_copy_desc *copies, int n)
+{
+	struct knod_sdma *sdma = &knod->sdma[idx];
+	u32 capacity = sdma->sdma->size / 4;
+	u32 completed, inflight;
+	int i;
+
+	/*
+	 * knod_sdma_copy has no overflow guard (sdma->idx is the dword write
+	 * cursor that just wraps), so drop the whole batch once the ring is
+	 * near full.  @completed is the last fenced cursor (knod_sdma_kick
+	 * writes sdma->idx into the signal), so @inflight is the unprocessed
+	 * span; reserve room for these @n copies (7 dwords each) plus a fence.
+	 * Signed compare so a stale signal ahead of the cursor (e.g. at
+	 * startup) reads as "negative" inflight rather than a false full.
+	 */
+	completed = (u32)READ_ONCE(((struct amd_signal *)
+				    sdma->queue_signal->kaddr)->value);
+	inflight = (u32)sdma->idx - completed;
+	if ((s32)(inflight + n * 7) >= (s32)(capacity - 64))
+		return 0;
+
+	for (i = 0; i < n; i++)
+		knod_sdma_copy(knod, copies[i].dst, copies[i].src, idx,
+			       copies[i].len);
+
+	return (u32)sdma->idx;
+}
+EXPORT_SYMBOL(knod_sdma_submit);
+
+void knod_sdma_kick(struct knod *knod, int idx)
+{
+	struct knod_sdma *sdma = &knod->sdma[idx];
+	u64 fence_addr;
+
+	fence_addr = sdma->queue_signal->gaddr +
+		     offsetof(struct amd_signal, value);
+	knod_sdma_fence(knod, fence_addr, (u32)sdma->idx, idx);
+	knod_sdma_doorbell(knod, idx);
+}
+EXPORT_SYMBOL(knod_sdma_kick);
+
+int knod_gart_map(struct amdgpu_device *adev, u64 npages,
+		  dma_addr_t *addr, u64 *gart_addr, u64 flags)
+{
+	struct amdgpu_ring *ring =
+		to_amdgpu_ring(adev->mman.buffer_funcs_scheds[0]);
+	struct amdgpu_job *job;
+	unsigned int num_dw, num_bytes;
+	struct dma_fence *fence;
+	u64 src_addr, dst_addr;
+	u64 pte_flags;
+	void *cpu_addr;
+	int r;
+
+	/* use gart window 0 */
+	*gart_addr = adev->gmc.gart_start;
+
+	num_dw = ALIGN(adev->mman.buffer_funcs->copy_num_dw, 8);
+	num_bytes = npages * 8;
+
+	r = amdgpu_job_alloc_with_ib(adev, &adev->mman.default_entity.base,
+				     AMDGPU_FENCE_OWNER_UNDEFINED,
+				     num_dw * 4 + num_bytes,
+				     AMDGPU_IB_POOL_DELAYED,
+				     &job,
+				     AMDGPU_KERNEL_JOB_ID_TTM_MAP_BUFFER);
+	if (r)
+		return r;
+
+	src_addr = num_dw * 4;
+	src_addr += job->ibs[0].gpu_addr;
+
+	dst_addr = amdgpu_bo_gpu_offset(adev->gart.bo);
+	amdgpu_emit_copy_buffer(adev, &job->ibs[0], src_addr,
+				dst_addr, num_bytes, 0);
+
+	amdgpu_ring_pad_ib(ring, &job->ibs[0]);
+	WARN_ON(job->ibs[0].length_dw > num_dw);
+
+	pte_flags = AMDGPU_PTE_VALID | AMDGPU_PTE_READABLE;
+	pte_flags |= AMDGPU_PTE_SYSTEM | AMDGPU_PTE_SNOOPED;
+	if (!(flags & KFD_IOCTL_SVM_FLAG_GPU_RO))
+		pte_flags |= AMDGPU_PTE_WRITEABLE;
+	pte_flags |= adev->gart.gart_pte_flags;
+
+	cpu_addr = &job->ibs[0].ptr[num_dw];
+
+	amdgpu_gart_map(adev, 0, npages, addr, pte_flags, cpu_addr);
+	fence = amdgpu_job_submit(job);
+	dma_fence_put(fence);
+
+	return r;
+}
+
+static void knod_init_aql_queue(struct knod *knod, int qid, int idx)
+{
+	struct amd_queue *amd_queue =
+		(struct amd_queue *)knod->kaql[idx].queue->kaddr;
+	union knod_aql_rsrc1 rsrc1;
+	u32 scratch_gaddr;
+	int i;
+
+	for (i = 0; i < knod->nr_aql_ring; i++)
+		knod_setup_invalidate(knod, i, idx);
+
+	amd_queue->hsa_queue.type = HSA_QUEUE_TYPE_MULTI;
+	amd_queue->hsa_queue.features = HSA_QUEUE_FEATURE_KERNEL_DISPATCH;
+	amd_queue->hsa_queue.base_address = (void *)knod->kaql[idx].aql->gaddr;
+	amd_queue->hsa_queue.doorbell_signal =
+		knod->kaql[idx].queue_signal->gaddr;
+	amd_queue->hsa_queue.size = knod->nr_aql_ring;
+	amd_queue->hsa_queue.reserved1 = 0;
+	amd_queue->hsa_queue.id = qid;
+
+	amd_queue->write_dispatch_id = 0; /* id is index */
+	amd_queue->group_segment_aperture_base_hi = 0;
+	amd_queue->private_segment_aperture_base_hi = 0;
+	amd_queue->max_cu_id = -1;
+	amd_queue->max_wave_id = -1;
+	amd_queue->max_legacy_doorbell_dispatch_id_plus_1 = 0;
+	amd_queue->legacy_doorbell_lock = 0;
+	amd_queue->read_dispatch_id = 0; /* id is index */
+	amd_queue->read_dispatch_id_field_base_byte_offset = 0x80;
+	/* scratch resource descriptor - use allocated scratch BO address */
+	scratch_gaddr = (u32)(knod->kaql[idx].scratch->gaddr & 0xffffffff);
+	amd_queue->scratch_resource_descriptor[0] = scratch_gaddr;
+	scratch_gaddr = (u32)((knod->kaql[idx].scratch->gaddr >> 32) & 0xffff);
+	rsrc1.base_address_hi = scratch_gaddr;
+	rsrc1.stride = 0;
+	rsrc1.cache_swizzle = 0;
+	rsrc1.swizzle_enable = 1;
+	memcpy(&amd_queue->scratch_resource_descriptor[1], &rsrc1, sizeof(u32));
+	amd_queue->scratch_resource_descriptor[2] =
+		knod->kaql[idx].scratch->size;
+	amd_queue->scratch_resource_descriptor[3] = 0x00ffffff;
+	amd_queue->scratch_backing_memory_location =
+		knod->kaql[idx].scratch->gaddr;
+	amd_queue->scratch_backing_memory_byte_size =
+		knod->kaql[idx].scratch->size;
+	amd_queue->scratch_wave64_lane_byte_size = 64;
+
+	amd_queue->queue_properties.is_ptr64 = 1;
+	amd_queue->queue_properties.enable_trap_handler_debug_sgprs = 0;
+}
+
+static void knod_init_sdma_queue(struct knod *knod, int qid, int idx)
+{
+}
+
+static void knod_init_queue(struct knod *knod, int qid, int idx, int type)
+{
+	if (type == KFD_IOC_QUEUE_TYPE_COMPUTE_AQL)
+		knod_init_aql_queue(knod, qid, idx);
+	else if (type == KFD_IOC_QUEUE_TYPE_SDMA)
+		knod_init_sdma_queue(knod, qid, idx);
+}
+
+static void stop_umh(pid_t umh_pid)
+{
+	struct pid *p = find_get_pid(umh_pid);
+
+	if (!p)
+		return;
+
+	kill_pid(p, SIGKILL, 1);
+	put_pid(p);
+}
+
+static int umh_init(struct subprocess_info *info, struct cred *new)
+{
+	struct umh_data *d = info->data;
+
+	d->pid = current->pid;
+
+	return 0;
+}
+
+static int launch_and_get_pid(void)
+{
+	static const char * const argv[] = { "/bin/sleep", "2147483647", NULL };
+	static const char * const envp[] = { "HOME=/", "PATH=/sbin:/bin", NULL };
+	struct umh_data data = { .pid = -1 };
+	struct subprocess_info *info;
+	int ret;
+
+	info = call_usermodehelper_setup(argv[0], (char **)argv,
+					 (char **)envp, GFP_KERNEL,
+					 umh_init, NULL, &data);
+	if (!info) {
+		pr_err("UMH setup failed\n");
+		return 0;
+	}
+
+	ret = call_usermodehelper_exec(info, UMH_WAIT_EXEC);
+	if (ret < 0) {
+		pr_err("UMH exec failed: %d\n", ret);
+		return 0;
+	}
+
+	return data.pid;
+}
+
+static int knod_set_isa(struct knod *knod)
+{
+	enum amd_asic_type asic_type;
+
+	asic_type = knod->process->pdds[0]->dev->adev->asic_type;
+	if (knod->isa_version != 9 && knod->isa_version != 10)
+		knod->isa_version = 0;
+	if (knod->isa_version == 0) {
+		if (asic_type <= CHIP_VEGAM) {
+			pr_err("Not supported chip");
+			return -EOPNOTSUPP;
+		} else if (asic_type == CHIP_VEGA10 ||
+				asic_type == CHIP_VEGA12 ||
+				asic_type == CHIP_VEGA20 ||
+				asic_type == CHIP_RAVEN ||
+				asic_type == CHIP_RENOIR) {
+			pr_debug("GCN5 is detected");
+			knod->isa_version = 9;
+		} else if (asic_type == CHIP_NAVI10 ||
+				asic_type == CHIP_NAVI12 ||
+				asic_type == CHIP_NAVI14 ||
+				asic_type == CHIP_CYAN_SKILLFISH) {
+			pr_err("RDNA1 is not supported yet");
+			return -EOPNOTSUPP;
+		} else if (asic_type == CHIP_SIENNA_CICHLID ||
+				asic_type == CHIP_NAVY_FLOUNDER ||
+				asic_type == CHIP_DIMGREY_CAVEFISH ||
+				asic_type == CHIP_BEIGE_GOBY ||
+				asic_type == CHIP_YELLOW_CARP) {
+			pr_debug("RDNA2 is detected");
+			knod->isa_version = 10;
+		} else if (asic_type == CHIP_ARCTURUS ||
+				asic_type == CHIP_ALDEBARAN) {
+			pr_err("CDNA is not supported yet");
+			return -EOPNOTSUPP;
+		} else if (asic_type == CHIP_IP_DISCOVERY) {
+			pr_err("Can't detect chip version, firmware update may be needed");
+			return -EOPNOTSUPP;
+		} else {
+			pr_err("Not supported chip");
+			return -EOPNOTSUPP;
+		}
+	}
+
+	knod->igpu = false;
+	if (asic_type == CHIP_RENOIR || asic_type == CHIP_RAVEN ||
+	    asic_type == CHIP_CYAN_SKILLFISH || asic_type == CHIP_BEIGE_GOBY ||
+	    asic_type == CHIP_YELLOW_CARP) {
+		pr_debug("iGPU is detected");
+		knod->igpu = true;
+	}
+
+	return 0;
+}
+
+static void knod_destroy_one_queue(struct knod *knod, int idx)
+{
+	/* Destroy queue and event BEFORE freeing BOs.
+	 * Queue holds references to BO VAs; freeing BOs first causes
+	 * NULL deref in kfd_queue_unref_bo_vas when UMH exits.
+	 * Use the created flag (not queue_id value) since KFD IDR can
+	 * assign queue_id=0 to the first queue.
+	 */
+	if (knod->aql_queue_created[idx]) {
+		int ret;
+
+		ret = knod_destroy_queue(knod->process,
+					 knod->aql_queue_id[idx]);
+		if (ret)
+			pr_err("knod: destroy_queue failed: %d (queue_id=%u)\n",
+			       ret, knod->aql_queue_id[idx]);
+		knod->aql_queue_created[idx] = false;
+		knod->aql_queue_id[idx] = 0;
+	}
+	if (knod->aql_event[idx].id) {
+		knod_destroy_event(knod->process, knod->aql_event[idx].id);
+		knod->aql_event[idx].id = 0;
+	}
+
+	knod_free_mem(knod, knod->kaql[idx].aql);
+	knod_free_mem(knod, knod->kaql[idx].queue);
+	knod_free_mem(knod, knod->kaql[idx].eop);
+	knod_free_mem(knod, knod->kaql[idx].ctx);
+	knod_free_mem(knod, knod->kaql[idx].queue_signal);
+	knod_free_mem(knod, knod->kaql[idx].amd_queue);
+	knod_free_mem(knod, knod->kaql[idx].scratch);
+}
+
+static int knod_alloc_one_queue(struct knod *knod, int idx,
+				struct kfd_topology_device *topo_dev,
+				struct kfd_process_device *pdd, void *ptr)
+{
+	int buf_flags = KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE |
+			KFD_IOC_ALLOC_MEM_FLAGS_COHERENT |
+			KFD_IOC_ALLOC_MEM_FLAGS_VRAM;
+	int flags = KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE |
+		    KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE |
+		    KFD_IOC_ALLOC_MEM_FLAGS_COHERENT |
+		    KFD_IOC_ALLOC_MEM_FLAGS_GTT;
+	struct amd_signal *queue_signal;
+	struct amd_queue *amd_queue;
+	struct queue_properties qp;
+	u32 total_cwsr_size;
+	int err;
+
+	knod->kaql[idx].aql = knod_alloc_mem(knod,
+		(NR_AQL_RING *
+		 sizeof(struct hsa_kernel_dispatch_packet) * 2),
+		KFD_IOC_ALLOC_MEM_FLAGS_GTT |
+		KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE |
+		KFD_IOC_ALLOC_MEM_FLAGS_COHERENT |
+		KFD_IOC_ALLOC_MEM_FLAGS_UNCACHED |
+		KFD_IOC_ALLOC_MEM_FLAGS_AQL_QUEUE_MEM |
+		KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE);
+	if (IS_ERR(knod->kaql[idx].aql)) {
+		err = PTR_ERR(knod->kaql[idx].aql);
+		knod->kaql[idx].aql = NULL;
+		goto err_mem;
+	}
+
+	knod->kaql[idx].queue = knod_alloc_mem(knod, PAGE_SIZE, flags);
+	if (IS_ERR(knod->kaql[idx].queue)) {
+		err = PTR_ERR(knod->kaql[idx].queue);
+		knod->kaql[idx].queue = NULL;
+		goto err_mem;
+	}
+
+	knod->kaql[idx].eop = knod_alloc_mem(knod, PAGE_SIZE, flags);
+	if (IS_ERR(knod->kaql[idx].eop)) {
+		err = PTR_ERR(knod->kaql[idx].eop);
+		knod->kaql[idx].eop = NULL;
+		goto err_mem;
+	}
+
+	total_cwsr_size = (topo_dev->node_props.cwsr_size +
+			   topo_dev->node_props.debug_memory_size)
+			  * NUM_XCC(pdd->dev->xcc_mask);
+	total_cwsr_size = ALIGN(total_cwsr_size, PAGE_SIZE);
+
+	knod->kaql[idx].ctx = knod_alloc_mem(knod, total_cwsr_size, buf_flags);
+	if (IS_ERR(knod->kaql[idx].ctx)) {
+		err = PTR_ERR(knod->kaql[idx].ctx);
+		knod->kaql[idx].ctx = NULL;
+		goto err_mem;
+	}
+
+	knod->kaql[idx].queue_signal = knod_alloc_mem(knod,
+		PAGE_SIZE << 5,
+		KFD_IOC_ALLOC_MEM_FLAGS_GTT |
+		KFD_IOC_ALLOC_MEM_FLAGS_COHERENT |
+		KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE |
+		KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE);
+	if (IS_ERR(knod->kaql[idx].queue_signal)) {
+		err = PTR_ERR(knod->kaql[idx].queue_signal);
+		knod->kaql[idx].queue_signal = NULL;
+		goto err_mem;
+	}
+
+	knod->kaql[idx].amd_queue = knod_alloc_mem(knod,
+		PAGE_SIZE << 5,
+		KFD_IOC_ALLOC_MEM_FLAGS_GTT |
+		KFD_IOC_ALLOC_MEM_FLAGS_COHERENT);
+	if (IS_ERR(knod->kaql[idx].amd_queue)) {
+		err = PTR_ERR(knod->kaql[idx].amd_queue);
+		knod->kaql[idx].amd_queue = NULL;
+		goto err_mem;
+	}
+
+	knod->kaql[idx].scratch = knod_alloc_mem(knod,
+		PAGE_SIZE << 5,
+		KFD_IOC_ALLOC_MEM_FLAGS_VRAM |
+		KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE |
+		KFD_IOC_ALLOC_MEM_FLAGS_COHERENT);
+	if (IS_ERR(knod->kaql[idx].scratch)) {
+		err = PTR_ERR(knod->kaql[idx].scratch);
+		knod->kaql[idx].scratch = NULL;
+		goto err_mem;
+	}
+
+	knod->kaql[idx].idx = 0;
+
+	/*
+	 * KNOD dispatchers poll queue_signal->value directly. Do not attach a
+	 * KFD signal event to every AQL completion; at high packet rates the
+	 * unused event notifications overflow the KFD interrupt/event path and
+	 * add latency without contributing to completion detection.
+	 */
+	knod->aql_event[idx].id = 0;
+	knod->aql_event[idx].slot = 0;
+	knod->signal_eid = 0;
+
+	/* Zero ALL queue-related GPU memory BEFORE creating the HW queue.
+	 * kfd_ioctl_create_queue loads HQD immediately - if the memory
+	 * contains stale data from a previous session (write/read pointers,
+	 * AQL dispatch packets, doorbell values), the GPU starts processing
+	 * garbage packets and goes to 100%.
+	 */
+	memset(knod->kaql[idx].queue->kaddr, 0, knod->kaql[idx].queue->size);
+	/* AQL ring is allocated 2x with AQL_QUEUE_MEM flag; the second
+	 * half is a GPU-side mirror for wrap-around.  kaddr only covers
+	 * the first half.
+	 */
+	memset(knod->kaql[idx].aql->kaddr, 0, knod->kaql[idx].aql->size / 2);
+	memset(knod->kaql[idx].eop->kaddr, 0, knod->kaql[idx].eop->size);
+	memset(knod->kaql[idx].amd_queue->kaddr, 0,
+	       knod->kaql[idx].amd_queue->size);
+	knod_init_queue(knod, idx, idx, KFD_IOC_QUEUE_TYPE_COMPUTE_AQL);
+
+	memset(&qp, 0, sizeof(qp));
+	qp.type = KFD_QUEUE_TYPE_COMPUTE;
+	qp.format = KFD_QUEUE_FORMAT_AQL;
+	qp.queue_percent = 100;
+	qp.priority = 15;
+	qp.ctl_stack_size = topo_dev->node_props.ctl_stack_size;
+	qp.ctx_save_restore_area_size = topo_dev->node_props.cwsr_size;
+	qp.ctx_save_restore_area_address = (u64)knod->kaql[idx].ctx->gaddr;
+	qp.queue_address = (u64)knod->kaql[idx].aql->gaddr;
+	qp.queue_size = knod->kaql[idx].aql->size / 2;
+	qp.write_ptr =
+		(void __user *)((u64)knod->kaql[idx].queue->gaddr + 0x38);
+	qp.read_ptr = (void __user *)((u64)knod->kaql[idx].queue->gaddr + 0x80);
+	qp.eop_ring_buffer_address = (u64)knod->kaql[idx].eop->gaddr;
+	qp.eop_ring_buffer_size = knod->kaql[idx].eop->size;
+
+	err = knod_create_event(knod->process, KFD_IOC_EVENT_SIGNAL, true, 1,
+				&knod->aql_event[idx]);
+	if (err) {
+		knod_err(" failed to create AQL event[%d] err=%d\n", idx, err);
+		goto err_mem;
+	}
+
+	queue_signal = (struct amd_signal *)knod->kaql[idx].queue_signal->kaddr;
+	queue_signal->kind = AMD_SIGNAL_KIND_USER;
+	if (queue_signal->kind == AMD_SIGNAL_KIND_DOORBELL) {
+		queue_signal->value = 0;
+		queue_signal->event_id = 0;
+		queue_signal->queue_ptr = (void *)knod->kaql[idx].queue->gaddr;
+		queue_signal->hardware_doorbell_ptr = ptr;
+		queue_signal->event_mailbox_ptr = 0;
+	} else if (queue_signal->kind == AMD_SIGNAL_KIND_USER) {
+		queue_signal->event_id = knod->aql_event[idx].id;
+		queue_signal->queue_ptr = 0;
+		queue_signal->event_mailbox_ptr = knod->mailbox->gaddr +
+			knod->aql_event[idx].slot * 8;
+		/* value is in union with hardware_doorbell_ptr - set last */
+		queue_signal->value = 0xffffffffffffffff;
+	}
+	knod_dbg(" pre-create-queue signal=%lld kaddr=%p gaddr=0x%llx\n",
+		 READ_ONCE(queue_signal->value), queue_signal,
+		 knod->kaql[idx].queue_signal->gaddr);
+
+	err = knod_create_queue(knod->process, &qp, pdd->user_gpu_id,
+				&knod->aql_queue_id[idx],
+				&knod->aql_doorbell_offset[idx]);
+	if (err) {
+		knod_err(" failed to create queue, err = %d\n", err);
+		goto err_mem;
+	}
+	knod->aql_queue_created[idx] = true;
+
+	knod_dbg(" post-create-queue signal=%lld queue_id=%d\n",
+		 READ_ONCE(queue_signal->value), knod->aql_queue_id[idx]);
+
+	/* Update queue id with the actual KFD-assigned value */
+	amd_queue = (struct amd_queue *)knod->kaql[idx].queue->kaddr;
+	amd_queue->hsa_queue.id = knod->aql_queue_id[idx];
+	return 0;
+
+err_mem:
+	if (knod->aql_event[idx].id) {
+		knod_destroy_event(knod->process, knod->aql_event[idx].id);
+		knod->aql_event[idx].id = 0;
+	}
+	knod_free_mem(knod, knod->kaql[idx].aql);
+	knod_free_mem(knod, knod->kaql[idx].queue);
+	knod_free_mem(knod, knod->kaql[idx].eop);
+	knod_free_mem(knod, knod->kaql[idx].ctx);
+	knod_free_mem(knod, knod->kaql[idx].queue_signal);
+	knod_free_mem(knod, knod->kaql[idx].amd_queue);
+	knod_free_mem(knod, knod->kaql[idx].scratch);
+	memset(&knod->kaql[idx], 0, sizeof(knod->kaql[idx]));
+	return err;
+}
+
+/*
+ * Write a minimal no-op kernel into the kernel BO so that any dispatch
+ * before a real shader (BPF/IPsec/MACsec/WG) is loaded executes a
+ * harmless s_endpgm instead of faulting on uninitialised VRAM.
+ *
+ * Layout:
+ *   [0..63]     kernel_descriptor  (kernel_code_entry_byte_offset = 256)
+ *   [256..259]  s_endpgm           (0xBF810000)
+ *   [260..1023] s_code_end padding (GFX10 SQC prefetch safety)
+ */
+#define KNOD_DEFAULT_KD_ENTRY_OFFSET	256
+#define KNOD_S_ENDPGM			0xBF810000u
+#define KNOD_S_CODE_END			0xBF9F0000u
+
+static void knod_init_default_kernel(struct knod *knod)
+{
+	struct kernel_descriptor *kd = knod->kernels[0]->kaddr;
+	u32 *code = (u32 *)((u8 *)knod->kernels[0]->kaddr +
+			    KNOD_DEFAULT_KD_ENTRY_OFFSET);
+	int i;
+
+	memset(kd, 0, sizeof(*kd));
+	kd->kernel_code_entry_byte_offset = KNOD_DEFAULT_KD_ENTRY_OFFSET;
+	kd->compute_pgm_rsrc1.granulated_workitem_vgpr_count = 0;
+	kd->compute_pgm_rsrc1.granulated_wavefront_sgpr_count = 0;
+	kd->compute_pgm_rsrc1.float_denorm_mode_32 = 3;
+	kd->compute_pgm_rsrc1.float_denorm_mode_16_64 = 3;
+	kd->compute_pgm_rsrc1.enable_dx10_clamp = 1;
+	kd->compute_pgm_rsrc1.enable_ieee_mode = 1;
+	if (knod->isa_version >= 10)
+		kd->compute_pgm_rsrc1.mem_ordered = 1;
+	kd->compute_pgm_rsrc2.enable_sgpr_workgroup_id_x = 1;
+
+	code[0] = KNOD_S_ENDPGM;
+	for (i = 1; i < (1024 - KNOD_DEFAULT_KD_ENTRY_OFFSET) / 4; i++)
+		code[i] = KNOD_S_CODE_END;
+}
+
+#define KNOD_DEFAULT_BATCH	64
+
+static struct knod_accel_xdp_ops *registered_xdp_ops;
+
+/*
+ * feature=none has no per-feature ops: the default worker stamps XDP_PASS and
+ * the NIC act handler delivers via knod_d2h_copy / knod_d2h_drain.
+ */
+static struct knod_accel_xdp_ops default_xdp_ops = {
+};
+
+static int knod_default_worker(void *arg)
+{
+	struct knod *knod = arg;
+	struct spsc_bd *bds[KNOD_DEFAULT_BATCH];
+	struct knod_dev *knodev;
+	unsigned int cnt;
+	int qi, i;
+
+	while (!kthread_should_stop()) {
+		knodev = READ_ONCE(knod->accel->knodev);
+		if (!knodev || !knodev->started) {
+			usleep_range(1000, 2000);
+			continue;
+		}
+
+		for (qi = 0; qi < knod->channels; qi++) {
+			struct knod_work_priv *wpriv = &knodev->wpriv[qi];
+
+			if (!wpriv->napi)
+				continue;
+
+			if (spsc_peek(&wpriv->spsc_bds, (void **)bds,
+				      KNOD_DEFAULT_BATCH, &cnt))
+				continue;
+
+			/*
+			 * feature=none has no program, so every packet passes.
+			 * Stamp the verdict before advancing the consumer
+			 * cursor: spsc_acquire() publishes the window with a
+			 * release barrier the act handler pairs with, so the
+			 * verdict has to be written first or the act handler
+			 * races a POISON read.  The NIC act handler does the
+			 * device->host delivery via knod_d2h_copy.
+			 */
+			for (i = 0; i < cnt; i++)
+				WRITE_ONCE(bds[i]->act, XDP_PASS);
+
+			spsc_acquire(&wpriv->spsc_bds, NULL, cnt, NULL);
+			knod_napi_kick(wpriv);
+		}
+
+		usleep_range(100, 200);
+	}
+	return 0;
+}
+
+static int knod_start_default_worker(struct knod *knod)
+{
+	struct task_struct *p;
+
+	knod->worker_fn = knod_default_worker;
+	knod->flush_fn = NULL;
+	knod->worker_ctx = knod;
+	p = kthread_run(knod_default_worker, knod, "knod_dflt_%d",
+			knod->accel->id);
+	if (IS_ERR(p))
+		return PTR_ERR(p);
+
+	get_task_struct(p);
+	knod->worker = p;
+	return 0;
+}
+
+static void knod_stop_worker(struct knod *knod)
+{
+	if (!knod->worker)
+		return;
+
+	if (knod->flush_fn)
+		knod->flush_fn(knod->worker_ctx);
+
+	kthread_stop(knod->worker);
+	put_task_struct(knod->worker);
+	knod->worker = NULL;
+	knod->worker_fn = NULL;
+	knod->flush_fn = NULL;
+	knod->worker_ctx = NULL;
+}
+
+int knod_register_worker(struct knod *knod, knod_worker_fn_t fn,
+			 knod_flush_fn_t flush, void *ctx)
+{
+	struct task_struct *p;
+
+	knod_stop_worker(knod);
+
+	knod->worker_fn = fn;
+	knod->flush_fn = flush;
+	knod->worker_ctx = ctx;
+	p = kthread_run(fn, ctx, "knod_%d", knod->accel->id);
+	if (IS_ERR(p)) {
+		knod->worker_fn = NULL;
+		knod->flush_fn = NULL;
+		knod->worker_ctx = NULL;
+		return PTR_ERR(p);
+	}
+
+	get_task_struct(p);
+	knod->worker = p;
+	return 0;
+}
+
+void knod_unregister_worker(struct knod *knod)
+{
+	knod_stop_worker(knod);
+	knod_start_default_worker(knod);
+}
+
+int knod_wait_on_events(struct kfd_process *p, u32 num_events,
+			void __user *data, bool all, u32 *user_timeout_ms,
+			u32 *wait_result)
+{
+	return kfd_wait_on_events_kernel(p, num_events, data, all,
+					 user_timeout_ms, wait_result);
+}
+EXPORT_SYMBOL(knod_wait_on_events);
+
+static int knod_alloc_ctx_init(struct knod *knod, int id, void **doorbell,
+			       struct kfd_topology_device **out_topo_dev,
+			       struct kfd_process_device **out_pdd)
+{
+	struct kfd_topology_device *topo_dev;
+	struct kfd_process_device *pdd;
+	struct file *drm_file;
+	size_t mem_size;
+	char path[64];
+	int err;
+
+	sprintf(path, "/dev/dri/renderD%d", id);
+
+	drm_file = filp_open(path, O_RDWR, 0);
+	if (IS_ERR(drm_file))
+		return PTR_ERR(drm_file);
+
+	knod->drm_file = drm_file;
+
+	knod->process = kfd_create_process(knod->umh_task);
+	if (IS_ERR(knod->process)) {
+		err = PTR_ERR(knod->process);
+		goto err_filp_close;
+	}
+	kref_get(&knod->process->ref);
+
+	err = knod_set_isa(knod);
+	if (err < 0)
+		goto err_unref_process;
+
+	pdd = knod->process->pdds[0];
+
+	topo_dev = kfd_topology_device_by_id(pdd->dev->id);
+	if (!topo_dev) {
+		pr_err("knod: can't find topo_dev for dev id %u\n",
+		       pdd->dev->id);
+		err = -ENODEV;
+		goto err_unref_process;
+	}
+
+	/* Acquire VM directly without installing an fd into any process's
+	 * fdtable.  get_file() provides the ref that pdd->drm_file will own.
+	 * kfd_process_destroy_pdds will fput it during process cleanup.
+	 */
+	get_file(drm_file);
+	mutex_lock(&knod->process->mutex);
+	err = kfd_process_device_init_vm(pdd, drm_file);
+	mutex_unlock(&knod->process->mutex);
+	if (err) {
+		fput(drm_file);
+		goto err_unref_process;
+	}
+
+	*doorbell = kfd_kernel_doorbell_mmap(pdd->dev, knod->process);
+	if (!*doorbell) {
+		err = -ENOMEM;
+		goto err_unref_process;
+	}
+
+	knod->pool = gen_pool_create(PAGE_SHIFT,
+				     dev_to_node(pdd->dev->adev->dev));
+	if (!knod->pool) {
+		err = -ENOMEM;
+		goto err_release_doorbell;
+	}
+
+	knod->dev = pdd->dev;
+	knod->reserved_addr = pdd->gpuvm_base << 2;
+	knod->limit_addr = pdd->gpuvm_limit;
+	mem_size = knod->limit_addr - knod->reserved_addr;
+
+	err = gen_pool_add(knod->pool, knod->reserved_addr, mem_size,
+			   dev_to_node(pdd->dev->adev->dev));
+	if (err)
+		goto err_gen_pool_destroy;
+
+	knod->kernels[0] = knod_alloc_mem(knod, PAGE_SIZE << 10,
+				      KFD_IOC_ALLOC_MEM_FLAGS_VRAM |
+				      KFD_IOC_ALLOC_MEM_FLAGS_COHERENT |
+				      KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE |
+				      KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE);
+	if (IS_ERR(knod->kernels[0])) {
+		err = PTR_ERR(knod->kernels[0]);
+		knod->kernels[0] = NULL;
+		goto err_gen_pool_destroy;
+	}
+	knod_init_default_kernel(knod);
+
+	/*
+	 * Second dispatch slot for the BPF ping-pong swap, allocated right
+	 * after knod->kernels[0] so both kernel BOs sit in the same low VA
+	 * region.
+	 * Allocating it later (at feature activate, past the RX buffers) put it
+	 * at a high VA, and switching the dispatch kernel_object to that BO
+	 * mid-stream wedged the compute queue.
+	 */
+	knod->kernels[1] = knod_alloc_mem(knod, PAGE_SIZE << 10,
+					  KFD_IOC_ALLOC_MEM_FLAGS_VRAM |
+					  KFD_IOC_ALLOC_MEM_FLAGS_COHERENT |
+					  KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE |
+					  KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE);
+	if (IS_ERR(knod->kernels[1])) {
+		err = PTR_ERR(knod->kernels[1]);
+		knod->kernels[1] = NULL;
+		goto err_free_kernel;
+	}
+
+	knod->mailbox = knod_alloc_mem(knod, PAGE_SIZE << 5,
+				       KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE |
+				       KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE |
+				       KFD_IOC_ALLOC_MEM_FLAGS_COHERENT |
+				       KFD_IOC_ALLOC_MEM_FLAGS_VRAM);
+	if (IS_ERR(knod->mailbox)) {
+		err = PTR_ERR(knod->mailbox);
+		knod->mailbox = NULL;
+		goto err_free_kernel;
+	}
+
+	*out_topo_dev = topo_dev;
+	*out_pdd = pdd;
+	return 0;
+
+err_free_kernel:
+	if (knod->kernels[1])
+		knod_free_mem(knod, knod->kernels[1]);
+	knod->kernels[1] = NULL;
+	knod_free_mem(knod, knod->kernels[0]);
+	knod->kernels[0] = NULL;
+err_gen_pool_destroy:
+	gen_pool_destroy(knod->pool);
+	knod->pool = NULL;
+err_release_doorbell:
+	iounmap(*doorbell);
+	*doorbell = NULL;
+err_unref_process:
+	kfd_unref_process(knod->process);
+err_filp_close:
+	fput(knod->drm_file);
+	return err;
+}
+
+struct knod *knod_alloc_ctx(struct knod_dev *knodev, int queue_cnt, int id,
+			    int channels)
+{
+	int buf_flags = KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE |
+			KFD_IOC_ALLOC_MEM_FLAGS_COHERENT |
+			KFD_IOC_ALLOC_MEM_FLAGS_VRAM;
+	struct kfd_topology_device *topo_dev;
+	struct kfd_process_device *pdd;
+	struct process_queue_node *pqn;
+	struct knod_work_priv *wpriv;
+	struct task_struct *task;
+	unsigned int sdma_cnt;
+	void *ptr = NULL;
+	struct knod *knod;
+	struct knod_mem *buf;
+	struct pid *spid;
+	struct queue *q;
+	int err, size;
+	pid_t pid;
+	int idx;
+
+	pid = launch_and_get_pid();
+	if (!pid) {
+		knod_err(" failed to create UMH\n");
+		return ERR_PTR(-EINVAL);
+	}
+
+	spid = find_get_pid(pid);
+	if (!spid) {
+		knod_err(" failed to find pid\n");
+		return ERR_PTR(-ESRCH);
+	}
+
+	task = get_pid_task(spid, PIDTYPE_PID);
+	put_pid(spid);
+	if (!task) {
+		knod_err(" no task found\n");
+		return ERR_PTR(-ESRCH);
+	}
+
+	knod = kzalloc(sizeof(struct knod), GFP_KERNEL);
+	if (!knod) {
+		put_task_struct(task);
+		return ERR_PTR(-ENOMEM);
+	}
+
+	knod->queue_cnt = queue_cnt;
+	knod->sdma_cnt = queue_cnt;
+	knod->channels = channels;
+	knod->umh_pid = pid;
+	knod->umh_task = task;
+	knod->nr_aql_ring = NR_AQL_RING;
+	INIT_LIST_HEAD(&knod->list);
+	INIT_LIST_HEAD(&knod->active_list);
+
+	err = knod_alloc_ctx_init(knod, id, &ptr, &topo_dev, &pdd);
+	if (err)
+		goto err_free_knod;
+
+	sdma_cnt = topo_dev->node_props.num_sdma_engines *
+		   topo_dev->node_props.num_sdma_queues_per_engine;
+	if (!sdma_cnt)
+		sdma_cnt = queue_cnt;
+	knod->sdma_cnt = min_t(int, queue_cnt, sdma_cnt);
+	if (!knod->sdma_cnt)
+		knod->sdma_cnt = 1;
+	pr_info("knod: AQL queues=%d SDMA queues=%d channels=%d\n",
+		queue_cnt, knod->sdma_cnt, channels);
+
+	knod->doorbell_base = ptr;
+
+	knod->buf = kmalloc_array(channels, sizeof(struct knod_mem *),
+				  GFP_KERNEL | __GFP_ZERO);
+	if (!knod->buf) {
+		err = -ENOMEM;
+		goto err_free_mailbox;
+	}
+
+	if (knod->igpu)
+		size = PAGE_SIZE << MAX_PAGE_ORDER;
+	else
+		size = PAGE_SIZE << 14;
+
+	for (idx = 0; idx < channels; idx++) {
+		wpriv = &knodev->wpriv[idx];
+		buf = __knod_alloc_mem(knod, size, buf_flags);
+		if (IS_ERR(buf)) {
+			err = PTR_ERR(buf);
+			goto err_free_bufs;
+		}
+		if (__knod_export_dma_buf(knod, buf)) {
+			knod_free_mem(knod, buf);
+			err = -ENOMEM;
+			goto err_free_bufs;
+		}
+		if (__knod_map_kaddr(knod, buf)) {
+			knod_free_mem(knod, buf);
+			err = -ENOMEM;
+			goto err_free_bufs;
+		}
+		if (__knod_map_mem(knod, buf)) {
+			knod_free_mem(knod, buf);
+			err = -ENOMEM;
+			goto err_free_bufs;
+		}
+		wpriv->dmabuf = buf->mem->dmabuf;
+		wpriv->index = idx;
+		knod->buf[idx] = buf;
+	}
+
+	/*
+	 * GPU->host delivery buffer + per-queue page_pools are owned by the
+	 * NOD framework (knod_pass_attach), allocated via accel_ops->alloc_mem.
+	 */
+
+	memset(knod->mailbox->kaddr, 0, knod->mailbox->size);
+
+	err = kfd_event_page_set(knod->process,
+				 knod->mailbox->kaddr,
+				 KFD_SIGNAL_EVENT_LIMIT * 8,
+				 knod->mailbox->gaddr);
+	if (err < 0)
+		goto err_free_bufs;
+
+	for (idx = 0; idx < queue_cnt; idx++) {
+		err = knod_alloc_one_queue(knod, idx, topo_dev, pdd, ptr);
+		if (err)
+			goto err_free_queues;
+	}
+
+	/* Allocate only the SDMA queues the device can actually provide. BPF
+	 * XDP_TX does not use SDMA for the verdict path, while PASS/d2h maps RX
+	 * queues onto the available SDMA queues modulo sdma_cnt.
+	 */
+	for (idx = 0; idx < knod->sdma_cnt; idx++) {
+		struct amd_signal *sdma_signal;
+		struct queue_properties qp;
+		long *sdma_lptr;
+		int sdma_flags = KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE |
+				 KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE |
+				 KFD_IOC_ALLOC_MEM_FLAGS_COHERENT |
+				 KFD_IOC_ALLOC_MEM_FLAGS_GTT;
+
+		knod->sdma[idx].sdma = knod_alloc_mem(knod, PAGE_SIZE << 4,
+			KFD_IOC_ALLOC_MEM_FLAGS_GTT |
+			KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE |
+			KFD_IOC_ALLOC_MEM_FLAGS_COHERENT |
+			KFD_IOC_ALLOC_MEM_FLAGS_UNCACHED |
+			KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE);
+		if (IS_ERR(knod->sdma[idx].sdma)) {
+			err = PTR_ERR(knod->sdma[idx].sdma);
+			knod->sdma[idx].sdma = NULL;
+			goto err_free_sdma;
+		}
+		knod->sdma[idx].queue = knod_alloc_mem(knod, PAGE_SIZE,
+						       sdma_flags);
+		if (IS_ERR(knod->sdma[idx].queue)) {
+			err = PTR_ERR(knod->sdma[idx].queue);
+			knod->sdma[idx].queue = NULL;
+			goto err_free_sdma;
+		}
+		knod->sdma[idx].queue_signal = knod_alloc_mem(knod,
+			PAGE_SIZE << 5,
+			KFD_IOC_ALLOC_MEM_FLAGS_GTT |
+			KFD_IOC_ALLOC_MEM_FLAGS_COHERENT |
+			KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE |
+			KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE);
+		if (IS_ERR(knod->sdma[idx].queue_signal)) {
+			err = PTR_ERR(knod->sdma[idx].queue_signal);
+			knod->sdma[idx].queue_signal = NULL;
+			goto err_free_sdma;
+		}
+		knod->sdma[idx].idx = 0;
+
+		err = knod_create_event(knod->process, KFD_IOC_EVENT_SIGNAL,
+					true, 1, &knod->sdma_event[idx]);
+		if (err) {
+			knod_err(" failed to create SDMA event[%d] err=%d\n",
+				 idx, err);
+			goto err_free_sdma;
+		}
+
+		memset(&qp, 0, sizeof(qp));
+		qp.type = KFD_QUEUE_TYPE_SDMA;
+		qp.format = KFD_QUEUE_FORMAT_PM4;
+		qp.queue_percent = 100;
+		qp.priority = 15;
+		qp.queue_address = (u64)knod->sdma[idx].sdma->gaddr;
+		qp.queue_size = knod->sdma[idx].sdma->size;
+		qp.write_ptr = (void __user *)
+			((u64)knod->sdma[idx].queue->gaddr + 0x08);
+		qp.read_ptr = (void __user *)
+			((u64)knod->sdma[idx].queue->gaddr + 0x10);
+
+		sdma_signal = (struct amd_signal *)
+			knod->sdma[idx].queue_signal->kaddr;
+		sdma_signal->kind = AMD_SIGNAL_KIND_USER;
+		sdma_signal->event_id = knod->sdma_event[idx].id;
+		sdma_signal->queue_ptr = 0;
+		sdma_signal->event_mailbox_ptr = knod->mailbox->gaddr +
+			(knod->sdma_event[idx].slot * 8);
+		/* SDMA fence writes u32 to low 32 bits of value.
+		 * Init to 0 (not -1 like AQL) so comparison works.
+		 * value is in union with hardware_doorbell_ptr - set last.
+		 */
+		sdma_signal->value = 0;
+		sdma_lptr = knod->mailbox->kaddr + (idx * 8);
+		sdma_lptr[knod->sdma_event[idx].slot] = 0;
+
+		err = knod_create_queue(knod->process, &qp, pdd->user_gpu_id,
+					&knod->sdma_queue_id[idx],
+					&knod->sdma_doorbell_offset[idx]);
+		if (err) {
+			knod_err(" failed to create SDMA queue[%d] err=%d\n",
+				 idx, err);
+			goto err_free_sdma;
+		}
+		knod->sdma_queue_created[idx] = true;
+
+		knod_init_queue(knod, knod->sdma_queue_id[idx], idx,
+				KFD_IOC_QUEUE_TYPE_SDMA);
+		knod->sdma[idx].doorbell =
+			(u64 *)((u8 *)ptr +
+				(u32)knod->sdma_doorbell_offset[idx]);
+		knod_dbg(" SDMA queue[%d] created: id=%d doorbell_offset=0x%x doorbell=%p\n",
+			 idx, knod->sdma_queue_id[idx],
+			 (u32)knod->sdma_doorbell_offset[idx],
+			 knod->sdma[idx].doorbell);
+	}
+
+	list_for_each_entry(pqn, &knod->process->pqm.queues,
+			    process_queue_list) {
+		if (!pqn->q)
+			continue;
+		q = pqn->q;
+		for (idx = 0; idx < queue_cnt; idx++) {
+			if (q->properties.queue_id == idx) {
+				knod->kaql[idx].doorbell =
+					(u64 *)((u8 *)ptr +
+					(u32)knod->aql_doorbell_offset[idx]);
+				q->properties.doorbell_ptr =
+					knod->kaql[idx].doorbell;
+				knod_dbg(" doorbell: idx=%d qid=%d door_off=0x%x doorbell=%p\n",
+					 idx, q->properties.queue_id,
+					 (u32)knod->aql_doorbell_offset[idx],
+					 knod->kaql[idx].doorbell);
+			}
+		}
+	}
+
+	list_add_tail(&knod->list, &ctx_list);
+
+	/* Create shared debugfs directory under dri/<N>/knod/ */
+	{
+		struct drm_minor *minor = adev_to_drm(pdd->dev->adev)->render;
+
+		if (minor && minor->debugfs_root) {
+			struct dentry *dir;
+
+			dir = debugfs_lookup("knod", minor->debugfs_root);
+			if (!dir) {
+				dir = debugfs_create_dir("knod",
+							 minor->debugfs_root);
+				if (IS_ERR(dir))
+					dir = NULL;
+			}
+			knod->debug_dir = dir;
+		}
+	}
+
+	return knod;
+
+err_free_sdma:
+	{
+		int j;
+
+		for (j = idx; j >= 0; j--) {
+			if (knod->sdma_queue_created[j]) {
+				knod_destroy_queue(knod->process,
+						   knod->sdma_queue_id[j]);
+				knod->sdma_queue_created[j] = false;
+			}
+			if (knod->sdma_event[j].id) {
+				knod_destroy_event(knod->process,
+						   knod->sdma_event[j].id);
+				knod->sdma_event[j].id = 0;
+			}
+			knod_free_mem(knod, knod->sdma[j].sdma);
+			knod_free_mem(knod, knod->sdma[j].queue);
+			knod_free_mem(knod, knod->sdma[j].queue_signal);
+		}
+	}
+	idx = queue_cnt;
+err_free_queues:
+	while (idx-- > 0)
+		knod_destroy_one_queue(knod, idx);
+err_free_bufs:
+	for (idx = 0; idx < channels; idx++)
+		knod_free_mem(knod, knod->buf[idx]);
+	kfree(knod->buf);
+err_free_mailbox:
+	knod_free_mem(knod, knod->mailbox);
+	knod_free_mem(knod, knod->kernels[1]);
+	knod_free_mem(knod, knod->kernels[0]);
+	gen_pool_destroy(knod->pool);
+	iounmap(ptr);
+	kfd_unref_process(knod->process);
+	fput(knod->drm_file);
+err_free_knod:
+	put_task_struct(knod->umh_task);
+	stop_umh(knod->umh_pid);
+	kfree(knod);
+	return ERR_PTR(err);
+}
+EXPORT_SYMBOL(knod_alloc_ctx);
+
+void knod_release_ctx(struct knod *knod)
+{
+	int idx;
+
+	list_del(&knod->list);
+
+	debugfs_remove_recursive(knod->debug_dir);
+
+	/* Release SDMA queues - destroy queue/event BEFORE freeing BOs.
+	 * Same ordering as knod_destroy_one_queue() for AQL: the queue
+	 * holds references to BO VAs, so freeing BOs first causes the
+	 * queue destroy to fail and leaves internal BOs in kfd_bo_list.
+	 */
+	for (idx = 0; idx < knod->sdma_cnt; idx++) {
+		knod_destroy_queue(knod->process, knod->sdma_queue_id[idx]);
+		knod_destroy_event(knod->process, knod->sdma_event[idx].id);
+		knod_free_mem(knod, knod->sdma[idx].sdma);
+		knod_free_mem(knod, knod->sdma[idx].queue);
+		knod_free_mem(knod, knod->sdma[idx].queue_signal);
+	}
+
+	for (idx = 0; idx < knod->queue_cnt; idx++)
+		knod_destroy_one_queue(knod, idx);
+
+	for (idx = 0; idx < knod->channels; idx++)
+		knod_free_mem(knod, knod->buf[idx]);
+	kfree(knod->buf);
+
+	knod_free_mem(knod, knod->mailbox);
+	knod_free_mem(knod, knod->kernels[1]);
+	knod_free_mem(knod, knod->kernels[0]);
+	gen_pool_destroy(knod->pool);
+	iounmap(knod->doorbell_base);
+
+	/* Tear down the KFD process.  It holds 3 refs:
+	 *   ref 1: "open" ref from kfd_create_process (normally dropped
+	 *           by kfd_release when /dev/kfd is closed - KNOD never
+	 *           opens /dev/kfd so we must drop this ourselves)
+	 *   ref 2: KNOD's explicit kref_get in knod_alloc_ctx
+	 *   ref 3: mmu_notifier alloc_notifier ref (dropped by
+	 *           free_notifier callback via SRCU after mmu_notifier_put)
+	 *
+	 * kfd_process_notifier_release_internal removes from hash,
+	 * destroys queues, and calls mmu_notifier_put which schedules
+	 * the SRCU callback to drop ref 3.  We drop refs 1 and 2 here.
+	 * After SRCU fires, ref reaches 0 -> kfd_process_wq_release
+	 * runs (sysfs removal, BO/PDD/doorbell/event cleanup, kfree).
+	 *
+	 * kfd_process_destroy_pdds fput's pdd->drm_file (the get_file
+	 * ref from init_vm), bringing the DRM file ref from 2->1.
+	 * The file/VM stays alive until we fput the filp_open ref below.
+	 */
+	kfd_process_notifier_release_internal(knod->process);
+	kfd_unref_process(knod->process);
+	kfd_unref_process(knod->process);
+	mmu_notifier_synchronize();
+	kfd_process_flush_wq();
+
+	/* Drop the filp_open ref (file ref 1->0).  kfd_process_destroy_pdds
+	 * already fput'd the get_file ref during wq_release above.
+	 */
+	fput(knod->drm_file);
+	flush_delayed_fput();
+
+	stop_umh(knod->umh_pid);
+	put_task_struct(knod->umh_task);
+	kfree(knod);
+}
+EXPORT_SYMBOL(knod_release_ctx);
+
+/* ---- feature control (knod genetlink) ---- */
+
+static void *knod_feature_ops(enum knod_feature feat)
+{
+	switch (feat) {
+	case KNOD_FEATURE_BPF:
+		return registered_xdp_ops;
+	case KNOD_FEATURE_IPSEC:
+		return accel_ops.ipsec_ops;
+	default:
+		return NULL;
+	}
+}
+
+/*
+ * Feature lifecycle across three independent axes:
+ *   init/exit            attach/detach      permanent per-attach state
+ *   activate/deactivate  feature select     feature GPU resources
+ *   start/stop           interface up/down  worker + GPU in-flight drain
+ *
+ * The worker only does useful work while (interface up AND a feature is
+ * selected), so a feature switch is bracketed stop -> change -> start: the
+ * worker is stopped and its GPU dispatch drained before deactivate() frees
+ * the resources it used, then restarted afterwards.  start/stop never touch
+ * the framework/NIC-owned RX SPSC ring (mlx5 co-owns it via rq->knodev), so
+ * they are safe to run while the interface is up and traffic is flowing.
+ */
+static void knod_feature_stop(struct knod *knod)
+{
+	struct knod_dev *knodev = knod->accel->knodev;
+
+	knod_stop_worker(knod);
+
+	switch (knod->active_feature) {
+	case KNOD_FEATURE_BPF:
+		if (registered_xdp_ops && registered_xdp_ops->stop)
+			registered_xdp_ops->stop(knodev);
+		break;
+	case KNOD_FEATURE_IPSEC:
+		if (accel_ops.ipsec_ops && accel_ops.ipsec_ops->stop)
+			accel_ops.ipsec_ops->stop(knodev);
+		break;
+	default:
+		break;
+	}
+}
+
+static void knod_feature_start(struct knod *knod)
+{
+	struct knod_dev *knodev = knod->accel->knodev;
+
+	switch (knod->active_feature) {
+	case KNOD_FEATURE_BPF:
+		if (registered_xdp_ops && registered_xdp_ops->start)
+			registered_xdp_ops->start(knodev);
+		break;
+	case KNOD_FEATURE_IPSEC:
+		if (accel_ops.ipsec_ops && accel_ops.ipsec_ops->start)
+			accel_ops.ipsec_ops->start(knodev);
+		break;
+	default:
+		knod_start_default_worker(knod);
+		break;
+	}
+}
+
+static void knod_feature_deactivate(struct knod *knod)
+{
+	struct knod_dev *knodev = knod->accel->knodev;
+
+	switch (knod->active_feature) {
+	case KNOD_FEATURE_BPF:
+		/*
+		 * Phase 1: unregister the offload dev - this force-frees any
+		 * user XDP progs/maps still bound.  The map-free ndo routes
+		 * back through accel_ops.xdp_ops->xdp_install, so xdp_ops must
+		 * still point at the BPF ops (and priv must be alive) here.
+		 */
+		if (registered_xdp_ops &&
+		    registered_xdp_ops->xdp_offload_uninit)
+			registered_xdp_ops->xdp_offload_uninit(knodev);
+		/*
+		 * Phase 2: repoint RX delivery at the default drain_pass and
+		 * wait out the in-flight NAPI readers.
+		 */
+		WRITE_ONCE(accel_ops.xdp_ops, &default_xdp_ops);
+		synchronize_net();
+		knod_stop_worker(knod);
+		if (registered_xdp_ops && registered_xdp_ops->deactivate)
+			registered_xdp_ops->deactivate(knodev);
+		break;
+	case KNOD_FEATURE_IPSEC:
+		/*
+		 * knod_ipsec_detach() clears the netdev xfrm flags and runs
+		 * ->deactivate(), which NULLs ipsec_priv and does its own
+		 * synchronize_net() before freeing.
+		 */
+		knod_ipsec_detach(knodev);
+		break;
+	default:
+		break;
+	}
+}
+
+static int knod_feature_activate(struct knod *knod)
+{
+	struct knod_dev *knodev = knod->accel->knodev;
+	int err;
+
+	switch (knod->active_feature) {
+	case KNOD_FEATURE_BPF:
+		if (!registered_xdp_ops)
+			return -ENODEV;
+		/* Phase A: GPU compute buffers. */
+		if (registered_xdp_ops->activate) {
+			err = registered_xdp_ops->activate(knodev);
+			if (err)
+				return err;
+		}
+		/*
+		 * Publish xdp_ops, then phase B: register the offload dev so
+		 * user XDP progs/maps can bind (their install/free route
+		 * through accel_ops.xdp_ops->xdp_install).
+		 */
+		WRITE_ONCE(accel_ops.xdp_ops, registered_xdp_ops);
+		if (registered_xdp_ops->xdp_offload_init) {
+			err = registered_xdp_ops->xdp_offload_init(knodev);
+			if (err) {
+				WRITE_ONCE(accel_ops.xdp_ops, &default_xdp_ops);
+				synchronize_net();
+				knod_stop_worker(knod);
+				if (registered_xdp_ops->deactivate)
+					registered_xdp_ops->deactivate(knodev);
+				return err;
+			}
+		}
+		return 0;
+	case KNOD_FEATURE_IPSEC:
+		if (!accel_ops.ipsec_ops)
+			return -ENODEV;
+		/* knod_ipsec_attach() runs ->activate() + sets netdev flags. */
+		return knod_ipsec_attach(knodev);
+	case KNOD_FEATURE_NONE:
+		return 0;
+	default:
+		return -EINVAL;
+	}
+}
+
+static int knod_accel_feature_get(struct knod_accel *accel,
+				u32 *ena, u32 *cap)
+{
+	struct knod *knod = READ_ONCE(accel->priv);
+	u32 mask = 0;
+	int i;
+
+	/* A registered-but-not-attached accel has no context yet. */
+	*ena = knod ? knod->active_feature : KNOD_FEATURE_NONE;
+	for (i = 0; i < KNOD_FEATURE_MAX; i++)
+		if (i == KNOD_FEATURE_NONE || knod_feature_ops(i))
+			mask |= BIT(i);
+	*cap = mask;
+	return 0;
+}
+
+static int knod_accel_feature_set(struct knod_accel *accel, u32 feature,
+				struct netlink_ext_ack *extack)
+{
+	struct knod *knod = READ_ONCE(accel->priv);
+	struct knod_dev *knodev;
+	bool started;
+	int err = 0;
+
+	if (feature >= KNOD_FEATURE_MAX) {
+		NL_SET_ERR_MSG(extack, "unknown feature");
+		return -EINVAL;
+	}
+	if (feature != KNOD_FEATURE_NONE && !knod_feature_ops(feature)) {
+		NL_SET_ERR_MSG(extack, "feature not available on this accelerator");
+		return -ENOENT;
+	}
+	knodev = knod ? READ_ONCE(knod->accel->knodev) : NULL;
+	if (!knod || !knodev) {
+		NL_SET_ERR_MSG(extack, "accelerator not attached to a NIC");
+		return -ENODEV;
+	}
+	if (feature == knod->active_feature)
+		return 0;
+
+	/*
+	 * Refuse to leave BPF while a user XDP prog or offloaded map is still
+	 * bound - tearing the offload down underneath them is unsafe (the GPU
+	 * dispatch keeps running against freed state).  The user must detach
+	 * first, e.g. "ip link set dev <if> xdp off".
+	 */
+	if (knod->active_feature == KNOD_FEATURE_BPF && registered_xdp_ops &&
+	    registered_xdp_ops->busy && registered_xdp_ops->busy(knodev)) {
+		NL_SET_ERR_MSG(extack, "detach the XDP program/maps first");
+		return -EBUSY;
+	}
+	if (knod->active_feature == KNOD_FEATURE_IPSEC && accel_ops.ipsec_ops &&
+	    accel_ops.ipsec_ops->busy && accel_ops.ipsec_ops->busy(knodev)) {
+		NL_SET_ERR_MSG(extack, "remove the offloaded xfrm SAs first");
+		return -EBUSY;
+	}
+
+	/*
+	 * stop -> change -> start.  The worker only runs while the interface
+	 * is up; bracket the resource change with worker stop/start in that
+	 * case.  stop() drains the GPU in-flight so deactivate() frees safely.
+	 */
+	started = READ_ONCE(knodev->started);
+	if (started)
+		knod_feature_stop(knod);
+	knod_feature_deactivate(knod);
+	knod->active_feature = KNOD_FEATURE_NONE;
+
+	if (feature != KNOD_FEATURE_NONE) {
+		knod->active_feature = feature;
+		err = knod_feature_activate(knod);
+		if (err) {
+			knod->active_feature = KNOD_FEATURE_NONE;
+			NL_SET_ERR_MSG(extack, "failed to activate feature");
+		}
+	}
+	if (started)
+		knod_feature_start(knod);
+	return err;
+}
+
+static int knod_attach(struct knod_dev *knodev)
+{
+	struct knod_accel *accel = knodev->accel;
+	int render_idx = accel->id / KNOD_MAX_AQL;
+	struct net_device *netdev = knodev->netdev;
+	struct amdgpu_device *adev;
+	struct knod *knod;
+
+	{
+		int q_cnt = clamp(READ_ONCE(knod_requested_queue_cnt), 1,
+				      KNOD_MAX_QUEUE_CNT);
+
+		knod = knod_alloc_ctx(knodev, q_cnt, render_idx,
+				      min(netdev->num_rx_queues,
+					  KNOD_SPSC_MAX));
+	}
+	if (IS_ERR(knod)) {
+		pr_err("knod: Failed to allocate context\n");
+		return -EINVAL;
+	}
+
+	accel->priv = knod;
+	knod->accel = accel;
+
+	/*
+	 * Keep GFX engine out of GFXOFF while a KNOD accel is attached.
+	 * On RDNA2 (tested RX6600 / gfx1032) the very first AQL dispatch
+	 * after boot hangs with signal=-1 when GFXOFF is active - SMU exit
+	 * from GFXOFF races with the doorbell ring and the completion
+	 * signal is never decremented. Forcing GFXOFF off for the attached
+	 * lifetime avoids the race entirely.
+	 *
+	 * Pin MCLK soft-min to HW max. KNOD's RX path is bandwidth-bound
+	 * on VRAM (NIC->GPU p2pdma delivers frames directly to VRAM, then
+	 * the shader streams them back out) but the individual dispatches
+	 * are too short for SMU's activity monitor to react - MCLK sticks
+	 * at the lowest DPM (~96 MHz on RX6600) in AUTO mode and caps
+	 * throughput far below what the compute path can sustain. Switching
+	 * to the COMPUTE power profile did not help on RDNA2: the COMPUTE
+	 * DpmActivityMonitor coefficients tune GFX upclock aggressively
+	 * but leave memory activity detection conservative.
+	 *
+	 * Setting soft_min via SetSoftMinByFreq is independent of
+	 * pp_power_profile_mode and of power_dpm_force_performance_level,
+	 * so AUTO governance stays in effect for SCLK/voltage - we get the
+	 * same 30W full-throughput state the user reaches via "force
+	 * performance = manual + echo 3 > pp_dpm_mclk", without the 75W
+	 * voltage pin that PROFILE_PEAK imposes. Passing 0xFFFF MHz lets
+	 * SMU firmware clamp to the actual hardware max.
+	 */
+	adev = knod->process->pdds[0]->dev->adev;
+	amdgpu_gfx_off_ctrl_immediate(adev, false);
+	amdgpu_dpm_set_soft_freq_range(adev, PP_MCLK, 0xFFFF, 0xFFFF);
+
+	/*
+	 * Attach settles in KNOD_FEATURE_NONE with no worker running.  The
+	 * worker is started by the NIC driver bringing the interface up
+	 * (knod_dev_start -> ->dev_start), and each feature's GPU resources
+	 * are allocated when the feature is selected (->activate).
+	 */
+	knod->active_feature = KNOD_FEATURE_NONE;
+
+	/*
+	 * Permanent per-attach feature state (e.g. the BPF bpf_offload_dev,
+	 * which must outlive feature switches so user XDP progs/maps survive).
+	 */
+	if (registered_xdp_ops && registered_xdp_ops->init)
+		registered_xdp_ops->init(knodev);
+	return 0;
+}
+
+static void knod_pre_detach(struct knod_dev *knodev)
+{
+	struct knod *knod = knodev->accel->priv;
+
+	/*
+	 * Detach requires the interface down, so the worker is already
+	 * stopped; stop again defensively, free the active feature's
+	 * resources, then tear down the permanent per-attach state.
+	 */
+	knod_feature_stop(knod);
+	knod_feature_deactivate(knod);
+	knod->active_feature = KNOD_FEATURE_NONE;
+
+	if (registered_xdp_ops && registered_xdp_ops->exit)
+		registered_xdp_ops->exit(knodev);
+}
+
+static void knod_dev_start_worker(struct knod_dev *knodev)
+{
+	struct knod *knod = knodev->accel->priv;
+
+	/* Interface up: start the current feature's worker. */
+	if (knod)
+		knod_feature_start(knod);
+}
+
+static void knod_dev_stop_worker(struct knod_dev *knodev)
+{
+	struct knod *knod = knodev->accel->priv;
+
+	/* Interface down: stop the worker + drain the GPU in-flight. */
+	if (knod)
+		knod_feature_stop(knod);
+}
+
+static void knod_detach(struct knod_dev *knodev)
+{
+	struct knod_accel *accel = knodev->accel;
+	struct knod *knod = accel->priv;
+	struct amdgpu_device *adev = knod->process->pdds[0]->dev->adev;
+
+	knod_stop_worker(knod);
+	knod_release_ctx(knod);
+	WRITE_ONCE(accel->priv, NULL);
+
+	/*
+	 * Restore default MCLK range. min=1 triggers SetSoftMinByFreq (the
+	 * API skips the call when min==0) and SMU clamps to HW min; max
+	 * 0xFFFF clamps to HW max. Together this matches the pre-attach
+	 * "no soft constraint" state so DPM can idle MCLK back down.
+	 */
+	amdgpu_dpm_set_soft_freq_range(adev, PP_MCLK, 1, 0xFFFF);
+	amdgpu_gfx_off_ctrl(adev, true);
+}
+
+static void *knod_accel_alloc_mem(struct knod_dev *knodev, size_t size,
+				  u64 *gaddr, struct page ***pages, void **priv)
+{
+	struct knod_accel *accel = knodev->accel;
+	struct knod *knod = accel->priv;
+	struct knod_mem *mem;
+	struct ttm_tt *tt;
+	u32 flags;
+
+	/* SPSC rings are hot producer/consumer control data. Keep them plain
+	 * GTT; only host-read delivery buffers need coherent CPU visibility.
+	 */
+	flags = KFD_IOC_ALLOC_MEM_FLAGS_GTT | KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE;
+	if (pages)
+		flags |= KFD_IOC_ALLOC_MEM_FLAGS_COHERENT;
+
+	mem = knod_alloc_mem(knod, size, flags);
+	if (IS_ERR(mem))
+		return NULL;
+
+	if (pages) {
+		tt = mem->mem->bo ? mem->mem->bo->tbo.ttm : NULL;
+		if (!tt || !tt->pages) {
+			knod_free_mem(knod, mem);
+			return NULL;
+		}
+		*pages = tt->pages;
+	}
+
+	*gaddr = mem->gaddr;
+	*priv = mem;
+	return mem->kaddr;
+}
+
+static void knod_accel_free_mem(struct knod_dev *knodev, void *priv)
+{
+	struct knod_accel *accel = knodev->accel;
+	struct knod *knod = accel->priv;
+	struct knod_mem *mem = priv;
+
+	knod_free_mem(knod, mem);
+}
+
+/*
+ * Device->host copy primitives for the common knod_d2h_copy/knod_d2h_drain
+ * path.  Thin wrappers over the SDMA engine (sdma[0]); the framework owns the
+ * pending ring, fence counter and dst pool.
+ */
+static u32 knod_accel_d2h_submit(struct knod_dev *knodev, u64 dst, int queue,
+				 u32 page_idx, u16 off, u32 len)
+{
+	struct knod *knod = knodev->accel->priv;
+	struct knod_sdma_copy_desc c;
+
+	/* The netmem page_pool dma_addr is the NIC address; SDMA needs
+	 * the GPU VM address, so derive the source from buf[queue].
+	 */
+	c.src = knod->buf[queue]->gaddr + ((u64)page_idx << PAGE_SHIFT) + off;
+	c.dst = dst;
+	c.len = len;
+	return knod_sdma_submit(knod, 0, &c, 1);
+}
+
+static void knod_accel_d2h_kick(struct knod_dev *knodev)
+{
+	knod_sdma_kick(knodev->accel->priv, 0);
+}
+
+static u32 knod_accel_d2h_fence(struct knod_dev *knodev, int sdma_idx)
+{
+	struct knod *knod = READ_ONCE(knodev->accel->priv);
+	struct knod_mem *signal;
+
+	if (!knod || sdma_idx < 0 || sdma_idx >= READ_ONCE(knod->sdma_cnt))
+		return 0;
+	signal = READ_ONCE(knod->sdma[sdma_idx].queue_signal);
+	if (!signal || !signal->kaddr)
+		return 0;
+
+	return (u32)READ_ONCE(((struct amd_signal *)signal->kaddr)->value);
+}
+
+static int knod_accel_mp_map(struct knod_dev *knodev)
+{
+	struct knod *knod = knodev->accel->priv;
+	int i;
+
+	if (!knod)
+		return -ENODEV;
+
+	for (i = 0; i < knod->channels; i++)
+		if (__knod_map_mem(knod, knod->buf[i]))
+			knod_err(" mp_map failed ch %d\n", i);
+
+	return 0;
+}
+
+static struct knod_accel_ops accel_ops = {
+	.attach = knod_attach,
+	.pre_detach = knod_pre_detach,
+	.detach = knod_detach,
+	.dev_start = knod_dev_start_worker,
+	.dev_stop = knod_dev_stop_worker,
+	.alloc_mem = knod_accel_alloc_mem,
+	.free_mem = knod_accel_free_mem,
+	.mp_map = knod_accel_mp_map,
+	.d2h_submit = knod_accel_d2h_submit,
+	.d2h_kick = knod_accel_d2h_kick,
+	.d2h_fence = knod_accel_d2h_fence,
+	.xdp_ops = &default_xdp_ops,
+	.feature_get = knod_accel_feature_get,
+	.feature_set = knod_accel_feature_set,
+};
+
+/*
+ * Tear a feature down on every accel still using it before its ops pointer
+ * is cleared (module unload).  Mirrors the feature_set transition to NONE.
+ */
+static void knod_feature_force_none(enum knod_feature feat)
+{
+	int i;
+
+	for (i = 0; i < nr_accels; i++) {
+		struct knod_dev *knodev;
+		struct knod *knod;
+		bool started;
+
+		if (!accels[i])
+			continue;
+		knod = READ_ONCE(accels[i]->priv);
+		knodev = knod ? READ_ONCE(accels[i]->knodev) : NULL;
+		if (!knod || !knodev || knod->active_feature != feat)
+			continue;
+		started = READ_ONCE(knodev->started);
+		if (started)
+			knod_feature_stop(knod);
+		knod_feature_deactivate(knod);
+		knod->active_feature = KNOD_FEATURE_NONE;
+		if (started)
+			knod_feature_start(knod);
+	}
+}
+
+void knod_accel_xdp_register(struct knod_accel_xdp_ops *xdp_ops)
+{
+	int i;
+
+	/*
+	 * Module load advertises the feature and sets up the permanent
+	 * per-attach state (bpf_offload_dev) on already-attached accels;
+	 * GPU compute resources wait for ->activate() on feature select.
+	 */
+	WRITE_ONCE(registered_xdp_ops, xdp_ops);
+	for (i = 0; i < nr_accels; i++) {
+		struct knod_dev *knodev;
+
+		if (!accels[i])
+			continue;
+		knodev = READ_ONCE(accels[i]->knodev);
+		if (knodev && xdp_ops->init)
+			xdp_ops->init(knodev);
+	}
+}
+EXPORT_SYMBOL(knod_accel_xdp_register);
+
+void knod_accel_xdp_unregister(void)
+{
+	int i;
+
+	/* Force any active BPF feature off, then drop the permanent state. */
+	knod_feature_force_none(KNOD_FEATURE_BPF);
+	for (i = 0; i < nr_accels; i++) {
+		struct knod_dev *knodev;
+
+		if (!accels[i])
+			continue;
+		knodev = READ_ONCE(accels[i]->knodev);
+		if (knodev && registered_xdp_ops && registered_xdp_ops->exit)
+			registered_xdp_ops->exit(knodev);
+	}
+	WRITE_ONCE(registered_xdp_ops, NULL);
+}
+EXPORT_SYMBOL(knod_accel_xdp_unregister);
+
+void knod_accel_ipsec_register(struct knod_accel_ipsec_ops *ipsec_ops)
+{
+	/* Advertise only; resources are allocated by ->activate() on select. */
+	WRITE_ONCE(accel_ops.ipsec_ops, ipsec_ops);
+}
+EXPORT_SYMBOL(knod_accel_ipsec_register);
+
+void knod_accel_ipsec_unregister(void)
+{
+	knod_feature_force_none(KNOD_FEATURE_IPSEC);
+	WRITE_ONCE(accel_ops.ipsec_ops, NULL);
+}
+EXPORT_SYMBOL(knod_accel_ipsec_unregister);
+
+/*
+ * Accel modules call this from their module_init before NOD attach to
+ * request the minimum AQL/SDMA queue pair count their dispatcher
+ * infrastructure needs. The value is a high-water mark: multiple
+ * callers raise it but never lower it, so whichever module needs the
+ * most queues wins.
+ *
+ * Must be called before knod_attach() runs - i.e. before any
+ * `echo X,0 > /sys/kernel/debug/knod_dev/attach` - otherwise
+ * the already-created knod context keeps its old queue_cnt and the
+ * caller must unbind/rebind to pick up the new value.
+ *
+ * Clamped to [1, KNOD_MAX_QUEUE_CNT]. Values outside that range are ignored
+ * (the internal cap can grow later without ABI break).
+ */
+void knod_request_queue_cnt(int n)
+{
+	int cur;
+
+	if (n < 1 || n > KNOD_MAX_QUEUE_CNT)
+		return;
+
+	do {
+		cur = READ_ONCE(knod_requested_queue_cnt);
+		if (n <= cur)
+			return;
+	} while (cmpxchg(&knod_requested_queue_cnt, cur, n) != cur);
+}
+EXPORT_SYMBOL(knod_request_queue_cnt);
+
+int knod_init(struct amdgpu_device *adev)
+{
+	int base_id = adev_to_drm(adev)->render->index * KNOD_MAX_AQL;
+	struct knod_accel *accel;
+	int i, n;
+
+	n = KNOD_MAX_AQL;
+
+	for (i = 0; i < n; i++) {
+		accel = kzalloc_obj(struct knod_accel, GFP_KERNEL);
+		if (!accel) {
+			pr_err("knod: Failed to allocate accel node %d\n", i);
+			goto err;
+		}
+
+		INIT_LIST_HEAD(&accel->list);
+		accel->type = KNOD_TYPE_GPU;
+		accel->accel_ops = &accel_ops;
+		accel->owner = THIS_MODULE;
+		accel->id = base_id + i;
+		snprintf(accel->name, sizeof(accel->name), "amdgpu-%d", i);
+		knod_accel_register(accel);
+		accels[i] = accel;
+	}
+	nr_accels = n;
+
+	return 0;
+
+err:
+	while (i-- > 0) {
+		knod_accel_unregister(accels[i]);
+		kfree(accels[i]);
+		accels[i] = NULL;
+	}
+	return 0;
+}
+
+void knod_fini(struct amdgpu_device *adev)
+{
+	int base_id = adev_to_drm(adev)->render->index * KNOD_MAX_AQL;
+	int i;
+
+	for (i = 0; i < nr_accels; i++) {
+		if (!accels[i] ||
+		    accels[i]->id / KNOD_MAX_AQL != base_id / KNOD_MAX_AQL)
+			continue;
+		knod_accel_unregister(accels[i]);
+		kfree(accels[i]);
+		accels[i] = NULL;
+	}
+}
+
+void knod_exit(void)
+{
+	struct knod *knod, *tmp;
+	int i;
+
+	for (i = 0; i < nr_accels; i++) {
+		if (!accels[i])
+			continue;
+		knod_accel_unregister(accels[i]);
+		kfree(accels[i]);
+		accels[i] = NULL;
+	}
+	nr_accels = 0;
+
+	list_for_each_entry_safe(knod, tmp, &ctx_list, list) {
+		knod_release_ctx(knod);
+	}
+}
diff --git a/drivers/gpu/drm/amd/amdkfd/knod/kfd_knod.h b/drivers/gpu/drm/amd/amdkfd/knod/kfd_knod.h
new file mode 100644
index 000000000000..a1eafb13a410
--- /dev/null
+++ b/drivers/gpu/drm/amd/amdkfd/knod/kfd_knod.h
@@ -0,0 +1,270 @@
+/* SPDX-License-Identifier: GPL-2.0-or-later */
+/* Copyright (c) 2021 Taehee Yoo <ap420073@gmail.com>
+ * Copyright (c) 2021 Hoyeon Lee <hoyeon.rhee@gmail.com>
+ */
+
+#ifndef KFD_KNOD_H_
+#define KFD_KNOD_H_
+#include <linux/kfd_ioctl.h>
+#include <net/knod.h>
+#include <linux/genalloc.h>
+#include "kfd_hsa.h"
+#include <linux/completion.h>
+
+/*
+ * knod_dbg() is a pr_debug(), so it is off by default and toggled with
+ * dynamic debug; knod_err() always fires.
+ */
+#define knod_dbg(fmt, ...)						\
+	pr_debug("knod %s:%d" fmt, __func__, __LINE__, ##__VA_ARGS__)
+#define knod_err(fmt, ...)						\
+	pr_err("knod %s:%d" fmt, __func__, __LINE__, ##__VA_ARGS__)
+
+struct page_pool;
+
+struct knod_mem {
+	struct list_head list;
+	struct kgd_mem *mem;
+	void *kaddr;
+	u32 flags;
+	u32 size;
+	u32 order;
+	u64 gaddr;
+};
+
+union knod_aql_rsrc1 {
+	struct {
+#if defined(__LITTLE_ENDIAN)
+		unsigned int base_address_hi : 16;
+		unsigned int stride : 14;
+		unsigned int cache_swizzle : 1;
+		unsigned int swizzle_enable : 1;
+#elif defined(__LITTLE_ENDIAN)
+		unsigned int swizzle_enable : 1;
+		unsigned int cache_swizzle : 1;
+		unsigned int stride : 14;
+		unsigned int base_address_hi : 16;
+#endif
+	};
+};
+
+struct knod_aql {
+	struct knod_mem *aql;
+	struct knod_mem *ctx;
+	struct knod_mem *queue;
+	struct knod_mem *scratch;
+	struct knod_mem *eop;
+	struct knod_mem *queue_signal;
+	struct knod_mem *tba;
+	struct knod_mem *tma;
+	struct knod_mem *amd_queue;
+	u64 *doorbell;
+	int idx;
+};
+
+struct knod_sdma {
+	struct knod_mem *sdma;
+	struct knod_mem *queue;
+	struct knod_mem *queue_signal;
+	u64 *doorbell;
+	int idx;
+};
+
+/* KFD event handle as knod tracks it: signal event id + its slot index. */
+struct knod_event {
+	u32 id;
+	u32 slot;
+};
+
+typedef int (*knod_worker_fn_t)(void *ctx);
+typedef void (*knod_flush_fn_t)(void *ctx);
+
+enum knod_feature {
+	KNOD_FEATURE_NONE = 0,
+	KNOD_FEATURE_BPF,
+	KNOD_FEATURE_IPSEC,
+	KNOD_FEATURE_MAX,
+};
+
+/*
+ * One accel per GPU: NIC:GPU is fixed 1:1 so a single NIC owns the whole
+ * device. The accel id is just the DRM render index (stride 1). Pipeline
+ * depth within the one accel is provided by KNOD_MAX_QUEUE_CNT HW queues,
+ * unrelated to this stride.
+ */
+#define KNOD_MAX_AQL  1
+
+/* Internal AQL/SDMA queue pairs per attached knod context. BPF can keep these
+ * queues in flight independently while the public accel-id ABI remains stable.
+ */
+#define KNOD_MAX_QUEUE_CNT  32
+
+#define NR_AQL_RING 16384
+#define AQL_STRUCT_SIZE 128
+struct knod {
+	struct list_head list;
+	struct list_head active_list;
+
+	struct gen_pool *pool;
+
+	struct hsa_kernel_dispatch_packet *dp;
+	pid_t umh_pid;
+	struct task_struct *umh_task;
+	u32 nr_aql_ring;
+	/* AQLs */
+	int queue_cnt;
+	int sdma_cnt;
+	int igpu;
+	int isa_version;
+	/* NAPIs */
+	int channels;
+	struct kfd_process *process;
+	struct mm_struct *mm;
+	struct kfd_node *dev;
+	struct file *drm_file;
+	void __iomem *doorbell_base;
+	u64 reserved_addr;
+	u64 limit_addr;
+	struct mutex lock;
+	struct hsa_event *event;
+	struct knod_mem *kernels[2];	/* dispatch slots: [0] default/pass, [1] BPF alt */
+	struct knod_mem *mailbox;
+	/* packet data path buf */
+	struct knod_mem **buf;
+
+	u32 signal_eid;
+	u32 completion_eid;
+	struct knod_aql kaql[NR_CPUS];
+	struct knod_sdma sdma[NR_CPUS];
+	struct knod_event aql_event[NR_CPUS];
+	struct knod_event sdma_event[NR_CPUS];
+	u32 aql_queue_id[NR_CPUS];
+	u32 sdma_queue_id[NR_CPUS];
+	u64 aql_doorbell_offset[NR_CPUS];
+	u64 sdma_doorbell_offset[NR_CPUS];
+	bool aql_queue_created[NR_CPUS];
+	bool sdma_queue_created[NR_CPUS];
+	struct kfd_event_data *event_data;
+	struct knod_accel *accel;
+	struct dentry *debug_dir;
+	/* Worker callback - one active worker at a time */
+	enum knod_feature active_feature;
+	knod_worker_fn_t worker_fn;
+	knod_flush_fn_t flush_fn;
+	void *worker_ctx;
+	struct task_struct *worker;
+};
+
+struct knod_dispatch_params {
+	u16 workgroup_size_x;
+	u32 grid_size_x;
+	u32 grid_size_y;
+	u32 private_segment_size;
+	u32 group_segment_size;
+	u64 kernel_object;
+	u64 kernarg_address;
+};
+
+static inline void
+knod_setup_invalidate(struct knod *knod, int idx, int q_idx)
+{
+	struct hsa_kernel_dispatch_packet *dp = knod->kaql[q_idx].aql->kaddr;
+
+	dp += idx;
+	dp->header = HSA_PACKET_TYPE_INVALID << HSA_PACKET_HEADER_TYPE;
+}
+
+static inline void
+knod_setup_dispatch(struct knod *knod, int idx,
+		    const struct knod_dispatch_params *p, int q_idx)
+{
+	struct hsa_kernel_dispatch_packet *dp = knod->kaql[q_idx].aql->kaddr;
+
+	dp += idx;
+	dp->setup = 2;
+	dp->workgroup_size_x = p->workgroup_size_x;
+	dp->workgroup_size_y = 1;
+	dp->workgroup_size_z = 1;
+	dp->grid_size_x = p->grid_size_x;
+	dp->grid_size_y = p->grid_size_y;
+	dp->grid_size_z = 1;
+	dp->private_segment_size = p->private_segment_size;
+	dp->group_segment_size = p->group_segment_size;
+	dp->kernel_object = p->kernel_object;
+	dp->kernarg_address = (void *)p->kernarg_address;
+	dp->completion_signal = knod->kaql[q_idx].queue_signal->gaddr;
+	/* publish the packet body before the valid header (WRITE_ONCE below) */
+	wmb();
+	WRITE_ONCE(dp->header,
+		   (HSA_PACKET_TYPE_KERNEL_DISPATCH <<
+		    HSA_PACKET_HEADER_TYPE) |
+		   (HSA_FENCE_SCOPE_SYSTEM <<
+		    HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
+		   (HSA_FENCE_SCOPE_SYSTEM <<
+		    HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE));
+}
+
+static inline void
+knod_setup_header(struct knod *knod,
+		  const struct knod_dispatch_params *p, int q_idx)
+{
+	struct amd_queue *amd_queue = (struct amd_queue *)knod->kaql[q_idx].amd_queue->kaddr;
+	int curr_idx = knod->kaql[q_idx].idx;
+	int next_idx = curr_idx + 1;
+	u64 *ptr = knod->kaql[q_idx].doorbell;
+
+	knod_setup_invalidate(knod, next_idx % knod->nr_aql_ring, q_idx);
+	knod_setup_dispatch(knod, curr_idx % knod->nr_aql_ring, p, q_idx);
+	WRITE_ONCE(amd_queue->write_dispatch_id, curr_idx);
+	writeq(curr_idx, ptr);
+	knod->kaql[q_idx].idx = next_idx;
+}
+
+#define KNOD_NR_AQL_DEFAULT   1
+struct knod *knod_alloc_ctx(struct knod_dev *knodev, int queue_cnt, int id,
+			    int channels);
+void knod_release_ctx(struct knod *knod);
+void knod_accel_xdp_register(struct knod_accel_xdp_ops *xdp_ops);
+void knod_accel_xdp_unregister(void);
+void knod_accel_ipsec_register(struct knod_accel_ipsec_ops *ipsec_ops);
+void knod_accel_ipsec_unregister(void);
+void knod_request_queue_cnt(int n);
+struct knod_mem *knod_alloc_mem(struct knod *knod, size_t size, int flags);
+struct knod_mem *__knod_alloc_mem(struct knod *knod, size_t size, int flags);
+int __knod_map_mem(struct knod *knod, struct knod_mem *mem);
+int __knod_export_dma_buf(struct knod *knod, struct knod_mem *mem);
+int __knod_map_kaddr(struct knod *knod, struct knod_mem *mem);
+void knod_free_mem(struct knod *pknod, struct knod_mem *mem);
+void knod_sdma_copy(struct knod *knod, u64 dst_gart_addr, u64 src_gart_addr,
+		    int idx, int size);
+void knod_sdma_fence(struct knod *knod, u64 fence_addr, u32 fence_val,
+		     int idx);
+void knod_sdma_trap(struct knod *knod, int idx);
+void knod_sdma_doorbell(struct knod *knod, int idx);
+
+/* One linear GPU->host SDMA copy (GPU VM addresses). */
+struct knod_sdma_copy_desc {
+	u64 dst;
+	u64 src;
+	u32 len;
+};
+
+/*
+ * Emit @n copies on sdma[@idx] as one batch with ring backpressure.
+ * Returns the post-batch ring position to fence/await, or 0 if the ring
+ * is too full (caller drops the whole batch).  Pair with knod_sdma_kick().
+ */
+u32 knod_sdma_submit(struct knod *knod, int idx,
+		     const struct knod_sdma_copy_desc *copies, int n);
+void knod_sdma_kick(struct knod *knod, int idx);
+
+int knod_gart_map(struct amdgpu_device *adev, u64 npages,
+		  dma_addr_t *addr, u64 *gart_addr, u64 flags);
+int knod_register_worker(struct knod *knod, knod_worker_fn_t fn,
+			 knod_flush_fn_t flush, void *ctx);
+void knod_unregister_worker(struct knod *knod);
+int knod_wait_on_events(struct kfd_process *p, u32 num_events,
+			void __user *data, bool all, u32 *user_timeout_ms,
+			u32 *wait_result);
+
+#endif /* KFD_KNOD_H_ */
-- 
2.43.0