Adds nvcc as a RawKernel backend

[sjperkins]

Adds nvcc as a backend for RawKernel (issue #1928). The nvcc.py is a cut and paste of much of the functionality in cupy_setup_build.py and install/*.py -- I've avoided refactoring this at the moment as I'd like to invite input as to whether you think this PR is useful.

[niboshi]

Thank you for the PR and I'm sorry for late review.

Does it have to rely on distutils? How about implementing a class with the same interface as _NVRTCProgram (cupy/cuda/compiler.py)?

[kmaehashi]

I discussed with @okuta and agreed to support NVCC in RawKernel.

cupy_setup_build.py and install/*.py has lots of clutters. Originally CuPy was using NVCC, so I think it's better to borrow some code from here, instead of using above files: https://github.com/cupy/cupy/blob/v1/cupy/cuda/compiler.py

[sjperkins]

Thanks for the suggestions. I'll try find some time to incorporate them in the next while.

[sjperkins]

I think this is now almost ready. One thing I'm still not quite sure how to handle with the nvcc route is the _preprocess call. compile_with_nvcc directly produces a cubin (a bytes type) while compile_with_nvrtc produces unicode from which a bytes type cubin is produced. I haven't succeeded in plugging the output of compile_using_nvcc into the preprocessing code, but this could be because I'm not experienced with unicode. I've tried running a .decode('utf-8') on the cubin but this fails on translation.

Any suggestions on how to handle this?

[niboshi]

Sorry for late review. LGTM except this comment.

[niboshi]

Jenkins, test this please

[chainer-ci]

Jenkins CI test (for commit a63b2a8, target branch master) failed with status FAILURE.

[sjperkins]

Will take a look at the test case failures.

[niboshi]

Those failures in TEST=chainer-py3/py35 (such as TestBias) are irrelevant. Please ignore them. I'm sorry for inconvenience.

Those in TEST=cupy-py??? seem relevant.

[sjperkins]

I think this is now almost ready. One thing I'm still not quite sure how to handle with the nvcc route is the _preprocess call. compile_with_nvcc directly produces a cubin (a bytes type) while compile_with_nvrtc produces unicode from which a bytes type cubin is produced. I haven't succeeded in plugging the output of compile_using_nvcc into the preprocessing code, but this could be because I'm not experienced with unicode. I've tried running a .decode('utf-8') on the cubin but this fails on translation.

OK so the test cases are failing because of a problem that I ran into previously. It doesn't seem possible to decode the test case cubin bytes produced by nvcc into unicode, as there's some byte that causes a failure in translation. So the following

with codecs.open(cubin_path, 'rb', encoding='utf-8') as bin_file:
    return bin_file.read()

produces:

UnicodeDecodeError: 'utf8' codec can't decode byte 0xbe in position 18: invalid start byte

Without this conversion, it's not possible for the cubin to interact with other unicode variables in compiler.py.

Could someone with more unicode experience advise on how to handle this situation?

[leofang]

I am not familiar with unicode either, but does the follow piece of code help? It describes how a unicode object is generated from the ptx:

cupy/cupy/cuda/nvrtc.pyx

Lines 111 to 127 in f85374d

	cpdef unicode getPTX(size_t prog):
	cdef size_t ptxSizeRet
	cdef bytes ptx
	cdef char* ptx_ptr
	with nogil:
	status = nvrtcGetPTXSize(<Program>prog, &ptxSizeRet)
	check_status(status)
	ptx = b' ' * ptxSizeRet
	ptx_ptr = ptx
	with nogil:
	status = nvrtcGetPTX(<Program>prog, ptx_ptr)
	check_status(status)
	# Strip the trailing NULL.
	assert ptx.endswith(b'\x00')
	ptx = ptx[:-1]
	return ptx.decode('UTF-8')

[sjperkins]

I am not familiar with unicode either, but does the follow piece of code help? It describes how a unicode object is generated from the ptx:

Thanks for the suggestion of removing the trailing null, but unfortunately it still fails on the byte at position 18.

[sjperkins]

Having thought about this a bit more, I think the issue is that compile_using_nvrtc returns text PTX assembly, while compile_using_nvcc produces a binary cubin. I probably need to modify the nvcc command within _preprocess to just return the PTX... Let me see what I can do.

[sjperkins]

OK, got it working.

[niboshi]

Thanks.
Jenkins, test this please

[chainer-ci]

Jenkins CI test (for commit 5b205f3, target branch master) failed with status FAILURE.

[niboshi]

The failure seems relevant.
Could you check again?

[sjperkins]

@niboshi I've been taking a look at the logs. Somewhat difficult to diagnose as it's failing on python2 but succeeding on python3. It also succeeds on python2 on my machine.

The error is:

CUDADriverError: CUDA_ERROR_INVALID_SOURCE: device kernel image is invalid

This seems to be a somewhat broad range of errors. For example:

1. It may be that the LD_LIBRARY_PATH needs to be set up to point at appropriate CUDA and NVIDIA drivers on the test machine. See NVRTC_ERROR: device kernel image is invalid fireice-uk/xmr-stak#2319 (comment) for instance.
2. I see there's also a stackoverflow question for chainer on this issue. Are there multiple GPUs on the test system, for instance?

[niboshi]

OK.
Jenkins, test this please

[pfn-ci-bot]

Successfully created a job for commit d4ff126:

• Dashboard for commit d4ff126

[chainer-ci]

Jenkins CI test (for commit d4ff126, target branch master) failed with status FAILURE.

[leofang]

@niboshi Looks good, thanks. The two build test errors are to be fixed in #2514.

[leofang]

In 54ad318 I fixed TypeError: Expected unicode, got str observed in Jenkins.

Jenkins confirmed this fix btw.

[leofang]

@niboshi I did four things:

1. First revert back to d4ff126 as the -cubin flag was already added
2. Enforce the correct backend selection in the tests
3. Clear the preprocessing cache in tearDown()
4. Add backend name to CompileException

I suspect there're some subtle cubin shadowing in the PY2 failure. Not sure if the changes 2 & 3 would fix it, but they're necessary anyway.

Could you kick off Jenkins again for me please? Sorry for keeping bugging you. Thanks.

[niboshi]

chainer/chainer-test#523 has been merged.

@leofang Sure, sorry for delay.
Jenkins, test this please

[pfn-ci-bot]

Successfully created a job for commit 45534ae:

• Dashboard for commit 45534ae

[chainer-ci]

Jenkins CI test (for commit 45534ae, target branch master) succeeded!

[leofang]

Ah, too bad PY2 is disabled! I'd love to know where it was fixed...Anyway, thanks for moving forward @niboshi.

[niboshi]

Thank you. LGTM except some cosmetics.

[leofang]

Thank you, @niboshi. Just committed.

[sjperkins]

@leofang Thanks for taking over the PR!

[niboshi]

Thanks!

[niboshi]

Jenkins, test this please

[pfn-ci-bot]

Successfully created a job for commit a3f1703:

• Dashboard for commit a3f1703

[chainer-ci]

Jenkins CI test (for commit a3f1703, target branch master) succeeded!

[leofang]

Thanks to @sjperkins for laying the groundwork and to @niboshi for your help and patience!