51 lines · plain
1// REQUIRES: amdgpu-registered-target2// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-feature +wavefrontsize32 -emit-llvm -o - %s | FileCheck -enable-var-scope %s3// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -emit-llvm -o - %s | FileCheck -enable-var-scope %s4// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -target-feature +wavefrontsize32 -emit-llvm -o - %s | FileCheck -enable-var-scope %s5// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -target-feature +wavefrontsize32 -emit-llvm -o - %s | FileCheck -enable-var-scope %s6 7typedef unsigned int uint;8 9 10// CHECK-LABEL: @test_ballot_wave32(11// CHECK: call i32 @llvm.amdgcn.ballot.i32(i1 %{{.+}})12void test_ballot_wave32(global uint* out, int a, int b)13{14 *out = __builtin_amdgcn_ballot_w32(a == b);15}16 17// CHECK: declare i32 @llvm.amdgcn.ballot.i32(i1) #[[$NOUNWIND_READONLY:[0-9]+]]18 19// CHECK-LABEL: @test_ballot_wave32_target_attr(20// CHECK: call i32 @llvm.amdgcn.ballot.i32(i1 %{{.+}})21__attribute__((target("wavefrontsize32")))22void test_ballot_wave32_target_attr(global uint* out, int a, int b)23{24 *out = __builtin_amdgcn_ballot_w32(a == b);25}26 27// CHECK-LABEL: @test_inverse_ballot_wave32(28// CHECK: call i1 @llvm.amdgcn.inverse.ballot.i32(i32 %{{.+}})29void test_inverse_ballot_wave32(global bool* out, int a)30{31 *out = __builtin_amdgcn_inverse_ballot_w32(a);32}33 34// CHECK-LABEL: @test_read_exec(35// CHECK: call i32 @llvm.amdgcn.ballot.i32(i1 true)36void test_read_exec(global uint* out) {37 *out = __builtin_amdgcn_read_exec();38}39 40// CHECK-LABEL: @test_read_exec_lo(41// CHECK: call i32 @llvm.amdgcn.ballot.i32(i1 true)42void test_read_exec_lo(global uint* out) {43 *out = __builtin_amdgcn_read_exec_lo();44}45 46// CHECK-LABEL: @test_read_exec_hi(47// CHECK: store i32 0, ptr addrspace(1) %out48void test_read_exec_hi(global uint* out) {49 *out = __builtin_amdgcn_read_exec_hi();50}51