CUDA, OpenCL ン 実習
濱田 剛
長崎大学工学部
超高速 コ コン ュ ン セン
hamada@cis.nagasaki-u.ac.jp
14：45～16：15 •CUDA, OpenCL ン 実習CUDA OpenCL環境構築
•CUDA入門
•CUDA ッ 使い方
•OpenCL入門
Contents
実習環境 い
C/C++ 言語＊ 64-bit Ｊ inux (CentOS 5)
ン用端 シン (80 )
Core2Duo
GPU い い ！
Happy Hacking Keyboard
計算用 シン
GPU ス (128 1 式 )
Core2Quad
GeForce9800GTX+ x 2GPUs
CUDA ( ), OpenCL
CUDA: (Compute unified device architecture)
NVIDIA 提供 NVIDIA 製 GPU 開発環境
無料
C/C++ 文法 GPU ン 能
ン SPMD
ン セス , gather/scatter operation
OpenCL
AMD, NVIDIA 製 GPU, Cell, コ CPU 利用 能 開発
環境
ＡＵＢＡ う 汎用的 ン 環境 NVIDIA 製 GPU 以外
利用 う 群
当然対象 階層 深い 複雑
始 前
GPGPU習得 奥義
い
ぶ 挫折
(奥義そ 1)
ュ 最初 べ
ぶ 挫折
(奥義そ 2)
SDK サン コ
最初 べ
CUDA 学習 質 関係
(奥義そ 3)
理屈 学ぶ前
あえ 動 ．
CUDA 質 理解 コ
手
CUDA学習用コ ： ／
例題 500個 全 sqrt
// 関数 (GPU 処理部 )
__global__
void calc_on_gpu (float *y)
{
int i = threadIdx.x;
y[i] = sqrt(y[i]);
CUDA学習用コ ： ０
例題 500個 全 sqrt
// ス 処理 int main() {
int n = 500;
int nb = sizeof(float) * n; float *x = (float*) malloc(nb); float *y;
cudaMalloc( (void**) &y, nb); for(int i=0; i<n; i++) x[i] = i;
cudaMemcpy( y, x, nb, cudaMemcpyHostToDevice);
calc_on_gpu <<<1, 500 >>> ( y ); // 関数 (500個 ス ッ 処理) cudaMemcpy(x, y, nb, cudaMemcpyDeviceToHost);
return 0; }
CUDAコ コン 実行
CUDA SDK CUDA Toolkit ン
& ンス ＆後 明 )
例題 保存 (test.cu)
例題 以 場所 あ
~/GPGPU-seminar-CUDA/sample1
コン
nvcc test.cu
コ 何 い ？
＆動作 仕組 い 明＇
CUDA 基
移動 ス ッ 並列処理
必 GPU 移動 処理
そ 処理 関数 呼ぶ
Host GPU GPU
CUDA学習用コ
// ス 処理 int main() {
int n = 500;
int nb = sizeof(float) * n; float *x = (float*) malloc(nb); float *y;
cudaMalloc( (void**) &y, nb); for(int i=0; i<n; i++) x[i] = i;
cudaMemcpy( y, x, nb, cudaMemcpyHostToDevice); calc_on_gpu <<<1, 500 >>> ( y );
cudaMemcpy(x, y, nb, cudaMemcpyDeviceToHost); return 0;
}
// ~/GPGPU-seminar-CUDA/sample1/test.cu
ＧＮＵ 確保
転送 ( ス → GPU
転送
(GPU → ス )
500個 ス ッ
関数 処理
ス ッ 番号
// 関数 (GPU 処理部 )
__global__
void calc_on_gpu (float *y)
{
int i = threadIdx.x; // 各ス ッ 毎 番号
y[i] = sqrt(y[i]);
CUDA学習用コ
３．．個 ス ッ 関数(calc_on_gpu) 並列実行！！
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ ス ッ
ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ ス ッ
ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ ス ッ
ス ッ
ス ッ ス ッ
ス ッ ス ッ
ス ッ ス ッ ス ッ
ス ッ
ス ッ ス ッ ス ッ
ス ッ
ス ッ ス ッ
Block( ッ ) い
// ス 処理 int main() {
int n = 500;
int nb = sizeof(float) * n; float *x = (float*) malloc(nb); float *y;
cudaMalloc( (void**) &y, nb); for(int i=0; i<n; i++) x[i] = i;
cudaMemcpy( y, x, nb, cudaMemcpyHostToDevice);
<<<1, 500 >>>
// ~/GPGPU-seminar-CUDA/sample1/test.cu
何？？？
……… 答え = block
GPU
1 ～ 2 個 GPU
構成
GPU
30個 SM
(Streaming Multiprocessor) 構成
SM, SP
8個 SP
(Streaming processor) 構成．
各SP 浮動小数点
積和算器
NVIDIA GT200
SP SP SP SP
SP SP SP SP Shared Memory
SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM
GPU
DRAM DRAM DRAM
×
）
<<<1, 500>>> 1 ッ 500ス
ッ いう意味
学習コ
500 個 ス ッ 1 個 SM(Stream
Multiprocessor) 中 け 動作 い
SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM
多数 ッ 使用 方法
// ス 処理 int main() {
int n = 60000; // 120 block , 500 thread/block; int nb = sizeof(float) * n;
float *x = (float*) malloc(nb); float *y;
cudaMalloc( (void**) &y, nb); for(int i=0; i<n; i++) x[i] = i;
cudaMemcpy( y, x, nb, cudaMemcpyHostToDevice); calc_on_gpu
<<<120, 500 >>>
( y );cudaMemcpy(x, y, nb, cudaMemcpyDeviceToHost); return 0;
}
// ~/GPGPU-seminar-CUDA/sample2/test.cu
計算 増や
→ そ 前 GPU う い 明
ッ 数 増や
若干修正
__global__
void calc_on_gpu (float *y)
{
int tid = threadIdx.x; // 各ス ッ 毎 番号
int bid = blockIdx.x; // 各 ッ 毎 番号
int i = 500 * bid + tid;
y[i] = sqrt(y[i]);
<<<120, 500>>>
120 ッ 500ス ッ いう意味
SM あ 6 ッ 替え 実行
GPU 全体 ／４．．．ス ッ 動作
SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM
SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM SM
CUDA/OpenCL
GPU ンス
(実演)
自身
GPGPU環境 構築
CUDA /OpenCL ンス
Initlevel 3 変更
念 Xサ 落
CUDA /OpenCL ンス
ンス 起動
CUDA /OpenCL ンス
CUDA /OpenCL ンス
CUDA /OpenCL ンス
CUDA /OpenCL ンス
CUDA /OpenCL ンス
CUDA /OpenCL ンス
CUDA /OpenCL ンス
CUDA /OpenCL ンス
以 GPU ンス 終了 ．
明 簡略化 ＊ 再起動
CUDA
Toolkit & SDK
ンス
CUDA Toolkit ンス
CUDA SDK ンス 起動
Root
CUDA ンス
環境設定 開く
以 う CUDA 必要 環境変数 追加
CUDA SDK ンス
CUDA SDK ンス 起動
追加 環境変数 効
CUDA SDK ンス
Enter (標準 場所 ンス 場合)
Enter (標準 場所 CUDA Toolkit ンス い 場合)
CUDA SDK ンス
ンス 終了 様子
CUDA
SDKサン コ
コン
サン コ ＊ＳＢＩ
コン
多少無駄 あ
cd ~NVIDIA_CUDA_SDK/
make
＊サン コ コン
べ 自動 行わ ．
最低限必要 コン 作業
Common
CUDA ッ
(CUDA-GDB)
入門
CUDA-GDB ？
GDB CUDA コ ッ う 機
能拡張
Host GPU 側 コ 対象
CUDA-GDB 2.2 Linux 64 bit 版 対応
CUDA-GDB 2.2 用 ュ ン 公開
コン 起動
ッ ション付 コン
nvcc -g -G foo.cu -o foo
CUDA-GDB 起動
cuda-gdb foo
動
(cuda-gdb) break myFirstKernel
(cuda-gdb) run
(cuda-gdb) next
(cuda-gdb) print idx
(cuda-gdb) finish
(cuda-gdb) quit
GDB 拡張コ ン
現在 focus い ス ッ 表示
(cuda-gdb) thread
中 ス ッ 替え
例 : ッ (0, 0) ス ッ (1, 0, 0) 替え
(cuda-gdb) thread <<<(0, 0), (1, 0, 0)>>>
全ス ッ 状態一覧
(cuda-gdb) info cuda threads all
起動 breakpoint設定 実行
ス ッ 情報表示 ス ッ 実行
ス ッ 替え ス ッ 実行
nvcc ッ 用 ション
--debug (-g)
Generate debug information for host code.
--device-debug <level> (-G)
Generate debug information for device code, plus
also specify the optimization level for the device
code in order to control its 'debuggability.
Allowed values for this option: 0,1,2,3.
OpenCL 入門
OpenCL
OpenCL
AMD, NVIDIA 製 GPU, Cell, コ CPU 利用 能 開発
環境
ン 業界標準 API 策定 行う非営利団体 OpenCLKhronos
Group 仕様策定
現在 Nvidia 実装 公開 & 無料 利用 能
ＡＵＢＡ う 汎用的 ン 環境 NVIDIA 製 GPU 以外
利用 う 群
当然対象 階層 深い 複雑
＊考え方 CUDA 同
＊ く ス ッ
CUDA OpenCL 対応
CUDA OpenCL 同 点
ス GPU 移動 必要 あ
多数 ス ッ 実行
ス ッ 識別番号 用い ス ッ 動作 指定
CUDA OpenCL 異 点
OpenCL 方 細 手続 多い
CUDA OpenCL 対応
＆ス ッ 識別＇
CUDA OpenCL
threadIdx.x get_local_id(0)
blockIdx.x get_group_id(0)
blockDim.x get_local_size(0)
blockIdx.x * blockDim.x + threadIdx.x get_global_size(0)
CUDA OpenCL 似 う 識別子 取得
⇒単一 複数 ス ッ 並列実行
基 的 ン CUDA 同
サン 明
各 work-item 自 識別子ＩＢ 使 配
列 ＩＢ番目 要素 ＩＢ 値 入 サン
→ く単純 動作
＆以 実演 交え 明 ＇
ＭｐｅｎＡＪ関係 変数
// OpenCL Vars
cl_context cxGPUContext; // OpenCL context
cl_command_queue cqCommandQue; // OpenCL command que cl_device_id* cdDevices; // OpenCL device list
cl_program cpProgram; // OpenCL program cl_kernel ckKernel; // OpenCL kernel
cl_mem cmDevSrcA; // OpenCL device source buffer A cl_mem cmDevSrcB; // OpenCL device source buffer B cl_mem cmDevDst; // OpenCL device destination buffer size_t szGlobalWorkSize; // 1D var for Total # of work items
size_t szLocalWorkSize; // 1D var for # of work items in the work group size_t szParmDataBytes; // Byte size of context information
size_t szKernelLength; // Byte size of kernel code cl_int ciErr1, ciErr2; // Error code var
char* cPathAndName = NULL; // var for full paths to data, src, etc. char* cSourceCL = NULL; // Buffer to hold source for compilation
work-item数 指定
// Length of float arrays to process int iNumElements = 11444777;
// set and log Global and Local work size dimensions SzLocalWorkSize = 256;
// rounded up to the nearest multiple of the LocalWorkSize
szGlobalWorkSize = shrRoundUp((int)szLocalWorkSize, iNumElements);
処理内容
// OpenCL Kernel Function for element by element vector addition
__kernel void VectorAdd(__global float* c, __global int iNumElements) {
// get index into global data array int iGID = get_global_id(0);
// bound check (equivalent to the limit on a 'for' loop for standard/serial C code
if (iGID >= iNumElements) {
return; }
// add the vector elements c[iGID] = iGID;
}
Host側 確保
// Allocate and initialize host arrays
dst = (void *)malloc(sizeof(cl_float) * szGlobalWorkSize);
cl_mem_flags Description
CL_MEM_READ_WRITE カ ー ネ っ て 読み 書き さ CL_MEM_WRITE_ONLY カ ー ネ write only CL_MEM_READ_ONLY カ ー ネ read only
CL_MEM_USE_HOST_PTR ー タ の 格納先と し て ホ ス の を 指定
スコン ス コ ン ュ
作成
// Create the OpenCL context on a GPU device
cxGPUContext = clCreateContextFromType(0, CL_DEVICE_TYPE_GPU, NULL, NULL, &ciErr1);
// Get the list of GPU devices associated with context
clGetContextInfo( cxGPUContext, CL_CONTEXT_DEVICES, 0, NULL, &szParmDataBytes);
cdDevices = (cl_device_id*)malloc(szParmDataBytes);
clGetContextInfo( cxGPUContext, CL_CONTEXT_DEVICES,
szParmDataBytes, cdDevices, NULL);
// Create a command-queue CqCommandQue =
clCreateCommandQueue(cxGPUContext, cdDevices[0], 0, &ciErr1);
ス側 確保
// Allocate the OpenCL buffer memory objects on the device GMEM. cmDevDst = clCreateBuffer(cxGPUContext, CL_MEM_WRITE_ONLY,
sizeof(cl_float) * szGlobalWorkSize, NULL, &ciErr2);
込 コン
// Name of the file with the source code for the computation kernel const char* cSourceFile = "VectorInit.cl";
// Read the OpenCL kernel in from source file
cPathAndName = shrFindFilePath(cSourceFile, argv[0]);
cSourceCL = oclLoadProgSource(cPathAndName, "", &szKernelLength);
// Create the program CpProgram =
ClCreateProgramWithSource( cxGPUContext, 1, (const char **)&cSourceCL,
&szKernelLength, &ciErr1);
// Build the program
clBuildProgram(cpProgram, 0, NULL, NULL, NULL, NULL); // Create the kernel
ckKernel = clCreateKernel(cpProgram, "VectorAdd", &ciErr1);
引数 指定
起動
// Set the Argument values
clSetKernelArg(ckKernel, 0, sizeof(cl_mem), (void*)&cmDevDst); clSetKernelArg(ckKernel, 1, sizeof(cl_int), (void*)&iNumElements);
// Launch kernel
clEnqueueNDRangeKernel(cqCommandQue, ckKernel, 1, NULL,
&szGlobalWorkSize, &szLocalWorkSize, 0, NULL, NULL);
// Synchronous/blocking read of results, and check accumulated errors ClEnqueueReadBuffer(cqCommandQue, cmDevDst, CL_TRUE, 0,
sizeof(cl_float) * szGlobalWorkSize, dst,
OpenCL入門
OpenCL CUDA 比べ 非常 柔軟 仕様
言い方 え 複雑
ン 考え方 CUDA 同 ．
CUDA 学習時間 無駄 ．
う 必要
CUDA 一通 学習 後
OpenCL 取 掛 近道 思い ．
CUDA/OpenCL
く く基 的 入門 行い
機会 あ
高速化編 行い い
思い
Enjoy the CUDA and OpenCL !
長崎大学 う そ！
付録 :
CUDA-GDB ンス
CUDA-GDB ンス
付録 :
OpenCL ンス
注意事項
OpenCL 64-bit Linux
正規 方法 動作
(Ubuntu8.1以外 )
以 方法
CentOS5.3-x86_64 動
裏技 入 い ！！
