PyCUDA/Thrust interop

thrust_demo.py

import pycuda
import pycuda.autoinit
import pycuda.gpuarray as gpuarray
import numpy as np

from codepy.cgen import *
from codepy.bpl import BoostPythonModule
from codepy.cuda import CudaModule

#Make a host_module, compiled for CPU
host_mod = BoostPythonModule()

#Make a device module, compiled with NVCC
nvcc_mod = CudaModule(host_mod)

#Describe device module code
#NVCC includes
nvcc_includes = [
    'thrust/sort.h',
    'thrust/device_vector.h',
    'cuda.h',
    ]
#Add includes to module
nvcc_mod.add_to_preamble([Include(x) for x in nvcc_includes])

#NVCC function
nvcc_function = FunctionBody(
    FunctionDeclaration(Value('void', 'my_sort'),
                        [Value('CUdeviceptr', 'input_ptr'),
                         Value('int', 'length')]),
    Block([Statement('thrust::device_ptr<float> thrust_ptr((float*)input_ptr)'),
           Statement('thrust::sort(thrust_ptr, thrust_ptr+length)')]))

#Add declaration to nvcc_mod
#Adds declaration to host_mod as well
nvcc_mod.add_function(nvcc_function)

host_includes = [
    'boost/python/extract.hpp',
    ]
#Add host includes to module
host_mod.add_to_preamble([Include(x) for x in host_includes])

host_namespaces = [
    'using namespace boost::python',
    ]

#Add BPL using statement
host_mod.add_to_preamble([Statement(x) for x in host_namespaces])


host_statements = [
    #Extract information from PyCUDA GPUArray
    #Get length
    'tuple shape = extract<tuple>(gpu_array.attr("shape"))',
    'int length = extract<int>(shape[0])',
    #Get data pointer
    'CUdeviceptr ptr = extract<CUdeviceptr>(gpu_array.attr("gpudata"))',
    #Call Thrust routine, compiled into the CudaModule
    'my_sort(ptr, length)',
    #Return result
    'return gpu_array',
    ]

host_mod.add_function(
    FunctionBody(
        FunctionDeclaration(Value('object', 'host_entry'),
                            [Value('object', 'gpu_array')]),
        Block([Statement(x) for x in host_statements])))

#Print out generated code, to see what we're actually compiling
print("---------------------- Host code ----------------------")
print(host_mod.generate())
print("--------------------- Device code ---------------------")
print(nvcc_mod.generate())
print("-------------------------------------------------------")
      


#Compile modules
import codepy.jit, codepy.toolchain
gcc_toolchain = codepy.toolchain.guess_toolchain()
nvcc_toolchain = codepy.toolchain.guess_nvcc_toolchain()

module = nvcc_mod.compile(gcc_toolchain, nvcc_toolchain, debug=True)



length = 100
a = np.array(np.random.rand(length), dtype=np.float32)
print("---------------------- Unsorted -----------------------")
print(a)
b = gpuarray.to_gpu(a)
# Call Thrust!!
c = module.host_entry(b)
print("----------------------- Sorted ------------------------")
print c.get()
print("-------------------------------------------------------")

I get this error when trying your example:

CompileError Traceback (most recent call last)
in
4 nvcc_toolchain = codepy.toolchain.guess_nvcc_toolchain()
5
----> 6 module = nvcc_mod.compile(gcc_toolchain, nvcc_toolchain, debug=True)
7
8

~/anaconda3/envs/test/lib/python3.6/site-packages/codepy/cuda.py in compile(self, host_toolchain, nvcc_toolchain, host_kwargs, nvcc_kwargs, **kwargs)
81 host_checksum, host_mod_name, host_object, host_compiled = compile_from_string(
82 host_toolchain, self.boost_module.name, host_code,
---> 83 object=True, **local_host_kwargs)
84 device_checksum, device_mod_name, device_object, device_compiled = compile_from_string(
85 nvcc_toolchain, 'gpu', device_code, 'gpu.cu',

~/anaconda3/envs/test/lib/python3.6/site-packages/codepy/jit.py in compile_from_string(toolchain, name, source_string, source_name, cache_dir, debug, wait_on_error, debug_recompile, object, source_is_binary, sleep_delay)
429
430 if object:
--> 431 toolchain.build_object(ext_file, source_paths, debug=debug)
432 else:
433 toolchain.build_extension(ext_file, source_paths, debug=debug)

~/anaconda3/envs/test/lib/python3.6/site-packages/codepy/toolchain.py in build_object(self, ext_file, source_files, debug)
190 print("FAILED compiler invocation:" +
191 " ".join(cc_cmdline), file=sys.stderr)
--> 192 raise CompileError("module compilation failed")
193
194 def build_extension(self, ext_file, source_files, debug=False):

CompileError: module compilation failed

Any suggestion is welcome

Hi Bryan,

In trying to run this example, I am stuck on this issue for a bit, anyway, you can help me or point me as to where and what I need to fix?

TIA

Traceback (most recent call last):
File "/home/john/anaconda3/lib/python3.8/site-packages/pytools/prefork.py", line 48, in call_capture_output
popen = Popen(cmdline, cwd=cwd, stdin=PIPE, stdout=PIPE,
File "/home/john/anaconda3/lib/python3.8/subprocess.py", line 858, in init
self._execute_child(args, executable, preexec_fn, close_fds,
File "/home/john/anaconda3/lib/python3.8/subprocess.py", line 1706, in _execute_child
raise child_exception_type(errno_num, err_msg, err_filename)
FileNotFoundError: [Errno 2] No such file or directory: 'x86_64-conda_cos6-linux-gnu-c++'

During handling of the above exception, another exception occurred:

Traceback (most recent call last):
File "thrustcudatest.py", line 82, in
gcc_toolchain = codepy.toolchain.guess_toolchain()
File "/home/john/anaconda3/lib/python3.8/site-packages/codepy/toolchain.py", line 443, in guess_toolchain
result, version, stderr = call_capture_output([kwargs["cc"], "--version"])
File "/home/john/anaconda3/lib/python3.8/site-packages/codepy/toolchain.py", line 433, in call_capture_output
result, stdout, stderr = call_capture_output(*args)
File "/home/john/anaconda3/lib/python3.8/site-packages/pytools/prefork.py", line 226, in call_capture_output
return forker.call_capture_output(cmdline, cwd, error_on_nonzero)
File "/home/john/anaconda3/lib/python3.8/site-packages/pytools/prefork.py", line 59, in call_capture_output
raise ExecError("error invoking '%s': %s"
pytools.prefork.ExecError: error invoking 'x86_64-conda_cos6-linux-gnu-c++ --version': [Errno 2] No such file or directory: 'x86_64-conda_cos6-linux-gnu-c++'