clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl

   1 // RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -D__AMDGCN_WAVEFRONT_SIZE=32 -target-feature +wavefrontsize32 -S -emit-llvm -o - %s | FileCheck -enable-var-scope %s
   2 // RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck -enable-var-scope %s
   3 // RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -target-feature +wavefrontsize32 -S -emit-llvm -o - %s | FileCheck -enable-var-scope %s
   4 // RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -target-feature +wavefrontsize32 -S -emit-llvm -o - %s | FileCheck -enable-var-scope %s
   5
   6 typedef unsigned int uint;
   7
   8
   9 // CHECK-LABEL: @test_ballot_wave32(
  10 // CHECK: call i32 @llvm.amdgcn.ballot.i32(i1 %{{.+}})
  11 void test_ballot_wave32(global uint* out, int a, int b)
  12 {
  13   *out = __builtin_amdgcn_ballot_w32(a == b);
  14 }
  15
  16 // CHECK: declare i32 @llvm.amdgcn.ballot.i32(i1) #[[$NOUNWIND_READONLY:[0-9]+]]
  17
  18 // CHECK-LABEL: @test_ballot_wave32_target_attr(
  19 // CHECK: call i32 @llvm.amdgcn.ballot.i32(i1 %{{.+}})
  20 __attribute__((target("wavefrontsize32")))
  21 void test_ballot_wave32_target_attr(global uint* out, int a, int b)
  22 {
  23   *out = __builtin_amdgcn_ballot_w32(a == b);
  24 }
  25
  26 // CHECK-LABEL: @test_read_exec(
  27 // CHECK: call i64 @llvm.amdgcn.ballot.i64(i1 true)
  28 void test_read_exec(global uint* out) {
  29   *out = __builtin_amdgcn_read_exec();
  30 }
  31
  32 // CHECK: declare i64 @llvm.amdgcn.ballot.i64(i1) #[[$NOUNWIND_READONLY:[0-9]+]]
  33
  34 // CHECK-LABEL: @test_read_exec_lo(
  35 // CHECK: call i32 @llvm.amdgcn.ballot.i32(i1 true)
  36 void test_read_exec_lo(global uint* out) {
  37   *out = __builtin_amdgcn_read_exec_lo();
  38 }
  39
  40 // CHECK-LABEL: @test_read_exec_hi(
  41 // CHECK: call i64 @llvm.amdgcn.ballot.i64(i1 true)
  42 // CHECK: lshr i64 [[A:%.*]], 32
  43 // CHECK: trunc i64 [[B:%.*]] to i32
  44 void test_read_exec_hi(global uint* out) {
  45   *out = __builtin_amdgcn_read_exec_hi();
  46 }
  47
  48 #if __AMDGCN_WAVEFRONT_SIZE != 32
  49 #error Wrong wavesize detected
  50 #endif