CUDA

【CUDA】初学者の勉強の仕方

CUDAを本格的に勉強し始めて(1ヶ月くらい?),気づいたこととして, 日本語資料の少なさがあります. 特に最新の情報はほとんど手に入りません.自分の思う最高効率の勉強は,CUDA by Example www.amazon.co.jpを一通り見て,概要を掴んで, 次は,CUDA Too…

【CUDA】__shfl_xorによるbutterfly exchange

warp shuffleの関数の中で一番直観的でないのが, このint __shfl_xor(int value, int laneMask)です. https://people.maths.ox.ac.uk/gilesm/cuda/lecs/lec4.pdfより引用がわかりやすくて,指定されたlaneMaskが,どのlaneIDのビットをflipさせるかを決め…

【CUDA】__shflによるwrap内の要素ずらし

__shfl_upや__shfl_downでは,端の値を循環させることが出来ない. そこで,__shflを使えば,端の値を循環させることができる. #include <stdio.h> __global__ void shfl_wrap(int *out_dev, int *in_dev, int const offset){ int value = in_dev[threadIdx.x]; valu</stdio.h>…

【CUDA】Warp shuffleの__shfl()_down, __shfl()_upサンプル

warp shuffleの__shfl_up()や__shfl_down()は、指定オフセット分だけずれたスレッドから値を読み出します。 #include <stdio.h> __global__ void shfl_test(int arg){ int laneId = threadIdx.x & 0x1f; int value; if(laneId == 3){ value = arg; }else{ value = 0; </stdio.h>…

【CUDA】Warp shuffleの__shfl()サンプル

warp shuffleは,kepler世代のcc3.x以上から使える, shared memoryを用いずに,warp内のthread間で値を交換することができる機能です.GPGPUでは,shared memoryをいじるのが当然なのですが,それをせずにさらに高速化することができるということで,使える…

【CUDA】cudaMallocHostとページング

CUDA by Exampleを読んでます. Amazon.co.jp: CUDA by Example 汎用GPUプログラミング入門: Jason Sanders, Edward Kandrot, 株式会社クイープ: 本その中で,ホスト側のメモリ確保のためにc標準のmallocと, CUDAに用意されているcudaMallocHostの違いが説…

【CUDA】Shared memoryによる行列積の高速化

Shared memoryによる行列積の高速化を試みてみました. CUDA公式のprogramming guideにも載っている基本課題ですが, 正直,全然(英語が)理解できなかったので, 挫折しかかっていたのですが,http://www.ccn.yamanashi.ac.jp/~stomo/%E3%82%B5%E3%82%A4%E…

【CUDA】Shared memoryの動的な確保

Shared memoryは,オンチップでキャッシュされないグローバルメモリにアクセスするより100倍速い. Shared memoryはブロックごとに配置されているおり,同じブロックのスレッドは,Shared memoryにアクセスできる. そのため,競合を防ぐために__syncthreads…

【CUDA】Dynamic parallelismによるQuicksort

http://blogs.nvidia.com/blog/2012/09/12/how-tesla-k20-speeds-up-quicksort-a-familiar-comp-sci-code/CC3.5以上で使えるDynamic parallelismは,デバイス内で再帰的にカーネルを呼び出すことできるので, Quicksortなどが簡易に実装できる. 【CUDA】Dyn…

【CUDA】Dynamic Parallelism

Dynamic parallelismとは、カーネルの中からカーネルを呼び出すことができるsm_35以降の機能です。ちょっと具体的にどういった場面に使えるかよくわからないのですが、 Adaptive Parallel Computation with CUDA Dynamic Parallelism | Parallel Forall によ…

【CUDA】SM,Warp,Occupancyなどの概念

CUDAを勉強しようとして,まずつまずくのが, ThreadやBlock, Gridなどのソフトウェア上の概念と, Streaming Multiprocessor(SM),CUDA Core,Warp, Occupancyなどのハードウェアの概念がごっちゃになる点だ.今回はじめて理解できた(気がする)ので,忘れ…

【CUDA】グリッド,ブロック,スレッド

CUDA勉強中 CUDA by Exampleからのメモ. #define N 1024 __global__ void add(int *a, int *b, int *c){ //int tid = threadIdx.x; if(tid < N){ c[tid] = a[tid] + b[tid]; } } といったベクトルの加算をデバイスにさせたいときに、どうブロックやスレッド…