45 lines · plain
1// RUN: %clang_cc1 -target-sdk-version=8.0 -emit-llvm %s -o - \2// RUN: | FileCheck %s --check-prefixes=CUDA-OLD,CHECK3// RUN: %clang_cc1 -target-sdk-version=9.2 -emit-llvm %s -o - \4// RUN: | FileCheck %s --check-prefixes=CUDA-NEW,CHECK5// RUN: %clang_cc1 -target-sdk-version=9.2 -emit-llvm %s -o - \6// RUN: -fgpu-default-stream=per-thread -DCUDA_API_PER_THREAD_DEFAULT_STREAM \7// RUN: | FileCheck %s --check-prefixes=CUDA-PTH,CHECK8// RUN: %clang_cc1 -x hip -emit-llvm %s -o - \9// RUN: | FileCheck %s --check-prefixes=HIP-OLD,CHECK10// RUN: %clang_cc1 -fhip-new-launch-api -x hip -emit-llvm %s -o - \11// RUN: | FileCheck %s --check-prefixes=HIP-NEW,LEGACY,CHECK12// RUN: %clang_cc1 -fhip-new-launch-api -x hip -emit-llvm %s -o - \13// RUN: -fgpu-default-stream=legacy \14// RUN: | FileCheck %s --check-prefixes=HIP-NEW,LEGACY,CHECK15// RUN: %clang_cc1 -fhip-new-launch-api -x hip -emit-llvm %s -o - \16// RUN: -fgpu-default-stream=per-thread -DHIP_API_PER_THREAD_DEFAULT_STREAM \17// RUN: | FileCheck %s --check-prefixes=HIP-NEW,PTH,CHECK18 19#include "Inputs/cuda.h"20 21// CHECK-LABEL: define{{.*}}g122// HIP-OLD: call{{.*}}hipSetupArgument23// HIP-OLD: call{{.*}}hipLaunchByPtr24// HIP-NEW: call{{.*}}__hipPopCallConfiguration25// LEGACY: call{{.*}}hipLaunchKernel26// PTH: call{{.*}}hipLaunchKernel_spt27// CUDA-OLD: call{{.*}}cudaSetupArgument28// CUDA-OLD: call{{.*}}cudaLaunch29// CUDA-NEW: call{{.*}}__cudaPopCallConfiguration30// CUDA-NEW: call{{.*}}cudaLaunchKernel31// CUDA-PTH: call{{.*}}cudaLaunchKernel_ptsz32__global__ void g1(int x) {}33 34// CHECK-LABEL: define{{.*}}main35int main(void) {36 // HIP-OLD: call{{.*}}hipConfigureCall37 // HIP-NEW: call{{.*}}__hipPushCallConfiguration38 // CUDA-OLD: call{{.*}}cudaConfigureCall39 // CUDA-NEW: call{{.*}}__cudaPushCallConfiguration40 // CHECK: icmp41 // CHECK: br42 // CHECK: call{{.*}}g143 g1<<<1, 1>>>(42);44}45