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 -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,CHECK
   9
  10 #include "Inputs/cuda.h"
  11
  12 // CHECK-LABEL: define{{.*}}g1
  13 // HIP-OLD: call{{.*}}hipSetupArgument
  14 // HIP-OLD: call{{.*}}hipLaunchByPtr
  15 // HIP-NEW: call{{.*}}__hipPopCallConfiguration
  16 // HIP-NEW: call{{.*}}hipLaunchKernel
  17 // CUDA-OLD: call{{.*}}cudaSetupArgument
  18 // CUDA-OLD: call{{.*}}cudaLaunch
  19 // CUDA-NEW: call{{.*}}__cudaPopCallConfiguration
  20 // CUDA-NEW: call{{.*}}cudaLaunchKernel
  21 __global__ void g1(int x) {}
  22
  23 // CHECK-LABEL: define{{.*}}main
  24 int main(void) {
  25   // HIP-OLD: call{{.*}}hipConfigureCall
  26   // HIP-NEW: call{{.*}}__hipPushCallConfiguration
  27   // CUDA-OLD: call{{.*}}cudaConfigureCall
  28   // CUDA-NEW: call{{.*}}__cudaPushCallConfiguration
  29   // CHECK: icmp
  30   // CHECK: br
  31   // CHECK: call{{.*}}g1
  32   g1<<<1, 1>>>(42);
  33 }