def METHOD_NAME(ctype_str, dtype, ctypes_ctor):
    try:
        q = dpctl.SyclQueue("opencl", property="enable_profiling")
    except dpctl.SyclQueueCreationError:
        pytest.skip("OpenCL queue could not be created")
    if dtype == dpt.dtype("f8") and q.sycl_device.has_aspect_fp64 is False:
        pytest.skip(
            "Device does not support double precision floating point type"
        )
    # OpenCL conventions for indexing global_id is opposite to
    # that of SYCL (and DPCTL)
    oclSrc = (
        "kernel void axpy("
        "   global " + ctype_str + " *a, global " + ctype_str + " *b,"
        "   global " + ctype_str + " *c, " + ctype_str + " d) {"
        "   size_t index = get_global_id(0);"
        "   c[index] = d * a[index] + b[index];"
        "}"
    )
    prog = dpctl_prog.create_program_from_source(q, oclSrc)
    axpyKernel = prog.get_sycl_kernel("axpy")
    n_elems = 1024 * 512
    lws = 128
    a = dpt.arange(n_elems, dtype=dtype, sycl_queue=q)
    b = dpt.arange(n_elems, stop=0, step=-1, dtype=dtype, sycl_queue=q)
    c = dpt.zeros(n_elems, dtype=dtype, sycl_queue=q)
    d = 2
    args = [a.usm_data, b.usm_data, c.usm_data, ctypes_ctor(d)]
    assert n_elems % lws == 0
    b_np = dpt.asnumpy(b)
    a_np = dpt.asnumpy(a)
    for r in (
        [
            n_elems,
        ],
        [2, n_elems],
        [2, 2, n_elems],
    ):
        c[:] = 0
        timer = dpctl.SyclTimer()
        with timer(q):
            q.submit(axpyKernel, args, r).wait()
            ref_c = a_np * np.array(d, dtype=dtype) + b_np
        host_dt, device_dt = timer.dt
        assert type(host_dt) is float and type(device_dt) is float
        assert np.allclose(dpt.asnumpy(c), ref_c), "Failed for {}".format(r)
    for gr, lr in (
        (
            [
                n_elems,
            ],
            [lws],
        ),
        ([2, n_elems], [2, lws // 2]),
        ([2, 2, n_elems], [2, 2, lws // 4]),
    ):
        c[:] = 0
        timer = dpctl.SyclTimer()
        with timer(q):
            q.submit(axpyKernel, args, gr, lr, [dpctl.SyclEvent()]).wait()
            ref_c = a_np * np.array(d, dtype=dtype) + b_np
        host_dt, device_dt = timer.dt
        assert type(host_dt) is float and type(device_dt) is float
        assert np.allclose(dpt.asnumpy(c), ref_c), "Failed for {}, {}".formatg(
            r, lr
        )