head	1.1;
branch	1.1.1;
access;
symbols
	netbsd-10-2-RELEASE:1.1.1.2
	netbsd-9-5-RELEASE:1.1.1.1
	netbsd-11-0-RELEASE:1.1.1.2
	netbsd-11-0-RC7:1.1.1.2
	netbsd-11-0-RC6:1.1.1.2
	netbsd-11-0-RC5:1.1.1.2
	netbsd-11-0-RC4:1.1.1.2
	netbsd-11-0-RC3:1.1.1.2
	netbsd-11-0-RC2:1.1.1.2
	netbsd-11-0-RC1:1.1.1.2
	perseant-exfatfs-base-20250801:1.1.1.2
	netbsd-11:1.1.1.2.0.10
	netbsd-11-base:1.1.1.2
	netbsd-10-1-RELEASE:1.1.1.2
	perseant-exfatfs-base-20240630:1.1.1.2
	perseant-exfatfs:1.1.1.2.0.8
	perseant-exfatfs-base:1.1.1.2
	netbsd-9-4-RELEASE:1.1.1.1
	netbsd-10-0-RELEASE:1.1.1.2
	netbsd-10-0-RC6:1.1.1.2
	netbsd-10-0-RC5:1.1.1.2
	netbsd-10-0-RC4:1.1.1.2
	netbsd-10-0-RC3:1.1.1.2
	netbsd-10-0-RC2:1.1.1.2
	netbsd-10-0-RC1:1.1.1.2
	netbsd-10:1.1.1.2.0.6
	netbsd-10-base:1.1.1.2
	netbsd-9-3-RELEASE:1.1.1.1
	cjep_sun2x:1.1.1.2.0.4
	cjep_sun2x-base:1.1.1.2
	cjep_staticlib_x-base1:1.1.1.2
	netbsd-9-2-RELEASE:1.1.1.1
	cjep_staticlib_x:1.1.1.2.0.2
	cjep_staticlib_x-base:1.1.1.2
	netbsd-9-1-RELEASE:1.1.1.1
	phil-wifi-20200421:1.1.1.2
	phil-wifi-20200411:1.1.1.2
	phil-wifi-20200406:1.1.1.2
	netbsd-9-0-RELEASE:1.1.1.1
	netbsd-9-0-RC2:1.1.1.1
	netbsd-9-0-RC1:1.1.1.1
	netbsd-9:1.1.1.1.0.6
	netbsd-9-base:1.1.1.1
	phil-wifi:1.1.1.1.0.4
	phil-wifi-20190609:1.1.1.1
	pgoyette-compat-merge-20190127:1.1.1.1.2.2
	pgoyette-compat-20190127:1.1.1.1
	pgoyette-compat-20190118:1.1.1.1
	pgoyette-compat-1226:1.1.1.1
	pgoyette-compat-1126:1.1.1.1
	pgoyette-compat-1020:1.1.1.1
	pgoyette-compat-0930:1.1.1.1
	pgoyette-compat-0906:1.1.1.1
	pgoyette-compat:1.1.1.1.0.2
	pgoyette-compat-0728:1.1.1.1
	clang-337282:1.1.1.1
	LLVM:1.1.1;
locks; strict;
comment	@// @;


1.1
date	2018.07.17.18.31.40;	author joerg;	state Exp;
branches
	1.1.1.1;
next	;
commitid	wDzL46ALjrCZgwKA;

1.1.1.1
date	2018.07.17.18.31.40;	author joerg;	state Exp;
branches
	1.1.1.1.2.1
	1.1.1.1.4.1;
next	1.1.1.2;
commitid	wDzL46ALjrCZgwKA;

1.1.1.2
date	2019.11.13.22.22.44;	author joerg;	state dead;
branches;
next	;
commitid	QD8YATxuNG34YJKB;

1.1.1.1.2.1
date	2018.07.17.18.31.40;	author pgoyette;	state dead;
branches;
next	1.1.1.1.2.2;
commitid	1UP1xAIUxv1ZgRLA;

1.1.1.1.2.2
date	2018.07.28.04.34.06;	author pgoyette;	state Exp;
branches;
next	;
commitid	1UP1xAIUxv1ZgRLA;

1.1.1.1.4.1
date	2018.07.17.18.31.40;	author christos;	state dead;
branches;
next	1.1.1.1.4.2;
commitid	jtc8rnCzWiEEHGqB;

1.1.1.1.4.2
date	2019.06.10.21.46.28;	author christos;	state Exp;
branches;
next	1.1.1.1.4.3;
commitid	jtc8rnCzWiEEHGqB;

1.1.1.1.4.3
date	2020.04.13.07.50.12;	author martin;	state dead;
branches;
next	;
commitid	X01YhRUPVUDaec4C;


desc
@@


1.1
log
@Initial revision
@
text
@// Only test codegen on target side, as private clause does not require any action on the host side
// Test target codegen - host bc file has to be created first.
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32

// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}

// expected-no-diagnostics
#ifndef HEADER
#define HEADER

template<typename tx, typename ty>
struct TT{
  tx X;
  ty Y;
  TT<tx, ty> operator*(const TT<tx, ty> &) { return *this; }
};

// TCHECK: [[S1:%.+]] = type { double }

int foo(int n) {
  int a = 0;
  short aa = 0;
  float b[10];
  float bn[n];
  double c[5][10];
  double cn[5][n];
  TT<long long, char> d;

  #pragma omp target reduction(*:a)
  {
  }

  // TCHECK: define weak void @@__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}})
  // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: store {{.+}}, {{.+}} [[A]],
  // TCHECK: load i32*, i32** [[A]],
  // TCHECK: ret void

#pragma omp target reduction(+:a)
  {
    a = 1;
  }

  // TCHECK:  define weak void @@__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}})
  // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: store {{.+}}, {{.+}} [[A]],
  // TCHECK: [[REF:%.+]] = load i32*, i32** [[A]],
  // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[REF]],
  // TCHECK: ret void

  #pragma omp target reduction(-:a, aa)
  {
    a = 1;
    aa = 1;
  }

  // TCHECK:  define weak void @@__omp_offloading_{{.+}}(i32*{{.+}} [[A:%.+]], i16*{{.+}} [[AA:%.+]])
  // TCHECK:  [[A:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK:  [[AA:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: store {{.+}}, {{.+}} [[A]],
  // TCHECK: store {{.+}}, {{.+}} [[AA]],
  // TCHECK: [[A_REF:%.+]] = load i32*, i32** [[A]],
  // TCHECK: [[AA_REF:%.+]] = load i16*, i16** [[AA]],
  // TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A_REF]],
  // TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[AA_REF]],
  // TCHECK:  ret void

  return a;
}


template<typename tx>
tx ftemplate(int n) {
  tx a = 0;
  short aa = 0;
  tx b[10];

#pragma omp target reduction(+:a,aa,b)
  {
    a = 1;
    aa = 1;
    b[2] = 1;
  }

  return a;
}

static
int fstatic(int n) {
  int a = 0;
  short aa = 0;
  char aaa = 0;
  int b[10];

#pragma omp target reduction(-:a,aa,aaa,b)
  {
    a = 1;
    aa = 1;
    aaa = 1;
    b[2] = 1;
  }

  return a;
}

// TCHECK: define weak void @@__omp_offloading_{{.+}}(i32*{{.+}}, i16*{{.+}}, i8*{{.+}}, [10 x i32]*{{.+}})
// TCHECK:  [[A:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK:  [[A2:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK:  [[A3:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK:  [[B:%.+]] = alloca [10 x i{{[0-9]+}}]*,
// TCHECK: store {{.+}}, {{.+}} [[A]],
// TCHECK: store {{.+}}, {{.+}} [[A2]],
// TCHECK: store {{.+}}, {{.+}} [[A3]],
// TCHECK: store {{.+}}, {{.+}} [[B]],
// TCHECK: [[A_REF:%.+]] = load i32*, i32** [[A]],
// TCHECK: [[AA_REF:%.+]] = load i16*, i16** [[AA]],
// TCHECK: [[A3_REF:%.+]] = load i8*, i8** [[A3]],
// TCHECK: [[B_REF:%.+]] = load {{.+}}*, {{.+}}** [[B]],
// TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A_REF]],
// TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[AA_REF]],
// TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A3_REF]],
// TCHECK:  [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
// TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
// TCHECK:  ret void

struct S1 {
  double a;

  int r1(int n){
    int b = n+1;
    short int c[2][n];

#pragma omp target reduction(max:b,c)
    {
      this->a = (double)b + 1.5;
      c[1][1] = ++a;
    }

    return c[1][1] + (int)b;
  }

  // TCHECK: define weak void @@__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i32*{{.+}}, i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]], i16*{{.+}})
  // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
  // TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
  // TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}},
  // TCHECK: [[C_ADDR:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]],
  // TCHECK: store i{{[0-9]+}}* {{.+}}, i{{[0-9]+}}** [[B_ADDR]],
  // TCHECK: store i{{[0-9]+}} [[VLA]], i{{[0-9]+}}* [[VLA_ADDR]],
  // TCHECK: store i{{[0-9]+}} [[VLA1]], i{{[0-9]+}}* [[VLA_ADDR2]],
  // TCHECK: store i{{[0-9]+}}* {{.+}}, i{{[0-9]+}}** [[C_ADDR]],
  // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]],
  // TCHECK: [[B_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[B_ADDR]],
  // TCHECK: [[VLA_ADDR_REF:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR]],
  // TCHECK: [[VLA_ADDR_REF2:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR2]],
  // TCHECK: [[C_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[C_ADDR]],

  // this->a = (double)b + 1.5;
  // TCHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B_REF]],
  // TCHECK: [[B_CONV:%.+]] = sitofp i{{[0-9]+}} [[B_VAL]] to double
  // TCHECK: [[NEW_A_VAL:%.+]] = fadd double [[B_CONV]], 1.5{{.+}}+00
  // TCHECK: [[A_FIELD:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
  // TCHECK: store double [[NEW_A_VAL]], double* [[A_FIELD]],

  // c[1][1] = ++a;
  // TCHECK: [[A_FIELD4:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
  // TCHECK: [[A_FIELD4_VAL:%.+]] = load double, double* [[A_FIELD4]],
  // TCHECK: [[A_FIELD_INC:%.+]] = fadd double [[A_FIELD4_VAL]], 1.0{{.+}}+00
  // TCHECK: store double [[A_FIELD_INC]], double* [[A_FIELD4]],
  // TCHECK: [[A_FIELD_INC_CONV:%.+]] = fptosi double [[A_FIELD_INC]] to i{{[0-9]+}}
  // TCHECK: [[C_IND:%.+]] = mul{{.+}} i{{[0-9]+}} 1, [[VLA_ADDR_REF2]]
  // TCHECK: [[C_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, i{{[0-9]+}}* [[C_REF]], i{{[0-9]+}} [[C_IND]]
  // TCHECK: [[C_1_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, i{{[0-9]+}}* [[C_1_REF]], i{{[0-9]+}} 1
  // TCHECK: store i{{[0-9]+}} [[A_FIELD_INC_CONV]], i{{[0-9]+}}* [[C_1_1_REF]],

  // finish
  // TCHECK: ret void
};


int bar(int n){
  int a = 0;
  a += foo(n);
  S1 S;
  a += S.r1(n);
  a += fstatic(n);
  a += ftemplate<int>(n);

  return a;
}

// template
// TCHECK: define weak void @@__omp_offloading_{{.+}}(i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}}, [10 x i32]*{{.+}})
// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}]*,
// TCHECK: store {{.+}}, {{.+}} [[A]],
// TCHECK: store {{.+}}, {{.+}} [[A2]],
// TCHECK: store {{.+}}, {{.+}} [[B]],
// TCHECK: [[A_REF:%.+]] = load i32*, i32** [[A]],
// TCHECK: [[AA_REF:%.+]] = load i16*, i16** [[AA]],
// TCHECK: [[B_REF:%.+]] = load {{.+}}*, {{.+}}** [[B]],
// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A_REF]],
// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[AA_REF]],
// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
// TCHECK: ret void

#endif
@


1.1.1.1
log
@Import clang r337282 from trunk
@
text
@@


1.1.1.2
log
@Mark old LLVM instance as dead.
@
text
@@


1.1.1.1.4.1
log
@file target_reduction_codegen.cpp was added on branch phil-wifi on 2019-06-10 21:46:28 +0000
@
text
@d1 225
@


1.1.1.1.4.2
log
@Sync with HEAD
@
text
@a0 225
// Only test codegen on target side, as private clause does not require any action on the host side
// Test target codegen - host bc file has to be created first.
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32

// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}

// expected-no-diagnostics
#ifndef HEADER
#define HEADER

template<typename tx, typename ty>
struct TT{
  tx X;
  ty Y;
  TT<tx, ty> operator*(const TT<tx, ty> &) { return *this; }
};

// TCHECK: [[S1:%.+]] = type { double }

int foo(int n) {
  int a = 0;
  short aa = 0;
  float b[10];
  float bn[n];
  double c[5][10];
  double cn[5][n];
  TT<long long, char> d;

  #pragma omp target reduction(*:a)
  {
  }

  // TCHECK: define weak void @@__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}})
  // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: store {{.+}}, {{.+}} [[A]],
  // TCHECK: load i32*, i32** [[A]],
  // TCHECK: ret void

#pragma omp target reduction(+:a)
  {
    a = 1;
  }

  // TCHECK:  define weak void @@__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}})
  // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: store {{.+}}, {{.+}} [[A]],
  // TCHECK: [[REF:%.+]] = load i32*, i32** [[A]],
  // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[REF]],
  // TCHECK: ret void

  #pragma omp target reduction(-:a, aa)
  {
    a = 1;
    aa = 1;
  }

  // TCHECK:  define weak void @@__omp_offloading_{{.+}}(i32*{{.+}} [[A:%.+]], i16*{{.+}} [[AA:%.+]])
  // TCHECK:  [[A:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK:  [[AA:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: store {{.+}}, {{.+}} [[A]],
  // TCHECK: store {{.+}}, {{.+}} [[AA]],
  // TCHECK: [[A_REF:%.+]] = load i32*, i32** [[A]],
  // TCHECK: [[AA_REF:%.+]] = load i16*, i16** [[AA]],
  // TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A_REF]],
  // TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[AA_REF]],
  // TCHECK:  ret void

  return a;
}


template<typename tx>
tx ftemplate(int n) {
  tx a = 0;
  short aa = 0;
  tx b[10];

#pragma omp target reduction(+:a,aa,b)
  {
    a = 1;
    aa = 1;
    b[2] = 1;
  }

  return a;
}

static
int fstatic(int n) {
  int a = 0;
  short aa = 0;
  char aaa = 0;
  int b[10];

#pragma omp target reduction(-:a,aa,aaa,b)
  {
    a = 1;
    aa = 1;
    aaa = 1;
    b[2] = 1;
  }

  return a;
}

// TCHECK: define weak void @@__omp_offloading_{{.+}}(i32*{{.+}}, i16*{{.+}}, i8*{{.+}}, [10 x i32]*{{.+}})
// TCHECK:  [[A:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK:  [[A2:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK:  [[A3:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK:  [[B:%.+]] = alloca [10 x i{{[0-9]+}}]*,
// TCHECK: store {{.+}}, {{.+}} [[A]],
// TCHECK: store {{.+}}, {{.+}} [[A2]],
// TCHECK: store {{.+}}, {{.+}} [[A3]],
// TCHECK: store {{.+}}, {{.+}} [[B]],
// TCHECK: [[A_REF:%.+]] = load i32*, i32** [[A]],
// TCHECK: [[AA_REF:%.+]] = load i16*, i16** [[AA]],
// TCHECK: [[A3_REF:%.+]] = load i8*, i8** [[A3]],
// TCHECK: [[B_REF:%.+]] = load {{.+}}*, {{.+}}** [[B]],
// TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A_REF]],
// TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[AA_REF]],
// TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A3_REF]],
// TCHECK:  [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
// TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
// TCHECK:  ret void

struct S1 {
  double a;

  int r1(int n){
    int b = n+1;
    short int c[2][n];

#pragma omp target reduction(max:b,c)
    {
      this->a = (double)b + 1.5;
      c[1][1] = ++a;
    }

    return c[1][1] + (int)b;
  }

  // TCHECK: define weak void @@__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i32*{{.+}}, i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]], i16*{{.+}})
  // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
  // TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
  // TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}},
  // TCHECK: [[C_ADDR:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]],
  // TCHECK: store i{{[0-9]+}}* {{.+}}, i{{[0-9]+}}** [[B_ADDR]],
  // TCHECK: store i{{[0-9]+}} [[VLA]], i{{[0-9]+}}* [[VLA_ADDR]],
  // TCHECK: store i{{[0-9]+}} [[VLA1]], i{{[0-9]+}}* [[VLA_ADDR2]],
  // TCHECK: store i{{[0-9]+}}* {{.+}}, i{{[0-9]+}}** [[C_ADDR]],
  // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]],
  // TCHECK: [[B_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[B_ADDR]],
  // TCHECK: [[VLA_ADDR_REF:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR]],
  // TCHECK: [[VLA_ADDR_REF2:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR2]],
  // TCHECK: [[C_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[C_ADDR]],

  // this->a = (double)b + 1.5;
  // TCHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B_REF]],
  // TCHECK: [[B_CONV:%.+]] = sitofp i{{[0-9]+}} [[B_VAL]] to double
  // TCHECK: [[NEW_A_VAL:%.+]] = fadd double [[B_CONV]], 1.5{{.+}}+00
  // TCHECK: [[A_FIELD:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
  // TCHECK: store double [[NEW_A_VAL]], double* [[A_FIELD]],

  // c[1][1] = ++a;
  // TCHECK: [[A_FIELD4:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
  // TCHECK: [[A_FIELD4_VAL:%.+]] = load double, double* [[A_FIELD4]],
  // TCHECK: [[A_FIELD_INC:%.+]] = fadd double [[A_FIELD4_VAL]], 1.0{{.+}}+00
  // TCHECK: store double [[A_FIELD_INC]], double* [[A_FIELD4]],
  // TCHECK: [[A_FIELD_INC_CONV:%.+]] = fptosi double [[A_FIELD_INC]] to i{{[0-9]+}}
  // TCHECK: [[C_IND:%.+]] = mul{{.+}} i{{[0-9]+}} 1, [[VLA_ADDR_REF2]]
  // TCHECK: [[C_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, i{{[0-9]+}}* [[C_REF]], i{{[0-9]+}} [[C_IND]]
  // TCHECK: [[C_1_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, i{{[0-9]+}}* [[C_1_REF]], i{{[0-9]+}} 1
  // TCHECK: store i{{[0-9]+}} [[A_FIELD_INC_CONV]], i{{[0-9]+}}* [[C_1_1_REF]],

  // finish
  // TCHECK: ret void
};


int bar(int n){
  int a = 0;
  a += foo(n);
  S1 S;
  a += S.r1(n);
  a += fstatic(n);
  a += ftemplate<int>(n);

  return a;
}

// template
// TCHECK: define weak void @@__omp_offloading_{{.+}}(i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}}, [10 x i32]*{{.+}})
// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}]*,
// TCHECK: store {{.+}}, {{.+}} [[A]],
// TCHECK: store {{.+}}, {{.+}} [[A2]],
// TCHECK: store {{.+}}, {{.+}} [[B]],
// TCHECK: [[A_REF:%.+]] = load i32*, i32** [[A]],
// TCHECK: [[AA_REF:%.+]] = load i16*, i16** [[AA]],
// TCHECK: [[B_REF:%.+]] = load {{.+}}*, {{.+}}** [[B]],
// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A_REF]],
// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[AA_REF]],
// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
// TCHECK: ret void

#endif
@


1.1.1.1.4.3
log
@Mostly merge changes from HEAD upto 20200411
@
text
@@


1.1.1.1.2.1
log
@file target_reduction_codegen.cpp was added on branch pgoyette-compat on 2018-07-28 04:34:06 +0000
@
text
@d1 225
@


1.1.1.1.2.2
log
@Sync with HEAD
@
text
@a0 225
// Only test codegen on target side, as private clause does not require any action on the host side
// Test target codegen - host bc file has to be created first.
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32
// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
// RUN: %clang_cc1 -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32

// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}

// expected-no-diagnostics
#ifndef HEADER
#define HEADER

template<typename tx, typename ty>
struct TT{
  tx X;
  ty Y;
  TT<tx, ty> operator*(const TT<tx, ty> &) { return *this; }
};

// TCHECK: [[S1:%.+]] = type { double }

int foo(int n) {
  int a = 0;
  short aa = 0;
  float b[10];
  float bn[n];
  double c[5][10];
  double cn[5][n];
  TT<long long, char> d;

  #pragma omp target reduction(*:a)
  {
  }

  // TCHECK: define weak void @@__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}})
  // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: store {{.+}}, {{.+}} [[A]],
  // TCHECK: load i32*, i32** [[A]],
  // TCHECK: ret void

#pragma omp target reduction(+:a)
  {
    a = 1;
  }

  // TCHECK:  define weak void @@__omp_offloading_{{.+}}(i32*{{.+}} %{{.+}})
  // TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: store {{.+}}, {{.+}} [[A]],
  // TCHECK: [[REF:%.+]] = load i32*, i32** [[A]],
  // TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[REF]],
  // TCHECK: ret void

  #pragma omp target reduction(-:a, aa)
  {
    a = 1;
    aa = 1;
  }

  // TCHECK:  define weak void @@__omp_offloading_{{.+}}(i32*{{.+}} [[A:%.+]], i16*{{.+}} [[AA:%.+]])
  // TCHECK:  [[A:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK:  [[AA:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: store {{.+}}, {{.+}} [[A]],
  // TCHECK: store {{.+}}, {{.+}} [[AA]],
  // TCHECK: [[A_REF:%.+]] = load i32*, i32** [[A]],
  // TCHECK: [[AA_REF:%.+]] = load i16*, i16** [[AA]],
  // TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A_REF]],
  // TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[AA_REF]],
  // TCHECK:  ret void

  return a;
}


template<typename tx>
tx ftemplate(int n) {
  tx a = 0;
  short aa = 0;
  tx b[10];

#pragma omp target reduction(+:a,aa,b)
  {
    a = 1;
    aa = 1;
    b[2] = 1;
  }

  return a;
}

static
int fstatic(int n) {
  int a = 0;
  short aa = 0;
  char aaa = 0;
  int b[10];

#pragma omp target reduction(-:a,aa,aaa,b)
  {
    a = 1;
    aa = 1;
    aaa = 1;
    b[2] = 1;
  }

  return a;
}

// TCHECK: define weak void @@__omp_offloading_{{.+}}(i32*{{.+}}, i16*{{.+}}, i8*{{.+}}, [10 x i32]*{{.+}})
// TCHECK:  [[A:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK:  [[A2:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK:  [[A3:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK:  [[B:%.+]] = alloca [10 x i{{[0-9]+}}]*,
// TCHECK: store {{.+}}, {{.+}} [[A]],
// TCHECK: store {{.+}}, {{.+}} [[A2]],
// TCHECK: store {{.+}}, {{.+}} [[A3]],
// TCHECK: store {{.+}}, {{.+}} [[B]],
// TCHECK: [[A_REF:%.+]] = load i32*, i32** [[A]],
// TCHECK: [[AA_REF:%.+]] = load i16*, i16** [[AA]],
// TCHECK: [[A3_REF:%.+]] = load i8*, i8** [[A3]],
// TCHECK: [[B_REF:%.+]] = load {{.+}}*, {{.+}}** [[B]],
// TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A_REF]],
// TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[AA_REF]],
// TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A3_REF]],
// TCHECK:  [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
// TCHECK:  store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
// TCHECK:  ret void

struct S1 {
  double a;

  int r1(int n){
    int b = n+1;
    short int c[2][n];

#pragma omp target reduction(max:b,c)
    {
      this->a = (double)b + 1.5;
      c[1][1] = ++a;
    }

    return c[1][1] + (int)b;
  }

  // TCHECK: define weak void @@__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i32*{{.+}}, i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]], i16*{{.+}})
  // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
  // TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: [[VLA_ADDR:%.+]] = alloca i{{[0-9]+}},
  // TCHECK: [[VLA_ADDR2:%.+]] = alloca i{{[0-9]+}},
  // TCHECK: [[C_ADDR:%.+]] = alloca i{{[0-9]+}}*,
  // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]],
  // TCHECK: store i{{[0-9]+}}* {{.+}}, i{{[0-9]+}}** [[B_ADDR]],
  // TCHECK: store i{{[0-9]+}} [[VLA]], i{{[0-9]+}}* [[VLA_ADDR]],
  // TCHECK: store i{{[0-9]+}} [[VLA1]], i{{[0-9]+}}* [[VLA_ADDR2]],
  // TCHECK: store i{{[0-9]+}}* {{.+}}, i{{[0-9]+}}** [[C_ADDR]],
  // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]],
  // TCHECK: [[B_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[B_ADDR]],
  // TCHECK: [[VLA_ADDR_REF:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR]],
  // TCHECK: [[VLA_ADDR_REF2:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[VLA_ADDR2]],
  // TCHECK: [[C_REF:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[C_ADDR]],

  // this->a = (double)b + 1.5;
  // TCHECK: [[B_VAL:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[B_REF]],
  // TCHECK: [[B_CONV:%.+]] = sitofp i{{[0-9]+}} [[B_VAL]] to double
  // TCHECK: [[NEW_A_VAL:%.+]] = fadd double [[B_CONV]], 1.5{{.+}}+00
  // TCHECK: [[A_FIELD:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
  // TCHECK: store double [[NEW_A_VAL]], double* [[A_FIELD]],

  // c[1][1] = ++a;
  // TCHECK: [[A_FIELD4:%.+]] = getelementptr inbounds [[S1]], [[S1]]* [[TH_ADDR_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
  // TCHECK: [[A_FIELD4_VAL:%.+]] = load double, double* [[A_FIELD4]],
  // TCHECK: [[A_FIELD_INC:%.+]] = fadd double [[A_FIELD4_VAL]], 1.0{{.+}}+00
  // TCHECK: store double [[A_FIELD_INC]], double* [[A_FIELD4]],
  // TCHECK: [[A_FIELD_INC_CONV:%.+]] = fptosi double [[A_FIELD_INC]] to i{{[0-9]+}}
  // TCHECK: [[C_IND:%.+]] = mul{{.+}} i{{[0-9]+}} 1, [[VLA_ADDR_REF2]]
  // TCHECK: [[C_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, i{{[0-9]+}}* [[C_REF]], i{{[0-9]+}} [[C_IND]]
  // TCHECK: [[C_1_1_REF:%.+]] = getelementptr inbounds i{{[0-9]+}}, i{{[0-9]+}}* [[C_1_REF]], i{{[0-9]+}} 1
  // TCHECK: store i{{[0-9]+}} [[A_FIELD_INC_CONV]], i{{[0-9]+}}* [[C_1_1_REF]],

  // finish
  // TCHECK: ret void
};


int bar(int n){
  int a = 0;
  a += foo(n);
  S1 S;
  a += S.r1(n);
  a += fstatic(n);
  a += ftemplate<int>(n);

  return a;
}

// template
// TCHECK: define weak void @@__omp_offloading_{{.+}}(i{{[0-9]+}}*{{.+}}, i{{[0-9]+}}*{{.+}}, [10 x i32]*{{.+}})
// TCHECK: [[A:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK: [[A2:%.+]] = alloca i{{[0-9]+}}*,
// TCHECK: [[B:%.+]] = alloca [10 x i{{[0-9]+}}]*,
// TCHECK: store {{.+}}, {{.+}} [[A]],
// TCHECK: store {{.+}}, {{.+}} [[A2]],
// TCHECK: store {{.+}}, {{.+}} [[B]],
// TCHECK: [[A_REF:%.+]] = load i32*, i32** [[A]],
// TCHECK: [[AA_REF:%.+]] = load i16*, i16** [[AA]],
// TCHECK: [[B_REF:%.+]] = load {{.+}}*, {{.+}}** [[B]],
// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[A_REF]],
// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[AA_REF]],
// TCHECK: [[B_GEP:%.+]] = getelementptr inbounds [10 x i{{[0-9]+}}], [10 x i{{[0-9]+}}]* [[B_REF]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
// TCHECK: store i{{[0-9]+}} 1, i{{[0-9]+}}* [[B_GEP]],
// TCHECK: ret void

#endif
@