cuda undefined reference to constant array

Question

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.

@rps corrected it, was a typo in my post (corrected it) but it's correct in my code. — soriak, Feb 02 '13 at 17:35
@talonmies thanks for the tip I tried to not declare it extern, but then I get a problem with multiple definitions. I don't understand why because I only define it in header file and want to access it from a .cu file which contains the host-based initialization and an other one which is the one above to access the values in my kernel. Any idea how to achieve what I need? Thanks for the help so far. — soriak, Feb 02 '13 at 17:43
You must *define* it somewhere, exactly once. When you had it declared `extern` it wasn't being defined anywhere. Without it being `extern` you are defining it everywhere. — talonmies, Feb 03 '13 at 10:05
@talonmies I removed the `extern` I figured out why this is wrong. I get a multiple definition error for each file which includes the .h-file. I tried to add `namespace Constants{ extern __device__ __constant__ Type1 array[10] }` at the beginning a file including the .h. I use the other two namespaces in all device files so this shouldn't be the problem. Unfortunately this doesn't fix it, neither does it result in a new error. — soriak, Feb 03 '13 at 12:29
Are you sure you understand what the difference between a declaration and a definition is in C++? A \_\_constant\_\_ variable must be (statically) *defined* once somewhere in your code, just like a subroutine or a class method. — talonmies, Feb 03 '13 at 12:36
@talonmies turned out I didn't understand the difference for variables. But after finding this problem I read some [info](http://www.cprogramming.com/declare_vs_define.html) on this and thought that my solution in the previous comment should work. As it turns out it doesn't but I don't understand why because I define it in the .h-file and declare it using `extern` in the file I use it. **EDIT:** I think that this wrong, I shouldn't need any extern if I include the.h file right? — soriak, Feb 03 '13 at 12:52
let us [continue this discussion in chat](http://chat.stackoverflow.com/rooms/23855/discussion-between-soriak-and-talonmies) — soriak, Feb 03 '13 at 12:55

score 1 · Answer 1 · edited May 23 '17 at 11:56

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];
      }
    }
  }
}

cuda undefined reference to __constant__ array

1 Answers1

cuda undefined reference to constant array