commit 61a8495e1e0485bd40f613004bf29f8a925ab37c
author Andrew Top <aabtop@google.com> Tue Apr 30 15:07:33 2019 -0700
committer Andrew Top <aabtop@google.com> Thu May 02 17:06:51 2019 -0700
tree 8b4860759d8854c2c58b6db4cf3ad2bf3f68aa5c
parent 193dc3d54bfe1fcc7a4d95650f8e9f2c8d9e7d1e

Import Cobalt 19.master.0.203780

Includes the following patches:
  https://cobalt-review.googlesource.com/c/cobalt/+/5210
    by errong.leng@samsung.com
  https://cobalt-review.googlesource.com/c/cobalt/+/5270
    by linus.wang@samsung.com

src/third_party/llvm-project/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err-clamp.cl

diff --git a/src/third_party/llvm-project/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err-clamp.cl b/src/third_party/llvm-project/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err-clamp.cl
new file mode 100644
index 0000000..f8356fe
--- /dev/null
+++ b/src/third_party/llvm-project/clang/test/CodeGenOpenCL/builtins-amdgcn-dl-insts-err-clamp.cl
@@ -0,0 +1,25 @@
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx906 -verify -S -emit-llvm -o - %s
+
+typedef unsigned int uint;
+typedef half __attribute__((ext_vector_type(2))) half2;
+typedef short __attribute__((ext_vector_type(2))) short2;
+typedef unsigned short __attribute__((ext_vector_type(2))) ushort2;
+
+kernel void builtins_amdgcn_dl_insts_err(
+    global float *fOut, global int *siOut, global uint *uiOut,
+    half2 v2hA, half2 v2hB, float fC,
+    short2 v2ssA, short2 v2ssB, int siA, int siB, int siC,
+    ushort2 v2usA, ushort2 v2usB, uint uiA, uint uiB, uint uiC, uint isClamp) {
+  fOut[0] = __builtin_amdgcn_fdot2(v2hA, v2hB, fC, isClamp == 0 ? false : true);     // expected-error {{'__builtin_amdgcn_fdot2' must be a constant integer}}
+
+  siOut[0] = __builtin_amdgcn_sdot2(v2ssA, v2ssB, siC, isClamp == 0 ? false : true); // expected-error {{'__builtin_amdgcn_sdot2' must be a constant integer}}
+  uiOut[0] = __builtin_amdgcn_udot2(v2usA, v2usB, uiC, isClamp == 0 ? false : true); // expected-error {{'__builtin_amdgcn_udot2' must be a constant integer}}
+
+  siOut[1] = __builtin_amdgcn_sdot4(siA, siB, siC, isClamp == 0 ? false : true);     // expected-error {{'__builtin_amdgcn_sdot4' must be a constant integer}}
+  uiOut[1] = __builtin_amdgcn_udot4(uiA, uiB, uiC, isClamp == 0 ? false : true);     // expected-error {{'__builtin_amdgcn_udot4' must be a constant integer}}
+
+  siOut[2] = __builtin_amdgcn_sdot8(siA, siB, siC, isClamp == 0 ? false : true);     // expected-error {{'__builtin_amdgcn_sdot8' must be a constant integer}}
+  uiOut[2] = __builtin_amdgcn_udot8(uiA, uiB, uiC, isClamp == 0 ? false : true);     // expected-error {{'__builtin_amdgcn_udot8' must be a constant integer}}
+}