CODEWITHSUNDEEP
c++cuda
user2924261
user2924261

Reputation: 127

CUDA 5.5 cudaMemcpyToSymbol, __constant__ and out of scope error

I'm trying to compile a CUDA example which has;

cuda.cu:

__constant__ unsigned VERTICES;
__constant__ unsigned TRIANGLES;

and the corresponding code in main.cpp;

cudaMemcpyToSymbol(VERTICES, &verticesNo, sizeof(int));
cudaMemcpyToSymbol(TRIANGLES, &trianglesNo, sizeof(int));

How to avoid "VERTICES not declared in this scope" error when compiling the main.cpp?

TIA.

cheers,

Upvotes: 1

Views: 1319

Answers (2)

opetrenko
opetrenko

Reputation: 346

CUDA is defined by the nvcc compiler which itself is an extension of C99. It sounds like what you really want to do is separate out CUDA so you have a generic header file. You could then use it from C or C++. I prefer to work with data in C++ personally, and as such I have found best way to do this is the following files and include paths:

               WRAPPER_HEADER.h          CUDA_HEADER.cuh
                 /          \               /       \
                /            \             /         \
            #include      #include    #include     #include
              /                \         /             \
             /                  \       /               \
            /                    \     /                 \
  CALL.cpp/CALL.c            CUDA_WRAPPER.cu         CUDA_KERNEL.cu

CALL.c/CALL.cpp is C/C++ containing whatever you want that will call the wrapper function

CUDA_WRAPPER.cu is the wrapper function which:

  • uses cudaMalloc/cudaMemcpy/cudaMemcpyToSymbol/cudaFree to create/free-up device memory and to manage I/O from the original calling function called from C or C++.
  • calls the kernel itself in the kernel<<>>(...) format

WRAPPER_HEADER.h contains a C version of:

  • declaration for the wrapper function (which must be written solely in C99)
  • constants in their generic C form
  • input that is different for each thread or for each block
  • pointers for where the results of the calculations end up
  • number of threads per block
  • number of blocks

CUDA_HEADER.cuh contains:

  • declaration of __constant__ memory that the wrapper can write to via cudaMemcpyToSymbol(...)
  • declaration for the kernel function with the __global__ specifier

CUDA_KERNEL.cu contains:

  • the implementation of the kernel __global__ void kernel(...) function
  • declaration and implementation of device functions declared with __device__ specifier
  • declaration of __shared__ memory (it only has lifetime of a block so cannot be called from a wrapper based on what I can tell... feel free to correct this anyone)

There is some of this demonstrated in the CUDA literature and I like it because it really separates out CUDA C as the specialized language it is. It is only necessary when you are dealing with setting up and running the kernels.

Upvotes: -1

hubs
hubs

Reputation: 1809

CUDA __constant__ variables have a file scope linkage. That means that the cudaMemcpyToSymbol have to be in the same .cu file where the __constant__ variable is defined.

You can add a wrapper function to the .cu file and call this one from your .cpp file.

sample for cuda.cu:

__constant__ unsigned VERTICES;
__constant__ unsigned TRIANGLES;

void wrapper_fn(unsigned *verticesNo, unsigned *trianglesNo)
{
  cudaMemcpyToSymbol(VERTICES, verticesNo, sizeof(unsigned));
  cudaMemcpyToSymbol(TRIANGLES, trianglesNo, sizeof(unsigned));
}

Then only call wrapper_fn in your main.cpp.

Upvotes: 4

Related Questions