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 %s
5
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 %s
10
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 %s
15
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 %s
19
20 #if __HIP__
21 #include "Inputs/cuda.h"
22 #endif
23
24 // Check local struct 'Op' uses Itanium mangling number instead of MSVC mangling
25 // number in device side name mangling. It is the same in device and host
26 // 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__
34 template<typename T>
kernel()35 __attribute__((global)) void kernel()
36 {
37 }
38 #endif
39
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"()
45 template<typename T>
fun()46 void fun()
47 {
48 }
49
50 template <typename T>
51 class Test {
52 public:
run()53 void run()
54 {
55 struct Op
56 {
57 };
58 #if __HIP__
59 kernel<Op><<<1, 1>>>();
60 #endif
61 fun<Op>();
62 }
63 };
64
main()65 int main() {
66 Test<int> A;
67 A.run();
68 }
69