19ecbb34eSYaxun (Sam) Liu // REQUIRES: x86-registered-target 2*9466b491SNikita Popov // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \ 39ecbb34eSYaxun (Sam) Liu // RUN: -x hip -emit-llvm-bc %s -o %t.hip.bc 4*9466b491SNikita Popov // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu \ 59ecbb34eSYaxun (Sam) Liu // RUN: -mlink-bitcode-file %t.hip.bc -DHIP_PLATFORM -emit-llvm \ 65cf2a37fSYaxun (Sam) Liu // RUN: %s -o - | FileCheck %s 75cf2a37fSYaxun (Sam) Liu 85cf2a37fSYaxun (Sam) Liu #include "Inputs/cuda.h" 95cf2a37fSYaxun (Sam) Liu 10*9466b491SNikita Popov // CHECK: @_Z2g1i = constant ptr @_Z17__device_stub__g1i, align 8 115cf2a37fSYaxun (Sam) Liu #if __HIP__ g1(int x)125cf2a37fSYaxun (Sam) Liu__global__ void g1(int x) {} 135cf2a37fSYaxun (Sam) Liu #else 145cf2a37fSYaxun (Sam) Liu extern void g1(int x); 155cf2a37fSYaxun (Sam) Liu 165cf2a37fSYaxun (Sam) Liu // CHECK: call i32 @hipLaunchKernel{{.*}}@_Z2g1i test()175cf2a37fSYaxun (Sam) Liuvoid test() { 185cf2a37fSYaxun (Sam) Liu hipLaunchKernel((void*)g1, 1, 1, nullptr, 0, 0); 195cf2a37fSYaxun (Sam) Liu } 205cf2a37fSYaxun (Sam) Liu 215cf2a37fSYaxun (Sam) Liu // CHECK: __hipRegisterFunction{{.*}}@_Z2g1i 225cf2a37fSYaxun (Sam) Liu #endif 23