8000 Allow vectorization of ternaries, indices · Issue #870 · inducer/loopy · GitHub
[go: up one dir, main page]
More Web Proxy on the site http://driver.im/
Skip to content
Allow vectorization of ternaries, indices #870
Open
@nkoskelo

Description

@nkoskelo
knl = lp.make_kernel( " { [i] : 0 <= i  < 8}"," out[i] = a if i == 0 else (b if i == 1 else c)")
knl = lp.tag_inames(knl, {"i": "vec"})
from loopy.kernel.array import VectorArrayDimTag
knl = lp.tag_array_axes(knl, "out", [VectorArrayDimTag()])
knl = lp.add_and_infer_dtypes(knl, {"a": np.float32, "b": np.float32, "c": np.float32})
lp.generate_code_v2(knl).device_code()

fails due to there being a dependence on i within the store operation due to the presence of an if statement.

Removing the tag_array_axes call and trying again will result in a warning but device code will be generated.

knl = lp.make_kernel( " { [i] : 0 <= i  < 8}"," out[i] = a if i == 0 else (b if i == 1 else c)")

knl = lp.tag_inames(knl, {"i": "vec"})
from loopy.kernel.array import VectorArrayDimTag
lp.generate_code_v2(knl).device_code()

will generate the following device code:

__kernel void __attribute__  (( reqn_work_group_size(1,1,1))) loopy_kernel(float const a, float const b, float const c, __global float *__restrict__ out)
{
out[0] = ((0 == 0) ? a : ((0 == 1) ? b : c));
out[1] = ((1 == 0) ? a : ((1 == 1) ? b : c));
out[2] = ((2 == 0) ? a : ((2 == 1) ? b : c));
out[3] = ((3 == 0) ? a : ((3 == 1) ? b : c));
out[4] = ((4 == 0) ? a : ((4 == 1) ? b : c));
out[5] = ((5 == 0) ? a : ((5 == 1) ? b : c));
out[6] = ((6 == 0) ? a : ((6 == 1) ? b : c));
out[7] = ((7 == 0) ? a : ((7 == 1) ? b : c));
}

Metadata

Metadata

Assignees

No one assigned

    Labels

    No labels
    No labels

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions

      0