33 lines · plain
1// RUN: %clang_cc1 -triple nvptx -fcuda-is-device -emit-llvm -o - %s \2// RUN: | FileCheck -check-prefix=NORDC %s3// RUN: %clang_cc1 -triple nvptx -fcuda-is-device -emit-llvm -o - %s \4// RUN: | FileCheck -check-prefix=NORDC-NEG %s5// RUN: %clang_cc1 -triple nvptx -fcuda-is-device -fgpu-rdc -emit-llvm -o - %s \6// RUN: | FileCheck -check-prefix=RDC %s7// RUN: %clang_cc1 -triple nvptx -fcuda-is-device -fgpu-rdc -emit-llvm -o - %s \8// RUN: | FileCheck -check-prefix=RDC-NEG %s9 10#include "Inputs/cuda.h"11 12template <typename T> __device__ void func() {}13template <typename T> __global__ void kernel() {}14 15template __device__ void func<int>();16// NORDC: define internal void @_Z4funcIiEvv()17// RDC: define weak_odr void @_Z4funcIiEvv()18 19template __global__ void kernel<int>();20// NORDC: define ptx_kernel void @_Z6kernelIiEvv()21// RDC: define weak_odr ptx_kernel void @_Z6kernelIiEvv()22 23// Ensure that unused static device function is eliminated24static __device__ void static_func() {}25// NORDC-NEG-NOT: define{{.*}} void @_ZL13static_funcv()26// RDC-NEG-NOT: define{{.*}} void @_ZL13static_funcv[[FILEID:.*]]()27 28// Ensure that kernel function has external or weak_odr29// linkage regardless static specifier30static __global__ void static_kernel() {}31// NORDC: define ptx_kernel void @_ZL13static_kernelv()32// RDC: define weak_odr ptx_kernel void @_ZL13static_kernelv[[FILEID:.*]]()33