69 lines · plain
1// RUN: %clang_cc1 -emit-llvm -o - -aux-triple x86_64-pc-windows-msvc \2// RUN: -fms-extensions -triple amdgcn-amd-amdhsa \3// RUN: -target-cpu gfx1030 -fcuda-is-device -x hip %s \4// RUN: | FileCheck -check-prefix=DEV %s5 6// RUN: %clang_cc1 -emit-llvm -o - -triple x86_64-pc-windows-msvc \7// RUN: -fms-extensions -aux-triple amdgcn-amd-amdhsa \8// RUN: -aux-target-cpu gfx1030 -x hip %s \9// RUN: | FileCheck -check-prefix=HOST %s10 11// RUN: %clang_cc1 -emit-llvm -o - -triple x86_64-pc-windows-msvc \12// RUN: -fms-extensions -aux-triple amdgcn-amd-amdhsa \13// RUN: -aux-target-cpu gfx1030 -x hip %s \14// RUN: | FileCheck -check-prefix=HOST-NEG %s15 16// RUN: %clang_cc1 -emit-llvm -o - -triple x86_64-pc-windows-msvc \17// RUN: -fms-extensions -x c++ %s \18// RUN: | FileCheck -check-prefix=CPP %s19 20#if __HIP__21#include "Inputs/cuda.h"22#endif23 24// Check local struct 'Op' uses Itanium mangling number instead of MSVC mangling25// number in device side name mangling. It is the same in device and host26// compilation.27 28// DEV: define amdgpu_kernel void @_Z6kernelIZN4TestIiE3runEvE2OpEvv(29 30// HOST-DAG: @{{.*}} = {{.*}}c"_Z6kernelIZN4TestIiE3runEvE2OpEvv\00"31 32// HOST-NEG-NOT: @{{.*}} = {{.*}}c"_Z6kernelIZN4TestIiE3runEvE2Op_1Evv\00"33#if __HIP__34template<typename T>35__attribute__((global)) void kernel()36{37}38#endif39 40// Check local struct 'Op' uses MSVC mangling number in host function name mangling.41// It is the same when compiled as HIP or C++ program.42 43// HOST-DAG: call void @"??$fun@UOp@?2??run@?$Test@H@@QEAAXXZ@@@YAXXZ"()44// CPP: call void @"??$fun@UOp@?2??run@?$Test@H@@QEAAXXZ@@@YAXXZ"()45template<typename T>46void fun()47{48}49 50template <typename T>51class Test {52public:53 void run()54 {55 struct Op56 {57 };58#if __HIP__59 kernel<Op><<<1, 1>>>();60#endif61 fun<Op>();62 }63};64 65int main() {66 Test<int> A;67 A.run();68}69