Compiler error when building simple .cu file with thrust, -G, and --expt-relaxed-constexpr

We recently upgraded from cuda_10.0.130 to cuda_10.1.168, but a few of our .cu files were getting the following error when being build in debug only:

CrashCompiler.cu
ptxas C:/Users/xxx/AppData/Local/Temp/tmpxft_00008658_00000000-10_CrashCompiler.compute_60.ptx, line 7553; error   : Alignment of argument does not match formal parameter '_'
ptxas fatal   : Ptx assembly aborted due to errors

The filename has been renamed from the original to “CrashCompiler.cu”. After a LOT of debugging I was able to reduce the original file (which included hundreds of headers and had thousands of lines of code) to this one simple repro:

#include <thrust/device_vector.h>
#include <algorithm>

void someFunctionThatCallsStdMin() {
  std::min(16, 32);
}

void someOtherFuncThatCreatesADeviceVector()
{
  thrust::device_vector<float4> foo(32);
}

Note that the BOTH of the offending lines (std::min and thrust::device_vector) must be in the file. They can be in the same function, but are split here to show that that isn’t required.

This only occurs when both -G and --expt-relaxed-constexpr are passed as flags. An entire ps script that can be used to reproduce the error is pasted below. --expt-relaxed-constexpr isn’t needed in this particular example, but IS required in other code in the actual file.

Is this known? Is there a fix?

Example powershell script that causes the issue:

third-party/cuda/cuda_10.1.168/x64-windows/bin/nvcc.exe `
 -ccbin `
 'C:\open\bravo\third-party\toolchains\vs2017_15.9\BuildTools\VC\Tools\MSVC\14.16.27023\bin\Hostx64\x64\cl.exe' `
 -'gencode=arch=compute_60,code=sm_60' `
 -'gencode=arch=compute_75,code=sm_75' `
 -G `
 '-Ithird-party/toolchains/vm' `
 '-Ithird-party/toolchains/windows10sdk/10.0.17763.0/Include/10.0.17763.0/shared' `
 '-Ithird-party/toolchains/windows10sdk/10.0.17763.0/Include/10.0.17763.0/ucrt' `
 '-Ithird-party/toolchains/windows10sdk/10.0.17763.0/Include/10.0.17763.0/um' `
 '-Ithird-party/toolchains/windows10sdk/10.0.17763.0/Include/10.0.17763.0/winrt' `
 '-Ithird-party/toolchains/windows10sdk/10.0.17763.0/Include/10.0.17763.0/cppwinrt' `
 '-Ithird-party/toolchains/windows10sdk/10.0.18362.0/Include/10.0.18362.0/km' `
 '-Ithird-party/toolchains/windows10sdk/10.0.18362.0/Include/10.0.18362.0/shared' `
 '-Ithird-party/toolchains/windows10sdk/10.0.18362.0/Include/10.0.18362.0/ucrt' `
 '-Ithird-party/toolchains/windows10sdk/10.0.18362.0/Include/10.0.18362.0/um' `
 '-Ithird-party/toolchains/windows10sdk/10.0.18362.0/Include/10.0.18362.0/winrt' `
 '-Ithird-party/toolchains/windows10sdk/10.0.18362.0/Include/10.0.18362.0/cppwinrt' `
 --expt-relaxed-constexpr `
 '-Ithird-party\cuda\cuda_10.1.168\x64-windows\include' `
 -Xcompiler `
 /MTd `
 -Xcompiler `
 /Od `
 -Xcompiler `
 /bigobj `
 -Xcompiler `
 /EHs `
 -Xcompiler `
 /GF `
 -Xcompiler `
 /Gw `
 -Xcompiler `
 /Gy `
 -Xcompiler `
 /Oi `
 -Xcompiler `
 /Z7 `
 -Xcompiler `
 /FC `
 -Xcompiler `
 /Zc:throwingNew `
 -Xcompiler `
 /Brepro `
 -Xcompiler `
 /d2threads1 `
 'CrashCompiler.cu'

I’m able to reproduce the issue on CUDA 10.1.243 also.

My suggestion would be to test the latest CUDA 10.2.89. If the problem still occurs there, file a bug according to the instructions linked to a sticky post at the top of this forum.

As time permits, I will probably try to do the same thing, but for bugs that I file internally, I won’t be able to provide any sort of regular status updates or anything like that. If you’d like bug-oriented communication, I suggest filing your own bug.

Thanks Robert,

I submitted a bug: https://developer.nvidia.com/nvidia_bug/2807504

I’ll try upgrading to 10.2 tomorrow however I don’t think that is a very easy solution for us.

Thanks,
Joe

I don’t think 10.2 will fix the issue.

The development team has had a look and confirmed they see the issue.

I expect a fix to come in a future CUDA release. At this time I can’t give any specifics regarding which release, schedule, etc.

I also have no information on any suggested workarounds, other than those that seem obvious (e.g. don’t compile with -G). I acknowledge that such obvious workarounds are probably not actually useful.

As best I can tell, the issue does not occur on linux.