clang/test/CodeGenCUDA/kernel-call.cu

   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 -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,CHECK
   8 // RUN: %clang_cc1 -x hip -emit-llvm %s -o - \
   9 // RUN: | FileCheck %s --check-prefixes=HIP-OLD,CHECK
  10 // RUN: %clang_cc1 -fhip-new-launch-api -x hip -emit-llvm %s -o - \
  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=legacy \
  14 // RUN:   | FileCheck %s --check-prefixes=HIP-NEW,LEGACY,CHECK
  15 // 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,CHECK
  18
  19 #include "Inputs/cuda.h"
  20
  21 // CHECK-LABEL: define{{.*}}g1
  22 // HIP-OLD: call{{.*}}hipSetupArgument
  23 // HIP-OLD: call{{.*}}hipLaunchByPtr
  24 // HIP-NEW: call{{.*}}__hipPopCallConfiguration
  25 // LEGACY: call{{.*}}hipLaunchKernel
  26 // PTH: call{{.*}}hipLaunchKernel_spt
  27 // CUDA-OLD: call{{.*}}cudaSetupArgument
  28 // CUDA-OLD: call{{.*}}cudaLaunch
  29 // CUDA-NEW: call{{.*}}__cudaPopCallConfiguration
  30 // CUDA-NEW: call{{.*}}cudaLaunchKernel
  31 // CUDA-PTH: call{{.*}}cudaLaunchKernel_ptsz
  32 __global__ void g1(int x) {}
  33
  34 // CHECK-LABEL: define{{.*}}main
  35 int main(void) {
  36   // HIP-OLD: call{{.*}}hipConfigureCall
  37   // HIP-NEW: call{{.*}}__hipPushCallConfiguration
  38   // CUDA-OLD: call{{.*}}cudaConfigureCall
  39   // CUDA-NEW: call{{.*}}__cudaPushCallConfiguration
  40   // CHECK: icmp
  41   // CHECK: br
  42   // CHECK: call{{.*}}g1
  43   g1<<<1, 1>>>(42);
  44 }