/**
 *
 * bash版キャリーチェーンのC言語版のGPU/CUDA移植版
 *
 詳しい説明はこちらをどうぞ
 https://suzukiiichiro.github.io/search/?keyword=Ｎクイーン問題
 *
アーキテクチャの指定（なくても問題なし、あれば高速）
-arch=sm_13 or -arch=sm_61

CPUの再帰での実行
$ nvcc -O3 -arch=sm_61 05CUDA_CarryChain.cu && ./a.out -r

CPUの非再帰での実行
$ nvcc -O3 -arch=sm_61 05CUDA_CarryChain.cu && ./a.out -c

GPUのシングルスレッド
$ nvcc -O3 -arch=sm_61 05CUDA_CarryChain.cu && ./a.out -g

GPUのマルチスレッド
$ nvcc -O3 -arch=sm_61 05CUDA_CarryChain.cu && ./a.out -n
*/
#include <iostream>
#include <vector>
#include <stdio.h>
#include <stdlib.h>
#include <stdbool.h>
#include <math.h>
#include <string.h>
#include <time.h>
#include <sys/time.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#define THREAD_NUM		96
#define MAX 27
// システムによって以下のマクロが必要であればコメントを外してください。
//#define UINT64_C(c) c ## ULL
//
// グローバル変数
unsigned long TOTAL=0; 
unsigned long UNIQUE=0;
// キャリーチェーン 非再帰版
// 構造体
typedef struct
{
  unsigned int size;
  unsigned int pres_a[930]; 
  unsigned int pres_b[930];
  // uint64_t COUNTER[3];      
  // //カウンター配列
  // unsigned int COUNT2;
  // unsigned int COUNT4;
  // unsigned int COUNT8;
}Global; Global g;
// 構造体
typedef struct Board
{
  uint64_t row;
  uint64_t down;
  uint64_t left;
  uint64_t right;
  long long x[MAX];
}Board ;
typedef struct Local
{
  unsigned int size;
  struct Board B;
  struct Board nB;
  struct Board eB;
  struct Board sB;
  struct Board wB;
  unsigned n;
  unsigned e;
  unsigned s;
  unsigned w;
  uint64_t dimx;
  uint64_t dimy;
  uint64_t COUNTER[3];      
  //カウンター配列
  unsigned int COUNT2;
  unsigned int COUNT4;
  unsigned int COUNT8;
  unsigned int STEPS;
}Local;
/**
  CPU/CPUR 再帰・非再帰共通
  */
// チェーンのリストを作成
void listChain()
{
  unsigned int idx=0;
  for(unsigned int a=0;a<(unsigned)g.size;++a){
    for(unsigned int b=0;b<(unsigned)g.size;++b){
      if(((a>=b)&&(a-b)<=1)||((b>a)&&(b-a)<=1)){ continue; }
      g.pres_a[idx]=a;
      g.pres_b[idx]=b;
      ++idx;
    }
  }
}
/**
  CPU 非再帰
*/
// クイーンの効きをチェック
bool placement(void* args)
{
  Local *l=(Local *)args;
  if(l->B.x[l->dimx]==l->dimy){ return true;  }  
  if (l->B.x[0]==0){
    if (l->B.x[1]!=(uint64_t)-1){
      if((l->B.x[1]>=l->dimx)&&(l->dimy==1)){ return false; }
    }
  }else{
    if( (l->B.x[0]!=(uint64_t)-1) ){
      if(( (l->dimx<l->B.x[0]||l->dimx>=g.size-l->B.x[0])
        && (l->dimy==0 || l->dimy==g.size-1)
      )){ return 0; } 
      if ((  (l->dimx==g.size-1)&&((l->dimy<=l->B.x[0])||
          l->dimy>=g.size-l->B.x[0]))){
        return 0;
      } 
    }
  }
  l->B.x[l->dimx]=l->dimy;                    //xは行 yは列
  uint64_t row=UINT64_C(1)<<l->dimx;
  uint64_t down=UINT64_C(1)<<l->dimy;
  uint64_t left=UINT64_C(1)<<(g.size-1-l->dimx+l->dimy); //右上から左下
  uint64_t right=UINT64_C(1)<<(l->dimx+l->dimy);       // 左上から右下
  if((l->B.row&row)||(l->B.down&down)||(l->B.left&left)||(l->B.right&right)){ return false; }     
  l->B.row|=row; l->B.down|=down; l->B.left|=left; l->B.right|=right;
  return true;
}
//非再帰
uint64_t solve(int size,int current,uint64_t row,uint64_t left,uint64_t down,uint64_t right)
{
  uint64_t row_a[MAX];
  uint64_t right_a[MAX];
  uint64_t left_a[MAX];
  uint64_t down_a[MAX];
  uint64_t bitmap_a[MAX];
  for (int i=0;i<size;i++){
    row_a[i]=0;
    left_a[i]=0;
    down_a[i]=0;
    right_a[i]=0;
    bitmap_a[i]=0;
  }
  row_a[current]=row;
  left_a[current]=left;
  down_a[current]=down;
  right_a[current]=right;
  uint64_t bitmap=bitmap_a[current]=~(left_a[current]|down_a[current]|right_a[current]);
  uint64_t total=0;
  uint64_t bit;

  while(current>-1){
    if((bitmap!=0||row&1)&&current<size){
      if(!(down+1)){
        total++;
        current--;
        row=row_a[current];
        left=left_a[current];
        right=right_a[current];
        down=down_a[current];
        bitmap=bitmap_a[current];
        continue;
      }else if(row&1){
        while( row&1 ){
          row>>=1;
          left<<=1;
          right>>=1;
        }
        bitmap=~(left|down|right);  //再帰に必要な変数は必ず定義する必要があります。
        continue;
      }else{
        bit=-bitmap&bitmap;
        bitmap=bitmap^bit;
        if(current<size){
          row_a[current]=row;
          left_a[current]=left;
          down_a[current]=down;
          right_a[current]=right;
          bitmap_a[current]=bitmap;
          current++;
        }
        row>>=1;      //１行下に移動する
        left=(left|bit)<<1;
        down=down|bit;
        right=(right|bit)>>1;
        bitmap=~(left|down|right);  //再帰に必要な変数は必ず定義する必要があります。
      }
    }else{
      current--;
      row=row_a[current];
      left=left_a[current];
      right=right_a[current];
      down=down_a[current];
      bitmap=bitmap_a[current];
    }
  }
  return total;
}
//非再帰 対称解除法
void carryChain_symmetry(void* args)
{
  Local *l=(Local *)args;
  // 対称解除法
  unsigned const int ww=(g.size-2)*(g.size-1)-1-l->w;
  unsigned const int w2=(g.size-2)*(g.size-1)-1;
  // # 対角線上の反転が小さいかどうか確認する
  if((l->s==ww)&&(l->n<(w2-l->e))){ return ; }
  // # 垂直方向の中心に対する反転が小さいかを確認
  if((l->e==ww)&&(l->n>(w2-l->n))){ return; }
  // # 斜め下方向への反転が小さいかをチェックする
  if((l->n==ww)&&(l->e>(w2-l->s))){ return; }
  // 枝刈り １行目が角の場合回転対称チェックせずCOUNT8にする
  if(l->B.x[0]==0){
    l->COUNTER[l->COUNT8]+=solve(g.size,0,l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return ;
  }
  // n,e,s==w の場合は最小値を確認する。右回転で同じ場合は、
  // w=n=e=sでなければ値が小さいのでskip  w=n=e=sであれば90度回転で同じ可能性
  if(l->s==l->w){ if((l->n!=l->w)||(l->e!=l->w)){ return; }
    l->COUNTER[l->COUNT2]+=solve(g.size,0,l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return;
  }
  // e==wは180度回転して同じ 180度回転して同じ時n>=sの時はsmaller?
  if((l->e==l->w)&&(l->n>=l->s)){ if(l->n>l->s){ return; }
    l->COUNTER[l->COUNT4]+=solve(g.size,0,l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return;
  }
  l->COUNTER[l->COUNT8]+=solve(g.size,0,l->B.row>>2,
  l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
  return;
}
//非再帰  pthread run()
void thread_run(void* args)
{
  Local *l=(Local *)args;

  // memcpy(&l->B,&l->wB,sizeof(Board));       // B=wB;
  l->B=l->wB;
  l->dimx=0; l->dimy=g.pres_a[l->w];
  //if(!placement(l)){ continue; }
  if(!placement(l)){ return; }
  l->dimx=1; l->dimy=g.pres_b[l->w];
  // if(!placement(l)){ continue; }
  if(!placement(l)){ return; }
  //２ 左２行に置く
  // memcpy(&l->nB,&l->B,sizeof(Board));       // nB=B;
  l->nB=l->B;
  for(l->n=l->w;l->n<(g.size-2)*(g.size-1)-l->w;++l->n){
    // memcpy(&l->B,&l->nB,sizeof(Board));     // B=nB;
    l->B=l->nB;
    l->dimx=g.pres_a[l->n]; l->dimy=g.size-1;
    if(!placement(l)){ continue; }
    l->dimx=g.pres_b[l->n]; l->dimy=g.size-2;
    if(!placement(l)){ continue; }
    // ３ 下２行に置く
    // memcpy(&l->eB,&l->B,sizeof(Board));     // eB=B;
    l->eB=l->B;
    for(l->e=l->w;l->e<(g.size-2)*(g.size-1)-l->w;++l->e){
      // memcpy(&l->B,&l->eB,sizeof(Board));   // B=eB;
      l->B=l->eB;
      l->dimx=g.size-1; l->dimy=g.size-1-g.pres_a[l->e];
      if(!placement(l)){ continue; }
      l->dimx=g.size-2; l->dimy=g.size-1-g.pres_b[l->e];
      if(!placement(l)){ continue; }
      // ４ 右２列に置く
      // memcpy(&l->sB,&l->B,sizeof(Board));   // sB=B;
      l->sB=l->B;
      for(l->s=l->w;l->s<(g.size-2)*(g.size-1)-l->w;++l->s){
        // memcpy(&l->B,&l->sB,sizeof(Board)); // B=sB;
        l->B=l->sB;
        l->dimx=g.size-1-g.pres_a[l->s]; l->dimy=0;
        if(!placement(l)){ continue; }
        l->dimx=g.size-1-g.pres_b[l->s]; l->dimy=1;
        if(!placement(l)){ continue; }
        // 対称解除法
        carryChain_symmetry(l);
      } //w
    } //e
  } //n
}
//非再帰  チェーンのビルド
void buildChain()
{
  Local l[(g.size/2)*(g.size-3)];

  // カウンターの初期化
  l->COUNT2=0; l->COUNT4=1; l->COUNT8=2;
  l->COUNTER[l->COUNT2]=l->COUNTER[l->COUNT4]=l->COUNTER[l->COUNT8]=0;
  // Board の初期化 nB,eB,sB,wB;
  l->B.row=l->B.down=l->B.left=l->B.right=0;
  // Board x[]の初期化
  for(unsigned int i=0;i<g.size;++i){ l->B.x[i]=-1; }
  //１ 上２行に置く
  // memcpy(&l->wB,&l->B,sizeof(Board));         // wB=B;
  l->wB=l->B;
  for(l->w=0;l->w<=(unsigned)(g.size/2)*(g.size-3);++l->w){
    thread_run(&l);
  } //w
  /**
   * 集計
   */
  UNIQUE= l->COUNTER[l->COUNT2]+
          l->COUNTER[l->COUNT4]+
          l->COUNTER[l->COUNT8];
  TOTAL=  l->COUNTER[l->COUNT2]*2+
          l->COUNTER[l->COUNT4]*4+
          l->COUNTER[l->COUNT8]*8;
}
//非再帰  キャリーチェーン
void carryChain()
{
  listChain();  //チェーンのリストを作成
  buildChain(); // チェーンのビルド
  // calcChain(&l);  // 集計
}
/**
  CPUR 再帰
  */
//再帰 ボード外側２列を除く内側のクイーン配置処理
uint64_t solveR(uint64_t row,uint64_t left,uint64_t down,uint64_t right)
{
  if(down+1==0){ return  1; }
  while((row&1)!=0) { 
    row>>=1;
    left<<=1;
    right>>=1;
  }
  row>>=1;
  uint64_t total=0;
  for(uint64_t carryChain=~(left|down|right);carryChain!=0;){
    uint64_t const bit=carryChain&-carryChain;
    total+=solveR(row,(left|bit)<<1,down|bit,(right|bit)>>1);
    carryChain^=bit;
  }
  return total;
} 
//再帰 対称解除法
void carryChain_symmetryR(void* args)
{
  Local *l=(Local *)args;
  // 対称解除法
  unsigned const int ww=(g.size-2)*(g.size-1)-1-l->w;
  unsigned const int w2=(g.size-2)*(g.size-1)-1;
  // # 対角線上の反転が小さいかどうか確認する
  if((l->s==ww)&&(l->n<(w2-l->e))){ return ; }
  // # 垂直方向の中心に対する反転が小さいかを確認
  if((l->e==ww)&&(l->n>(w2-l->n))){ return; }
  // # 斜め下方向への反転が小さいかをチェックする
  if((l->n==ww)&&(l->e>(w2-l->s))){ return; }
  // 枝刈り １行目が角の場合回転対称チェックせずCOUNT8にする
  if(l->B.x[0]==0){
    l->COUNTER[l->COUNT8]+=solveR(l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return ;
  }
  // n,e,s==w の場合は最小値を確認する。右回転で同じ場合は、
  // w=n=e=sでなければ値が小さいのでskip  w=n=e=sであれば90度回転で同じ可能性
  if(l->s==l->w){ if((l->n!=l->w)||(l->e!=l->w)){ return; }
    l->COUNTER[l->COUNT2]+=solveR(l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return;
  }
  // e==wは180度回転して同じ 180度回転して同じ時n>=sの時はsmaller?
  if((l->e==l->w)&&(l->n>=l->s)){ if(l->n>l->s){ return; }
    l->COUNTER[l->COUNT4]+=solveR(l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return;
  }
  l->COUNTER[l->COUNT8]+=solveR(l->B.row>>2,
  l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
  return;
}
//再帰  pthread run()
void thread_runR(void* args)
{
  Local *l=(Local *)args;

  // memcpy(&l->B,&l->wB,sizeof(Board));       // B=wB;
  l->B=l->wB;
  l->dimx=0; l->dimy=g.pres_a[l->w];
  //if(!placement(l)){ continue; }
  if(!placement(l)){ return; }
  l->dimx=1; l->dimy=g.pres_b[l->w];
  // if(!placement(l)){ continue; }
  if(!placement(l)){ return; }
  //２ 左２行に置く
  // memcpy(&l->nB,&l->B,sizeof(Board));       // nB=B;
  l->nB=l->B;
  for(l->n=l->w;l->n<(g.size-2)*(g.size-1)-l->w;++l->n){
    // memcpy(&l->B,&l->nB,sizeof(Board));     // B=nB;
    l->B=l->nB;
    l->dimx=g.pres_a[l->n]; l->dimy=g.size-1;
    if(!placement(l)){ continue; }
    l->dimx=g.pres_b[l->n]; l->dimy=g.size-2;
    if(!placement(l)){ continue; }
    // ３ 下２行に置く
    // memcpy(&l->eB,&l->B,sizeof(Board));     // eB=B;
    l->eB=l->B;
    for(l->e=l->w;l->e<(g.size-2)*(g.size-1)-l->w;++l->e){
      // memcpy(&l->B,&l->eB,sizeof(Board));   // B=eB;
      l->B=l->eB;
      l->dimx=g.size-1; l->dimy=g.size-1-g.pres_a[l->e];
      if(!placement(l)){ continue; }
      l->dimx=g.size-2; l->dimy=g.size-1-g.pres_b[l->e];
      if(!placement(l)){ continue; }
      // ４ 右２列に置く
      // memcpy(&l->sB,&l->B,sizeof(Board));   // sB=B;
      l->sB=l->B;
      for(l->s=l->w;l->s<(g.size-2)*(g.size-1)-l->w;++l->s){
        // memcpy(&l->B,&l->sB,sizeof(Board)); // B=sB;
        l->B=l->sB;
        l->dimx=g.size-1-g.pres_a[l->s]; l->dimy=0;
        if(!placement(l)){ continue; }
        l->dimx=g.size-1-g.pres_b[l->s]; l->dimy=1;
        if(!placement(l)){ continue; }
        // 対称解除法
        carryChain_symmetryR(l);
      } //w
    } //e
  } //n
}
//再帰  チェーンのビルド
void buildChainR()
{
  Local l[(g.size/2)*(g.size-3)];

  // カウンターの初期化
  l->COUNT2=0; l->COUNT4=1; l->COUNT8=2;
  l->COUNTER[l->COUNT2]=l->COUNTER[l->COUNT4]=l->COUNTER[l->COUNT8]=0;
  // Board の初期化 nB,eB,sB,wB;
  l->B.row=l->B.down=l->B.left=l->B.right=0;
  // Board x[]の初期化
  for(unsigned int i=0;i<g.size;++i){ l->B.x[i]=-1; }
  //１ 上２行に置く
  // memcpy(&l->wB,&l->B,sizeof(Board));         // wB=B;
  l->wB=l->B;
  for(l->w=0;l->w<=(unsigned)(g.size/2)*(g.size-3);++l->w){
    thread_runR(&l);
  } //w
  /**
   * 集計
   */
  UNIQUE= l->COUNTER[l->COUNT2]+
          l->COUNTER[l->COUNT4]+
          l->COUNTER[l->COUNT8];
  TOTAL=  l->COUNTER[l->COUNT2]*2+
          l->COUNTER[l->COUNT4]*4+
          l->COUNTER[l->COUNT8]*8;
}
//再帰  キャリーチェーン
void carryChainR()
{
  listChain();  //チェーンのリストを作成
  buildChainR(); // チェーンのビルド
  // calcChain(&l);  // 集計
}
/**
  GPU 
 */
// GPU クイーンの効きをチェック
bool GPU_placement(void* args)
{
  Local *l=(Local *)args;
  if(l->B.x[l->dimx]==l->dimy){ return true;  }  
  if (l->B.x[0]==0){
    if (l->B.x[1]!=(uint64_t)-1){
      if((l->B.x[1]>=l->dimx)&&(l->dimy==1)){ return false; }
    }
  }else{
    if( (l->B.x[0]!=(uint64_t)-1) ){
      if(( (l->dimx<l->B.x[0]||l->dimx>=g.size-l->B.x[0])
        && (l->dimy==0 || l->dimy==g.size-1)
      )){ return 0; } 
      if ((  (l->dimx==g.size-1)&&((l->dimy<=l->B.x[0])||
          l->dimy>=g.size-l->B.x[0]))){
        return 0;
      } 
    }
  }
  l->B.x[l->dimx]=l->dimy;                    //xは行 yは列
  uint64_t row=UINT64_C(1)<<l->dimx;
  uint64_t down=UINT64_C(1)<<l->dimy;
  uint64_t left=UINT64_C(1)<<(g.size-1-l->dimx+l->dimy); //右上から左下
  uint64_t right=UINT64_C(1)<<(l->dimx+l->dimy);       // 左上から右下
  if((l->B.row&row)||(l->B.down&down)||(l->B.left&left)||(l->B.right&right)){ return false; }     
  l->B.row|=row; l->B.down|=down; l->B.left|=left; l->B.right|=right;
  return true;
}
//GPU 再帰 ボード外側２列を除く内側のクイーン配置処理
uint64_t GPU_solveR(uint64_t row,uint64_t left,uint64_t down,uint64_t right)
{
  if(down+1==0){ return  1; }
  while((row&1)!=0) { 
    row>>=1;
    left<<=1;
    right>>=1;
  }
  row>>=1;
  uint64_t total=0;
  for(uint64_t carryChain=~(left|down|right);carryChain!=0;){
    uint64_t const bit=carryChain&-carryChain;
    total+=solveR(row,(left|bit)<<1,down|bit,(right|bit)>>1);
    carryChain^=bit;
  }
  return total;
} 
//GPU 再帰 対称解除法
void GPU_carryChain_symmetryR(void* args)
{
  Local *l=(Local *)args;
  // 対称解除法
  unsigned const int ww=(g.size-2)*(g.size-1)-1-l->w;
  unsigned const int w2=(g.size-2)*(g.size-1)-1;
  // # 対角線上の反転が小さいかどうか確認する
  if((l->s==ww)&&(l->n<(w2-l->e))){ return ; }
  // # 垂直方向の中心に対する反転が小さいかを確認
  if((l->e==ww)&&(l->n>(w2-l->n))){ return; }
  // # 斜め下方向への反転が小さいかをチェックする
  if((l->n==ww)&&(l->e>(w2-l->s))){ return; }
  // 枝刈り １行目が角の場合回転対称チェックせずCOUNT8にする
  if(l->B.x[0]==0){
    l->COUNTER[l->COUNT8]+=solveR(l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return ;
  }
  // n,e,s==w の場合は最小値を確認する。右回転で同じ場合は、
  // w=n=e=sでなければ値が小さいのでskip  w=n=e=sであれば90度回転で同じ可能性
  if(l->s==l->w){ if((l->n!=l->w)||(l->e!=l->w)){ return; }
    l->COUNTER[l->COUNT2]+=solveR(l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return;
  }
  // e==wは180度回転して同じ 180度回転して同じ時n>=sの時はsmaller?
  if((l->e==l->w)&&(l->n>=l->s)){ if(l->n>l->s){ return; }
    l->COUNTER[l->COUNT4]+=solveR(l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return;
  }
  l->COUNTER[l->COUNT8]+=solveR(l->B.row>>2,
  l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
  return;
}
//GPU 再帰  pthread run()
void GPU_thread_runR(void* args)
{
  Local *l=(Local *)args;

  // memcpy(&l->B,&l->wB,sizeof(Board));       // B=wB;
  l->B=l->wB;
  l->dimx=0; l->dimy=g.pres_a[l->w];
  //if(!GPU_placement(l)){ continue; }
  if(!GPU_placement(l)){ return; }
  l->dimx=1; l->dimy=g.pres_b[l->w];
  // if(!GPU_placement(l)){ continue; }
  if(!GPU_placement(l)){ return; }
  //２ 左２行に置く
  // memcpy(&l->nB,&l->B,sizeof(Board));       // nB=B;
  l->nB=l->B;
  for(l->n=l->w;l->n<(g.size-2)*(g.size-1)-l->w;++l->n){
    // memcpy(&l->B,&l->nB,sizeof(Board));     // B=nB;
    l->B=l->nB;
    l->dimx=g.pres_a[l->n]; l->dimy=g.size-1;
    if(!GPU_placement(l)){ continue; }
    l->dimx=g.pres_b[l->n]; l->dimy=g.size-2;
    if(!GPU_placement(l)){ continue; }
    // ３ 下２行に置く
    // memcpy(&l->eB,&l->B,sizeof(Board));     // eB=B;
    l->eB=l->B;
    for(l->e=l->w;l->e<(g.size-2)*(g.size-1)-l->w;++l->e){
      // memcpy(&l->B,&l->eB,sizeof(Board));   // B=eB;
      l->B=l->eB;
      l->dimx=g.size-1; l->dimy=g.size-1-g.pres_a[l->e];
      if(!GPU_placement(l)){ continue; }
      l->dimx=g.size-2; l->dimy=g.size-1-g.pres_b[l->e];
      if(!GPU_placement(l)){ continue; }
      // ４ 右２列に置く
      // memcpy(&l->sB,&l->B,sizeof(Board));   // sB=B;
      l->sB=l->B;
      for(l->s=l->w;l->s<(g.size-2)*(g.size-1)-l->w;++l->s){
        // memcpy(&l->B,&l->sB,sizeof(Board)); // B=sB;
        l->B=l->sB;
        l->dimx=g.size-1-g.pres_a[l->s]; l->dimy=0;
        if(!GPU_placement(l)){ continue; }
        l->dimx=g.size-1-g.pres_b[l->s]; l->dimy=1;
        if(!GPU_placement(l)){ continue; }
        // 対称解除法
        carryChain_symmetryR(l);
      } //w
    } //e
  } //n
}
//GPU 再帰  チェーンのビルド
void GPU_buildChainR(const unsigned int size,unsigned int STEPS)
{
  Local l[(g.size/2)*(g.size-3)];
  l->STEPS=STEPS;
  l->size=size;
  Local lDevice[(g.size/2)*(g.size-3)];
  cudaMallocHost((void**) &l,   sizeof(struct Local)*l->STEPS);
  cudaMalloc((void**) &lDevice, sizeof(struct Local)*l->STEPS);

  // カウンターの初期化
  l->COUNT2=0; l->COUNT4=1; l->COUNT8=2;
  l->COUNTER[l->COUNT2]=l->COUNTER[l->COUNT4]=l->COUNTER[l->COUNT8]=0;
  // Board の初期化 nB,eB,sB,wB;
  l->B.row=l->B.down=l->B.left=l->B.right=0;
  // Board x[]の初期化
  for(unsigned int i=0;i<g.size;++i){ l->B.x[i]=-1; }
  //１ 上２行に置く
  // memcpy(&l->wB,&l->B,sizeof(Board));         // wB=B;
  l->wB=l->B;
  unsigned int limit=(unsigned)(g.size/2)*(g.size-3);
  cudaMemcpy(lDevice,l,
      sizeof(struct Local)*limit,cudaMemcpyHostToDevice);
  for(l->w=0;l->w<=(unsigned)(g.size/2)*(g.size-3);++l->w){
    thread_runR(&l);
    //GPU_thread_runR<<<l->STEPS/THREAD_NUM,THREAD_NUM>>>(&l);
    //GPU_thread_runR<<<l->STEPS/THREAD_NUM,THREAD_NUM>>>(&l);
  } //w
  cudaMemcpy(l,lDevice,
      sizeof(struct Local)*limit,cudaMemcpyDeviceToHost);
  /**
   * 集計
   */
  UNIQUE= l->COUNTER[l->COUNT2]+
          l->COUNTER[l->COUNT4]+
          l->COUNTER[l->COUNT8];
  TOTAL=  l->COUNTER[l->COUNT2]*2+
          l->COUNTER[l->COUNT4]*4+
          l->COUNTER[l->COUNT8]*8;
}
//GPU 再帰  キャリーチェーン
void GPU_carryChainR(const unsigned int size,unsigned int STEPS)
{
  listChain();  //チェーンのリストを作成
  GPU_buildChainR(size,STEPS); // チェーンのビルド
  // calcChain(&l);  // 集計
}
// CUDA 初期化
bool InitCUDA()
{
  int count;
  cudaGetDeviceCount(&count);
  if(count==0){fprintf(stderr,"There is no device.\n");return false;}
  int i;
  for(i=0;i<count;i++){
    struct cudaDeviceProp prop;
    if(cudaGetDeviceProperties(&prop,i)==cudaSuccess){if(prop.major>=1){break;} }
  }
  if(i==count){fprintf(stderr,"There is no device supporting CUDA 1.x.\n");return false;}
  cudaSetDevice(i);
  return true;
}
//メイン
int main(int argc,char** argv)
{
  bool cpu=false,cpur=false,gpu=false,gpuNodeLayer=false;
  int argstart=2;
  if(argc>=2&&argv[1][0]=='-'){
    if(argv[1][1]=='c'||argv[1][1]=='C'){cpu=true;}
    else if(argv[1][1]=='r'||argv[1][1]=='R'){cpur=true;}
    else if(argv[1][1]=='c'||argv[1][1]=='C'){cpu=true;}
    else if(argv[1][1]=='g'||argv[1][1]=='G'){gpu=true;}
    else if(argv[1][1]=='n'||argv[1][1]=='N'){gpuNodeLayer=true;}
    else{ gpuNodeLayer=true; } //デフォルトをgpuとする
    argstart=2;
  }
  if(argc<argstart){
    printf("Usage: %s [-c|-g|-r|-s] n steps\n",argv[0]);
    printf("  -r: CPU 再帰\n");
    printf("  -c: CPU 非再帰\n");
    printf("  -g: GPU 再帰\n");
    printf("  -n: GPU キャリーチェーン\n");
  }
  if(cpur){ printf("\n\nCPU キャリーチェーン 再帰 \n"); }
  else if(cpu){ printf("\n\nCPU キャリーチェーン 非再帰 \n"); }
  else if(gpu){ printf("\n\nGPU キャリーチェーン シングルスレッド\n"); }
  else if(gpuNodeLayer){ printf("\n\nGPU キャリーチェーン マルチスレッド\n"); }
  if(cpu||cpur)
  {
    int min=4; 
    int targetN=17;
    struct timeval t0;
    struct timeval t1;
    printf("%s\n"," N:        Total      Unique      dd:hh:mm:ss.ms");
    for(int size=min;size<=targetN;size++){
      TOTAL=UNIQUE=0;
      gettimeofday(&t0, NULL);//計測開始
      if(cpur){ //再帰
        g.size=size;
        carryChainR();
      }
      if(cpu){ //非再帰
        g.size=size;
        carryChain();
      }
      //
      gettimeofday(&t1, NULL);//計測終了
      int ss;int ms;int dd;
      if(t1.tv_usec<t0.tv_usec) {
        dd=(t1.tv_sec-t0.tv_sec-1)/86400;
        ss=(t1.tv_sec-t0.tv_sec-1)%86400;
        ms=(1000000+t1.tv_usec-t0.tv_usec+500)/10000;
      }else {
        dd=(t1.tv_sec-t0.tv_sec)/86400;
        ss=(t1.tv_sec-t0.tv_sec)%86400;
        ms=(t1.tv_usec-t0.tv_usec+500)/10000;
      }//end if
      int hh=ss/3600;
      int mm=(ss-hh*3600)/60;
      ss%=60;
      printf("%2d:%13ld%12ld%8.2d:%02d:%02d:%02d.%02d\n",size,TOTAL,UNIQUE,dd,hh,mm,ss,ms);
    } //end for
  }//end if
  if(gpu||gpuNodeLayer)
  {
    if(!InitCUDA()){return 0;}
    int STEPS=24576;
    int min=4;
    int targetN=21;
    struct timeval t0;
    struct timeval t1;
    printf("%s\n"," N:        Total      Unique      dd:hh:mm:ss.ms");
    for(int size=min;size<=targetN;size++){
      gettimeofday(&t0,NULL);   // 計測開始
      if(gpu){
        TOTAL=UNIQUE=0;
        g.size=size;
        GPU_carryChainR(size,STEPS); //キャリーチェーン
      }else if(gpuNodeLayer){
        TOTAL=UNIQUE=0;
        g.size=size;
        GPU_carryChainR(size,STEPS); // キャリーチェーン
      }
      gettimeofday(&t1,NULL);   // 計測終了
      int ss;int ms;int dd;
      if (t1.tv_usec<t0.tv_usec) {
        dd=(int)(t1.tv_sec-t0.tv_sec-1)/86400;
        ss=(t1.tv_sec-t0.tv_sec-1)%86400;
        ms=(1000000+t1.tv_usec-t0.tv_usec+500)/10000;
      } else {
        dd=(int)(t1.tv_sec-t0.tv_sec)/86400;
        ss=(t1.tv_sec-t0.tv_sec)%86400;
        ms=(t1.tv_usec-t0.tv_usec+500)/10000;
      }//end if
      int hh=ss/3600;
      int mm=(ss-hh*3600)/60;
      ss%=60;
      printf("%2d:%13ld%12ld%8.2d:%02d:%02d:%02d.%02d\n",size,TOTAL,UNIQUE,dd,hh,mm,ss,ms);
    }//end for
  }//end if
  return 0;
}

	.file	"tmpxft_00381faf_00000000-6_05CUDA_CarryChain_0801-1_suzuki.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4020:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4020:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z9listChainv
	.type	_Z9listChainv, @function
_Z9listChainv:
.LFB3998:
	.cfi_startproc
	endbr64
	movl	g(%rip), %esi
	xorl	%eax, %eax
	xorl	%ecx, %ecx
	leaq	g(%rip), %rdi
.L3:
	cmpl	%eax, %esi
	je	.L14
	xorl	%edx, %edx
.L7:
	cmpl	%edx, %eax
	jb	.L4
	movl	%eax, %r8d
	subl	%edx, %r8d
	cmpl	$1, %r8d
	ja	.L6
	jmp	.L5
.L4:
	movl	%edx, %r8d
	subl	%eax, %r8d
	decl	%r8d
	je	.L5
.L6:
	movl	%ecx, %r8d
	incl	%ecx
	movl	%eax, 4(%rdi,%r8,4)
	movl	%edx, 3724(%rdi,%r8,4)
.L5:
	incl	%edx
	cmpl	%edx, %esi
	jne	.L7
	incl	%eax
	jmp	.L3
.L14:
	ret
	.cfi_endproc
.LFE3998:
	.size	_Z9listChainv, .-_Z9listChainv
	.globl	_Z9placementPv
	.type	_Z9placementPv, @function
_Z9placementPv:
.LFB3999:
	.cfi_startproc
	endbr64
	movq	1264(%rdi), %rsi
	movq	1272(%rdi), %rdx
	cmpq	%rdx, 40(%rdi,%rsi,8)
	je	.L39
	movq	%rdi, %rax
	movq	40(%rdi), %rdi
	leaq	4(%rsi), %r8
	testq	%rdi, %rdi
	jne	.L17
	movq	48(%rax), %rcx
	cmpq	$-1, %rcx
	je	.L18
	cmpq	%rsi, %rcx
	jb	.L18
	cmpq	$1, %rdx
	jne	.L18
	jmp	.L40
.L17:
	cmpq	$-1, %rdi
	je	.L18
	cmpq	%rdi, %rsi
	jb	.L20
	movl	g(%rip), %ecx
	subq	%rdi, %rcx
	cmpq	%rcx, %rsi
	jb	.L21
.L20:
	testq	%rdx, %rdx
	je	.L40
	movl	g(%rip), %ecx
	decl	%ecx
	cmpq	%rcx, %rdx
	je	.L40
.L21:
	movl	g(%rip), %ecx
	leal	-1(%rcx), %r9d
	cmpq	%r9, %rsi
	jne	.L18
	cmpq	%rdx, %rdi
	jnb	.L40
	subq	%rdi, %rcx
	cmpq	%rcx, %rdx
	jnb	.L40
.L18:
	movl	$1, %edi
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	%esi, %ecx
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rbp
	movq	8(%rax), %r10
	salq	%cl, %rbp
	movq	%rdx, 8(%rax,%r8,8)
	movl	g(%rip), %r8d
	testq	%rbp, %r10
	je	.L43
.L22:
	xorl	%eax, %eax
	jmp	.L15
.L43:
	movq	%rdi, %rbx
	movl	%edx, %ecx
	movq	16(%rax), %r9
	salq	%cl, %rbx
	testq	%rbx, %r9
	jne	.L22
	leal	-1(%r8,%rdx), %ecx
	movq	%rdi, %r11
	movq	24(%rax), %r8
	subl	%esi, %ecx
	salq	%cl, %r11
	testq	%r11, %r8
	jne	.L22
	leal	(%rdx,%rsi), %ecx
	movq	32(%rax), %rdx
	salq	%cl, %rdi
	testq	%rdi, %rdx
	jne	.L22
	orq	%rbp, %r10
	orq	%rbx, %r9
	orq	%r11, %r8
	orq	%rdi, %rdx
	movq	%r10, 8(%rax)
	movq	%r9, 16(%rax)
	movq	%r8, 24(%rax)
	movq	%rdx, 32(%rax)
	movb	$1, %al
.L15:
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
.L40:
	.cfi_restore 3
	.cfi_restore 6
	xorl	%eax, %eax
	ret
.L39:
	movb	$1, %al
	ret
	.cfi_endproc
.LFE3999:
	.size	_Z9placementPv, .-_Z9placementPv
	.globl	_Z5solveiimmmm
	.type	_Z5solveiimmmm, @function
_Z5solveiimmmm:
.LFB4000:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdx, %r10
	movq	%rcx, %rdx
	subq	$1088, %rsp
	.cfi_def_cfa_offset 1104
	movq	%fs:40, %rax
	movq	%rax, 1080(%rsp)
	xorl	%eax, %eax
.L45:
	cmpl	%eax, %edi
	jle	.L63
	xorl	%ecx, %ecx
	movq	%rcx, (%rsp,%rax,8)
	movq	%rcx, 432(%rsp,%rax,8)
	movq	%rcx, 648(%rsp,%rax,8)
	movq	%rcx, 216(%rsp,%rax,8)
	movq	%rcx, 864(%rsp,%rax,8)
	incq	%rax
	jmp	.L45
.L63:
	movq	%r8, %rax
	movslq	%esi, %rcx
	xorl	%r11d, %r11d
	orq	%r9, %rax
	movq	%r10, (%rsp,%rcx,8)
	orq	%rdx, %rax
	movq	%rdx, 432(%rsp,%rcx,8)
	notq	%rax
	movq	%r8, 648(%rsp,%rcx,8)
	movq	%r9, 216(%rsp,%rcx,8)
	movq	%rax, 864(%rsp,%rcx,8)
.L47:
	testl	%esi, %esi
	js	.L64
	movq	%r10, %rcx
	andl	$1, %ecx
	movq	%rcx, %rbx
	orq	%rax, %rbx
	je	.L48
	cmpl	%edi, %esi
	jge	.L48
	cmpq	$-1, %r8
	jne	.L49
	incq	%r11
	jmp	.L48
.L49:
	testq	%rcx, %rcx
	je	.L60
.L51:
	testl	$1, %r10d
	je	.L62
	shrq	%r10
	addq	%rdx, %rdx
	shrq	%r9
	jmp	.L51
.L60:
	movq	%rax, %rcx
	movslq	%esi, %rbx
	incl	%esi
	negq	%rcx
	movq	%rdx, 432(%rsp,%rbx,8)
	andq	%rax, %rcx
	movq	%r8, 648(%rsp,%rbx,8)
	xorq	%rcx, %rax
	orq	%rcx, %rdx
	orq	%rcx, %r8
	orq	%r9, %rcx
	movq	%rax, 864(%rsp,%rbx,8)
	shrq	%rcx
	addq	%rdx, %rdx
	movq	%r10, (%rsp,%rbx,8)
	shrq	%r10
	movq	%r9, 216(%rsp,%rbx,8)
	movq	%rcx, %r9
.L62:
	movq	%rdx, %rax
	orq	%r9, %rax
	orq	%r8, %rax
	notq	%rax
	jmp	.L47
.L48:
	decl	%esi
	movslq	%esi, %rax
	movq	(%rsp,%rax,8), %r10
	movq	432(%rsp,%rax,8), %rdx
	movq	216(%rsp,%rax,8), %r9
	movq	648(%rsp,%rax,8), %r8
	movq	864(%rsp,%rax,8), %rax
	jmp	.L47
.L64:
	movq	1080(%rsp), %rax
	subq	%fs:40, %rax
	je	.L55
	call	__stack_chk_fail@PLT
.L55:
	addq	$1088, %rsp
	.cfi_def_cfa_offset 16
	movq	%r11, %rax
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4000:
	.size	_Z5solveiimmmm, .-_Z5solveiimmmm
	.globl	_Z19carryChain_symmetryPv
	.type	_Z19carryChain_symmetryPv, @function
_Z19carryChain_symmetryPv:
.LFB4001:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rdi, %rbx
	movl	g(%rip), %edi
	movl	1260(%rbx), %r8d
	movl	1256(%rbx), %ecx
	leal	-1(%rdi), %edx
	leal	-2(%rdi), %eax
	movl	1252(%rbx), %esi
	imull	%edx, %eax
	movl	1248(%rbx), %edx
	decl	%eax
	movl	%eax, %r9d
	subl	%r8d, %r9d
	cmpl	%r9d, %ecx
	jne	.L66
	movl	%eax, %r10d
	subl	%esi, %r10d
	cmpl	%r10d, %edx
	jb	.L65
.L66:
	cmpl	%esi, %r9d
	jne	.L68
	movl	%eax, %r10d
	subl	%edx, %r10d
	cmpl	%edx, %r10d
	jb	.L65
.L68:
	cmpl	%r9d, %edx
	jne	.L69
	subl	%ecx, %eax
	cmpl	%esi, %eax
	jb	.L65
.L69:
	cmpq	$0, 40(%rbx)
	je	.L72
	cmpl	%ecx, %r8d
	jne	.L71
	cmpl	%esi, %r8d
	jne	.L65
	cmpl	%edx, %r8d
	jne	.L65
	leal	-4(%rdi), %ecx
	movq	16(%rbx), %rax
	orl	$-1, %edx
	leal	-5(%rdi), %r8d
	sall	%cl, %edx
	movq	32(%rbx), %r9
	movq	24(%rbx), %rsi
	movl	%r8d, %ecx
	movslq	%edx, %rdx
	shrq	$2, %rax
	shrq	$4, %rsi
	orq	%rdx, %rax
	movq	8(%rbx), %rdx
	shrq	$4, %r9
	incq	%rax
	salq	%cl, %r9
	salq	%cl, %rax
	shrq	$2, %rdx
	movq	%rsi, %rcx
	xorl	%esi, %esi
	leaq	-1(%rax), %r8
	call	_Z5solveiimmmm
	movq	%rax, %rdx
	movl	1304(%rbx), %eax
	jmp	.L80
.L71:
	cmpl	%ecx, %edx
	jb	.L72
	cmpl	%esi, %r8d
	jne	.L72
	cmpl	%edx, %ecx
	jb	.L65
	leal	-4(%rdi), %ecx
	movq	16(%rbx), %rax
	orl	$-1, %edx
	leal	-5(%rdi), %r8d
	sall	%cl, %edx
	movq	32(%rbx), %r9
	movq	24(%rbx), %rsi
	movl	%r8d, %ecx
	movslq	%edx, %rdx
	shrq	$2, %rax
	shrq	$4, %rsi
	orq	%rdx, %rax
	movq	8(%rbx), %rdx
	shrq	$4, %r9
	incq	%rax
	salq	%cl, %r9
	salq	%cl, %rax
	shrq	$2, %rdx
	movq	%rsi, %rcx
	xorl	%esi, %esi
	leaq	-1(%rax), %r8
	call	_Z5solveiimmmm
	movq	%rax, %rdx
	movl	1308(%rbx), %eax
	jmp	.L80
.L72:
	leal	-4(%rdi), %ecx
	movq	16(%rbx), %rax
	orl	$-1, %edx
	leal	-5(%rdi), %r8d
	sall	%cl, %edx
	movq	32(%rbx), %r9
	movq	24(%rbx), %rsi
	movl	%r8d, %ecx
	movslq	%edx, %rdx
	shrq	$2, %rax
	shrq	$4, %rsi
	orq	%rdx, %rax
	movq	8(%rbx), %rdx
	shrq	$4, %r9
	incq	%rax
	salq	%cl, %r9
	salq	%cl, %rax
	shrq	$2, %rdx
	movq	%rsi, %rcx
	xorl	%esi, %esi
	leaq	-1(%rax), %r8
	call	_Z5solveiimmmm
	movq	%rax, %rdx
	movl	1312(%rbx), %eax
.L80:
	addq	%rdx, 1280(%rbx,%rax,8)
.L65:
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4001:
	.size	_Z19carryChain_symmetryPv, .-_Z19carryChain_symmetryPv
	.globl	_Z10thread_runPv
	.type	_Z10thread_runPv, @function
_Z10thread_runPv:
.LFB4002:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$62, %ecx
	xorl	%edx, %edx
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdi, %r12
	addq	$8, %rdi
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	leaq	1000(%r12), %rsi
	leaq	g(%rip), %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$72, %rsp
	.cfi_def_cfa_offset 128
	movl	1260(%r12), %ebx
	rep movsl
	movq	%rdx, 1264(%r12)
	movq	%r12, %rdi
	movl	4(%rbp,%rbx,4), %eax
	movq	%rax, 1272(%r12)
	call	_Z9placementPv
	testb	%al, %al
	leaq	8(%r12), %rax
	movq	%rax, 8(%rsp)
	je	.L81
	movl	3724(%rbp,%rbx,4), %eax
	movq	%r12, %rdi
	movq	$1, 1264(%r12)
	movq	%rax, 1272(%r12)
	call	_Z9placementPv
	testb	%al, %al
	je	.L81
	leaq	256(%r12), %rdi
	movl	$62, %ecx
	movl	%ebx, 16(%rsp)
	leaq	8(%r12), %rsi
	rep movsl
	leaq	256(%r12), %rax
	movq	%rbx, %r13
	movq	%rax, 40(%rsp)
	leaq	504(%r12), %rax
	movq	%rax, 48(%rsp)
.L84:
	movl	0(%rbp), %ebx
	movl	16(%rsp), %eax
	leal	-2(%rbx), %r14d
	movl	%eax, 1248(%r12)
	decl	%ebx
	movl	%r14d, %eax
	imull	%ebx, %eax
	subl	%r13d, %eax
	movl	%eax, 20(%rsp)
	cmpl	%eax, 16(%rsp)
	jnb	.L81
	movl	16(%rsp), %r15d
	movq	8(%rsp), %rdi
	movl	$62, %ecx
	movq	40(%rsp), %rsi
	rep movsl
	movl	4(%rbp,%r15,4), %eax
	movq	%r12, %rdi
	movq	%rax, 1264(%r12)
	movl	%ebx, %eax
	movq	%rax, 1272(%r12)
	movq	%rax, 32(%rsp)
	call	_Z9placementPv
	testb	%al, %al
	jne	.L85
.L86:
	incl	16(%rsp)
	jmp	.L84
.L85:
	movl	3724(%rbp,%r15,4), %eax
	movq	%r12, %rdi
	movq	%rax, 1264(%r12)
	movl	%r14d, %eax
	movq	%rax, 1272(%r12)
	movq	%rax, 24(%rsp)
	call	_Z9placementPv
	testb	%al, %al
	je	.L86
	movq	8(%rsp), %rsi
	leaq	504(%r12), %rdi
	movl	$62, %ecx
	movl	%r13d, %r14d
	rep movsl
.L87:
	movl	20(%rsp), %eax
	movl	%r14d, 1252(%r12)
	cmpl	%eax, %r14d
	jnb	.L86
	movq	32(%rsp), %rax
	movq	8(%rsp), %rdi
	movl	$62, %ecx
	movl	%r14d, %r15d
	movq	48(%rsp), %rsi
	rep movsl
	movq	%rax, 1264(%r12)
	movl	%ebx, %eax
	movq	%r12, %rdi
	subl	4(%rbp,%r15,4), %eax
	movq	%rax, 1272(%r12)
	call	_Z9placementPv
	testb	%al, %al
	jne	.L88
.L89:
	incl	%r14d
	jmp	.L87
.L88:
	movq	24(%rsp), %rax
	movq	%r12, %rdi
	movq	%rax, 1264(%r12)
	movl	%ebx, %eax
	subl	3724(%rbp,%r15,4), %eax
	movq	%rax, 1272(%r12)
	call	_Z9placementPv
	testb	%al, %al
	je	.L89
	movq	8(%rsp), %rsi
	leaq	752(%r12), %rdi
	movl	$62, %ecx
	movl	%r13d, %r15d
	rep movsl
.L90:
	movl	20(%rsp), %eax
	movl	%r15d, 1256(%r12)
	cmpl	%eax, %r15d
	jnb	.L89
	movq	8(%rsp), %rdi
	leaq	752(%r12), %rsi
	movl	$62, %ecx
	movl	%ebx, %eax
	rep movsl
	movl	%r15d, %ecx
	movq	%r12, %rdi
	subl	4(%rbp,%rcx,4), %eax
	movq	%rcx, 56(%rsp)
	movq	%rax, 1264(%r12)
	xorl	%eax, %eax
	movq	%rax, 1272(%r12)
	call	_Z9placementPv
	testb	%al, %al
	je	.L92
	movq	56(%rsp), %rcx
	movl	%ebx, %eax
	movq	%r12, %rdi
	subl	3724(%rbp,%rcx,4), %eax
	movq	$1, 1272(%r12)
	movq	%rax, 1264(%r12)
	call	_Z9placementPv
	testb	%al, %al
	je	.L92
	movq	%r12, %rdi
	call	_Z19carryChain_symmetryPv
.L92:
	incl	%r15d
	jmp	.L90
.L81:
	addq	$72, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4002:
	.size	_Z10thread_runPv, .-_Z10thread_runPv
	.globl	_Z10buildChainv
	.type	_Z10buildChainv, @function
_Z10buildChainv:
.LFB4003:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	subq	$16, %rsp
	.cfi_offset 14, -24
	.cfi_offset 13, -32
	.cfi_offset 12, -40
	.cfi_offset 3, -48
	movl	g(%rip), %edx
	movq	%fs:40, %rax
	movq	%rax, -40(%rbp)
	xorl	%eax, %eax
	movq	%rsp, %rcx
	movl	%edx, %r14d
	leal	-3(%rdx), %eax
	shrl	%r14d
	imull	%eax, %r14d
	movl	%r14d, %eax
	imulq	$1320, %rax, %rax
	addq	$15, %rax
	shrq	$4, %rax
	salq	$4, %rax
	movq	%rax, %rsi
	andq	$-4096, %rsi
	subq	%rsi, %rcx
.L111:
	cmpq	%rcx, %rsp
	je	.L112
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L111
.L112:
	andl	$4095, %eax
	subq	%rax, %rsp
	testq	%rax, %rax
	je	.L113
	orq	$0, -8(%rsp,%rax)
.L113:
	leaq	7(%rsp), %r12
	xorl	%eax, %eax
	xorl	%ecx, %ecx
	movq	%r12, %rbx
	andq	$-8, %r12
	shrq	$3, %rbx
	movq	%rax, 1280(,%rbx,8)
	movq	%rax, 1288(,%rbx,8)
	movq	%rax, 1296(,%rbx,8)
	movl	$1, %eax
	movl	$2, 1312(,%rbx,8)
	salq	$32, %rax
	movq	%rcx, 32(,%rbx,8)
	movq	%rcx, 24(,%rbx,8)
	movq	%rcx, 16(,%rbx,8)
	movq	%rcx, 8(,%rbx,8)
	movq	%rax, 1304(,%rbx,8)
	xorl	%eax, %eax
.L114:
	cmpl	%edx, %eax
	jnb	.L123
	movq	$-1, 40(%r12,%rax,8)
	incq	%rax
	jmp	.L114
.L123:
	leaq	1000(%r12), %rdi
	leaq	8(%r12), %rsi
	movl	$62, %ecx
	xorl	%r13d, %r13d
	rep movsl
.L116:
	movl	%r13d, 1260(%r12)
	cmpl	%r13d, %r14d
	jb	.L124
	movq	%r12, %rdi
	incl	%r13d
	call	_Z10thread_runPv
	jmp	.L116
.L124:
	movq	1280(,%rbx,8), %rcx
	movq	1288(,%rbx,8), %rsi
	movq	1296(,%rbx,8), %rax
	leaq	(%rcx,%rsi), %rdx
	addq	%rax, %rdx
	addq	%rax, %rax
	addq	%rsi, %rax
	movq	%rdx, UNIQUE(%rip)
	addq	%rax, %rax
	addq	%rcx, %rax
	addq	%rax, %rax
	movq	%rax, TOTAL(%rip)
	movq	-40(%rbp), %rax
	subq	%fs:40, %rax
	je	.L118
	call	__stack_chk_fail@PLT
.L118:
	leaq	-32(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE4003:
	.size	_Z10buildChainv, .-_Z10buildChainv
	.globl	_Z10carryChainv
	.type	_Z10carryChainv, @function
_Z10carryChainv:
.LFB4004:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z9listChainv
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	_Z10buildChainv
	.cfi_endproc
.LFE4004:
	.size	_Z10carryChainv, .-_Z10carryChainv
	.globl	_Z6solveRmmmm
	.type	_Z6solveRmmmm, @function
_Z6solveRmmmm:
.LFB4005:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$1, %r15d
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	cmpq	$-1, %rdx
	je	.L127
	movq	%rdi, %r12
	movq	%rsi, %rbp
	movq	%rdx, %r8
	movq	%rcx, %r13
.L129:
	movq	%r12, %r15
	shrq	%r12
	andl	$1, %r15d
	testq	%r15, %r15
	je	.L136
	addq	%rbp, %rbp
	shrq	%r13
	jmp	.L129
.L136:
	movq	%rbp, %rbx
	orq	%r13, %rbx
	orq	%r8, %rbx
	notq	%rbx
.L131:
	testq	%rbx, %rbx
	je	.L127
	movq	%rbx, %r14
	movq	%r13, %rcx
	movq	%rbp, %rsi
	movq	%r8, %rdx
	negq	%r14
	movq	%r12, %rdi
	movq	%r8, 8(%rsp)
	andq	%rbx, %r14
	orq	%r14, %rcx
	orq	%r14, %rsi
	orq	%r14, %rdx
	xorq	%r14, %rbx
	shrq	%rcx
	addq	%rsi, %rsi
	call	_Z6solveRmmmm
	movq	8(%rsp), %r8
	addq	%rax, %r15
	jmp	.L131
.L127:
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	movq	%r15, %rax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4005:
	.size	_Z6solveRmmmm, .-_Z6solveRmmmm
	.globl	_Z20carryChain_symmetryRPv
	.type	_Z20carryChain_symmetryRPv, @function
_Z20carryChain_symmetryRPv:
.LFB4006:
	.cfi_startproc
	endbr64
	movl	g(%rip), %eax
	movl	1260(%rdi), %r8d
	movq	%rdi, %r9
	movl	1256(%rdi), %esi
	leal	-1(%rax), %ecx
	leal	-2(%rax), %edx
	imull	%ecx, %edx
	movl	1248(%rdi), %ecx
	movl	1252(%rdi), %edi
	decl	%edx
	movl	%edx, %r10d
	subl	%r8d, %r10d
	cmpl	%r10d, %esi
	jne	.L138
	movl	%edx, %r11d
	subl	%edi, %r11d
	cmpl	%r11d, %ecx
	jb	.L152
.L138:
	cmpl	%edi, %r10d
	jne	.L140
	movl	%edx, %r11d
	subl	%ecx, %r11d
	cmpl	%ecx, %r11d
	jb	.L152
.L140:
	cmpl	%r10d, %ecx
	jne	.L141
	subl	%esi, %edx
	cmpl	%edi, %edx
	jb	.L152
.L141:
	pushq	%rdx
	.cfi_def_cfa_offset 16
	cmpq	$0, 40(%r9)
	jne	.L142
	movq	32(%r9), %rsi
	leal	-5(%rax), %r8d
	movq	16(%r9), %rdx
	orl	$-1, %edi
	movl	%r8d, %ecx
	shrq	$4, %rsi
	shrq	$2, %rdx
	salq	%cl, %rsi
	leal	-4(%rax), %ecx
	sall	%cl, %edi
	movl	%r8d, %ecx
	movslq	%edi, %rdi
	orq	%rdi, %rdx
	incq	%rdx
	jmp	.L155
.L142:
	cmpl	%esi, %r8d
	jne	.L143
	cmpl	%edi, %r8d
	jne	.L137
	cmpl	%ecx, %r8d
	jne	.L137
	movq	32(%r9), %rsi
	leal	-5(%rax), %r8d
	movq	16(%r9), %rdx
	orl	$-1, %edi
	movl	%r8d, %ecx
	shrq	$4, %rsi
	shrq	$2, %rdx
	salq	%cl, %rsi
	leal	-4(%rax), %ecx
	movq	24(%r9), %rax
	sall	%cl, %edi
	movl	%r8d, %ecx
	movslq	%edi, %rdi
	shrq	$4, %rax
	orq	%rdi, %rdx
	movq	8(%r9), %rdi
	incq	%rdx
	salq	%cl, %rdx
	shrq	$2, %rdi
	movq	%rsi, %rcx
	movq	%rax, %rsi
	decq	%rdx
	call	_Z6solveRmmmm
	movq	%rax, %rdx
	movl	1304(%r9), %eax
	jmp	.L154
.L143:
	cmpl	%esi, %ecx
	jb	.L144
	cmpl	%edi, %r8d
	jne	.L144
	cmpl	%ecx, %esi
	jb	.L137
	movq	32(%r9), %rsi
	leal	-5(%rax), %r8d
	movq	16(%r9), %rdx
	orl	$-1, %edi
	movl	%r8d, %ecx
	shrq	$4, %rsi
	shrq	$2, %rdx
	salq	%cl, %rsi
	leal	-4(%rax), %ecx
	movq	24(%r9), %rax
	sall	%cl, %edi
	movl	%r8d, %ecx
	movslq	%edi, %rdi
	shrq	$4, %rax
	orq	%rdi, %rdx
	movq	8(%r9), %rdi
	incq	%rdx
	salq	%cl, %rdx
	shrq	$2, %rdi
	movq	%rsi, %rcx
	movq	%rax, %rsi
	decq	%rdx
	call	_Z6solveRmmmm
	movq	%rax, %rdx
	movl	1308(%r9), %eax
	jmp	.L154
.L144:
	movq	32(%r9), %rsi
	leal	-5(%rax), %edi
	movq	16(%r9), %rdx
	movl	%edi, %ecx
	shrq	$4, %rsi
	shrq	$2, %rdx
	salq	%cl, %rsi
	leal	-4(%rax), %ecx
	orl	$-1, %eax
	sall	%cl, %eax
	movl	%edi, %ecx
	cltq
	orq	%rax, %rdx
	incq	%rdx
.L155:
	movq	24(%r9), %rax
	movq	8(%r9), %rdi
	salq	%cl, %rdx
	movq	%rsi, %rcx
	decq	%rdx
	shrq	$4, %rax
	shrq	$2, %rdi
	movq	%rax, %rsi
	call	_Z6solveRmmmm
	movq	%rax, %rdx
	movl	1312(%r9), %eax
.L154:
	addq	%rdx, 1280(%r9,%rax,8)
.L137:
	popq	%rax
	.cfi_def_cfa_offset 8
	ret
.L152:
	ret
	.cfi_endproc
.LFE4006:
	.size	_Z20carryChain_symmetryRPv, .-_Z20carryChain_symmetryRPv
	.globl	_Z11thread_runRPv
	.type	_Z11thread_runRPv, @function
_Z11thread_runRPv:
.LFB4007:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	movl	$62, %ecx
	xorl	%edx, %edx
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdi, %r12
	addq	$8, %rdi
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	leaq	1000(%r12), %rsi
	leaq	g(%rip), %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$56, %rsp
	.cfi_def_cfa_offset 112
	movl	1260(%r12), %ebx
	rep movsl
	movq	%rdx, 1264(%r12)
	movq	%r12, %rdi
	movl	4(%rbp,%rbx,4), %eax
	movq	%rax, 1272(%r12)
	call	_Z9placementPv
	testb	%al, %al
	leaq	8(%r12), %rax
	movq	%rax, 8(%rsp)
	je	.L156
	movl	3724(%rbp,%rbx,4), %eax
	movq	%r12, %rdi
	movq	$1, 1264(%r12)
	movq	%rax, 1272(%r12)
	call	_Z9placementPv
	testb	%al, %al
	je	.L156
	leaq	256(%r12), %rdi
	movl	$62, %ecx
	movl	%ebx, 16(%rsp)
	leaq	8(%r12), %rsi
	rep movsl
	leaq	256(%r12), %rax
	movq	%rbx, %r13
	movq	%rax, 24(%rsp)
	leaq	504(%r12), %rax
	movq	%rax, 32(%rsp)
.L159:
	movl	0(%rbp), %ebx
	movl	16(%rsp), %eax
	leal	-2(%rbx), %r14d
	movl	%eax, 1248(%r12)
	decl	%ebx
	movl	%r14d, %eax
	imull	%ebx, %eax
	subl	%r13d, %eax
	movl	%eax, 20(%rsp)
	cmpl	%eax, 16(%rsp)
	jnb	.L156
	movl	16(%rsp), %r15d
	movq	8(%rsp), %rdi
	movl	$62, %ecx
	movd	%ebx, %xmm0
	movq	24(%rsp), %rsi
	rep movsl
	movl	4(%rbp,%r15,4), %eax
	movq	%r12, %rdi
	movq	%xmm0, 1272(%r12)
	movq	%rax, 1264(%r12)
	call	_Z9placementPv
	testb	%al, %al
	jne	.L160
.L161:
	incl	16(%rsp)
	jmp	.L159
.L160:
	movl	3724(%rbp,%r15,4), %eax
	movd	%r14d, %xmm1
	movq	%r12, %rdi
	movq	%xmm1, 1272(%r12)
	movq	%rax, 1264(%r12)
	call	_Z9placementPv
	testb	%al, %al
	je	.L161
	movq	8(%rsp), %rsi
	leaq	504(%r12), %rdi
	movl	$62, %ecx
	movl	%r13d, %r14d
	rep movsl
.L162:
	movl	20(%rsp), %eax
	movl	%r14d, 1252(%r12)
	cmpl	%eax, %r14d
	jnb	.L161
	movq	8(%rsp), %rdi
	movq	32(%rsp), %rsi
	movl	%r14d, %r15d
	movl	%ebx, %eax
	movl	$62, %ecx
	subl	4(%rbp,%r15,4), %eax
	rep movsl
	movq	%rax, 1272(%r12)
	movq	%r12, %rdi
	movq	%xmm0, 1264(%r12)
	call	_Z9placementPv
	testb	%al, %al
	jne	.L163
.L164:
	incl	%r14d
	jmp	.L162
.L163:
	movq	%xmm1, 1264(%r12)
	movl	%ebx, %eax
	movq	%r12, %rdi
	subl	3724(%rbp,%r15,4), %eax
	movq	%rax, 1272(%r12)
	call	_Z9placementPv
	testb	%al, %al
	je	.L164
	movq	8(%rsp), %rsi
	leaq	752(%r12), %rdi
	movl	$62, %ecx
	movl	%r13d, %r15d
	rep movsl
.L165:
	movl	20(%rsp), %eax
	movl	%r15d, 1256(%r12)
	cmpl	%eax, %r15d
	jnb	.L164
	movq	8(%rsp), %rdi
	leaq	752(%r12), %rsi
	movl	$62, %ecx
	movl	%ebx, %eax
	rep movsl
	movl	%r15d, %ecx
	movq	%r12, %rdi
	subl	4(%rbp,%rcx,4), %eax
	movq	%rcx, 40(%rsp)
	movq	%rax, 1264(%r12)
	xorl	%eax, %eax
	movq	%rax, 1272(%r12)
	call	_Z9placementPv
	testb	%al, %al
	je	.L167
	movq	40(%rsp), %rcx
	movl	%ebx, %eax
	movq	%r12, %rdi
	subl	3724(%rbp,%rcx,4), %eax
	movq	$1, 1272(%r12)
	movq	%rax, 1264(%r12)
	call	_Z9placementPv
	testb	%al, %al
	je	.L167
	movq	%r12, %rdi
	call	_Z20carryChain_symmetryRPv
.L167:
	incl	%r15d
	jmp	.L165
.L156:
	addq	$56, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4007:
	.size	_Z11thread_runRPv, .-_Z11thread_runRPv
	.globl	_Z11buildChainRv
	.type	_Z11buildChainRv, @function
_Z11buildChainRv:
.LFB4008:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	subq	$16, %rsp
	.cfi_offset 14, -24
	.cfi_offset 13, -32
	.cfi_offset 12, -40
	.cfi_offset 3, -48
	movl	g(%rip), %edx
	movq	%fs:40, %rax
	movq	%rax, -40(%rbp)
	xorl	%eax, %eax
	movq	%rsp, %rcx
	movl	%edx, %r14d
	leal	-3(%rdx), %eax
	shrl	%r14d
	imull	%eax, %r14d
	movl	%r14d, %eax
	imulq	$1320, %rax, %rax
	addq	$15, %rax
	shrq	$4, %rax
	salq	$4, %rax
	movq	%rax, %rsi
	andq	$-4096, %rsi
	subq	%rsi, %rcx
.L186:
	cmpq	%rcx, %rsp
	je	.L187
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L186
.L187:
	andl	$4095, %eax
	subq	%rax, %rsp
	testq	%rax, %rax
	je	.L188
	orq	$0, -8(%rsp,%rax)
.L188:
	leaq	7(%rsp), %r12
	xorl	%eax, %eax
	xorl	%ecx, %ecx
	movq	%r12, %rbx
	andq	$-8, %r12
	shrq	$3, %rbx
	movq	%rax, 1280(,%rbx,8)
	movq	%rax, 1288(,%rbx,8)
	movq	%rax, 1296(,%rbx,8)
	movl	$1, %eax
	movl	$2, 1312(,%rbx,8)
	salq	$32, %rax
	movq	%rcx, 32(,%rbx,8)
	movq	%rcx, 24(,%rbx,8)
	movq	%rcx, 16(,%rbx,8)
	movq	%rcx, 8(,%rbx,8)
	movq	%rax, 1304(,%rbx,8)
	xorl	%eax, %eax
.L189:
	cmpl	%edx, %eax
	jnb	.L198
	movq	$-1, 40(%r12,%rax,8)
	incq	%rax
	jmp	.L189
.L198:
	leaq	1000(%r12), %rdi
	leaq	8(%r12), %rsi
	movl	$62, %ecx
	xorl	%r13d, %r13d
	rep movsl
.L191:
	movl	%r13d, 1260(%r12)
	cmpl	%r13d, %r14d
	jb	.L199
	movq	%r12, %rdi
	incl	%r13d
	call	_Z11thread_runRPv
	jmp	.L191
.L199:
	movq	1280(,%rbx,8), %rcx
	movq	1288(,%rbx,8), %rsi
	movq	1296(,%rbx,8), %rax
	leaq	(%rcx,%rsi), %rdx
	addq	%rax, %rdx
	addq	%rax, %rax
	addq	%rsi, %rax
	movq	%rdx, UNIQUE(%rip)
	addq	%rax, %rax
	addq	%rcx, %rax
	addq	%rax, %rax
	movq	%rax, TOTAL(%rip)
	movq	-40(%rbp), %rax
	subq	%fs:40, %rax
	je	.L193
	call	__stack_chk_fail@PLT
.L193:
	leaq	-32(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE4008:
	.size	_Z11buildChainRv, .-_Z11buildChainRv
	.globl	_Z11carryChainRv
	.type	_Z11carryChainRv, @function
_Z11carryChainRv:
.LFB4009:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z9listChainv
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	_Z11buildChainRv
	.cfi_endproc
.LFE4009:
	.size	_Z11carryChainRv, .-_Z11carryChainRv
	.globl	_Z13GPU_placementPv
	.type	_Z13GPU_placementPv, @function
_Z13GPU_placementPv:
.LFB4663:
	.cfi_startproc
	endbr64
	jmp	_Z9placementPv
	.cfi_endproc
.LFE4663:
	.size	_Z13GPU_placementPv, .-_Z13GPU_placementPv
	.globl	_Z10GPU_solveRmmmm
	.type	_Z10GPU_solveRmmmm, @function
_Z10GPU_solveRmmmm:
.LFB4011:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movl	$1, %r12d
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	pushq	%r8
	.cfi_def_cfa_offset 48
	cmpq	$-1, %rdx
	je	.L203
	movq	%rdi, %r11
	movq	%rsi, %r10
	movq	%rdx, %r13
	movq	%rcx, %rbx
.L205:
	movq	%r11, %r12
	shrq	%r11
	andl	$1, %r12d
	testq	%r12, %r12
	je	.L212
	addq	%r10, %r10
	shrq	%rbx
	jmp	.L205
.L212:
	movq	%r10, %r9
	orq	%rbx, %r9
	orq	%r13, %r9
	notq	%r9
.L207:
	testq	%r9, %r9
	je	.L203
	movq	%r9, %rbp
	movq	%rbx, %rcx
	movq	%r10, %rsi
	movq	%r13, %rdx
	negq	%rbp
	movq	%r11, %rdi
	andq	%r9, %rbp
	orq	%rbp, %rcx
	orq	%rbp, %rsi
	orq	%rbp, %rdx
	xorq	%rbp, %r9
	shrq	%rcx
	addq	%rsi, %rsi
	call	_Z6solveRmmmm
	addq	%rax, %r12
	jmp	.L207
.L203:
	popq	%rdx
	.cfi_def_cfa_offset 40
	movq	%r12, %rax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4011:
	.size	_Z10GPU_solveRmmmm, .-_Z10GPU_solveRmmmm
	.globl	_Z24GPU_carryChain_symmetryRPv
	.type	_Z24GPU_carryChain_symmetryRPv, @function
_Z24GPU_carryChain_symmetryRPv:
.LFB4667:
	.cfi_startproc
	endbr64
	jmp	_Z20carryChain_symmetryRPv
	.cfi_endproc
.LFE4667:
	.size	_Z24GPU_carryChain_symmetryRPv, .-_Z24GPU_carryChain_symmetryRPv
	.globl	_Z15GPU_thread_runRPv
	.type	_Z15GPU_thread_runRPv, @function
_Z15GPU_thread_runRPv:
.LFB4665:
	.cfi_startproc
	endbr64
	jmp	_Z11thread_runRPv
	.cfi_endproc
.LFE4665:
	.size	_Z15GPU_thread_runRPv, .-_Z15GPU_thread_runRPv
	.globl	_Z15GPU_buildChainRjj
	.type	_Z15GPU_buildChainRjj, @function
_Z15GPU_buildChainRjj:
.LFB4014:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	%edi, %ecx
	movq	%rsp, %rbp
	.cfi_def_cfa_register 6
	pushq	%r15
	pushq	%r14
	pushq	%r13
	pushq	%r12
	pushq	%rbx
	subq	$24, %rsp
	.cfi_offset 15, -24
	.cfi_offset 14, -32
	.cfi_offset 13, -40
	.cfi_offset 12, -48
	.cfi_offset 3, -56
	movl	g(%rip), %edx
	movq	%fs:40, %rax
	movq	%rax, -56(%rbp)
	xorl	%eax, %eax
	movq	%rsp, %rdi
	movl	%edx, %eax
	subl	$3, %edx
	shrl	%eax
	imull	%edx, %eax
	imulq	$1320, %rax, %rax
	addq	$15, %rax
	movq	%rax, %rdx
	shrq	$4, %rdx
	salq	$4, %rdx
	movq	%rdx, %r8
	andq	$-4096, %r8
	subq	%r8, %rdi
.L216:
	cmpq	%rdi, %rsp
	je	.L217
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L216
.L217:
	andl	$4095, %edx
	subq	%rdx, %rsp
	testq	%rdx, %rdx
	je	.L218
	orq	$0, -8(%rsp,%rdx)
.L218:
	leaq	7(%rsp), %rbx
	shrq	$4, %rax
	movq	%rsp, %rdx
	movq	%rbx, %r12
	salq	$4, %rax
	andq	$-8, %rbx
	shrq	$3, %r12
	movl	%ecx, 0(,%r12,8)
	movq	%rax, %rcx
	movl	%esi, 1316(,%r12,8)
	andq	$-4096, %rcx
	subq	%rcx, %rdx
.L219:
	cmpq	%rdx, %rsp
	je	.L220
	subq	$4096, %rsp
	orq	$0, 4088(%rsp)
	jmp	.L219
.L220:
	andl	$4095, %eax
	subq	%rax, %rsp
	testq	%rax, %rax
	je	.L221
	orq	$0, -8(%rsp,%rax)
.L221:
	movl	%esi, %esi
	movq	%rbx, %rdi
	movq	%rsp, %r15
	imulq	$1320, %rsi, %rsi
	call	cudaMallocHost@PLT
	movl	1316(,%r12,8), %esi
	movq	%r15, %rdi
	imulq	$1320, %rsi, %rsi
	call	cudaMalloc@PLT
	xorl	%eax, %eax
	xorl	%edx, %edx
	movl	$2, 1312(,%r12,8)
	movq	%rax, 1280(,%r12,8)
	movq	%rax, 1288(,%r12,8)
	movq	%rax, 1296(,%r12,8)
	movl	$1, %eax
	salq	$32, %rax
	movq	%rdx, 32(,%r12,8)
	movq	%rax, 1304(,%r12,8)
	movl	g(%rip), %eax
	movq	%rdx, 24(,%r12,8)
	movq	%rdx, 16(,%r12,8)
	movq	%rdx, 8(,%r12,8)
	xorl	%edx, %edx
.L222:
	cmpl	%eax, %edx
	jnb	.L234
	movq	$-1, 40(%rbx,%rdx,8)
	incq	%rdx
	jmp	.L222
.L234:
	movl	%eax, %r13d
	subl	$3, %eax
	leaq	1000(%rbx), %rdi
	xorl	%r14d, %r14d
	shrl	%r13d
	leaq	8(%rbx), %rsi
	movl	$62, %ecx
	rep movsl
	imull	%eax, %r13d
	movl	$1, %ecx
	movq	%rbx, %rsi
	movq	%r15, %rdi
	imulq	$1320, %r13, %r13
	movq	%r13, %rdx
	call	cudaMemcpy@PLT
.L224:
	movl	g(%rip), %eax
	movl	%r14d, 1260(%rbx)
	movl	%eax, %edx
	subl	$3, %eax
	shrl	%edx
	imull	%edx, %eax
	cmpl	%r14d, %eax
	jb	.L235
	movq	%rbx, %rdi
	incl	%r14d
	call	_Z11thread_runRPv
	jmp	.L224
.L235:
	movl	$2, %ecx
	movq	%r13, %rdx
	movq	%r15, %rsi
	movq	%rbx, %rdi
	call	cudaMemcpy@PLT
	movl	1304(,%r12,8), %eax
	movq	1280(%rbx,%rax,8), %rcx
	movl	1308(,%r12,8), %eax
	movq	1280(%rbx,%rax,8), %rsi
	movl	1312(,%r12,8), %eax
	movq	1280(%rbx,%rax,8), %rax
	leaq	(%rcx,%rsi), %rdx
	addq	%rax, %rdx
	addq	%rax, %rax
	addq	%rsi, %rax
	movq	%rdx, UNIQUE(%rip)
	addq	%rax, %rax
	addq	%rcx, %rax
	addq	%rax, %rax
	movq	%rax, TOTAL(%rip)
	movq	-56(%rbp), %rax
	subq	%fs:40, %rax
	je	.L226
	call	__stack_chk_fail@PLT
.L226:
	leaq	-40(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r13
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa 7, 8
	ret
	.cfi_endproc
.LFE4014:
	.size	_Z15GPU_buildChainRjj, .-_Z15GPU_buildChainRjj
	.globl	_Z15GPU_carryChainRjj
	.type	_Z15GPU_carryChainRjj, @function
_Z15GPU_carryChainRjj:
.LFB4015:
	.cfi_startproc
	endbr64
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	%edi, 12(%rsp)
	movl	%esi, 8(%rsp)
	call	_Z9listChainv
	movl	8(%rsp), %esi
	movl	12(%rsp), %edi
	addq	$24, %rsp
	.cfi_def_cfa_offset 8
	jmp	_Z15GPU_buildChainRjj
	.cfi_endproc
.LFE4015:
	.size	_Z15GPU_carryChainRjj, .-_Z15GPU_carryChainRjj
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"There is no device.\n"
.LC1:
	.string	"There is no device supporting CUDA 1.x.\n"
	.text
	.globl	_Z8InitCUDAv
	.type	_Z8InitCUDAv, @function
_Z8InitCUDAv:
.LFB4016:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	subq	$1064, %rsp
	.cfi_def_cfa_offset 1088
	movq	%fs:40, %rax
	movq	%rax, 1048(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %rdi
	call	cudaGetDeviceCount@PLT
	cmpl	$0, 12(%rsp)
	leaq	.LC0(%rip), %rdx
	je	.L252
	xorl	%ebx, %ebx
	leaq	16(%rsp), %rbp
.L239:
	cmpl	%ebx, 12(%rsp)
	jle	.L243
	movl	%ebx, %esi
	movq	%rbp, %rdi
	call	cudaGetDeviceProperties_v2@PLT
	testl	%eax, %eax
	jne	.L241
	cmpl	$0, 376(%rsp)
	jg	.L243
.L241:
	incl	%ebx
	jmp	.L239
.L243:
	cmpl	%ebx, 12(%rsp)
	jne	.L245
	leaq	.LC1(%rip), %rdx
.L252:
	movq	stderr(%rip), %rdi
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	xorl	%eax, %eax
	jmp	.L238
.L245:
	movl	%ebx, %edi
	call	cudaSetDevice@PLT
	movb	$1, %al
.L238:
	movq	1048(%rsp), %rdx
	subq	%fs:40, %rdx
	je	.L247
	call	__stack_chk_fail@PLT
.L247:
	addq	$1064, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4016:
	.size	_Z8InitCUDAv, .-_Z8InitCUDAv
	.section	.rodata.str1.1
.LC2:
	.string	"Usage: %s [-c|-g|-r|-s] n steps\n"
.LC3:
	.string	"  -r: CPU \345\206\215\345\270\260\n"
.LC4:
	.string	"  -c: CPU \351\235\236\345\206\215\345\270\260\n"
.LC5:
	.string	"  -g: GPU \345\206\215\345\270\260\n"
.LC6:
	.string	"  -n: GPU \343\202\255\343\203\243\343\203\252\343\203\274\343\203\201\343\202\247\343\203\274\343\203\263\n"
.LC7:
	.string	"\n\nCPU \343\202\255\343\203\243\343\203\252\343\203\274\343\203\201\343\202\247\343\203\274\343\203\263 \345\206\215\345\270\260 \n"
.LC8:
	.string	"\n\nCPU \343\202\255\343\203\243\343\203\252\343\203\274\343\203\201\343\202\247\343\203\274\343\203\263 \351\235\236\345\206\215\345\270\260 \n"
.LC9:
	.string	"\n\nGPU \343\202\255\343\203\243\343\203\252\343\203\274\343\203\201\343\202\247\343\203\274\343\203\263 \343\202\267\343\203\263\343\202\260\343\203\253\343\202\271\343\203\254\343\203\203\343\203\211\n"
.LC10:
	.string	" N:        Total      Unique      dd:hh:mm:ss.ms"
.LC11:
	.string	"%s\n"
.LC12:
	.string	"%2d:%13ld%12ld%8.2d:%02d:%02d:%02d.%02d\n"
.LC13:
	.string	"\n\nGPU \343\202\255\343\203\243\343\203\252\343\203\274\343\203\201\343\202\247\343\203\274\343\203\263 \343\203\236\343\203\253\343\203\201\343\202\271\343\203\254\343\203\203\343\203\211\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB4017:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	subq	$56, %rsp
	.cfi_def_cfa_offset 96
	movq	%fs:40, %rax
	movq	%rax, 40(%rsp)
	xorl	%eax, %eax
	decl	%edi
	jle	.L254
	movq	8(%rsi), %rax
	cmpb	$45, (%rax)
	jne	.L261
	movb	1(%rax), %cl
	subl	$67, %ecx
	cmpb	$47, %cl
	ja	.L256
	movabsq	$4294967297, %rdx
	movl	$1, %eax
	salq	%cl, %rax
	testq	%rdx, %rax
	jne	.L276
	movabsq	$140737488388096, %rdx
	testq	%rdx, %rax
	jne	.L277
	movabsq	$68719476752, %rdx
	leaq	.LC9(%rip), %rsi
	testq	%rdx, %rax
	jne	.L294
	jmp	.L256
.L254:
	movq	(%rsi), %rdx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC2(%rip), %rsi
	call	__printf_chk@PLT
	leaq	.LC3(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC4(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC5(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	leaq	.LC6(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	jmp	.L261
.L276:
	xorl	%r12d, %r12d
	movb	$1, %bpl
	leaq	.LC8(%rip), %rsi
.L291:
	movl	$2, %edi
	xorl	%eax, %eax
	movl	$4, %ebx
	call	__printf_chk@PLT
	leaq	.LC10(%rip), %rdx
	leaq	.LC11(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	leaq	.LC12(%rip), %r13
	call	__printf_chk@PLT
.L266:
	xorl	%edx, %edx
	xorl	%esi, %esi
	leaq	8(%rsp), %rdi
	movq	%rdx, UNIQUE(%rip)
	movq	%rdx, TOTAL(%rip)
	call	gettimeofday@PLT
	testb	%r12b, %r12b
	je	.L262
	movl	%ebx, g(%rip)
	call	_Z11carryChainRv
.L262:
	testb	%bpl, %bpl
	je	.L263
	movl	%ebx, g(%rip)
	call	_Z10carryChainv
.L263:
	leaq	24(%rsp), %rdi
	xorl	%esi, %esi
	call	gettimeofday@PLT
	movq	32(%rsp), %rcx
	movq	16(%rsp), %rdi
	movq	24(%rsp), %rax
	cmpq	%rdi, %rcx
	jge	.L264
	subq	8(%rsp), %rax
	movl	$86400, %r8d
	subq	%rdi, %rcx
	decq	%rax
	addq	$1000500, %rcx
	cqto
	idivq	%r8
	movq	%rax, %r9
	movl	%edx, %esi
.L292:
	movq	%rcx, %rax
	movl	$10000, %edi
	movq	UNIQUE(%rip), %r8
	cqto
	idivq	%rdi
	movl	$3600, %edi
	movq	%rax, %rcx
	movl	%esi, %eax
	cltd
	pushq	%rcx
	.cfi_def_cfa_offset 104
	movl	$60, %ecx
	idivl	%edi
	movl	%eax, %edi
	movl	%esi, %eax
	cltd
	idivl	%ecx
	imull	$-3600, %edi, %eax
	addl	%esi, %eax
	movq	%r13, %rsi
	pushq	%rdx
	.cfi_def_cfa_offset 112
	cltd
	idivl	%ecx
	movq	TOTAL(%rip), %rcx
	movl	%ebx, %edx
	incl	%ebx
	pushq	%rax
	.cfi_def_cfa_offset 120
	xorl	%eax, %eax
	pushq	%rdi
	.cfi_def_cfa_offset 128
	movl	$2, %edi
	call	__printf_chk@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 96
	cmpl	$18, %ebx
	jne	.L266
.L261:
	movq	40(%rsp), %rax
	subq	%fs:40, %rax
	je	.L274
	call	__stack_chk_fail@PLT
.L277:
	movb	$1, %r12b
	xorl	%ebp, %ebp
	leaq	.LC7(%rip), %rsi
	jmp	.L291
.L264:
	subq	8(%rsp), %rax
	movl	$86400, %r8d
	subq	%rdi, %rcx
	cqto
	addq	$500, %rcx
	idivq	%r8
	movq	%rax, %r9
	movl	%edx, %esi
	jmp	.L292
.L256:
	leaq	.LC13(%rip), %rsi
.L294:
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	call	_Z8InitCUDAv
	testb	%al, %al
	je	.L261
	leaq	.LC10(%rip), %rdx
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$4, %ebx
	leaq	.LC12(%rip), %rbp
.L269:
	xorl	%esi, %esi
	leaq	8(%rsp), %rdi
	call	gettimeofday@PLT
	xorl	%eax, %eax
	movl	$24576, %esi
	movl	%ebx, %edi
	movq	%rax, UNIQUE(%rip)
	movq	%rax, TOTAL(%rip)
	movl	%ebx, g(%rip)
	call	_Z15GPU_carryChainRjj
	xorl	%esi, %esi
	leaq	24(%rsp), %rdi
	call	gettimeofday@PLT
	movq	32(%rsp), %rcx
	movq	16(%rsp), %rdi
	movq	24(%rsp), %rsi
	movq	8(%rsp), %r8
	cmpq	%rdi, %rcx
	jge	.L267
	leal	-1(%rsi), %eax
	movl	$86400, %r9d
	subq	%r8, %rsi
	subq	%rdi, %rcx
	subl	%r8d, %eax
	movl	$86400, %r8d
	addq	$1000500, %rcx
	cltd
	idivl	%r9d
	movl	%eax, %r9d
	leaq	-1(%rsi), %rax
	cqto
	idivq	%r8
	movl	%edx, %esi
	jmp	.L293
.L267:
	movl	%esi, %eax
	movl	$86400, %r10d
	subq	%rdi, %rcx
	subl	%r8d, %eax
	addq	$500, %rcx
	cltd
	idivl	%r10d
	movl	%eax, %r9d
	movq	%rsi, %rax
	subq	%r8, %rax
	movl	$86400, %r8d
	cqto
	idivq	%r8
	movl	%edx, %esi
.L293:
	movq	%rcx, %rax
	movl	$10000, %edi
	movq	UNIQUE(%rip), %r8
	cqto
	idivq	%rdi
	movl	$3600, %edi
	movq	%rax, %rcx
	movl	%esi, %eax
	cltd
	pushq	%rcx
	.cfi_def_cfa_offset 104
	movl	$60, %ecx
	idivl	%edi
	movl	%eax, %edi
	movl	%esi, %eax
	cltd
	idivl	%ecx
	imull	$-3600, %edi, %eax
	addl	%esi, %eax
	movq	%rbp, %rsi
	pushq	%rdx
	.cfi_def_cfa_offset 112
	cltd
	idivl	%ecx
	movq	TOTAL(%rip), %rcx
	movl	%ebx, %edx
	incl	%ebx
	pushq	%rax
	.cfi_def_cfa_offset 120
	xorl	%eax, %eax
	pushq	%rdi
	.cfi_def_cfa_offset 128
	movl	$2, %edi
	call	__printf_chk@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 96
	cmpl	$22, %ebx
	jne	.L269
	jmp	.L261
.L274:
	addq	$56, %rsp
	.cfi_def_cfa_offset 40
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 32
	popq	%rbp
	.cfi_def_cfa_offset 24
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4017:
	.size	main, .-main
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4043:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE4043:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.globl	g
	.bss
	.align 32
	.type	g, @object
	.size	g, 7444
g:
	.zero	7444
	.globl	UNIQUE
	.align 8
	.type	UNIQUE, @object
	.size	UNIQUE, 8
UNIQUE:
	.zero	8
	.globl	TOTAL
	.align 8
	.type	TOTAL, @object
	.size	TOTAL, 8
TOTAL:
	.zero	8
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

/**
 *
 * bash版キャリーチェーンのC言語版のGPU/CUDA移植版
 *
 詳しい説明はこちらをどうぞ
 https://suzukiiichiro.github.io/search/?keyword=Ｎクイーン問題
 *
アーキテクチャの指定（なくても問題なし、あれば高速）
-arch=sm_13 or -arch=sm_61

CPUの再帰での実行
$ nvcc -O3 -arch=sm_61 05CUDA_CarryChain.cu && ./a.out -r

CPUの非再帰での実行
$ nvcc -O3 -arch=sm_61 05CUDA_CarryChain.cu && ./a.out -c

GPUのシングルスレッド
$ nvcc -O3 -arch=sm_61 05CUDA_CarryChain.cu && ./a.out -g

GPUのマルチスレッド
$ nvcc -O3 -arch=sm_61 05CUDA_CarryChain.cu && ./a.out -n
*/
#include <iostream>
#include <vector>
#include <stdio.h>
#include <stdlib.h>
#include <stdbool.h>
#include <math.h>
#include <string.h>
#include <time.h>
#include <sys/time.h>
#include <hip/hip_runtime.h>


#define THREAD_NUM		96
#define MAX 27
// システムによって以下のマクロが必要であればコメントを外してください。
//#define UINT64_C(c) c ## ULL
//
// グローバル変数
unsigned long TOTAL=0; 
unsigned long UNIQUE=0;
// キャリーチェーン 非再帰版
// 構造体
typedef struct
{
  unsigned int size;
  unsigned int pres_a[930]; 
  unsigned int pres_b[930];
  // uint64_t COUNTER[3];      
  // //カウンター配列
  // unsigned int COUNT2;
  // unsigned int COUNT4;
  // unsigned int COUNT8;
}Global; Global g;
// 構造体
typedef struct Board
{
  uint64_t row;
  uint64_t down;
  uint64_t left;
  uint64_t right;
  long long x[MAX];
}Board ;
typedef struct Local
{
  unsigned int size;
  struct Board B;
  struct Board nB;
  struct Board eB;
  struct Board sB;
  struct Board wB;
  unsigned n;
  unsigned e;
  unsigned s;
  unsigned w;
  uint64_t dimx;
  uint64_t dimy;
  uint64_t COUNTER[3];      
  //カウンター配列
  unsigned int COUNT2;
  unsigned int COUNT4;
  unsigned int COUNT8;
  unsigned int STEPS;
}Local;
/**
  CPU/CPUR 再帰・非再帰共通
  */
// チェーンのリストを作成
void listChain()
{
  unsigned int idx=0;
  for(unsigned int a=0;a<(unsigned)g.size;++a){
    for(unsigned int b=0;b<(unsigned)g.size;++b){
      if(((a>=b)&&(a-b)<=1)||((b>a)&&(b-a)<=1)){ continue; }
      g.pres_a[idx]=a;
      g.pres_b[idx]=b;
      ++idx;
    }
  }
}
/**
  CPU 非再帰
*/
// クイーンの効きをチェック
bool placement(void* args)
{
  Local *l=(Local *)args;
  if(l->B.x[l->dimx]==l->dimy){ return true;  }  
  if (l->B.x[0]==0){
    if (l->B.x[1]!=(uint64_t)-1){
      if((l->B.x[1]>=l->dimx)&&(l->dimy==1)){ return false; }
    }
  }else{
    if( (l->B.x[0]!=(uint64_t)-1) ){
      if(( (l->dimx<l->B.x[0]||l->dimx>=g.size-l->B.x[0])
        && (l->dimy==0 || l->dimy==g.size-1)
      )){ return 0; } 
      if ((  (l->dimx==g.size-1)&&((l->dimy<=l->B.x[0])||
          l->dimy>=g.size-l->B.x[0]))){
        return 0;
      } 
    }
  }
  l->B.x[l->dimx]=l->dimy;                    //xは行 yは列
  uint64_t row=UINT64_C(1)<<l->dimx;
  uint64_t down=UINT64_C(1)<<l->dimy;
  uint64_t left=UINT64_C(1)<<(g.size-1-l->dimx+l->dimy); //右上から左下
  uint64_t right=UINT64_C(1)<<(l->dimx+l->dimy);       // 左上から右下
  if((l->B.row&row)||(l->B.down&down)||(l->B.left&left)||(l->B.right&right)){ return false; }     
  l->B.row|=row; l->B.down|=down; l->B.left|=left; l->B.right|=right;
  return true;
}
//非再帰
uint64_t solve(int size,int current,uint64_t row,uint64_t left,uint64_t down,uint64_t right)
{
  uint64_t row_a[MAX];
  uint64_t right_a[MAX];
  uint64_t left_a[MAX];
  uint64_t down_a[MAX];
  uint64_t bitmap_a[MAX];
  for (int i=0;i<size;i++){
    row_a[i]=0;
    left_a[i]=0;
    down_a[i]=0;
    right_a[i]=0;
    bitmap_a[i]=0;
  }
  row_a[current]=row;
  left_a[current]=left;
  down_a[current]=down;
  right_a[current]=right;
  uint64_t bitmap=bitmap_a[current]=~(left_a[current]|down_a[current]|right_a[current]);
  uint64_t total=0;
  uint64_t bit;

  while(current>-1){
    if((bitmap!=0||row&1)&&current<size){
      if(!(down+1)){
        total++;
        current--;
        row=row_a[current];
        left=left_a[current];
        right=right_a[current];
        down=down_a[current];
        bitmap=bitmap_a[current];
        continue;
      }else if(row&1){
        while( row&1 ){
          row>>=1;
          left<<=1;
          right>>=1;
        }
        bitmap=~(left|down|right);  //再帰に必要な変数は必ず定義する必要があります。
        continue;
      }else{
        bit=-bitmap&bitmap;
        bitmap=bitmap^bit;
        if(current<size){
          row_a[current]=row;
          left_a[current]=left;
          down_a[current]=down;
          right_a[current]=right;
          bitmap_a[current]=bitmap;
          current++;
        }
        row>>=1;      //１行下に移動する
        left=(left|bit)<<1;
        down=down|bit;
        right=(right|bit)>>1;
        bitmap=~(left|down|right);  //再帰に必要な変数は必ず定義する必要があります。
      }
    }else{
      current--;
      row=row_a[current];
      left=left_a[current];
      right=right_a[current];
      down=down_a[current];
      bitmap=bitmap_a[current];
    }
  }
  return total;
}
//非再帰 対称解除法
void carryChain_symmetry(void* args)
{
  Local *l=(Local *)args;
  // 対称解除法
  unsigned const int ww=(g.size-2)*(g.size-1)-1-l->w;
  unsigned const int w2=(g.size-2)*(g.size-1)-1;
  // # 対角線上の反転が小さいかどうか確認する
  if((l->s==ww)&&(l->n<(w2-l->e))){ return ; }
  // # 垂直方向の中心に対する反転が小さいかを確認
  if((l->e==ww)&&(l->n>(w2-l->n))){ return; }
  // # 斜め下方向への反転が小さいかをチェックする
  if((l->n==ww)&&(l->e>(w2-l->s))){ return; }
  // 枝刈り １行目が角の場合回転対称チェックせずCOUNT8にする
  if(l->B.x[0]==0){
    l->COUNTER[l->COUNT8]+=solve(g.size,0,l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return ;
  }
  // n,e,s==w の場合は最小値を確認する。右回転で同じ場合は、
  // w=n=e=sでなければ値が小さいのでskip  w=n=e=sであれば90度回転で同じ可能性
  if(l->s==l->w){ if((l->n!=l->w)||(l->e!=l->w)){ return; }
    l->COUNTER[l->COUNT2]+=solve(g.size,0,l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return;
  }
  // e==wは180度回転して同じ 180度回転して同じ時n>=sの時はsmaller?
  if((l->e==l->w)&&(l->n>=l->s)){ if(l->n>l->s){ return; }
    l->COUNTER[l->COUNT4]+=solve(g.size,0,l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return;
  }
  l->COUNTER[l->COUNT8]+=solve(g.size,0,l->B.row>>2,
  l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
  return;
}
//非再帰  pthread run()
void thread_run(void* args)
{
  Local *l=(Local *)args;

  // memcpy(&l->B,&l->wB,sizeof(Board));       // B=wB;
  l->B=l->wB;
  l->dimx=0; l->dimy=g.pres_a[l->w];
  //if(!placement(l)){ continue; }
  if(!placement(l)){ return; }
  l->dimx=1; l->dimy=g.pres_b[l->w];
  // if(!placement(l)){ continue; }
  if(!placement(l)){ return; }
  //２ 左２行に置く
  // memcpy(&l->nB,&l->B,sizeof(Board));       // nB=B;
  l->nB=l->B;
  for(l->n=l->w;l->n<(g.size-2)*(g.size-1)-l->w;++l->n){
    // memcpy(&l->B,&l->nB,sizeof(Board));     // B=nB;
    l->B=l->nB;
    l->dimx=g.pres_a[l->n]; l->dimy=g.size-1;
    if(!placement(l)){ continue; }
    l->dimx=g.pres_b[l->n]; l->dimy=g.size-2;
    if(!placement(l)){ continue; }
    // ３ 下２行に置く
    // memcpy(&l->eB,&l->B,sizeof(Board));     // eB=B;
    l->eB=l->B;
    for(l->e=l->w;l->e<(g.size-2)*(g.size-1)-l->w;++l->e){
      // memcpy(&l->B,&l->eB,sizeof(Board));   // B=eB;
      l->B=l->eB;
      l->dimx=g.size-1; l->dimy=g.size-1-g.pres_a[l->e];
      if(!placement(l)){ continue; }
      l->dimx=g.size-2; l->dimy=g.size-1-g.pres_b[l->e];
      if(!placement(l)){ continue; }
      // ４ 右２列に置く
      // memcpy(&l->sB,&l->B,sizeof(Board));   // sB=B;
      l->sB=l->B;
      for(l->s=l->w;l->s<(g.size-2)*(g.size-1)-l->w;++l->s){
        // memcpy(&l->B,&l->sB,sizeof(Board)); // B=sB;
        l->B=l->sB;
        l->dimx=g.size-1-g.pres_a[l->s]; l->dimy=0;
        if(!placement(l)){ continue; }
        l->dimx=g.size-1-g.pres_b[l->s]; l->dimy=1;
        if(!placement(l)){ continue; }
        // 対称解除法
        carryChain_symmetry(l);
      } //w
    } //e
  } //n
}
//非再帰  チェーンのビルド
void buildChain()
{
  Local l[(g.size/2)*(g.size-3)];

  // カウンターの初期化
  l->COUNT2=0; l->COUNT4=1; l->COUNT8=2;
  l->COUNTER[l->COUNT2]=l->COUNTER[l->COUNT4]=l->COUNTER[l->COUNT8]=0;
  // Board の初期化 nB,eB,sB,wB;
  l->B.row=l->B.down=l->B.left=l->B.right=0;
  // Board x[]の初期化
  for(unsigned int i=0;i<g.size;++i){ l->B.x[i]=-1; }
  //１ 上２行に置く
  // memcpy(&l->wB,&l->B,sizeof(Board));         // wB=B;
  l->wB=l->B;
  for(l->w=0;l->w<=(unsigned)(g.size/2)*(g.size-3);++l->w){
    thread_run(&l);
  } //w
  /**
   * 集計
   */
  UNIQUE= l->COUNTER[l->COUNT2]+
          l->COUNTER[l->COUNT4]+
          l->COUNTER[l->COUNT8];
  TOTAL=  l->COUNTER[l->COUNT2]*2+
          l->COUNTER[l->COUNT4]*4+
          l->COUNTER[l->COUNT8]*8;
}
//非再帰  キャリーチェーン
void carryChain()
{
  listChain();  //チェーンのリストを作成
  buildChain(); // チェーンのビルド
  // calcChain(&l);  // 集計
}
/**
  CPUR 再帰
  */
//再帰 ボード外側２列を除く内側のクイーン配置処理
uint64_t solveR(uint64_t row,uint64_t left,uint64_t down,uint64_t right)
{
  if(down+1==0){ return  1; }
  while((row&1)!=0) { 
    row>>=1;
    left<<=1;
    right>>=1;
  }
  row>>=1;
  uint64_t total=0;
  for(uint64_t carryChain=~(left|down|right);carryChain!=0;){
    uint64_t const bit=carryChain&-carryChain;
    total+=solveR(row,(left|bit)<<1,down|bit,(right|bit)>>1);
    carryChain^=bit;
  }
  return total;
} 
//再帰 対称解除法
void carryChain_symmetryR(void* args)
{
  Local *l=(Local *)args;
  // 対称解除法
  unsigned const int ww=(g.size-2)*(g.size-1)-1-l->w;
  unsigned const int w2=(g.size-2)*(g.size-1)-1;
  // # 対角線上の反転が小さいかどうか確認する
  if((l->s==ww)&&(l->n<(w2-l->e))){ return ; }
  // # 垂直方向の中心に対する反転が小さいかを確認
  if((l->e==ww)&&(l->n>(w2-l->n))){ return; }
  // # 斜め下方向への反転が小さいかをチェックする
  if((l->n==ww)&&(l->e>(w2-l->s))){ return; }
  // 枝刈り １行目が角の場合回転対称チェックせずCOUNT8にする
  if(l->B.x[0]==0){
    l->COUNTER[l->COUNT8]+=solveR(l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return ;
  }
  // n,e,s==w の場合は最小値を確認する。右回転で同じ場合は、
  // w=n=e=sでなければ値が小さいのでskip  w=n=e=sであれば90度回転で同じ可能性
  if(l->s==l->w){ if((l->n!=l->w)||(l->e!=l->w)){ return; }
    l->COUNTER[l->COUNT2]+=solveR(l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return;
  }
  // e==wは180度回転して同じ 180度回転して同じ時n>=sの時はsmaller?
  if((l->e==l->w)&&(l->n>=l->s)){ if(l->n>l->s){ return; }
    l->COUNTER[l->COUNT4]+=solveR(l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return;
  }
  l->COUNTER[l->COUNT8]+=solveR(l->B.row>>2,
  l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
  return;
}
//再帰  pthread run()
void thread_runR(void* args)
{
  Local *l=(Local *)args;

  // memcpy(&l->B,&l->wB,sizeof(Board));       // B=wB;
  l->B=l->wB;
  l->dimx=0; l->dimy=g.pres_a[l->w];
  //if(!placement(l)){ continue; }
  if(!placement(l)){ return; }
  l->dimx=1; l->dimy=g.pres_b[l->w];
  // if(!placement(l)){ continue; }
  if(!placement(l)){ return; }
  //２ 左２行に置く
  // memcpy(&l->nB,&l->B,sizeof(Board));       // nB=B;
  l->nB=l->B;
  for(l->n=l->w;l->n<(g.size-2)*(g.size-1)-l->w;++l->n){
    // memcpy(&l->B,&l->nB,sizeof(Board));     // B=nB;
    l->B=l->nB;
    l->dimx=g.pres_a[l->n]; l->dimy=g.size-1;
    if(!placement(l)){ continue; }
    l->dimx=g.pres_b[l->n]; l->dimy=g.size-2;
    if(!placement(l)){ continue; }
    // ３ 下２行に置く
    // memcpy(&l->eB,&l->B,sizeof(Board));     // eB=B;
    l->eB=l->B;
    for(l->e=l->w;l->e<(g.size-2)*(g.size-1)-l->w;++l->e){
      // memcpy(&l->B,&l->eB,sizeof(Board));   // B=eB;
      l->B=l->eB;
      l->dimx=g.size-1; l->dimy=g.size-1-g.pres_a[l->e];
      if(!placement(l)){ continue; }
      l->dimx=g.size-2; l->dimy=g.size-1-g.pres_b[l->e];
      if(!placement(l)){ continue; }
      // ４ 右２列に置く
      // memcpy(&l->sB,&l->B,sizeof(Board));   // sB=B;
      l->sB=l->B;
      for(l->s=l->w;l->s<(g.size-2)*(g.size-1)-l->w;++l->s){
        // memcpy(&l->B,&l->sB,sizeof(Board)); // B=sB;
        l->B=l->sB;
        l->dimx=g.size-1-g.pres_a[l->s]; l->dimy=0;
        if(!placement(l)){ continue; }
        l->dimx=g.size-1-g.pres_b[l->s]; l->dimy=1;
        if(!placement(l)){ continue; }
        // 対称解除法
        carryChain_symmetryR(l);
      } //w
    } //e
  } //n
}
//再帰  チェーンのビルド
void buildChainR()
{
  Local l[(g.size/2)*(g.size-3)];

  // カウンターの初期化
  l->COUNT2=0; l->COUNT4=1; l->COUNT8=2;
  l->COUNTER[l->COUNT2]=l->COUNTER[l->COUNT4]=l->COUNTER[l->COUNT8]=0;
  // Board の初期化 nB,eB,sB,wB;
  l->B.row=l->B.down=l->B.left=l->B.right=0;
  // Board x[]の初期化
  for(unsigned int i=0;i<g.size;++i){ l->B.x[i]=-1; }
  //１ 上２行に置く
  // memcpy(&l->wB,&l->B,sizeof(Board));         // wB=B;
  l->wB=l->B;
  for(l->w=0;l->w<=(unsigned)(g.size/2)*(g.size-3);++l->w){
    thread_runR(&l);
  } //w
  /**
   * 集計
   */
  UNIQUE= l->COUNTER[l->COUNT2]+
          l->COUNTER[l->COUNT4]+
          l->COUNTER[l->COUNT8];
  TOTAL=  l->COUNTER[l->COUNT2]*2+
          l->COUNTER[l->COUNT4]*4+
          l->COUNTER[l->COUNT8]*8;
}
//再帰  キャリーチェーン
void carryChainR()
{
  listChain();  //チェーンのリストを作成
  buildChainR(); // チェーンのビルド
  // calcChain(&l);  // 集計
}
/**
  GPU 
 */
// GPU クイーンの効きをチェック
bool GPU_placement(void* args)
{
  Local *l=(Local *)args;
  if(l->B.x[l->dimx]==l->dimy){ return true;  }  
  if (l->B.x[0]==0){
    if (l->B.x[1]!=(uint64_t)-1){
      if((l->B.x[1]>=l->dimx)&&(l->dimy==1)){ return false; }
    }
  }else{
    if( (l->B.x[0]!=(uint64_t)-1) ){
      if(( (l->dimx<l->B.x[0]||l->dimx>=g.size-l->B.x[0])
        && (l->dimy==0 || l->dimy==g.size-1)
      )){ return 0; } 
      if ((  (l->dimx==g.size-1)&&((l->dimy<=l->B.x[0])||
          l->dimy>=g.size-l->B.x[0]))){
        return 0;
      } 
    }
  }
  l->B.x[l->dimx]=l->dimy;                    //xは行 yは列
  uint64_t row=UINT64_C(1)<<l->dimx;
  uint64_t down=UINT64_C(1)<<l->dimy;
  uint64_t left=UINT64_C(1)<<(g.size-1-l->dimx+l->dimy); //右上から左下
  uint64_t right=UINT64_C(1)<<(l->dimx+l->dimy);       // 左上から右下
  if((l->B.row&row)||(l->B.down&down)||(l->B.left&left)||(l->B.right&right)){ return false; }     
  l->B.row|=row; l->B.down|=down; l->B.left|=left; l->B.right|=right;
  return true;
}
//GPU 再帰 ボード外側２列を除く内側のクイーン配置処理
uint64_t GPU_solveR(uint64_t row,uint64_t left,uint64_t down,uint64_t right)
{
  if(down+1==0){ return  1; }
  while((row&1)!=0) { 
    row>>=1;
    left<<=1;
    right>>=1;
  }
  row>>=1;
  uint64_t total=0;
  for(uint64_t carryChain=~(left|down|right);carryChain!=0;){
    uint64_t const bit=carryChain&-carryChain;
    total+=solveR(row,(left|bit)<<1,down|bit,(right|bit)>>1);
    carryChain^=bit;
  }
  return total;
} 
//GPU 再帰 対称解除法
void GPU_carryChain_symmetryR(void* args)
{
  Local *l=(Local *)args;
  // 対称解除法
  unsigned const int ww=(g.size-2)*(g.size-1)-1-l->w;
  unsigned const int w2=(g.size-2)*(g.size-1)-1;
  // # 対角線上の反転が小さいかどうか確認する
  if((l->s==ww)&&(l->n<(w2-l->e))){ return ; }
  // # 垂直方向の中心に対する反転が小さいかを確認
  if((l->e==ww)&&(l->n>(w2-l->n))){ return; }
  // # 斜め下方向への反転が小さいかをチェックする
  if((l->n==ww)&&(l->e>(w2-l->s))){ return; }
  // 枝刈り １行目が角の場合回転対称チェックせずCOUNT8にする
  if(l->B.x[0]==0){
    l->COUNTER[l->COUNT8]+=solveR(l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return ;
  }
  // n,e,s==w の場合は最小値を確認する。右回転で同じ場合は、
  // w=n=e=sでなければ値が小さいのでskip  w=n=e=sであれば90度回転で同じ可能性
  if(l->s==l->w){ if((l->n!=l->w)||(l->e!=l->w)){ return; }
    l->COUNTER[l->COUNT2]+=solveR(l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return;
  }
  // e==wは180度回転して同じ 180度回転して同じ時n>=sの時はsmaller?
  if((l->e==l->w)&&(l->n>=l->s)){ if(l->n>l->s){ return; }
    l->COUNTER[l->COUNT4]+=solveR(l->B.row>>2,
    l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
    return;
  }
  l->COUNTER[l->COUNT8]+=solveR(l->B.row>>2,
  l->B.left>>4,((((l->B.down>>2)|(~0<<(g.size-4)))+1)<<(g.size-5))-1,(l->B.right>>4)<<(g.size-5));
  return;
}
//GPU 再帰  pthread run()
void GPU_thread_runR(void* args)
{
  Local *l=(Local *)args;

  // memcpy(&l->B,&l->wB,sizeof(Board));       // B=wB;
  l->B=l->wB;
  l->dimx=0; l->dimy=g.pres_a[l->w];
  //if(!GPU_placement(l)){ continue; }
  if(!GPU_placement(l)){ return; }
  l->dimx=1; l->dimy=g.pres_b[l->w];
  // if(!GPU_placement(l)){ continue; }
  if(!GPU_placement(l)){ return; }
  //２ 左２行に置く
  // memcpy(&l->nB,&l->B,sizeof(Board));       // nB=B;
  l->nB=l->B;
  for(l->n=l->w;l->n<(g.size-2)*(g.size-1)-l->w;++l->n){
    // memcpy(&l->B,&l->nB,sizeof(Board));     // B=nB;
    l->B=l->nB;
    l->dimx=g.pres_a[l->n]; l->dimy=g.size-1;
    if(!GPU_placement(l)){ continue; }
    l->dimx=g.pres_b[l->n]; l->dimy=g.size-2;
    if(!GPU_placement(l)){ continue; }
    // ３ 下２行に置く
    // memcpy(&l->eB,&l->B,sizeof(Board));     // eB=B;
    l->eB=l->B;
    for(l->e=l->w;l->e<(g.size-2)*(g.size-1)-l->w;++l->e){
      // memcpy(&l->B,&l->eB,sizeof(Board));   // B=eB;
      l->B=l->eB;
      l->dimx=g.size-1; l->dimy=g.size-1-g.pres_a[l->e];
      if(!GPU_placement(l)){ continue; }
      l->dimx=g.size-2; l->dimy=g.size-1-g.pres_b[l->e];
      if(!GPU_placement(l)){ continue; }
      // ４ 右２列に置く
      // memcpy(&l->sB,&l->B,sizeof(Board));   // sB=B;
      l->sB=l->B;
      for(l->s=l->w;l->s<(g.size-2)*(g.size-1)-l->w;++l->s){
        // memcpy(&l->B,&l->sB,sizeof(Board)); // B=sB;
        l->B=l->sB;
        l->dimx=g.size-1-g.pres_a[l->s]; l->dimy=0;
        if(!GPU_placement(l)){ continue; }
        l->dimx=g.size-1-g.pres_b[l->s]; l->dimy=1;
        if(!GPU_placement(l)){ continue; }
        // 対称解除法
        carryChain_symmetryR(l);
      } //w
    } //e
  } //n
}
//GPU 再帰  チェーンのビルド
void GPU_buildChainR(const unsigned int size,unsigned int STEPS)
{
  Local l[(g.size/2)*(g.size-3)];
  l->STEPS=STEPS;
  l->size=size;
  Local lDevice[(g.size/2)*(g.size-3)];
  hipHostMalloc((void**) &l,   sizeof(struct Local)*l->STEPS, hipHostMallocDefault);
  hipMalloc((void**) &lDevice, sizeof(struct Local)*l->STEPS);

  // カウンターの初期化
  l->COUNT2=0; l->COUNT4=1; l->COUNT8=2;
  l->COUNTER[l->COUNT2]=l->COUNTER[l->COUNT4]=l->COUNTER[l->COUNT8]=0;
  // Board の初期化 nB,eB,sB,wB;
  l->B.row=l->B.down=l->B.left=l->B.right=0;
  // Board x[]の初期化
  for(unsigned int i=0;i<g.size;++i){ l->B.x[i]=-1; }
  //１ 上２行に置く
  // memcpy(&l->wB,&l->B,sizeof(Board));         // wB=B;
  l->wB=l->B;
  unsigned int limit=(unsigned)(g.size/2)*(g.size-3);
  hipMemcpy(lDevice,l,
      sizeof(struct Local)*limit,hipMemcpyHostToDevice);
  for(l->w=0;l->w<=(unsigned)(g.size/2)*(g.size-3);++l->w){
    thread_runR(&l);
    //GPU_thread_runR<<<l->STEPS/THREAD_NUM,THREAD_NUM>>>(&l);
    //GPU_thread_runR<<<l->STEPS/THREAD_NUM,THREAD_NUM>>>(&l);
  } //w
  hipMemcpy(l,lDevice,
      sizeof(struct Local)*limit,hipMemcpyDeviceToHost);
  /**
   * 集計
   */
  UNIQUE= l->COUNTER[l->COUNT2]+
          l->COUNTER[l->COUNT4]+
          l->COUNTER[l->COUNT8];
  TOTAL=  l->COUNTER[l->COUNT2]*2+
          l->COUNTER[l->COUNT4]*4+
          l->COUNTER[l->COUNT8]*8;
}
//GPU 再帰  キャリーチェーン
void GPU_carryChainR(const unsigned int size,unsigned int STEPS)
{
  listChain();  //チェーンのリストを作成
  GPU_buildChainR(size,STEPS); // チェーンのビルド
  // calcChain(&l);  // 集計
}
// CUDA 初期化
bool InitCUDA()
{
  int count;
  hipGetDeviceCount(&count);
  if(count==0){fprintf(stderr,"There is no device.\n");return false;}
  int i;
  for(i=0;i<count;i++){
    struct hipDeviceProp_t prop;
    if(hipGetDeviceProperties(&prop,i)==hipSuccess){if(prop.major>=1){break;} }
  }
  if(i==count){fprintf(stderr,"There is no device supporting CUDA 1.x.\n");return false;}
  hipSetDevice(i);
  return true;
}
//メイン
int main(int argc,char** argv)
{
  bool cpu=false,cpur=false,gpu=false,gpuNodeLayer=false;
  int argstart=2;
  if(argc>=2&&argv[1][0]=='-'){
    if(argv[1][1]=='c'||argv[1][1]=='C'){cpu=true;}
    else if(argv[1][1]=='r'||argv[1][1]=='R'){cpur=true;}
    else if(argv[1][1]=='c'||argv[1][1]=='C'){cpu=true;}
    else if(argv[1][1]=='g'||argv[1][1]=='G'){gpu=true;}
    else if(argv[1][1]=='n'||argv[1][1]=='N'){gpuNodeLayer=true;}
    else{ gpuNodeLayer=true; } //デフォルトをgpuとする
    argstart=2;
  }
  if(argc<argstart){
    printf("Usage: %s [-c|-g|-r|-s] n steps\n",argv[0]);
    printf("  -r: CPU 再帰\n");
    printf("  -c: CPU 非再帰\n");
    printf("  -g: GPU 再帰\n");
    printf("  -n: GPU キャリーチェーン\n");
  }
  if(cpur){ printf("\n\nCPU キャリーチェーン 再帰 \n"); }
  else if(cpu){ printf("\n\nCPU キャリーチェーン 非再帰 \n"); }
  else if(gpu){ printf("\n\nGPU キャリーチェーン シングルスレッド\n"); }
  else if(gpuNodeLayer){ printf("\n\nGPU キャリーチェーン マルチスレッド\n"); }
  if(cpu||cpur)
  {
    int min=4; 
    int targetN=17;
    struct timeval t0;
    struct timeval t1;
    printf("%s\n"," N:        Total      Unique      dd:hh:mm:ss.ms");
    for(int size=min;size<=targetN;size++){
      TOTAL=UNIQUE=0;
      gettimeofday(&t0, NULL);//計測開始
      if(cpur){ //再帰
        g.size=size;
        carryChainR();
      }
      if(cpu){ //非再帰
        g.size=size;
        carryChain();
      }
      //
      gettimeofday(&t1, NULL);//計測終了
      int ss;int ms;int dd;
      if(t1.tv_usec<t0.tv_usec) {
        dd=(t1.tv_sec-t0.tv_sec-1)/86400;
        ss=(t1.tv_sec-t0.tv_sec-1)%86400;
        ms=(1000000+t1.tv_usec-t0.tv_usec+500)/10000;
      }else {
        dd=(t1.tv_sec-t0.tv_sec)/86400;
        ss=(t1.tv_sec-t0.tv_sec)%86400;
        ms=(t1.tv_usec-t0.tv_usec+500)/10000;
      }//end if
      int hh=ss/3600;
      int mm=(ss-hh*3600)/60;
      ss%=60;
      printf("%2d:%13ld%12ld%8.2d:%02d:%02d:%02d.%02d\n",size,TOTAL,UNIQUE,dd,hh,mm,ss,ms);
    } //end for
  }//end if
  if(gpu||gpuNodeLayer)
  {
    if(!InitCUDA()){return 0;}
    int STEPS=24576;
    int min=4;
    int targetN=21;
    struct timeval t0;
    struct timeval t1;
    printf("%s\n"," N:        Total      Unique      dd:hh:mm:ss.ms");
    for(int size=min;size<=targetN;size++){
      gettimeofday(&t0,NULL);   // 計測開始
      if(gpu){
        TOTAL=UNIQUE=0;
        g.size=size;
        GPU_carryChainR(size,STEPS); //キャリーチェーン
      }else if(gpuNodeLayer){
        TOTAL=UNIQUE=0;
        g.size=size;
        GPU_carryChainR(size,STEPS); // キャリーチェーン
      }
      gettimeofday(&t1,NULL);   // 計測終了
      int ss;int ms;int dd;
      if (t1.tv_usec<t0.tv_usec) {
        dd=(int)(t1.tv_sec-t0.tv_sec-1)/86400;
        ss=(t1.tv_sec-t0.tv_sec-1)%86400;
        ms=(1000000+t1.tv_usec-t0.tv_usec+500)/10000;
      } else {
        dd=(int)(t1.tv_sec-t0.tv_sec)/86400;
        ss=(t1.tv_sec-t0.tv_sec)%86400;
        ms=(t1.tv_usec-t0.tv_usec+500)/10000;
      }//end if
      int hh=ss/3600;
      int mm=(ss-hh*3600)/60;
      ss%=60;
      printf("%2d:%13ld%12ld%8.2d:%02d:%02d:%02d.%02d\n",size,TOTAL,UNIQUE,dd,hh,mm,ss,ms);
    }//end for
  }//end if
  return 0;
}

	.text
	.file	"05CUDA_CarryChain_0801-1_suzuki.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z9listChainv                   # -- Begin function _Z9listChainv
	.type	_Z9listChainv,@function
_Z9listChainv:                          # @_Z9listChainv
	.cfi_startproc
# %bb.0:
	movl	g(%rip), %eax
	testl	%eax, %eax
	je	.LBB0_10
# %bb.1:                                # %.preheader.preheader
	xorl	%ecx, %ecx
	xorl	%edx, %edx
	xorl	%esi, %esi
.LBB0_2:                                # %.preheader
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB0_3 Depth 2
	movl	%edx, %edi
	xorl	%r8d, %r8d
.LBB0_3:                                #   Parent Loop BB0_2 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	cmpl	%r8d, %edx
	jb	.LBB0_5
# %bb.4:                                #   in Loop: Header=BB0_3 Depth=2
	cmpl	$2, %edi
	jb	.LBB0_8
.LBB0_5:                                #   in Loop: Header=BB0_3 Depth=2
	cmpl	%edx, %r8d
	jbe	.LBB0_7
# %bb.6:                                #   in Loop: Header=BB0_3 Depth=2
	leal	(%rcx,%r8), %r9d
	cmpl	$2, %r9d
	jb	.LBB0_8
.LBB0_7:                                #   in Loop: Header=BB0_3 Depth=2
	movl	%esi, %r9d
	movl	%edx, g+4(,%r9,4)
	movl	%r8d, g+3724(,%r9,4)
	incl	%esi
.LBB0_8:                                #   in Loop: Header=BB0_3 Depth=2
	incl	%r8d
	decl	%edi
	cmpl	%r8d, %eax
	jne	.LBB0_3
# %bb.9:                                #   in Loop: Header=BB0_2 Depth=1
	incl	%edx
	decl	%ecx
	cmpl	%eax, %edx
	jne	.LBB0_2
.LBB0_10:                               # %._crit_edge
	retq
.Lfunc_end0:
	.size	_Z9listChainv, .Lfunc_end0-_Z9listChainv
	.cfi_endproc
                                        # -- End function
	.globl	_Z9placementPv                  # -- Begin function _Z9placementPv
	.type	_Z9placementPv,@function
_Z9placementPv:                         # @_Z9placementPv
	.cfi_startproc
# %bb.0:
	movq	1264(%rdi), %rsi
	movq	1272(%rdi), %rdx
	movb	$1, %al
	cmpq	%rdx, 40(%rdi,%rsi,8)
	je	.LBB1_1
# %bb.2:
	movq	40(%rdi), %rcx
	cmpq	$-1, %rcx
	je	.LBB1_15
# %bb.3:
	testq	%rcx, %rcx
	jne	.LBB1_7
# %bb.4:
	cmpq	$1, %rdx
	jne	.LBB1_15
# %bb.5:
	movq	48(%rdi), %rax
	cmpq	$-1, %rax
	je	.LBB1_15
# %bb.6:
	cmpq	%rsi, %rax
	jae	.LBB1_19
	jmp	.LBB1_15
.LBB1_7:
	cmpq	%rcx, %rsi
	jb	.LBB1_10
# %bb.8:
	movl	g(%rip), %r8d
	movq	%r8, %rax
	subq	%rcx, %rax
	cmpq	%rax, %rsi
	jae	.LBB1_10
# %bb.9:                                # %._crit_edge
	leal	-1(%r8), %eax
	jmp	.LBB1_12
.LBB1_10:
	testq	%rdx, %rdx
	je	.LBB1_19
# %bb.11:
	movl	g(%rip), %r8d
	leal	-1(%r8), %eax
	cmpq	%rax, %rdx
	je	.LBB1_19
.LBB1_12:
	cmpq	%rax, %rsi
	jne	.LBB1_15
# %bb.13:
	xorl	%eax, %eax
	cmpq	%rcx, %rdx
	jbe	.LBB1_1
# %bb.14:
	movl	%r8d, %r8d
	subq	%rcx, %r8
	cmpq	%r8, %rdx
	jae	.LBB1_1
.LBB1_15:
	movq	%rdx, 40(%rdi,%rsi,8)
	movl	$1, %eax
	movl	%esi, %ecx
	shlq	%cl, %rax
	movl	$1, %r8d
	movl	%edx, %ecx
	shlq	%cl, %r8
	movl	%edx, %ecx
	subl	%esi, %ecx
	addl	g(%rip), %ecx
	movl	$1, %r9d
	decb	%cl
	movl	$1, %r10d
                                        # kill: def $cl killed $cl killed $ecx
	shlq	%cl, %r10
	addl	%esi, %edx
	movl	%edx, %ecx
	shlq	%cl, %r9
	movq	8(%rdi), %rcx
	btq	%rsi, %rcx
	jb	.LBB1_19
# %bb.16:
	movq	16(%rdi), %rdx
	testq	%r8, %rdx
	jne	.LBB1_19
# %bb.17:
	movq	24(%rdi), %rsi
	testq	%r10, %rsi
	jne	.LBB1_19
# %bb.18:
	movq	32(%rdi), %r11
	testq	%r9, %r11
	je	.LBB1_20
.LBB1_19:
	xorl	%eax, %eax
.LBB1_1:
                                        # kill: def $al killed $al killed $eax
	retq
.LBB1_20:
	orq	%rax, %rcx
	movq	%rcx, 8(%rdi)
	orq	%r8, %rdx
	movq	%rdx, 16(%rdi)
	orq	%r10, %rsi
	movq	%rsi, 24(%rdi)
	orq	%r9, %r11
	movq	%r11, 32(%rdi)
	movb	$1, %al
                                        # kill: def $al killed $al killed $eax
	retq
.Lfunc_end1:
	.size	_Z9placementPv, .Lfunc_end1-_Z9placementPv
	.cfi_endproc
                                        # -- End function
	.globl	_Z5solveiimmmm                  # -- Begin function _Z5solveiimmmm
	.type	_Z5solveiimmmm,@function
_Z5solveiimmmm:                         # @_Z5solveiimmmm
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$1128, %rsp                     # imm = 0x468
	.cfi_def_cfa_offset 1184
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%r9, %r13
	movq	%r8, %r12
	movq	%rcx, %r15
	movq	%rdx, %r14
	movl	%esi, %ebx
	testl	%edi, %edi
	jle	.LBB2_2
# %bb.1:                                # %.lr.ph.preheader
	movl	%edi, %ebp
	shlq	$3, %rbp
	movl	%edi, 12(%rsp)                  # 4-byte Spill
	leaq	912(%rsp), %rdi
	xorl	%esi, %esi
	movq	%rbp, %rdx
	callq	memset@PLT
	leaq	464(%rsp), %rdi
	xorl	%esi, %esi
	movq	%rbp, %rdx
	callq	memset@PLT
	leaq	240(%rsp), %rdi
	xorl	%esi, %esi
	movq	%rbp, %rdx
	callq	memset@PLT
	leaq	688(%rsp), %rdi
	xorl	%esi, %esi
	movq	%rbp, %rdx
	callq	memset@PLT
	leaq	16(%rsp), %rdi
	xorl	%esi, %esi
	movq	%rbp, %rdx
	callq	memset@PLT
	movl	12(%rsp), %edi                  # 4-byte Reload
.LBB2_2:                                # %._crit_edge
	movslq	%ebx, %rax
	movq	%r14, 912(%rsp,%rax,8)
	movq	%r15, 464(%rsp,%rax,8)
	movq	%r12, 240(%rsp,%rax,8)
	movq	%r13, 688(%rsp,%rax,8)
	movq	%r12, %rcx
	orq	%r15, %rcx
	orq	%r13, %rcx
	notq	%rcx
	movq	%rcx, 16(%rsp,%rax,8)
	testl	%eax, %eax
	js	.LBB2_3
# %bb.4:                                # %.lr.ph123.lr.ph.preheader
	xorl	%eax, %eax
.LBB2_7:                                # %.lr.ph123
                                        # =>This Loop Header: Depth=1
                                        #     Child Loop BB2_5 Depth 2
                                        #       Child Loop BB2_10 Depth 3
	cmpl	%edi, %ebx
	jge	.LBB2_12
# %bb.8:                                # %.lr.ph123.split.preheader
                                        #   in Loop: Header=BB2_7 Depth=1
	movl	%r14d, %edx
	andl	$1, %edx
	movq	%rdx, %rsi
	orq	%rcx, %rsi
	je	.LBB2_12
.LBB2_5:                                # %.lr.ph
                                        #   Parent Loop BB2_7 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB2_10 Depth 3
	cmpq	$-1, %r12
	je	.LBB2_6
# %bb.9:                                #   in Loop: Header=BB2_5 Depth=2
	testb	$1, %dl
	je	.LBB2_15
.LBB2_10:                               # %.lr.ph115
                                        #   Parent Loop BB2_7 Depth=1
                                        #     Parent Loop BB2_5 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movq	%r14, %rcx
	shrq	%r14
	addq	%r15, %r15
	shrq	%r13
	testb	$2, %cl
	jne	.LBB2_10
# %bb.11:                               # %._crit_edge116
                                        #   in Loop: Header=BB2_5 Depth=2
	movq	%r15, %rcx
	orq	%r12, %rcx
	orq	%r13, %rcx
	xorl	%edx, %edx
	cmpq	$-1, %rcx
	notq	%rcx
	jne	.LBB2_5
.LBB2_12:                               # %.split
                                        #   in Loop: Header=BB2_7 Depth=1
	decl	%ebx
	movslq	%ebx, %rcx
	movq	912(%rsp,%rcx,8), %r14
	movq	464(%rsp,%rcx,8), %r15
	movq	688(%rsp,%rcx,8), %r13
	movq	240(%rsp,%rcx,8), %r12
	movq	16(%rsp,%rcx,8), %rcx
.LBB2_13:                               # %.outer
                                        #   in Loop: Header=BB2_7 Depth=1
	testl	%ebx, %ebx
	jns	.LBB2_7
	jmp	.LBB2_14
.LBB2_6:                                # %.outer.outer
                                        #   in Loop: Header=BB2_7 Depth=1
	incq	%rax
	decl	%ebx
	movslq	%ebx, %rdx
	movq	912(%rsp,%rdx,8), %r14
	movq	464(%rsp,%rdx,8), %r15
	movq	688(%rsp,%rdx,8), %r13
	movq	240(%rsp,%rdx,8), %r12
	movq	16(%rsp,%rdx,8), %rcx
	testl	%edx, %edx
	jns	.LBB2_7
	jmp	.LBB2_14
.LBB2_15:                               #   in Loop: Header=BB2_7 Depth=1
	movq	%rcx, %rdx
	negq	%rdx
	andq	%rcx, %rdx
	xorq	%rdx, %rcx
	movl	%ebx, %esi
	movq	%r14, 912(%rsp,%rsi,8)
	movq	%r15, 464(%rsp,%rsi,8)
	movq	%r12, 240(%rsp,%rsi,8)
	movq	%r13, 688(%rsp,%rsi,8)
	movq	%rcx, 16(%rsp,%rsi,8)
	incl	%ebx
	shrq	%r14
	orq	%rdx, %r15
	addq	%r15, %r15
	orq	%rdx, %r12
	orq	%r13, %rdx
	shrq	%rdx
	movq	%r15, %rcx
	orq	%r12, %rcx
	orq	%rdx, %rcx
	notq	%rcx
	movq	%rdx, %r13
	jmp	.LBB2_13
.LBB2_3:
	xorl	%eax, %eax
.LBB2_14:                               # %.outer._crit_edge
	addq	$1128, %rsp                     # imm = 0x468
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	_Z5solveiimmmm, .Lfunc_end2-_Z5solveiimmmm
	.cfi_endproc
                                        # -- End function
	.globl	_Z19carryChain_symmetryPv       # -- Begin function _Z19carryChain_symmetryPv
	.type	_Z19carryChain_symmetryPv,@function
_Z19carryChain_symmetryPv:              # @_Z19carryChain_symmetryPv
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	%rdi, %rbx
	movl	g(%rip), %edi
	leal	-2(%rdi), %eax
	leal	-1(%rdi), %esi
	imull	%eax, %esi
	decl	%esi
	movl	1256(%rbx), %ecx
	movl	1260(%rbx), %eax
	movl	%esi, %r9d
	subl	%eax, %r9d
	cmpl	%r9d, %ecx
	jne	.LBB3_1
# %bb.2:
	movl	1252(%rbx), %edx
	movl	%esi, %r8d
	subl	%edx, %r8d
	cmpl	%r8d, 1248(%rbx)
	jae	.LBB3_3
	jmp	.LBB3_18
.LBB3_1:                                # %._crit_edge
	movl	1252(%rbx), %edx
.LBB3_3:
	movl	1248(%rbx), %r8d
	cmpl	%r9d, %edx
	jne	.LBB3_5
# %bb.4:
	movl	%esi, %r10d
	subl	%r8d, %r10d
	cmpl	%r10d, %r8d
	ja	.LBB3_18
.LBB3_5:                                # %._crit_edge60
	subl	%ecx, %esi
	cmpl	%esi, %edx
	jbe	.LBB3_7
# %bb.6:                                # %._crit_edge60
	cmpl	%r9d, %r8d
	je	.LBB3_18
.LBB3_7:
	cmpq	$0, 40(%rbx)
	je	.LBB3_16
# %bb.8:
	cmpl	%eax, %ecx
	jne	.LBB3_12
# %bb.9:
	cmpl	%eax, %edx
	jne	.LBB3_18
# %bb.10:
	cmpl	%eax, %r8d
	jne	.LBB3_18
# %bb.11:
	movq	8(%rbx), %rdx
	movq	16(%rbx), %rsi
	shrq	$2, %rdx
	movq	24(%rbx), %rax
	shrq	$2, %rsi
	leal	-4(%rdi), %ecx
	movl	$-1, %r8d
                                        # kill: def $cl killed $cl killed $ecx
	shll	%cl, %r8d
	movslq	%r8d, %r8
	orq	%rsi, %r8
	incq	%r8
	leal	-5(%rdi), %ecx
	shlq	%cl, %r8
	shrq	$4, %rax
	decq	%r8
	movq	32(%rbx), %r9
	shrq	$4, %r9
                                        # kill: def $cl killed $cl killed $ecx
	shlq	%cl, %r9
                                        # kill: def $edi killed $edi killed $rdi
	xorl	%esi, %esi
	movq	%rax, %rcx
	callq	_Z5solveiimmmm
	leaq	1304(%rbx), %rcx
	jmp	.LBB3_17
.LBB3_12:
	cmpl	%eax, %edx
	jne	.LBB3_16
# %bb.13:
	cmpl	%ecx, %r8d
	jb	.LBB3_16
# %bb.14:
	ja	.LBB3_18
# %bb.15:
	movq	8(%rbx), %rdx
	movq	16(%rbx), %rsi
	shrq	$2, %rdx
	movq	24(%rbx), %rax
	shrq	$2, %rsi
	leal	-4(%rdi), %ecx
	movl	$-1, %r8d
                                        # kill: def $cl killed $cl killed $ecx
	shll	%cl, %r8d
	movslq	%r8d, %r8
	orq	%rsi, %r8
	incq	%r8
	leal	-5(%rdi), %ecx
	shlq	%cl, %r8
	shrq	$4, %rax
	decq	%r8
	movq	32(%rbx), %r9
	shrq	$4, %r9
                                        # kill: def $cl killed $cl killed $ecx
	shlq	%cl, %r9
                                        # kill: def $edi killed $edi killed $rdi
	xorl	%esi, %esi
	movq	%rax, %rcx
	callq	_Z5solveiimmmm
	leaq	1308(%rbx), %rcx
	jmp	.LBB3_17
.LBB3_16:
	movq	8(%rbx), %rdx
	movq	16(%rbx), %rsi
	shrq	$2, %rdx
	movq	24(%rbx), %rax
	shrq	$2, %rsi
	leal	-4(%rdi), %ecx
	movl	$-1, %r8d
                                        # kill: def $cl killed $cl killed $ecx
	shll	%cl, %r8d
	movslq	%r8d, %r8
	orq	%rsi, %r8
	incq	%r8
	leal	-5(%rdi), %ecx
	shlq	%cl, %r8
	shrq	$4, %rax
	decq	%r8
	movq	32(%rbx), %r9
	shrq	$4, %r9
                                        # kill: def $cl killed $cl killed $ecx
	shlq	%cl, %r9
                                        # kill: def $edi killed $edi killed $rdi
	xorl	%esi, %esi
	movq	%rax, %rcx
	callq	_Z5solveiimmmm
	leaq	1312(%rbx), %rcx
.LBB3_17:                               # %.sink.split
	movl	(%rcx), %ecx
	addq	%rax, 1280(%rbx,%rcx,8)
.LBB3_18:
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end3:
	.size	_Z19carryChain_symmetryPv, .Lfunc_end3-_Z19carryChain_symmetryPv
	.cfi_endproc
                                        # -- End function
	.globl	_Z10thread_runPv                # -- Begin function _Z10thread_runPv
	.type	_Z10thread_runPv,@function
_Z10thread_runPv:                       # @_Z10thread_runPv
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdi, %rbx
	leaq	1000(%rdi), %rsi
	leaq	8(%rdi), %r14
	movl	$248, %edx
	movq	%r14, %rdi
	callq	memcpy@PLT
	movq	$0, 1264(%rbx)
	movl	1260(%rbx), %eax
	movl	g+4(,%rax,4), %eax
	movq	%rax, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	testb	%al, %al
	je	.LBB4_20
# %bb.1:
	movq	$1, 1264(%rbx)
	movl	1260(%rbx), %eax
	movl	g+3724(,%rax,4), %eax
	movq	%rax, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	testb	%al, %al
	je	.LBB4_20
# %bb.2:
	leaq	256(%rbx), %rdi
	movl	$248, %edx
	movq	%rdi, 16(%rsp)                  # 8-byte Spill
	movq	%r14, %rsi
	callq	memcpy@PLT
	movl	1260(%rbx), %r15d
	movl	%r15d, 1248(%rbx)
	movl	g(%rip), %ebp
	leal	-2(%rbp), %eax
	decl	%ebp
	imull	%ebp, %eax
	subl	%r15d, %eax
	cmpl	%eax, %r15d
	jae	.LBB4_20
# %bb.3:                                # %.lr.ph71
	leaq	504(%rbx), %rax
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	leaq	752(%rbx), %r13
.LBB4_4:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_9 Depth 2
                                        #       Child Loop BB4_14 Depth 3
	movl	$248, %edx
	movq	%r14, %rdi
	movq	16(%rsp), %rsi                  # 8-byte Reload
	callq	memcpy@PLT
	movl	%r15d, %eax
	movl	g+4(,%rax,4), %eax
	movq	%rax, 1264(%rbx)
	movl	%ebp, %eax
	movq	%rax, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	movl	g(%rip), %ebp
	testb	%al, %al
	je	.LBB4_5
# %bb.6:                                #   in Loop: Header=BB4_4 Depth=1
	movl	1248(%rbx), %eax
	movl	g+3724(,%rax,4), %eax
	movq	%rax, 1264(%rbx)
	addl	$-2, %ebp
	movq	%rbp, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	movl	g(%rip), %ebp
	movl	1260(%rbx), %r15d
	testb	%al, %al
	je	.LBB4_7
# %bb.8:                                #   in Loop: Header=BB4_4 Depth=1
	movl	$248, %edx
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	%r14, %rsi
	callq	memcpy@PLT
	movl	%r15d, 1252(%rbx)
	leal	-2(%rbp), %eax
	leal	-1(%rbp), %r12d
	imull	%r12d, %eax
	subl	%r15d, %eax
	cmpl	%eax, %r15d
	jae	.LBB4_7
.LBB4_9:                                # %.lr.ph69
                                        #   Parent Loop BB4_4 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB4_14 Depth 3
	movl	$248, %edx
	movq	%r14, %rdi
	movq	8(%rsp), %rsi                   # 8-byte Reload
	callq	memcpy@PLT
	movl	%r12d, %eax
	movq	%rax, 1264(%rbx)
	movl	%r15d, %eax
	movl	g+4(,%rax,4), %eax
	notl	%eax
	addl	%ebp, %eax
	movq	%rax, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	movl	g(%rip), %ebp
	testb	%al, %al
	je	.LBB4_10
# %bb.11:                               #   in Loop: Header=BB4_9 Depth=2
	leal	-2(%rbp), %eax
	movq	%rax, 1264(%rbx)
	movl	1252(%rbx), %eax
	movl	g+3724(,%rax,4), %eax
	notl	%eax
	addl	%ebp, %eax
	movq	%rax, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	movl	g(%rip), %ebp
	movl	1260(%rbx), %r15d
	testb	%al, %al
	je	.LBB4_12
# %bb.13:                               #   in Loop: Header=BB4_9 Depth=2
	movl	$248, %edx
	movq	%r13, %rdi
	movq	%r14, %rsi
	callq	memcpy@PLT
	movl	%r15d, 1256(%rbx)
	leal	-2(%rbp), %eax
	leal	-1(%rbp), %ecx
	imull	%eax, %ecx
	subl	%r15d, %ecx
	cmpl	%ecx, %r15d
	jae	.LBB4_12
.LBB4_14:                               # %.lr.ph
                                        #   Parent Loop BB4_4 Depth=1
                                        #     Parent Loop BB4_9 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movl	$248, %edx
	movq	%r14, %rdi
	movq	%r13, %rsi
	callq	memcpy@PLT
	movl	%r15d, %eax
	movl	g+4(,%rax,4), %eax
	notl	%eax
	addl	%ebp, %eax
	movq	%rax, 1264(%rbx)
	movq	$0, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	testb	%al, %al
	je	.LBB4_17
# %bb.15:                               #   in Loop: Header=BB4_14 Depth=3
	movl	1256(%rbx), %eax
	movl	g+3724(,%rax,4), %eax
	notl	%eax
	addl	g(%rip), %eax
	movq	%rax, 1264(%rbx)
	movq	$1, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	testb	%al, %al
	je	.LBB4_17
# %bb.16:                               #   in Loop: Header=BB4_14 Depth=3
	movq	%rbx, %rdi
	callq	_Z19carryChain_symmetryPv
.LBB4_17:                               #   in Loop: Header=BB4_14 Depth=3
	movl	1256(%rbx), %r15d
	movl	1260(%rbx), %eax
	incl	%r15d
	movl	%r15d, 1256(%rbx)
	movl	g(%rip), %ebp
	leal	-2(%rbp), %ecx
	leal	-1(%rbp), %edx
	imull	%ecx, %edx
	subl	%eax, %edx
	cmpl	%edx, %r15d
	jb	.LBB4_14
	jmp	.LBB4_18
.LBB4_10:                               # %.lr.ph69..loopexit_crit_edge
                                        #   in Loop: Header=BB4_9 Depth=2
	movl	1260(%rbx), %eax
	jmp	.LBB4_18
.LBB4_12:                               #   in Loop: Header=BB4_9 Depth=2
	movl	%r15d, %eax
.LBB4_18:                               # %.loopexit
                                        #   in Loop: Header=BB4_9 Depth=2
	movl	1252(%rbx), %r15d
	incl	%r15d
	movl	%r15d, 1252(%rbx)
	leal	-2(%rbp), %ecx
	leal	-1(%rbp), %r12d
	imull	%r12d, %ecx
	subl	%eax, %ecx
	cmpl	%ecx, %r15d
	jb	.LBB4_9
	jmp	.LBB4_19
.LBB4_5:                                # %..loopexit65_crit_edge
                                        #   in Loop: Header=BB4_4 Depth=1
	movl	1260(%rbx), %eax
	jmp	.LBB4_19
.LBB4_7:                                #   in Loop: Header=BB4_4 Depth=1
	movl	%r15d, %eax
.LBB4_19:                               # %.loopexit65
                                        #   in Loop: Header=BB4_4 Depth=1
	movl	1248(%rbx), %r15d
	incl	%r15d
	movl	%r15d, 1248(%rbx)
	leal	-2(%rbp), %ecx
	decl	%ebp
	imull	%ebp, %ecx
	subl	%eax, %ecx
	cmpl	%ecx, %r15d
	jb	.LBB4_4
.LBB4_20:                               # %.loopexit66
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end4:
	.size	_Z10thread_runPv, .Lfunc_end4-_Z10thread_runPv
	.cfi_endproc
                                        # -- End function
	.globl	_Z10buildChainv                 # -- Begin function _Z10buildChainv
	.type	_Z10buildChainv,@function
_Z10buildChainv:                        # @_Z10buildChainv
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset %rbp, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register %rbp
	pushq	%r15
	pushq	%r14
	pushq	%rbx
	pushq	%rax
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	movl	g(%rip), %edx
	movl	%edx, %eax
	shrl	%eax
	leal	-3(%rdx), %r15d
	imull	%eax, %r15d
	imulq	$1320, %r15, %rax               # imm = 0x528
	addq	$15, %rax
	andq	$-16, %rax
	movq	%rsp, %rcx
	movq	%rcx, %rbx
	subq	%rax, %rbx
	movq	%rbx, %rsp
	leaq	8(%rbx), %r14
	negq	%rax
	movabsq	$4294967296, %rsi               # imm = 0x100000000
	movq	%rsi, 1304(%rbx)
	movl	$2, 1312(%rbx)
	xorps	%xmm0, %xmm0
	movups	%xmm0, 8(%rbx)
	movups	%xmm0, 24(%rbx)
	movaps	%xmm0, 1280(%rbx)
	movq	$0, 1296(%rcx,%rax)
	testq	%rdx, %rdx
	je	.LBB5_2
# %bb.1:                                # %.lr.ph.preheader
	leaq	40(%rbx), %rdi
	shlq	$3, %rdx
	movl	$255, %esi
	callq	memset@PLT
.LBB5_2:                                # %._crit_edge
	leaq	1000(%rbx), %rdi
	movl	$248, %edx
	movq	%r14, %rsi
	callq	memcpy@PLT
	movl	$0, 1260(%rbx)
.LBB5_3:                                # =>This Inner Loop Header: Depth=1
	movq	%rbx, %rdi
	callq	_Z10thread_runPv
	movl	1260(%rbx), %eax
	incl	%eax
	movl	%eax, 1260(%rbx)
	cmpl	%r15d, %eax
	jbe	.LBB5_3
# %bb.4:
	movl	1304(%rbx), %eax
	movl	1308(%rbx), %ecx
	movq	1280(%rbx,%rax,8), %rax
	movq	1280(%rbx,%rcx,8), %rcx
	leaq	(%rcx,%rax), %rdx
	movl	1312(%rbx), %esi
	movq	1280(%rbx,%rsi,8), %rsi
	addq	%rsi, %rdx
	movq	%rdx, UNIQUE(%rip)
	addq	%rax, %rax
	leaq	(%rax,%rcx,4), %rax
	leaq	(%rax,%rsi,8), %rax
	movq	%rax, TOTAL(%rip)
	leaq	-24(%rbp), %rsp
	popq	%rbx
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa %rsp, 8
	retq
.Lfunc_end5:
	.size	_Z10buildChainv, .Lfunc_end5-_Z10buildChainv
	.cfi_endproc
                                        # -- End function
	.globl	_Z10carryChainv                 # -- Begin function _Z10carryChainv
	.type	_Z10carryChainv,@function
_Z10carryChainv:                        # @_Z10carryChainv
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	_Z9listChainv
	popq	%rax
	.cfi_def_cfa_offset 8
	jmp	_Z10buildChainv                 # TAILCALL
.Lfunc_end6:
	.size	_Z10carryChainv, .Lfunc_end6-_Z10carryChainv
	.cfi_endproc
                                        # -- End function
	.globl	_Z6solveRmmmm                   # -- Begin function _Z6solveRmmmm
	.type	_Z6solveRmmmm,@function
_Z6solveRmmmm:                          # @_Z6solveRmmmm
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	cmpq	$-1, %rdx
	je	.LBB7_1
# %bb.2:                                # %.preheader
	movq	%rcx, %r15
	movq	%rdx, %rbx
	movq	%rsi, %rcx
	testb	$1, %dil
	jne	.LBB7_4
# %bb.3:
	movq	%rdi, %r12
	jmp	.LBB7_6
.LBB7_1:
	movl	$1, %r13d
	jmp	.LBB7_10
.LBB7_4:                                # %.lr.ph.preheader
	movq	%rdi, %r12
.LBB7_5:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	shrq	%r12
	addq	%rcx, %rcx
	shrq	%r15
	testb	$2, %dil
	movq	%r12, %rdi
	jne	.LBB7_5
.LBB7_6:                                # %._crit_edge
	movq	%rbx, %r14
	orq	%r15, %r14
	orq	%rcx, %r14
	cmpq	$-1, %r14
	je	.LBB7_7
# %bb.8:                                # %.lr.ph40.preheader
	shrq	%r12
	notq	%r14
	xorl	%r13d, %r13d
	movq	%rcx, (%rsp)                    # 8-byte Spill
.LBB7_9:                                # %.lr.ph40
                                        # =>This Inner Loop Header: Depth=1
	movq	%r14, %rbp
	negq	%rbp
	andq	%r14, %rbp
	movq	%rbp, %rsi
	orq	%rcx, %rsi
	addq	%rsi, %rsi
	movq	%rbp, %rdx
	orq	%rbx, %rdx
	movq	%rbp, %rcx
	orq	%r15, %rcx
	shrq	%rcx
	movq	%r12, %rdi
	callq	_Z6solveRmmmm
	movq	(%rsp), %rcx                    # 8-byte Reload
	addq	%rax, %r13
	xorq	%rbp, %r14
	jne	.LBB7_9
	jmp	.LBB7_10
.LBB7_7:
	xorl	%r13d, %r13d
.LBB7_10:                               # %.loopexit
	movq	%r13, %rax
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end7:
	.size	_Z6solveRmmmm, .Lfunc_end7-_Z6solveRmmmm
	.cfi_endproc
                                        # -- End function
	.globl	_Z20carryChain_symmetryRPv      # -- Begin function _Z20carryChain_symmetryRPv
	.type	_Z20carryChain_symmetryRPv,@function
_Z20carryChain_symmetryRPv:             # @_Z20carryChain_symmetryRPv
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	%rdi, %rbx
	movl	g(%rip), %eax
	leal	-2(%rax), %ecx
	leal	-1(%rax), %edi
	imull	%ecx, %edi
	decl	%edi
	movl	1256(%rbx), %edx
	movl	1260(%rbx), %ecx
	movl	%edi, %r9d
	subl	%ecx, %r9d
	cmpl	%r9d, %edx
	jne	.LBB8_1
# %bb.2:
	movl	1252(%rbx), %esi
	movl	%edi, %r8d
	subl	%esi, %r8d
	cmpl	%r8d, 1248(%rbx)
	jae	.LBB8_3
	jmp	.LBB8_18
.LBB8_1:                                # %._crit_edge
	movl	1252(%rbx), %esi
.LBB8_3:
	movl	1248(%rbx), %r8d
	cmpl	%r9d, %esi
	jne	.LBB8_5
# %bb.4:
	movl	%edi, %r10d
	subl	%r8d, %r10d
	cmpl	%r10d, %r8d
	ja	.LBB8_18
.LBB8_5:                                # %._crit_edge60
	subl	%edx, %edi
	cmpl	%edi, %esi
	jbe	.LBB8_7
# %bb.6:                                # %._crit_edge60
	cmpl	%r9d, %r8d
	je	.LBB8_18
.LBB8_7:
	cmpq	$0, 40(%rbx)
	je	.LBB8_16
# %bb.8:
	cmpl	%ecx, %edx
	jne	.LBB8_12
# %bb.9:
	cmpl	%ecx, %esi
	jne	.LBB8_18
# %bb.10:
	cmpl	%ecx, %r8d
	jne	.LBB8_18
# %bb.11:
	movq	8(%rbx), %rdi
	movq	16(%rbx), %r8
	shrq	$2, %rdi
	movq	24(%rbx), %rsi
	shrq	$2, %r8
	leal	-4(%rax), %ecx
	movl	$-1, %edx
                                        # kill: def $cl killed $cl killed $ecx
	shll	%cl, %edx
	movslq	%edx, %rdx
	orq	%r8, %rdx
	incq	%rdx
	addb	$-5, %al
	movl	%eax, %ecx
	shlq	%cl, %rdx
	shrq	$4, %rsi
	decq	%rdx
	movq	32(%rbx), %r8
	shrq	$4, %r8
	shlq	%cl, %r8
	movq	%r8, %rcx
	callq	_Z6solveRmmmm
	leaq	1304(%rbx), %rcx
	jmp	.LBB8_17
.LBB8_12:
	cmpl	%ecx, %esi
	jne	.LBB8_16
# %bb.13:
	cmpl	%edx, %r8d
	jb	.LBB8_16
# %bb.14:
	ja	.LBB8_18
# %bb.15:
	movq	8(%rbx), %rdi
	movq	16(%rbx), %r8
	shrq	$2, %rdi
	movq	24(%rbx), %rsi
	shrq	$2, %r8
	leal	-4(%rax), %ecx
	movl	$-1, %edx
                                        # kill: def $cl killed $cl killed $ecx
	shll	%cl, %edx
	movslq	%edx, %rdx
	orq	%r8, %rdx
	incq	%rdx
	addb	$-5, %al
	movl	%eax, %ecx
	shlq	%cl, %rdx
	shrq	$4, %rsi
	decq	%rdx
	movq	32(%rbx), %r8
	shrq	$4, %r8
	shlq	%cl, %r8
	movq	%r8, %rcx
	callq	_Z6solveRmmmm
	leaq	1308(%rbx), %rcx
	jmp	.LBB8_17
.LBB8_16:
	movq	8(%rbx), %rdi
	movq	16(%rbx), %r8
	shrq	$2, %rdi
	movq	24(%rbx), %rsi
	shrq	$2, %r8
	leal	-4(%rax), %ecx
	movl	$-1, %edx
                                        # kill: def $cl killed $cl killed $ecx
	shll	%cl, %edx
	movslq	%edx, %rdx
	orq	%r8, %rdx
	incq	%rdx
	addb	$-5, %al
	movl	%eax, %ecx
	shlq	%cl, %rdx
	shrq	$4, %rsi
	decq	%rdx
	movq	32(%rbx), %r8
	shrq	$4, %r8
	shlq	%cl, %r8
	movq	%r8, %rcx
	callq	_Z6solveRmmmm
	leaq	1312(%rbx), %rcx
.LBB8_17:                               # %.sink.split
	movl	(%rcx), %ecx
	addq	%rax, 1280(%rbx,%rcx,8)
.LBB8_18:
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end8:
	.size	_Z20carryChain_symmetryRPv, .Lfunc_end8-_Z20carryChain_symmetryRPv
	.cfi_endproc
                                        # -- End function
	.globl	_Z11thread_runRPv               # -- Begin function _Z11thread_runRPv
	.type	_Z11thread_runRPv,@function
_Z11thread_runRPv:                      # @_Z11thread_runRPv
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdi, %rbx
	leaq	1000(%rdi), %rsi
	leaq	8(%rdi), %r14
	movl	$248, %edx
	movq	%r14, %rdi
	callq	memcpy@PLT
	movq	$0, 1264(%rbx)
	movl	1260(%rbx), %eax
	movl	g+4(,%rax,4), %eax
	movq	%rax, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	testb	%al, %al
	je	.LBB9_20
# %bb.1:
	movq	$1, 1264(%rbx)
	movl	1260(%rbx), %eax
	movl	g+3724(,%rax,4), %eax
	movq	%rax, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	testb	%al, %al
	je	.LBB9_20
# %bb.2:
	leaq	256(%rbx), %rdi
	movl	$248, %edx
	movq	%rdi, 16(%rsp)                  # 8-byte Spill
	movq	%r14, %rsi
	callq	memcpy@PLT
	movl	1260(%rbx), %r15d
	movl	%r15d, 1248(%rbx)
	movl	g(%rip), %ebp
	leal	-2(%rbp), %eax
	decl	%ebp
	imull	%ebp, %eax
	subl	%r15d, %eax
	cmpl	%eax, %r15d
	jae	.LBB9_20
# %bb.3:                                # %.lr.ph71
	leaq	504(%rbx), %rax
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	leaq	752(%rbx), %r13
.LBB9_4:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB9_9 Depth 2
                                        #       Child Loop BB9_14 Depth 3
	movl	$248, %edx
	movq	%r14, %rdi
	movq	16(%rsp), %rsi                  # 8-byte Reload
	callq	memcpy@PLT
	movl	%r15d, %eax
	movl	g+4(,%rax,4), %eax
	movq	%rax, 1264(%rbx)
	movl	%ebp, %eax
	movq	%rax, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	movl	g(%rip), %ebp
	testb	%al, %al
	je	.LBB9_5
# %bb.6:                                #   in Loop: Header=BB9_4 Depth=1
	movl	1248(%rbx), %eax
	movl	g+3724(,%rax,4), %eax
	movq	%rax, 1264(%rbx)
	addl	$-2, %ebp
	movq	%rbp, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	movl	g(%rip), %ebp
	movl	1260(%rbx), %r15d
	testb	%al, %al
	je	.LBB9_7
# %bb.8:                                #   in Loop: Header=BB9_4 Depth=1
	movl	$248, %edx
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	%r14, %rsi
	callq	memcpy@PLT
	movl	%r15d, 1252(%rbx)
	leal	-2(%rbp), %eax
	leal	-1(%rbp), %r12d
	imull	%r12d, %eax
	subl	%r15d, %eax
	cmpl	%eax, %r15d
	jae	.LBB9_7
.LBB9_9:                                # %.lr.ph69
                                        #   Parent Loop BB9_4 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB9_14 Depth 3
	movl	$248, %edx
	movq	%r14, %rdi
	movq	8(%rsp), %rsi                   # 8-byte Reload
	callq	memcpy@PLT
	movl	%r12d, %eax
	movq	%rax, 1264(%rbx)
	movl	%r15d, %eax
	movl	g+4(,%rax,4), %eax
	notl	%eax
	addl	%ebp, %eax
	movq	%rax, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	movl	g(%rip), %ebp
	testb	%al, %al
	je	.LBB9_10
# %bb.11:                               #   in Loop: Header=BB9_9 Depth=2
	leal	-2(%rbp), %eax
	movq	%rax, 1264(%rbx)
	movl	1252(%rbx), %eax
	movl	g+3724(,%rax,4), %eax
	notl	%eax
	addl	%ebp, %eax
	movq	%rax, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	movl	g(%rip), %ebp
	movl	1260(%rbx), %r15d
	testb	%al, %al
	je	.LBB9_12
# %bb.13:                               #   in Loop: Header=BB9_9 Depth=2
	movl	$248, %edx
	movq	%r13, %rdi
	movq	%r14, %rsi
	callq	memcpy@PLT
	movl	%r15d, 1256(%rbx)
	leal	-2(%rbp), %eax
	leal	-1(%rbp), %ecx
	imull	%eax, %ecx
	subl	%r15d, %ecx
	cmpl	%ecx, %r15d
	jae	.LBB9_12
.LBB9_14:                               # %.lr.ph
                                        #   Parent Loop BB9_4 Depth=1
                                        #     Parent Loop BB9_9 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movl	$248, %edx
	movq	%r14, %rdi
	movq	%r13, %rsi
	callq	memcpy@PLT
	movl	%r15d, %eax
	movl	g+4(,%rax,4), %eax
	notl	%eax
	addl	%ebp, %eax
	movq	%rax, 1264(%rbx)
	movq	$0, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	testb	%al, %al
	je	.LBB9_17
# %bb.15:                               #   in Loop: Header=BB9_14 Depth=3
	movl	1256(%rbx), %eax
	movl	g+3724(,%rax,4), %eax
	notl	%eax
	addl	g(%rip), %eax
	movq	%rax, 1264(%rbx)
	movq	$1, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z9placementPv
	testb	%al, %al
	je	.LBB9_17
# %bb.16:                               #   in Loop: Header=BB9_14 Depth=3
	movq	%rbx, %rdi
	callq	_Z20carryChain_symmetryRPv
.LBB9_17:                               #   in Loop: Header=BB9_14 Depth=3
	movl	1256(%rbx), %r15d
	movl	1260(%rbx), %eax
	incl	%r15d
	movl	%r15d, 1256(%rbx)
	movl	g(%rip), %ebp
	leal	-2(%rbp), %ecx
	leal	-1(%rbp), %edx
	imull	%ecx, %edx
	subl	%eax, %edx
	cmpl	%edx, %r15d
	jb	.LBB9_14
	jmp	.LBB9_18
.LBB9_10:                               # %.lr.ph69..loopexit_crit_edge
                                        #   in Loop: Header=BB9_9 Depth=2
	movl	1260(%rbx), %eax
	jmp	.LBB9_18
.LBB9_12:                               #   in Loop: Header=BB9_9 Depth=2
	movl	%r15d, %eax
.LBB9_18:                               # %.loopexit
                                        #   in Loop: Header=BB9_9 Depth=2
	movl	1252(%rbx), %r15d
	incl	%r15d
	movl	%r15d, 1252(%rbx)
	leal	-2(%rbp), %ecx
	leal	-1(%rbp), %r12d
	imull	%r12d, %ecx
	subl	%eax, %ecx
	cmpl	%ecx, %r15d
	jb	.LBB9_9
	jmp	.LBB9_19
.LBB9_5:                                # %..loopexit65_crit_edge
                                        #   in Loop: Header=BB9_4 Depth=1
	movl	1260(%rbx), %eax
	jmp	.LBB9_19
.LBB9_7:                                #   in Loop: Header=BB9_4 Depth=1
	movl	%r15d, %eax
.LBB9_19:                               # %.loopexit65
                                        #   in Loop: Header=BB9_4 Depth=1
	movl	1248(%rbx), %r15d
	incl	%r15d
	movl	%r15d, 1248(%rbx)
	leal	-2(%rbp), %ecx
	decl	%ebp
	imull	%ebp, %ecx
	subl	%eax, %ecx
	cmpl	%ecx, %r15d
	jb	.LBB9_4
.LBB9_20:                               # %.loopexit66
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end9:
	.size	_Z11thread_runRPv, .Lfunc_end9-_Z11thread_runRPv
	.cfi_endproc
                                        # -- End function
	.globl	_Z11buildChainRv                # -- Begin function _Z11buildChainRv
	.type	_Z11buildChainRv,@function
_Z11buildChainRv:                       # @_Z11buildChainRv
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset %rbp, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register %rbp
	pushq	%r15
	pushq	%r14
	pushq	%rbx
	pushq	%rax
	.cfi_offset %rbx, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	movl	g(%rip), %edx
	movl	%edx, %eax
	shrl	%eax
	leal	-3(%rdx), %r15d
	imull	%eax, %r15d
	imulq	$1320, %r15, %rax               # imm = 0x528
	addq	$15, %rax
	andq	$-16, %rax
	movq	%rsp, %rcx
	movq	%rcx, %rbx
	subq	%rax, %rbx
	movq	%rbx, %rsp
	leaq	8(%rbx), %r14
	negq	%rax
	movabsq	$4294967296, %rsi               # imm = 0x100000000
	movq	%rsi, 1304(%rbx)
	movl	$2, 1312(%rbx)
	xorps	%xmm0, %xmm0
	movups	%xmm0, 8(%rbx)
	movups	%xmm0, 24(%rbx)
	movaps	%xmm0, 1280(%rbx)
	movq	$0, 1296(%rcx,%rax)
	testq	%rdx, %rdx
	je	.LBB10_2
# %bb.1:                                # %.lr.ph.preheader
	leaq	40(%rbx), %rdi
	shlq	$3, %rdx
	movl	$255, %esi
	callq	memset@PLT
.LBB10_2:                               # %._crit_edge
	leaq	1000(%rbx), %rdi
	movl	$248, %edx
	movq	%r14, %rsi
	callq	memcpy@PLT
	movl	$0, 1260(%rbx)
.LBB10_3:                               # =>This Inner Loop Header: Depth=1
	movq	%rbx, %rdi
	callq	_Z11thread_runRPv
	movl	1260(%rbx), %eax
	incl	%eax
	movl	%eax, 1260(%rbx)
	cmpl	%r15d, %eax
	jbe	.LBB10_3
# %bb.4:
	movl	1304(%rbx), %eax
	movl	1308(%rbx), %ecx
	movq	1280(%rbx,%rax,8), %rax
	movq	1280(%rbx,%rcx,8), %rcx
	leaq	(%rcx,%rax), %rdx
	movl	1312(%rbx), %esi
	movq	1280(%rbx,%rsi,8), %rsi
	addq	%rsi, %rdx
	movq	%rdx, UNIQUE(%rip)
	addq	%rax, %rax
	leaq	(%rax,%rcx,4), %rax
	leaq	(%rax,%rsi,8), %rax
	movq	%rax, TOTAL(%rip)
	leaq	-24(%rbp), %rsp
	popq	%rbx
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa %rsp, 8
	retq
.Lfunc_end10:
	.size	_Z11buildChainRv, .Lfunc_end10-_Z11buildChainRv
	.cfi_endproc
                                        # -- End function
	.globl	_Z11carryChainRv                # -- Begin function _Z11carryChainRv
	.type	_Z11carryChainRv,@function
_Z11carryChainRv:                       # @_Z11carryChainRv
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	_Z9listChainv
	popq	%rax
	.cfi_def_cfa_offset 8
	jmp	_Z11buildChainRv                # TAILCALL
.Lfunc_end11:
	.size	_Z11carryChainRv, .Lfunc_end11-_Z11carryChainRv
	.cfi_endproc
                                        # -- End function
	.globl	_Z13GPU_placementPv             # -- Begin function _Z13GPU_placementPv
	.type	_Z13GPU_placementPv,@function
_Z13GPU_placementPv:                    # @_Z13GPU_placementPv
	.cfi_startproc
# %bb.0:
	movq	1264(%rdi), %rsi
	movq	1272(%rdi), %rdx
	movb	$1, %al
	cmpq	%rdx, 40(%rdi,%rsi,8)
	je	.LBB12_1
# %bb.2:
	movq	40(%rdi), %rcx
	cmpq	$-1, %rcx
	je	.LBB12_15
# %bb.3:
	testq	%rcx, %rcx
	jne	.LBB12_7
# %bb.4:
	cmpq	$1, %rdx
	jne	.LBB12_15
# %bb.5:
	movq	48(%rdi), %rax
	cmpq	$-1, %rax
	je	.LBB12_15
# %bb.6:
	cmpq	%rsi, %rax
	jae	.LBB12_19
	jmp	.LBB12_15
.LBB12_7:
	cmpq	%rcx, %rsi
	jb	.LBB12_10
# %bb.8:
	movl	g(%rip), %r8d
	movq	%r8, %rax
	subq	%rcx, %rax
	cmpq	%rax, %rsi
	jae	.LBB12_10
# %bb.9:                                # %._crit_edge
	leal	-1(%r8), %eax
	jmp	.LBB12_12
.LBB12_10:
	testq	%rdx, %rdx
	je	.LBB12_19
# %bb.11:
	movl	g(%rip), %r8d
	leal	-1(%r8), %eax
	cmpq	%rax, %rdx
	je	.LBB12_19
.LBB12_12:
	cmpq	%rax, %rsi
	jne	.LBB12_15
# %bb.13:
	xorl	%eax, %eax
	cmpq	%rcx, %rdx
	jbe	.LBB12_1
# %bb.14:
	movl	%r8d, %r8d
	subq	%rcx, %r8
	cmpq	%r8, %rdx
	jae	.LBB12_1
.LBB12_15:
	movq	%rdx, 40(%rdi,%rsi,8)
	movl	$1, %eax
	movl	%esi, %ecx
	shlq	%cl, %rax
	movl	$1, %r8d
	movl	%edx, %ecx
	shlq	%cl, %r8
	movl	%edx, %ecx
	subl	%esi, %ecx
	addl	g(%rip), %ecx
	movl	$1, %r9d
	decb	%cl
	movl	$1, %r10d
                                        # kill: def $cl killed $cl killed $ecx
	shlq	%cl, %r10
	addl	%esi, %edx
	movl	%edx, %ecx
	shlq	%cl, %r9
	movq	8(%rdi), %rcx
	btq	%rsi, %rcx
	jb	.LBB12_19
# %bb.16:
	movq	16(%rdi), %rdx
	testq	%r8, %rdx
	jne	.LBB12_19
# %bb.17:
	movq	24(%rdi), %rsi
	testq	%r10, %rsi
	jne	.LBB12_19
# %bb.18:
	movq	32(%rdi), %r11
	testq	%r9, %r11
	je	.LBB12_20
.LBB12_19:
	xorl	%eax, %eax
.LBB12_1:
                                        # kill: def $al killed $al killed $eax
	retq
.LBB12_20:
	orq	%rax, %rcx
	movq	%rcx, 8(%rdi)
	orq	%r8, %rdx
	movq	%rdx, 16(%rdi)
	orq	%r10, %rsi
	movq	%rsi, 24(%rdi)
	orq	%r9, %r11
	movq	%r11, 32(%rdi)
	movb	$1, %al
                                        # kill: def $al killed $al killed $eax
	retq
.Lfunc_end12:
	.size	_Z13GPU_placementPv, .Lfunc_end12-_Z13GPU_placementPv
	.cfi_endproc
                                        # -- End function
	.globl	_Z10GPU_solveRmmmm              # -- Begin function _Z10GPU_solveRmmmm
	.type	_Z10GPU_solveRmmmm,@function
_Z10GPU_solveRmmmm:                     # @_Z10GPU_solveRmmmm
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	pushq	%rax
	.cfi_def_cfa_offset 64
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	cmpq	$-1, %rdx
	je	.LBB13_1
# %bb.2:                                # %.preheader
	movq	%rcx, %r15
	movq	%rdx, %rbx
	movq	%rsi, %rcx
	testb	$1, %dil
	jne	.LBB13_4
# %bb.3:
	movq	%rdi, %r12
	jmp	.LBB13_6
.LBB13_1:
	movl	$1, %r13d
	jmp	.LBB13_10
.LBB13_4:                               # %.lr.ph.preheader
	movq	%rdi, %r12
.LBB13_5:                               # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	shrq	%r12
	addq	%rcx, %rcx
	shrq	%r15
	testb	$2, %dil
	movq	%r12, %rdi
	jne	.LBB13_5
.LBB13_6:                               # %._crit_edge
	movq	%rbx, %r14
	orq	%r15, %r14
	orq	%rcx, %r14
	cmpq	$-1, %r14
	je	.LBB13_7
# %bb.8:                                # %.lr.ph40.preheader
	shrq	%r12
	notq	%r14
	xorl	%r13d, %r13d
	movq	%rcx, (%rsp)                    # 8-byte Spill
.LBB13_9:                               # %.lr.ph40
                                        # =>This Inner Loop Header: Depth=1
	movq	%r14, %rbp
	negq	%rbp
	andq	%r14, %rbp
	movq	%rbp, %rsi
	orq	%rcx, %rsi
	addq	%rsi, %rsi
	movq	%rbp, %rdx
	orq	%rbx, %rdx
	movq	%rbp, %rcx
	orq	%r15, %rcx
	shrq	%rcx
	movq	%r12, %rdi
	callq	_Z6solveRmmmm
	movq	(%rsp), %rcx                    # 8-byte Reload
	addq	%rax, %r13
	xorq	%rbp, %r14
	jne	.LBB13_9
	jmp	.LBB13_10
.LBB13_7:
	xorl	%r13d, %r13d
.LBB13_10:                              # %.loopexit
	movq	%r13, %rax
	addq	$8, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end13:
	.size	_Z10GPU_solveRmmmm, .Lfunc_end13-_Z10GPU_solveRmmmm
	.cfi_endproc
                                        # -- End function
	.globl	_Z24GPU_carryChain_symmetryRPv  # -- Begin function _Z24GPU_carryChain_symmetryRPv
	.type	_Z24GPU_carryChain_symmetryRPv,@function
_Z24GPU_carryChain_symmetryRPv:         # @_Z24GPU_carryChain_symmetryRPv
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset %rbx, -16
	movq	%rdi, %rbx
	movl	g(%rip), %eax
	leal	-2(%rax), %ecx
	leal	-1(%rax), %edi
	imull	%ecx, %edi
	decl	%edi
	movl	1256(%rbx), %edx
	movl	1260(%rbx), %ecx
	movl	%edi, %r9d
	subl	%ecx, %r9d
	cmpl	%r9d, %edx
	jne	.LBB14_1
# %bb.2:
	movl	1252(%rbx), %esi
	movl	%edi, %r8d
	subl	%esi, %r8d
	cmpl	%r8d, 1248(%rbx)
	jae	.LBB14_3
	jmp	.LBB14_18
.LBB14_1:                               # %._crit_edge
	movl	1252(%rbx), %esi
.LBB14_3:
	movl	1248(%rbx), %r8d
	cmpl	%r9d, %esi
	jne	.LBB14_5
# %bb.4:
	movl	%edi, %r10d
	subl	%r8d, %r10d
	cmpl	%r10d, %r8d
	ja	.LBB14_18
.LBB14_5:                               # %._crit_edge60
	subl	%edx, %edi
	cmpl	%edi, %esi
	jbe	.LBB14_7
# %bb.6:                                # %._crit_edge60
	cmpl	%r9d, %r8d
	je	.LBB14_18
.LBB14_7:
	cmpq	$0, 40(%rbx)
	je	.LBB14_16
# %bb.8:
	cmpl	%ecx, %edx
	jne	.LBB14_12
# %bb.9:
	cmpl	%ecx, %esi
	jne	.LBB14_18
# %bb.10:
	cmpl	%ecx, %r8d
	jne	.LBB14_18
# %bb.11:
	movq	8(%rbx), %rdi
	movq	16(%rbx), %r8
	shrq	$2, %rdi
	movq	24(%rbx), %rsi
	shrq	$2, %r8
	leal	-4(%rax), %ecx
	movl	$-1, %edx
                                        # kill: def $cl killed $cl killed $ecx
	shll	%cl, %edx
	movslq	%edx, %rdx
	orq	%r8, %rdx
	incq	%rdx
	addb	$-5, %al
	movl	%eax, %ecx
	shlq	%cl, %rdx
	shrq	$4, %rsi
	decq	%rdx
	movq	32(%rbx), %r8
	shrq	$4, %r8
	shlq	%cl, %r8
	movq	%r8, %rcx
	callq	_Z6solveRmmmm
	leaq	1304(%rbx), %rcx
	jmp	.LBB14_17
.LBB14_12:
	cmpl	%ecx, %esi
	jne	.LBB14_16
# %bb.13:
	cmpl	%edx, %r8d
	jb	.LBB14_16
# %bb.14:
	ja	.LBB14_18
# %bb.15:
	movq	8(%rbx), %rdi
	movq	16(%rbx), %r8
	shrq	$2, %rdi
	movq	24(%rbx), %rsi
	shrq	$2, %r8
	leal	-4(%rax), %ecx
	movl	$-1, %edx
                                        # kill: def $cl killed $cl killed $ecx
	shll	%cl, %edx
	movslq	%edx, %rdx
	orq	%r8, %rdx
	incq	%rdx
	addb	$-5, %al
	movl	%eax, %ecx
	shlq	%cl, %rdx
	shrq	$4, %rsi
	decq	%rdx
	movq	32(%rbx), %r8
	shrq	$4, %r8
	shlq	%cl, %r8
	movq	%r8, %rcx
	callq	_Z6solveRmmmm
	leaq	1308(%rbx), %rcx
	jmp	.LBB14_17
.LBB14_16:
	movq	8(%rbx), %rdi
	movq	16(%rbx), %r8
	shrq	$2, %rdi
	movq	24(%rbx), %rsi
	shrq	$2, %r8
	leal	-4(%rax), %ecx
	movl	$-1, %edx
                                        # kill: def $cl killed $cl killed $ecx
	shll	%cl, %edx
	movslq	%edx, %rdx
	orq	%r8, %rdx
	incq	%rdx
	addb	$-5, %al
	movl	%eax, %ecx
	shlq	%cl, %rdx
	shrq	$4, %rsi
	decq	%rdx
	movq	32(%rbx), %r8
	shrq	$4, %r8
	shlq	%cl, %r8
	movq	%r8, %rcx
	callq	_Z6solveRmmmm
	leaq	1312(%rbx), %rcx
.LBB14_17:                              # %.sink.split
	movl	(%rcx), %ecx
	addq	%rax, 1280(%rbx,%rcx,8)
.LBB14_18:
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end14:
	.size	_Z24GPU_carryChain_symmetryRPv, .Lfunc_end14-_Z24GPU_carryChain_symmetryRPv
	.cfi_endproc
                                        # -- End function
	.globl	_Z15GPU_thread_runRPv           # -- Begin function _Z15GPU_thread_runRPv
	.type	_Z15GPU_thread_runRPv,@function
_Z15GPU_thread_runRPv:                  # @_Z15GPU_thread_runRPv
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$24, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	movq	%rdi, %rbx
	leaq	1000(%rdi), %rsi
	leaq	8(%rdi), %r14
	movl	$248, %edx
	movq	%r14, %rdi
	callq	memcpy@PLT
	movq	$0, 1264(%rbx)
	movl	1260(%rbx), %eax
	movl	g+4(,%rax,4), %eax
	movq	%rax, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z13GPU_placementPv
	testb	%al, %al
	je	.LBB15_20
# %bb.1:
	movq	$1, 1264(%rbx)
	movl	1260(%rbx), %eax
	movl	g+3724(,%rax,4), %eax
	movq	%rax, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z13GPU_placementPv
	testb	%al, %al
	je	.LBB15_20
# %bb.2:
	leaq	256(%rbx), %rdi
	movl	$248, %edx
	movq	%rdi, 16(%rsp)                  # 8-byte Spill
	movq	%r14, %rsi
	callq	memcpy@PLT
	movl	1260(%rbx), %r15d
	movl	%r15d, 1248(%rbx)
	movl	g(%rip), %ebp
	leal	-2(%rbp), %eax
	decl	%ebp
	imull	%ebp, %eax
	subl	%r15d, %eax
	cmpl	%eax, %r15d
	jae	.LBB15_20
# %bb.3:                                # %.lr.ph71
	leaq	504(%rbx), %rax
	movq	%rax, 8(%rsp)                   # 8-byte Spill
	leaq	752(%rbx), %r13
.LBB15_4:                               # =>This Loop Header: Depth=1
                                        #     Child Loop BB15_9 Depth 2
                                        #       Child Loop BB15_14 Depth 3
	movl	$248, %edx
	movq	%r14, %rdi
	movq	16(%rsp), %rsi                  # 8-byte Reload
	callq	memcpy@PLT
	movl	%r15d, %eax
	movl	g+4(,%rax,4), %eax
	movq	%rax, 1264(%rbx)
	movl	%ebp, %eax
	movq	%rax, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z13GPU_placementPv
	movl	g(%rip), %ebp
	testb	%al, %al
	je	.LBB15_5
# %bb.6:                                #   in Loop: Header=BB15_4 Depth=1
	movl	1248(%rbx), %eax
	movl	g+3724(,%rax,4), %eax
	movq	%rax, 1264(%rbx)
	addl	$-2, %ebp
	movq	%rbp, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z13GPU_placementPv
	movl	g(%rip), %ebp
	movl	1260(%rbx), %r15d
	testb	%al, %al
	je	.LBB15_7
# %bb.8:                                #   in Loop: Header=BB15_4 Depth=1
	movl	$248, %edx
	movq	8(%rsp), %rdi                   # 8-byte Reload
	movq	%r14, %rsi
	callq	memcpy@PLT
	movl	%r15d, 1252(%rbx)
	leal	-2(%rbp), %eax
	leal	-1(%rbp), %r12d
	imull	%r12d, %eax
	subl	%r15d, %eax
	cmpl	%eax, %r15d
	jae	.LBB15_7
.LBB15_9:                               # %.lr.ph69
                                        #   Parent Loop BB15_4 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB15_14 Depth 3
	movl	$248, %edx
	movq	%r14, %rdi
	movq	8(%rsp), %rsi                   # 8-byte Reload
	callq	memcpy@PLT
	movl	%r12d, %eax
	movq	%rax, 1264(%rbx)
	movl	%r15d, %eax
	movl	g+4(,%rax,4), %eax
	notl	%eax
	addl	%ebp, %eax
	movq	%rax, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z13GPU_placementPv
	movl	g(%rip), %ebp
	testb	%al, %al
	je	.LBB15_10
# %bb.11:                               #   in Loop: Header=BB15_9 Depth=2
	leal	-2(%rbp), %eax
	movq	%rax, 1264(%rbx)
	movl	1252(%rbx), %eax
	movl	g+3724(,%rax,4), %eax
	notl	%eax
	addl	%ebp, %eax
	movq	%rax, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z13GPU_placementPv
	movl	g(%rip), %ebp
	movl	1260(%rbx), %r15d
	testb	%al, %al
	je	.LBB15_12
# %bb.13:                               #   in Loop: Header=BB15_9 Depth=2
	movl	$248, %edx
	movq	%r13, %rdi
	movq	%r14, %rsi
	callq	memcpy@PLT
	movl	%r15d, 1256(%rbx)
	leal	-2(%rbp), %eax
	leal	-1(%rbp), %ecx
	imull	%eax, %ecx
	subl	%r15d, %ecx
	cmpl	%ecx, %r15d
	jae	.LBB15_12
.LBB15_14:                              # %.lr.ph
                                        #   Parent Loop BB15_4 Depth=1
                                        #     Parent Loop BB15_9 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	movl	$248, %edx
	movq	%r14, %rdi
	movq	%r13, %rsi
	callq	memcpy@PLT
	movl	%r15d, %eax
	movl	g+4(,%rax,4), %eax
	notl	%eax
	addl	%ebp, %eax
	movq	%rax, 1264(%rbx)
	movq	$0, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z13GPU_placementPv
	testb	%al, %al
	je	.LBB15_17
# %bb.15:                               #   in Loop: Header=BB15_14 Depth=3
	movl	1256(%rbx), %eax
	movl	g+3724(,%rax,4), %eax
	notl	%eax
	addl	g(%rip), %eax
	movq	%rax, 1264(%rbx)
	movq	$1, 1272(%rbx)
	movq	%rbx, %rdi
	callq	_Z13GPU_placementPv
	testb	%al, %al
	je	.LBB15_17
# %bb.16:                               #   in Loop: Header=BB15_14 Depth=3
	movq	%rbx, %rdi
	callq	_Z20carryChain_symmetryRPv
.LBB15_17:                              #   in Loop: Header=BB15_14 Depth=3
	movl	1256(%rbx), %r15d
	movl	1260(%rbx), %eax
	incl	%r15d
	movl	%r15d, 1256(%rbx)
	movl	g(%rip), %ebp
	leal	-2(%rbp), %ecx
	leal	-1(%rbp), %edx
	imull	%ecx, %edx
	subl	%eax, %edx
	cmpl	%edx, %r15d
	jb	.LBB15_14
	jmp	.LBB15_18
.LBB15_10:                              # %.lr.ph69..loopexit_crit_edge
                                        #   in Loop: Header=BB15_9 Depth=2
	movl	1260(%rbx), %eax
	jmp	.LBB15_18
.LBB15_12:                              #   in Loop: Header=BB15_9 Depth=2
	movl	%r15d, %eax
.LBB15_18:                              # %.loopexit
                                        #   in Loop: Header=BB15_9 Depth=2
	movl	1252(%rbx), %r15d
	incl	%r15d
	movl	%r15d, 1252(%rbx)
	leal	-2(%rbp), %ecx
	leal	-1(%rbp), %r12d
	imull	%r12d, %ecx
	subl	%eax, %ecx
	cmpl	%ecx, %r15d
	jb	.LBB15_9
	jmp	.LBB15_19
.LBB15_5:                               # %..loopexit65_crit_edge
                                        #   in Loop: Header=BB15_4 Depth=1
	movl	1260(%rbx), %eax
	jmp	.LBB15_19
.LBB15_7:                               #   in Loop: Header=BB15_4 Depth=1
	movl	%r15d, %eax
.LBB15_19:                              # %.loopexit65
                                        #   in Loop: Header=BB15_4 Depth=1
	movl	1248(%rbx), %r15d
	incl	%r15d
	movl	%r15d, 1248(%rbx)
	leal	-2(%rbp), %ecx
	decl	%ebp
	imull	%ebp, %ecx
	subl	%eax, %ecx
	cmpl	%ecx, %r15d
	jb	.LBB15_4
.LBB15_20:                              # %.loopexit66
	addq	$24, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end15:
	.size	_Z15GPU_thread_runRPv, .Lfunc_end15-_Z15GPU_thread_runRPv
	.cfi_endproc
                                        # -- End function
	.globl	_Z15GPU_buildChainRjj           # -- Begin function _Z15GPU_buildChainRjj
	.type	_Z15GPU_buildChainRjj,@function
_Z15GPU_buildChainRjj:                  # @_Z15GPU_buildChainRjj
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset %rbp, -16
	movq	%rsp, %rbp
	.cfi_def_cfa_register %rbp
	pushq	%r15
	pushq	%r14
	pushq	%r12
	pushq	%rbx
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	movl	g(%rip), %eax
	movl	%eax, %ecx
	shrl	%ecx
	addl	$-3, %eax
	imull	%ecx, %eax
	imulq	$1320, %rax, %rax               # imm = 0x528
	addq	$15, %rax
	andq	$-16, %rax
	movq	%rsp, %rbx
	subq	%rax, %rbx
	movq	%rbx, %rsp
	movl	%esi, 1316(%rbx)
	movl	%edi, (%rbx)
	movq	%rsp, %r14
	subq	%rax, %r14
	movq	%r14, %rsp
	movl	%esi, %eax
	imulq	$1320, %rax, %rsi               # imm = 0x528
	movq	%rbx, %rdi
	xorl	%edx, %edx
	callq	hipHostMalloc
	movl	1316(%rbx), %eax
	imulq	$1320, %rax, %rsi               # imm = 0x528
	movq	%r14, %rdi
	callq	hipMalloc
	movabsq	$4294967296, %rax               # imm = 0x100000000
	movq	%rax, 1304(%rbx)
	movl	$2, 1312(%rbx)
	leaq	8(%rbx), %r15
	xorps	%xmm0, %xmm0
	movups	%xmm0, 8(%rbx)
	movups	%xmm0, 24(%rbx)
	movaps	%xmm0, 1280(%rbx)
	movq	$0, 1296(%rbx)
	movl	g(%rip), %r12d
	testq	%r12, %r12
	je	.LBB16_2
# %bb.1:                                # %.lr.ph.preheader
	leaq	40(%rbx), %rdi
	leaq	(,%r12,8), %rdx
	movl	$255, %esi
	callq	memset@PLT
.LBB16_2:                               # %._crit_edge
	leaq	1000(%rbx), %rdi
	movl	$248, %edx
	movq	%r15, %rsi
	callq	memcpy@PLT
	movl	%r12d, %eax
	shrl	%eax
	addl	$-3, %r12d
	imull	%eax, %r12d
	imulq	$1320, %r12, %r15               # imm = 0x528
	movq	%r14, %rdi
	movq	%rbx, %rsi
	movq	%r15, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movl	g(%rip), %r12d
	movl	%r12d, %eax
	shrl	%eax
	addl	$-3, %r12d
	imull	%eax, %r12d
	movl	$0, 1260(%rbx)
.LBB16_3:                               # =>This Inner Loop Header: Depth=1
	movq	%rbx, %rdi
	callq	_Z11thread_runRPv
	movl	1260(%rbx), %eax
	incl	%eax
	movl	%eax, 1260(%rbx)
	cmpl	%r12d, %eax
	jbe	.LBB16_3
# %bb.4:
	movq	%rbx, %rdi
	movq	%r14, %rsi
	movq	%r15, %rdx
	movl	$2, %ecx
	callq	hipMemcpy
	movl	1304(%rbx), %eax
	movl	1308(%rbx), %ecx
	movq	1280(%rbx,%rax,8), %rax
	movq	1280(%rbx,%rcx,8), %rcx
	leaq	(%rcx,%rax), %rdx
	movl	1312(%rbx), %esi
	movq	1280(%rbx,%rsi,8), %rsi
	addq	%rsi, %rdx
	movq	%rdx, UNIQUE(%rip)
	addq	%rax, %rax
	leaq	(%rax,%rcx,4), %rax
	leaq	(%rax,%rsi,8), %rax
	movq	%rax, TOTAL(%rip)
	leaq	-32(%rbp), %rsp
	popq	%rbx
	popq	%r12
	popq	%r14
	popq	%r15
	popq	%rbp
	.cfi_def_cfa %rsp, 8
	retq
.Lfunc_end16:
	.size	_Z15GPU_buildChainRjj, .Lfunc_end16-_Z15GPU_buildChainRjj
	.cfi_endproc
                                        # -- End function
	.globl	_Z15GPU_carryChainRjj           # -- Begin function _Z15GPU_carryChainRjj
	.type	_Z15GPU_carryChainRjj,@function
_Z15GPU_carryChainRjj:                  # @_Z15GPU_carryChainRjj
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %rbp, -16
	movl	%esi, %ebx
	movl	%edi, %ebp
	callq	_Z9listChainv
	movl	%ebp, %edi
	movl	%ebx, %esi
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	jmp	_Z15GPU_buildChainRjj           # TAILCALL
.Lfunc_end17:
	.size	_Z15GPU_carryChainRjj, .Lfunc_end17-_Z15GPU_carryChainRjj
	.cfi_endproc
                                        # -- End function
	.globl	_Z8InitCUDAv                    # -- Begin function _Z8InitCUDAv
	.type	_Z8InitCUDAv,@function
_Z8InitCUDAv:                           # @_Z8InitCUDAv
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	subq	$1480, %rsp                     # imm = 0x5C8
	.cfi_def_cfa_offset 1504
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	leaq	4(%rsp), %rbx
	movq	%rbx, %rdi
	callq	hipGetDeviceCount
	movl	(%rbx), %ecx
	testl	%ecx, %ecx
	je	.LBB18_3
# %bb.1:                                # %.preheader
	jle	.LBB18_2
# %bb.5:                                # %.lr.ph
	xorl	%ebx, %ebx
	leaq	8(%rsp), %r14
.LBB18_6:                               # =>This Inner Loop Header: Depth=1
	movq	%r14, %rdi
	movl	%ebx, %esi
	callq	hipGetDevicePropertiesR0600
	movl	4(%rsp), %ecx
	testl	%eax, %eax
	jne	.LBB18_8
# %bb.7:                                #   in Loop: Header=BB18_6 Depth=1
	cmpl	$0, 368(%rsp)
	jg	.LBB18_9
.LBB18_8:                               #   in Loop: Header=BB18_6 Depth=1
	incl	%ebx
	cmpl	%ecx, %ebx
	jl	.LBB18_6
	jmp	.LBB18_9
.LBB18_2:
	xorl	%ebx, %ebx
.LBB18_9:                               # %._crit_edge
	cmpl	%ecx, %ebx
	je	.LBB18_10
# %bb.11:
	movl	%ebx, %edi
	callq	hipSetDevice
	movb	$1, %al
.LBB18_12:
                                        # kill: def $al killed $al killed $eax
	addq	$1480, %rsp                     # imm = 0x5C8
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.LBB18_3:
	.cfi_def_cfa_offset 1504
	movq	stderr(%rip), %rcx
	movl	$.L.str, %edi
	movl	$20, %esi
	jmp	.LBB18_4
.LBB18_10:
	movq	stderr(%rip), %rcx
	movl	$.L.str.1, %edi
	movl	$40, %esi
.LBB18_4:
	movl	$1, %edx
	callq	fwrite@PLT
	xorl	%eax, %eax
	jmp	.LBB18_12
.Lfunc_end18:
	.size	_Z8InitCUDAv, .Lfunc_end18-_Z8InitCUDAv
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$40, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	cmpl	$2, %edi
	jl	.LBB19_17
# %bb.1:
	movq	8(%rsi), %rax
	cmpb	$45, (%rax)
	jne	.LBB19_18
# %bb.2:
	movzbl	1(%rax), %eax
	addl	$-67, %eax
	cmpl	$47, %eax
	ja	.LBB19_13
# %bb.3:
	movabsq	$4294967297, %rcx               # imm = 0x100000001
	btq	%rax, %rcx
	jb	.LBB19_7
# %bb.4:
	movl	$.Lstr.5, %edi
	movabsq	$68719476752, %rcx              # imm = 0x1000000010
	btq	%rax, %rcx
	jb	.LBB19_14
# %bb.5:
	movabsq	$140737488388096, %rcx          # imm = 0x800000008000
	btq	%rax, %rcx
	jae	.LBB19_13
# %bb.6:
	movb	$1, %bpl
	movl	$.Lstr.7, %edi
	jmp	.LBB19_8
.LBB19_17:                              # %.thread174
	movq	(%rsi), %rsi
	movl	$.L.str.2, %edi
	xorl	%eax, %eax
	callq	printf
	movl	$.Lstr, %edi
	callq	puts@PLT
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	jmp	.LBB19_18
.LBB19_7:
	movl	$.Lstr.6, %edi
	xorl	%ebp, %ebp
.LBB19_8:                               # %.thread186
	callq	puts@PLT
	movl	$.L.str.12, %edi
	callq	puts@PLT
	movl	$4, %ebx
	xorl	%r12d, %r12d
	leaq	8(%rsp), %r15
	movabsq	$3777893186295716171, %r14      # imm = 0x346DC5D63886594B
.LBB19_9:                               # =>This Inner Loop Header: Depth=1
	movq	%r12, UNIQUE(%rip)
	movq	%r12, TOTAL(%rip)
	leaq	24(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movl	%ebx, g(%rip)
	callq	_Z9listChainv
	testb	%bpl, %bpl
	je	.LBB19_11
# %bb.10:                               #   in Loop: Header=BB19_9 Depth=1
	callq	_Z11buildChainRv
	jmp	.LBB19_12
.LBB19_11:                              # %.critedge
                                        #   in Loop: Header=BB19_9 Depth=1
	callq	_Z10buildChainv
.LBB19_12:                              #   in Loop: Header=BB19_9 Depth=1
	movq	%r15, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	16(%rsp), %rax
	movq	24(%rsp), %rcx
	movq	32(%rsp), %rdi
	leaq	1000000(%rax), %rsi
	movq	%rcx, %rdx
	negq	%rdx
	cmpq	%rdi, %rax
	notq	%rcx
	cmovgeq	%rdx, %rcx
	cmovgeq	%rax, %rsi
	addq	8(%rsp), %rcx
	movq	%rcx, %rax
	movabsq	$1749024623285053783, %rdx      # imm = 0x1845C8A0CE512957
	imulq	%rdx
	movq	%rdx, %r8
	movq	%rdx, %rax
	shrq	$63, %rax
	sarq	$13, %r8
	addq	%rax, %r8
	subq	%rdi, %rsi
	addq	$500, %rsi                      # imm = 0x1F4
	movq	%rsi, %rax
	imulq	%r14
	movq	%rdx, %r10
	movq	%rdx, %rax
	shrq	$63, %rax
	shrq	$11, %r10
	addl	%eax, %r10d
	imulq	$86400, %r8, %rax               # imm = 0x15180
	subq	%rax, %rcx
	movslq	%ecx, %r11
	imulq	$-1851608123, %r11, %r9         # imm = 0x91A2B3C5
	shrq	$32, %r9
	addl	%r11d, %r9d
	movl	%r9d, %eax
	shrl	$31, %eax
	sarl	$11, %r9d
	addl	%eax, %r9d
	imull	$-3600, %r9d, %eax              # imm = 0xF1F0
	addl	%r11d, %eax
	movslq	%eax, %r13
	imulq	$-2004318071, %r13, %rax        # imm = 0x88888889
	shrq	$32, %rax
	addl	%eax, %r13d
	movl	%r13d, %eax
	shrl	$31, %eax
	sarl	$5, %r13d
	addl	%eax, %r13d
	imulq	$-2004318071, %r11, %rax        # imm = 0x88888889
	shrq	$32, %rax
	addl	%r11d, %eax
	movl	%eax, %ecx
	shrl	$31, %ecx
	sarl	$5, %eax
	addl	%ecx, %eax
	imull	$60, %eax, %eax
	subl	%eax, %r11d
	movq	TOTAL(%rip), %rdx
	movq	UNIQUE(%rip), %rcx
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movl	$.L.str.13, %edi
	movl	%ebx, %esi
                                        # kill: def $r8d killed $r8d killed $r8
                                        # kill: def $r9d killed $r9d killed $r9
	xorl	%eax, %eax
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	%r11
	.cfi_adjust_cfa_offset 8
	pushq	%r13
	.cfi_adjust_cfa_offset 8
	callq	printf
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
	incl	%ebx
	cmpl	$18, %ebx
	jne	.LBB19_9
	jmp	.LBB19_18
.LBB19_13:
	movl	$.Lstr.4, %edi
.LBB19_14:
	callq	puts@PLT
	callq	_Z8InitCUDAv
	testb	%al, %al
	je	.LBB19_18
# %bb.15:
	movl	$.L.str.12, %edi
	callq	puts@PLT
	movl	$4, %ebx
	xorl	%r12d, %r12d
	leaq	8(%rsp), %r15
	movabsq	$3777893186295716171, %r13      # imm = 0x346DC5D63886594B
	movabsq	$1749024623285053783, %rbp      # imm = 0x1845C8A0CE512957
.LBB19_16:                              # %.sink.split
                                        # =>This Inner Loop Header: Depth=1
	leaq	24(%rsp), %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	%r12, UNIQUE(%rip)
	movq	%r12, TOTAL(%rip)
	movl	%ebx, g(%rip)
	callq	_Z9listChainv
	movl	%ebx, %edi
	movl	$24576, %esi                    # imm = 0x6000
	callq	_Z15GPU_buildChainRjj
	movq	%r15, %rdi
	xorl	%esi, %esi
	callq	gettimeofday
	movq	16(%rsp), %rdx
	movq	24(%rsp), %rcx
	movq	32(%rsp), %rsi
	leaq	1000000(%rdx), %rax
	movq	%rcx, %rdi
	negq	%rdi
	cmpq	%rsi, %rdx
	notq	%rcx
	cmovgeq	%rdi, %rcx
	cmovgeq	%rdx, %rax
	addq	8(%rsp), %rcx
	movslq	%ecx, %r8
	imulq	$-1037155065, %r8, %rdx         # imm = 0xC22E4507
	shrq	$32, %rdx
	addl	%edx, %r8d
	movl	%r8d, %edx
	shrl	$31, %edx
	sarl	$16, %r8d
	addl	%edx, %r8d
	subq	%rsi, %rax
	addq	$500, %rax                      # imm = 0x1F4
	imulq	%r13
	movq	%rdx, %r10
	movq	%rdx, %rax
	shrq	$63, %rax
	shrq	$11, %r10
	addl	%eax, %r10d
	movq	%rcx, %rax
	imulq	%rbp
	movq	%rdx, %rax
	shrq	$63, %rax
	sarq	$13, %rdx
	addq	%rax, %rdx
	imulq	$86400, %rdx, %rax              # imm = 0x15180
	subq	%rax, %rcx
	movslq	%ecx, %r11
	imulq	$-1851608123, %r11, %r9         # imm = 0x91A2B3C5
	shrq	$32, %r9
	addl	%r11d, %r9d
	movl	%r9d, %eax
	shrl	$31, %eax
	sarl	$11, %r9d
	addl	%eax, %r9d
	imull	$-3600, %r9d, %eax              # imm = 0xF1F0
	addl	%r11d, %eax
	movslq	%eax, %r14
	imulq	$-2004318071, %r14, %rax        # imm = 0x88888889
	shrq	$32, %rax
	addl	%eax, %r14d
	movl	%r14d, %eax
	shrl	$31, %eax
	sarl	$5, %r14d
	addl	%eax, %r14d
	imulq	$-2004318071, %r11, %rax        # imm = 0x88888889
	shrq	$32, %rax
	addl	%r11d, %eax
	movl	%eax, %ecx
	shrl	$31, %ecx
	sarl	$5, %eax
	addl	%ecx, %eax
	imull	$60, %eax, %eax
	subl	%eax, %r11d
	movq	TOTAL(%rip), %rdx
	movq	UNIQUE(%rip), %rcx
	subq	$8, %rsp
	.cfi_adjust_cfa_offset 8
	movl	$.L.str.13, %edi
	movl	%ebx, %esi
                                        # kill: def $r8d killed $r8d killed $r8
                                        # kill: def $r9d killed $r9d killed $r9
	xorl	%eax, %eax
	pushq	%r10
	.cfi_adjust_cfa_offset 8
	pushq	%r11
	.cfi_adjust_cfa_offset 8
	pushq	%r14
	.cfi_adjust_cfa_offset 8
	callq	printf
	addq	$32, %rsp
	.cfi_adjust_cfa_offset -32
	incl	%ebx
	cmpl	$22, %ebx
	jne	.LBB19_16
.LBB19_18:                              # %.thread
	xorl	%eax, %eax
	addq	$40, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end19:
	.size	main, .Lfunc_end19-main
	.cfi_endproc
                                        # -- End function
	.type	TOTAL,@object                   # @TOTAL
	.bss
	.globl	TOTAL
	.p2align	3, 0x0
TOTAL:
	.quad	0                               # 0x0
	.size	TOTAL, 8

	.type	UNIQUE,@object                  # @UNIQUE
	.globl	UNIQUE
	.p2align	3, 0x0
UNIQUE:
	.quad	0                               # 0x0
	.size	UNIQUE, 8

	.type	g,@object                       # @g
	.globl	g
	.p2align	2, 0x0
g:
	.zero	7444
	.size	g, 7444

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"There is no device.\n"
	.size	.L.str, 21

	.type	.L.str.1,@object                # @.str.1
.L.str.1:
	.asciz	"There is no device supporting CUDA 1.x.\n"
	.size	.L.str.1, 41

	.type	.L.str.2,@object                # @.str.2
.L.str.2:
	.asciz	"Usage: %s [-c|-g|-r|-s] n steps\n"
	.size	.L.str.2, 33

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	" N:        Total      Unique      dd:hh:mm:ss.ms"
	.size	.L.str.12, 49

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"%2d:%13ld%12ld%8.2d:%02d:%02d:%02d.%02d\n"
	.size	.L.str.13, 41

	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.type	.Lstr,@object                   # @str
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr:
	.asciz	"  -r: CPU \345\206\215\345\270\260"
	.size	.Lstr, 17

	.type	.Lstr.1,@object                 # @str.1
.Lstr.1:
	.asciz	"  -c: CPU \351\235\236\345\206\215\345\270\260"
	.size	.Lstr.1, 20

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"  -g: GPU \345\206\215\345\270\260"
	.size	.Lstr.2, 17

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"  -n: GPU \343\202\255\343\203\243\343\203\252\343\203\274\343\203\201\343\202\247\343\203\274\343\203\263"
	.size	.Lstr.3, 35

	.type	.Lstr.4,@object                 # @str.4
.Lstr.4:
	.asciz	"\n\nGPU \343\202\255\343\203\243\343\203\252\343\203\274\343\203\201\343\202\247\343\203\274\343\203\263 \343\203\236\343\203\253\343\203\201\343\202\271\343\203\254\343\203\203\343\203\211"
	.size	.Lstr.4, 53

	.type	.Lstr.5,@object                 # @str.5
.Lstr.5:
	.asciz	"\n\nGPU \343\202\255\343\203\243\343\203\252\343\203\274\343\203\201\343\202\247\343\203\274\343\203\263 \343\202\267\343\203\263\343\202\260\343\203\253\343\202\271\343\203\254\343\203\203\343\203\211"
	.size	.Lstr.5, 56

	.type	.Lstr.6,@object                 # @str.6
.Lstr.6:
	.asciz	"\n\nCPU \343\202\255\343\203\243\343\203\252\343\203\274\343\203\201\343\202\247\343\203\274\343\203\263 \351\235\236\345\206\215\345\270\260 "
	.size	.Lstr.6, 42

	.type	.Lstr.7,@object                 # @str.7
.Lstr.7:
	.asciz	"\n\nCPU \343\202\255\343\203\243\343\203\252\343\203\274\343\203\201\343\202\247\343\203\274\343\203\263 \345\206\215\345\270\260 "
	.size	.Lstr.7, 39

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
