Hey Nicholas, On Freitag 30 Januar 2009, you wrote: > Attached is a slightly hacky patch for pycuda to enable reading of > source c files. Currently, includes must be in the same directory... I'm > not sure what the best policy is for dependencies; gcc -P -xc++ seems to > fail when encountering CUDA directives.
I like the sentiment of the patch, but I'd prefer if we could get away without the explicit CudaSourceFile, and instead default to automatic include processing. nvcc has an -M (or --generate-dependencies) flag. That could be used instead of makedepend, which is likely not around on Win32. Opinions? Andreas PS: Cc'ed the PyCuda mailing list. You patch reattached for the ML to see.
diff --git a/src/python/driver.py b/src/python/driver.py
index a56e7fb..c0936fa 100644
--- a/src/python/driver.py
+++ b/src/python/driver.py
@@ -418,14 +418,16 @@ def _get_nvcc_version(nvcc):
-def _do_compile(source, options, keep, nvcc, cache_dir):
+def _do_compile(source, options, keep, nvcc, cache_dir, file_root, include_files):
from os.path import join
if cache_dir:
- import md5
- checksum = md5.new()
+ import hashlib
+ checksum = hashlib.md5()
checksum.update(source)
+ for fname in include_files:
+ checksum.update(open(fname, "r").read())
for option in options:
checksum.update(option)
checksum.update(_get_nvcc_version(nvcc))
@@ -439,8 +441,10 @@ def _do_compile(source, options, keep, nvcc, cache_dir):
pass
from tempfile import mkdtemp
+ import shutil
file_dir = mkdtemp()
- file_root = "kernel"
+ for fname in include_files:
+ shutil.copy(fname, file_dir)
cu_file_name = file_root + ".cu"
cu_file_path = join(file_dir, cu_file_name)
@@ -487,12 +491,39 @@ def _do_compile(source, options, keep, nvcc, cache_dir):
+class CudaSourceFile(object):
+ def __init__(self, filename, calcdeps = True, include_files = []):
+ import os
+ filename = os.path.realpath(filename)
+ self.text = "#line 1 \"%s\"\n" %(filename) + open(filename, "r").read()
+ self.filename = filename
+ self.basename_root = os.path.basename(filename).replace(".cu", "")
+ self.include_files = include_files
+ if calcdeps:
+ import subprocess, re
+ # note: possible problems with awkward filenames with ": "
+ mdout = subprocess.Popen(["makedepend", "-Y", "-f-", filename],
+ stdout=subprocess.PIPE).communicate()[0]
+ self.include_files += re.findall(r"^[\w\d \-_\./]+: (.+)$",
+ mdout, flags=re.M)
+
+ def __str__(self):
+ return self.text
+
class SourceModule(object):
def __init__(self, source, nvcc="nvcc",
options=[], keep=False,
no_extern_c=False, arch=None, code=None,
cache_dir=None):
+ if isinstance(source, CudaSourceFile):
+ tmp_file_root = source.basename_root
+ include_files = source.include_files
+ source = str(source)
+ else:
+ tmp_file_root = "kernel"
+ include_files = []
+
if not no_extern_c:
source = 'extern "C" {\n%s\n}\n' % source
@@ -519,7 +550,8 @@ class SourceModule(object):
if code is not None:
options.extend(["-code", code])
- cubin = _do_compile(source, options, keep, nvcc, cache_dir)
+ cubin = _do_compile(source, options, keep, nvcc, cache_dir,
+ tmp_file_root, include_files)
def failsafe_extract(key, cubin):
pattern = r"%s\s*=\s*([0-9]+)" % key
signature.asc
Description: This is a digitally signed message part.
_______________________________________________ PyCuda mailing list [email protected] http://tiker.net/mailman/listinfo/pycuda_tiker.net
