diff --git a/src/realm/cuda/cuda_internal.cc b/src/realm/cuda/cuda_internal.cc index 2c52a7496d2..5c7558e1bf6 100644 --- a/src/realm/cuda/cuda_internal.cc +++ b/src/realm/cuda/cuda_internal.cc @@ -21,6 +21,8 @@ #include "realm/cuda/cuda_memcpy.h" #include "realm/realm_assert.h" +#include + namespace Realm { extern Logger log_xd; @@ -302,8 +304,6 @@ namespace Realm { return in_lstride > in_pstride || out_lstride > out_pstride; } - // Calculates the maximum alignment native type alignment the GPU supports that will - // work with the given size. static size_t calculate_type_alignment(size_t v) { // We don't need a full log2 here @@ -314,45 +314,127 @@ namespace Realm { return 1; // Unfortunately this can only be byte aligned :( } - static size_t populate_affine_copy_info(AffineCopyInfo<3> ©_infos, - size_t &min_align, - MemcpyTransposeInfo &transpose_info, - AddressListCursor &in_alc, uintptr_t in_base, - GPU *in_gpu, AddressListCursor &out_alc, - uintptr_t out_base, GPU *out_gpu, - size_t bytes_left) + size_t GPUXferDes::read_address_entry(AffineCopyInfo<3> ©_infos, + size_t &min_align, + MemcpyTransposeInfo &transpose_info, + AddressListCursor &src_cur, uintptr_t in_base, + AddressListCursor &dst_cur, uintptr_t out_base, + size_t bytes_left, size_t max_xfer_fields, + size_t &fields_total) { AffineCopyPair<3> ©_info = copy_infos.subrects[copy_infos.num_rects++]; - uintptr_t in_offset = in_alc.get_offset(); - uintptr_t out_offset = out_alc.get_offset(); - // the reported dim is reduced for partially consumed address - // ranges - whatever we get can be assumed to be regular - int in_dim = in_alc.get_dim(); - int out_dim = out_alc.get_dim(); - size_t icount = in_alc.remaining(0); - size_t ocount = out_alc.remaining(0); - // contig bytes is always the min of the first dimensions - size_t contig_bytes = std::min(std::min(icount, ocount), bytes_left); + + using std::max; + using std::min; + + // --------------------------------------------------------------------------- + // helpers + // --------------------------------------------------------------------------- + + const auto attach_fields = [&](AddressListCursor &c, AffineSubRect<3> &subr) { + if(c.field_block()) { + subr.num_fields = min(max_xfer_fields, c.remaining_fields()); + subr.fields = c.fields_data(); + fields_total = max(fields_total, subr.num_fields); + } + }; + + size_t icount = src_cur.remaining(0); + size_t ocount = dst_cur.remaining(0); + + const uintptr_t in_offset = src_cur.get_offset(); + const uintptr_t out_offset = dst_cur.get_offset(); + + int in_dim = src_cur.get_dim(); + int out_dim = dst_cur.get_dim(); + + const size_t contig_bytes = min({icount, ocount, bytes_left}); + + // After we know the volume of the rectangle, decide how many + // fields we can really move without exceeding bytes_left and + // update both sub-rects’ num_fields accordingly. + // + // The lambda returns the final field‐count so callers can pass + // it straight to `AddressListCursor::advance`. + const auto final_field_count = [&](size_t rect_volume, size_t left_bytes, size_t ic, + size_t oc) -> size_t { + size_t rect_fields = + max(1, max(copy_info.src.num_fields, copy_info.dst.num_fields)); + + // TODO(apryakhin:): Use left_bytes + size_t left = bytes_left; + + if(src_cur.field_block()) { + left = min(left, src_cur.partial + ? ic + : ic * max(size_t(1), ic * copy_info.src.num_fields)); + } else { + left = min(left, ic); + } + + if(dst_cur.field_block()) { + left = min(left, dst_cur.partial + ? oc + : oc * max(size_t(1), oc * copy_info.dst.num_fields)); + } else { + left = min(left, oc); + } + + size_t max_by_bytes = (rect_volume == 0) ? 0 : (left / rect_volume); + rect_fields = max(size_t(1), min(rect_fields, max_by_bytes)); + + if(copy_info.src.num_fields) { + copy_info.src.num_fields = min(copy_info.src.num_fields, rect_fields); + } + + if(copy_info.dst.num_fields) { + copy_info.dst.num_fields = min(copy_info.dst.num_fields, rect_fields); + } + + return rect_fields; + }; + + auto advance = [&](AddressListCursor &c, int dim, size_t amount, size_t scale, + size_t fields) { + if(c.field_block()) + c.advance(dim, amount * scale, fields); + else + c.advance(dim, amount * scale * fields); + }; log_gpudma.info() << "IN: " << in_dim << ' ' << icount << ' ' << in_offset << ' ' << contig_bytes; log_gpudma.info() << "OUT: " << out_dim << ' ' << ocount << ' ' << out_offset << ' ' - << contig_bytes; + << contig_bytes << " left:" << bytes_left; - assert(in_dim > 0); - assert(out_dim > 0); + assert(in_dim > 0 && out_dim > 0); - copy_info.src.addr = static_cast(in_base + in_offset); - copy_info.dst.addr = static_cast(out_base + out_offset); + copy_info.src.addr = in_base + in_offset; + copy_info.dst.addr = out_base + out_offset; copy_info.extents[1] = 1; copy_info.extents[2] = 1; - min_align = std::min(min_align, calculate_type_alignment(copy_info.src.addr)); - min_align = std::min(min_align, calculate_type_alignment(copy_info.dst.addr)); - // Calculate the minimum alignment for contig bytes - min_align = std::min(min_align, calculate_type_alignment(contig_bytes)); + min_align = min(min_align, calculate_type_alignment(copy_info.src.addr)); + min_align = min(min_align, calculate_type_alignment(copy_info.dst.addr)); + min_align = min(min_align, calculate_type_alignment(contig_bytes)); - // catch simple 1D case first + attach_fields(src_cur, copy_info.src); + attach_fields(dst_cur, copy_info.dst); + + if(src_cur.field_block() && dst_cur.field_block()) { + copy_info.src.field_stride = src_cur.addrlist->full_field_bytes(); + copy_info.dst.field_stride = dst_cur.addrlist->full_field_bytes(); + } else if(src_cur.field_block()) { + copy_info.src.field_stride = copy_info.dst.field_stride = + src_cur.addrlist->full_field_bytes(); + } else if(dst_cur.field_block()) { + copy_info.dst.field_stride = copy_info.src.field_stride = + dst_cur.addrlist->full_field_bytes(); + } + + // --------------------------------------------------------------------------- + // fast path – pure 1‑D + // --------------------------------------------------------------------------- if((contig_bytes == bytes_left) || ((contig_bytes == icount) && (in_dim == 1)) || ((contig_bytes == ocount) && (out_dim == 1))) { copy_info.extents[0] = contig_bytes; @@ -360,12 +442,16 @@ namespace Realm { copy_info.dst.strides[0] = contig_bytes; copy_info.volume = contig_bytes; - in_alc.advance(0, contig_bytes); - out_alc.advance(0, contig_bytes); - return contig_bytes; + fields_total = final_field_count(copy_info.src.field_stride, copy_info.volume, + icount, ocount); + advance(src_cur, 0, contig_bytes, 1, fields_total); + advance(dst_cur, 0, contig_bytes, 1, fields_total); + return contig_bytes * fields_total; } - // grow to a 2D copy + // --------------------------------------------------------------------------- + // grow to 2‑D (id/od == sub‑dimension chosen for “lines") + // --------------------------------------------------------------------------- int id; size_t iscale; uintptr_t in_lstride; @@ -381,8 +467,8 @@ namespace Realm { } else { assert(in_dim > 1); id = 1; - icount = in_alc.remaining(id); - in_lstride = in_alc.get_stride(id); + icount = src_cur.remaining(id); + in_lstride = src_cur.get_stride(id); iscale = 1; } @@ -401,23 +487,19 @@ namespace Realm { } else { assert(out_dim > 1); od = 1; - ocount = out_alc.remaining(od); - out_lstride = out_alc.get_stride(od); + ocount = dst_cur.remaining(od); + out_lstride = dst_cur.get_stride(od); oscale = 1; } - size_t lines = std::min(std::min(icount, ocount), bytes_left / contig_bytes); + const size_t lines = min(min(icount, ocount), bytes_left / contig_bytes); - // *_lstride is the number of bytes for each line, so recalculate - // the minimum alignment to make sure the alignment matches the - // byte alignment across all lines. - min_align = std::min(min_align, calculate_type_alignment(in_lstride)); - min_align = std::min(min_align, calculate_type_alignment(out_lstride)); + min_align = min(min_align, calculate_type_alignment(in_lstride)); + min_align = min(min_align, calculate_type_alignment(out_lstride)); - // see if we need to stop at 2D if(((contig_bytes * lines) == bytes_left) || - ((lines == icount) && (id == (in_dim - 1))) || - ((lines == ocount) && (od == (out_dim - 1)))) { + ((lines == icount) && (id == in_dim - 1)) || + ((lines == ocount) && (od == out_dim - 1))) { copy_info.src.strides[0] = in_lstride; copy_info.src.strides[1] = lines; copy_info.dst.strides[0] = out_lstride; @@ -426,12 +508,16 @@ namespace Realm { copy_info.extents[1] = lines; copy_info.volume = lines * contig_bytes; - in_alc.advance(id, lines * iscale); - out_alc.advance(od, lines * oscale); - return lines * contig_bytes; + fields_total = final_field_count(copy_info.src.field_stride, copy_info.volume, + in_lstride * icount, ocount * out_lstride); + advance(src_cur, id, lines, iscale, fields_total); + advance(dst_cur, od, lines, oscale, fields_total); + return copy_info.volume * fields_total; } - // Grow to a 3D copy + // --------------------------------------------------------------------------- + // need full 3‑D or transpose + // --------------------------------------------------------------------------- uintptr_t in_pstride; if(lines < icount) { // third input dim comes from splitting current @@ -443,8 +529,8 @@ namespace Realm { } else { id++; assert(in_dim > id); - icount = in_alc.remaining(id); - in_pstride = in_alc.get_stride(id); + icount = src_cur.remaining(id); + in_pstride = src_cur.get_stride(id); iscale = 1; } @@ -459,45 +545,44 @@ namespace Realm { } else { od++; assert(out_dim > od); - ocount = out_alc.remaining(od); - out_pstride = out_alc.get_stride(od); + ocount = dst_cur.remaining(od); + out_pstride = dst_cur.get_stride(od); oscale = 1; } const size_t planes = - std::min(std::min(icount, ocount), (bytes_left / (contig_bytes * lines))); - - if(needs_transpose(in_lstride, in_pstride, out_lstride, out_pstride)) { - transpose_info.src = static_cast(in_base + in_offset); - transpose_info.dst = static_cast(out_base + out_offset); + min(min(icount, ocount), (bytes_left / (contig_bytes * lines))); + const bool do_transpose = (in_lstride > in_pstride) || (out_lstride > out_pstride); + if(do_transpose) { + transpose_info.src = in_base + in_offset; + transpose_info.dst = out_base + out_offset; transpose_info.src_strides[0] = in_lstride; transpose_info.src_strides[1] = in_pstride; - transpose_info.dst_strides[0] = out_lstride; transpose_info.dst_strides[1] = out_pstride; - transpose_info.extents[0] = contig_bytes; transpose_info.extents[1] = lines; transpose_info.extents[2] = planes; copy_infos.num_rects--; } else { + copy_info.src.strides[0] = in_lstride; + copy_info.src.strides[1] = in_pstride / in_lstride; copy_info.dst.strides[0] = out_lstride; copy_info.dst.strides[1] = out_pstride / out_lstride; - copy_info.extents[0] = contig_bytes; copy_info.extents[1] = lines; copy_info.extents[2] = planes; - - copy_info.src.strides[0] = in_lstride; - copy_info.src.strides[1] = in_pstride / in_lstride; - copy_info.volume = planes * lines * contig_bytes; } - in_alc.advance(id, planes * iscale); - out_alc.advance(od, planes * oscale); - return planes * lines * contig_bytes; + fields_total = + final_field_count(copy_info.src.field_stride, planes * lines * contig_bytes, + in_pstride * icount, ocount * out_pstride); + + advance(src_cur, id, planes, iscale, fields_total); + advance(dst_cur, od, planes, oscale, fields_total); + return contig_bytes * lines * planes * fields_total; } bool GPU::is_accessible_host_mem(const MemoryImpl *mem) const @@ -543,17 +628,6 @@ namespace Realm { bool GPUXferDes::progress_xd(GPUChannel *channel, TimeLimit work_until) { - // Mininum amount to transfer in a single quantum before returning in order to - // ensure forward progress - // TODO: make controllable - const size_t MIN_XFER_SIZE = 4 << 20; - // Maximum amount to transfer in a single quantum in order to ensure other requests - // have a chance to make forward progress. This should be large enough that the - // overhead of splitting the copy shouldn't be noticable in terms of latency (4GiB - // should be good here for most purposes) - // TODO: make controllable - const size_t flow_control_bytes = 4ULL * 1024ULL * 1024ULL * 1024ULL; - ReadSequenceCache rseqcache(this, 2 << 20); WriteSequenceCache wseqcache(this, 2 << 20); GPUStream *stream = 0; @@ -568,12 +642,13 @@ namespace Realm { memset(©_infos, 0, sizeof(copy_infos)); // The general algorithm here can be described in three loops: - // 1) Outer loop - iterates over all the addresses for each request. This typically - // corresponds to each rectangle in an index space transfer. 2) Batch loop - Map the - // address list that can be a mix of different rectangle sizes to a batch of copies - // that can be pushed in a single launch (either kernel or cuMemcpy call) - // 2.a) At this point advancing the address list commits us to submitting the copy - // in #3, thus flow control happens here. + // 1) Outer loop - iterates over all the addresses for each request. This + // typically corresponds to each rectangle in an index space transfer. 2) Batch + // loop - Map the address list that can be a mix of different rectangle sizes to a + // batch of copies that can be pushed in a single launch (either kernel or + // cuMemcpy call) + // 2.a) At this point advancing the address list commits us to submitting the + // copy in #3, thus flow control happens here. // 3) Copy loop - Based on the batch, descide the best copy to push. // 1) Outer loop - iterate over all the addresses for each request @@ -588,15 +663,18 @@ namespace Realm { const InstanceLayoutPieceBase *in_nonaffine, *out_nonaffine; - if(((total_bytes >= MIN_XFER_SIZE) && work_until.is_expired()) || - (total_bytes >= flow_control_bytes)) { + if(((total_bytes >= min_xfer_size) && work_until.is_expired()) || + (total_bytes >= max_xfer_size)) { log_gpudma.info() << "Flow control hit, copied " << total_bytes - << " leave the rest for later!"; + << " max_xfer_size:" << max_xfer_size + << " min_xfer_size:" << min_xfer_size << " xd=" << std::hex + << guid << std::dec; break; } const size_t max_bytes = - get_addresses(MIN_XFER_SIZE, &rseqcache, in_nonaffine, out_nonaffine); + get_addresses(min_xfer_size, &rseqcache, in_nonaffine, out_nonaffine); + if(max_bytes == 0) { break; } @@ -619,10 +697,11 @@ namespace Realm { out_is_ipc = dst_is_ipc[output_control.current_io_port]; } - // We need a kernel copy if this is a H2H copy, as CUDA forces the calling thread - // to perform H2H copies synchronously with cuMemcpyAsync. We have decided that - // the GPU is the best one to do this (either via a faster inter-connect than one - // CPU thread can saturate or the fact it can be done asynchronously) + // We need a kernel copy if this is a H2H copy, as CUDA forces the calling + // thread to perform H2H copies synchronously with cuMemcpyAsync. We have + // decided that the GPU is the best one to do this (either via a faster + // inter-connect than one CPU thread can saturate or the fact it can be done + // asynchronously) needs_kernel_copy = (in_port != nullptr) && (in_gpu->is_accessible_host_mem(in_port->mem)) && (out_port != nullptr) && (out_gpu->is_accessible_host_mem(out_port->mem)); @@ -663,7 +742,7 @@ namespace Realm { size_t copy_info_total = 0; size_t min_align = 16; // Hope for the highest type alignment we can get, 16 bytes copy_infos.num_rects = 0; - size_t bytes_left = std::min(flow_control_bytes - total_bytes, max_bytes); + size_t bytes_left = std::min(max_xfer_size - total_bytes, max_bytes); if(cuda_copy.WidthInBytes != 0) { memset(&cuda_copy, 0, sizeof(cuda_copy)); @@ -672,26 +751,32 @@ namespace Realm { memset(&transpose_copy, 0, sizeof(transpose_copy)); } + bool needs_fast_multifield = false; + size_t fields_total = 1; + // 2) Batch loop - Collect all the rectangles for this inport/outport pair by // iterating the address list cursor for each and figure out what copy we can do // that best fits the layout of the source and destinations - while(bytes_left > 0 && copy_infos.num_rects < AffineCopyInfo<3>::MAX_NUM_RECTS) { - AddressListCursor &in_alc = in_port->addrcursor; - AddressListCursor &out_alc = out_port->addrcursor; + while(bytes_left > 0 && copy_infos.num_rects < AffineCopyInfo<3>::MAX_NUM_RECTS && + !needs_fast_multifield) { + AddressListCursor &src_cur = in_port->addrcursor; + AddressListCursor &dst_cur = out_port->addrcursor; + if(!in_nonaffine && !out_nonaffine) { - log_gpudma.info() << "Affine -> Affine"; - // limit transfer size for host<->device copies - // this is because CUDA stages these copies through a staging buffer and is a - // blocking call. Thus to limit the amount of time spent within the cuda - // driver and to allow us to time out early if needed, split these larger - // copies into smaller ones ourselves + // log_gpudma.info() << "Affine -> Affine"; + // limit transfer size for host<->device copies + // this is because CUDA stages these copies through a staging buffer and is + // a blocking call. Thus to limit the amount of time spent within the cuda + // driver and to allow us to time out early if needed, split these larger + // copies into smaller ones ourselves if(!in_gpu || (!out_gpu && !out_is_ipc)) { bytes_left = std::min(bytes_left, (size_t)(4U << 20U)); } - const size_t bytes_to_copy = populate_affine_copy_info( - copy_infos, min_align, transpose_copy, in_alc, in_base, in_gpu, out_alc, - out_base, out_gpu, bytes_left); + size_t bytes_to_copy = 0; + bytes_to_copy = read_address_entry(copy_infos, min_align, transpose_copy, + src_cur, in_base, dst_cur, out_base, + bytes_left, max_xfer_fields, fields_total); // Either src or dst can't be accessed with a kernel, so just break out and // perform a standard cuMemcpy @@ -701,16 +786,23 @@ namespace Realm { } log_gpudma.info() << "\tAdded " << bytes_to_copy - << " Bytes left= " << (bytes_left - bytes_to_copy); + << " Bytes left= " << (bytes_left - bytes_to_copy) + << " xd=" << std::hex << guid << std::dec; + assert(bytes_to_copy <= bytes_left); copy_info_total += bytes_to_copy; bytes_left -= bytes_to_copy; + + if(src_cur.field_block() || dst_cur.field_block()) { + needs_fast_multifield = true; + break; + } } else { // Non-affine transfers AddressInfoCudaArray ainfo; if(in_nonaffine) { assert(!out_nonaffine); - log_gpudma.info() << "Array -> Affine"; + /// log_gpudma.info() << "Array -> Affine"; size_t bytes = in_port->iter->step_custom(bytes_left, ainfo, false); if(bytes == 0) break; // flow control or end of array @@ -721,9 +813,9 @@ namespace Realm { cuda_copy.srcZ = ainfo.pos[2]; cuda_copy.dstMemoryType = CU_MEMORYTYPE_UNIFIED; cuda_copy.dstDevice = - static_cast(out_base + out_alc.get_offset()); + static_cast(out_base + dst_cur.get_offset()); get_nonaffine_strides(cuda_copy.dstPitch, cuda_copy.dstHeight, ainfo, - out_alc, bytes); + dst_cur, bytes); } else { assert(!in_nonaffine); log_gpudma.info() << "Affine -> Array"; @@ -737,9 +829,9 @@ namespace Realm { cuda_copy.dstZ = ainfo.pos[2]; cuda_copy.srcMemoryType = CU_MEMORYTYPE_UNIFIED; cuda_copy.srcDevice = - static_cast(in_base + in_alc.get_offset()); + static_cast(in_base + src_cur.get_offset()); get_nonaffine_strides(cuda_copy.srcPitch, cuda_copy.srcHeight, ainfo, - in_alc, bytes); + src_cur, bytes); } cuda_copy.WidthInBytes = ainfo.width_in_bytes; cuda_copy.Height = ainfo.height; @@ -751,10 +843,11 @@ namespace Realm { } } // 3) Copy loop - Actually perform the copies enumerated earlier and track their - // completion This logic will determine which path was ultimately chosen based on - // the enumeration logic and should only be one launch API call, regardless of the - // size of the batch. These copies *must* be submitted and cannot be interrupted, - // as we've already updated the addresslistcursor and committed to submitting them + // completion This logic will determine which path was ultimately chosen based + // on the enumeration logic and should only be one launch API call, regardless + // of the size of the batch. These copies *must* be submitted and cannot be + // interrupted, as we've already updated the addresslistcursor and committed to + // submitting them size_t bytes_to_fence = 0; if(cuda_copy.WidthInBytes != 0) { // First the non-affine copies @@ -809,23 +902,35 @@ namespace Realm { transpose_copy.extents[2]; } - if((copy_infos.num_rects > 1) || needs_kernel_copy) { - // Adjust all the rectangles' sizes to account for the element size based on the - // calculated alignment + // if(needs_fast_multifield) { + if((copy_infos.num_rects > 1) || needs_kernel_copy || needs_fast_multifield) { + // Adjust all the rectangles' sizes to account for the element size based on + // the calculated alignment for(size_t i = 0; (min_align > 1) && (i < copy_infos.num_rects); i++) { copy_infos.subrects[i].dst.strides[0] /= min_align; copy_infos.subrects[i].src.strides[0] /= min_align; + if(copy_infos.subrects[i].src.field_stride) { + copy_infos.subrects[i].src.field_stride /= min_align; + } + if(copy_infos.subrects[i].dst.field_stride) { + copy_infos.subrects[i].dst.field_stride /= min_align; + } copy_infos.subrects[i].extents[0] /= min_align; copy_infos.subrects[i].volume /= min_align; } - // TODO: add some heuristics here, like if some rectangles are very large, do a - // cuMemcpy instead, possibly utilizing the copy engines or better optimized + + // TODO: add some heuristics here, like if some rectangles are very large, do + // a cuMemcpy instead, possibly utilizing the copy engines or better optimized // kernels log_gpudma.info() << "\tLaunching kernel for rects=" << copy_infos.num_rects + << " xd=" << std::hex << guid << std::dec << " bytes=" << copy_info_total - << " out_is_ipc=" << out_is_ipc; - stream->get_gpu()->launch_batch_affine_kernel( - ©_infos, 3, min_align, copy_info_total / min_align, stream); + << " out_is_ipc=" << out_is_ipc << " fields=" << fields_total + << " needs_multi:" << needs_fast_multifield; + + stream->get_gpu()->launch_batch_affine_kernel(©_infos, 3, min_align, + (copy_info_total / min_align), + needs_fast_multifield, stream); bytes_to_fence += copy_info_total; } else if(copy_infos.num_rects == 1) { // Then the affine copies to/from the device @@ -1019,8 +1124,8 @@ namespace Realm { size_t addr_size = 0; - AddressListCursor &in_alc = in_port->addrcursor; - AddressListCursor &out_alc = out_port->addrcursor; + AddressListCursor &src_cur = in_port->addrcursor; + AddressListCursor &dst_cur = out_port->addrcursor; size_t write_ind_bytes = 0; uintptr_t dst_ind_base = 0; @@ -1034,9 +1139,9 @@ namespace Realm { out_port->iter->get_base_offset(), 0)); out_base += addr_info.base_offset; - dst_ind_base += (out_alc.get_offset() / addr_info.bytes_per_chunk) * addr_size; + dst_ind_base += (dst_cur.get_offset() / addr_info.bytes_per_chunk) * addr_size; } else { - out_base += out_alc.get_offset(); + out_base += dst_cur.get_offset(); } size_t read_ind_bytes = 0; @@ -1051,9 +1156,9 @@ namespace Realm { in_port->iter->get_base_offset(), 0)); in_base += addr_info.base_offset; - src_ind_base += (in_alc.get_offset() / addr_info.bytes_per_chunk) * addr_size; + src_ind_base += (src_cur.get_offset() / addr_info.bytes_per_chunk) * addr_size; } else { - in_base += in_alc.get_offset(); + in_base += src_cur.get_offset(); } log_gpudma.info() << "cuda gathe/scatter bytes_per_chunk=" @@ -1100,8 +1205,8 @@ namespace Realm { max_bytes, strides, in_base, out_base, src_ind_base, dst_ind_base); - in_alc.advance(0, max_bytes); - out_alc.advance(0, max_bytes); + src_cur.advance(0, max_bytes); + dst_cur.advance(0, max_bytes); // TODO(apryakhin@): Add control flow total_bytes += max_bytes; @@ -1553,8 +1658,8 @@ namespace Realm { case GPU_FB_MEM: continue; default: - add_path(local_gpu_mems, static_cast(i), /*src_global=*/false, - bw, latency, frag_overhead, XFER_GPU_FROM_FB) + add_path(local_gpu_mems, static_cast(i), + /*src_global=*/false, bw, latency, frag_overhead, XFER_GPU_FROM_FB) .set_max_dim(2); break; } @@ -1675,6 +1780,69 @@ namespace Realm { return 0; } + RemoteChannelInfo *GPUChannel::construct_remote_info() const + { + return new GPURemoteChannelInfo(node, kind, reinterpret_cast(this), + paths); + } + + //////////////////////////////////////////////////////////////////////// + // + // class GPURemoteChannelInfo + // + + GPURemoteChannelInfo::GPURemoteChannelInfo( + NodeID _owner, XferDesKind _kind, uintptr_t _remote_ptr, + const std::vector &_paths) + : SimpleRemoteChannelInfo(_owner, _kind, _remote_ptr, _paths) + {} + + RemoteChannel *GPURemoteChannelInfo::create_remote_channel() + { + GPURemoteChannel *rc = new GPURemoteChannel(remote_ptr); + rc->node = owner; + rc->kind = kind; + rc->paths.swap(paths); + return rc; + } + + // these templates can go here because they're only used by the helper below + template + bool GPURemoteChannelInfo::serialize(S &serializer) const + { + return ((serializer << owner) && (serializer << kind) && + (serializer << remote_ptr) && (serializer << paths)); + } + + template + /*static*/ RemoteChannelInfo *GPURemoteChannelInfo::deserialize_new(S &deserializer) + { + NodeID owner; + XferDesKind kind; + uintptr_t remote_ptr; + std::vector paths; + + if((deserializer >> owner) && (deserializer >> kind) && + (deserializer >> remote_ptr) && (deserializer >> paths)) { + return new GPURemoteChannelInfo(owner, kind, remote_ptr, paths); + } else { + return 0; + } + } + + /*static*/ Serialization::PolymorphicSerdezSubclass + GPURemoteChannelInfo::serdez_subclass; + + //////////////////////////////////////////////////////////////////////// + // + // class GPURemoteChannel + // + + GPURemoteChannel::GPURemoteChannel(uintptr_t _remote_ptr) + : RemoteChannel(_remote_ptr) + {} + //////////////////////////////////////////////////////////////////////// // // class GPUCompletionEvent @@ -1865,17 +2033,17 @@ namespace Realm { fill_info.num_rects = 0; while(total_bytes < max_bytes) { - AddressListCursor &out_alc = out_port->addrcursor; + AddressListCursor &dst_cur = out_port->addrcursor; - uintptr_t out_offset = out_alc.get_offset(); + uintptr_t out_offset = dst_cur.get_offset(); // the reported dim is reduced for partially consumed address // ranges - whatever we get can be assumed to be regular - int out_dim = out_alc.get_dim(); + int out_dim = dst_cur.get_dim(); if((reduced_fill_size < sizeof(fill_info.fill_value)) && ((reduced_fill_size & (reduced_fill_size - 1)) == 0)) { - const size_t bytes = std::min(out_alc.remaining(0), max_bytes); - size_t lines = (out_dim > 1 ? out_alc.remaining(1) : 1); + const size_t bytes = std::min(dst_cur.remaining(0), max_bytes); + size_t lines = (out_dim > 1 ? dst_cur.remaining(1) : 1); if((lines * bytes) > max_bytes) { lines = std::max(1, max_bytes / bytes); } @@ -1888,7 +2056,7 @@ namespace Realm { bytes / reduced_fill_size; fill_info.subrects[fill_info.num_rects].extents[1] = lines; fill_info.subrects[fill_info.num_rects].strides[0] = - (out_dim > 1 ? out_alc.get_stride(1) : bytes) / reduced_fill_size; + (out_dim > 1 ? dst_cur.get_stride(1) : bytes) / reduced_fill_size; fill_info.num_rects++; total_info_bytes += bytes; @@ -1904,12 +2072,12 @@ namespace Realm { } total_bytes += bytes * lines; - out_alc.advance((out_dim == 1 ? 0 : 1), (out_dim == 1 ? bytes : lines)); + dst_cur.advance((out_dim == 1 ? 0 : 1), (out_dim == 1 ? bytes : lines)); } else { // more general approach - use strided 2d copies to fill the first // line, and then we can use logarithmic doublings to deal with // multiple lines and/or planes - size_t bytes = out_alc.remaining(0); + size_t bytes = dst_cur.remaining(0); size_t elems = bytes / reduced_fill_size; #ifdef DEBUG_REALM assert((bytes % reduced_fill_size) == 0); @@ -1972,11 +2140,11 @@ namespace Realm { if(out_dim == 1) { // all done - out_alc.advance(0, bytes); + dst_cur.advance(0, bytes); total_bytes += bytes; } else { - size_t lines = out_alc.remaining(1); - size_t lstride = out_alc.get_stride(1); + size_t lines = dst_cur.remaining(1); + size_t lstride = dst_cur.get_stride(1); CUDA_MEMCPY2D copy2d; copy2d.srcMemoryType = CU_MEMORYTYPE_DEVICE; @@ -2004,11 +2172,11 @@ namespace Realm { } if(out_dim == 2) { - out_alc.advance(1, lines); + dst_cur.advance(1, lines); total_bytes += bytes * lines; } else { - size_t planes = out_alc.remaining(2); - size_t pstride = out_alc.get_stride(2); + size_t planes = dst_cur.remaining(2); + size_t pstride = dst_cur.get_stride(2); // logarithmic version requires that pstride be a multiple of // lstride @@ -2046,7 +2214,7 @@ namespace Realm { planes_done += todo; } - out_alc.advance(2, planes); + dst_cur.advance(2, planes); total_bytes += bytes * lines * planes; } else { // plane-at-a-time fallback - can reuse most of copy2d @@ -2059,7 +2227,7 @@ namespace Realm { CHECK_CU(CUDA_DRIVER_FNPTR(cuMemcpy2DAsync)(©2d, stream->get_stream())); } - out_alc.advance(2, planes); + dst_cur.advance(2, planes); total_bytes += bytes * lines * planes; } } diff --git a/src/realm/cuda/cuda_internal.h b/src/realm/cuda/cuda_internal.h index 96b8446daa7..f1d750765fb 100644 --- a/src/realm/cuda/cuda_internal.h +++ b/src/realm/cuda/cuda_internal.h @@ -416,7 +416,8 @@ namespace Realm { void launch_batch_affine_fill_kernel(void *fill_info, size_t dim, size_t elemSize, size_t volume, GPUStream *stream); void launch_batch_affine_kernel(void *copy_info, size_t dim, size_t elemSize, - size_t volume, GPUStream *stream); + size_t volume, bool multified_optimized, + GPUStream *stream); void launch_transpose_kernel(MemcpyTransposeInfo ©_info, size_t elemSize, GPUStream *stream); @@ -470,6 +471,8 @@ namespace Realm { GPUFuncInfo indirect_copy_kernels[REALM_MAX_DIM][CUDA_MEMCPY_KERNEL_MAX2_LOG2_BYTES] [CUDA_MEMCPY_KERNEL_MAX2_LOG2_BYTES]; GPUFuncInfo batch_affine_kernels[REALM_MAX_DIM][CUDA_MEMCPY_KERNEL_MAX2_LOG2_BYTES]; + GPUFuncInfo multi_batch_affine_kernels[REALM_MAX_DIM] + [CUDA_MEMCPY_KERNEL_MAX2_LOG2_BYTES]; GPUFuncInfo batch_fill_affine_kernels[REALM_MAX_DIM] [CUDA_MEMCPY_KERNEL_MAX2_LOG2_BYTES]; GPUFuncInfo fill_affine_large_kernels[REALM_MAX_DIM] @@ -800,9 +803,28 @@ namespace Realm { bool progress_xd(GPUChannel *channel, TimeLimit work_until); + static size_t read_address_entry(AffineCopyInfo<3> ©_infos, size_t &min_align, + MemcpyTransposeInfo &transpose_info, + AddressListCursor &in_alc, uintptr_t in_base, + AddressListCursor &out_alc, uintptr_t out_base, + size_t bytes_left, size_t max_xfer_fields, + size_t &fields_total); + private: std::vector src_gpus, dst_gpus; std::vector dst_is_ipc; + + // Mininum amount to transfer in a single quantum before returning in order to + // ensure forward progress + // TODO: make controllable + static constexpr size_t min_xfer_size = 4 << 20; + // Maximum amount to transfer in a single quantum in order to ensure other requests + // have a chance to make forward progress. This should be large enough that the + // overhead of splitting the copy shouldn't be noticable in terms of latency (4GiB + // should be good here for most purposes) + // TODO: make controllable + static constexpr size_t max_xfer_size = 4ULL * 1024ULL * 1024ULL * 1024ULL; + static constexpr size_t max_xfer_fields = 2000; }; class GPUIndirectChannel; @@ -919,9 +941,46 @@ namespace Realm { long submit(Request **requests, long nr); GPU *get_gpu() const { return src_gpu; } + virtual RemoteChannelInfo *construct_remote_info() const; + + virtual bool support_idindexed_fields(Memory src_mem, Memory dst_mem) const + { + return true; + } + private: GPU *src_gpu; - // std::deque pending_copies; + }; + + class GPURemoteChannelInfo : public SimpleRemoteChannelInfo { + public: + GPURemoteChannelInfo(NodeID _owner, XferDesKind _kind, uintptr_t _remote_ptr, + const std::vector &_paths); + + virtual RemoteChannel *create_remote_channel(); + + template + bool serialize(S &serializer) const; + + template + static RemoteChannelInfo *deserialize_new(S &deserializer); + + protected: + static Serialization::PolymorphicSerdezSubclass + serdez_subclass; + }; + + class GPURemoteChannel : public RemoteChannel { + friend class GPURemoteChannelInfo; + + GPURemoteChannel(uintptr_t _remote_ptr); + + public: + virtual bool support_idindexed_fields(Memory src_mem, Memory dst_mem) const + { + return true; + } }; class GPUfillChannel; diff --git a/src/realm/cuda/cuda_memcpy.cu b/src/realm/cuda/cuda_memcpy.cu index 7b8b44ad007..0e274b56d6f 100644 --- a/src/realm/cuda/cuda_memcpy.cu +++ b/src/realm/cuda/cuda_memcpy.cu @@ -139,8 +139,73 @@ memcpy_kernel_transpose(Realm::Cuda::MemcpyTransposeInfo info, T *tile template static __device__ inline void -memcpy_affine_batch(Realm::Cuda::AffineCopyPair *info, - size_t nrects, size_t start_offset = 0) +memcpy_multi_affine_batch(Realm::Cuda::AffineCopyPair *info, size_t nrects, + size_t start_offset = 0) +{ + const Offset_t blk_stride = blockDim.x; + const Offset_t tid_global = threadIdx.x; + + /* -------- iterate over copy rectangles -------- */ + for(size_t r = 0; r < nrects; ++r) { + auto &cp = info[r]; + Offset_t v = cp.volume; // elements in one field + Offset_t n = max(size_t(1), max(cp.src.num_fields, cp.dst.num_fields)); + + const T *__restrict__ src = reinterpret_cast(cp.src.addr); + T *__restrict__ dst = reinterpret_cast(cp.dst.addr); + + /* -------- iterate over fields handled by this block -------- */ + for(Offset_t f = blockIdx.x; f < n; f += gridDim.x) { + + Offset_t src_field_base = cp.src.num_fields > 0 + ? cp.src.fields[f] * cp.src.field_stride + : f * cp.src.field_stride; + + Offset_t dst_field_base = cp.dst.num_fields > 0 + ? cp.dst.fields[f] * cp.dst.field_stride + : f * cp.dst.field_stride; + + Offset_t off = tid_global; + + while(off < v) { + /* -------- issue up to MAX_UNROLL loads ---------- */ + T buf[MAX_UNROLL]; + unsigned loaded = 0; + +#pragma unroll + for(unsigned i = 0; i < MAX_UNROLL; ++i) { + Offset_t idx = off + i * blk_stride; + if(idx >= v) { + break; + } + + Offset_t src_coords[N]; + index_to_coords(src_coords, idx, cp.extents); + Offset_t src_lin = coords_to_index(src_coords, cp.src.strides); + buf[i] = src[src_field_base + src_lin]; + ++loaded; + } + +/* -------- corresponding stores ------------------ */ +#pragma unroll + for(unsigned i = 0; i < loaded; ++i) { + Offset_t idx = off + i * blk_stride; + Offset_t dst_coords[N]; + index_to_coords(dst_coords, idx, cp.extents); + Offset_t dst_lin = coords_to_index(dst_coords, cp.dst.strides); + dst[dst_field_base + dst_lin] = buf[i]; + } + + off += loaded * blk_stride; + } // while off < v + } // for each field handled by this block + } // for each rect +} + +template +static __device__ inline void +memcpy_affine_batch(Realm::Cuda::AffineCopyPair *info, size_t nrects, + size_t start_offset = 0) { Offset_t offset = blockIdx.x * blockDim.x + threadIdx.x - start_offset; const unsigned grid_stride = gridDim.x * blockDim.x; @@ -170,8 +235,7 @@ memcpy_affine_batch(Realm::Cuda::AffineCopyPair *info, for(unsigned j = 0; j < i; j++) { Offset_t dst_coords[N]; - index_to_coords(dst_coords, - (offset + j * grid_stride), + index_to_coords(dst_coords, (offset + j * grid_stride), current_info.extents); const size_t dst_idx = @@ -247,7 +311,7 @@ memcpy_indirect_points(Realm::Cuda::MemcpyIndirectInfo<3, Offset_t> info) template static __device__ inline void -memfill_affine_batch(const Realm::Cuda::AffineFillInfo& info) +memfill_affine_batch(const Realm::Cuda::AffineFillInfo &info) { Offset_t offset = blockIdx.x * blockDim.x + threadIdx.x; const unsigned grid_stride = gridDim.x * blockDim.x; @@ -277,15 +341,24 @@ memfill_affine_batch(const Realm::Cuda::AffineFillInfo& info) } } -#define MEMCPY_TEMPLATE_INST(type, dim, offt, name) \ - extern "C" __global__ __launch_bounds__(256, 4) void \ - memcpy_affine_batch##name(Realm::Cuda::AffineCopyInfo info) { \ - memcpy_affine_batch(info.subrects, info.num_rects); \ +#define MEMCPY_MULTI_TEMPLATE_INST(type, dim, offt, name) \ + extern "C" __global__ __launch_bounds__(256, 4) void multi_affine_batch##name( \ + Realm::Cuda::AffineCopyInfo info) \ + { \ + memcpy_multi_affine_batch(info.subrects, info.num_rects); \ + } + +#define MEMCPY_TEMPLATE_INST(type, dim, offt, name) \ + extern "C" __global__ __launch_bounds__(256, 4) void memcpy_affine_batch##name( \ + Realm::Cuda::AffineCopyInfo info) \ + { \ + memcpy_affine_batch(info.subrects, info.num_rects); \ } #define FILL_TEMPLATE_INST(type, dim, offt, name) \ extern "C" __global__ void fill_affine_batch##name( \ - Realm::Cuda::AffineFillInfo info) { \ + Realm::Cuda::AffineFillInfo info) \ + { \ memfill_affine_batch(info); \ } @@ -311,16 +384,17 @@ memfill_affine_batch(const Realm::Cuda::AffineFillInfo& info) #define INST_TEMPLATES(type, sz, dim, off) \ MEMCPY_TEMPLATE_INST(type, dim, off, dim##D_##sz) \ + MEMCPY_MULTI_TEMPLATE_INST(type, dim, off, dim##D_##sz) \ FILL_TEMPLATE_INST(type, dim, off, dim##D_##sz) \ FILL_LARGE_TEMPLATE_INST(type, dim, off, dim##D_##sz) \ MEMCPY_INDIRECT_TEMPLATE_INST(int, type, dim, off, dim##D_##sz##32) \ MEMCPY_INDIRECT_TEMPLATE_INST(long long, type, dim, off, dim##D_##sz##64) -#define INST_TEMPLATES_FOR_TYPES(dim, off) \ - INST_TEMPLATES(unsigned char, 8, dim, off) \ - INST_TEMPLATES(unsigned short, 16, dim, off) \ - INST_TEMPLATES(unsigned int, 32, dim, off) \ - INST_TEMPLATES(unsigned long long, 64, dim, off) \ +#define INST_TEMPLATES_FOR_TYPES(dim, off) \ + INST_TEMPLATES(unsigned char, 8, dim, off) \ + INST_TEMPLATES(unsigned short, 16, dim, off) \ + INST_TEMPLATES(unsigned int, 32, dim, off) \ + INST_TEMPLATES(unsigned long long, 64, dim, off) \ INST_TEMPLATES(uint4, 128, dim, off) #define INST_TEMPLATES_FOR_DIMS() \ diff --git a/src/realm/cuda/cuda_memcpy.h b/src/realm/cuda/cuda_memcpy.h index 7573909ba72..4c6c4f92394 100644 --- a/src/realm/cuda/cuda_memcpy.h +++ b/src/realm/cuda/cuda_memcpy.h @@ -31,10 +31,15 @@ namespace Realm { template struct alignas(8) AffineSubRect { + using FieldID = int; // Extent of the ND array Offset_t strides[N - 1]; // Address of the ND array uintptr_t addr; + + const FieldID *fields; + size_t num_fields; + Offset_t field_stride; }; template diff --git a/src/realm/cuda/cuda_module.cc b/src/realm/cuda/cuda_module.cc index a92cba58f60..4991b3a1c6c 100644 --- a/src/realm/cuda/cuda_module.cc +++ b/src/realm/cuda/cuda_module.cc @@ -1142,7 +1142,8 @@ namespace Realm { } void GPU::launch_batch_affine_kernel(void *copy_info, size_t dim, size_t elem_size, - size_t volume, GPUStream *stream) + size_t volume, bool mutlfield_optimized, + GPUStream *stream) { size_t log_elem_size = std::min(static_cast(ctz(elem_size)), CUDA_MEMCPY_KERNEL_MAX2_LOG2_BYTES - 1); @@ -1151,10 +1152,13 @@ namespace Realm { assert(dim <= REALM_MAX_DIM); assert(dim >= 1); - // TODO: probably replace this - // with a better data-structure - GPUFuncInfo &func_info = batch_affine_kernels[dim - 1][log_elem_size]; - launch_kernel(func_info, copy_info, volume, stream); + if(!mutlfield_optimized) { + GPUFuncInfo &func_info = batch_affine_kernels[dim - 1][log_elem_size]; + launch_kernel(func_info, copy_info, volume, stream); + } else { + GPUFuncInfo &func_info = multi_batch_affine_kernels[dim - 1][log_elem_size]; + launch_kernel(func_info, copy_info, volume, stream); + } } const GPU::CudaIpcMapping *GPU::find_ipc_mapping(Memory mem) const @@ -2095,6 +2099,15 @@ namespace Realm { 0)); batch_affine_kernels[d - 1][log_bit_sz] = func_info; + std::snprintf(name, sizeof(name), "multi_affine_batch%uD_%u", d, bit_sz); + CHECK_CU(CUDA_DRIVER_FNPTR(cuModuleGetFunction)(&func_info.func, device_module, + name)); + + CHECK_CU(CUDA_DRIVER_FNPTR(cuOccupancyMaxPotentialBlockSize)( + &func_info.occ_num_blocks, &func_info.occ_num_threads, func_info.func, 0, 0, + 0)); + multi_batch_affine_kernels[d - 1][log_bit_sz] = func_info; + std::snprintf(name, sizeof(name), "fill_affine_large%uD_%u", d, bit_sz); CHECK_CU(CUDA_DRIVER_FNPTR(cuModuleGetFunction)(&func_info.func, device_module, name)); diff --git a/src/realm/inst_layout.h b/src/realm/inst_layout.h index a67b535790b..b7635cab5d3 100644 --- a/src/realm/inst_layout.h +++ b/src/realm/inst_layout.h @@ -227,6 +227,7 @@ namespace Realm { using FieldMap = std::map; FieldMap fields; + bool idindexed_fields{false}; }; REALM_PUBLIC_API @@ -421,6 +422,9 @@ namespace Realm { IndexSpace space; std::vector> piece_lists; + // Pre-computed dimension ordering for idindexed_fields + std::vector preferred_dim_order; + static Serialization::PolymorphicSerdezSubclass> serdez_subclass; diff --git a/src/realm/inst_layout.inl b/src/realm/inst_layout.inl index 0ee4db69603..e7597aa6bb2 100644 --- a/src/realm/inst_layout.inl +++ b/src/realm/inst_layout.inl @@ -138,6 +138,9 @@ namespace Realm { // track that std::map, size_t> pl_indexes, pl_starts, pl_sizes; + size_t field_stride = 0; + layout->idindexed_fields = true; + // reserve space so that we don't have to copy piece lists as we grow layout->piece_lists.reserve(ilc.field_groups.size()); for(size_t i = 0; i < ilc.field_groups.size(); i++) { @@ -147,6 +150,7 @@ namespace Realm { // pieces size_t gsize = 0; size_t galign = 1; + // we can't set field offsets in a single pass because we don't know // the whole group's alignment until we look at every field std::map field_offsets; @@ -166,6 +170,7 @@ namespace Realm { } // increase size and alignment if needed gsize = max(gsize, offset + it2->size); + if((it2->alignment > 1) && ((galign % it2->alignment) != 0)) galign = lcm(galign, size_t(it2->alignment)); field_offsets[it2->field_id] = offset; @@ -182,8 +187,10 @@ namespace Realm { std::pair pl_key(gsize, galign); std::map, size_t>::const_iterator it = pl_indexes.find(pl_key); + size_t li; size_t reuse_offset; + if(it != pl_indexes.end()) { li = it->second; size_t piece_start = round_up(layout->bytes_used, galign); @@ -229,6 +236,9 @@ namespace Realm { // final value of stride is total bytes used by piece - use that // to set new instance footprint layout->bytes_used = piece_start + stride; + if(field_stride == 0) { + field_stride = stride; + } pl.pieces.push_back(piece); } @@ -242,6 +252,10 @@ namespace Realm { it2 != field_offsets.end(); ++it2) { // should not have seen this field before assert(layout->fields.count(it2->first) == 0); + + layout->idindexed_fields &= + (static_cast(it2->first * field_stride) == reuse_offset); + InstanceLayoutGeneric::FieldLayout &fl = layout->fields[it2->first]; fl.list_idx = li; fl.rel_offset = /*group_offset +*/ it2->second + reuse_offset; @@ -249,6 +263,61 @@ namespace Realm { } } + // Compute preferred dimension ordering for idindexed_fields + if(layout->idindexed_fields) { + layout->preferred_dim_order.clear(); + std::vector preferred; + preferred.reserve(N); + + // Consider all fields + for(InstanceLayoutGeneric::FieldMap::const_iterator it = layout->fields.begin(); + it != layout->fields.end(); ++it) { + const InstancePieceList &ipl = layout->piece_lists[it->second.list_idx]; + + for(typename std::vector *>::const_iterator it2 = + ipl.pieces.begin(); + it2 != ipl.pieces.end(); ++it2) { + + const AffineLayoutPiece *affine = + static_cast *>(*it2); + + preferred.clear(); + for(int d = 0; d < N; d++) { + preferred.push_back(d); + } + + // Sort dimensions by stride + std::sort(preferred.begin(), preferred.end(), [&](int a, int b) { + return affine->strides[a] < affine->strides[b]; + }); + + // Reconcile dimensions orders + if(preferred.size() > layout->preferred_dim_order.size()) { + if(std::equal(layout->preferred_dim_order.begin(), + layout->preferred_dim_order.end(), preferred.begin())) { + layout->preferred_dim_order = preferred; + } + } + + preferred.clear(); + } + } + + // If we didn't end up choosing all the dimensions, add the rest back in + // in ascending order + if(layout->preferred_dim_order.size() != N) { + std::vector present(N, false); + for(size_t i = 0; i < layout->preferred_dim_order.size(); i++) { + present[layout->preferred_dim_order[i]] = true; + } + for(int i = 0; i < N; i++) { + if(!present[i]) { + layout->preferred_dim_order.push_back(i); + } + } + } + } + return layout; } @@ -511,7 +580,8 @@ namespace Realm { { InstanceLayout *il = new InstanceLayout; if((s >> il->bytes_used) && (s >> il->alignment_reqd) && (s >> il->fields) && - (s >> il->space) && (s >> il->piece_lists)) { + (s >> il->idindexed_fields) && (s >> il->space) && (s >> il->piece_lists) && + (s >> il->preferred_dim_order)) { return il; } else { delete il; @@ -530,7 +600,9 @@ namespace Realm { copy->bytes_used = bytes_used; copy->alignment_reqd = alignment_reqd; copy->fields = fields; + copy->idindexed_fields = idindexed_fields; copy->space = space; + copy->preferred_dim_order = preferred_dim_order; copy->piece_lists.resize(piece_lists.size()); for(size_t i = 0; i < piece_lists.size(); i++) { copy->piece_lists[i].pieces.resize(piece_lists[i].pieces.size()); @@ -596,8 +668,9 @@ namespace Realm { template inline bool InstanceLayout::serialize(S &s) const { - return ((s << bytes_used) && (s << alignment_reqd) && (s << fields) && (s << space) && - (s << piece_lists)); + return ((s << bytes_used) && (s << alignment_reqd) && (s << fields) && + (s << idindexed_fields) && (s << space) && (s << piece_lists) && + (s << preferred_dim_order)); } //////////////////////////////////////////////////////////////////////// diff --git a/src/realm/runtime_impl.cc b/src/realm/runtime_impl.cc index 50de5eaed09..59282320ce4 100644 --- a/src/realm/runtime_impl.cc +++ b/src/realm/runtime_impl.cc @@ -758,9 +758,11 @@ namespace Realm { config_map.insert({"pin_util_procs", &pin_util_procs}); config_map.insert({"use_ext_sysmem", &use_ext_sysmem}); config_map.insert({"regmem", ®_mem_size}); + config_map.insert({"ib_regmem", ®_ib_mem_size}); config_map.insert({"report_sparsity_leaks", &report_sparsity_leaks}); config_map.insert({"barrier_broadcast_radix", &barrier_broadcast_radix}); config_map.insert({"diskmem", &disk_mem_size}); + config_map.insert({"dma_multi_field", &dma_multi_field}); resource_map.insert({"cpu", &res_num_cpus}); resource_map.insert({"sysmem", &res_sysmem_size}); @@ -815,6 +817,7 @@ namespace Realm { .add_option_int_units("-ll:ib_rsize", reg_ib_mem_size, 'm') .add_option_int_units("-ll:dsize", disk_mem_size, 'm') .add_option_int("-ll:dma", dma_worker_threads) + .add_option_int("-ll:dma_multi_field", dma_multi_field) .add_option_bool("-ll:pin_dma", pin_dma_threads) .add_option_int("-ll:dummy_rsrv_ok", dummy_reservation_ok) .add_option_bool("-ll:show_rsrv", show_reservations) diff --git a/src/realm/runtime_impl.h b/src/realm/runtime_impl.h index 72b72c75feb..8624cb433b2 100644 --- a/src/realm/runtime_impl.h +++ b/src/realm/runtime_impl.h @@ -145,6 +145,7 @@ namespace Realm { size_t reg_mem_size = 0; size_t disk_mem_size = 0; unsigned dma_worker_threads = 0; // unused - warning on application use + bool dma_multi_field = true; #ifdef EVENT_TRACING size_t event_trace_block_size = 1 << 20; double event_trace_exp_arrv_rate = 1e3; diff --git a/src/realm/transfer/channel.cc b/src/realm/transfer/channel.cc index 9f2c609abf1..adb7b6afb9b 100644 --- a/src/realm/transfer/channel.cc +++ b/src/realm/transfer/channel.cc @@ -4195,7 +4195,7 @@ namespace Realm { { unsigned bw = 5000; // HACK - estimate at 5 GB/s unsigned latency = 2000; // HACK - estimate at 2 us - unsigned frag_overhead = 1000; // HACK - estimate at 1 us + unsigned frag_overhead = 2000; // HACK - estimate at 2 us // any combination of SYSTEM/REGDMA/Z_COPY/SOCKET_MEM // for(size_t i = 0; i < num_cpu_mem_kinds; i++) // add_path(cpu_mem_kinds[i], false, diff --git a/src/realm/transfer/channel.h b/src/realm/transfer/channel.h index e693d0ff01a..ad886b82969 100644 --- a/src/realm/transfer/channel.h +++ b/src/realm/transfer/channel.h @@ -56,7 +56,7 @@ namespace Realm { typedef unsigned long long XferDesID; -// clang-format off + // clang-format off #define REALM_XFERDES_KINDS(__op__) \ __op__(XFER_NONE) \ __op__(XFER_DISK_READ) \ @@ -739,6 +739,11 @@ namespace Realm { virtual bool supports_redop(ReductionOpID redop_id) const; + virtual bool support_idindexed_fields(Memory src_mem, Memory dst_mem) const + { + return false; + }; + // attempt to make progress on the specified xferdes virtual long progress_xd(XferDes *xd, long max_nr); diff --git a/src/realm/transfer/transfer.cc b/src/realm/transfer/transfer.cc index 54020a174fb..c999b8839cc 100644 --- a/src/realm/transfer/transfer.cc +++ b/src/realm/transfer/transfer.cc @@ -817,6 +817,253 @@ namespace Realm { return true; } + //////////////////////////////////////////////////////////////////////// + // + // class IDIndexedFieldsIterator + // + + template + IDIndexedFieldsIterator::IDIndexedFieldsIterator( + const int _dim_order[N], const std::vector &_fields, size_t _field_size, + RegionInstanceImpl *_inst_impl, const IndexSpace &_is, + ReplicatedHeap *_repl_heap) + : TransferIteratorBase(_inst_impl, _dim_order) + , is(_is) + , repl_heap(_repl_heap) + , field_block(FieldBlock::create(*_repl_heap, _fields.data(), _fields.size())) + { + if(is.is_valid()) { + reset_internal(); + } else { + iter_init_deferred = true; + } + + if(iter_init_deferred || iter.valid) { + fields = _fields; + field_size = _field_size; + inst_layout = + checked_cast *>(this->inst_impl->metadata.layout); + } + } + + template + template + /*static*/ TransferIterator * + IDIndexedFieldsIterator::deserialize_new(S &deserializer) + { + IndexSpace is; + RegionInstance inst; + std::vector fields; + size_t field_size; + int dim_order[N]; + + if(!((deserializer >> is) && (deserializer >> inst) && (deserializer >> fields) && + (deserializer >> field_size))) { + return 0; + } + + for(int i = 0; i < N; i++) { + if(!(deserializer >> dim_order[i])) { + return 0; + } + } + + IDIndexedFieldsIterator *tiis = new IDIndexedFieldsIterator( + dim_order, fields, field_size, get_runtime()->get_instance_impl(inst), is, + &get_runtime()->repl_heap); + + return tiis; + } + + template + IDIndexedFieldsIterator::~IDIndexedFieldsIterator(void) + { + if(field_block) { + repl_heap->free_obj(field_block); + } + } + + template + Event IDIndexedFieldsIterator::request_metadata(void) + { + Event e = TransferIteratorBase::request_metadata(); + + if(iter_init_deferred) { + e = Event::merge_events(e, is.make_valid()); + } + + return e; + } + + template + void IDIndexedFieldsIterator::reset(void) + { + TransferIteratorBase::reset(); + rect_idx = 0; + reset_internal(); + } + + template + void IDIndexedFieldsIterator::reset_internal(void) + { + // assert(!iter_init_deferred); + if(!sparsity_impl) { + assert(is.is_valid()); + iter.reset(is); + } else { + iter.reset(is.bounds, is.bounds, + reinterpret_cast *>(sparsity_impl)); + } + iter_init_deferred = false; + this->is_done = !iter.valid; + } + + template + bool + IDIndexedFieldsIterator::get_addresses(AddressList &addrlist, + const InstanceLayoutPieceBase *&nonaffine) + { + nonaffine = 0; + + if(rect_idx == 0) { + addrlist.attach_field_block(field_block); + } + + while(!this->done()) { + if(!this->have_rect) { + return false; + } + + const InstancePieceList &piece_list = + inst_layout->piece_lists[inst_layout->fields.begin()->second.list_idx]; + const InstanceLayoutPiece *layout_piece = + piece_list.find_piece(this->cur_point); + + assert(layout_piece->layout_type == PieceLayoutTypes::AffineLayoutType); + + Rect target_subrect; + this->have_rect = + compute_target_subrect(layout_piece->bounds, this->cur_rect, this->cur_point, + target_subrect, &this->dim_order[0]); + + size_t contig_bytes = 0; + size_t total_bytes = 0; + + const AffineLayoutPiece *affine = + static_cast *>(layout_piece); + + // assert(this->inst_impl->metadata.is_valid()); + size_t base_offset = this->inst_impl->metadata.inst_offset + affine->offset + + affine->strides.dot(target_subrect.lo); //+ field_rel_offset; + +#ifdef DEBUG_REALM + assert(layout_piece->bounds.contains(target_subrect)); +#endif + + assert(layout_piece->layout_type == PieceLayoutTypes::AffineLayoutType); + + // TODO(apryakhin@): If that's the same rec consider caching it + std::unordered_map> count_strides; + int ndims = compact_affine_dims( + static_cast *>(layout_piece), target_subrect, + this->dim_order, field_size, total_bytes, contig_bytes, count_strides); + assert(ndims > 0); + assert(total_bytes > 0); + assert(contig_bytes > 0); + + bool commited = addrlist.append_entry(ndims, contig_bytes, total_bytes, base_offset, + count_strides); + assert(commited); + } + + return true; + } + + template + size_t IDIndexedFieldsIterator::step(size_t max_bytes, + TransferIterator::AddressInfo &info, + unsigned flags, bool tentative /*= false*/) + { + // NOT SUPPORTED + assert(0); + return 0; + } + + template + size_t + IDIndexedFieldsIterator::step_custom(size_t max_bytes, + TransferIterator::AddressInfoCustom &info, + bool tentative /*= false*/) + { + // NOT SUPPORTED + assert(0); + return 0; + } + + template + void IDIndexedFieldsIterator::confirm_step(void) + { + // NOT SUPPORTED + assert(0); + } + + template + void IDIndexedFieldsIterator::cancel_step(void) + { + // NOT SUPPORTED + assert(0); + } + + template + bool IDIndexedFieldsIterator::get_next_rect(Rect &r, FieldID &fid, + size_t &offset, size_t &fsize) + { + if(iter_init_deferred) { + reset_internal(); + if(!iter.valid) { + this->is_done = true; + return false; + } + } + + if(this->is_done) { + return false; + } + + r = iter.rect; + rect_idx++; + + iter.step(); + if(!iter.valid) { + reset_internal(); + this->is_done = true; + } + return true; + } + + template + /*static*/ Serialization::PolymorphicSerdezSubclass> + IDIndexedFieldsIterator::serdez_subclass; + + template + template + bool IDIndexedFieldsIterator::serialize(S &serializer) const + { + if(!((serializer << iter.space) && (serializer << this->inst_impl->me) && + (serializer << fields) && (serializer << field_size))) { + return false; + } + + for(int i = 0; i < N; i++) { + if(!(serializer << this->dim_order[i])) { + return false; + } + } + + return true; + } + //////////////////////////////////////////////////////////////////////// // // class WrappingTransferIteratorIndirect @@ -1622,6 +1869,7 @@ namespace Realm { class TransferDomainIndexSpace : public TransferDomain { public: TransferDomainIndexSpace(IndexSpace _is); + ~TransferDomainIndexSpace(); template static TransferDomain *deserialize_new(S &deserializer); @@ -1648,7 +1896,8 @@ namespace Realm { const std::vector &dim_order, const std::vector &fields, const std::vector &fld_offsets, - const std::vector &fld_sizes) const; + const std::vector &fld_sizes, + bool idindexed_fields = false) const; virtual TransferIterator *create_iterator(RegionInstance inst, RegionInstance peer, const std::vector &fields, @@ -1664,6 +1913,8 @@ namespace Realm { template bool serialize(S &serializer) const; + static constexpr size_t MIN_IDINDEXED_FIELDS = 2; + // protected: IndexSpace is; }; @@ -1673,6 +1924,10 @@ namespace Realm { : is(_is) {} + template + TransferDomainIndexSpace::~TransferDomainIndexSpace() + {} + template template /*static*/ TransferDomain * @@ -1824,6 +2079,69 @@ namespace Realm { std::vector preferred; preferred.reserve(N); + // Fast path for dst and src with idindexed_fields and pre-computed ordering + RegionInstanceImpl *dst_impl = nullptr; + const InstanceLayout *dst_layout = nullptr; + RegionInstanceImpl *src_impl = nullptr; + const InstanceLayout *src_layout = nullptr; + + // Check destination + if((dsts[0].field_id != FieldID(-1)) && dsts[0].inst.exists() && + dsts[0].indirect_index == -1) { + dst_impl = get_runtime()->get_instance_impl(dsts[0].inst); + dst_layout = checked_cast *>(dst_impl->metadata.layout); + } + + // Check source + if((srcs[0].field_id != FieldID(-1)) && srcs[0].inst.exists() && + srcs[0].indirect_index == -1) { + src_impl = get_runtime()->get_instance_impl(srcs[0].inst); + src_layout = checked_cast *>(src_impl->metadata.layout); + } + + // Check if both layouts support the fast path optimization + bool layouts_support_fastpath = + (dst_layout && dst_layout->idindexed_fields && + !dst_layout->preferred_dim_order.empty() && src_layout && + src_layout->idindexed_fields && !src_layout->preferred_dim_order.empty()); + + if(layouts_support_fastpath) { + + for(int d : dst_layout->preferred_dim_order) { + if(!trivial[d]) { + dim_order.push_back(d); + } + } + + preferred.clear(); + for(int d : src_layout->preferred_dim_order) { + if(!trivial[d]) { + preferred.push_back(d); + } + } + + reconcile_dim_orders(dim_order, preferred); + + // if we didn't end up choosing all the dimensions, add the rest back in + // in ascending order + if(dim_order.size() != N) { + std::vector present(N, false); + for(size_t i = 0; i < dim_order.size(); i++) { + present[dim_order[i]] = true; + } + for(int i = 0; i < N; i++) { + if(!present[i]) { + dim_order.push_back(i); + } + } +#ifdef DEBUG_REALM + assert(dim_order.size() == N); +#endif + } + + return; + } + // consider destinations first for(size_t i = 0; i < dsts.size(); i++) { if((dsts[i].field_id != FieldID(-1)) && dsts[i].inst.exists()) { @@ -1831,7 +2149,6 @@ namespace Realm { max_stride); reconcile_dim_orders(dim_order, preferred); preferred.clear(); - continue; } // TODO: ask opinion of indirections? @@ -1844,7 +2161,6 @@ namespace Realm { max_stride); reconcile_dim_orders(dim_order, preferred); preferred.clear(); - continue; } // TODO: ask opinion of indirections? @@ -1925,10 +2241,30 @@ namespace Realm { fragments.assign(N + 2, 0); - for(size_t i = 0; i < fields.size(); i++) { - FieldID fid = fields[i]; - size_t field_size = fld_sizes[i]; + // Determine processing strategy: bulk (for idindexed_fields) or individual + bool use_bulk_processing = false; + size_t fields_to_process = fields.size(); // Process one field per iteration; + if(inst_layout->idindexed_fields && fields.size() >= 2) { + use_bulk_processing = true; + fields_to_process = 1; // Process all fields at once + } + + for(size_t field_batch = 0; field_batch < fields_to_process; field_batch++) { + // Determine field info for this batch + FieldID fid; + size_t field_size, effective_field_count; + if(use_bulk_processing) { + fid = fields[0]; + field_size = fld_sizes[0]; + effective_field_count = fields.size(); + } else { + fid = fields[field_batch]; + field_size = fld_sizes[field_batch]; + effective_field_count = 1; + } + + // Get layout pieces for representative field const InstancePieceList *ipl; { InstanceLayoutGeneric::FieldMap::const_iterator it = @@ -1944,18 +2280,18 @@ namespace Realm { assert(layout_piece != 0); if(layout_piece->bounds.contains(is)) { - // easy case: one piece covers our entire domain and the iteration order + // Easy case: one piece covers our entire domain and the iteration order // doesn't impact the fragment count if(layout_piece->layout_type == PieceLayoutTypes::AffineLayoutType) { const AffineLayoutPiece *affine = static_cast *>(layout_piece); do { - add_fragments_for_rect(isi.rect, field_size, 1 /*field count*/, + add_fragments_for_rect(isi.rect, field_size, effective_field_count, affine->strides, dim_order, fragments); isi.step(); } while(isi.valid); } else { - // not affine - add one fragment for each rectangle + // Not affine - add one fragment for each rectangle size_t num_rects; if(is.dense()) { num_rects = 1; @@ -1973,7 +2309,7 @@ namespace Realm { do { Point next_start = isi.rect.lo; while(true) { - // look up new piece if needed + // Look up new piece if needed if(!layout_piece->bounds.contains(next_start)) { layout_piece = ipl->find_piece(next_start); assert(layout_piece != 0); @@ -1985,7 +2321,7 @@ namespace Realm { if(layout_piece->layout_type == PieceLayoutTypes::AffineLayoutType) { const AffineLayoutPiece *affine = static_cast *>(layout_piece); - add_fragments_for_rect(isi.rect, field_size, 1 /*field count*/, + add_fragments_for_rect(isi.rect, field_size, effective_field_count, affine->strides, dim_order, fragments); } else { non_affine_rects++; @@ -2012,12 +2348,24 @@ namespace Realm { TransferIterator *TransferDomainIndexSpace::create_iterator( RegionInstance inst, const std::vector &dim_order, const std::vector &fields, const std::vector &fld_offsets, - const std::vector &fld_sizes) const + const std::vector &fld_sizes, bool idindexed_fields) const { assert(dim_order.size() == N); RegionInstanceImpl *impl = get_runtime()->get_instance_impl(inst); - return new TransferIteratorIndexSpace(dim_order.data(), fields, fld_offsets, - fld_sizes, impl, is); + const InstanceLayout *inst_layout = + checked_cast *>(impl->metadata.layout); + if(idindexed_fields && inst_layout->idindexed_fields && is.dense() && + fields.size() >= MIN_IDINDEXED_FIELDS) { + // TODO(apryakhin@): There is an untested path where either src or dst + // might not have idindexed_fields layout. There aren't any checks that prevents + // this from running either. + return new IDIndexedFieldsIterator(dim_order.data(), fields, + fld_sizes.front(), impl, is, + &get_runtime()->repl_heap); + } else { + return new TransferIteratorIndexSpace(dim_order.data(), fields, fld_offsets, + fld_sizes, impl, is); + } } template @@ -3540,16 +3888,16 @@ namespace Realm { << " srcs=" << srcs.size() << " dsts=" << dsts.size(); if(log_xplan.want_debug()) { for(size_t i = 0; i < srcs.size(); i++) { - log_xplan.debug() << "created: plan=" << (void *)this << " srcs[" << i - << "]=" << srcs[i]; + log_xplan.info() << "created: plan=" << (void *)this << " srcs[" << i + << "]=" << srcs[i]; } for(size_t i = 0; i < dsts.size(); i++) { - log_xplan.debug() << "created: plan=" << (void *)this << " dsts[" << i - << "]=" << dsts[i]; + log_xplan.info() << "created: plan=" << (void *)this << " dsts[" << i + << "]=" << dsts[i]; } for(size_t i = 0; i < indirects.size(); i++) { - log_xplan.debug() << "created: plan=" << (void *)this << " indirects[" << i - << "]=" << *indirects[i]; + log_xplan.info() << "created: plan=" << (void *)this << " indirects[" << i + << "]=" << *indirects[i]; } } @@ -3629,8 +3977,13 @@ namespace Realm { min_granularity = combined_field_size; } + size_t max_ib_size = 0; + RealmStatus status = + get_runtime()->get_module_config("core")->get_property("ib_regmem", max_ib_size); + assert(status == REALM_SUCCESS); + size_t ib_size = domain_size * element_size + serdez_pad; - const size_t IB_MAX_SIZE = 16 << 20; // 16MB + const size_t IB_MAX_SIZE = max_ib_size; // 16 << 20; // 16MB if(ib_size > IB_MAX_SIZE) { // take up to IB_MAX_SIZE, respecting the min granularity if(min_granularity > 1) { @@ -3724,8 +4077,9 @@ namespace Realm { // for now, pick a global dimension ordering // TODO: allow this to vary for independent subgraphs (or dependent ones // with transposes in line) - domain->choose_dim_order(dim_order, srcs, dsts, indirects, - false /*!force_fortran_order*/, 65536 /*max_stride*/); + + domain->choose_dim_order(dim_order, srcs, dsts, indirects, (domain->volume() == 1), + 65536 /*max_stride*/); src_fields.resize(srcs.size()); dst_fields.resize(dsts.size()); @@ -4032,6 +4386,17 @@ namespace Realm { xdn.gather_control_input = -1; xdn.scatter_control_input = -1; xdn.target_node = path_info.xd_channels[j]->node; + + bool enable_multi_field = false; + RealmStatus success = + get_runtime()->get_module_config("core")->get_property( + "dma_multi_field", enable_multi_field); + assert(success == REALM_SUCCESS); + + xdn.idindexed_fields = + enable_multi_field && + path_info.xd_channels[j]->support_idindexed_fields(src_mem, dst_mem); + xdn.channel = path_info.xd_channels[j]; xdn.inputs.resize(1); xdn.inputs[0] = ((j == 0) ? TransferGraph::XDTemplate::mk_inst( @@ -4198,49 +4563,49 @@ namespace Realm { IBAllocOrderSorter(graph.ib_edges)); } - if(log_xplan.want_debug()) { - log_xplan.debug() << "analysis: plan=" << (void *)this - << " dim_order=" << PrettyVector(dim_order) - << " xds=" << graph.xd_nodes.size() - << " ibs=" << graph.ib_edges.size(); - - for(size_t i = 0; i < graph.xd_nodes.size(); i++) { - if(graph.xd_nodes[i].redop.id != 0) { - log_xplan.debug() - << "analysis: plan=" << (void *)this << " xds[" << i - << "]: target=" << graph.xd_nodes[i].target_node << " inputs=" - << PrettyVector(graph.xd_nodes[i].inputs) - << " outputs=" - << PrettyVector(graph.xd_nodes[i].outputs) - << " channel=" - << ((graph.xd_nodes[i].channel) ? graph.xd_nodes[i].channel->kind : -1) - << " redop=(" << graph.xd_nodes[i].redop.id << "," - << graph.xd_nodes[i].redop.is_fold << "," - << graph.xd_nodes[i].redop.in_place << ")"; - } else { - log_xplan.debug() - << "analysis: plan=" << (void *)this << " xds[" << i - << "]: target=" << graph.xd_nodes[i].target_node << " inputs=" - << PrettyVector(graph.xd_nodes[i].inputs) - << " outputs=" - << PrettyVector(graph.xd_nodes[i].outputs) - << " channel=" - << ((graph.xd_nodes[i].channel) ? graph.xd_nodes[i].channel->kind : -1); - } + // if(log_xplan.want_debug()) { + log_xplan.info() << "analysis: plan=" << (void *)this + << " dim_order=" << PrettyVector(dim_order) + << " xds=" << graph.xd_nodes.size() + << " ibs=" << graph.ib_edges.size(); + + for(size_t i = 0; i < graph.xd_nodes.size(); i++) { + if(graph.xd_nodes[i].redop.id != 0) { + log_xplan.info() + << "analysis: plan=" << (void *)this << " xds[" << i + << "]: target=" << graph.xd_nodes[i].target_node << " inputs=" + << PrettyVector(graph.xd_nodes[i].inputs) + << " outputs=" + << PrettyVector(graph.xd_nodes[i].outputs) + << " channel=" + << ((graph.xd_nodes[i].channel) ? graph.xd_nodes[i].channel->kind : -1) + << " redop=(" << graph.xd_nodes[i].redop.id << "," + << graph.xd_nodes[i].redop.is_fold << "," << graph.xd_nodes[i].redop.in_place + << ")"; + } else { + log_xplan.info() + << "analysis: plan=" << (void *)this << " xds[" << i + << "]: target=" << graph.xd_nodes[i].target_node << " inputs=" + << PrettyVector(graph.xd_nodes[i].inputs) + << " outputs=" + << PrettyVector(graph.xd_nodes[i].outputs) + << " channel=" + << ((graph.xd_nodes[i].channel) ? graph.xd_nodes[i].channel->kind : -1); } + } - for(size_t i = 0; i < graph.ib_edges.size(); i++) { - log_xplan.debug() << "analysis: plan=" << (void *)this << " ibs[" << i - << "]: memory=" << graph.ib_edges[i].memory << ":" - << graph.ib_edges[i].memory.kind() - << " size=" << graph.ib_edges[i].size; - } + for(size_t i = 0; i < graph.ib_edges.size(); i++) { + log_xplan.info() << "analysis: plan=" << (void *)this << " ibs[" << i + << "]: memory=" << graph.ib_edges[i].memory << ":" + << graph.ib_edges[i].memory.kind() + << " size=" << graph.ib_edges[i].size; + } - if(!graph.ib_edges.empty()) { - log_xplan.debug() << "analysis: plan=" << (void *)this - << " ib_alloc=" << PrettyVector(graph.ib_alloc_order); - } + if(!graph.ib_edges.empty()) { + log_xplan.info() << "analysis: plan=" << (void *)this + << " ib_alloc=" << PrettyVector(graph.ib_alloc_order); } + //} // mark that the analysis is complete and see if there are any pending // ops that can start allocating ibs @@ -4672,7 +5037,8 @@ namespace Realm { src_sizes[k] = desc.src_fields[xdn.inputs[j].inst.fld_start + k].size; } ii.iter = desc.domain->create_iterator(xdn.inputs[j].inst.inst, desc.dim_order, - src_fields, src_offsets, src_sizes); + src_fields, src_offsets, src_sizes, + xdn.idindexed_fields); // use first field's serdez - they all have to be the same ii.serdez_id = desc.src_fields[xdn.inputs[j].inst.fld_start].serdez_id; ii.ib_offset = 0; @@ -4802,7 +5168,8 @@ namespace Realm { dst_sizes[k] = desc.dst_fields[xdn.outputs[j].inst.fld_start + k].size; } oi.iter = desc.domain->create_iterator(xdn.outputs[j].inst.inst, desc.dim_order, - dst_fields, dst_offsets, dst_sizes); + dst_fields, dst_offsets, dst_sizes, + xdn.idindexed_fields); // use first field's serdez - they all have to be the same oi.serdez_id = desc.dst_fields[xdn.outputs[j].inst.fld_start].serdez_id; oi.ib_offset = 0; @@ -5055,6 +5422,7 @@ namespace Realm { const ProfilingRequestSet &, Event, int) const; \ template class TransferIteratorIndexSpace; \ template class TransferIteratorIndirect; \ + template class IDIndexedFieldsIterator; \ template class WrappingTransferIteratorIndirect; \ template class TransferIteratorIndirectRange; \ template class AddressSplitXferDesFactory; \ diff --git a/src/realm/transfer/transfer.h b/src/realm/transfer/transfer.h index 8a68a240c1c..e9bae8c71d5 100644 --- a/src/realm/transfer/transfer.h +++ b/src/realm/transfer/transfer.h @@ -217,6 +217,68 @@ namespace Realm { size_t field_idx{0}; }; + //////////////////////////////////////////////////////////////////////// + // + // class IDIndexedFieldsIterator + // + + template + class IDIndexedFieldsIterator : public TransferIteratorBase { + protected: + IDIndexedFieldsIterator(void); + + public: + IDIndexedFieldsIterator(const int _dim_order[N], const std::vector &_fields, + size_t _field_size, RegionInstanceImpl *_inst_impl, + const IndexSpace &_is, ReplicatedHeap *_repl_heap); + + template + static TransferIterator *deserialize_new(S &deserializer); + + virtual ~IDIndexedFieldsIterator(void); + + Event request_metadata(void) override; + void reset(void) override; + + static Serialization::PolymorphicSerdezSubclass> + serdez_subclass; + + template + bool serialize(S &serializer) const; + + bool get_addresses(AddressList &addrlist, + const InstanceLayoutPieceBase *&nonaffine) override; + + size_t step(size_t max_bytes, TransferIterator::AddressInfo &info, unsigned flags, + bool tentative = false) override; + size_t step_custom(size_t max_bytes, TransferIterator::AddressInfoCustom &info, + bool tentative = false) override; + void confirm_step(void) override; + void cancel_step(void) override; + + protected: + void reset_internal(void); + + bool get_next_rect(Rect &r, FieldID &fid, size_t &offset, + size_t &fsize) override; + + IndexSpace is; + SparsityMapImpl *sparsity_impl{nullptr}; + IndexSpaceIterator iter; + bool iter_init_deferred{false}; + std::vector fields; + + // const InstanceLayoutPiece *layout_piece{nullptr}; + const InstanceLayout *inst_layout{nullptr}; + + size_t field_size{0}; + size_t rect_idx{0}; + + ReplicatedHeap *repl_heap{nullptr}; + FieldBlock *field_block{nullptr}; + }; + template class TransferIteratorIndirect : public TransferIteratorBase { protected: @@ -294,11 +356,12 @@ namespace Realm { const std::vector &fld_sizes, std::vector &fragments) const = 0; - virtual TransferIterator * - create_iterator(RegionInstance inst, const std::vector &dim_order, - const std::vector &fields, - const std::vector &fld_offsets, - const std::vector &fld_sizes) const = 0; + virtual TransferIterator *create_iterator(RegionInstance inst, + const std::vector &dim_order, + const std::vector &fields, + const std::vector &fld_offsets, + const std::vector &fld_sizes, + bool idindexed_fields = false) const = 0; virtual TransferIterator * create_iterator(RegionInstance inst, RegionInstance peer, @@ -324,6 +387,7 @@ namespace Realm { int scatter_control_input; XferDesRedopInfo redop; Channel *channel = nullptr; + bool idindexed_fields = false; enum IOType { diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index a6213d8b46f..bf0adf14fec 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -111,6 +111,7 @@ list( repl_heap_test.cc nodeset_test.cc transfer_iterator_test.cc + idindexed_fields_iterator_test.cc lowlevel_dma_test.cc circ_queue_test.cc gather_scatter_test.cc @@ -428,6 +429,10 @@ if(TEST_USE_GPU) target_link_libraries(cuda_memcpy_test CUDA::cudart) target_link_libraries(cuda_scatter_test CUDA::cudart) target_link_libraries(test_cuhook CUDA::cudart) + set(multifield_transfer_ARGS -ll:gpu 1) + set(multifield_transfer_RESOURCE_LOCK gpu) + add_integration_test(multifield_transfer "${REALM_TEST_DIR}/multifield_transfer.cc") + target_link_libraries(multifield_transfer CUDA::cudart) endif() set(transpose_test_gpu_ARGS -ll:gpu 1) set(transpose_test_gpu_RESOURCE_LOCK gpu) diff --git a/tests/multifield_transfer.cc b/tests/multifield_transfer.cc new file mode 100644 index 00000000000..6aed49989ae --- /dev/null +++ b/tests/multifield_transfer.cc @@ -0,0 +1,809 @@ +/* + * Copyright 2025 Stanford University, NVIDIA Corporation + * SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include +#include +#include + +#include +#include +#include + +using namespace Realm; + +const size_t MAX_DIM = 2; +// typedef size_t ElementType; +typedef unsigned short ElementType; + +typedef Realm::IndexSpace CopyIndexSpace; + +Logger log_app("app"); + +enum +{ + BENCH_TIMING_TASK = Processor::TASK_ID_FIRST_AVAILABLE + 0, + UPDATE_OP_TIMING_TASK, + TOP_LEVEL_TASK, + REMOTE_COPY_TASK, +}; + +namespace TestConfig { + bool verify = false; + bool enable_profiling = true; + bool enable_remote_copy = false; + bool graphviz = false; + int graph_type = 2; // 0 means isolated dense, 1 means concurrent dense + size_t num_iterations = 2; + size_t num_samples = 2; + size_t max_ops = 4; + size_t max_copy_fields = 32000; + size_t num_fields = 1; + size_t field_size = sizeof(ElementType); + size_t size = 4ULL * 256ULL; // 1024ULL * 4ULL;//1024ULL; +}; // namespace TestConfig + +template +inline void copy(RegionInstance src_inst, RegionInstance dst_inst, + const std::vector &fields, IndexSpace index_space) +{ + std::vector srcs(fields.size()), dsts(fields.size()); + for(size_t i = 0; i < fields.size(); i++) { + srcs[i].set_field(src_inst, fields[i], sizeof(FT)); + dsts[i].set_field(dst_inst, fields[i], sizeof(FT)); + } + index_space.copy(srcs, dsts, ProfilingRequestSet()).wait(); +} + +template +static void dump_and_verify(RegionInstance inst, RegionInstance proxy_inst, + const std::vector &fields, + const IndexSpace &is, size_t row_size, + const std::vector
&values, bool verbose = false) +{ + copy(inst, proxy_inst, fields, is); + size_t max_fail_count = 50; + size_t fail_count = 0; + int index = 0; + for(FieldID fid : fields) { + GenericAccessor acc(proxy_inst, fid); + size_t i = 0; + DT value = values[index++]; + for(IndexSpaceIterator it(is); it.valid; it.step()) { + for(PointInRectIterator it2(it.rect); it2.valid; it2.step()) { + DT v = acc[it2.p]; + if(value != DT(-1) && v != value) { + std::cout << "bad v:" << int(v) << " p:" << it2.p << " fid:" << fid + << " expv:" << value << std::endl; + fail_count++; + } else { + // std::cout << "good v:" << int(v) << " p:" << it2.p << " fid:" << fid + //<< std::endl; + } + if(verbose) { + if((i++) % row_size == 0) + std::cout << std::endl; + std::cout << it2.p << ": " << v << " "; + } + if(fail_count >= max_fail_count) + break; + } + if(verbose) + std::cout << "\n"; + } + } + assert(fail_count == 0); +} + +class Stat { +public: + Stat() + : count(0) + , mean(0.0) + , sum(0.0) + , square_sum(0.0) + , smallest(std::numeric_limits::max()) + , largest(-std::numeric_limits::max()) + {} + void reset() { *this = Stat(); } + void sample(double s) + { + count++; + if(s < smallest) + smallest = s; + if(s > largest) + largest = s; + sum += s; + double delta0 = s - mean; + mean += delta0 / count; + double delta1 = s - mean; + square_sum += delta0 * delta1; + } + unsigned get_count() const { return count; } + double get_average() const { return mean; } + double get_sum() const { return sum; } + double get_stddev() const + { + return get_variance() > 0.0 ? std::sqrt(get_variance()) : 0.0; + } + double get_variance() const { return square_sum / (count > 2 ? 1 : count - 1); } + double get_smallest() const { return smallest; } + double get_largest() const { return largest; } + + friend std::ostream &operator<<(std::ostream &os, const Stat &s); + +private: + unsigned count; + double mean; + double sum; + double square_sum; + double smallest; + double largest; +}; + +std::ostream &operator<<(std::ostream &os, const Stat &s) +{ + return os << std::scientific << std::setprecision(2) + << s.get_average() /*<< "(+/-" << s.get_stddev() << ')'*/ + << ", MIN=" << s.get_smallest() << ", MAX=" << s.get_largest() + << ", N=" << s.get_count(); +} + +struct CopyOperation; +class TestGraphFactory; + +struct RemoteCopyTaskArgs { + CopyOperation *op; + Realm::UserEvent profiling_event; + Realm::UserEvent remote_copy_event; + Realm::Event wait_on; + CopyIndexSpace index_space; + Realm::Processor profile_proc; + Realm::CopySrcDstField dsts; + Realm::CopySrcDstField srcs; +}; + +struct UpdateOpTimingTaskArgs { + CopyOperation *op; + Realm::UserEvent profiling_event; +}; + +static const char *mem_kind_to_string(Realm::Memory::Kind kind) +{ + switch(kind) { + case Realm::Memory::Kind::DISK_MEM: + return "DISK"; + case Realm::Memory::Kind::FILE_MEM: + return "FILE"; + case Realm::Memory::Kind::GLOBAL_MEM: + return "GLOBAL"; + case Realm::Memory::Kind::GPU_DYNAMIC_MEM: + return "GPU_DYN"; + case Realm::Memory::Kind::GPU_FB_MEM: + return "GPU_FB"; + case Realm::Memory::Kind::GPU_MANAGED_MEM: + return "GPU_MANAGED"; + case Realm::Memory::Kind::HDF_MEM: + return "HDF_MEM"; + case Realm::Memory::Kind::LEVEL1_CACHE: + return "L1$"; + case Realm::Memory::Kind::LEVEL2_CACHE: + return "L2$"; + case Realm::Memory::Kind::LEVEL3_CACHE: + return "L3$"; + case Realm::Memory::Kind::REGDMA_MEM: + return "REGDMA"; + case Realm::Memory::Kind::SOCKET_MEM: + return "SOCKET"; + case Realm::Memory::Kind::SYSTEM_MEM: + return "SYSTEM"; + case Realm::Memory::Kind::Z_COPY_MEM: + return "Z_COPY"; + default: + return "Unknown"; + } +} + +static void display_memory_info(Memory m) +{ + std::vector affinities; + affinities.clear(); + Realm::Machine::get_machine().get_mem_mem_affinity(affinities, m, Memory::NO_MEMORY, + false); + log_app.print() << "Memory: " << m << " kind: " << mem_kind_to_string(m.kind()) + << " size: " << static_cast(m.capacity()) / (1024.0 * 1024.0) + << "MiB"; + for(Machine::MemoryMemoryAffinity &m2m : affinities) { + log_app.print() << "\t" << m2m.m2 << " est-bw: " << std::fixed << std::setprecision(2) + << (m2m.bandwidth / 1000.0) << "GB/s est-lat: " << m2m.latency + << "ns"; + } +} + +struct CopyOperation { + CopyIndexSpace index_space; + std::vector owned_instances; + std::vector dsts; + std::vector srcs; + std::vector dependencies; + Realm::Memory src_mem; + bool is_dependency = false; + Realm::Event current_event = Realm::Event::NO_EVENT; + Stat measured_time; // nanoseconds + Stat lag_time; // nanoseconds + CopyOperation(CopyIndexSpace i, std::vector &d, + std::vector &s, Realm::Memory src) + : index_space(i) + , dsts(d) + , srcs(s) + , src_mem(src) + {} + ~CopyOperation() + { + for(Realm::RegionInstance &inst : owned_instances) { + inst.destroy(); + } + } + size_t get_total_size() const + { + size_t total = 0; + for(size_t i = 0; i < dsts.size(); i++) { + total += dsts[i].size; + } + return index_space.volume() * total; + } + void add_dependency(CopyOperation *op) + { + dependencies.push_back(op); + op->is_dependency = true; + } + + template + static size_t get_total_size(FwdIter begin, FwdIter end) + { + size_t total = 0; + for(; begin != end; ++begin) { + CopyOperation &op = *begin; + total += op.get_total_size(); + } + return total; + } +}; + +class TestGraphFactory { +public: + virtual ~TestGraphFactory() {} + virtual void create(std::vector &graph) = 0; + virtual void verify() = 0; +}; + +template +class RandomPicker { +public: + // Constructor takes a vector of memories to pick from + RandomPicker(const std::vector &_elements) + : elements(_elements) + , gen(std::random_device{}()) + , dist(0, elements.size() - 1) + {} + + std::pair operator()(size_t i) + { + size_t random_index = dist(gen); + return {random_index, elements[random_index]}; + } + +protected: + std::vector elements; + std::mt19937 gen; + std::uniform_int_distribution<> dist; +}; + +class MultiFieldTestGraphFactory : public TestGraphFactory { +public: + std::vector memories_to_test; + size_t size; + std::map fields; + std::vector> inst_to_fields; + std::vector, + std::vector>> + validate_queue; + + MultiFieldTestGraphFactory(std::vector &_mems, size_t _sz, + std::map _fields) + : memories_to_test(_mems) + , size(_sz) + , fields(_fields) + {} + + void verify() override + { + Realm::Machine::MemoryQuery mq(Realm::Machine::get_machine()); + mq.only_kind(Memory::SYSTEM_MEM).has_capacity(1); + std::vector memories(mq.begin(), mq.end()); + + for(const auto &[index_space, inst, value, fields] : validate_queue) { + RegionInstance validate_instance; + std::map src_fields; + assert(value.empty() == false); + for(FieldID fid : fields) + src_fields[fid] = sizeof(value.front()); + Realm::RegionInstance::create_instance(validate_instance, *mq.begin(), index_space, + src_fields, 0, ProfilingRequestSet()) + .wait(); + dump_and_verify(inst, validate_instance, fields, + index_space, 1, value); + } + } + + /*virtual*/ + void create(std::vector &graph) override + { + graph.clear(); + + constexpr ElementType value = 9; + + Realm::Point start_pnt(0, 0); + Realm::Point end_pnt(1, TestConfig::size); + + // Realm::Point start_pnt(0); + // Realm::Point end_pnt(TestConfig::size); + CopyIndexSpace is(Rect{start_pnt, end_pnt}); + + std::map src_fields; + for(const auto [field_id, field_size] : fields) { + src_fields[field_id] = field_size; + } + + std::map dst_fields; + for(const auto [field_id, field_size] : fields) { + dst_fields[field_id] = field_size; + } + + assert(src_fields.empty() == false); + assert(dst_fields.empty() == false); + + std::map values; + + std::vector instances; + std::vector src_instances; + std::vector dst_instances; + + for(size_t i = 0; i < memories_to_test.size(); i++) { + for(size_t j = 0; j < memories_to_test.size(); j++) { + Realm::RegionInstance src_inst, dst_inst; + Realm::RegionInstance::create_instance(src_inst, memories_to_test[i], is, + src_fields, 0, ProfilingRequestSet()) + .wait(); + Realm::RegionInstance::create_instance(dst_inst, memories_to_test[j], is, + dst_fields, 0, ProfilingRequestSet()) + .wait(); + + if(i == j) { + int index = 1; + std::vector fill_events; + for(const auto &[field_id, field_size] : src_fields) { + std::vector srcs(1), dsts(1); + ElementType fill_value = value + index++; + values[field_id] = fill_value; + srcs[0].set_fill(fill_value); // this needs conversion + assert(field_size == sizeof(ElementType)); + dsts[0].set_field(src_inst, field_id, field_size); + fill_events.emplace_back(is.copy(srcs, dsts, ProfilingRequestSet())); + } + + Event::merge_events(fill_events).wait(); + } + + instances.push_back(src_inst); + src_instances.push_back(src_inst); + inst_to_fields.push_back(src_fields); + instances.push_back(dst_inst); + dst_instances.push_back(dst_inst); + inst_to_fields.push_back(dst_fields); + } + } + + const size_t max_concurrent_ops = TestConfig::max_ops; + + RandomPicker inst_picker(instances); + + for(size_t i = 0; i < max_concurrent_ops; i++) { + auto src_inst = src_instances[0]; + auto dst_inst = + dst_instances[std::max(size_t(0), dst_instances.size() - 1)]; // inst_picker(i); + + // auto src_fields = src_fields;//inst_to_fields[0]; + // auto dst_fields = dst_fields;//inst_to_fields[dst_inst.first]; + + const size_t max_fields = + std::min(TestConfig::max_copy_fields, fields.size()); + std::vector srcs(max_fields), dsts(max_fields); + + std::vector src_field_ids; + for(const auto &[field_id, field_size] : src_fields) { + src_field_ids.push_back(field_id); + } + + std::vector dst_field_ids; + for(const auto &[field_id, field_size] : dst_fields) { + dst_field_ids.push_back(field_id); + } + + // Shuffle to select random fields for source and destination + std::shuffle(src_field_ids.begin(), src_field_ids.end(), + std::mt19937(std::random_device()())); + std::shuffle(dst_field_ids.begin(), dst_field_ids.end(), + std::mt19937(std::random_device()())); + + std::vector verify_values; + + // Pick `max_fields` random fields for the source + size_t field_index = 0; + for(size_t i = 0; i < max_fields && field_index < src_field_ids.size(); i++) { + srcs[i].set_field(src_inst, src_field_ids[field_index], + src_fields[src_field_ids[field_index]]); + verify_values.push_back(values[src_field_ids[field_index]]); + field_index++; + } + + // Pick `max_fields` random fields for the destination + field_index = 0; + for(size_t i = 0; i < max_fields && field_index < dst_field_ids.size(); i++) { + dsts[i].set_field(dst_inst, dst_field_ids[field_index], + dst_fields[dst_field_ids[field_index]]); + field_index++; + } + + auto sub = + std::vector(dst_field_ids.begin(), dst_field_ids.begin() + max_fields); + validate_queue.emplace_back(std::make_tuple(is, dst_inst, verify_values, sub)); + + graph.emplace_back(is, dsts, srcs, src_inst.get_location()); + } + + graph[0].owned_instances = instances; + } +}; + +static void display_node_data(std::vector &graph) +{ + if(TestConfig::graphviz) { + std::cout << "digraph g {" << std::endl; + } + // Node information + for(size_t i = 0; i < graph.size(); i++) { + // Assume instances are the same across all src and all dst operations + Memory src = graph[i].srcs[0].inst.get_location(); + Memory dst = graph[i].dsts[0].inst.get_location(); + if(!src.exists()) { + src = dst; + } + std::vector affinity; + if(Realm::Machine::get_machine().get_mem_mem_affinity(affinity, src, dst, false) == + 0) { + Realm::Machine::MemoryMemoryAffinity fake_aff; + fake_aff.m1 = src; + fake_aff.m2 = dst; + fake_aff.bandwidth = 1; + fake_aff.latency = UINT_MAX; + affinity.push_back(fake_aff); + } + const double bw = + (graph[i].get_total_size() * 1000ULL) / graph[i].measured_time.get_average(); + const double lag = graph[i].lag_time.get_average(); + if(TestConfig::graphviz) { + std::cout << "node_" << i << "[label=<" << i << "
"; + } + std::cout << src << '(' << mem_kind_to_string(src.kind()) << ") : " << dst << '(' + << mem_kind_to_string(dst.kind()) << ")"; + if(TestConfig::graphviz) { + std::cout << "
"; + } + std::cout << " sz: " << graph[i].get_total_size() / (1024ULL * 1024ULL) << "MiB" + << " bw: " << std::fixed << std::setprecision(2) << bw / 1000.0 << "GB/s (" + << std::fixed << 100.0 * bw / affinity[0].bandwidth << "%)" + << " lag: " << lag << "ns"; + if(TestConfig::graphviz) { + std::cout << ">];"; + } + std::cout << std::endl; + } + // Links + if(TestConfig::graphviz) { + for(size_t i = 0; i < graph.size(); i++) { + for(size_t j = 0; j < graph[i].dependencies.size(); j++) { + const size_t dependency_idx = graph[i].dependencies[j] - graph.data(); + std::cout << "node_" << dependency_idx << " -> node_" << i << ';' << std::endl; + } + } + std::cout << '}' << std::endl; + } +} + +static void remote_copy_task(const void *args, size_t arglen, const void *userdata, + size_t userlen, Processor p) +{ + const RemoteCopyTaskArgs &self_args = + *reinterpret_cast(args); + assert(arglen == sizeof(RemoteCopyTaskArgs)); + + Realm::ProfilingRequestSet prs; + if(TestConfig::enable_profiling) { + UpdateOpTimingTaskArgs prof_args; + prof_args.op = self_args.op; + prof_args.profiling_event = self_args.profiling_event; + prs.add_request(self_args.profile_proc, UPDATE_OP_TIMING_TASK, &prof_args, + sizeof(prof_args)) + .add_measurement(ProfilingMeasurements::OperationTimeline::ID); + } + std::vector srcs(1, self_args.srcs), dsts(1, self_args.dsts); + Realm::Event event = self_args.index_space.copy(srcs, dsts, prs, self_args.wait_on); + self_args.remote_copy_event.trigger(event); + + Realm::RegionInstance src_inst = self_args.srcs.inst; + Realm::RegionInstance dst_inst = self_args.dsts.inst; + log_app.debug("Remote Copy(%p) from src:%llx(%llx) to dst:%llx(%llx) is issued on " + "processor %llx", + self_args.op, src_inst.id, src_inst.get_location().id, dst_inst.id, + dst_inst.get_location().id, p.id); +} + +static void issue_copy_from_remote(std::vector &finish_events, + CopyOperation &op, Realm::Event wait_on, + Processor local_p, Processor remote_p) +{ + UserEvent profiling_event = UserEvent::NO_USER_EVENT; + UserEvent remote_copy_event = UserEvent::create_user_event(); + if(TestConfig::enable_profiling) { + profiling_event = UserEvent::create_user_event(); + finish_events.push_back(profiling_event); + } + + RemoteCopyTaskArgs remote_args; + remote_args.profiling_event = profiling_event; + remote_args.remote_copy_event = remote_copy_event; + remote_args.wait_on = wait_on; + remote_args.op = + &op; // profiling task is executed on local proc, so it is OK to pass a pointer + remote_args.profile_proc = local_p; + remote_args.index_space = op.index_space; + remote_args.srcs = op.srcs[0]; + remote_args.dsts = op.dsts[0]; + + Realm::Event remote_task_event = + remote_p.spawn(REMOTE_COPY_TASK, &remote_args, sizeof(RemoteCopyTaskArgs)); + finish_events.push_back(remote_task_event); + + // TODO: Add some validation of the copy here + // This is a dangling node (a node without children), so make sure to + // capture it's finish event as one that marks the graph as complete + if(!op.is_dependency) { + finish_events.push_back(remote_copy_event); + } +} + +static void issue_copy_from_local(std::vector &finish_events, + CopyOperation &op, Realm::Event wait_on, Processor p) +{ + // Queue up the copy! + // TODO: Use the profiling request set to accumulate the times of the + // individual copies for bandwidth verification + Realm::ProfilingRequestSet prs; + if(TestConfig::enable_profiling) { + UpdateOpTimingTaskArgs prof_args; + UserEvent profiling_event = UserEvent::create_user_event(); + finish_events.push_back(profiling_event); + prof_args.op = &op; + prof_args.profiling_event = profiling_event; + prs.add_request(p, UPDATE_OP_TIMING_TASK, &prof_args, sizeof(prof_args)) + .add_measurement(ProfilingMeasurements::OperationTimeline::ID); + } + op.current_event = op.index_space.copy(op.srcs, op.dsts, prs, wait_on); + + Realm::RegionInstance src_inst = op.srcs[0].inst; + Realm::RegionInstance dst_inst = op.dsts[0].inst; + log_app.debug("Local Copy(%p) from src:%llx(%llx) to dst:%llx(%llx) is issued on " + "processor %llx", + &op, src_inst.id, src_inst.get_location().id, dst_inst.id, + dst_inst.get_location().id, p.id); + + // TODO: Add some validation of the copy here + // This is a dangling node (a node without children), so make sure to + // capture it's finish event as one that marks the graph as complete + if(!op.is_dependency) { + finish_events.push_back(op.current_event); + } +} + +static Realm::Event run_graph(std::vector &graph, Realm::Event start_event, + Realm::Processor p) +{ + // build a map of processor + std::map proc_map; + if(TestConfig::enable_remote_copy) { + for(realm_address_space_t i = 0; + i < Realm::Machine::get_machine().get_address_space_count(); i++) { + proc_map[i] = Realm::Processor::NO_PROC; + } + for(Machine::ProcessorQuery::iterator it = + Realm::Machine::ProcessorQuery(Realm::Machine::get_machine()) + .only_kind(Realm::Processor::LOC_PROC) + .begin(); + it; ++it) { + Processor proc = *it; + if(proc_map[proc.address_space()] == Realm::Processor::NO_PROC) { + proc_map[proc.address_space()] = proc; + } + } + } + std::vector finish_events; + // TODO: Add fill operations for validation purposes + // This implementation assumes a topologically sorted graph from the graph + // generator + for(CopyOperation &op : graph) { + Realm::Event wait_on = start_event; + if(op.dependencies.size() > 0) { + std::vector events(op.dependencies.size() + 1, + Realm::Event::NO_EVENT); + for(size_t i = 0; i < op.dependencies.size(); i++) { + events[i] = op.dependencies[i]->current_event; + } + events.back() = start_event; + wait_on = Event::merge_events(events); + } + + realm_address_space_t src_rank = op.src_mem.address_space(); + if(!TestConfig::enable_remote_copy || src_rank == p.address_space()) { + issue_copy_from_local(finish_events, op, wait_on, p); + } else { + issue_copy_from_remote(finish_events, op, wait_on, p, proc_map[src_rank]); + } + } + // And the final event signaling this graph is complete + return Realm::Event::merge_events(finish_events); +} + +static void update_operation_time(const void *args, size_t arglen, const void *userdata, + size_t userlen, Realm::Processor p) +{ + ProfilingResponse resp(args, arglen); + assert(resp.user_data_size() == sizeof(UpdateOpTimingTaskArgs)); + const UpdateOpTimingTaskArgs &self_args = + *static_cast(resp.user_data()); + ProfilingMeasurements::OperationTimeline timeline; + + if(resp.get_measurement(timeline)) { + self_args.op->measured_time.sample(timeline.complete_time - timeline.start_time); + self_args.op->lag_time.sample(timeline.start_time - timeline.ready_time); + } else { + assert(0 && "Failed to get timeline measurement"); + } + self_args.profiling_event.trigger(); + log_app.debug("Profile Copy:%p is done on processor %llx", self_args.op, p.id); +} + +static void bench_timing_task(const void *args, size_t arglen, const void *userdata, + size_t userlen, Processor p) +{ + log_app.print("=== Memory Info ==="); + Realm::Machine::MemoryQuery mq(Realm::Machine::get_machine()); + mq.only_kind(Memory::GPU_FB_MEM).has_capacity(1); + std::vector memories(mq.begin(), mq.end()); + for(Memory m : memories) { + display_memory_info(m); + } + log_app.print("==================="); + + // mq = mq.has_capacity(pow(TestConfig::size, MAX_DIM) * sizeof(ElementType)); + memories.assign(mq.begin(), mq.end()); + if(memories.size() == 0) { + abort(); + } + + TestGraphFactory *test_factory = nullptr; + std::map fields; + for(size_t i = 0; i < TestConfig::num_fields; i++) { + fields[i] = TestConfig::field_size; + } + test_factory = new MultiFieldTestGraphFactory(memories, TestConfig::size, fields); + + std::vector graph; + test_factory->create(graph); + + Stat graph_time; + + for(size_t sample_iter = 0; sample_iter < TestConfig::num_samples; sample_iter++) { + Realm::UserEvent trigger_event = Realm::UserEvent::create_user_event(); + Realm::Event current_graph_event = trigger_event; + for(size_t graph_iter = 0; graph_iter < TestConfig::num_iterations; graph_iter++) { + current_graph_event = run_graph(graph, current_graph_event, p); + } + + size_t start_time = Clock::current_time_in_microseconds(); + trigger_event.trigger(); // Start the graphs + current_graph_event.wait(); // Wait for them to finish + size_t end_time = Clock::current_time_in_microseconds(); + + if(TestConfig::verify) { + test_factory->verify(); + } + + if(sample_iter != 0) { + graph_time.sample(double(end_time - start_time) / TestConfig::num_iterations); + } + log_app.info() << "\tGraph sample (us): " << end_time - start_time; + } + + size_t total_size_bytes = CopyOperation::get_total_size(graph.begin(), graph.end()); + + log_app.print() << "Graph total transfer size: " + << total_size_bytes / (1024ULL * 1024ULL) << "MiB"; + log_app.print() << "Graph time (us): " << graph_time.get_average(); + log_app.print() << "Graph bandwidth (GB/s): " + << total_size_bytes / (1000.0 * graph_time.get_average()); + if(TestConfig::enable_profiling) { + display_node_data(graph); + } + graph.clear(); // Should destroy all the created instances + + delete test_factory; + + usleep(100000); +} + +int main(int argc, char **argv) +{ + Runtime r; + + bool ok = r.init(&argc, &argv); + assert(ok); + + CommandLineParser cp; + cp.add_option_int("-profile", TestConfig::enable_profiling) + .add_option_int("-remote-copy", TestConfig::enable_remote_copy) + .add_option_int("-iter", TestConfig::num_iterations) + .add_option_int("-samples", TestConfig::num_samples) + .add_option_int("-size", TestConfig::size) + .add_option_int("-copy_fields", TestConfig::max_copy_fields) + .add_option_int("-num_fields", TestConfig::num_fields) + .add_option_int("-field_size", TestConfig::field_size) + .add_option_int("-graphviz", TestConfig::graphviz) + .add_option_int("-max_ops", TestConfig::max_ops) + .add_option_int("-verify", TestConfig::verify); + ok = cp.parse_command_line(argc, (const char **)argv); + assert(ok); + + r.register_task(BENCH_TIMING_TASK, bench_timing_task); + r.register_task(UPDATE_OP_TIMING_TASK, update_operation_time); + r.register_task(REMOTE_COPY_TASK, remote_copy_task); + + Processor p = Machine::ProcessorQuery(Machine::get_machine()) + .only_kind(Processor::LOC_PROC) + .first(); + + // collective launch of a single task - everybody gets the same finish event + Event e = r.collective_spawn(p, BENCH_TIMING_TASK, 0, 0); + + // request shutdown once that task is complete + r.shutdown(e); + + // now sleep this thread until that shutdown actually happens + r.wait_for_shutdown(); + + return 0; +} diff --git a/tests/unit_tests/idindexed_fields_iterator_test.cc b/tests/unit_tests/idindexed_fields_iterator_test.cc new file mode 100644 index 00000000000..c0071714220 --- /dev/null +++ b/tests/unit_tests/idindexed_fields_iterator_test.cc @@ -0,0 +1,198 @@ +/* + * Copyright 2025 Stanford University, NVIDIA Corporation + * SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "realm/transfer/transfer.h" +#include "realm/inst_layout.h" +#include "test_common.h" +#include +#include + +using namespace Realm; + +// Base class for parameterized test cases +struct BaseIDIndexedIteratorTestCaseData { + virtual ~BaseIDIndexedIteratorTestCaseData() = default; + virtual int get_dim() const = 0; +}; + +template +struct IDIndexedIteratorTestCaseData { + // The iteration space + Rect domain; + // The expected subrects from get_addresses (for uniform fields it's just the full + // domain) + std::vector> expected; + // Dimension traversal order + std::vector dim_order; + // The list of field IDs (all fields have the same uniform size) + std::vector fields; + // Uniform field size in bytes + size_t field_size; +}; + +template +struct WrappedIDIndexedIteratorTestCaseData : public BaseIDIndexedIteratorTestCaseData { + IDIndexedIteratorTestCaseData data; + explicit WrappedIDIndexedIteratorTestCaseData(IDIndexedIteratorTestCaseData d) + : data(std::move(d)) + {} + int get_dim() const override { return N; } +}; + +class IDIndexedFieldsIteratorGetAddressesTest + : public ::testing::TestWithParam { +protected: + void TearDown() override { delete GetParam(); } +}; + +struct MockHeap : public ReplicatedHeap { + void *alloc_obj(std::size_t bytes, std::size_t align = 16) override + { + void *ptr = nullptr; +#ifdef REALM_ON_WINDOWS + ptr = _aligned_malloc(bytes, align); +#else + int ret = posix_memalign(&ptr, align, bytes); + if(ret != 0) + ptr = nullptr; +#endif + assert(ptr != nullptr); + return ptr; + } + + void free_obj(void *ptr) override + { +#ifdef REALM_ON_WINDOWS + _aligned_free(ptr); +#else + free(ptr); +#endif + } +}; + +template +void run_uniform_test_case(const IDIndexedIteratorTestCaseData &tc) +{ + using T = int; + + // Create an instance with one entry per field, all uniformly sized + std::vector field_sizes(tc.fields.size(), tc.field_size); + RegionInstanceImpl *inst_impl = + create_inst(nullptr, tc.domain, tc.fields, field_sizes); + + MockHeap mock_heap; + + // Build the IDIndexedFieldsIterator + auto it = std::make_unique>( + tc.dim_order.data(), tc.fields, tc.field_size, inst_impl, tc.domain, &mock_heap); + + const InstanceLayoutPieceBase *nonaffine = nullptr; + AddressList addrlist; + AddressListCursor cursor; + + // Invoke get_addresses + bool ok = it->get_addresses(addrlist, nonaffine); + ASSERT_TRUE(ok) << "get_addresses() failed"; + ASSERT_TRUE(it->done()) << "Iterator should be done after get_addresses"; + + // Check total bytes pending + size_t total_volume = 0; + for(auto &r : tc.expected) + total_volume += r.volume(); + size_t expected_bytes = total_volume * tc.field_size * tc.fields.size(); + ASSERT_EQ(addrlist.bytes_pending(), expected_bytes) << "bytes_pending mismatch"; + + // If 1D, walk through with the cursor and drain the bytes + cursor.set_addrlist(&addrlist); + if(expected_bytes > 0 && cursor.get_dim() == 1) { + int dim = cursor.get_dim() - 1; + for(size_t f = 0; f < tc.fields.size(); f++) { + for(auto &r : tc.expected) { + size_t rem = cursor.remaining(dim); + ASSERT_EQ(rem, r.volume() * tc.field_size) + << "Unexpected remaining bytes for field " << f; + cursor.advance(dim, rem); + } + } + ASSERT_EQ(addrlist.bytes_pending(), 0u) << "All bytes should have been consumed"; + } +} + +TEST_P(IDIndexedFieldsIteratorGetAddressesTest, Base) +{ + BaseIDIndexedIteratorTestCaseData const *base = GetParam(); + dispatch_for_dimension( + base->get_dim(), + [&](auto Dim) { + constexpr int N = Dim; + auto const &tc = + static_cast const *>(base)->data; + run_uniform_test_case(tc); + }, + std::make_index_sequence{}); +} + +INSTANTIATE_TEST_SUITE_P(UniformFieldsCases, IDIndexedFieldsIteratorGetAddressesTest, + ::testing::Values( + // 1D: empty domain + new WrappedIDIndexedIteratorTestCaseData<1>({ + /* domain */ Rect<1, int>::make_empty(), + /* expected */ {}, + /* dim_order */ {0}, + /* fields */ {0}, + /* field_size */ sizeof(int), + }), + // 1D: single field, full domain + new WrappedIDIndexedIteratorTestCaseData<1>({ + /* domain */ Rect<1, int>(0, 14), + /* expected */ {Rect<1, int>(0, 14)}, + /* dim_order */ {0}, + /* fields */ {0}, + /* field_size */ sizeof(int), + }), + // 1D: two uniform fields, full domain + new WrappedIDIndexedIteratorTestCaseData<1>({ + /* domain */ Rect<1, int>(0, 14), + /* expected */ {Rect<1, int>(0, 14)}, + /* dim_order */ {0}, + /* fields */ {0, 1}, + /* field_size */ sizeof(int), + }), + // 2D: single field, full rectangular domain + new WrappedIDIndexedIteratorTestCaseData<2>({ + /* domain */ Rect<2, int>({0, 0}, {10, 10}), + /* expected */ {Rect<2, int>({0, 0}, {10, 10})}, + /* dim_order */ {0, 1}, + /* fields */ {0}, + /* field_size */ sizeof(int), + }), + // 2D: single field, reverse dims + new WrappedIDIndexedIteratorTestCaseData<2>({ + /* domain */ Rect<2, int>({0, 0}, {10, 10}), + /* expected */ {Rect<2, int>({0, 0}, {10, 10})}, + /* dim_order */ {1, 0}, + /* fields */ {0}, + /* field_size */ sizeof(int), + }), + // 3D: single field, small cube + new WrappedIDIndexedIteratorTestCaseData<3>({ + /* domain */ Rect<3, int>({0, 0, 0}, {1, 1, 1}), + /* expected */ {Rect<3, int>({0, 0, 0}, {1, 1, 1})}, + /* dim_order */ {0, 1, 2}, + /* fields */ {0}, + /* field_size */ sizeof(int), + }))); diff --git a/tests/unit_tests/read_address_entry_test.cc b/tests/unit_tests/read_address_entry_test.cc new file mode 100644 index 00000000000..c2cc0d41ebb --- /dev/null +++ b/tests/unit_tests/read_address_entry_test.cc @@ -0,0 +1,699 @@ +/* + * Copyright 2025 Stanford University, NVIDIA Corporation + * SPDX-License-Identifier: Apache-2.0 + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include "realm/transfer/address_list.h" +#include "realm/cuda/cuda_memcpy.h" +#include "realm/cuda/cuda_internal.h" +#include +#include +#include +#include +#include +#include +#include +#include + +using namespace Realm; +using namespace Realm::Cuda; + +namespace { + + // === Utility helpers for the tests + // ============================================================== + + struct MockHeap { + void *alloc_obj(std::size_t bytes, std::size_t align = 16) + { + void *ptr = nullptr; +#ifdef REALM_ON_WINDOWS + ptr = _aligned_malloc(bytes, align); +#else + int ret = posix_memalign(&ptr, align, bytes); + if(ret != 0) + ptr = nullptr; +#endif + assert(ptr != nullptr); + return ptr; + } + + void free_obj(void *ptr) + { +#ifdef REALM_ON_WINDOWS + _aligned_free(ptr); +#else + free(ptr); +#endif + } + }; + + static void append_entry_1d(AddressList &al, size_t contig_bytes, size_t base = 0, + bool wrap_mode = true) + { + bool ret = + al.append_entry(/*dims=*/1, contig_bytes, contig_bytes, base, {}, wrap_mode); + assert(ret); + } + + static void append_entry_2d(AddressList &al, size_t contig_bytes, size_t lines, + size_t base = 0, bool wrap_mode = true) + { + std::unordered_map> count_strides; + // dim 1 : {count, stride} + count_strides[1] = {lines, contig_bytes}; + bool ret = al.append_entry(/*dims=*/2, contig_bytes, contig_bytes * lines, base, + count_strides, wrap_mode); + assert(ret); + } + + static void append_entry_3d(AddressList &al, size_t contig_bytes, size_t lines, + size_t planes, size_t base = 0, bool wrap_mode = true) + { + std::unordered_map> count_strides; + count_strides[1] = {lines, contig_bytes}; + count_strides[2] = {planes, contig_bytes * lines}; + bool ret = al.append_entry(/*dims=*/3, contig_bytes, contig_bytes * lines * planes, + base, count_strides, wrap_mode); + assert(ret); + } + +} // anonymous namespace + +// ================================================================================================ +// TESTS +// ================================================================================================ + +TEST(ReadAddressEntryTests, FastPath1D_NoFields) +{ + constexpr size_t CONTIG = 64; + AddressList in_al, out_al; + append_entry_1d(in_al, CONTIG); + append_entry_1d(out_al, CONTIG); + + AddressListCursor in_cur, out_cur; + in_cur.set_addrlist(&in_al); + out_cur.set_addrlist(&out_al); + + AffineCopyInfo<3> copy_infos{}; + MemcpyTransposeInfo transpose_info{}; + size_t min_align = 16; + size_t fields_tot = 0; + const size_t bytes_left = CONTIG; + const size_t max_fields = 8; + + size_t ret_bytes = GPUXferDes::read_address_entry( + copy_infos, min_align, transpose_info, in_cur, + /*in_base=*/0, out_cur, /*out_base=*/0, bytes_left, max_fields, fields_tot); + + ASSERT_EQ(ret_bytes, CONTIG); + ASSERT_EQ(copy_infos.num_rects, 1u); + const auto &ci = copy_infos.subrects[0]; + EXPECT_EQ(ci.extents[0], CONTIG); + EXPECT_EQ(ci.volume, CONTIG); + EXPECT_EQ(ci.src.strides[0], CONTIG); + EXPECT_EQ(ci.dst.strides[0], CONTIG); + EXPECT_EQ(fields_tot, 1u); + EXPECT_EQ(in_al.bytes_pending(), 0u); + EXPECT_EQ(out_al.bytes_pending(), 0u); +} + +TEST(ReadAddressEntryTests, TwoDimensionalSplit_FirstDim) +{ + constexpr size_t CONTIG = 32; + constexpr size_t LINES = 4; + const size_t kTotalBytes = CONTIG * LINES; + + AddressList in_al, out_al; + append_entry_2d(in_al, CONTIG, LINES); + append_entry_2d(out_al, CONTIG, LINES); + + AddressListCursor in_cur, out_cur; + in_cur.set_addrlist(&in_al); + out_cur.set_addrlist(&out_al); + + AffineCopyInfo<3> copy_infos{}; + MemcpyTransposeInfo transpose_info{}; + size_t min_align = 16; + size_t fields_tot = 0; + const size_t max_fields = 8; + + size_t ret_bytes = + GPUXferDes::read_address_entry(copy_infos, min_align, transpose_info, in_cur, 0, + out_cur, 0, kTotalBytes, max_fields, fields_tot); + + ASSERT_EQ(ret_bytes, kTotalBytes); + ASSERT_EQ(copy_infos.num_rects, 1u); + const auto &ci = copy_infos.subrects[0]; + EXPECT_EQ(ci.extents[0], CONTIG); + EXPECT_EQ(ci.extents[1], LINES); + EXPECT_EQ(ci.extents[2], 1u); + EXPECT_EQ(ci.volume, kTotalBytes); + EXPECT_EQ(ci.src.strides[0], CONTIG); + EXPECT_EQ(ci.src.strides[1], LINES); + EXPECT_EQ(ci.dst.strides[0], CONTIG); + EXPECT_EQ(ci.dst.strides[1], LINES); + EXPECT_EQ(fields_tot, 1); // TODO: fix +} + +TEST(ReadAddressEntryTests, ThreeDimensional_NoTranspose) +{ + constexpr size_t CONTIG = 16; + constexpr size_t LINES = 4; + constexpr size_t PLANES = 2; + const size_t kTotalBytes = CONTIG * LINES * PLANES; + + AddressList in_al, out_al; + append_entry_3d(in_al, CONTIG, LINES, PLANES); + append_entry_3d(out_al, CONTIG, LINES, PLANES); + + AddressListCursor in_cur, out_cur; + in_cur.set_addrlist(&in_al); + out_cur.set_addrlist(&out_al); + + AffineCopyInfo<3> copy_infos{}; + MemcpyTransposeInfo transpose_info{}; + size_t min_align = 16; + size_t fields_tot = 0; + + const size_t max_fields = 8; + size_t ret_bytes = + GPUXferDes::read_address_entry(copy_infos, min_align, transpose_info, in_cur, 0, + out_cur, 0, kTotalBytes, max_fields, fields_tot); + + ASSERT_EQ(ret_bytes, kTotalBytes); + ASSERT_EQ(copy_infos.num_rects, 1u); + const auto &ci = copy_infos.subrects[0]; + EXPECT_EQ(ci.extents[0], CONTIG); + EXPECT_EQ(ci.extents[1], LINES); + EXPECT_EQ(ci.extents[2], PLANES); + EXPECT_EQ(ci.volume, kTotalBytes); + EXPECT_EQ(ci.src.strides[0], CONTIG); + EXPECT_EQ(ci.src.strides[1], LINES); + EXPECT_EQ(ci.dst.strides[0], CONTIG); + EXPECT_EQ(ci.dst.strides[1], LINES); + EXPECT_EQ(fields_tot, 1); +} + +TEST(ReadAddressEntryTests, FieldBlock_LimitedTransfer) +{ + constexpr size_t CONTIG = 32; + constexpr int kTotalFields = 5; + constexpr int kMaxFieldsPerXfer = 2; + + AddressList in_al, out_al; + // Attach field blocks BEFORE appending the entry so that the byte accounting + // inside AddressList includes all fields. + std::vector field_ids(kTotalFields); + std::iota(field_ids.begin(), field_ids.end(), 0); + + MockHeap heap; + auto *fb_in = FieldBlock::create(heap, field_ids.data(), field_ids.size()); + auto *fb_out = FieldBlock::create(heap, field_ids.data(), field_ids.size()); + + in_al.attach_field_block(fb_in); + out_al.attach_field_block(fb_out); + + append_entry_1d(in_al, CONTIG); + append_entry_1d(out_al, CONTIG); + + AddressListCursor in_cur, out_cur; + in_cur.set_addrlist(&in_al); + out_cur.set_addrlist(&out_al); + + AffineCopyInfo<3> copy_infos{}; + MemcpyTransposeInfo transpose_info{}; + size_t min_align = 16; + size_t fields_tot = 0; + + size_t ret_bytes = GPUXferDes::read_address_entry( + copy_infos, min_align, transpose_info, in_cur, 0, out_cur, 0, CONTIG * kTotalFields, + kMaxFieldsPerXfer, fields_tot); + + ASSERT_EQ(copy_infos.num_rects, 1u); + const auto &ci = copy_infos.subrects[0]; + + EXPECT_EQ(fields_tot, static_cast(kMaxFieldsPerXfer)); + EXPECT_EQ(ci.src.num_fields, static_cast(kMaxFieldsPerXfer)); + EXPECT_EQ(ci.dst.num_fields, static_cast(kMaxFieldsPerXfer)); + EXPECT_EQ(ret_bytes, CONTIG * kMaxFieldsPerXfer); + EXPECT_EQ(in_al.bytes_pending(), CONTIG * (kTotalFields - kMaxFieldsPerXfer)); + EXPECT_EQ(out_al.bytes_pending(), CONTIG * (kTotalFields - kMaxFieldsPerXfer)); + + heap.free_obj(fb_in); + heap.free_obj(fb_out); +} + +// ───────────────────────────────────────────────────────────────────────────── +// (A) Partial-field loop: move at most Fmax fields per iteration until done +// ───────────────────────────────────────────────────────────────────────────── +TEST(ReadAddressEntry, Loop_PartialFieldConsumption) +{ + constexpr size_t C = 64; // bytes per rectangle + constexpr int Ftot = 7; // total number of fields + constexpr int Fmax = 3; // copy at most this many per iteration + + // Build identical src/dst AddressLists and attach FieldBlocks + AddressList in_al, out_al; + + std::vector ids(Ftot); + std::iota(ids.begin(), ids.end(), 0); + MockHeap h; + auto *fb_in = FieldBlock::create(h, ids.data(), Ftot); + auto *fb_out = FieldBlock::create(h, ids.data(), Ftot); + in_al.attach_field_block(fb_in); + out_al.attach_field_block(fb_out); + + append_entry_1d(in_al, C); + append_entry_1d(out_al, C); + + AddressListCursor ic, oc; + ic.set_addrlist(&in_al); + oc.set_addrlist(&out_al); + + size_t bytes_moved_total = 0; + int fields_remaining = Ftot; + + while(in_al.bytes_pending() > 0) { + AffineCopyInfo<3> info{}; + MemcpyTransposeInfo tr{}; + size_t min_align = 16, fields_total = 0; + + // Allow up to `C * Fmax` bytes this round — enough for at most Fmax fields + size_t moved = GPUXferDes::read_address_entry(info, min_align, tr, ic, 0, oc, 0, + C * Fmax, Fmax, fields_total); + + ASSERT_GT(moved, 0u); + EXPECT_EQ(fields_total, static_cast(std::min(Fmax, fields_remaining))); + + bytes_moved_total += moved; + fields_remaining -= static_cast(fields_total); + } + + EXPECT_EQ(bytes_moved_total, C * Ftot); + EXPECT_EQ(fields_remaining, 0); + EXPECT_EQ(in_al.bytes_pending(), 0u); + EXPECT_EQ(out_al.bytes_pending(), 0u); + + h.free_obj(fb_in); + h.free_obj(fb_out); +} + +// ───────────────────────────────────────────────────────────────────────────── +// (B) Partial-rect loop: move *bytes_left_step* at a time within ONE field +// ───────────────────────────────────────────────────────────────────────────── +TEST(ReadAddressEntry, Loop_PartialRectConsumption) +{ + constexpr size_t C = 64; // full rectangle size + constexpr size_t STEP = 16; // copy only 16 bytes per iteration + + AddressList in_al, out_al; + append_entry_1d(in_al, C); + append_entry_1d(out_al, C); + + AddressListCursor ic, oc; + ic.set_addrlist(&in_al); + oc.set_addrlist(&out_al); + + size_t bytes_moved_total = 0; + + while(in_al.bytes_pending() > 0) { + AffineCopyInfo<3> info{}; + MemcpyTransposeInfo tr{}; + size_t min_align = 16, fields_total = 0; + + size_t moved = GPUXferDes::read_address_entry(info, min_align, tr, ic, 0, oc, 0, + STEP, // bytes_left -> only STEP bytes + 8, // max_xfer_fields + fields_total); + + ASSERT_GT(moved, 0u); + EXPECT_LE(moved, STEP); + EXPECT_EQ(fields_total, 1u); // only one field in play + + bytes_moved_total += moved; + } + + EXPECT_EQ(bytes_moved_total, C); // whole rectangle eventually moved + EXPECT_EQ(in_al.bytes_pending(), 0u); + EXPECT_EQ(out_al.bytes_pending(), 0u); +} + +// ───────────────────────────────────────────────────────────────────────────── +// (C) 3-D partial–rect loop : copy one plane at a time +// ───────────────────────────────────────────────────────────────────────────── +TEST(ReadAddressEntry, Loop_PartialRectConsumption3D) +{ + constexpr size_t CONTIG = 16; // bytes per contiguous chunk (X) + constexpr size_t LINES = 4; // Y + constexpr size_t PLANES = 5; // Z → total volume = 16 * 4 * 5 = 320 + constexpr size_t STEP = CONTIG * LINES; // 64 bytes ⇒ one full plane + + AddressList in_al, out_al; + append_entry_3d(in_al, CONTIG, LINES, PLANES); + append_entry_3d(out_al, CONTIG, LINES, PLANES); + + AddressListCursor ic, oc; + ic.set_addrlist(&in_al); + oc.set_addrlist(&out_al); + + size_t total_moved = 0; + + while(in_al.bytes_pending() > 0) { + AffineCopyInfo<3> info{}; + MemcpyTransposeInfo tr{}; + size_t min_align = 16, fields_total = 0; + + size_t moved = GPUXferDes::read_address_entry(info, min_align, tr, ic, 0, oc, 0, + STEP, // allow only one plane’s bytes + 8, // unlimited fields (1 here) + fields_total); + + ASSERT_GT(moved, 0u); + EXPECT_LE(moved, STEP); // never more than a plane + EXPECT_EQ(fields_total, 1); // TODO: fix + total_moved += moved; + } + + EXPECT_EQ(total_moved, CONTIG * LINES * PLANES); // all bytes copied + EXPECT_EQ(in_al.bytes_pending(), 0u); + EXPECT_EQ(out_al.bytes_pending(), 0u); +} + +// ───────────────────────────────────────────────────────────────────────────── +// 2-D entry: source row 64 B, destination row 32 B +// bytes_left = 64 → contig_bytes becomes 32 (< icount 64) so id = od = 0 path +// ───────────────────────────────────────────────────────────────────────────── +TEST(ReadAddressEntry, TwoDimensional_SplitDim0_OneShot) +{ + constexpr size_t SRC_CONTIG = 64; // src row width (icount starts at 64) + constexpr size_t DST_CONTIG = 32; // dst row width (ocount 32) + constexpr size_t LINES = 2; // second dimension + constexpr size_t BYTES_LEFT = 64; // > dst row, < src row ⇒ split path + + // Build AddressLists + AddressList in_al, out_al; + append_entry_2d(in_al, SRC_CONTIG, LINES); // 64-byte rows + append_entry_2d(out_al, DST_CONTIG, LINES); // 32-byte rows + + AddressListCursor ic, oc; + ic.set_addrlist(&in_al); + oc.set_addrlist(&out_al); + + AffineCopyInfo<3> info{}; + MemcpyTransposeInfo tr{}; + size_t min_align = 16, fields_total = 0; + + // Call read_address_entry – should copy *one* 32-byte line + size_t moved = GPUXferDes::read_address_entry(info, min_align, tr, ic, 0, oc, 0, + BYTES_LEFT, // flow-control limit + 8 /*max fields*/, // no FieldBlocks + fields_total); + + // ─── Checks ─────────────────────────────────────────────────────────────── + EXPECT_EQ(moved, SRC_CONTIG); // 64 bytes copied + EXPECT_EQ(info.num_rects, 1u); + + const auto &rect = info.subrects[0]; + EXPECT_EQ(rect.extents[0], DST_CONTIG); // contig reduced to 32 + EXPECT_EQ(rect.extents[1], LINES); + + EXPECT_EQ(rect.src.strides[0], DST_CONTIG); + EXPECT_EQ(rect.src.strides[1], LINES); + EXPECT_EQ(rect.dst.strides[0], DST_CONTIG); + EXPECT_EQ(rect.dst.strides[1], LINES); + + EXPECT_EQ(rect.volume, SRC_CONTIG); + EXPECT_EQ(fields_total, 1u); // TODO: fix + + // Remaining bytes: (64-32) * 2 lines = 64 bytes still pending + EXPECT_EQ(in_al.bytes_pending(), (SRC_CONTIG - DST_CONTIG) * LINES); + EXPECT_EQ(out_al.bytes_pending(), 0); +} + +TEST(ReadAddressEntryTests, SrcWithFields_DstNoFields_Partial) +{ + constexpr size_t kContig = 64; + constexpr int kSrcFields = 5; + constexpr int kMoveFields = 2; + constexpr int kLeft = 3; + constexpr size_t kBytesLeft = kContig * 5; // allow exactly one field worth of bytes + + // Build address lists + AddressList in_al, out_al; + + // Attach field block to **source** only + std::vector src_field_ids(kSrcFields); + std::iota(src_field_ids.begin(), src_field_ids.end(), 0); + MockHeap heap; + auto *fb_src = FieldBlock::create(heap, src_field_ids.data(), src_field_ids.size()); + in_al.attach_field_block(fb_src); + + append_entry_1d(in_al, kContig); + append_entry_1d(out_al, kContig * kMoveFields + kLeft, 0); + + size_t bytes_left = + std::min(kBytesLeft, std::min(in_al.bytes_pending(), out_al.bytes_pending())); + + AddressListCursor ic, oc; + ic.set_addrlist(&in_al); + oc.set_addrlist(&out_al); + + AffineCopyInfo<3> infos{}; + MemcpyTransposeInfo tr{}; + size_t min_align = 16, fields_total = 0; + + const size_t max_fields = 8; // not the limiting factor here + + EXPECT_EQ(ic.get_offset(), 0); + EXPECT_EQ(oc.get_offset(), 0); + + { + size_t moved = GPUXferDes::read_address_entry(infos, min_align, tr, ic, 0, oc, 0, + bytes_left, max_fields, fields_total); + // Expectations + ASSERT_EQ(infos.num_rects, 1u); + const auto &ci = infos.subrects[0]; + EXPECT_EQ(ci.extents[0], kContig); + EXPECT_EQ(ci.volume, kContig); + + // Only one field can be moved because bytes_left == contig + EXPECT_EQ(fields_total, kMoveFields); + EXPECT_EQ(ci.src.num_fields, kMoveFields); + EXPECT_EQ(ci.dst.num_fields, 0u); + EXPECT_EQ(moved, kContig * kMoveFields); + + // Pending bytes: src started with 3*64, dst with 64 + EXPECT_EQ(in_al.bytes_pending(), kContig * (kSrcFields - kMoveFields)); + EXPECT_EQ(out_al.bytes_pending(), kLeft); + + EXPECT_EQ(ic.get_offset(), 0); + EXPECT_EQ(oc.get_offset(), kContig * kMoveFields); + } + + infos.num_rects = 0; + bytes_left = + std::min(kBytesLeft, std::min(in_al.bytes_pending(), out_al.bytes_pending())); + + { + size_t moved = GPUXferDes::read_address_entry(infos, min_align, tr, ic, 0, oc, 0, + kLeft, max_fields, fields_total); + + // Expectations + ASSERT_EQ(infos.num_rects, 1u); + const auto &ci = infos.subrects[0]; + EXPECT_EQ(ci.extents[0], kLeft); + EXPECT_EQ(ci.volume, kLeft); + + EXPECT_EQ(fields_total, 1); + EXPECT_EQ(ci.src.num_fields, 1); + EXPECT_EQ(ci.dst.num_fields, 0u); + EXPECT_EQ(moved, kLeft); + + EXPECT_EQ(in_al.bytes_pending(), kContig * (kSrcFields - kMoveFields) - kLeft); + EXPECT_EQ(out_al.bytes_pending(), 0); + + EXPECT_EQ(ic.get_offset(), kLeft); + EXPECT_EQ(oc.get_offset(), 0); // WrapAround + EXPECT_EQ(oc.remaining(0), kContig * kMoveFields + kLeft); + } + + heap.free_obj(fb_src); +} + +// ───────────────────────────────────────────────────────────────────────────── +// 2-D branch with field blocks on both sides, bytes_left moves two fields +// ───────────────────────────────────────────────────────────────────────────── +TEST(ReadAddressEntry, FieldBlock_2D_MoveTwoFields) +{ + constexpr size_t CONTIG = 32; + constexpr size_t LINES = 4; // volume per field = 128 + constexpr int FIELDS = 4; + constexpr size_t BYTES_LEFT = CONTIG * LINES * 2; // allow exactly 2 fields (256) + + AddressList in_al, out_al; + std::vector ids(FIELDS); + std::iota(ids.begin(), ids.end(), 0); + MockHeap h; + auto *fb_in = FieldBlock::create(h, ids.data(), ids.size()); + auto *fb_out = FieldBlock::create(h, ids.data(), ids.size()); + in_al.attach_field_block(fb_in); + out_al.attach_field_block(fb_out); + + append_entry_2d(in_al, CONTIG, LINES); + append_entry_2d(out_al, CONTIG, LINES); + + AddressListCursor ic, oc; + ic.set_addrlist(&in_al); + oc.set_addrlist(&out_al); + AffineCopyInfo<3> info{}; + MemcpyTransposeInfo tr{}; + size_t min_align = 16, fields_total = 0; + + size_t moved = GPUXferDes::read_address_entry( + info, min_align, tr, ic, 0, oc, 0, BYTES_LEFT, /*max_fields=*/8, fields_total); + + // We should hit the 2-D branch (one rectangle) and move exactly two fields + ASSERT_EQ(info.num_rects, 1u); + const auto &ci = info.subrects[0]; + EXPECT_EQ(ci.extents[0], CONTIG); + EXPECT_EQ(ci.extents[1], LINES); + EXPECT_EQ(ci.volume, CONTIG * LINES); + EXPECT_EQ(fields_total, 2u); + EXPECT_EQ(ci.src.num_fields, 2u); + EXPECT_EQ(ci.dst.num_fields, 2u); + EXPECT_EQ(moved, BYTES_LEFT); + + const size_t BYTES_PER_FIELD = CONTIG * LINES; // 128 + EXPECT_EQ(in_al.bytes_pending(), BYTES_PER_FIELD * (FIELDS - 2)); + EXPECT_EQ(out_al.bytes_pending(), BYTES_PER_FIELD * (FIELDS - 2)); + + h.free_obj(fb_in); + h.free_obj(fb_out); +} + +// ───────────────────────────────────────────────────────────────────────────── +// 3-D branch with destination field block only, moves three fields +// ───────────────────────────────────────────────────────────────────────────── +TEST(ReadAddressEntry, DstFieldBlock_3D_MoveThreeFields) +{ + constexpr size_t CONTIG = 16; + constexpr size_t LINES = 2; + constexpr size_t PLANES = 3; // volume per field = 96 + constexpr size_t DST_FIELDS = 5; + constexpr size_t MOVE_FIELDS = 3; + constexpr size_t BYTES_LEFT = + CONTIG * LINES * PLANES * MOVE_FIELDS; // 288 bytes → 3 fields + + AddressList in_al, out_al; + std::vector ids(DST_FIELDS); + std::iota(ids.begin(), ids.end(), 0); + MockHeap h; + auto *fb_dst = FieldBlock::create(h, ids.data(), ids.size()); + out_al.attach_field_block(fb_dst); + + append_entry_3d(in_al, CONTIG, LINES, PLANES * MOVE_FIELDS); // large enough incoming IB + append_entry_3d(out_al, CONTIG, LINES, PLANES); + + AddressListCursor ic, oc; + ic.set_addrlist(&in_al); + oc.set_addrlist(&out_al); + AffineCopyInfo<3> info{}; + MemcpyTransposeInfo tr{}; + size_t min_align = 16, fields_total = 0; + + size_t moved = GPUXferDes::read_address_entry( + info, min_align, tr, ic, 0, oc, 0, BYTES_LEFT, /*max_fields=*/8, fields_total); + + // Expect 3-D branch (one rect) transferring three destination fields + ASSERT_EQ(info.num_rects, 1u); + const auto &ci = info.subrects[0]; + EXPECT_EQ(ci.extents[0], CONTIG); + EXPECT_EQ(ci.extents[1], LINES); + EXPECT_EQ(ci.extents[2], PLANES); + EXPECT_EQ(ci.volume, CONTIG * LINES * PLANES); + EXPECT_EQ(fields_total, 3u); + EXPECT_EQ(ci.src.num_fields, 0u); + EXPECT_EQ(ci.dst.num_fields, 3u); + EXPECT_EQ(moved, BYTES_LEFT); + + const size_t VOL = CONTIG * LINES * PLANES; // 96 + EXPECT_EQ(in_al.bytes_pending(), 0u); + EXPECT_EQ(out_al.bytes_pending(), VOL * (DST_FIELDS - 3)); + + h.free_obj(fb_dst); +} + +// ───────────────────────────────────────────────────────────────────────────── +// Mis-aligned base addresses must force min_align down to 1 +// ───────────────────────────────────────────────────────────────────────────── +TEST(ReadAddressEntry, MisalignedBase_ByteAlignment) +{ + constexpr size_t CONTIG = 64; // 64-byte 1-D rectangle (16-aligned size) + // 1-byte-mis-aligned bases + constexpr uintptr_t IN_BASE = 1; // 0x…01 + constexpr uintptr_t OUT_BASE = 3; // 0x…03 + + AddressList in_al, out_al; + append_entry_1d(in_al, CONTIG); + append_entry_1d(out_al, CONTIG); + + AddressListCursor ic, oc; + ic.set_addrlist(&in_al); + oc.set_addrlist(&out_al); + + AffineCopyInfo<3> info{}; + MemcpyTransposeInfo tr{}; + size_t min_align = 16; // starts optimistic + size_t fields_total = 0; + + size_t moved = + GPUXferDes::read_address_entry(info, min_align, tr, ic, IN_BASE, // <-- mis-aligned + oc, OUT_BASE, // <-- mis-aligned + CONTIG, // bytes_left + /*max_fields*/ 8, fields_total); + + ASSERT_EQ(moved, CONTIG); // full rectangle copied + ASSERT_EQ(info.num_rects, 1u); // fast path 1-D + EXPECT_EQ(min_align, 1u); // **critical check** + EXPECT_EQ(fields_total, 1u); + EXPECT_EQ(in_al.bytes_pending(), 0u); + EXPECT_EQ(out_al.bytes_pending(), 0u); +} + +TEST(ReadAddressEntryTests, DISABLED_Misaligned3D_NoFieldBlock) +{ + constexpr size_t CONTIG = 48; // not 16-byte aligned + constexpr size_t LINES = 3; + constexpr size_t PLANES = 2; + AddressList src_al, dst_al; + append_entry_3d(src_al, CONTIG, LINES, PLANES); + append_entry_3d(dst_al, CONTIG, LINES, PLANES); + + AddressListCursor sc, dc; + sc.set_addrlist(&src_al); + dc.set_addrlist(&dst_al); + + AffineCopyInfo<3> ci{}; + MemcpyTransposeInfo tr{}; + size_t min_align = 16; + size_t fields_tot = 0; + + // this *must* return 0 because CONTIG is not 16-aligned + size_t moved = GPUXferDes::read_address_entry(ci, min_align, tr, sc, 0, dc, 0, + CONTIG * LINES * PLANES, + /*max_fields*/ 1, fields_tot); + EXPECT_EQ(moved, 0u); +}