while compiling a CUDA kernel, I get the compiler error:
Assertion failure at line 590 of …/…/be/cg/NVISA/cgemit_targ.cxx:
Compiler Error in file /tmp/tmpxft_000038d2_00000000-29_traceKernel.cpp3.i during Assembly phase:
stack params not supported
nvopencc INTERNAL ERROR: /opt/nvidia/cuda/open64/lib//be returned non-zero status 1
This error is preceeded by about 20 advisories of the form:
/tmp/tmpxft_000038d2_00000000-29_traceKernel.cpp3.i(13): Advisory: Cannot tell what pointer points to, assuming global memory space
I’ve attached the source file, any help would be very appreciated. All of the advisories occur in the first function __sphere_intersect_kd at lines 13, 17, 21, 22, 25, 26, 28, 34, 37, 47, 48, 51, 52.
Information on exactly what the messages mean and how people have fixed these problems in the past would be very much appreciated, I haven’t found any information on the stack params compiler error at all.
Thanks in advance
You did not attach the source file. For some reason this forum doesn’t allow files of the type .c, .cpp and .cu as attachments to the posts. However the error message can easily be overlooked.
The “Cannot tel what pointer points to …” advisories are really just advisories and usually don’t cause problems. Well, as long as the compilers guess is correct and they actually point to global memory. However I have never seen them cause the compiler to crash.
I can reproduce this problem:
/*!
\file ValueCompare.h
\date Wednesday May 27, 2009
\author Gregory Diamos <gregory.diamos@gatech.edu>
\brief The header file for the ValueCompare class.
*/
#ifndef VALUE_COMPARE_H_INCLUDED
#define VALUE_COMPARE_H_INCLUDED
#include <functional>
namespace hydrazine
{
/*!
\brief A class for comparing key/value pairs based on the key only
*/
template< typename Compare, typename Container >
class ValueCompare :
public std::binary_function< typename Container::value_type,
typename Container::value_type, bool >,
public std::binary_function< typename Container::value_type,
typename Container::key_type, bool >,
public std::binary_function< typename Container::key_type,
typename Container::value_type, bool >
{
public:
typedef size_t size_type;
typedef typename Container::key_type key_type;
typedef typename Container::mapped_type mapped_type;
typedef typename Container::value_type value_type;
typedef ptrdiff_t difference_type;
typedef value_type* pointer;
typedef value_type& reference;
typedef const value_type* const_pointer;
typedef const value_type& const_reference;
protected:
Compare _compare;
public:
const Compare& compare() const
{
return _compare;
}
public:
explicit ValueCompare( const Compare& c ) : _compare( c ) {}
ValueCompare() {}
public:
bool operator()( const_reference x, const_reference y ) const
{
return _compare( x.first, y.first );
}
bool operator()( const key_type& x, const_reference y ) const
{
return _compare( x, y.first );
}
bool operator()( const_reference x, const key_type& y ) const
{
return _compare( x.first, y );
}
};
}
#endif
#include <thrust/sort.h>
#include <thrust/device_vector.h>
#include <hydrazine/interface/ValueCompare.h>
#include <datalog/interface/Map.h>
#include <map>
#define SIZE 169
typedef std::map< unsigned int, unsigned int > Map;
int main()
{
thrust::device_vector< std::pair< unsigned int, unsigned int > >
vector( SIZE );
thrust::sort( vector.begin(), vector.end(),
hydrazine::ValueCompare< std::less< unsigned int >, Map >() );
}
The compiler dies with the same message:
Assertion failure at line 590 of …/…/be/cg/NVISA/cgemit_targ.cxx:
Compiler Error in file /tmp/tmpxft_000038d2_00000000-29_traceKernel.cpp3.i during Assembly phase:
stack params not supported
This code works if I use std::less< std::pair< unsigned int, unsigned int > >() instead of hydrazine::ValueCompare< std::less< unsigned int >, Map >() …