Visual Studio 2013 の CUDA で、カーネルからカーネルを呼び出す (Dynamic Parallelism) には、次のページに書かれている設定を行う必要がある。
Compiling CUDA Projects with Dynamic Parallelism (VS 2012/13) | Viral F#:
http://viralfsharp.com/2014/08/17/compiling-cuda-projects-with-dynamic-parallelism-vs-201213/
Geforce GTX 970 で CUDA カーネル呼び出しにかかるオーバーヘッドを計測した。
(1) Dynamic Parallelism なし： ホストからカーネルを呼び出し
(2) Dynamic Parallelism あり： ホストからカーネルを呼び出し
(3) Dynamic Parallelism あり： カーネルからカーネルを呼び出し
結果は次のとおり。
(1) 3.80 µsec (マイクロ秒)
(2) 7.56 µsec
(3) 7.19 µsec
Dynamic Parallelism を使用にすると、コンパイルオプションの [Generate Relocatable Device Code] の影響なのか分からないが、ホストからカーネルを呼び出す場合もオーバーヘッドが増加している。
CPU から GPU の呼び出しを GPU から GPU の呼び出しに単純に変更しても、速度は速くならないようだ。
[参考]
CUDA の Kernel 呼び出しのオーバーヘッドと引数の数について ( 周辺機器 ) - 正統納豆天国ブログ - Yahoo!ブログ:
http://blogs.yahoo.co.jp/natto_heaven/33360615.html
CUDA の Kernel 呼び出しオーバーヘッド: Dynamic Parallelism 編 ( 周辺機器 ) - 正統納豆天国ブログ - Yahoo!ブログ:
http://blogs.yahoo.co.jp/natto_heaven/33395337.html
2015-03-13
2015-03-09
CUDA カーネルでの malloc とメモリ使用量
CUDA カーネルで malloc に大きいサイズを指定するとヒープ不足で失敗する。
事前に下記の API を呼び出して余裕を持たせたヒープサイズを指定する必要がある。
Geforce GTX 970 メモリ 4GB で、cudaDeviceSetLimit(cudaLimitMallocHeapSize) に 2GB を指定する。そして、カーネルから malloc(1) を呼び出して 1バイトのメモリを確保する。その後、cudaMemGetInfo API が返す free は 2GB へ減少する。malloc されるたびにメモリが確保されるのではなく、一度に最大ヒープサイズのメモリが確保されるようだ。
事前に下記の API を呼び出して余裕を持たせたヒープサイズを指定する必要がある。
cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)Geforce GTX 970 メモリ 4GB で、cudaDeviceSetLimit(cudaLimitMallocHeapSize) に 2GB を指定する。そして、カーネルから malloc(1) を呼び出して 1バイトのメモリを確保する。その後、cudaMemGetInfo API が返す free は 2GB へ減少する。malloc されるたびにメモリが確保されるのではなく、一度に最大ヒープサイズのメモリが確保されるようだ。
2015-03-07
Thrust: CUDA の C++ template ライブラリ
CUDA SDK に Thrust という CUDA の C++ template ライブラリが
標準で入っているのにさっき気付いた。
知らずに host_vector とかを自分で実装してしまった…。
Thrust :: CUDA Toolkit Documentation:
http://docs.nvidia.com/cuda/thrust/
C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v6.5/include/thrust
にあるライブラリのバージョンは 1.7.2 のようだ。
GitHub の Thrust project page の最新版は 1.8.0 となっている。
標準で入っているのにさっき気付いた。
知らずに host_vector とかを自分で実装してしまった…。
Thrust :: CUDA Toolkit Documentation:
http://docs.nvidia.com/cuda/thrust/
C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v6.5/include/thrust
にあるライブラリのバージョンは 1.7.2 のようだ。
GitHub の Thrust project page の最新版は 1.8.0 となっている。
2015-03-06
CUDA で分岐中の __syncthreads の動作
CUDA で分岐中に __syncthreads() を実行させた場合にデッドロックが起きると
インターネット上で見かけた文書に書かれていた。
そこで実際に GTX 970 で実験してみたが、デッドロックは起きなかった。
my_kernel<<<1, 2>>>(); でカーネル関数を呼び出す。
実行結果は次のとおり。
デッドロックはしなかったが、予期しない結果になった。
また、上記の関数で分岐の片方の __syncthreads() をコメントアウトしても
デッドロックはしなかった。
分岐中での __syncthreads() の動作は未定義ということだが、
これをテストした環境では何もしないという動作のようだ。
なお、上記のコードから if 文を取り除けば、
次のように正しい結果になる。
インターネット上で見かけた文書に書かれていた。
そこで実際に GTX 970 で実験してみたが、デッドロックは起きなかった。
__global__
void my_kernel()
{
__shared__ int shared[2];
shared[0] = -1;
shared[1] = -1;
int val;
if (threadIdx.x == 0)
{
shared[1 - threadIdx.x] = threadIdx.x;
__syncthreads();
val = shared[threadIdx.x];
}
else
{
shared[1 - threadIdx.x] = threadIdx.x;
__syncthreads();
val = shared[threadIdx.x];
}
printf("threadIdx.x=%d, val=%d.\n", threadIdx.x, val);
}
my_kernel<<<1, 2>>>(); でカーネル関数を呼び出す。
実行結果は次のとおり。
threadIdx.x=0, val=1.
threadIdx.x=1, val=-1.
デッドロックはしなかったが、予期しない結果になった。
また、上記の関数で分岐の片方の __syncthreads() をコメントアウトしても
デッドロックはしなかった。
分岐中での __syncthreads() の動作は未定義ということだが、
これをテストした環境では何もしないという動作のようだ。
なお、上記のコードから if 文を取り除けば、
次のように正しい結果になる。
threadIdx.x=0, val=1.
threadIdx.x=1, val=0.
Visual Studio で CUDA のソースコード編集時の赤い波線を消す
[注意]
この方法は副作用があるようなので、おすすめはしない。
Visual Studio で CUDA のソースコードを編集していると、
構文間違いではないのに赤い波線のエラーが表示される。
例えば、
__syncthreads();
や
my_cuda_kernel_func<<<1, 1>>>();
などで赤い波線が表示される。
これを解消するには、ソースコードの先頭に
を挿入する。
[参考]
visual studio 2010 - CUDA __syncthreads() compiles fine but is underlined with red - Stack Overflow:
http://stackoverflow.com/questions/13893919/cuda-syncthreads-compiles-fine-but-is-underlined-with-red
[追記]
これをすると、IntelliSense が効かなくなるようだ。
この方法は副作用があるようなので、おすすめはしない。
Visual Studio で CUDA のソースコードを編集していると、
構文間違いではないのに赤い波線のエラーが表示される。
例えば、
__syncthreads();
や
my_cuda_kernel_func<<<1, 1>>>();
などで赤い波線が表示される。
これを解消するには、ソースコードの先頭に
#ifndef __CUDACC__
#define __CUDACC__
#endif
を挿入する。
[参考]
visual studio 2010 - CUDA __syncthreads() compiles fine but is underlined with red - Stack Overflow:
http://stackoverflow.com/questions/13893919/cuda-syncthreads-compiles-fine-but-is-underlined-with-red
[追記]
これをすると、IntelliSense が効かなくなるようだ。
2015-03-05
CUDA の block と warp、threadIdx と thread ID について
1つの block が 31*31 のスレッドで構成される場合、
warpSize = 32
31 * 31 / warpSize = 30
31 * 31 % warpSize = 1
となり、block は 31個の warp に分割される。
また、thread ID は、
(thread ID) = threadIdx.x + threadIdx.y * blockDim.x
[参考]
cuda - How is the 2D thread blocks padded for warp scheduling? - Stack Overflow:
http://stackoverflow.com/questions/15044671/how-is-the-2d-thread-blocks-padded-for-warp-scheduling
warpSize = 32
31 * 31 / warpSize = 30
31 * 31 % warpSize = 1
となり、block は 31個の warp に分割される。
また、thread ID は、
(thread ID) = threadIdx.x + threadIdx.y * blockDim.x
[参考]
cuda - How is the 2D thread blocks padded for warp scheduling? - Stack Overflow:
http://stackoverflow.com/questions/15044671/how-is-the-2d-thread-blocks-padded-for-warp-scheduling
2015-03-04
CUDA で高速に配列の合計値を計算する方法
GPU を使って配列の合計値を計算する（parallel reductions）には、共有メモリとスレッド間の同期をとるためのバリアを使う方法が一般的だ。
CUDA には、warp という32個のスレッドのまとまりがある。warp 内のスレッドは常に同期しているので、バリアが不要になり、その分だけ高速化できる。
Kepler 以降（__CUDA_ARCH__ >= 300）の場合は __shfl_down という命令を使ってさらに高速化できる。この命令は、あるスレッドが同じ warp 内の別のスレッドのレジスタを直接参照できるので、共有メモリを使わずに warp 内のレジスタの合計を計算することができる。
詳しくは下記のページを参照されたい。
Faster Parallel Reductions on Kepler | Parallel Forall:
http://devblogs.nvidia.com/parallelforall/faster-parallel-reductions-kepler/
また、CUDA SDK に含まれるサンプルソースコード reduction_kernel.cu も参考になる。CUDA SDK 6.5 の場合は下記のパスにある。
C:\ProgramData\NVIDIA Corporation\CUDA Samples\v6.5\6_Advanced\reduction
CUDA には、warp という32個のスレッドのまとまりがある。warp 内のスレッドは常に同期しているので、バリアが不要になり、その分だけ高速化できる。
Kepler 以降（__CUDA_ARCH__ >= 300）の場合は __shfl_down という命令を使ってさらに高速化できる。この命令は、あるスレッドが同じ warp 内の別のスレッドのレジスタを直接参照できるので、共有メモリを使わずに warp 内のレジスタの合計を計算することができる。
詳しくは下記のページを参照されたい。
Faster Parallel Reductions on Kepler | Parallel Forall:
http://devblogs.nvidia.com/parallelforall/faster-parallel-reductions-kepler/
また、CUDA SDK に含まれるサンプルソースコード reduction_kernel.cu も参考になる。CUDA SDK 6.5 の場合は下記のパスにある。
C:\ProgramData\NVIDIA Corporation\CUDA Samples\v6.5\6_Advanced\reduction
2015-03-03
[Geforce GTX 970] cudaDeviceProp
CUDA Device
Name GeForce GTX 970
Driver WDDM
DeviceIndex 0
GPU Family GM204-A
RmGpuId 256
Compute Major 5
Compute Minor 2
MAX_THREADS_PER_BLOCK 1024
MAX_BLOCK_DIM_X 1024
MAX_BLOCK_DIM_Y 1024
MAX_BLOCK_DIM_Z 64
MAX_GRID_DIM_X 2147483647
MAX_GRID_DIM_Y 65535
MAX_GRID_DIM_Z 65535
MAX_SHARED_MEMORY_PER_BLOCK 49152
TOTAL_CONSTANT_MEMORY 65536
WARP_SIZE 32
MAX_PITCH 2147483647
MAX_REGISTERS_PER_BLOCK 65536
CLOCK_RATE 1253000
TEXTURE_ALIGNMENT 512
GPU_OVERLAP 1
MULTIPROCESSOR_COUNT 13
KERNEL_EXEC_TIMEOUT 1
INTEGRATED 0
CAN_MAP_HOST_MEMORY 1
COMPUTE_MODE 0
MAXIMUM_TEXTURE1D_WIDTH 65536
MAXIMUM_TEXTURE2D_WIDTH 65536
MAXIMUM_TEXTURE2D_HEIGHT 65536
MAXIMUM_TEXTURE3D_WIDTH 4096
MAXIMUM_TEXTURE3D_HEIGHT 4096
MAXIMUM_TEXTURE3D_DEPTH 4096
MAXIMUM_TEXTURE2D_LAYERED_WIDTH 16384
MAXIMUM_TEXTURE2D_LAYERED_HEIGHT 16384
MAXIMUM_TEXTURE2D_LAYERED_LAYERS 2048
SURFACE_ALIGNMENT 512
CONCURRENT_KERNELS 1
ECC_ENABLED 0
PCI_BUS_ID 1
PCI_DEVICE_ID 0
TCC_DRIVER 0
MEMORY_CLOCK_RATE 3505000
GLOBAL_MEMORY_BUS_WIDTH 256
L2_CACHE_SIZE 1835008
MAX_THREADS_PER_MULTIPROCESSOR 2048
ASYNC_ENGINE_COUNT 2
UNIFIED_ADDRESSING 1
MAXIMUM_TEXTURE1D_LAYERED_WIDTH 16384
MAXIMUM_TEXTURE1D_LAYERED_LAYERS 2048
CAN_TEX2D_GATHER 1
MAXIMUM_TEXTURE2D_GATHER_WIDTH 16384
MAXIMUM_TEXTURE2D_GATHER_HEIGHT 16384
MAXIMUM_TEXTURE3D_WIDTH_ALTERNATE 2048
MAXIMUM_TEXTURE3D_HEIGHT_ALTERNATE 2048
MAXIMUM_TEXTURE3D_DEPTH_ALTERNATE 16384
PCI_DOMAIN_ID 0
TEXTURE_PITCH_ALIGNMENT 32
MAXIMUM_TEXTURECUBEMAP_WIDTH 16384
MAXIMUM_TEXTURECUBEMAP_LAYERED_WIDTH 16384
MAXIMUM_TEXTURECUBEMAP_LAYERED_LAYERS 2046
MAXIMUM_SURFACE1D_WIDTH 65536
MAXIMUM_SURFACE2D_WIDTH 65536
MAXIMUM_SURFACE2D_HEIGHT 32768
MAXIMUM_SURFACE3D_WIDTH 65536
MAXIMUM_SURFACE3D_HEIGHT 32768
MAXIMUM_SURFACE3D_DEPTH 2048
MAXIMUM_SURFACE1D_LAYERED_WIDTH 65536
MAXIMUM_SURFACE1D_LAYERED_LAYERS 2048
MAXIMUM_SURFACE2D_LAYERED_WIDTH 65536
MAXIMUM_SURFACE2D_LAYERED_HEIGHT 32768
MAXIMUM_SURFACE2D_LAYERED_LAYERS 2048
MAXIMUM_SURFACECUBEMAP_WIDTH 32768
MAXIMUM_SURFACECUBEMAP_LAYERED_WIDTH 32768
MAXIMUM_SURFACECUBEMAP_LAYERED_LAYERS 2046
MAXIMUM_TEXTURE1D_LINEAR_WIDTH 134217728
MAXIMUM_TEXTURE2D_LINEAR_WIDTH 65000
MAXIMUM_TEXTURE2D_LINEAR_HEIGHT 65000
MAXIMUM_TEXTURE2D_LINEAR_PITCH 1048544
MAXIMUM_TEXTURE2D_MIPMAPPED_WIDTH 16384
MAXIMUM_TEXTURE2D_MIPMAPPED_HEIGHT 16384
MAXIMUM_TEXTURE1D_MIPMAPPED_WIDTH 16384
STREAM_PRIORITIES_SUPPORTED 0
GLOBAL_L1_CACHE_SUPPORTED 0
LOCAL_L1_CACHE_SUPPORTED 0
MAX_SHARED_MEMORY_PER_MULTIPROCESSOR 98304
MAX_REGISTERS_PER_MULTIPROCESSOR 65536
MANAGED_MEMORY 1
MULTI_GPU_BOARD 0
MULTI_GPU_BOARD_GROUP_ID 0
DISPLAY_NAME GeForce GTX 970
COMPUTE_CAPABILITY_MAJOR 5
COMPUTE_CAPABILITY_MINOR 2
TOTAL_MEMORY 4294967296
RAM_TYPE 8
RAM_LOCATION 1
GPU_PCI_DEVICE_ID 331485406
GPU_PCI_SUB_SYSTEM_ID 2231898179
GPU_PCI_REVISION_ID 161
GPU_PCI_EXT_DEVICE_ID 5058
GPU_PCI_EXT_GEN 2
GPU_PCI_EXT_GPU_GEN 2
GPU_PCI_EXT_GPU_LINK_RATE 8000
GPU_PCI_EXT_GPU_LINK_WIDTH 16
GPU_PCI_EXT_DOWNSTREAM_LINK_RATE 8000
GPU_PCI_EXT_DOWNSTREAM_LINK_WIDTH 16
[GeForce GTX 970] deviceQuery
deviceQuery.exe Starting...
CUDA Device Query (Runtime API) version (CUDART static linking)
Detected 1 CUDA Capable device(s)
Device 0: "GeForce GTX 970"
CUDA Driver Version / Runtime Version 6.5 / 6.5
CUDA Capability Major/Minor version number: 5.2
Total amount of global memory: 4096 MBytes (4294967296 bytes)
(13) Multiprocessors, (128) CUDA Cores/MP: 1664 CUDA Cores
GPU Clock rate: 1253 MHz (1.25 GHz)
Memory Clock rate: 3505 Mhz
Memory Bus Width: 256-bit
L2 Cache Size: 1835008 bytes
Maximum Texture Dimension Size (x,y,z) 1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
Maximum Layered 1D Texture Size, (num) layers 1D=(16384), 2048 layers
Maximum Layered 2D Texture Size, (num) layers 2D=(16384, 16384), 2048 layers
Total amount of constant memory: 65536 bytes
Total amount of shared memory per block: 49152 bytes
Total number of registers available per block: 65536
Warp size: 32
Maximum number of threads per multiprocessor: 2048
Maximum number of threads per block: 1024
Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
Max dimension size of a grid size (x,y,z): (2147483647, 65535, 65535)
Maximum memory pitch: 2147483647 bytes
Texture alignment: 512 bytes
Concurrent copy and kernel execution: Yes with 2 copy engine(s)
Run time limit on kernels: Yes
Integrated GPU sharing Host Memory: No
Support host page-locked memory mapping: Yes
Alignment requirement for Surfaces: Yes
Device has ECC support: Disabled
CUDA Device Driver Mode (TCC or WDDM): WDDM (Windows Display Driver Model)
Device supports Unified Addressing (UVA): Yes
Device PCI Bus ID / PCI location ID: 1 / 0
Compute Mode:
< Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >
deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 6.5, CUDA Runtime Version = 6.5, NumDevs = 1, Device0 = GeForce GTX 970
Result = PASS2015-03-01
OpenCL のソースコード
CUDA の SDK に含まれている OpenCL のバージョンは 1.1 のようだが、
cl.hpp ファイルが見つからない。
なお、cl.hpp のソースコードは次のページにある。
Khronos OpenCL Registry:
https://www.khronos.org/registry/cl/
cl.hpp ファイルが見つからない。
なお、cl.hpp のソースコードは次のページにある。
Khronos OpenCL Registry:
https://www.khronos.org/registry/cl/
Subscribe to:
Posts (Atom)