CUDA API & Multi-GPU zЛекторы: yБоресков А.В. (ВМиК МГУ)Боресков А.В. (ВМиК МГУ) yХарламов А. (NVidia)Харламов А. (NVidia) yМикушин Д. (НИВЦ) - презентация

1 CUDA API & Multi-GPU zЛекторы: yБоресков А.В. (ВМиК МГУ)Боресков А.В. (ВМиК МГУ) yХарламов А. (NVidia)Харламов А. (NVidia) yМикушин Д. (НИВЦ)

2 CUDA API zCUDA C zCUDA Driver API zOpenCL zDirectX Compute ATIs Compute Solution

3 CUDA C (Runtime API) zРасширение языка C zCUDA API: yРасширения языка C yRuntime библиотека

4 CUDA C Runtime float * a = new float [N]; float * dev = NULL; cudaMalloc( (void**)&dev, N * sizeof ( float ) ); dim3 threads = dim3( 512, 1 ); dim3 blocks = dim3( N / threads.x, 1 ); kernel >> ( dev ); cudaThreadSynchronize(); cudaMemcpy(a, dev, N*sizeof(float), cudaMemcpyDeviceToHost); cudaFree( dev );

5 CUDA C Runtime __global__ void kernel ( float * data ) { int idx = blockIdx.x * blockDim.x + threadIdx.x ; data [idx] = idx; }

6 CUDA C Runtime zNVCC z.ptx y-keep

7 CUDA C Runtime zNVCC z.ptx y-keep __global__ void kernel ( float * data ) { int idx = blockIdx.x * blockDim.x + threadIdx.x ; data [idx] = idx; }

8 CUDA C Runtime.entry _Z6kernelPf (.param.u32 __cudaparm__Z6kernelPf_data) {.reg.u16 %rh ;.reg.u32 %r ;.reg.f32 %f ;.loc1460 $LBB1__Z6kernelPf:.loc14100 mov.u16 %rh1, %ctaid.x; // mov.u16 %rh2, %ntid.x; // mul.wide.u16 %r1, %rh1, %rh2;// cvt.u32.u16 %r2, %tid.x; // add.u32 %r3, %r2, %r1; // cvt.rn.f32.s32 %f1, %r3; // ld.param.u32 %r4, [__cudaparm__Z6kernelPf_data]; // id:14 mul.lo.u32 %r5, %r3, 4; // add.u32 %r6, %r4, %r5; // st.global.f32 [%r6+0], %f1; // id:15.loc14110 exit; $LDWend__Z6kernelPf: } // _Z6kernelPf

9 CUDA C Driver CUdevice device; CUcontext context; CUmodule module; CUfunction function; CUdeviceptr pData; float * pHostData = new float[N]; cuInit(0); cuDeviceGetCount(&device_count); cuDeviceGet( &device, 0 ); cuCtxCreate( &context, 0, device ); cuModuleLoad( &module, "hello.cuda_runtime.ptx" ); cuModuleGetFunction( &function, module, "_Z6kernelPf" ); cuMemAlloc( &pData, N * sizeof(float) ); //...

10 CUDA C Driver //... cuFuncSetBlockShape( function, N, 1, 1 ); cuParamSeti( function, 0, pData ); cuParamSetSize( function, sizeof(void *) ); cuLaunchGrid( function, 1, 1 ); cuMemcpyDtoH( pHostData, pData, N * sizeof( float) ); cuMemFree( pData );

11 OpenCL zКроссплатформенный стандарт yGPU, CPU, Cell, … zПроблема: функциональность, но не производительность yРазный код для разных платформ yРазные расширения openGL-style

12 CUDA vs OpenCL Терминология zCUDA C yПоток (thread) yБлок потоков (thread block) yСеть (grid) yЯдро z OpenCL yЭлемент работы (work-item) y Группа работы (work-group) y N-мерное пространство индексов (ND-Range index space) y Ядро

13 CUDA vs OpenCL Спецификаторы функций zCUDA C y__global__ y__host__ y__device__ z OpenCL y__kernel yn/a

14 CUDA vs OpenCL Пространство памяти zCUDA C y__device__ y__shared__ y__constant__ ylocal z OpenCL y__global y__local y__constant y__private

15 OpenCL cl_context ctx; cl_command_queue cmd_q; cl_program program; cl_kernel kernel; cl_device_id * pDevId = NULL; ctx = clCreateContextFromType(0,CL_DEVICE_TYPE_GPU,0,0,0); clGetContextInfo(ctx,CL_CONTEXT_DEVICES,0,0,&dev_cnt); clGetContextInfo(ctx,CL_CONTEXT_DEVICES,dev_cnt,pDevId,0); cmd_q= clCreateCommandQueue(ctx,pDevId[0],0,0); program = clCreateProgramWithSource(ctx,1,pText,0,0); clBuildProgram(program, 0,0,0,0,0); kernel = clCreateKernel(program, "simple", 0);

16 OpenCL cl_mem mem = clCreateBuffer(ctx,CL_MEM_WRITE_ONLY, N*sizeof(float),0,0); clSetKernelArg(kernel, 0, sizeof(cl_mem),(void*) &mem); clSetKernelArg(kernel, 1, sizeof(int), (void*) &N); clEnqueueNDRangeKernel(cmd_q,kernel,1,0,&N,&N,0,0,0); clEnqueueReadBuffer(cmd_q, mem, CL_TRUE, 0, N*sizeof(float), pData,0,0,0); clReleaseMemObject(mem); clReleaseKernel(kernel); clReleaseProgram(program); clReleaseCommandQueue(cmd_q); clReleaseContext(ctx);

17 DirectX Compute zMicrosoft API zТесно интегрирован с Direct3D zДоступен yCS 4.x: DirectX 10 HW yCS 5.x: DirectX 11 HW

18 DirectX zID3D11Device yID3D11Resource yID3D11View zID3D11DeviceContext

19 CUDA vs DirectX Спецификаторы функций zCUDA C y__global__ y__host__ y__device__ z OpenCL yCompute Shader yn/a

20 CUDA vs DirectX Compute Пространство памяти zCUDA C y__device__ y__shared__ y__constant__ ylocal z OpenCL y[Structured]Buffer ygroupshared yConstant Buffer yn/a

21 DirectX zID3D11Device yID3D11Resource yID3D11View zID3D11DeviceContext zID3D11AsynchronousID3D11Asynchronous yID3D11Query zID3D11ComputeShader zID3DX11Effect

22 DirectX zID3D11Device yID3D11Resource xBuffer xStructuredBuffer xTexture yID3D11View xShaderResourceView xUnorderedAccessView xRenderTargetView

23 DirectX zID3D11DeviceContext yDispatch(bx, by, bz) yDispatchIndirect(pBuffer, offset) yEnd(pQuery) yGetData(g_pQuerry, NULL, 0, 0 )

24 DirectX ID3D11ComputeShader z ConstantBuffer z ShaderResourceView z UnorderedAccessView ID3D11Effect z ConstantBuffer z ShaderResourceView z UnorderedAccessView

25 DirectX ID3D11ComputeShader pContext->CSSetShader(pCS, NULL, 0); pContext->CSSetUnorderedAccessViews(0, 1, &pRWBufUAV, NULL); ID3D11Effect pEffect->GetVariableByName(tSimple)->AsUnorderedAccessView()- >SetUnorderedAccessView(pRWBufUAV); pEffect->GetTechniqueByName(tSimple)->GetPassByName(pSimple")- >Apply(0, pContext); Apply(0, pContext);">

26

27 DirectX Compute

28

29 Multi-GPU zCUDA zOpenMP zMPI zOS Threads

30 Multi-GPU CPUCPU GPUGPU

31 CUDA

32 Кластер

33 Пример main() thread0 thread1 malloc() shared resource BARRIERBARRIER thread0 thread1 malloc() shared resource

34 Пример main() thread0 thread1 malloc() shared resource BARRIERBARRIER thread0 thread1 malloc() shared resource

35 Модельная задача main() thread0 thread1 malloc() shared resource

36 Модельная задача main() thread0 thread1 malloc() shared resource printf() malloc() shared resource BARRIERBARRIER

37 CUDA zcudaGetDeviceCount() zcudaSetDevice()

38 OpenMP zomp_set_num_threads() zomp_get_thread_num() zomp_get_num_threads() z#pragma omp parallel z#pragma omp barrier

39 Модельная задача

40

41

42 MPI zMPICH2MPICH2 zИнструкция по установкеИнструкция по установке zProgramming GuideProgramming Guide

43 Запуск MPI

44 Модельная задача

45

46 MPI

47

48 Ссылки zhttp://iproc.ru/programming/openmp- visual-studio/ visual-studio/ zhttp://iproc.ru/programming/mpich- windows/ windows/ zhttp:// /mpich2/ /mpich2/

49 Ресурсы нашего курса zCUDA.CS.MSU.SUCUDA.CS.MSU.SU yМесто для вопросов и дискуссий yМесто для материалов нашего курса yМесто для ваших статей! xЕсли вы нашли какой-то интересный подход! xИли исследовали производительность разных подходов и знаете, какой из них самый быстрый! xИли знаете способы сделать работу с CUDA проще! z s teps3d.narod.rus teps3d.narod.ru z