bug in generating ptx

Hello every CUDAers :D

I dont know if this is a well-known bug or not, so I just post it here if anyone of you ran into it. :ph34r:

I am using Linux 64bit nvcc on Ubuntu (CUDA 1.1); not sure if it’s the same case on Windows but you can still try so. Basically the following code snippet cannot be compiled due to the total variable length in ptx files filled up the ptxas buffer and cause ptxas terminated:

typedef struct

{

   int id; // add whatever field you like

} Type1;

typedef struct

{

   int id; // add whatever field you like

} Type2;

typedef struct

{

   int id; // add whatever field you like

} Type3;

typedef struct

{

   int id; // add whatever field you like

} Type4;

__global__ static void WHATEVER_LONG_FUNCTION_NAME_XXXXXXXXXXXXXXXXXXXXXXXXXXXXXXXXXX(

Type1 *pType1_WITH_WHATEVER_LONG_VARIABLE_NAME_AAAA, 

Type2 *pType2_SOME_OTHER_LONG_VARIABLE_NAME_BBBB,

Type3 *pType3_SOME_OTHER_LONG_LONG_VARIABLE_NAME_CCCC,

Type4 *pType4_SOME_OTHER_LONG_LONG_LONG_VARIABLE_NAME_DDDD

//you can add more types and variables here if you want to

)

{

 // do nothing here

}

then you will get:

This is pretty depressing because it just cannot compile!! So I changed the declaration in several ways and finally, after digging into the generated ptx file, I found that those variable names are using all types as identifier in their parameter declaration like this:

   .entry 

_Z4<function_name>P3<type_name_1>P3<type_name_2>P3<type_name_3>P3<type_name­_4> 

    { 

    .param .u64 __cudaparm__Z4<function_name>P3<type_name_1>P3<type_name_2>P3<type_name_3>P­3<type_name_4>_<var_name_1>; 

    .param .u64 __cudaparm__Z4<function_name>P3<type_name_1>P3<type_name_2>P3<type_name_3>P­3<type_name_4>_<var_name_2>; 

    .param .u64 __cudaparm__Z4<function_name>P3<type_name_1>P3<type_name_2>P3<type_name_3>P­3<type_name_4>_<var_name_3>; 

    .param .u64 __cudaparm__Z4<function_name>P3<type_name_1>P3<type_name_2>P3<type_name_3>P­3<type_name_4>_<var_name_4>; 

    ... 

   }  

well, i know it makes sense here to use all types as a part of the function declaration because of function overloading…but is that necessary for parameters? ptx parameters are all declared in global scope? :unsure:

and I guess somehow ptxas has a small buffer that just can not hold those long variable names and finally got buffer overflow and screw up the runtime stack.

So the trick here would be making some dummy macro to re-define those type names and function names to make it shorter. It works for me!

Hope this can be fixed upon next release. Thanks! :thumbup:

This issue was also reported by others and will be fixed in the next release.

Thanks!

Mark

Would you explain,how should i define such macros?

Sorry,i am new in this. :">

I’ve got it now.It’s works! :)

Thanks,anyway.

You just saved the day + the weekend. Thanks!