SWDEV-179954 - OpenCL/LC - Merge branch amd-master into amd-common

Change-Id: I2ecbb97ee0ed8fe015c92c25ebd3c213cb21f11d

include/clang/Basic/BuiltinsAMDGPU.def

@@ -108,6 +108,16 @@ BUILTIN(__builtin_amdgcn_ds_fminf, "ff*3fIiIiIb", "n")
 BUILTIN(__builtin_amdgcn_ds_fmaxf, "ff*3fIiIiIb", "n")
 BUILTIN(__builtin_amdgcn_ds_append, "ii*3", "n")
 BUILTIN(__builtin_amdgcn_ds_consume, "ii*3", "n")
+BUILTIN(__builtin_amdgcn_alignbit, "UiUiUiUi", "nc")
+BUILTIN(__builtin_amdgcn_alignbyte, "UiUiUiUi", "nc")
+BUILTIN(__builtin_amdgcn_ubfe, "UiUiUiUi", "nc")
+BUILTIN(__builtin_amdgcn_sbfe, "UiUiUiUi", "nc")
+BUILTIN(__builtin_amdgcn_cvt_pkrtz, "E2hff", "nc")
+BUILTIN(__builtin_amdgcn_cvt_pknorm_i16, "E2sff", "nc")
+BUILTIN(__builtin_amdgcn_cvt_pknorm_u16, "E2Usff", "nc")
+BUILTIN(__builtin_amdgcn_cvt_pk_i16, "E2sii", "nc")
+BUILTIN(__builtin_amdgcn_cvt_pk_u16, "E2UsUiUi", "nc")
+BUILTIN(__builtin_amdgcn_cvt_pk_u8_f32, "UifUiUi", "nc")
 
 //===----------------------------------------------------------------------===//
 // CI+ only builtins.
@@ -162,6 +172,13 @@ TARGET_BUILTIN(__builtin_amdgcn_udot4, "UiUiUiUiIb", "nc", "dot2-insts")
 TARGET_BUILTIN(__builtin_amdgcn_sdot8, "SiSiSiSiIb", "nc", "dot1-insts")
 TARGET_BUILTIN(__builtin_amdgcn_udot8, "UiUiUiUiIb", "nc", "dot2-insts")
 
+//===----------------------------------------------------------------------===//
+// GFX10+ only builtins.
+//===----------------------------------------------------------------------===//
+TARGET_BUILTIN(__builtin_amdgcn_permlane16, "UiUiUiUiUiIbIb", "nc", "gfx10-insts")
+TARGET_BUILTIN(__builtin_amdgcn_permlanex16, "UiUiUiUiUiIbIb", "nc", "gfx10-insts")
+TARGET_BUILTIN(__builtin_amdgcn_mov_dpp8, "UiUiIUi", "nc", "gfx10-insts")
+
 //===----------------------------------------------------------------------===//
 // Special builtins.
 //===----------------------------------------------------------------------===//

include/clang/Basic/BuiltinsWebAssembly.def

@@ -29,6 +29,9 @@ BUILTIN(__builtin_wasm_memory_grow, "zIiz", "n")
 TARGET_BUILTIN(__builtin_wasm_memory_init, "vIUiIUiv*UiUi", "", "bulk-memory")
 TARGET_BUILTIN(__builtin_wasm_data_drop, "vIUi", "", "bulk-memory")
 
+// Thread-local storage
+TARGET_BUILTIN(__builtin_wasm_tls_size, "z", "nc", "bulk-memory")
+
 // Floating point min/max
 BUILTIN(__builtin_wasm_min_f32, "fff", "nc")
 BUILTIN(__builtin_wasm_max_f32, "fff", "nc")

lib/CodeGen/CGBuiltin.cpp

@@ -12679,6 +12679,8 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
 
   case AMDGPU::BI__builtin_amdgcn_ds_swizzle:
     return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_ds_swizzle);
+  case AMDGPU::BI__builtin_amdgcn_mov_dpp8:
+    return emitBinaryBuiltin(*this, E, Intrinsic::amdgcn_mov_dpp8);
   case AMDGPU::BI__builtin_amdgcn_mov_dpp:
   case AMDGPU::BI__builtin_amdgcn_update_dpp: {
     llvm::SmallVector<llvm::Value *, 6> Args;
@@ -12744,6 +12746,10 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_fract);
   case AMDGPU::BI__builtin_amdgcn_lerp:
     return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_lerp);
+  case AMDGPU::BI__builtin_amdgcn_ubfe:
+    return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_ubfe);
+  case AMDGPU::BI__builtin_amdgcn_sbfe:
+    return emitTernaryBuiltin(*this, E, Intrinsic::amdgcn_sbfe);
   case AMDGPU::BI__builtin_amdgcn_uicmp:
   case AMDGPU::BI__builtin_amdgcn_uicmpl:
   case AMDGPU::BI__builtin_amdgcn_sicmp:
@@ -13913,6 +13919,11 @@ Value *CodeGenFunction::EmitWebAssemblyBuiltinExpr(unsigned BuiltinID,
     Function *Callee = CGM.getIntrinsic(Intrinsic::wasm_data_drop);
     return Builder.CreateCall(Callee, {Arg});
   }
+  case WebAssembly::BI__builtin_wasm_tls_size: {
+    llvm::Type *ResultType = ConvertType(E->getType());
+    Function *Callee = CGM.getIntrinsic(Intrinsic::wasm_tls_size, ResultType);
+    return Builder.CreateCall(Callee);
+  }
   case WebAssembly::BI__builtin_wasm_throw: {
     Value *Tag = EmitScalarExpr(E->getArg(0));
     Value *Obj = EmitScalarExpr(E->getArg(1));

lib/CodeGen/CGExpr.cpp

@@ -3892,6 +3892,23 @@ LValue CodeGenFunction::EmitLValueForLambdaField(const FieldDecl *Field) {
   return EmitLValueForField(LambdaLV, Field);
 }
 
+/// Get the field index in the debug info. The debug info structure/union
+/// will ignore the unnamed bitfields.
+unsigned CodeGenFunction::getDebugInfoFIndex(const RecordDecl *Rec,
+                                             unsigned FieldIndex) {
+  unsigned I = 0, Skipped = 0;
+
+  for (auto F : Rec->getDefinition()->fields()) {
+    if (I == FieldIndex)
+      break;
+    if (F->isUnnamedBitfield())
+      Skipped++;
+    I++;
+  }
+
+  return FieldIndex - Skipped;
+}
+
 /// Get the address of a zero-sized field within a record. The resulting
 /// address doesn't necessarily have the right type.
 static Address emitAddrOfZeroSizeField(CodeGenFunction &CGF, Address Base,
@@ -3931,7 +3948,7 @@ static Address emitPreserveStructAccess(CodeGenFunction &CGF, Address base,
       CGF.CGM.getTypes().getCGRecordLayout(rec).getLLVMFieldNo(field);
 
   return CGF.Builder.CreatePreserveStructAccessIndex(
-      base, idx, field->getFieldIndex(), DbgInfo);
+      base, idx, CGF.getDebugInfoFIndex(rec, field->getFieldIndex()), DbgInfo);
 }
 
 static bool hasAnyVptr(const QualType Type, const ASTContext &Context) {
@@ -4048,7 +4065,7 @@ LValue CodeGenFunction::EmitLValueForField(LValue base,
           getContext().getRecordType(rec), rec->getLocation());
       addr = Address(
           Builder.CreatePreserveUnionAccessIndex(
-              addr.getPointer(), field->getFieldIndex(), DbgInfo),
+              addr.getPointer(), getDebugInfoFIndex(rec, field->getFieldIndex()), DbgInfo),
           addr.getAlignment());
     }
   } else {

lib/CodeGen/CodeGenFunction.h

@@ -2652,6 +2652,9 @@ class CodeGenFunction : public CodeGenTypeCache {
   /// Converts Location to a DebugLoc, if debug information is enabled.
   llvm::DebugLoc SourceLocToDebugLoc(SourceLocation Location);
 
+  /// Get the record field index as represented in debug info.
+  unsigned getDebugInfoFIndex(const RecordDecl *Rec, unsigned FieldIndex);
+
 
   //===--------------------------------------------------------------------===//
   //                            Declaration Emission

lib/Driver/ToolChains/Arch/X86.cpp

@@ -135,7 +135,7 @@ void x86::getX86TargetFeatures(const Driver &D, const llvm::Triple &Triple,
     if (ArchType == llvm::Triple::x86_64) {
       Features.push_back("+sse4.2");
       Features.push_back("+popcnt");
-      Features.push_back("+mcx16");
+      Features.push_back("+cx16");
     } else
       Features.push_back("+ssse3");
   }

lib/Format/TokenAnnotator.cpp

@@ -388,6 +388,10 @@ class AnnotatingParser {
   bool isCpp11AttributeSpecifier(const FormatToken &Tok) {
     if (!Style.isCpp() || !Tok.startsSequence(tok::l_square, tok::l_square))
       return false;
+    // The first square bracket is part of an ObjC array literal
+    if (Tok.Previous && Tok.Previous->is(tok::at)) {
+      return false;
+    }
     const FormatToken *AttrTok = Tok.Next->Next;
     if (!AttrTok)
       return false;
@@ -400,7 +404,7 @@ class AnnotatingParser {
     while (AttrTok && !AttrTok->startsSequence(tok::r_square, tok::r_square)) {
       // ObjC message send. We assume nobody will use : in a C++11 attribute
       // specifier parameter, although this is technically valid:
-      // [[foo(:)]]
+      // [[foo(:)]].
       if (AttrTok->is(tok::colon) ||
           AttrTok->startsSequence(tok::identifier, tok::identifier) ||
           AttrTok->startsSequence(tok::r_paren, tok::identifier))

test/CodeGen/builtin-preserve-access-index.c

@@ -0,0 +1,177 @@
+// RUN: %clang -target x86_64 -emit-llvm -S -g %s -o - | FileCheck %s
+
+#define _(x) (__builtin_preserve_access_index(x))
+
+const void *unit1(const void *arg) {
+  return _(arg);
+}
+// CHECK: define dso_local i8* @unit1
+// CHECK-NOT: llvm.preserve.array.access.index
+// CHECK-NOT: llvm.preserve.struct.access.index
+// CHECK-NOT: llvm.preserve.union.access.index
+
+const void *unit2(void) {
+  return _((const void *)0xffffffffFFFF0000ULL);
+}
+// CHECK: define dso_local i8* @unit2
+// CHECK-NOT: llvm.preserve.array.access.index
+// CHECK-NOT: llvm.preserve.struct.access.index
+// CHECK-NOT: llvm.preserve.union.access.index
+
+const void *unit3(const int *arg) {
+  return _(arg + 1);
+}
+// CHECK: define dso_local i8* @unit3
+// CHECK-NOT: llvm.preserve.array.access.index
+// CHECK-NOT: llvm.preserve.struct.access.index
+// CHECK-NOT: llvm.preserve.union.access.index
+
+const void *unit4(const int *arg) {
+  return _(&arg[1]);
+}
+// CHECK: define dso_local i8* @unit4
+// CHECK-NOT: getelementptr
+// CHECK: call i32* @llvm.preserve.array.access.index.p0i32.p0i32(i32* %{{[0-9a-z]+}}, i32 0, i32 1)
+
+const void *unit5(const int *arg[5]) {
+  return _(&arg[1][2]);
+}
+// CHECK: define dso_local i8* @unit5
+// CHECK-NOT: getelementptr
+// CHECK: call i32** @llvm.preserve.array.access.index.p0p0i32.p0p0i32(i32** %{{[0-9a-z]+}}, i32 0, i32 1)
+// CHECK-NOT: getelementptr
+// CHECK: call i32* @llvm.preserve.array.access.index.p0i32.p0i32(i32* %{{[0-9a-z]+}}, i32 0, i32 2)
+
+struct s1 {
+  char a;
+  int b;
+};
+
+struct s2 {
+  char a1:1;
+  char a2:1;
+  int b;
+};
+
+struct s3 {
+  char a1:1;
+  char a2:1;
+  char :6;
+  int b;
+};
+
+const void *unit6(struct s1 *arg) {
+  return _(&arg->a);
+}
+// CHECK: define dso_local i8* @unit6
+// CHECK-NOT: getelementptr
+// CHECK: call i8* @llvm.preserve.struct.access.index.p0i8.p0s_struct.s1s(%struct.s1* %{{[0-9a-z]+}}, i32 0, i32 0), !dbg !{{[0-9]+}}, !llvm.preserve.access.index ![[STRUCT_S1:[0-9]+]]
+
+const void *unit7(struct s1 *arg) {
+  return _(&arg->b);
+}
+// CHECK: define dso_local i8* @unit7
+// CHECK-NOT: getelementptr
+// CHECK: call i32* @llvm.preserve.struct.access.index.p0i32.p0s_struct.s1s(%struct.s1* %{{[0-9a-z]+}}, i32 1, i32 1), !dbg !{{[0-9]+}}, !llvm.preserve.access.index ![[STRUCT_S1]]
+
+const void *unit8(struct s2 *arg) {
+  return _(&arg->b);
+}
+// CHECK: define dso_local i8* @unit8
+// CHECK-NOT: getelementptr
+// CHECK: call i32* @llvm.preserve.struct.access.index.p0i32.p0s_struct.s2s(%struct.s2* %{{[0-9a-z]+}}, i32 1, i32 2), !dbg !{{[0-9]+}}, !llvm.preserve.access.index ![[STRUCT_S2:[0-9]+]]
+
+const void *unit9(struct s3 *arg) {
+  return _(&arg->b);
+}
+// CHECK: define dso_local i8* @unit9
+// CHECK-NOT: getelementptr
+// CHECK: call i32* @llvm.preserve.struct.access.index.p0i32.p0s_struct.s3s(%struct.s3* %{{[0-9a-z]+}}, i32 1, i32 2), !dbg !{{[0-9]+}}, !llvm.preserve.access.index ![[STRUCT_S3:[0-9]+]]
+
+union u1 {
+  char a;
+  int b;
+};
+
+union u2 {
+  char a;
+  int :32;
+  int b;
+};
+
+const void *unit10(union u1 *arg) {
+  return _(&arg->a);
+}
+// CHECK: define dso_local i8* @unit10
+// CHECK-NOT: getelementptr
+// CHECK: call %union.u1* @llvm.preserve.union.access.index.p0s_union.u1s.p0s_union.u1s(%union.u1* %{{[0-9a-z]+}}, i32 0), !dbg !{{[0-9]+}}, !llvm.preserve.access.index ![[UNION_U1:[0-9]+]]
+
+const void *unit11(union u1 *arg) {
+  return _(&arg->b);
+}
+// CHECK: define dso_local i8* @unit11
+// CHECK-NOT: getelementptr
+// CHECK: call %union.u1* @llvm.preserve.union.access.index.p0s_union.u1s.p0s_union.u1s(%union.u1* %{{[0-9a-z]+}}, i32 1), !dbg !{{[0-9]+}}, !llvm.preserve.access.index ![[UNION_U1]]
+
+const void *unit12(union u2 *arg) {
+  return _(&arg->b);
+}
+// CHECK: define dso_local i8* @unit12
+// CHECK-NOT: getelementptr
+// CHECK: call %union.u2* @llvm.preserve.union.access.index.p0s_union.u2s.p0s_union.u2s(%union.u2* %{{[0-9a-z]+}}, i32 1), !dbg !{{[0-9]+}}, !llvm.preserve.access.index ![[UNION_U2:[0-9]+]]
+
+struct s4 {
+  char d;
+  union u {
+    int b[4];
+    char a;
+  } c;
+};
+
+union u3 {
+  struct s {
+    int b[4];
+  } c;
+  char a;
+};
+
+const void *unit13(struct s4 *arg) {
+  return _(&arg->c.b[2]);
+}
+// CHECK: define dso_local i8* @unit13
+// CHECK: call %union.u* @llvm.preserve.struct.access.index.p0s_union.us.p0s_struct.s4s(%struct.s4* %{{[0-9a-z]+}}, i32 1, i32 1), !dbg !{{[0-9]+}}, !llvm.preserve.access.index ![[STRUCT_S4:[0-9]+]]
+// CHECK: call %union.u* @llvm.preserve.union.access.index.p0s_union.us.p0s_union.us(%union.u* %{{[0-9a-z]+}}, i32 0), !dbg !{{[0-9]+}}, !llvm.preserve.access.index ![[UNION_I_U:[0-9]+]]
+// CHECK: call i32* @llvm.preserve.array.access.index.p0i32.p0a4i32([4 x i32]* %{{[0-9a-z]+}}, i32 1, i32 2)
+
+const void *unit14(union u3 *arg) {
+  return _(&arg->c.b[2]);
+}
+// CHECK: define dso_local i8* @unit14
+// CHECK: call %union.u3* @llvm.preserve.union.access.index.p0s_union.u3s.p0s_union.u3s(%union.u3* %{{[0-9a-z]+}}, i32 0), !dbg !{{[0-9]+}}, !llvm.preserve.access.index ![[UNION_U3:[0-9]+]]
+// CHECK: call [4 x i32]* @llvm.preserve.struct.access.index.p0a4i32.p0s_struct.ss(%struct.s* %{{[0-9a-z]+}}, i32 0, i32 0), !dbg !{{[0-9]+}}, !llvm.preserve.access.index ![[STRUCT_I_S:[0-9]+]]
+// CHECK: call i32* @llvm.preserve.array.access.index.p0i32.p0a4i32([4 x i32]* %{{[0-9a-z]+}}, i32 1, i32 2)
+
+const void *unit15(struct s4 *arg) {
+  return _(&arg[2].c.a);
+}
+// CHECK: define dso_local i8* @unit15
+// CHECK: call %struct.s4* @llvm.preserve.array.access.index.p0s_struct.s4s.p0s_struct.s4s(%struct.s4* %{{[0-9a-z]+}}, i32 0, i32 2)
+// CHECK: call %union.u* @llvm.preserve.struct.access.index.p0s_union.us.p0s_struct.s4s(%struct.s4* %{{[0-9a-z]+}}, i32 1, i32 1), !dbg !{{[0-9]+}}, !llvm.preserve.access.index ![[STRUCT_S4]]
+// CHECK: call %union.u* @llvm.preserve.union.access.index.p0s_union.us.p0s_union.us(%union.u* %{{[0-9a-z]+}}, i32 1), !dbg !{{[0-9]+}}, !llvm.preserve.access.index ![[UNION_I_U]]
+
+const void *unit16(union u3 *arg) {
+  return _(&arg[2].a);
+}
+// CHECK: define dso_local i8* @unit16
+// CHECK: call %union.u3* @llvm.preserve.array.access.index.p0s_union.u3s.p0s_union.u3s(%union.u3* %{{[0-9a-z]+}}, i32 0, i32 2)
+// CHECK: call %union.u3* @llvm.preserve.union.access.index.p0s_union.u3s.p0s_union.u3s(%union.u3* %{{[0-9a-z]+}}, i32 1), !dbg !{{[0-9]+}}, !llvm.preserve.access.index ![[UNION_U3]]
+
+// CHECK: ![[STRUCT_S1]] = distinct !DICompositeType(tag: DW_TAG_structure_type, name: "s1"
+// CHECK: ![[STRUCT_S2]] = distinct !DICompositeType(tag: DW_TAG_structure_type, name: "s2"
+// CHECK: ![[STRUCT_S3]] = distinct !DICompositeType(tag: DW_TAG_structure_type, name: "s3"
+// CHECK: ![[UNION_U1]] = distinct !DICompositeType(tag: DW_TAG_union_type, name: "u1"
+// CHECK: ![[UNION_U2]] = distinct !DICompositeType(tag: DW_TAG_union_type, name: "u2"
+// CHECK: ![[STRUCT_S4]] = distinct !DICompositeType(tag: DW_TAG_structure_type, name: "s4"
+// CHECK: ![[UNION_I_U]] = distinct !DICompositeType(tag: DW_TAG_union_type, name: "u"
+// CHECK: ![[UNION_U3]] = distinct !DICompositeType(tag: DW_TAG_union_type, name: "u3"
+// CHECK: ![[STRUCT_I_S]] = distinct !DICompositeType(tag: DW_TAG_structure_type, name: "s"

test/CodeGen/builtins-wasm.c

@@ -38,6 +38,12 @@ void data_drop() {
   // WEBASSEMBLY64: call void @llvm.wasm.data.drop(i32 3)
 }
 
+__SIZE_TYPE__ tls_size() {
+  return __builtin_wasm_tls_size();
+  // WEBASSEMBLY32: call i32 @llvm.wasm.tls.size.i32()
+  // WEBASSEMBLY64: call i64 @llvm.wasm.tls.size.i64()
+}
+
 void throw(void *obj) {
   return __builtin_wasm_throw(0, obj);
   // WEBASSEMBLY32: call void @llvm.wasm.throw(i32 0, i8* %{{.*}})

test/CodeGenOpenCL/builtins-amdgcn-gfx10.cl

@@ -0,0 +1,24 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1011 -S -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1012 -S -emit-llvm -o - %s | FileCheck %s
+
+typedef unsigned int uint;
+
+// CHECK-LABEL: @test_permlane16(
+// CHECK: call i32 @llvm.amdgcn.permlane16(i32 %a, i32 %b, i32 %c, i32 %d, i1 true, i1 true)
+void test_permlane16(global uint* out, uint a, uint b, uint c, uint d) {
+  *out = __builtin_amdgcn_permlane16(a, b, c, d, 1, 1);
+}
+
+// CHECK-LABEL: @test_permlanex16(
+// CHECK: call i32 @llvm.amdgcn.permlanex16(i32 %a, i32 %b, i32 %c, i32 %d, i1 true, i1 true)
+void test_permlanex16(global uint* out, uint a, uint b, uint c, uint d) {
+  *out = __builtin_amdgcn_permlanex16(a, b, c, d, 1, 1);
+}
+
+// CHECK-LABEL: @test_mov_dpp8(
+// CHECK: call i32 @llvm.amdgcn.mov.dpp8.i32(i32 %a, i32 1)
+void test_mov_dpp8(global uint* out, uint a) {
+  *out = __builtin_amdgcn_mov_dpp8(a, 1);
+}