clang/test/CodeGenOpenCL/builtins-amdgcn-wmma-w64.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 -S -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-GFX1100
   4
   5 typedef float  v4f   __attribute__((ext_vector_type(4)));
   6 typedef float  v8f   __attribute__((ext_vector_type(8)));
   7 typedef half   v8h   __attribute__((ext_vector_type(8)));
   8 typedef half   v16h  __attribute__((ext_vector_type(16)));
   9 typedef int    v2i   __attribute__((ext_vector_type(2)));
  10 typedef int    v4i   __attribute__((ext_vector_type(4)));
  11 typedef int    v8i   __attribute__((ext_vector_type(8)));
  12 typedef short  v8s   __attribute__((ext_vector_type(8)));
  13 typedef short  v16s  __attribute__((ext_vector_type(16)));
  14
  15 #ifdef WMMA_GFX1100_TESTS
  16
  17 // Wave64
  18
  19 //
  20 // amdgcn_wmma_f32_16x16x16_f16
  21 //
  22
  23 // CHECK-GFX1100-LABEL: @test_amdgcn_wmma_f32_16x16x16_f16_w64(
  24 // CHECK-GFX1100-NEXT:  entry:
  25 // CHECK-GFX1100-NEXT:    [[TMP0:%.*]] = tail call <4 x float> @llvm.amdgcn.wmma.f32.16x16x16.f16.v4f32(<16 x half> [[A:%.*]], <16 x half> [[B:%.*]], <4 x float> [[C:%.*]])
  26 // CHECK-GFX1100-NEXT:    store <4 x float> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 16, !tbaa [[TBAA4:![0-9]+]]
  27 // CHECK-GFX1100-NEXT:    ret void
  28 //
  29 void test_amdgcn_wmma_f32_16x16x16_f16_w64(global v4f* out, v16h a, v16h b, v4f c)
  30 {
  31   *out = __builtin_amdgcn_wmma_f32_16x16x16_f16_w64(a, b, c);
  32 }
  33
  34 //
  35 // amdgcn_wmma_f32_16x16x16_bf16
  36 //
  37
  38 // CHECK-GFX1100-LABEL: @test_amdgcn_wmma_f32_16x16x16_bf16_w64(
  39 // CHECK-GFX1100-NEXT:  entry:
  40 // CHECK-GFX1100-NEXT:    [[TMP0:%.*]] = tail call <4 x float> @llvm.amdgcn.wmma.f32.16x16x16.bf16.v4f32(<16 x i16> [[A:%.*]], <16 x i16> [[B:%.*]], <4 x float> [[C:%.*]])
  41 // CHECK-GFX1100-NEXT:    store <4 x float> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 16, !tbaa [[TBAA4]]
  42 // CHECK-GFX1100-NEXT:    ret void
  43 //
  44 void test_amdgcn_wmma_f32_16x16x16_bf16_w64(global v4f* out, v16s a, v16s b, v4f c)
  45 {
  46   *out = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w64(a, b, c);
  47 }
  48
  49 //
  50 // amdgcn_wmma_f16_16x16x16_f16
  51 //
  52
  53 // CHECK-GFX1100-LABEL: @test_amdgcn_wmma_f16_16x16x16_f16_w64(
  54 // CHECK-GFX1100-NEXT:  entry:
  55 // CHECK-GFX1100-NEXT:    [[TMP0:%.*]] = tail call <8 x half> @llvm.amdgcn.wmma.f16.16x16x16.f16.v8f16(<16 x half> [[A:%.*]], <16 x half> [[B:%.*]], <8 x half> [[C:%.*]], i1 true)
  56 // CHECK-GFX1100-NEXT:    store <8 x half> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 16, !tbaa [[TBAA4]]
  57 // CHECK-GFX1100-NEXT:    ret void
  58 //
  59 void test_amdgcn_wmma_f16_16x16x16_f16_w64(global v8h* out, v16h a, v16h b, v8h c)
  60 {
  61   *out = __builtin_amdgcn_wmma_f16_16x16x16_f16_w64(a, b, c, true);
  62 }
  63
  64 //
  65 // amdgcn_wmma_bf16_16x16x16_bf16
  66 //
  67
  68 // CHECK-GFX1100-LABEL: @test_amdgcn_wmma_bf16_16x16x16_bf16_w64(
  69 // CHECK-GFX1100-NEXT:  entry:
  70 // CHECK-GFX1100-NEXT:    [[TMP0:%.*]] = tail call <8 x i16> @llvm.amdgcn.wmma.bf16.16x16x16.bf16.v8i16(<16 x i16> [[A:%.*]], <16 x i16> [[B:%.*]], <8 x i16> [[C:%.*]], i1 true)
  71 // CHECK-GFX1100-NEXT:    store <8 x i16> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 16, !tbaa [[TBAA4]]
  72 // CHECK-GFX1100-NEXT:    ret void
  73 //
  74 void test_amdgcn_wmma_bf16_16x16x16_bf16_w64(global v8s* out, v16s a, v16s b, v8s c)
  75 {
  76   *out = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w64(a, b, c, true);
  77 }
  78
  79 //
  80 // amdgcn_wmma_i32_16x16x16_iu8
  81 //
  82
  83 // CHECK-GFX1100-LABEL: @test_amdgcn_wmma_i32_16x16x16_iu8_w64(
  84 // CHECK-GFX1100-NEXT:  entry:
  85 // CHECK-GFX1100-NEXT:    [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.wmma.i32.16x16x16.iu8.v4i32(i1 true, <4 x i32> [[A:%.*]], i1 true, <4 x i32> [[B:%.*]], <4 x i32> [[C:%.*]], i1 false)
  86 // CHECK-GFX1100-NEXT:    store <4 x i32> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 16, !tbaa [[TBAA4]]
  87 // CHECK-GFX1100-NEXT:    ret void
  88 //
  89 void test_amdgcn_wmma_i32_16x16x16_iu8_w64(global v4i* out, v4i a, v4i b, v4i c)
  90 {
  91   *out = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w64(true, a, true, b, c, false);
  92 }
  93
  94 //
  95 // amdgcn_wmma_i32_16x16x16_iu4
  96 //
  97
  98 // CHECK-GFX1100-LABEL: @test_amdgcn_wmma_i32_16x16x16_iu4_w64(
  99 // CHECK-GFX1100-NEXT:  entry:
 100 // CHECK-GFX1100-NEXT:    [[TMP0:%.*]] = tail call <4 x i32> @llvm.amdgcn.wmma.i32.16x16x16.iu4.v4i32(i1 true, <2 x i32> [[A:%.*]], i1 true, <2 x i32> [[B:%.*]], <4 x i32> [[C:%.*]], i1 false)
 101 // CHECK-GFX1100-NEXT:    store <4 x i32> [[TMP0]], ptr addrspace(1) [[OUT:%.*]], align 16, !tbaa [[TBAA4]]
 102 // CHECK-GFX1100-NEXT:    ret void
 103 //
 104 void test_amdgcn_wmma_i32_16x16x16_iu4_w64(global v4i* out, v2i a, v2i b, v4i c)
 105 {
 106   *out = __builtin_amdgcn_wmma_i32_16x16x16_iu4_w64(true, a, true, b, c, false);
 107 }
 108
 109 #endif