83 lines · plain
1// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \2// RUN: -fgpu-rdc -std=c++11 -emit-llvm -o - -target-cpu gfx906 | FileCheck %s3 4// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \5// RUN: -fgpu-rdc -std=c++11 -emit-llvm -o - -target-cpu gfx906 \6// RUN: | FileCheck -check-prefix=NEG %s7 8// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \9// RUN: -std=c++11 -emit-llvm -o - -target-cpu gfx906 \10// RUN: | FileCheck -check-prefixes=NEG,NORDC %s11 12// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -x hip %s \13// RUN: -fgpu-rdc -std=c++11 -emit-llvm -o - \14// RUN: | FileCheck -check-prefix=HOST-NEG %s15 16 17#include "Inputs/cuda.h"18 19// CHECK-LABEL: @__clang_gpu_used_external = internal {{.*}}global20// CHECK-DAG: @_Z7kernel1v21// CHECK-DAG: @_Z7kernel4v22// CHECK-DAG: @var123// CHECK-LABEL: @llvm.compiler.used = {{.*}} @__clang_gpu_used_external24 25// NEG-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel2v26// NEG-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel3v27// NEG-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel5v28// NEG-NOT: @__clang_gpu_used_external = {{.*}} @var229// NEG-NOT: @__clang_gpu_used_external = {{.*}} @var330// NEG-NOT: @__clang_gpu_used_external = {{.*}} @ext_shvar31// NEG-NOT: @__clang_gpu_used_external = {{.*}} @shvar32// NORDC-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel1v33// NORDC-NOT: @__clang_gpu_used_external = {{.*}} @_Z7kernel4v34// NORDC-NOT: @__clang_gpu_used_external = {{.*}} @var135// HOST-NEG-NOT: call void @__hipRegisterVar({{.*}}, ptr @ext_shvar36// HOST-NEG-NOT: call void @__hipRegisterVar({{.*}}, ptr @shvar37__global__ void kernel1();38 39// kernel2 is not marked as used since it is a definition.40__global__ void kernel2() {}41 42// kernel3 is not marked as used since it is not called by host function.43__global__ void kernel3();44 45// kernel4 is marked as used even though it is not called.46__global__ void kernel4();47 48// kernel5 is not marked as used since it is called by host function49// with weak_odr linkage, which may be dropped by linker.50__global__ void kernel5();51 52extern __device__ int var1;53 54__device__ int var2;55 56extern __device__ int var3;57 58void use(int *p);59 60void test() {61 kernel1<<<1, 1>>>();62 void *p = (void*)kernel4;63 use(&var1);64}65 66__global__ void test_lambda_using_extern_shared() {67 extern __shared__ int ext_shvar[];68 __shared__ int shvar[10];69 auto lambda = [&]() {70 ext_shvar[0] = 1;71 shvar[0] = 2;72 };73 lambda();74}75 76template<class T>77void template_caller() {78 kernel5<<<1, 1>>>();79 var1 = 1;80}81 82template void template_caller<int>();83