CODEWITHSUNDEEP
c++cuda
emher
emher

Reputation: 6024

CUDA / C++ : error C3861: 'uint2float': identifier not found

I am experiencing this issue in Visual Studio 2012 using CUDA 5.5 with in the file "FLOAT_UTIL_DEVICE.HCU", which is provided by Nvidia. According to similar issues, it might be a matter of code structure, but i can't see any problems:

#include <vector_functions.h>
#include <device_functions.h>

//// ADDED BY ME FOR TEST PURPOSES
//inline __device__ float uint2float( unsigned int a )
//{
//  return (float) a;
//}
//// END

inline __device__ float2 uintd_to_floatd( uint2 a )
{
    return make_float2( uint2float(a.x), uint2float(a.y) );
}

inline __device__ float3 uintd_to_floatd( uint3 a )
{
    return make_float3( uint2float(a.x), uint2float(a.y), uint2float(a.z) );
}

inline __device__ float4 uintd_to_floatd( uint4 a )
{
    return make_float4( uint2float(a.x), uint2float(a.y), uint2float(a.z), uint2float(a.w) );
}

The methods in question should be defined in "device_functions.h" according to Nvidia documentation. If i uncomment the test code (which defines the function which is missing), i get a new error: "more than one instance of the overloaded function "uint2float" matches the argument list"; hence somehow it IS defined already. What am i missing?

Upvotes: 0

Views: 1312

Answers (1)

talonmies
talonmies

Reputation: 72349

I can't reproduce this problem in CUDA 5.0 using gcc. If I take a complete repro case using your device functions:

#include <vector_functions.h>
#include <device_functions.h>

inline __device__ float2 uintd_to_floatd( uint2 a )
{
    return make_float2( uint2float(a.x), uint2float(a.y) );
}

inline __device__ float3 uintd_to_floatd( uint3 a )
{
    return make_float3( uint2float(a.x), uint2float(a.y), 
                uint2float(a.z) );

}

inline __device__ float4 uintd_to_floatd( uint4 a )
{
    return make_float4( uint2float(a.x), uint2float(a.y), 
                uint2float(a.z), uint2float(a.w) );
}


template<typename Tin, typename Tout>
__global__
void kernel(Tin *in, Tout *out) {
    out[threadIdx.x] = uintd_to_floatd(in[threadIdx.x]);
}

template __global__ void kernel<uint2,float2>(uint2 *, float2 *);
template __global__ void kernel<uint3,float3>(uint3 *, float3 *);
template __global__ void kernel<uint4,float4>(uint4 *, float4 *);

and compile it:

$ nvcc -c -arch=sm_20 -Xptxas="-v" uint2float.cu 
ptxas info    : 0 bytes gmem
ptxas info    : Compiling entry function '_Z6kernelI5uint36float3EvPT_PT0_' for 'sm_20'
ptxas info    : Function properties for _Z6kernelI5uint36float3EvPT_PT0_
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 7 registers, 40 bytes cmem[0]
ptxas info    : Compiling entry function '_Z6kernelI5uint46float4EvPT_PT0_' for 'sm_20'
ptxas info    : Function properties for _Z6kernelI5uint46float4EvPT_PT0_
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 8 registers, 40 bytes cmem[0]
ptxas info    : Compiling entry function '_Z6kernelI5uint26float2EvPT_PT0_' for 'sm_20'
ptxas info    : Function properties for _Z6kernelI5uint26float2EvPT_PT0_
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 6 registers, 40 bytes cmem[0]

it builds without any compile errors. This either means that there is another error in some code that you haven't shown us, or this is a problem specific to Visual Studio or the MS C++ compiler. Thrust code using certain vector types is known to break when compiled with the VS toolchain. It might be that you are seeing a symptom of the same issue. If you are desperate for a short term fix, you might try defining your own version of the vector types and re-writing your __device__ functions to work with those types instead.

Upvotes: 1

Related Questions