1// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --check-attributes --check-globals 2// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -fno-ident -fcuda-is-device \ 3// RUN: -emit-llvm -o - %s | FileCheck -check-prefix=OPTNONE %s 4 5#define __device__ __attribute__((device)) 6#define __global__ __attribute__((global)) 7 8//. 9// OPTNONE: @__hip_cuid_ = addrspace(1) global i8 0 10// OPTNONE: @llvm.compiler.used = appending addrspace(1) global [1 x ptr] [ptr addrspacecast (ptr addrspace(1) @__hip_cuid_ to ptr)], section "llvm.metadata" 11// OPTNONE: @__oclc_ABI_version = weak_odr hidden local_unnamed_addr addrspace(4) constant i32 500 12//. 13__device__ void extern_func(); 14 15// OPTNONE: Function Attrs: convergent mustprogress noinline nounwind optnone 16// OPTNONE-LABEL: define {{[^@]+}}@_Z4funcv 17// OPTNONE-SAME: () #[[ATTR0:[0-9]+]] { 18// OPTNONE-NEXT: entry: 19// OPTNONE-NEXT: call void @_Z11extern_funcv() #[[ATTR3:[0-9]+]] 20// OPTNONE-NEXT: ret void 21// 22__device__ void func() { 23 extern_func(); 24} 25 26// OPTNONE: Function Attrs: convergent mustprogress noinline norecurse nounwind optnone 27// OPTNONE-LABEL: define {{[^@]+}}@_Z6kernelv 28// OPTNONE-SAME: () #[[ATTR2:[0-9]+]] { 29// OPTNONE-NEXT: entry: 30// OPTNONE-NEXT: call void @_Z11extern_funcv() #[[ATTR3]] 31// OPTNONE-NEXT: ret void 32// 33__global__ void kernel() { 34 extern_func(); 35} 36//. 37// OPTNONE: attributes #[[ATTR0]] = { convergent mustprogress noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" } 38// OPTNONE: attributes #[[ATTR1:[0-9]+]] = { convergent nounwind "no-trapping-math"="true" "stack-protector-buffer-size"="8" } 39// OPTNONE: attributes #[[ATTR2]] = { convergent mustprogress noinline norecurse nounwind optnone "amdgpu-flat-work-group-size"="1,1024" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" } 40// OPTNONE: attributes #[[ATTR3]] = { convergent nounwind } 41//. 42// OPTNONE: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 500} 43// OPTNONE: [[META1:![0-9]+]] = !{i32 1, !"amdgpu_printf_kind", !"hostcall"} 44// OPTNONE: [[META2:![0-9]+]] = !{i32 1, !"wchar_size", i32 4} 45//. 46