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'