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

   1 // REQUIRES: amdgpu-registered-target
   2 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
   3 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1011 -S -emit-llvm -o - %s | FileCheck %s
   4 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck %s
   5
   6 typedef unsigned int uint;
   7 typedef unsigned long ulong;
   8
   9 // CHECK-LABEL: @test_permlane16(
  10 // CHECK: call i32 @llvm.amdgcn.permlane16(i32 %a, i32 %b, i32 %c, i32 %d, i1 false, i1 false)
  11 void test_permlane16(global uint* out, uint a, uint b, uint c, uint d) {
  12   *out = __builtin_amdgcn_permlane16(a, b, c, d, 0, 0);
  13 }
  14
  15 // CHECK-LABEL: @test_permlanex16(
  16 // CHECK: call i32 @llvm.amdgcn.permlanex16(i32 %a, i32 %b, i32 %c, i32 %d, i1 false, i1 false)
  17 void test_permlanex16(global uint* out, uint a, uint b, uint c, uint d) {
  18   *out = __builtin_amdgcn_permlanex16(a, b, c, d, 0, 0);
  19 }
  20
  21 // CHECK-LABEL: @test_mov_dpp8(
  22 // CHECK: call i32 @llvm.amdgcn.mov.dpp8.i32(i32 %a, i32 1)
  23 void test_mov_dpp8(global uint* out, uint a) {
  24   *out = __builtin_amdgcn_mov_dpp8(a, 1);
  25 }
  26
  27 // CHECK-LABEL: @test_s_memtime
  28 // CHECK: call i64 @llvm.amdgcn.s.memtime()
  29 void test_s_memtime(global ulong* out)
  30 {
  31   *out = __builtin_amdgcn_s_memtime();
  32 }
  33
  34 // CHECK-LABEL: @test_groupstaticsize
  35 // CHECK: call i32 @llvm.amdgcn.groupstaticsize()
  36 void test_groupstaticsize(global uint* out)
  37 {
  38   *out = __builtin_amdgcn_groupstaticsize();
  39 }
  40
  41 // CHECK-LABEL: @test_ballot_wave32(
  42 // CHECK: call i32 @llvm.amdgcn.ballot.i32(i1 %{{.+}})
  43 void test_ballot_wave32(global uint* out, int a, int b)
  44 {
  45   *out = __builtin_amdgcn_ballot_w32(a == b);
  46 }