llvm/clang/test/CodeGenCUDA/device-fun-linkage.cu

// RUN: %clang_cc1 -triple nvptx -fcuda-is-device -emit-llvm -o - %s \
// RUN:   | FileCheck -check-prefix=NORDC %s
// RUN: %clang_cc1 -triple nvptx -fcuda-is-device -emit-llvm -o - %s \
// RUN:   | FileCheck -check-prefix=NORDC-NEG %s
// RUN: %clang_cc1 -triple nvptx -fcuda-is-device -fgpu-rdc -emit-llvm -o - %s \
// RUN:   | FileCheck -check-prefix=RDC %s
// RUN: %clang_cc1 -triple nvptx -fcuda-is-device -fgpu-rdc -emit-llvm -o - %s \
// RUN:   | FileCheck -check-prefix=RDC-NEG %s

#include "Inputs/cuda.h"

template <typename T> __device__ void func() {}
template <typename T> __global__ void kernel() {}

template __device__ void func<int>();
// NORDC:     define internal void @_Z4funcIiEvv()
// RDC:       define weak_odr void @_Z4funcIiEvv()

template __global__ void kernel<int>();
// NORDC:     define void @_Z6kernelIiEvv()
// RDC:       define weak_odr void @_Z6kernelIiEvv()

// Ensure that unused static device function is eliminated
static __device__ void static_func() {}
// NORDC-NEG-NOT: define{{.*}} void @_ZL13static_funcv()
// RDC-NEG-NOT:   define{{.*}} void @_ZL13static_funcv[[FILEID:.*]]()

// Ensure that kernel function has external or weak_odr
// linkage regardless static specifier
static __global__ void static_kernel() {}
// NORDC:     define void @_ZL13static_kernelv()
// RDC:       define weak_odr void @_ZL13static_kernelv[[FILEID:.*]]()