| // 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=AMDGCN-GFX900 %s | |
| // RUN: %clang_cc1 -triple amdgpu10.10-amd-amdhsa -emit-llvm %s -o - | FileCheck --check-prefix=AMDGCN-GFX1010 %s | |
| // RUN: %clang_cc1 -triple spirv64-amd-amdhsa -emit-llvm %s -o - | FileCheck --check-prefix=AMDGCNSPIRV %s | |
| // Test that, depending on triple and, if applicable, target-cpu, one of three | |
| // things happens: | |
| // 1) for gfx900 we emit an empty kernel (concrete target, lacks feature) | |
| // 2) for gfx1010 we emit a call to trap (concrete target, has feature) | |
| // 3) for AMDGCNSPIRV we emit a boolean specialisation constant, via a call | |
| // to __spirv_SpecConstant, with the id of UINT32_MAX, and the boolean | |
| // value of false, which will yield an OpSpecConstantFalse in SPIR-V | |
| // AMDGCN-GFX900-LABEL: define dso_local void @foo( | |
| // AMDGCN-GFX900-SAME: ) #[[ATTR0:[0-9]+]] { | |
| // AMDGCN-GFX900-NEXT: [[ENTRY:.*:]] | |
| // AMDGCN-GFX900-NEXT: ret void | |
| // | |
| // AMDGCN-GFX1010-LABEL: define dso_local void @foo( | |
| // AMDGCN-GFX1010-SAME: ) #[[ATTR0:[0-9]+]] { | |
| // AMDGCN-GFX1010-NEXT: [[ENTRY:.*:]] | |
| // AMDGCN-GFX1010-NEXT: call void @llvm.trap() #[[ATTR2:[0-9]+]] | |
| // AMDGCN-GFX1010-NEXT: unreachable | |
| // AMDGCN-GFX1010: [[BB0:.*:]] | |
| // AMDGCN-GFX1010-NEXT: ret void | |
| // | |
| // AMDGCNSPIRV-LABEL: define spir_func void @foo( | |
| // AMDGCNSPIRV-SAME: ) addrspace(4) #[[ATTR0:[0-9]+]] { | |
| // AMDGCNSPIRV-NEXT: [[ENTRY:.*:]] | |
| // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = call addrspace(4) i1 @llvm.spv.named.boolean.spec.constant(i32 -1, i1 false, metadata [[META2:![0-9]+]]) | |
| // AMDGCNSPIRV-NEXT: [[TOBOOL:%.*]] = icmp ne i1 [[TMP0]], false | |
| // AMDGCNSPIRV-NEXT: br i1 [[TOBOOL]], label %[[IF_THEN:.*]], label %[[LOR_LHS_FALSE:.*]] | |
| // AMDGCNSPIRV: [[LOR_LHS_FALSE]]: | |
| // AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = call addrspace(4) i1 @llvm.spv.named.boolean.spec.constant(i32 -1, i1 false, metadata [[META3:![0-9]+]]) | |
| // AMDGCNSPIRV-NEXT: [[TOBOOL1:%.*]] = icmp ne i1 [[TMP1]], false | |
| // AMDGCNSPIRV-NEXT: br i1 [[TOBOOL1]], label %[[IF_THEN]], label %[[LOR_LHS_FALSE2:.*]] | |
| // AMDGCNSPIRV: [[LOR_LHS_FALSE2]]: | |
| // AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = call addrspace(4) i1 @llvm.spv.named.boolean.spec.constant(i32 -1, i1 false, metadata [[META4:![0-9]+]]) | |
| // AMDGCNSPIRV-NEXT: [[TOBOOL3:%.*]] = icmp ne i1 [[TMP2]], false | |
| // AMDGCNSPIRV-NEXT: br i1 [[TOBOOL3]], label %[[IF_THEN]], label %[[IF_END:.*]] | |
| // AMDGCNSPIRV: [[IF_THEN]]: | |
| // AMDGCNSPIRV-NEXT: call addrspace(4) void @llvm.trap() #[[ATTR3:[0-9]+]] | |
| // AMDGCNSPIRV-NEXT: unreachable | |
| // AMDGCNSPIRV: [[BB3:.*:]] | |
| // AMDGCNSPIRV-NEXT: br label %[[IF_END]] | |
| // AMDGCNSPIRV: [[IF_END]]: | |
| // AMDGCNSPIRV-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(); | |
| } | |
| //. | |
| // AMDGCN-GFX900: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" } | |
| //. | |
| // AMDGCN-GFX1010: attributes #[[ATTR0]] = { convergent noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" } | |
| // AMDGCN-GFX1010: attributes #[[ATTR1:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) } | |
| // AMDGCN-GFX1010: attributes #[[ATTR2]] = { noreturn nounwind } | |
| //. | |
| // AMDGCNSPIRV: attributes #[[ATTR0]] = { noinline nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" } | |
| // AMDGCNSPIRV: attributes #[[ATTR1:[0-9]+]] = { nounwind } | |
| // AMDGCNSPIRV: attributes #[[ATTR2:[0-9]+]] = { cold noreturn nounwind memory(inaccessiblemem: write) } | |
| // AMDGCNSPIRV: attributes #[[ATTR3]] = { noreturn nounwind } | |
| //. | |
| // AMDGCN-GFX900: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600} | |
| // AMDGCN-GFX900: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} | |
| //. | |
| // AMDGCN-GFX1010: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600} | |
| // AMDGCN-GFX1010: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} | |
| //. | |
| // AMDGCNSPIRV: [[META0:![0-9]+]] = !{i32 1, !"amdhsa_code_object_version", i32 600} | |
| // AMDGCNSPIRV: [[META1:![0-9]+]] = !{!"{{.*}}clang version {{.*}}"} | |
| // AMDGCNSPIRV: [[META2]] = !{!"has.gfx10-insts"} | |
| // AMDGCNSPIRV: [[META3]] = !{!"has.gfx12-insts"} | |
| // AMDGCNSPIRV: [[META4]] = !{!"has.ashr-pk-insts"} | |
| //. |