Nクイーン問題(62)第七章 並列処理 対称解除法 ノードレイヤー NVIDIA CUDA編


【参考リンク】Nクイーン問題 過去記事一覧はこちらから

エイト・クイーンのプログラムアーカイブ
Bash、Lua、C、Java、Python、CUDAまで!
https://github.com/suzukiiichiro/N-Queens

CUDAセットアップ

手順など

今どきは少なくなったわけですが、CUDAをインストールから始めたい人はこちら

EC2 G4インスタンスのAmazon Linux 2にNVIDIA CUDAをインストールしてみた
https://dev.classmethod.jp/articles/install-nvidia-cuda-on-ec2-g4-amazon-linux-2/

今どきは、GPUインスタンスを使えば最初からnvccが実行できます。

AWSのGPUインスタンスでCUDAを動かす
https://qiita.com/navitime_tech/items/0a8073347c21800f1cad

料金やインスタンスの選定、起動までが詳しく書かれているページはこちら

AWS/EC2でGPUマシンを起動してみる
https://motomura-kazushi.net/2021/08/15/aws-ec2%E3%81%A7gpu%E3%83%9E%E3%82%B7%E3%83%B3%E3%82%92%E8%B5%B7%E5%8B%95%E3%81%97%E3%81%A6%E6%B7%B1%E5%B1%A4%E5%AD%A6%E7%BF%92%E3%82%92%E5%9B%9E%E3%81%97%E3%81%A6%E3%81%BF%E3%82%8B/

確認方法

GPUインスタンスでCUDAが使えるようになったかな?と、確認するには以下のコマンドを実行します。

$ nvcc --version

以下の内容が表示されればCUDAは実行可能な状態にあります。

NQueens2$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2023 NVIDIA Corporation
Built on Tue_Jul_11_02:31:28_PDT_2023
Cuda compilation tools, release 12.2, V12.2.128
Build cuda_12.2.r12.2/compiler.33053471_0
NQueens2$

よくありがちな質問ですが、topコマンドでは、GPUインスタンスの負荷を調べることはできませんので、以下のコマンドでGPUインスタンスのCPU不可やメモリ使用状況を確認します。

$ watch nvidia-smi

ミラー CUDA部分についての解説

実行

実行は以下の4種類です。

1.シングルスレッドで再帰

// 再帰 対称解除法
void symmetry_R(unsigned int size,struct local* l)
$ nvcc 03CUDA_Symmetry_NodeLayer.cu && ./a.out -r

2.シングルスレッドで非再帰

// 非再帰 対称解除法
void symmetry_NR(unsigned int size,struct local* l)
$ nvcc 03CUDA_Symmetry_NodeLayer.cu && ./a.out -c

3.シングルスレッドのGPU
こちらはInitCUDA()を通過後、以下のメソッドを呼び出します。
シングルスレッドで動作します。

// GPU 対称解除法 -g の実行時のみ呼び出されます
__host__ __device__
void GPU_symmetry_R(unsigned int size,struct local* hostLocal)

$ nvcc 03CUDA_Symmetry_NodeLayer.cu && ./a.out -g

3の必要性は、4を開発している間、問題点を局所化するためです。
3の解がきちんと出力していれば、他の箇所に問題があると特定することが目的です。

4.マルチスレッドのGPU
こちらはInitCUDA()を通過後、以下のメソッドを呼び出します。
マルチスレッドで動作します。

// GPU -n ノードレイヤーの作成
void symmetry_build_nodeLayer(unsigned int size)
$ nvcc 03CUDA_Symmetry_NodeLayer.cu && ./a.out -n

次の項では、4について具体的に説明します。

関数の説明

InitCUDA()を通過後、以下のメソッドを順次辿って実行されます。

CUDA 初期化

bool InitCUDA()

CUDAデバイスが搭載されているか、GPUの数などを確認します。
GPUが搭載されていない場合は、returnします。

GPU -n ノードレイヤーの作成

void symmetry_build_nodeLayer(unsigned int size)

以下でレイヤーの数を指定します。
Nが増えればレイヤーは枯渇します。
Nが16まではレイヤーは4で足りますが、以降、レイヤーは、5,6と増やす必要があり、レイヤーが増えることによって、速度は加速度的に遅くなります。
ノードレイヤーの考え方はスマートではありますが、Nの最大化と高速化を求める場合は限界がまもなくおとずれるロジックです。

今の段階では、レイヤー4で進めていきます。

  // ツリーの3番目のレイヤーにあるノード
  //(それぞれ連続する3つの数字でエンコードされる)のベクトル。
  // レイヤー2以降はノードの数が均等なので対称性を利用できる。
  // レイヤ4には十分なノードがある(N16の場合、9844)。
  // ここではレイヤーを5に設定、Nに併せて増やしていく
  std::vector<local>L;
  // NodeLayerは、N18でabortします。
  std::vector<long> nodes=kLayer_nodeLayer(size,5,L); 

以下の部分で nodeSizeを取得します。
また、hostNodesとdeviceNodesを宣言しメモリ領域を確保します。
併せて、hostNodes配列の先頭に空の値を格納し、cudaMemcpy()でデバイス側に配列をコピー(転送)します。

  // デバイスにはクラスがないので、
  // 最初の要素を指定してからデバイスにコピーする。
  size_t nodeSize=nodes.size() * sizeof(long);
  long* hostNodes=(long*)malloc(nodeSize);
  hostNodes=&nodes[0];
  long* deviceNodes=NULL;
  cudaMalloc((void**)&deviceNodes,nodeSize);
  cudaMemcpy(deviceNodes,hostNodes,nodeSize,cudaMemcpyHostToDevice);

この部分では、構造体 Localをデバイスに転送して利用するために、hostLocaldeviceLocalを作成します。
併せて、hostLocal構造体の配列の先頭にからの値を格納し、cudaMemcpy()でデバイス側に配列をコピー(転送)します。

  // host/device Local
  //size_t localSize=numSolutions * sizeof(struct local);
  size_t localSize=L.size() * sizeof(struct local);
  struct local* hostLocal=(local*)malloc(localSize);
  hostLocal=&L[0];
  local* deviceLocal=NULL;
  cudaMalloc((void**)&deviceLocal,localSize);
  cudaMemcpy(deviceLocal, hostLocal, localSize, cudaMemcpyHostToDevice);

以下は、解を格納する deviceSolutions を宣言します。
ミラーでは、必要なノードは半分で済みます。
3の整数というのは、一つのノードにleft down right が格納されるからです。

  // デバイス出力の割り当て
  // 必要なのはノードの半分だけで、
  // 各ノードは3つの整数で符号化される。
  long* deviceSolutions=NULL;
  int numSolutions=nodes.size() / 3; 
  size_t solutionSize=numSolutions * sizeof(long);
  cudaMalloc((void**)&deviceSolutions,solutionSize);

以下でCUDAカーネルを起動します。
起動は

起動したい関数 <<< 並列処理の数 >>>( 関数に渡すパラメータ)

となります。

  // CUDAカーネルを起動する。
  int threadsPerBlock=256;
  int blocksPerGrid=(numSolutions+threadsPerBlock-1)/threadsPerBlock;
  dim_nodeLayer <<<blocksPerGrid,threadsPerBlock>>>(size,deviceNodes,deviceSolutions,numSolutions,deviceLocal);

以下では、デバイス側で処理した結果がdeviceSolusionsに格納され、その配列を cudaMemcpy()でホスト側にコピー(転送)します。

  // 結果をホストにコピー
  long* hostSolutions=(long*)malloc(solutionSize);
  cudaMemcpy(hostSolutions,deviceSolutions,solutionSize,cudaMemcpyDeviceToHost);

デバイスからホストへ転送する場合は

cudaMemcpyDeviceToHost

ホストからデバイスへ転送する場合は

cudaMemcpyHostToDevice

となります。

以下は、スレッドごとに格納された解をforで回して集計します。

  // 部分解を加算し、結果を表示する。
  unsigned long solutions=0;
  for(unsigned long i=0;i<numSolutions;i++){
       solutions += hostSolutions[i]; // Symmetry
  }
  // 出力
  TOTAL=solutions;

GPU -n Kレイヤー k 番目のレイヤのすべてのノードを含むベクトルを返す

std::vector<long> kLayer_nodeLayer(unsigned int size,unsigned int k,std::vector<local>& L)

こちらの関数は、kLayer_nodeLayer()を使って、必要なノードを埋める処理をします。ようするに、Nクイーンの当たり判定をすべて行いノードを作成するということになります。

ノードレイヤーが遅い理由は、Nクイーンの処理を、当たり判定と本番確定の2回行っていることです。

メモ
GPUで並列実行するためのleft,right,downを作成する

kLayer_nodeLayer(size,4)
第2引数の4は4行目までnqueenを実行し、それまでのleft,down,rightをnodes配列に格納する

nodesはベクター配列で構造体でもなんでも格納できる
push_backで追加。

nodes配列は3個で1セットleft,dwon,rightの情報を同じ配列に格納する
[0]left[1]down[2]right

GPU -n Kレイヤー 再帰 角にQがないときのバックトラック

unsigned long kLayer_nodeLayer_backTrack(int size,std::vector<long>& nodes,unsigned int k,unsigned long left,unsigned long down,unsigned long right,std::vector<local>& L,struct local* l)

こちらの関数は、「ノードレイヤーにおける」ボードの最上部の角にQがないときのバックトラックです。枝刈り処理が含まれています。

    if(row<l->BOUND1){   //枝刈り
      bitmap=bitmap|2;
      bitmap=bitmap^2;
    }

GPU -n Kレイヤー 角にQがあるときのバックトラック

unsigned long kLayer_nodeLayer_backTrack_corner(unsigned int size,std::vector<long>& nodes,unsigned int k,unsigned long left,unsigned long down,unsigned long right,std::vector<local>& L,struct local* l)

こちらの関数は、「ノードレイヤーにおける」ボードの最上部の角にQがないときのバックトラックです。枝刈り処理が含まれています。

    if(row<l->BOUND1){
      bitmap=bitmap|l->SIDEMASK;
      bitmap=bitmap^l->SIDEMASK;
    }else{
      if(row==l->BOUND2){
        if((down&l->SIDEMASK)==0){
          return 0;
        }
        if( (down&l->SIDEMASK)!=l->SIDEMASK){
          bitmap=bitmap&l->SIDEMASK;
        }
      }
    }

GPU -n ノードレイヤー i 番目のメンバを i 番目の部分木の解で埋める

__global__ 
void dim_nodeLayer(unsigned int size,long* nodes,long* solutions,unsigned int numElements,struct local* l)

ここでは、解を導き出すために、ノードにより、クイーンの移動の候補を列挙する処理を行います。
ノードレイヤーは、解を導き出すための処理に加えて、クイーンの移動候補を列挙するための処理を行うため、事実上、2回エイトクイーンを行っていることが、速度低下の最大のボトルネックと言えます。

ボード最上部の角にQがあるかないかの判定でGPU_symmetry_solve_nodeLayer_corner()を呼び出すかGPU_symmetry_solve_nodeLayer()を呼び出すかの判定処理を行っています。

GPU -n ノードレイヤーによる対称解除法 -n の実行時に呼び出される

__host__ __device__ 
long GPU_symmetry_solve_nodeLayer(unsigned int size,unsigned long left,unsigned long down,unsigned long right,struct local* l)

ここでは、ボードの最上部の角にQがないときの処理を行います。
「ノードレイヤーにおける」ではありません。実際の解を求める処理となります。

GPU -n ノードレイヤーによる対称解除法 -n の実行時に呼び出される

__host__ __device__ 
long GPU_symmetry_solve_nodeLayer_corner(unsigned int size,unsigned long left,unsigned long down,unsigned long right,struct local* l)

ここでは、ボードの最上部の角にQがあるときの処理を行います。
「ノードレイヤーにおける」ではありません。実際の解を求める処理となります。

0以外のbitをカウント

__host__ __device__
unsigned int countBits_nodeLayer(unsigned long n)

ここでは、ビットで表現されたボード情報の右端にあるゼロ以外の数字を削除する処理を行います。

GPU -n の 対称解除法

__host__ __device__
long GPUN_symmetryOps(unsigned int size,struct local* l)

ここが対称解除法の肝の部分です。
事項で詳細に説明します。

対象解除法について

まず、Nが小さな盤面で考えていきます。
以下、順に見て理解を深めてもらえればと思います。

##対象解除法について(N2版)
ひとつの解には、盤面を90度・180度・270度回転、及びそれらの鏡像の合計8個の対称解が存在します。

  原型   90度  180度   270度
  12   41   34   23
  43   32   21   14
  
      上の行を左右反転                     
  21   14   43   32
  34   23   12   41

上図左上がユニーク解です。

  原型(ユニーク解)
  12 
  43 

1行目はユニーク解を90度、180度、270度回転したものです。

  原型   90度  180度   270度
  12   41   34   23
  43   32   21   14

2行目は1行目のそれぞれを左右反転したものです。

      上の行を左右反転                     
  21   14   43   32
  34   23   12   41

2行目はユニーク解を左右反転、対角反転、上下反転、逆対角反転したものとも解釈できます。
ただし、回転・線対称な解もあるので注意が必要です。

対称的な解を除去し、ユニーク解のみを求める方法はいくつかあります。
ここでは、一つの解には8つのパターンが存在することを覚えておいてください。

対象解を導き出す2つの方法

対象解を導き出す方法として、
1.解が見つかるたびに回転・対象をチェック
2.解のすべてを保存して、保存された解の回転・対象をチェック

最小解選択法
解がひとつみつかるとすべての対称解を生成、 状態を数値とみなして最も小さいもののみを解とする方法。
最も最適と言われています。

最小選択法よりも劣る方法
解全てを保存しておき、新しい解が発見されるたびに対称形が無いかどうかを調べる方法。
保存領域を必要とするし、比較時に対称形を生成する必要があります。

対象解除法(N5版)

全探索によって得られたある1つの解が、回転・反転などによる本質的に変わることのない変換によって他の解と同型となるものが存在する場合、それを別の解とはしないとする解の数え方で得られる解を「ユニーク解」といいます。

ユニーク解とは、全解の中から回転・反転などによる変換によって同型になるものどうしをグループ化することを意味しています。

グループ1:
+-+-+-+-+-+  +-+-+-+-+-+ 
| | | |Q| |  | |Q| | | |
+-+-+-+-+-+  +-+-+-+-+-+
|Q| | | | |  | | | | |Q|
+-+-+-+-+-+  +-+-+-+-+-+
| | |Q| | |  | | |Q| | |
+-+-+-+-+-+  +-+-+-+-+-+
| | | | |Q|  |Q| | | | |
+-+-+-+-+-+  +-+-+-+-+-+
| |Q| | | |  | | | |Q| |
+-+-+-+-+-+  +-+-+-+-+-+

グループ2:
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  
| | | | |Q|  |Q| | | | |  | | |Q| | |  | | |Q| | |  | | | |Q| |  | |Q| | | |  |Q| | | | |  | | | | |Q|  
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  
| | |Q| | |  | | |Q| | |  |Q| | | | |  | | | | |Q|  | |Q| | | |  | | | |Q| |  | | | |Q| |  | |Q| | | |  
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  
|Q| | | | |  | | | | |Q|  | | | |Q| |  | |Q| | | |  | | | | |Q|  |Q| | | | |  | |Q| | | |  | | | |Q| |  
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  
| | | |Q| |  | |Q| | | |  | |Q| | | |  | | | |Q| |  | | |Q| | |  | | |Q| | |  | | | | |Q|  |Q| | | | |  
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  
| |Q| | | |  | | | |Q| |  | | | | |Q|  |Q| | | | |  |Q| | | | |  | | | | |Q|  | | |Q| | |  | | |Q| | |  
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  

ユニーク解はその「個数のみ」に着目され、この解はユニーク解であり、この解はユニーク解ではないという定まった判定方法はありません。
ユニーク解であるかどうかの判断はユニーク解の個数を数える目的の為だけにあります。
どのような定義をしたとしてもユニーク解の個数それ自体は変わりません。

Nクイーン問題は、正方形のボードで形成されるので、回転・反転による変換パターンはぜんぶで8通りあります。

だからといって「全解数=ユニーク解数×8」と単純にはいきません。
ひとつのグループの要素数が必ず8個あるとは限らないのです。

N=5の例では2つのグループがあり、要素数が2個のグループ、要素数が8個のグループがあります。
結論から言うと、N=5の全解は10個で、ユニーク解は2個です。

ユニーク解を判定するための定義

各行のクイーンが左から何番目にあるかを調べて、最上段の行(row)から下の行(row)へ順番に列挙します。
そしてそれをN桁の数値として見た場合に「最大値」となるものをユニーク解とします。
このN桁の数を以後は「ユニーク判定値」と呼ぶことにします。

 0 1 2 3 4 
+-+-+-+-+-+  
| | | | |Q|    4
+-+-+-+-+-+  
| | |Q| | |    2 
+-+-+-+-+-+  
|Q| | | | |    0  ----> 4 2 0 3 1 (ユニーク判定値)
+-+-+-+-+-+             数が大きい方をユニークとみなす 
| | | |Q| |    3
+-+-+-+-+-+  
| |Q| | | |    1 
+-+-+-+-+-+  

 0 1 2 3 4     左右反転!
+-+-+-+-+-+  
|Q| | | | |    0
+-+-+-+-+-+  
| | |Q| | |    2 
+-+-+-+-+-+  
| | | | |Q|    4  ----> 0 2 4 1 3 
+-+-+-+-+-+            数が小さいのでユニーク解とはしません 
| |Q| | | |    1
+-+-+-+-+-+  
| | | |Q| |    3 
+-+-+-+-+-+  

探索によって得られたある1つの解(オリジナル)がユニーク解であるかどうかを判定するには、「8通りの変換を試み、その中でオリジナルのユニーク判定値が最小であるかを調べる」ことになります。

N2の場合の8通りの変換

  原型   90度  180度   270度
  12   41   34   23
  43   32   21   14
  
      上の行を左右反転                     
  21   14   43   32
  34   23   12   41

結論から先にいえば、ユニーク解とは成り得ないことが明確なパターンを探索中に切り捨てる「枝刈り」を組み込むことにより、3通りの変換を試みるだけでユニーク解の判定が可能になります。

枝刈り

先ず最上段の行(row)のクイーンの位置に着目します。
その位置が左半分の領域にあればユニーク解には成り得ません。
何故なら左右反転によって得られるパターンのユニーク判定値の方が確実に小さくなるからです。

 0 1 2 3 4     
+-+-+-+-+-+  
|Q| | | | |    0
+-+-+-+-+-+  
| | |Q| | |    2 
+-+-+-+-+-+  
| | | | |Q|    4  ----> 0 2 4 1 3 
+-+-+-+-+-+            数が小さいのでユニークとはしません
| |Q| | | |    1
+-+-+-+-+-+  
| | | |Q| |    3 
+-+-+-+-+-+  

 0 1 2 3 4     左右反転!
+-+-+-+-+-+  
| | | | |Q|    4
+-+-+-+-+-+  
| | |Q| | |    2 
+-+-+-+-+-+  
|Q| | | | |    0  ----> 4 2 0 3 1 (ユニーク判定値)
+-+-+-+-+-+             数が大きい方をユニークとみなす 
| | | |Q| |    3
+-+-+-+-+-+  
| |Q| | | |    1 
+-+-+-+-+-+  

Nが奇数の場合に中央にあった場合はどうでしょう。

+-+-+-+-+-+  
| | |Q| | |    2
+-+-+-+-+-+  
| | | | | |      
+-+-+-+-+-+  
| | | | | |       
+-+-+-+-+-+             
| | | | | |     
+-+-+-+-+-+  
| | | | | |      
+-+-+-+-+-+  

これもユニーク解には成り得ません。
何故なら仮に中央にあった場合、それがユニーク解であるためには少なくとも他の3辺におけるクイーンの位置も中央になければならず、それは互いの効き筋にあたるので解になりえないからです。

+-+-+-+-+-+  
| | |Q| | |    
+-+-+-+-+-+  
| | | | | |      
+-+-+-+-+-+  
|Q| | | |Q|       
+-+-+-+-+-+             
| | | | | |     
+-+-+-+-+-+  
| | |Q| | |      
+-+-+-+-+-+  

最上段の行のクイーンの位置は中央を除く右側の領域に限定されます。(ただし、N ≧ 2)

+-+-+-+-+-+  
| | | |Q|Q|    
+-+-+-+-+-+  
| | | | | |      
+-+-+-+-+-+  
| | | | | |       
+-+-+-+-+-+             
| | | | | |     
+-+-+-+-+-+  
| | | | | |      
+-+-+-+-+-+  

次にその中でも一番右端(右上の角)にクイーンがある場合を考えてみます。

+-+-+-+-+-+  
| | | | |Q|    
+-+-+-+-+-+  
| | | | | |      
+-+-+-+-+-+  
| | | | | |       
+-+-+-+-+-+             
| | | | | |     
+-+-+-+-+-+  
| | | | | |      
+-+-+-+-+-+  

他の3つの角にクイーンを置くことはできないので、

+-+-+-+-+-+  
|Q| | | |Q|    
+-+-+-+-+-+  
| | | | | |      
+-+-+-+-+-+  
| | | | | |       
+-+-+-+-+-+             
| | | | | |     
+-+-+-+-+-+  
|Q| | | |Q|      
+-+-+-+-+-+  

ユニーク解であるかどうかを判定するには、右上角から左下角を通る斜軸で反転させたパターンとの比較だけになります。

+-+-+-+-+-+  
| | | | |Q|    
+-+-+-+-+-+  
| | | |/| |      
+-+-+-+-+-+  
| | |/| | |       
+-+-+-+-+-+             
| |/| | | |     
+-+-+-+-+-+  
|/| | | | |      
+-+-+-+-+-+  

突き詰めれば、上から2行目(row1)のクイーンの位置が右から何番目にあるかと、

+-+-+-+-+-+  
| | | | |Q|    
+-+-+-+-+-+  
| |Q| |/| |  ←
+-+-+-+-+-+  
| | |/| | |       
+-+-+-+-+-+             
| |/| | | |     
+-+-+-+-+-+  
|/| | | | |      
+-+-+-+-+-+  

右から2列目のクイーンの位置が上から何番目にあるかを比較するだけで判定することができます。

+-+-+-+-+-+  
| | | | |Q|    
+-+-+-+-+-+  
| |Q| |/| |      
+-+-+-+-+-+  
| | |/| | |       
+-+-+-+-+-+             
| |/| |Q| |  ←    
+-+-+-+-+-+  
|/| | | | |      
+-+-+-+-+-+  

この2つの値が同じになることはないからです。

       3 0
+-+-+-+-+-+  
| | | | |Q| 0
+-+-+-+-+-+  
| |Q| |/| | 3    上から2行目のクイーンの位置が右から4番目にある。
+-+-+-+-+-+  
| | |/| | |       
+-+-+-+-+-+             
| |/| |Q| |  ←  右から2列目のクイーンの位置が上から4番目にある。
+-+-+-+-+-+      しかし、互いの効き筋にあたるのでこれは有り得ない。 
|/| | | | |      
+-+-+-+-+-+  

結局、再帰探索中において下図の X への配置を禁止する枝刈りを入れておけば、得られる解は総てユニーク解であることが保証されます。

+-+-+-+-+-+  
| | | |X|Q| 
+-+-+-+-+-+  
| |Q| |X| | 
+-+-+-+-+-+  
| | | |X| |       
+-+-+-+-+-+             
| | | |Q| | 
+-+-+-+-+-+ 
| | | | | |      
+-+-+-+-+-+  

次に右端以外にクイーンがある場合を考えてみます。
オリジナルがユニーク解であるためには先ず下図の X への配置は禁止されます。
よって、その枝刈りを先ず入れておきます。

+-+-+-+-+-+-+-+-+  
|X|X| | | |Q|X|X| 
+-+-+-+-+-+-+-+-+  
|X| | | | | | |X| 
+-+-+-+-+-+-+-+-+  
| | | | | | | | |       
+-+-+-+-+-+-+-+-+             
| | | | | | | | | 
+-+-+-+-+-+-+-+-+ 
| | | | | | | | |      
+-+-+-+-+-+-+-+-+  
| | | | | | | | |
+-+-+-+-+-+-+-+-+
|X| | | | | | |X|
+-+-+-+-+-+-+-+-+
|X|X| | | | |X|X|
+-+-+-+-+-+-+-+-+

次にクイーンの利き筋を辿っていくと、結局、オリジナルがユニーク解ではない可能性があるのは、下図の A,B,C の位置のどこかにクイーンがある場合に限られます。

+-+-+-+-+-+-+-+-+  
|X|X| | | |Q|X|X| 
+-+-+-+-+-+-+-+-+  
|X| | | |x|x|x|X| 
+-+-+-+-+-+-+-+-+  
|C| | |x| |x| |x|       
+-+-+-+-+-+-+-+-+             
| | |x| | |x| | | 
+-+-+-+-+-+-+-+-+ 
| |x| | | |x| | |      
+-+-+-+-+-+-+-+-+  
|x| | | | |x| |A|
+-+-+-+-+-+-+-+-+
|X| | | | |x| |X|
+-+-+-+-+-+-+-+-+
|X|X|B| | |x|X|X|
+-+-+-+-+-+-+-+-+

従って、90度回転、180度回転、270度回転の3通りの変換パターンだけを調べれはよいことになります。

ユニーク解を数える

これまでの考察はユニーク解の個数を求めるためのものでした。
全解数を求めるにはユニーク解を求めるための枝刈りを取り除いて全探索する必要があります。
したがって探索時間を犠牲にしてしまうことになります。
そこで「ユニーク解の個数から全解数を導いてしまおう」という試みが考えられます。
これは、左右反転(前章のミラー)によるパターンの探索を省略して最後に結果を2倍するというアイデアの拡張版といえるものです。
そしてそれを実現させるには「ユニーク解が属するグループの要素数はいくつあるのか」という考察が必要になってきます。

最初に、クイーンが右上角にあるユニーク解を考えます。

+-+-+-+-+-+  
| | | | |Q| 
+-+-+-+-+-+  
| | | | | | 
+-+-+-+-+-+  
| | | | | |       
+-+-+-+-+-+             
| | | | | | 
+-+-+-+-+-+ 
| | | | | |      
+-+-+-+-+-+  

斜軸で反転したパターンがオリジナルと同型になることは有り得ないことと(×2)、

+-+-+-+-+-+  +-+-+-+-+-+ 
| | | | |Q|  | | | | |Q| 
+-+-+-+-+-+  +-+-+-+-+-+ 
| | |Q|/| |  | | | |/| | 
+-+-+-+-+-+  +-+-+-+-+-+ 
| | |/| | |  | | |/|Q| | 
+-+-+-+-+-+  +-+-+-+-+-+            
| |/| | | |  | |/| | | | 
+-+-+-+-+-+  +-+-+-+-+-+ 
|/| | | | |  |/| | | | |     
+-+-+-+-+-+  +-+-+-+-+-+ 

右上角のクイーンを他の3つの角に写像させることができるので(×4)、

+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+ 
| | | | |Q|  | | | | | |  | | | | | |  |Q| | | | | 
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+ 
| | | | | |  | | | | | |  | | | | | |  | | | | | | 
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+ 
| | | | | |  | | | | | |  | | | | | |  | | | | | |      
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+            
| | | | | |  | | | | | |  | | | | | |  | | | | | | 
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+ 
| | | | | |  | | | | |Q|  |Q| | | | |  | | | | | |     
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+ 

このユニーク解が属するグループの要素数は必ず8個(=2×4)になります。

次に、クイーンが右上角以外にある場合は少し複雑になります。

(1) 90度回転させてオリジナルと同型になる場合、さらに90度回転(オリジナルから180度回転)させても、さらに90度回転(オリジナルから270度回転)させてもオリジナルと同型になる。

90度回転させても同じ場合は

+-+-+-+-+-+  +-+-+-+-+-+ 
| | | |Q| |  | | | |Q| | 
+-+-+-+-+-+  +-+-+-+-+-+ 
|Q| | | | |  |Q| | | | | 
+-+-+-+-+-+  +-+-+-+-+-+ 
| | |Q| | |  | | |Q| | | 
+-+-+-+-+-+  +-+-+-+-+-+            
| | | | |Q|  | | | | |Q| 
+-+-+-+-+-+  +-+-+-+-+-+ 
| |Q| | | |  | |Q| | | |     
+-+-+-+-+-+  +-+-+-+-+-+ 

さらに90度回転(原型から180度回転)させても

                 さらに90度回転
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+ 
| | | |Q| |  | | | |Q| |  | | | |Q| | 
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+ 
|Q| | | | |  |Q| | | | |  |Q| | | | | 
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+ 
| | |Q| | |  | | |Q| | |  | | |Q| | | 
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+            
| | | | |Q|  | | | | |Q|  | | | | |Q| 
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+ 
| |Q| | | |  | |Q| | | |  | |Q| | | |     
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+ 

さらに90度回転(原型から270度回転)させても同じ!

                               さらに90度回転
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+ 
| | | |Q| |  | | | |Q| |  | | | |Q| |  | | | |Q| | 
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+ 
|Q| | | | |  |Q| | | | |  |Q| | | | |  |Q| | | | | 
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+ 
| | |Q| | |  | | |Q| | |  | | |Q| | |  | | |Q| | | 
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+            
| | | | |Q|  | | | | |Q|  | | | | |Q|  | | | | |Q| 
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+ 
| |Q| | | |  | |Q| | | |  | |Q| | | |  | |Q| | | |     
+-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+  +-+-+-+-+-+ 


(2) 90度回転させてオリジナルと異なる場合は、270度回転させても必ずオリジナルとは異なる。
ただし、180度回転させた場合はオリジナルと同型になることも有り得る。

ユニーク数から全解を求める

1.クイーンが右上角にある場合、ユニーク解が属するグループの要素数は必ず8個(=2×4)

2.クイーンが右上角以外にある場合、
(1) 90度回転させてオリジナルと同型になる場合、さらに90度回転(オリジナルから180度回転)させても、
さらに90度回転(オリジナルから270度回転)させてもオリジナルと同型になる。
こちらに該当するユニーク解が属するグループの要素数は、左右反転させたパターンを加えて2個しかありません。

(2) 90度回転させてオリジナルと異なる場合は、270度回転させても必ずオリジナルとは異なる。
ただし、180度回転させた場合はオリジナルと同型になることも有り得る。
こちらに該当するユニーク解が属するグループの要素数は、180度回転させて同型になる場合は4個(左右反転×縦横回転)

(3)180度回転させてもオリジナルと異なる場合は、8個(左右反転×縦横回転×上下反転)

1の場合は、解の数x8
2−(1)の場合は、解の数x2
2−(2)の場合は、解の数x4
2−(3)の場合は、解の数x8
それぞれのグループの解の数にグループのパターン数をかけ合わせた数の総計が全解となります。
ひとつひとつのユニーク解が上のどの種類に該当するのかを調べることにより全解数を計算で導き出すことができるのです。

対象解除法 4つの関数

対象解除法のソースは4つの関数で構成されています。

1.printRecord()
解を発見するたびにボードレイアウトを表示します。
以降、新しい枝刈りポイントを発見するためにも使います。
board[]を使ってQの位置を確認するのに役立ちます。

2.symmetryOps()
対象解除を行うロジックメソッドです。
以下のロジックが同梱されています。

1.クイーンが右上角にある場合、ユニーク解が属するグループの要素数は必ず8個(=2×4)
2.クイーンが右上角以外にある場合、

(1) 90度回転させてオリジナルと同型になる場合、さらに90度回転(オリジナルから180
度回転)させても、さらに90度回転(オリジナルから270度回転)させてもオリジナルと同
型になる。
こちらに該当するユニーク解が属するグループの要素数は、左右反転させたパター ンを
加えて2個しかありません。
(2) 90度回転させてオリジナルと異なる場合は、270度回転させても必ずオリジナルとは
  異なる。ただし、180度回転させた場合はオリジナルと同型になることも有り得る。
こちらに該当するユニーク解が属するグループの要素数は、180度回転させて同型になる
  場合は4個(左右反転×縦横回転)
(3)180度回転させてもオリジナルと異なる場合は、8個(左右反転×縦横回転×上下反転)

3.backTrack()
再帰ロジックで作られています。
右上角にQが配置されているかいないかで判定処理が分岐されます。
その判定処理に従って、見合った枝刈りが行われます。

4.symmetry()
本体メソッドです。
こちらのメソッドでは、Qが角に配置されている場合、そうでない場合を分岐して3のbackTrack()に渡します。

symmetry()メソッド

角にQがある時の処理

symmetry()メソッドは、Qが角にあるかどうかを判定し、適切にbackTrack()に処理を渡します。

では上から順番に見ていきます。
とても大切なことなのですが、グローバル変数の初期化を行います。

  TOTAL=UNIQUE=COUNT2=COUNT4=COUNT8=0;

角にQがある時のボードのイメージは以下のとおりです。

  角にQがある時の処理
    +-+-+-+-+-+  
    | | | | |Q| 
    +-+-+-+-+-+  
    | | | | | | 
    +-+-+-+-+-+  
    | | | | | |       
    +-+-+-+-+-+             
    | | | | | | 
    +-+-+-+-+-+ 
    | | | | | |      
    +-+-+-+-+-+  

MASKについてですが、こちらはrow全体にビットを立てて配置済みとするフィルタのことです。

N5の場合、 10進数では31、2進数では 11111 となります。

MASK=(1<<size)-1;
    +-+-+-+-+-+  
    |M|M|M|M|M|  MASK=(1<<size)-1
    +-+-+-+-+-+  $(( (1<<5)-1 )):31
    | | | | | |  $(( 2#11111 )) 
    +-+-+-+-+-+  $ 31
    | | | | | |       
    +-+-+-+-+-+             
    | | | | | | 
    +-+-+-+-+-+ 
    | | | | | |      
    +-+-+-+-+-+  

bashでは以下のようにすると2進数を10進数に変換して表示することができます。

$ echo $(( 2#11111 ))
$ 31

TOPBITは左上角にビットを立てたものです。

  TOPBIT=1<<(size-1);
    +-+-+-+-+-+  
    |T| | | | |  TOPBIT=1<<(size-1);
    +-+-+-+-+-+  $(( (1<<(5-1) )):16 
    | | | | | |  $(( 2#10000 )) 
    +-+-+-+-+-+  $ 16 
    | | | | | |       
    +-+-+-+-+-+             
    | | | | | | 
    +-+-+-+-+-+ 
    | | | | | |      
    +-+-+-+-+-+  

ENDBITはTOPBITの1つ右にビットを立てたものです。現在はまだ初期値0のままです。
後述しますがENDBITは、

ENDBIT=TOPBIT>>1;

となります。

  ENDBIT=LASTMASK=SIDEMASK=0;
    は何も置かない(なにもはじまってない)
    +-+-+-+-+-+  
    |T|E| | | |  0#01000
    +-+-+-+-+-+  
    | | | | | | 
    +-+-+-+-+-+  
    | | | | | |       
    +-+-+-+-+-+             
    | | | | | | 
    +-+-+-+-+-+ 
    | | | | | |      
    +-+-+-+-+-+  

配列BOUND1とBOUND2を初期化して
Qを右上角に配置します。

  BOUND1=2; 
  BOUND2=0;
    +-+-+-+-+-+  右角から1つ左から開始
    | | | |B| |  BOUND1は右から左へシフト
    +-+-+-+-+-+   
    | | | | | | 
    +-+-+-+-+-+  
    | | | | | |       
    +-+-+-+-+-+             
    | | | | | | 
    +-+-+-+-+-+ 
    | | | | | |      
    +-+-+-+-+-+  
    +-+-+-+-+-+  左角から1つ右から開始
    | |B| | | |  BOUND2は左から右へシフトします。
    +-+-+-+-+-+  
    | | | | | | 
    +-+-+-+-+-+  
    | | | | | |       
    +-+-+-+-+-+             
    | | | | | | 
    +-+-+-+-+-+ 
    | | | | | |      
    +-+-+-+-+-+  

ここはすごく重要で、row[0]の1というのは右上の角のことを意味します。ですので、「Qを右上角に配置した」ということになります。

  board[0]=1;       # Qを右上角に配置

ここで、rowは2行目の処理に入ります。
というのも、1行目は右上角にQを配置したので次に進んだわけです。

その後、以下の行は「2行目の真ん中にQを配置する」という意味となります。

      bit=1<<BOUND1;

イメージでは以下のとおりです。

       の場合、BOUND1は2と3になる
        1<<BOUND1  行目は真ん中に置く
        +-+-+-+-+-+  
        | | | | |Q|  1#00001
        +-+-+-+-+-+  
        | | |B| | |  BOUND1=2 4#00100
        +-+-+-+-+-+  $(( 1<<2 ))
        | | | | | |  $ 4
        +-+-+-+-+-+  1->2->4               
        | | | | | | 
        +-+-+-+-+-+ 
        | | | | | |      
        +-+-+-+-+-+  

さらに、処理の中でBOUND1は左に1つシフトします。

        1<<BOUND1  行目は真ん中の次に置く
        +-+-+-+-+-+  
        | | | | |Q|  1#00001
        +-+-+-+-+-+  
        | |B| | | |  BOUND1=3 8#01000
        +-+-+-+-+-+  $(( 1<<3 ))
        | | | | | |  $ 8
        +-+-+-+-+-+  1->2->4->8             
        | | | | | | 
        +-+-+-+-+-+ 
        | | | | | |      
        +-+-+-+-+-+  

角にQがない時の処理

TOPBITは先程も書いたとおり、左上角にビットを立てた状態を示します。

  TOPBIT=1<<(size-1)
    +-+-+-+-+-+  
    |T| | | | | TOPBIT=1<<(size-1)
    +-+-+-+-+-+ $(( 2#10000 )) 
    | | | | | | $ 16 
    +-+-+-+-+-+  
    | | | | | |       
    +-+-+-+-+-+             
    | | | | | | 
    +-+-+-+-+-+ 
    | | | | | |      
    +-+-+-+-+-+  

ENDBITは、TOPBITの1つ右にビットを立てた状態です。

  ENDBIT=TOPBIT>>1
    +-+-+-+-+-+  
    |T|E| | | | ENDBIT=TOPBIT>>1
    +-+-+-+-+-+ $(( 2#01000 )) 
    | | | | | | $ 8
    +-+-+-+-+-+  
    | | | | | |       
    +-+-+-+-+-+             
    | | | | | | 
    +-+-+-+-+-+ 
    | | | | | |      
    +-+-+-+-+-+  

SIDEMASKは行の両サイドの角にビットを立てた状態を示します。

  SIDEMASK=TOPBIT|1
    +-+-+-+-+-+  
    |S| | | |S| SIDEMASK=TOPBIT|1
    +-+-+-+-+-+ $(( 2#10001 )) 
    | | | | | | $ 17
    +-+-+-+-+-+  
    | | | | | |       
    +-+-+-+-+-+             
    | | | | | | 
    +-+-+-+-+-+ 
    | | | | | |      
    +-+-+-+-+-+  

LASTMASKもSIDEMASK同様左右の角のビットを立てた状態を示します。ただ使いみちとしては、上記の状態でも使えるわけですが、

  LASTMASK=TOPBIT|1
    +-+-+-+-+-+  
    |L| | | |L| LASTMASK=TOPBIT|1
    +-+-+-+-+-+ $(( 2#10001 )) 
    | | | | | | $ 17
    +-+-+-+-+-+  
    | | | | | |       
    +-+-+-+-+-+             
    | | | | | | 
    +-+-+-+-+-+ 
    | | | | | |      
    +-+-+-+-+-+  

こういうシチュエーションで使います。

  LASTMASK=TOPBIT|1
    +-+-+-+-+-+  
    | | | | | | LASTMASK=TOPBIT|1
    +-+-+-+-+-+ $(( 2#10001 )) 
    | | | | | | $ 17
    +-+-+-+-+-+  
    | | | | | |       
    +-+-+-+-+-+             
    | | | | | | 
    +-+-+-+-+-+ 
    |L| | | |L|      
    +-+-+-+-+-+  

ENDBITは左角の1つ手前から右へシフトするビットです。

    ENDBIT=ENDBIT>>1
      +-+-+-+-+-+  
      |T|E| | | | ENDBIT=TOPBIT>>1
      +-+-+-+-+-+ $(( 2#01000 )) 
      | | | | | | $ 8
      +-+-+-+-+-+  
      | | | | | |       
      +-+-+-+-+-+             
      | | | | | | 
      +-+-+-+-+-+ 
      | | | | | |      
      +-+-+-+-+-+  
           
      +-+-+-+-+-+  
      |T| |E| | | ENDBIT=ENDBIT>>1
      +-+-+-+-+-+ $(( 2#00100 )) 
      | | | | | | $ 4
      +-+-+-+-+-+  
      | | | | | |       
      +-+-+-+-+-+             
      | | | | | | 
      +-+-+-+-+-+ 
      | | | | | |      
      +-+-+-+-+-+  

LASTMASKに代入されるビットは以下のイメージとなります。使いみちはLASTMASKと似ていますが、

    +-+-+-+-+-+  
    |L|L| |L|L| LASTMASK=LASTMASK<<1 | LASTMASK | LASTMASK>>1
    +-+-+-+-+-+ $(( 2#11011 )) 
    | | | | | | $ 27
    +-+-+-+-+-+  
    | | | | | |       
    +-+-+-+-+-+             
    | | | | | | 
    +-+-+-+-+-+ 
    | | | | | |      
    +-+-+-+-+-+  

こういったシチュエーションにも使います。

    +-+-+-+-+-+  
    | | | | | | LASTMASK=LASTMASK<<1 | LASTMASK | LASTMASK>>1
    +-+-+-+-+-+ $(( 2#11011 )) 
    | | | | | | $ 27
    +-+-+-+-+-+  
    | | | | | |       
    +-+-+-+-+-+             
    | | | | | | 
    +-+-+-+-+-+ 
    |L|L| |L|L|      
    +-+-+-+-+-+  

backTrack()メソッド

角にQがある時の処理

このメソッドでは、角にQがある場合の処理、ない場合の処理の分岐があり、その分岐の中で、((COUNT8++))であったり、symmetryOpsといった対象解除の処理へ進んだり、はたまた「枝刈り」をしたりと、濃厚なメソッドです。

では、順番に見ていきます。

ここは、Qが角にある場合の処理で、rowが最終行までたどり着いた(配置できた)場合の処理となります。グループの要素数は必ず8あるわけですので、あとから8倍するCOUNT8をインクリメントします。

      .クイーンが右上角にある場合、ユニーク解が属する
      グループの要素数は必ず(×)
      if(row==(size-1)){
        if(bitmap){
          board[row]=bitmap;
          COUNT8++;
        }
      }else{
  次は、Qが角にある場合の処理で、
  上から2行目のクイーンの位置が左から何番目にあるかと、
  右から2列目のクイーンの位置が上から何番目にあるかを、
  比較するだけで判定します。
  具体的には、2行目と2列目の位置を数値とみなし、
  2行目<2列目という条件を課せばよい
  という枝仮ルールを実現している処理となります。
			}else{
				if(row<BOUND1){   //枝刈り
					bitmap=bitmap|2;
					bitmap=bitmap^2;
				}
				while(bitmap){
        : '
        上から行目のクイーンの位置が左から何番目にあるかと、
        右から列目のクイーンの位置が上から何番目にあるかを、
        比較するだけで判定します。
        具体的には、行目と列目の位置を数値とみなし、
        行目<列目という条件を課せばよい
        結論: 以下の図では、を枝刈りを入れる
          
          +-+-+-+-+-+  
          | | | |X|Q| 
          +-+-+-+-+-+  
          | |Q| |X| |  8(左から数えて
          +-+-+-+-+-+  
          | | | |X| |       
          +-+-+-+-+-+             
          | | | |Q| |  8(上から数えて 
          +-+-+-+-+-+ 
          | | | | | |      
          +-+-+-+-+-+  
        ';
      fi

たったこれだけの2行ですごいですね。

         bitmap=bitmap|2;
         bitmap=bitmap^2;

角にQがない時の処理

ここは少しわかりにくいですね。丁寧に説明文を入れましたので、じっくりと理解してみてください。

    }else{                            # Qが角にない
      : '
      オリジナルがユニーク解であるためには先ず、
      前提:symmetryOpsは回転・鏡像変換により得られる状態の
      ユニーク値を比較し最小のものだけがユニーク解となるようにしている。
      Qができるだけ右に置かれている方がユニーク値は小さい
      例えば1行目の2列目にQが置かれている方が3列目に置かれているより
      ユニーク値は小さくユニーク解に近い。
      1行目のクイーンの位置が同じなら2行目のクイーンの位置がより右の
      列におかれているものがユニーク値は小さくユニーク解に近い。

      下図の X への配置は禁止されます。
      Qの位置より右位置の8対象位置X)にクイーンを置くことはできない。
      置いた場合回転・鏡像変換したユニーク値が最小にならなくなり、symmetryOps
      で負けるので枝刈りをする


      1行目のクイーンが3列目に置かれている場合
      +-+-+-+-+-+-+-+-+  
      |X|X| | | |Q|X|X| 
      +-+-+-+-+-+-+-+-+  
      |X| | | | | | |X| 
      +-+-+-+-+-+-+-+-+  
      | | | | | | | | |       
      +-+-+-+-+-+-+-+-+             
      | | | | | | | | | 
      +-+-+-+-+-+-+-+-+ 
      | | | | | | | | |      
      +-+-+-+-+-+-+-+-+  
      | | | | | | | | |
      +-+-+-+-+-+-+-+-+
      |X| | | | | | |X|
      +-+-+-+-+-+-+-+-+
      |X|X| | | | |X|X|
      +-+-+-+-+-+-+-+-+

      1行目のクイーンが4列目に置かれている場合
      +-+-+-+-+-+-+-+-+  
      |X|X|X| |Q|X|X|X| 
      +-+-+-+-+-+-+-+-+  
      |X| | | | | | |X| 
      +-+-+-+-+-+-+-+-+  
      |X| | | | | | |X|       
      +-+-+-+-+-+-+-+-+             
      | | | | | | | | | 
      +-+-+-+-+-+-+-+-+ 
      | | | | | | | | |      
      +-+-+-+-+-+-+-+-+  
      |X| | | | | | |X|
      +-+-+-+-+-+-+-+-+
      |X| | | | | | |X|
      +-+-+-+-+-+-+-+-+
      |X|X|X| | |X|X|X|
      +-+-+-+-+-+-+-+-+

      プログラムではこの枝刈を上部サイド枝刈り、下部サイド枝刈り、最下段枝刈り
      3か所で行っている
      それぞれ、1,2,3の数字で表すと以下の通り

      1行目のクイーンが3列目に置かれている場合
      +-+-+-+-+-+-+-+-+  
      |X|X| | | |Q|X|X| 
      +-+-+-+-+-+-+-+-+  
      |1| | | | | | |1| 
      +-+-+-+-+-+-+-+-+  
      | | | | | | | | |       
      +-+-+-+-+-+-+-+-+             
      | | | | | | | | | 
      +-+-+-+-+-+-+-+-+ 
      | | | | | | | | |      
      +-+-+-+-+-+-+-+-+  
      | | | | | | | | |
      +-+-+-+-+-+-+-+-+
      |2| | | | | | |2|
      +-+-+-+-+-+-+-+-+
      |2|3| | | | |3|2|
      +-+-+-+-+-+-+-+-+
      1行目にXが残っているが当然Qの効き筋なので枝刈する必要はない
      ';
			if(row<BOUND1){
				bitmap=bitmap|SIDEMASK;
				bitmap=bitmap^SIDEMASK;
			}else{
				if(row==BOUND2){
					if((down&SIDEMASK)==0){
						return;
					}
					if( (down&SIDEMASK)!=SIDEMASK){
						bitmap=bitmap&SIDEMASK;
					}
				}
			}

symmetryOps()メソッド

ここもロジックを理解するところから初めて見ると良いと思います。
あまり細かいことを考えるよりも、大枠でどんなときにこのメソッドが呼ばれ、カウンターがいつインクリメントするのか、などを掴んだほうが良いと思います。

symmetryOps()では以下を実現しています。

2.クイーンが右上角以外にある場合、
(1) 90度回転させてオリジナルと同型になる場合、さらに90度回転(オリジナルか
ら180度回転)させても、さらに90度回転(オリジナルから270度回転)させてもオリ
ジナルと同型になる。
こちらに該当するユニーク解が属するグループの要素数は、左右反転させたパター
ンを加えて2個しかありません。
2.クイーンが右上角以外にある場合、
(2) 90度回転させてオリジナルと異なる場合は、270度回転させても必ずオリジナル
とは異なる。ただし、180度回転させた場合はオリジナルと同型になることも有り得
る。こちらに該当するユニーク解が属するグループの要素数は、180度回転させて同
型になる場合は4個(左右反転×縦横回転)
2.クイーンが右上角以外にある場合、
(3)180度回転させてもオリジナルと異なる場合は、8個(左右反転×縦横回転×上下反転)

クイーンの利き筋を辿っていくと、オリジナルがユニーク解ではない可能性があり、
それは下図の A,B,C の位置のどこかにクイーンがある場合に限られます。
symmetryOpsは、以下の図のA,B,CにQが置かれた場合にユニーク解かを判定します。
原型と、90,180,270回転させたもののユニーク値を比較します。

     0 1 2 3 4 
    +-+-+-+-+-+  
    | | | | |Q|    4
    +-+-+-+-+-+  
    | | |Q| | |    2 
    +-+-+-+-+-+  
    |Q| | | | |    0  ----> 4 2 0 3 1 (ユニーク判定値)
    +-+-+-+-+-+             数が大きい方をユニークとみなす 
    | | | |Q| |    3
    +-+-+-+-+-+  
    | |Q| | | |    1 
    +-+-+-+-+-+  

     0 1 2 3 4     左右反転!
    +-+-+-+-+-+  
    |Q| | | | |    0
    +-+-+-+-+-+  
    | | |Q| | |    2 
    +-+-+-+-+-+  
    | | | | |Q|    4  ----> 0 2 4 1 3 
    +-+-+-+-+-+            数が小さいのでユニーク解とはしません 
    | |Q| | | |    1
    +-+-+-+-+-+  
    | | | |Q| |    3 
    +-+-+-+-+-+  

Qができるだけ右に置かれている方がユニーク値は大きくなります。
例えば1行目の2列目にQが置かれている方が、
3列目に置かれているよりユニーク値は大きくユニーク解に近い。
1行目のクイーンの位置が同じなら2行目のクイーンの位置がより右の列におかれてい
るものがユニーク値は大きくユニーク解に近くなります。
それ以外はユニーク解なのでCOUNT8にする

   +-+-+-+-+-+-+-+-+  
   |X|X| | | |Q|X|X| 
   +-+-+-+-+-+-+-+-+  
   |X| | | |x|x|x|X| 
   +-+-+-+-+-+-+-+-+  
   |C| | |x| |x| |x|       
   +-+-+-+-+-+-+-+-+             
   | | |x| | |x| | | 
   +-+-+-+-+-+-+-+-+ 
   | |x| | | |x| | |      
   +-+-+-+-+-+-+-+-+  
   |x| | | | |x| |A|
   +-+-+-+-+-+-+-+-+
   |X| | | | |x| |X|
   +-+-+-+-+-+-+-+-+
   |X|X|B| | |x|X|X|
   +-+-+-+-+-+-+-+-+

Aの場合 右90度回転 board[BOUND2]==1
Bの場合 右180度回転 board[size-1]==ENDBIT
Cの場合 右270度回転 board[BOUND1]==TOPBIT

CUDA ノードレイヤー

・std::vector kLayer_nodeLayer(unsigned int size,unsigned int k,std::vector& L)
5行分backtrack1,2を実行し、実行結果をnodes,Lに格納する
Lは、vector 構造体 l を格納する
lはBOUND1,BOUND2,TOPBIT,ENDBIT,SIDEMASK,LASTMASK,board[MAX]を格納する
(COUNT2,4,8 TOTAL,UNIQUEは今回は不要)

CUDA ソースコード

ソースコードは以下のとおりです。

/**
 *
 * bash版対称解除法のC言語版のGPU/CUDA移植版
 *
 詳しい説明はこちらをどうぞ
 https://suzukiiichiro.github.io/search/?keyword=Nクイーン問題

非再帰でのコンパイルと実行
$ nvcc -O3 -arch=sm_61 03CUDA_Symmetry_NodeLayer.cu && ./a.out -c

再帰でのコンパイルと実行
$ nvcc -O3 -arch=sm_61 03CUDA_Symmetry_NodeLayer.cu && ./a.out -r

GPU で並列処理せずに実行
$ nvcc -O3 -arch=sm_61 03CUDA_Symmetry_NodeLayer.cu && ./a.out -n

GPU で並列処理で実行(ノードレイヤー)
$ nvcc -O3 -arch=sm_61 03CUDA_Symmetry_NodeLayer.cu && ./a.out -n

対称解除法 GPUノードレイヤー
 N:        Total      Unique      dd:hh:mm:ss.ms
 4:            0           0      00:00:00:00.13
 5:            0           0      00:00:00:00.00
 6:            0           0      00:00:00:00.00
 7:           40           0      00:00:00:00.00
 8:           92           0      00:00:00:00.00
 9:          352           0      00:00:00:00.00
10:          724           0      00:00:00:00.00
11:         2680           0      00:00:00:00.00
12:        14200           0      00:00:00:00.00
13:        73712           0      00:00:00:00.03
14:       365596           0      00:00:00:00.15
15:      2279184           0      00:00:00:00.86
16:     14772512           0      00:00:00:05.50
17:     95815104           0      00:00:00:38.87
18:    666090624           0      00:00:04:46.16

以降はバーストします。

コメント
・std::vector<long> kLayer_nodeLayer(unsigned int size,unsigned int k,std::vector<local>& L)
 5行分backtrack1,2を実行し、実行結果をnodes,Lに格納する
 Lは、vector 構造体 l を格納する
 lはBOUND1,BOUND2,TOPBIT,ENDBIT,SIDEMASK,LASTMASK,board[MAX]を格納する
 (COUNT2,4,8 TOTAL,UNIQUEは今回は不要)

*
*/
#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 MAX 27
#define THREAD_NUM (1<<MAX)-1
using std::cout; using std::endl;
using std::vector; using std::string;
// システムによって以下のマクロが必要であればコメントを外してください。
//#define UINT64_C(c) c ## ULL
//
// グローバル変数
unsigned long TOTAL=0;
unsigned long UNIQUE=0;
//GPU で使うローカル構造体
typedef struct local
{
  unsigned int BOUND1,BOUND2;
  unsigned int TOPBIT,ENDBIT,SIDEMASK,LASTMASK;
  unsigned long board[MAX];
  unsigned long COUNT2,COUNT4,COUNT8,TOTAL,UNIQUE;
  unsigned long TYPE;
}local;
// CPU 再帰/非再帰共通 対称解除法
void symmetryOps(unsigned int size,struct local* l)
{
  /**
  2.クイーンが右上角以外にある場合、
  (1) 90度回転させてオリジナルと同型になる場合、さらに90度回転(オリジナルか
  ら180度回転)させても、さらに90度回転(オリジナルから270度回転)させてもオリ
  ジナルと同型になる。
  こちらに該当するユニーク解が属するグループの要素数は、左右反転させたパター
  ンを加えて2個しかありません。
  */
  if(l->board[l->BOUND2]==1){
    unsigned int ptn;
    unsigned int own;
    for(ptn=2,own=1;own<size;++own,ptn<<=1){
      unsigned int bit;
      unsigned int you;
      for(bit=1,you=size-1;(l->board[you]!=ptn)&&l->board[own]>=bit;--you){
        bit<<=1;
      }
      if(l->board[own]>bit){
        return ;
      }
      if(l->board[own]<bit){
        break;
      }
    }//end for
    // 90度回転して同型なら180度回転しても270度回転しても同型である
    if(own>size-1){
      l->COUNT2++;
      return ;
    }//end if
  }//end if
  /**
  2.クイーンが右上角以外にある場合、
    (2) 90度回転させてオリジナルと異なる場合は、270度回転させても必ずオリジナル
    とは異なる。ただし、180度回転させた場合はオリジナルと同型になることも有り得
    る。こちらに該当するユニーク解が属するグループの要素数は、180度回転させて同
    型になる場合は4個(左右反転×縦横回転)
   */
  //180度回転
  if(l->board[size-1]==l->ENDBIT){
    unsigned int you;
    unsigned int own;
    for(you=size-1-1,own=1;own<=size-1;++own,--you){
      unsigned int bit;
      unsigned int ptn;
      for(bit=1,ptn=l->TOPBIT;(ptn!=l->board[you])&&(l->board[own]>=bit);ptn>>=1){
        bit<<=1;
      }
      if(l->board[own]>bit){
        return ;
      }
      if(l->board[own]<bit){
        break;
      }
    }//end for
    //90度回転が同型でなくても180度回転が同型であることもある
    if(own>size-1){
      l->COUNT4++;
      return ;
    }
  }//end if
  /**
  2.クイーンが右上角以外にある場合、
    (3)180度回転させてもオリジナルと異なる場合は、8個(左右反転×縦横回転×上下反転)
  */
  //270度回転
  if(l->board[l->BOUND1]==l->TOPBIT){
    unsigned int ptn;
    unsigned int own;
    unsigned int you;
    unsigned int bit;
    for(ptn=l->TOPBIT>>1,own=1;own<=size-1;++own,ptn>>=1){
      for(bit=1,you=0;(l->board[you]!=ptn)&&(l->board[own]>=bit);++you){
        bit<<=1;
      }
      if(l->board[own]>bit){
        return ;
      }
      if(l->board[own]<bit){
        break;
      }
    }//end for
  }//end if
  l->COUNT8++;
}
/**
  CPU -c
  */
// 非再帰 角にQがないときのバックトラック
void symmetry_backTrack_NR(unsigned int size,unsigned int row,unsigned int _left,unsigned int _down,unsigned int _right,struct local *l)
{
  unsigned int mask=(1<<size)-1;
  unsigned int down[size];
  unsigned int left[size];
  unsigned int right[size];
  unsigned int bitmap[size];
  left[row]=_left;
  down[row]=_down;
  right[row]=_right;
  bitmap[row]=mask&~(left[row]|down[row]|right[row]);
  while(row>0){
    if(bitmap[row]>0){
      if(row<l->BOUND1){ //上部サイド枝刈り
        bitmap[row]|=l->SIDEMASK;
        bitmap[row]^=l->SIDEMASK;
      }else if(row==l->BOUND2){ //下部サイド枝刈り
        if((down[row]&l->SIDEMASK)==0){
          row--; 
        }
        if((down[row]&l->SIDEMASK)!=l->SIDEMASK){
          bitmap[row]&=l->SIDEMASK;
        }
      }
      unsigned int save_bitmap=bitmap[row];
      unsigned int bit=-bitmap[row]&bitmap[row];
      bitmap[row]^=bit;
      l->board[row]=bit; //Qを配置
      if((bit&mask)!=0){
        if(row==(size-1)){
          if( (save_bitmap&l->LASTMASK)==0){
            symmetryOps(size,l);  //対称解除法
          }
          row--;
        }else{
          unsigned int n=row++;
          left[row]=(left[n]|bit)<<1;
          down[row]=(down[n]|bit);
          right[row]=(right[n]|bit)>>1;
          bitmap[row]=mask&~(left[row]|down[row]|right[row]);
        }
      }else{
        row--;
      }
    }else{
      row--;
    }
  }//end while
}
// 非再帰 角にQがあるときのバックトラック
void symmetry_backTrack_corner_NR(unsigned int size,unsigned int row,unsigned int _left,unsigned int _down,unsigned int _right,struct local *l)
{
  unsigned int mask=(1<<size)-1;
  unsigned int bit=0;
  unsigned int down[size];
  unsigned int left[size];
  unsigned int right[size];
  unsigned int bitmap[size];
  left[row]=_left;
  down[row]=_down;
  right[row]=_right;
  bitmap[row]=mask&~(left[row]|down[row]|right[row]);
  while(row>=2){
    if(row<l->BOUND1){
      // bitmap[row]=bitmap[row]|2;
      // bitmap[row]=bitmap[row]^2;
      bitmap[row]&=~2;
    }
    if(bitmap[row]>0){
      bit=-bitmap[row]&bitmap[row];
      bitmap[row]^=bit;
      if(row==(size-1)){
        l->COUNT8++;
        row--;
      }else{
        unsigned int n=row++;
        left[row]=(left[n]|bit)<<1;
        down[row]=(down[n]|bit);
        right[row]=(right[n]|bit)>>1;
        l->board[row]=bit; //Qを配置
        //クイーンが配置可能な位置を表す
        bitmap[row]=mask&~(left[row]|down[row]|right[row]);
      }
    }else{
      row--;
    }
  }//end while
}
// 非再帰 対称解除法
void symmetry_NR(unsigned int size,struct local* l)
{
  l->TOTAL=l->UNIQUE=l->COUNT2=l->COUNT4=l->COUNT8=0;
  unsigned int bit=0;
  l->TOPBIT=1<<(size-1);
  l->ENDBIT=l->SIDEMASK=l->LASTMASK=0;
  l->BOUND1=2;
  l->BOUND2=0;
  l->board[0]=1;
  while(l->BOUND1>1&&l->BOUND1<size-1){
    if(l->BOUND1<size-1){
      bit=1<<l->BOUND1;
      l->board[1]=bit;   //2行目にQを配置
      //角にQがあるときのバックトラック
      symmetry_backTrack_corner_NR(size,2,(2|bit)<<1,1|bit,(2|bit)>>1,l);
    }
    l->BOUND1++;
  }
  l->TOPBIT=1<<(size-1);
  l->ENDBIT=l->TOPBIT>>1;
  l->SIDEMASK=l->TOPBIT|1;
  l->LASTMASK=l->TOPBIT|1;
  l->BOUND1=1;
  l->BOUND2=size-2;
  while(l->BOUND1>0 && l->BOUND2<size-1 && l->BOUND1<l->BOUND2){
    if(l->BOUND1<l->BOUND2){
      bit=1<<l->BOUND1;
      l->board[0]=bit;   //Qを配置
      //角にQがないときのバックトラック
      symmetry_backTrack_NR(size,1,bit<<1,bit,bit>>1,l);
    }
    l->BOUND1++;
    l->BOUND2--;
    l->ENDBIT=l->ENDBIT>>1;
    l->LASTMASK=l->LASTMASK<<1|l->LASTMASK|l->LASTMASK>>1;
  }//ene while
  UNIQUE=l->COUNT2+l->COUNT4+l->COUNT8;
  TOTAL=l->COUNT2*2+l->COUNT4*4+l->COUNT8*8;
}
/**
  CPU -r
  */
// 再帰 角にQがないときのバックトラック
void symmetry_backTrack(unsigned int size,unsigned int row,unsigned int left,unsigned int down,unsigned int right,struct local* l)
{
  unsigned int mask=(1<<size)-1;
  unsigned int bitmap=mask&~(left|down|right);
  if(row==(size-1)){
    if(bitmap){
      if( (bitmap&l->LASTMASK)==0){
        l->board[row]=bitmap;  //Qを配置
        symmetryOps(size,l);    //対称解除
      }
    }
  }else{
    if(row<l->BOUND1){
      bitmap=bitmap|l->SIDEMASK;
      bitmap=bitmap^l->SIDEMASK;
    }else{
      if(row==l->BOUND2){
        if((down&l->SIDEMASK)==0){
          return;
        }
        if( (down&l->SIDEMASK)!=l->SIDEMASK){
          bitmap=bitmap&l->SIDEMASK;
        }
      }
    }
    while(bitmap){
      unsigned int bit=-bitmap&bitmap;
      bitmap=bitmap^bit;
      l->board[row]=bit;
      symmetry_backTrack(size,row+1,(left|bit)<<1,down|bit,(right|bit)>>1,l);
    }
  }
}
// 再帰 角にQがあるときのバックトラック
void symmetry_backTrack_corner(unsigned int size,unsigned int row,unsigned int left,unsigned int down,unsigned int right,struct local* l)
{
  unsigned int mask=(1<<size)-1;
  unsigned int bitmap=mask&~(left|down|right);
  unsigned int bit=0;
  if(row==(size-1)){
    if(bitmap){
      l->board[row]=bitmap;
      l->COUNT8++;
    }
  }else{
    if(row<l->BOUND1){   //枝刈り
      bitmap=bitmap|2;
      bitmap=bitmap^2;
    }
    while(bitmap){
      bit=-bitmap&bitmap;
      bitmap=bitmap^bit;
      l->board[row]=bit;   //Qを配置
      symmetry_backTrack_corner(size,row+1,(left|bit)<<1,down|bit,(right|bit)>>1,l);
    }
  }
}
// 再帰 対称解除法
void symmetry_R(unsigned int size,struct local* l)
{
  l->TOTAL=l->UNIQUE=l->COUNT2=l->COUNT4=l->COUNT8=0;
  unsigned int bit=0;
  l->TOPBIT=1<<(size-1);
  l->ENDBIT=l->LASTMASK=l->SIDEMASK=0;
  l->BOUND1=2;
  l->BOUND2=0;
  l->board[0]=1;
  while(l->BOUND1>1 && l->BOUND1<size-1){
    if(l->BOUND1<size-1){
      bit=1<<l->BOUND1;
      l->board[1]=bit;   //2行目にQを配置
      //角にQがあるときのバックトラック
      symmetry_backTrack_corner(size,2,(2|bit)<<1,1|bit,(2|bit)>>1,l);
    }
    l->BOUND1++;
  }//end while
  l->TOPBIT=1<<(size-1);
  l->ENDBIT=l->TOPBIT>>1;
  l->SIDEMASK=l->TOPBIT|1;
  l->LASTMASK=l->TOPBIT|1;
  l->BOUND1=1;
  l->BOUND2=size-2;
  while(l->BOUND1>0 && l->BOUND2<size-1 && l->BOUND1<l->BOUND2){
    if(l->BOUND1<l->BOUND2){
      bit=1<<l->BOUND1;
      l->board[0]=bit;   //Qを配置
      //角にQがないときのバックトラック
      symmetry_backTrack(size,1,bit<<1,bit,bit>>1,l);
    }
    l->BOUND1++;
    l->BOUND2--;
    l->ENDBIT=l->ENDBIT>>1;
    l->LASTMASK=l->LASTMASK<<1|l->LASTMASK|l->LASTMASK>>1;
  }//ene while
  UNIQUE=l->COUNT2+l->COUNT4+l->COUNT8;
  TOTAL=l->COUNT2*2+l->COUNT4*4+l->COUNT8*8;
}
/**
  GPU -g
  */
__device__
struct dlocal{
  unsigned int BOUND1,BOUND2;
  unsigned int TOPBIT,ENDBIT,SIDEMASK,LASTMASK;
  unsigned long board[MAX];
  unsigned long COUNT2,COUNT4,COUNT8,TOTAL,UNIQUE;
  unsigned long TYPE;
}dlocal;
__device__ struct dlocal gdl[9999];
// GPU 対称解除法
__host__ __device__
long GPU_symmetryOps(unsigned int size,struct dlocal* l)
{
  /**
  2.クイーンが右上角以外にある場合、
  (1) 90度回転させてオリジナルと同型になる場合、さらに90度回転(オリジナルか
  ら180度回転)させても、さらに90度回転(オリジナルから270度回転)させてもオリ
  ジナルと同型になる。
  こちらに該当するユニーク解が属するグループの要素数は、左右反転させたパター
  ンを加えて2個しかありません。
  */
  if(l->board[l->BOUND2]==1){
    unsigned int ptn;
    unsigned int own;
    for(ptn=2,own=1;own<size;++own,ptn<<=1){
      unsigned int bit;
      unsigned int you;
      for(bit=1,you=size-1;(l->board[you]!=ptn)&& l->board[own]>=bit;--you){
        bit<<=1;
      }
      if(l->board[own]>bit){
        return 0;
      }
      if(l->board[own]<bit){
        break;
      }
    }//end for
    // 90度回転して同型なら180度回転しても270度回転しても同型である
    if(own>size-1){
      l->COUNT2++;
      return 2;
    }//end if
  }//end if
  /**
  2.クイーンが右上角以外にある場合、
    (2) 90度回転させてオリジナルと異なる場合は、270度回転させても必ずオリジナル
    とは異なる。ただし、180度回転させた場合はオリジナルと同型になることも有り得
    る。こちらに該当するユニーク解が属するグループの要素数は、180度回転させて同
    型になる場合は4個(左右反転×縦横回転)
   */
  //180度回転
  if(l->board[size-1]==l->ENDBIT){
    unsigned int you;
    unsigned int own;
    for(you=size-1-1,own=1;own<=size-1;++own,--you){
      unsigned int bit;
      unsigned int ptn;
      for(bit=1,ptn=l->TOPBIT;(ptn!=l->board[you])&&(l->board[own]>=bit);ptn>>=1){
        bit<<=1;
      }
      if(l->board[own]>bit){
        return 0;
      }
      if(l->board[own]<bit){
        break;
      }
    }//end for
    //90度回転が同型でなくても180度回転が同型であることもある
    if(own>size-1){
      l->COUNT4++;
      return 4;
    }
  }//end if
  /**
  2.クイーンが右上角以外にある場合、
    (3)180度回転させてもオリジナルと異なる場合は、8個(左右反転×縦横回転×上下反転)
  */
  //270度回転
  if(l->board[l->BOUND1]==l->TOPBIT){
    unsigned int ptn;
    unsigned int own;
    unsigned int you;
    unsigned int bit;
    for(ptn=l->TOPBIT>>1,own=1;own<=size-1;++own,ptn>>=1){
      for(bit=1,you=0;(l->board[you]!=ptn)&&(l->board[own]>=bit);++you){
        bit<<=1;
      }
      if(l->board[own]>bit){
        return 0;
      }
      if(l->board[own]<bit){
        break;
      }
    }//end for
  }//end if
  l->COUNT8++;
  return 8;
}
// GPU 角にQがないときのバックトラック
__host__ __device__
long GPU_symmetry_backTrack(unsigned int size,unsigned int row,unsigned int left,unsigned int down,unsigned int right,struct dlocal* l)
{
  unsigned long counter=0;
  unsigned int mask=(1<<size)-1;
  unsigned int bitmap=mask&~(left|down|right);
  if(row==(size-1)){
    if(bitmap){
      if( (bitmap& l->LASTMASK)==0){
        l->board[row]=bitmap;  //Qを配置
        counter+=GPU_symmetryOps(size,l);    //対称解除
      }
    }
  }else{
    if(row<l->BOUND1){
      bitmap=bitmap|l->SIDEMASK;
      bitmap=bitmap^l->SIDEMASK;
    }else{
      if(row==l->BOUND2){
        if((down&l->SIDEMASK)==0){
          return 0;
        }
        if( (down&l->SIDEMASK)!=l->SIDEMASK){
          bitmap=bitmap&l->SIDEMASK;
        }
      }
    }
    while(bitmap){
      unsigned int bit=-bitmap&bitmap;
      bitmap=bitmap^bit;
      l->board[row]=bit;
      counter+=GPU_symmetry_backTrack(size,row+1,(left|bit)<<1,down|bit,(right|bit)>>1,l);
    }
  }
  return counter;
}
// GPU 角にQがあるときのバックトラック
__host__ __device__
long GPU_symmetry_backTrack_corner(unsigned int size,unsigned int row,unsigned int left,unsigned int down,unsigned int right,struct dlocal* l)
{
  unsigned long counter=0;
  unsigned int mask=(1<<size)-1;
  unsigned int bitmap=mask&~(left|down|right);
  unsigned int bit=0;
  if(row==(size-1)){
    if(bitmap){
      l->board[row]=bitmap;
      l->COUNT8++;
      counter+=8;
    }
  }else{
    if(row<l->BOUND1){   //枝刈り
      bitmap=bitmap|2;
      bitmap=bitmap^2;
    }
    while(bitmap){
      bit=-bitmap&bitmap;
      bitmap=bitmap^bit;
      l->board[row]=bit;   //Qを配置
      counter+=GPU_symmetry_backTrack_corner(size,row+1,(left|bit)<<1,down|bit,(right|bit)>>1,l);
    }
  }
  return counter;
}
// GPU 対称解除法 -g の実行時のみ呼び出されます
__host__ __device__
void GPU_symmetry_R(unsigned int size,struct local* hostLocal)
{
  // GPU内部で使うための dlocal構造体
  struct dlocal l;
  l.TOTAL=l.UNIQUE=l.COUNT2=l.COUNT4=l.COUNT8=0;
  unsigned int bit=0;
  l.TOPBIT=1<<(size-1);
  l.ENDBIT=l.LASTMASK=l.SIDEMASK=0;
  l.BOUND1=2;
  l.BOUND2=0;
  l.board[0]=1;
  while(l.BOUND1>1 && l.BOUND1<size-1){
    if(l.BOUND1<size-1){
      bit=1<<l.BOUND1;
      l.board[1]=bit;   //2行目にQを配置
      //角にQがあるときのバックトラック
      GPU_symmetry_backTrack_corner(size,2,(2|bit)<<1,1|bit,(2|bit)>>1,&l);
    }
    l.BOUND1++;
  }//end while
  l.TOPBIT=1<<(size-1);
  l.ENDBIT=l.TOPBIT>>1;
  l.SIDEMASK=l.TOPBIT|1;
  l.LASTMASK=l.TOPBIT|1;
  l.BOUND1=1;
  l.BOUND2=size-2;
  while(l.BOUND1>0 && l.BOUND2<size-1 && l.BOUND1<l.BOUND2){
    if(l.BOUND1<l.BOUND2){
      bit=1<<l.BOUND1;
      l.board[0]=bit;   //Qを配置
      //角にQがないときのバックトラック
      GPU_symmetry_backTrack(size,1,bit<<1,bit,bit>>1,&l);
    }
    l.BOUND1++;
    l.BOUND2--;
    l.ENDBIT=l.ENDBIT>>1;
    l.LASTMASK=l.LASTMASK<<1|l.LASTMASK|l.LASTMASK>>1;
  }//ene while
  // 集計値は hostLocalへ代入
  hostLocal->UNIQUE=l.COUNT2+l.COUNT4+l.COUNT8;
  hostLocal->TOTAL=l.COUNT2*2+l.COUNT4*4+l.COUNT8*8;
}
/**
  GPU -n
  */
// GPU -n の 対称解除法
__host__ __device__
long GPUN_symmetryOps(unsigned int size,struct local* l)
{
  /**
  2.クイーンが右上角以外にある場合、
  (1) 90度回転させてオリジナルと同型になる場合、さらに90度回転(オリジナルか
  ら180度回転)させても、さらに90度回転(オリジナルから270度回転)させてもオリ
  ジナルと同型になる。
  こちらに該当するユニーク解が属するグループの要素数は、左右反転させたパター
  ンを加えて2個しかありません。
  */
  if(l->board[l->BOUND2]==1){
    unsigned int ptn;
    unsigned int own;
    for(ptn=2,own=1;own<size;++own,ptn<<=1){
      unsigned int bit;
      unsigned int you;
      for(bit=1,you=size-1;(l->board[you]!=ptn)&& l->board[own]>=bit;--you){
        bit<<=1;
      }
      if(l->board[own]>bit){
        return 0;
      }
      if(l->board[own]<bit){
        break;
      }
    }//end for
    // 90度回転して同型なら180度回転しても270度回転しても同型である
    if(own>size-1){
      l->COUNT2++;
      return 2;
    }//end if
  }//end if
  /**
  2.クイーンが右上角以外にある場合、
    (2) 90度回転させてオリジナルと異なる場合は、270度回転させても必ずオリジナル
    とは異なる。ただし、180度回転させた場合はオリジナルと同型になることも有り得
    る。こちらに該当するユニーク解が属するグループの要素数は、180度回転させて同
    型になる場合は4個(左右反転×縦横回転)
   */
  //180度回転
  if(l->board[size-1]==l->ENDBIT){
    unsigned int you;
    unsigned int own;
    for(you=size-1-1,own=1;own<=size-1;++own,--you){
      unsigned int bit;
      unsigned int ptn;
      for(bit=1,ptn=l->TOPBIT;(ptn!=l->board[you])&&(l->board[own]>=bit);ptn>>=1){
        bit<<=1;
      }
      if(l->board[own]>bit){
        return 0;
      }
      if(l->board[own]<bit){
        break;
      }
    }//end for
    //90度回転が同型でなくても180度回転が同型であることもある
    if(own>size-1){
      l->COUNT4++;
      return 4;
    }
  }//end if
  /**
  2.クイーンが右上角以外にある場合、
    (3)180度回転させてもオリジナルと異なる場合は、8個(左右反転×縦横回転×上下反転)
  */
  //270度回転
  if(l->board[l->BOUND1]==l->TOPBIT){
    unsigned int ptn;
    unsigned int own;
    unsigned int you;
    unsigned int bit;
    for(ptn=l->TOPBIT>>1,own=1;own<=size-1;++own,ptn>>=1){
      for(bit=1,you=0;(l->board[you]!=ptn)&&(l->board[own]>=bit);++you){
        bit<<=1;
      }
      if(l->board[own]>bit){
        return 0;
      }
      if(l->board[own]<bit){
        break;
      }
    }//end for
  }//end if
  l->COUNT8++;
  return 8;
}
// 0以外のbitをカウント
__host__ __device__
unsigned int countBits_nodeLayer(unsigned long n)
{
  unsigned int counter=0;
  while(n){ n&=(n-1); counter++; }
  return counter;
}
// GPU -n ノードレイヤーによる対称解除法 -n の実行時に呼び出される
__host__ __device__ 
long GPU_symmetry_solve_nodeLayer_corner(unsigned int size,unsigned long left,unsigned long down,unsigned long right,struct local* l)
{
  unsigned long counter = 0;
  unsigned long mask=(1<<size)-1;
  unsigned long bitmap=mask&~(left|down|right);
  unsigned long bit=0;
  unsigned int row=countBits_nodeLayer(down);
  if(row==(size-1)){
    if(bitmap){
      l->board[row]=bitmap;
      return 8;
    }
  }else{
    if(row<l->BOUND1){   //枝刈り
      bitmap=bitmap|2;
      bitmap=bitmap^2;
    }
    while(bitmap){
      bit=-bitmap&bitmap;
      bitmap=bitmap^bit;
      l->board[row]=bit;   //Qを配置
      counter += GPU_symmetry_solve_nodeLayer_corner(size,(left|bit)<<1,(down|bit),(right|bit)>>1,l);
    }
  }
  return counter;
}
// GPU -n ノードレイヤーによる対称解除法 -n の実行時に呼び出される
__host__ __device__ 
long GPU_symmetry_solve_nodeLayer(unsigned int size,unsigned long left,unsigned long down,unsigned long right,struct local* l)
{
  unsigned long counter = 0;
  unsigned long mask=(1<<size)-1;
  unsigned long bitmap=mask&~(left|down|right);
  unsigned int row=countBits_nodeLayer(down);
  if(row==(size-1)){
    if(bitmap){
      if( (bitmap& l->LASTMASK)==0){
        l->board[row]=bitmap;  //Qを配置
        return GPUN_symmetryOps(size,l);    //対称解除;
      }
    }
  }else{
    if(row<l->BOUND1){
      bitmap=bitmap|l->SIDEMASK;
      bitmap=bitmap^l->SIDEMASK;
    }else{
      if(row==l->BOUND2){
        if((down&l->SIDEMASK)==0){
          return 0;
        }
        if( (down&l->SIDEMASK)!=l->SIDEMASK){
          bitmap=bitmap&l->SIDEMASK;
        }
      }
    }
    while(bitmap){
      unsigned long bit=-bitmap&bitmap;
      bitmap=bitmap^bit;
      l->board[row]=bit;
      counter+=GPU_symmetry_solve_nodeLayer(size,(left|bit)<<1,down|bit,(right|bit)>>1,l);
    }
  }
  return counter;
}
// GPU -n ノードレイヤー i 番目のメンバを i 番目の部分木の解で埋める
__global__ 
void dim_nodeLayer(unsigned int size,long* nodes,long* solutions,unsigned int numElements,struct local* l)
{
  unsigned int i=blockDim.x*blockIdx.x+threadIdx.x;
  if(i<numElements){
    if(l[i].TYPE==0){
      solutions[i]=GPU_symmetry_solve_nodeLayer_corner(size,nodes[3*i],nodes[3*i+1],nodes[3*i+2],&l[i]);
    }else{
      solutions[i]=GPU_symmetry_solve_nodeLayer(size,nodes[3*i],nodes[3*i+1],nodes[3*i+2],&l[i]);
    }
  }
}
// GPU -n Kレイヤー 再帰 角にQがないときのバックトラック
unsigned long kLayer_nodeLayer_backTrack(int size,std::vector<long>& nodes,unsigned int k,unsigned long left,unsigned long down,unsigned long right,std::vector<local>& L,struct local* l)
{
  unsigned long counter=0;
  unsigned long mask=(1<<size)-1;
  unsigned long bitmap=mask&~(left|down|right);
  unsigned int row=countBits_nodeLayer(down);
  if(row==k) {
      nodes.push_back(left);
      nodes.push_back(down);
      nodes.push_back(right);
      L.push_back(*l);
      return 1;
  }else{
    if(row<l->BOUND1){
      bitmap=bitmap|l->SIDEMASK;
      bitmap=bitmap^l->SIDEMASK;
    }else{
      if(row==l->BOUND2){
        if((down&l->SIDEMASK)==0){
          return 0;
        }
        if( (down&l->SIDEMASK)!=l->SIDEMASK){
          bitmap=bitmap&l->SIDEMASK;
        }
      }
    }
    while(bitmap){
      unsigned long bit=-bitmap&bitmap;
      bitmap=bitmap^bit;
      l->board[row]=bit;
      counter+=kLayer_nodeLayer_backTrack(size,nodes,k,(left|bit)<<1,(down|bit),(right|bit)>>1,L,l); 
    }
  }
  return counter;
}
// GPU -n Kレイヤー 角にQがあるときのバックトラック
unsigned long kLayer_nodeLayer_backTrack_corner(unsigned int size,std::vector<long>& nodes,unsigned int k,unsigned long left,unsigned long down,unsigned long right,std::vector<local>& L,struct local* l)
{
  unsigned long counter=0;
  unsigned long mask=(1<<size)-1;
  unsigned long bitmap=mask&~(left|down|right);
  unsigned long bit=0;
  int row=countBits_nodeLayer(down);
    if(row==k) {
      nodes.push_back(left);
      nodes.push_back(down);
      nodes.push_back(right);
      L.push_back(*l);
    }
    if(row<l->BOUND1){   //枝刈り
      bitmap=bitmap|2;
      bitmap=bitmap^2;
    }
    while(bitmap){
      bit=-bitmap&bitmap;
      bitmap=bitmap^bit;
      l->board[row]=bit;   //Qを配置
      counter+=kLayer_nodeLayer_backTrack_corner(size,nodes,k,(left|bit)<<1,(down|bit),(right|bit)>>1,L,l); 
    }
  return counter;
}
// GPU -n Kレイヤー k 番目のレイヤのすべてのノードを含むベクトルを返す
std::vector<long> kLayer_nodeLayer(unsigned int size,unsigned int k,std::vector<local>& L)
{
  std::vector<long> nodes{};
  unsigned int bit=0;
  struct local l;
  l.TOTAL=l.UNIQUE=l.COUNT2=l.COUNT4=l.COUNT8=0;
  l.TOPBIT=1<<(size-1);
  l.ENDBIT=l.LASTMASK=l.SIDEMASK=0;
  l.BOUND1=2;
  l.BOUND2=0;
  l.board[0]=1;
  while(l.BOUND1>1 && l.BOUND1<size-1){
    if(l.BOUND1<size-1){
      bit=1<<l.BOUND1;
      l.board[1]=bit;   //2行目にQを配置
      //角にQがあるときのバックトラック
      l.TYPE=0;
      kLayer_nodeLayer_backTrack_corner(size,nodes,k,(2|bit)<<1,1|bit,(2|bit)>>1,L,&l);
    }
    l.BOUND1++;
  }
  l.TOPBIT=1<<(size-1);
  l.ENDBIT=l.TOPBIT>>1;
  l.SIDEMASK=l.TOPBIT|1;
  l.LASTMASK=l.TOPBIT|1;
  l.BOUND1=1;
  l.BOUND2=size-2;
  while(l.BOUND1>0 && l.BOUND2<size-1 && l.BOUND1<l.BOUND2){
    if(l.BOUND1<l.BOUND2){
      bit=1<<l.BOUND1;
      l.board[0]=bit;   //Qを配置
      //角にQがないときのバックトラック
      l.TYPE=1;
      kLayer_nodeLayer_backTrack(size,nodes,k,bit<<1,bit,bit>>1,L,&l);
    }
    l.BOUND1++;
    l.BOUND2--;
    l.ENDBIT=l.ENDBIT>>1;
    l.LASTMASK=l.LASTMASK<<1|l.LASTMASK|l.LASTMASK>>1;
  }
  return nodes;
}
// GPU -n ノードレイヤーの作成
void symmetry_build_nodeLayer(unsigned int size)
{
  // ツリーの3番目のレイヤーにあるノード
  //(それぞれ連続する3つの数字でエンコードされる)のベクトル。
  // レイヤー2以降はノードの数が均等なので対称性を利用できる。
  // レイヤ4には十分なノードがある(N16の場合、9844)。
  // ここではレイヤーを5に設定、Nに併せて増やしていく
  std::vector<local>L;
  // NodeLayerは、N18でabortします。
  std::vector<long> nodes=kLayer_nodeLayer(size,5,L); 

  // デバイスにはクラスがないので、
  // 最初の要素を指定してからデバイスにコピーする。
  size_t nodeSize=nodes.size() * sizeof(long);
  long* hostNodes=(long*)malloc(nodeSize);
  hostNodes=&nodes[0];
  long* deviceNodes=NULL;
  cudaMalloc((void**)&deviceNodes,nodeSize);
  cudaMemcpy(deviceNodes,hostNodes,nodeSize,cudaMemcpyHostToDevice);

  // host/device Local
  //size_t localSize=numSolutions * sizeof(struct local);
  size_t localSize=L.size() * sizeof(struct local);
  struct local* hostLocal=(local*)malloc(localSize);
  hostLocal=&L[0];
  local* deviceLocal=NULL;
  cudaMalloc((void**)&deviceLocal,localSize);
  cudaMemcpy(deviceLocal, hostLocal, localSize, cudaMemcpyHostToDevice);

  // デバイス出力の割り当て
  // 必要なのはノードの半分だけで、
  // 各ノードは3つの整数で符号化される。
  long* deviceSolutions=NULL;
  int numSolutions=nodes.size() / 3; 
  size_t solutionSize=numSolutions * sizeof(long);
  cudaMalloc((void**)&deviceSolutions,solutionSize);

  // CUDAカーネルを起動する。
  int threadsPerBlock=256;
  int blocksPerGrid=(numSolutions+threadsPerBlock-1)/threadsPerBlock;
  dim_nodeLayer <<<blocksPerGrid,threadsPerBlock>>>(size,deviceNodes,deviceSolutions,numSolutions,deviceLocal);

  // 結果をホストにコピー
  long* hostSolutions=(long*)malloc(solutionSize);
  cudaMemcpy(hostSolutions,deviceSolutions,solutionSize,cudaMemcpyDeviceToHost);
  
  // 部分解を加算し、結果を表示する。
  unsigned long solutions=0;
  for(unsigned long i=0;i<numSolutions;i++){
       solutions += hostSolutions[i]; // Symmetry
  }
  // 出力
  TOTAL=solutions;
}
// CUDA 初期化
bool InitCUDA()
{
  int count;
  cudaGetDeviceCount(&count);
  if(count==0){fprintf(stderr,"There is no device.\n");return false;}
  unsigned 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;
  unsigned 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\n対称解除法 再帰 \n"); }
  else if(cpu){ printf("\n\n対称解除法 非再帰 \n"); }
  else if(gpu){ printf("\n\n対称解除法 GPU\n"); }
  else if(gpuNodeLayer){ printf("\n\n対称解除法 GPUノードレイヤー \n"); }
  if(cpu||cpur)
  {
    unsigned int min=4; 
    unsigned int targetN=17;
    struct timeval t0;
    struct timeval t1;
    printf("%s\n"," N:        Total      Unique      dd:hh:mm:ss.ms");
    for(unsigned int size=min;size<=targetN;size++){
      local l;
      gettimeofday(&t0,NULL);//計測開始
      if(cpur){ //再帰
        symmetry_R(size,&l);
      }
      if(cpu){ //非再帰
        symmetry_NR(size,&l);
      }
      //
      gettimeofday(&t1,NULL);//計測終了
      unsigned int ss;
      unsigned int ms;
      unsigned 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
      unsigned int hh=ss/3600;
      unsigned 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; */
    unsigned int min=4;
    unsigned int targetN=21;
    struct timeval t0;
    struct timeval t1;
    printf("%s\n"," N:        Total      Unique      dd:hh:mm:ss.ms");
    for(unsigned int size=min;size<=targetN;size++){
      gettimeofday(&t0,NULL);
      if(gpu){
        TOTAL=UNIQUE=0;
        local l[MAX];
        GPU_symmetry_R(size,&l[0]);
        TOTAL=l->TOTAL;
        UNIQUE=l->UNIQUE;
      }else if(gpuNodeLayer){
        TOTAL=UNIQUE=0;
        symmetry_build_nodeLayer(size);
      }
      gettimeofday(&t1,NULL);
      unsigned int ss;
      unsigned int ms;
      unsigned 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
      unsigned int hh=ss/3600;
      unsigned 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);
    }
  }
  return 0;
}

CUDA 実行結果

GPU で並列処理で実行(ノードレイヤー)
$ nvcc -O3 -arch=sm_61 03CUDA_Symmetry_NodeLayer.cu && ./a.out -n

対称解除法 GPUノードレイヤー
 N:        Total      Unique      dd:hh:mm:ss.ms
 4:            0           0      00:00:00:00.13
 5:            0           0      00:00:00:00.00
 6:            0           0      00:00:00:00.00
 7:           40           0      00:00:00:00.00
 8:           92           0      00:00:00:00.00
 9:          352           0      00:00:00:00.00
10:          724           0      00:00:00:00.00
11:         2680           0      00:00:00:00.00
12:        14200           0      00:00:00:00.00
13:        73712           0      00:00:00:00.03
14:       365596           0      00:00:00:00.15
15:      2279184           0      00:00:00:00.86
16:     14772512           0      00:00:00:05.50
17:     95815104           0      00:00:00:38.87
18:    666090624           0      00:00:04:46.16

以降はバーストします。

次回はビットボードに行きます!

参考リンク

以下の詳細説明を参考にしてください。
【参考リンク】Nクイーン問題 過去記事一覧
【Github】エイト・クイーンのソース置き場 BashもJavaもPythonも!

Nクイーン問題(50)第七章 マルチプロセス Python編
https://suzukiiichiro.github.io/posts/2023-06-21-04-n-queens-suzuki/
Nクイーン問題(49)第七章 マルチスレッド Python編
https://suzukiiichiro.github.io/posts/2023-06-21-03-n-queens-suzuki/
Nクイーン問題(48)第七章 シングルスレッド Python編
https://suzukiiichiro.github.io/posts/2023-06-21-02-n-queens-suzuki/
Nクイーン問題(47)第七章 クラス Python編
https://suzukiiichiro.github.io/posts/2023-06-21-01-n-queens-suzuki/
Nクイーン問題(46)第七章 ステップNの実装 Python編
https://suzukiiichiro.github.io/posts/2023-06-16-02-n-queens-suzuki/
Nクイーン問題(45)第七章 キャリーチェーン Python編
https://suzukiiichiro.github.io/posts/2023-06-16-01-n-queens-suzuki/
Nクイーン問題(44)第七章 対象解除法 Python編
https://suzukiiichiro.github.io/posts/2023-06-14-02-n-queens-suzuki/
Nクイーン問題(43)第七章 ミラー Python編
https://suzukiiichiro.github.io/posts/2023-06-14-01-n-queens-suzuki/
Nクイーン問題(42)第七章 ビットマップ Python編
https://suzukiiichiro.github.io/posts/2023-06-13-05-n-queens-suzuki/
Nクイーン問題(41)第七章 配置フラグ Python編
https://suzukiiichiro.github.io/posts/2023-06-13-04-n-queens-suzuki/
Nクイーン問題(40)第七章 バックトラック Python編
https://suzukiiichiro.github.io/posts/2023-06-13-03-n-queens-suzuki/
Nクイーン問題(39)第七章 バックトラック準備編 Python編
https://suzukiiichiro.github.io/posts/2023-06-13-02-n-queens-suzuki/
Nクイーン問題(38)第七章 ブルートフォース Python編
https://suzukiiichiro.github.io/posts/2023-06-13-01-n-queens-suzuki/
Nクイーン問題(37)第六章 C言語移植 その17 pthread並列処理完成
https://suzukiiichiro.github.io/posts/2023-05-30-17-n-queens-suzuki/
Nクイーン問題(36)第六章 C言語移植 その16 pthreadの実装
https://suzukiiichiro.github.io/posts/2023-05-30-16-n-queens-suzuki/
Nクイーン問題(35)第六章 C言語移植 その15 pthread実装直前版完成
https://suzukiiichiro.github.io/posts/2023-05-30-15-n-queens-suzuki/
Nクイーン問題(34)第六章 C言語移植 その14
https://suzukiiichiro.github.io/posts/2023-05-30-14-n-queens-suzuki/
Nクイーン問題(33)第六章 C言語移植 その13
https://suzukiiichiro.github.io/posts/2023-05-30-13-n-queens-suzuki/
Nクイーン問題(32)第六章 C言語移植 その12
https://suzukiiichiro.github.io/posts/2023-05-30-12-n-queens-suzuki/
Nクイーン問題(31)第六章 C言語移植 その11
https://suzukiiichiro.github.io/posts/2023-05-30-11-n-queens-suzuki/
Nクイーン問題(30)第六章 C言語移植 その10
https://suzukiiichiro.github.io/posts/2023-05-30-10-n-queens-suzuki/
Nクイーン問題(29)第六章 C言語移植 その9
https://suzukiiichiro.github.io/posts/2023-05-30-09-n-queens-suzuki/
Nクイーン問題(28)第六章 C言語移植 その8
https://suzukiiichiro.github.io/posts/2023-05-30-08-n-queens-suzuki/
Nクイーン問題(27)第六章 C言語移植 その7
https://suzukiiichiro.github.io/posts/2023-05-30-07-n-queens-suzuki/
Nクイーン問題(26)第六章 C言語移植 その6
https://suzukiiichiro.github.io/posts/2023-05-30-06-n-queens-suzuki/
Nクイーン問題(25)第六章 C言語移植 その5
https://suzukiiichiro.github.io/posts/2023-05-30-05-n-queens-suzuki/
Nクイーン問題(24)第六章 C言語移植 その4
https://suzukiiichiro.github.io/posts/2023-05-30-04-n-queens-suzuki/
Nクイーン問題(23)第六章 C言語移植 その3
https://suzukiiichiro.github.io/posts/2023-05-30-03-n-queens-suzuki/
Nクイーン問題(22)第六章 C言語移植 その2
https://suzukiiichiro.github.io/posts/2023-05-30-02-n-queens-suzuki/
Nクイーン問題(21)第六章 C言語移植 その1
N-Queens問://suzukiiichiro.github.io/posts/2023-05-30-01-n-queens-suzuki/
Nクイーン問題(20)第五章 並列処理
https://suzukiiichiro.github.io/posts/2023-05-23-02-n-queens-suzuki/
Nクイーン問題(19)第五章 キャリーチェーン
https://suzukiiichiro.github.io/posts/2023-05-23-01-n-queens-suzuki/
Nクイーン問題(18)第四章 エイト・クイーンノスタルジー
https://suzukiiichiro.github.io/posts/2023-04-25-01-n-queens-suzuki/
Nクイーン問題(17)第四章 偉人のソースを読む「N24を発見 Jeff Somers」
https://suzukiiichiro.github.io/posts/2023-04-21-01-n-queens-suzuki/
Nクイーン問題(16)第三章 対象解除法 ソース解説
https://suzukiiichiro.github.io/posts/2023-04-18-01-n-queens-suzuki/
Nクイーン問題(15)第三章 対象解除法 ロジック解説
https://suzukiiichiro.github.io/posts/2023-04-13-02-nqueens-suzuki/
Nクイーン問題(14)第三章 ミラー
https://suzukiiichiro.github.io/posts/2023-04-13-01-nqueens-suzuki/
Nクイーン問題(13)第三章 ビットマップ
https://suzukiiichiro.github.io/posts/2023-04-05-01-nqueens-suzuki/
Nクイーン問題(12)第二章 まとめ
https://suzukiiichiro.github.io/posts/2023-03-17-02-n-queens-suzuki/
Nクイーン問題(11)第二章 配置フラグの再帰・非再帰
https://suzukiiichiro.github.io/posts/2023-03-17-01-n-queens-suzuki/
Nクイーン問題(10)第二章 バックトラックの再帰・非再帰
https://suzukiiichiro.github.io/posts/2023-03-16-01-n-queens-suzuki/
Nクイーン問題(9)第二章 ブルートフォースの再帰・非再帰
https://suzukiiichiro.github.io/posts/2023-03-14-01-n-queens-suzuki/
Nクイーン問題(8)第一章 まとめ
https://suzukiiichiro.github.io/posts/2023-03-09-01-n-queens-suzuki/
Nクイーン問題(7)第一章 ブルートフォース再び
https://suzukiiichiro.github.io/posts/2023-03-08-01-n-queens-suzuki/
Nクイーン問題(6)第一章 配置フラグ
https://suzukiiichiro.github.io/posts/2023-03-07-01-n-queens-suzuki/
Nクイーン問題(5)第一章 進捗表示テーブルの作成
https://suzukiiichiro.github.io/posts/2023-03-06-01-n-queens-suzuki/
Nクイーン問題(4)第一章 バックトラック
https://suzukiiichiro.github.io/posts/2023-02-21-01-n-queens-suzuki/
Nクイーン問題(3)第一章 バックトラック準備編
https://suzukiiichiro.github.io/posts/2023-02-14-03-n-queens-suzuki/
Nクイーン問題(2)第一章 ブルートフォース
https://suzukiiichiro.github.io/posts/2023-02-14-02-n-queens-suzuki/
Nクイーン問題(1)第一章 エイトクイーンについて
https://suzukiiichiro.github.io/posts/2023-02-14-01-n-queens-suzuki/

書籍の紹介

Nクイーン問題(63)第七章 並列処理 対称解除法 ビットボード NVIDIA CUDA編

Nクイーン問題(63)第七章 並列処理 対称解除法 ビットボード NVIDIA CUDA編

Nクイーン問題(61)第七章 並列処理 ミラー NVIDIA CUDA編

Nクイーン問題(61)第七章 並列処理 ミラー NVIDIA CUDA編