untitled
|
|
|
- ゆいと にかどり
- 5 years ago
- Views:
Transcription
1 CUDA Vol II: CUDA : NVIDIA Q2 2008
2 CUDA...1 CUDA...3 CUDA
3 CUDA CUDA ( ) CUDA 2 NVIDIA Corporation
4 CUDA CUDA GPU GPU CPU NVIDIA Corporation
5 V V d call put 1 2 = S CND( d ) X e = X e log( d = log( = S X rt S X 1 CND( d 2 v ) + ( r + ) T 2 v T 2 v ) + ( r ) T 2 v T CND( d) = 1 CND( d) 2 rt CND( d ) S CND( d ) 2 ) 1 S X CND r v NVIDIA Corporation u 1 x 2 N ( x) = e du 2 (Hull ) 5 6 host device float CND(float d) { float K = 1.0f / (1.0f f * fabsf(d)); float CND = RSQRT2PI * expf(- 0.5f * d * d) * (K * (A1 + K * (A2 + K * (A3 + K * (A4 + K * A5))))); if(d > 0) CND = 1.0f - CND; return CND; float 2 2 NVIDIA Corporation
6 1. : hoptprice(n) hoptstrike (N).. 2. : doptprice(n) doptstrike (N) GPU 6. GPU 7. GPU NVIDIA Corporation /* */ float *hoptprice, *hoptstrike, *hoptyear; hoptprice = (float *) malloc(sizeof(float*n); hoptstrike = (float *) malloc(sizeof(float*n); hoptyear = (float *) malloc(sizeof(float*n); /* cudamalloc GPU */ float *doptprice, *doptstrike, *doptyear; cudamalloc( (void **) &doptprice, sizeof(float)*n); cudamalloc( (void **) &doptstrike, sizeof(float)*n); cudamalloc( (void **) &doptyear, sizeof(float)*n); /* hoptprice hoptstrike hoptyear */ NVIDIA Corporation
7 4 5 /* cudamemcpy (target, source, size, direction)*/ cudamemcpy (doptprice, hoptprice, sizeof(float)*n, cudamemcpyhosttodevice); cudamemcpy (doptstrike, hoptstrike, sizeof(float)*n, cudamemcpyhosttodevice); cudamemcpy (doptyears, hoptyears, sizeof(float)*n, cudamemcpyhosttodevice); /* GPU <<<, >>>*/ / BlackScholesGPU<<<128, 256>>>( dcallresult, dputresult, doptionstrike, doptionprice, doptionyears, RISKFREE, VOLATILITY, OPT_N); NVIDIA Corporation /* cudamemcpy(target, source, size, direction)*/ cudamemcpy (hcallresult, dcallresult, sizeof(float)*n, cudamemcpydevicetohost); cudamemcpy (hputresult, dputresult, sizeof(float)*n, cudamemcpydevicetohost); /* CPU */ BlackScholesCPU( ); /* */ /* */ / free( hoptprice); cudafree(doptprice); NVIDIA Corporation
8 BlackScholesGPU OptN? = 65,536 1 = 512 OptN 33M : { global void BlackScholes (float *., int OptN) const int const int tid = blockdim.x * blockidx.x + threadidx.x; THREAD_N = blockdim.x * griddim.x; for(int opt = tid; opt < OptN; opt += THREAD_N) BlackScholesBody( d_callresult[opt], d_putresult[opt], d_optionprice[opt], d_optionstrike[opt],d_optionyears[opt], Riskfree,Volatility ); NVIDIA Corporation BlackScholes.cu nvcc O3 o BlackScholes BlacScholes.cu \\ -I../../common/inc/ -L../../lib/ -lcutil -lgl lglut : BlackScholes (1,000,000 ) GPU GPU GPU CPU : GPU : : CPU : L1 : E-08 : E-05 NVIDIA Corporation
9 PCIe PCIe GB/s CPU 3 GB/s NVIDIA LinkBoost 4 GB CUDA NVIDIA Corporation pinned cudamalloc cudamallochost cudafree cudafreehost /* malloc */ / cudamallochost ((void **) &h_callresultgpu, OPT_SZ); /* free cudafreehost(h_callresultgpu); NVIDIA Corporation
10 BlackScholesPinned.cu nvcc O3 o BlackScholesPinned BlacScholesPinned.cu\\ -I../../common/inc/ -L../../lib/ -lcutil -lgl lglut ( ) ) : BlackScholesPinned 1,000,000 GPU GPU GPU : GPU : : CPU CPU : L1 : E-08 : E-05 NVIDIA Corporation CUDA
11 CUDA NVIDIA Corporation i. ii. iii. % MATLAB % An Introduction to Financial Option Valuation: Mathematics, Stochastics % and Computation D. Higham S = 2; E = 1; r = 0.05; 05 sigma = 0.25; T = 3; M = 1e6; Svals = S*exp((r-0.5*sigma^2)*T + sigma*sqrt(t)*randn(m,1)); Pvals = exp(-r*t)*max(svals-e,0); Pmean = mean(pvals) width = 1.96*std(Pvals)/sqrt(M); conf = [Pmean - width, Pmean + width] NVIDIA Corporation
12 i. M :M=200,000,000 i. MT ii. ii. N : N=128 iii. iv. : MT GPU CUDA SDK v1.1 MonteCarloMultiGPU NVIDIA Corporation RNG DCMT RandomGPU<<<32,128>>>( d_random, N_PER_RNG, seed); : 4096 Tesla C870 GPU !!!! NVIDIA Corporation
13 BoxMullerGPU<<<32,128>>>( d_random, N_PER_RNG, RNG seed); Tesla C870 GPU #define PI f device void BoxMuller(float& u1, float& u2){ float r = sqrtf(-2.0f * logf(u1)); float phi = 2 * PI * u2; u1 = r * cosf(phi); u2 = r * sinf(phi); Beasley-Springer-Moro NVIDIA Corporation void MonteCarloGPU(d_Random,.) { // ,384 MonteCarloKernelGPU<<<64, 256, 0>>>(d_Random); // cudamemcpy(h_sum, d_sum, ACCUM_SZ, cudamemcpydevicetohost) ; cudamemcpy(h_sum2, d_sum2, ACCUM_SZ, cudamemcpydevicetohost) ; // 2 double dblsum = 0, dblsum2 = 0; for(int i = 0; i < ACCUM_N; i++){ dblsum += h_sum[i]; dblsum2 += h_sum2[i]; NVIDIA Corporation
14 global void MonteCarloKernelGPU( ) { const int tid = blockdim.x * blockidx.x + threadidx.x; const int threadn = blockdim.x * griddim.x; //... for(int iaccum = tid; iaccum < accumn; iaccum += threadn) { float sum = 0, sum2 = 0; for(int ipath = iaccum; ipath < pathn; ipath += accumn) { float r = d_random[ipath]; //... sum += endoptionprice; sum2 += endoptionprice * endoptionprice; d_sum[iaccum] = sum; d_sum2[iaccum] = sum2; NVIDIA Corporation N a i S 0 =0 S i =S i-1 +a 1 S=S n Wilkinson 1963 N 2 NVIDIA Corporation
15 Montecarlo.cu : nvcc O3 o Montecarlo Montecarlo.cu \\ -I../../common/inc/ -L../../lib/ -lcutil -lgl lglut ( ) : Montecarlo 200,000,000 GPU GPU : ; : : e-05; : e-05; : ; : : e-05; : e-06; NVIDIA Corporation :256 16M : K NVIDIA Corporation
16 1 NVIDIA Corporation NVIDIA Corporation
17 1 NVIDIA Corporation NVIDIA Corporation
18 GPU CPU GPU CPU 1 : 64 NVIDIA Corporation GPU GPU
19 NVIDIA Corporation
20 1 1 CPU cudamemcpy : 1 1 NVIDIA Corporation ? NVIDIA bool multiblock = ((numpaths / numoptions >= 8192); multiblock false NVIDIA Corporation
21 1 1 NVIDIA Corporation CUDA SDK 1.1 MonteCarlo GPU MonteCarloMultiGPU NVIDIA Corporation
22 CUDA NVIDIA Corporation
23 FFT GPU GPU NVIDIA Corporation φ = r FFT ( k 2 x + k 2 y ) ˆ φ = rˆ 1. 2 FFT r(k) k 2 FFT k 2. (k) u(k) ˆ rˆ φ = 2 ( k + 2 x k y ) 3. 2 FFT u(k) u NVIDIA Corporation
24 MATLAB % N = 64; % L = 1; % f << 1 sig = 0.1; % k = (2*pi/L)*[0:(N/2-1) (-N/2):(-1)]; % (m,n) % (x,y) [KX KY] = meshgrid(k,k); % delsq = -(KX.^2 + KY.^2); % (0,0) % % 0 delsq(1,1) = 1; % h = L/N; x = (0:(N-1))*h ; y = (0:(N-1))*h; [X Y] = meshgrid(x,y); % RHS f(x,y) rsq = (X-0.5*L).^2 + (Y-0.5*L).^2; sigsq = sig^2; f = exp(-rsq/(2*sigsq)).* (rsq - 2*sigsq)/(sigsq^2); % fhat = fft2(f); u = real(ifft2(fhat./delsq)); % u = 0 % u = u - u(1,1); % L2 Linf uex = exp(-rsq/(2*sigsq)); errmax = norm(u(:)-uex(:),inf); errmax2 = norm(u(:)-uex(:),2)/(n*n); % L2 Linf fprintf('n=%d n',n); fprintf('solution f(' i at (%d,%d):%d) ',N/2,N/2); fprintf('computed=%10.6f reference = %10.6f n',u(n/2,n/2), uex(n/2,n/2)); fprintf('linf err=%10.6e L2 norm err = %10.6e n',errmax, errmax2); NVIDIA Corporation : r (NxN) u (NxN) kx (N) ky (N) 2. : r_d u_d kx_d ky_d d d ky d 3. kx ky 4. FFT FFT FFT GPU C2C NVIDIA Corporation
25 1 3 /* */ float *kx, *ky, *r; kx = (float *) malloc(sizeof(float*n); ky = (float *) malloc(sizeof(float*n); r = (float *) malloc(sizeof(float*n*n); /* cudamalloc GPU */ float *kx_d, *ky_d, *r_d; cudamalloc( (void **) &kx_d, sizeof(cufftcomplex)*n); cudamalloc( (void **) &ky_d, sizeof(cufftcomplex)*n); cudamalloc( (void **) &r_d, sizeof(cufftcomplex)*n*n); cufftcomplex *r_complex_d; cudamalloc( (void **) &r_complex_d, sizeof(cufftcomplex)*n*n); NVIDIA Corporation /* r kx ky */ /* cudamemcpy cpy (target, source, size, direction)*/ cudamemcpy (kx_d, kx, sizeof(float)*n, cudamemcpyhosttodevice); cudamemcpy (ky_d, ky, sizeof(float)*n, cudamemcpyhosttodevice); cudamemcpy (r_d, r, sizeof(float)*n*n, cudamemcpyhosttodevice); /* CUDA FFT (FFTW ) */ cuffthandle plan; cufftplan2d( &plan, N, N, CUFFT_C2C); NVIDIA Corporation
26 5 /* : block_size_x*block_size_y = G */ dim3 dimblock(block_size_x, block_size_y); dim3 dimgrid (N/dimBlock.x, N/dimBlock.y); /* N block_size_x block_size_y */ if (N % block_size_x!=0 ) dimgrid.x+=1; if (N % block_size_y!=0 ) dimgrid.y+=1 NVIDIA Corporation /* */ real2complex<<<dimgrid, dimblock>>> (r_d, r_complex_d, N); /* FFT */ cufftexecc2c (plan, r_complex_d, r_complex_d, CUFFT_FORWARD); FORWARD) /* */ solve_poisson<<<dimgrid, dimblock>>> (r_complex_d, kx_d, ky_d,n); /* FFT */ cufftexecc2c (plan, r_complex_d, r_complex_d, CUFFT_INVERSE); /* FFT ifft */ scale = 1.f / ( (float N * (float) N ); complex2real_scaled<<<dimgrid, dimblock>>> (r_d, r_complex_d, N, scale); NVIDIA Corporation
27 11 /* cudamemcpy(target, source, size, direction)*/ cudamemcpy (r, r_d, sizeof(float)*n*n, cudamemcpydevicetohost); /* */ cufftdestroy( plan); cudafree(r_complex_d); cudafree(kx_d); NVIDIA Corporation real2complex /* */ global void real2complex (float *a, cufftcomplex *c, int N) { /* NxN idx idy */ int idx = blockid.x*blockdim.x+threadidx.x; int idy = blockid.y*blockdim.y+threadidx.y; if ( idx < N && idy <N) { int index = idx + idy*n; c[index].x = a[index]; c[index].y = 0.f; NVIDIA Corporation
28 solve_poisson global void solve_poisson (cufftcomplex *c, float *kx, float *ky, int N) { /* N N idx idy */ int idx = blockid.x*blockdim.x+threadidx.x; int idy = blockid.y*blockdim.y+threadidx.y; if ( idx < N && idy <N) { int index = idx + idy*n; float scale = - ( kx[idx]*kx[idx] + ky[idy]*ky[idy] ); if ( idx ==0 && idy == 0 ) scale =1.f; scale = 1.f / scale; c[index].x *= scale; c[index].y *= scale; ˆ rˆ φ = 2 2 ( k x + k y ) NVIDIA Corporation complex2real_scaled /* */ global void complex2real_scaled (cufftcomplex *c, float *a, int N, float scale) { /* N N idx idy */ int idx = blockid.x*blockdim.x+threadidx.x; int idy = blockid.y*blockdim.y+threadidx.y; if ( idx < N && idy <N) { int index = idx + idy*n; a[index] = scale*c[index].x ; NVIDIA Corporation
29 poisson_1 poisson_1.cu nvcc O3 o poisson_1 poisson_1.cu \\ -I/usr/local/cuda/include L/usr/local/cuda/lib -lcufft lcudart./poisson_1 -N dimblock dimgrid 2 4 L e-08: : I/O ( ): (32,32) 32) = = MATLAB N=64 (32,32) : = = Linf = e-05 L2 = e-08 NVIDIA Corporation CUDA CUDA_PROFILE: 1 0 CUDA_PROFILE_LOG: filename cuda_profile.log CUDA_PROFILE_CSV: 1 0 NVIDIA Corporation
30 Poisson_1./poisson_1 N1024 method=[ memcopy ] gputime=[ ] method=[ memcopy ] gputime=[ ] method=[ memcopy ] gputime=[ ] method=[ real2complex ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_radix4 ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ transpose ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_radix4 ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_transpose ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ solve_poisson] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_radix4 ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_transpose] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_radix4 ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_transpose] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ complex2real_scaled ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ memcopy ] gputime=[ ] NVIDIA Corporation pinned CPU GPU #ifdef PINNED cudamallochost((void **) &r,sizeof(float)*n*n); // rhs 2 #else r = (float *) malloc(sizeof(float)*n*n); // rhs 2 #endif $./poisson_1 1,024 1,024 : (ms) : (ms) I/O : ( ) (ms) $./poisson_1_pinned pinned 1,024 1,024 : (ms) : (ms) I/O : ( ) (ms) NVIDIA Corporation
31 solve_poisson kx ky ( umul24) NVIDIA Corporation solve_poisson global void solve_poisson (cufftcomplex *c, float *kx, float *ky, int N) { unsigned int idx = umul24(blockidx.x,blockdim.x)+threadidx.x; unsigned int idy = umul24(blockidx.y,blockdim.y)+threadidx.y; // k shared float kx_s[block_width], ky_s[block_height] if (threadix.x < 1) kx_s[threadidx.x] = kx[idx]; if (threadix.y < 1) ky_s[threadidx.y] = ky[idy]; syncthreads(); if ( idx < N && idy <N) { unsigned int index = idx + umul24(idy,n); float scale = - ( kx_s[threadidx.x]*kx_s[threadidx.x] + ky_s[threadidy.y]*ky_s[threadidy.y] ); if ( idx ==0 && idy == 0 ) scale =1.f; scale = 1.f / scale; c[index].x *= scale; c[index].y*= scale; NVIDIA Corporation
32 Poisson_2./poisson_2 N1024 x16 y16 method=[ memcopy ] gputime=[ ] method=[ memcopy ] gputime=[ ] method=[ memcopy ] gputime=[ ] method=[ real2complex ] gputime=[ ] cputime=[ ] occupancy=[ ] (was 1654) method=[ c2c_radix4 ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_transpose ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_radix4 ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_transpose] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ solve_poisson] gputime=[ ] cputime=[ ] occupancy=[ ] (was 6389) method=[ c2c_radix4 ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_transpose] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_radix4 ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_transpose] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ complex2real_scaled ] gputime=[ ] cputime=[ ] occupancy=[ ]??????? method=[ memcopy ] gputime=[ ] NVIDIA Corporation complex2real_scaled ( ) global void complex2real_scaled (cufftcomplex *c, float *a, int N, floatscale) { /* NxN idx idy */ int idx = blockid.x*blockdim.x+threadidx.x; kdi did int idy = blockid.y*blockdim.y+threadidx.y; volatile float2 c2; if ( idx < N && idy <N) { int index = idx + idy*n; c2.x= c[index].x; c2.y= c[index].y; a[index] = scale*c2.x c2.x ; ptx NVIDIA Corporation
33 Poisson_3 method=[ memcopy ] gputime=[ ] method=[ memcopy ] gputime=[ ] method=[ memcopy ] gputime=[ ] method=[ real2complex] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_radix4] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_transpose] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_radix4 ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_transpose] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ solve_poisson ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_radix4 ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_transpose] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_radix4 ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ c2c_transpose] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ complex2real_scaled ] gputime=[ ] cputime=[ ] occupancy=[ ] method=[ memcopy ] gputime=[ ] NVIDIA Corporation pinned pinned (r2c c2r) 67ms (10.8ms) 63ms ms (7.1ms) 59.4ms +c2r 62.1ms 58.2ms (5.8ms) Tesla C870 pinned : 10.4ms NVIDIA Corporation
34 CUDA 7 NVIDIA Corporation
35 GPU? NVIDIA Corporation ? : GPU GPU M b B >M* b GPU NVIDIA Corporation
36 : NVIDIA Corporation ? GPU GFLOP/s: : G80 GPU MHz DDR 384 * 1800 / 8 = 86.4 GB/s NVIDIA Corporation
![1: global void reduce0(int *g_idata, int *g_odata) { extern shared int sdata[]; // 1 1 unsigned int tid = threadidx.x; unsigned int i = blockidx.x*blockdim.x + threadidx.](/docs-images/93/118017109/images/37-0.jpg "x; sdata[tid] = g_idata[i]; syncthreads(); // for(unsigned int s=1; s < blockdim.x; s *= 2) { if (tid % (2*s) == 0) { sdata[tid] += sdata[tid + s]; syncthreads(); // if (tid == 0) g_odata[blockidx.")
37 1: global void reduce0(int *g_idata, int *g_odata) { extern shared int sdata[]; // 1 1 unsigned int tid = threadidx.x; unsigned int i = blockidx.x*blockdim.x + threadidx.x; sdata[tid] = g_idata[i]; syncthreads(); // for(unsigned int s=1; s < blockdim.x; s *= 2) { if (tid % (2*s) == 0) { sdata[tid] += sdata[tid + s]; syncthreads(); // if (tid == 0) g_odata[blockidx.x] = sdata[0]; NVIDIA Corporation : 1 == 1 2 == 2 3 == 4 4 == 8 ID ID ID ID NVIDIA Corporation
38 1: global void reduce1(int *g_idata, int *g_odata) { extern shared int sdata[]; // 1 unsigned int tid = threadidx.x; unsigned int i = blockidx.x*blockdim.x + threadidx.x; sdata[tid] = g_idata[i]; syncthreads(); // for (unsigned int s=1; s < blockdim.x; s *= 2) { if (tid % (2*s) == 0) { sdata[tid] += sdata[tid + s]; : syncthreads(); // if (tid == 0) g_odata[blockidx.x] = sdata[0]; NVIDIA Corporation M 1: ms GB/s : 128 NVIDIA Corporation
39 2: for (unsigned int s=1; s < blockdim.x; s *= 2) { if (tid %(2*s) == 0) { sdata[tid] += sdata[tid + s]; syncthreads(); for (unsigned int s=1; s < blockdim.x; s *= 2) { int index = 2 * s * tid; if (index < blockdim.x) { sdata[index] += sdata[index + s]; syncthreads(); : NVIDIA Corporation M 1: 2: 2 22 int ms GB/s ms GB/s 2.33 NVIDIA Corporation
40 : 1 == 8 2 == 4 3 == 2 4 == 1 ID ID ID ID NVIDIA Corporation : for (unsigned int s=1; s < blockdim.x; s *= 2) { int index = 2 * s * tid; if (index < blockdim.x) { sdata[index] += sdata[index + s]; syncthreads(); ID for (unsigned int s=blockdim.x/2; s>0; s>>=1) { if (tid < s) { sdata[tid] += sdata[tid + s]; syncthreads(); NVIDIA Corporation
41 4M 1: 2: 3: 2 22 int ms GB/s ms GB/s ms GB/s 2.01 NVIDIA Corporation : for (unsigned int s=blockdim.x/2; s>0; s>>=1) { if (tid <s){ sdata[tid] += sdata[tid + s]; syncthreads();! NVIDIA Corporation
42 4: 1 // 1 unsigned int tid = threadidx.x; x; unsigned int i = blockidx.x*blockdim.x + threadidx.x; sdata[tid] = g_idata[i]; syncthreads(); 2 // // unsigned int tid = threadidx.x; unsigned int i = blockidx.x*(blockdim.x*2) + threadidx.x; sdata[tid] = g_idata[i] + g_idata[i+blockdim.x]; syncthreads(); NVIDIA Corporation M 1: 2: 3: 4: 2 22 int ms GB/s ms GB/s ms GB/s ms GB/s 1.78 NVIDIA Corporation
43 17 GB/s : NVIDIA Corporation s <= 32 1 SIMD s <= 32 syncthreads() if (tid < s) 6 6 NVIDIA Corporation
44 5: for (unsigned int s=blockdim.x/2; s>32; s>>=1) { if (tid < s) sdata[tid] += sdata[tid + s]; syncthreads(); if (tid < 32) { sdata[tid] += sdata[tid + 32]; sdata[tid] += sdata[tid + 16]; sdata[tid] += sdata[tid + 8]; sdata[tid] += sdata[tid + 4]; sdata[tid] += sdata[tid + 2]; sdata[tid] += sdata[tid + 1]; for if NVIDIA Corporation M 1: 2: 3: 4: 2 22 int ms GB/s ms GB/s ms GB/s ms GB/s : ms GB/s 1.8 NVIDIA Corporation
45 GPU :? CUDA C++ NVIDIA Corporation template <unsigned int blocksize> global void reduce5(int *g_idata, int *g_odata) NVIDIA Corporation
46 6: if (blocksize >= 512) { if (tid < 256) { sdata[tid] += sdata[tid + 256]; syncthreads(); if (blocksize >= 256) { if (tid < 128) { sdata[tid] += sdata[tid + 128]; syncthreads(); if (blocksize >= 128) { if (tid < 64) { sdata[tid] += sdata[tid + 64]; syncthreads(); if (tid < 32) { if (blocksize >= 64) sdata[tid] += sdata[tid + 32]; if (blocksize >= 32) sdata[tid] += sdata[tid + 16]; if (blocksize >= 16) sdata[tid] += sdata[tid + 8]; if (blocksize >= 8) sdata[tid] += sdata[tid + 4]; if (blocksize >= 4) sdata[tid] += sdata[tid + 2]; if (blocksize >= 2) sdata[tid] += sdata[tid + 1]; : NVIDIA Corporation ? 10 switch switch (threads) { case 512: reduce5<512><<< dimgrid, dimblock, smemsize >>>(d_idata, d_odata); break; case 256: reduce5<256><<< dimgrid, dimblock, smemsize >>>(d_idata, d_odata); break; case 128: reduce5<128><<< dimgrid, dimblock, smemsize >>>(d_idata, d_odata); break; case 64: reduce5< 64><<< dimgrid, dimblock, smemsize >>>(d_idata, d_odata); break; case 32: reduce5< 32><<< dimgrid, dimblock, smemsize >>>(d_idata, d_odata); break; case 16: reduce5< 16><<< dimgrid, dimblock, smemsize >>>(d_idata, d_odata); break; case 8: reduce5< 8><<< dimgrid, dimblock, smemsize >>>(d_idata, d_odata); break; case 4: reduce5< 4><<< dimgrid, dimblock, smemsize >>>(d_idata, d_odata); break; case 2: reduce5< 2><<< dimgrid, dimblock, smemsize >>>(d_idata, d_odata); break; case 1: reduce5< 1><<< dimgrid, dimblock, smemsize >>>(d_idata, d_odata); break; NVIDIA Corporation
47 4M 1: 2: 3: 4: 2 22 int ms GB/s ms GB/s ms GB/s ms GB/s : ms GB/s 1.8 6: ms GB/s 1.41 NVIDIA Corporation Log(N) S N/2 O(log N) N=2 D S [1..D] 2 D-S = N-1 O(N): P P O(N/P + log N) O(N) N=P O(log N) NVIDIA Corporation
48 x : O(N) O(log N) O(N log N): O(N/log N) O(log N) O(N/log N) O(log N) = O((N/log N) * log N) = O(N) NVIDIA Corporation O(log n) 1 1,024 2, G ,024 4,096 NVIDIA Corporation
49 7: 2 unsigned int tid = threadidx.x; unsigned int i = blockidx.x*(blockdim.x*2) + threadidx.x; sdata[tid] = g_idata[i] + g_idata[i+blockdim.x]; syncthreads(); while unsigned int tid = threadidx.x; unsigned int i = blockidx.x*(blocksize*2) + threadidx.x; unsigned int gridsize = blocksize*2*griddim.x; sdata[tid] = 0; do { sdata[tid] += g_idata[i] + g_idata[i+blocksize]; i += gridsize; while (i < n); syncthreads(); NVIDIA Corporation M 1: 2: 3: 4: 2 22 int ms GB/s ms GB/s ms GB/s ms GB/s : ms GB/s 1.8 6: 7: ms GB/s ms GB/s M 7: 72 GB/s! : 30! NVIDIA Corporation
![template <unsigned int blocksize> global void reduce6(int *g_idata, int *g_odata, unsigned int n) { extern shared int sdata[]; unsigned int tid = threadidx.x; unsigned int i = blockidx.](/docs-images/93/118017109/images/50-1.jpg "x*(blocksize*2) + tid; unsigned int gridsize = blocksize*2*griddim.")
50 template <unsigned int blocksize> global void reduce6(int *g_idata, int *g_odata, unsigned int n) { extern shared int sdata[]; unsigned int tid = threadidx.x; unsigned int i = blockidx.x*(blocksize*2) + tid; unsigned int gridsize = blocksize*2*griddim.x; sdata[tid] = 0; do { sdata[tid] += g_idata[i] + g_idata[i+blocksize]; i += gridsize; while (i < n); syncthreads(); if (blocksize >= 512) { if (tid < 256) { sdata[tid] += sdata[tid + 256]; syncthreads(); if (blocksize >= 256) { if (tid < 128) { sdata[tid] += sdata[tid + 128]; syncthreads(); if (blocksize >= 128) { if (tid < 64) { sdata[tid] += sdata[tid + 64]; syncthreads(); if (tid < 32) { if (blocksize >= 64) sdata[tid] += sdata[tid + 32]; if (blocksize >= 32) sdata[tid] += sdata[tid + 16]; if (blocksize >= 16) sdata[tid] += sdata[tid + 8]; if (blocksize >= 8) sdata[tid] += sdata[tid + 4]; if (blocksize >= 4) sdata[tid] += sdata[tid + 2]; if (blocksize >= 2) sdata[tid] += sdata[tid + 1]; if (tid == 0) g_odata[blockidx.x] = sdata[0]; NVIDIA Corporation NVIDIA Corporation
51
52 NVIDIA NVIDIA NVIDIA Corporation NVIDIA Corporation NVIDIA Corporation NVIDIA Corporation NVIDIA NVIDIA CUDA Tesla NVIDIA Corporation Copyright 2008 NVIDIA Corporation.All rights reserved. NVIDIA Corporation 2701 San Tomas Expressway Santa Clara, CA
GPU CUDA CUDA 2010/06/28 1
GPU CUDA CUDA 2010/06/28 1 GPU NVIDIA Mark Harris, Optimizing Parallel Reduction in CUDA http://developer.download.nvidia.com/ compute/cuda/1_1/website/data- Parallel_Algorithms.html#reduction CUDA SDK
07-二村幸孝・出口大輔.indd
GPU Graphics Processing Units HPC High Performance Computing GPU GPGPU General-Purpose computation on GPU CPU GPU GPU *1 Intel Quad-Core Xeon E5472 3.0 GHz 2 6 MB L2 cache 1600 MHz FSB 80 GFlops 1 nvidia
untitled
GPGPU NVIDACUDA Learn More about CUDA - NVIDIA http://www.nvidia.co.jp/object/cuda_education_jp.html NVIDIA CUDA programming Guide CUDA http://www.sintef.no/upload/ikt/9011/simoslo/evita/2008/seland.pdf
( CUDA CUDA CUDA CUDA ( NVIDIA CUDA I
GPGPU (II) GPGPU CUDA 1 GPGPU CUDA(CUDA Unified Device Architecture) CUDA NVIDIA GPU *1 C/C++ (nvcc) CUDA NVIDIA GPU GPU CUDA CUDA 1 CUDA CUDA 2 CUDA NVIDIA GPU PC Windows Linux MaxOSX CUDA GPU CUDA NVIDIA
Slide 1
CUDA プログラミングの基本 パート II - カーネル CUDA の基本の概要 パート I CUDAのソフトウェアスタックとコンパイル GPUのメモリ管理 パート II カーネルの起動 GPUコードの具体像 注 : 取り上げているのは基本事項のみです そのほか多数の API 関数についてはプログラミングガイドを ご覧ください GPU 上でのコードの実行 カーネルは C 関数 + 多少の制約 ホストメモリはアクセスできない戻り値型は
CUDA 連携とライブラリの活用 2
1 09:30-10:00 受付 10:00-12:00 Reedbush-H ログイン GPU 入門 13:30-15:00 OpenACC 入門 15:15-16:45 OpenACC 最適化入門と演習 17:00-18:00 OpenACC の活用 (CUDA 連携とライブラリの活用 ) CUDA 連携とライブラリの活用 2 3 OpenACC 簡単にGPUプログラムが作成できる それなりの性能が得られる
1 4 1.1........................................... 4 1.2.................................. 4 1.3................................... 4 2 5 2.1 GPU.....
CPU GPU N Q07-065 2011 2 17 1 1 4 1.1........................................... 4 1.2.................................. 4 1.3................................... 4 2 5 2.1 GPU...........................................
NUMAの構成
GPU のプログラム 天野 アクセラレータとは? 特定の性質のプログラムを高速化するプロセッサ 典型的なアクセラレータ GPU(Graphic Processing Unit) Xeon Phi FPGA(Field Programmable Gate Array) 最近出て来た Deep Learning 用ニューロチップなど Domain Specific Architecture 1GPGPU:General
概要 目的 CUDA Fortran の利用に関する基本的なノウハウを提供する 本チュートリアル受講後は Web 上で公開されている資料等を参照しながら独力で CUDA Fortran が利用できることが目標 対象 CUDA Fortran の利用に興味を抱いている方 前提とする知識 Fortran
CUDA Fortran チュートリアル 2010 年 9 月 29 日 NEC 概要 目的 CUDA Fortran の利用に関する基本的なノウハウを提供する 本チュートリアル受講後は Web 上で公開されている資料等を参照しながら独力で CUDA Fortran が利用できることが目標 対象 CUDA Fortran の利用に興味を抱いている方 前提とする知識 Fortran を用いた Linux
1. GPU コンピューティング GPU コンピューティング GPUによる 汎用コンピューティング GPU = Graphics Processing Unit CUDA Compute Unified Device Architecture NVIDIA の GPU コンピューティング環境 Lin
Windows で始める CUDA 入門 GTC 2013 チュートリアル エヌビディアジャパン CUDA エンジニア森野慎也 1. GPU コンピューティング GPU コンピューティング GPUによる 汎用コンピューティング GPU = Graphics Processing Unit CUDA Compute Unified Device Architecture NVIDIA の GPU コンピューティング環境
CUDA を用いた画像処理 画像処理を CUDA で並列化 基本的な並列化の考え方 目標 : 妥当な Naïve コードが書ける 最適化の初歩がわかる ブロックサイズ メモリアクセスパターン
CUDA 画像処理入門 エヌビディアジャパン CUDA エンジニア森野慎也 GTC Japan 2014 CUDA を用いた画像処理 画像処理を CUDA で並列化 基本的な並列化の考え方 目標 : 妥当な Naïve コードが書ける 最適化の初歩がわかる ブロックサイズ メモリアクセスパターン RGB Y( 輝度 ) 変換 カラー画像から グレイスケールへの変換 Y = 0.299 R + 0.587
GPU Computing on Business
GPU Computing on Business 2010 Numerical Technologies Incorporated http://www.numtech.com/ 1 2 3 4 5 6 7 8 9 GPU Computing $$$ Revenue Total Cost low BEP Quantity 10 11 12 13 14 15 GPU Computing $$$ Revenue
‚æ4›ñ
( ) ( ) ( ) A B C D E F G H I J K L M N O P Q R S T U V W X Y Z a b c d e f g h i j k l m n o p q r s t u v w x y z 0 1 2 3 4 5 6 7 8 9 (OUS) 9 26 1 / 28 ( ) ( ) ( ) A B C D Z a b c d z 0 1 2 9 (OUS) 9
MPI または CUDA を用いた将棋評価関数学習プログラムの並列化 2009/06/30
MPI または CUDA を用いた将棋評価関数学習プログラムの並列化 2009/06/30 目次 1. まえがき 3 2. 計算方法 4 3. MPI を用いた並列化 6 4. CUDA を用いた並列化 11 5. 計算結果 20 6. まとめ 24 2 1. まえがき 目的将棋の評価関数を棋譜から学習するボナンザメソッドの簡易版を作成し それを MPI または CUDA を用いて並列化し 計算時間を短縮することを目的とする
CUDA基礎１
CUDA 基礎 1 東京工業大学 学術国際情報センター 黄遠雄 2016/6/27 第 20 回 GPU コンピューティング講習会 1 ヘテロジニアス コンピューティング ヘテロジニアス コンピューティング (CPU + GPU) は広く使われている Financial Analysis Scientific Simulation Engineering Simulation Data Intensive
1206_Cray_PE_Overview+Roadmap_JPN
Cray プログラミング環境の開発と現状 寺西慶太 Cray Inc. Cray のプログラミング環境へのビジョン Crayはアプリケーション性能を最大限に向アプリケーション性能を最大限に向上させることを目標にプログラミング環境を研究開発 コンパイラ ライブラリ ツールの統合されたプログラミング環境で HPC プログラミングの複雑さを克服 スケーラービリティ 機能の拡張と自動化による使いやすさの向上
2 09:00-09:30 受付 09:30-12:00 GPU 入門,CUDA 入門 13:00-14:30 OpenACC 入門 + HA-PACS ログイン 14:45-16:15 OpenACC 最適化入門と演習 16:30-18:00 CUDA 最適化入門と演習
担当 大島聡史 ( 助教 ) ohshima@cc.u-tokyo.ac.jp 星野哲也 ( 助教 ) hoshino@cc.u-tokyo.ac.jp 質問やサンプルプログラムの提供についてはメールでお問い合わせください 1 2016 年 6 月 8 日 ( 水 ) 東京大学情報基盤センター 2 09:00-09:30 受付 09:30-12:00 GPU 入門,CUDA 入門 13:00-14:30
TSUBAME2.0におけるGPUの 活用方法
GPU プログラミング 基礎編 東京工業大学学術国際情報センター 1. GPU コンピューティングと TSUBAME2.0 スーパーコンピュータ GPU コンピューティングとは グラフィックプロセッサ (GPU) は グラフィック ゲームの画像計算のために 進化を続けてきた 現在 CPU のコア数は 2~12 個に対し GPU 中には数百コア その GPU を一般アプリケーションの高速化に利用! GPGPU
double float
2015 3 13 1 2 2 3 2.1.......................... 3 2.2............................. 3 3 4 3.1............................... 4 3.2 double float......................... 5 3.3 main.......................
Slide 1
CUDA プログラミングの基本 パート I - ソフトウェアスタックとメモリ管理 CUDA の基本の概要 パート I CUDAのソフトウェアスタックとコンパイル GPUのメモリ管理 パートII カーネルの起動 GPUコードの具体項目 注 : 取り上げているのは基本事項のみです そのほか多数の API 関数についてはプログラミングガイドを ご覧ください CUDA インストレーション CUDA インストレーションの構成
GPU GPU CPU CPU CPU GPU GPU N N CPU ( ) 1 GPU CPU GPU 2D 3D CPU GPU GPU GPGPU GPGPU 2 nvidia GPU CUDA 3 GPU 3.1 GPU Core 1
GPU 4 2010 8 28 1 GPU CPU CPU CPU GPU GPU N N CPU ( ) 1 GPU CPU GPU 2D 3D CPU GPU GPU GPGPU GPGPU 2 nvidia GPU CUDA 3 GPU 3.1 GPU Core 1 Register & Shared Memory ( ) CPU CPU(Intel Core i7 965) GPU(Tesla
３次多項式パラメタ推定計算の CUDAを用いた実装 （CUDAプログラミングの練習として） Implementation of the Estimation of the parameters of 3rd-order-Polynomial with CUDA
3 次多項式パラメタ推定計算の CUDA を用いた実装 (CUDA プログラミングの練習として ) Estimating the Parameters of 3rd-order-Polynomial with CUDA ISS 09/11/12 問題の選択 目的 CUDA プログラミングを経験 ( 試行錯誤と習得 ) 実際に CPU のみの場合と比べて高速化されることを体験 問題 ( インプリメントする内容
CudaWaveField
CudaWaveField 2012 3 22 2 CudaWaveField Rel 1.0.0 Rel 1.0 CudaWaveField ( cwfl) / cwfl cwfl http://www.laser.ee.kansai-u.ac.jp/wavefieldtools Note Acrobat Reader 3 I CudaWaveField 9 1 11 1.1 CudaWaveField......................
GPGPUクラスタの性能評価
2008 年度理研 HPC シンポジウム第 3 世代 PC クラスタ GPGPU クラスタの性能評価 2009 年 3 月 12 日 富士通研究所成瀬彰 発表の概要 背景 GPGPU による高速化 CUDA の概要 GPU のメモリアクセス特性調査 姫野 BMT の高速化 GPGPU クラスタによる高速化 GPU Host 間のデータ転送 GPU-to-GPU の通信性能 GPGPU クラスタ上での姫野
:30 12:00 I. I VI II. III. IV. a d V. VI
2018 2018 08 02 10:30 12:00 I. I VI II. III. IV. a d V. VI. 80 100 60 1 I. Backus-Naur BNF N N y N x N xy yx : yxxyxy N N x, y N (parse tree) (1) yxyyx (2) xyxyxy (3) yxxyxyy (4) yxxxyxxy N y N x N yx
「産業上利用することができる発明」の審査の運用指針（案）
1 1.... 2 1.1... 2 2.... 4 2.1... 4 3.... 6 4.... 6 1 1 29 1 29 1 1 1. 2 1 1.1 (1) (2) (3) 1 (4) 2 4 1 2 2 3 4 31 12 5 7 2.2 (5) ( a ) ( b ) 1 3 2 ( c ) (6) 2. 2.1 2.1 (1) 4 ( i ) ( ii ) ( iii ) ( iv)
IPSJ SIG Technical Report Vol.2013-HPC-138 No /2/21 GPU CRS 1,a) 2,b) SpMV GPU CRS SpMV GPU NVIDIA Kepler CUDA5.0 Fermi GPU Kepler Kepler Tesla
GPU CRS 1,a),b) SpMV GPU CRS SpMV GPU NVIDIA Kepler CUDA5.0 Fermi GPU Kepler Kepler Tesla K0 CUDA5.0 cusparse CRS SpMV 00 1.86 177 1. SpMV SpMV CRS Compressed Row Storage *1 SpMV GPU GPU NVIDIA Kepler
untitled
II yacc 005 : 1, 1 1 1 %{ int lineno=0; 3 int wordno=0; 4 int charno=0; 5 6 %} 7 8 %% 9 [ \t]+ { charno+=strlen(yytext); } 10 "\n" { lineno++; charno++; } 11 [^ \t\n]+ { wordno++; charno+=strlen(yytext);}
DO 時間積分 START 反変速度の計算 contravariant_velocity 移流項の計算 advection_adams_bashforth_2nd DO implicit loop( 陰解法 ) 速度勾配, 温度勾配の計算 gradient_cell_center_surface 速
1 1, 2 1, 2 3 2, 3 4 GP LES ASUCA LES NVIDIA CUDA LES 1. Graphics Processing Unit GP General-Purpose SIMT Single Instruction Multiple Threads 1 2 3 4 1),2) LES Large Eddy Simulation 3) ASUCA 4) LES LES
(MIRU2010) NTT Graphic Processor Unit GPU graphi
(MIRU2010) 2010 7 889 2192 1-1 905 2171 905 NTT 243 0124 3-1 E-mail: ac094608@edu.okinawa-ct.ac.jp, akisato@ieee.org Graphic Processor Unit GPU graphic processor unit CUDA Fully automatic extraction of
Microsoft PowerPoint - GPGPU実践基礎工学（web）.pptx
GPU のプログラム構造 長岡技術科学大学電気電子情報工学専攻出川智啓 今回の内容 GPU プログラミング環境 (CUDA) GPU プログラムの実行の流れ CUDA によるプログラムの記述 カーネル (GPU で処理する関数 ) の構造 記述方法とその理由 GPU 固有のパラメータの確認 405 GPU(Graphics Processing Unit) とは 画像処理専用のハードウェア 具体的には画像処理用のチップ
a n a n ( ) (1) a m a n = a m+n (2) (a m ) n = a mn (3) (ab) n = a n b n (4) a m a n = a m n ( m > n ) m n 4 ( ) 552
3 3.0 a n a n ( ) () a m a n = a m+n () (a m ) n = a mn (3) (ab) n = a n b n (4) a m a n = a m n ( m > n ) m n 4 ( ) 55 3. (n ) a n n a n a n 3 4 = 8 8 3 ( 3) 4 = 8 3 8 ( ) ( ) 3 = 8 8 ( ) 3 n n 4 n n
I. Backus-Naur BNF S + S S * S S x S +, *, x BNF S (parse tree) : * x + x x S * S x + S S S x x (1) * x x * x (2) * + x x x (3) + x * x + x x (4) * *
2015 2015 07 30 10:30 12:00 I. I VI II. III. IV. a d V. VI. 80 100 60 1 I. Backus-Naur BNF S + S S * S S x S +, *, x BNF S (parse tree) : * x + x x S * S x + S S S x x (1) * x x * x (2) * + x x x (3) +
untitled
CUDA Vol I: CUDA : NVIDIA Q2 2008 GPU...1 CUDA...10 CUDA...25...56 G8x...62...67...105...115 CUDA...128 CUBLAS...130 CUFFT...142 CUDA...157 CUDA...159 CUDA Fortran...168 CUDA API...173...174 CUDA...176
GPU 画像 動画処理用ハードウェア 低性能なプロセッサがたくさん詰まっている ピーク性能が非常に高い GPUを数値計算に用いるのがGPGPU Graphics Processing Unit General Purpose GPU TSUBAME2.0: GPUスパコン 本演習ではNVIDIA社の
演習II (連続系アルゴリズム) 第2回: GPGPU 須田研究室 M1 本谷 徹 motoya@is.s.u-tokyo.ac.jp 2012/10/19 GPU 画像 動画処理用ハードウェア 低性能なプロセッサがたくさん詰まっている ピーク性能が非常に高い GPUを数値計算に用いるのがGPGPU Graphics Processing Unit General Purpose GPU TSUBAME2.0:
44 4 I (1) ( ) (10 15 ) ( 17 ) ( 3 1 ) (2)
(1) I 44 II 45 III 47 IV 52 44 4 I (1) ( ) 1945 8 9 (10 15 ) ( 17 ) ( 3 1 ) (2) 45 II 1 (3) 511 ( 451 1 ) ( ) 365 1 2 512 1 2 365 1 2 363 2 ( ) 3 ( ) ( 451 2 ( 314 1 ) ( 339 1 4 ) 337 2 3 ) 363 (4) 46
i ii i iii iv 1 3 3 10 14 17 17 18 22 23 28 29 31 36 37 39 40 43 48 59 70 75 75 77 90 95 102 107 109 110 118 125 128 130 132 134 48 43 43 51 52 61 61 64 62 124 70 58 3 10 17 29 78 82 85 102 95 109 iii
Microsoft Word - C.....u.K...doc
C uwêííôöðöõ Ð C ÔÖÐÖÕ ÐÊÉÌÊ C ÔÖÐÖÕÊ C ÔÖÐÖÕÊ Ç Ê Æ ~ if eíè ~ for ÒÑÒ ÌÆÊÉÉÊ ~ switch ÉeÍÈ ~ while ÒÑÒ ÊÍÍÔÖÐÖÕÊ ~ 1 C ÔÖÐÖÕ ÐÊÉÌÊ uê~ ÏÒÏÑ Ð ÓÏÖ CUI Ô ÑÊ ÏÒÏÑ ÔÖÐÖÕÎ d ÈÍÉÇÊ ÆÒ Ö ÒÐÑÒ ÊÔÎÏÖÎ d ÉÇÍÊ
Slides: TimeGraph: GPU Scheduling for Real-Time Multi-Tasking Environments
計算機アーキテクチャ第 11 回 マルチプロセッサ 本資料は授業用です 無断で転載することを禁じます 名古屋大学 大学院情報科学研究科 准教授加藤真平 デスクトップ ジョブレベル並列性 スーパーコンピュータ 並列処理プログラム プログラムの並列化 for (i = 0; i < N; i++) { x[i] = a[i] + b[i]; } プログラムの並列化 x[0] = a[0] + b[0];
For_Beginners_CAPL.indd
CAPL Vector Japan Co., Ltd. 目次 1 CAPL 03 2 CAPL 03 3 CAPL 03 4 CAPL 04 4.1 CAPL 4.2 CAPL 4.3 07 5 CAPL 08 5.1 CANoe 5.2 CANalyzer 6 CAPL 10 7 CAPL 11 7.1 CAPL 7.2 CAPL 7.3 CAPL 7.4 CAPL 16 7.5 18 8 CAPL
8 if switch for while do while 2
(Basic Theory of Information Processing) ( ) if for while break continue 1 8 if switch for while do while 2 8.1 if (p.52) 8.1.1 if 1 if ( ) 2; 3 1 true 2 3 false 2 3 3 8.1.2 if-else (p.54) if ( ) 1; else
Java演習(4) -- 変数と型 --
50 20 20 5 (20, 20) O 50 100 150 200 250 300 350 x (reserved 50 100 y 50 20 20 5 (20, 20) (1)(Blocks1.java) import javax.swing.japplet; import java.awt.graphics; (reserved public class Blocks1 extends
EGunGPU
Super Computing in Accelerator simulations - Electron Gun simulation using GPGPU - K. Ohmi, KEK-Accel Accelerator Physics seminar 2009.11.19 Super computers in KEK HITACHI SR11000 POWER5 16 24GB 16 134GFlops,
GPGPUイントロダクション
大島聡史 ( 並列計算分科会主査 東京大学情報基盤センター助教 ) GPGPU イントロダクション 1 目的 昨今注目を集めている GPGPU(GPU コンピューティング ) について紹介する GPGPU とは何か? 成り立ち 特徴 用途 ( ソフトウェアや研究例の紹介 ) 使い方 ( ライブラリ 言語 ) CUDA GPGPU における課題 2 GPGPU とは何か? GPGPU General-Purpose
! 行行 CPUDSP PPESPECell/B.E. CPUGPU 行行 SIMD [SSE, AltiVec] 用 HPC CPUDSP PPESPE (Cell/B.E.) SPE CPUGPU GPU CPU DSP DSP PPE SPE SPE CPU DSP SPE 2
![! 行行 CPUDSP PPESPECell/B.E. CPUGPU 行行 SIMD [SSE, AltiVec] 用 HPC CPUDSP PPESPE (Cell/B.E.) SPE CPUGPU GPU CPU DSP DSP PPE SPE SPE CPU DSP SPE 2](/thumbs/88/115445009.jpg "! 行行 CPUDSP PPESPECell/B.E. CPUGPU 行行 SIMD [SSE, AltiVec] 用 HPC CPUDSP PPESPE (Cell/B.E.) SPE CPUGPU GPU CPU DSP DSP PPE SPE SPE CPU DSP SPE 2")
! OpenCL [Open Computing Language] 言 [OpenCL C 言 ] CPU, GPU, Cell/B.E.,DSP 言 行行 [OpenCL Runtime] OpenCL C 言 API Khronos OpenCL Working Group AMD Broadcom Blizzard Apple ARM Codeplay Electronic Arts Freescale
2011 年 3 月 3 日 GPGPU ハンズオンプログラミング演習 株式会社クロスアビリティ ability.jp 3 Mar 2011 Copyright (C) 2011 X-Ability Co.,Ltd. All rights reserved.
2011 年 3 月 3 日 GPGPU ハンズオンプログラミング演習 株式会社クロスアビリティ rkoga@x ability.jp 講師 : 古賀良太 / 古川祐貴 取締役計算化学ソルバー XA CHEM SUITE の開発 コンサルティングパートナー 並列ソフトウェアの開発 ビルド サーバ販売 ソフトウェア代理店 会社紹介 社名株式会社クロスアビリティ (X Ability Co.,Ltd)
PC Windows 95, Windows 98, Windows NT, Windows 2000, MS-DOS, UNIX CPU
1. 1.1. 1.2. 1 PC Windows 95, Windows 98, Windows NT, Windows 2000, MS-DOS, UNIX CPU 2. 2.1. 2 1 2 C a b N: PC BC c 3C ac b 3 4 a F7 b Y c 6 5 a ctrl+f5) 4 2.2. main 2.3. main 2.4. 3 4 5 6 7 printf printf
解きながら学ぶC++入門編
!... 38!=... 35 "... 112 " "... 311 " "... 4, 264 #... 371 #define... 126, 371 #endif... 369 #if... 369 #ifndef... 369 #include... 3, 311 #undef... 371 %... 17, 18 %=... 85 &... 222 &... 203 &&... 40 &=...
I. Backus-Naur BNF : N N 0 N N N N N N 0, 1 BNF N N 0 11 (parse tree) 11 (1) (2) (3) (4) II. 0(0 101)* (
2016 2016 07 28 10:30 12:00 I. I VI II. III. IV. a d V. VI. 80 100 60 1 I. Backus-Naur BNF : 11011 N N 0 N N 11 1001 N N N N 0, 1 BNF N N 0 11 (parse tree) 11 (1) 1100100 (2) 1111011 (3) 1110010 (4) 1001011
untitled
A = QΛQ T A n n Λ Q A = XΛX 1 A n n Λ X GPGPU A 3 T Q T AQ = T (Q: ) T u i = λ i u i T {λ i } {u i } QR MR 3 v i = Q u i A {v i } A n = 9000 Quad Core Xeon 2 LAPACK (4/3) n 3 O(n 2 ) O(n 3 ) A {v i }
TSUBAMEのTeslaを用いたGPGPU（CUDAの基礎）
GPU コンピューティング (CUDA) 講習会 CUDA プログラミング基礎 丸山直也 2010/09/13 1 はじめに 本講習では時間の関係上ごくごく基礎的な内容のみをとりあげます ただし 資料の後半にはメモリアクセスなどに関するチューニングに向けた情報をのせてあります それらは講習時間内には取り上げません チューニングやよりアドバンストな内容の講習会は別途開催しております 本講習で取り上げる概念等は基礎的なものに限られるため
Java Java Java Java Java 4 p * *** ***** *** * Unix p a,b,c,d 100,200,250,500 a*b = a*b+c = a*b+c*d = (a+b)*(c+d) = 225
Java Java Java Java Java 4 p35 4-2 * *** ***** *** * Unix p36 4-3 a,b,c,d 100,200,250,500 a*b = 20000 a*b+c = 20250 a*b+c*d = 145000 (a+b)*(c+d) = 225000 a+b*c+d = 50600 b/a+d/c = 4 p38 4-4 (1) mul = 1
卒業論文
卒業論文 GPU を用いた液晶用ガラスの欠損検出画像処理の高速化 (Ⅲ) 氏名 : 原田健史学籍番号 :2260060085-4 指導教員 : 山崎勝弘教授提出日 :2011 年 2 月 17 日 立命館大学理工学部電子情報デザイン学科 1 内容梗概 液晶用ガラスは近年のデジタルテレビやスマートフォンの普及により採用数が激増している液晶パネルの主要部品である 液晶用ガラスはその加工工程において高温処理が必要とされるため
(search: ) [1] ( ) 2 (linear search) (sequential search) 1
![(search: ) [1] ( ) 2 (linear search) (sequential search) 1](/thumbs/97/131825822.jpg "(search: ) [1] ( ) 2 (linear search) (sequential search) 1")
2005 11 14 1 1.1 2 1.2 (search:) [1] () 2 (linear search) (sequential search) 1 2.1 2.1.1 List 2-1(p.37) 1 1 13 n
23 Fig. 2: hwmodulev2 3. Reconfigurable HPC 3.1 hw/sw hw/sw hw/sw FPGA PC FPGA PC FPGA HPC FPGA FPGA hw/sw hw/sw hw- Module FPGA hwmodule hw/sw FPGA h
23 FPGA CUDA Performance Comparison of FPGA Array with CUDA on Poisson Equation (lijiang@sekine-lab.ei.tuat.ac.jp), (kazuki@sekine-lab.ei.tuat.ac.jp), (takahashi@sekine-lab.ei.tuat.ac.jp), (tamukoh@cc.tuat.ac.jp),
FFTSS Library Version 3.0 User's Guide
: 19 10 31 FFTSS 3.0 Copyright (C) 2002-2007 The Scalable Software Infrastructure Project, (CREST),,. http://www.ssisc.org/ Contents 1 4 2 (DFT) 4 3 4 3.1 UNIX............................................
1 GPU GPGPU GPU CPU 2 GPU 2007 NVIDIA GPGPU CUDA[3] GPGPU CUDA GPGPU CUDA GPGPU GPU GPU GPU Graphics Processing Unit LSI LSI CPU ( ) DRAM GPU LSI GPU
![1 GPU GPGPU GPU CPU 2 GPU 2007 NVIDIA GPGPU CUDA[3] GPGPU CUDA GPGPU CUDA GPGPU GPU GPU GPU Graphics Processing Unit LSI LSI CPU ( ) DRAM GPU LSI GPU](/thumbs/89/99132402.jpg "1 GPU GPGPU GPU CPU 2 GPU 2007 NVIDIA GPGPU CUDA[3] GPGPU CUDA GPGPU CUDA GPGPU GPU GPU GPU Graphics Processing Unit LSI LSI CPU ( ) DRAM GPU LSI GPU")
GPGPU (I) GPU GPGPU 1 GPU(Graphics Processing Unit) GPU GPGPU(General-Purpose computing on GPUs) GPU GPGPU GPU ( PC ) PC PC GPU PC PC GPU GPU 2008 TSUBAME NVIDIA GPU(Tesla S1070) TOP500 29 [1] 2009 AMD
Microsoft PowerPoint - GPGPU実践基礎工学（web）.pptx
GPU のメモリ階層 長岡技術科学大学電気電子情報工学専攻出川智啓 今回の内容 GPU のメモリ階層 グローバルメモリ 共有メモリ モザイク処理への適用 コンスタントメモリ 空間フィルタへの適用 577 GPU の主要部品 基盤 GPU( チップ )+ 冷却部品 画面出力端子 電源入力端子 メモリ 特性の把握が重要 電源入力端子 画面出力端子 メモリ チップ PCI Ex 端子 http://www.geforce.com/whats
CPU Levels in the memory hierarchy Level 1 Level 2... Increasing distance from the CPU in access time Level n Size of the memory at each level 1: 2.2
FFT 1 Fourier fast Fourier transform FFT FFT FFT 1 FFT FFT 2 Fourier 2.1 Fourier FFT Fourier discrete Fourier transform DFT DFT n 1 y k = j=0 x j ω jk n, 0 k n 1 (1) x j y k ω n = e 2πi/n i = 1 (1) n DFT
untitled
A = QΛQ T A n n Λ Q A = XΛX 1 A n n Λ X GPGPU A 3 T Q T AQ = T (Q: ) T u i = λ i u i T {λ i } {u i } QR MR 3 v i = Q u i A {v i } A n = 9000 Quad Core Xeon 2 LAPACK (4/3) n 3 O(n 2 ) O(n 3 ) A {v i }
O(N) ( ) log 2 N
2005 11 21 1 1.1 2 O(N) () log 2 N 1.2 2 1 List 3-1 List 3-3 List 3-4? 3 3.1 3.1.1 List 2-1(p.70) 1 1 10 1 3.1.2 List 3-1(p.70-71) 1 1 2 1 2 2 1: 1 3 3.1.3 1 List 3-1(p.70-71) 2 #include stdlib.h
untitled
23 12 10 12:55 ~ 18:45 KKR Tel0557-85-2000 FAX0557-85-6604 12:55~13:00 13:00~13:38 I 1) 13:00~13:12 2) 13:13~13:25 3) 13:26~13:38 13:39~14:17 II 4) 13:39~13:51 5) 13:52 ~ 14:04 6) 14:05 ~ 14:17 14:18 ~
活用ガイド (ソフトウェア編)
ii iii iv NEC Corporation 1998 v vi PA RT 1 vii PA RT 2 viii PA RT 3 PA RT 4 ix P A R T 1 2 3 1 4 5 1 1 2 1 2 3 4 6 1 2 3 4 5 7 1 6 7 8 1 9 1 10 1 2 3 4 5 6 7 8 9 10 11 11 1 12 12 1 13 1 1 14 2 3 4 5 1
(1) (2) (1) (2) 2 3 {a n } a 2 + a 4 + a a n S n S n = n = S n
. 99 () 0 0 0 () 0 00 0 350 300 () 5 0 () 3 {a n } a + a 4 + a 6 + + a 40 30 53 47 77 95 30 83 4 n S n S n = n = S n 303 9 k d 9 45 k =, d = 99 a d n a n d n a n = a + (n )d a n a n S n S n = n(a + a n
ohp11.dvi
19 11 ( ) 2019.4.20 1 / ( ) n O(n 2 ) O(n 2 ) ( ) 1 d n 1 n logn O(nlogn) n ( n logn C ) 2 ( ) ( merge) 2 1 1 3 1 4 5 4 2 3 7 9 7 1 2 3 4 5 7 9 1: 2 ivec merge 3 ( ) (2) int *ivec_new(int size) { int *a
r11.dvi
19 11 ( ) 2019.4.20 1 / 1.1 ( n n O(n 2 O(n 2 ) ( 1 d n 1 n logn O(nlogn n ( n logn C 1.2 ( ( merge 2 1 1 3 1 4 5 4 2 3 7 9 7 1 2 3 4 5 7 9 1: 2 ivec merge int *ivec_new(int size) { int *a = (int*)malloc((size+1)
入門ガイド
ii iii iv NEC Corporation 1998 v P A R 1 P A R 2 P A R 3 T T T vi P A R T 4 P A R T 5 P A R T 6 P A R T 7 vii 1P A R T 1 2 2 1 3 1 4 1 1 5 2 3 6 4 1 7 1 2 3 8 1 1 2 3 9 1 2 10 1 1 2 11 3 12 1 2 1 3 4 13
Copyright 2008 All Rights Reserved 2
Copyright 2008 All Rights Reserved 1 Copyright 2008 All Rights Reserved 2 Copyright 2008 All Rights Reserved 3 Copyright 2008 All Rights Reserved 4 Copyright 2008 All Rights Reserved 5 Copyright 2008 All
ハピタス のコピー.pages
Copyright (C) All Rights Reserved. 10 12,500 () ( ) ()() 1 : 2 : 3 : 2 4 : 5 : Copyright (C) All Rights Reserved. Copyright (C) All Rights Reserved. Copyright (C) All Rights Reserved. Copyright (C) All
Java updated
Java 2003.07.14 updated 3 1 Java 5 1.1 Java................................. 5 1.2 Java..................................... 5 1.3 Java................................ 6 1.3.1 Java.......................
Copyright 2008 NIFTY Corporation All rights reserved. 2
Copyright 2008 NIFTY Corporation All rights reserved. 2 Copyright 2008 NIFTY Corporation All rights reserved. 3 Copyright 2008 NIFTY Corporation All rights reserved. 4 Copyright 2008 NIFTY Corporation
SystemC言語概論
SystemC CPU S/W 2004/01/29 4 SystemC 1 SystemC 2.0.1 CPU S/W 3 ISS SystemC Co-Simulation 2004/01/29 4 SystemC 2 ISS SystemC Co-Simulation GenericCPU_Base ( ) GenericCPU_ISS GenericCPU_Prog GenericCPU_CoSim
Java学習教材
Java 2016/4/17 Java 1 Java1 : 280 : (2010/1/29) ISBN-10: 4798120987 ISBN-13: 978-4798120980 2010/1/29 1 Java 1 Java Java Java class FirstExample { public static void main(string[] args) { System.out.println("
1. A0 A B A0 A : A1,...,A5 B : B1,...,B
1. A0 A B A0 A : A1,...,A5 B : B1,...,B12 2. 3. 4. 5. A0 A B f : A B 4 (i) f (ii) f (iii) C 2 g, h: C A f g = f h g = h (iv) C 2 g, h: B C g f = h f g = h 4 (1) (i) (iii) (2) (iii) (i) (3) (ii) (iv) (4)
いまからはじめる組み込みGPU実装
いまからはじめる組み込み GPU 実装 ~ コンピュータービジョン ディープラーニング編 ~ MathWorks Japan アプリケーションエンジニアリング部シニアアプリケーションエンジニア大塚慶太郎 2017 The MathWorks, Inc. 1 コンピュータービジョン ディープラーニングによる 様々な可能性 自動運転 ロボティクス 予知保全 ( 製造設備 ) セキュリティ 2 転移学習を使った画像分類
XACCの概要
2 global void kernel(int a[max], int llimit, int ulimit) {... } : int main(int argc, char *argv[]){ MPI_Int(&argc, &argc); MPI_Comm_rank(MPI_COMM_WORLD, &rank); MPI_Comm_size(MPI_COMM_WORLD, &size); dx
=
2. 2.1 2.2 kuri@ice.uec.ac.jp ( 2007/10/30/16:46) 1 . 1. 1 + 2 = 5. 2. 180. 3. 3 3. 4.. 5.. 2 2.1 1.,,,,. 2., ( ) ( ).,,,, 3.,. 4.,,,. 3 1.,. 1. 1 + 2 = 5. (, ) 2. 180. (, ) 3. 3, 3. (, ) 4.. (, ) 5..
i
14 i ii iii iv v vi 14 13 86 13 12 28 14 16 14 15 31 (1) 13 12 28 20 (2) (3) 2 (4) (5) 14 14 50 48 3 11 11 22 14 15 10 14 20 21 20 (1) 14 (2) 14 4 (3) (4) (5) 12 12 (6) 14 15 5 6 7 8 9 10 7
Microsoft PowerPoint _OpenCAE並列計算分科会.pptx
地球流体力学に関する GPGPU を用いた数値計算 神戸大学惑星科学研究センター西澤誠也 地球流体力学とは 地球 惑星に関連がある流体の力学 回転, 重力の影響 e.g. 大気, 海洋, マントル 数値計算は天気予報 & 弾道軌道予測から始まった ベクトル計算機 地球流体の計算はベクトル長が長いものが多い ベクトル計算機の凋落 某社の次世代スパコンからの撤退 個人的スパコンの将来予想 個々の演算器はシンプルに
- - http://168iroha.net 018 10 14 i 1 1 1.1.................................................... 1 1.................................................... 7.1................................................
GPU のアーキテクチャとプログラム構造 長岡技術科学大学電気電子情報工学専攻出川智啓
GPU のアーキテクチャとプログラム構造 長岡技術科学大学電気電子情報工学専攻出川智啓 今回の内容 GPU のアーキテクチャ CUDA CUDA によるプログラミング 58 GPU(Graphics Processing Unit) とは 画像処理専用のハードウェア 具体的には画像処理用のチップ チップ単体では販売されていない PCI Ex カードで販売 ( チップ単体と区別せずに GPU と呼ぶことも多い
1. マシンビジョンにおける GPU の活用
CUDA 画像処理入門 GTC 213 チュートリアル エヌビディアジャパン CUDA エンジニア森野慎也 1. マシンビジョンにおける GPU の活用 1. 医用画像処理における GPU の活用 CT や MRI から画像を受信して三次元画像の構築をするシステム 2 次元スキャンデータから 3 次元 4 次元イメージの高速生成 CUDA 化により画像処理速度を約 2 倍に高速化 1. CUDA で画像処理
rank ”«‘‚“™z‡Ì GPU ‡É‡æ‡éŁÀŠñ›»
rank GPU ERATO 2011 11 1 1 / 26 GPU rank/select wavelet tree balanced parenthesis GPU rank 2 / 26 GPU rank/select wavelet tree balanced parenthesis GPU rank 2 / 26 GPU rank/select wavelet tree balanced
Microsoft PowerPoint - CproNt02.ppt [互換モード]
![Microsoft PowerPoint - CproNt02.ppt [互換モード]](/thumbs/95/124866096.jpg "Microsoft PowerPoint - CproNt02.ppt [互換モード]")
第 2 章 C プログラムの書き方 CPro:02-01 概要 C プログラムの構成要素は関数 ( プログラム = 関数の集まり ) 関数は, ヘッダと本体からなる 使用する関数は, プログラムの先頭 ( 厳密には, 使用場所より前 ) で型宣言 ( プロトタイプ宣言 ) する 関数は仮引数を用いることができる ( なくてもよい ) 関数には戻り値がある ( なくてもよい void 型 ) コメント
課題
float[] xball; float[] yball; int numberofballs = (a) ; int radius=10; size(400,400); xball = (b) (c) [numberofballs]; yball = (d) (e) [numberofballs]; xball[i] = random(radius,width-radius); yball[i]
熊本大学学術リポジトリ Kumamoto University Repositor Title GPGPU による高速演算について Author(s) 榎本, 昌一 Citation Issue date Type URL Presentation
熊本大学学術リポジトリ Kumamoto University Repositor Title GPGPU による高速演算について Author(s) 榎本, 昌一 Citation Issue date 2011-03-17 Type URL Presentation http://hdl.handle.net/2298/23539 Right GPGPU による高速演算について 榎本昌一 東京大学大学院工学系研究科システム創成学専攻
ストリーミング SIMD 拡張命令2 (SSE2) を使用した、倍精度浮動小数点ベクトルの最大/最小要素とそのインデックスの検出
SIMD 2(SSE2) / 2.0 2000 7 : 248602J-001 01/10/30 1 305-8603 115 Fax: 0120-47-8832 * Copyright Intel Corporation 1999-2001 01/10/30 2 1...5 2...5 2.1...5 2.1.1...5 2.1.2...8 3...9 3.1...9 3.2...9 4...9
2 ColorSpace DepthSpace CameraSpace Kinect V2 Kinect V2 BOdyIndex 3. NtKinect Kinect V2 C++ NtKinect [4] NtKinect = Kinect SDK + + STL(C++) + OpenCV +
![2 ColorSpace DepthSpace CameraSpace Kinect V2 Kinect V2 BOdyIndex 3. NtKinect Kinect V2 C++ NtKinect [4] NtKinect = Kinect SDK + + STL(C++) + OpenCV +](/thumbs/78/77559809.jpg "2 ColorSpace DepthSpace CameraSpace Kinect V2 Kinect V2 BOdyIndex 3. NtKinect Kinect V2 C++ NtKinect [4] NtKinect = Kinect SDK + + STL(C++) + OpenCV +")
NtKinect: C++ Class Library for Kinect V2 1,a) Kinect for Windows V2 C++ NtKinect NtKinect DLL Kinect V2 Kinect V2, C++, DLL, Unity NtKinect: C++ Class Library for Kinect V2 Nitta Yoshihisa 1,a) Abstract:
ex01.dvi
,. 0. 0.0. C () /******************************* * $Id: ex_0_0.c,v.2 2006-04-0 3:37:00+09 naito Exp $ * * 0. 0.0 *******************************/ #include int main(int argc, char **argv) double
実際の株価データを用いたオプション料の計算
2002 2 20 1 1 3 2 3 2.1 : : : : : : : : : : : : : : : : : : : : : : : : : : : : 5 2.1.1 : : : : : : : : : : : : : : : : : : : : 5 2.1.2 : : : : : : : : : : : : : : : : : : : : 6 2.2 : : : : : : : : : :
Cell/B.E. BlockLib
Cell/B.E. BlockLib 17 17115080 21 2 10 i Cell/B.E. BlockLib SIMD CELL SIMD Cell Cell BlockLib BlockLib NestStep libspe1 Cell SDK 3.1 libspe2 BlockLib Cell SDK 3.1 NestStep libspe2 BlockLib BlockLib libspe1
課題
int[] scores; PFont font; int[] scores = { (a) ; PFont font; size(300,400); scores = (a); scores[0] = 10000; scores[1] = 9000; scores[2] = 5000; scores[3] = 1000; scores[4] = 30; font = loadfont("serif-48.vlw");