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// 11bool device_function() { 12 return __nvvm_reflect("__CUDA_ARCH") >= 700; 13} 14 15// CHECK-LABEL: define dso_local ptx_kernel void @kernel_function( 16// CHECK-SAME: ptr addrspace(1) noundef align 4 [[I:%.*]]) #[[ATTR2:[0-9]+]] !kernel_arg_addr_space [[META3:![0-9]+]] !kernel_arg_access_qual [[META4:![0-9]+]] !kernel_arg_type [[META5:![0-9]+]] !kernel_arg_base_type [[META5]] !kernel_arg_type_qual [[META6:![0-9]+]] { 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} 29//. 30// CHECK: [[META3]] = !{i32 1} 31// CHECK: [[META4]] = !{!"none"} 32// CHECK: [[META5]] = !{!"int*"} 33// CHECK: [[META6]] = !{!""} 34//. 35