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-LABEL: @test_ballot_wave32_target_attr(
  17 // CHECK: call i32 @llvm.amdgcn.ballot.i32(i1 %{{.+}})
  18 __attribute__((target("wavefrontsize32")))
  19 void test_ballot_wave32_target_attr(global uint* out, int a, int b)
  20 {
  21   *out = __builtin_amdgcn_ballot_w32(a == b);
  22 }
  23
  24 #if __AMDGCN_WAVEFRONT_SIZE != 32
  25 #error Wrong wavesize detected
  26 #endif