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
g1(int x)28 __global__ void g1(int x) {}
29
30 // CHECK-LABEL: define{{.*}}main
main(void)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