# NVCC Compile Shared Library

**URL:** <https://forums.developer.nvidia.com/t/nvcc-compile-shared-library/179393>\
**Category:** CUDA Programming and Performance\
**Created:** [May 31, 2021, 4:56am UTC](https://forums.developer.nvidia.com/t/nvcc-compile-shared-library/179393 "2021-05-31T04:56:40Z")\
**Posts on this page:** 11\
**Page:** 1

<div class="post-metadata">

**Author:** ![Daniel\_Wong](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/daniel_wong/32/60786_2.png) [@Daniel\_Wong](https://forums.developer.nvidia.com/u/Daniel_Wong)\
**Post date:** [May 31, 2021, 4:56am UTC](https://forums.developer.nvidia.com/t/nvcc-compile-shared-library/179393/1 "2021-05-31T04:56:40Z")

</div>

I want to use NVCC to compile the shared library first then link it with main program.

```auto
nvcc --compiler-options '-fPIC' -o liblibrary.so --shared library.cu
nvcc -Xcompiler -I./ main.cu -L. -llibrary
nvcc -I./ main.cu -L. -llibrary
export LD_LIBRARY_PATH=$PWD:$LD_LIBRARY_PATH 
./a.out

```

There are several files.   
  
`library.cu`

```auto
#ifndef LIB1_H_INCLUDED
#define LIB1_H_INCLUDED

void print_value ( int x );
__global__ void cuda_hello();

#endif /* LIB1_H_INCLUDED */

```

`library.cuh`

```auto
#include "library.cuh"
#include <stdio.h>

void print_value( int x )
{
    printf("%d\n", x);
}

__global__ void cuda_hello(){
    printf("Hello World from GPU!\n");
}

```

`main.cu`

```auto
#include <stdio.h>
#include "library.cuh"

int main ( void )
{
    print_value(10);
    cuda_hello<<<1,1>>>(); 
    return 0;
}

```

However, **it will not print any output, is there anything wrong?** while `print_value` can print the output correctly.

---

<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:** [June 1, 2021, 9:28pm UTC](https://forums.developer.nvidia.com/t/nvcc-compile-shared-library/179393/2 "2021-06-01T21:28:39Z")

</div>

after:

> [@Daniel\_Wong](#):
>
> ` cuda_hello<<<1,1>>>();`

you should have:

```
cudaDeviceSynchronize();

```

Please also use proper CUDA error checking (google that, take the first hit, modify your code).

---

<div class="post-metadata">

**Author:** ![Daniel\_Wong](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/daniel_wong/32/60786_2.png) [@Daniel\_Wong](https://forums.developer.nvidia.com/u/Daniel_Wong)\
**Post date:** [June 2, 2021, 12:29am UTC](https://forums.developer.nvidia.com/t/nvcc-compile-shared-library/179393/3 "2021-06-02T00:29:07Z")

</div>

Hi, @Robert_Crovella  
This is the code I tried, but still no output for the error message.

```auto
int main ( void )
{
    print_value(10);
    cuda_hello<<<1,1>>>(); 
    gpuErrchk( cudaDeviceSynchronize());
    gpuErrchk( cudaGetLastError());
    return 0;
}

```

where `gpuErrchk` is defined as:

```auto
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__ , __LINE__ ); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

```

---

<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:** [June 2, 2021, 2:35am UTC](https://forums.developer.nvidia.com/t/nvcc-compile-shared-library/179393/4 "2021-06-02T02:35:43Z")

</div>

I recommend that you keep CUDA syntax out of the shared library interface.

```
$ cat library.cuh
#ifndef LIB1_H_INCLUDED
#define LIB1_H_INCLUDED

void print_value ( int x );
void cuda_hello();

#endif /* LIB1_H_INCLUDED */
$ cat library.cu
#include "library.cuh"
#include <stdio.h>

void print_value( int x )
{
    printf("%d\n", x);
}

__global__ void cuda_hello_kernel(){
    printf("Hello World from GPU!\n");
}

void cuda_hello(){
     cuda_hello_kernel<<<1,1>>>();
     cudaDeviceSynchronize();
}

$ cat main.cu
#include <stdio.h>
#include "library.cuh"

int main ( void )
{
    print_value(10);
    cuda_hello();
    return 0;
}
$ nvcc -Xcompiler -fPIC -o liblibrary.so --shared library.cu
$ nvcc -I. main.cu -L. -llibrary
$ LD_LIBRARY_PATH=. cuda-memcheck ./a.out
========= CUDA-MEMCHECK
10
Hello World from GPU!
========= ERROR SUMMARY: 0 errors
$

```

---

<div class="post-metadata">

**Author:** ![Daniel\_Wong](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/daniel_wong/32/60786_2.png) [@Daniel\_Wong](https://forums.developer.nvidia.com/u/Daniel_Wong)\
**Post date:** [June 2, 2021, 2:52am UTC](https://forums.developer.nvidia.com/t/nvcc-compile-shared-library/179393/5 "2021-06-02T02:52:36Z")

</div>

Hi @Robert_Crovella Thanks a lot!  
Just have one more follow up question.

> when should we use **`-rdc=true`** during the compilation?

---

<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:** [June 2, 2021, 1:20pm UTC](https://forums.developer.nvidia.com/t/nvcc-compile-shared-library/179393/6 "2021-06-02T13:20:10Z")

</div>

when device code in one compilation unit is referenceing or calling device code in another compilation unit. You have no examples of that here.

---

<div class="post-metadata">

**Author:** ![Daniel\_Wong](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/daniel_wong/32/60786_2.png) [@Daniel\_Wong](https://forums.developer.nvidia.com/u/Daniel_Wong)\
**Post date:** [June 2, 2021, 3:50pm UTC](https://forums.developer.nvidia.com/t/nvcc-compile-shared-library/179393/7 "2021-06-02T15:50:48Z")

</div>

Does this apply to the case where I have the definition of a ` __device__ ` function in one `.cu` file and compile it to `.so`, then link this `.so` with another `.cu` file in the later compilation?

---

<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:** [June 2, 2021, 3:59pm UTC](https://forums.developer.nvidia.com/t/nvcc-compile-shared-library/179393/8 "2021-06-02T15:59:45Z")

</div>

No, that won’t work, generally. It is a [stated limitation](https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#libraries) that you cannot device-link across a `.so` boundary. You can device link across a static library boundary.

You can device link within a .so, just not across the boundary/interface.

---

<div class="post-metadata">

**Author:** ![Daniel\_Wong](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/daniel_wong/32/60786_2.png) [@Daniel\_Wong](https://forums.developer.nvidia.com/u/Daniel_Wong)\
**Post date:** [June 2, 2021, 5:31pm UTC](https://forums.developer.nvidia.com/t/nvcc-compile-shared-library/179393/9 "2021-06-02T17:31:31Z")

</div>

Got it, Now, I understand why CUDA-based libraries such as `cuDNN, cuBLAS, NCCL`, only offer host CPU API instead of the ` __global__ ` function API.

However, I found there is one exception is for `NVSHMEM`, they also offer device-level API for a thread/warp/block that can directly be called from a ` __global__ ` kernel, which is quite different from those above CUDA libraries.

---

<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:** [June 2, 2021, 5:37pm UTC](https://forums.developer.nvidia.com/t/nvcc-compile-shared-library/179393/10 "2021-06-02T17:37:54Z")

</div>

nvshmem is implemented via a static library. As we’ve already discussed, you can do more-or-less anything with a static library.

---

<div class="post-metadata">

**Author:** ![TomNVIDIA](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/tomnvidia/32/14181_2.png) [@TomNVIDIA](https://forums.developer.nvidia.com/u/TomNVIDIA)\
**Post date:** [October 12, 2021, 4:04pm UTC](https://forums.developer.nvidia.com/t/nvcc-compile-shared-library/179393/11 "2021-10-12T16:04:10Z")

</div>

Closed on October 12, 2021, 4:04pm UTC