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

   1 // REQUIRES: amdgpu-registered-target
   2 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -S -emit-llvm -o - %s | FileCheck %s
   3 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1101 -S -emit-llvm -o - %s | FileCheck %s
   4 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1102 -S -emit-llvm -o - %s | FileCheck %s
   5 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1103 -S -emit-llvm -o - %s | FileCheck %s
   6 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1150 -S -emit-llvm -o - %s | FileCheck %s
   7 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1151 -S -emit-llvm -o - %s | FileCheck %s
   8
   9 typedef unsigned int uint;
  10 typedef unsigned long ulong;
  11 typedef uint uint2 __attribute__((ext_vector_type(2)));
  12 typedef uint uint4 __attribute__((ext_vector_type(4)));
  13
  14 // CHECK-LABEL: @test_s_sendmsg_rtn(
  15 // CHECK: call i32 @llvm.amdgcn.s.sendmsg.rtn.i32(i32 0)
  16 void test_s_sendmsg_rtn(global uint* out) {
  17   *out = __builtin_amdgcn_s_sendmsg_rtn(0);
  18 }
  19
  20 // CHECK-LABEL: @test_s_sendmsg_rtnl(
  21 // CHECK: call i64 @llvm.amdgcn.s.sendmsg.rtn.i64(i32 0)
  22 void test_s_sendmsg_rtnl(global ulong* out) {
  23   *out = __builtin_amdgcn_s_sendmsg_rtnl(0);
  24 }
  25
  26 // CHECK-LABEL: @test_ds_bvh_stack_rtn(
  27 // CHECK: %0 = tail call { i32, i32 } @llvm.amdgcn.ds.bvh.stack.rtn(i32 %addr, i32 %data, <4 x i32> %data1, i32 128)
  28 // CHECK: %1 = extractvalue { i32, i32 } %0, 0
  29 // CHECK: %2 = extractvalue { i32, i32 } %0, 1
  30 // CHECK: %3 = insertelement <2 x i32> poison, i32 %1, i64 0
  31 // CHECK: %4 = insertelement <2 x i32> %3, i32 %2, i64 1
  32 void test_ds_bvh_stack_rtn(global uint2* out, uint addr, uint data, uint4 data1)
  33 {
  34   *out = __builtin_amdgcn_ds_bvh_stack_rtn(addr, data, data1, 128);
  35 }
  36
  37 // CHECK-LABEL: @test_permlane64(
  38 // CHECK: call i32 @llvm.amdgcn.permlane64(i32 %a)
  39 void test_permlane64(global uint* out, uint a) {
  40   *out = __builtin_amdgcn_permlane64(a);
  41 }
  42
  43 // CHECK-LABEL: @test_s_wait_event_export_ready
  44 // CHECK: call void @llvm.amdgcn.s.wait.event.export.ready
  45 void test_s_wait_event_export_ready() {
  46   __builtin_amdgcn_s_wait_event_export_ready();
  47 }
  48
  49 // CHECK-LABEL: @test_global_add_f32
  50 // CHECK: call float @llvm.amdgcn.global.atomic.fadd.f32.p1.f32(ptr addrspace(1) %{{.*}}, float %{{.*}})
  51 void test_global_add_f32(float *rtn, global float *addr, float x) {
  52   *rtn = __builtin_amdgcn_global_atomic_fadd_f32(addr, x);
  53 }