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> 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> 46 void fun() 47 { 48 } 49 50 template <typename T> 51 class Test { 52 public: 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 65 int main() { 66 Test<int> A; 67 A.run(); 68 } 69