390 likes | 544 Views
TROCHĘ HISTORII. PO CO UŻYWAĆ OPENCL?. AMD Radeon HD7870 OC 1,1 GHz. Intel Core i5 2550k OC 4,2 GHz. Intel Core i5 2550k 4 rdzenie. AMD Radeon HD7870 1280 rdzeni. DLACZEGO NIE UŻYWAĆ OPENCL?. CZEGO NAM TRZEBA?. CZAS NA NAJLEPSZE!. KOD!. TRADYCYJNE PODEJŚCIE.
E N D
TROCHĘ HISTORII Marek Zając
PO CO UŻYWAĆ OPENCL? Marek Zając
AMD Radeon HD7870 OC 1,1 GHz Intel Core i5 2550k OC 4,2 GHz Marek Zając
Intel Core i5 2550k 4 rdzenie AMD Radeon HD7870 1280 rdzeni Marek Zając
DLACZEGO NIE UŻYWAĆ OPENCL? Marek Zając
CZEGO NAM TRZEBA? Marek Zając
CZAS NA NAJLEPSZE! Marek Zając
KOD! Marek Zając
TRADYCYJNE PODEJŚCIE Marek Zając
double* tab = newdouble[100000000]; //preparearray... //... for(int i = 0; i < 100000000; i++) { tab[i] *= 2.5; } Marek Zając
OPENCL Marek Zając
double* tab = newdouble[100000000]; //preparearray... //... … Marek Zając
… cl_interr; cl::vector<cl::Platform> platformList; cl::Platform::get(&platformList); … Marek Zając
… cl_context_propertiescprops[3] = { CL_CONTEXT_PLATFORM, (cl_context_properties) (platformList[0])(), 0 }; cl::Context context(CL_DEVICE_TYPE_ALL, cprops, NULL, NULL, &err); cl::vector<cl::Device> devices; devices = context.getInfo<CL_CONTEXT_DEVICES>(); … Marek Zając
… std::ifstream file("func.cl"); std::string prog(std::istreambuf_iterator<char>(file), (std::istreambuf_iterator<char>())); cl::Program::Sources source(1, std::make_pair(prog.c_str(), prog.length() + 1)); cl::Program program(context, source); err = program.build(devices, ""); cl::Kernel kernel(program, "func", &err); … Marek Zając
… JUŻ PRAWIE, JESZCZE TYLKO … Marek Zając
… cl::Buffer in(context, CL_MEM_READ_WRITE, sizeof(double) * 100000000); cl::CommandQueue queue(context, devices[0], 0, NULL); constsize_tlocal_ws = 256; constsize_tglobal_ws = ceil((double)local_ws / 100000000); cl::Eventevent; … Marek Zając
… queue.enqueueWriteBuffer(in, CL_TRUE, 0, sizeof(double) * 100000000, tab); kernel.setArg(0, in); kernel.setArg(1, 100000000); queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(global_ws), cl::NDRange(local_ws), NULL, &event); event.wait(); queue.enqueueReadBuffer(in, CL_TRUE, 0, sizeof(double) * 100000000, tab); Marek Zając
CZYŻBY KONIEC? Marek Zając
__kernelvoidfunc(__globaldouble* in, int n) { size_tid = get_global_id(0); in[id] *= 2.5; } Marek Zając
TO MOŻE JAKIŚ TEST? „ŻĄDAMY DOWODÓW!” Marek Zając
__kernelvoidfunc(__globaldouble* in, double x, int n) { size_tid = get_global_id(0); if(id < n) { in[id] *= x; in[id] += 10.0 * x; in[id] = sqrt(in[i]); } } Marek Zając
K = 10000000 K * 200 Marek Zając
WYNIK TO… Marek Zając
CPU 11,731 sek. GPU 6,068 sek. Marek Zając
double => float CPU 12,963 sek. GPU 2,979 sek. Marek Zając
DZIĘKUJĘ ZA UWAGĘ ŹRÓDŁA OBRAZKÓW: 99% KRADZIONE Z INTERNETU, 1% WŁAŚNE Marek Zając