1// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx908 -verify -S -o - %s 2 3#pragma OPENCL EXTENSION cl_khr_fp64:enable 4 5typedef float v4f __attribute__((ext_vector_type(4))); 6typedef float v16f __attribute__((ext_vector_type(16))); 7typedef float v32f __attribute__((ext_vector_type(32))); 8typedef half v4h __attribute__((ext_vector_type(4))); 9typedef half v16h __attribute__((ext_vector_type(16))); 10typedef half v32h __attribute__((ext_vector_type(32))); 11typedef int v4i __attribute__((ext_vector_type(4))); 12typedef int v16i __attribute__((ext_vector_type(16))); 13typedef int v32i __attribute__((ext_vector_type(32))); 14typedef short v2s __attribute__((ext_vector_type(2))); 15typedef short v4s __attribute__((ext_vector_type(4))); 16typedef short v16s __attribute__((ext_vector_type(16))); 17typedef short v32s __attribute__((ext_vector_type(32))); 18typedef double v4d __attribute__((ext_vector_type(4))); 19 20 21void test_mfma_f32_32x32x1f32(global v32f* out, float a, float b, v32f c, int d) 22{ 23 *out = __builtin_amdgcn_mfma_f32_32x32x1f32(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x1f32' must be a constant integer}} 24 *out = __builtin_amdgcn_mfma_f32_32x32x1f32(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x1f32' must be a constant integer}} 25 *out = __builtin_amdgcn_mfma_f32_32x32x1f32(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x1f32' must be a constant integer}} 26} 27 28void test_mfma_f32_16x16x1f32(global v16f* out, float a, float b, v16f c, int d) 29{ 30 *out = __builtin_amdgcn_mfma_f32_16x16x1f32(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x1f32' must be a constant integer}} 31 *out = __builtin_amdgcn_mfma_f32_16x16x1f32(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x1f32' must be a constant integer}} 32 *out = __builtin_amdgcn_mfma_f32_16x16x1f32(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x1f32' must be a constant integer}} 33} 34 35void test_mfma_f32_4x4x1f32(global v4f* out, float a, float b, v4f c, int d) 36{ 37 *out = __builtin_amdgcn_mfma_f32_4x4x1f32(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_4x4x1f32' must be a constant integer}} 38 *out = __builtin_amdgcn_mfma_f32_4x4x1f32(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_4x4x1f32' must be a constant integer}} 39 *out = __builtin_amdgcn_mfma_f32_4x4x1f32(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_4x4x1f32' must be a constant integer}} 40} 41 42void test_mfma_f32_32x32x2f32(global v16f* out, float a, float b, v16f c, int d) 43{ 44 *out = __builtin_amdgcn_mfma_f32_32x32x2f32(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x2f32' must be a constant integer}} 45 *out = __builtin_amdgcn_mfma_f32_32x32x2f32(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x2f32' must be a constant integer}} 46 *out = __builtin_amdgcn_mfma_f32_32x32x2f32(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x2f32' must be a constant integer}} 47} 48 49void test_mfma_f32_16x16x4f32(global v4f* out, float a, float b, v4f c, int d) 50{ 51 *out = __builtin_amdgcn_mfma_f32_16x16x4f32(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x4f32' must be a constant integer}} 52 *out = __builtin_amdgcn_mfma_f32_16x16x4f32(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x4f32' must be a constant integer}} 53 *out = __builtin_amdgcn_mfma_f32_16x16x4f32(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x4f32' must be a constant integer}} 54} 55 56void test_mfma_f32_32x32x4f16(global v32f* out, v4h a, v4h b, v32f c, int d) 57{ 58 *out = __builtin_amdgcn_mfma_f32_32x32x4f16(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x4f16' must be a constant integer}} 59 *out = __builtin_amdgcn_mfma_f32_32x32x4f16(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x4f16' must be a constant integer}} 60 *out = __builtin_amdgcn_mfma_f32_32x32x4f16(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x4f16' must be a constant integer}} 61} 62 63void test_mfma_f32_16x16x4f16(global v16f* out, v4h a, v4h b, v16f c, int d) 64{ 65 *out = __builtin_amdgcn_mfma_f32_16x16x4f16(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x4f16' must be a constant integer}} 66 *out = __builtin_amdgcn_mfma_f32_16x16x4f16(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x4f16' must be a constant integer}} 67 *out = __builtin_amdgcn_mfma_f32_16x16x4f16(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x4f16' must be a constant integer}} 68} 69 70void test_mfma_f32_4x4x4f16(global v4f* out, v4h a, v4h b, v4f c, int d) 71{ 72 *out = __builtin_amdgcn_mfma_f32_4x4x4f16(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_4x4x4f16' must be a constant integer}} 73 *out = __builtin_amdgcn_mfma_f32_4x4x4f16(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_4x4x4f16' must be a constant integer}} 74 *out = __builtin_amdgcn_mfma_f32_4x4x4f16(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_4x4x4f16' must be a constant integer}} 75} 76 77void test_mfma_f32_32x32x8f16(global v16f* out, v4h a, v4h b, v16f c, int d) 78{ 79 *out = __builtin_amdgcn_mfma_f32_32x32x8f16(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x8f16' must be a constant integer}} 80 *out = __builtin_amdgcn_mfma_f32_32x32x8f16(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x8f16' must be a constant integer}} 81 *out = __builtin_amdgcn_mfma_f32_32x32x8f16(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x8f16' must be a constant integer}} 82} 83 84void test_mfma_f32_16x16x16f16(global v4f* out, v4h a, v4h b, v4f c, int d) 85{ 86 *out = __builtin_amdgcn_mfma_f32_16x16x16f16(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x16f16' must be a constant integer}} 87 *out = __builtin_amdgcn_mfma_f32_16x16x16f16(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x16f16' must be a constant integer}} 88 *out = __builtin_amdgcn_mfma_f32_16x16x16f16(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x16f16' must be a constant integer}} 89} 90 91void test_mfma_i32_32x32x4i8(global v32i* out, int a, int b, v32i c, int d) 92{ 93 *out = __builtin_amdgcn_mfma_i32_32x32x4i8(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_i32_32x32x4i8' must be a constant integer}} 94 *out = __builtin_amdgcn_mfma_i32_32x32x4i8(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_i32_32x32x4i8' must be a constant integer}} 95 *out = __builtin_amdgcn_mfma_i32_32x32x4i8(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_i32_32x32x4i8' must be a constant integer}} 96} 97 98void test_mfma_i32_16x16x4i8(global v16i* out, int a, int b, v16i c, int d) 99{ 100 *out = __builtin_amdgcn_mfma_i32_16x16x4i8(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_i32_16x16x4i8' must be a constant integer}} 101 *out = __builtin_amdgcn_mfma_i32_16x16x4i8(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_i32_16x16x4i8' must be a constant integer}} 102 *out = __builtin_amdgcn_mfma_i32_16x16x4i8(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_i32_16x16x4i8' must be a constant integer}} 103} 104 105void test_mfma_i32_4x4x4i8(global v4i* out, int a, int b, v4i c, int d) 106{ 107 *out = __builtin_amdgcn_mfma_i32_4x4x4i8(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_i32_4x4x4i8' must be a constant integer}} 108 *out = __builtin_amdgcn_mfma_i32_4x4x4i8(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_i32_4x4x4i8' must be a constant integer}} 109 *out = __builtin_amdgcn_mfma_i32_4x4x4i8(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_i32_4x4x4i8' must be a constant integer}} 110} 111 112void test_mfma_i32_32x32x8i8(global v16i* out, int a, int b, v16i c, int d) 113{ 114 *out = __builtin_amdgcn_mfma_i32_32x32x8i8(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_i32_32x32x8i8' must be a constant integer}} 115 *out = __builtin_amdgcn_mfma_i32_32x32x8i8(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_i32_32x32x8i8' must be a constant integer}} 116 *out = __builtin_amdgcn_mfma_i32_32x32x8i8(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_i32_32x32x8i8' must be a constant integer}} 117} 118 119void test_mfma_i32_16x16x16i8(global v4i* out, int a, int b, v4i c, int d) 120{ 121 *out = __builtin_amdgcn_mfma_i32_16x16x16i8(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_i32_16x16x16i8' must be a constant integer}} 122 *out = __builtin_amdgcn_mfma_i32_16x16x16i8(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_i32_16x16x16i8' must be a constant integer}} 123 *out = __builtin_amdgcn_mfma_i32_16x16x16i8(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_i32_16x16x16i8' must be a constant integer}} 124} 125 126void test_mfma_f32_32x32x2bf16(global v32f* out, v2s a, v2s b, v32f c, int d) 127{ 128 *out = __builtin_amdgcn_mfma_f32_32x32x2bf16(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x2bf16' must be a constant integer}} 129 *out = __builtin_amdgcn_mfma_f32_32x32x2bf16(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x2bf16' must be a constant integer}} 130 *out = __builtin_amdgcn_mfma_f32_32x32x2bf16(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x2bf16' must be a constant integer}} 131} 132 133void test_mfma_f32_16x16x2bf16(global v16f* out, v2s a, v2s b, v16f c, int d) 134{ 135 *out = __builtin_amdgcn_mfma_f32_16x16x2bf16(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x2bf16' must be a constant integer}} 136 *out = __builtin_amdgcn_mfma_f32_16x16x2bf16(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x2bf16' must be a constant integer}} 137 *out = __builtin_amdgcn_mfma_f32_16x16x2bf16(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x2bf16' must be a constant integer}} 138} 139 140void test_mfma_f32_4x4x2bf16(global v4f* out, v2s a, v2s b, v4f c, int d) 141{ 142 *out = __builtin_amdgcn_mfma_f32_4x4x2bf16(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_4x4x2bf16' must be a constant integer}} 143 *out = __builtin_amdgcn_mfma_f32_4x4x2bf16(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_4x4x2bf16' must be a constant integer}} 144 *out = __builtin_amdgcn_mfma_f32_4x4x2bf16(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_4x4x2bf16' must be a constant integer}} 145} 146 147void test_mfma_f32_32x32x4bf16(global v16f* out, v2s a, v2s b, v16f c, int d) 148{ 149 *out = __builtin_amdgcn_mfma_f32_32x32x4bf16(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x4bf16' must be a constant integer}} 150 *out = __builtin_amdgcn_mfma_f32_32x32x4bf16(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x4bf16' must be a constant integer}} 151 *out = __builtin_amdgcn_mfma_f32_32x32x4bf16(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_32x32x4bf16' must be a constant integer}} 152} 153 154void test_mfma_f32_16x16x8bf16(global v4f* out, v2s a, v2s b, v4f c, int d) 155{ 156 *out = __builtin_amdgcn_mfma_f32_16x16x8bf16(a, b, c, d, 0, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x8bf16' must be a constant integer}} 157 *out = __builtin_amdgcn_mfma_f32_16x16x8bf16(a, b, c, 0, d, 0); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x8bf16' must be a constant integer}} 158 *out = __builtin_amdgcn_mfma_f32_16x16x8bf16(a, b, c, 0, 0, d); // expected-error{{argument to '__builtin_amdgcn_mfma_f32_16x16x8bf16' must be a constant integer}} 159} 160