![コンスタントメモリへの
デ`タ僕
__constant__ int foo;
__constant__ int bar[100];
int func() {
int i = 100;
int j[100];
cudaMemcpyToSymbol(foo, &i, sizeof(i));
cudaMemcpyToSymbol(bar, j, sizeof(j));
}
塘双でもそうじゃなくても&&はいらない
18](https://image.slidesharecdn.com/cudamemo-150416140105-conversion-gate01/85/CUDA-18-320.jpg)
![グロ`バルメモリの旋喘
__kernel__ void func(int *arg, int n) {
for (int i = 0; i < n; i++) {
printf(＾%dn￣, arg[d]);
}
}
int func() {
int val[3] = {1, 3, 5};
int *d_val;
cudaMalloc((void**)&d_val, sizeof(val));
cudaMemcpy(d_val, val, sizeof(val), cudaMemcpyHostToDevice);
func<<<1, 1>>>(d_val, 3);
}
19](https://image.slidesharecdn.com/cudamemo-150416140105-conversion-gate01/85/CUDA-19-320.jpg)
遺雨禽粥メモ
- 2. MMaaxxwweellllのブロック GM107 GPU には、1 児の GPC、5 児の Maxwell ストリ`ミング?プロセッサSMM、2 児の 64 ビッ ト?メモリ?コントロ`ラ┷腕 128 ビットが喜dされています。これは、揖チップのフルg廾であり、 GeForce GTX 750 Ti での竃塞rと揖じ撹になっています。 v弼の膨叔がCCUUDDAA CCoorree CCUUDDAA CCoorreeの嶄で }方スレッドが揖rにg佩される 2
- 4. カ`ネルv方協x カ`ネルv方 GGPPUUにg佩させるv方 ____gglloobbaall____ vvooiidd kkeerrnneell(()) {{ }} ホストから柵び竃し辛嬬なv方 ____ddeevviiccee____ vvooiidd ffuunncc(()) {{ }} デバイスからのみ柵び竃し辛嬬なv方 4
- 5. カ`ネルv方柵び竃し kkeerrnneell<<<<<<ggiirrddDDiimm,, bblloocckkDDiimm>>>>>>((哈方));; ggrriiddDDiimm:: グリッドの寄きさ┘屮蹈奪の方? 11 oorr 22肝圷で峺協 bblloocckkDDIImm:: ブロックの寄きさ┘好譽奪匹諒? 11?33肝圷で峺協 ggrriidd bblloocckk tthhrreeaadd ggrriiddDDiimm xx bblloocckkDDiimmの方だけ スレッドが伏撹される 方嵐゛方為嵐スレッド 5
- 6. スレッドIIDD麿 スレッドIIDD tthhrreeaaddIIddxxで函誼tthhrreeaaddIIddxx..xx,, tthhrreeaaddIIddxx..yy,, tthhrreeaaddIIddxx..zz ブロックIIDD bblloocckkIIddxxで函誼bblloocckkIIddxx..xx,, bblloocckkIIddxx..yy ブロックの寄きさ11ブロックあたりのスレッド方 bblloocckkDDiimmで函誼bblloocckkDDiimm..xx,, bblloocckkDDiimm..yy,, bblloocckkDDiimm..zz グリッドの寄きさ11グリッドあたりのブロック方 ggrriiddDDiimmで函誼ggrriiddDDiimm..xx,, ggrriiddDDiimm..yy 6
- 7. HHeelllloo WWoorrlldd #include <stdio.h> __kernel__ void hello() { if (threadIdx.x == 0 & blockIdx.x == 0) { printf(＾Hello World!￣); } } void main() { hello<<<2, 10>>>(); cudaThreadSynchronize(); // カ`ネルv方のK阻を棋つ } 7
- 9. WWAARRPP ベクトル處麻グル`プ SSPP坪では3322スレッドg了でSSIIMMTTg佩 SSIIMMTT ssiinnggllee iinnssttrruuccttiioonn mmuullttiippllee tthhrreeaadd aadddd aadddd aadddd aadddd aadddd aadddd aadddd aadddd mmuull mmuull mmuull mmuull mmuull mmuull mmuull mmuull sseettpp sseettpp sseettpp sseettpp sseettpp sseettpp sseettpp sseettpp bbrraa LL11 bbrraa LL11 bbrraa LL11 bbrraa LL11 bbrraa LL11 bbrraa LL11 bbrraa LL11 bbrraa LL11 ssuubb ssuubb ssuubb bbrraa LL22 bbrraa LL22 bbrraa LL22 aadddd aadddd aadddd aadddd aadddd mmoovv mmoovv mmoovv mmoovv mmoovv mmoovv mmoovv mmoovv tthhrreeaadd ttiimmee LL11 LL22 9
- 11. aattoommiiccI尖 aattoommiiccAAdddd,, aattoommiiccSSuubb,, aattoommiiccCCAASS 吉 和返に聞うとWくなるので廣吭 方嵐スレッドで揖豚I尖がg佩さ れてしまう バリア揖豚とMみ栽わせて旋喘さ れることが謹い 11
- 12. カ`ネルv方 柵び竃し堀業 方マイクロ昼 22..55 GGHHzz CCPPUUクロックQ麻で? 方認クロックのオ`バ`ヘッド 12
- 13. GGPPUUのメモリレイアウト CPU Memory GPU Global Memory L2 Cache Shared Memory texture memory Constant memory L1/texture cache Shared Memory L1/texture cache local memory register register GPU CPU SM 13
- 14. 慌嗤メモリ LL11キャッシュ 11ブロックあたりの否楚 4488 KKBB 慌嗤メモリ1166 KKBB キャッシュ 1166 KKBB 慌嗤メモリ4488 KKBB キャッシュ アクセス堀業┘譽ぅ謄鵐沓 グロ`バルメモリの110000蔚殻業互堀 hhttttpp::////ddeevvbbllooggss..nnvviiddiiaa..ccoomm//ppaarraalllleellffoorraallll// uussiinngg--sshhaarreedd--mmeemmoorryy--ccuuddaa--cccc// 14
- 15. メモリアクセス堀業 ┘譽ぅ謄鵐沓 グロ`バルメモリ 11,,000000 クロック念瘁 LL22キャッシュ 330000 クロック念瘁 LL11キャッシュ,, テクスチャメモリコンスタ ントキャッシュ 方噴クロック 15
- 17. コンスタントメモリと 慌嗤メモリ ____ccoonnssttaanntt____ iinntt ffoooo;; コンスタントメモリ ____sshhaarreedd____ iinntt bbaarr;; 慌嗤メモリ デバイス坪のv方のみで協x辛嬬 17
- 18. コンスタントメモリへの デ`タ僕 __constant__ int foo; __constant__ int bar[100]; int func() { int i = 100; int j[100]; cudaMemcpyToSymbol(foo, &i, sizeof(i)); cudaMemcpyToSymbol(bar, j, sizeof(j)); } 塘双でもそうじゃなくても&&はいらない 18
- 19. グロ`バルメモリの旋喘 __kernel__ void func(int *arg, int n) { for (int i = 0; i < n; i++) { printf(＾%dn￣, arg[d]); } } int func() { int val[3] = {1, 3, 5}; int *d_val; cudaMalloc((void**)&d_val, sizeof(val)); cudaMemcpy(d_val, val, sizeof(val), cudaMemcpyHostToDevice); func<<<1, 1>>>(d_val, 3); } 19
- 20. 書指乾れなかったこと テクスチャメモリ コアとメモリgの鉦xに餓があるメモリ イメ`ジ議にはNNUUMMAA議なメモリ ストリ`ムと掲揖豚g佩 GGPPUUへのメモリコピ`とGGPPUUでの麻を揖rに佩える いわゆるパイプライン晒 マルチGGPPUU より互堀に麻できる 謹方のメモリのN ppiinnnneedd mmeemmoorryy,, mmaappppeedd mmeemmoorryy,, uunniiffiieedd mmeemmoorryy 20