CODEWITHSUNDEEP
linkercudaundefined-reference
soriak
soriak

Reputation: 672

cuda undefined reference to __constant__ array

i got the following problem using nvcc. I use separate compilation (finally got it running with cmake) and got a problem with the declaration of a __device__ __constant__ (extern) type1<type2> var[length] array.

Here is my header:

#include <type.h>
#include <type2.h>

#ifndef GUARD_H_
#define GUARD_H_

namespace NameSpace1{
  namespace NameSpace2{

    namespace Constants{
      __device__ __constant__ extern Type1<Type2> array[10];
      __device__ Type1<Type2> accessConstantX(size_t index);
    }
  }
}
#endif

And her my .cu-file:

#include <header.h>
#include <assert.h>

namespace NameSpace1{
  namespace NameSpace2{

    namespace Constants{
      __device__ Type1<Type2> accessConstantX(size_t index){
    assert(index <= 9);
    return array[index];
      }
    }
  }
}

I get the following error from intermediate separate compilation step:

nvlink error   : Undefined reference to '_ZN12NameSpace115NameSpace29Constants17arrayE'

This results from the access in the .cu-file. Thanks for suggestions.

Upvotes: 0

Views: 1098

Answers (1)

soriak
soriak

Reputation: 672

I found out what my mistake was, be reading this post. It turned out that I didn't understand how to correctly declare an global variable in a header file. My problem had nothing to do with nvcc. Thanks @talonmies for your help it pointed my in the direction to look for the answer.

Here the solution to my problem:

Here is my header:

#include <type.h>
#include <type2.h>

#ifndef GUARD_H_
#define GUARD_H_

namespace NameSpace1{
  namespace NameSpace2{

    namespace Constants{
      extern __device__ __constant__ extern Type1<Type2> array[10];
      __device__ Type1<Type2> accessConstantX(size_t index);
    }
  }
}
#endif

And her my .cu-file:

#include <header.h>
#include <assert.h>

namespace NameSpace1{
  namespace NameSpace2{

    namespace Constants{
      __device__ __constant__ Type1<Type2> array[10];
      __device__ Type1<Type2> accessConstantX(size_t index){
        assert(index <= 9);
        return array[index];
      }
    }
  }
}

Upvotes: 1

Related Questions