M3 및 A17 Pro의 GPU 개선 사항 알아보기

안녕하세요, 저는 Jedd Haberstro입니다 Apple의 GPU 그래픽, 디스플레이 소프트웨어 그룹의 엔지니어입니다

iPhone 15 Pro와 새로운 Mac 모델의 핵심인 A17 Pro 및 M3 칩 제품군의 새 Apple Family 9 GPU 아키텍처를 소개합니다

Apple의 제품 라인 전반에서 GPU는 많은 사용자들에게 사랑받는 다양하고 풍부한 경험을 구현합니다 새 iPhone 15 Pro로 이동 중에 게임을 즐길 때도

새 iMac에서 매끄러운 UI 애니메이션을 볼 때도

새 MacBook Pro에서 머신 러닝을 활용하여 고도의 비디오 및 이미지 처리 작업을 할 때도

GPU는 이러한 앱의 구동에 중요한 역할을 합니다

Metal API는 Apple GPU의 컴퓨팅 기능을 구동하는 데 사용됩니다 이러한 앱들은 전부 다양한 Metal 셰이딩 언어 프로그램을 실행합니다 이러한 셰이더 프로그램은 단 몇 줄의 코드를 실행하는 저용량의 간단한 셰이더부터 수많은 코드, 프레임워크 라이브러리를 사용하는 복잡한 대용량 셰이더까지 다양합니다 이 셰이더들의 공통점은 방대한 데이터 병렬성으로 이는 앱을 병렬로 실행하여 성능을 크게 향상할 수 있는 기회입니다 이러한 병렬성은 Metal 셰이더 프로그램을 다양한 입력을 통해 병렬로 반복 실행하는 데서 달성됩니다 3D 렌더링된 장면의 각 버텍스나 화면의 각 픽셀 등이 해당됩니다

궁극적으로, 이러한 셰이더를 병렬로 실행하는 건 GPU입니다

모든 GPU의 중심에는 셰이더 코어가 있습니다 각 셰이더 코어는 수천 개의 스레드를 병렬로 실행합니다 또한 성능을 더욱 확장하기 위해 GPU에는 병렬로 실행할 수 있는 많은 셰이더 코어가 있습니다 따라서 앱에 실행을 위한 수만 개의 병렬 스레드를 제공합니다

이미 기존 iPhone iPad, Mac에 있는 GPU는 놀라운 성능을 낼 수 있습니다

또한 앱 개발자들이 GPU의 잠재력을 극대화할 수 있게 해주는 강력한 도구도 있죠 하지만 새로운 Apple Family 9 GPU는 성능을 전례 없던 수준으로 끌어올리는 새로운 개선 사항을 선보입니다

첫 번째는 새로운 코어 아키텍처입니다 기존 앱의 성능과 전력 효율을 향상하여 앱이 사용자에게 제공하는 경험을 개선하며 차세대 앱을 개발할 때 발생하는 까다로운 과제를 해결합니다

하드웨어 가속 레이 트레이싱은 Metal의 레이 트레이싱 API를 사용하는 앱에 유용하며 풍부한 렌더링 효과와 탁월한 성능을 위한 레이 트레이싱 기능을 사용할 기회를 넓힙니다

하드웨어 가속 메시 셰이딩은 앱이 고도의 지오메트리 프로세싱 파이프라인을 빌드할 수 있게 해줍니다

이러한 내용을 자세히 살펴보기 전에 아무런 변경 사항 없이 새로운 Apple Family 9 GPU로 앱에서 실현 가능한 성능을 알아보겠습니다 화면의 게임은 Larian Studios의 ‘발더스 게이트 3’입니다 위 화면은 M3 시스템 MacBook Pro 아래는 M2 시스템 MacBook Pro입니다 각 화면은 울트라 비디오 품질(1800p)로 렌더링되었습니다 M3 Mac은 더 높은 스레드 점유율로 게임의 Metal 셰이더를 실행하는 차세대 셰이더 코어 덕분에 비약적으로 향상된 성능을 선보입니다 화면은 Blender가 Cycles Path Tracer를 사용하여 바버샵 장면을 렌더링하는 모습입니다 M3 Mac에서 Metal 레이 트레이싱을 활용하는 툴입니다 두 렌더링은 동시에 시작되었지만 하드웨어 가속 레이 트레이싱과 차세대 셰이더 코어 덕분에 M3 Mac의 렌더링이 훨씬 빠릅니다

이것은 ‘토이 스토리 4’ 골동품 가게 USD의 실시간 시각화입니다 Pixar의 Hydra Storm으로 렌더링했습니다 Hydra Storm은 M3 Mac에서 Metal 메시 셰이딩을 사용하는데 하드웨어 가속 메시 셰이딩과 합치면 그 어느 때보다도 빠릅니다

이제 각 기능들을 자세히 알아보겠습니다 먼저 차세대 셰이더 코어입니다

Apple Family 9 GPU에는 여러 구성 요소가 있습니다 컴퓨팅 및 버텍스 명령 프로세서는 Metal 명령 버퍼를 분석합니다 래스터라이저는 실행을 위해 프래그먼트 셰이더를 디스패치합니다

캐시 계층도 있습니다 GPU LLC(Last Level Cache)는 모든 GPU 메모리 트래픽을 서비스합니다

그러나 모든 GPU의 핵심은 셰이더 코어입니다 셰이더 코어는 앱의 Metal 셰이더를 실행하는 구성 요소입니다

셰이더 코어는 텍스처 유닛과도 쌍을 이루어 텍스처 리소스는 물론, 광선 충돌(Intersection) 요청을 가속하는 새로운 레이 트레이싱 유닛을 샘플링 및 작성할 수 있습니다

셰이더 코어의 구성 부분은 더 세분화될 수 있습니다

각 셰이더 코어에는 FP32 FP16 및 정수 연산 등 다양한 유형의 명령어를 실행하는 일련의 실행 파이프라인이 있으며 이러한 명령어는 Metal 데이터 유형이 float, half, int 등인 셰이더 변수에 대한 작업에 해당합니다 텍스처와 버퍼에 읽기 및 쓰기 작업을 하기 위한 메모리 파이프라인도 있습니다

이러한 실행 파이프라인을 모두 사용하는 데는 일반적으로 여러 SIMD 그룹의 명령어를 실행해야 합니다 따라서 셰이더 코어에서 실행 중인 SIMD 그룹을 추적하는 풀과 다음에 명령어를 실행할 SIMD 그룹을 선택하는 스케줄러가 있습니다

일반적으로 셰이더 프로그램이 사용할 수 있는 다양한 유형의 데이터를 저장하기 위한 온칩 메모리도 몇 개 있습니다 변수 값을 저장하는 레지스터 컴퓨트 스레드 그룹 전반에서 공유되는 데이터나 타일 전반에서 공유되는 색상 첨부 데이터를 저장하는 스레드 그룹 및 타일 메모리

스택과 버퍼에 대한 액세스 성능을 향상시키는 캐시가 있습니다

지금까지 알아본 셰이더 코어의 역할과 그 구성 요소를 바탕으로 Apple Family 9 GPU 셰이더 코어의 새로운 개선 사항 세 가지를 설명하겠습니다 이러한 개선 사항 덕분에 앱을 변경하지 않고도 셰이더의 성능이 향상됩니다 하지만 이 새로운 셰이더 코어의 작동 방식을 이해하면 훨씬 더 많은 이점을 얻을 수 있습니다

첫 번째 변경 사항은 동적 셰이더 코어 메모리입니다 이를 통해 앱은 스레드 점유율을 개선하고 그 결과 성능도 향상될 수 있습니다

두 번째 변경 사항은 유연한 온칩 메모리입니다 이를 통해 셰이더가 버퍼 스택, 스레드 그룹 타일 메모리에 액세스할 때 효율성이 향상됩니다

마지막 변경 사항은 셰이더 코어의 고성능 ALU 파이프라인으로 병렬 실행 기능이 향상되었습니다 이렇게 하면 부동 소수점 또는 정수 연산의 조합을 수행하는 앱의 성능이 향상됩니다

이 새로운 기능을 더 자세히 살펴보기 전에 셰이더 코어가 실행 파이프라인을 사용하는 방식과, 이 과정에서 스레드 점유율이 갖는 중요성에 대해 알아보겠습니다

Metal 셰이더가 ALU 파이프라인을 사용하여 일부 수학 연산을 실행한 후 버퍼를 읽고 그 결과를 바로 사용한다고 가정해 보겠습니다 버퍼에 액세스하려면 기기 메모리까지 이동해야 할 수 있으며 이는 지연 시간이 긴 작업입니다 이 시간 동안 SIMD 그룹은 다른 작업을 실행할 수 없으며 ALU 파이프라인이 사용되지 않게 됩니다

이를 해결하기 위해 셰이더 코어는 다른 SIMD 그룹에서 고유한 ALU 명령어가 포함할 수 있는 명령어를 실행할 수 있습니다 이렇게 하면 ALU가 이동하여 사용되는 시간이 줄고 SIMD 그룹이 병렬로 실행될 수 있으므로 성능이 향상됩니다

셰이더 코어를 실행 중인 추가 SIMD 그룹이 있을 경우 ALU 및 기타 실행 파이프라인이 실행할 명령어를 계속 받을 때까지 이 작업을 반복할 수 있습니다

셰이더 코어에서 동시에 실행되는 SIMD 그룹의 수를 스레드 점유율이라고 합니다

그런데 셰이더 코어에서 동시에 실행될 SIMD 그룹의 수는 어떻게 결정될까요?

그 질문에 대답하기 위해 예제를 하나 보겠습니다 이것은 레이 트레이싱 컴퓨트 커널 프로토타입입니다 가속 스트럭처가 있는 광선과 충돌하고

충돌 결과를 검사하며

이전의 충돌한 머티리얼에 따라 다른 셰이딩 함수를 실행합니다 이 예제에서는 유리, 가죽 머티리얼 셰이딩을 지원합니다

각 코드 라인은 프로그램의 변수를 저장하기 위해 일정량의 레지스터를 사용합니다 프로그램의 다양한 포인트에서 코드의 기능에 따라 레지스터의 사용량이 결정됩니다 이 예제에서는 shadeGlass 함수의 구현은 프로그램의 다른 부분보다 더 많은 레지스터를 사용합니다

Apple Family 9 GPU 이전에는 SIMD 그룹이 온칩 레지스터 파일에서 레지스터를 할당할 때까지는 셰이더 코어에서 실행을 시작할 수 없었습니다 할당된 양은 프로그램의 모든 포인트에서 최대 레지스터 사용량과 동일합니다 SIMD 그룹은 전체 사용 시간 동안 많은 레지스터를 할당 받은 채로 있지만 이러한 레지스터의 대부분은 프로그램의 많은 섹션에서 사용되지 않을 수 있습니다 따라서 최대 레지스터 사용량에 따라 예를 들어, 셰이더 코어에서 한 번에 4개의 SIMD그룹만 실행할 수 있습니다 더 사용하려면 보유한 온칩 레지스터 파일 메모리로는 부족하기 때문입니다 Apple Family 9 GPU는 새로운 동적 셰이더 코어 메모리 기능을 제공하여 더 이상 최대 레지스터 사용량이 실행 가능한 SIMD 그룹의 수를 제한하지 않습니다 이제 온칩 레지스터 메모리는 프로그램의 각 부분이 실제 사용량에 따라 동적으로 할당되고, 셰이더의 사용 시간이 끝나면 할당이 해제됩니다 이를 통해 공간이 확보되어 SIMD 그룹이 온칩 레지스터 파일을 훨씬 더 효율적으로 사용할 수 있습니다 이는 동시에 실행되는 SIMD 그룹의 수를 늘려 앱의 스레드 점유율에 중요한 영향을 주고 결과적으로 성능을 향상합니다

설명드린 내용처럼 이제 레지스터가 SIMD 그룹의 사용 시간 동안 동적으로 할당 및 할당 해제됩니다 이것은 부분적으로는 레지스터 파일이 영구적으로 저장되지 않고 캐시로 존재하기 때문에 가능한 것입니다 즉, 칩이 저장할 수 있는 양보다 더 많은 레지스터를 사용할 수 있습니다

유연한 온칩 메모리 기능은 이 처리를 나머지 셰이더 코어의 메모리 유형에도 적용합니다 스레드 그룹, 타일 메모리 캐시 생성 등에도 적용합니다

이제 레지스터, 스레드 그룹, 타일, 스택 버퍼 데이터가 모두 칩에 캐싱되므로 이러한 모든 메모리 유형을 서비스하는 온칩 메모리 캐시의 수는 줄어들고 용량은 커졌습니다 이러한 유연성은 각 메모리 유형을 과하게 사용하지 않는 셰이더에 도움이 됩니다 과거에는 컴퓨트 커널이 스레드 그룹 메모리 등을 사용하지 않은 경우 해당 온칩 스토리지는 전혀 사용되지 않았습니다 이제 온칩 스토리지는 셰이더에서 사용하는 메모리 유형에 동적으로 할당되어 이전보다 더 많은 온칩 스토리지를 제공하며 궁극적으로는 성능이 향상됩니다

예를 들어 레지스터 사용량이 많은 셰이더는 점유율이 높을 수 있습니다

버퍼 데이터의 대규모 작업 세트에 반복적으로 액세스하는 셰이더의 경우 캐시 적중률이 향상되고 버퍼 액세스 지연 시간이 낮아져 성능이 향상됩니다 함수 포인터 표시되는 함수 테이블 동적으로 연결된 셰이더 라이브러리 등 인라인이 아닌 함수를 많이 사용하는 앱의 경우 함수 매개변수를 전달하기 위한 온칩 스택 공간이 늘어나 더 빠른 함수 호출이 가능합니다

하지만 앱이 여전히 온칩 스토리지보다 더 많은 메모리를 사용한다면 어떻게 될까요? 그대로 두면 해당 데이터는 다음 캐시 레벨, 심지어 메인 메모리로 유출됩니다 셰이더 코어는 셰이더의 동작을 동적으로 모니터링하고 점유율을 조정하여 이러한 일이 발생하지 않도록 합니다 덕분에 데이터는 칩에 유지되며 궁극적으로는 실행 파이프라인이 계속 가동됩니다

하지만 동적 레지스터 사용은 물론 셰이더의 액세스 스레드 그룹 타일, 스택 버퍼 메모리가 셰이더의 점유율에 영향을 줄 것입니다

이러한 새로운 하드웨어 기능은 많은 앱의 점유율을 향상하므로 개발자의 점유율 최적화 작업이 전보다 훨씬 줄어듭니다 그렇지만 Apple Family 9 GPU에서 점유율을 더 최적화해야 하는 경우를 위해 개발된 프로파일링 도구가 있습니다 점유율의 진단 및 최적화 방법을 자세히 알아보려면 이 세션들을 시청하시기 바랍니다

Apple Family 9 GPU 셰이더 코어에서 소개해드릴 마지막 기능은 고성능 ALU 파이프라인입니다

Apple GPU 셰이더 코어에는 FP16 명령어를 포함하여 다양한 명령어 유형에 대한 별도의 ALU 파이프라인이 있습니다 Apple GPU는 FP16 산술 실행을 위해 최적화되었습니다 가능하면 FP16 데이터 유형을 사용하는 것이 좋습니다

FP16 수학 명령어는 최대 처리량에서 실행되며

FP32에 비해 레지스터를 적게 사용합니다 버퍼가 기본적으로 FP16으로 데이터를 저장하는 경우, 메모리 대역폭이 줄어듭니다 그리고 수학 연산의 소스나 대상 변수가 아직 FP16이 아니라면 간단하게 변환할 수 있습니다

그러나 앱에서 FP32 및 정수 등 다른 수학 연산을 수행해야 하는 경우 Apple Family 9 GPU 셰이더 코어는 그 어느 때보다 우수한 병렬 처리 성능으로 세 데이터 유형의 명령어를 모두 실행할 수 있습니다 이전 Apple GPU에 비해 ​​최대 2배의 ALU 성능을 제공할 수 있습니다 이러한 탁월한 병렬성을 활용하려면 명령어가 여러 SIMD 그룹에서 실행되어야 합니다 즉, 점유율을 늘리면 ALU 파이프라인의 활용을 향상할 수 있습니다 예를 들어 보겠습니다 동시 실행 중인 두 개의 SIMD 그룹이 ALU 명령을 실행하고 있습니다 이전에는 이러한 SIMD 그룹이 차례대로 실행되어야 했을 것입니다

이제는 다른 시점에 실행할 FP32 및 FP16 명령어가 있는 경우 이를 동시에 실행하여 병렬성과 성능을 향상할 수 있습니다

차세대 셰이더 코어의 새로운 기능을 요약하자면 셰이더의 사용 시간 동안 레지스터를 동적으로 할당 및 할당 해제하여 스레드 점유율을 향상합니다

레지스터, 스레드 그룹 타일, 스택, 버퍼 메모리를 서비스하는 대용량 온칩 캐시가 있어 이러한 메모리 유형에 대한 액세스 성능이 향상됩니다

셰이더 코어는 점유율을 동적으로 조정하여, 데이터를 칩에 유지하고 실행 파이프라인을 계속 가동합니다

마지막으로 FP16 FP32 및 정수 연산의 병렬 실행 성능이 그 어느 때보다 높아져 ALU 성능이 향상됩니다

다음으로 하드웨어 가속 레이 트레이싱입니다

Metal 레이 트레이싱을 사용하면 앱은 Apple GPU의 막대한 병렬성을 활용하여 장면 지오메트리와 충돌하는 광선을 구현할 수 있습니다 Metal 레이 트레이싱에 대해 자세히 알아보려면 ‘Metal 레이 트레이싱 가이드’와 ‘Metal 레이 트레이싱으로 앱 향상하기’를 시청하세요

Metal 레이 트레이싱 API의 중심에는 인터섹터 오브젝트가 있습니다 이 오브젝트는 가속 스트럭처 안에 있는 프리미티브로 광선의 충돌 포인트를 지정합니다 레이 트레이싱 앱의 GPU 함수 즉, 셰이더는 이 오브젝트를 자주 반복 호출하므로 앱 성능의 핵심이 됩니다

이 GPU 함수는 아까 레이 트레이싱 커널의 레지스터 사용과 함께 보여드렸습니다 이 함수는 인터섹터 객체를 생성하고 해당 객체의 intersect 메서드를 호출하여 충돌을 찾습니다

인터섹터는 충돌 포인트를 결정하기 위해 주요 단계를 거칩니다 먼저 가속 스트럭처를 순회하여 후보 프리미티브를 찾습니다 다음으로 앱에서 제공할 수 있는 충돌 함수를 호출하여 광선이 프리미티브와 충돌하는지 확인합니다

충돌하는 경우 해당 충돌을 이전 충돌과 비교하며 가장 가까운 포인트를 찾을 때까지 이 과정을 반복합니다

그런 다음 앱에서 처리하기 위해 가장 가까운 충돌이 호출했던 GPU 함수로 반환됩니다

Apple Family 9 GPU의 새로운 기능인 인터섹터 객체 구현은 매우 중요한 작업으로, 하드웨어 가속으로 그 성능이 크게 향상됩니다

하드웨어 가속 충돌은 GPU 함수와 인라인으로 실행되지 않습니다 따라서 광선과 광선-충돌 간 광선 페이로드의 커뮤니케이션을 용이하게 하기 위해 온칩 메모리에서 데이터를 읽고 씁니다 이는 새로운 Xcode의 RT 스크래치 성능 카운터에서 관측할 수 있습니다

인터섹터의 역할과 작업 범위를 알아보았습니다 이제 예제를 사용하여 인터섹터 루프를 따라가 보며 성능 특성을 분석해 보겠습니다

예제의 앱은 4개의 광선을 가속 스트럭처로 충돌시키려는 2개의 SIMD 그룹을 각각 실행하고 있습니다

이 예제에서 가속 스트럭처에는 직육면체 객체 하나 구체 객체 하나가 포함된 일반적인 Cornell 박스가 있습니다

광선은 intersect 메서드를 호출하고 여기에 광선, 가속 구조 충돌 함수 테이블을 전달하여 장면에 반영됩니다 각 SIMD 그룹은 직육면체와 충돌하는 광선 2개 구체와 충돌하는 광선 2개가 있습니다 이 예제에서 직육면체는 불투명 삼각 프리미티브로 정의됩니다 MTLAccelerationStructure

TriangleGeometryDescriptor에서 opaque 속성을 YES로 설정하면 광선 충돌은 Metal에 내장된 충돌 함수를 사용하여 계산됩니다

하지만 구체의 충돌은 사용자 설정 BoundingBoxIntersection 함수를 반드시 호출하여 단계적으로 정의됩니다

사용자 설정 BoundingBoxIntersection 함수는 bounding_box 매개변수와 함께 intersection 속성을 사용하여 선언됩니다

앞서 언급했듯이 가속 스트럭처에서 광선을 테스트하는 각 스레드는 intersect 메서드를 호출합니다 따라서 이 예제를 염두에 두고 기존 구현에서 각 충돌 호출 순회와 충돌 테스트가 실행되는 방식을 살펴보겠습니다

일반적인 사용 시 모든 순회에서 광선을 테스트할 프리미티브를 찾는 데 동일한 시간이 걸리는 건 아닙니다 따라서 실행 분기(Execution Divergence)가 생성됩니다 SIMD 그룹의 각 스레드는 다음으로 진행하기 전에 해당 SIMD 그룹에서 가장 긴 순회를 기다리게 됩니다

결과적으로 충돌 함수를 실행할 때도 동일한 오버헤드가 복합됩니다 실행 분기로 인해 각 충돌 함수 유형이 한 번에 하나씩 실행되어 병렬성이 더 감소합니다 각 스레드는 두 단계에서 모두 유휴 런타임의 상당 부분을 소비하며 SIMD 그룹의 다른 스레드가 완료되기를 기다려야 하므로 중대한 성능 병목 현상이 발생합니다

기존 구조를 염두에 두고 하드웨어 가속 레이 트레이싱이 이러한 비효율을 어떻게 최적화하는지 보겠습니다

첫 번째 주요 개선점은 하드웨어 인터섹터가 고정 기능 하드웨어를 사용하여 각 순회를 독립적으로 실행할 수 있게 된 것입니다 이것이 가능한 이유는, 부분적으로 해당 배열이 GPU 함수와 인라인으로 실행되는 대신 처리를 위해 하드웨어 인터섹터로 전송되기 때문입니다 이는 순회에 소요되는 시간을 크게 줄이고 기존 순회 실행 분기의 오버헤드를 제거합니다

반면 충돌 함수는 Metal 셰이딩 언어 코드입니다 따라서 셰이더 코어에서 실행하려면 SIMD 그룹으로 그룹화되어야 합니다 하지만 하드웨어 인터섹터는 각 광선을 독립적으로 실행하므로 별도의 SIMD 그룹에서 발생한 광선의 충돌 함수 호출을 자유롭게 그룹화할 수 있습니다

이것이 리오더(Reorder) 단계의 역할입니다 광선이 가까운 거리와 시간 내에 이 단계에 도달하면 충돌 함수 호출은 해당하는 SIMD 그룹으로 그룹화됩니다 기존 구현에서 보였던 실행 분기는 감소하거나, 심지어 완전히 사라집니다

지금까지 하드웨어 가속 레이 트레이싱이 앱의 인터섹터 호출 성능을 어떻게 향상할 수 있는지 보여드렸습니다 이제 그 이점을 극대화하기 위해 앱에 구현할 수 있는 모범 사례를 살펴보겠습니다

첫 번째로, 가능하면 인터섹터 객체 API를 사용하는 것이 좋습니다 Metal은 충돌 쿼리 API를 사용하는 레이 트레이싱도 지원하지만 이 API는 읽기 및 쓰기를 요하는 임시 레이 트레이싱 메모리의 양을 증가시키며 리오더 단계도 비활성화됩니다

두 번째로 사용자 설정 교차 함수를 작성할 때 논리적으로 다양한 충돌 루틴을 실행할 수 있는 단일 우버 함수를 생성하지 않는 것이 좋습니다 대신 논리적 충돌 루틴마다 하나의 Metal 충돌 함수를 생성하세요 이렇게 하면 리오더 단계의 이점이 늘어납니다

인터섹터 객체로 전달 및 반환되는 광선 페이로드 스트럭처의 크기를 최소화하는 것도 중요합니다 이렇게 하면 셰이더의 지연 시간이 줄어들고 스레드 점유율이 늘어날 가능성이 있습니다

레이 트레이싱 앱을 최적화하는 방법에 관한 자세한 내용과 지침은 화면의 세션에서 확인하세요

요약하자면 Apple Family 9 GPU는 고정 기능 순회 블록과 충돌 함수 리오더 단계를 갖춘 새로운 하드웨어 가속을 통해 레이 트레이싱 성능을 크게 향상했습니다

이 새로운 하드웨어는 모든 Metal 레이 트레이싱 앱의 성능을 향상시켜 GPU가 제공하는 이점을 극대화하지만 가능하면 충돌 쿼리 API 대신 충돌 API를 사용하는 것이 가장 좋습니다

Apple Family 9 GPU에 관해 소개할 마지막 개선 사항은 하드웨어 가속 메시 셰이딩입니다

메시 셰이딩은 렌더링 파이프라인에서 GPU를 기반으로 하는 유연한 지오메트리 프로세싱 단계로, 기존의 버텍스 셰이더 단계가 컴퓨트와 유사한 2개의 셰이더로 변경되었습니다

객체 셰이더는 첫 번째 단계에서 실행되며 전체 메시 객체와 같이 앱에 특정된 입력의 대략적인 프로세싱에 사용할 수 있습니다 각 객체 스레드 그룹은 메시 그룹을 생성하여 이후의 세부적인 프로세싱을 수행하도록 선택할 수 있습니다 메시 셰이더는 여기서 두 번째 단계를 구성합니다 일반적으로 메시 스레드 그룹은 종종 메시렛(Meshlet)이라고 하는 상위 객체의 구성 부분을 처리합니다

메시 스레드 그룹의 출력은 Metal 메시 객체로 기존 그래픽 파이프라인의 나머지 부분에서 처리할 버텍스 및 프리미티브의 목록을 캡슐화합니다

메시 셰이딩을 응용하는 방법은 다양합니다 세부 지오메트리 컬링

단계적 지오메트리 생성 압축 포맷과 같은 사용자 설정 앱 특정 지오메트리 표현 등이 있습니다 다른 그래픽 API에서 지오메트리 및 테셀레이션 셰이더를 포팅할 때도 사용됩니다

Metal 메시 셰이딩에 대한 자세한 내용은 화면의 두 세션에서 확인해 보세요

Apple Family 9 GPU의 하드웨어 가속 메시 셰이딩에서 가장 눈에 띄는 개선 사항은 기존 메시 셰이딩 코드의 성능이 훨씬 향상된다는 것입니다

Apple Family 9 GPU는 객체와 메시 스레드 그룹을 더 효율적으로 스케줄링하여 중간 메시렛 데이터를 온칩으로 유지하므로 메모리 트래픽이 감소합니다

새로운 하드웨어는 향상된 Metal API 기능도 갖추고 있습니다 첫 번째로 메시 드로우 명령을 간접 명령 버퍼로 인코딩하는 것을 지원합니다 따라서 GPU 기반 렌더링 파이프라인이 기존의 버텍스 셰이더는 물론 메시 셰이딩을 활용할 수 있습니다

두 번째 API 개선 사항은 메시 그리드당 최대 스레드 그룹 수가 1,024에서 1백만 이상으로 확대된 것입니다

이제 최적의 메시 셰이딩 성능을 위한 모범 사례를 검토해 보겠습니다

메시 스레드 그룹이 출력한 metal::mesh 객체에는 여러 템플릿 매개변수가 있는데 그 크기를 최대한 작게 유지하는 게 중요합니다

메시의 버텍스 및 프리미티브 데이터 유형의 경우 이를 관련 없는 다른 버텍스 또는 메시 함수와 공유하여 생긴, 사용하지 않는 속성을 제거하여 크기를 줄일 수 있습니다 메시 유형은 출력될 수 있는 프리미티브 및 버텍스의 최대 개수도 지정해야 합니다 앱의 지오메트리 파이프라인 애셋에 실제 필요한 수보다 더 높게 설정하면 안 됩니다 이 크기를 작게 유지하면 메모리 트래픽이 줄며 점유율을 높일 수 있습니다

메시 셰이더에서 프리미티브별 컬링을 수행하는 경우 메시 객체에 단지 하드웨어의 후속 호출 단계에서 호출하려는 목적으로 버텍스 위치를 작성하지 않는 것이 좋습니다 대신 그러한 프리미티브를 완전히 생략하세요 나머지 하드웨어 지오메트리 프로세싱 단계에서 처리 시간이 크게 절약됩니다

Apple Family 9 GPU에 관한 내용을 요약하겠습니다

차세대 셰이더 코어는 레지스터 스토리지를 동적으로 할당하고 다양한 메모리 유형에서 온칩 메모리를 공유하고 활용도를 증가시켜 스레드 점유율과 성능을 향상시킵니다

하드웨어 가속 레이 트레이싱은 새로운 고품질 시각 효과를 실현하는 Metal 레이 트레이싱 API로 앱의 성능을 크게 향상시킵니다 마지막으로 하드웨어 가속으로 메시 셰이딩 성능이 크게 향상되어 더 많은 앱에서 지오메트리 프로세싱 파이프라인을 사용자 설정할 수 있게 되었습니다

시청해 주셔서 감사합니다