clang/test/CodeGenOpenCL/reflect.cl

   1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 4
   2 // RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -emit-llvm -O0 -o - | FileCheck %s
   3
   4 // CHECK-LABEL: define dso_local zeroext i1 @device_function(
   5 // CHECK-SAME: ) #[[ATTR0:[0-9]+]] {
   6 // CHECK-NEXT:  entry:
   7 // CHECK-NEXT:    [[TMP0:%.*]] = call i32 @llvm.nvvm.reflect(ptr addrspacecast (ptr addrspace(4) @.str to ptr))
   8 // CHECK-NEXT:    [[CMP:%.*]] = icmp uge i32 [[TMP0]], 700
   9 // CHECK-NEXT:    ret i1 [[CMP]]
  10 //
  11 bool device_function() {
  12   return __nvvm_reflect("__CUDA_ARCH") >= 700;
  13 }
  14
  15 // CHECK-LABEL: define dso_local spir_kernel void @kernel_function(
  16 // CHECK-SAME: ptr addrspace(1) noundef align 4 [[I:%.*]]) #[[ATTR2:[0-9]+]] !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 {
  17 // CHECK-NEXT:  entry:
  18 // CHECK-NEXT:    [[I_ADDR:%.*]] = alloca ptr addrspace(1), align 4
  19 // CHECK-NEXT:    store ptr addrspace(1) [[I]], ptr [[I_ADDR]], align 4
  20 // CHECK-NEXT:    [[CALL:%.*]] = call zeroext i1 @device_function() #[[ATTR3:[0-9]+]]
  21 // CHECK-NEXT:    [[CONV:%.*]] = zext i1 [[CALL]] to i32
  22 // CHECK-NEXT:    [[TMP0:%.*]] = load ptr addrspace(1), ptr [[I_ADDR]], align 4
  23 // CHECK-NEXT:    store i32 [[CONV]], ptr addrspace(1) [[TMP0]], align 4
  24 // CHECK-NEXT:    ret void
  25 //
  26 __kernel void kernel_function(__global int *i) {
  27   *i = device_function();
  28 }