CodeGen: Cast temporary variable to proper address space

[yaxunl]

Summary

In C++ all variables are in default address space. Previously change has been made to cast automatic variables to default address space. However that is not sufficient since all temporary variables need to be casted to default address space.

This patch casts all temporary variables to default address space except those for passing indirect arguments since they are only used for load/store.

This patch only affects target having non-zero alloca address space.

[yaxunl]

Fix comments.

[Anastasia]

lib/CodeGen/CGDecl.cpp

1107 ↗

(On Diff #100819)

Is this code still needed here, considering that CreateTempAlloca is doing this already?

1109 ↗

(On Diff #100819)

Why removing the comment?

lib/CodeGen/CodeGenFunction.h

1930 ↗

(On Diff #100819)

I am still confused, what is the case the alloca AS is different from the language AS?

[yaxunl]

lib/CodeGen/CGDecl.cpp

1107 ↗

(On Diff #100819)

I have planed but forgot to remove this. This is also the reason I removed the comments. I will remove it in the update.

lib/CodeGen/CodeGenFunction.h

1930 ↗

(On Diff #100819)

In C or C++ language where all variables without explicit address space attribute are assumed to be in the default address space.

[yaxunl]

Remove redundant code.

[rjmccall]

Please split this up into two patches:

• one for your target-specific changes to the global address space and so on
• one for the generic fix to CreateTempAlloca etc.

lib/CodeGen/CGExpr.cpp

76 ↗

(On Diff #100899)

These arguments should both be AST address spaces, not lowered address spaces. That is, this second argument should be LangAS::Default, not DestAddrSpace.

...we should really make proper types for these instead of always passing them around as unsigneds.

lib/CodeGen/CodeGenFunction.h

1931 ↗

(On Diff #100899)

"On certain targets, the alloca address space is not the same as the generic address space.
Since the address of a temporary is often exposed to the program as a pointer or reference
in the generic address space, this function will convert the returned address into the generic
address space by default. That conversion may be suppressed by passing false as
\p CastToGenericAddrSpace."

And please rename DoCast to CastToGenericAddrSpace.

1973 ↗

(On Diff #100899)

"generic" instead of "expected", please, and rename DoCast to CastToGenericAddrSpace.

[yaxunl]

Keep changes for temporary variables only.

[rjmccall]

lib/CodeGen/CGCall.cpp

3805 ↗

(On Diff #101070)

How about "indirect-arg-temp" as the name here?

3835 ↗

(On Diff #101070)

"byval-temp"

lib/CodeGen/CodeGenFunction.h

1937 ↗

(On Diff #101070)

You have an extra newline here.

Hmm. I see what you're getting at with your addition to the comment, but I think the
real problem is that my original recommendation is just wrong. Let's fix that.

"LangAS::Default is the address space of pointers to local variables and temporaries,
as exposed in the source language. In certain configurations, this is not the same as
the alloca address space, and a cast is needed to lift the pointer from the alloca AS
into LangAS::Default. This can happen when the target uses a restricted address space
for the stack but the source language requires LangAS::Default to be a generic address
space. The latter condition is common for most programming languages; OpenCL is an
exception in that LangAS::Default is the private address space, which naturally maps
to the stack.

Because the address of a temporary is often exposed to the program in various ways,
this function will perform the cast by default. The cast may be avoided by passing false
as \p CastToDefaultAddrSpace; this is more efficient if the caller knows that the address
will not be exposed."

And you should rename the parameter to CastToDefaultAddrSpace here and the other
places.

[yaxunl]

Revised by John's comments.

[rjmccall]

Thanks. One tweak and then LGTM.

lib/CodeGen/CodeGenFunction.h

1937 ↗

(On Diff #101221)

This line should have ///.

[yaxunl]

Fix a comment.

[rjmccall]

Thanks!

cfe/trunk/lib/CodeGen/CGCall.cpp

case ABIArgInfo::InAlloca: {
       }
       break;
     }
 
     case ABIArgInfo::Indirect: {
       assert(NumIRArgs == 1);
       if (RV.isScalar() || RV.isComplex()) {
         // Make a temporary alloca to pass the argument.
-        Address Addr = CreateMemTemp(I->Ty, ArgInfo.getIndirectAlign());
+        Address Addr = CreateMemTemp(I->Ty, ArgInfo.getIndirectAlign(),
+                                     "indirect-arg-temp", false);
         IRCallArgs[FirstIRArg] = Addr.getPointer();
 
         LValue argLV = MakeAddrLValue(Addr, I->Ty);
         EmitInitStoreOfNonAggregate(*this, RV, argLV);
       } else {
         // We want to avoid creating an unnecessary temporary+copy here;
         // however, we need one in three cases:
         // 1. If the argument is not byval, and we are required to copy the
case ABIArgInfo::Indirect: {
                  : 0);
         if ((!ArgInfo.getIndirectByVal() && I->NeedsCopy) ||
             (ArgInfo.getIndirectByVal() && Addr.getAlignment() < Align &&
              llvm::getOrEnforceKnownAlignment(Addr.getPointer(),
                                               Align.getQuantity(), *TD)
                < Align.getQuantity()) ||
             (ArgInfo.getIndirectByVal() && (RVAddrSpace != ArgAddrSpace))) {
           // Create an aligned temporary, and copy to it.
-          Address AI = CreateMemTemp(I->Ty, ArgInfo.getIndirectAlign());
+          Address AI = CreateMemTemp(I->Ty, ArgInfo.getIndirectAlign(),
+                                     "byval-temp", false);
           IRCallArgs[FirstIRArg] = AI.getPointer();
           EmitAggregateCopy(AI, Addr, I->Ty, RV.isVolatileQualified());
         } else {
           // Skip the extra memcpy call.
           IRCallArgs[FirstIRArg] = Addr.getPointer();
         }
       }
       break;

cfe/trunk/lib/CodeGen/CGDecl.cpp

void CodeGenFunction::EmitLifetimeEnd(llvm::Value *Size, llvm::Value *Addr) {
   C->setDoesNotThrow();
 }
 
 /// EmitAutoVarAlloca - Emit the alloca and debug information for a
 /// local variable.  Does not emit initialization or destruction.
 CodeGenFunction::AutoVarEmission
 CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) {
   QualType Ty = D.getType();
+  assert(Ty.getAddressSpace() == LangAS::Default);
 
   AutoVarEmission emission(D);
 
   bool isByRef = D.hasAttr<BlocksAttr>();
   emission.IsByRef = isByRef;
 
   CharUnits alignment = getContext().getDeclAlign(&D);
 
if (NRVO) {
       } else {
         allocaTy = ConvertTypeForMem(Ty);
         allocaAlignment = alignment;
       }
 
       // Create the alloca.  Note that we set the name separately from
       // building the instruction so that it's there even in no-asserts
       // builds.
-      address = CreateTempAlloca(allocaTy, allocaAlignment);
-      address.getPointer()->setName(D.getName());
+      address = CreateTempAlloca(allocaTy, allocaAlignment, D.getName());
 
       // Don't emit lifetime markers for MSVC catch parameters. The lifetime of
       // the catch parameter starts in the catchpad instruction, and we can't
       // insert code in those basic blocks.
       bool IsMSCatchParam =
           D.isExceptionVariable() && getTarget().getCXXABI().isMicrosoft();
 
       // Emit a lifetime intrinsic if meaningful. There's no point in doing this
if (Ty->isConstantSizeType()) {
 
     llvm::Value *elementCount;
     QualType elementType;
     std::tie(elementCount, elementType) = getVLASize(Ty);
 
     llvm::Type *llvmTy = ConvertTypeForMem(elementType);
 
     // Allocate memory for the array.
-    llvm::AllocaInst *vla = Builder.CreateAlloca(llvmTy, elementCount, "vla");
-    vla->setAlignment(alignment.getQuantity());
-
-    address = Address(vla, alignment);
+    address = CreateTempAlloca(llvmTy, alignment, "vla", elementCount);
   }
 
-  // Alloca always returns a pointer in alloca address space, which may
-  // be different from the type defined by the language. For example,
-  // in C++ the auto variables are in the default address space. Therefore
-  // cast alloca to the expected address space when necessary.
-  auto T = D.getType();
-  assert(T.getAddressSpace() == LangAS::Default);
-  if (getASTAllocaAddressSpace() != LangAS::Default) {
-    auto *Addr = getTargetHooks().performAddrSpaceCast(
-        *this, address.getPointer(), getASTAllocaAddressSpace(),
-        T.getAddressSpace(),
-        address.getElementType()->getPointerTo(
-            getContext().getTargetAddressSpace(T.getAddressSpace())),
-        /*non-null*/ true);
-    address = Address(Addr, address.getAlignment());
-  }
   setAddrOfLocalVar(&D, address);
   emission.Addr = address;
 
   // Emit debug info for local var declaration.
   if (HaveInsertPoint())
     if (CGDebugInfo *DI = getDebugInfo()) {
       if (CGM.getCodeGenOpts().getDebugInfo() >=
           codegenoptions::LimitedDebugInfo) {

cfe/trunk/lib/CodeGen/CGExpr.cpp

llvm::Value *CodeGenFunction::EmitCastToVoidPtr(llvm::Value *value) {
 
   if (value->getType() == destType) return value;
   return Builder.CreateBitCast(value, destType);
 }
 
 /// CreateTempAlloca - This creates a alloca and inserts it into the entry
 /// block.
 Address CodeGenFunction::CreateTempAlloca(llvm::Type *Ty, CharUnits Align,
-                                          const Twine &Name) {
-  auto Alloca = CreateTempAlloca(Ty, Name);
+                                          const Twine &Name,
+                                          llvm::Value *ArraySize,
+                                          bool CastToDefaultAddrSpace) {
+  auto Alloca = CreateTempAlloca(Ty, Name, ArraySize);
   Alloca->setAlignment(Align.getQuantity());
-  return Address(Alloca, Align);
-}
-
-/// CreateTempAlloca - This creates a alloca and inserts it into the entry
-/// block.
+  llvm::Value *V = Alloca;
+  // Alloca always returns a pointer in alloca address space, which may
+  // be different from the type defined by the language. For example,
+  // in C++ the auto variables are in the default address space. Therefore
+  // cast alloca to the default address space when necessary.
+  if (CastToDefaultAddrSpace && getASTAllocaAddressSpace() != LangAS::Default) {
+    auto DestAddrSpace = getContext().getTargetAddressSpace(LangAS::Default);
+    V = getTargetHooks().performAddrSpaceCast(
+        *this, V, getASTAllocaAddressSpace(), LangAS::Default,
+        Ty->getPointerTo(DestAddrSpace), /*non-null*/ true);
+  }
+
+  return Address(V, Align);
+}
+
+/// CreateTempAlloca - This creates an alloca and inserts it into the entry
+/// block if \p ArraySize is nullptr, otherwise inserts it at the current
+/// insertion point of the builder.
 llvm::AllocaInst *CodeGenFunction::CreateTempAlloca(llvm::Type *Ty,
-                                                    const Twine &Name) {
+                                                    const Twine &Name,
+                                                    llvm::Value *ArraySize) {
+  if (ArraySize)
+    return Builder.CreateAlloca(Ty, ArraySize, Name);
   return new llvm::AllocaInst(Ty, CGM.getDataLayout().getAllocaAddrSpace(),
-                              nullptr, Name, AllocaInsertPt);
+                              ArraySize, Name, AllocaInsertPt);
 }
 
 /// CreateDefaultAlignTempAlloca - This creates an alloca with the
 /// default alignment of the corresponding LLVM type, which is *not*
 /// guaranteed to be related in any way to the expected alignment of
 /// an AST type that might have been lowered to Ty.
 Address CodeGenFunction::CreateDefaultAlignTempAlloca(llvm::Type *Ty,
                                                       const Twine &Name) {
void CodeGenFunction::InitTempAlloca(Address Var, llvm::Value *Init) {
   Block->getInstList().insertAfter(AllocaInsertPt->getIterator(), Store);
 }
 
 Address CodeGenFunction::CreateIRTemp(QualType Ty, const Twine &Name) {
   CharUnits Align = getContext().getTypeAlignInChars(Ty);
   return CreateTempAlloca(ConvertType(Ty), Align, Name);
 }
 
-Address CodeGenFunction::CreateMemTemp(QualType Ty, const Twine &Name) {
+Address CodeGenFunction::CreateMemTemp(QualType Ty, const Twine &Name,
+                                       bool CastToDefaultAddrSpace) {
   // FIXME: Should we prefer the preferred type alignment here?
-  return CreateMemTemp(Ty, getContext().getTypeAlignInChars(Ty), Name);
+  return CreateMemTemp(Ty, getContext().getTypeAlignInChars(Ty), Name,
+                       CastToDefaultAddrSpace);
 }
 
 Address CodeGenFunction::CreateMemTemp(QualType Ty, CharUnits Align,
-                                       const Twine &Name) {
-  return CreateTempAlloca(ConvertTypeForMem(Ty), Align, Name);
+                                       const Twine &Name,
+                                       bool CastToDefaultAddrSpace) {
+  return CreateTempAlloca(ConvertTypeForMem(Ty), Align, Name, nullptr,
+                          CastToDefaultAddrSpace);
 }
 
 /// EvaluateExprAsBool - Perform the usual unary conversions on the specified
 /// expression and compare the result against zero, returning an Int1Ty value.
 llvm::Value *CodeGenFunction::EvaluateExprAsBool(const Expr *E) {
   PGO.setCurrentStmt(E);
   if (const MemberPointerType *MPT = E->getType()->getAs<MemberPointerType>()) {
     llvm::Value *MemPtr = EmitScalarExpr(E);

cfe/trunk/lib/CodeGen/CodeGenFunction.h

#endif
   Address EmitLoadOfReference(Address Ref, const ReferenceType *RefTy,
                               LValueBaseInfo *BaseInfo = nullptr);
   LValue EmitLoadOfReferenceLValue(Address Ref, const ReferenceType *RefTy);
 
   Address EmitLoadOfPointer(Address Ptr, const PointerType *PtrTy,
                             LValueBaseInfo *BaseInfo = nullptr);
   LValue EmitLoadOfPointerLValue(Address Ptr, const PointerType *PtrTy);
 
-  /// CreateTempAlloca - This creates a alloca and inserts it into the entry
-  /// block. The caller is responsible for setting an appropriate alignment on
+  /// CreateTempAlloca - This creates an alloca and inserts it into the entry
+  /// block if \p ArraySize is nullptr, otherwise inserts it at the current
+  /// insertion point of the builder. The caller is responsible for setting an
+  /// appropriate alignment on
   /// the alloca.
-  llvm::AllocaInst *CreateTempAlloca(llvm::Type *Ty,
-                                     const Twine &Name = "tmp");
+  ///
+  /// \p ArraySize is the number of array elements to be allocated if it
+  ///    is not nullptr.
+  ///
+  /// LangAS::Default is the address space of pointers to local variables and
+  /// temporaries, as exposed in the source language. In certain
+  /// configurations, this is not the same as the alloca address space, and a
+  /// cast is needed to lift the pointer from the alloca AS into
+  /// LangAS::Default. This can happen when the target uses a restricted
+  /// address space for the stack but the source language requires
+  /// LangAS::Default to be a generic address space. The latter condition is
+  /// common for most programming languages; OpenCL is an exception in that
+  /// LangAS::Default is the private address space, which naturally maps
+  /// to the stack.
+  ///
+  /// Because the address of a temporary is often exposed to the program in
+  /// various ways, this function will perform the cast by default. The cast
+  /// may be avoided by passing false as \p CastToDefaultAddrSpace; this is
+  /// more efficient if the caller knows that the address will not be exposed.
+  llvm::AllocaInst *CreateTempAlloca(llvm::Type *Ty, const Twine &Name = "tmp",
+                                     llvm::Value *ArraySize = nullptr);
   Address CreateTempAlloca(llvm::Type *Ty, CharUnits align,
-                           const Twine &Name = "tmp");
+                           const Twine &Name = "tmp",
+                           llvm::Value *ArraySize = nullptr,
+                           bool CastToDefaultAddrSpace = true);
 
   /// CreateDefaultAlignedTempAlloca - This creates an alloca with the
   /// default ABI alignment of the given LLVM type.
   ///
   /// IMPORTANT NOTE: This is *not* generally the right alignment for
   /// any given AST type that happens to have been lowered to the
   /// given IR type.  This should only ever be used for function-local,
   /// IR-driven manipulations like saving and restoring a value.  Do
#endif
   /// PHI construction), but the type is the IR type, not the type appropriate
   /// for storing in memory.
   ///
   /// That is, this is exactly equivalent to CreateMemTemp, but calling
   /// ConvertType instead of ConvertTypeForMem.
   Address CreateIRTemp(QualType T, const Twine &Name = "tmp");
 
   /// CreateMemTemp - Create a temporary memory object of the given type, with
-  /// appropriate alignment.
-  Address CreateMemTemp(QualType T, const Twine &Name = "tmp");
-  Address CreateMemTemp(QualType T, CharUnits Align, const Twine &Name = "tmp");
+  /// appropriate alignment. Cast it to the default address space if
+  /// \p CastToDefaultAddrSpace is true.
+  Address CreateMemTemp(QualType T, const Twine &Name = "tmp",
+                        bool CastToDefaultAddrSpace = true);
+  Address CreateMemTemp(QualType T, CharUnits Align, const Twine &Name = "tmp",
+                        bool CastToDefaultAddrSpace = true);
 
   /// CreateAggTemp - Create a temporary memory object for the given
   /// aggregate type.
   AggValueSlot CreateAggTemp(QualType T, const Twine &Name = "tmp") {
     return AggValueSlot::forAddr(CreateMemTemp(T, Name),
                                  T.getQualifiers(),
                                  AggValueSlot::IsNotDestructed,
                                  AggValueSlot::DoesNotNeedGCBarriers,

cfe/trunk/test/CodeGen/address-space.c

-// RUN: %clang_cc1 -triple x86_64-apple-darwin -emit-llvm < %s | FileCheck -check-prefixes=CHECK,GIZ %s
+// RUN: %clang_cc1 -triple x86_64-apple-darwin -emit-llvm < %s | FileCheck -check-prefixes=CHECK,X86,GIZ %s
 // RUN: %clang_cc1 -triple amdgcn -emit-llvm < %s | FileCheck -check-prefixes=CHECK,PIZ %s
-// RUN: %clang_cc1 -triple amdgcn---amdgiz -emit-llvm < %s | FileCheck -check-prefixes=CHECK,GIZ %s
+// RUN: %clang_cc1 -triple amdgcn---amdgiz -emit-llvm < %s | FileCheck -check-prefixes=CHECK,AMDGIZ,GIZ %s
 
 // CHECK: @foo = common addrspace(1) global
 int foo __attribute__((address_space(1)));
 
 // CHECK: @ban = common addrspace(1) global
 int ban[10] __attribute__((address_space(1)));
 
 // CHECK: @a = common global

cfe/trunk/test/CodeGen/default-address-space.c

 int *B;
 
 // COM-LABEL: define i32 @test1()
 // PIZ: load i32, i32 addrspace(4)* @foo
 // CHECK: load i32, i32* @foo
 int test1() { return foo; }
 
 // COM-LABEL: define i32 @test2(i32 %i)
-// PIZ: load i32, i32 addrspace(4)*
+// COM: %[[addr:.*]] = getelementptr
+// PIZ: load i32, i32 addrspace(4)* %[[addr]]
 // PIZ-NEXT: ret i32
-// CHECK: load i32, i32*
+// CHECK: load i32, i32* %[[addr]]
 // CHECK-NEXT: ret i32
 int test2(int i) { return ban[i]; }
 
 // COM-LABEL: define void @test3()
 // PIZ: load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* @B
 // PIZ: load i32, i32 addrspace(4)*
 // PIZ: load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* @A
 // PIZ: store i32 {{.*}}, i32 addrspace(4)*
 // CHECK: load i32*, i32** @B
 // CHECK: load i32, i32*
 // CHECK: load i32*, i32** @A
 // CHECK: store i32 {{.*}}, i32*
 void test3() {
   *A = *B;
 }
 
 // PIZ-LABEL: define void @test4(i32 addrspace(4)* %a)
-// PIZ: %[[a_addr:.*]] = alloca i32 addrspace(4)*
-// PIZ: store i32 addrspace(4)* %a, i32 addrspace(4)** %[[a_addr]]
-// PIZ: %[[r0:.*]] = load i32 addrspace(4)*, i32 addrspace(4)** %[[a_addr]]
+// PIZ: %[[alloca:.*]] = alloca i32 addrspace(4)*
+// PIZ: %[[a_addr:.*]] = addrspacecast{{.*}} %[[alloca]] to i32 addrspace(4)* addrspace(4)*
+// PIZ: store i32 addrspace(4)* %a, i32 addrspace(4)* addrspace(4)* %[[a_addr]]
+// PIZ: %[[r0:.*]] = load i32 addrspace(4)*, i32 addrspace(4)* addrspace(4)* %[[a_addr]]
 // PIZ: %[[arrayidx:.*]] = getelementptr inbounds i32, i32 addrspace(4)* %[[r0]]
 // PIZ: store i32 0, i32 addrspace(4)* %[[arrayidx]]
 // CHECK-LABEL: define void @test4(i32* %a)
-// CHECK: %[[a_addr:.*]] = alloca i32*, align 4, addrspace(5)
-// CHECK: store i32* %a, i32* addrspace(5)* %[[a_addr]]
-// CHECK: %[[r0:.*]] = load i32*, i32* addrspace(5)* %[[a_addr]]
+// CHECK: %[[alloca:.*]] = alloca i32*, align 4, addrspace(5)
+// CHECK: %[[a_addr:.*]] = addrspacecast{{.*}} %[[alloca]] to i32**
+// CHECK: store i32* %a, i32** %[[a_addr]]
+// CHECK: %[[r0:.*]] = load i32*, i32** %[[a_addr]]
 // CHECK: %[[arrayidx:.*]] = getelementptr inbounds i32, i32* %[[r0]]
 // CHECK: store i32 0, i32* %[[arrayidx]]
 void test4(int *a) {
   a[0] = 0;
 }

cfe/trunk/test/CodeGen/x86_64-arguments.c

 
 void test54_helper(__m256, ...);
 __m256 x54;
 void test54() {
   test54_helper(x54, x54, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0i);
   test54_helper(x54, x54, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0i);
 }
 // AVX: @test54_helper(<8 x float> {{%[a-zA-Z0-9]+}}, <8 x float> {{%[a-zA-Z0-9]+}}, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double {{%[a-zA-Z0-9]+}}, double {{%[a-zA-Z0-9]+}})
-// AVX: @test54_helper(<8 x float> {{%[a-zA-Z0-9]+}}, <8 x float> {{%[a-zA-Z0-9]+}}, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, { double, double }* byval align 8 {{%[a-zA-Z0-9]+}})
+// AVX: @test54_helper(<8 x float> {{%[a-zA-Z0-9]+}}, <8 x float> {{%[a-zA-Z0-9]+}}, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, { double, double }* byval align 8 {{%[^)]+}})
 
 typedef float __m512 __attribute__ ((__vector_size__ (64)));
 typedef struct {
   __m512 m;
 } s512;
 
 s512 x55;
 __m512 x56;

 // AVX512-LABEL: define void @f63
 // AVX512-NOT: br i1
 // AVX512: ret void
 void f63(__m512 *m, __builtin_va_list argList) {
   *m = __builtin_va_arg(argList, __m512);
 }
 
 // AVX512: @f64_helper(<16 x float> {{%[a-zA-Z0-9]+}}, <16 x float> {{%[a-zA-Z0-9]+}}, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double {{%[a-zA-Z0-9]+}}, double {{%[a-zA-Z0-9]+}})
-// AVX512: @f64_helper(<16 x float> {{%[a-zA-Z0-9]+}}, <16 x float> {{%[a-zA-Z0-9]+}}, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, { double, double }* byval align 8 {{%[a-zA-Z0-9]+}})
+// AVX512: @f64_helper(<16 x float> {{%[a-zA-Z0-9]+}}, <16 x float> {{%[a-zA-Z0-9]+}}, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, double 1.000000e+00, { double, double }* byval align 8 {{%[^)]+}})
 void f64_helper(__m512, ...);
 __m512 x64;
 void f64() {
   f64_helper(x64, x64, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0i);
   f64_helper(x64, x64, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0i);
 }
 
 struct t65 {
   __m256 m;
   int : 0;
 };
 // SSE-LABEL: @f65(%struct.t65* byval align 32 %{{[^,)]+}})
 // AVX: @f65(<8 x float> %{{[^,)]+}})
 void f65(struct t65 a0) {
 }

cfe/trunk/test/CodeGenCXX/amdgcn-automatic-variable.cpp

 // RUN: %clang_cc1 -O0 -triple amdgcn---amdgiz -emit-llvm %s -o - | FileCheck %s
 
 // CHECK-LABEL: define void @_Z5func1Pi(i32* %x)
 void func1(int *x) {
   // CHECK: %[[x_addr:.*]] = alloca i32*{{.*}}addrspace(5)
-  // CHECK: store i32* %x, i32* addrspace(5)* %[[x_addr]]
-  // CHECK: %[[r0:.*]] = load i32*, i32* addrspace(5)* %[[x_addr]]
-  // CHECK: store i32 1, i32* %[[r0]]
+  // CHECK: %[[r0:.*]] = addrspacecast i32* addrspace(5)* %[[x_addr]] to i32**
+  // CHECK: store i32* %x, i32** %[[r0]]
+  // CHECK: %[[r1:.*]] = load i32*, i32** %[[r0]]
+  // CHECK: store i32 1, i32* %[[r1]]
   *x = 1;
 }
 
 // CHECK-LABEL: define void @_Z5func2v()
 void func2(void) {
   // CHECK: %lv1 = alloca i32, align 4, addrspace(5)
   // CHECK: %lv2 = alloca i32, align 4, addrspace(5)
   // CHECK: %la = alloca [100 x i32], align 4, addrspace(5)

 // CHECK-LABEL: define void @_Z5func3v
 void func3() {
   // CHECK: %[[a:.*]] = alloca %class.A, align 4, addrspace(5)
   // CHECK: %[[r0:.*]] = addrspacecast %class.A addrspace(5)* %[[a]] to %class.A*
   // CHECK: call void @_ZN1AC1Ev(%class.A* %[[r0]])
   // CHECK: call void @_ZN1AD1Ev(%class.A* %[[r0]])
   A a;
 }
+
+// CHECK-LABEL: define void @_Z5func4i
+void func4(int x) {
+  // CHECK: %[[x_addr:.*]] = alloca i32, align 4, addrspace(5)
+  // CHECK: %[[r0:.*]] = addrspacecast i32 addrspace(5)* %[[x_addr]] to i32*
+  // CHECK: store i32 %x, i32* %[[r0]], align 4
+  // CHECK: call void @_Z5func1Pi(i32* %[[r0]])
+  func1(&x);
+}