#pragma omp declare target

__attribute__ ((noipa))
int
gfx900 (void)
{
  return 0x900;
}

__attribute__ ((noipa))
int
gfx902 (void)
{
  return 0x902;
}

__attribute__ ((noipa))
int
gfx904 (void)
{
  return 0x904;
}

__attribute__ ((noipa))
int
gfx906 (void)
{
  return 0x906;
}

__attribute__ ((noipa))
int
gfx908 (void)
{
  return 0x908;
}

__attribute__ ((noipa))
int
gfx909 (void)
{
  return 0x909;
}

__attribute__ ((noipa))
int
gfx90a (void)
{
  return 0x90a;
}

__attribute__ ((noipa))
int
gfx90c (void)
{
  return 0x90c;
}

__attribute__ ((noipa))
int
gfx942 (void)
{
  return 0x942;
}

__attribute__ ((noipa))
int
gfx950 (void)
{
  return 0x950;
}

__attribute__ ((noipa))
int
gfx1030 (void)
{
  return 0x1030;
}

__attribute__ ((noipa))
int
gfx1031 (void)
{
  return 0x1031;
}

__attribute__ ((noipa))
int
gfx1032 (void)
{
  return 0x1032;
}

__attribute__ ((noipa))
int
gfx1033 (void)
{
  return 0x1033;
}

__attribute__ ((noipa))
int
gfx1034 (void)
{
  return 0x1034;
}

__attribute__ ((noipa))
int
gfx1035 (void)
{
  return 0x1035;
}

__attribute__ ((noipa))
int
gfx1036 (void)
{
  return 0x1036;
}

__attribute__ ((noipa))
int
gfx1100 (void)
{
  return 0x1100;
}

__attribute__ ((noipa))
int
gfx1101 (void)
{
  return 0x1101;
}

__attribute__ ((noipa))
int
gfx1102 (void)
{
  return 0x1102;
}

__attribute__ ((noipa))
int
gfx1103 (void)
{
  return 0x1103;
}

__attribute__ ((noipa))
int
gfx1150 (void)
{
  return 0x1150;
}

__attribute__ ((noipa))
int
gfx1151 (void)
{
  return 0x1151;
}

__attribute__ ((noipa))
int
gfx1152 (void)
{
  return 0x1152;
}

__attribute__ ((noipa))
int
gfx1153 (void)
{
  return 0x1153;
}

__attribute__ ((noipa))
int
gfx9_generic (void)
{
  return 0x90ff;
}

__attribute__ ((noipa))
int
gfx9_4_generic (void)
{
  return 0x94ff;
}

__attribute__ ((noipa))
int
gfx10_3_generic (void)
{
  return 0x103ff;
}

__attribute__ ((noipa))
int
gfx11_generic (void)
{
  return 0x110ff;
}


#pragma omp declare variant(gfx900) match(device = {isa("gfx900")})
#pragma omp declare variant(gfx902) match(device = {isa("gfx902")})
#pragma omp declare variant(gfx904) match(device = {isa("gfx904")})
#pragma omp declare variant(gfx906) match(device = {isa("gfx906")})
#pragma omp declare variant(gfx908) match(device = {isa("gfx908")})
#pragma omp declare variant(gfx909) match(device = {isa("gfx909")})
#pragma omp declare variant(gfx90a) match(device = {isa("gfx90a")})
#pragma omp declare variant(gfx90c) match(device = {isa("gfx90c")})
#pragma omp declare variant(gfx942) match(device = {isa("gfx942")})
#pragma omp declare variant(gfx950) match(device = {isa("gfx950")})
#pragma omp declare variant(gfx1030) match(device = {isa("gfx1030")})
#pragma omp declare variant(gfx1031) match(device = {isa("gfx1031")})
#pragma omp declare variant(gfx1032) match(device = {isa("gfx1032")})
#pragma omp declare variant(gfx1033) match(device = {isa("gfx1033")})
#pragma omp declare variant(gfx1034) match(device = {isa("gfx1034")})
#pragma omp declare variant(gfx1035) match(device = {isa("gfx1035")})
#pragma omp declare variant(gfx1036) match(device = {isa("gfx1036")})
#pragma omp declare variant(gfx1100) match(device = {isa("gfx1100")})
#pragma omp declare variant(gfx1101) match(device = {isa("gfx1101")})
#pragma omp declare variant(gfx1102) match(device = {isa("gfx1102")})
#pragma omp declare variant(gfx1103) match(device = {isa("gfx1103")})
#pragma omp declare variant(gfx1150) match(device = {isa("gfx1150")})
#pragma omp declare variant(gfx1151) match(device = {isa("gfx1151")})
#pragma omp declare variant(gfx1152) match(device = {isa("gfx1152")})
#pragma omp declare variant(gfx1153) match(device = {isa("gfx1153")})
#pragma omp declare variant(gfx9_generic) match(device = {isa("gfx9-generic")})
#pragma omp declare variant(gfx9_4_generic) match(device = {isa("gfx9-4-generic")})
#pragma omp declare variant(gfx10_3_generic) match(device = {isa("gfx10-3-generic")})
#pragma omp declare variant(gfx11_generic) match(device = {isa("gfx11-generic")})
__attribute__ ((noipa))
int
f (void)
{
  return 0;
}

#pragma omp end declare target

int
main (void)
{
  int v = 0;

#pragma omp target map(from : v)
  v = f ();

#ifdef OFFLOAD_DEVICE_GCN
  if (v == 0)
    __builtin_abort ();

  __builtin_printf ("AMDGCN accelerator: gfx%x\n", v);
#else
  if (v != 0)
    __builtin_abort ();
#endif

  return 0;
}