A quick guide to writing OpenCL kernels for PowerVR Rogue GPUs (by Doug Watt)

PowerVR Rogue GPU에서  OpenCL 커널 작성을 위한 간단한 가이드라인

(PowerVR 사의 Rogue GPU에 특화된 내용이 많지만, OpenCL과 관련된 fundamental한 내용들이 좋아서 번역하였다)

This article is written in Korean which translates an article written by Doug Watt. 

source : http://blog.imgtec.com/powervr/a-quick-guide-to-writing-opencl-kernels-for-rogue

This provides the background to understand the programming guidelines for the Rogue architecturewhich are illustrated by using a case study of an image filtering program.

이 문서는 이미지 필터링 프로그램의 적용 사례를 설명한 “Rogue 아키텍처(GPU)에서의 프로그래밍 가이드라인” 을 이해하기 위한 배경 지식을 설명하고 있다.

OpenCL overview

Consider a simple C program that multiplies the corresponding elements in two 2-dimensional arrays:

다음과 같이 2차원 배열의 값들을 곱하는 간단한 C 프로그램이 있다

void matrix_mul_cpu(const float *a, const float *b, float *c, int l, int h)

{

   int i;

   for (i=0; i<l; i++)

   for (j=0; j<h; j++)

   c[i][j] = a[i][j] * b[i][j];

}

The for loop consists of two parts: the vector range (two dimensions containing l*h elements), and a multiplication operation. All loop iterations are independent in that there are no data-dependencies between any iterations of the loops. This means that the multiplication operations can be executed in parallel on the GPU’s numerous threads.

For 루프문은 이중으로 구성되어 있다 벡터 영역과 곱 연산. 모든 루프의 과정은 각각의 루프 내에서 데이터들 사이에 의존성이 없다. 이는 곱 연산이 GPU의 수많은 스레드들 위에서 병렬로 처리가 가능하다는 뜻이다.

OpenCL programs consist of two parts: one that runs on the host CPU, the other that runs on the GPU device. To execute code on the GPU, you first define an OpenCL kernel, which is written in a variant of C with some additional keywords and data types. Continuing with the above example, an OpenCL kernel that implements a single multiplication is as follows:

OpenCL 프로그램들은 두 개의 파트로 이루어져 있다. 하나는 host CPU 에서 동작하고, 다른 하나는 GPU 디바이스에서 동작한다.(사실 GPU 뿐만 아니라 다른 병렬 처리를 위한 벡터 디바이스를 모두 지원하는걸 목표로 한다. 역자주) 코드를 GPU 위에서 실행하기 위해 가장 먼저 할일은 OpenCL 커널을 정의하는 것이다. 이 커널은 C 언어를 확장한 것으로 여러 키워드들과 데이터 타입이 추가가 된 언어이다(OpenCL C 언어를 말한다. 역자주) 위의 예제로 돌아가서 OpenCL 커널은 아래의 코드처럼 단 한 줄의 곱 연산으로만 구현된다.

_kernel void vector_mul(__global const float *a, __global const float *b, __global float *c)

{

  size_t i = get_global_id(0);

  size_t j = get_global_id(1);

  c[i][j] = a[i][j] * b[i][j];

}

The host program launches (or enqueues) this kernel, spawning a virtual grid of concurrent work-items as illustrated below for a 2D grid of size 512 (32×16). The grid is an N-dimensional range where N=1, 2 or 3 (or NDRange). A single work-item executes for each point on the NDRange; each work-item executes the same kernel code but has a unique ID that can be queried and used to determine the data to be processed. This ID is obtained using the built-in function get_global_id.

호스트 프로그램은 이 커널을 실행(또는 enqueuer 한다)하고, 아래 그림에 표현된 것처럼 512(32x16) 크기의 2D grid 처럼 동시 다발적인 work-item의 가상 그리드를 생성한다. 이 grid는 N-차원 영역 이다.(N= 1, 2, 또는 3, (또는 NDRange)). 하나의 단일 work-item은 ND Range 위의 각각의 지점을 수행한다. 각각의 work-item은 동일한 커널 코드를 수행하지만, 연산을 해야할 대상 데이터를 지정하는데 사용하는 특정한 ID를 갖는다. 이 ID는 get_global_id() 라는 내장 함수에서 사용된다.

In OpenCL, multiple work-items are grouped together to form workgroups. In the figure above, each workgroup size is 8×4 comprising a total of 32 work-items. Work-items in a workgroup can synchronize with one another and share data using local memory (to be explained in a later article).

OpenCL 에서는 여러 개의 work-item 들이 하나의 workgroup을 생성한다. 위의 그림에서 각각의 workgroup 사이즈는 8x4 이고 총 32개의 work-item들을 묶어놓은 것이다. 하나의 workgroup 안의 work-item 들은 각각의 다른 work-item 들과 동기화를 하는데 이때 데이터를 공유하는데 사용되는 것이 local memory이다.(추후 설명하겠다)

OpenCL execution on the PowerVR Rogue architecture

PowerVR Rogue 아키텍쳐(GPU)에서 OpenCL 실행하기

As shown below, from an OpenCL perspective, PowerVR GPUs are built around scalable arrays of multithreaded processors called Unified Shading Clusters (USCs). When a program running on the CPU enqueues an OpenCL kernel, all work-items in the NDRange are enumerated. The workgroup IDs and work-item IDs are enqueued sequentially in row-major order. The Coarse Grain Scheduler assembles the work-items into groups of up to 32 threads, termed warps, and streams these warps to the multiprocessors with available execution capacity. The warps containing the work-items for a complete workgroup are always allocated to a single multiprocessor.

아래 그림에 보여지듯이 OpenCL 관점에서는 PowerVR GPU들은 확장 가능한 멀티 스레드 프로세서들의 연속이다. (Unified Shading Clusters 라고 부른다. USCs)

하나의 프로그램이 CPU위에서 동작하고 OpenCL 커널을 enqueue 하였을 때, NDRange 내부의 모든 work-item들이 쭉 나열된다. Workgroup ID 들과 work-item ID들은 row 방향부터 순차적으로  enqueue 된다. Coarse Grain Scheduler라고 불리는 스케줄러는 work-item들을 최대 32 개의 스레드(termed warps) 그룹 안에 모으고, 이 warp 들을 실행 가능한 멀티 프로세서들에 보낸다. 이 warp 들은 완전한 하나의 workgroup에 해당하는 work-item들을 포함하고 있으며, 항상 하나의 멀티프로세서에 할당된다.

The multiprocessors execute the warps obtained from the Coarse Grain Scheduler, with the assistance of one or more Texture Processing Unit (TPU), L1 and L2 cache units. The precise number and grouping of these hardware blocks is product-specific.

멀티프로세서들은 Coarse Grain Scheduler 로부터 포함된 warp들을 수행한다. 이때 Texture Processing Unit(TPU), L1, L2 캐시 유닛들이 이를 지원한다. 이러한 하드웨어 그룹과 정밀도 숫자들은 제품마다 다르다.

As shown below, each multiprocessor contains a number residency slots that at any time are either empty (illustrated as a grid with a dotted line) or occupied by a warp (illustrated as a grid with a solid line and colour fill). Each warp can be in one of three states: active (green fill) which means it is currently running on the execution unit; ready (orange fill) which means the scheduler may run the warp on the execution unit after executing the active warp; and blocked (red fill) which means that one or more of the work-items in the warp is awaiting a memory or barrier operation to complete. In this example, the multiprocessor has six warps to schedule: the first is active, four are ready to be executed and a single warp is blocked on a memory or barrier operation.

아래 그림에서 보여지듯, 각각의 멀티프로세서들은 수많은 residency slot들을 포함하며(그림에서 점으로 표시된 영역) 이 슬롯들은 때로는 비어 있거나, warp에 의해 선점 당한다.(그림에서 색깔로 칠해진 영역). 각각의 warp는 3가지의 상태를 가질수 있다. Active(초록)는 현재 execution unit(실행 유닛)위에서 동작 중. Ready(오렌지색)은 active warp의 실행이 끝난 다음에 스케줄러가 execution unit 위에서 실행 시킬 수 있는 대기 상태. Blocked(빨강)는 warp 안의 하나 또는 그 이상의 work-item들이 작업 완료를 위해 메모리 연산이나 barrier 연산을 기다리고 있는 상태이다.

이 그림의 예시에서, 멀티프로세서들은 스케줄링 할 수 있는 여섯 개의 warp들을 갖는다. 첫 번째는 active, 네 개는 ready, 하나는 block된 상태이다.

A multiprocessor can execute at most one warp on its execution unit at any one time. The work-items in the active warp are usually executed together in lock-step parallel as shown below, and continue executing either until completion or until the work-items become blocked on a memory or barrier operation. In this example, all 32 work-items in the active warp execute the first kernel statement together and then all progress on to the second kernel statement together.

하나의 멀티프로세서는 프로세서 내의 execution unit 위에 한번에 단지 하나의 warp를 실행할 수 있다. active warp 안의 Work-item 들은 아래 그림에서처럼 lock-step 안에서 함께 실행된다. 그리고 work-item 들이 메모리 연산이나 barrier 연산에 의해 block되거나, 작업이 끝날 때까지 실행 된다. 아래 그림 예시에서 active warp 안의 모든 32개의 work-item 들은 첫 번째 커널 statement를 함께 수행한 후, 계속해서 두 번째 statement를 수행한다.

Full lock-step parallelism is achieved when all threads in a warp follow the same control path. If the threads diverge via a conditional branch, as illustrated in the figure below, the work-items stop executing together lock-step and instead the hardware serially executes each branch path taken. In this example, all work-items in the warp follow separate paths (work-item 0 executes case 0, work-item 1 executes case 1 and so on), resulting in a 32-way branch divergence. A more common scenario is an if-else statement whereby some of the work-items follow the if-statement and the other work-items follow the else-statement, resulting in a two-way divergence. In this case the compiler and hardware can minimise the impact of the divergence by using hardware predication. The compiler translates the code sequence whereby the branch condition is first calculated (true orfalse) and then both the if and else target instructions are executed in sequence using the branch condition as an instruction predicate. In this way, all of the instructions are executed, but only the instructions with a true predicate write results.

Full lock-step 병렬화는 하나의 warp 안에 모든 스레드들이 같은 연산 로직을 따를 때 가능하다. 만약 스레드들이 조건문에 의해 분기가 되면(아래 그림 참조), work-item들은 lock-step을 수행하고 멈춘다. 그리고 나서 하드웨어가 순서대로 각각의 branch를 수행한다. 이 예시에서,  warp 안의 모든 work-item들은 별도의 path(work-item 0는 case 0을 수행, work-item1은 case 1을 수행)를 수행한다. 그리고 그 결과들은 32 가지의 다양한 결과를 내놓는다.  더 일반적인 시나리오는 if-else 문인데 몇몇 work-item들이 if문을 따르고, 또 다른 몇몇 work-item들은 else문을 수행할 때, 두 가지의 분기를 생성한다.  이 때엔 컴파일러와 하드웨어는 특정 하드웨어 명령(hardware predication)을 이용해 분기의 영향을 최소화 하는 최적화 작업을 한다. 컴파일러가 분기문 영역의 코드 시퀀스를 먼저 계산하고(true or false), 그 다음 if 와 else 모두 instruction 셋에 의해 분기문을 순차적으로 실행한다. 이 방법으로 모든 instruction들이 수행되지만 true 에 해당하는 instruction만 결과를 내놓는다.

Whenever a warp is descheduled, for example due to completion, a memory or barrier operation, the multiprocessor selects another resident warp that is ready to execute. In the figure below, the first warp has reached a statement that reads from an array allocated to system memory (mem), preventing it from continuing until some point in the future when the data has been fetched; the multiprocessor is able to continue performing useful work during this time by selecting one of the other four ready warps to execute.

하나의 warp가 스케줄링이 되지 않을 때마다(예를 들면 연산이 종료되거나, 메모리 접근 연산 또는 barrier 연산에 의해), 멀티프로세서는 실행 가능한 다른 잔여 warp를 선택한다. 아래 그림에서, 첫 번째 warp는 system memory에 할당된 배열로부터 데이터를 읽는 statement에 도달하였고, 나중에 데이터 fetch가 완료될 때까지 더 이상의 진행을 멈춘다. 이 시간 동안 멀티프로세서는 다른 네 개의 ready 상태의 warp를 선택해서 작업을 계속 수행할 수 있다.

The execution state for each work-item (and, therefore, each warp) is maintained on-chip for the entire lifetime of the work-item in the unified store; shared local memory is held in the common-store(both shown in above). As a result, a multiprocessor can context switch between warps with zero cost. The figure below shows an example timeline for a multiprocessor scheduling between four warps. The scheduler starts executing warp 0, which runs until it reaches a blocking operation such a read or write to system memory. At this time the scheduler de-schedules warp 0 and starts executing warp 1. The work-items in all warps implement the same kernel, and therefore have similar performance characteristics, reflected in the similar (phase-adjusted) timelines in the example. At a time around half way into execution of warp 2, the operation that warp 0 was blocked on completes, so that warp 0 returns to a ready state. Later, after the scheduler has executed the first statements in warp 3, the scheduler re-schedules warp 0. The effect of this concurrent scheduling is that the impact of memory latency when processing each warp is hidden. In this scenario, the system is able to completely hide the memory latency.

각각의 work-item의 실행 상태는 unified store 안에 work-item의 총 라이프 타임 동안에 한번에 관리가 된다. 공유된 local memory는 common-store에 묶인다(둘 다 위에서 보았다.. 뭔소리지?). 결론적으로, 하나의 멀티프로세서는 별도의 비용 없이 warp 사이에 context switch를 할 수 있다. 아래 그림은 네 개의 warp 들 간의 멀티프로세서 스케줄링 타임라인을 보여준다. 스케줄러를 보면 맨 처음 warp0를 수행하고 시스템 메모리를 읽거나 쓰는 연산을 수행하면서 block 될 때까지 스케줄링 된다. 이 때에 스케줄러는 warp0 스케줄링을 취소하고 warp1을 수행한다. 모든 warp들에 work-item들은 같은 커널을 구현하였다. 따라서 비슷한 성능 특성을 보여주고 이 예제에서 비슷한 타임라인을 보여준다. Warp 2의 실행 되는 구간의 약 절반 정도 되는 부분에서 warp 0의 연산의 block 이 끝났기 때문에 warp 0는 다시 ready 상태로 돌아간다.  나중에 스케줄러가 warp 3의 첫 번째 실행을 수행하고 마친 이후에, warp 0를 다시 스케줄링 한다. 이 동시다발적인 스케줄링 시나리오 그림은 사실 각각의 warp가 수행될 때 메모리 지연(memory latency)이 무시되었다 .이 시나리오의 시스템에서 메모리 지연(memory latency)은 완전히 무시한다고 가정한다.

The figure below shows a similar scenario but where the round-trip-time to memory is larger. In this case, by the time warp 3 blocks on its first memory operation, warp 0, 1 and 2 are all still waiting for their memory operations to complete. All warps are therefore all blocked at the same time, with a corresponding reducing in utilization and performance.

아래 그림에서는 비슷한 시나리오를 보여주지만 메모리가 더 큰 round-trip-time 의 예시이다. 여기에서 warp 3 가 첫 번째 메모리 연산에서 블락이 되었을 때, warp 0, 1, 2가 여전히 메모리 연산이 끝나길 기다리고 있다. 따라서 모든 warp들이 동시에 블록 되고, utilization과 성능이 모두 감소한다.

In this case, utilization could be improved if there were more than four resident warps on the multiprocessor. The total number of warps that a multiprocessor can maintain on-chip depends on factors such as the time it takes to execute each warp and the memory requirements of each warp. If a work-item allocates a large amount of private memory, the total number of work-items (and, therefore, warps) that can reside on a multiprocessor may be limited by the size of the unified store.

이 상황에서 만약 멀티 프로세서에 여분의 warp들이 더 존재한다면 utilization은 향상될 수 있다. 멀티프로세서가 하나의 칩에서 관리할 수 있는 Warp의 총 개수는 각각의 warp를 수행하는데 드는 시간이나, 메모리 요구사항 등에 달려 있다. 만약 work-item이 아주 큰 private 메모리를 할당하면, 멀티프로세서 위에 남아 있을 수 있는 총 work-item 수는 unified store의 크기에 제한될 수도 있다.

Similarly, if a workgroup allocates a large amount of shared local memory, the total number of workgroups that can reside on a multiprocessor may be limited by the size of the common store.

비슷한 예로, 만약 workgroup 이 큰 사이즈의 shared local memory를 할당하면, 멀티프로세서에 남아있는 총 workgroup의 수 역시 common store의 사이즈에 제한될 수 있다.

Programming guidelines

프로그래밍 가이드라인

RECOMMENDATION: If possible use a workgroup size of 32.

가능하면 workgroup 크기를 32개로 유지하라

If permitted by a kernel’s logical structuring, you are advised to use a workgroup size of 32. In this case, each warp contains a complete workgroup and all synchronization costs are free (they all happen together on the same cycle in lock-step). For larger workgroup sizes, each workgroup is implemented using multiple warps, and barrier synchronization requires the multiprocessor to perform context switches between these warps. For smaller workgroup sizes such as 4, 8 and 16, if the workgroup size is specified at compile-time, multiple workgroups are packed into a warp. For non-standard workgroup sizes, alignment restrictions on how the hardware can pack the workgroups into warps means that the warps are underpopulated with idle threads, reducing efficiency.

만약 커널에 로직 구조상 허용된다면, workgroup의 사이즈를 32개로 정하는 것을 권장한다. 이렇게 하면 각각의 warp들은 하나의 완전한 workgroup을 포함하며, 모든 warp들의 동기화 비용이 없다(lock-stop 내에서 같은 사이클에 동시에 수행된다). 더 큰 workgroup 크기를 정하면 각각의 workgroup이 다중 warp로 구현되며, 이럴 경우 barrier 동기화 연산은 이 warp들 간의 context switch를 수행하게 한다.  4, 8, 16과 같이 더 작은 workgroup 사이즈일 경우(workgroup 크기가 컴파일 시에 정해질 때) 다중 workgroup들이 하나의 warp에 묶여진다. 비 표준 workgroup 사이즈들을 사용하게 되면, 하드웨어가 workgroup을 warp 안에 묶을 수 있는 alignment 제한들이 생기게 되는데, 결국 warp의 수가 감소하게 되고 이는 warp들이 idle thread나 효율성의 감소를 야기한다.

RECOMMENDATION: Specify your workgroup size at compile-time.

workgroup 사이즈를 컴파일 시에 정하라

For a given kernel, you can specify the workgroup size for which it can be enqueued by specifying the following attribute.

하나의 커널에서 프로그래머는 아래와 같이 attribute를 정의함으로써 한번에 enqueue 되는 workgroup 사이즈를 정의할 수 있다.

__attribute__((reqd_work_group_size(X, Y, Z)))

The workgroup size is the product of X, Y and Z. Specifying a fixed workgroup size allows a number of additional compile-time optimizations that improve performance such as reducing the number of instructions required for barrier synchronization.

Workgroup 사이즈는 X, Y, Z에 의해 정해진다. 고정된 workgroup size를 미리 정의하게 되면 컴파일 시에 몇몇 최적화 효과를 얻을 수 있는데, 예를 들면 barrier 동기화 연산에 사용되는 명령어(instruction)셋의 수를 줄일 수가 있다.

RECOMMENDATION: For image processing tasks, declare your kernel parameters as image and sampler data types (as opposed to character arrays) to benefit from TPU acceleration.

이미지를 처리하는 태스크에서는 커널 파라미터를 image 와 sampler 데이터 타입으로 선언해라.(문자열 배열은 사용하지 말라). 이는 TPU 가속의 효과를 얻을 수 있다.

As discussed previously, the TPU provides efficient sampling of image data, image interpolation, border-pixel handling. Specifying border-pixel handling in the sampler avoids the need for special conditional code within the kernel to handle these edge conditions, reducing divergence. Using image parameters also enables your kernel to be efficiently incorporated as part of a larger heterogeneous program.

앞에서 언급 했듯이 TPU는 image interpolation, border-pixel handler 같은 효율적인 이미지 데이터의 sampling 효과를 제공한다. sampler에서 border-pixel handling 을 정의하면 커널 내에서 edge condition들을 처리하기 위한 특정한 조건문 코드 사용을 피할 수 있고, 분기를 줄일 수 있다. 이미지 파라미터를 사용하는 것 또한 커널이 더 큰 이기종 프로그램의 일부로써 효율적으로 포함되도록 할 수 있다.

RECOMMENDATION: Use float data types inside your kernel to maximize ALU throughput.

ALU 유닛의 대역폭을 최대한 활용하기 위해 커널 내에서 float 데이터 타입을 사용하라

On Rogue devices, each thread can perform floating-point operations on up to two datasets per cycle, or it can perform integer operations on a single dataset. In practice, there is usually sufficient instruction-level parallelism in the kernel code for the compiler to generate dual-issue floating-point instructions. To maximize arithmetic throughput, you should therefore perform arithmetic on float data types wherever possible.

Rogue GPU 에서는 각각의 스레드들이 한 사이클에 두 개의 데이터 셋으로 floating-point 연산을 수행한다. 하지만 정수 연산에서는 하나의 데이터 셋으로 수행한다. 실제로는 컴파일러가 Floating-point 명령어의 dual-issue를 만들기 위해 대부분 커널 내에서 충분한 명령어 레벨에서의 병렬화가 이루이지긴 한다. 따라서 산술 연산의 성능을 최대화 하기 위해 프로그래머는 가능한 float 데이터 타입의 연산을 하도록 커널을 작성해야 한다.

 Use the built-in function read_imagef to sample a pixel from an OpenCL image. This function uses hardware to fetch the pixel’s constituent values (either r/b/g or y/u/v values) and write these values into elements of the returned float4 vector.

OpenCL 이미지의 pixel을 접근할 때 read_imagef()내장 함수를 사용해라. 이 함수는 하드웨어를 사용해서 픽셀의 값들을 접근하고 float4 타입 벡터에 값을 쓴다.

RECOMMENDATION: Trade precision for speed wherever possible. Use the compiler flag
-cl-fast-relaxed-math to enable arithmetic optimizations that trade precision for speed.

가능하다면 속도를 위해 정밀도를 조절하라. 컴파일러 flag인 –cl-fast-relaxed-math 를 이용하면 정밀도를 조절할 수 있으며 산술 연산 최적화가 가능하다.

Depending on the precision requirements of the application, you can often improve speed by limiting the use of arithmetic operations with long execution times.
어플리케이션에서 요구되는 산술 정밀도(배정밀도 등)는 경우에 따라 다르므로 실행 시간이 오래 걸리는 산술 연산 사용에 제한을 둠으로써 속도 향상을 얻을 수 있다.

·         Use the -cl-fast-relaxed-math build option, which enables many aggressive compiler optimizations for floating-point arithmetic.

·         -cl-fast-relaxed-math 옵션을 사용하여 빌드를 하면, floating-point 산술 연산들에 대해 상당히 강력한 컴파일러 최적화가 이루어진다.

·         Use native_* and half_* math built-ins, which have lower precision but are faster than their un-prefixed variants.

·         native_* 와 half_* 로 시작하는 내장 함수들을 사용하면 낮은 정밀도를 사용하지만 훨씬 빠른 속도를 얻을 수 있다.

RECOMMENDATION: Minimize conditional code.

조건문을 최소화 하라

In general you should minimise conditional code, especially with regards to nested conditional statements. Flow control statements (if, switch, do, for, while) can reduce instruction throughput by causing threads to diverge in their execution paths requiring the hardware to serialize execution of work-items in a warp. The compiler attempts to reduce the effect of divergence using hardware predication.
일반적으로 프로그래머는 조건문을 최대한 적게 사용해야 한다. 특히 if-else 가 난무하는 nested 조건 문은 지양하라. Flow control statement들은(if, switch, do, for, while) 각각의 스레드들이 하나의 warp안에서 work-item 의 순차 실행을 위해 하드웨어에 요구하는 명령들이 많아지고 다양해 지기 때문에 명령어의 성능을 저하시킨다. 컴파일러는 hardware predication를 이용해 이 분기를 최적화 하려고 한다.

RECOMMENDATION: Avoid short kernels with low arithmetic intensity.

산술 연산이 적은 짧은 커널 작성을 피하라

In general, the smaller the ratio of the number of arithmetic to memory instructions (or arithmetic intensity), the higher the occupancy is required to hide the memory latency. If this ratio is 50, for example, then to hide a memory access latency of about 500 clock cycles about 10 warps are required. If the multiprocessor contains 16 residency slots, this translates to a required occupancy of around 63%.

일반적으로 메모리 연산(또는 산술 집중도. Arithmetic intensity라고 한다)에 대한 산술 연산의 비중이 작을수록 메모리 지연을 죽이기 위한 비용이 높아진다. 만약 이 비율이 50%라고 하면, 10개의 warp들에 대해 약 500 clock 사이클이 메모리 접근 지연을 처리하기 위해 요구된다. 만약 멀티프로세서가 16개의 잔여 slot을 포함하고 있다면 이는 약 63%의 점유율을 요구한다.

Barrier instructions can also force the multiprocessor to idle as more and more warps wait for other warps in the same workgroup to complete execution of instructions prior to the barrier. Having multiple resident workgroups in the multiprocessor can help reduce idling in this case, as warps from different workgroups do not wait for each other at barriers. To increase the number of resident workgroups per multiprocessor, consider reducing your workgroup size.Barrier 명령어 또한 barrier 명령 이전에 연산을 처리하기 위해 workgroup 내의 다른 warp를 기다리도록 강제할 수 있다. 이 멀티프로세서 안에 여러 개의 여분의  workgroup이 있다면 다른 workgroup의 warp들이 각각의 barrier 연산을 기다릴 필요 없기 때문에 idling을 줄일 수 있다. 멀티프로세서당 잔여 workgroup 수를 늘리기 위해서는 workgroup size를 줄여야 한다.