В функции numba замените cuda.popc() эквивалентом ЦП, если его нет в CUDA. ⇐ Python

[Anonymous]

Я пишу общий код, который поддерживает как numba-jitting на CPU, так и numba.cuda-jitting на GPU.
Все работает хорошо, за исключением того, что глубоко внутри общего кода я хотел бы использовать встроенную инструкцию, которая подсчитывает количество бит в целом числе. Это cuda.popc() для пути CUDA и вспомогательная функция cpu_popc() для пути CPU. К сожалению, cuda.popc действителен только в ядре графического процессора с числом.cuda, а cpu_popc действителен только в функции ЦП с числом.
Есть ли есть ли способ реализовать как cpu_compute, так и gpu_compute, если не считать дублирования всего общего кода?
Вот простая структура для проверки этого:
# CPU equivalent ctpop, from https://stackoverflow.com/a/77103233
@numba.extending.intrinsic
def popc_helper(typing_context, src):
def codegen(context, builder, signature, args):
return numba.cpython.mathimpl.call_fp_intrinsic(builder, "llvm.ctpop.i64", args)
return numba.uint64(numba.uint64), codegen

@numba.njit(numba.uint64(numba.uint64))
def cpu_popc(x):
"""Return the (population) count of set bits in an integer."""
return popc_helper(x)

@numba.njit
def common_function(x):
# ...
# some_long_code_that_should_not_get_duplicated.
# ...
# return cpu_popc(x) # This works on the CPU path.
return cuda.popc(x) # This works on the GPU path.

@numba.njit
def cpu_compute(n=5):
array_in = np.arange(n)
array_out = np.empty_like(array_in)
for i, value in enumerate(array_in):
array_out = common_function(value)
return array_out

@cuda.jit
def gpu_kernel(array_in, array_out):
thread_index = cuda.grid(1)
if thread_index < len(array_in):
array_out[thread_index] = common_function(array_in[thread_index])

def gpu_compute(n=5):
array_in = np.arange(n)
array_out = cuda.device_array_like(array_in)
gpu_kernel[1, len(array_in)](cuda.to_device(array_in), array_out)
return array_out.copy_to_host()

# print(cpu_compute())
print(gpu_compute())


Подробнее здесь: https://stackoverflow.com/questions/793 ... ot-in-cuda