• トップ
  • G-DEPについて
  • ご購入ガイド
  • サポート
  • お問い合わせ

G-DEPトップ  >  G-DEPの高速演算記  >  高速演算記 第25回 「Kepler解説その2 〜Kepler世代の新機能〜」

高速演算記 第25回 「Kepler解説その2 〜Kepler世代の新機能〜」

 

前回の記事ではサンプルプログラムを通じてK20の優れた性能を見ました。今回は具体的なカーネルを例に挙げて、K20のハイパフォーマンス性と新機能を利用したチューニング例をご紹介します。今回扱うカーネルは疎行列-ベクトル積 y=Ax、いわゆるSpMVです。疎行列とはその成分のほとんどが零の行列の事で、SpMVでは疎行列の非零要素だけを取り出して上手く処理する事が、メモリアクセスと省メモリの観点から重要となります。なお、本記事で扱うRead-OnlyデータキャッシュとWarp Shuffle命令を利用したSpMVのK20向けチューニングは、2013年1月に開催されたNVIDIA Manufacturing Day 2013にてプロメテック・ソフトウェア社の北岡氏が講演内で紹介されたものです。

 

SpMVのカーネル

SpMVのCUDAでの実装例のいくつかは、有名な論文『Efficient Sparse Matrix-Vector Multiplication on CUDA』(Nathan Bell and Michael Garland著) に掲載されています。今回はその論文中のFigure 22に掲載されているCSR形式のSpMVカーネル(vector kernel)を取り上げます。それは次のようなカーネルです。(※世代を考慮してvolatileを加えています):
 
  __global__   void   spmv_csr_vector_kernel ( const int           num_rows,
                                                         const int           *ptr,
                                                         const int           *indices,
                                                         const double  *data,
                                                         const double  *x,
                                                            double             *y)
 {
    extern  __shared__  volatile  double  vals[];
 
    int  thread_id = blockDim.x * blockIdx.x + threadIdx.x;   // global thread index
    int  warp_id    = thread_id / 32;                             // global warp index
    int  lane          = thread_id & (32-1);                         // thread index within the warp
 
    // one warp per row
    int row = warp_id;
 
    if (row < num_rows)
    {
       int  row_start  = ptr[row];
       int  row_end   = ptr[row+1];
 
       // compute running sum per thread
       double  sum = 0.0;
       for (int  jj=row_start+lane;  jj<row_end;  jj+=32)
          sum += data[jj] * x[ indices[jj] ];
 
       // parallel reduction in shared memory
    // ( slightly faster than naive code, vals[threadIdx.x] += vals[threadIdx.x + ... ] )
       vals[threadIdx.x] = sum; 
       vals[threadIdx.x] = sum = sum + vals[threadIdx.x + 16];
       vals[threadIdx.x] = sum = sum + vals[threadIdx.x +   8];
       vals[threadIdx.x] = sum = sum + vals[threadIdx.x +   4];
       vals[threadIdx.x] = sum = sum + vals[threadIdx.x +   2];
       sum = sum + vals[threadIdx.x+ 1];
 
       // first thread writes the result
       if (lane == 0)
           y[row] = sum;
   }
 }
 
 
このカーネルの内容を簡単に確認しておきます。データ構造はCSR形式になっていて、y=AxというSpMVに対し、num_rowsには疎行列Aの行数、data[ i ]  ( i = 0,1, ...) には疎行列Aの Row-major での i番目の非零要素、indices[ i ] には data [ i ] の列番号、ptr[ i ] (i=0,1, ..., num_rows-1) には疎行列Aの第 i 行の最初の非零要素が配列 data の何番目に格納されているか、がそれぞれ格納されています。ptr [num_rows] には非零要素の総数が格納されています。以下に 4×5疎行列での例を挙げておきます:
 

          num_rows  =  4
               data      =  [ 1.0,  4.0,  2.0,  3.0,  5.0,  7.0,  8.0,  9.0,  6.0 ]
             indices    =  [   0,     1,     1,     2,     0,     3,     4,     2,    4   ]
                ptr        =  [   0,             2,             4,                     7,    9   ]
 
このカーネルでは1 warp、すなわちthreadIdxの連続する32スレッドがSpMVの1つの内積を担当しています。1 warpで1つの内積を並列処理し、最後にwarp-wideな並列リダクションを行っています。同じwarp内のスレッドは __syncthreads() がなくとも自動的に同期するという性質を活かした形です。

このカーネルをK20向けにチューニングする前に、まずはこのままC2075とK20の両方で実行して性能比較をしておきましょう。その準備として、次にこのカーネルに最適なグリッド構成 <<<Dg, Db, (Db + 16) * sizeof(double)>>> を考えてみます。

 

K20での最適なグリッド構成は?

上記のSpMVカーネルの実行に最適なグリッド構成 <<<Dg, Db, (Db + 16) * sizeof(double)>>> はどのようなものでしょうか? それはGPUのハードウェア・スペックによって異なります。NVIDIA GPUのハードウェア・スペックはCompute Capability (CC) というラベルで区別されるので、つまりはCC毎に異なります。Fermi世代TeslaであるC2075等はCC2.0であり、Kepler世代TeslaであるK20とK20XはCC3.5なので、以下ではこの2つのCCそれぞれでSpMVに最適なブロック構成を探してみます。ちなみに、CCの1の位は世代を定義していて、2がFermi世代、3がKepler世代です。CCの少数点部はマイナー番号で、例えば同じ世代のTeslaとGeForceでもアーキテクチャに微妙な違いがある事を区別しています。なお、CC2.0では Dg ≦ 65535という制限がありましたが、CC3.5ではそれは Dg ≦ 231-1 にまで緩和されました。Db ≦ 1024という制限と1 warp = 32スレッドである事は両世代に共通なので、両世代とも1ブロックは1 warp (= 32スレッド) 〜 32 warp (= 1024スレッド) の範囲に組まなければなりません。CC毎の違いの一覧表はこちらにあります。
 
  CC2.0 (C2075等)  CC3.5 (K20, K20X)
 gridDim.x の最大値 65535 231-1
 gridDim.y, gridDIm.z の最大値 65535 65535
 1ブロック内の最大スレッド数   1024 1024
 1ブロック内の最大warp数 32 32
 
さて、最適なグリッド構成の目安となるのはTeslaが最大限 busyになる構成です。それを見つけるには、Tesla内部に複数搭載されている "SMX" の以下の4つのハードウェア・スペックに注目する必要があります:
 
                 ・1 SMXが同時管理できる最大warp数
                 ・1 SMXが同時管理できる最大ブロック数
                 ・1 SMXあたりの共有メモリ容量
                 ・1 SMXあたりの32-bit レジスタファイル数
 
この4つの制限下でなるべく高いOccupancy(= カーネル実行中に1 SMXが同時管理しているwarp数 / 1 SMXが同時管理できる最大warp数)を実現させたいわけです。まずは話を簡単にするために、最初の2つの制限だけの下で最適なグリッド構成を考えてみましょう。
 
 
上の写真は、前回の記事でもご紹介した、CC3.5のTeslaの中にあるチップGK110です。右の設計図に記されているように、ここにSMXが並んでいます。GK110にはSMXが #0〜#14の15個搭載されていますが、K20ではその内の13個しか動作しないように制御がかかっています。よりハイエンドなTesla K20Xでは14個が動作しますし、将来的には15個全てが動作するTeslaが発売されるかもしれません。
 
CUDAプログラミングモデルにおけるブロックはこのSMX内で管理されます。1つのブロックが異なるSMXにまたがって管理される事はありません。CUDAプログラミングモデルにおけるスレッドはSMX内のコアに対応していて、スレッドの処理は、1つのコアが一人二役、三役…と働いて遂行されていきます。複数のブロックが1 SMX内で同時に管理されますが、1 SMXが同時に管理できるwarp数・ブロック数には上限があります。先の4つの制限の内の最初の2つはこれです。この上限を引き出す、Occupancyの高いグリッド構成を選ぶ事で、Teslaをbusyにできるわけです。
 
C2075等のCC2.0では1 SMが同時管理できる最大warp数は48、最大ブロック数は8でした。このため、共有メモリとレジスタ数が制限を与えなければ、1ブロックのwarp数を 6 (Db = 192),  8 (Db = 256),  12 (Db = 384),  16 (Db = 512),  24 (Db = 768) のいずれかに設定する事で、1 SMが同時管理するブロック数はそれぞれ8,  6,  4,  3,  2となり、48/48 = 100%の Occupancyが達成されます。(※ C2075ならSMは14個搭載されているので、Dgを8×14=112の倍数かもしくは112より十分大きくしない限り、Tail Effectにより効率が下がるためり、通常は1000程度あるいはそれよりも多く取ります。)
 
一方、K20等のCC3.5では1 SMXが同時管理できる最大warp数は64に、最大ブロック数は16に強化されています。こちらは、共有メモリとレジスタ数が制限を与えなければ、1ブロックのwarp数を 4 (Db = 128),  8 (Db = 256),  16 (Db = 512),  32 (Db = 1024) のいずれかに設定する事で、1 SMXが同時管理するブロック数はそれぞれ16,  8,  4,  2となり、64/64 = 100%のOccupancyが達成されます。(※こちらは、K20ならSMXは13個搭載されているので、Dgを16×13=208の倍数かもしくは208より十分大きくしない限り、Tail Effectにより効率は下がります。)共有メモリとレジスタ数の制限がなければ、こういう結果を得るわけです。共有メモリとレジスタ数も含めて比較したのが下図です。
 
  CC2.0 (C2075等) CC3.5 (K20, K20X)
 1SMX (SM) が同時管理できる最大warp数 48 64
 1SMX (SM) が同時管理できる最大ブロック数 8 16
 1SMX (SM) あたりの共有メモリ容量 49152 Byte 49152 Byte
 1SMX (SM) あたりの 32-bit レジスタファイル数 32768 65536
 
しかし実際には、CC2.0ではレジスタ数がボトルネックとなって100%のOccupancyを達成できない事が多々有りました。カーネル実行時に1スレッドあたりが使用するレジスタ数は、オプション『--ptxas-options=-v』をつけてnvccコンパイルする事で簡単に確認できます。あるいはCUDA5のnvprofでも簡単に確認できます。CC2.0において100%のOccupancyを達成するためには、この1 スレッドが使用するレジスタ数が32768 / (48×32) = 21 以下である必要があります…が、実はCC2.0ではレジスタのアロケーションはwarp毎に64単位で行われるため、1スレッドが21レジスタを使用しようとすると各warpが確保するレジスタ数が21×32 = 672(64の倍数でない)から704(64の倍数)に切り上げされてしまい、1 SMが48 warpを同時管理するのに必要なレジスタ数は704×48 = 33792となって制限値32768を超えてしまいます。このため、CC2.0で100%のOccupancyを達成するためには 1スレッドが使用するレジスタ数が20以下である必要があり、これは経験上、中々厳しい条件でした。

これに対して、K20等のCC3.5では1 スレッドが使用するレジスタ数が65536 / (64×32) = 32 以下であれば100%のOccupancyを達成できるというように制限が緩和されています。こちらのレジスタのアロケーションはwarp毎に256単位で行われ、不運な切り上げもありません。勿論、折角増えたレジスタを余らせてしまっては、それはそれで勿体ないのですが。

さて、今のSpMVカーネルで1スレッドが使用するレジスタ数を確認すると、CC2.0(-arch=sm_20でコンパイル)では21となり、残念ながら100%のOccupancyを達成できない事がわかります。レジスタ数が足りず、1 SMが48warpも同時管理できないからです。この場合に達成可能な最大のOccupancyは、例えば 1ブロックを 6 warp (Db = 192) で組んで 1 SMが7ブロックを同時管理する 42/48 = 88%です。一方、K20等のCC3.5(-arch=sm_35でコンパイル)では1スレッドが使用するレジスタ数は27で、100%のOccupancyを達成できる事がわかります。1ブロックのwarp数を 4 (Db = 128)、8 (Db =256)、16 (Db = 512)、32 (Db = 1024) のいずれかに設定する事で100%のOccupancyを達成できます。今のSpMVカーネルをK20で実行する場合、これらのグリッド構成が最適化へ向けた第一候補となるわけです。このカーネルはまさに、Kepler世代で増強された(最大warp数とのバランスが良くなった)レジスタ数の恩恵を受けるわけです。(幸い、今は共有メモリ容量は制限を与えません。)

ちなみに、1ブロックあたりのwarp数に応じたOccupancyの変化はCUDA Occupancy Calculatorで簡単に計算できます。このツールを使ってKepler世代Teslaで上記SpMVを実行する際のOccupancyを計算して図にしたのが下図です。上記の通り、1ブロックを 4 warp、8 warp、16 warp、32 warpで組んだ場合に一番上の100% Occupancyラインに達する事がわかります。1ブロックが 7 warp、9 warp、21 warpの場合は 63/64 = 98%であり、これらもほぼ100%に達しています。

Kepler世代TeslaでのSpMVのOccupancy。横軸はブロックサイズ Db、縦軸は1 SMXの同時管理するwarp数。
赤いポイントは1ブロックを4warpで組んでOccupancy 100%を得るところ。  
                 
なお、Occupancyの他にも様々要因があるため、Occupancyを最大にするグリッド構成が必ずしも最適とは限りませんが、Occupancyは最初に意識考すべき目安となります。最適化には高いOccupancyを追求するThread-level Parallelism (TLP) だけでなく、1スレッドあたりの使用レジスタ数を高めて活かすInstruction-level Parallelism (ILP) もあります。後者に関する有名なプレゼンテーションはこちらにあります。        

 

K20での実行時間

では先のSpMVの実行時間をC2075とK20で測定して比較してみましょう。ここではフロリダ大学のSparse Matrix Collectionのmark3jac140(行列サイズは(60K)2、非零要素数は400K)でのSpMVの実行時間を測定しました。1ブロックを1 warp 〜32 warpで組んだ各場合に1000回測定し、平均実行時間をプロットした結果が下図です。Dgは必要最小限の整数値(≒ 60K/Db)に設定しています。共有メモリとL1キャッシュの比率はデフォルト設定にしています。誤差は小さく、グラフの見やすさのためにあえて省略しています。(テスト環境は前回の記事と同じです。

 

 

この結果を見ると、Occupancyと実行時間の相関がよく見てとれます。最速値は、C2075では1ブロックを6 warp(Db = 192)で組んだ場合の 0.5088msとなり、K20では1ブロックを4 warp(Db = 128)で組んだ場合の 0.3045msとなりました。少なくともこの例では、全く同じカーネルをC2075とK20で実行してK20で約1.7倍の高速化がなされた事になります。

また、C2075ではDg ≦ 65535という制限があるため、このままでは(65535×32)2より大きなサイズの行列を扱えません。そのためには、各warpが複数の内積を処理するようにカーネルを書き換える必要があります。しかし、K20ではその制限はDg ≦231-1にまで緩和されたため、((231-1)×32)2 ものサイズまでであれば今のカーネルのまま扱う事ができます。

 

K20でのチューニング例

Kepler世代TeslaはFermi世代に比べて単に各種スペックが向上しているだけではなく、いくつかの新機能も備えています。次はこれらを利用して先のSpMVカーネルを最適化してみましょう。ここでは北岡氏の講演と同じく、Kepler世代Teslaの新機能である『Read-Onlyデータキャッシュ』と『Warp Shuffle命令』を利用してみます。なお、これらを使用するためにはCUDA 5も必要です。

Read-Onlyデータキャッシュとは、1 SMXあたりに48KB搭載されている読み取り専用のキャッシュです。実は、このRead-OnlyデータキャッシュはFermi世代以前にテクスチャキャッシュと呼ばれていたものです。K20ではテクスチャキャッシュをレジスタに直結するパスが新設され、テクスチャユニットを経由せずに高速にレジスタへデータを転送できるようになったため、テクスチャキャッシュはRead-Onlyデータキャッシュへと名称変更されたわけです。このRead-Onlyデータキャッシュは非常に簡単な方法で使用する事ができます。そこにキャッシュさせたいデータをカーネルに渡す際、const修飾子に加えて“__restrict__” で修飾し、CUDA5にてオプション『-arch=sm_35』を付けてコンパイルするだけです。この簡単な手順で、そのデータは自動的にRead-Onlyデータキャッシュにキャッシュされるようになります。ここにはSpMVのベクトルxを置いてみます。

共有メモリ、L1キャッシュに並ぶ Read-Onlyデータキャッシュ

 

もう1つの新機能であるWarp Shuffle命令は、同じwarp内のスレッド間でレジスタレベルでデータ交換を行う命令です。この命令を使えば共有メモリを介さすにデータ交換を行えます。CUDA C Programming Guideにも掲載されているように、このWarp Shuffle命令は並列リダクションに応用できます。

 

Kepler世代ではこれら4パターンのWarp Shuffle命令を使える(※up, downでは図のような境界の循環は起こりません)
 

Read-Onlyデータキャッシュの効果

では、順を追って最適化を行っていきましょう。まずはRead-Onlyデータキャッシュを適用してみます。比較のため、テクスチャユニットを利用した場合もテストしてみます。テクスチャユニットを利用したカーネルは以下のようになります:
 

   texture<int2, cudaTextureType1D, cudaReadModeElementType> texture_;

   __global__   void   spmv_csr_vector_kernel_texture ( const int          num_rows,
                                        const int         *ptr,
                                                             const int         *indices,
                                                             const double *data,
                                                             const double *x,
                                                                double            *y)
  {
     extern  __shared__  volatile  double  vals[];
 
     int  thread_id = blockDim.x * blockIdx.x + threadIdx.x;   // global thread index
     int  warp_id    = thread_id / 32;                                // global warp index
     int  lane          = thread_id & (32-1);                            // thread index within the warp
 
     // one warp per row
     int  row = warp_id;
 
     if (row < num_rows)
     {
        int  row_start  = ptr[row];
        int  row_end   = ptr[row+1];
 
        // compute running sum per thread
        double sum = 0.0;
        for (int jj=row_start+lane;  jj<row_end;  jj+=32)
        {
           int2 const v = tex1Dfetch( texture_,  indices[jj] );
           sum += data[jj] * __hiloint2double(v.y, v.x);
        }
 
        // parallel reduction in shared memory
     // ( slightly faster than naive code, vals[threadIdx.x] += vals[threadIdx.x + ... ] )
        vals[threadIdx.x] = sum;
        vals[threadIdx.x] = sum = sum + vals[threadIdx.x + 16];
        vals[threadIdx.x] = sum = sum + vals[threadIdx.x +   8];
        vals[threadIdx.x] = sum = sum + vals[threadIdx.x +   4];
        vals[threadIdx.x] = sum = sum + vals[threadIdx.x +   2];
        sum = sum + vals[threadIdx.x+ 1];
 
        // first thread writes the result
        if (lane == 0)
            y[row] = sum;
     }
   }
 
 
そして、Read-Onlyデータキャッシュを利用したカーネルが以下です。このコラムの最初のカーネルの引数 xの前に__restrict__を追記しただけです。但し、CUDA5でオプション -arch=sm_35 を指定してコンパイルする必要が有ります:
 
   
   __global__   void   spmv_csr_vector_kernel_ReadOnly ( const int          num_rows,
                                                             const int         *ptr,
                                                             const int         *indices,
                                                             const double *data,
                                                             const double * __restrict__ x,
                                                                double            *y)
   {
     extern  __shared__  volatile  double  vals[];
 
     int  thread_id = blockDim.x * blockIdx.x + threadIdx.x;   // global thread index
     int  warp_id    = thread_id / 32;                                   // global warp index
     int  lane          = thread_id & (32-1);                               // thread index within the warp
 
     // one warp per row
     int row = warp_id;
 
     if ( row < num_rows )
     {
        int  row_start  = ptr[row];
        int  row_end   = ptr[row+1];
 
        // compute running sum per thread
        double  sum = 0.0;
        for (int jj=row_start+lane;  jj<row_end;  jj+=32)
          sum += data[jj] * x[ indices[jj] ];
 
        // parallel reduction in shared memory
     // ( slightly faster than naive code, vals[threadIdx.x] += vals[threadIdx.x + ... ] )
        vals[threadIdx.x] = sum; 
        vals[threadIdx.x] = sum = sum + vals[threadIdx.x + 16];
        vals[threadIdx.x] = sum = sum + vals[threadIdx.x +   8];
        vals[threadIdx.x] = sum = sum + vals[threadIdx.x +   4];
        vals[threadIdx.x] = sum = sum + vals[threadIdx.x +   2];
        sum = sum + vals[threadIdx.x+ 1];
 
        // first thread writes the result
        if (lane == 0)
            y[row] = sum;
     }
   }
 
 
これらのカーネルを実行した結果は下図のようになりました。テクスチャユニットを利用するカーネルはC2075とK20の両方で実行しています。カーネルの変更に伴い1スレッドが使用するレジスタ数は増えますが、K20には大量のレジスタが搭載されているためOccupancyを制限するに至りませんでした。そのためK20では、Occupancyは最初のカーネルの場合と全く同じです。そして再びOccupancyと実行時間がよく相関しています。K20でテクスチャユニットを使用した際の最速値は 0.2712ms、Read-Onlyデータキャッシュを使用した際の最速値は 0.2581msとなっており、いずれも1ブロックを4 warpで組んだ場合です。このように、Read-Onlyデータキャッシュは手軽に使えてしかも有用な新機能になっています。

 

 

Warp Shuffle命令の効果

では最後に、Read-Onlyデータキャッシュに加えてWarp-Shuffle命令も利用してみましょう。それが以下のカーネルです。共有メモリはもはや全く使用していません。
 
   
   __device__   __inline__    double  shfl_xor ( double value,  int const lane )
   {
       return  __hiloint2double( __shfl_xor(__double2hiint(value),lane),
                                                    __shfl_xor(__double2loint(value),lane)); 
   }
 
 
   __global__   void   spmv_csr_vector_kernel_K20 ( const int           num_rows,
                                                           const int          *ptr,
                                                           const int          *indices,
                                                           const double  *data,
                                                           const double  *   __restrict__   x,
                                                              double             *y)
   {
      int  thread_id = blockDim.x * blockIdx.x + threadIdx.x;   // global thread index
      int  warp_id    = thread_id / 32;                                // global warp index
      int  lane          = thread_id & (32-1);                            // thread index within the warp
 
      // one warp per row
      int row = warp_id;
 
      if (row < num_rows)
      {
         int  row_start  = ptr[row];
         int  row_end   = ptr[row+1];
 
         // compute running sum per thread
         double sum = 0.0;
         for ( int jj=row_start+lane;  jj<row_end;  jj+=32 )
            sum += data[jj] * x[indices[jj]];
 
         // parallel reduction in registers
         sum += shfl_xor(sum, 16);
         sum += shfl_xor(sum,   8);
         sum += shfl_xor(sum,   4);
         sum += shfl_xor(sum,   2);
         sum += shfl_xor(sum,   1);
 
         // first thread writes the result
         if ( lane == 0 )
             y[row] = sum;
      }
   }
 
 
このカーネル実行時のOccupancyも最初のカーネルと全く同じです。これを実行した結果も加えれば、最終的に下図を得ます。このカーネルでも最速値は1ブロックを4 warpで組んだ場合に得られ、0.2264msでした。最初のC2075の実行時間の約2.2倍となっています。少なくともこれらの例では、同じカーネルでは Db = 128 〜 256 で安定的に良い性能が出てているようです。こうしてKepler世代Teslaの優れた新機能を確認する事ができました:

 

 

その他の疎行列

同程度のサイズの他の疎行列としてvenkat01(非零要素数は1.7M)、water_tank(非零要素数は2M)を使ったテストもしたところ、以下のように同様の結果を得ました。venkat01でのはC2075での最速が 0.4951ms、K20での最速が 0.2341msです。water_tankではC2075での最速が 0.5362ms、K20での最速が 0.2913msです。勿論、より複雑な構造の疎行列であれば異なる結果を得る可能性はあります:

 

 

 

なお、高速なSpMVを追求するという意味では、CUDAに標準搭載されているCUSPARSEライブラリを使用する事が推奨されます。SpMVをK20にてCUSPARSEで実行したところ、mark3jac140では 0.1484ms、venkat01では 0.3564ms、water_tankでは 0.4011msという平均時間で処理する事が出来ました。これらの行列規模では先の最適化されたカーネルに負けたりもしていますが、行列サイズがより大きくなるとCUSPARSEの方が概ね倍以上高速になります。例えばCUSPARSEの実行時間は、rajat16(サイズは(100K)2、非零要素数は500K)では 0.4231ms、netherlands_osm(サイズは(2M)2、非零要素数は5M)では 11.01ms、memchip(サイズは(2.7M)2、非零要素数は13M)では 3.269msとなりました。これらはそれぞれ、WarpShuffle命令を使った先のカーネルに対して2.1倍、5.6倍、2.7倍高速でした。nvprofで確認したところ、CUSPARSEは使用レジスタ数が多く、ILPによる最適化が行われていると予想されます。

 

K20のその他の新機能

最後に、K20で強化されたその他の主なポイントを挙げておきます。それは以下のような点です:

  • ・atomic演算が高速化された
  • ・1 SMX内で共有メモリ32KB + L1キャッシュ32KBという設定が可能になった
  • ・L2キャッシュがFermi世代の2倍の1536KBになった
  • ・ECCのon/offでの性能差が削減された

この内のatomic演算の高速化を実際に確認してみるべくatomicAddによるint、float、doubleの3パターンの和を試してみると、C2075からK20にかけてそれぞれ約2.5倍、2.4倍、24倍に高速化されました。足し上げは<<<1,1024>>>というグリッド構成で行い、double用のatomicAddはCUDA C Programming Guideに記載されているコードをそのまま使用した結果です。K20でatomic演算が高速化されている事がよくわかります。

この結果を受けて、atomicAddを先のSpMV (mark3jac140) のリダクション部分に適用して最速値を確認したところ、C2075での 37.07msに対し、K20では10倍以上高速な 3.340msとなりました。これだけatomic演算が高速化されていれば、カーネルの内部ループでの使用に耐える事もあるかもしれません。

(G-DEP 技術本部 チーフエンジニア 河井博紀)