NVIDIA CUDA Compute Unified Device Architecture プログラミング ガイド( 日 本 語 版 ) Version 1.1 3/2/2008
ii CUDA Programming Guide Version 1.1
目 次 Chapter 1. CUDAの 紹 介...1 1.1 データ 並 列 演 算 デバイスとしてのグラフィック プロセッサ ユニット...1 1.2 CUDA: GPU での 演 算 のための 新 しいアーキテクチャ...3 1.3 本 書 の 構 成...6 Chapter 2. プログラミング モデル...7 2.1 高 度 なマルチスレッド コプロセッサ...7 2.2 スレッドの 集 合...エラー! ブックマークが 定 義 されていません 2.2.1 スレッド ビロック...7 2.2.2 スレッド ブロックのグリッド...8 2.3 メモリ モデル... 10 Chapter 3. ハードウェア 実 装... 13 3.1 オンチップ シェアード メモリ 付 SIMD マルチ プロセッサのセット... 13 3.2 実 行 モデル... 14 3.3 演 算 能 力... 15 3.4 マルチ デバイス... 16 3.5 モード スイッチ... 16 Chapter 4. アプリケーション プログラミング インターフェイス(API)... 17 4.1 C 言 語 での 拡 張... 17 4.2 言 語 の 拡 張... 17 4.2.1 関 数 型 修 飾 子... 18 4.2.1.1 device... 18 4.2.1.2 global... 18 4.2.1.3 host... 18 4.2.1.4 制 限...エラー! ブックマークが 定 義 されていません 4.2.2 修 飾 子 の 変 数 型... 19 4.2.2.1 device... 19 CUDA Programming Guide Version 1.1 iii
4.2.2.2 constant... 19 4.2.2.3 shared... 19 4.2.2.4 Restrictions...エラー! ブックマークが 定 義 されていません 4.2.3 実 行 コンフィグレーション...エラー! ブックマークが 定 義 されていません 4.2.4 組 み 込 み 変 数... 21 4.2.4.1 griddim... 21 4.2.4.2 blockidx... 22 4.2.4.3 blockdim... 22 4.2.4.4 threadidx... 22 4.2.4.5 制 限...エラー! ブックマークが 定 義 されていません 4.2.5 NVCC を 伴 うコンパイル... 22 4.2.5.1 noinline... 22 4.2.5.2 #pragma unroll... 23 4.3 共 通 ランタイム コンポーネント... 23 4.3.1 組 み 込 みベクター 型... 23 4.3.1.1 char1, uchar1, char2, uchar2, char3, uchar3, char4, uchar4, short1, ushort1, short2, ushort2, short3, ushort3, short4, ushort4, int1, uint1, int2, uint2, int3, uint3, int4, uint4, long1, ulong1, long2, ulong2, long3, ulong3, long4, ulong4, float1, float2, float3, float4... 23 4.3.1.2 dim3 型... 23 4.3.2 数 学 的 関 数...エラー! ブックマークが 定 義 されていません 4.3.3 時 間 関 数... 24 4.3.4 テクスチャ 型... 24 4.3.4.1 テクスチャ レファレンスの 宣 言... 24 4.3.4.2 ランタイム テクスチャ 参 照 属 性... 25 4.3.4.3 リニア メモリ 対 CUDA 行 列 によるテクスチャ... 25 4.4 デバイス ランタイム コンポーネント... 26 4.4.1 数 学 関 数... 26 4.4.2 同 期 の 関 数... 26 4.4.3 型 変 換 関 数... 27 4.4.4 型 キャスト 関 数... 27 4.4.5 テクスチャ 関 数... 27 iv CUDA Programming Guide Version 1.1
4.4.5.1 デバイス メモリからのテクスチャリング... 27 4.4.5.2 CUDA 行 列 からのテクスチャリング... 28 4.4.6 原 子 関 数... 28 4.5 ホスト ランタイム コンポーネント...エラー! ブックマークが 定 義 されていません 4.5.1 共 通 概 念... 29 4.5.1.1 デバイス...エラー! ブックマークが 定 義 されていません 4.5.1.2 メモリ...エラー! ブックマークが 定 義 されていません 4.5.1.3 OpenGL 相 互 運 用 性... 30 4.5.1.4 Direct3D 相 互 運 用 性... 30 4.5.1.5 コンカレント 実 行 の 非 同 期... 31 4.5.2 ランタイム API... 32 4.5.2.1 初 期 化...エラー! ブックマークが 定 義 されていません 4.5.2.2 デバイス 管 理... 32 4.5.2.3 メモリ 管 理... 32 4.5.2.4 ストリーム 管 理... 34 4.5.2.5 イベント 管 理... 34 4.5.2.6 テクスチャ 参 照 管 理... 35 4.5.2.7 OpenGL 相 互 運 用 性... 37 4.5.2.8 Direct3D 相 互 運 用 性... 37 4.5.2.9 デバイス エミュレーション モードを 使 ったデバッグ... 37 4.5.3 ドライバ API... 39 4.5.3.1 初 期 化...エラー! ブックマークが 定 義 されていません 4.5.3.2 デバイス 管 理... 39 4.5.3.3 コンテクスト 管 理...エラー! ブックマークが 定 義 されていません 4.5.3.4 モジュール 管 理... 40 4.5.3.5 実 行 制 御... 40 4.5.3.6 メモリ 管 理...エラー! ブックマークが 定 義 されていません 4.5.3.7 ストリーム 管 理... 42 4.5.3.8 イベント 管 理... 43 4.5.3.9 テクスチャ 参 照 管 理... 44 4.5.3.10 OpenGL 相 互 運 用 性... 44 CUDA Programming Guide Version 1.1 v
4.5.3.11 Direct3D 相 互 運 用 性... 44 Chapter 5. 性 能 ガイドライン... 47 5.1 性 能 命 令... 47 5.1.1 命 令 スループット... 47 5.1.1.1 演 算 命 令... 47 5.1.1.2 フロー 命 令 の 制 御... 48 5.1.1.3 メモリ 命 令... 49 5.1.1.4 同 期 命 令...49 5.1.2 メモリ 帯 域 幅... 49 5.1.2.1 グローバル メモリ...エラー! ブックマークが 定 義 されていません 5.1.2.2 定 数 メモリ... 55 5.1.2.3 テクスチャ メモリ... 55 5.1.2.4 シェアード メモリ...エラー! ブックマークが 定 義 されていません 5.1.2.5 レジスタ...エラー! ブックマークが 定 義 されていません 5.2 ブロックあたりのスレッドの 数... 62 5.3 ホストとデバイス 間 のデータ 転 送... 63 5.4 テクスチャ フェッチ 対 グローバルまたは 定 数 メモリ 読 出 し... 63 5.5 総 合 的 な 性 能 の 最 適 化 戦 略... 64 Chapter 6. 行 列 乗 算 の 例... 67 6.1 概 要...エラー! ブックマークが 定 義 されていません 6.2 ソース コードのリスト... 69 6.3 ソース コード ウォークスルー... 71 6.3.1 Mul()... 71 6.3.2 Muld()... 71 Appendix A. 技 術 仕 様... エラー! ブックマークが 定 義 されていません A.1 一 般 仕 様...エラー! ブックマークが 定 義 されていません A.2 標 準 浮 動 小 数 点... 74 Appendix B. 数 学 関 数... 77 B.1 共 通 ランタイム コンポーネント... 77 B.2 デバイス ランタイム コンポーネント... 80 Appendix C. 原 子 関 数... 83 vi CUDA Programming Guide Version 1.1
C.1 算 術 関 数...エラー! ブックマークが 定 義 されていません C.1.1 atomicadd()... 83 C.1.2 atomicsub()... 83 C.1.3 atomicexch()... 83 C.1.4 atomicmin()... 84 C.1.5 atomicmax()... 84 C.1.6 atomicinc()... 84 C.1.7 atomicdec()... 84 C.1.8 atomiccas()... 84 C.2 ビット 単 位 関 数... 85 C.2.1 atomicand()... 85 C.2.2 atomicor()... 85 C.2.3 atomicxor()... 85 Appendix D. ランタイム API 参 照... 87 D.1 デバイス 管 理... 87 D.1.1 cudagetdevicecount()... 87 D.1.2 cudasetdevice()... 87 D.1.3 cudagetdevice()... 87 D.1.4 cudagetdeviceproperties()... 88 D.1.5 cudachoosedevice()... 89 D.2 スレッド 管 理... 89 D.2.1 cudathreadsynchronize()... 89 D.2.2 cudathreadexit()... 89 D.3 ストリーム 管 理... 89 D.3.1 cudastreamcreate()... 89 D.3.2 cudastreamquery()... 89 D.3.3 cudastreamsynchronize()... 89 D.3.4 cudastreamdestroy()... 89 D.4 イベント 管 理...エラー! ブックマークが 定 義 されていません D.4.1 cudaeventcreate()... 90 D.4.2 cudaeventrecord()... 90 CUDA Programming Guide Version 1.1 vii
D.4.3 cudaeventquery()... 90 D.4.4 cudaeventsynchronize()... 90 D.4.5 cudaeventdestroy()... 90 D.4.6 cudaeventelapsedtime()... 90 D.5 メモリ 管 理...エラー! ブックマークが 定 義 されていません D.5.1 cudamalloc()... 91 D.5.2 cudamallocpitch()... 91 D.5.3 cudafree()... 91 D.5.4 cudamallocarray()... 92 D.5.5 cudafreearray()... 92 D.5.6 cudamallochost()... 92 D.5.7 cudafreehost()... 92 D.5.8 cudamemset()... 92 D.5.9 cudamemset2d()... 92 D.5.10 cudamemcpy()... 93 D.5.11 cudamemcpy2d()... 93 D.5.12 cudamemcpytoarray()... 94 D.5.13 cudamemcpy2dtoarray()... 94 D.5.14 cudamemcpyfromarray()... 95 D.5.15 cudamemcpy2dfromarray()... 95 D.5.16 cudamemcpyarraytoarray()... 96 D.5.17 cudamemcpy2darraytoarray()... 96 D.5.18 cudamemcpytosymbol()... 96 D.5.19 cudamemcpyfromsymbol()... 96 D.5.20 cudagetsymboladdress()... 97 D.5.21 cudagetsymbolsize()... 97 D.6 テクスチャ 参 照 管 理... 97 D.6.1 低 レベル API... 97 D.6.1.1 cudacreatechanneldesc()... 97 D.6.1.2 cudagetchanneldesc()... 97 D.6.1.3 cudagettexturereference()... 97 viii CUDA Programming Guide Version 1.1
D.6.1.4 cudabindtexture()... 98 D.6.1.5 cudabindtexturetoarray()... 98 D.6.1.6 cudaunbindtexture()... 98 D.6.1.7 cudagettexturealignmentoffset()... 98 D.6.2 高 レベル API... 98 D.6.2.1 cudacreatechanneldesc()... 98 D.6.2.2 cudabindtexture()... 99 D.6.2.3 cudabindtexturetoarray()... 99 D.6.2.4 cudaunbindtexture()... 99 D.7 実 行 制 御...エラー! ブックマークが 定 義 されていません D.7.1 cudaconfigurecall()...100 D.7.2 cudalaunch()...100 D.7.3 cudasetupargument()...100 D.8 OpenGL 相 互 運 用 性...100 D.8.1 cudaglregisterbufferobject()...100 D.8.2 cudaglmapbufferobject()...101 D.8.3 cudaglunmapbufferobject()...101 D.8.4 cudaglunregisterbufferobject()...101 D.9 Direct3D 相 互 運 用 性...101 D.9.1 cudad3d9begin()...101 D.9.2 cudad3d9end()...101 D.9.3 cudad3d9registervertexbuffer()...101 D.9.4 cudad3d9mapvertexbuffer()...101 D.9.5 cudad3d9unmapvertexbuffer()...102 D.9.6 cudad3d9unregistervertexbuffer()...102 D.9.7 cudad3d9getdevice()...102 D.10 エラーの 取 り 扱 い...102 D.10.1 cudagetlasterror()...102 D.10.2 cudageterrorstring()...102 Appendix E. ドライバ API 参 照...103 E.1 初 期 化...エラー! ブックマークが 定 義 されていません CUDA Programming Guide Version 1.1 ix
E.1.1 cuinit()...103 E.2 デバイス 管 理...103 E.2.1 cudevicegetcount()...103 E.2.2 cudeviceget()...103 E.2.3 cudevicegetname()...103 E.2.4 cudevicetotalmem()...104 E.2.5 cudevicecomputecapability()...104 E.2.6 cudevicegetattribute()...104 E.2.7 cudevicegetproperties()...105 E.3 コンテクスト 管 理...エラー! ブックマークが 定 義 されていません E.3.1 cuctxcreate()...106 E.3.2 cuctxattach()...106 E.3.3 cuctxdetach()...106 E.3.4 cuctxgetdevice()...106 E.3.5 cuctxsynchronize()...106 E.4 モジュール 管 理...エラー! ブックマークが 定 義 されていません E.4.1 cumoduleload()...106 E.4.2 cumoduleloaddata()...107 E.4.3 cumoduleloadfatbinary()...107 E.4.4 cumoduleunload()...107 E.4.5 cumodulegetfunction()...107 E.4.6 cumodulegetglobal()...107 E.4.7 cumodulegettexref()...108 E.5 ストリーム 管 理...108 E.5.1 custreamcreate()...108 E.5.2 custreamquery()...108 E.5.3 custreamsynchronize()...108 E.5.4 custreamdestroy()...108 E.6 イベント 管 理...エラー! ブックマークが 定 義 されていません E.6.1 cueventcreate()...108 E.6.2 cueventrecord()...108 x CUDA Programming Guide Version 1.1
E.6.3 cueventquery()...109 E.6.4 cueventsynchronize()...109 E.6.5 cueventdestroy()...109 E.6.6 cueventelapsedtime()...109 E.7 実 行 制 御...エラー! ブックマークが 定 義 されていません E.7.1 cufuncsetblockshape()...109 E.7.2 cufuncsetsharedsize()...110 E.7.3 cuparamsetsize()...110 E.7.4 cuparamseti()...110 E.7.5 cuparamsetf()...110 E.7.6 cuparamsetv()...110 E.7.7 cuparamsettexref()...110 E.7.8 culaunch()...111 E.7.9 culaunchgrid()...111 E.8 メモリ 管 理...111 E.8.1 cumemgetinfo()...111 E.8.2 cumemalloc()...111 E.8.3 cumemallocpitch()...111 E.8.4 cumemfree()...112 E.8.5 cumemallochost()...112 E.8.6 cumemfreehost()...112 E.8.7 cumemgetaddressrange()...112 E.8.8 cuarraycreate()...113 E.8.9 cuarraygetdescriptor()...114 E.8.10 cuarraydestroy()...114 E.8.11 cumemset()...114 E.8.12 cumemset2d()...114 E.8.13 cumemcpyhtod()...115 E.8.14 cumemcpydtoh()...115 E.8.15 cumemcpydtod()...115 E.8.16 cumemcpydtoa()...116 CUDA Programming Guide Version 1.1 xi
E.8.17 cumemcpyatod()...116 E.8.18 cumemcpyatoh()...116 E.8.19 cumemcpyhtoa()...116 E.8.20 cumemcpyatoa()...117 E.8.21 cumemcpy2d()...117 E.9 テクスチ 参 照 管 理...119 E.9.1 cutexrefcreate()...119 E.9.2 cutexrefdestroy()...119 E.9.3 cutexrefsetarray()...119 E.9.4 cutexrefsetaddress()...120 E.9.5 cutexrefsetformat()...120 E.9.6 cutexrefsetaddressmode()...120 E.9.7 cutexrefsetfiltermode()...120 E.9.8 cutexrefsetflags()...121 E.9.9 cutexrefgetaddress()...121 E.9.10 cutexrefgetarray()...121 E.9.11 cutexrefgetaddressmode()...121 E.9.12 cutexrefgetfiltermode()...121 E.9.13 cutexrefgetformat()...122 E.9.14 cutexrefgetflags()...122 E.10 OpenGL 相 互 運 用 性...122 E.10.1 cuglinit()...122 E.10.2 cuglregisterbufferobject()...122 E.10.3 cuglmapbufferobject()...122 E.10.4 cuglunmapbufferobject()...122 E.10.5 cuglunregisterbufferobject()...123 E.11 Direct3D 相 互 運 用 性...123 E.11.1 cud3d9begin()...123 E.11.2 cud3d9end()...123 E.11.3 cud3d9registervertexbuffer()...123 E.11.4 cud3d9mapvertexbuffer()...123 xii CUDA Programming Guide Version 1.1
E.11.5 cud3d9unmapvertexbuffer()...123 E.11.6 cud3d9unregistervertexbuffer()...123 E.11.7 cud3d9getdevice()...124 Appendix F. テクスチャ フェッチ...125 F.1 直 近 ポイントのサンプリング...126 F.2 リニア フィルタリング...エラー! ブックマークが 定 義 されていません F.3 参 照 テーブル...128 CUDA Programming Guide Version 1.1 xiii
図 表 リスト Figure 1-1. CPUとGPUの 浮 動 小 数 点 演 算 能 力...1 Figure 1-2. GPUはデータ 処 理 用 に 多 くのトランジスタを 割 当 てられる...2 Figure 1-3. CUDAのソフトウェア スタック...3 Figure 1-4. ギャザーとスキャッターのメモリ 動 作...4 Figure 1-5. シェアード メモリはALUにより 緊 密 にデータを 持 ち 込 む...5 Figure 2-1. スレッドの 集 合...9 Figure 2-2. メモリ モデル... 11 Figure 3-1. ハードウェア モデル... 14 Figure 5-1. 結 合 したグローバル メモリ アクセス パターンの 例... 52 Figure 5-2. 非 結 合 グローバル メモリ パターンの 例... 53 Figure 5-3. 非 結 合 グローバル メモリ アクセス パターンの 例. エラー! ブックマークが 定 義 さ れていません Figure 5-4. バンク 競 合 のシェアード メモリ アクセスパターンの 例. エラー! ブックマークが 定 義 されていません Figure 5-5. バンク 競 合 のないシェアード メモリ アクセスパターン 例... 59 Figure 5-6. バンク 競 合 のシェアード メモリ アクセス パターンの 例 エラー! ブックマークが 定 義 されていません Figure 5-7. ブロードキャストのシェアード メモリ 読 出 しアクセス パターンの 例... 61 Figure 6-1. 行 列 乗 法... 68 xiv CUDA Programming Guide Version 1.1
Chapter 1. CUDA の 紹 介 1.1 データ 並 列 演 算 デバイスとしてのグラフィック プロセッサ ユニット わずか 数 年 間 の 事 態 で プログラマブルグラフィックプロセッサユニットはFigure 1-1によって 示 すように 明 確 にコンピューティングの 主 力 製 品 に 発 展 しました マルチ コアが 非 常 に 高 いメモリ 帯 域 幅 によって 動 作 されている 状 態 で 今 日 のGPUはグラフィックスと 非 グラフィックス 処 理 の 両 方 のための 信 じられないリソースを 提 案 します GFLOPS G80GL = Quadro 5600 FX G80 = GeForce 8800 GTX G71 = GeForce 7900 GTX G70 = GeForce 7800 GTX NV40 = GeForce 6800 Ultra NV35 = GeForce FX 5950 Ultra NV30 = GeForce FX 5800 G80GL Figure 1-1. CPUとGPUの 浮 動 小 数 点 演 算 能 力 そのような 発 展 の 主 な 背 景 はGPUが 演 算 集 約 のために 特 化 されるということです 高 度 並 列 演 算 -まさにグラフィックス レンダリングに 対 するように- のために 設 計 されたようなものです より 多 くのトランジスタがデータ キャッシュやフロー 制 御 よりも 図 表 1-2のようにデータ 処 理 用 に 専 念 されています CUDA Programming Guide Version 1.1 1
Chapter 1. Introduction to CUDA Control Cache ALU ALU ALU ALU DRAM DRAM CPU GPU Figure 1-2. GPUはデータ 処 理 用 に 多 くのトランジスタを 割 当 てられる GPU のメモリ 処 理 用 計 算 命 令 に 使 える 面 積 比 率 は 特 にデータ 並 列 演 算 を 表 現 するアドレスの 問 題 に 適 合 します ( 同 一 のプログラムは 高 強 度 の 計 算 を 伴 う 並 列 内 の 多 くのデータ 要 素 を 実 行 します) なぜなら 同 一 プログラムは 高 度 で 洗 練 されたフロー 制 御 を 必 要 としない 各 データ 要 素 を 実 行 します GPUは 強 力 な 計 算 力 を 有 し 処 理 は 多 くのデータ 要 素 上 で 実 行 されます メモリ アクセスの 遅 延 は 大 きなデータ キャッシュの 代 わりの 演 算 能 力 により 表 面 化 しません データ 並 列 処 理 は 並 列 処 理 スレッドへデータ 要 素 を 割 り 当 てます 多 くのアプリケーションは 行 列 のような 大 きなデータ セットを 処 理 し 演 算 処 理 向 上 のためのデータ 並 列 プログラミング モデル を 活 用 できます 3 次 元 のレンダリング 用 の 大 きなピクセルと 頂 点 セットは 並 列 スレッドに 割 り 当 てられます レンダリングした 画 像 のポスト プロセス ビデオ エンコード 画 像 スケーリングや 立 体 視 のような 画 像 とメディア 処 理 アプリケーションに 似 ています パターン 認 識 では 画 像 ブロッ クとピクセルを 並 列 処 理 スレッドに 割 り 当 てることができます つまり 画 像 レンダリング フィール ドの 外 部 の 多 くのアルゴリズムと 処 理 は 一 般 的 な 信 号 処 理 物 理 シミュレ-ション 財 務 予 測 や 生 物 学 的 計 算 などをデータ 並 列 処 理 化 により 加 速 されるのです ところで かねてより 演 算 能 力 はGPU 内 に 潜 在 的 にあったのですが 非 グラフィックス アプリケ ーションのために 効 率 的 に その 演 算 パワーを 扱 うのが 困 難 だったのです GPUは 初 心 者 への 高 い 学 習 カーブと 不 十 分 なAPIしかない 非 グラフィックス アプリケ ーショを 使 ってしかプログラミングできなかったのです GPU DRAM の 一 般 的 読 み 込 み 方 法 として GPU プログラムは DRAM のあらゆる 部 分 からデータ 要 素 を 集 めることができましたが GPU プログラムは 一 般 的 に 書 き 出 すkと ができませんでした GPU プログラムは 如 何 なる DRAM の 部 分 へもスキャッタできない などの CPU で 容 易 に 利 用 可 能 な 多 くのプログラミングの 柔 軟 性 が 欠 けていました いくつかのアプリケーションがGPUの 演 算 能 力 を 利 用 している 時 に DRAM メモリの 帯 域 幅 がボトルネックとなっていました このドキュメントは これらの 課 題 についての 直 接 的 な 回 答 をすべく 真 のジェネリック データ 並 列 演 算 デバイスとしての 目 新 しいGPUハードウェアとプログラミングモデルについて 記 述 してい ます 2 CUDA Programming Guide Version 1.1
Chapter 1. Introduction to CUDA 1.2 CUDA: GPUでの 演 算 のための 新 しいアーキテ クチャ CUDA は Compute Unified Device Architecture の 省 略 で データ 並 列 処 理 デバイスとして 画 像 データ 割 り 当 てのためのAPIを 除 く GPU での 演 算 管 理 をするための 新 しいハードウェアとソフト ウェア アーキテクチャです それは GeForce8 Series Tesla 及 び Quadro で 利 用 可 能 です( 詳 細 に 関 しては Appendix A を 参 照 ください) オペレーティング システムの 多 重 タスキング メカニ ズムは 同 時 に 稼 働 するいくつかの CUDA とグラフィックス アプリケーションが GPU へアクセス するのを 管 理 します CUDA ソフトウェア スタックは 図 表 1-3 のように 複 数 レイヤーから 成 ります:ハードウェア ドライ バ API そのランタイムと2つの 上 位 のレイヤーで 共 通 に 使 う 数 学 ライブラリ CUFFT CUBLAS 等 については 別 のドキュメントに 記 述 しています このハードウェアは 高 性 能 をもたらす 軽 量 なド ライバーやランタイム レイヤーをサポートするように 設 計 されてきました CPU Application CUDA Libraries CUDA Runtime CUDA Driver GPU Figure 1-3. CUDAのソフトウェア スタック CUDA API は 最 小 のラーニングカーブのために 拡 張 C プログラミング 言 語 を 包 含 して います( 第 4 章 を 参 照 してください) CUDA Programming Guide Version 1.1 3
Chapter 1. Introduction to CUDA CUDA は より 多 くの 柔 軟 なプログラミングのために 図 表 1-4 に 示 すような 一 般 的 なスキャッタや ギャザー 両 方 の DRAM メモリ アドレッシングを 提 供 します プログラミングの 見 地 からは まさし く CPU などのように DRAM のどんな 位 置 でもデータを 読 み 書 きする 転 送 できます Control Cache ALU ALU ALU... Control Cache ALU ALU ALU... DRAM d 0 d 1 d 2 d 3 d 4 d 5 d 6 d 7 ギャザー Control Cache ALU ALU ALU... Control Cache ALU ALU ALU... DRAM d 0 d 1 d 2 d 3 d 4 d 5 d 6 d 7 スキャッター Figure 1-4. ギャザーとスキャッターのメモリ 動 作 4 CUDA Programming Guide Version 1.1
Chapter 1. Introduction to CUDA CUDAは 互 いのシェアー データを 使 うスレッドが 並 列 データ キャッシュか 一 般 的 オン チップ シ ェアード メモリをとても 高 速 に 読 み 書 きすることを 特 色 とします( 第 3 章 を 参 照 してください) 図 表 1-5 に 示 すように アプリケーションはオーバー フェッチとラウンド トリップを DRAM 用 に 最 小 化 します したがって DRAM メモリ 帯 域 幅 に 依 存 せずに 利 用 することができます Control Cache ALU ALU ALU... Control Cache ALU ALU ALU... DRAM d 0 d 1 d 2 d 3 d 4 d 5 d 6 d 7 シェアード メモリを 伴 わずに Control Cache ALU ALU ALU... Control Cache ALU ALU ALU... Shared memory Shared memory d 0 d 1 d 2 d 3 d 4 d 5 d 6 d 7 DRAM d 0 d 1 d 2 d 3 d 4 d 5 d 6 d 7 シェアード メモリを 伴 って Figure 1-5. シェアード メモリはALUにより 緊 密 にデータを 持 ち 込 む CUDA Programming Guide Version 1.1 5
Chapter 1. Introduction to CUDA 1.3 本 書 の 構 成 本 書 は 下 記 の 章 から 成 り 立 っています 第 1 章 CUDA の 紹 介. 第 2 章 プログラミング モデルの 概 要 第 3 章 ハードウェアの 実 装 の 記 述 第 4 章 CUDA APIとランタイムの 記 述 第 5 章 どのように 最 高 性 能 を 引 き 出 すかのガイダンス 第 6 章 幾 つかの 簡 単 なサンプル コードでのウォーキング スルーによる 前 章 の 図 解 エラー! 参 照 元 が 見 つかりません 幾 つかのデバイスの 技 術 仕 様 を 示 します Appendix B CUDA でサポートしている 数 学 演 算 子 のリスト Appendix C CUDA でサポートしている 原 子 演 算 子 のリスト Appendix D CUDA ランタイム API レファレンス Appendix E CUDA ドライバーAPIレファレンス Appendix F より 詳 細 なテクスチャ フェッチ 6 CUDA Programming Guide Version 1.1
Chapter 2. プログラミング モデル 2.1 高 度 なマルチスレッド コプロセッサ CUDA を 通 してプログラムされると GPUは 主 CPUのコプロセッサ 並 列 スレッドのとても 多 い 数 を 実 行 する 演 算 デバイスとみなします それはメイン CPU かホストのコプロセッサとして 作 動 しま す 言 い 換 えれば ホストで 動 くアプリケーションが 使 う 並 列 データ 演 算 集 約 的 な 部 分 は GPU へ 任 せます 幾 度 も 繰 り 返 し 実 行 したアプリケーションの 部 分 で しかし 異 なるデータ 上 の 独 立 した1つの 機 能 に 分 離 できたものは 幾 つかの 異 なるスレッドのように このデバイス 上 で 実 行 できる その 趣 旨 で デバイスの 命 令 セットにそのような 機 能 をコンパイルします そして カーネルと 呼 ばれる 結 果 としてのプログラムをデバイスにダウンロードします ホストとそのデバイスの 両 方 は それぞれホストメモリとデバイスメモ リと 呼 ばれたそれ 自 身 の DRAM を 維 持 します あるデータはデバイスの 高 性 能 Direct Memory Access(DMA)エンジンを 活 用 した API の 呼 び 出 しで 1 つの DRAM から 他 方 へデータをコピーすることができます. 2.2 スレッドの 集 合 スレッドの 集 まりは Section2.2.1 と 2.2.2 で 記 述 し また 図 表 2-1 に 表 しているスレッド ブロック のグリッドとして 整 理 されたカーネルを 実 行 します 2.2.1 スレッド ブロック スレッド ブロックはメモリ アクセスを 調 整 するために ある 速 い 共 有 メモリを 通 して 効 率 的 にデ ータを 共 有 して それらの 実 行 を 同 時 にさせることによって 同 期 できるスレッドの 集 まりです ある ものはカーネルで 同 期 ポイントを 指 定 することができます そこでは 同 期 ポイントに 達 するまで ブロックのスレッドがサスペンドしています 各 スレッドはスレッド ID によって 特 定 されます (それは ブロックの 中 のスレッド 番 号 です) スレ ッドIDに 基 づく 複 雑 なアドレッシングを 助 けるために アプリケーションは 任 意 にサイズの2か3 次 元 行 列 としてブロックを 指 定 でき 2-3のコンポーネント インデックスに 代 えて 使 った 各 スレッド を 認 識 します サイズ (D x, D y )の 2 次 元 ブロック 用 のインデックス(x, y)のスレッドのスレッドidは CUDA Programming Guide Version 1.1 7
Chapter 2. Programming Model (x + y D x )ですし サイズ(D x, D y, D z )の3 次 元 ブロック 用 のインデックス(x, y, z)のスレッドのスレッド IDは(x + y D x + z D x D y )です 2.2.2 スレッド ブロックのグリッド ひとつのブロックが 持 てる 最 大 のスレッド 数 には 制 限 があります ところで 同 じ 次 数 とサイズの ブロックの 場 合 は 複 数 のブロックから 一 つのブロックに 集 めた 同 じカーネルで 実 行 します これ による 一 つのカーネル 呼 び 出 しで 立 ち 上 げることができるスレッドの 総 数 はとても 大 きいのです これは 減 少 したスレッドの 協 力 を 犠 牲 にしています なぜなら 同 じグリッドからの 異 なるスレッド のブロック 内 のスレッドは 相 互 の 交 信 と 同 期 ができないからです このモデルでは 異 なる 並 列 能 力 を 伴 う 様 々なデバイス 上 で リコンパイルなしにカーネルが 効 率 的 に 動 作 ができます:そのデ バイスに もしほんの 少 しか 膨 大 な 並 列 能 力 あるいは その 両 方 があれば 通 常 はグリッドの 全 てのブロックを 連 続 稼 動 するかも 知 れません 各 ブロックはグリッド 内 のブロック 番 号 であり それ 自 身 のブロックIDにより 認 識 されます ブロッ クIDに 基 づく 複 雑 なアドレッシングを 助 けるため アプリケーションは 任 意 のサイズの 2 次 元 行 列 としての 一 つのグリッドを 特 定 できます また 2コンポーネント インデックスに 代 えて 使 った 各 ブ ロックを 認 識 します サイズ (D x, D y )の 2 次 元 ブロック 用 のインデックス(x, y)のブロックのブロックi Dは (x + y D x )です 8 CUDA Programming Guide Version 1.1
Chapter 2. Programming Model Host Device Grid 1 Kernel 1 Block (0, 0) Block (1, 0) Block (2, 0) Block (0, 1) Block (1, 1) Block (2, 1) Grid 2 Kernel 2 Block (1, 1) Thread (0, 0) Thread (1, 0) Thread (2, 0) Thread (3, 0) Thread (4, 0) Thread (0, 1) Thread (1, 1) Thread (2, 1) Thread (3, 1) Thread (4, 1) Thread (0, 2) Thread (1, 2) Thread (2, 2) Thread (3, 2) Thread (4, 2) ホストはデバイスにカーネル 呼 び 出 しの 継 続 を 発 行 します スレッドの 集 まりが 幾 つか のスレッド ブロックを 一 つのグリッドとして 整 理 されたので 各 カーネルは 実 行 されま す Figure 2-1. スレッドの 集 合 CUDA Programming Guide Version 1.1 9
Chapter 2. Programming Model 2.3 メモリ モデル スレッドは 図 表 2-2 に 示 した デバイス 上 のそのデバイスが 持 つDRAMと 下 記 のメモリ 空 間 を 経 由 したオンチップ メモリへアクセスする 手 段 のみを 持 っています: スレッド レジスタ 毎 の 読 み 込 み/ 書 き 出 し スレッド ローカル メモリ 毎 の 読 み 込 み/ 書 き 出 し シェアード メモリ ブロック 毎 の 読 み 込 み/ 書 き 出 し グローバル メモリ グリッド 毎 の 読 み 込 み/ 書 き 出 し コンスタント メモリ グリッド 毎 の 読 み 込 み/ 書 き 出 し テクスチャ メモリ グリッド 毎 の 読 み 込 み/ 書 き 出 し グローバル コンスタント 及 びテクスチャ メモリ 空 間 はホストによる 読 み 込 み 或 いは 書 き 出 しがで きます また これらは 同 じアプリケーションにより 永 続 的 にいたるところのカーネルが 起 動 しま す グローバル コンスタント 及 びテクスチャ メモリ 空 間 は 異 なるメモリ 使 用 量 のために 最 適 化 され ます( 図 表 5.1.2.1 5.1.2.2 及 び 5.1.2.3 を 参 照 下 さい) テクスチャ メモリはまた 幾 つかの 特 定 データ フォーマット 用 にデータ フィルタリングや 異 なるアドレッシング モデルを 提 供 します (Section 4.3.4 を 参 照 下 さい) 10 CUDA Programming Guide Version 1.1
Chapter 2. Programming Model Grid Block (0, 0) Block (1, 0) Shared Memory Shared Memory Registers Registers Registers Registers Thread (0, 0) Thread (1, 0) Thread (0, 0) Thread (1, 0) Local Memory Local Memory Local Memory Local Memory Global Memory Constant Memory テクスチャ Memory 1つのスレッドは 様 々な 範 囲 のメモリ 空 間 を 経 由 して そのデバイス 自 身 ののDRAMとオンチッ プ メモリへアクセスします Figure 2-2. メモリ モデル l CUDA Programming Guide Version 1.1 11
Chapter 3. ハードウェア 実 装 3.1 オンチップ シェアード メモリ 付 きSIMDマルチ プロセッサのセット デバイスは Figure 3-1 に 示 すマルチプロセッサのセットとして 実 装 されます 各 マルチプロセッサ は 単 一 命 令 複 データ アーキテクチャ(SIMD)を 持 ちます:あらゆる 与 えられたクロック 周 期 マル チプロセッサの 各 プロセッサは 同 じ 指 示 を 実 行 しますが 異 なったデータを 操 作 します 各 マルチプロセッサは 次 の4つの 型 のオン チップ メモリを 持 ちます: プロセッサあたりのローカル 32 ビット レジスタ を1セット 持 ちます 全 てのプロセッサでシェアーされ シェアード メモリ 空 間 で 実 装 されるパラレル デー タ キャッシュまたはシェアード メモリ 全 てのプロセッサによりシェアーされ 定 数 メモリ 空 間 から 読 出 しが 高 速 になるリード オンリー 定 数 キャッシュ でこれはデバイス メモリのリード オンリー 区 域 として 実 装 され ています 全 てのプロセッサによりシェアーされ テクスチャ メモリ 空 間 から 読 出 しが 高 速 になるリ ード オンリーテクスチャ キャッシュ でこれはデバイス メモリのリード オンリー 区 域 と して 実 装 されています ローカル 及 びグローバル メモリ 空 間 はデバイス メモリのリード-ライト 区 域 として 実 装 され キ ャッシュされません Section2.3 で 言 及 した 様 々なアドレッシング モードとデータ フィルタリング を 与 えるテクスチャ ユニットで 各 マルチプロセッサはテクスチャ キャッシュにアクセスします CUDA Programming Guide Version 1.1 13
Chapter 4. Application Programming Interface Device Multiprocessor N Multiprocessor 2 Multiprocessor 1 Shared Memory Registers Processor 1 Registers Processor 2 Registers Processor M Instruction Unit Constant Cache テクスチャ Cache Device Memory オンチップ メモリ 付 SIMDマルチプロセッサのセット Figure 3-1. ハードウェア モデル 3.2 実 行 モデル スレッド ブロックの1つのグリッドは マルチプロセッサ 上 の 実 行 用 スケジューリング ブロックに よりデバイス 上 で 実 行 されます 各 マルチプロセッサはブロックの 集 まりの1つのブロックを 次 々と 処 理 します 1つのブロックは ただ1つのマルチプロセッサにより 処 理 されます そしてこれは とても 高 速 なメモリへ 読 み 込 む オンチップ シェアード メモリ 内 にシェアード メモリ 空 間 を 備 えています 14 CUDA Programming Guide Version 1.1
Chapter 4. Application Programming Interface 各 マルチプロセッサが 1 つの 集 まりで 幾 つのブロックを 処 理 することができるかは マルチプロセ ッサのレジスタと 共 有 メモリがブロックの 集 まりの 全 てのスレッドの 中 で 分 けられるのでいくつの 1 スレッドあたりのレジスタと どのくらいの 1 ブロックあたりの 共 有 メモリが 与 えられたカーネルに 必 要 であるかよります 少 なくとも 1 ブロックを 処 理 するために 十 分 なレジスタか 1 マルチプロセ ッサあたり 利 用 可 能 な 共 有 メモリがないと カーネルは 起 動 しないでしょう 1 つのマルチプロセッサによって 1 つの 集 まりで 処 理 されるブロックはアクティブであると 呼 ばれ ます それぞれのアクティブなブロックはワープと 呼 ばれるスレッドの SIMD グループに 分 けられ ます:それぞれのこれらのワープは ワープ サイズと 呼 ばれるスレッドの 同 じ 数 を 含 んでいて マルチプロセッサにより SIMD 方 式 で 実 行 されます アクティブ ワープ- 例 えば 全 てのアクティブ ブロックからの 全 てのワープ-はタイム スライス されます:スレッド スケジューラは マルチプロセッサのコンピュータのリソースの 使 用 を 最 大 に するために 定 期 的 に 1 ヶ 所 のワープから 別 のものに 切 り 替 わります ハーフ ワープはワープ の 前 半 か 後 半 のどちらかです ブロックがワープに 分 けられる 方 法 はいつも 同 じです; 各 ワープは スレッド0を 含 む 最 初 のワー プを 伴 うスレッドIDが インクリメントし 連 続 するスレッドを 含 みます Section2.2.1 ではスレッド の ID がブロックでどうスレッドのインデックス リストに 関 連 するかを 説 明 します ブロックの 中 のワープの 発 行 順 序 は 未 定 義 ですが それらの 実 行 は 同 時 にすることができます グローバルか 共 有 メモリ アクセスを 調 整 するために Section2.2.1 で 言 及 します スレッド ブロックのグリッドの 中 のブロックの 発 行 順 序 は 未 定 義 であり ブロック 間 の 同 期 メカニ ズムが 全 くないので 同 じグリッドの 2 つの 異 なったブロックからのスレッドはグリッドの 実 行 の 間 グローバルなメモリを 通 して 安 全 に 互 いに 通 信 することができません もし 非 アトミック 命 令 が ワープの1つ 以 上 のスレッド 用 の グローバルかシェアード メモリ 内 の 同 じロケーションに 書 き 込 むワープにより 実 行 された 時 は それらの 出 現 は 未 定 義 ですが 書 き 込 みの1つは 成 功 するように 保 証 されていて ロケーションと 順 序 へ 出 現 する 順 番 を 書 き 込 みま す もし アトミック 命 令 (Section 1.11.6 を 参 照 下 さい)がグローバル メモリ 内 (このグローバル メモ リとは 全 て 順 序 付 けされ 出 現 したロケーションへワープ 各 読 み 込 み 編 集 書 き 込 みのうちの 1つ 以 上 のスレッドのためのもの)へのワープ 読 み 込 み 編 集 書 き 込 みにより 実 行 されたら そ こに 出 現 したその 順 序 は 未 定 義 になります 3.3 演 算 能 力 デバイスの 演 算 能 力 はメジャー レビジョン 番 号 とマイナー レビジョン 番 号 により 定 義 されます メジャー レビジョン 番 号 付 きデバイスは 同 じコア アーキテクチャです 追 補 Aに 掲 載 したデバイ スは 全 て 演 算 能 力 1.x です(それらのメジャー レビジョン 番 号 は1ですから) マイナー レビジョン 番 号 はコア アーキテクチャの 改 訂 番 号 や 新 機 能 を 含 む 可 能 性 のあるもの に 対 応 しています 様 々な 演 算 能 力 の 技 術 仕 様 は 追 補 Aに 説 明 ある 方 式 で 与 えられます CUDA Programming Guide Version 1.1 15
Chapter 4. Application Programming Interface 3.4 マルチ デバイス 複 数 のGPUの 使 用 はマルチプルGPUシステム 上 の 稼 動 アプリケーションによるCUDAデバイ スとして それらのGPUが 同 じタイプで 動 作 する 時 だけ 保 証 されます もし そのシステムがSLIモードの 場 合 は 全 てのGPUはドライバー スタック 内 の 最 下 位 でフュ ーズしますので 1つのGPUしかCUDAデバイスとして 使 えません 各 GPUを 独 立 したものとして 見 えるためには SLIモードをCUDAのためにコントロール パネル を オフ にしておく 必 要 があります 3.5 モード スイッチ GPUはプライマリ サーフェスと 呼 ばれる 幾 つかのDRAMメモリに 専 念 します プライマリ サー フェスは ユーザーによる 表 示 出 力 している 際 に 表 示 装 置 のリフレッシュに 使 います ユーザーがディスプレイの 解 像 度 やビットの 深 さの 切 り 替 えによりディスプレイのモード スイッチ を 起 動 した 際 に(NVIDIA コントロール パネルや Windows のディスプレイ コントロール パネルを 使 って) 相 当 なメモリをプライマリ サーフェスの 変 更 用 に 必 要 とします 例 えば ユーザーがデ ィスプレイの 解 像 度 を from 1280x1024x32-bit to 1600x1200x32-bit へ 変 更 した 場 合 システム はプライマリ サーフェス 用 に 5.24MB より 多 くの 7.68MB を 専 念 させなくてはなりません (アン チ エイリアシングを 伴 うフル スクリーン グラフィックス アプリケーションの 場 合 は 更 に 多 くのデ ィスプレイ メモリをプライマリ サーフェスに 必 要 とします ) Windows において フル スクリーン DirectX アプリケーションの 起 動 や コンピュータをロックするための Ctrl+Alt+Del 操 作 を 含 むデ ィスプレイ モードの 切 り 替 えの 起 動 という 他 のイベントをする 場 合 も 同 様 です もし モード スイッチがプライマリ サーフェス 用 に 必 要 な 相 当 量 のメモリを 増 加 させるなら シス テムはCUDAアプリケーションに 専 念 しているメモリを 奪 い 合 わなければなりませんので 結 果 と してそれらのアプリケーションはクラッシュするかも 知 れません 16 CUDA Programming Guide Version 1.1
Chapter 4. Application Programming Interface Chapter 4. アフ リケーション フ ロク ラミンク インターフェイス(API) 4.1 C 言 語 での 拡 張 CUDAプログラミング インターフェイスの 目 的 は ユーザーが 容 易 にCに 近 いプログラミング 言 語 で デバイスによる 実 行 用 のプログラミング 記 述 のための 比 較 的 簡 単 なパスを 提 供 することで す これは 下 記 から 成 り 立 っています: 最 小 限 のC 言 語 への 拡 張 セットは Section4.2 に 述 べています:プログラマはデバイス 上 で 実 行 するためのソース コードの 部 分 を 対 象 とします; ランタイム ライブラリは 以 下 へ 分 割 します: ホスト コンポーネント(Section4.5 に 述 べています)はホストで 稼 動 し ホストからの1つ 以 上 の 演 算 デバイスにアクセスし 制 御 する 機 能 を 提 供 します; デバイス コンポーネント(Section4.4.で 述 べています)はデバイスで 稼 動 し デバイスに 特 化 した 機 能 を 提 供 します; 共 通 コンポーネント(Section4.3 に 述 べています)はビルト イン ベクター 型 とホストとデバ イス コードの 両 方 でサポートしているC 標 準 ライブラリのサブセットを 提 供 します これらは 共 通 ランタイム コンポーネントにより 提 供 された 関 数 であり デバイスで 稼 動 するのに サポートしている C 標 準 ライブラリからの 関 数 のみに 重 視 されるべきです 4.2 言 語 の 拡 張 Cプログラミング 言 語 の 拡 張 は 下 記 の4つです: ホストまたはデバイスで 実 行 するか 無 関 係 に さらにホストまたはデバイスから 呼 び 出 せるかに 関 係 なく 指 定 する 関 数 型 修 飾 詞 (Section4.2.1); 変 数 のデバイス 上 でのメモリ ロケーションを 指 定 する 変 数 型 (Section4.2.2); CUDA Programming Guide Version 1.1 17
Chapter 4. Application Programming Interface カーネルがホストからのデバイス 上 でどのように 稼 動 するかを 指 定 する 新 ディテクティ ブ(Section4.2.3); グリッド ブロックの 次 数 ブロック 及 びスレッド インデックスの4つの 組 み 込 み 変 数 (Section4.2.4) これらの 拡 張 子 を 含 む 各 ソース ファイルは Section4.2.5 に 概 略 の 述 べている CUDA コンパイ ラ nvcc を 伴 ってコンパイルしなければなりません nvcc の 詳 細 な 記 述 は 別 のドキュメントで 読 むことができます これらの 拡 張 子 のそれぞれは 以 下 の Section 毎 に 制 限 を 記 述 しています nvcc はこれらの 制 限 の 同 じ 警 告 上 のエラーまたはワーニングを 与 えるでしょうが それらのいくつかは 発 覚 できま せん 4.2.1 関 数 型 修 飾 子 4.2.1.1 device device 修 飾 子 は 次 の 機 能 を 宣 言 します: デバイスでの 実 行 デバイスからのみ 呼 び 出 し 可 能 4.2.1.2 global global 修 飾 子 は 存 在 としてのカーネルの 機 能 を 宣 言 します その 機 能 とは; デバイスでの 実 行 ホストからのみ 呼 び 出 し 可 能 4.2.1.3 host host 修 飾 子 は 次 の 機 能 を 宣 言 します; デバイスでの 実 行 ホストからのみ 呼 び 出 し 可 能 それは host 修 飾 子 のみを 伴 う 機 能 を 宣 言 するか またはあらゆる host, device か global 修 飾 子 を 伴 わないのと 等 価 なものです いずれの 場 合 の 機 能 もホストだけのためにコン パイルされます ところで host 修 飾 子 は device 修 飾 子 と 組 み 合 わせで 使 うことができます この 場 合 の 関 数 はホストとデバイスの 両 方 に 対 してコンパイルされます 4.2.1.4 制 限 device と global 関 数 は 帰 納 をサポートしません device と global 関 数 はそれらの 本 体 内 の 静 的 変 数 を 宣 言 できません device と global 関 数 は 引 数 の 変 数 番 号 を 保 有 することはできません device 関 数 はそれらのアドレスを 持 つことはできません 他 方 で global 関 数 への 関 数 ポイ ンターはサポートされます 18 CUDA Programming Guide Version 1.1
Chapter 4. Application Programming Interface global と host 修 飾 子 は 一 緒 に 使 用 できません global 関 数 はボイド 復 帰 型 を 持 っていなければなりません global 関 数 を 呼 ぶあらゆるものは Section4.2.3 に 述 べている 実 行 コンフィギュレーションを 指 定 しなければなりません global 関 数 を 呼 ぶものは デバイスが 実 行 の 完 了 を 意 味 する 非 同 期 です. global 関 数 パラメータは 現 在 デバイスへのシェアード メモリを 経 由 して 送 られ 256 バイトの 制 限 があります 4.2.2 修 飾 子 の 変 数 型 4.2.2.1 device device 修 飾 子 は 変 数 がデバイスに 存 在 することを 宣 言 します 他 の 型 の 修 飾 子 は 変 数 が 属 するメモリ 空 間 で 更 に 指 定 し device を 伴 って 使 う 際 に 次 の3つ の Section で 定 義 します もし それらの 表 示 がないなら 変 数 は: グローバル メモリ 空 間 に 存 在 しています アプリケーションのライフタイムを 保 有 しています グリッド 内 部 の 全 てのスレッドとランタイム ライブラリを 経 由 してホストからからアクセス できます 4.2.2.2 constant constant 修 飾 子 は device をオプションとして 一 緒 に 使 って 変 数 を 宣 言 します: メモリ 空 間 定 数 に 存 在 しています アプリケーションのライフタイムを 保 有 しています グリッド 内 部 の 全 てのスレッドとランタイム ライブラリを 経 由 してホストからからアクセス できます 4.2.2.3 shared shared 修 飾 子 は device をオプションとして 一 緒 に 使 って 変 数 を 宣 言 します: スレッド ブロックのシェアード メモリ 空 間 に 存 在 します ブロックのライフ タイムを 保 有 します ブロック 内 部 の 全 てのスレッドからのみアクセスできます これらはスレッドを 渡 る 弛 緩 順 序 付 けであっても スレッド 内 部 のシェアード 変 数 の 完 全 順 次 整 合 性 です syncthreads()(section4.4.2)の 実 行 後 でのみ 見 えるように 保 証 した 他 のスレッドから 書 き 込 みます 変 数 が 揮 発 性 として 宣 言 でもしない 限 り コンパイラは 読 み 出 しを 最 適 化 するの に 自 由 であり 前 の 宣 言 が 満 たされる 限 りシェアード メモリに 書 き 込 みます: 変 数 を 次 の 外 部 行 列 のようにシェアード メモリ 内 で 宣 言 している 際 に extern shared float shared[]; 行 列 サイズは 起 動 時 に 決 定 しています(Section4.2.3) 全 ての 変 数 はメモリの 同 じアドレスで 開 始 し この 方 式 で 宣 言 します CUDA Programming Guide Version 1.1 19
Chapter 4. Application Programming Interface 4.2.2.4 制 限 行 列 内 の 変 数 の 配 置 はオフセットを 経 由 して 明 示 的 に 管 理 しなければなりません 例 えば もし それが 以 下 の 方 程 式 を 望 んだ 場 合 short array0[128]; float array1[64]; int array2[256]; ダイナミックに 割 り 当 てたシェアード メモリでは それは 以 下 の 方 法 で 行 列 を 宣 言 して 初 期 化 す るかも 知 れません extern shared char array[]; device void func() // device or global function { short* array0 = (short*)array; float* array1 = (float*)&array0[128]; int* array2 = (int*)&array1[64]; } これらの 修 飾 子 は struct と union メンバーで 正 式 なパラメータで そしてホストで 実 行 する 関 数 内 部 のローカル 変 数 では 許 可 されません shared と constant 変 数 は 暗 黙 の 静 的 ストレージを 保 有 します device, shared 及 び constant 変 数 は extern キーワードを 使 った 外 部 として 宣 言 できま せん device と constant 変 数 はファイル 有 効 範 囲 でのみ 許 可 されます constant 変 数 はホスト ランタイム 関 数 を 経 由 してのホストからのみとデバイスから 割 り 当 て できません(Section4.5.2.3 と 4.5.3.6) shared 変 数 はそれらの 宣 言 子 の 一 部 として 初 期 化 を 持 つことができません 一 般 的 にこれらのあらゆる 修 飾 子 がなくとも デバイス コードで 宣 言 された 自 動 変 数 は レジス タにあります また 一 方 幾 つかの 場 合 コンパイラはローカル メモリにそれを 置 くかも 知 れませ ん これはしばしば 余 りに 多 くのレジスタ 空 間 を 費 やす 大 きな 構 造 か 行 列 やコンパイラが 定 数 量 付 きインデックスを 決 定 できない 行 列 を 示 します ptx アッセンブリ コードの 検 査 (t ptx か-keep 付 きコンパイリングにより 取 得 したもの)は ld.local と st.local ニーモニックを 使 用 して 変 数 が 宣 言 されますから 最 初 のコンパイル 段 階 の 間 local ニーモニックを 使 用 することでローカル メモリに 置 かれ アクセスできるかどうかを 伝 えるで しょう もっとも それが 対 象 としたアーキテクチャのために 余 りに 多 くのレジスター 空 間 を 費 やすことが 判 明 したなら その 後 のコンパイル 段 階 はそのまま 他 の 方 法 を 決 めるかも 知 れません これはロ ーカル メモリ 使 用 量 (lmem)を 報 告 する--ptxas-options=-v オプションでコンパイルすることによ りチェックできます デバイスで 実 行 するコードのポインターは コンパイラがそれらをシェアード メモリ 空 間 をしてい るか 否 かに 関 係 なく 解 決 することができ グローバルなメモリ 空 間 であるかぎりサポートされま す さもなければ それらはグローバル メモリ 空 間 で 割 り 当 てるか または 宣 言 するメモリを 示 す だけのために 制 限 されます コード 内 のグローバルかシェアード メモリに 修 飾 子 参 照 するポインターは 大 抵 は 分 離 の 失 敗 かアプリケーションの 終 了 で 未 定 義 の 挙 動 内 のデバイス 結 果 を 実 行 した ホストかコード 内 の ホスト メモリで 実 行 します a device, shared か constant 変 数 のアドレスで 取 得 したアド レスはデバイス コードでのみ 使 用 することができます その device か constant 変 数 は 20 CUDA Programming Guide Version 1.1
Chapter 4. Application Programming Interface Section4.5.2.3 で 述 べている cudagetsymboladdress() 経 由 して 取 得 したアドレスはホスト コード でのみ 使 用 できます 4.2.3 実 行 コンフィグレーション global 関 数 を 呼 ぶすべてのものは それを 呼 び 出 すために 実 行 コンフィグレーションを 指 定 し なければなりません 実 行 コンフィグレーションは デバイスで 実 行 する 関 数 を 使 うのにグリッドとブロックの 次 数 を 定 義 します 同 様 にや 関 連 するストリーム(Section4.5.1.5 にストリームについて 述 べています)でも それは 挿 入 している<<< Dg, Db, Ns, S >>>の 関 数 名 から 括 弧 内 の 引 数 の 間 からくる 式 によって 指 定 されます ここで: Dim3(Section4.3.1.2 を 参 照 下 さい) 型 には Dg があり 起 動 開 始 ブロックの 数 と Dg.x * Dg.y が 同 等 なようにグリッドのサイズと 次 数 を 指 定 します Dg.z は 未 使 用 です; Dim3(Section4.3.1.2 を 参 照 下 さい) 型 には Db があり ブロック 当 りのスレッド 数 と Db.x * Db..y * Db.z が 同 等 なように 各 ブロックのサイズと 次 数 を 指 定 します; size_t 型 には Ns があり 静 的 に 割 り 当 てられたメモリに 加 えるこの 呼 び 出 しのために ブロック 毎 に 動 的 に 割 り 当 てるシェアード メモリ 内 のバイト 数 を 指 定 します この 動 的 に 割 り 当 てられたメモリは Section4.2.2.3 に 言 及 している 外 部 行 列 として 宣 言 した 変 数 の 全 てによって 使 用 されたものです;Ns は0をデフォルトとするオプション 引 数 です; cudastream_t には S があり 関 連 するストリームを 指 定 します S は0をデフォルトとす るオプション 引 数 です 宣 言 された 関 数 の 例 global void Func(float* parameter); このように 呼 び 出 さなくてはなりません: Func<<< Dg, Db, Ns >>>(parameter); 実 行 コンフィグレーションのための 引 数 は 実 関 数 引 数 の 前 に 関 数 引 数 のように 評 価 され 現 状 で はそれはデバイスへのシェアード メモリ 経 由 でパスされます 関 数 呼 び 出 しは もし Dg か Db が Appendix A.1 にて 指 定 したデバイス 用 に 許 された 最 大 サイズ より 大 きいか または Ns が 静 的 割 り 当 て 関 数 引 数 や 実 行 コンフィグレーションに 必 要 なシェア ード メモリの 容 量 を 差 し 引 いたデバイスで 可 能 なシェアード メモリの 最 大 容 量 より 大 きいと 失 敗 します 4.2.4 組 み 込 み 変 数 4.2.4.1 griddim この 変 数 は dim3 型 (Section4.3.1.2 を 参 照 )で グリッドの 次 数 を 含 んでいます CUDA Programming Guide Version 1.1 21
Chapter 4. Application Programming Interface. 4.2.4.2 blockidx 4.2.4.3 blockdim 4.2.4.4 threadidx 4.2.4.5 制 限 この 変 数 は uint3 型 (Section4.3.1.3 を 参 照 )で グリッド 内 部 のブロック インデックスを 含 んでい ます この 変 数 は dim3 型 (Section4.3.1.2 を 参 照 )で ブロックの 次 数 を 含 んでいます この 変 数 は uint3 型 (Section4.3.1.3 を 参 照 )で ブロック 内 部 のスレッド インデックスを 含 んでい ます あらゆる 組 み 込 み 変 数 のアドレスの 取 得 を 許 可 しません あらゆる 組 み 込 み 変 数 への 値 の 割 り 当 てを 許 可 しません 4.2.5 NVCC を 伴 うコンパイル nvcc はコンパイルしているCUDAコードの 処 理 を 単 純 にするコンパイラ ドライバです: 簡 単 で 身 近 なコマンド ライン オプションを 提 供 し 異 なるコンパイル 段 階 を 実 装 するツールの 収 集 を 呼 び 出 すことで 実 行 します nvcc の 基 本 的 流 れはホスト コードからデバイス コードを 分 離 するのと バイナリ フォームまた は cubin オブジェクトへコンパイルしているデバイス コードをから 成 り 立 ちます 生 成 したホスト コードは 別 のツールを 使 用 してコンパイルした 残 りか 最 終 コンパイル 段 階 にてホスト コンパイ ラが 直 接 呼 び 出 したオブジェクト コードとしての Cコードとしての 出 力 です アプリケーションは 生 成 したホスト コードを 無 視 するか CUDAドライバ API(Section4.5.3 を 参 照 下 さい)を 使 った デバイス 上 の cubin オブジェクトのロードか 実 行 のどちらもできます または それらはグローバルに 初 期 化 したデータ 行 列 としての cubin オブジェクトを 含 んんでいたり 必 要 なCUDAランタイム スタートアップ コードからロード 及 び 起 動 した 各 コンパイル カーネル (Section4.5.2 を 参 照 下 さい)へ Section4.2.3 に 述 べている 実 行 コンフィグレーション 構 文 の 変 換 を 含 んで 生 成 したホスト コードにリンクすることができます C++の 構 文 ルールによるとコンパイラのフロント エンドは CUDA ソースファイルを 処 理 します フ ル C++はホスト コード 用 にサポートされます また 一 方 C++ののサブセット C だけはデバイス コードを 全 てサポートします;C++の 基 本 ブロック 内 部 変 数 のクラス 継 承 または 宣 言 子 などの 特 定 の 機 能 は 違 います C++ 構 文 ルールを 使 用 した 帰 結 として 無 効 ポインター( 例 ;malloc()によ る 返 し)は 型 に 嵌 っていない 非 無 効 ポインターへの 割 り 当 てができません nvcc の 流 れの 詳 細 な 記 述 とコマンド オプションはこれ 以 外 のドキュメントで 読 めます nvcc は 以 下 の Section で 説 明 している2つのコンパイラ 指 示 文 で 紹 介 しています 4.2.5.1 noinline デフォルトで device 関 数 はいつもインラインされています 22 CUDA Programming Guide Version 1.1
Chapter 4. Application Programming Interface noinline 関 数 修 飾 子 はできれば 関 数 についてどんなインラインのためでなく コンパイラのた めならヒントとして 使 えます 関 数 本 体 はそれが 呼 ばれたのと 同 じファイル 内 に まだあるに 違 い ありません コンパイラは noinline 修 飾 子 をポインター パラメータ 付 き 関 数 用 と 大 きなパラメータ リスト 付 き 関 数 用 に 引 き 受 けません 4.2.5.2 #pragma unroll デフォルトでコンパイラは 既 知 のトリップ カウント 付 きの 小 さなループを 展 開 します すべての 与 えられたループを 展 開 しながら 制 御 するのに#pragma unroll 命 令 を 使 用 できます それはループ の 直 前 に 置 かれなくてはならなく ループに 対 してだけ 適 用 されます その 数 字 はオプションでル ープを 何 回 展 開 しなければならないかをという 指 示 することになります 例 として このコード サンプル 内 で: #pragma unroll 5 for (int i = 0; i < n; ++i) このループは 5 回 展 開 します それはプログラマー 次 第 で その 展 開 はプログラムの 正 当 性 に 影 響 を 与 えないでしょう( 上 記 の 例 でもし n が 5 より 小 さい 場 合 はそうかも 知 れません) #pragma unroll 1 はコンパイラがループを 展 開 するのを 防 止 するでしょう もしトリップ カウンターが 定 数 で 数 値 が 全 く#pragma unroll の 後 に 指 定 されないなら ループは 展 開 でき さもなければ 全 く 展 開 できません 4.3 共 通 ランタイム コンポーネント 共 通 ランタイム コンポーネントはホストとデバイス 関 数 の 両 方 により 使 用 できます 4.3.1 組 み 込 みベクター 型 4.3.1.1 char1, uchar1, char2, uchar2, char3, uchar3, char4, uchar4, short1, ushort1, short2, ushort2, short3, ushort3, short4, ushort4, int1, uint1, int2, uint2, int3, uint3, int4, uint4, long1, ulong1, long2, ulong2, long3, ulong3, long4, ulong4, float1, float2, float3, float4 4.3.1.2 dim3 型 これらはから 基 本 整 数 と 浮 動 小 数 点 型 から 生 成 されたベクター 型 です それらは 構 造 体 で 1 番 目 から 4 番 目 のコンポーネントはフィールド x,y,z と w それぞれを 経 由 してアクセスできます こ れらは 全 てフォーム make_<type name>の 構 造 体 関 数 とともに 来 ます 例 として int2 make_int2(int x, int y); 値 (x, y)を 伴 う int2 型 のベクターを 生 成 します この 型 は 次 数 を 指 定 するのに 使 用 される uint3 に 基 づく 整 数 ベクター 型 です dim3 型 の 変 数 を 定 義 する 際 に 不 特 定 の 状 態 で 残 っている どんなコンポーネントも1に 初 期 化 します CUDA Programming Guide Version 1.1 23
Chapter 4. Application Programming Interface 4.3.2 数 学 的 関 数 デバイスで 実 行 されると Table B-1 には 各 エラー 領 域 と 共 に 現 在 サポートしている 数 学 的 関 数 C/C++ 標 準 ライブラリの 総 覧 を 収 めています ホスト コードで 実 行 されると 与 えられた 関 数 は 可 能 ならば C ランタイム 実 装 を 使 います 4.3.3 時 間 関 数 clock_t clock(); これは 各 クロック 周 期 で 増 加 されるカウンターの 値 を 返 します カーネルの 最 初 と 最 後 のカウンターを 抽 出 し 2つのサンプルの 違 いを 取 得 し 完 全 にスレッドを 実 行 したデバイスにより 取 得 した クロック サイクルの 数 の 各 スレッドあたりの 測 定 結 果 を 記 録 します しかし デバイスで 実 行 したスレッド 命 令 を 費 やしたクロック サイクルの 数 ではありませ ん 前 の 数 はタイム スライスした 最 後 のスレッドよりも 大 きいです 4.3.4 テクスチャ 型 CUDAはテクスチャ メモリにアクセスするグラフィックス 用 GPUのテクスチャリング ハードウェア のサブセットをサポートします グローバル メモリの 代 わりにテクスチャ メモリから 読 み 込 んだデータは Section5.4 に 記 述 した 幾 つかの 性 能 利 得 を 得 ることができます テクスチャ メモリは Section4.4.5 に 記 述 したテクスチャ フェッチと 呼 ばれるデバイス 関 数 を 使 っ てカーネルから 読 み 込 みます 最 初 のテクスチャ フェッチのパラメータはテクスチャ レファレン スと 呼 ばれるオブジェクトを 指 定 します テクスチャ レファレンスはテクスチャメモリの 部 分 をフェッチするか 定 義 します それはホスト ラ ンタイム 関 数 (Section4.5.2.6 お 呼 び 4.5.3.9)を 経 てテクスチャと 呼 んでいるメモリのある 領 域 に 結 合 されなければなりません これ 以 前 に それはカーネルにより 使 用 できます 幾 つかの 特 殊 テ クスチャ レファレンスは 同 じテクスチャかメモリ 内 の 重 複 したテクスチへ 結 合 されるかも 知 れませ ん テクスチャ レファレンスは 幾 つかの 属 性 を 保 有 しています それらの 一 つは テクスチャがテ クスチャ 座 標 を 使 用 する1 次 元 行 列 か 2 次 元 行 列 として 記 述 するかどうかを 2つのテクスチャを 使 用 することで 指 定 する 次 元 数 が 調 整 されます 行 列 の 要 素 は テクスチャ エレメント を 短 縮 してテクセルと 呼 ばれます 他 の 属 性 は どのように 入 力 座 標 が 割 り 込 まれて 処 理 されたかとい うのと 同 様 に テクスチャ フェッチの 入 出 力 データ 型 を 定 義 します 4.3.4.1 テクスチャ レファレンスの 宣 言 幾 つかのテクスチャ レファレンスの 属 性 は 不 変 で コンパイルする 時 に 分 かっていなければなら なく それらはテクスチャ レファレンスを 宣 言 した 時 に 指 定 されます テクスチャ レファレンスは テクスチャ 型 の 変 数 としてファイル スコープで 宣 言 されます: 24 CUDA Programming Guide Version 1.1
Chapter 4. Application Programming Interface テクスチャ<Type, Dim, ReadMode> texref; ここで: Type はテクスチャをフェッチしたときに 返 されるデータ 型 です;Type は Section4.3.1.1 に 記 述 している 基 本 的 整 数 浮 動 小 数 点 型 と 1-, 2- 及 び 4-コンンポーネント ベクタ ー 型 に 制 限 します; Dim はテクスチャ 参 照 の 1 か2の 次 数 を 指 定 します;Dim はデフォルトを1とするオプシ ョン 引 数 です; ReadMode は cudareadmodenormalizedfloat か cudareadmodeelementtype と 同 等 です;もし それが cudareadmodenormalizerfloat で Type が16ビットか8ビット 整 数 型 で その 値 は 最 大 範 囲 の 符 号 なし 整 数 型 用 の[0.0, 1.0]と 符 号 付 整 数 型 用 の[-1.0, 1.0]にマップされたときに 実 際 に 浮 動 小 数 点 として 値 を 返 します; 例 えば 値 0xff を 伴 う 符 号 なし8ビットテクスチャ 要 素 を1と 読 む もし それが cudareadmodeelementtype なら 変 換 は 実 行 しません ReadMode は cudareadelementtype をデフォルトとするオ プション 引 数 です 4.3.4.2 ランタイム テクスチャ 参 照 属 性 テクスチャ 参 照 の 他 の 属 性 は 易 変 で ホストからホスト ランタイムのときに 変 えることができます (ランタイム API については Section4.5.2.6 ドライバーAPI いついては 4.5.3.9) それらはテクス チャ 座 標 が 正 規 化 かどうか アドレッシング モード 及 びテクスチャ フィルタリングを 以 下 のように 詳 細 に 指 定 します デフォルトでテクスチャは[0, N]の 範 囲 内 の 浮 動 小 数 点 座 標 を 使 って 参 照 されます このとき N は 座 標 に 対 応 する 寸 法 のテクスチャのサイズです 例 えば サイズの 64x32 にあるテクスチャ は x と y 次 数 の 座 標 が[0, 63]と[0, 31]に 参 照 されます 正 規 化 テクスチャ 座 標 で [0, N]の 代 わ りに[0.0, 1.0]の 範 囲 に 座 標 を 指 定 します そして 同 じ 64x32 のテクスチャは x, y 両 座 標 の[0, 1] 範 囲 内 に 正 規 化 してアドレスされます 正 規 化 テクスチャ 座 標 は もしそれが テクスチャ サイ ズ 如 何 に 関 係 なくテクスチャ 座 標 が 望 ましいなら 幾 つかのアプリケーションの 要 求 に 自 然 に 適 合 します アドレッシング モードはテクスチャ 座 標 が 範 囲 外 のときにどうなるかを 定 義 します 非 正 規 化 テ クスチャ 座 標 を 使 用 するとき [0, N] 範 囲 外 のテクスチャ 座 標 はクランプされています 値 が 0 未 満 は 0 に 設 定 され N 以 上 は N-1 に 設 定 されます クランピングはまた 正 規 化 テクスチャ 座 標 を 使 用 しているときはデフォルト アドレッシング モードです:0.0 未 満 か 1.0 を 超 える 値 は[0.0, 1.0]の 範 囲 にクランプされます 正 規 化 座 標 として ラップ アドレッシング モードも 指 定 される かも 知 れません テクスチャが 周 期 的 信 号 を 含 むときに 通 常 はラップ アドレッシングが 使 用 さ れます それはテクスチャ 座 標 の 断 片 的 部 分 だけを 使 用 します 例 えば 1.25 は 0.25 に -1.25 は 0.75 と 同 じように 扱 われるということです リニア テクスチャ フィルタリングは 浮 動 小 数 点 データを 返 すための 構 成 されるテクスチャ 用 にだけ 実 行 されるかも 知 れません それは 隣 接 の テクセル 間 の 低 い 精 度 の 補 間 をします 可 能 になると テクスチャ フェッチ 位 置 を 囲 むテクセル が 読 まれ テクスチャ フェッチのリターン 値 はテクスチャ 座 標 がテクセルの 間 で 落 下 したところに 基 づいた 状 態 で 補 間 されます 簡 単 なリニア 補 間 は 一 次 元 テクスチャのために 実 行 されます そ して バイ リニアの 補 間 は 二 次 元 テクスチャのために 実 行 されます Appendix F はテクスチャ フェッチングのより 詳 細 なことを 記 述 しています 4.3.4.3 リニア メモリ 対 CUDA 行 列 によるテクスチャ CUDA Programming Guide Version 1.1 25
Chapter 4. Application Programming Interface テクスチャはリニア メモリか CUDA 行 列 のどんな 領 域 にあるかもしれません(Section4.5.1.2 を 参 照 下 さい) テクスチャはリニア メモリに 割 り 当 てられます: 次 数 =1の 時 のみ 実 行 できます: テクスチャ フィルタリングはサポートしません; 非 正 規 化 整 数 テクスチャ 座 標 をしたときのみアドレスできます; 前 のアドレッシング モードはサポートしません; 範 囲 外 のときにテクスチャ アクセスは 0を 返 します ハードウェアは 整 列 要 求 にテクスチャ ベース アドレスで 実 行 します プログラマから 行 列 要 求 を 抽 出 するために この 関 数 はデバイス メモリ 上 へテクスチャ 参 照 を 拘 束 し ます デバイス メモリにテクスチャ 参 照 を 拘 束 する 関 数 は 必 要 なメモリから 読 むために テクスチャ フェッチに 適 用 するパス バックされた 1 バイトを 戻 します CUDA の 配 分 ルーチンで 返 されたベース ポインタはこの 整 列 規 制 に 一 致 しています そしてアプリ ケーションは 割 り 当 てられたポインタを cudabind テクスチャ()/cuTexRefSetAddress() に 通 過 することによって 全 体 でオフセットを 避 けることができます 4.4 デバイス ランタイム コンポーネント デバイス ランタイム コンポーネントはデバイス 関 数 でのみ 使 用 できます 4.4.1 数 学 関 数 Table B-1 の 幾 つかの 関 数 は それほど 正 確 ではありませんが より 速 いバージョンはデバイ ス ランタイム コンポーネントに 存 在 しています; それで ( sin(x)のような)と 共 に 同 じ 名 前 を 前 に 置 いています それらの 組 み 込 み 関 数 は それらの 各 エラー 結 合 と 共 に Table B-2 に 列 記 しています コンパイラに 存 在 しているなら あらゆる 関 数 にそれほど 正 確 でないカウンターパー トにコンパイルさせるオプション(-use_fast_math)があります 4.4.2 関 数 の 同 期 void syncthreads(); ブロック 内 のすべてのスレッドを 同 期 します すべてのスレッドが 一 旦 このポイントに 達 すると 通 常 は 実 行 を 再 開 します syncthreads()は 同 じブロックのスレッドのコミュニケーションを 調 整 するのに 使 用 されます ブ ロックの 中 のいくつかのスレッドが 共 有 されたかグローバルなメモリの 同 じアドレスにアクセスす るとき それらは 潜 在 的 なリード-アフター-ライト ライト-アフター-リードまたはライト-アフター- ライトのメモリ アクセスの 危 険 性 があります これらがアクセスする 中 間 のスレッドを 連 動 させる ことによって これらのデータ 危 険 を 避 けることができます syncthreads() は 条 件 が 全 体 のスレッド ブロックにわたり 完 全 に 同 じと 評 価 した 時 に 条 件 付 コ ードを 許 可 します さもなければコード 実 行 は 故 意 でない 副 作 用 に 掛 かるか ハングを 発 生 しそ うです 26 CUDA Programming Guide Version 1.1
Chapter 4. Application Programming Interface 4.4.3 型 変 換 関 数 以 下 での 関 数 における 接 尾 語 は IEEE-754 丸 めモードを 示 します: rn は round-to-nearest-even のことです rz は round-towards-zero のことです ru は round-up のことです ( 正 の 無 限 大 へ), rd は round-down のことです( 負 の 無 限 大 へ). int float2int_[rn,rz,ru,rd](float); 指 定 した 丸 めモードを 使 用 して 浮 動 小 数 点 の 引 数 を 整 数 に 変 換 します unsigned int float2uint_[rn,rz,ru,rd](float); 指 定 した 丸 めモードを 使 用 して 浮 動 小 数 点 の 引 数 を 符 号 なし 整 数 に 変 換 します float int2float_[rn,rz,ru,rd](int); 指 定 した 丸 めモードを 使 用 して 整 数 の 引 数 を 浮 動 小 数 点 に 変 換 します float uint2float_[rn,rz,ru,rd](unsigned int); 指 定 した 丸 めモードを 使 用 して 符 号 なし 整 数 の 引 数 を 浮 動 小 数 点 に 変 換 します 4.4.4 型 キャスト 関 数 float int_as_float(int); 値 を 変 更 せず 整 数 の 引 数 に 浮 動 小 数 点 の 型 キャストを 実 行 します 例 えば int_as_float(0xc0000000)は-2 と 等 しいです performs a floating-point type cast on the integer argument, leaving the value unchanged. For example, int_as_float(0xc0000000) is equal to -2. int float_as_int(float); 値 を 変 更 せず 浮 動 小 数 点 の 引 数 に 整 数 の 型 キャストを 実 行 します 例 えば float_as_int(1.0f)は to 0x3f800000 と 等 しいです 4.4.5 テクスチャ 関 数 4.4.5.1 デバイス メモリからのテクスチャリング デバイス メモリからのテクスチャリングのときに テクスチャは tex1dfetch() 関 数 群 と 共 にアクセスされます; 例 として: template<class Type> Type tex1dfetch( texture<type, 1, cudareadmodeelementtype> texref, int x); float tex1dfetch( texture<unsigned char, 1, cudareadmodenormalizedfloat> texref, int x); float tex1dfetch( texture<signed char, 1, cudareadmodenormalizedfloat> texref, CUDA Programming Guide Version 1.1 27
Chapter 4. Application Programming Interface int x); float tex1dfetch( texture<unsigned short, 1, cudareadmodenormalizedfloat> texref, int x); float tex1dfetch( texture<signed short, 1, cudareadmodenormalizedfloat> texref, int x); これらの 関 数 は テクスチャ 座 標 x を 使 用 することでテクスチャ 参 照 texref に 拘 束 されたリニア メモリの 範 囲 をとって 来 ます テクスチャ フィルタリングとアドレッシング モードは 全 くサポートさ れません 整 数 型 のために これらの 関 数 は 整 数 を 32 ビットの 浮 動 小 数 点 に 任 意 にプロモート するかもしれません そのうえ 上 に 示 された 関 数 2 と 4 倍 はサポートされます; 例 えば: float4 tex1dfetch( texture<uchar4, 1, cudareadmodenormalizedfloat> texref, int x); テクスチャ 座 標 x を 使 用 することでテクスチャ 参 照 texref に 拘 束 したリニア メモリをとって 来 ま す 4.4.5.2 CUDA 行 列 からのテクスチャリング CUDA 行 列 からテクスチャリングするとき テクスチャは tex1d() か tex2d() と 共 にアクセスされます: template<class Type, enum cudatexturereadmode readmode> Type tex1d(texture<type, 1, readmode> texref, float x); template<class Type, enum cudatexturereadmode readmode> Type tex2d(texture<type, 2, readmode> texref, float x, float y); これらの 関 数 は テクスチャ 座 標 x と y を 使 用 することでテクスチャ 参 照 texref に 縛 られた CUDA 行 列 をとって 来 ます テクスチャ 参 照 の 不 変 (コンパイル 時 )の そして 可 変 (ランタイム) の 属 性 の 組 み 合 わせは 座 標 がどんな 処 理 がテクスチャ フェッチの 間 で 起 こるか 返 し 値 はテ クスチャ フェッチにより 配 信 したか 読 み 取 るかを 決 定 します (Section 4.3.4.1 と 4.3.4.2 を 参 照 下 さい) 4.4.6 原 子 関 数 原 子 関 数 は 演 算 能 力 1.1 のデバイス 用 でのみ 可 能 です それらは Appendix C にリストされています 原 子 関 数 はグローバル メモリに 存 在 する 1 つの 32 ビ ット ワードの 読 込 み- 編 集 - 書 込 み 原 子 演 算 子 を 実 行 します 例 えば atomicadd()はグローバル メモリの 同 じアドレスにある 32 ビット ワードを 読 込 み 整 数 をアドします そして 同 じアドレスに 結 果 を 書 込 みます その 演 算 子 はセンス 内 の 原 子 で 他 のスレッドから 干 渉 なく 実 行 することを 保 証 されます 言 い 換 えれば 演 算 子 が 完 全 になるまで 他 のどんなスレッドもこのアドレスにアクセスすることができません 原 子 演 算 子 は 32 ビット 符 号 付 か 符 号 なし 整 数 を 伴 うときのみ 稼 動 します 28 CUDA Programming Guide Version 1.1
Chapter 4. Application Programming Interface 4.5 ホスト ランタイム コンポーネント ホスト ランタイム コンポーネントはホスト 関 数 によってのみ 使 用 できます それは 操 作 するため に 下 記 の 関 数 を 提 供 します: デバイス 管 理 コンテキスト 管 理 メモリ 管 理 コード モジュール 管 理 実 行 制 御 テクスチャ 参 照 管 理 OpenGL と Direct3D の 可 搬 性 それは2つのAPIで 構 成 されます: CUDA driver API と 呼 ばれている 低 レベルの API CUDA runtime API と 呼 ばれている CUDA ドライバ API の 上 の 実 装 されている 高 レベ ルの API それらの API は 相 互 に 排 他 的 です:アプリケーションは1つか 他 方 を 使 用 しなければなりません CUDA ランタイムは 暗 黙 の 初 期 化 コンテキスト 管 理 及 びモジュール 管 理 を 提 供 することによっ て デバイスコード 管 理 を 容 易 にします Nvcc により 生 成 されたCホスト コードは CUDA ランタ イムに 依 存 します(Section4.2.5 を 参 照 下 さい) そして アプリケーションは CUDA ランタイム API を 使 用 しなければならない このコードにリンクする コントラストに 於 いて CUDA ドライバ API はさらにコードを 要 求 します このことはプログラミングやデバッギングには 難 しいが cubin オブジェクトに 対 処 するだけであるので より 良 い 管 理 水 準 を 提 供 して 言 語 に 依 存 していません (Section4.2.5 を 参 照 下 さい) CUDA ドライバーAPI を 使 用 するカーネルを 構 成 して 始 動 する のは 特 に 難 しいです 明 白 なファンクション コールが Section4.2.3 で 説 明 された 実 行 構 成 構 文 の 代 わりにある 状 態 で 実 行 構 成 とカーネル パラメタを 指 定 しなければなりませんから また デバイス エミュレーション(Section4.5.2.7 を 参 照 下 さい)は CUDA ドライバ API では 動 作 しませ ん CUDA ドライバ API は cuda ダイナミック ライブラリを 経 由 して 配 信 され それらの 全 てのエ ントリー ポイントは cu の 接 頭 語 です CUDA ランタイム API は cudart ダイナミック ライブラリを 経 由 して 配 信 され それらの 全 てのエントリー ポイントは cuda の 接 頭 語 です 4.5.1 共 通 概 念 4.5.1.1 デバイス 両 方 の API は カーネル 実 行 のためにシステムの 上 で 利 用 可 能 なデバイスを 列 挙 するために 関 数 を 提 供 して 彼 らの 特 性 について 問 い 合 わせし それらの 1 つを 選 択 します(ランタイム API の ために Section4.5.2.2 とドライバ API のために Section4.5.3.2 を 参 照 下 さい) 幾 つかのホスト スレッドは 同 一 デバイスでデバイス コードを 実 行 できます しかし 設 計 により 1 つのホスト スレッドはデバイス コードを1つのデバイスでしか 実 行 できません 結 果 として 複 数 のホストスレッドが 複 数 のデバイスでデバイス コードを 実 行 するのに 必 要 となります さらに 別 のホスト スレッドからのランタイムはランタイムを 経 由 して 1 つのホスト スレッドで 作 成 された どんな CUDA リソースも 使 用 することはできません CUDA Programming Guide Version 1.1 29
Chapter 4. Application Programming Interface 4.5.1.2 メモリ デバイス メモリはリニア メモリまたは CUDA 行 列 のいずれかとして 割 り 当 てできます リニア メモリは32ビット アドレス 空 間 のデバイスに 存 在 します 例 えば 別 々に 割 り 当 てられたエンティ ティーは 二 分 木 のポインタ 経 由 でお 互 いに 参 照 できます CUDA 行 列 はテクスチャ フェッチするために 最 適 化 された 不 透 明 なメモリ レイアウトです (Section4.3.4 を 参 照 下 さい) それらは1 次 元 か2 次 元 で エレメントの 集 合 です それぞれは 1, 2か4のコンポーネントを 保 有 し おそらく 符 号 付 か 符 号 なし 8-, 16-または 32 ビット 整 数 16 ビ ット( 現 在 ドライバ API を 経 由 してのみサポートしています)か 32 ビット 浮 動 です CUDA 行 列 は 単 にカーネルでテクスチャのフェッチすることで 読 込 み 可 能 であり 同 じ 数 の 詰 まっ ているコンポーネントでテクスチャ 参 照 に 拘 束 されるだけかもしれません リニア メモリと CUDA 行 列 の 両 方 は Section4.5.2.3 及 び 4.5.3.6 に 記 述 していているメモリ コピー 関 数 経 由 でホストに より 読 込 み 可 能 で 書 込 み 可 能 です また ホスト ランタイムは malloc()によって 割 り 当 てられた 通 常 のページ- 可 能 なホストメモリと 対 照 的 に 割 り 当 てる 関 数 とフリー ページ - 固 定 ホスト メモ リ- を 提 供 します(ランタイム API のための SectionD.5.6 及 び D.5.7 とドライバ API のための E8.5 及 び E8.6 を 参 照 下 さい) ページ 固 定 メモリの1つの 優 位 点 は もしホスト スレッドによるデータ 交 換 を 実 行 するためにのみ ホスト メモリに 割 り 当 てられたものがページ 固 定 に 割 り 当 てられた ら ホスト メモリとデバイス メモリ 間 のバンド 幅 が 高 いことです ページ 固 定 メモリは 希 少 リソー スです それでページ 固 定 されたメモリにおける 配 分 はページ- 可 能 なメモリにおける 配 分 のずっ と 前 に 失 敗 し 始 めるでしょう ページング 用 にオペレーティングシステムに 利 用 可 能 な 物 理 的 なメ モリの 量 を 減 少 させることで あまりに 多 くのページ 固 定 メモリを 割 り 当 てると 総 合 システム 性 能 は 抑 えられます 4.5.1.3 OpenGL の 相 互 運 用 性 OpenGL バッファ オブジェクトは CUDA のアドレッシング 内 にマップされるでしょう CUDA が OpenGL によって 記 述 されたデータを 読 出 すか または CUDA が OpenGL で 費 やされるために データを 書 込 むのを 可 能 にするどちらかで Section4.5.2.7 でどのようにランタイム API で 実 行 されるかを 記 述 し Section4.5.3.10 ではドライバ API について 説 明 します 4.5.1.4 Direct3D の 相 互 運 用 性 Direct3D 9.0 頂 点 バッファは CUDA のアドレス 空 間 内 部 へマップされるでしょう Direct3D により 記 述 された DUDA がデータを 読 出 すか Direct3D により 費 やされるために CUDA がデータを 書 き 込 むのを 可 能 のするどちらかで Section4.5.2.8 でどのようにランタイム API で 実 行 されるか を 記 述 し Section4.5.2.8 ではドライバ API について 説 明 します CUDA コンテクストは 一 度 の1つだけの Direct3D デバイスを 伴 い 相 互 運 用 するでしょう 最 初 / 最 後 の 関 数 を 呼 び 出 すことは Section4.5.2.8 及 び 4.5.3.11 に 記 述 しています CUDA コンテクストtpDirect3D デバイスは 同 じ GPU で 生 成 されなければなりません これはラン タイム API 用 の cudad3d9getdevice()(sectiond.9.7 を 参 照 下 さい) またはドライバ API 用 の cud3d9getdevice()(sectione.11.7 を 参 照 下 さい)を 使 用 して Direct3D によって 使 用 されるアダプターに 対 応 する CUDA デバイスについて 問 い 合 わせすることで 確 実 にすることが できます Direct3D デバイスは D3DCREATE_HARDWARE_VERTEXPROCESSING フラッグ 付 で 生 成 されな ければなりません CUDA 以 下 を 未 だサポートしていません 30 CUDA Programming Guide Version 1.1
Chapter 4. Application Programming Interface Direct3D 9.0 以 外 のバージョン 頂 点 バッファ 以 外 の Direct3D オブジェクト また Direct3D ドライバと CUDA コンテクストが 異 なったドライバで Direct3D と CUDA の 負 荷 バ ランスが 相 互 運 用 性 より 好 まれる 場 合 に 作 成 されるのを 保 証 するために cudad3d9getdevice() か cud3d9getdevice を 使 用 できます 4.5.1.5 コンカレント 実 行 の 非 同 期 ホストとデバイス 間 のコンカレントの 実 行 を 容 易 にするための 幾 つかのランタイ ム 関 数 は 非 同 期 です: デバイスが 要 求 されたタスクを 完 了 する 前 に 制 御 をア プリケーションに 返 します それらは: カーネルは global 関 数 または cugridlaunch() 及 び cugridlaunchasync()を 経 由 して 起 動 します; メモリコピーを 実 行 して Async で 接 尾 される 関 数 ; デバイスとデバイスの 双 方 向 でのメモリコピーを 実 行 する 関 数 ; メモリをセットする 関 数 ; また 幾 つかのデバイスはページ 固 定 したホスト メモリとデバイス メモリ 間 でカーネル 実 行 を 伴 うコンカレントにコピーを 実 行 できます アプリケーションは CU_DEVICE_ATTRIBUTE_GPU_ OVERLAP 付 きの cudevicegetattribute()を 呼 び 出 すことで この 機 能 を 問 い 合 わせするでしょう (それぞれ SectionE.2.6 を 参 照 下 さい) この 機 能 は 現 在 では cudamallocpitch() (Section4.5.2.3 を 参 照 下 さい)か cumemallocpitch()(section4.5.3.6 を 参 照 下 さい)を 経 由 して 割 り 当 てられた CUDA 行 列 か 2D 行 列 にかかわらないメモリコピーのためだけにサポー トされます アプリケーションはストリームを 経 由 してコンカレントに 管 理 します ストリームは その 命 令 で 実 行 する 関 数 の 順 序 です 他 方 で 異 なったストリームは 個 別 の 順 序 外 の 関 数 をもう 他 方 かコンカレ ントに 実 行 するでしょう ストリームはストリ-ミング オブジェクトを 生 成 することで 定 義 され ストリーム パラメータとして カーネル 起 動 の 順 序 とホストとデバイスの 双 方 向 のコピーを 指 定 します Section4.5.2.4 ではこ れをランタイム API と 共 に Section4.5.3.7 ではドライバ API と 共 にどのように 実 行 したかを 記 述 しています すべての 先 行 関 数 の 後 でのみ ゼロ ストリーム パラメタで 指 定 した あらゆるカーネルの 起 動 メモリのセット またはメモリーのコピーが 始 まります ストリームの 一 部 の 関 数 を 含 み 後 続 でな い 関 数 はそれが 完 了 するまで 始 まります ランタイム API 用 の cudastreamquery() 及 びドライバ API 用 の custreamquery()(sectiond.3.2 及 び E.5.2 のそれぞれを 参 照 下 さい)はストリーム 内 の 全 ての 先 行 関 数 が 完 結 しているなら そ れを 知 るための 方 法 をアプリケーションに 提 供 します ランタイム API 用 の cudastreamsynchronize()とドライバ API 用 の custreamsynchronize() (SectionE.5.2 及 び E.5.3 のそれぞれを 参 照 下 さい)はストリーム 内 の 全 ての 先 行 関 数 が 完 結 ま で 待 つ ランタイムを 明 示 的 に 強 制 するための 方 法 を 提 供 します ランタイム API 用 の cudathreadsynchronize()とドライバ API 用 の cuctxsynchronize() (SectionD.2.1 及 び E.3.5 のそれぞれを 参 照 下 さい)アプリケーションはストリーム 内 の 全 ての 先 行 タスクが 完 結 するまで 待 つ ランタイムを 強 制 できます 不 要 なスローダウンを 避 けるために タイミング 目 的 や 起 動 の 隔 離 やメモリ コピーが 失 敗 しているときに これらの 関 数 を 使 用 するの は 最 も 良 いです CUDA Programming Guide Version 1.1 31
Chapter 4. Application Programming Interface ランタイムもまたデバイスの 進 捗 を 密 接 にモニタして アプリケーションでそれらのイベントが 記 録 されたときに プログラムとクエリーのあらゆるポイントのイベントを 非 同 期 に 記 録 することを 送 出 することで 正 確 なタイミングを 実 行 する 方 法 を 提 供 します イベントはイベントが 完 結 する 前 な ら 全 てのタスク(または 全 ての 関 数 はストリームを 与 えたもの)を 記 録 します Section4.5.2.5 は ランタイム API で そして Section4.5.3.8 はドライバ API で これをどのように 実 行 するかを 記 述 しています 異 なるストリームからの2つの 関 数 は もし ページ 固 定 したホスト メモリの 割 り 当 てか デバイ ス メモリ 割 り 当 てか デバイス メモリ セットか デバイス/デバイス 間 の 双 方 向 メモリ コピーの いずれかはコンカレントに 動 作 できません また イベントはそれらの 間 の 記 録 を 発 生 します プログラマは CUDA_LAUNCH_BLOCKING 環 境 変 数 を1に 設 定 することにより システムで 動 作 する 全 ての CUDA アプリケーションのための 非 同 期 実 行 を グローバルに 無 効 にすることができ ます この 機 能 をデバッグ 目 的 だけに 供 給 すべきであり 決 してプロダクション ソフトウェアを 確 実 に 動 作 させる 方 法 として 使 用 するべきではありません 4.5.2 ランタイム API 4.5.2.1 初 期 化 ランタイム API 用 の 明 白 な 初 期 化 関 数 はありません; ランタイム 関 数 が 呼 ぶ 1 回 目 を 初 期 化 し ます それはランタイム 関 数 が 呼 んだタイミングと 最 初 のランタイムへの 呼 び 出 しエラー コード を 解 釈 したときを 記 憶 しておく 必 要 があります 4.5.2.2 デバイス 管 理 SectionD.1 の 関 数 はシステム 内 のデバイス プレゼントを 管 理 するのに 使 います cudagetdevicecount() 及 び cudagetdeviceproperties() はデバイスを 数 えて それらの 特 性 を 検 索 するの 方 法 を 提 供 します: int devicecount; cudagetdevicecount(&devicecount); int device; for (device = 0; device < devicecount; ++device) { cudadeviceprop deviceprop; cudagetdeviceproperties(&deviceprop, device); } cudasetdevice() はホスト スレッド 関 連 のデバイスを 選 択 するのに 使 います cudasetdevice(device); あらゆる global 関 数 や Appendix D からのどんな 関 数 も 呼 ばれる 前 に デバイスを 選 択 しなけ ればなりません cudasetdevice()への 明 白 なコールでこれをしないなら 自 動 的 にデバイス 0 を 選 択 します そして その 後 の cudasetdevice()へのどんな 明 白 なコールも 効 果 はないでしょう 4.5.2.3 メモリ 管 理 SectionD.5 の 関 数 はデバイス メモリの 割 り 当 てや 開 放 とメモリがホストとデバイス メモリの 間 で グローバルなメモリ 空 間 および 転 送 データで 宣 言 された あらゆる 変 数 のためにも 割 り 当 てたア クセスに 使 われます リニア メモリは cudamalloc() や cudamallocpitch()を 使 った 割 り 当 てや 32 CUDA Programming Guide Version 1.1