J3ff/IQ.Pilot
1
0
Fork 0
forked from IQ.Lvbs/IQ.Pilot
Code Pull Requests Activity

IQ.Pilot Release Commit @ 5bc9cd3

Browse Source
This commit is contained in:
IQ.Lvbs history cleanup
2026-08-22 23:42:42 -05:00
commit d037016281
4504 changed files with 1129827 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1740
tinygrad_repo/extra/hip_gpu_driver/amdgpu_drm.h Normal file
View File

File diff suppressed because it is too large Load Diff

13609
tinygrad_repo/extra/hip_gpu_driver/gc_10_3_0_offset.h Normal file
View File

File diff suppressed because it is too large Load Diff

11685
tinygrad_repo/extra/hip_gpu_driver/gc_11_0_0_offset.h Normal file
View File

File diff suppressed because it is too large Load Diff

117
tinygrad_repo/extra/hip_gpu_driver/hip_ioctl.py Normal file
View File

@@ -0,0 +1,117 @@
# type: ignore
import ctypes, ctypes.util, struct, platform, pathlib, re, time, os
start = time.perf_counter()
# *** ioctl lib ***
libc = ctypes.CDLL(ctypes.util.find_library("c"))
# platform.processor calls `uname -p` which can return `unknown` on some systems
processor = os.getenv("IOCTL_PROCESSOR") or platform.processor() or platform.machine()
IOCTL_SYSCALL = {"aarch64": 0x1d, "x86_64":16}[processor]
def get_struct(argp, stype):
return ctypes.cast(ctypes.c_void_p(argp), ctypes.POINTER(stype)).contents
def format_struct(s):
sdats = []
for field_name, *_ in s._real_fields_:
dat = getattr(s, field_name)
if isinstance(dat, int): sdats.append(f"{field_name}:0x{dat:X}")
else: sdats.append(f"{field_name}:{dat}")
return sdats
def install_hook(c_function, python_function):
python_function_addr = ctypes.cast(ctypes.byref(python_function), ctypes.POINTER(ctypes.c_ulong)).contents.value
# AARCH64 trampoline to ioctl
if processor == "aarch64":
# 0x0000000000000000: 70 00 00 10 adr x16, #0xc
# 0x0000000000000004: 10 02 40 F9 ldr x16, [x16]
# 0x0000000000000008: 00 02 1F D6 br x16
tramp = b"\x70\x00\x00\x10\x10\x02\x40\xf9\x00\x02\x1f\xd6"
tramp += struct.pack("Q", python_function_addr)
elif processor == "x86_64":
# 0x0000000000000000: 49 B8 aa aa aa aa aa aa aa aa movabs r8, <address>
# 0x000000000000000a: 41 FF E0 jmp r8
tramp = b"\x49\xB8" + struct.pack("Q", python_function_addr) + b"\x41\xFF\xE0"
else:
raise Exception(f"processor {processor} not supported")
# get real ioctl address
ioctl_address = ctypes.cast(ctypes.byref(c_function), ctypes.POINTER(ctypes.c_ulong))
# hook ioctl
ret = libc.mprotect(ctypes.c_ulong((ioctl_address.contents.value//0x1000)*0x1000), 0x2000, 7)
assert ret == 0
libc.memcpy(ioctl_address.contents, ctypes.create_string_buffer(tramp), len(tramp))
# *** ioctl lib end ***
import tinygrad.runtime.autogen.kfd as kfd_ioctl
import tinygrad.runtime.autogen.hsa as hsa
def print_aql_queue(read_pointer_address):
rptr_offset = getattr(hsa.amd_queue_v2_t, 'read_dispatch_id').offset
queue_base = read_pointer_address - rptr_offset
queue = hsa.amd_queue_v2_t.from_address(queue_base)
print(f" AQL Queue @ 0x{queue_base:X}:")
for field_name, *_ in hsa.amd_queue_v2_t._real_fields_:
val = getattr(queue, field_name)
if isinstance(val, int): print(f" {field_name}: 0x{val:X}")
elif hasattr(val, '_length_'):
arr_vals = [f"{format_struct(v)}" if hasattr(v, '_real_fields_') else f"{v:#X}" for v in val]
print(f" {field_name}: [{', '.join(arr_vals)}]")
elif hasattr(val, '_real_fields_'): print(f" {field_name}: {format_struct(val)}")
else: print(f" {field_name}: {val}")
def ioctls_from_header():
hdr = (pathlib.Path(__file__).parent / "kfd_ioctl.h").read_text().replace("\\\n", "")
pattern = r'#define\s+(AMDKFD_IOC_[A-Z0-9_]+)\s+AMDKFD_IOW?R?\((0x[0-9a-fA-F]+),\s+struct\s([A-Za-z0-9_]+)\)'
matches = re.findall(pattern, hdr, re.MULTILINE)
return {int(nr, 0x10):(name, getattr(kfd_ioctl, "struct_"+sname, None)) for name, nr, sname in matches}
nrs = ioctls_from_header()
@ctypes.CFUNCTYPE(ctypes.c_int, ctypes.c_int, ctypes.c_ulong, ctypes.c_void_p)
def ioctl(fd, request, argp):
st = time.perf_counter()
ret = libc.syscall(IOCTL_SYSCALL, ctypes.c_int(fd), ctypes.c_ulong(request), ctypes.c_void_p(argp))
et = time.perf_counter()-st
idir, size, itype, nr = (request>>30), (request>>16)&0x3FFF, (request>>8)&0xFF, request&0xFF
if nr in nrs and itype == 75:
# /dev/kfd
name, stype = nrs[nr]
s = get_struct(argp, stype)
print(f"{(st-start)*1000:7.2f} ms +{et*1000.:7.2f} ms : {ret:2d} = {name:40s}", ' '.join(format_struct(s)))
if name == "AMDKFD_IOC_SVM":
out = ctypes.cast(s.attrs, ctypes.POINTER(kfd_ioctl.struct_kfd_ioctl_svm_attribute))
for i in range(s.nattr): print(f"{i}: {kfd_ioctl.enum_kfd_ioctl_svm_attr_type.get(out[i].type):40s}: {out[i].value:#x}")
if name == "AMDKFD_IOC_CREATE_QUEUE" and s.queue_type == kfd_ioctl.KFD_IOC_QUEUE_TYPE_COMPUTE_AQL: print_aql_queue(s.read_pointer_address)
else:
print(f"{(st-start)*1000:7.2f} ms +{et*1000.:7.2f} ms : ioctl",
f"{idir=} {size=} {itype=} {nr=} {fd=} {ret=}", os.readlink(f"/proc/self/fd/{fd}") if fd >= 0 else "")
return ret
install_hook(libc.ioctl, ioctl)
# AMD_LOG_LEVEL=4 HSAKMT_DEBUG_LEVEL=7
if __name__ == "__main__":
print("***** import tinygrad")
from tinygrad import Tensor, Device, TinyJit
print("***** access HIP")
dev = Device["HIP"]
print("***** create tensor a")
a = Tensor([1.,2.]*1024*1024, device="HIP").realize()
print("***** create tensor b")
b = Tensor([3.,4.]*1024*1024, device="HIP").realize()
@TinyJit
def add(a, b): return (a+b).realize()
for i in range(4):
print(f"***** add tensors {i}")
c = add(a, b)
#dev.synchronize()
c = add(b, a)
dev.synchronize()
print(f"***** copyout")
nc = c.numpy()
print(f"***** delete")
del add, a, b, c, dev
print(f"***** done")
os._exit(0)

1889
tinygrad_repo/extra/hip_gpu_driver/kfd_ioctl.h Normal file
View File

File diff suppressed because it is too large Load Diff

247
tinygrad_repo/extra/hip_gpu_driver/kfd_pm4_headers_ai.h Normal file
View File

@@ -0,0 +1,247 @@
/* SPDX-License-Identifier: GPL-2.0 OR MIT */
/*
* Copyright 2016-2022 Advanced Micro Devices, Inc.
*
* 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.
*
*/
#ifndef F32_MES_PM4_PACKETS_H
#define F32_MES_PM4_PACKETS_H
#define uint32_t unsigned int
#define int32_t int
#ifndef PM4_MES_HEADER_DEFINED
#define PM4_MES_HEADER_DEFINED
union PM4_MES_TYPE_3_HEADER {
struct {
uint32_t reserved1 : 8; /* < reserved */
uint32_t opcode : 8; /* < IT opcode */
uint32_t count : 14;/* < number of DWORDs - 1 in the
* information body.
*/
uint32_t type : 2; /* < packet identifier.
* It should be 3 for type 3 packets
*/
};
uint32_t u32All;
};
#endif /* PM4_MES_HEADER_DEFINED */
#ifndef PM4_MEC_RELEASE_MEM_DEFINED
#define PM4_MEC_RELEASE_MEM_DEFINED
enum mec_release_mem_event_index_enum {
event_index__mec_release_mem__end_of_pipe = 5,
event_index__mec_release_mem__shader_done = 6
};
enum mec_release_mem_cache_policy_enum {
cache_policy__mec_release_mem__lru = 0,
cache_policy__mec_release_mem__stream = 1
};
enum mec_release_mem_pq_exe_status_enum {
pq_exe_status__mec_release_mem__default = 0,
pq_exe_status__mec_release_mem__phase_update = 1
};
enum mec_release_mem_dst_sel_enum {
dst_sel__mec_release_mem__memory_controller = 0,
dst_sel__mec_release_mem__tc_l2 = 1,
dst_sel__mec_release_mem__queue_write_pointer_register = 2,
dst_sel__mec_release_mem__queue_write_pointer_poll_mask_bit = 3
};
enum mec_release_mem_int_sel_enum {
int_sel__mec_release_mem__none = 0,
int_sel__mec_release_mem__send_interrupt_only = 1,
int_sel__mec_release_mem__send_interrupt_after_write_confirm = 2,
int_sel__mec_release_mem__send_data_after_write_confirm = 3,
int_sel__mec_release_mem__unconditionally_send_int_ctxid = 4,
int_sel__mec_release_mem__conditionally_send_int_ctxid_based_on_32_bit_compare = 5,
int_sel__mec_release_mem__conditionally_send_int_ctxid_based_on_64_bit_compare = 6
};
enum mec_release_mem_data_sel_enum {
data_sel__mec_release_mem__none = 0,
data_sel__mec_release_mem__send_32_bit_low = 1,
data_sel__mec_release_mem__send_64_bit_data = 2,
data_sel__mec_release_mem__send_gpu_clock_counter = 3,
data_sel__mec_release_mem__send_cp_perfcounter_hi_lo = 4,
data_sel__mec_release_mem__store_gds_data_to_memory = 5
};
struct pm4_mec_release_mem {
union {
union PM4_MES_TYPE_3_HEADER header; /*header */
unsigned int ordinal1;
};
union {
struct {
unsigned int event_type:6;
unsigned int reserved1:2;
enum mec_release_mem_event_index_enum event_index:4;
unsigned int tcl1_vol_action_ena:1;
unsigned int tc_vol_action_ena:1;
unsigned int reserved2:1;
unsigned int tc_wb_action_ena:1;
unsigned int tcl1_action_ena:1;
unsigned int tc_action_ena:1;
uint32_t reserved3:1;
uint32_t tc_nc_action_ena:1;
uint32_t tc_wc_action_ena:1;
uint32_t tc_md_action_ena:1;
uint32_t reserved4:3;
enum mec_release_mem_cache_policy_enum cache_policy:2;
uint32_t reserved5:2;
enum mec_release_mem_pq_exe_status_enum pq_exe_status:1;
uint32_t reserved6:2;
} bitfields2;
unsigned int ordinal2;
};
union {
struct {
uint32_t reserved7:16;
enum mec_release_mem_dst_sel_enum dst_sel:2;
uint32_t reserved8:6;
enum mec_release_mem_int_sel_enum int_sel:3;
uint32_t reserved9:2;
enum mec_release_mem_data_sel_enum data_sel:3;
} bitfields3;
unsigned int ordinal3;
};
union {
struct {
uint32_t reserved10:2;
unsigned int address_lo_32b:30;
} bitfields4;
struct {
uint32_t reserved11:3;
uint32_t address_lo_64b:29;
} bitfields4b;
uint32_t reserved12;
unsigned int ordinal4;
};
union {
uint32_t address_hi;
uint32_t reserved13;
uint32_t ordinal5;
};
union {
uint32_t data_lo;
uint32_t cmp_data_lo;
struct {
uint32_t dw_offset:16;
uint32_t num_dwords:16;
} bitfields6c;
uint32_t reserved14;
uint32_t ordinal6;
};
union {
uint32_t data_hi;
uint32_t cmp_data_hi;
uint32_t reserved15;
uint32_t reserved16;
uint32_t ordinal7;
};
uint32_t int_ctxid;
};
#endif
#ifndef PM4_MEC_WRITE_DATA_DEFINED
#define PM4_MEC_WRITE_DATA_DEFINED
enum WRITE_DATA_dst_sel_enum {
dst_sel___write_data__mem_mapped_register = 0,
dst_sel___write_data__tc_l2 = 2,
dst_sel___write_data__gds = 3,
dst_sel___write_data__memory = 5,
dst_sel___write_data__memory_mapped_adc_persistent_state = 6,
};
enum WRITE_DATA_addr_incr_enum {
addr_incr___write_data__increment_address = 0,
addr_incr___write_data__do_not_increment_address = 1
};
enum WRITE_DATA_wr_confirm_enum {
wr_confirm___write_data__do_not_wait_for_write_confirmation = 0,
wr_confirm___write_data__wait_for_write_confirmation = 1
};
enum WRITE_DATA_cache_policy_enum {
cache_policy___write_data__lru = 0,
cache_policy___write_data__stream = 1
};
struct pm4_mec_write_data_mmio {
union {
union PM4_MES_TYPE_3_HEADER header; /*header */
unsigned int ordinal1;
};
union {
struct {
unsigned int reserved1:8;
unsigned int dst_sel:4;
unsigned int reserved2:4;
unsigned int addr_incr:1;
unsigned int reserved3:2;
unsigned int resume_vf:1;
unsigned int wr_confirm:1;
unsigned int reserved4:4;
unsigned int cache_policy:2;
unsigned int reserved5:5;
} bitfields2;
unsigned int ordinal2;
};
union {
struct {
unsigned int dst_mmreg_addr:18;
unsigned int reserved6:14;
} bitfields3;
unsigned int ordinal3;
};
uint32_t reserved7;
uint32_t data;
};
#endif
enum {
CACHE_FLUSH_AND_INV_TS_EVENT = 0x00000014
};
#endif

4886
tinygrad_repo/extra/hip_gpu_driver/navi10_sdma_pkt_open.h Normal file
View File

File diff suppressed because it is too large Load Diff

470
tinygrad_repo/extra/hip_gpu_driver/nvd.h Normal file
View File

@@ -0,0 +1,470 @@
/*
* Copyright 2019 Advanced Micro Devices, Inc.
*
* 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.
*
*/
#ifndef NVD_H
#define NVD_H
/**
* Navi's PM4 definitions
*/
#define PACKET_TYPE0 0
#define PACKET_TYPE1 1
#define PACKET_TYPE2 2
#define PACKET_TYPE3 3
#define CP_PACKET_GET_TYPE(h) (((h) >> 30) & 3)
#define CP_PACKET_GET_COUNT(h) (((h) >> 16) & 0x3FFF)
#define CP_PACKET0_GET_REG(h) ((h) & 0xFFFF)
#define CP_PACKET3_GET_OPCODE(h) (((h) >> 8) & 0xFF)
#define PACKET0(reg, n) ((PACKET_TYPE0 << 30) | \
((reg) & 0xFFFF) | \
((n) & 0x3FFF) << 16)
#define CP_PACKET2 0x80000000
#define PACKET2_PAD_SHIFT 0
#define PACKET2_PAD_MASK (0x3fffffff << 0)
#define PACKET2(v) (CP_PACKET2 | REG_SET(PACKET2_PAD, (v)))
#define PACKET3(op, n) ((PACKET_TYPE3 << 30) | \
(((op) & 0xFF) << 8) | \
((n) & 0x3FFF) << 16)
#define PACKET3_COMPUTE(op, n) (PACKET3(op, n) | 1 << 1)
/* Packet 3 types */
#define PACKET3_NOP 0x10
#define PACKET3_SET_BASE 0x11
#define PACKET3_BASE_INDEX(x) ((x) << 0)
#define CE_PARTITION_BASE 3
#define PACKET3_CLEAR_STATE 0x12
#define PACKET3_INDEX_BUFFER_SIZE 0x13
#define PACKET3_DISPATCH_DIRECT 0x15
#define PACKET3_DISPATCH_INDIRECT 0x16
#define PACKET3_INDIRECT_BUFFER_END 0x17
#define PACKET3_INDIRECT_BUFFER_CNST_END 0x19
#define PACKET3_ATOMIC_GDS 0x1D
#define PACKET3_ATOMIC_MEM 0x1E
#define PACKET3_OCCLUSION_QUERY 0x1F
#define PACKET3_SET_PREDICATION 0x20
#define PACKET3_REG_RMW 0x21
#define PACKET3_COND_EXEC 0x22
#define PACKET3_PRED_EXEC 0x23
#define PACKET3_DRAW_INDIRECT 0x24
#define PACKET3_DRAW_INDEX_INDIRECT 0x25
#define PACKET3_INDEX_BASE 0x26
#define PACKET3_DRAW_INDEX_2 0x27
#define PACKET3_CONTEXT_CONTROL 0x28
#define PACKET3_INDEX_TYPE 0x2A
#define PACKET3_DRAW_INDIRECT_MULTI 0x2C
#define PACKET3_DRAW_INDEX_AUTO 0x2D
#define PACKET3_NUM_INSTANCES 0x2F
#define PACKET3_DRAW_INDEX_MULTI_AUTO 0x30
#define PACKET3_INDIRECT_BUFFER_PRIV 0x32
#define PACKET3_INDIRECT_BUFFER_CNST 0x33
#define PACKET3_COND_INDIRECT_BUFFER_CNST 0x33
#define PACKET3_STRMOUT_BUFFER_UPDATE 0x34
#define PACKET3_DRAW_INDEX_OFFSET_2 0x35
#define PACKET3_DRAW_PREAMBLE 0x36
#define PACKET3_WRITE_DATA 0x37
#define WRITE_DATA_DST_SEL(x) ((x) << 8)
/* 0 - register
* 1 - memory (sync - via GRBM)
* 2 - gl2
* 3 - gds
* 4 - reserved
* 5 - memory (async - direct)
*/
#define WR_ONE_ADDR (1 << 16)
#define WR_CONFIRM (1 << 20)
#define WRITE_DATA_CACHE_POLICY(x) ((x) << 25)
/* 0 - LRU
* 1 - Stream
*/
#define WRITE_DATA_ENGINE_SEL(x) ((x) << 30)
/* 0 - me
* 1 - pfp
* 2 - ce
*/
#define PACKET3_DRAW_INDEX_INDIRECT_MULTI 0x38
#define PACKET3_MEM_SEMAPHORE 0x39
# define PACKET3_SEM_USE_MAILBOX (0x1 << 16)
# define PACKET3_SEM_SEL_SIGNAL_TYPE (0x1 << 20) /* 0 = increment, 1 = write 1 */
# define PACKET3_SEM_SEL_SIGNAL (0x6 << 29)
# define PACKET3_SEM_SEL_WAIT (0x7 << 29)
#define PACKET3_DRAW_INDEX_MULTI_INST 0x3A
#define PACKET3_COPY_DW 0x3B
#define PACKET3_WAIT_REG_MEM 0x3C
#define WAIT_REG_MEM_FUNCTION(x) ((x) << 0)
/* 0 - always
* 1 - <
* 2 - <=
* 3 - ==
* 4 - !=
* 5 - >=
* 6 - >
*/
#define WAIT_REG_MEM_MEM_SPACE(x) ((x) << 4)
/* 0 - reg
* 1 - mem
*/
#define WAIT_REG_MEM_OPERATION(x) ((x) << 6)
/* 0 - wait_reg_mem
* 1 - wr_wait_wr_reg
*/
#define WAIT_REG_MEM_ENGINE(x) ((x) << 8)
/* 0 - me
* 1 - pfp
*/
#define PACKET3_INDIRECT_BUFFER 0x3F
#define INDIRECT_BUFFER_VALID (1 << 23)
#define INDIRECT_BUFFER_CACHE_POLICY(x) ((x) << 28)
/* 0 - LRU
* 1 - Stream
* 2 - Bypass
*/
#define INDIRECT_BUFFER_PRE_ENB(x) ((x) << 21)
#define INDIRECT_BUFFER_PRE_RESUME(x) ((x) << 30)
#define PACKET3_COND_INDIRECT_BUFFER 0x3F
#define PACKET3_COPY_DATA 0x40
#define PACKET3_CP_DMA 0x41
#define PACKET3_PFP_SYNC_ME 0x42
#define PACKET3_SURFACE_SYNC 0x43
#define PACKET3_ME_INITIALIZE 0x44
#define PACKET3_COND_WRITE 0x45
#define PACKET3_EVENT_WRITE 0x46
#define EVENT_TYPE(x) ((x) << 0)
#define EVENT_INDEX(x) ((x) << 8)
/* 0 - any non-TS event
* 1 - ZPASS_DONE, PIXEL_PIPE_STAT_*
* 2 - SAMPLE_PIPELINESTAT
* 3 - SAMPLE_STREAMOUTSTAT*
* 4 - *S_PARTIAL_FLUSH
*/
#define PACKET3_EVENT_WRITE_EOP 0x47
#define PACKET3_EVENT_WRITE_EOS 0x48
#define PACKET3_RELEASE_MEM 0x49
#define PACKET3_RELEASE_MEM_EVENT_TYPE(x) ((x) << 0)
#define PACKET3_RELEASE_MEM_EVENT_INDEX(x) ((x) << 8)
#define PACKET3_RELEASE_MEM_GCR_GLM_WB (1 << 12)
#define PACKET3_RELEASE_MEM_GCR_GLM_INV (1 << 13)
#define PACKET3_RELEASE_MEM_GCR_GLV_INV (1 << 14)
#define PACKET3_RELEASE_MEM_GCR_GL1_INV (1 << 15)
#define PACKET3_RELEASE_MEM_GCR_GL2_US (1 << 16)
#define PACKET3_RELEASE_MEM_GCR_GL2_RANGE (1 << 17)
#define PACKET3_RELEASE_MEM_GCR_GL2_DISCARD (1 << 19)
#define PACKET3_RELEASE_MEM_GCR_GL2_INV (1 << 20)
#define PACKET3_RELEASE_MEM_GCR_GL2_WB (1 << 21)
#define PACKET3_RELEASE_MEM_GCR_SEQ (1 << 22)
#define PACKET3_RELEASE_MEM_CACHE_POLICY(x) ((x) << 25)
/* 0 - cache_policy__me_release_mem__lru
* 1 - cache_policy__me_release_mem__stream
* 2 - cache_policy__me_release_mem__noa
* 3 - cache_policy__me_release_mem__bypass
*/
#define PACKET3_RELEASE_MEM_EXECUTE (1 << 28)
#define PACKET3_RELEASE_MEM_DATA_SEL(x) ((x) << 29)
/* 0 - discard
* 1 - send low 32bit data
* 2 - send 64bit data
* 3 - send 64bit GPU counter value
* 4 - send 64bit sys counter value
*/
#define PACKET3_RELEASE_MEM_INT_SEL(x) ((x) << 24)
/* 0 - none
* 1 - interrupt only (DATA_SEL = 0)
* 2 - interrupt when data write is confirmed
*/
#define PACKET3_RELEASE_MEM_DST_SEL(x) ((x) << 16)
/* 0 - MC
* 1 - TC/L2
*/
#define PACKET3_PREAMBLE_CNTL 0x4A
# define PACKET3_PREAMBLE_BEGIN_CLEAR_STATE (2 << 28)
# define PACKET3_PREAMBLE_END_CLEAR_STATE (3 << 28)
#define PACKET3_DMA_DATA 0x50
/* 1. header
* 2. CONTROL
* 3. SRC_ADDR_LO or DATA [31:0]
* 4. SRC_ADDR_HI [31:0]
* 5. DST_ADDR_LO [31:0]
* 6. DST_ADDR_HI [7:0]
* 7. COMMAND [31:26] | BYTE_COUNT [25:0]
*/
/* CONTROL */
# define PACKET3_DMA_DATA_ENGINE(x) ((x) << 0)
/* 0 - ME
* 1 - PFP
*/
# define PACKET3_DMA_DATA_SRC_CACHE_POLICY(x) ((x) << 13)
/* 0 - LRU
* 1 - Stream
*/
# define PACKET3_DMA_DATA_DST_SEL(x) ((x) << 20)
/* 0 - DST_ADDR using DAS
* 1 - GDS
* 3 - DST_ADDR using L2
*/
# define PACKET3_DMA_DATA_DST_CACHE_POLICY(x) ((x) << 25)
/* 0 - LRU
* 1 - Stream
*/
# define PACKET3_DMA_DATA_SRC_SEL(x) ((x) << 29)
/* 0 - SRC_ADDR using SAS
* 1 - GDS
* 2 - DATA
* 3 - SRC_ADDR using L2
*/
# define PACKET3_DMA_DATA_CP_SYNC (1 << 31)
/* COMMAND */
# define PACKET3_DMA_DATA_CMD_SAS (1 << 26)
/* 0 - memory
* 1 - register
*/
# define PACKET3_DMA_DATA_CMD_DAS (1 << 27)
/* 0 - memory
* 1 - register
*/
# define PACKET3_DMA_DATA_CMD_SAIC (1 << 28)
# define PACKET3_DMA_DATA_CMD_DAIC (1 << 29)
# define PACKET3_DMA_DATA_CMD_RAW_WAIT (1 << 30)
#define PACKET3_CONTEXT_REG_RMW 0x51
#define PACKET3_GFX_CNTX_UPDATE 0x52
#define PACKET3_BLK_CNTX_UPDATE 0x53
#define PACKET3_INCR_UPDT_STATE 0x55
#define PACKET3_ACQUIRE_MEM 0x58
/* 1. HEADER
* 2. COHER_CNTL [30:0]
* 2.1 ENGINE_SEL [31:31]
* 2. COHER_SIZE [31:0]
* 3. COHER_SIZE_HI [7:0]
* 4. COHER_BASE_LO [31:0]
* 5. COHER_BASE_HI [23:0]
* 7. POLL_INTERVAL [15:0]
* 8. GCR_CNTL [18:0]
*/
#define PACKET3_ACQUIRE_MEM_GCR_CNTL_GLI_INV(x) ((x) << 0)
/*
* 0:NOP
* 1:ALL
* 2:RANGE
* 3:FIRST_LAST
*/
#define PACKET3_ACQUIRE_MEM_GCR_CNTL_GL1_RANGE(x) ((x) << 2)
/*
* 0:ALL
* 1:reserved
* 2:RANGE
* 3:FIRST_LAST
*/
#define PACKET3_ACQUIRE_MEM_GCR_CNTL_GLM_WB(x) ((x) << 4)
#define PACKET3_ACQUIRE_MEM_GCR_CNTL_GLM_INV(x) ((x) << 5)
#define PACKET3_ACQUIRE_MEM_GCR_CNTL_GLK_WB(x) ((x) << 6)
#define PACKET3_ACQUIRE_MEM_GCR_CNTL_GLK_INV(x) ((x) << 7)
#define PACKET3_ACQUIRE_MEM_GCR_CNTL_GLV_INV(x) ((x) << 8)
#define PACKET3_ACQUIRE_MEM_GCR_CNTL_GL1_INV(x) ((x) << 9)
#define PACKET3_ACQUIRE_MEM_GCR_CNTL_GL2_US(x) ((x) << 10)
#define PACKET3_ACQUIRE_MEM_GCR_CNTL_GL2_RANGE(x) ((x) << 11)
/*
* 0:ALL
* 1:VOL
* 2:RANGE
* 3:FIRST_LAST
*/
#define PACKET3_ACQUIRE_MEM_GCR_CNTL_GL2_DISCARD(x) ((x) << 13)
#define PACKET3_ACQUIRE_MEM_GCR_CNTL_GL2_INV(x) ((x) << 14)
#define PACKET3_ACQUIRE_MEM_GCR_CNTL_GL2_WB(x) ((x) << 15)
#define PACKET3_ACQUIRE_MEM_GCR_CNTL_SEQ(x) ((x) << 16)
/*
* 0: PARALLEL
* 1: FORWARD
* 2: REVERSE
*/
#define PACKET3_ACQUIRE_MEM_GCR_RANGE_IS_PA (1 << 18)
#define PACKET3_REWIND 0x59
#define PACKET3_INTERRUPT 0x5A
#define PACKET3_GEN_PDEPTE 0x5B
#define PACKET3_INDIRECT_BUFFER_PASID 0x5C
#define PACKET3_PRIME_UTCL2 0x5D
#define PACKET3_LOAD_UCONFIG_REG 0x5E
#define PACKET3_LOAD_SH_REG 0x5F
#define PACKET3_LOAD_CONFIG_REG 0x60
#define PACKET3_LOAD_CONTEXT_REG 0x61
#define PACKET3_LOAD_COMPUTE_STATE 0x62
#define PACKET3_LOAD_SH_REG_INDEX 0x63
#define PACKET3_SET_CONFIG_REG 0x68
#define PACKET3_SET_CONFIG_REG_START 0x00002000
#define PACKET3_SET_CONFIG_REG_END 0x00002c00
#define PACKET3_SET_CONTEXT_REG 0x69
#define PACKET3_SET_CONTEXT_REG_START 0x0000a000
#define PACKET3_SET_CONTEXT_REG_END 0x0000a400
#define PACKET3_SET_CONTEXT_REG_INDEX 0x6A
#define PACKET3_SET_VGPR_REG_DI_MULTI 0x71
#define PACKET3_SET_SH_REG_DI 0x72
#define PACKET3_SET_CONTEXT_REG_INDIRECT 0x73
#define PACKET3_SET_SH_REG_DI_MULTI 0x74
#define PACKET3_GFX_PIPE_LOCK 0x75
#define PACKET3_SET_SH_REG 0x76
#define PACKET3_SET_SH_REG_START 0x00002c00
#define PACKET3_SET_SH_REG_END 0x00003000
#define PACKET3_SET_SH_REG_OFFSET 0x77
#define PACKET3_SET_QUEUE_REG 0x78
#define PACKET3_SET_UCONFIG_REG 0x79
#define PACKET3_SET_UCONFIG_REG_START 0x0000c000
#define PACKET3_SET_UCONFIG_REG_END 0x0000c400
#define PACKET3_SET_UCONFIG_REG_INDEX 0x7A
#define PACKET3_FORWARD_HEADER 0x7C
#define PACKET3_SCRATCH_RAM_WRITE 0x7D
#define PACKET3_SCRATCH_RAM_READ 0x7E
#define PACKET3_LOAD_CONST_RAM 0x80
#define PACKET3_WRITE_CONST_RAM 0x81
#define PACKET3_DUMP_CONST_RAM 0x83
#define PACKET3_INCREMENT_CE_COUNTER 0x84
#define PACKET3_INCREMENT_DE_COUNTER 0x85
#define PACKET3_WAIT_ON_CE_COUNTER 0x86
#define PACKET3_WAIT_ON_DE_COUNTER_DIFF 0x88
#define PACKET3_SWITCH_BUFFER 0x8B
#define PACKET3_DISPATCH_DRAW_PREAMBLE 0x8C
#define PACKET3_DISPATCH_DRAW_PREAMBLE_ACE 0x8C
#define PACKET3_DISPATCH_DRAW 0x8D
#define PACKET3_DISPATCH_DRAW_ACE 0x8D
#define PACKET3_GET_LOD_STATS 0x8E
#define PACKET3_DRAW_MULTI_PREAMBLE 0x8F
#define PACKET3_FRAME_CONTROL 0x90
# define FRAME_TMZ (1 << 0)
# define FRAME_CMD(x) ((x) << 28)
/*
* x=0: tmz_begin
* x=1: tmz_end
*/
#define PACKET3_INDEX_ATTRIBUTES_INDIRECT 0x91
#define PACKET3_WAIT_REG_MEM64 0x93
#define PACKET3_COND_PREEMPT 0x94
#define PACKET3_HDP_FLUSH 0x95
#define PACKET3_COPY_DATA_RB 0x96
#define PACKET3_INVALIDATE_TLBS 0x98
# define PACKET3_INVALIDATE_TLBS_DST_SEL(x) ((x) << 0)
# define PACKET3_INVALIDATE_TLBS_ALL_HUB(x) ((x) << 4)
# define PACKET3_INVALIDATE_TLBS_PASID(x) ((x) << 5)
#define PACKET3_AQL_PACKET 0x99
#define PACKET3_DMA_DATA_FILL_MULTI 0x9A
#define PACKET3_SET_SH_REG_INDEX 0x9B
#define PACKET3_DRAW_INDIRECT_COUNT_MULTI 0x9C
#define PACKET3_DRAW_INDEX_INDIRECT_COUNT_MULTI 0x9D
#define PACKET3_DUMP_CONST_RAM_OFFSET 0x9E
#define PACKET3_LOAD_CONTEXT_REG_INDEX 0x9F
#define PACKET3_SET_RESOURCES 0xA0
/* 1. header
* 2. CONTROL
* 3. QUEUE_MASK_LO [31:0]
* 4. QUEUE_MASK_HI [31:0]
* 5. GWS_MASK_LO [31:0]
* 6. GWS_MASK_HI [31:0]
* 7. OAC_MASK [15:0]
* 8. GDS_HEAP_SIZE [16:11] | GDS_HEAP_BASE [5:0]
*/
# define PACKET3_SET_RESOURCES_VMID_MASK(x) ((x) << 0)
# define PACKET3_SET_RESOURCES_UNMAP_LATENTY(x) ((x) << 16)
# define PACKET3_SET_RESOURCES_QUEUE_TYPE(x) ((x) << 29)
#define PACKET3_MAP_PROCESS 0xA1
#define PACKET3_MAP_QUEUES 0xA2
/* 1. header
* 2. CONTROL
* 3. CONTROL2
* 4. MQD_ADDR_LO [31:0]
* 5. MQD_ADDR_HI [31:0]
* 6. WPTR_ADDR_LO [31:0]
* 7. WPTR_ADDR_HI [31:0]
*/
/* CONTROL */
# define PACKET3_MAP_QUEUES_QUEUE_SEL(x) ((x) << 4)
# define PACKET3_MAP_QUEUES_VMID(x) ((x) << 8)
# define PACKET3_MAP_QUEUES_QUEUE(x) ((x) << 13)
# define PACKET3_MAP_QUEUES_PIPE(x) ((x) << 16)
# define PACKET3_MAP_QUEUES_ME(x) ((x) << 18)
# define PACKET3_MAP_QUEUES_QUEUE_TYPE(x) ((x) << 21)
# define PACKET3_MAP_QUEUES_ALLOC_FORMAT(x) ((x) << 24)
# define PACKET3_MAP_QUEUES_ENGINE_SEL(x) ((x) << 26)
# define PACKET3_MAP_QUEUES_NUM_QUEUES(x) ((x) << 29)
/* CONTROL2 */
# define PACKET3_MAP_QUEUES_CHECK_DISABLE(x) ((x) << 1)
# define PACKET3_MAP_QUEUES_DOORBELL_OFFSET(x) ((x) << 2)
#define PACKET3_UNMAP_QUEUES 0xA3
/* 1. header
* 2. CONTROL
* 3. CONTROL2
* 4. CONTROL3
* 5. CONTROL4
* 6. CONTROL5
*/
/* CONTROL */
# define PACKET3_UNMAP_QUEUES_ACTION(x) ((x) << 0)
/* 0 - PREEMPT_QUEUES
* 1 - RESET_QUEUES
* 2 - DISABLE_PROCESS_QUEUES
* 3 - PREEMPT_QUEUES_NO_UNMAP
*/
# define PACKET3_UNMAP_QUEUES_QUEUE_SEL(x) ((x) << 4)
# define PACKET3_UNMAP_QUEUES_ENGINE_SEL(x) ((x) << 26)
# define PACKET3_UNMAP_QUEUES_NUM_QUEUES(x) ((x) << 29)
/* CONTROL2a */
# define PACKET3_UNMAP_QUEUES_PASID(x) ((x) << 0)
/* CONTROL2b */
# define PACKET3_UNMAP_QUEUES_DOORBELL_OFFSET0(x) ((x) << 2)
/* CONTROL3a */
# define PACKET3_UNMAP_QUEUES_DOORBELL_OFFSET1(x) ((x) << 2)
/* CONTROL3b */
# define PACKET3_UNMAP_QUEUES_RB_WPTR(x) ((x) << 0)
/* CONTROL4 */
# define PACKET3_UNMAP_QUEUES_DOORBELL_OFFSET2(x) ((x) << 2)
/* CONTROL5 */
# define PACKET3_UNMAP_QUEUES_DOORBELL_OFFSET3(x) ((x) << 2)
#define PACKET3_QUERY_STATUS 0xA4
/* 1. header
* 2. CONTROL
* 3. CONTROL2
* 4. ADDR_LO [31:0]
* 5. ADDR_HI [31:0]
* 6. DATA_LO [31:0]
* 7. DATA_HI [31:0]
*/
/* CONTROL */
# define PACKET3_QUERY_STATUS_CONTEXT_ID(x) ((x) << 0)
# define PACKET3_QUERY_STATUS_INTERRUPT_SEL(x) ((x) << 28)
# define PACKET3_QUERY_STATUS_COMMAND(x) ((x) << 30)
/* CONTROL2a */
# define PACKET3_QUERY_STATUS_PASID(x) ((x) << 0)
/* CONTROL2b */
# define PACKET3_QUERY_STATUS_DOORBELL_OFFSET(x) ((x) << 2)
# define PACKET3_QUERY_STATUS_ENG_SEL(x) ((x) << 25)
#define PACKET3_RUN_LIST 0xA5
#define PACKET3_MAP_PROCESS_VM 0xA6
/* GFX11 */
#define PACKET3_SET_Q_PREEMPTION_MODE 0xF0
# define PACKET3_SET_Q_PREEMPTION_MODE_IB_VMID(x) ((x) << 0)
# define PACKET3_SET_Q_PREEMPTION_MODE_INIT_SHADOW_MEM (1 << 0)
#endif

571
tinygrad_repo/extra/hip_gpu_driver/sdma_registers.h Normal file
View File

@@ -0,0 +1,571 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2020, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.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 with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS OR COPYRIGHT HOLDERS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef HSA_RUNTIME_CORE_INC_SDMA_REGISTERS_H_
#define HSA_RUNTIME_CORE_INC_SDMA_REGISTERS_H_
#include <stddef.h>
#include <stdint.h>
namespace rocr {
namespace AMD {
// SDMA packet for VI device.
// Reference: http://people.freedesktop.org/~agd5f/dma_packets.txt
const unsigned int SDMA_OP_COPY = 1;
const unsigned int SDMA_OP_FENCE = 5;
const unsigned int SDMA_OP_TRAP = 6;
const unsigned int SDMA_OP_POLL_REGMEM = 8;
const unsigned int SDMA_OP_ATOMIC = 10;
const unsigned int SDMA_OP_CONST_FILL = 11;
const unsigned int SDMA_OP_TIMESTAMP = 13;
const unsigned int SDMA_OP_GCR = 17;
const unsigned int SDMA_SUBOP_COPY_LINEAR = 0;
const unsigned int SDMA_SUBOP_COPY_LINEAR_RECT = 4;
const unsigned int SDMA_SUBOP_TIMESTAMP_GET_GLOBAL = 2;
const unsigned int SDMA_SUBOP_USER_GCR = 1;
const unsigned int SDMA_ATOMIC_ADD64 = 47;
typedef struct SDMA_PKT_COPY_LINEAR_TAG {
union {
struct {
unsigned int op : 8;
unsigned int sub_op : 8;
unsigned int extra_info : 16;
};
unsigned int DW_0_DATA;
} HEADER_UNION;
union {
struct {
unsigned int count : 22;
unsigned int reserved_0 : 10;
};
unsigned int DW_1_DATA;
} COUNT_UNION;
union {
struct {
unsigned int reserved_0 : 16;
unsigned int dst_swap : 2;
unsigned int reserved_1 : 6;
unsigned int src_swap : 2;
unsigned int reserved_2 : 6;
};
unsigned int DW_2_DATA;
} PARAMETER_UNION;
union {
struct {
unsigned int src_addr_31_0 : 32;
};
unsigned int DW_3_DATA;
} SRC_ADDR_LO_UNION;
union {
struct {
unsigned int src_addr_63_32 : 32;
};
unsigned int DW_4_DATA;
} SRC_ADDR_HI_UNION;
union {
struct {
unsigned int dst_addr_31_0 : 32;
};
unsigned int DW_5_DATA;
} DST_ADDR_LO_UNION;
union {
struct {
unsigned int dst_addr_63_32 : 32;
};
unsigned int DW_6_DATA;
} DST_ADDR_HI_UNION;
static const size_t kMaxSize_ = 0x3fffe0;
} SDMA_PKT_COPY_LINEAR;
// linear sub-window
typedef struct SDMA_PKT_COPY_LINEAR_RECT_TAG {
static const unsigned int pitch_bits = 19;
static const unsigned int slice_bits = 28;
static const unsigned int rect_xy_bits = 14;
static const unsigned int rect_z_bits = 11;
union {
struct {
unsigned int op : 8;
unsigned int sub_op : 8;
unsigned int reserved : 13;
unsigned int element : 3;
};
unsigned int DW_0_DATA;
} HEADER_UNION;
union {
struct {
unsigned int src_addr_31_0 : 32;
};
unsigned int DW_1_DATA;
} SRC_ADDR_LO_UNION;
union {
struct {
unsigned int src_addr_63_32 : 32;
};
unsigned int DW_2_DATA;
} SRC_ADDR_HI_UNION;
union {
struct {
unsigned int src_offset_x : 14;
unsigned int reserved_1 : 2;
unsigned int src_offset_y : 14;
unsigned int reserved_2 : 2;
};
unsigned int DW_3_DATA;
} SRC_PARAMETER_1_UNION;
union {
struct {
unsigned int src_offset_z : 11;
unsigned int reserved_1 : 2;
unsigned int src_pitch : pitch_bits;
};
unsigned int DW_4_DATA;
} SRC_PARAMETER_2_UNION;
union {
struct {
unsigned int src_slice_pitch : slice_bits;
unsigned int reserved_1 : 4;
};
unsigned int DW_5_DATA;
} SRC_PARAMETER_3_UNION;
union {
struct {
unsigned int dst_addr_31_0 : 32;
};
unsigned int DW_6_DATA;
} DST_ADDR_LO_UNION;
union {
struct {
unsigned int dst_addr_63_32 : 32;
};
unsigned int DW_7_DATA;
} DST_ADDR_HI_UNION;
union {
struct {
unsigned int dst_offset_x : 14;
unsigned int reserved_1 : 2;
unsigned int dst_offset_y : 14;
unsigned int reserved_2 : 2;
};
unsigned int DW_8_DATA;
} DST_PARAMETER_1_UNION;
union {
struct {
unsigned int dst_offset_z : 11;
unsigned int reserved_1 : 2;
unsigned int dst_pitch : pitch_bits;
};
unsigned int DW_9_DATA;
} DST_PARAMETER_2_UNION;
union {
struct {
unsigned int dst_slice_pitch : slice_bits;
unsigned int reserved_1 : 4;
};
unsigned int DW_10_DATA;
} DST_PARAMETER_3_UNION;
union {
struct {
unsigned int rect_x : rect_xy_bits;
unsigned int reserved_1 : 2;
unsigned int rect_y : rect_xy_bits;
unsigned int reserved_2 : 2;
};
unsigned int DW_11_DATA;
} RECT_PARAMETER_1_UNION;
union {
struct {
unsigned int rect_z : rect_z_bits;
unsigned int reserved_1 : 5;
unsigned int dst_swap : 2;
unsigned int reserved_2 : 6;
unsigned int src_swap : 2;
unsigned int reserved_3 : 6;
};
unsigned int DW_12_DATA;
} RECT_PARAMETER_2_UNION;
} SDMA_PKT_COPY_LINEAR_RECT;
typedef struct SDMA_PKT_CONSTANT_FILL_TAG {
union {
struct {
unsigned int op : 8;
unsigned int sub_op : 8;
unsigned int sw : 2;
unsigned int reserved_0 : 12;
unsigned int fillsize : 2;
};
unsigned int DW_0_DATA;
} HEADER_UNION;
union {
struct {
unsigned int dst_addr_31_0 : 32;
};
unsigned int DW_1_DATA;
} DST_ADDR_LO_UNION;
union {
struct {
unsigned int dst_addr_63_32 : 32;
};
unsigned int DW_2_DATA;
} DST_ADDR_HI_UNION;
union {
struct {
unsigned int src_data_31_0 : 32;
};
unsigned int DW_3_DATA;
} DATA_UNION;
union {
struct {
unsigned int count : 22;
unsigned int reserved_0 : 10;
};
unsigned int DW_4_DATA;
} COUNT_UNION;
static const size_t kMaxSize_ = 0x3fffe0;
} SDMA_PKT_CONSTANT_FILL;
typedef struct SDMA_PKT_FENCE_TAG {
union {
struct {
unsigned int op : 8;
unsigned int sub_op : 8;
unsigned int mtype : 3;
unsigned int gcc : 1;
unsigned int sys : 1;
unsigned int pad1 : 1;
unsigned int snp : 1;
unsigned int gpa : 1;
unsigned int l2_policy : 2;
unsigned int reserved_0 : 6;
};
unsigned int DW_0_DATA;
} HEADER_UNION;
union {
struct {
unsigned int addr_31_0 : 32;
};
unsigned int DW_1_DATA;
} ADDR_LO_UNION;
union {
struct {
unsigned int addr_63_32 : 32;
};
unsigned int DW_2_DATA;
} ADDR_HI_UNION;
union {
struct {
unsigned int data : 32;
};
unsigned int DW_3_DATA;
} DATA_UNION;
} SDMA_PKT_FENCE;
typedef struct SDMA_PKT_POLL_REGMEM_TAG {
union {
struct {
unsigned int op : 8;
unsigned int sub_op : 8;
unsigned int reserved_0 : 10;
unsigned int hdp_flush : 1;
unsigned int reserved_1 : 1;
unsigned int func : 3;
unsigned int mem_poll : 1;
};
unsigned int DW_0_DATA;
} HEADER_UNION;
union {
struct {
unsigned int addr_31_0 : 32;
};
unsigned int DW_1_DATA;
} ADDR_LO_UNION;
union {
struct {
unsigned int addr_63_32 : 32;
};
unsigned int DW_2_DATA;
} ADDR_HI_UNION;
union {
struct {
unsigned int value : 32;
};
unsigned int DW_3_DATA;
} VALUE_UNION;
union {
struct {
unsigned int mask : 32;
};
unsigned int DW_4_DATA;
} MASK_UNION;
union {
struct {
unsigned int interval : 16;
unsigned int retry_count : 12;
unsigned int reserved_0 : 4;
};
unsigned int DW_5_DATA;
} DW5_UNION;
} SDMA_PKT_POLL_REGMEM;
typedef struct SDMA_PKT_ATOMIC_TAG {
union {
struct {
unsigned int op : 8;
unsigned int sub_op : 8;
unsigned int l : 1;
unsigned int reserved_0 : 8;
unsigned int operation : 7;
};
unsigned int DW_0_DATA;
} HEADER_UNION;
union {
struct {
unsigned int addr_31_0 : 32;
};
unsigned int DW_1_DATA;
} ADDR_LO_UNION;
union {
struct {
unsigned int addr_63_32 : 32;
};
unsigned int DW_2_DATA;
} ADDR_HI_UNION;
union {
struct {
unsigned int src_data_31_0 : 32;
};
unsigned int DW_3_DATA;
} SRC_DATA_LO_UNION;
union {
struct {
unsigned int src_data_63_32 : 32;
};
unsigned int DW_4_DATA;
} SRC_DATA_HI_UNION;
union {
struct {
unsigned int cmp_data_31_0 : 32;
};
unsigned int DW_5_DATA;
} CMP_DATA_LO_UNION;
union {
struct {
unsigned int cmp_data_63_32 : 32;
};
unsigned int DW_6_DATA;
} CMP_DATA_HI_UNION;
union {
struct {
unsigned int loop_interval : 13;
unsigned int reserved_0 : 19;
};
unsigned int DW_7_DATA;
} LOOP_UNION;
} SDMA_PKT_ATOMIC;
typedef struct SDMA_PKT_TIMESTAMP_TAG {
union {
struct {
unsigned int op : 8;
unsigned int sub_op : 8;
unsigned int reserved_0 : 16;
};
unsigned int DW_0_DATA;
} HEADER_UNION;
union {
struct {
unsigned int addr_31_0 : 32;
};
unsigned int DW_1_DATA;
} ADDR_LO_UNION;
union {
struct {
unsigned int addr_63_32 : 32;
};
unsigned int DW_2_DATA;
} ADDR_HI_UNION;
} SDMA_PKT_TIMESTAMP;
typedef struct SDMA_PKT_TRAP_TAG {
union {
struct {
unsigned int op : 8;
unsigned int sub_op : 8;
unsigned int reserved_0 : 16;
};
unsigned int DW_0_DATA;
} HEADER_UNION;
union {
struct {
unsigned int int_ctx : 28;
unsigned int reserved_1 : 4;
};
unsigned int DW_1_DATA;
} INT_CONTEXT_UNION;
} SDMA_PKT_TRAP;
// HDP flush packet, no parameters.
typedef struct SDMA_PKT_HDP_FLUSH_TAG {
unsigned int DW_0_DATA;
unsigned int DW_1_DATA;
unsigned int DW_2_DATA;
unsigned int DW_3_DATA;
unsigned int DW_4_DATA;
unsigned int DW_5_DATA;
// Version of gfx9 sDMA microcode introducing SDMA_PKT_HDP_FLUSH
static const uint16_t kMinVersion_ = 0x1A5;
} SDMA_PKT_HDP_FLUSH;
static const SDMA_PKT_HDP_FLUSH hdp_flush_cmd = {0x8, 0x0, 0x80000000, 0x0, 0x0, 0x0};
typedef struct SDMA_PKT_GCR_TAG {
union {
struct {
unsigned int op : 8;
unsigned int sub_op : 8;
unsigned int : 16;
};
unsigned int DW_0_DATA;
} HEADER_UNION;
union {
struct {
unsigned int : 7;
unsigned int BaseVA_LO : 25;
};
unsigned int DW_1_DATA;
} WORD1_UNION;
union {
struct {
unsigned int BaseVA_HI : 16;
unsigned int GCR_CONTROL_GLI_INV : 2;
unsigned int GCR_CONTROL_GL1_RANGE : 2;
unsigned int GCR_CONTROL_GLM_WB : 1;
unsigned int GCR_CONTROL_GLM_INV : 1;
unsigned int GCR_CONTROL_GLK_WB : 1;
unsigned int GCR_CONTROL_GLK_INV : 1;
unsigned int GCR_CONTROL_GLV_INV : 1;
unsigned int GCR_CONTROL_GL1_INV : 1;
unsigned int GCR_CONTROL_GL2_US : 1;
unsigned int GCR_CONTROL_GL2_RANGE : 2;
unsigned int GCR_CONTROL_GL2_DISCARD : 1;
unsigned int GCR_CONTROL_GL2_INV : 1;
unsigned int GCR_CONTROL_GL2_WB : 1;
};
unsigned int DW_2_DATA;
} WORD2_UNION;
union {
struct {
unsigned int GCR_CONTROL_RANGE_IS_PA : 1;
unsigned int GCR_CONTROL_SEQ : 2;
unsigned int : 4;
unsigned int LimitVA_LO : 25;
};
unsigned int DW_3_DATA;
} WORD3_UNION;
union {
struct {
unsigned int LimitVA_HI : 16;
unsigned int : 8;
unsigned int VMID : 4;
unsigned int : 4;
};
unsigned int DW_4_DATA;
} WORD4_UNION;
} SDMA_PKT_GCR;
} // namespace amd
} // namespace rocr
#endif // HSA_RUNTIME_CORE_INC_SDMA_REGISTERS_H_

5664
tinygrad_repo/extra/hip_gpu_driver/sdma_v6_0_0_pkt_open.h Normal file
View File

File diff suppressed because it is too large Load Diff

1166
tinygrad_repo/extra/hip_gpu_driver/sienna_cichlid_ip_offset.h Normal file
View File

File diff suppressed because it is too large Load Diff

444
tinygrad_repo/extra/hip_gpu_driver/soc15d.h Normal file
View File

@@ -0,0 +1,444 @@
/*
* Copyright 2014 Advanced Micro Devices, Inc.
*
* 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.
*
*/
#ifndef SOC15_H
#define SOC15_H
#define GFX9_NUM_GFX_RINGS 1
#define GFX9_NUM_COMPUTE_RINGS 8
/*
* PM4
*/
#define PACKET_TYPE0 0
#define PACKET_TYPE1 1
#define PACKET_TYPE2 2
#define PACKET_TYPE3 3
#define CP_PACKET_GET_TYPE(h) (((h) >> 30) & 3)
#define CP_PACKET_GET_COUNT(h) (((h) >> 16) & 0x3FFF)
#define CP_PACKET0_GET_REG(h) ((h) & 0xFFFF)
#define CP_PACKET3_GET_OPCODE(h) (((h) >> 8) & 0xFF)
#define PACKET0(reg, n) ((PACKET_TYPE0 << 30) | \
((reg) & 0xFFFF) | \
((n) & 0x3FFF) << 16)
#define CP_PACKET2 0x80000000
#define PACKET2_PAD_SHIFT 0
#define PACKET2_PAD_MASK (0x3fffffff << 0)
#define PACKET2(v) (CP_PACKET2 | REG_SET(PACKET2_PAD, (v)))
#define PACKET3(op, n) ((PACKET_TYPE3 << 30) | \
(((op) & 0xFF) << 8) | \
((n) & 0x3FFF) << 16)
#define PACKET3_COMPUTE(op, n) (PACKET3(op, n) | 1 << 1)
#define PACKETJ_CONDITION_CHECK0 0
#define PACKETJ_CONDITION_CHECK1 1
#define PACKETJ_CONDITION_CHECK2 2
#define PACKETJ_CONDITION_CHECK3 3
#define PACKETJ_CONDITION_CHECK4 4
#define PACKETJ_CONDITION_CHECK5 5
#define PACKETJ_CONDITION_CHECK6 6
#define PACKETJ_CONDITION_CHECK7 7
#define PACKETJ_TYPE0 0
#define PACKETJ_TYPE1 1
#define PACKETJ_TYPE2 2
#define PACKETJ_TYPE3 3
#define PACKETJ_TYPE4 4
#define PACKETJ_TYPE5 5
#define PACKETJ_TYPE6 6
#define PACKETJ_TYPE7 7
#define PACKETJ(reg, r, cond, type) ((reg & 0x3FFFF) | \
((r & 0x3F) << 18) | \
((cond & 0xF) << 24) | \
((type & 0xF) << 28))
#define CP_PACKETJ_NOP 0x60000000
#define CP_PACKETJ_GET_REG(x) ((x) & 0x3FFFF)
#define CP_PACKETJ_GET_RES(x) (((x) >> 18) & 0x3F)
#define CP_PACKETJ_GET_COND(x) (((x) >> 24) & 0xF)
#define CP_PACKETJ_GET_TYPE(x) (((x) >> 28) & 0xF)
/* Packet 3 types */
#define PACKET3_NOP 0x10
#define PACKET3_SET_BASE 0x11
#define PACKET3_BASE_INDEX(x) ((x) << 0)
#define CE_PARTITION_BASE 3
#define PACKET3_CLEAR_STATE 0x12
#define PACKET3_INDEX_BUFFER_SIZE 0x13
#define PACKET3_DISPATCH_DIRECT 0x15
#define PACKET3_DISPATCH_INDIRECT 0x16
#define PACKET3_ATOMIC_GDS 0x1D
#define PACKET3_ATOMIC_MEM 0x1E
#define PACKET3_OCCLUSION_QUERY 0x1F
#define PACKET3_SET_PREDICATION 0x20
#define PACKET3_REG_RMW 0x21
#define PACKET3_COND_EXEC 0x22
#define PACKET3_PRED_EXEC 0x23
#define PACKET3_DRAW_INDIRECT 0x24
#define PACKET3_DRAW_INDEX_INDIRECT 0x25
#define PACKET3_INDEX_BASE 0x26
#define PACKET3_DRAW_INDEX_2 0x27
#define PACKET3_CONTEXT_CONTROL 0x28
#define PACKET3_INDEX_TYPE 0x2A
#define PACKET3_DRAW_INDIRECT_MULTI 0x2C
#define PACKET3_DRAW_INDEX_AUTO 0x2D
#define PACKET3_NUM_INSTANCES 0x2F
#define PACKET3_DRAW_INDEX_MULTI_AUTO 0x30
#define PACKET3_INDIRECT_BUFFER_CONST 0x33
#define PACKET3_STRMOUT_BUFFER_UPDATE 0x34
#define PACKET3_DRAW_INDEX_OFFSET_2 0x35
#define PACKET3_DRAW_PREAMBLE 0x36
#define PACKET3_WRITE_DATA 0x37
#define WRITE_DATA_DST_SEL(x) ((x) << 8)
/* 0 - register
* 1 - memory (sync - via GRBM)
* 2 - gl2
* 3 - gds
* 4 - reserved
* 5 - memory (async - direct)
*/
#define WR_ONE_ADDR (1 << 16)
#define WR_CONFIRM (1 << 20)
#define WRITE_DATA_CACHE_POLICY(x) ((x) << 25)
/* 0 - LRU
* 1 - Stream
*/
#define WRITE_DATA_ENGINE_SEL(x) ((x) << 30)
/* 0 - me
* 1 - pfp
* 2 - ce
*/
#define PACKET3_DRAW_INDEX_INDIRECT_MULTI 0x38
#define PACKET3_MEM_SEMAPHORE 0x39
# define PACKET3_SEM_USE_MAILBOX (0x1 << 16)
# define PACKET3_SEM_SEL_SIGNAL_TYPE (0x1 << 20) /* 0 = increment, 1 = write 1 */
# define PACKET3_SEM_SEL_SIGNAL (0x6 << 29)
# define PACKET3_SEM_SEL_WAIT (0x7 << 29)
#define PACKET3_WAIT_REG_MEM 0x3C
#define WAIT_REG_MEM_FUNCTION(x) ((x) << 0)
/* 0 - always
* 1 - <
* 2 - <=
* 3 - ==
* 4 - !=
* 5 - >=
* 6 - >
*/
#define WAIT_REG_MEM_MEM_SPACE(x) ((x) << 4)
/* 0 - reg
* 1 - mem
*/
#define WAIT_REG_MEM_OPERATION(x) ((x) << 6)
/* 0 - wait_reg_mem
* 1 - wr_wait_wr_reg
*/
#define WAIT_REG_MEM_ENGINE(x) ((x) << 8)
/* 0 - me
* 1 - pfp
*/
#define PACKET3_INDIRECT_BUFFER 0x3F
#define INDIRECT_BUFFER_VALID (1 << 23)
#define INDIRECT_BUFFER_CACHE_POLICY(x) ((x) << 28)
/* 0 - LRU
* 1 - Stream
* 2 - Bypass
*/
#define INDIRECT_BUFFER_PRE_ENB(x) ((x) << 21)
#define INDIRECT_BUFFER_PRE_RESUME(x) ((x) << 30)
#define PACKET3_COPY_DATA 0x40
#define PACKET3_PFP_SYNC_ME 0x42
#define PACKET3_COND_WRITE 0x45
#define PACKET3_EVENT_WRITE 0x46
#define EVENT_TYPE(x) ((x) << 0)
#define EVENT_INDEX(x) ((x) << 8)
/* 0 - any non-TS event
* 1 - ZPASS_DONE, PIXEL_PIPE_STAT_*
* 2 - SAMPLE_PIPELINESTAT
* 3 - SAMPLE_STREAMOUTSTAT*
* 4 - *S_PARTIAL_FLUSH
*/
#define PACKET3_RELEASE_MEM 0x49
#define EVENT_TYPE(x) ((x) << 0)
#define EVENT_INDEX(x) ((x) << 8)
#define EOP_TCL1_VOL_ACTION_EN (1 << 12)
#define EOP_TC_VOL_ACTION_EN (1 << 13) /* L2 */
#define EOP_TC_WB_ACTION_EN (1 << 15) /* L2 */
#define EOP_TCL1_ACTION_EN (1 << 16)
#define EOP_TC_ACTION_EN (1 << 17) /* L2 */
#define EOP_TC_NC_ACTION_EN (1 << 19)
#define EOP_TC_MD_ACTION_EN (1 << 21) /* L2 metadata */
#define EOP_EXEC (1 << 28) /* For Trailing Fence */
#define DATA_SEL(x) ((x) << 29)
/* 0 - discard
* 1 - send low 32bit data
* 2 - send 64bit data
* 3 - send 64bit GPU counter value
* 4 - send 64bit sys counter value
*/
#define INT_SEL(x) ((x) << 24)
/* 0 - none
* 1 - interrupt only (DATA_SEL = 0)
* 2 - interrupt when data write is confirmed
*/
#define DST_SEL(x) ((x) << 16)
/* 0 - MC
* 1 - TC/L2
*/
#define PACKET3_PREAMBLE_CNTL 0x4A
# define PACKET3_PREAMBLE_BEGIN_CLEAR_STATE (2 << 28)
# define PACKET3_PREAMBLE_END_CLEAR_STATE (3 << 28)
#define PACKET3_DMA_DATA 0x50
/* 1. header
* 2. CONTROL
* 3. SRC_ADDR_LO or DATA [31:0]
* 4. SRC_ADDR_HI [31:0]
* 5. DST_ADDR_LO [31:0]
* 6. DST_ADDR_HI [7:0]
* 7. COMMAND [30:21] | BYTE_COUNT [20:0]
*/
/* CONTROL */
# define PACKET3_DMA_DATA_ENGINE(x) ((x) << 0)
/* 0 - ME
* 1 - PFP
*/
# define PACKET3_DMA_DATA_SRC_CACHE_POLICY(x) ((x) << 13)
/* 0 - LRU
* 1 - Stream
*/
# define PACKET3_DMA_DATA_DST_SEL(x) ((x) << 20)
/* 0 - DST_ADDR using DAS
* 1 - GDS
* 3 - DST_ADDR using L2
*/
# define PACKET3_DMA_DATA_DST_CACHE_POLICY(x) ((x) << 25)
/* 0 - LRU
* 1 - Stream
*/
# define PACKET3_DMA_DATA_SRC_SEL(x) ((x) << 29)
/* 0 - SRC_ADDR using SAS
* 1 - GDS
* 2 - DATA
* 3 - SRC_ADDR using L2
*/
# define PACKET3_DMA_DATA_CP_SYNC (1 << 31)
/* COMMAND */
# define PACKET3_DMA_DATA_CMD_SAS (1 << 26)
/* 0 - memory
* 1 - register
*/
# define PACKET3_DMA_DATA_CMD_DAS (1 << 27)
/* 0 - memory
* 1 - register
*/
# define PACKET3_DMA_DATA_CMD_SAIC (1 << 28)
# define PACKET3_DMA_DATA_CMD_DAIC (1 << 29)
# define PACKET3_DMA_DATA_CMD_RAW_WAIT (1 << 30)
#define PACKET3_ACQUIRE_MEM 0x58
/* 1. HEADER
* 2. COHER_CNTL [30:0]
* 2.1 ENGINE_SEL [31:31]
* 3. COHER_SIZE [31:0]
* 4. COHER_SIZE_HI [7:0]
* 5. COHER_BASE_LO [31:0]
* 6. COHER_BASE_HI [23:0]
* 7. POLL_INTERVAL [15:0]
*/
/* COHER_CNTL fields for CP_COHER_CNTL */
#define PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_TC_NC_ACTION_ENA(x) ((x) << 3)
#define PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_TC_WC_ACTION_ENA(x) ((x) << 4)
#define PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_TC_INV_METADATA_ACTION_ENA(x) ((x) << 5)
#define PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_TCL1_VOL_ACTION_ENA(x) ((x) << 15)
#define PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_TC_WB_ACTION_ENA(x) ((x) << 18)
#define PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_TCL1_ACTION_ENA(x) ((x) << 22)
#define PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_TC_ACTION_ENA(x) ((x) << 23)
#define PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_CB_ACTION_ENA(x) ((x) << 25)
#define PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_DB_ACTION_ENA(x) ((x) << 26)
#define PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_SH_KCACHE_ACTION_ENA(x) ((x) << 27)
#define PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_SH_KCACHE_VOL_ACTION_ENA(x) ((x) << 28)
#define PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_SH_ICACHE_ACTION_ENA(x) ((x) << 29)
#define PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_SH_KCACHE_WB_ACTION_ENA(x) ((x) << 30)
#define PACKET3_REWIND 0x59
#define PACKET3_LOAD_UCONFIG_REG 0x5E
#define PACKET3_LOAD_SH_REG 0x5F
#define PACKET3_LOAD_CONFIG_REG 0x60
#define PACKET3_LOAD_CONTEXT_REG 0x61
#define PACKET3_SET_CONFIG_REG 0x68
#define PACKET3_SET_CONFIG_REG_START 0x00002000
#define PACKET3_SET_CONFIG_REG_END 0x00002c00
#define PACKET3_SET_CONTEXT_REG 0x69
#define PACKET3_SET_CONTEXT_REG_START 0x0000a000
#define PACKET3_SET_CONTEXT_REG_END 0x0000a400
#define PACKET3_SET_CONTEXT_REG_INDIRECT 0x73
#define PACKET3_SET_SH_REG 0x76
#define PACKET3_SET_SH_REG_START 0x00002c00
#define PACKET3_SET_SH_REG_END 0x00003000
#define PACKET3_SET_SH_REG_OFFSET 0x77
#define PACKET3_SET_QUEUE_REG 0x78
#define PACKET3_SET_UCONFIG_REG 0x79
#define PACKET3_SET_UCONFIG_REG_START 0x0000c000
#define PACKET3_SET_UCONFIG_REG_END 0x0000c400
#define PACKET3_SET_UCONFIG_REG_INDEX_TYPE (2 << 28)
#define PACKET3_SCRATCH_RAM_WRITE 0x7D
#define PACKET3_SCRATCH_RAM_READ 0x7E
#define PACKET3_LOAD_CONST_RAM 0x80
#define PACKET3_WRITE_CONST_RAM 0x81
#define PACKET3_DUMP_CONST_RAM 0x83
#define PACKET3_INCREMENT_CE_COUNTER 0x84
#define PACKET3_INCREMENT_DE_COUNTER 0x85
#define PACKET3_WAIT_ON_CE_COUNTER 0x86
#define PACKET3_WAIT_ON_DE_COUNTER_DIFF 0x88
#define PACKET3_SWITCH_BUFFER 0x8B
#define PACKET3_FRAME_CONTROL 0x90
# define FRAME_TMZ (1 << 0)
# define FRAME_CMD(x) ((x) << 28)
/*
* x=0: tmz_begin
* x=1: tmz_end
*/
#define PACKET3_INVALIDATE_TLBS 0x98
# define PACKET3_INVALIDATE_TLBS_DST_SEL(x) ((x) << 0)
# define PACKET3_INVALIDATE_TLBS_ALL_HUB(x) ((x) << 4)
# define PACKET3_INVALIDATE_TLBS_PASID(x) ((x) << 5)
# define PACKET3_INVALIDATE_TLBS_FLUSH_TYPE(x) ((x) << 29)
#define PACKET3_SET_RESOURCES 0xA0
/* 1. header
* 2. CONTROL
* 3. QUEUE_MASK_LO [31:0]
* 4. QUEUE_MASK_HI [31:0]
* 5. GWS_MASK_LO [31:0]
* 6. GWS_MASK_HI [31:0]
* 7. OAC_MASK [15:0]
* 8. GDS_HEAP_SIZE [16:11] | GDS_HEAP_BASE [5:0]
*/
# define PACKET3_SET_RESOURCES_VMID_MASK(x) ((x) << 0)
# define PACKET3_SET_RESOURCES_UNMAP_LATENTY(x) ((x) << 16)
# define PACKET3_SET_RESOURCES_QUEUE_TYPE(x) ((x) << 29)
#define PACKET3_MAP_QUEUES 0xA2
/* 1. header
* 2. CONTROL
* 3. CONTROL2
* 4. MQD_ADDR_LO [31:0]
* 5. MQD_ADDR_HI [31:0]
* 6. WPTR_ADDR_LO [31:0]
* 7. WPTR_ADDR_HI [31:0]
*/
/* CONTROL */
# define PACKET3_MAP_QUEUES_QUEUE_SEL(x) ((x) << 4)
# define PACKET3_MAP_QUEUES_VMID(x) ((x) << 8)
# define PACKET3_MAP_QUEUES_QUEUE(x) ((x) << 13)
# define PACKET3_MAP_QUEUES_PIPE(x) ((x) << 16)
# define PACKET3_MAP_QUEUES_ME(x) ((x) << 18)
# define PACKET3_MAP_QUEUES_QUEUE_TYPE(x) ((x) << 21)
# define PACKET3_MAP_QUEUES_ALLOC_FORMAT(x) ((x) << 24)
# define PACKET3_MAP_QUEUES_ENGINE_SEL(x) ((x) << 26)
# define PACKET3_MAP_QUEUES_NUM_QUEUES(x) ((x) << 29)
/* CONTROL2 */
# define PACKET3_MAP_QUEUES_CHECK_DISABLE(x) ((x) << 1)
# define PACKET3_MAP_QUEUES_DOORBELL_OFFSET(x) ((x) << 2)
#define PACKET3_UNMAP_QUEUES 0xA3
/* 1. header
* 2. CONTROL
* 3. CONTROL2
* 4. CONTROL3
* 5. CONTROL4
* 6. CONTROL5
*/
/* CONTROL */
# define PACKET3_UNMAP_QUEUES_ACTION(x) ((x) << 0)
/* 0 - PREEMPT_QUEUES
* 1 - RESET_QUEUES
* 2 - DISABLE_PROCESS_QUEUES
* 3 - PREEMPT_QUEUES_NO_UNMAP
*/
# define PACKET3_UNMAP_QUEUES_QUEUE_SEL(x) ((x) << 4)
# define PACKET3_UNMAP_QUEUES_ENGINE_SEL(x) ((x) << 26)
# define PACKET3_UNMAP_QUEUES_NUM_QUEUES(x) ((x) << 29)
/* CONTROL2a */
# define PACKET3_UNMAP_QUEUES_PASID(x) ((x) << 0)
/* CONTROL2b */
# define PACKET3_UNMAP_QUEUES_DOORBELL_OFFSET0(x) ((x) << 2)
/* CONTROL3a */
# define PACKET3_UNMAP_QUEUES_DOORBELL_OFFSET1(x) ((x) << 2)
/* CONTROL3b */
# define PACKET3_UNMAP_QUEUES_RB_WPTR(x) ((x) << 0)
/* CONTROL4 */
# define PACKET3_UNMAP_QUEUES_DOORBELL_OFFSET2(x) ((x) << 2)
/* CONTROL5 */
# define PACKET3_UNMAP_QUEUES_DOORBELL_OFFSET3(x) ((x) << 2)
#define PACKET3_QUERY_STATUS 0xA4
/* 1. header
* 2. CONTROL
* 3. CONTROL2
* 4. ADDR_LO [31:0]
* 5. ADDR_HI [31:0]
* 6. DATA_LO [31:0]
* 7. DATA_HI [31:0]
*/
/* CONTROL */
# define PACKET3_QUERY_STATUS_CONTEXT_ID(x) ((x) << 0)
# define PACKET3_QUERY_STATUS_INTERRUPT_SEL(x) ((x) << 28)
# define PACKET3_QUERY_STATUS_COMMAND(x) ((x) << 30)
/* CONTROL2a */
# define PACKET3_QUERY_STATUS_PASID(x) ((x) << 0)
/* CONTROL2b */
# define PACKET3_QUERY_STATUS_DOORBELL_OFFSET(x) ((x) << 2)
# define PACKET3_QUERY_STATUS_ENG_SEL(x) ((x) << 25)
#define PACKET3_RUN_CLEANER_SHADER 0xD2
/* 1. header
* 2. RESERVED [31:0]
*/
#define VCE_CMD_NO_OP 0x00000000
#define VCE_CMD_END 0x00000001
#define VCE_CMD_IB 0x00000002
#define VCE_CMD_FENCE 0x00000003
#define VCE_CMD_TRAP 0x00000004
#define VCE_CMD_IB_AUTO 0x00000005
#define VCE_CMD_SEMAPHORE 0x00000006
#define VCE_CMD_IB_VM 0x00000102
#define VCE_CMD_WAIT_GE 0x00000106
#define VCE_CMD_UPDATE_PTB 0x00000107
#define VCE_CMD_FLUSH_TLB 0x00000108
#define VCE_CMD_REG_WRITE 0x00000109
#define VCE_CMD_REG_WAIT 0x0000010a
#define HEVC_ENC_CMD_NO_OP 0x00000000
#define HEVC_ENC_CMD_END 0x00000001
#define HEVC_ENC_CMD_FENCE 0x00000003
#define HEVC_ENC_CMD_TRAP 0x00000004
#define HEVC_ENC_CMD_IB_VM 0x00000102
#define HEVC_ENC_CMD_REG_WRITE 0x00000109
#define HEVC_ENC_CMD_REG_WAIT 0x0000010a
#endif

22477
tinygrad_repo/extra/hip_gpu_driver/soc21_enum.h Normal file
View File

File diff suppressed because it is too large Load Diff

207
tinygrad_repo/extra/hip_gpu_driver/test_kfd_2.py Normal file
View File

@@ -0,0 +1,207 @@
import os, ctypes, pathlib, re, fcntl, functools, mmap, time
import tinygrad.runtime.autogen.kfd as kfd
from tinygrad.helpers import to_mv, getenv
from extra.hip_gpu_driver import hip_ioctl
import tinygrad.runtime.autogen.hsa as hsa
from hexdump import hexdump
libc = ctypes.CDLL("libc.so.6")
libc.memset.argtypes = [ctypes.c_void_p, ctypes.c_char, ctypes.c_int]
libc.mmap.argtypes = [ctypes.c_void_p, ctypes.c_size_t, ctypes.c_int, ctypes.c_int, ctypes.c_int, ctypes.c_long]
libc.mmap.restype = ctypes.c_void_p
MAP_NORESERVE = 0x4000
MAP_FIXED = 0x10
def kfd_ioctl(idir, nr, user_struct, fd, **kwargs):
made = user_struct(**kwargs)
ret = fcntl.ioctl(fd, (idir<<30) | (ctypes.sizeof(user_struct)<<16) | (ord('K')<<8) | nr, made)
if ret != 0: raise RuntimeError(f"ioctl returned {ret}")
return made
def format_struct(s):
sdats = []
for field_name, field_type in s._fields_:
dat = getattr(s, field_name)
if isinstance(dat, int): sdats.append(f"{field_name}:0x{dat:X}")
else: sdats.append(f"{field_name}:{dat}")
return sdats
idirs = {"IOW": 1, "IOR": 2, "IOWR": 3}
def ioctls_from_header():
hdr = pathlib.Path("/usr/include/linux/kfd_ioctl.h").read_text().replace("\\\n", "")
pattern = r'#define\s+(AMDKFD_IOC_[A-Z0-9_]+)\s+AMDKFD_(IOW?R?)\((0x[0-9a-fA-F]+),\s+struct\s([A-Za-z0-9_]+)\)'
matches = re.findall(pattern, hdr, re.MULTILINE)
fxns = {}
for name, idir, nr, sname in matches:
fxns[name.replace("AMDKFD_IOC_", "").lower()] = functools.partial(kfd_ioctl, idirs[idir], int(nr, 0x10), getattr(kfd, "struct_"+sname))
return type("KIO", (object, ), fxns)
kio = ioctls_from_header()
# sudo su -c "echo 'file drivers/gpu/drm/amd/* +p' > /sys/kernel/debug/dynamic_debug/control"
def gpu_alloc_userptr(fd, size, flags):
addr = libc.mmap(0, size, mmap.PROT_READ|mmap.PROT_WRITE, mmap.MAP_SHARED|mmap.MAP_ANONYMOUS, -1, 0)
assert addr != 0xffffffffffffffff
mem = kio.alloc_memory_of_gpu(fd, va_addr=addr, size=size, gpu_id=GPU_ID, flags=flags, mmap_offset=addr)
return mem
def gpu_alloc(fd, size, flags):
addr = libc.mmap(0, size, 0, mmap.MAP_PRIVATE|mmap.MAP_ANONYMOUS|MAP_NORESERVE, -1, 0)
assert addr != 0xffffffffffffffff
mem = kio.alloc_memory_of_gpu(fd, va_addr=addr, size=size, gpu_id=GPU_ID, flags=flags)
buf = libc.mmap(mem.va_addr, mem.size, mmap.PROT_READ|mmap.PROT_WRITE, mmap.MAP_SHARED|MAP_FIXED, drm_fd, mem.mmap_offset)
assert buf != 0xffffffffffffffff
assert addr == buf == mem.va_addr
return mem
if __name__ == "__main__":
fd = os.open("/dev/kfd", os.O_RDWR)
gpu_num = getenv("GPU", 0)
drm_fd = os.open(f"/dev/dri/renderD{128+gpu_num}", os.O_RDWR)
with open(f"/sys/devices/virtual/kfd/kfd/topology/nodes/{1+gpu_num}/gpu_id", "r") as f: GPU_ID = int(f.read())
#ver = kio.get_version(fd)
st = kio.acquire_vm(fd, drm_fd=drm_fd, gpu_id=GPU_ID)
#exit(0)
# 0xF0000001 = KFD_IOC_ALLOC_MEM_FLAGS_VRAM | KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE | KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE | KFD_IOC_ALLOC_MEM_FLAGS_PUBLIC | KFD_IOC_ALLOC_MEM_FLAGS_NO_SUBSTITUTE
# 0xD6000002 = KFD_IOC_ALLOC_MEM_FLAGS_GTT | KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE | KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE | KFD_IOC_ALLOC_MEM_FLAGS_NO_SUBSTITUTE
# 0xD6000004 = KFD_IOC_ALLOC_MEM_FLAGS_USERPTR | KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE | KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE | KFD_IOC_ALLOC_MEM_FLAGS_NO_SUBSTITUTE
# 0x94000010 = KFD_IOC_ALLOC_MEM_FLAGS_MMIO_REMAP | KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE | KFD_IOC_ALLOC_MEM_FLAGS_NO_SUBSTITUTE
#addr = libc.mmap(0, 0x1000, mmap.PROT_READ|mmap.PROT_WRITE, mmap.MAP_PRIVATE|mmap.MAP_ANONYMOUS, -1, 0)
#addr = libc.mmap(0, 0x1000, mmap.PROT_READ|mmap.PROT_WRITE, mmap.MAP_SHARED|mmap.MAP_ANONYMOUS, -1, 0)
#mem = kio.AMDKFD_IOC_ALLOC_MEMORY_OF_GPU(fd, va_addr=addr, size=0x1000, gpu_id=GPU_ID, flags=0xD6000004)
#mem = gpu_alloc(fd, 0x1000, kfd.KFD_IOC_ALLOC_MEM_FLAGS_VRAM |
# kfd.KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE | kfd.KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE |
# kfd.KFD_IOC_ALLOC_MEM_FLAGS_PUBLIC | kfd.KFD_IOC_ALLOC_MEM_FLAGS_NO_SUBSTITUTE)
#arr = (ctypes.c_int32 * 1)(GPU_ID)
#stm = kio.map_memory_to_gpu(fd, handle=mem.handle, device_ids_array_ptr=ctypes.addressof(arr), n_devices=1)
arr = (ctypes.c_int32 * 1)(GPU_ID)
rw_ptr = gpu_alloc(fd, 0x1000, kfd.KFD_IOC_ALLOC_MEM_FLAGS_GTT | kfd.KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE |
kfd.KFD_IOC_ALLOC_MEM_FLAGS_COHERENT | kfd.KFD_IOC_ALLOC_MEM_FLAGS_UNCACHED |
kfd.KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE | kfd.KFD_IOC_ALLOC_MEM_FLAGS_NO_SUBSTITUTE)
stm = kio.map_memory_to_gpu(fd, handle=rw_ptr.handle, device_ids_array_ptr=ctypes.addressof(arr), n_devices=1)
assert stm.n_success == 1
event_page = gpu_alloc(fd, 0x8000, kfd.KFD_IOC_ALLOC_MEM_FLAGS_GTT | kfd.KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE |
kfd.KFD_IOC_ALLOC_MEM_FLAGS_COHERENT | kfd.KFD_IOC_ALLOC_MEM_FLAGS_UNCACHED |
kfd.KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE | kfd.KFD_IOC_ALLOC_MEM_FLAGS_NO_SUBSTITUTE)
stm = kio.map_memory_to_gpu(fd, handle=event_page.handle, device_ids_array_ptr=ctypes.addressof(arr), n_devices=1)
assert stm.n_success == 1
ring_base = gpu_alloc_userptr(fd, 0x1000, kfd.KFD_IOC_ALLOC_MEM_FLAGS_USERPTR | kfd.KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE |
kfd.KFD_IOC_ALLOC_MEM_FLAGS_COHERENT | kfd.KFD_IOC_ALLOC_MEM_FLAGS_UNCACHED |
kfd.KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE | kfd.KFD_IOC_ALLOC_MEM_FLAGS_NO_SUBSTITUTE)
stm = kio.map_memory_to_gpu(fd, handle=ring_base.handle, device_ids_array_ptr=ctypes.addressof(arr), n_devices=1)
assert stm.n_success == 1
signals = gpu_alloc_userptr(fd, 0x1000, kfd.KFD_IOC_ALLOC_MEM_FLAGS_USERPTR | kfd.KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE |
kfd.KFD_IOC_ALLOC_MEM_FLAGS_COHERENT | kfd.KFD_IOC_ALLOC_MEM_FLAGS_UNCACHED |
kfd.KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE | kfd.KFD_IOC_ALLOC_MEM_FLAGS_NO_SUBSTITUTE)
stm = kio.map_memory_to_gpu(fd, handle=signals.handle, device_ids_array_ptr=ctypes.addressof(arr), n_devices=1)
assert stm.n_success == 1
eop_buffer = gpu_alloc(fd, 0x1000, kfd.KFD_IOC_ALLOC_MEM_FLAGS_VRAM |
kfd.KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE | kfd.KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE |
kfd.KFD_IOC_ALLOC_MEM_FLAGS_NO_SUBSTITUTE)
stm = kio.map_memory_to_gpu(fd, handle=eop_buffer.handle, device_ids_array_ptr=ctypes.addressof(arr), n_devices=1)
assert stm.n_success == 1
ctx_save_restore_address = gpu_alloc(fd, 0x2C02000, kfd.KFD_IOC_ALLOC_MEM_FLAGS_VRAM |
kfd.KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE | kfd.KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE |
kfd.KFD_IOC_ALLOC_MEM_FLAGS_NO_SUBSTITUTE)
stm = kio.map_memory_to_gpu(fd, handle=ctx_save_restore_address.handle, device_ids_array_ptr=ctypes.addressof(arr), n_devices=1)
assert stm.n_success == 1
#113.00 ms + 0.00 ms : 0 = AMDKFD_IOC_CREATE_QUEUE ring_base_address:0x797465200000 write_pointer_address:0x79751C068038 read_pointer_address:0x79751C068080 doorbell_offset:0x0 ring_size:0x800000 gpu_id:0x433D queue_type:0x2 queue_per
#centage:0x64 queue_priority:0x7 queue_id:0x0 eop_buffer_address:0x79751C064000 eop_buffer_size:0x1000 ctx_save_restore_address:0x796E52400000 ctx_save_restore_size:0x2BEA000 ctl_stack_size:0xA000
#113.84 ms + 0.59 ms : 0 = AMDKFD_IOC_CREATE_QUEUE ring_base_address:0x71AC3F600000 write_pointer_address:0x71B302AB0038 read_pointer_address:0x71B302AB0080 doorbell_offset:0xD0CF400000000008 ring_size:0x800000 gpu_id:0x433D queue_typ
#e:0x2 queue_percentage:0x64 queue_priority:0x7 queue_id:0x1 eop_buffer_address:0x71B302AAC000 eop_buffer_size:0x1000 ctx_save_restore_address:0x71AC3C800000 ctx_save_restore_size:0x2BEA000 ctl_stack_size:0xA000
#define KFD_MMAP_TYPE_SHIFT 62
#define KFD_MMAP_TYPE_DOORBELL (0x3ULL << KFD_MMAP_TYPE_SHIFT)
evt = kio.create_event(fd, event_page_offset=event_page.handle, auto_reset=1)
nq = kio.create_queue(fd, ring_base_address=ring_base.va_addr, ring_size=0x1000, gpu_id=GPU_ID,
queue_type=kfd.KFD_IOC_QUEUE_TYPE_COMPUTE_AQL, queue_percentage=kfd.KFD_MAX_QUEUE_PERCENTAGE,
queue_priority=kfd.KFD_MAX_QUEUE_PRIORITY,
eop_buffer_address=eop_buffer.va_addr, eop_buffer_size=0x1000,
ctx_save_restore_address=ctx_save_restore_address.va_addr, ctx_save_restore_size=0x2C02000,
ctl_stack_size = 0xa000,
# write_pointer_address and read_pointer_address are on GART
#write_pointer_address=0xaaaabbbb, read_pointer_address=0xaaaacccc)
write_pointer_address=rw_ptr.va_addr+0, read_pointer_address=rw_ptr.va_addr+0x8)
doorbell = libc.mmap(0, 8192, mmap.PROT_READ|mmap.PROT_WRITE, mmap.MAP_SHARED, fd, nq.doorbell_offset)
print("doorbell", hex(doorbell))
to_mv(signals.va_addr, 0x40)
"""
hexdump(to_mv(event_page.va_addr, 0x40))
kio.set_event(fd, event_id=evt.event_id)
hexdump(to_mv(event_page.va_addr, 0x40))
kio.reset_event(fd, event_id=evt.event_id)
hexdump(to_mv(event_page.va_addr, 0x40))
"""
# KFD_EVENT_TYPE_SIGNAL
BARRIER_HEADER = 1 << hsa.HSA_PACKET_HEADER_BARRIER
BARRIER_HEADER |= hsa.HSA_FENCE_SCOPE_SYSTEM << hsa.HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE
BARRIER_HEADER |= hsa.HSA_FENCE_SCOPE_SYSTEM << hsa.HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE
BARRIER_HEADER |= hsa.HSA_PACKET_TYPE_BARRIER_AND << hsa.HSA_PACKET_HEADER_TYPE
AQL_PACKET_SIZE = ctypes.sizeof(hsa.hsa_kernel_dispatch_packet_t)
EMPTY_SIGNAL = hsa.hsa_signal_t()
ds = to_mv(rw_ptr.va_addr, 0x100).cast("Q")
ds[0] = 1 #ring_base.va_addr + AQL_PACKET_SIZE
ds[1] = 0 #ring_base.va_addr
#libc.memset(rw_ptr.va_addr, 0xaa, 0x100)
#hexdump(to_mv(rw_ptr.va_addr, 0x100))
#packet = hsa.hsa_barrier_and_packet_t.from_address(rw_ptr.va_addr+0x38)
packet = hsa.hsa_barrier_and_packet_t.from_address(ring_base.va_addr)
packet.reserved0 = 0
packet.reserved1 = 0
for i in range(5): packet.dep_signal[i] = EMPTY_SIGNAL
#packet.dep_signal[0] = hsa.hsa_signal_t(evt.event_id)
packet.reserved2 = 0
#packet.completion_signal = EMPTY_SIGNAL
packet.completion_signal = hsa.hsa_signal_t(signals.va_addr)
packet.header = BARRIER_HEADER
hexdump(to_mv(ring_base.va_addr, AQL_PACKET_SIZE))
# _HsaEventData
to_mv(signals.va_addr, 0x40).cast("Q")[0] = 1
to_mv(signals.va_addr, 0x40).cast("Q")[1] = 1
#to_mv(signals.va_addr, 0x40).cast("Q")[2] = event_page
to_mv(signals.va_addr, 0x40).cast("Q")[2] = event_page.va_addr + evt.event_slot_index*8 # HWData2=HWAddress
to_mv(signals.va_addr, 0x40).cast("Q")[3] = evt.event_trigger_data # HWData3=HWData
print(hex(ds[0]), hex(ds[1]), hex(ds[2]))
hexdump(to_mv(signals.va_addr, 0x40))
# 10 08 49 3E 46 77 00 00
# ring doorbell
print(hex(to_mv(doorbell, 0x10).cast("I")[0]))
#to_mv(doorbell, 0x10).cast("I")[0] = 0xffffffff
to_mv(doorbell, 0x10).cast("I")[0] = 0
evt_arr = (kfd.struct_kfd_event_data * 1)()
evt_arr[0].event_id = evt.event_id
kio.wait_events(fd, events_ptr=ctypes.addressof(evt_arr), num_events=1, wait_for_all=0, timeout=1000)
print(hex(ds[0]), hex(ds[1]), hex(ds[2]))
hexdump(to_mv(signals.va_addr, 0x40))
#nq = kio.create_queue(fd, ring_base_address=buf, ring_size=0x1000, gpu_id=GPU_ID,
# queue_type=kfd.KFD_IOC_QUEUE_TYPE_COMPUTE_AQL, queue_percentage=kfd.KFD_MAX_QUEUE_PERCENTAGE,
# queue_priority=kfd.KFD_MAX_QUEUE_PRIORITY, write_pointer_address=buf+8, read_pointer_address=buf+0x10)
#print(nq)
#mv = to_mv(buf, 0x1000)
#addr = libc.mmap(0, 0x1000, mmap.PROT_READ|mmap.PROT_WRITE, mmap.MAP_PRIVATE|mmap.MAP_ANONYMOUS, -1, 0)
#print('\n'.join(format_struct(ver)))
#print('\n'.join(format_struct(st)))

156
tinygrad_repo/extra/hip_gpu_driver/test_pm4.py Normal file
View File

@@ -0,0 +1,156 @@
import time
from hexdump import hexdump
from tinygrad import Tensor, Device
import tinygrad.runtime.autogen.amd_gpu as amd_gpu
import tinygrad.runtime.autogen.kfd as kfd
import tinygrad.runtime.autogen.hsa as hsa
from tinygrad.runtime.ops_amd import kio, AMDProgram
from tinygrad.helpers import to_mv
DISPATCH_INIT_VALUE = 0x21 | 0x8000
#mmCOMPUTE_START_X = 0x2e04
#mmCOMPUTE_PGM_LO = 0x2e0c
BASE_ADDR = 0x00001260
PACKET3_SET_SH_REG_START = 0x2c00
SUB = PACKET3_SET_SH_REG_START - BASE_ADDR
regCOMPUTE_PGM_LO = 0x1bac - SUB
regCOMPUTE_START_X = 0x1ba4 - SUB
regCOMPUTE_NUM_THREAD_X = 0x1ba7 - SUB
regCOMPUTE_USER_DATA_0 = 0x1be0 - SUB
regCOMPUTE_USER_DATA_8 = 0x1be8 - SUB
regCOMPUTE_PGM_RSRC1 = 0x1bb2 - SUB
regCOMPUTE_PGM_RSRC2 = 0x1bb3 - SUB
# DEBUG=6 python3 extra/hip_gpu_driver/test_pm4.py
# sudo umr -i 1 -s amd744c.gfx1100 --sbank 1 1 2 | grep regCOMPUTE
# 0x00009025
COMPUTE_SHADER_EN = 1
USE_THREAD_DIMENSIONS = 1 << 5
CS_W32_EN = 1 << 15
def format_struct(s):
sdats = []
for field_name, field_type in s._fields_:
dat = getattr(s, field_name)
if isinstance(dat, int): sdats.append(f"{field_name}:0x{dat:X}")
else: sdats.append(f"{field_name}:{dat}")
return sdats
if __name__ == "__main__":
dev = Device["KFD"]
a = Tensor([0.,1.,2.], device="KFD").realize()
b = a + 7
b.uop.buffer.allocate()
si = b.schedule()[-1]
runner = dev.get_runner(*si.ast)
prg: AMDProgram = runner.clprg
print("device initted")
# Compute Queue
gart_compute = dev._gpu_alloc(0x1000, kfd.KFD_IOC_ALLOC_MEM_FLAGS_GTT, uncached=True)
eop_buffer = dev._gpu_alloc(0x1000, kfd.KFD_IOC_ALLOC_MEM_FLAGS_VRAM)
compute_ring = dev._gpu_alloc(0x800000, kfd.KFD_IOC_ALLOC_MEM_FLAGS_GTT, uncached=True)
ctx_save_restore_address = dev._gpu_alloc(0x2C02000, kfd.KFD_IOC_ALLOC_MEM_FLAGS_VRAM)
compute_queue = kio.create_queue(dev.kfd, ring_base_address=compute_ring.va_addr, ring_size=compute_ring.size, gpu_id=dev.gpu_id,
queue_type=kfd.KFD_IOC_QUEUE_TYPE_COMPUTE, queue_percentage=kfd.KFD_MAX_QUEUE_PERCENTAGE, queue_priority=kfd.KFD_MAX_QUEUE_PRIORITY,
#eop_buffer_address=eop_buffer.va_addr, eop_buffer_size=eop_buffer.size,
#ctx_save_restore_address=ctx_save_restore_address.va_addr, ctx_save_restore_size=ctx_save_restore_address.size,
#ctl_stack_size = 0xa000,
write_pointer_address=gart_compute.va_addr, read_pointer_address=gart_compute.va_addr+8)
compute_doorbell = to_mv(dev.doorbells + compute_queue.doorbell_offset - dev.doorbells_base, 4).cast("I")
#scratch = dev._gpu_alloc(0x10000, kfd.KFD_IOC_ALLOC_MEM_FLAGS_VRAM)
ka = to_mv(dev.kernargs_ptr, 0x10).cast("Q")
ka[0] = b.uop.buffer._buf.va_addr
ka[1] = a.uop.buffer._buf.va_addr
compute_read_pointer = to_mv(compute_queue.read_pointer_address, 8).cast("Q")
compute_write_pointer = to_mv(compute_queue.write_pointer_address, 8).cast("Q")
hexdump(to_mv(prg.handle, 0x40))
code = hsa.amd_kernel_code_t.from_address(prg.handle)
#print(format_struct(code))
#print("code")
#hexdump(to_mv(code_ptr, 0x100))
#runner.local_size = [2,1,1]
print(runner.local_size, runner.global_size)
#pm4_cmd += [amd_gpu.PACKET3(amd_gpu.PACKET3_SET_SH_REG, 6), mmCOMPUTE_PGM_LO,
# prg.handle&0xFFFFFFFF, prg.handle>>32, 0, 0, (scratch.va_addr>>8)&0xFFFFFFFF, scratch.va_addr>>40]
code_ptr = (prg.handle + code.kernel_code_entry_byte_offset) >> 8
pm4_cmd = [amd_gpu.PACKET3(amd_gpu.PACKET3_SET_SH_REG, 6), regCOMPUTE_PGM_LO, code_ptr&0xFFFFFFFF, code_ptr>>32, 0, 0, 0, 0]
pm4_cmd += [amd_gpu.PACKET3(amd_gpu.PACKET3_SET_SH_REG, 2), regCOMPUTE_PGM_RSRC1, code.compute_pgm_rsrc1, code.compute_pgm_rsrc2]
pm4_cmd += [amd_gpu.PACKET3(amd_gpu.PACKET3_SET_SH_REG, 2), regCOMPUTE_USER_DATA_0, dev.kernargs_ptr&0xFFFFFFFF, dev.kernargs_ptr>>32]
#pm4_cmd += [amd_gpu.PACKET3(amd_gpu.PACKET3_SET_SH_REG, 2), regCOMPUTE_USER_DATA_0, 0, 0]
pm4_cmd += [amd_gpu.PACKET3(amd_gpu.PACKET3_SET_SH_REG, 8), regCOMPUTE_START_X, 0,0,0,
runner.local_size[0],runner.local_size[1],runner.local_size[2],0,0]
# disabled USE_THREAD_DIMENSIONS
pm4_cmd += [amd_gpu.PACKET3(amd_gpu.PACKET3_DISPATCH_DIRECT, 3),
runner.global_size[0],runner.global_size[1],runner.global_size[2], CS_W32_EN | COMPUTE_SHADER_EN]
#pm4_cmd = [amd_gpu.PACKET3(amd_gpu.PACKET3_NOP, 0x3fff)]*0x200
"""
addr=0x0
sz=(1 << 64)-1
gli=0
glv=0
glk=0
gl1=0
gl2=0
pm4_cmd = [amd_gpu.PACKET3(amd_gpu.PACKET3_ACQUIRE_MEM, 6), 0,
sz & 0xffffffff, (sz >> 32) & 0xff, addr & 0xffffffff, (addr >> 32) & 0xffffff, 0,
amd_gpu.PACKET3_ACQUIRE_MEM_GCR_CNTL_GLI_INV(gli) | amd_gpu.PACKET3_ACQUIRE_MEM_GCR_CNTL_GLK_INV(glk) | \
amd_gpu.PACKET3_ACQUIRE_MEM_GCR_CNTL_GLV_INV(glv) | amd_gpu.PACKET3_ACQUIRE_MEM_GCR_CNTL_GL1_INV(gl1) | \
amd_gpu.PACKET3_ACQUIRE_MEM_GCR_CNTL_GL2_INV(gl2)]
print(pm4_cmd)
"""
wptr = 0
pm4_buffer_view = to_mv(compute_ring.va_addr, compute_ring.size).cast("I")
for j in range(0x80000):
for i, value in enumerate(pm4_cmd): pm4_buffer_view[wptr+i] = value
wptr += len(pm4_cmd)
compute_write_pointer[0] = wptr
compute_doorbell[0] = wptr
for k in range(10):
done = compute_read_pointer[0] == compute_write_pointer[0]
print(compute_read_pointer[0], compute_write_pointer[0], done)
if done: break
time.sleep(0.01)
break
#break
#print(compute_read_pointer[0])
#time.sleep(0.05)
#print(compute_read_pointer[0])
#time.sleep(100)
print(a.numpy())
print(b.numpy())
exit(0)
#pm4_cmd = [amd_gpu.PACKET3(amd_gpu.PACKET3_SET_SH_REG, 8), mmCOMPUTE_PGM_LO, 0,0,0,1,1,1,0,0]
#pm4_cmd += [amd_gpu.PACKET3(amd_gpu.PACKET3_DISPATCH_DIRECT, )]
#pm4_cmd = [amd_gpu.PACKET3(amd_gpu.PACKET3_ACQUIRE_MEM, 6), 0,
# sz & 0xffffffff, (sz >> 32) & 0xff, addr & 0xffffffff, (addr >> 32) & 0xffffff, 0,
# amd_gpu.PACKET3_ACQUIRE_MEM_GCR_CNTL_GLI_INV(gli) | amd_gpu.PACKET3_ACQUIRE_MEM_GCR_CNTL_GLK_INV(glk) | \
# amd_gpu.PACKET3_ACQUIRE_MEM_GCR_CNTL_GLV_INV(glv) | amd_gpu.PACKET3_ACQUIRE_MEM_GCR_CNTL_GL1_INV(gl1) | \
# amd_gpu.PACKET3_ACQUIRE_MEM_GCR_CNTL_GL2_INV(gl2)]

43
tinygrad_repo/extra/hip_gpu_driver/test_sdma_fun.py Normal file
View File

@@ -0,0 +1,43 @@
import ctypes, mmap, time
from tinygrad.runtime.ops_amd import AMDDevice, kio, sdma_pkts, libc
import tinygrad.runtime.autogen.amd_sdma as amd_sdma
import tinygrad.runtime.autogen.kfd as kfd
from tinygrad.helpers import to_mv
if __name__ == "__main__":
dev = AMDDevice()
sdma_ring = dev._gpu_alloc(1 << 22, kfd.KFD_IOC_ALLOC_MEM_FLAGS_USERPTR, uncached=True)
gart = dev._gpu_alloc(0x1000, kfd.KFD_IOC_ALLOC_MEM_FLAGS_GTT, uncached=True)
sdma_queue = kio.create_queue(AMDDevice.kfd,
ring_base_address=sdma_ring.va_addr, ring_size=sdma_ring.size, gpu_id=dev.gpu_id,
queue_type=kfd.KFD_IOC_QUEUE_TYPE_SDMA, queue_percentage=kfd.KFD_MAX_QUEUE_PERCENTAGE, queue_priority=kfd.KFD_MAX_QUEUE_PRIORITY,
write_pointer_address=gart.va_addr + 0x100, read_pointer_address=gart.va_addr + 0x108)
doorbells_base = sdma_queue.doorbell_offset & (~0xfff)
doorbells = libc.mmap(0, 8192, mmap.PROT_READ|mmap.PROT_WRITE, mmap.MAP_SHARED, AMDDevice.kfd, doorbells_base)
sdma_read_pointer = to_mv(sdma_queue.read_pointer_address, 8).cast("Q")
sdma_write_pointer = to_mv(sdma_queue.write_pointer_address, 8).cast("Q")
sdma_doorbell = to_mv(doorbells + sdma_queue.doorbell_offset - doorbells_base, 4).cast("I")
test_write_page = dev._gpu_alloc(0x1000, kfd.KFD_IOC_ALLOC_MEM_FLAGS_USERPTR, uncached=True)
cmd = sdma_pkts.timestamp(op=amd_sdma.SDMA_OP_TIMESTAMP, sub_op=amd_sdma.SDMA_SUBOP_TIMESTAMP_GET_GLOBAL, addr=test_write_page.va_addr)
sdma_doorbell_value = 0
def blit_sdma_command(cmd):
ctypes.memmove(sdma_ring.va_addr + (sdma_doorbell_value % sdma_ring.size), ctypes.addressof(cmd), sz:=ctypes.sizeof(cmd))
return sz
while True:
sdma_doorbell_value += blit_sdma_command(cmd)
sdma_write_pointer[0] = sdma_doorbell_value
sdma_doorbell[0] = sdma_doorbell_value
while sdma_read_pointer[0] != sdma_write_pointer[0]: continue
tm = to_mv(test_write_page.va_addr, 0x1000).cast("Q")[0]/1e8
print(f"{tm:.3f} s @ 0x{sdma_ring.va_addr + (sdma_doorbell_value % sdma_ring.size):X} R:0x{sdma_queue.read_pointer_address:X} W:0x{sdma_queue.write_pointer_address:X}")
time.sleep(0.01)

3335
tinygrad_repo/extra/hip_gpu_driver/vega10_sdma_pkt_open.h Normal file
View File

File diff suppressed because it is too large Load Diff