Expose the kfd internals the knod provider needs to build and run GPU queues on behalf of the kernel: kernel-side event waiting, GPUVM allocation for VRAM/GTT, doorbell access, and the HSA queue/packet layout (kfd_hsa.h). No functional change for existing user-mode queue users; the provider itself is added in a following patch.
Signed-off-by: Taehee Yoo <[email protected]> (cherry picked from commit 7d2c7cbfbb2a4a80052b5793841f6a229271973d) --- drivers/gpu/drm/amd/amdgpu/amdgpu_amdkfd.h | 13 +- .../gpu/drm/amd/amdgpu/amdgpu_amdkfd_gpuvm.c | 65 +++ drivers/gpu/drm/amd/amdgpu/amdgpu_drv.c | 3 + drivers/gpu/drm/amd/amdkfd/kfd_chardev.c | 117 ++--- drivers/gpu/drm/amd/amdkfd/kfd_doorbell.c | 20 +- drivers/gpu/drm/amd/amdkfd/kfd_events.c | 112 ++++- drivers/gpu/drm/amd/amdkfd/kfd_events.h | 4 + drivers/gpu/drm/amd/amdkfd/kfd_hsa.h | 451 ++++++++++++++++++ drivers/gpu/drm/amd/amdkfd/kfd_module.c | 1 + drivers/gpu/drm/amd/amdkfd/kfd_priv.h | 34 ++ drivers/gpu/drm/amd/amdkfd/kfd_process.c | 26 +- 11 files changed, 778 insertions(+), 68 deletions(-) create mode 100644 drivers/gpu/drm/amd/amdkfd/kfd_hsa.h diff --git a/drivers/gpu/drm/amd/amdgpu/amdgpu_amdkfd.h b/drivers/gpu/drm/amd/amdgpu/amdgpu_amdkfd.h index e443a7277299..cf906b32eacb 100644 --- a/drivers/gpu/drm/amd/amdgpu/amdgpu_amdkfd.h +++ b/drivers/gpu/drm/amd/amdgpu/amdgpu_amdkfd.h @@ -328,6 +328,8 @@ int amdgpu_amdkfd_gpuvm_sync_memory( struct amdgpu_device *adev, struct kgd_mem *mem, bool intr); int amdgpu_amdkfd_gpuvm_map_gtt_bo_to_kernel(struct kgd_mem *mem, void **kptr, uint64_t *size); +int amdgpu_amdkfd_gpuvm_map_vram_bo_to_kernel(struct kgd_mem *mem, + void **kptr, uint64_t *size); void amdgpu_amdkfd_gpuvm_unmap_gtt_bo_from_kernel(struct kgd_mem *mem); int amdgpu_amdkfd_map_gtt_bo_to_gart(struct amdgpu_bo *bo, struct amdgpu_bo **bo_gart); @@ -446,7 +448,6 @@ bool kgd2kfd_vmfault_fast_path(struct amdgpu_device *adev, struct amdgpu_iv_entr bool retry_fault); void kgd2kfd_lock_kfd(void); void kgd2kfd_teardown_processes(struct amdgpu_device *adev); - #else static inline int kgd2kfd_init(void) { @@ -576,5 +577,15 @@ static inline void kgd2kfd_teardown_processes(struct amdgpu_device *adev) { } +#endif + +#if defined(CONFIG_HSA_AMD_KNOD) +int knod_init(struct amdgpu_device *adev); +void knod_fini(struct amdgpu_device *adev); +void knod_exit(void); +#else +static inline int knod_init(struct amdgpu_device *adev) { return 0; } +static inline void knod_fini(struct amdgpu_device *adev) { } +static inline void knod_exit(void) { } #endif #endif /* AMDGPU_AMDKFD_H_INCLUDED */ diff --git a/drivers/gpu/drm/amd/amdgpu/amdgpu_amdkfd_gpuvm.c b/drivers/gpu/drm/amd/amdgpu/amdgpu_amdkfd_gpuvm.c index 35fe2c974699..79fd0ce7fc24 100644 --- a/drivers/gpu/drm/amd/amdgpu/amdgpu_amdkfd_gpuvm.c +++ b/drivers/gpu/drm/amd/amdgpu/amdgpu_amdkfd_gpuvm.c @@ -2336,6 +2336,71 @@ int amdgpu_amdkfd_gpuvm_map_gtt_bo_to_kernel(struct kgd_mem *mem, return ret; } +/** amdgpu_amdkfd_gpuvm_map_vram_bo_to_kernel() - Map a VRAM BO for kernel CPU access + * + * @mem: Buffer object to be mapped for CPU access + * @kptr[out]: pointer in kernel CPU address space + * @size[out]: size of the buffer + * + * Pins the BO and maps it for kernel CPU access. The eviction fence is removed + * from the BO, since pinned BOs cannot be evicted. The bo must remain on the + * validate_list, so the GPU mapping can be restored after a page table was + * evicted. + * + * Return: 0 on success, error code on failure + */ +int amdgpu_amdkfd_gpuvm_map_vram_bo_to_kernel(struct kgd_mem *mem, + void **kptr, uint64_t *size) +{ + int ret; + struct amdgpu_bo *bo = mem->bo; + + if (amdgpu_ttm_tt_get_usermm(bo->tbo.ttm)) { + pr_err("userptr can't be mapped to kernel\n"); + return -EINVAL; + } + + mutex_lock(&mem->process_info->lock); + + ret = amdgpu_bo_reserve(bo, true); + if (ret) { + pr_err("Failed to reserve bo. ret %d\n", ret); + goto bo_reserve_failed; + } + + ret = amdgpu_bo_pin(bo, AMDGPU_GEM_DOMAIN_VRAM); + if (ret) { + pr_err("Failed to pin bo. ret %d\n", ret); + goto pin_failed; + } + + ret = amdgpu_bo_kmap(bo, kptr); + if (ret) { + pr_err("Failed to map bo to kernel. ret %d\n", ret); + goto kmap_failed; + } + + amdgpu_amdkfd_remove_eviction_fence( + bo, mem->process_info->eviction_fence); + + if (size) + *size = amdgpu_bo_size(bo); + + amdgpu_bo_unreserve(bo); + + mutex_unlock(&mem->process_info->lock); + return 0; + +kmap_failed: + amdgpu_bo_unpin(bo); +pin_failed: + amdgpu_bo_unreserve(bo); +bo_reserve_failed: + mutex_unlock(&mem->process_info->lock); + + return ret; +} + /** amdgpu_amdkfd_gpuvm_map_gtt_bo_to_kernel() - Unmap a GTT BO for kernel CPU access * * @mem: Buffer object to be unmapped for CPU access diff --git a/drivers/gpu/drm/amd/amdgpu/amdgpu_drv.c b/drivers/gpu/drm/amd/amdgpu/amdgpu_drv.c index 4c0c77eafbd1..aa974929d22a 100644 --- a/drivers/gpu/drm/amd/amdgpu/amdgpu_drv.c +++ b/drivers/gpu/drm/amd/amdgpu/amdgpu_drv.c @@ -2480,6 +2480,8 @@ static int amdgpu_pci_probe(struct pci_dev *pdev, drm_client_setup(adev_to_drm(adev), format); } + knod_init(adev); + ret = amdgpu_debugfs_init(adev); if (ret) DRM_ERROR("Creating debugfs files failed (%d).\n", ret); @@ -2540,6 +2542,7 @@ amdgpu_pci_remove(struct pci_dev *pdev) struct drm_device *dev = pci_get_drvdata(pdev); struct amdgpu_device *adev = drm_to_adev(dev); + knod_fini(adev); amdgpu_ras_eeprom_check_and_recover(adev); amdgpu_xcp_dev_unplug(adev); amdgpu_gmc_prepare_nps_mode_change(adev); diff --git a/drivers/gpu/drm/amd/amdkfd/kfd_chardev.c b/drivers/gpu/drm/amd/amdkfd/kfd_chardev.c index c7edebd2fd8a..7879a6ffaa76 100644 --- a/drivers/gpu/drm/amd/amdkfd/kfd_chardev.c +++ b/drivers/gpu/drm/amd/amdkfd/kfd_chardev.c @@ -335,32 +335,27 @@ static int set_queue_properties_from_user(struct queue_properties *q_properties, return 0; } -static int kfd_ioctl_create_queue(struct file *filep, struct kfd_process *p, - void *data) +/* + * Create a queue from a kernel-filled queue_properties. @q_properties is + * owned by the caller and must already be populated (the ioctl handler does + * this from user args via set_queue_properties_from_user(); in-kernel callers + * fill it directly). On success *@queue_id_out and *@doorbell_offset_out are + * set. + */ +int kfd_create_queue(struct kfd_process *p, struct queue_properties *q_properties, + u32 gpu_id, u32 *queue_id_out, u64 *doorbell_offset_out) { - struct kfd_ioctl_create_queue_args *args = data; struct kfd_node *dev; int err = 0; unsigned int queue_id; struct kfd_process_device *pdd; - struct queue_properties q_properties; uint32_t doorbell_offset_in_process = 0; - memset(&q_properties, 0, sizeof(struct queue_properties)); - - pr_debug("Creating queue ioctl\n"); - - err = set_queue_properties_from_user(&q_properties, args); - if (err) - return err; - - pr_debug("Looking for gpu id 0x%x\n", args->gpu_id); - mutex_lock(&p->mutex); - pdd = kfd_process_device_data_by_id(p, args->gpu_id); + pdd = kfd_process_device_data_by_id(p, gpu_id); if (!pdd) { - pr_debug("Could not find gpu id 0x%x\n", args->gpu_id); + pr_debug("Could not find gpu id 0x%x\n", gpu_id); err = -EINVAL; goto err_pdd; } @@ -372,14 +367,14 @@ static int kfd_ioctl_create_queue(struct file *filep, struct kfd_process *p, goto err_bind_process; } - if (q_properties.type == KFD_QUEUE_TYPE_SDMA_BY_ENG_ID) { + if (q_properties->type == KFD_QUEUE_TYPE_SDMA_BY_ENG_ID) { int max_sdma_eng_id = kfd_get_num_sdma_engines(dev) + kfd_get_num_xgmi_sdma_engines(dev) - 1; - if (q_properties.sdma_engine_id > max_sdma_eng_id) { + if (q_properties->sdma_engine_id > max_sdma_eng_id) { err = -EINVAL; pr_err("sdma_engine_id %i exceeds maximum id of %i\n", - q_properties.sdma_engine_id, max_sdma_eng_id); + q_properties->sdma_engine_id, max_sdma_eng_id); goto err_sdma_engine_id; } } @@ -392,7 +387,7 @@ static int kfd_ioctl_create_queue(struct file *filep, struct kfd_process *p, } } - err = kfd_queue_acquire_buffers(pdd, &q_properties); + err = kfd_queue_acquire_buffers(pdd, q_properties); if (err) { pr_debug("failed to acquire user queue buffers\n"); goto err_acquire_queue_buf; @@ -402,42 +397,32 @@ static int kfd_ioctl_create_queue(struct file *filep, struct kfd_process *p, p->lead_thread->pid, dev->id); - err = pqm_create_queue(&p->pqm, dev, &q_properties, &queue_id, + err = pqm_create_queue(&p->pqm, dev, q_properties, &queue_id, NULL, NULL, NULL, &doorbell_offset_in_process); if (err != 0) goto err_create_queue; - args->queue_id = queue_id; - + *queue_id_out = queue_id; /* Return gpu_id as doorbell offset for mmap usage */ - args->doorbell_offset = KFD_MMAP_TYPE_DOORBELL; - args->doorbell_offset |= KFD_MMAP_GPU_ID(args->gpu_id); + *doorbell_offset_out = KFD_MMAP_TYPE_DOORBELL; + *doorbell_offset_out |= KFD_MMAP_GPU_ID(gpu_id); if (KFD_IS_SOC15(dev)) /* On SOC15 ASICs, include the doorbell offset within the * process doorbell frame, which is 2 pages. */ - args->doorbell_offset |= doorbell_offset_in_process; + *doorbell_offset_out |= doorbell_offset_in_process; mutex_unlock(&p->mutex); - pr_debug("Queue id %d was created successfully\n", args->queue_id); - - pr_debug("Ring buffer address == 0x%016llX\n", - args->ring_base_address); - - pr_debug("Read ptr address == 0x%016llX\n", - args->read_pointer_address); - - pr_debug("Write ptr address == 0x%016llX\n", - args->write_pointer_address); + pr_debug("Queue id %d was created successfully\n", queue_id); kfd_dbg_ev_raise(KFD_EC_MASK(EC_QUEUE_NEW), p, dev, queue_id, false, NULL, 0); return 0; err_create_queue: - kfd_queue_unref_bo_vas(pdd, &q_properties); - kfd_queue_release_buffers(pdd, &q_properties); + kfd_queue_unref_bo_vas(pdd, q_properties); + kfd_queue_release_buffers(pdd, q_properties); err_acquire_queue_buf: err_sdma_engine_id: err_bind_process: @@ -446,6 +431,23 @@ static int kfd_ioctl_create_queue(struct file *filep, struct kfd_process *p, return err; } +static int kfd_ioctl_create_queue(struct file *filep, struct kfd_process *p, + void *data) +{ + struct kfd_ioctl_create_queue_args *args = data; + struct queue_properties q_properties; + int err; + + memset(&q_properties, 0, sizeof(q_properties)); + + err = set_queue_properties_from_user(&q_properties, args); + if (err) + return err; + + return kfd_create_queue(p, &q_properties, args->gpu_id, + &args->queue_id, &args->doorbell_offset); +} + static int kfd_ioctl_destroy_queue(struct file *filp, struct kfd_process *p, void *data) { @@ -643,7 +645,7 @@ static int kfd_ioctl_set_memory_policy(struct file *filep, return err; } -static int kfd_ioctl_set_trap_handler(struct file *filep, +int kfd_ioctl_set_trap_handler(struct file *filep, struct kfd_process *p, void *data) { struct kfd_ioctl_set_trap_handler_args *args = data; @@ -726,7 +728,7 @@ static int kfd_ioctl_get_clock_counters(struct file *filep, static int kfd_ioctl_get_process_apertures(struct file *filp, - struct kfd_process *p, void *data) + struct kfd_process *p, void *data) { struct kfd_ioctl_get_process_apertures_args *args = data; struct kfd_process_device_apertures *pAperture; @@ -859,7 +861,7 @@ static int kfd_ioctl_get_process_apertures_new(struct file *filp, } static int kfd_ioctl_create_event(struct file *filp, struct kfd_process *p, - void *data) + void *data) { struct kfd_ioctl_create_event_args *args = data; int err; @@ -894,24 +896,24 @@ static int kfd_ioctl_destroy_event(struct file *filp, struct kfd_process *p, return kfd_event_destroy(p, args->event_id); } -static int kfd_ioctl_set_event(struct file *filp, struct kfd_process *p, - void *data) +int kfd_ioctl_set_event(struct file *filp, struct kfd_process *p, + void *data) { struct kfd_ioctl_set_event_args *args = data; return kfd_set_event(p, args->event_id); } -static int kfd_ioctl_reset_event(struct file *filp, struct kfd_process *p, - void *data) +int kfd_ioctl_reset_event(struct file *filp, struct kfd_process *p, + void *data) { struct kfd_ioctl_reset_event_args *args = data; return kfd_reset_event(p, args->event_id); } -static int kfd_ioctl_wait_events(struct file *filp, struct kfd_process *p, - void *data) +int kfd_ioctl_wait_events(struct file *filp, struct kfd_process *p, + void *data) { struct kfd_ioctl_wait_events_args *args = data; @@ -920,8 +922,9 @@ static int kfd_ioctl_wait_events(struct file *filp, struct kfd_process *p, (args->wait_for_all != 0), &args->timeout, &args->wait_result); } -static int kfd_ioctl_set_scratch_backing_va(struct file *filep, - struct kfd_process *p, void *data) + +int kfd_ioctl_set_scratch_backing_va(struct file *filep, + struct kfd_process *p, void *data) { struct kfd_ioctl_set_scratch_backing_va_args *args = data; struct kfd_process_device *pdd; @@ -1003,8 +1006,8 @@ static int kfd_ioctl_get_tile_config(struct file *filep, return 0; } -static int kfd_ioctl_acquire_vm(struct file *filep, struct kfd_process *p, - void *data) +int kfd_ioctl_acquire_vm(struct file *filep, struct kfd_process *p, + void *data) { struct kfd_ioctl_acquire_vm_args *args = data; struct kfd_process_device *pdd; @@ -1078,8 +1081,8 @@ static int kfd_ioctl_get_available_memory(struct file *filep, return 0; } -static int kfd_ioctl_alloc_memory_of_gpu(struct file *filep, - struct kfd_process *p, void *data) +int kfd_ioctl_alloc_memory_of_gpu(struct file *filep, + struct kfd_process *p, void *data) { struct kfd_ioctl_alloc_memory_of_gpu_args *args = data; struct kfd_process_device *pdd; @@ -1279,8 +1282,8 @@ static int kfd_ioctl_free_memory_of_gpu(struct file *filep, return ret; } -static int kfd_ioctl_map_memory_to_gpu(struct file *filep, - struct kfd_process *p, void *data) +int kfd_ioctl_map_memory_to_gpu(struct file *filep, + struct kfd_process *p, void *data) { struct kfd_ioctl_map_memory_to_gpu_args *args = data; struct kfd_process_device *pdd, *peer_pdd; @@ -1624,8 +1627,8 @@ static int kfd_ioctl_import_dmabuf(struct file *filep, return r; } -static int kfd_ioctl_export_dmabuf(struct file *filep, - struct kfd_process *p, void *data) +int kfd_ioctl_export_dmabuf(struct file *filep, + struct kfd_process *p, void *data) { struct kfd_ioctl_export_dmabuf_args *args = data; struct kfd_process_device *pdd; diff --git a/drivers/gpu/drm/amd/amdkfd/kfd_doorbell.c b/drivers/gpu/drm/amd/amdkfd/kfd_doorbell.c index fdcf7f2d1b5b..78b7e956f590 100644 --- a/drivers/gpu/drm/amd/amdkfd/kfd_doorbell.c +++ b/drivers/gpu/drm/amd/amdkfd/kfd_doorbell.c @@ -145,10 +145,28 @@ int kfd_doorbell_mmap(struct kfd_node *dev, struct kfd_process *process, vma->vm_page_prot); } +void __iomem *kfd_kernel_doorbell_mmap(struct kfd_node *dev, + struct kfd_process *process) +{ + phys_addr_t address; + struct kfd_process_device *pdd; + + pdd = kfd_get_process_device_data(dev, process); + if (!pdd) + return NULL; + + /* Calculate physical address of doorbell */ + address = kfd_get_process_doorbells(pdd); + if (!address) + return NULL; + + return ioremap(address, kfd_doorbell_process_slice(dev->kfd)); +} + /* get kernel iomem pointer for a doorbell */ void __iomem *kfd_get_kernel_doorbell(struct kfd_dev *kfd, - unsigned int *doorbell_off) + unsigned int *doorbell_off) { u32 inx; diff --git a/drivers/gpu/drm/amd/amdkfd/kfd_events.c b/drivers/gpu/drm/amd/amdkfd/kfd_events.c index 81900b49d9d5..91ef4c61a8bc 100644 --- a/drivers/gpu/drm/amd/amdkfd/kfd_events.c +++ b/drivers/gpu/drm/amd/amdkfd/kfd_events.c @@ -326,7 +326,7 @@ static bool event_can_be_cpu_signaled(const struct kfd_event *ev) return ev->type == KFD_EVENT_TYPE_SIGNAL; } -static int kfd_event_page_set(struct kfd_process *p, void *kernel_address, +int kfd_event_page_set(struct kfd_process *p, void *kernel_address, uint64_t size, uint64_t user_handle) { struct kfd_signal_page *page; @@ -1068,6 +1068,116 @@ int kfd_wait_on_events(struct kfd_process *p, return ret; } +int kfd_wait_on_events_kernel(struct kfd_process *p, + uint32_t num_events, void __user *data, + bool all, uint32_t *user_timeout_ms, + uint32_t *wait_result) +{ + struct kfd_event_data *events = (struct kfd_event_data *) data; + uint32_t i; + int ret = 0; + + struct kfd_event_waiter *event_waiters = NULL; + long timeout = user_timeout_to_jiffies(*user_timeout_ms); + + event_waiters = alloc_event_waiters(num_events); + if (!event_waiters) { + ret = -ENOMEM; + goto out; + } + + /* Use p->event_mutex here to protect against concurrent creation and + * destruction of events while we initialize event_waiters. + */ + mutex_lock(&p->event_mutex); + + for (i = 0; i < num_events; i++) { + struct kfd_event_data event_data; + + memcpy(&event_data, &events[i], sizeof(struct kfd_event_data)); + ret = init_event_waiter(p, &event_waiters[i], &event_data); + if (ret) + goto out_unlock; + } + + /* Check condition once. */ + *wait_result = test_event_condition(all, num_events, event_waiters); + if (*wait_result == KFD_IOC_WAIT_RESULT_COMPLETE) { + ret = copy_signaled_event_data(num_events, + event_waiters, events); + goto out_unlock; + } else if (WARN_ON(*wait_result == KFD_IOC_WAIT_RESULT_FAIL)) { + /* This should not happen. Events shouldn't be + * destroyed while we're holding the event_mutex + */ + goto out_unlock; + } + + mutex_unlock(&p->event_mutex); + + while (true) { + if (fatal_signal_pending(current)) { + ret = -EINTR; + break; + } + + if (signal_pending(current)) { + ret = -ERESTARTSYS; + if (*user_timeout_ms != KFD_EVENT_TIMEOUT_IMMEDIATE && + *user_timeout_ms != KFD_EVENT_TIMEOUT_INFINITE) + *user_timeout_ms = jiffies_to_msecs( + max(0l, timeout-1)); + break; + } + + /* Set task state to interruptible sleep before + * checking wake-up conditions. A concurrent wake-up + * will put the task back into runnable state. In that + * case schedule_timeout will not put the task to + * sleep and we'll get a chance to re-check the + * updated conditions almost immediately. Otherwise, + * this race condition would lead to a soft hang or a + * very long sleep. + */ + set_current_state(TASK_INTERRUPTIBLE); + + *wait_result = test_event_condition(all, num_events, + event_waiters); + if (*wait_result != KFD_IOC_WAIT_RESULT_TIMEOUT) + break; + + if (timeout <= 0) + break; + + timeout = schedule_timeout(timeout); + } + __set_current_state(TASK_RUNNING); + + mutex_lock(&p->event_mutex); + /* copy_signaled_event_data may sleep. So this has to happen + * after the task state is set back to RUNNING. + * + * The event may also have been destroyed after signaling. So + * copy_signaled_event_data also must confirm that the event + * still exists. Therefore this must be under the p->event_mutex + * which is also held when events are destroyed. + */ + if (!ret && *wait_result == KFD_IOC_WAIT_RESULT_COMPLETE) + ret = copy_signaled_event_data(num_events, + event_waiters, events); + +out_unlock: + free_waiters(num_events, event_waiters, ret == -ERESTARTSYS); + mutex_unlock(&p->event_mutex); +out: + if (ret) + *wait_result = KFD_IOC_WAIT_RESULT_FAIL; + else if (*wait_result == KFD_IOC_WAIT_RESULT_FAIL) + ret = -EIO; + + return ret; +} + int kfd_event_mmap(struct kfd_process *p, struct vm_area_struct *vma) { unsigned long pfn; diff --git a/drivers/gpu/drm/amd/amdkfd/kfd_events.h b/drivers/gpu/drm/amd/amdkfd/kfd_events.h index 1dc21c13833b..aa6ead122084 100644 --- a/drivers/gpu/drm/amd/amdkfd/kfd_events.h +++ b/drivers/gpu/drm/amd/amdkfd/kfd_events.h @@ -87,5 +87,9 @@ struct kfd_event { extern void kfd_signal_event_interrupt(u32 pasid, uint32_t partial_id, uint32_t valid_id_bits, bool signal_mailbox_updated); +int kfd_wait_on_events_kernel(struct kfd_process *p, + uint32_t num_events, void __user *data, + bool all, uint32_t *user_timeout_ms, + uint32_t *wait_result); #endif diff --git a/drivers/gpu/drm/amd/amdkfd/kfd_hsa.h b/drivers/gpu/drm/amd/amdkfd/kfd_hsa.h new file mode 100644 index 000000000000..794688c615f0 --- /dev/null +++ b/drivers/gpu/drm/amd/amdkfd/kfd_hsa.h @@ -0,0 +1,451 @@ +/* SPDX-License-Identifier: GPL-2.0-or-later */ +#ifndef KFD_HSA_CODE_H_ +#define KFD_HSA_CODE_H_ + +struct amd_queue_properties { + u32 enable_trap_handler:1, + is_ptr64:1, + enable_trap_handler_debug_sgprs:1, + enable_profiling:1, + use_scratch_once:1, + reserved1:27; +}; + +enum hsa_packet_type { + /* + * Vendor-specific packet. + */ + HSA_PACKET_TYPE_VENDOR_SPECIFIC = 0, + /* + * The packet has been processed in the past, but has not been reassigned to + * the packet processor. A packet processor must not process a packet of this + * type. All queues support this packet type. + */ + HSA_PACKET_TYPE_INVALID = 1, + /* + * Packet used by agents for dispatching jobs to kernel agents. Not all + * queues support packets of this type (see ::hsa_queue_feature_t). + */ + HSA_PACKET_TYPE_KERNEL_DISPATCH = 2, + /* + * Packet used by agents to delay processing of subsequent packets, and to + * express complex dependencies between multiple packets. All queues support + * this packet type. + */ + HSA_PACKET_TYPE_BARRIER_AND = 3, + /* + * Packet used by agents for dispatching jobs to agents. Not all + * queues support packets of this type (see ::hsa_queue_feature_t). + */ + HSA_PACKET_TYPE_AGENT_DISPATCH = 4, + /* + * Packet used by agents to delay processing of subsequent packets, and to + * express complex dependencies between multiple packets. All queues support + * this packet type. + */ + HSA_PACKET_TYPE_BARRIER_OR = 5 +}; + +enum hsa_packet_header { + /* + * Packet type. The value of this sub-field must be one of + * ::hsa_packet_type_t. If the type is ::HSA_PACKET_TYPE_VENDOR_SPECIFIC, the + * packet layout is vendor-specific. + */ + HSA_PACKET_HEADER_TYPE = 0, + /* + * Barrier bit. If the barrier bit is set, the processing of the current + * packet only launches when all preceding packets (within the same queue) are + * complete. + */ + HSA_PACKET_HEADER_BARRIER = 8, + /* + * Acquire fence scope. The value of this sub-field determines the scope and + * type of the memory fence operation applied before the packet enters the + * active phase. An acquire fence ensures that any subsequent global segment + * or image loads by any unit of execution that belongs to a dispatch that has + * not yet entered the active phase on any queue of the same kernel agent, + * sees any data previously released at the scopes specified by the acquire + * fence. The value of this sub-field must be one of ::hsa_fence_scope_t. + */ + HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE = 9, + /* + * Release fence scope, The value of this sub-field determines the scope and + * type of the memory fence operation applied after kernel completion but + * before the packet is completed. A release fence makes any global segment or + * image data that was stored by any unit of execution that belonged to a + * dispatch that has completed the active phase on any queue of the same + * kernel agent visible in all the scopes specified by the release fence. The + * value of this sub-field must be one of ::hsa_fence_scope_t. + */ + HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE = 11 +}; + +struct code_properties { + /* 4-sgprs */ + u16 enable_sgpr_private_segment_buffer:1, + /* 2-sgprs */ + enable_sgpr_dispatch_ptr:1, + /* 2-sgprs */ + enable_sgpr_queue_ptr:1, + /* 2-sgprs */ + enable_sgpr_kernarg_segment_ptr:1, + /* 2-sgprs */ + enable_sgpr_dispatch_id:1, + /* 2-sgprs */ + enable_sgpr_flat_scratch_init:1, + /* 2-sgprs */ + enable_sgpr_private_segment_size:1, + reserved0:3, + enable_wavefront_size32:1, // gfx10+ + uses_dynamic_stack:1, + reserved1:4; +}; + +struct compute_pgm_rsrc1 { + u32 granulated_workitem_vgpr_count:6, + granulated_wavefront_sgpr_count:4, + priority:2, + float_round_mode_32:2, + float_round_mode_16_64:2, + float_denorm_mode_32:2, + float_denorm_mode_16_64:2, + priv:1, + enable_dx10_clamp:1, + debug_mode:1, + enable_ieee_mode:1, + bulky:1, + cdbg_user:1, + fp16_ovfl:1, /* gfx9+ */ + reserved0:2, + wgp_mode:1, /* gfx10+ */ + mem_ordered:1, /* gfx10+ */ + fwd_progress:1; /* gfx10+ */ +}; + +struct compute_pgm_rsrc2 { + /* 1-sgprs + * enable_sgpr_private_segment_wavefront_offset + */ + u32 enable_private_segment:1, + /* More or equal to enable_sgpr_*. + * Exceed 16 will be ignored. + */ + user_sgpr_count:5, + enable_trap_handler:1, + /* 1-sgpr */ + enable_sgpr_workgroup_id_x:1, + /* 1-sgpr */ + enable_sgpr_workgroup_id_y:1, + /* 1-sgpr */ + enable_sgpr_workgroup_id_z:1, + /* 1-sgpr */ + enable_sgpr_workgroup_info:1, + enable_vgpr_workitem_id:2, + enable_exception_address_watch:1, + enable_exception_memory:1, + granulated_lds_size:9, + enable_exception_ieee_754_fp_invalid_operation:1, + enable_exception_fp_denormal_source:1, + enable_exception_ieee_754_fp_division_by_zero:1, + enable_exception_ieee_754_fp_overflow:1, + enable_exception_ieee_754_fp_underflow:1, + enable_exception_ieee_754_fp_inexact:1, + enable_exception_int_divide_by_zero:1, + reserved0:1; +}; + +struct compute_pgm_rsrc3 { + u32 accum_offset:6, + reserved0:10, + tg_split:1, + reserved1:15; +}; + +struct kernel_descriptor { + u32 group_segment_fixed_size; + u32 private_segment_fixed_size; + u32 kernarg_size; + u8 reserved0[4]; + s64 kernel_code_entry_byte_offset; + u8 reserved1[20]; + struct compute_pgm_rsrc3 compute_pgm_rsrc3; /* GFX10+ and GFX90A+ */ + struct compute_pgm_rsrc1 compute_pgm_rsrc1; + struct compute_pgm_rsrc2 compute_pgm_rsrc2; + struct code_properties code_properties; + u8 reserved2[6]; +}; + +enum hsa_queue_type { + HSA_QUEUE_TYPE_MULTI = 0, + HSA_QUEUE_TYPE_SINGLE = 1 +}; + +enum hsa_queue_feature { + HSA_QUEUE_FEATURE_KERNEL_DISPATCH = 1, + HSA_QUEUE_FEATURE_AGENT_DISPATCH = 2 +}; + +struct hsa_kernel_dispatch_packet { + u16 header; + u16 setup; + u16 workgroup_size_x; + u16 workgroup_size_y; + u16 workgroup_size_z; + u16 reserved0; + u32 grid_size_x; + u32 grid_size_y; + u32 grid_size_z; + u32 private_segment_size; + u32 group_segment_size; + u64 kernel_object; + void *kernarg_address; + u64 reserved2; + u64 completion_signal; +}; + +enum amd_signal_kind { + AMD_SIGNAL_KIND_INVALID = 0, + AMD_SIGNAL_KIND_USER = 1, + AMD_SIGNAL_KIND_DOORBELL = -1, + AMD_SIGNAL_KIND_LEGACY_DOORBELL = -2 +}; + +/* An AMD Signal object must always be 64 byte aligned to ensure it cannot + * span a page boundary. This is required by CP microcode which optimizes + * access to the structure by only doing a single SUA (System Uniform Address) + * translation when accessing signal fields. This optimization is used in GFX8. + */ +struct amd_signal { + u64 kind; + union { + volatile s64 value; + volatile u32 *legacy_hardware_doorbell_ptr; + volatile u64 *hardware_doorbell_ptr; + }; + /* For AMD_SIGNAL_KIND_USER: mailbox address for event notification + * in Signal operations. + */ + u64 event_mailbox_ptr; + /* For AMD_SIGNAL_KIND_USER: event id for event notification in Signal + * operations. + */ + u32 event_id; + u32 reserved1; + /* Start of the AQL packet timestamp, when profiled. */ + u64 start_ts; + /* End of the AQL packet timestamp, when profiled. */ + u64 end_ts; + union { + /* For AMD_SIGNAL_KIND_*DOORBELL: the address of the associated + * amd_queue, otherwise reserved and must be 0. + */ + void *queue_ptr; + u64 reserved2; + }; + u32 reserved3[2]; +} __aligned(64); + +struct hsa_queue { + u32 type; + + u32 features; + + void *base_address; + /* + * Signal object used by the application to indicate the ID of a packet that + * is ready to be processed. The HSA runtime manages the doorbell signal. If + * the application tries to replace or destroy this signal, the behavior is + * undefined. + * + * If @a type is ::HSA_QUEUE_TYPE_SINGLE, the doorbell signal value must be + * updated in a monotonically increasing fashion. If @a type is + * ::HSA_QUEUE_TYPE_MULTI, the doorbell signal value can be updated with any + * value. + */ + u64 doorbell_signal; + + /* + * Maximum number of packets the queue can hold. Must be a power of 2. + */ + u32 size; + /* Reserved. Must be 0. */ + u32 reserved1; + /* + * Queue identifier, which is unique over the lifetime of the application. + */ + u64 id; + +}; + +enum hsa_fence_scope { + HSA_FENCE_SCOPE_NONE = 0, + HSA_FENCE_SCOPE_AGENT = 1, + HSA_FENCE_SCOPE_SYSTEM = 2 +}; + +struct amd_queue { + struct hsa_queue hsa_queue; + u32 reserved1[4]; + volatile u64 write_dispatch_id; + u32 group_segment_aperture_base_hi; + u32 private_segment_aperture_base_hi; + u32 max_cu_id; + u32 max_wave_id; + volatile u64 max_legacy_doorbell_dispatch_id_plus_1; + volatile u32 legacy_doorbell_lock; + u32 reserved2[9]; + volatile u64 read_dispatch_id; + u32 read_dispatch_id_field_base_byte_offset; + u32 compute_tmpring_size; + u32 scratch_resource_descriptor[4]; + u64 scratch_backing_memory_location; + u64 scratch_backing_memory_byte_size; + u32 scratch_wave64_lane_byte_size; + struct amd_queue_properties queue_properties; + u32 reserved3[2]; + u64 queue_inactive_signal; + u32 reserved4[14]; +} __aligned(64); + +struct hsa_sync_var { + union { + /* pointer to user mode data */ + void *user_data; + /* 64bit compatibility of value */ + u64 user_data_ptr_value; + }; + u64 SyncVarSize; +}; + +enum hsa_event_type { + /* user-mode generated GPU signal */ + HSA_EVENTTYPE_SIGNAL = 0, + /* HSA node change (attach/detach) */ + HSA_EVENTTYPE_NODECHANGE = 1, + /* HSA device state change( start/stop ) */ + HSA_EVENTTYPE_DEVICESTATECHANGE = 2, + /* GPU shader exception event */ + HSA_EVENTTYPE_HW_EXCEPTION = 3, + /* GPU SYSCALL with parameter info */ + HSA_EVENTTYPE_SYSTEM_EVENT = 4, + /* GPU signal for debugging */ + HSA_EVENTTYPE_DEBUG_EVENT = 5, + /* GPU signal for profiling */ + HSA_EVENTTYPE_PROFILE_EVENT = 6, + /* GPU signal queue idle state (EOP pm4) */ + HSA_EVENTTYPE_QUEUE_EVENT = 7, + /* GPU signal for signaling memory access faults and memory subsystem issues */ + HSA_EVENTTYPE_MEMORY = 8, + /* ... */ + HSA_EVENTTYPE_MAXID, + HSA_EVENTTYPE_TYPE_SIZE = 0xffffffff +}; + +enum hsa_eventtype_nodechange_flags { + HSA_EVENTTYPE_NODECHANGE_ADD = 0, + HSA_EVENTTYPE_NODECHANGE_REMOVE = 1, + HSA_EVENTTYPE_NODECHANGE_SIZE = 0xffffffff +}; + +struct hsa_node_change { + /* HSA node added/removed on the platform */ + enum hsa_eventtype_nodechange_flags flags; +}; + +enum hsa_device { + HSA_DEVICE_CPU = 0, + HSA_DEVICE_GPU = 1, + MAX_HSA_DEVICE = 2 +}; + +enum hsa_eventtype_devicestatechange_flags { + /* device started (and available) */ + HSA_EVENTTYPE_DEVICESTATUSCHANGE_START = 0, + /* device stopped (i.e. unavailable) */ + HSA_EVENTTYPE_DEVICESTATUSCHANGE_STOP = 1, + HSA_EVENTTYPE_DEVICESTATUSCHANGE_SIZE = 0xffffffff +}; + +struct hsa_device_state_change { + /* F-NUMA node that contains the device */ + u32 node_id; + /* device type: GPU or CPU */ + enum hsa_device device; + /* event flags */ + enum hsa_eventtype_devicestatechange_flags flags; +}; + +struct hsa_access_attribute_failure { + /* Page not present or supervisor privilege */ + unsigned int not_present:1; + /* Write access to a read-only page */ + unsigned int readonly:1; + /* Execute access to a page marked NX */ + unsigned int no_execute:1; + /* Host access only */ + unsigned int gpu_access:1; + /* RAS ECC failure (notification of DRAM ECC - non-recoverable - + * error, if supported by HW) + */ + unsigned int ecc:1; + /* Can't determine the exact fault address */ + unsigned int imprecise:1; + /* Indicates RAS errors or other errors causing the access to GPU to fail + * 0 = no RAS error, + * 1 = ECC_SRAM, + * 2 = Link_SYNFLOOD (poison), + * 3 = GPU hang (not attributable to a specific cause), other values reserved + */ + unsigned int error_type:3; + /* must be 0 */ + unsigned int Reserved:23; +}; + +enum hsa_eventid_memory_flags { + /* access fault, recoverable after page adjustment */ + HSA_EVENTID_MEMORY_RECOVERABLE = 0, + /* memory access requires process context destruction, unrecoverable */ + HSA_EVENTID_MEMORY_FATAL_PROCESS = 1, + /* memory access requires all GPU VA context destruction, unrecoverable */ + HSA_EVENTID_MEMORY_FATAL_VM = 2, +}; + +struct hsa_memory_access_fault { + /* H-NUMA node that contains the device where the memory access occurred */ + u32 node_id; + /* virtual address this occurred on */ + u64 virtual_address; + /* failure attribute */ + struct hsa_access_attribute_failure failure; + /* event flags */ + enum hsa_eventid_memory_flags flags; +}; + +struct hsa_event_data { + enum hsa_event_type event_type; + + union { + /* return data associated with HSA_EVENTTYPE_SIGNAL and other events */ + struct hsa_sync_var sync_var; + /* data associated with HSA_EVENTTYPE_NODE_CHANGE */ + struct hsa_node_change node_change_state; + /* data associated with HSA_EVENTTYPE_DEVICE_STATE_CHANGE */ + struct hsa_device_state_change device_state; + /* data associated with HSA_EVENTTYPE_MEMORY */ + struct hsa_memory_access_fault memory_access_fault; + }; + + // the following data entries are internal to the KFD & thunk itself. + + u64 hw_data1; // internal thunk store for Event data (OsEventHandle) + u64 hw_data2; // internal thunk store for Event data (HWAddress) + u32 hw_data3; // internal thunk store for Event data (HWData) +}; + +struct hsa_event { + u32 event_id; + struct hsa_event_data event_data; +}; + +#endif /* KFD_HSA_CODE_H_ */ diff --git a/drivers/gpu/drm/amd/amdkfd/kfd_module.c b/drivers/gpu/drm/amd/amdkfd/kfd_module.c index 33aa23450b3f..2a2405db5b9c 100644 --- a/drivers/gpu/drm/amd/amdkfd/kfd_module.c +++ b/drivers/gpu/drm/amd/amdkfd/kfd_module.c @@ -77,6 +77,7 @@ static int kfd_init(void) static void kfd_exit(void) { + knod_exit(); kfd_cleanup_processes(); kfd_process_destroy_wq(); kfd_debugfs_fini(); diff --git a/drivers/gpu/drm/amd/amdkfd/kfd_priv.h b/drivers/gpu/drm/amd/amdkfd/kfd_priv.h index acd0e41e744c..4f5d4b0bfa89 100644 --- a/drivers/gpu/drm/amd/amdkfd/kfd_priv.h +++ b/drivers/gpu/drm/amd/amdkfd/kfd_priv.h @@ -1069,6 +1069,7 @@ bool kfd_dev_is_large_bar(struct kfd_node *dev); struct kfd_process *create_process(const struct task_struct *thread, bool primary); int kfd_process_create_wq(void); void kfd_process_destroy_wq(void); +void kfd_process_flush_wq(void); void kfd_cleanup_processes(void); struct kfd_process *kfd_create_process(struct task_struct *thread); int kfd_create_process_sysfs(struct kfd_process *process); @@ -1536,6 +1537,9 @@ void kfd_signal_hw_exception_event(u32 pasid); int kfd_set_event(struct kfd_process *p, uint32_t event_id); int kfd_reset_event(struct kfd_process *p, uint32_t event_id); int kfd_kmap_event_page(struct kfd_process *p, uint64_t event_page_offset); +int kfd_event_page_set(struct kfd_process *p, void *kernel_address, + uint64_t size, uint64_t user_handle); + int kfd_event_create(struct file *devkfd, struct kfd_process *p, uint32_t event_type, bool auto_reset, uint32_t node_id, @@ -1641,3 +1645,33 @@ static inline void kfd_debugfs_remove_process(struct kfd_process *p) {} #endif #endif + +int kfd_process_alloc_gpuvm(struct kfd_process_device *pdd, + uint64_t gpu_va, uint32_t size, + uint32_t flags, struct kgd_mem **mem, void **kptr); +void kfd_process_free_gpuvm(struct kgd_mem *mem, + struct kfd_process_device *pdd, void **kptr); +int kfd_create_queue(struct kfd_process *p, struct queue_properties *q_properties, + u32 gpu_id, u32 *queue_id_out, u64 *doorbell_offset_out); +int kfd_ioctl_set_event(struct file *filp, struct kfd_process *p, + void *data); +int kfd_ioctl_reset_event(struct file *filp, struct kfd_process *p, + void *data); +int kfd_ioctl_wait_events(struct file *filp, struct kfd_process *p, void *data); + +int kfd_ioctl_set_trap_handler(struct file *filep, + struct kfd_process *p, void *data); +int kfd_ioctl_set_scratch_backing_va(struct file *filep, + struct kfd_process *p, void *data); + + +int kfd_ioctl_acquire_vm(struct file *filep, struct kfd_process *p, void *data); +int kfd_ioctl_map_memory_to_gpu(struct file *filep, struct kfd_process *p, + void *data); +int kfd_ioctl_alloc_memory_of_gpu(struct file *filep, struct kfd_process *p, + void *data); +int kfd_ioctl_export_dmabuf(struct file *filep, + struct kfd_process *p, void *data); + +void __iomem *kfd_kernel_doorbell_mmap(struct kfd_node *dev, + struct kfd_process *process); diff --git a/drivers/gpu/drm/amd/amdkfd/kfd_process.c b/drivers/gpu/drm/amd/amdkfd/kfd_process.c index ca71fa726e32..3c9268241800 100644 --- a/drivers/gpu/drm/amd/amdkfd/kfd_process.c +++ b/drivers/gpu/drm/amd/amdkfd/kfd_process.c @@ -715,8 +715,14 @@ void kfd_process_destroy_wq(void) } } -static void kfd_process_free_gpuvm(struct kgd_mem *mem, - struct kfd_process_device *pdd, void **kptr) +void kfd_process_flush_wq(void) +{ + if (kfd_process_wq) + flush_workqueue(kfd_process_wq); +} + +void kfd_process_free_gpuvm(struct kgd_mem *mem, + struct kfd_process_device *pdd, void **kptr) { struct kfd_node *dev = pdd->dev; @@ -736,9 +742,9 @@ static void kfd_process_free_gpuvm(struct kgd_mem *mem, * to avoid concurrency. Because of that exclusiveness, we do * not need to take p->mutex. */ -static int kfd_process_alloc_gpuvm(struct kfd_process_device *pdd, - uint64_t gpu_va, uint32_t size, - uint32_t flags, struct kgd_mem **mem, void **kptr) +int kfd_process_alloc_gpuvm(struct kfd_process_device *pdd, + uint64_t gpu_va, uint32_t size, + uint32_t flags, struct kgd_mem **mem, void **kptr) { struct kfd_node *kdev = pdd->dev; int err; @@ -761,10 +767,14 @@ static int kfd_process_alloc_gpuvm(struct kfd_process_device *pdd, } if (kptr) { - err = amdgpu_amdkfd_gpuvm_map_gtt_bo_to_kernel( - (struct kgd_mem *)*mem, kptr, NULL); + if (flags & KFD_IOC_ALLOC_MEM_FLAGS_VRAM) + err = amdgpu_amdkfd_gpuvm_map_vram_bo_to_kernel( + (struct kgd_mem *)*mem, kptr, NULL); + else + err = amdgpu_amdkfd_gpuvm_map_gtt_bo_to_kernel( + (struct kgd_mem *)*mem, kptr, NULL); if (err) { - pr_debug("Map GTT BO to kernel failed\n"); + pr_debug("Map BO to kernel failed\n"); goto sync_memory_failed; } } -- 2.43.0
