Updated PPCG Code Generation for OpenCL compatibility

[PhilippSchaad]

Summary

Added a small change to the way pointer arguments are set in the kernel code generation. The way the pointer is retrieved now, specifically requests global address space to be annotated. This is necessary, if the IR should be run through NVPTX to generate OpenCL compatible PTX. The changes do not affect the PTX Strings generated for the CUDA target (nvptx64-nvidia-cuda), but are necessary for OpenCL (nvptx64-nvidia-nvcl).

Additionally, the data layout has been updated to what the NVPTX Backend requests/recommends.

[grosser]

Add Siddharth, as another review. Siddharth, could you take a look?

[bollu]

lib/CodeGen/PPCGCodeGeneration.cpp

1283

Please add a test case to check that both code paths are taken.

1306–1307

1. Change the magic constant 1 to a static const GlobalMemory = 1 or something of the sort.
2. This change seems unrelated to the initial change. Can this be split into two separate patches?
3. Please add a test case that makes sure that the correct code is generated.

[PhilippSchaad]

lib/CodeGen/PPCGCodeGeneration.cpp

1283

Not both branches are being taken at the moment. The only invocation of this function, at this time, is being done with constant true. The 32-bit part was already present pre-patch, I just adapted BOTH strings to fit the ones recommended by the NVPTX backend. (See NVPTX Usage)

1306–1307

1. What's the issue exactly?
2. The change should do exactly the same thing, just in one statement instead of two. Before, I used getInt8Ty, and then got a pointer to it, declaring address space 1 = globa, now I'm doing it in one step.
3. Sure, can do

[PhilippSchaad]

Updated tests to fit address space changes and correct NVPTX data layout.

Each GPGPU test case now carries the correct NVPTX data-layout requested for 64-bit execution. Additionally, all parameters have been adapted to be global pointers. Consequently, each occurrence of said pointers in the following kernel, now carries the correct address space annotation as well.

[jlebar]

Additionally, all parameters have been adapted to be global pointers.

FYI, the LLVM backend has a pass that converts kernel args to address space 1 -- it's not necessary for front-ends to do this. Depending on what you're doing in these tests it may or may not be necessary -- I can't check because this patch was uploaded without context.

[PhilippSchaad]

In D32215#734619, @jlebar wrote:

Additionally, all parameters have been adapted to be global pointers.

FYI, the LLVM backend has a pass that converts kernel args to address space 1 -- it's not necessary for front-ends to do this. Depending on what you're doing in these tests it may or may not be necessary -- I can't check because this patch was uploaded without context.

Oh, thanks for pointing that out! As long as the generated PTX kernel's arguments are annotated with .global it should do the trick.

PS: What I'm trying to do is building said PTX kernel with OpenCL, and OpenCL requires the memory-object pointers in the kernel arguments to be marked as .global. :-)

[PhilippSchaad]

@jlebar From what I understand, NVPTX only marks a kernel argument as .global if its address space in the IR has been marked with addrspace(1), and said conversion to global address space would then take place inside the kernel, instead of marking the argument itself with .global. Am I wrong on this?

[jlebar]

Am I wrong on this?

No, it looks like I was wrong and you are correct.

Sorry about that.

[PhilippSchaad]

In D32215#734796, @jlebar wrote:

Am I wrong on this?

No, it looks like I was wrong and you are correct.

Sorry about that.

No problem.

In that case a review would be appreciated.

[jlebar]

Since this patch was sent without sufficient context, I cannot really tell what it is dong. See http://llvm.org/docs/Phabricator.html#requesting-a-review-via-the-web-interface. I'm also not comfortable reviewing Polly code, which I am unfamiliar with.

But in principle what is described in the patch description sounds fine; perhaps that's all the LGTM you need.

[PhilippSchaad]

@grosser @bollu ping

[bollu]

lib/CodeGen/PPCGCodeGeneration.cpp

1306–1307

From what I understand, the constant 1 is used to refer to the global address space in the NVPTX backend (please correct me if I am wrong). Could you change the invocation to:

static const int UseGlobalMemory = 1;
Args.push_back(Builder.getInt8PtrTy(UseGlobalMemory));

I personally dislike magic constants in code. Documenting it this way should save the next person a google search.

[PhilippSchaad]

lib/CodeGen/PPCGCodeGeneration.cpp

1306–1307

Oh that's what you mean, yes, in that case I agree with you. This is a better solution. I will update that, thanks for pointing it out.

[PhilippSchaad]

Removed magic constant in favor of more doc-friendly 'explanatory-constant'.

[grosser]

Hi Philippe,

one comment regarding the data-layout changes:

test/GPGPU/double-parallel-loop.ll

215

Hi Philippe,

I am slightly surprised you change the data-layout of the input code. This is generally nothing we have control over. Is this change indeed needed to make your patch work?

[bollu]

LGTM, thanks for the patch!

[PhilippSchaad]

test/GPGPU/double-parallel-loop.ll

215

It is not, this was only done for consistency reasons, which does not make sense. My bad, I will revert those.

[PhilippSchaad]

Reverted unnecessary datalayout changes in tests.

[PhilippSchaad]

@bollu Could you maybe land this when @grosser agrees with the revert?

[grosser]

LGTM.

[PhilippSchaad]

@bollu ping: I would love it if you could commit this change please ;-)

lib/CodeGen/PPCGCodeGeneration.cpp

 ///
 /// @param is64Bit Are we looking for a 64 bit architecture?
 static std::string computeNVPTXDataLayout(bool is64Bit) {
-  std::string Ret = "e";
+  std::string Ret = "";
 
-  if (!is64Bit)
-    Ret += "-p:32:32";
-
-  Ret += "-i64:64-v16:16-v32:32-n16:32:64";
+  if (!is64Bit) {
+    Ret += "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:"
+           "64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:"
+           "64-v128:128:128-n16:32:64";
+  } else {
+    Ret += "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:"
+           "64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:"
+           "64-v128:128:128-n16:32:64";
+  }
 
   return Ret;
 }

       const ScopArrayInfo *SAI = ScopArrayInfo::getFromId(Id);
       Args.push_back(SAI->getElementType());
     } else {
-      Args.push_back(Builder.getInt8PtrTy());
+      static const int UseGlobalMemory = 1;
+      Args.push_back(Builder.getInt8PtrTy(UseGlobalMemory));
     }
   }
 

test/GPGPU/cuda-annotations.ll

 
 ; REQUIRES: pollyacc
 
-; KERNEL: define ptx_kernel void @kernel_0(i8* %MemRef_A, i64 %n) #0 {
+; KERNEL: define ptx_kernel void @kernel_0(i8 addrspace(1)* %MemRef_A, i64 %n) #0 {
 
 ; KERNEL: !nvvm.annotations = !{!0}
 
-; KERNEL: !0 = !{void (i8*, i64)* @kernel_0, !"maxntidx", i32 32, !"maxntidy", i32 1, !"maxntidz", i32 1}
+; KERNEL: !0 = !{void (i8 addrspace(1)*, i64)* @kernel_0, !"maxntidx", i32 32, !"maxntidy", i32 1, !"maxntidz", i32 1}
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @foo(i64* %A, i64 %n) {
 bb:

test/GPGPU/double-parallel-loop.ll

 ;          A[i][j] += i * j;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @double_parallel_loop([1024 x float]* %A) {
 bb:

test/GPGPU/host-control-flow.ll

 ; IR-NEXT:   %polly.loop_cond = icmp sle i64 %polly.indvar, 98
 ; IR-NEXT:   br i1 %polly.loop_cond, label %polly.loop_header, label %polly.loop_exit
 
-; KERNEL-IR: define ptx_kernel void @kernel_0(i8* %MemRef_A, i64 %c0)
+; KERNEL-IR: define ptx_kernel void @kernel_0(i8 addrspace(1)* %MemRef_A, i64 %c0)
 ; KERNEL-IR-LABEL: entry:
 ; KERNEL-IR-NEXT:   %0 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
 ; KERNEL-IR-NEXT:   %b0 = zext i32 %0 to i64

 ; KERNEL-IR-NEXT:   br label %polly.stmt.for.body3
 
 ; KERNEL-IR-LABEL: polly.stmt.for.body3:                             ; preds = %polly.then
-; KERNEL-IR-NEXT:   %polly.access.cast.MemRef_A = bitcast i8* %MemRef_A to float*
+; KERNEL-IR-NEXT:   %polly.access.cast.MemRef_A = bitcast i8 addrspace(1)* %MemRef_A to float addrspace(1)*
 ; KERNEL-IR-NEXT:   %pexp.pdiv_r = urem i64 %c0, 2
 ; KERNEL-IR-NEXT:   %polly.access.mul.MemRef_A = mul nsw i64 %pexp.pdiv_r, 100
 ; KERNEL-IR-NEXT:   %7 = mul nsw i64 32, %b0
 ; KERNEL-IR-NEXT:   %8 = add nsw i64 %7, %t0
 ; KERNEL-IR-NEXT:   %polly.access.add.MemRef_A = add nsw i64 %polly.access.mul.MemRef_A, %8
-; KERNEL-IR-NEXT:   %polly.access.MemRef_A = getelementptr float, float* %polly.access.cast.MemRef_A, i64 %polly.access.add.MemRef_A
-; KERNEL-IR-NEXT:   %tmp_p_scalar_ = load float, float* %polly.access.MemRef_A, align 4
-; KERNEL-IR-NEXT:   %polly.access.cast.MemRef_A1 = bitcast i8* %MemRef_A to float*
+; KERNEL-IR-NEXT:   %polly.access.MemRef_A = getelementptr float, float addrspace(1)* %polly.access.cast.MemRef_A, i64 %polly.access.add.MemRef_A
+; KERNEL-IR-NEXT:   %tmp_p_scalar_ = load float, float addrspace(1)* %polly.access.MemRef_A, align 4
+; KERNEL-IR-NEXT:   %polly.access.cast.MemRef_A1 = bitcast i8 addrspace(1)* %MemRef_A to float addrspace(1)*
 ; KERNEL-IR-NEXT:   %pexp.pdiv_r2 = urem i64 %c0, 2
 ; KERNEL-IR-NEXT:   %polly.access.mul.MemRef_A3 = mul nsw i64 %pexp.pdiv_r2, 100
 ; KERNEL-IR-NEXT:   %9 = mul nsw i64 32, %b0
 ; KERNEL-IR-NEXT:   %10 = add nsw i64 %9, %t0
 ; KERNEL-IR-NEXT:   %11 = add nsw i64 %10, 1
 ; KERNEL-IR-NEXT:   %polly.access.add.MemRef_A4 = add nsw i64 %polly.access.mul.MemRef_A3, %11
-; KERNEL-IR-NEXT:   %polly.access.MemRef_A5 = getelementptr float, float* %polly.access.cast.MemRef_A1, i64 %polly.access.add.MemRef_A4
-; KERNEL-IR-NEXT:   %tmp2_p_scalar_ = load float, float* %polly.access.MemRef_A5, align 4
+; KERNEL-IR-NEXT:   %polly.access.MemRef_A5 = getelementptr float, float addrspace(1)* %polly.access.cast.MemRef_A1, i64 %polly.access.add.MemRef_A4
+; KERNEL-IR-NEXT:   %tmp2_p_scalar_ = load float, float addrspace(1)* %polly.access.MemRef_A5, align 4
 ; KERNEL-IR-NEXT:   %p_add = fadd float %tmp_p_scalar_, %tmp2_p_scalar_
-; KERNEL-IR-NEXT:   %polly.access.cast.MemRef_A6 = bitcast i8* %MemRef_A to float*
+; KERNEL-IR-NEXT:   %polly.access.cast.MemRef_A6 = bitcast i8 addrspace(1)* %MemRef_A to float addrspace(1)*
 ; KERNEL-IR-NEXT:   %pexp.pdiv_r7 = urem i64 %c0, 2
 ; KERNEL-IR-NEXT:   %polly.access.mul.MemRef_A8 = mul nsw i64 %pexp.pdiv_r7, 100
 ; KERNEL-IR-NEXT:   %12 = mul nsw i64 32, %b0
 ; KERNEL-IR-NEXT:   %13 = add nsw i64 %12, %t0
 ; KERNEL-IR-NEXT:   %14 = add nsw i64 %13, 2
 ; KERNEL-IR-NEXT:   %polly.access.add.MemRef_A9 = add nsw i64 %polly.access.mul.MemRef_A8, %14
-; KERNEL-IR-NEXT:   %polly.access.MemRef_A10 = getelementptr float, float* %polly.access.cast.MemRef_A6, i64 %polly.access.add.MemRef_A9
-; KERNEL-IR-NEXT:   %tmp3_p_scalar_ = load float, float* %polly.access.MemRef_A10, align 4
+; KERNEL-IR-NEXT:   %polly.access.MemRef_A10 = getelementptr float, float addrspace(1)* %polly.access.cast.MemRef_A6, i64 %polly.access.add.MemRef_A9
+; KERNEL-IR-NEXT:   %tmp3_p_scalar_ = load float, float addrspace(1)* %polly.access.MemRef_A10, align 4
 ; KERNEL-IR-NEXT:   %p_add12 = fadd float %p_add, %tmp3_p_scalar_
-; KERNEL-IR-NEXT:   %polly.access.cast.MemRef_A11 = bitcast i8* %MemRef_A to float*
+; KERNEL-IR-NEXT:   %polly.access.cast.MemRef_A11 = bitcast i8 addrspace(1)* %MemRef_A to float addrspace(1)*
 ; KERNEL-IR-NEXT:   %15 = add nsw i64 %c0, 1
 ; KERNEL-IR-NEXT:   %pexp.pdiv_r12 = urem i64 %15, 2
 ; KERNEL-IR-NEXT:   %polly.access.mul.MemRef_A13 = mul nsw i64 %pexp.pdiv_r12, 100

 ; KERNEL-IR-NEXT:   %17 = add nsw i64 %16, %t0
 ; KERNEL-IR-NEXT:   %18 = add nsw i64 %17, 1
 ; KERNEL-IR-NEXT:   %polly.access.add.MemRef_A14 = add nsw i64 %polly.access.mul.MemRef_A13, %18
-; KERNEL-IR-NEXT:   %polly.access.MemRef_A15 = getelementptr float, float* %polly.access.cast.MemRef_A11, i64 %polly.access.add.MemRef_A14
-; KERNEL-IR-NEXT:   %tmp4_p_scalar_ = load float, float* %polly.access.MemRef_A15, align 4
+; KERNEL-IR-NEXT:   %polly.access.MemRef_A15 = getelementptr float, float addrspace(1)* %polly.access.cast.MemRef_A11, i64 %polly.access.add.MemRef_A14
+; KERNEL-IR-NEXT:   %tmp4_p_scalar_ = load float, float addrspace(1)* %polly.access.MemRef_A15, align 4
 ; KERNEL-IR-NEXT:   %p_add17 = fadd float %tmp4_p_scalar_, %p_add12
-; KERNEL-IR-NEXT:   %polly.access.cast.MemRef_A16 = bitcast i8* %MemRef_A to float*
+; KERNEL-IR-NEXT:   %polly.access.cast.MemRef_A16 = bitcast i8 addrspace(1)* %MemRef_A to float addrspace(1)*
 ; KERNEL-IR-NEXT:   %19 = add nsw i64 %c0, 1
 ; KERNEL-IR-NEXT:   %pexp.pdiv_r17 = urem i64 %19, 2
 ; KERNEL-IR-NEXT:   %polly.access.mul.MemRef_A18 = mul nsw i64 %pexp.pdiv_r17, 100

 ; KERNEL-IR-NEXT:   %21 = add nsw i64 %20, %t0
 ; KERNEL-IR-NEXT:   %22 = add nsw i64 %21, 1
 ; KERNEL-IR-NEXT:   %polly.access.add.MemRef_A19 = add nsw i64 %polly.access.mul.MemRef_A18, %22
-; KERNEL-IR-NEXT:   %polly.access.MemRef_A20 = getelementptr float, float* %polly.access.cast.MemRef_A16, i64 %polly.access.add.MemRef_A19
-; KERNEL-IR-NEXT:   store float %p_add17, float* %polly.access.MemRef_A20, align 4
+; KERNEL-IR-NEXT:   %polly.access.MemRef_A20 = getelementptr float, float addrspace(1)* %polly.access.cast.MemRef_A16, i64 %polly.access.add.MemRef_A19
+; KERNEL-IR-NEXT:   store float %p_add17, float addrspace(1)* %polly.access.MemRef_A20, align 4
 ; KERNEL-IR-NEXT:   br label %polly.merge
 
 ; KERNEL-IR-LABEL: polly.else:                                       ; preds = %polly.cond
 ; KERNEL-IR-NEXT:   br label %polly.merge
 ; KERNEL-IR-NEXT: }
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @foo([100 x float]* %A) {
 entry:

test/GPGPU/host-statement.ll

 
 ; REQUIRES: pollyacc
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 target triple = "x86_64-unknown-linux-gnu"
 
 declare void @llvm.lifetime.start(i64, i8* nocapture) #0

test/GPGPU/invalid-kernel.ll

 
 ; IR: br i1 false, label %polly.start, label %bb1
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @foo(i64* %A, i64* %B) {
 bb:

test/GPGPU/kernel-params-only-some-arrays.ll

 
 ; KERNEL: ; ModuleID = 'kernel_0'
 ; KERNEL-NEXT: source_filename = "kernel_0"
-; KERNEL-NEXT: target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
+; KERNEL-NEXT: target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 ; KERNEL-NEXT: target triple = "nvptx64-nvidia-cuda"
 
-; KERNEL: define ptx_kernel void @kernel_0(i8* %MemRef_A)
+; KERNEL: define ptx_kernel void @kernel_0(i8 addrspace(1)* %MemRef_A)
 ; KERNEL-NEXT:   entry:
 ; KERNEL-NEXT:     %0 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
 ; KERNEL-NEXT:     %b0 = zext i32 %0 to i64

 
 ; KERNEL: ; ModuleID = 'kernel_1'
 ; KERNEL-NEXT: source_filename = "kernel_1"
-; KERNEL-NEXT: target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
+; KERNEL-NEXT: target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 ; KERNEL-NEXT: target triple = "nvptx64-nvidia-cuda"
 
-; KERNEL: define ptx_kernel void @kernel_1(i8* %MemRef_B)
+; KERNEL: define ptx_kernel void @kernel_1(i8 addrspace(1)* %MemRef_B)
 ; KERNEL-NEXT:   entry:
 ; KERNEL-NEXT:     %0 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
 ; KERNEL-NEXT:     %b0 = zext i32 %0 to i64

 ; IR-NEXT:  [[DATA:%.*]] = bitcast i8** %polly_launch_1_param_0 to i8*
 ; IR-NEXT:  store i8* [[DATA]], i8** [[SLOT]]
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @kernel_params_only_some_arrays(float* %A, float* %B) {
 entry:

test/GPGPU/kernel-params-scop-parameter.ll

 ;        A[i] += 42;
 ;    }
 
-; KERNEL-IR: define ptx_kernel void @kernel_0(i8* %MemRef_A, i64 %n)
+; KERNEL-IR: define ptx_kernel void @kernel_0(i8 addrspace(1)* %MemRef_A, i64 %n)
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @kernel_params_scop_parameter(float* %A, i64 %n) {
 bb:

test/GPGPU/mostly-sequential.ll

 ; Verify that we identified this kernel as non-profitable.
 ; IR: br i1 false, label %polly.start, label %bb3
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @foo(float* %A) {
 bb:

test/GPGPU/non-read-only-scalars.ll

 ; CODE-NEXT: Stmt_bb20(c0);
 
 ; KERNEL-IR:       store float %p_tmp23, float* %sum.0.phiops
-; KERNEL-IR-NEXT:  [[REGA:%.+]] = bitcast i8* %MemRef_sum_0__phi to float*
+; KERNEL-IR-NEXT:  [[REGA:%.+]] = addrspacecast i8 addrspace(1)* %MemRef_sum_0__phi to float*
 ; KERNEL-IR-NEXT:  [[REGB:%.+]] = load float, float* %sum.0.phiops
 ; KERNEL-IR-NEXT:  store float [[REGB]], float* [[REGA]]
-; KERNEL-IR-NEXT:  [[REGC:%.+]] = bitcast i8* %MemRef_sum_0 to float*
+; KERNEL-IR-NEXT:  [[REGC:%.+]] = addrspacecast i8 addrspace(1)* %MemRef_sum_0 to float*
 ; KERNEL-IR-NEXT:  [[REGD:%.+]] = load float, float* %sum.0.s2a
 ; KERNEL-IR-NEXT:  store float [[REGD]], float* [[REGC]]
 ; KERNEL-IR-NEXT:  ret void
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 @.str = private unnamed_addr constant [4 x i8] c"%f\0A\00", align 1
 

test/GPGPU/non-zero-array-offset.ll

 ;    }
 ;    #endif
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @foo(float* %A, float* %B) {
 bb:

test/GPGPU/only-part-of-array-modified.ll

 ; CODE: cudaCheckReturn(cudaMemcpy(dev_MemRef_B, MemRef_B, (1024) * sizeof(i32), cudaMemcpyHostToDevice));
 ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (2047) * sizeof(i32), cudaMemcpyHostToDevice));
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @foo(float* %A, float* %B) {
 bb:

test/GPGPU/parametric-loop-bound.ll

 ; IR-NEXT: [[REGB:%.+]] = bitcast i64* %polly_launch_0_param_1 to i8*
 ; IR-NEXT: store i8* [[REGB]], i8** [[REGA]]
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @foo(i64* %A, i64 %n) {
 bb:

test/GPGPU/phi-nodes-in-kernel.ll

 
 ; REQUIRES: pollyacc
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 target triple = "x86_64-unknown-linux-gnu"
 
 ; CODE: # host

 ; KERNEL-IR: entry:
 ; KERNEL-IR-NEXT:   %out_l.055.s2a = alloca i32
 ; KERNEL-IR-NEXT:   %out_l.055.phiops = alloca i32
-; KERNEL-IR-NEXT:   %1 = bitcast i8* %MemRef_out_l_055__phi to i32*
+; KERNEL-IR-NEXT:   %1 = addrspacecast i8 addrspace(1)* %MemRef_out_l_055__phi to i32*
 ; KERNEL-IR-NEXT:   %2 = load i32, i32* %1
 ; KERNEL-IR-NEXT:   store i32 %2, i32* %out_l.055.phiops
-; KERNEL-IR-NEXT:   %3 = bitcast i8* %MemRef_out_l_055 to i32*
+; KERNEL-IR-NEXT:   %3 = addrspacecast i8 addrspace(1)* %MemRef_out_l_055 to i32*
 ; KERNEL-IR-NEXT:   %4 = load i32, i32* %3
 ; KERNEL-IR-NEXT:   store i32 %4, i32* %out_l.055.s2a
 

test/GPGPU/private-memory.ll

 
 ; KERNEL:   %polly.access.cast.private_array = bitcast [1 x float]* %private_array to float*
 ; KERNEL-NEXT:   %polly.access.private_array = getelementptr float, float* %polly.access.cast.private_array, i64 0
-; KERNEL-NEXT:   %polly.access.cast.MemRef_A = bitcast i8* %MemRef_A to float*
-; KERNEL-NEXT:   %polly.access.MemRef_A = getelementptr float, float* %polly.access.cast.MemRef_A, i64 %t0
-; KERNEL-NEXT:   %shared.read = load float, float* %polly.access.MemRef_A
+; KERNEL-NEXT:   %polly.access.cast.MemRef_A = bitcast i8 addrspace(1)* %MemRef_A to float addrspace(1)*
+; KERNEL-NEXT:   %polly.access.MemRef_A = getelementptr float, float addrspace(1)* %polly.access.cast.MemRef_A, i64 %t0
+; KERNEL-NEXT:   %shared.read = load float, float addrspace(1)* %polly.access.MemRef_A
 ; KERNEL-NEXT:   store float %shared.read, float* %polly.access.private_array
 
 ; KERNEL:   %polly.access.cast.private_array5 = bitcast [1 x float]* %private_array to float*
 ; KERNEL-NEXT:   %polly.access.private_array6 = getelementptr float, float* %polly.access.cast.private_array5, i64 0
-; KERNEL-NEXT:   %polly.access.cast.MemRef_A7 = bitcast i8* %MemRef_A to float*
-; KERNEL-NEXT:   %polly.access.MemRef_A8 = getelementptr float, float* %polly.access.cast.MemRef_A7, i64 %t0
+; KERNEL-NEXT:   %polly.access.cast.MemRef_A7 = bitcast i8 addrspace(1)* %MemRef_A to float addrspace(1)*
+; KERNEL-NEXT:   %polly.access.MemRef_A8 = getelementptr float, float addrspace(1)* %polly.access.cast.MemRef_A7, i64 %t0
 ; KERNEL-NEXT:   %shared.write = load float, float* %polly.access.private_array6
-; KERNEL-NEXT:   store float %shared.write, float* %polly.access.MemRef_A8
+; KERNEL-NEXT:   store float %shared.write, float addrspace(1)* %polly.access.MemRef_A8
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @add(float* %A) {
 bb:

test/GPGPU/region-stmt.ll

 ;    }
 ;
 source_filename = "/tmp/test.c"
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @foo(float* %A, float* %B) {
 entry:

test/GPGPU/remove-dead-instructions-in-stmt-2.ll

 
 ; REQUIRES: pollyacc
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 target triple = "x86_64-unknown-linux-gnu"
 
-; KERNEL-IR: store i32 0, i32* %polly.access.MemRef_sum_c, align 4
+; KERNEL-IR: store i32 0, i32 addrspace(1)* %polly.access.MemRef_sum_c, align 4
 ; KERNEL-IR-NEXT: br label %polly.merge
 
 define void @kernel_dynprog([50 x [50 x i32]]* %sum_c) {

test/GPGPU/remove-dead-instructions-in-stmt.ll

 ; condition. This code referred to CPU registers and consequently resulted
 ; in invalid bitcode.
 
-; KERNEL-IR: store i32 0, i32* %polly.access.MemRef_sum_c, align 4
+; KERNEL-IR: store i32 0, i32 addrspace(1)* %polly.access.MemRef_sum_c, align 4
 ; KERNEL-IR-NEXT: br label %polly.merge
 
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"

test/GPGPU/run-time-check.ll

 ; IR:       %tmp = icmp slt i64 %i.0, %n
 ; IR-NEXT:  br i1 %tmp, label %bb2, label %polly.merge_new_and_old
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @foo(i64 %n, [32 x float]* %A) {
 bb:

test/GPGPU/scalar-param-and-value-32-bit.ll

 ;        A[j + n] += 42;
 ;    }
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 ; CHECK: define ptx_kernel void @kernel_0(i8* %MemRef_A, i32 %n)
 

test/GPGPU/scalar-param-and-value-use.ll

 ;          A[i][j] += A[i + 1][j + 1];
 ;    }
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 ; This test case failed at some point as %n was only available in this kernel
 ; when referenced through an isl_id in an isl ast expression, but not when

test/GPGPU/scalar-parameter-fp128.ll

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @fp128(fp128* %A, fp128 %b) {
 bb:

test/GPGPU/scalar-parameter-half.ll

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @half(half* %A, half %b) {
 bb:

test/GPGPU/scalar-parameter-i120.ll

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @i120(i120* %A, i120 %b) {
 bb:

test/GPGPU/scalar-parameter-i128.ll

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @i128(i128* %A, i128 %b) {
 bb:

test/GPGPU/scalar-parameter-i3000.ll

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @i3000(i3000* %A, i3000 %b) {
 bb:

test/GPGPU/scalar-parameter-i80.ll

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @i80(i80* %A, i80 %b) {
 bb:

test/GPGPU/scalar-parameter-ppc_fp128.ll

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @ppc_fp128(ppc_fp128* %A, ppc_fp128 %b) {
 bb:

test/GPGPU/scalar-parameter-x86_fp80.ll

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @fp128(fp128* %A, fp128 %b) {
 bb:

test/GPGPU/scalar-parameter.ll

 
 ; REQUIRES: pollyacc,nvptx
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 ; KERNEL: define ptx_kernel void @kernel_0(i8* %MemRef_A, float %MemRef_b)
 

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @float(float* %A, float %b) {
 bb:

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @double(double* %A, double %b) {
 bb:

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @i1(i1* %A, i1 %b) {
 bb:

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @i3(i3* %A, i3 %b) {
 bb:

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @i8(i8* %A, i8 %b) {
 bb:

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @i32(i32* %A, i32 %b) {
 bb:

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @i60(i60* %A, i60 %b) {
 bb:

 ;        A[i] += b;
 ;    }
 ;
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @i64(i64* %A, i64 %b) {
 bb:

test/GPGPU/scheduler-timeout.ll

 
 ; REQUIRES: pollyacc
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 target triple = "x86_64-unknown-linux-gnu"
 
 ; This test case took at some point forever to schedule, as the isl scheduler

test/GPGPU/shared-memory-scalar.ll

 ; the scalar is not stored any more in shared memory. We still leave this
 ; test case as documentation if we every forget to mark scalars as read-only.
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @add(float* %A, float %alpha) {
 bb:

test/GPGPU/shared-memory-two-dimensional.ll

 
 ; KERNEL:        %polly.access.mul.MemRef_b = mul nsw i64 %polly.indvar, 8
 ; KERNEL-NEXT:   %polly.access.add.MemRef_b = add nsw i64 %polly.access.mul.MemRef_b, %t0
-; KERNEL-NEXT:   %polly.access.MemRef_b = getelementptr float, float* %polly.access.cast.MemRef_b, i64 %polly.access.add.MemRef_b
-; KERNEL-NEXT:   %shared.read = load float, float* %polly.access.MemRef_b
+; KERNEL-NEXT:   %polly.access.MemRef_b = getelementptr float, float addrspace(1)* %polly.access.cast.MemRef_b, i64 %polly.access.add.MemRef_b
+; KERNEL-NEXT:   %shared.read = load float, float addrspace(1)* %polly.access.MemRef_b
 ; KERNEL-NEXT:   store float %shared.read, float addrspace(3)* %polly.access.shared_MemRef_b
 
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @foo(float* %A, [8 x float]* %b) {
 bb:

test/GPGPU/shared-memory.ll

 ; KERNEL: @shared_MemRef_A = internal addrspace(3) global [32 x float] zeroinitializer, align 4
 
 ; KERNEL:   %polly.access.shared_MemRef_A = getelementptr float, float addrspace(3)* getelementptr inbounds ([32 x float], [32 x float] addrspace(3)* @shared_MemRef_A, i32 0, i32 0), i64 %t0
-; KERNEL-NEXT:   %polly.access.cast.MemRef_A = bitcast i8* %MemRef_A to float*
-; KERNEL-NEXT:   %polly.access.MemRef_A = getelementptr float, float* %polly.access.cast.MemRef_A, i64 %t0
-; KERNEL-NEXT:   %shared.read = load float, float* %polly.access.MemRef_A
+; KERNEL-NEXT:   %polly.access.cast.MemRef_A = bitcast i8 addrspace(1)* %MemRef_A to float addrspace(1)*
+; KERNEL-NEXT:   %polly.access.MemRef_A = getelementptr float, float addrspace(1)* %polly.access.cast.MemRef_A, i64 %t0
+; KERNEL-NEXT:   %shared.read = load float, float addrspace(1)* %polly.access.MemRef_A
 ; KERNEL-NEXT:   store float %shared.read, float addrspace(3)* %polly.access.shared_MemRef_A
 
 ; KERNEL:   %polly.access.shared_MemRef_A3 = getelementptr float, float addrspace(3)* getelementptr inbounds ([32 x float], [32 x float] addrspace(3)* @shared_MemRef_A, i32 0, i32 0), i64 %t0
-; KERNEL-NEXT:   %polly.access.cast.MemRef_A4 = bitcast i8* %MemRef_A to float*
-; KERNEL-NEXT:   %polly.access.MemRef_A5 = getelementptr float, float* %polly.access.cast.MemRef_A4, i64 %t0
+; KERNEL-NEXT:   %polly.access.cast.MemRef_A4 = bitcast i8 addrspace(1)* %MemRef_A to float addrspace(1)*
+; KERNEL-NEXT:   %polly.access.MemRef_A5 = getelementptr float, float addrspace(1)* %polly.access.cast.MemRef_A4, i64 %t0
 ; KERNEL-NEXT:   %shared.write = load float, float addrspace(3)* %polly.access.shared_MemRef_A3
-; KERNEL-NEXT:   store float %shared.write, float* %polly.access.MemRef_A5
+; KERNEL-NEXT:   store float %shared.write, float addrspace(1)* %polly.access.MemRef_A5
 
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 
 define void @add(float* %A) {
 bb:

test/GPGPU/size-cast.ll

 ; IR-NEXT:   mul i64
 ; IR-NEXT:   @polly_allocateMemoryForDevice
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 target triple = "x86_64-unknown-linux-gnu"
 
 define void @hoge(i32 %arg, i32 %arg1, [1000 x double]* %arg2, double* %arg3) {

test/GPGPU/untouched-arrays.ll

 ; CODE-NEXT: Stmt_bb33(t0, 0);
 
 
-target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 target triple = "x86_64-unknown-linux-gnu"
 
 %struct.hoge = type { [23 x i16], [22 x i16], [14 x i16], [13 x i16] }