"...composable_kernel_rocm.git" did not exist on "6d073d31bbc7d39d8b170d549f2af61970378150"
Commit e3b6b9f3 authored by Antoine Kaufmann's avatar Antoine Kaufmann
Browse files

i40e: draft for hmc fetch for direct access

parent 04400c02
......@@ -130,7 +130,7 @@ class host_mem_cache {
static const uint16_t MAX_SEGMENTS = 0x1000;
struct segment {
uint64_t pdir_addr;
uint64_t addr;
uint16_t pgcount;
bool valid;
bool direct;
......@@ -138,9 +138,18 @@ class host_mem_cache {
i40e_bm &dev;
segment segs[MAX_SEGMENTS];
public:
class mem_op : public dma_base {
public:
bool failed;
};
host_mem_cache(i40e_bm &dev);
void reg_updated(uint64_t addr);
// issue a hmc memory operation (address is in the context
void issue_mem_op(mem_op &op);
};
......
......@@ -15,7 +15,7 @@ host_mem_cache::host_mem_cache(i40e_bm &dev_)
: dev(dev_)
{
for (size_t i = 0; i < MAX_SEGMENTS; i++) {
segs[i].pdir_addr = 0;
segs[i].addr = 0;
segs[i].pgcount = 0;
segs[i].valid = false;
segs[i].direct = false;
......@@ -24,5 +24,86 @@ host_mem_cache::host_mem_cache(i40e_bm &dev_)
void host_mem_cache::reg_updated(uint64_t addr)
{
std::cerr << "hmc reg updated " << addr << std::endl;
if (addr == I40E_PFHMC_SDCMD) {
// read/write command for descriptor
uint32_t cmd = dev.regs.pfhmc_sdcmd;
uint16_t idx = (cmd & I40E_PFHMC_SDCMD_PMSDIDX_MASK) >>
I40E_PFHMC_SDCMD_PMSDIDX_SHIFT;
uint32_t lo = dev.regs.pfhmc_sddatalow;
uint32_t hi = dev.regs.pfhmc_sddatahigh;
if ((cmd & I40E_PFHMC_SDCMD_PMSDWR_MASK)) {
// write
std::cerr << "hmc: writing descriptor " << idx << std::endl;
segs[idx].addr = ((lo & I40E_PFHMC_SDDATALOW_PMSDDATALOW_MASK) >>
I40E_PFHMC_SDDATALOW_PMSDDATALOW_SHIFT) << 12;
segs[idx].addr |= ((uint64_t) hi) << 32;
segs[idx].pgcount = (lo & I40E_PFHMC_SDDATALOW_PMSDBPCOUNT_MASK) >>
I40E_PFHMC_SDDATALOW_PMSDBPCOUNT_SHIFT;
segs[idx].valid = !!(lo & I40E_PFHMC_SDDATALOW_PMSDVALID_MASK);
segs[idx].direct = !!(lo & I40E_PFHMC_SDDATALOW_PMSDTYPE_MASK);
std::cerr << " addr=" << segs[idx].addr << " pgcount=" <<
segs[idx].pgcount << " valid=" << segs[idx].valid <<
" direct=" << segs[idx].direct << std::endl;
} else {
// read
std::cerr << "hmc: reading descriptor " << idx << std::endl;
dev.regs.pfhmc_sddatalow = ((segs[idx].addr >> 12) <<
I40E_PFHMC_SDDATALOW_PMSDDATALOW_SHIFT) &
I40E_PFHMC_SDDATALOW_PMSDDATALOW_MASK;
dev.regs.pfhmc_sddatalow |= (segs[idx].pgcount <<
I40E_PFHMC_SDDATALOW_PMSDBPCOUNT_SHIFT) &
I40E_PFHMC_SDDATALOW_PMSDBPCOUNT_MASK;
if (segs[idx].valid)
dev.regs.pfhmc_sddatalow |= I40E_PFHMC_SDDATALOW_PMSDVALID_MASK;
if (segs[idx].direct)
dev.regs.pfhmc_sddatalow |= I40E_PFHMC_SDDATALOW_PMSDTYPE_MASK;
dev.regs.pfhmc_sddatahigh = segs[idx].addr >> 32;
}
}
}
void host_mem_cache::issue_mem_op(mem_op &op)
{
uint64_t addr = op.dma_addr;
uint16_t seg_idx = addr >> 21;
uint16_t seg_idx_last = (addr + op.len - 1) >> 21;
uint32_t dir_off = addr & ((1 << 21) - 1);
struct segment *seg = &segs[seg_idx];
if (seg_idx >= MAX_SEGMENTS) {
std::cerr << "hmc issue_mem_op: seg index too high " << seg_idx <<
std::endl;
abort();
}
if (!seg->valid) {
// TODO: errorinfo and data registers
std::cerr << "hmc issue_mem_op: segment invalid addr=" << addr <<
std::endl;
op.failed = true;
return;
}
if (seg_idx != seg_idx_last) {
std::cerr << "hmc issue_mem_op: operation crosses segs addr=" <<
addr << " len=" << op.len << std::endl;
abort();
}
if (!seg->direct) {
std::cerr << "hmc issue_mem_op: TODO paged ops addr=" << addr <<
std::endl;
abort();
}
op.failed = false;
op.dma_addr = seg->addr + dir_off;
std::cerr << "hmc issue_mem_op: hmc_addr=" << addr << " dma_addr=" <<
op.dma_addr << " len=" << op.len << std::endl;
runner->issue_dma(op);
}
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment