6630787695

Committed 24 Oct 2023 06:21PM UTC coverage: 78.3% (+0.07%) from 78.228%

Build # 6630787695

Build Type

Pull #1178

github

Committed by

web-flow

Commit Message

Merge aeb58f23c into 067cbbba6

Pull Request Pull Request #1178: An experimental kernel dispatcher for numba_dpex.kernel decorator

Run Details

1123 of 1704 branches covered (0.0%)

Branch coverage included in aggregate %.

347 of 406 new or added lines in 12 files covered. (85.47%)

10 existing lines in 1 file now uncovered.

4986 of 6098 relevant lines covered (81.76%)

0.82 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

82.69

/numba_dpex/experimental/launcher.py

# SPDX-FileCopyrightText: 2023 Intel Corporation
#
# SPDX-License-Identifier: Apache-2.0

from llvmlite import ir as llvmir
from numba.core import cgutils, cpu, types
from numba.extending import intrinsic, overload

from numba_dpex import config, dpjit
from numba_dpex.core.exceptions import UnreachableError
from numba_dpex.core.targets.kernel_target import DpexKernelTargetContext
from numba_dpex.core.types import DpnpNdArray, NdRangeType, RangeType
from numba_dpex.core.utils import kernel_launcher as kl
from numba_dpex.dpctl_iface import libsyclinterface_bindings as sycl
from numba_dpex.experimental.kernel_dispatcher import _KernelModule
from numba_dpex.utils import create_null_ptr


def _get_queue_ref_val(
    kernel_targetctx: DpexKernelTargetContext,
    builder: llvmir.IRBuilder,
    kernel_argtys: [types.Type, ...],
    kernel_args,
):
    """
    Get the sycl queue from the first DpnpNdArray argument. Prior passes
    before lowering make sure that compute-follows-data is enforceable
    for a specific call to a kernel. As such, at the stage of lowering
    the queue from the first DpnpNdArray argument can be extracted.
    """

    for arg_num, argty in enumerate(kernel_argtys):
        if isinstance(argty, DpnpNdArray):
            llvm_val = kernel_args[arg_num]
            datamodel = kernel_targetctx.data_model_manager.lookup(argty)
            sycl_queue_attr_pos = datamodel.get_field_position("sycl_queue")
            ptr_to_queue_ref = builder.extract_value(
                llvm_val, sycl_queue_attr_pos
            )
            break

    return ptr_to_queue_ref


def _create_kernel_bundle_from_spirv(
    builder: llvmir.IRBuilder,
    targetctx: cpu.CPUContext,
    queue_ref: llvmir.PointerType,
    kernel_bc: llvmir.Constant,
    kernel_bc_size_in_bytes: int,
):
    dref = sycl.dpctl_queue_get_device(builder, queue_ref)
    cref = sycl.dpctl_queue_get_context(builder, queue_ref)
    args = [
        cref,
        dref,
        kernel_bc,
        llvmir.Constant(llvmir.IntType(64), kernel_bc_size_in_bytes),
        builder.load(create_null_ptr(builder, targetctx)),
    ]
    kbref = sycl.dpctl_kernel_bundle_create_from_spirv(builder, *args)
    sycl.dpctl_context_delete(builder, cref)
    sycl.dpctl_device_delete(builder, dref)

    return kbref


def _get_num_flattened_kernel_args(
    kernel_targetctx: DpexKernelTargetContext,
    kernel_argtys: tuple[types.Type, ...],
):
    num_flattened_kernel_args = 0
    for arg_type in kernel_argtys:
        if isinstance(arg_type, DpnpNdArray):
            datamodel = kernel_targetctx.data_model_manager.lookup(arg_type)
            num_flattened_kernel_args += datamodel.flattened_field_count
        elif arg_type == types.complex64 or arg_type == types.complex128:
            num_flattened_kernel_args += 2
        else:
            num_flattened_kernel_args += 1

    return num_flattened_kernel_args


def _create_kernel_launcher_body(
    codegen_targetctx: cpu.CPUContext,
    kernel_targetctx: DpexKernelTargetContext,
    builder: llvmir.IRBuilder,
    indexer_argty: RangeType | NdRangeType,
    kernel_argtys: tuple[types.Type, ...],
    kernel_module: _KernelModule,
    index_space_arg: llvmir.BaseStructType,
    kernel_args: [llvmir.Instruction, ...],
):
    klbuilder = kl.KernelLaunchIRBuilder(kernel_targetctx, builder)

    if config.DEBUG_KERNEL_LAUNCHER:
        cgutils.printf(
            builder, "DPEX-DEBUG: Inside the kernel launcher function\n"
        )

    kernel_bc_byte_str: llvmir.Constant = codegen_targetctx.insert_const_bytes(
        builder.module,
        bytes=kernel_module.kernel_bitcode,
    )

    num_flattened_kernel_args = _get_num_flattened_kernel_args(
        kernel_targetctx=kernel_targetctx, kernel_argtys=kernel_argtys
    )

    # Create LLVM values for the kernel args list and kernel arg types list
    args_list = klbuilder.allocate_kernel_arg_array(num_flattened_kernel_args)
    args_ty_list = klbuilder.allocate_kernel_arg_ty_array(
        num_flattened_kernel_args
    )
    kernel_args_ptrs = []
    for arg in kernel_args:
        ptr = builder.alloca(arg.type)
        builder.store(arg, ptr)
        kernel_args_ptrs.append(ptr)

    # Populate the args_list and the args_ty_list LLVM arrays
    klbuilder.populate_kernel_args_and_args_ty_arrays(
        callargs_ptrs=kernel_args_ptrs,
        kernel_argtys=kernel_argtys,
        args_list=args_list,
        args_ty_list=args_ty_list,
        datamodel_mgr=kernel_targetctx.data_model_manager,
    )

    if config.DEBUG_KERNEL_LAUNCHER:
        cgutils.printf(
            builder, "DPEX-DEBUG: Populated kernel args and arg type arrays.\n"
        )

    qref = _get_queue_ref_val(
        kernel_targetctx=kernel_targetctx,
        builder=builder,
        kernel_argtys=kernel_argtys,
        kernel_args=kernel_args,
    )

    if config.DEBUG_KERNEL_LAUNCHER:
        cgutils.printf(
            builder,
            "DPEX-DEBUG: Extracted queue pointer from first dpnp array.\n",
        )

    kbref = _create_kernel_bundle_from_spirv(
        builder=builder,
        targetctx=codegen_targetctx,
        queue_ref=qref,
        kernel_bc=kernel_bc_byte_str,
        kernel_bc_size_in_bytes=len(kernel_module.kernel_bitcode),
    )

    if config.DEBUG_KERNEL_LAUNCHER:
        cgutils.printf(
            builder, "DPEX-DEBUG: Generated kernel_bundle from SPIR-V.\n"
        )

    # Get the pointer to the sycl::kernel object in the sycl::kernel_bundle
    kernel_name = codegen_targetctx.insert_const_string(
        builder.module, kernel_module.kernel_name
    )
    kref = sycl.dpctl_kernel_bundle_get_kernel(builder, kbref, kernel_name)

    # Submit synchronous kernel
    # FIXME: Needs to change once we support returning a SyclEvent back to
    # caller.
    if isinstance(indexer_argty, RangeType):
        range_ndim = indexer_argty.ndim
        range_extents = []
        datamodel = kernel_targetctx.data_model_manager.lookup(indexer_argty)
        for dim_num in range(range_ndim):
            dim_pos = datamodel.get_field_position("dim" + str(dim_num))
            range_extents.append(
                builder.extract_value(index_space_arg, dim_pos)
            )

        if config.DEBUG_KERNEL_LAUNCHER:
            cgutils.printf(builder, "DPEX-DEBUG: Submit sync range kernel.\n")

        eref = klbuilder.submit_sycl_kernel(
            sycl_kernel_ref=kref,
            sycl_queue_ref=qref,
            total_kernel_args=num_flattened_kernel_args,
            arg_list=args_list,
            arg_ty_list=args_ty_list,
            global_range=range_extents,
            local_range=[],
            wait_before_return=False,
        )

        if config.DEBUG_KERNEL_LAUNCHER:
            cgutils.printf(builder, "DPEX-DEBUG: Wait on event.\n")

        sycl.dpctl_event_wait(builder, eref)
        sycl.dpctl_event_delete(builder, eref)

    elif isinstance(indexer_argty, NdRangeType):
        ndrange_ndim = indexer_argty.ndim
        grange_extents = []
        lrange_extents = []
        datamodel = kernel_targetctx.data_model_manager.lookup(indexer_argty)
        for dim_num in range(ndrange_ndim):
            gdim_pos = datamodel.get_field_position("gdim" + str(dim_num))
            grange_extents.append(
                builder.extract_value(index_space_arg, gdim_pos)
            )
            ldim_pos = datamodel.get_field_position("ldim" + str(dim_num))
            lrange_extents.append(
                builder.extract_value(index_space_arg, ldim_pos)
            )

        eref = klbuilder.submit_sycl_kernel(
            sycl_kernel_ref=kref,
            sycl_queue_ref=qref,
            total_kernel_args=num_flattened_kernel_args,
            arg_list=args_list,
            arg_ty_list=args_ty_list,
            global_range=grange_extents,
            local_range=lrange_extents,
            wait_before_return=False,
        )
        if config.DEBUG_KERNEL_LAUNCHER:
            cgutils.printf(builder, "DPEX-DEBUG: Wait on event.\n")

        sycl.dpctl_event_wait(builder, eref)
        sycl.dpctl_event_delete(builder, eref)
    else:
        raise UnreachableError

    # Delete the kernel ref
    sycl.dpctl_kernel_delete(builder, kref)
    # Delete the kernel bundle pointer
    sycl.dpctl_kernel_bundle_delete(builder, kbref)


@intrinsic
def intrin_launch_trampoline(typingctx, kernel_fn, index_space, kernel_args):
    kernel_args_list = [arg for arg in kernel_args]
    # signature of this intrinsic
    sig = types.void(kernel_fn, index_space, kernel_args)
    # signature of the kernel_fn
    kernel_sig = types.void(*kernel_args_list)
    kmodule: _KernelModule = kernel_fn.dispatcher.compile(kernel_sig)
    kernel_targetctx = kernel_fn.dispatcher.targetctx

    def codegen(cgctx, builder, sig, llargs):
        kernel_argtys = kernel_sig.args
        kernel_args_unpacked = []
        for pos in range(len(kernel_args)):
            kernel_args_unpacked.append(builder.extract_value(llargs[2], pos))
        _create_kernel_launcher_body(
            codegen_targetctx=cgctx,
            kernel_targetctx=kernel_targetctx,
            builder=builder,
            indexer_argty=sig.args[1],
            kernel_argtys=kernel_argtys,
            kernel_module=kmodule,
            index_space_arg=llargs[1],
            kernel_args=kernel_args_unpacked,
        )

    return sig, codegen


def _launch_trampoline():
    pass


@overload(_launch_trampoline)
def _ol_launch_trampoline(kernel_fn, index_space, *kernel_args):
    def impl(kernel_fn, index_space, *kernel_args):
        intrin_launch_trampoline(kernel_fn, index_space, kernel_args)

    return impl


@dpjit
def call_kernel(kernel_fn, index_space, *kernel_args):
    """Calls a numba_dpex.kernel decorated function from CPython or from another
    dpjit function.

    Args:
        kernel_fn (numba_dpex.experimental.KernelDispatcher): A
        numba_dpex.kernel decorated function that is compiled to a
        KernelDispatcher by numba_dpex.
        index_space (Range | NdRange): A numba_dpex.Range or numba_dpex.NdRange
        type object that specifies the index space for the kernel.
        kernel_args : List of objects that are passed to the numba_dpex.kernel
        decorated function.
    """
    _launch_trampoline(kernel_fn, index_space, *kernel_args)

1	# SPDX-FileCopyrightText: 2023 Intel Corporation
2	#
3	# SPDX-License-Identifier: Apache-2.0
4
5	from llvmlite import ir as llvmir	1✔
6	from numba.core import cgutils, cpu, types	1✔
7	from numba.extending import intrinsic, overload	1✔
8
9	from numba_dpex import config, dpjit	1✔
10	from numba_dpex.core.exceptions import UnreachableError	1✔
11	from numba_dpex.core.targets.kernel_target import DpexKernelTargetContext	1✔
12	from numba_dpex.core.types import DpnpNdArray, NdRangeType, RangeType	1✔
13	from numba_dpex.core.utils import kernel_launcher as kl	1✔
14	from numba_dpex.dpctl_iface import libsyclinterface_bindings as sycl	1✔
15	from numba_dpex.experimental.kernel_dispatcher import _KernelModule	1✔
16	from numba_dpex.utils import create_null_ptr	1✔
17
18
19	def _get_queue_ref_val(	1✔
20	kernel_targetctx: DpexKernelTargetContext,
21	builder: llvmir.IRBuilder,
22	kernel_argtys: [types.Type, ...],
23	kernel_args,
24	):
25	"""
26	Get the sycl queue from the first DpnpNdArray argument. Prior passes
27	before lowering make sure that compute-follows-data is enforceable
28	for a specific call to a kernel. As such, at the stage of lowering
29	the queue from the first DpnpNdArray argument can be extracted.
30	"""
31
32	for arg_num, argty in enumerate(kernel_argtys):	1!
33	if isinstance(argty, DpnpNdArray):	1!
34	llvm_val = kernel_args[arg_num]	1✔
35	datamodel = kernel_targetctx.data_model_manager.lookup(argty)	1✔
36	sycl_queue_attr_pos = datamodel.get_field_position("sycl_queue")	1✔
37	ptr_to_queue_ref = builder.extract_value(	1✔
38	llvm_val, sycl_queue_attr_pos
39	)
40	break	1✔
41
42	return ptr_to_queue_ref	1✔
43
44
45	def _create_kernel_bundle_from_spirv(	1✔
46	builder: llvmir.IRBuilder,
47	targetctx: cpu.CPUContext,
48	queue_ref: llvmir.PointerType,
49	kernel_bc: llvmir.Constant,
50	kernel_bc_size_in_bytes: int,
51	):
52	dref = sycl.dpctl_queue_get_device(builder, queue_ref)	1✔
53	cref = sycl.dpctl_queue_get_context(builder, queue_ref)	1✔
54	args = [	1✔
55	cref,
56	dref,
57	kernel_bc,
58	llvmir.Constant(llvmir.IntType(64), kernel_bc_size_in_bytes),
59	builder.load(create_null_ptr(builder, targetctx)),
60	]
61	kbref = sycl.dpctl_kernel_bundle_create_from_spirv(builder, *args)	1✔
62	sycl.dpctl_context_delete(builder, cref)	1✔
63	sycl.dpctl_device_delete(builder, dref)	1✔
64
65	return kbref	1✔
66
67
68	def _get_num_flattened_kernel_args(	1✔
69	kernel_targetctx: DpexKernelTargetContext,
70	kernel_argtys: tuple[types.Type, ...],
71	):
72	num_flattened_kernel_args = 0	1✔
73	for arg_type in kernel_argtys:	1✔
74	if isinstance(arg_type, DpnpNdArray):	1!
75	datamodel = kernel_targetctx.data_model_manager.lookup(arg_type)	1✔
76	num_flattened_kernel_args += datamodel.flattened_field_count	1✔
NEW 77	elif arg_type == types.complex64 or arg_type == types.complex128:	×
NEW 78	num_flattened_kernel_args += 2	×
79	else:
NEW 80	num_flattened_kernel_args += 1	×
81
82	return num_flattened_kernel_args	1✔
83
84
85	def _create_kernel_launcher_body(	1✔
86	codegen_targetctx: cpu.CPUContext,
87	kernel_targetctx: DpexKernelTargetContext,
88	builder: llvmir.IRBuilder,
89	indexer_argty: RangeType \| NdRangeType,
90	kernel_argtys: tuple[types.Type, ...],
91	kernel_module: _KernelModule,
92	index_space_arg: llvmir.BaseStructType,
93	kernel_args: [llvmir.Instruction, ...],
94	):
95	klbuilder = kl.KernelLaunchIRBuilder(kernel_targetctx, builder)	1✔
96
97	if config.DEBUG_KERNEL_LAUNCHER:	1!
NEW 98	cgutils.printf(	×
99	builder, "DPEX-DEBUG: Inside the kernel launcher function\n"
100	)
101
102	kernel_bc_byte_str: llvmir.Constant = codegen_targetctx.insert_const_bytes(	1✔
103	builder.module,
104	bytes=kernel_module.kernel_bitcode,
105	)
106
107	num_flattened_kernel_args = _get_num_flattened_kernel_args(	1✔
108	kernel_targetctx=kernel_targetctx, kernel_argtys=kernel_argtys
109	)
110
111	# Create LLVM values for the kernel args list and kernel arg types list
112	args_list = klbuilder.allocate_kernel_arg_array(num_flattened_kernel_args)	1✔
113	args_ty_list = klbuilder.allocate_kernel_arg_ty_array(	1✔
114	num_flattened_kernel_args
115	)
116	kernel_args_ptrs = []	1✔
117	for arg in kernel_args:	1✔
118	ptr = builder.alloca(arg.type)	1✔
119	builder.store(arg, ptr)	1✔
120	kernel_args_ptrs.append(ptr)	1✔
121
122	# Populate the args_list and the args_ty_list LLVM arrays
123	klbuilder.populate_kernel_args_and_args_ty_arrays(	1✔
124	callargs_ptrs=kernel_args_ptrs,
125	kernel_argtys=kernel_argtys,
126	args_list=args_list,
127	args_ty_list=args_ty_list,
128	datamodel_mgr=kernel_targetctx.data_model_manager,
129	)
130
131	if config.DEBUG_KERNEL_LAUNCHER:	1!
NEW 132	cgutils.printf(	×
133	builder, "DPEX-DEBUG: Populated kernel args and arg type arrays.\n"
134	)
135
136	qref = _get_queue_ref_val(	1✔
137	kernel_targetctx=kernel_targetctx,
138	builder=builder,
139	kernel_argtys=kernel_argtys,
140	kernel_args=kernel_args,
141	)
142
143	if config.DEBUG_KERNEL_LAUNCHER:	1!
NEW 144	cgutils.printf(	×
145	builder,
146	"DPEX-DEBUG: Extracted queue pointer from first dpnp array.\n",
147	)
148
149	kbref = _create_kernel_bundle_from_spirv(	1✔
150	builder=builder,
151	targetctx=codegen_targetctx,
152	queue_ref=qref,
153	kernel_bc=kernel_bc_byte_str,
154	kernel_bc_size_in_bytes=len(kernel_module.kernel_bitcode),
155	)
156
157	if config.DEBUG_KERNEL_LAUNCHER:	1!
NEW 158	cgutils.printf(	×
159	builder, "DPEX-DEBUG: Generated kernel_bundle from SPIR-V.\n"
160	)
161
162	# Get the pointer to the sycl::kernel object in the sycl::kernel_bundle
163	kernel_name = codegen_targetctx.insert_const_string(	1✔
164	builder.module, kernel_module.kernel_name
165	)
166	kref = sycl.dpctl_kernel_bundle_get_kernel(builder, kbref, kernel_name)	1✔
167
168	# Submit synchronous kernel
169	# FIXME: Needs to change once we support returning a SyclEvent back to
170	# caller.
171	if isinstance(indexer_argty, RangeType):	1✔
172	range_ndim = indexer_argty.ndim	1✔
173	range_extents = []	1✔
174	datamodel = kernel_targetctx.data_model_manager.lookup(indexer_argty)	1✔
175	for dim_num in range(range_ndim):	1✔
176	dim_pos = datamodel.get_field_position("dim" + str(dim_num))	1✔
177	range_extents.append(	1✔
178	builder.extract_value(index_space_arg, dim_pos)
179	)
180
181	if config.DEBUG_KERNEL_LAUNCHER:	1!
NEW 182	cgutils.printf(builder, "DPEX-DEBUG: Submit sync range kernel.\n")	×
183
184	eref = klbuilder.submit_sycl_kernel(	1✔
185	sycl_kernel_ref=kref,
186	sycl_queue_ref=qref,
187	total_kernel_args=num_flattened_kernel_args,
188	arg_list=args_list,
189	arg_ty_list=args_ty_list,
190	global_range=range_extents,
191	local_range=[],
192	wait_before_return=False,
193	)
194
195	if config.DEBUG_KERNEL_LAUNCHER:	1!
NEW 196	cgutils.printf(builder, "DPEX-DEBUG: Wait on event.\n")	×
197
198	sycl.dpctl_event_wait(builder, eref)	1✔
199	sycl.dpctl_event_delete(builder, eref)	1✔
200
201	elif isinstance(indexer_argty, NdRangeType):	1!
202	ndrange_ndim = indexer_argty.ndim	1✔
203	grange_extents = []	1✔
204	lrange_extents = []	1✔
205	datamodel = kernel_targetctx.data_model_manager.lookup(indexer_argty)	1✔
206	for dim_num in range(ndrange_ndim):	1✔
207	gdim_pos = datamodel.get_field_position("gdim" + str(dim_num))	1✔
208	grange_extents.append(	1✔
209	builder.extract_value(index_space_arg, gdim_pos)
210	)
211	ldim_pos = datamodel.get_field_position("ldim" + str(dim_num))	1✔
212	lrange_extents.append(	1✔
213	builder.extract_value(index_space_arg, ldim_pos)
214	)
215
216	eref = klbuilder.submit_sycl_kernel(	1✔
217	sycl_kernel_ref=kref,
218	sycl_queue_ref=qref,
219	total_kernel_args=num_flattened_kernel_args,
220	arg_list=args_list,
221	arg_ty_list=args_ty_list,
222	global_range=grange_extents,
223	local_range=lrange_extents,
224	wait_before_return=False,
225	)
226	if config.DEBUG_KERNEL_LAUNCHER:	1!
NEW 227	cgutils.printf(builder, "DPEX-DEBUG: Wait on event.\n")	×
228
229	sycl.dpctl_event_wait(builder, eref)	1✔
230	sycl.dpctl_event_delete(builder, eref)	1✔
231	else:
NEW 232	raise UnreachableError	×
233
234	# Delete the kernel ref
235	sycl.dpctl_kernel_delete(builder, kref)	1✔
236	# Delete the kernel bundle pointer
237	sycl.dpctl_kernel_bundle_delete(builder, kbref)	1✔
238
239
240	@intrinsic	1✔
241	def intrin_launch_trampoline(typingctx, kernel_fn, index_space, kernel_args):	1✔
242	kernel_args_list = [arg for arg in kernel_args]	1✔
243	# signature of this intrinsic
244	sig = types.void(kernel_fn, index_space, kernel_args)	1✔
245	# signature of the kernel_fn
246	kernel_sig = types.void(*kernel_args_list)	1✔
247	kmodule: _KernelModule = kernel_fn.dispatcher.compile(kernel_sig)	1✔
248	kernel_targetctx = kernel_fn.dispatcher.targetctx	1✔
249
250	def codegen(cgctx, builder, sig, llargs):	1✔
251	kernel_argtys = kernel_sig.args	1✔
252	kernel_args_unpacked = []	1✔
253	for pos in range(len(kernel_args)):	1✔
254	kernel_args_unpacked.append(builder.extract_value(llargs[2], pos))	1✔
255	_create_kernel_launcher_body(	1✔
256	codegen_targetctx=cgctx,
257	kernel_targetctx=kernel_targetctx,
258	builder=builder,
259	indexer_argty=sig.args[1],
260	kernel_argtys=kernel_argtys,
261	kernel_module=kmodule,
262	index_space_arg=llargs[1],
263	kernel_args=kernel_args_unpacked,
264	)
265
266	return sig, codegen	1✔
267
268
269	def _launch_trampoline():	1✔
NEW 270	pass	×
271
272
273	@overload(_launch_trampoline)	1✔
274	def _ol_launch_trampoline(kernel_fn, index_space, *kernel_args):	1✔
275	def impl(kernel_fn, index_space, *kernel_args):	1✔
NEW 276	intrin_launch_trampoline(kernel_fn, index_space, kernel_args)	×
277
278	return impl	1✔
279
280
281	@dpjit	1✔
282	def call_kernel(kernel_fn, index_space, *kernel_args):	1✔
283	"""Calls a numba_dpex.kernel decorated function from CPython or from another
284	dpjit function.
285
286	Args:
287	kernel_fn (numba_dpex.experimental.KernelDispatcher): A
288	numba_dpex.kernel decorated function that is compiled to a
289	KernelDispatcher by numba_dpex.
290	index_space (Range \| NdRange): A numba_dpex.Range or numba_dpex.NdRange
291	type object that specifies the index space for the kernel.
292	kernel_args : List of objects that are passed to the numba_dpex.kernel
293	decorated function.
294	"""
NEW 295	_launch_trampoline(kernel_fn, index_space, *kernel_args)	×

IntelPython / numba-dpex / 6630787695

Source File Press 'n' to go to next uncovered line, 'b' for previous

Source File
Press 'n' to go to next uncovered line, 'b' for previous