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

Кластер

OpenMP

MultiThreading Hello World #include #include // для beginthread() void mtPrintf( void * pArg); int main() { int t0 = 0; int t1 = 1; _beginthread(mtPrintf, 0, (void*)&t0 ); mtPrintf( (void*)&t1); Sleep( 100 ); return 0; } void mtPrintf( void * pArg ) { int * pIntArg = (int *) pArg; printf( "The function was passed %d\n", (*pIntArg) ); }

MultiThreading Hello World // создание нового потока // необходимо указать: // entry point функцию, // размер стека, при 0 – OS выберет сама // (void *) – указатель на аргументы функции _beginthread(mtPrintf, 0, (void*)&t1 ); // напечатать из основного потока mtPrintf( (void*)&t0); // подождать 100 мс // создание потока windows требует времени // если основной поток закончит выполнение // то и все дочерние потоки будут прерваны Sleep( 100 );

SSE Hello World #include struct vec4 { union { float v[4]; __m128 v4; }; int main() { vec4 a = {5.0f, 2.0f, 1.0f, 3.0f}; vec4 b = {5.0f, 3.0f, 9.0f, 7.0f}; vec4 c; c.v4 = _mm_add_ps(a.v4, b.v4); printf("c = {%.3f, %.3f, %.3f, %.3f}\n", c.v[0], c.v[1], c.v[2], c.v[3]); return 0; }

CUDA Hello World #defineN(1024*1024) __global__ void kernel ( float * data ) { int idx = blockIdx.x * blockDim.x + threadIdx.x; float x = 2.0f * f * (float) idx / (float) N; data [idx] = sinf ( sqrtf ( x ) ); } int main ( int argc, char * argv [] ) { float a [N]; float * dev = NULL; cudaMalloc ( (void**)&dev, N * sizeof ( float ) ); kernel >> ( dev ); cudaMemcpy ( a, dev, N * sizeof ( float ), cudaMemcpyDeviceToHost ); cudaFree ( dev ); for (int idx = 0; idx < N; idx++) printf("a[%d] = %.5f\n", idx, a[idx]); return 0; }

CUDA Hello World __global__ void kernel ( float * data ) { int idx = blockIdx.x * blockDim.x + threadIdx.x; // номер текущей нити float x = 2.0f * f * (float) idx / (float) N; // значение аргумента data [idx] = sinf ( sqrtf ( x ) ); // найти значение и } // записать его в массив zДля каждого элемента массива (всего N) запускается отдельная нить, вычисляющая требуемое значение. zКаждая нить обладает уникальным id

CUDA Hello World float a [N]; float * dev = NULL; // выделить память на GPU под N элементов cudaMalloc ( (void**)&dev, N * sizeof ( float ) ); // запустить N нитей блоками по 512 нитей // выполняемая на нити функция - kernel // массив данных - dev kernel >> ( dev ); // скопировать результаты из памяти GPU (DRAM) в // память CPU (N элементов) cudaMemcpy ( a, dev, N * sizeof ( float ), cudaMemcpyDeviceToHost ); // освободить память GPU cudaFree ( dev );

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