# Cuda constant memory

**URL:** <https://forums.developer.nvidia.com/t/cuda-constant-memory/265584>\
**Category:** CUDA Programming and Performance\
**Tags:** cuda, kernel\
**Created:** [September 6, 2023, 2:45pm UTC](https://forums.developer.nvidia.com/t/cuda-constant-memory/265584 "2023-09-06T14:45:05Z")\
**Posts on this page:** 6\
**Page:** 1

<div class="post-metadata">

**Author:** ![cjl4dream](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@cjl4dream](https://forums.developer.nvidia.com/u/cjl4dream)\
**Post date:** [September 6, 2023, 2:45pm UTC](https://forums.developer.nvidia.com/t/cuda-constant-memory/265584/1 "2023-09-06T14:45:05Z")

</div>

I want to use constant memory, but I found a problem.  
If I define a constant symbol in the cu file and want to initialize it in cpp, I have to pass its size and ptr through the function.  
such as:

```auto
//kernel.cu
#include <cuda_runtime.h>
#include <iostream>
__constant__ int const_data;

extern size_t get_size() {
    size_t size;
    cudaGetSymbolSize(&size, const_data);
    return size;
}

extern void* get_pointer() {
    void* ptr;
    cudaGetSymbolAddress(&ptr, const_data);
    std::cout << "cu ptr:" << ptr << std::endl;
    return ptr;
}

__global__ void kernel_print_const() {
    printf("value = %d\n", const_data);
    return;
}

cudaError_t print_const() {
    kernel_print_const<<<1,1>>>();
    return cudaDeviceSynchronize();
}

```

and my main.cpp:

```auto
//main.cpp
#include <cuda_runtime.h>
#include <iostream>
extern cudaError_t print_const();
extern void* get_pointer();
extern size_t get_size();
int main() {
    std::cout << get_size() << std::endl;
    void* d_const_data = get_pointer();
    std::cout << "main ptr: " << d_const_data << std::endl;

    int h_const_data = 10;
    cudaMemcpy(d_const_data, &h_const_data, sizeof(int), cudaMemcpyHostToDevice);
    print_const();
    return 0;
}

```

how can I use cudaGetSymbolSize() and cudaMemcpyToSymbol() in my main.cpp?

---

<div class="post-metadata">

**Author:** ![cjl4dream](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@cjl4dream](https://forums.developer.nvidia.com/u/cjl4dream)\
**Post date:** [September 6, 2023, 2:55pm UTC](https://forums.developer.nvidia.com/t/cuda-constant-memory/265584/2 "2023-09-06T14:55:42Z")

</div>

By the way, if I define lots of constant symbol in the kernel.cu

```auto
//kernel.cu
__constant__ int a1;
__constant__ int a2;
__constant__ int a3;
__constant__ int a4;
...

```

Does I have to define lots of get\_point() and get\_size()?  
How can I use function parameters to define a function to solve this?  
such as :

```auto
template<class T>
void* get_size(T symbol) {
    void* ptr;
    cudaGetSymbolAddress(&ptr, symbol);
    std::cout << "cu ptr:" << ptr << std::endl;
    return ptr;
}

```

---

<div class="post-metadata">

**Author:** ![Robert\_Crovella](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/robert_crovella/32/14043_2.png) [@Robert\_Crovella](https://forums.developer.nvidia.com/u/Robert_Crovella)\
**Post date:** [September 6, 2023, 3:07pm UTC](https://forums.developer.nvidia.com/t/cuda-constant-memory/265584/3 "2023-09-06T15:07:17Z")

</div>

> [@cjl4dream](#):
>
> Does I have to define lots of get\_point() and get\_size()?

Yes. The ` __constant__ ` syntax is not something that can be processed by the host compiler. When you have a file in your project that is named filename **.cpp** , that file will be processed by the host compiler, by default. The host compiler will not understand ` __constant__ ` syntax.

If you want to switch to using a `main.cu` file, you can use the `extern` declaration in order to access the constant symbol in another compilation unit. For this you will need to compile with relocatable device code with device linking. There are [examples](https://stackoverflow.com/questions/15468997/cuda-5-0-generate-relocatable-device-code-leads-to-invalid-device-symbol-error/15480144#15480144) like this already published.

I haven’t tried using a constant symbol with templating. At first glance, I don’t think it will work. (A simple test suggests to me it will not work: `cudaErrorInvalidSymbol`).

---

<div class="post-metadata">

**Author:** ![njuffa](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@njuffa](https://forums.developer.nvidia.com/u/njuffa)\
**Post date:** [September 6, 2023, 7:28pm UTC](https://forums.developer.nvidia.com/t/cuda-constant-memory/265584/4 "2023-09-06T19:28:25Z")

</div>

It is not clear to me what exactly you are trying to achieve, but an alternative to consider may be passing the desired data as a kernel argument, which is placed in constant memory. With recent hardware and software, CUDA is able to pass a good-sized chunk of data this way. This avoids separate `cudaMemCpyFromSymbol()` calls.

> **[CUDA 12.1 Supports Large Kernel Parameters | NVIDIA Technical Blog](https://developer.nvidia.com/blog/cuda-12-1-supports-large-kernel-parameters/)**
>
> CUDA 12.1 offers you the option of passing up to 32,764 bytes using kernel parameters, which can be exploited to simplify applications as well as gain performance improvements.

> CUDA kernel function parameters are passed to the device through constant memory and have been limited to 4,096 bytes. CUDA 12.1 increases this parameter limit from 4,096 bytes to 32,764 bytes on all device architectures including [NVIDIA Volta](https://www.nvidia.com/en-us/data-center/volta-gpu-architecture/) and above.

---

<div class="post-metadata">

**Author:** ![cjl4dream](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@cjl4dream](https://forums.developer.nvidia.com/u/cjl4dream)\
**Post date:** [September 7, 2023, 5:02am UTC](https://forums.developer.nvidia.com/t/cuda-constant-memory/265584/6 "2023-09-07T05:02:32Z")

</div>

In fact, I want to encapsulate a cpp class for Cuda’s constant memory to implement the functions of upload, download, and memset. But I have to export its device pointer in the Cu file that defines the constant memory. If there are too many **constant** variables, this will require to be implemented get\_pointer() for each variable.

---

<div class="post-metadata">

**Author:** ![njuffa](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@njuffa](https://forums.developer.nvidia.com/u/njuffa)\
**Post date:** [September 7, 2023, 5:17am UTC](https://forums.developer.nvidia.com/t/cuda-constant-memory/265584/7 "2023-09-07T05:17:40Z")

</div>

> [@cjl4dream](#):
>
> I want to encapsulate a cpp class for Cuda’s constant memory to implement the functions of upload, download, and memset

I do not know your use case, but off-hand I fail to see the benefits of such an arrangement.

Note that there is not _the_ constant memory, but multiple banks of constant memory (the details differ by GPU architecture), one of which is used to back ` __constant__ `.

> [@cjl4dream](#):
>
> If there are too many **constant** variables,

You could always aggregate multiple variables into a single data object, a `struct` in the simplest case, to reduce the the number of distinct data objects in constant memory.