Visual Profilerとは？ CUDAを用いてハイパフォーマンスなプログラムを書くためには、 プログラムの特性を調べ、最もボトルネックとなっている部分を確認しながら、 適切な最適化テクニックを用いてコーディングを行う必要があります。 ここで言う「特性」には以下のようなものがあります。 プログラム全体の実行時間の内訳 Kernelにおけるレジスタやメモリの使用状況 Coalesced Access状況 Kernelの命令数 分岐によるWarp分割数 ここに挙げたのはほんの一例ですが、これらの特性を自分で算出するには、 本来行いたいプログラミングからは大きく離れた所に労力を費やさねばなりません。 そこで出番となるのが、CUDAプログラムのプロファイリングツール、CUDA Visual Profilerです。 もともとCUDAプログラムにはプロファイルをCSV形式で出力する機能が備わっ

このブログは、株式会社フィックスターズのエンジニアが、あらゆるテーマについて自由に書いているブログです。 ちょっとしたきっかけがあって、AMD用のGPUとNVIDIA用のGPU両方で高速化作業を行いました。 そのときに得られた知見を書いておきます。 AMD GCN isa と NVIDIA SASS GPUでプログラミングを行うときは、OpenCL、CUDA などを使うと思いますが、これらはより低レベルなGPU用の機械語にコンパイルされます。 AMD のGPUでは、これに”ISA”、NVIDIA の GPU では、これを”SASS”という名前が付けられています。ISAは一般的な単語なので、別に正しい名前があるかもしれませんが、ビルド時のオプションなどで指定する場合は、–isa 等を指定しているので、ここでは、GCN isa とします。 通常のGPUプログラミングでこれらを見ることはないかも

ようやく、CUDAのcoalesced accessというのが理解出来た。 単純なメモリコピーで、8GB/S程度しか出ていなかったのが、今日、20GB/Sまで向上した。 キモは、 連続したメモリに各スレッドが同時にアクセスする という事。 今までは、画像のフィルターを、ラインに分割して、1ラインを1スレッドにやらせていた。 これだと、全てのスレッドが、担当するラインのピクセルを読み込みに行く。 つまり、非連続なメモリをアクセスするので、読み込みの指令数分だけのアクセスが生じる。 この例で分かりやすく言えば、スレッド数分だけのアクセスが生じるわけだ。 これを、各スレッドが各ピクセルを処理させる様にする。 これがcoalescだ。 スレッド0番は、Pixel0を、スレッド1番は、Pixel1を...と、順番に並べてアクセスさせる。 次のループでは、スレッド数分先を読み込みに行く。 簡単に言えば

Shared Memoryのバンク衝突についての実験 作成日　2008/7/18 Shared Memoryのバンク衝突を回避するためにどのようなアドレッシングにすればよいのか、 また、バンク衝突が起きた場合、どの程度のパフォーマンスの低下が見られるのか、実験してみました。 まずは、ソースコードを。カーネル部分だけ示します。 もとは、SDKのtemplateサンプルとほぼ同じです。 #ifndef _SharedTest_KERNEL_H_ #define _SharedTest_KERNEL_H_ #include <stdio.h> #define SDATA( index)      CUT_BANK_CHECKER(sdata, index) //////////////////////////////////////////////////////////////////////

CUDAプログラミング 次に、CUDA SDKのmatrixMultという行列積のプログラムに沿って、CUDAのプログラミングを見て行こう。プログラミングガイドに例題が載っている行列積と基本的には同じプログラムであるが、CUDA SDKでは、通常実行用のGPUバイナリを生成する以外に、デバグ機能を組み込んだバイナリを生成したり、GPUのない環境でCPUで実行するエミュレーションバイナリを生成したりする機能があり、SDKの例題は、直接CUDAを呼び出すのではなく、cutil.hで定義された呼び出し形式で書かれているという点が異なっている。 このcutil.hは、例えば後述のCUDA_SAFE_CALLは、通常実行モードバイナリでは単純に引数として与えられたCUDA関数の呼び出しに置き換えるが、_DEBUGが定義されていると、引数の関数を呼び出し、その戻り値のエラーを検査してエラーの場合はエラ

この間のインタプリタをはじめから・・・で問題がありご指摘を受けました。 ですので次回より修正を行い、今回は GPGPU を使った計算についてやりたいと思います。 今回 nVidia の GeForce 8800 GTX というボードが手に入りましたので、専用の言語 （現在 GF8x 系のみで動作可能） であるCUDAを利用したいと思います。 CPUからの命令でGPUメモリのGPUのデータ領域のメモリを確保 CPUからGPUへメモリ内容をコピー GPUで演算処理、エラーの有無をチェック GPUからCPUへ出力用メモリ内容をコピー という流れで処理を行います。 注意しなければならない点としてGPUからCPUにメモリ内容をコピーする際、GPU内部で出力用メモリに書き出しが行われなかった場合、前回の出力結果とまったく同じものが出てきます。 (再起動してもフラッシュされない場合もありました。)

目次 >> CUDA >> メモリアクセス CUDA - メモリアクセス CUDAは、コンピュータ上のRAM、グラフィックカード上のRAMなどいくつかのメモリをもつ。 それぞれのメモリは、アクセス速度やアクセスできる範囲などが違う。 具体的には、コンピュータ上のメインメモリには直接アクセスはできない。そのため、cudaMemcpyを使って、メインメモリからグラフィックカード上のメモリへ転送してやる必要がある。 一方、グラフィックカード上のメモリにもいくつか種類があり、一つはグローバルメモリであるが、これは同一グリッド内であれば、どのブロック、どのスレッドからもアクセスできる。 一方、シェアードメモリは同一ブロック内のスレッドからしかアクセスできないものの、グローバルメモリに比べて、きわめて高速にアクセスできる。 今回は、これらの違いがどの程度実行速度に影響するか調べてみた。なお今回使用して

#include <stdio.h> #include <stdlib.h> #include <string.h> #include <cutil.h> #include <sys/time.h> #include <unistd.h> typedef struct __align__(8) { float real; float imaginary; } Complex; __device__ __host__ void initComplex( Complex *_arg_this, float _arg_real, float _arg_imaginary ) { _arg_this->real = _arg_real; _arg_this->imaginary = _arg_imaginary; } __device__ void addComplex( Complex *_ar

{{Framework.description ? Framework.description : 'Join the GeForce community. Browse categories, post your questions, or just chat with other members.'}}

NVIDIA CUDA Information Site — unleash the power of GPU / GPUのパワーを引き出そう — NVIDIA CUDA Information サイト（以下「本サイト」）は、マルチコアソリューションを提供するフィックスターズの技術者有志が運営する、NVIDIA CUDAの普及と利用促進を目的とする情報公開と情報交換のためのサイトです。 本サイトでは、NVIDIA社が発売するGPUアクセラレーターボードTesla C1060/S1070 (Tesla）上のCUDAアプリケーション開発に関する情報を中心として、GPUに関する各種の情報をお伝えします。

Release Notes CUDA Features Archive EULA Installation Guides Quick Start Guide Installation Guide Windows Installation Guide Linux Programming Guides Programming Guide Best Practices Guide Maxwell Compatibility Guide Pascal Compatibility Guide Volta Compatibility Guide Turing Compatibility Guide NVIDIA Ampere GPU Architecture Compatibility Guide Hopper Compatibility Guide Ada Compatibility Guide M

東工大のTSUBAMEですが、GPGPUを組み合わせた1.2にアップグレードするそうです。という話が先日のWindows HPC server 2008の発表会の時に松岡教授の口から出てきたのですが…… TSUBAMEは元々Opteronを大量採用した初期事例で有名なのにTSUBAME 1.2で使われるGPUはNVIDIA（実証環境はGeFroce 8800で、本環境はTesla T10Pだそうで）。 #ということで、下世話ながら表題のフレーズが頭をよぎったAC 興味深いモデが付いたので追記 現状のテストベッドはGeForce 8800GTSを使ったMS HPC-GPGPUクラスタは現在32ノードで、GPUを128枚使用、3ラック40U。Windows Compute Cluster 2003->Windows HPC 2008に移行中とのこと。開発環境はVisual Studio + C