clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w32.cl

   1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
   2 // REQUIRES: amdgpu-registered-target
   3 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -DWMMA_GFX1100_TESTS -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1100
   4
   5 typedef float  v8f   __attribute__((ext_vector_type(8)));
   6 typedef half   v16h  __attribute__((ext_vector_type(16)));
   7 typedef int    v2i   __attribute__((ext_vector_type(2)));
   8 typedef int    v4i   __attribute__((ext_vector_type(4)));
   9 typedef int    v8i   __attribute__((ext_vector_type(8)));
  10 typedef short  v16s  __attribute__((ext_vector_type(16)));
  11
  12 #ifdef WMMA_GFX1100_TESTS
  13
  14 // Wave32
  15
  16 //
  17 // amdgcn_wmma_f32_16x16x16_f16
  18 //
  19
  20 // CHECK-GFX1100-LABEL: @test_amdgcn_wmma_f32_16x16x16_f16_w32(
  21 // CHECK-GFX1100-NEXT:  entry:
  22 // CHECK-GFX1100-NEXT:    [[TMP0:%.*]] = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x16.f16.v8f32.v16f16(<16 x half> [[A:%.*]], <16 x half> [[B:%.*]], <8 x float> [[C:%.*]])
  23 // CHECK-GFX1100-NEXT:    store <8 x float> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 32, !tbaa [[TBAA4:![0-9]+]]
  24 // CHECK-GFX1100-NEXT:    ret void
  25 //
  26 void test_amdgcn_wmma_f32_16x16x16_f16_w32(global v8f* out, v16h a, v16h b, v8f c)
  27 {
  28   *out = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a, b, c);
  29 }
  30
  31 //
  32 // amdgcn_wmma_f32_16x16x16_bf16
  33 //
  34
  35 // CHECK-GFX1100-LABEL: @test_amdgcn_wmma_f32_16x16x16_bf16_w32(
  36 // CHECK-GFX1100-NEXT:  entry:
  37 // CHECK-GFX1100-NEXT:    [[TMP0:%.*]] = tail call <8 x float> @llvm.amdgcn.wmma.f32.16x16x16.bf16.v8f32.v16i16(<16 x i16> [[A:%.*]], <16 x i16> [[B:%.*]], <8 x float> [[C:%.*]])
  38 // CHECK-GFX1100-NEXT:    store <8 x float> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 32, !tbaa [[TBAA4]]
  39 // CHECK-GFX1100-NEXT:    ret void
  40 //
  41 void test_amdgcn_wmma_f32_16x16x16_bf16_w32(global v8f* out, v16s a, v16s b, v8f c)
  42 {
  43   *out = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32(a, b, c);
  44 }
  45
  46 //
  47 // amdgcn_wmma_f16_16x16x16_f16
  48 //
  49
  50 // CHECK-GFX1100-LABEL: @test_amdgcn_wmma_f16_16x16x16_f16_w32(
  51 // CHECK-GFX1100-NEXT:  entry:
  52 // CHECK-GFX1100-NEXT:    [[TMP0:%.*]] = tail call <16 x half> @llvm.amdgcn.wmma.f16.16x16x16.f16.v16f16.v16f16(<16 x half> [[A:%.*]], <16 x half> [[B:%.*]], <16 x half> [[C:%.*]], i1 true)
  53 // CHECK-GFX1100-NEXT:    store <16 x half> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 32, !tbaa [[TBAA4]]
  54 // CHECK-GFX1100-NEXT:    ret void
  55 //
  56 void test_amdgcn_wmma_f16_16x16x16_f16_w32(global v16h* out, v16h a, v16h b, v16h c)
  57 {
  58   *out = __builtin_amdgcn_wmma_f16_16x16x16_f16_w32(a, b, c, true);
  59 }
  60
  61 //
  62 // amdgcn_wmma_bf16_16x16x16_bf16
  63 //
  64
  65 // CHECK-GFX1100-LABEL: @test_amdgcn_wmma_bf16_16x16x16_bf16_w32(
  66 // CHECK-GFX1100-NEXT:  entry:
  67 // CHECK-GFX1100-NEXT:    [[TMP0:%.*]] = tail call <16 x i16> @llvm.amdgcn.wmma.bf16.16x16x16.bf16.v16i16.v16i16(<16 x i16> [[A:%.*]], <16 x i16> [[B:%.*]], <16 x i16> [[C:%.*]], i1 true)
  68 // CHECK-GFX1100-NEXT:    store <16 x i16> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 32, !tbaa [[TBAA4]]
  69 // CHECK-GFX1100-NEXT:    ret void
  70 //
  71 void test_amdgcn_wmma_bf16_16x16x16_bf16_w32(global v16s* out, v16s a, v16s b, v16s c)
  72 {
  73   *out = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32(a, b, c, true);
  74 }
  75
  76 //
  77 // amdgcn_wmma_f16_16x16x16_f16_tied
  78 //
  79
  80 // CHECK-GFX1100-LABEL: @test_amdgcn_wmma_f16_16x16x16_f16_tied_w32(
  81 // CHECK-GFX1100-NEXT:  entry:
  82 // CHECK-GFX1100-NEXT:    [[TMP0:%.*]] = tail call <16 x half> @llvm.amdgcn.wmma.f16.16x16x16.f16.tied.v16f16.v16f16(<16 x half> [[A:%.*]], <16 x half> [[B:%.*]], <16 x half> [[C:%.*]], i1 true)
  83 // CHECK-GFX1100-NEXT:    store <16 x half> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 32, !tbaa [[TBAA4]]
  84 // CHECK-GFX1100-NEXT:    ret void
  85 //
  86 void test_amdgcn_wmma_f16_16x16x16_f16_tied_w32(global v16h* out, v16h a, v16h b, v16h c)
  87 {
  88   *out = __builtin_amdgcn_wmma_f16_16x16x16_f16_tied_w32(a, b, c, true);
  89 }
  90
  91 //
  92 // amdgcn_wmma_bf16_16x16x16_bf16_tied
  93 //
  94
  95 // CHECK-GFX1100-LABEL: @test_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32(
  96 // CHECK-GFX1100-NEXT:  entry:
  97 // CHECK-GFX1100-NEXT:    [[TMP0:%.*]] = tail call <16 x i16> @llvm.amdgcn.wmma.bf16.16x16x16.bf16.tied.v16i16.v16i16(<16 x i16> [[A:%.*]], <16 x i16> [[B:%.*]], <16 x i16> [[C:%.*]], i1 true)
  98 // CHECK-GFX1100-NEXT:    store <16 x i16> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 32, !tbaa [[TBAA4]]
  99 // CHECK-GFX1100-NEXT:    ret void
 100 //
 101 void test_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32(global v16s* out, v16s a, v16s b, v16s c)
 102 {
 103   *out = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_tied_w32(a, b, c, true);
 104 }
 105
 106 //
 107 // amdgcn_wmma_i32_16x16x16_iu8
 108 //
 109
 110 // CHECK-GFX1100-LABEL: @test_amdgcn_wmma_i32_16x16x16_iu8_w32(
 111 // CHECK-GFX1100-NEXT:  entry:
 112 // CHECK-GFX1100-NEXT:    [[TMP0:%.*]] = tail call <8 x i32> @llvm.amdgcn.wmma.i32.16x16x16.iu8.v8i32.v4i32(i1 true, <4 x i32> [[A:%.*]], i1 true, <4 x i32> [[B:%.*]], <8 x i32> [[C:%.*]], i1 false)
 113 // CHECK-GFX1100-NEXT:    store <8 x i32> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 32, !tbaa [[TBAA4]]
 114 // CHECK-GFX1100-NEXT:    ret void
 115 //
 116 void test_amdgcn_wmma_i32_16x16x16_iu8_w32(global v8i* out, v4i a, v4i b, v8i c)
 117 {
 118   *out = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a, true, b, c, false);
 119 }
 120
 121 //
 122 // amdgcn_wmma_i32_16x16x16_iu4
 123 //
 124
 125 // CHECK-GFX1100-LABEL: @test_amdgcn_wmma_i32_16x16x16_iu4_w32(
 126 // CHECK-GFX1100-NEXT:  entry:
 127 // CHECK-GFX1100-NEXT:    [[TMP0:%.*]] = tail call <8 x i32> @llvm.amdgcn.wmma.i32.16x16x16.iu4.v8i32.v2i32(i1 true, <2 x i32> [[A:%.*]], i1 true, <2 x i32> [[B:%.*]], <8 x i32> [[C:%.*]], i1 false)
 128 // CHECK-GFX1100-NEXT:    store <8 x i32> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 32, !tbaa [[TBAA4]]
 129 // CHECK-GFX1100-NEXT:    ret void
 130 //
 131 void test_amdgcn_wmma_i32_16x16x16_iu4_w32(global v8i* out, v2i a, v2i b, v8i c)
 132 {
 133   *out = __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32(true, a, true, b, c, false);
 134 }
 135
 136 #endif