cross-project-tests/amdgpu/builtins-amdgcn-wmma-w32.cl

   1 // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -target-feature +wavefrontsize32 -DWMMA_GFX1100_TESTS -S -o - %s | FileCheck %s --check-prefix=CHECK-GFX1100
   2
   3 typedef float  v4f   __attribute__((ext_vector_type(4)));
   4 typedef float  v8f   __attribute__((ext_vector_type(8)));
   5 typedef half   v16h  __attribute__((ext_vector_type(16)));
   6 typedef int    v2i   __attribute__((ext_vector_type(2)));
   7 typedef int    v4i   __attribute__((ext_vector_type(4)));
   8 typedef int    v8i   __attribute__((ext_vector_type(8)));
   9 typedef short  v16s  __attribute__((ext_vector_type(16)));
  10
  11 #ifdef WMMA_GFX1100_TESTS
  12
  13 // Wave32
  14
  15
  16 // CHECK-GFX1100-LABEL: test_amdgcn_wmma_f32_16x16x16_f16_w32:
  17 // CHECK-GFX1100: v_wmma_f32_16x16x16_f16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}]
  18 //
  19 void test_amdgcn_wmma_f32_16x16x16_f16_w32(global v8f* out, v16h a, v16h b, v8f c)
  20 {
  21   *out = __builtin_amdgcn_wmma_f32_16x16x16_f16_w32(a, b, c);
  22 }
  23
  24
  25 // CHECK-GFX1100-LABEL: test_amdgcn_wmma_f32_16x16x16_bf16_w32:
  26 // CHECK-GFX1100: v_wmma_f32_16x16x16_bf16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}]
  27 //
  28 void test_amdgcn_wmma_f32_16x16x16_bf16_w32(global v8f* out, v16s a, v16s b, v8f c)
  29 {
  30   *out = __builtin_amdgcn_wmma_f32_16x16x16_bf16_w32(a, b, c);
  31 }
  32
  33
  34 // CHECK-GFX1100-LABEL: test_amdgcn_wmma_f16_16x16x16_f16_w32:
  35 // CHECK-GFX1100: v_wmma_f16_16x16x16_f16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}] op_sel:[0,0,1]
  36 //
  37 void test_amdgcn_wmma_f16_16x16x16_f16_w32(global v16h* out, v16h a, v16h b, v16h c)
  38 {
  39   *out = __builtin_amdgcn_wmma_f16_16x16x16_f16_w32(a, b, c, true);
  40 }
  41
  42
  43 // CHECK-GFX1100-LABEL: test_amdgcn_wmma_bf16_16x16x16_bf16_w32:
  44 // CHECK-GFX1100: v_wmma_bf16_16x16x16_bf16 v[{{.*}}], v[{{.*}} v[{{.*}}], v[{{.*}}] op_sel:[0,0,1]
  45 //
  46 void test_amdgcn_wmma_bf16_16x16x16_bf16_w32(global v16s* out, v16s a, v16s b, v16s c)
  47 {
  48   *out = __builtin_amdgcn_wmma_bf16_16x16x16_bf16_w32(a, b, c, true);
  49 }
  50
  51
  52 // CHECK-GFX1100-LABEL: test_amdgcn_wmma_i32_16x16x16_iu8_w32:
  53 // CHECK-GFX1100: v_wmma_i32_16x16x16_iu8 v[{{.*}}], v[{{.*}} v[{{.*}} v[{{.*}}] neg_lo:[1,1,0]
  54 //
  55 void test_amdgcn_wmma_i32_16x16x16_iu8_w32(global v8i* out, v4i a, v4i b, v8i c)
  56 {
  57   *out = __builtin_amdgcn_wmma_i32_16x16x16_iu8_w32(true, a, true, b, c, false);
  58 }
  59
  60
  61 // CHECK-GFX1100-LABEL: test_amdgcn_wmma_i32_16x16x16_iu4_w32:
  62 // CHECK-GFX1100: v_wmma_i32_16x16x16_iu4 v[{{.*}}, v[{{.*}} v[{{.*}} v[{{.*}} neg_lo:[1,1,0]
  63 void test_amdgcn_wmma_i32_16x16x16_iu4_w32(global v8i* out, v2i a, v2i b, v8i c)
  64 {
  65   *out = __builtin_amdgcn_wmma_i32_16x16x16_iu4_w32(true, a, true, b, c, false);
  66 }
  67
  68 #endif