host-device-calls-host.cu 897 B

1234567891011121314151617181920212223242526272829303132
  1. // RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-allow-host-calls-from-host-device -fcuda-is-device -Wno-cuda-compat -emit-llvm -o - | FileCheck %s
  2. #include "Inputs/cuda.h"
  3. extern "C"
  4. void host_function() {}
  5. // CHECK-LABEL: define void @hd_function_a
  6. extern "C"
  7. __host__ __device__ void hd_function_a() {
  8. // CHECK: call void @host_function
  9. host_function();
  10. }
  11. // CHECK: declare void @host_function
  12. // CHECK-LABEL: define void @hd_function_b
  13. extern "C"
  14. __host__ __device__ void hd_function_b(bool b) { if (b) host_function(); }
  15. // CHECK-LABEL: define void @device_function_b
  16. extern "C"
  17. __device__ void device_function_b() { hd_function_b(false); }
  18. // CHECK-LABEL: define void @global_function
  19. extern "C"
  20. __global__ void global_function() {
  21. // CHECK: call void @device_function_b
  22. device_function_b();
  23. }
  24. // CHECK: !{{[0-9]+}} = !{void ()* @global_function, !"kernel", i32 1}