| // 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 amdgpu10.10-amd-amdhsa -emit-llvm %s -o - | FileCheck --check-prefix=AMDGPU1010 %s |
| |
| // Test that __builtin_amdgcn_is_invocable resolves from the triple's subarch |
| // alone, with no -target-cpu: |
| // 1) for amdgpu9.00 (gfx900) we emit an empty kernel (lacks feature) |
| // 2) for amdgpu10.10 (gfx1010) we emit a call to trap (has feature) |
| // AMDGPU900-LABEL: define dso_local void @foo( |
| // AMDGPU900-SAME: ) #[[ATTR0:[0-9]+]] { |
| // AMDGPU900-NEXT: [[ENTRY:.*:]] |
| // AMDGPU900-NEXT: ret void |
| // |
| // AMDGPU1010-LABEL: define dso_local void @foo( |
| // AMDGPU1010-SAME: ) #[[ATTR0:[0-9]+]] { |
| // AMDGPU1010-NEXT: [[ENTRY:.*:]] |
| // AMDGPU1010-NEXT: call void @llvm.trap() #[[ATTR2:[0-9]+]] |
| // AMDGPU1010-NEXT: unreachable |
| // AMDGPU1010: [[BB0:.*:]] |
| // AMDGPU1010-NEXT: ret void |
| // |
| void foo() { |
| if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_permlanex16) || |
| (__builtin_amdgcn_is_invocable(__builtin_amdgcn_permlanex16_var)) || |
| (__builtin_amdgcn_is_invocable(__builtin_amdgcn_ashr_pk_i8_i32))) |
| return __builtin_trap(); |
| } |
| //. |
| // AMDGPU900: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" } |
| //. |
| // AMDGPU1010: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" } |
| // AMDGPU1010: attributes #[[ATTR1:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) } |
| // AMDGPU1010: attributes #[[ATTR2]] = { noreturn nounwind } |
| //. |
| // AMDGPU900: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600} |
| // AMDGPU900: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} |
| //. |
| // AMDGPU1010: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600} |
| // AMDGPU1010: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} |
| //. |
