diff --git a/doc/misc.rst b/doc/misc.rst index c43550c47d8bda6e59a30c502fd4778bba8614e7..97bac9fec35d1960f0b8dceb9489f8399b72520c 100644 --- a/doc/misc.rst +++ b/doc/misc.rst @@ -184,7 +184,7 @@ This list is always growing, but here are a few pointers: * Reorder loops - Use :func:`loopy.set_loop_priority`. + Use :func:`loopy.prioritize_loops`. * Precompute subexpressions: @@ -196,7 +196,7 @@ This list is always growing, but here are a few pointers: * Tile: Use :func:`loopy.split_iname` to produce enough loops, then use - :func:`loopy.set_loop_priority` to set the ordering. + :func:`loopy.prioritize_loops` to set the ordering. * Fix constants diff --git a/doc/tutorial.rst b/doc/tutorial.rst index 87daa9fc4fc01b0625066cfd7c934c046b546930..f8b2f2d8cfd1e1a78f2f51c2272fdc34b65b70f4 100644 --- a/doc/tutorial.rst +++ b/doc/tutorial.rst @@ -360,7 +360,7 @@ Let us take a look at the generated code for the above kernel: .. doctest:: >>> knl = lp.set_options(knl, "write_cl") - >>> knl = lp.set_loop_priority(knl, "i,j") + >>> knl = lp.prioritize_loops(knl, "i,j") >>> evt, (out,) = knl(queue, a=a_mat_dev) #define lid(N) ((int) get_local_id(N)) #define gid(N) ((int) get_group_id(N)) @@ -402,7 +402,7 @@ with identical bounds, for the use of the transpose: ... out[j,i] = a[i,j] {id=transpose} ... out[ii,jj] = 2*out[ii,jj] {dep=transpose} ... """) - >>> knl = lp.set_loop_priority(knl, "i,j,ii,jj") + >>> knl = lp.prioritize_loops(knl, "i,j,ii,jj") :func:`loopy.duplicate_inames` can be used to achieve the same goal. Now the intended code is generated and our test passes. @@ -459,9 +459,9 @@ The warning (and the nondeterminism it warns about) is easily resolved: .. doctest:: - >>> knl = lp.set_loop_priority(knl, "j,i") + >>> knl = lp.prioritize_loops(knl, "j,i") -:func:`loopy.set_loop_priority` indicates the textual order in which loops +:func:`loopy.prioritize_loops` indicates the textual order in which loops should be entered in the kernel code. Note that this priority has an advisory role only. If the kernel logically requires a different nesting, loop priority is ignored. Priority is only considered if loop nesting is @@ -507,7 +507,7 @@ is overwritten with the new kernel:: knl = lp.do_something(knl, arguments...) We've already seen an example of a transformation above: -For instance, :func:`set_loop_priority` fit the pattern. +For instance, :func:`prioritize_loops` fit the pattern. :func:`loopy.split_iname` is another fundamental (and useful) transformation. It turns one existing iname (recall that this is loopy's word for a 'loop @@ -526,7 +526,7 @@ Consider this example: ... "{ [i]: 0<=i>> knl = lp.split_iname(knl, "i", 16) - >>> knl = lp.set_loop_priority(knl, "i_outer,i_inner") + >>> knl = lp.prioritize_loops(knl, "i_outer,i_inner") >>> knl = lp.set_options(knl, "write_cl") >>> evt, (out,) = knl(queue, a=x_vec_dev) #define lid(N) ((int) get_local_id(N)) @@ -554,7 +554,12 @@ relation to loop nesting. For example, it's perfectly possible to request .. doctest:: - >>> knl = lp.set_loop_priority(knl, "i_inner,i_outer") + >>> knl = lp.make_kernel( + ... "{ [i]: 0<=i>> knl = lp.split_iname(knl, "i", 16) + >>> knl = lp.prioritize_loops(knl, "i_inner,i_outer") + >>> knl = lp.set_options(knl, "write_cl") >>> evt, (out,) = knl(queue, a=x_vec_dev) #define lid(N) ((int) get_local_id(N)) ... @@ -568,7 +573,7 @@ Notice how loopy has automatically generated guard conditionals to make sure the bounds on the old iname are obeyed. The combination of :func:`loopy.split_iname` and -:func:`loopy.set_loop_priority` is already good enough to implement what is +:func:`loopy.prioritize_loops` is already good enough to implement what is commonly called 'loop tiling': .. doctest:: @@ -579,7 +584,7 @@ commonly called 'loop tiling': ... assumptions="n mod 16 = 0 and n >= 1") >>> knl = lp.split_iname(knl, "i", 16) >>> knl = lp.split_iname(knl, "j", 16) - >>> knl = lp.set_loop_priority(knl, "i_outer,j_outer,i_inner") + >>> knl = lp.prioritize_loops(knl, "i_outer,j_outer,i_inner") >>> knl = lp.set_options(knl, "write_cl") >>> evt, (out,) = knl(queue, a=a_mat_dev) #define lid(N) ((int) get_local_id(N)) @@ -621,7 +626,7 @@ loop's tag to ``"unr"``: >>> orig_knl = knl >>> knl = lp.split_iname(knl, "i", 4) >>> knl = lp.tag_inames(knl, dict(i_inner="unr")) - >>> knl = lp.set_loop_priority(knl, "i_outer,i_inner") + >>> knl = lp.prioritize_loops(knl, "i_outer,i_inner") >>> knl = lp.set_options(knl, "write_cl") >>> evt, (out,) = knl(queue, a=x_vec_dev) #define lid(N) ((int) get_local_id(N)) @@ -743,7 +748,7 @@ assumption: >>> orig_knl = knl >>> knl = lp.split_iname(knl, "i", 4) >>> knl = lp.tag_inames(knl, dict(i_inner="unr")) - >>> knl = lp.set_loop_priority(knl, "i_outer,i_inner") + >>> knl = lp.prioritize_loops(knl, "i_outer,i_inner") >>> knl = lp.set_options(knl, "write_cl") >>> evt, (out,) = knl(queue, a=x_vec_dev) #define lid(N) ((int) get_local_id(N)) @@ -773,7 +778,7 @@ enabling some cost savings: >>> knl = orig_knl >>> knl = lp.split_iname(knl, "i", 4, slabs=(0, 1), inner_tag="unr") >>> knl = lp.set_options(knl, "write_cl") - >>> knl = lp.set_loop_priority(knl, "i_outer,i_inner") + >>> knl = lp.prioritize_loops(knl, "i_outer,i_inner") >>> evt, (out,) = knl(queue, a=x_vec_dev) #define lid(N) ((int) get_local_id(N)) ... diff --git a/loopy/__init__.py b/loopy/__init__.py index fad6e6eaba910d6a24927456b6bba27920ddc7e6..21a41b11c9b84d288aa2cbb5146db23538613688 100644 --- a/loopy/__init__.py +++ b/loopy/__init__.py @@ -61,7 +61,7 @@ from loopy.library.reduction import register_reduction_parser # {{{ import transforms from loopy.transform.iname import ( - set_loop_priority, + set_loop_priority, prioritize_loops, split_iname, chunk_iname, join_inames, tag_inames, duplicate_inames, rename_iname, remove_unused_inames, split_reduction_inward, split_reduction_outward, @@ -166,7 +166,7 @@ __all__ = [ # {{{ transforms - "set_loop_priority", + "set_loop_priority", "prioritize_loops", "split_iname", "chunk_iname", "join_inames", "tag_inames", "duplicate_inames", "rename_iname", "remove_unused_inames", diff --git a/loopy/check.py b/loopy/check.py index 2f48211da430b42e8000467f070b1633ab3a4f38..2556bee7b74bb1edcd88ca42fb2bc01567b472c2 100644 --- a/loopy/check.py +++ b/loopy/check.py @@ -78,9 +78,10 @@ def check_insn_attributes(kernel): def check_loop_priority_inames_known(kernel): - for iname in kernel.loop_priority: - if iname not in kernel.all_inames(): - raise LoopyError("unknown iname '%s' in loop priorities" % iname) + for prio in kernel.loop_priority: + for iname in prio: + if iname not in kernel.all_inames(): + raise LoopyError("unknown iname '%s' in loop priorities" % iname) def check_for_double_use_of_hw_axes(kernel): diff --git a/loopy/kernel/__init__.py b/loopy/kernel/__init__.py index e0ec0df31dac5cf3ee470d5e6337060cd84a5099..a24cdbb87bc25e3b82bb4336d09f3089cecca164 100644 --- a/loopy/kernel/__init__.py +++ b/loopy/kernel/__init__.py @@ -140,9 +140,10 @@ class LoopKernel(RecordWithoutPickling): .. attribute:: loop_priority - A list of inames. The earlier in the list the iname occurs, the earlier - it will be scheduled. (This applies to inames with non-parallel - implementation tags.) + A frozenset of priority constraints to the kernel. Each such constraint + is a tuple of inames. Inames occuring in such a tuple will be scheduled + earlier than any iname following in the tuple. This applies only to inames + with non-parallel implementation tags. .. attribute:: silenced_warnings @@ -184,7 +185,7 @@ class LoopKernel(RecordWithoutPickling): symbol_manglers=[], iname_slab_increments={}, - loop_priority=[], + loop_priority=frozenset(), silenced_warnings=[], applied_iname_rewrites=[], diff --git a/loopy/schedule/__init__.py b/loopy/schedule/__init__.py index d6eb1a32cdd32b2adaed08f92f4bb63d9501e5ce..21631a1581e5f7cdaf43de5071303336fdbc3a6b 100644 --- a/loopy/schedule/__init__.py +++ b/loopy/schedule/__init__.py @@ -355,6 +355,46 @@ def gen_dependencies_except(kernel, insn_id, except_insn_ids): for sub_dep_id in gen_dependencies_except(kernel, dep_id, except_insn_ids): yield sub_dep_id + +def get_priority_tiers(wanted, priorities): + # Get highest priority tier candidates: These are the first inames + # of all the given priority constraints + candidates = set(next(iter(p for p in prio if p in wanted)) + for prio in priorities + ) + + # Now shrink this set by removing those inames that are prohibited + # by other constraints + bad_candidates = [] + for c1 in candidates: + for c2 in candidates: + for prio in priorities: + try: + if prio.index(c1) < prio.index(c2): + bad_candidates.append(c2) + except ValueError: + # A ValueError in tuple.index just states that one of + # the candidates is not present in the priority constraint + pass + candidates = candidates - set(bad_candidates) + + if candidates: + # We found a valid priority tier! + yield candidates + else: + # If we did not, we stop the generator! + return + + # Now reduce the input data for recursion! + priorities = frozenset([tuple(i for i in prio if i not in candidates) + for prio in priorities + ]) - frozenset([()]) + wanted = wanted - candidates + + # Yield recursively! + for tier in get_priority_tiers(wanted, priorities): + yield tier + # }}} @@ -930,19 +970,27 @@ def generate_loop_schedules_internal( # Build priority tiers. If a schedule is found in the first tier, then # loops in the second are not even tried (and so on). - - loop_priority_set = set(sched_state.kernel.loop_priority) + loop_priority_set = set().union(*[set(prio) + for prio in + sched_state.kernel.loop_priority]) useful_loops_set = set(six.iterkeys(iname_to_usefulness)) useful_and_desired = useful_loops_set & loop_priority_set if useful_and_desired: - priority_tiers = [ - [iname] - for iname in sched_state.kernel.loop_priority - if iname in useful_and_desired - and iname not in sched_state.ilp_inames - and iname not in sched_state.vec_inames - ] + wanted = ( + useful_and_desired + - sched_state.ilp_inames + - sched_state.vec_inames + ) + priority_tiers = [t for t in + get_priority_tiers(wanted, + sched_state.kernel.loop_priority + ) + ] + + # Update the loop priority set, because some constraints may have + # have been contradictary. + loop_priority_set = set().union(*[set(t) for t in priority_tiers]) priority_tiers.append( useful_loops_set diff --git a/loopy/transform/fusion.py b/loopy/transform/fusion.py index 471c039080205015e95167ae6258edea5220351e..77c2d3adecb6db4c0b77c9eb32983c9c04067c43 100644 --- a/loopy/transform/fusion.py +++ b/loopy/transform/fusion.py @@ -267,9 +267,7 @@ def _fuse_two_kernels(knla, knlb): "iname slab increment", knla.iname_slab_increments, knlb.iname_slab_increments), - loop_priority=_ordered_merge_lists( - knla.loop_priority, - knlb.loop_priority), + loop_priority=knla.loop_priority.union(knlb.loop_priority), silenced_warnings=_ordered_merge_lists( knla.silenced_warnings, knlb.silenced_warnings), diff --git a/loopy/transform/iname.py b/loopy/transform/iname.py index 7fd56caa989d2bc06a676316a890ea5fd39e8897..4c2856746a4b82668635a5e49523f71fe7b16695 100644 --- a/loopy/transform/iname.py +++ b/loopy/transform/iname.py @@ -52,12 +52,12 @@ __doc__ = """ .. autofunction:: has_schedulable_iname_nesting +.. autofunction:: prioritize_loops + .. autofunction:: rename_iname .. autofunction:: remove_unused_inames -.. autofunction:: set_loop_priority - .. autofunction:: split_reduction_inward .. autofunction:: split_reduction_outward @@ -78,19 +78,39 @@ __doc__ = """ # {{{ set loop priority def set_loop_priority(kernel, loop_priority): + from warnings import warn + warn("set_loop_priority is deprecated. Use prioritize_loops instead." + "Attention: A call to set_loop_priority will overwrite any previously" + "set priorities!", DeprecationWarning, stacklevel=2) + + if isinstance(loop_priority, str): + loop_priority = tuple(s.strip() + for s in loop_priority.split(",") if s.strip()) + loop_priority = tuple(loop_priority) + + return kernel.copy(loop_priority=frozenset([loop_priority])) + + +def prioritize_loops(kernel, loop_priority): """Indicates the textual order in which loops should be entered in the kernel code. Note that this priority has an advisory role only. If the kernel logically requires a different nesting, priority is ignored. Priority is only considered if loop nesting is ambiguous. - :arg: an iterable of inames, or, for brevity, a comma-seaprated string of + prioritize_loops can be used multiple times. If you do so, each given + *loop_priority* specifies a scheduling constraint. The constraints from + all calls to prioritize_loops together establish a partial order on the + inames (see https://en.wikipedia.org/wiki/Partially_ordered_set). + + :arg: an iterable of inames, or, for brevity, a comma-separated string of inames """ - if isinstance(loop_priority, str): - loop_priority = [s.strip() for s in loop_priority.split(",") if s.strip()] + loop_priority = tuple(s.strip() + for s in loop_priority.split(",") if s.strip()) + loop_priority = tuple(loop_priority) - return kernel.copy(loop_priority=loop_priority) + return kernel.copy(loop_priority=kernel.loop_priority.union([loop_priority])) # }}} @@ -245,20 +265,22 @@ def _split_iname_backend(kernel, split_iname, iname_slab_increments = kernel.iname_slab_increments.copy() iname_slab_increments[outer_iname] = slabs - new_loop_priority = [] - for prio_iname in kernel.loop_priority: - if prio_iname == split_iname: - new_loop_priority.append(outer_iname) - new_loop_priority.append(inner_iname) - else: - new_loop_priority.append(prio_iname) + new_priorities = [] + for prio in kernel.loop_priority: + new_prio = () + for prio_iname in prio: + if prio_iname == split_iname: + new_prio = new_prio + (outer_iname, inner_iname) + else: + new_prio = new_prio + (prio_iname,) + new_priorities.append(new_prio) kernel = kernel.copy( domains=new_domains, iname_slab_increments=iname_slab_increments, instructions=new_insns, applied_iname_rewrites=applied_iname_rewrites, - loop_priority=new_loop_priority) + loop_priority=frozenset(new_priorities)) from loopy.match import parse_stack_match within = parse_stack_match(within) diff --git a/test/gnuma_loopy_transforms.py b/test/gnuma_loopy_transforms.py index 49a2a99bb8d4d518ac23134ba9d1a1ff3d6e8053..5a55f4a3d7d8bc1ff7580ebf1d79bcd1bd7cf7f6 100644 --- a/test/gnuma_loopy_transforms.py +++ b/test/gnuma_loopy_transforms.py @@ -38,7 +38,7 @@ def set_D_storage_format(kernel): def set_up_volume_loop(kernel, Nq): kernel = lp.fix_parameters(kernel, Nq=Nq) - kernel = lp.set_loop_priority(kernel, "e,k,j,i") + kernel = lp.prioritize_loops(kernel, "e,k,j,i") kernel = lp.tag_inames(kernel, dict(e="g.0", j="l.1", i="l.0")) kernel = lp.assume(kernel, "elements >= 1") return kernel diff --git a/test/test_apps.py b/test/test_apps.py index 840fa1f67240113f34ad6d5e88af8e90c76c770a..790a44f6acac72e4fa6fe04a45f32813e6204bb9 100644 --- a/test/test_apps.py +++ b/test/test_apps.py @@ -84,12 +84,12 @@ def test_convolution(ctx_factory): def variant_0(knl): #knl = lp.split_iname(knl, "im_x", 16, inner_tag="l.0") - knl = lp.set_loop_priority(knl, "iimg,im_x,im_y,ifeat,f_x,f_y") + knl = lp.prioritize_loops(knl, "iimg,im_x,im_y,ifeat,f_x,f_y") return knl def variant_1(knl): knl = lp.split_iname(knl, "im_x", 16, inner_tag="l.0") - knl = lp.set_loop_priority(knl, "iimg,im_x_outer,im_y,ifeat,f_x,f_y") + knl = lp.prioritize_loops(knl, "iimg,im_x_outer,im_y,ifeat,f_x,f_y") return knl def variant_2(knl): @@ -151,12 +151,12 @@ def test_convolution_with_nonzero_base(ctx_factory): def variant_0(knl): #knl = lp.split_iname(knl, "im_x", 16, inner_tag="l.0") - knl = lp.set_loop_priority(knl, "iimg,im_x,im_y,ifeat,f_x,f_y") + knl = lp.prioritize_loops(knl, "iimg,im_x,im_y,ifeat,f_x,f_y") return knl def variant_1(knl): knl = lp.split_iname(knl, "im_x", 16, inner_tag="l.0") - knl = lp.set_loop_priority(knl, "iimg,im_x_outer,im_y,ifeat,f_x,f_y") + knl = lp.prioritize_loops(knl, "iimg,im_x_outer,im_y,ifeat,f_x,f_y") return knl for variant in [ @@ -341,7 +341,7 @@ def test_stencil(ctx_factory): knl = lp.split_iname(knl, "i", 16, outer_tag="g.1", inner_tag="l.1") knl = lp.split_iname(knl, "j", 16, outer_tag="g.0", inner_tag="l.0") knl = lp.add_prefetch(knl, "a", ["i_inner", "j_inner"]) - knl = lp.set_loop_priority(knl, ["a_dim_0_outer", "a_dim_1_outer"]) + knl = lp.prioritize_loops(knl, ["a_dim_0_outer", "a_dim_1_outer"]) return knl def variant_2(knl): @@ -349,7 +349,7 @@ def test_stencil(ctx_factory): knl = lp.split_iname(knl, "j", 16, outer_tag="g.0", inner_tag="l.0") knl = lp.add_prefetch(knl, "a", ["i_inner", "j_inner"], fetch_bounding_box=True) - knl = lp.set_loop_priority(knl, ["a_dim_0_outer", "a_dim_1_outer"]) + knl = lp.prioritize_loops(knl, ["a_dim_0_outer", "a_dim_1_outer"]) return knl for variant in [ @@ -396,7 +396,7 @@ def test_stencil_with_overfetch(ctx_factory): slabs=(1, 1)) knl = lp.add_prefetch(knl, "a", ["i_inner", "j_inner"], fetch_bounding_box=True) - knl = lp.set_loop_priority(knl, ["a_dim_0_outer", "a_dim_1_outer"]) + knl = lp.prioritize_loops(knl, ["a_dim_0_outer", "a_dim_1_outer"]) return knl for variant in [variant_overfetch]: diff --git a/test/test_domain.py b/test/test_domain.py index 0631a08796eeda73376afcd1cd03a737a33241f6..531e792704d530d505a789e00344bf9271d0ced5 100644 --- a/test/test_domain.py +++ b/test/test_domain.py @@ -61,7 +61,7 @@ def test_assume(ctx_factory): [lp.GlobalArg("a", np.float32, shape="n"), "..."]) knl = lp.split_iname(knl, "i", 16) - knl = lp.set_loop_priority(knl, "i_outer,i_inner") + knl = lp.prioritize_loops(knl, "i_outer,i_inner") knl = lp.assume(knl, "n mod 16 = 0") knl = lp.assume(knl, "n > 10") knl = lp.preprocess_kernel(knl, ctx.devices[0]) diff --git a/test/test_fortran.py b/test/test_fortran.py index a5b1b830bc8834637d5f4c609fff8232ef7449e6..6e05aa6adba66ce0a1896527249d321de104c512 100644 --- a/test/test_fortran.py +++ b/test/test_fortran.py @@ -406,7 +406,7 @@ def test_fuse_kernels(ctx_factory): inner=(xd_line + "\n" + yd_line), name="xyderiv")) knl = lp.fuse_kernels((xderiv, yderiv)) - knl = lp.set_loop_priority(knl, "e,i,j,k") + knl = lp.prioritize_loops(knl, "e,i,j,k") assert len(knl.temporary_variables) == 2 diff --git a/test/test_linalg.py b/test/test_linalg.py index 7db8566084423ab85cc3870448ddefd60b0bc27b..96b95cf4b44c603e3ee1948c1d5eceecac277bce 100644 --- a/test/test_linalg.py +++ b/test/test_linalg.py @@ -285,7 +285,7 @@ def test_rank_one(ctx_factory): def variant_1(knl): knl = lp.add_prefetch(knl, "a") knl = lp.add_prefetch(knl, "b") - knl = lp.set_loop_priority(knl, ["i", "j"]) + knl = lp.prioritize_loops(knl, ["i", "j"]) knl = lp.add_inames_to_insn(knl, "i", "writes:b_fetch") return knl diff --git a/test/test_loopy.py b/test/test_loopy.py index d407b63fb95ba9ebf2ebf1381c91d61ecfe393c3..c6058f220b90e89ab635e3b9f3711070e5cb3c31 100644 --- a/test/test_loopy.py +++ b/test/test_loopy.py @@ -751,7 +751,7 @@ def test_ilp_loop_bound(ctx_factory): ref_knl = knl - knl = lp.set_loop_priority(knl, "j,i,k") + knl = lp.prioritize_loops(knl, "j,i,k") knl = lp.split_iname(knl, "k", 4, inner_tag="ilp") lp.auto_test_vs_ref(ref_knl, ctx, knl, @@ -788,7 +788,7 @@ def test_slab_decomposition_does_not_double_execute(ctx_factory): knl = ref_knl knl = lp.split_iname(knl, "i", 4, slabs=(0, 1), inner_tag="unr", outer_tag=outer_tag) - knl = lp.set_loop_priority(knl, "i_outer") + knl = lp.prioritize_loops(knl, "i_outer") a = cl.array.empty(queue, 20, np.float32) a.fill(17) @@ -1408,7 +1408,7 @@ def test_index_cse(ctx_factory): end """) knl = lp.tag_inames(knl, "l:unr") - knl = lp.set_loop_priority(knl, "i,j,k,l") + knl = lp.prioritize_loops(knl, "i,j,k,l") knl = lp.add_and_infer_dtypes(knl, {"a": np.float32, "b": np.float32}) knl = lp.fix_parameters(knl, n=5) print(lp.generate_code_v2(knl).device_code()) diff --git a/test/test_nbody.py b/test/test_nbody.py index 540a8f5c3fac4be02bd4bf53f8d5cc92a195c64a..e118b04b997020943d79ec1ba566eff85d56199a 100644 --- a/test/test_nbody.py +++ b/test/test_nbody.py @@ -77,7 +77,7 @@ def test_nbody(ctx_factory): ["x_fetch_j", "x_fetch_k"], default_tag=None) knl = lp.tag_inames(knl, dict(x_fetch_k="unr", x_fetch_j="l.0")) knl = lp.add_prefetch(knl, "x[i,k]", ["k"], default_tag=None) - knl = lp.set_loop_priority(knl, ["j_outer", "j_inner"]) + knl = lp.prioritize_loops(knl, ["j_outer", "j_inner"]) return knl n = 3000 diff --git a/test/test_numa_diff.py b/test/test_numa_diff.py index e4f303f78bd6442ce7d956b66b3d99bffc1d6252..0c304b7a854579007f57ba204cbff8f440aaf5fc 100644 --- a/test/test_numa_diff.py +++ b/test/test_numa_diff.py @@ -69,7 +69,7 @@ def test_gnuma_horiz_kernel(ctx_factory, ilp_multiple, Nq, opt_level): set_q_storage_format, set_D_storage_format) hsv = lp.fix_parameters(hsv, Nq=Nq) - hsv = lp.set_loop_priority(hsv, "e,k,j,i") + hsv = lp.prioritize_loops(hsv, "e,k,j,i") hsv = lp.tag_inames(hsv, dict(e="g.0", j="l.1", i="l.0")) hsv = lp.assume(hsv, "elements >= 1") diff --git a/test/test_reduction.py b/test/test_reduction.py index 29dda7419ae870e9a3d773aefeafb540567e94a1..68f6242440a14eeeb2144762f0d5175f2135ffa2 100644 --- a/test/test_reduction.py +++ b/test/test_reduction.py @@ -463,16 +463,16 @@ def test_poisson_fem(ctx_factory): ref_knl = knl - knl = lp.set_loop_priority(knl, ["c", "j", "i", "k"]) + knl = lp.prioritize_loops(knl, ["c", "j", "i", "k"]) def variant_1(knl): knl = lp.precompute(knl, "dpsi", "i,k,ell", default_tag='for') - knl = lp.set_loop_priority(knl, "c,i,j") + knl = lp.prioritize_loops(knl, "c,i,j") return knl def variant_2(knl): knl = lp.precompute(knl, "dpsi", "i,ell", default_tag='for') - knl = lp.set_loop_priority(knl, "c,i,j") + knl = lp.prioritize_loops(knl, "c,i,j") return knl def add_types(knl): diff --git a/test/test_target.py b/test/test_target.py index e8548ccca6732f2fac6ea84d98d5b06f1d472c6c..b656383e7bbe008892f45159faadd2d195d67a3b 100644 --- a/test/test_target.py +++ b/test/test_target.py @@ -133,7 +133,7 @@ def test_generate_c_snippet(): knl = lp.add_prefetch(knl, "q_v", "k,I", default_tag=None) knl = lp.split_iname(knl, "k", 4, inner_tag="unr", slabs=(0, 1)) - knl = lp.set_loop_priority(knl, "I,k_outer,k_inner") + knl = lp.prioritize_loops(knl, "I,k_outer,k_inner") knl = lp.preprocess_kernel(knl) knl = lp.get_one_scheduled_kernel(knl) @@ -215,7 +215,7 @@ def test_numba_cuda_target(): knl = lp.split_iname(knl, "j", 128, inner_tag='l.0', slabs=(0, 1)) knl = lp.add_prefetch(knl, "X[i,:]") knl = lp.fix_parameters(knl, N=3) - knl = lp.set_loop_priority(knl, "i_inner,j_outer") + knl = lp.prioritize_loops(knl, "i_inner,j_outer") knl = lp.tag_inames(knl, "k:unr") knl = lp.tag_array_axes(knl, "X", "N0,N1") diff --git a/test/test_transform.py b/test/test_transform.py index cf9564274d12b2889d5010bbd807c4bd94fa149f..ac5a26f6a5683bf9e86055a2729ea5ee995dee1e 100644 --- a/test/test_transform.py +++ b/test/test_transform.py @@ -63,7 +63,7 @@ def test_chunk_iname(ctx_factory): ref_knl = knl knl = lp.chunk_iname(knl, "i", 3, inner_tag="l.0") - knl = lp.set_loop_priority(knl, "i_outer, i_inner") + knl = lp.prioritize_loops(knl, "i_outer, i_inner") lp.auto_test_vs_ref(ref_knl, ctx, knl, parameters=dict(n=130)) @@ -369,7 +369,7 @@ def test_precompute_with_preexisting_inames(ctx_factory): knl = lp.precompute(knl, "D2_subst", "i,k", default_tag="for", precompute_inames="ii,jj") - knl = lp.set_loop_priority(knl, "ii,jj,e,j,k") + knl = lp.prioritize_loops(knl, "ii,jj,e,j,k") lp.auto_test_vs_ref( ref_knl, ctx, knl,