clang/test/CodeGenCUDA/struct-mangling-number.cu

   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 }