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

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

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

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 );

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

CUDA C Runtime zNVCC z.ptx y-keep

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

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

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) ); //...

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 );

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

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 Ядро

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

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

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);

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);

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

DirectX zID3D11Device yID3D11Resource yID3D11View zID3D11DeviceContext

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

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

DirectX zID3D11Device yID3D11Resource yID3D11View zID3D11DeviceContext zID3D11AsynchronousID3D11Asynchronous yID3D11Query zID3D11ComputeShader zID3DX11Effect

DirectX zID3D11Device yID3D11Resource xBuffer xStructuredBuffer xTexture yID3D11View xShaderResourceView xUnorderedAccessView xRenderTargetView

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

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

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);

DirectX Compute

Multi-GPU zCUDA zOpenMP zMPI zOS Threads

Multi-GPU CPUCPU GPUGPU

CUDA

Кластер

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

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

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

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

CUDA zcudaGetDeviceCount() zcudaSetDevice()

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

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

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

Запуск MPI

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

MPI

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

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