clang/test/SemaOpenCL/builtins-amdgcn-error-gfx90a-param.cl

   1 // RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx90a -verify -S -o - %s
   2
   3 #pragma OPENCL EXTENSION cl_khr_fp64:enable
   4
   5 typedef float  v4f   __attribute__((ext_vector_type(4)));
   6 typedef float  v16f  __attribute__((ext_vector_type(16)));
   7 typedef float  v32f  __attribute__((ext_vector_type(32)));
   8 typedef half   v4h   __attribute__((ext_vector_type(4)));
   9 typedef half   v16h  __attribute__((ext_vector_type(16)));
  10 typedef half   v32h  __attribute__((ext_vector_type(32)));
  11 typedef int    v4i   __attribute__((ext_vector_type(4)));
  12 typedef int    v16i  __attribute__((ext_vector_type(16)));
  13 typedef int    v32i  __attribute__((ext_vector_type(32)));
  14 typedef short  v2s   __attribute__((ext_vector_type(2)));
  15 typedef short  v4s   __attribute__((ext_vector_type(4)));
  16 typedef short  v16s  __attribute__((ext_vector_type(16)));
  17 typedef short  v32s  __attribute__((ext_vector_type(32)));
  18 typedef double v4d   __attribute__((ext_vector_type(4)));
  19
  20 void test_mfma_f32_32x32x4bf16_1k(global v32f* out, v4s a, v4s b, v32f c, int d)
  21 {
  22   *out = __builtin_amdgcn_mfma_f32_32x32x4bf16_1k(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x4bf16_1k' must be a constant integer}}
  23   *out = __builtin_amdgcn_mfma_f32_32x32x4bf16_1k(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x4bf16_1k' must be a constant integer}}
  24   *out = __builtin_amdgcn_mfma_f32_32x32x4bf16_1k(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x4bf16_1k' must be a constant integer}}
  25 }
  26
  27 void test_mfma_f32_16x16x4bf16_1k(global v16f* out, v4s a, v4s b, v16f c, int d)
  28 {
  29   *out = __builtin_amdgcn_mfma_f32_16x16x4bf16_1k(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x4bf16_1k' must be a constant integer}}
  30   *out = __builtin_amdgcn_mfma_f32_16x16x4bf16_1k(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x4bf16_1k' must be a constant integer}}
  31   *out = __builtin_amdgcn_mfma_f32_16x16x4bf16_1k(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x4bf16_1k' must be a constant integer}}
  32 }
  33
  34 void test_mfma_f32_4x4x4bf16_1k(global v4f* out, v4s a, v4s b, v4f c, int d)
  35 {
  36   *out = __builtin_amdgcn_mfma_f32_4x4x4bf16_1k(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_4x4x4bf16_1k' must be a constant integer}}
  37   *out = __builtin_amdgcn_mfma_f32_4x4x4bf16_1k(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_4x4x4bf16_1k' must be a constant integer}}
  38   *out = __builtin_amdgcn_mfma_f32_4x4x4bf16_1k(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_4x4x4bf16_1k' must be a constant integer}}
  39 }
  40
  41 void test_mfma_f32_32x32x8bf16_1k(global v16f* out, v4s a, v4s b, v16f c, int d)
  42 {
  43   *out = __builtin_amdgcn_mfma_f32_32x32x8bf16_1k(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x8bf16_1k' must be a constant integer}}
  44   *out = __builtin_amdgcn_mfma_f32_32x32x8bf16_1k(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x8bf16_1k' must be a constant integer}}
  45   *out = __builtin_amdgcn_mfma_f32_32x32x8bf16_1k(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x8bf16_1k' must be a constant integer}}
  46 }
  47
  48 void test_mfma_f32_16x16x16bf16_1k(global v4f* out, v4s a, v4s b, v4f c, int d)
  49 {
  50   *out = __builtin_amdgcn_mfma_f32_16x16x16bf16_1k(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x16bf16_1k' must be a constant integer}}
  51   *out = __builtin_amdgcn_mfma_f32_16x16x16bf16_1k(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x16bf16_1k' must be a constant integer}}
  52   *out = __builtin_amdgcn_mfma_f32_16x16x16bf16_1k(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x16bf16_1k' must be a constant integer}}
  53 }
  54
  55 void test_mfma_f64_16x16x4f64(global v4d* out, double a, double b, v4d c, int d)
  56 {
  57   *out = __builtin_amdgcn_mfma_f64_16x16x4f64(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f64_16x16x4f64' must be a constant integer}}
  58   *out = __builtin_amdgcn_mfma_f64_16x16x4f64(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f64_16x16x4f64' must be a constant integer}}
  59   *out = __builtin_amdgcn_mfma_f64_16x16x4f64(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f64_16x16x4f64' must be a constant integer}}
  60 }
  61
  62 void test_mfma_f64_4x4x4f64(global double* out, double a, double b, double c, int d)
  63 {
  64   *out = __builtin_amdgcn_mfma_f64_4x4x4f64(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f64_4x4x4f64' must be a constant integer}}
  65   *out = __builtin_amdgcn_mfma_f64_4x4x4f64(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f64_4x4x4f64' must be a constant integer}}
  66   *out = __builtin_amdgcn_mfma_f64_4x4x4f64(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f64_4x4x4f64' must be a constant integer}}
  67 }