이전됨(Moved). 이 드라이버는 이제 커널의 리포지토리에서 개발되고 있습니다. 히스토리는 다음과 같습니다:
hakios-xnu-kernel/hakios-metal (커널 2.3.3 이후)
커널을 체크아웃(checkout)하면 빌드되며, 수정 사항은 그곳에 반영됩니다. 이곳에는 새로운 내용이 올라오지 않으며, 이 리포지토리는 단지 해당 위치로 연결되는 링크를 위해 유지됩니다.
hakios의 Metal 드라이버가 C++로 구현되었습니다. 이는 휴대폰의 Adreno GPU(Adreno 7c+ Gen 3, chip id 0x06030500)에서 작동하는 Metal의 커맨드 모델을 커널 자체의 GPU 드라이버(HakiosGPU)를 통해 Mesa 없이 처리합니다.
현재 존재하는 내용과 각 부분이 검사되는 방식은 다음과 같습니다:
| 파트 | 기능 | 검사 주체 |
|---|
pm4.hpp | 커맨드 프로세서의 패킷(type 4 및 type 7, 패리티 비트 포함) | 모든 헤더는 Mesa의 pm4_pkt4_hdr / pm4_pkt7_hdr와 비교되며, 모든 128개의 opcode가 검사됩니다. opcode는 adreno_pm4.xml과 비교됩니다. |
| ... | | |
'fake GPU'가 수행하지 않는 것: 셰이더 명령어(shader instruction)를 실행하거나, fake_gpu.hpp에 명시된 검사를 넘어 레지스터의 값을 판단하는 것입니다. 이 드라이버가 수락하는 디스패치(dispatch)는 잘 구성되어 있으며, GPU가 할당받은 메모리만 이름을 지정합니다. 실제로 GPU가 이를 실행할지는 휴대폰이 결정할 문제입니다. 과거에는 'CP_EVENT_WRITE'의 TIMESTAMP 비트가 GPU에게 종료 대기(end-of-dispatch wait)가 값을 기대하는 곳에 타임스탬프를 쓰게 만들어, 모든 디스패치가 해당 비트가 클리어될 때까지 멈추는 문제가 있었습니다 (--bisect을 통해 발견되었으며, 부팅 시 한 번 발생). 이제 fake GPU는 이 비트에 대한 타임스탬프를 작성하므로, 스트림이 여기서도 실패하게 됩니다.
현재까지 Apple의 셰이더(shaders)는 다음과 같습니다: 프로브가 휴대폰 자체 파일 시스템에서 .metallib 파일을 읽고, 장치에서 해당 함수의 컴파일을 수행한 다음, 캐나리(canary)-채워진 버퍼들(--metallib FILE FUNCTION) 위에서 실행하며, 각 버퍼별로 기록된 워드 수, 전체 버퍼의 해시값, 그리고 처음 여덟 개의 워드를 출력합니다. 레퍼런스는 hakios의 tools/air/kernelref.py입니다: 동일한 커널을 같은 버퍼들 위에서 LLVM을 통해 리눅스 호스트(Linux host)에서 실행하고, 동일한 라인을 출력하므로 두 로그가 다릅니다(diff). 10월 28일(hakios logs/215528-run.log)에 휴대폰에서 컴파일되어 완료된 PortraitFilters의 여덟 가지 버퍼 전용 커널이 있었고, 이 여덟 개 모두 LLVM 실행 결과와 해시값을 포함하여 동일하게 나왔습니다 (MetalFaceMask_InitMinMax, MetalFaceMask_InitSpanTable, image_calcminmax0, image_calcminmax2, fbs_bistochastize_init, fbs_pcg_iter3, MetalFaceMask_PropagateParents, MetalFaceMask_CalcCubeInputScaling). 마지막 커널은 처음에는 세 개의 float에서 가장 마지막 자리에 GPU 차이로 한 단위가 읽혔는데, 이 단위는 호스트(host)의 것이었습니다: Apple의 IR이 포함하는 fast-math 플래그 하에서 LLVM의 x86 백엔드(backend)는 벡터를 역수 추정치로 나누고 뉴턴 스텝(Newton step)을 한 번 수행했는데, 여기서 Adreno의 역수 계산 및 32 곱하기가 올바르게 반올림된 부분은 레퍼런스가 정확하게 나눕니다. 텍스처(textures), 샘플러(samplers), 스레드그룹 메모리(threadgroup memory), 배리어(barriers) 또는 원자적 연산(atomics)을 사용하는 커널은 실행되지 않았습니다 (휴대폰에서는 이름으로 거부되었고, 호스트의 레퍼런스 범위 밖에 있습니다; 두 개의 InitMinMax/InitSpanTable 원자적 저장(atomic stores)은 일반 저장(plain stores)으로 비교되는데, 이는 값에 대해서는 무언가를 말해주지만 순서(ordering)에 대해서는 아무것도 말해주지 않습니다). 현재까지 Metal 자체 객체: hakios의 플러그인(userland/src/HakiosMetal.m: MTLDevice와 그것이 제공하는 객체들로, 각각 Metal의 비공개 기본 클래스(private base class)의 하위 클래스이며, hakios_metal.h를 통해 접근 가능)은 Metal 프로그램이 도달하는 지점이며, hakios의 userland/src/metal_compute_test.m은 이 레포지토리의 어떤 것도 연결하지 않고 오직 Metal.framework만을 통해 GPU 클라이언트로서 이를 구동합니다.
10월 28일 (hakios logs/222820-run.log, 양쪽 도로 모두 실패 없음) Metal 자체의 발견으로, 부트 인자(boot-arg) metal=1과 서비스에 있는 두 속성을 사용하여 MTLCreateSystemDefaultDevice()에서 우리 클래스를 반환했습니다. 이미지 라이브러리의 MetalFaceMask_InitMinMax와 메모리에서 빌드된 .metallib의 fill_id가 MTLComputeCommandEncoder를 통해 실행되었습니다. 1024개의 워드 중 1024개, 버퍼로부터의 상수 및 -setBytes:length:atIndex:로부터의 상수가 포함되었고, 동일한 명령 버퍼(command buffer) 내에서 블릿 복사(blit copy)와 채우기(fill)가 이루어졌으며, 절반 그리드에 걸쳐 -dispatchThreads:를 사용했습니다. 첫 번째 실행(logs/222249-run.log)은 GPU에 문제가 없었음에도 불구하고 두 번 실패했습니다: 드라이버가 ir3의 패딩된 64바이트 상수 범위(constant range)에 대한 4바이트 바인딩을 거부했고 (수정됨: 이 범위는 객체의 페이지를 기준으로 측정됨), 테스트 테이블은 프로브 라인의 재진술로부터 작성되었으며, 예상했던 네 개의 워드 대신 여섯 개가 측정되었습니다. 렌더 인코더(Render encoders), 텍스처(textures), 샘플러(samplers), 히프(heaps), 펜스(fences) 및 이벤트(events)는 이름으로 거부됩니다.
2D 컴퓨팅(compute)의 텍스처와 샘플러는 제공되었습니다 (샘플링된 텍스처, stib로 작성된 이미지, 경계 테이블을 통한 클램프-투-제로(clamp-to-zero)가 적용된 Metal의 샘플러). 이는 가짜 GPU에서 확인되었고, 10월 29일 (hakios logs/011104-run.log, 012127-run.log)에 tools/kernels/tex_fill.ll, tex_copy.ll, tex_sample_gid.ll 및 tex_fill_h.ll을 사용하여 휴대폰에서 측정되었습니다: IBO 테이블을 통해 작성된 이미지는 R32Uint의 모든 픽셀과 RGBA32Uint의 네 개 레인 전체에 기록되었으며 (2048개 중 2048개, 행 패딩은 그대로 유지되고 나머지 절반 그리드는 캐나리(canary)로 남겨짐), 인접한 텍스처 테이블을 통한 텍스처 읽기는 2048개 중 2048개를 복사했고, Metal의 샘플러는 가장 가까운 이웃 필터링(nearest filtering) 및 선형 필터링(linear filtering)으로 중심에서 정확히 텍셀을 반환했으며, 사분위 위치에서의 쌍선형 평균(bilinear mean)은 0.000123 내에 있었습니다 (예상대로 가장 가까운 이웃과 다른 64개 중 60개의 픽셀). 또한 절반 이미지는 2048개 중 2048개의 정확한 반 비트 패턴을 받았습니다.
이전에 발견된 버그가 hakios의 프론트엔드에서 발생했습니다: 호출의 void 반환 타입에 의해 모든 텍스처 쓰기가 타이핑되어, half 및 float 값이 정수(integer)로 저장되었고(stib.b.typed.2d.u16), 드라이버는 이를 읽어 커널 클래스를 기반으로 Apple의 MetalFaceMask_Clear에 대해 RGBA16Float을 거부했습니다. 이제 프론트엔드는 값 자체로 스토어를 타이핑하며, 명령어의 클래스와 불일치하는 선언은 컴파일 시간에 컴파일러와 픽스처 작성기(fixture writer) 모두에 의해 거부됩니다. 아직 지원하지 않는 기능으로는 2D가 아닌 다른 텍스처들, 하나의 레벨/하나의 슬라이스 (1D, 3D, cube, arrays, mipmaps, MSAA), 투명 검정색이 아닌 경계 색상, 비교 샘플러(compare samplers), 디스패치에서의 프론트엔드 null 텍스처(그 아래에서 그리는 draw가 이를 사용함), 버퍼 위의 텍스처(bytesPerRow는 호출자가 선택), private memory, threadgroup memory (각각 이름으로 거부되며 사용되지 않음) 등이 있습니다. 인덱싱된 constant 로드 ir3는 UBO 디스크립터를 통해 디스패치에서 draw와 마찬가지로 제공됩니다.
첫 번째 그리기(draw) 작업(render.hpp, 10월 29일)이 작성되었고, 가짜 GPU에서 검사되었으며, 같은 날 휴대폰에서 측정되었습니다 (hakios logs/051851-run.log, 첫 실행, 모든 라인 통과): 하나의 삼각형을 그리고, 다음으로 쌍(pair)을 그린 후, per-vertex 색상을 sysmem에 있는 64x64 RGBA8Unorm 텍스처로 기록했습니다 (직접 렌더링 모드). 이 값은 CPU가 다시 읽어와 프로브가 계산한 규칙과 픽셀별로 비교합니다 (아래 스테이지 9): 내부 픽셀 2016/2016이 색상을 가지며, 외부 픽셀 2016/2016은 캐나리아(canary)를 가집니다. 좌측 하단 삼각형에 의해 페인팅된 대각선 중심은 0개이고, 그 보완체에 의해 64/64가 페인팅되었습니다 (좌측 상단 규칙과 두 개가 합쳐져 타겟을 정확히 한 번 커버함: 4096/4096). 절반의 뷰포트 오른쪽(half viewport right), vertexStart 3은 두 번째 삼각형만 그립니다 (따라서 vertex_id는 VFD_INDEX_OFFSET을 포함합니다). 모든 4032개의 비대각선 픽셀에서 per-vertex 색상은 1/255 이내이며(최대 편차 0), 네 개의 스테이지가 호스트의 픽스처를 단어 그대로 컴파일했으며, 패스당 시간은 41~68 마이크로초(us)였습니다.
Metal 자체 객체들(hakios의 logs/051907-run.log, metal_render_test, 양쪽 경로 모두 0 실패)을 통해 동일한 드로우가 동일한 카운트로 작동합니다: 장치에 연결된 MTLRenderPipelineState, MTLRenderCommandEncoder, drawPrimitives: -- 그리고 렌더 사용처가 없는 텍스처는 해당 명령 버퍼를 오류로 종료시키며, 텍스처는 변경되지 않습니다. 셰이더는 저희가 작성한 것입니다 (tools/kernels/tri_flat_*.ll, tri_color_*.ll: 어느 단계에서도 버퍼가 없는, [[vertex_id]]에 의해 선택된 위치로 turnip의 자체 클리어 경로가 드로우하는 방식과 같습니다). 레지스터 시퀀스는 turnip의 sysmem 패스와 r3d 경로이며 fd6이 두 번째 증인 역할을 하고, 모든 레지스터는 해당 라인에서 인용되며, 각 그룹은 render::Options 스위치 아래에 위치합니다. 따라서 --bisect를 한 번 부팅하면 하드웨어가 어디를 원하는지, 또는 두 증인이나 커널 자체의 부팅 그림이 불일치하는지 알 수 있습니다 (RB_CCU_CNTL 비트 2, BIN_CONTROL의 LRZ 피드백 비트, RB_RENDER_CNTL, SP_FLOAT_CNTL, GRAS_UNKNOWN_8110, USE/IGNORE_VISIBILITY, LRZ_FLUSH, RB_DONE_TS). GRAS_CL_VPORT_YSCALE의 부호는 Metal의 관례(NDC y축이 위쪽, 0번째 행이 맨 위)에 따라 음수입니다: 프로브의 방향 라인이 이를 확인했습니다 (왼쪽 하단 절반이 드로우됨), 그리고 모든 스위치의 기본값 -- turnip과 fd6이 일치하는 값은 turnip의 값, 그렇지 않은 더 작은 세트에서는 그 값 -- 은 bisect 없이 드로우되었습니다.
아직 안 됨. 이 모든 것이 이름으로 거부됩니다(제공되지 않음): per-vertex Float..Float4, Half2, Half4 및 UChar4Normalized를 초과하는 정점 속성들 (이들은 정점 디스크립터: 'tri_attr_vs.ll'의 Float2 위치와 하나의 인터리브드 버퍼에서 가져온 UChar4Normalized 색상으로 제공됩니다. 휴대폰에서 정확히 4032/4032 픽셀입니다, 'logs/142841-run.log'); 고정 오프셋에서 읽는 'constant T&'를 초과하는 정점 또는 프래그먼트 스테이지의 버퍼들 (이것은 푸시된 범위이며 제공됩니다: 'tri_cflat_fs.ll'의 [[buffer(0)]]으로부터 가져온 색상으로, 휴대폰에서 정확히 두 드로우에 두 개의 색상이 있습니다, 'logs/142053-run.log'; 인덱싱된 하나는 UBO 디스크립터를 통해 제공되며, 'device' 관련은 3D IBO 테이블을 통해 제공됩니다, 아래); 정점 스테이지의 텍스처 및 샘플러, 드로우 내의 이미지들, Depth32Float를 제외한 스텐실(stencil), 깊이 포맷(depth formats), 프래그먼트 깊이 출력(fragment depth output), 이중 소스 블렌딩(dual-source blending), GMEM (타일링된) 렌더링, MSAA, 기본 인스턴스(base instance), 음수 기본 정점(negative base vertex), 점 및 선, 여러 색상 첨부(color attachments), 드로우가 제공하지 않는 것을 바인딩하는 스테이지를 가진 프로그램.
Clear의 로드 액션은 turnip이 단일 샘플링된 sysmem attachment를 지우는 방식과 같이 처리된다: 어떤 드로우 상태(clear_sysmem_attachment, tu_clear_blit.cc 내의 r2d_*)보다 먼저 2D 엔진을 사용하여 전체 타겟에 하나의 CP_BLIT 단색 채우기(solid fill)가 이루어지며, 색상은 r2d_clear_value로 패킹되고, 이후 색상은 CCU가 정리 및 무효화된다. 드로우가 없는 패스 역시 그 자체로 독립적인 패스가 된다(CommandBuffer::clear, hakios_metal_cb_clear)는 것이 기록한 내용이 아무것도 없는 인코더(encoder)가 필요한 것과 같다. 휴대폰 환경(hakios logs/145325-run.log, 프로브 라인 9i, 실패 0건)에서는 모든 픽셀이 패킹된 색상을 가지며 사각형 외부의 바이트는 변경되지 않는다. 이는 RGBA8Unorm, BGRA8Unorm (메모리에서 B가 먼저), RGBA8Unorm_sRGB (선형 0.5를 188로 저장하며 알파도 선형), RGBA16Float (서브노멀 하프까지 지원), R32Uint, RGBA32Float (-0 유지) 등이며, 48x20 타겟의 경우 행 패딩과 4행 패딩은 canary 상태를 유지한다. tri_flat 이전에 Clear가 발생하면 삼각형의 색상은 내부에는 남아있고 외부와 대각선 부분에만 Clear 색상이 남는다. 가짜 GPU는 2D 레지스터들이 제대로 작동하는 것을 확인한 후 채우기(fill) 내용을 메모리로 가져간다 (이것이 유일하게 가져가는 것). 따라서 호스트 테스트도 바이트를 비교하며, 아홉 개의 변형된 채우기 중 일부를 거부한다.
블렌딩(Blending)은 색상 첨부(colour attachment) 0에 대해 tu6_emit_blend가 기록하는 방식으로 제공됩니다: MTLBlendOperation의 숫자는 오퍼코드(opcodes)이며, MTLBlendFactor는 adreno_rb_blend_factor로 매핑되고, tu6_blend_factor는 Vulkan의 값을 매핑합니다. MTLColorWriteMask의 비트들은 컴포넌트 활성화(component enable)에 역순으로 기록되며, 블렌드 색상은 RB_BLEND_*_F32로 전달됩니다. Source1 팩터(dual source)와 정수 형식으로의 블렌딩은 이름상 거부됩니다. 휴대폰 환경에서 (hakios logs/150855-run.log, probe line 9j, 0 failures), tri_cflat의 반투명 색상이 지워진 타겟 위에 표시된 결과는 스트레이트(straight) 및 프리멀티플라이드 소스 오버(premultiplied source-over), 클램프 처리된 ReverseSubtract (0으로 제한), 알파에 대한 Max with Min, BlendColor와 BlendAlpha 팩터, 그리고 블렌딩이 없는 Red | Alpha 쓰기 마스크의 경우, 2016개 픽셀 모두에서 Metal의 방정식과 일치하며 최대 편차는 0입니다. 삼각형 외부에서는 지워진 색상이 정확하게 유지됩니다.
샘플링된 텍스처(Sampled textures)는 프래그먼트 단계(fragment stage)에서 제공됩니다. 텍스처 및 샘플러 디스크립터(texture and sampler descriptors)는 Gallium의 FS_TEX 그룹에 의해 패스 테이블(pass's table)에 로드됩니다 (SB6_FS_TEX, SP_FS_TEX_SAMP/CONST/COUNT, fd6_texture.cc). 변수(varying)인 좌표를 가진 샘플은 ir3가 첫 번째 명령어 이전에 프리페치(prefetch)로 전환하여 발행하며, 이는 A6XX용으로 fd6가 기록하는 대로 커맨드 슬롯에 담겨 SP_FS_PREFETCH_CMD로 나갑니다. 이제 스테이지 픽스처(stage fixtures)는 프리페치 명령을 가지고 (ir3-dump.patch), 드라이버 내 컴파일러가 동일한 방식으로 이를 채웁니다. 휴대폰 환경에서 (hakios logs/152752-run.log, probe line 9k, 0 failures), tri_tex_fs.ll은 tri_color_vs의 보간된 (red, green) 좌표에서 4x4 텍스처를 샘플링하며, 3840개 중 3840개의 판단된 픽셀이 정확히 해당 좌표 이름의 텍셀(texel)을 보유하고 있습니다. 192개의 픽셀은 텍셀 가장자리에서 1/64 떨어진 곳에 남아 판단되지 않았으며, 로우 플립된 텍셀은 발견되지 않았습니다.