33 lines · plain
1// Make sure that __global__ functions are emitted along with correct2// annotations and are added to @llvm.used to prevent their elimination.3// REQUIRES: nvptx-registered-target4//5// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s6 7#include "Inputs/cuda.h"8 9// CHECK-LABEL: define{{.*}} void @device_function10extern "C"11__device__ void device_function() {}12 13// CHECK-LABEL: define{{.*}} ptx_kernel void @global_function14extern "C"15__global__ void global_function() {16 // CHECK: call void @device_function17 device_function();18}19 20// Make sure host-instantiated kernels are preserved on device side.21template <typename T> __global__ void templated_kernel(T param) {}22// CHECK-DAG: define{{.*}} ptx_kernel void @_Z16templated_kernelIiEvT_(23 24namespace {25__global__ void anonymous_ns_kernel() {}26// CHECK-DAG: define{{.*}} void @_ZN12_GLOBAL__N_119anonymous_ns_kernelEv(27}28 29void host_function() {30 templated_kernel<<<0, 0>>>(0);31 anonymous_ns_kernel<<<0,0>>>();32}33