Compilation error in arch sm_20 but compiles well with -arch sm_13

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi all,

my app compiles well with -arch sm_12 and sm_13 but compilation fails with sm_20.

How can I be able to know what are the coding differences between these architectures, and how can I understand the error message?

1>### Assertion failure at line 2129 of ../../be/cg/NVISA/cgtarget.cxx:

1>### Compiler Error in file ....../AppData/Local/Temp/tmpxft_000007e4_00000000-15_rayTracer.compute_20.cpp3.i during Code_Expansion phase:

1>### asm m constraint must have simple variable

1>nvopencc ERROR: C:\Progra~2\NVIDIA~1\CUDA\v3.2\bin/../open64/lib//be.exe returned non-zero status 1

Thanks,

–pium

edit :

I found that on the thrust forum http://groups.google.com/group/thrust-users/browse_thread/thread/92e89d033c881127

And indeed, I am using texture as argument but not to a kernel, but as argument to a device function called in the kernel. Have I right to do so?

2

This kind of error message indicates an internal compiler error. The error information is designed for consumption by a compiler engineer and states the location inside the compiler at which the problem was deteced, i.e. not user actionable.

It would be helpful if you could file a bug for this. Thanks!

3

Arf…

My code is quite enourmous but I guess it is easy to make a simple example that have the same reaction.

As said in the edit part of my previous post, the problem came from the fact I passed a texture as a device function parameter.

I don’t know if I was taking advantage of a previous bug because it compiles and works in sm_13. Is it a clean way to code ? Because it is quite useful, not to have to copy code that do the same thing on several textures.

An important point is in the new implementation that copies the code for every texture is a bit faster.

template <class TexType>

__device__

myFunction( const TexType& tex )

{

...

}

__global__

void myKernel( )

{

    myFunction( myTexture );

}

I hope this can help you.

– pium

4

Compiler team indicates that they should be able to debug this from the compute_20.cpp3.i file that is causing trouble. If you are a registered developer, it would be helpful if you could file a bug for this, attaching the compute_20.cpp3.i file in question. If you are not yet a registered developer, I would encourage you to become one :-)

As a final option, you could send me the file as an attachment via a personal message through the forums and I will follow up with the compiler team. The benefit of filing a bug as a registered developer is being able to check on bug status. Internally filed bugs are not publicly visible.

5

done!

6

Thank you for you help. As far as I know, textures cannot be passed as arguments to device functions, but such usage should not give rise to an internal compiler error.