parallel processing - OpenCL barrier of finding max in a block -


i've found piece of opencl kernel sample code in nvidia's developer site purpose function maxoneblock find out biggest value of array maxvalue , store maxvalue[0].

i understand looping part, confused unroll part: why unroll part not need sync thread after each step done?

e.g: when 1 thread done comparison of localid , localid+32, how ensure other thread have stored result localid+16?

the kernel code:

void maxoneblock(__local float maxvalue[],                  __local int   maxind[]) {     uint localid   = get_local_id(0);     uint localsize = get_local_size(0);     int idx;     float m1, m2, m3;      (uint s = localsize/2; s > 32; s >>= 1)     {         if (localid < s)          {             m1 = maxvalue[localid];             m2 = maxvalue[localid+s];             m3 = (m1 >= m2) ? m1 : m2;             idx = (m1 >= m2) ? localid : localid + s;             maxvalue[localid] = m3;             maxind[localid] = maxind[idx];         }         barrier(clk_local_mem_fence);     }      // unroll final warp reduce loop , sync overheads     if (localid < 32)     {         m1 = maxvalue[localid];         m2 = maxvalue[localid+32];         m3 = (m1 > m2) ? m1 : m2;         idx = (m1 > m2) ? localid : localid + 32;         maxvalue[localid] = m3;         maxind[localid] = maxind[idx];           m1 = maxvalue[localid];         m2 = maxvalue[localid+16];         m3 = (m1 > m2) ? m1 : m2;         idx = (m1 > m2) ? localid : localid + 16;         maxvalue[localid] = m3;         maxind[localid] = maxind[idx];          m1 = maxvalue[localid];         m2 = maxvalue[localid+8];         m3 = (m1 > m2) ? m1 : m2;         idx = (m1 > m2) ? localid : localid + 8;         maxvalue[localid] = m3;         maxind[localid] = maxind[idx];          m1 = maxvalue[localid];         m2 = maxvalue[localid+4];         m3 = (m1 > m2) ? m1 : m2;         idx = (m1 > m2) ? localid : localid + 4;         maxvalue[localid] = m3;         maxind[localid] = maxind[idx];          m1 = maxvalue[localid];         m2 = maxvalue[localid+2];         m3 = (m1 > m2) ? m1 : m2;         idx = (m1 > m2) ? localid : localid + 2;         maxvalue[localid] = m3;         maxind[localid] = maxind[idx];          m1 = maxvalue[localid];         m2 = maxvalue[localid+1];         m3 = (m1 > m2) ? m1 : m2;         idx = (m1 > m2) ? localid : localid + 1;         maxvalue[localid] = m3;         maxind[localid] = maxind[idx];     } }

why unroll part not need sync thread after each step done?

the sample incorrect, barrier indeed required after each step.

it looks sample written in warp-synchronous style, way of exploiting underlying execution mechanism of warps on nvidia hardware, incorrect synchronization cause break if underlying execution mechanism changes or in presence of compiler optimizations.


Comments

Post a Comment

Popular posts from this blog

angularjs - ADAL JS Angular- WebAPI add a new role claim to the token -

i have created angular spa webapi backend using adal js authentication.as there no roles in ad, need manually add role claims in order give users access different api controllers. the roles stored in database. expecting inject claim through call webapi after authentication ad. webapi code might this. identity.addclaim(new claim("role", "user")); var ticket = new authenticationticket(identity, props); var accesstoken = startup.oauthbeareroptions.accesstokenformat.protect(ticket); is possible replace adal idtoken new token? is viable solution or there other better way handle this? as initial token generated azuread, possible edit token add new claim? appreciated. graph api supporting group claims. see here: http://justazure.com/azure-active-directory-part-4-group-claims/ if @ examples on page, users assigned groups , app can check group in claims. in current version of portal, need app manifest , modify it.
Read more

node.js - Using Node without global install -

i'm working on few node & electron (atom shell) projects , i'm curious how them work without user having install node globally. i've seen few applications including node directory in project contains node.exe binary. how go setting node project use binary? edit: thought should add, i've searched everywhere answer this, unrelated results.
Read more

php - CakePHP HttpSockets send array of paramms -

i problem when make request (i can request). $link = 'http://uri/path/?param1=1&sid=2&sid=3&sid=4' $http = $httpsocket = new httpsocket(); $res = $httpsocket->get($link); if check request pr($http); i see : [line] => /url/path/?param1=1&sid%5b0%5d=2&sid%5b1%5d=3&sid%5b2%5d=4 http/1.1 and i'm getting empty response because server doesn't know parameters sid%5b0%5d=2 when try send parameters array : $data = array('param1' => '1', 'sid' => array('2','3', '4')); i see same changes %5b0%5d - additional indexes. how fix it? unfortunately , can not change server side, if send request in browser, i'll normal response in json format. when make request, can parameters this: $param1 = $_get['param1']; $sid= $_get['sid']; tip: use diferente names parameters. have multiples sid. use sid1, sid2, sid3... update: can parameters this:...
Read more