Regces: Fill & Download for Free

Shared memory is used just like any other array but with a __shared__ in front of array. It does not synchronize data between threads in its thread block. Explicit synchronization is needed to see updated values, even within same thread.// declare __shared__ int sharedData[n];  // initialize using all threads sharedData[threadIdx.x]=someGlobalArray[index];   // reduce data on it for(int i=blockDim.x/2;i>0;i/=2) {  if(threadIdx.x < i)  sharedData[threadIdx.x] += sharedData[threadIdx.x + i];   // synchronize between threads  __syncthreads(); }  if(threadIdx.x == 0)  result[index] = sharedData[0];  It can be used as a coalescing layer between global memory and registers, it can also be used as a fast addressable memory. Coalesced global memory access is important for performance. For this, thread block’s threads do uniformly increasing index access between them on global memory and copy data to shared memory, then work on shared memory as random as it needs without limitations of global memory access patterns. Total shared memory bandwidth is on the order of terabytes per second so it yields high performance if data on it is re-used many times. It is also a fast way to share data between block threads as seen in above reduction example.There is instruction level parallelism in GPUs. This lets them move on next instruction if an instruction is busy doing something (like fetching data from shared memory) and if there is no dependency between them so synchronization commands should be used as less as possible because they hold execution of threads until all threads hit that barrier. If there are multiple independent shared memory operations, they should be grouped behind a common barrier command.Example for coalescing layer:matrix multiplicationloading square sub-matrices A and B as coalesced into shared memoryreading global memory line by line instead of element by elementdoing C += A x B in random-access from shared memoryreading element by element with not a bad performance hitstoring result C matrix into shared memorywriting element by elementstoring C from shared to global memorywriting line by line, coalescedmatrix transposeload matrix line by line to shared memorydo the transpose using element-wise operationsstore result matrix line by line from shared memory to global memorysortload sub-array into shared memory, coalescedsort, element-wise access on shared memorystore sub-array into global memory from shared memory, coalescedExample for fast memory:n-body algorithmload 1 particle into a registerload K particles into shared memoryiterate over all particles in shared memoryaccumulate force on a registertotal re-use per particle = number of threads iterating on shared dataimage filterload pixel into registerload K pixels into shared memoryiterate all neighbor pixels in shared memoryre-use ratio per pixel = K, on shared memorytree traversalload nodes into shared memorygo up or down on edgesre-use ratio per node = number of threads passing through itCoalescing layer does not fit every algorithm but it does help on some algorithms like matrix multiplication, even if you don’t use shared memory as a data re-use part(like using register-tiling only).Here is a working example kernel for matrix multiplication that achieves 500+ GFLOPS on a GT1030:#define _N_ 1024 #define _USE_Z_ORDER_INDEXING_ 1 const int block = 4; // size of register tiling const int dv = 32; // size of shared memory tiling  __global__ void multiplyBlock(const float * __restrict__ a, const float * __restrict__ b, float * __restrict__ c, const int * __restrict__ zorder) {     #ifdef _USE_Z_ORDER_INDEXING_  /* zorder indexing */  const int groupId0 = blockIdx.x + blockIdx.y * (_N_/dv);  const int groupIdZorder = loadUncached(&zorder[groupId0]);  const int blockIdxxZorder = groupIdZorder % (_N_/dv);  const int blockIdxyZorder = groupIdZorder / (_N_/dv);  const int idx = threadIdx.x + blockIdxxZorder * blockDim.x;  const int idy = threadIdx.y + blockIdxyZorder * blockDim.y;  const int bidx = blockIdxxZorder;  const int bidy = blockIdxyZorder; #else   /* normal indexing */  const int idx = threadIdx.x + blockIdx.x * blockDim.x;  const int idy = threadIdx.y + blockIdx.y * blockDim.y;  const int bidx = blockIdx.x;  const int bidy = blockIdx.y; #endif   const int indexX = bidx * dv;  const int indexY = bidy * dv;  const int localId = threadIdx.x + threadIdx.y*blockDim.x;  const int ratio = (dv/block);  bool dualLoad = ((blockDim.x * blockDim.y) == (2*dv));  float regA[block*block+1];  float regB[block*block+1];  float regC[block*block+1];   __shared__ float smA[1 + dv][1 + dv];  __shared__ float smB[1 + dv][1 + dv];    for(int i=0;i<block;i++)  for(int j=0;j<block;j++)  regC[i*block+j]=0.0f;  __syncwarp();  const int nStep = _N_ / dv;    // iterate all sub matrices A,B (dv x dv size)   for(int i=0;i<nStep;i++)  {  const int iBlock = i*dv;  const int iBlockNIndexXlocalId = iBlock*_N_+indexX+localId;  const int iBlockIndexYNlocalId = iBlock+indexY*_N_+localId;    // load matrix A,B into shared memory  // coalescing layer  if((!dualLoad) && (localId < dv))  {  #pragma unroll  for(int z =0; z< dv; z++)  {  smA[localId][z] = a[iBlockIndexYNlocalId + z*_N_];  smB[localId][z] = b[iBlockNIndexXlocalId + z*_N_];  }  }  else if(dualLoad)  {  if(localId<dv)  {  #pragma unroll  for(int z =0; z< dv; z++)  smA[localId][z] = a[iBlockIndexYNlocalId + z*_N_];  }   if(localId>=dv)  {  #pragma unroll  for(int z =0; z< dv; z++)  smB[localId-dv][z] = b[iBlockNIndexXlocalId-dv + z*_N_];  }  }  __syncthreads();      for(int j=0;j<ratio;j++)  {  const int jBlock = j*block;   // load sub-sub matrices X,Y (block x block size) from A,B   // data re-use layer by shared memory  #pragma unroll  for(int y=0;y<block;y++)  {  #pragma unroll  for(int x=0;x<block;x++)  {  regA[x + y*block] = smA[jBlock + x][threadIdx.y*block + y];  regB[x + y*block] = smB[threadIdx.x*block + x][y+jBlock];  }  }    //C += X*Y    regC[0]+=regA[0]*regB[0]; regC[0]+=regA[1]*regB[4]; regC[0]+=regA[2]*regB[8]; regC[0]+=regA[3]*regB[12]; regC[1]+=regA[0]*regB[1]; regC[1]+=regA[1]*regB[5]; regC[1]+=regA[2]*regB[9]; regC[1]+=regA[3]*regB[13]; regC[2]+=regA[0]*regB[2]; regC[2]+=regA[1]*regB[6]; regC[2]+=regA[2]*regB[10]; regC[2]+=regA[3]*regB[14]; regC[3]+=regA[0]*regB[3]; regC[3]+=regA[1]*regB[7]; regC[3]+=regA[2]*regB[11]; regC[3]+=regA[3]*regB[15]; regC[4]+=regA[4]*regB[0]; regC[4]+=regA[5]*regB[4]; regC[4]+=regA[6]*regB[8]; regC[4]+=regA[7]*regB[12]; regC[5]+=regA[4]*regB[1]; regC[5]+=regA[5]*regB[5]; regC[5]+=regA[6]*regB[9]; regC[5]+=regA[7]*regB[13]; regC[6]+=regA[4]*regB[2]; regC[6]+=regA[5]*regB[6]; regC[6]+=regA[6]*regB[10]; regC[6]+=regA[7]*regB[14]; regC[7]+=regA[4]*regB[3]; regC[7]+=regA[5]*regB[7]; regC[7]+=regA[6]*regB[11]; regC[7]+=regA[7]*regB[15]; regC[8]+=regA[8]*regB[0]; regC[8]+=regA[9]*regB[4]; regC[8]+=regA[10]*regB[8]; regC[8]+=regA[11]*regB[12]; regC[9]+=regA[8]*regB[1]; regC[9]+=regA[9]*regB[5]; regC[9]+=regA[10]*regB[9]; regC[9]+=regA[11]*regB[13]; regC[10]+=regA[8]*regB[2]; regC[10]+=regA[9]*regB[6]; regC[10]+=regA[10]*regB[10]; regC[10]+=regA[11]*regB[14]; regC[11]+=regA[8]*regB[3]; regC[11]+=regA[9]*regB[7]; regC[11]+=regA[10]*regB[11]; regC[11]+=regA[11]*regB[15]; regC[12]+=regA[12]*regB[0]; regC[12]+=regA[13]*regB[4]; regC[12]+=regA[14]*regB[8]; regC[12]+=regA[15]*regB[12]; regC[13]+=regA[12]*regB[1]; regC[13]+=regA[13]*regB[5]; regC[13]+=regA[14]*regB[9]; regC[13]+=regA[15]*regB[13]; regC[14]+=regA[12]*regB[2]; regC[14]+=regA[13]*regB[6]; regC[14]+=regA[14]*regB[10]; regC[14]+=regA[15]*regB[14]; regC[15]+=regA[12]*regB[3]; regC[15]+=regA[13]*regB[7]; regC[15]+=regA[14]*regB[11]; regC[15]+=regA[15]*regB[15];    }   __syncthreads();  }   // preparing coalescing layer for result-writes, shared mem  #pragma unroll  for(int j=0;j<ratio;j++)  {  #pragma unroll  for(int y =0; y<block ; y++)  {  #pragma unroll  for(int x =0; x<block ; x++)  {  smA[x + threadIdx.x*block][threadIdx.y*block + y]=regC[x + y*block];  }  }  }  __syncthreads();   if((!dualLoad) && (localId < dv))  {  #pragma unroll  for(int z =0; z< dv; z++)  {  c[indexY * _N_ + indexX + localId + z*_N_]=smA[localId][z];  }  }  else if(dualLoad)  {  if(localId<dv)  {  #pragma unroll  for(int z =0; z< dv; z+=2)  {  c[indexY * _N_ + indexX + localId + z*_N_]=smA[localId][z];  }  }    if(localId>=dv)  {  #pragma unroll  for(int z =1; z< dv; z+=2)  {  c[indexY * _N_ + indexX + localId-dv + z*_N_]=smA[localId-dv][z];  }  }   }   } This uses shared memory as a coalescing layer (global data load/store) and a data re-use source (sub-sub-matrix data load/store).It is called like this: dim3 dimBlock(dv/block,dv/block);  dim3 dimGrid(_N_/dv,_N_/dv);     multiplyBlock<<<dimGrid,dimBlock>>>(a,b,c,zorder);  cudaDeviceSynchronize(); “N” is square matrix size (this only works for square matrices).“block” is register tiling size.“dv” is shared memory tiling size. This is where shared memory data is re-used whenever register tiling needs to load sub-sub-matrices from sub-matrix selected.Zorder array has fully recursive z-order indices generated from a function like this (a quick stackoverflow plagiarism):// N/dv is given as size&N here so it z-orders threadblocks, not threads void zorder(const int * source, int y0, int x0, int size,  int * target, const int N, int * counter) {  if (size == 1) {  target[(*counter)++]=source[y0*N+x0];  } else {   int h = size/2;  zorder(source, y0, x0, h, target,N,counter);  zorder(source, y0, x0+h, h, target,N,counter);  zorder(source, y0+h, x0, h, target,N,counter);  zorder(source, y0+h, x0+h, h, target,N,counter);  } } // to let index values not disrupt L1 contents __device__ __forceinline__ int loadUncached(const int *addr) {  int result;  asm("ld.global.cg.s32 %0, [%1];" : "=r"(result) : "l"(addr));  return result; } used in z-ordering on block-level data (not thread-level) to better use unified texture/L1 cache of GT1030.Nvidia’s visual profiler gives below data for that matrix multiplication kernel at 1024x1024 size:Shared memory is a bottleneck when register tiling is not big enough.1721 MHz GPU with 384 CUDA pipelines = 1321 GFLOPS peak.1024 size matrix multiply in 4.2 milliseconds = 511 GFLOPS achieved (38% of peak)1024 size matrix multiply in 4.2 milliseconds = 2045 GB/s average bandwidth achievedJust using shared memory is not enough to achieve peak GFLOPS performance. It creates a memory dependency on bandwidth and stalls execution. Registers are needed. (This example used 4x4 tiles on registers, should be using 8x8 at least)Algorithms with unequal operator usage (+,*,/,..) this calculation does not apply and shared memory still be fast enough.

Officers on the full-career type of commission (the Regular Commission, RegC) have to retire by 60 or 35 years service, whichever comes first. Many officers are not on the full-career type of commission and serve for a far shorter period - eg a Short Service Commission is for 3 years (extendable to 8 years). Most officers overall serve for just 3 to 12 years.British Army - Recruiting Selection and Training - Officer Commissions - Armed Forces - a12a6Other Ranks usually have to retire after 22 years service, at the age of 40, although some can apply to serve for about another 3 years.

The Legal Officer in the British Army is just like any other commissioned officer — except that the L.O. is a qualified solicitor or barrister.Most people think the L.O.s are just prosecuting soldiers at courts martial. They’re more than that.They’re certainly more than Tom Cruise’s character in the 1992 movie A Few Good Men who reputedly cannot swim and had trouble buttoning up his service tunic.The L.O.s work in the Army Legal Services (ALS) branch of the Adjutant-General’s Corps (AGC), which is the army’s general administrative service.The ALS provides a range of crucial specialist military legal support — from service prosecutions to operational law (such as advising senior commanders on rules of engagement during overseas operations). The ALS also advise and train regular officers on discipline, administrative action and operational law.As a Legal Officer, this is the bird’s eye view:—Entry requirementsAge between 23 and 32 years old.If you’re over 26 years old, you can still apply online but it would be better to go through the Army recruitment centre.You must already be a qualified solicitor or barrister — you have completed your 2-year training contract or a 12-month pupillage.Qualified in the England and Wales, Scotland or Northern Ireland jurisdictions.Pay and rank progressionYou will join the Army as a captain.Starting minimum salary is £40,025 a year.30 days of paid annual leave, plus bank holidays.Extra time time on your return if you’ve been on operations.The Army pays all of the contributes towards your pension. If you serve more than 18 years, you’ll qualify for a tax-free lump sum from age 40.Most officers serve a “Short Service Commission” (3–12 years). You can apply to extend the service through through IRC (Intermediate Regular Commission) or RegC (Regular Commission). RegC allows service up to age 60.The maximum achievable rank for a Legal Officer is Major-General.Training for the roleAs a professionally qualified officer candidate, you will sit through an extra Selection Board to provide your specialist skills.You’ll spend 9 months of ordinary Army training first. During the first two weeks of your basic training, you will also do initial training and administration at the Directorate of Army Legal Services.After the above, you then go to the Royal Military Academy Sandhurst to do the short course Professionally Qualified Officers Commissioning Course (10 weeks). After this, you joint your regiment, corps or unit (ALS). See Note below.Post-training qualificationsWhile working in prosecutions at the Service Prosecutions Authority (SPA), you can obtain an LL.M. Higher Rights of Audience (Criminal Courts).Note about SandhurstIn simple terms, your training at Sandhurst counts towards the opportunity to achieve either a Bachelor in Science or a Master in Science in Leadership and Strategic Studies over 3–5 years of commissioned service.As a PQO at Sandhurst, your commissioning course is 10 weeks in duration and designed to be shorter than the one for regular officers. This short course still gives you a military training, teaching you core infantry skills such as drill, fieldcraft, officership and command.For comparison, regular officers do a longer commissioning course:—Junior Term (14 weeks) — Basic military skills and fundamentals of military leadership, focusing on fieldcraft and tactics. This includes two weeks on Adventurous Training (earning Adventurous Training Qualifications in caving, skiing, paragliding, diving, sailing, mountaineering, etc).Intermediate Term (14 weeks) — You transition from self-survival to taking responsibility for others. You will be trained how to operate up to company level. At the end of term, you go on two weeks’ adventurous training.Senior Term (14 weeks) — You will be trained in increasingly complex scenarios to refine your military skills, go on three more field exercise (one overseas) before final commissioning.If you’re over 26 or 30 years old (I can’t remember which), you may be put on the Late Entry Officer Course (LEOC) at Sandhurst.Final wordThe age band for candidates (23–32 years old) already tells you that the Army isn’t exactly looking for second-rate lawyers.If a person at age 18 starts Year 1 of the law degree to be a solicitor, he/she should be qualified normally by 25–26. That puts the person in the ‘older’ bracket of officer candidates, so the Army is looking for the ‘standard’ crop of intelligent and ‘performing’ type of lawyers.