| // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 6 | |
| // RUN: %clang_cc1 -triple amdgpu9.00-amd-amdhsa -emit-llvm %s -o - | FileCheck --check-prefix=AMDGPU900 %s | |
| // RUN: %clang_cc1 -triple amdgpu9-amd-amdhsa -emit-llvm %s -o - | FileCheck --check-prefix=AMDGPU9_NOCPU %s | |
| // RUN: %clang_cc1 -triple amdgpu9-amd-amdhsa -target-cpu gfx900 -emit-llvm %s -o - | FileCheck --check-prefix=AMDGPU900_CPU %s | |
| // Test that __builtin_amdgcn_processor_is works when the target is specified | |
| // only by the subarch without -target-cpu. | |
| // AMDGPU900-LABEL: define dso_local void @gfx900_name( | |
| // AMDGPU900-SAME: ) #[[ATTR0:[0-9]+]] { | |
| // AMDGPU900-NEXT: [[ENTRY:.*:]] | |
| // AMDGPU900-NEXT: call void @llvm.trap() #[[ATTR2:[0-9]+]] | |
| // AMDGPU900-NEXT: unreachable | |
| // AMDGPU900: [[BB0:.*:]] | |
| // AMDGPU900-NEXT: ret void | |
| // | |
| // AMDGPU9_NOCPU-LABEL: define dso_local void @gfx900_name( | |
| // AMDGPU9_NOCPU-SAME: ) #[[ATTR0:[0-9]+]] { | |
| // AMDGPU9_NOCPU-NEXT: [[ENTRY:.*:]] | |
| // AMDGPU9_NOCPU-NEXT: ret void | |
| // | |
| // AMDGPU900_CPU-LABEL: define dso_local void @gfx900_name( | |
| // AMDGPU900_CPU-SAME: ) #[[ATTR0:[0-9]+]] { | |
| // AMDGPU900_CPU-NEXT: [[ENTRY:.*:]] | |
| // AMDGPU900_CPU-NEXT: call void @llvm.trap() #[[ATTR2:[0-9]+]] | |
| // AMDGPU900_CPU-NEXT: unreachable | |
| // AMDGPU900_CPU: [[BB0:.*:]] | |
| // AMDGPU900_CPU-NEXT: ret void | |
| // | |
| void gfx900_name() { | |
| if (__builtin_amdgcn_processor_is("gfx900")) | |
| return __builtin_trap(); | |
| } | |
| //. | |
| // AMDGPU900: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" } | |
| // AMDGPU900: attributes #[[ATTR1:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) } | |
| // AMDGPU900: attributes #[[ATTR2]] = { noreturn nounwind } | |
| //. | |
| // AMDGPU9_NOCPU: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" } | |
| //. | |
| // AMDGPU900_CPU: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" "target-cpu"="gfx900" } | |
| // AMDGPU900_CPU: attributes #[[ATTR1:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) } | |
| // AMDGPU900_CPU: attributes #[[ATTR2]] = { noreturn nounwind } | |
| //. | |
| // AMDGPU900: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600} | |
| // AMDGPU900: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} | |
| //. | |
| // AMDGPU9_NOCPU: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600} | |
| // AMDGPU9_NOCPU: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} | |
| //. | |
| // AMDGPU900_CPU: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600} | |
| // AMDGPU900_CPU: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} | |
| //. |