1 // RUN: %clang_cc1 -target-sdk-version=8.0 -emit-llvm %s -o - \
2 // RUN: | FileCheck %s --check-prefixes=CUDA-OLD,CHECK
3 // RUN: %clang_cc1 -target-sdk-version=9.2  -emit-llvm %s -o - \
4 // RUN: | FileCheck %s --check-prefixes=CUDA-NEW,CHECK
5 // RUN: %clang_cc1 -x hip -emit-llvm %s -o - \
6 // RUN: | FileCheck %s --check-prefixes=HIP-OLD,CHECK
7 // RUN: %clang_cc1 -fhip-new-launch-api -x hip -emit-llvm %s -o - \
8 // RUN: | FileCheck %s --check-prefixes=HIP-NEW,LEGACY,CHECK
9 // RUN: %clang_cc1 -fhip-new-launch-api -x hip -emit-llvm %s -o - \
10 // RUN:   -fgpu-default-stream=legacy \
11 // RUN:   | FileCheck %s --check-prefixes=HIP-NEW,LEGACY,CHECK
12 // RUN: %clang_cc1 -fhip-new-launch-api -x hip -emit-llvm %s -o - \
13 // RUN:   -fgpu-default-stream=per-thread -DHIP_API_PER_THREAD_DEFAULT_STREAM \
14 // RUN:   | FileCheck %s --check-prefixes=HIP-NEW,PTH,CHECK
15 
16 #include "Inputs/cuda.h"
17 
18 // CHECK-LABEL: define{{.*}}g1
19 // HIP-OLD: call{{.*}}hipSetupArgument
20 // HIP-OLD: call{{.*}}hipLaunchByPtr
21 // HIP-NEW: call{{.*}}__hipPopCallConfiguration
22 // LEGACY: call{{.*}}hipLaunchKernel
23 // PTH: call{{.*}}hipLaunchKernel_spt
24 // CUDA-OLD: call{{.*}}cudaSetupArgument
25 // CUDA-OLD: call{{.*}}cudaLaunch
26 // CUDA-NEW: call{{.*}}__cudaPopCallConfiguration
27 // CUDA-NEW: call{{.*}}cudaLaunchKernel
28 __global__ void g1(int x) {}
29 
30 // CHECK-LABEL: define{{.*}}main
31 int main(void) {
32   // HIP-OLD: call{{.*}}hipConfigureCall
33   // HIP-NEW: call{{.*}}__hipPushCallConfiguration
34   // CUDA-OLD: call{{.*}}cudaConfigureCall
35   // CUDA-NEW: call{{.*}}__cudaPushCallConfiguration
36   // CHECK: icmp
37   // CHECK: br
38   // CHECK: call{{.*}}g1
39   g1<<<1, 1>>>(42);
40 }
41