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

   1 // REQUIRES: amdgpu-registered-target
   2 // RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu hawaii -emit-llvm -o - %s | FileCheck %s
   3 // RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu fiji -emit-llvm -o - %s | FileCheck %s
   4 // RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx906 -emit-llvm -o - %s | FileCheck %s
   5 // RUN: %clang_cc1 -cl-std=CL2.0 -O0 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -emit-llvm -o - %s | FileCheck %s
   6
   7 typedef unsigned int uint;
   8 typedef unsigned long ulong;
   9
  10 // CHECK-LABEL: @test_s_dcache_inv_vol
  11 // CHECK: call void @llvm.amdgcn.s.dcache.inv.vol(
  12 void test_s_dcache_inv_vol()
  13 {
  14   __builtin_amdgcn_s_dcache_inv_vol();
  15 }
  16
  17 // CHECK-LABEL: @test_buffer_wbinvl1_vol
  18 // CHECK: call void @llvm.amdgcn.buffer.wbinvl1.vol()
  19 void test_buffer_wbinvl1_vol()
  20 {
  21   __builtin_amdgcn_buffer_wbinvl1_vol();
  22 }
  23
  24 // CHECK-LABEL: @test_gws_sema_release_all(
  25 // CHECK: call void @llvm.amdgcn.ds.gws.sema.release.all(i32 %{{[0-9]+}})
  26 void test_gws_sema_release_all(uint id)
  27 {
  28   __builtin_amdgcn_ds_gws_sema_release_all(id);
  29 }
  30
  31 // CHECK-LABEL: @test_s_memtime
  32 // CHECK: call i64 @llvm.amdgcn.s.memtime()
  33 void test_s_memtime(global ulong* out)
  34 {
  35   *out = __builtin_amdgcn_s_memtime();
  36 }
  37
  38 // CHECK-LABEL: @test_is_shared(
  39 // CHECK: call i1 @llvm.amdgcn.is.shared(ptr %{{[0-9]+}}
  40 int test_is_shared(const int* ptr) {
  41   return __builtin_amdgcn_is_shared(ptr);
  42 }
  43
  44 // CHECK-LABEL: @test_is_private(
  45 // CHECK: call i1 @llvm.amdgcn.is.private(ptr %{{[0-9]+}}
  46 int test_is_private(const int* ptr) {
  47   return __builtin_amdgcn_is_private(ptr);
  48 }
  49
  50 // CHECK-LABEL: @test_is_shared_global(
  51 // CHECK: [[CAST:%[0-9]+]] = addrspacecast ptr addrspace(1) %{{[0-9]+}} to ptr
  52 // CHECK: call i1 @llvm.amdgcn.is.shared(ptr [[CAST]]
  53 int test_is_shared_global(const global int* ptr) {
  54   return __builtin_amdgcn_is_shared(ptr);
  55 }
  56
  57 // CHECK-LABEL: @test_is_private_global(
  58 // CHECK: [[CAST:%[0-9]+]] = addrspacecast ptr addrspace(1) %{{[0-9]+}} to ptr
  59 // CHECK: call i1 @llvm.amdgcn.is.private(ptr [[CAST]]
  60 int test_is_private_global(const global int* ptr) {
  61   return __builtin_amdgcn_is_private(ptr);
  62 }
  63
  64 // CHECK-LABEL: @test_groupstaticsize
  65 // CHECK: call i32 @llvm.amdgcn.groupstaticsize()
  66 void test_groupstaticsize(global uint* out) {
  67   *out = __builtin_amdgcn_groupstaticsize();
  68 }