[Particle Simulator] Compute Shader 사용해 gpu 연산하기

[Particle Simulator] cpu 연산 - gpu 랜더링 방법

기존 CPU 기반 구현의 한계

지난 글에서는 입자 시뮬레이션을 위해 프레임마다 각 입자에 대한 계산을 cpu에서 실행하고 그 결과를 gpu를 통해 렌더링하는 방식을 사용했습니다. 

private func updateSandSimulation() {
    // 아래에서 위로 순회 (중력 방향을 고려)
    for y in (0..<gridHeight).reversed() {
        for x in 0..<gridWidth {
            let index = y * gridWidth + x
            
            // 빈 공간이거나 이미 업데이트된 입자는 건너뛰기
            if particles[index].isEmpty || particles[index].hasBeenUpdated {
                continue
            }
            
            // 모래 입자 업데이트 로직
            let belowY = y + 1
            
            // 화면 바닥에 도달했는지 확인
            if belowY >= gridHeight {
                particles[index].hasBeenUpdated = true
                continue
            }
            
            // 1. 바로 아래가 비어있으면 아래로 떨어짐
            if particles[belowY * gridWidth + x].isEmpty {
                moveParticle(fromIndex: index, toIndex: belowY * gridWidth + x)
            }
            // 2. 왼쪽 아래가 비어있으면 왼쪽 아래로 이동
            else if x > 0 && particles[belowY * gridWidth + (x-1)].isEmpty {
                moveParticle(fromIndex: index, toIndex: belowY * gridWidth + (x-1))
            }
            // 3. 오른쪽 아래가 비어있으면 오른쪽 아래로 이동
            else if x < gridWidth-1 && particles[belowY * gridWidth + (x+1)].isEmpty {
                moveParticle(fromIndex: index, toIndex: belowY * gridWidth + (x+1))
            }
            // 4. 이동할 수 없으면 해당 위치에 정지
            else {
                particles[index].hasBeenUpdated = true
            }
        }
    }
    
    // 다음 프레임을 위해 업데이트 플래그 초기화
    for i in 0..<particles.count {
        particles[i].hasBeenUpdated = false
    }
}

이 방식은 구현이 직관적이지만 매 프레임마다 cpu가 모든 격자에 대해 순차적으로 새로운 상태를 계산하고 이를 업데이트하며 다음과 같은 문제가 발생했습니다. 

1. 성능 제한: 100x100 격자만 사용해도 CPU 사용량이 99%까지 치솟았습니다.
2. 확장성 부족: 입자 수를 늘리거나 격자를 키우면 성능이 급격히 저하되었습니다.
3. 발열 문제: 지속적인 고부하 CPU 작업으로 기기 발열이 심해졌습니다.

GPU로의 전환: 접근 방식 변경

GPU는 CPU와 달리 수많은 코어를 통한 병렬 처리에 최적화되어 있습니다. 이러한 특성은 CPU에게 부담이 되는 다량의 입자 상태 계산 및 업데이트에 이상적입니다. GPU는 원래 모니터와 같은 그래픽 출력 장치의 수많은 픽셀을 매 프레임마다 갱신하도록 설계되었기에, 입자 시뮬레이터처럼 무수한 객체가 동시에 계산되어야 하는 작업에 특히 적합합니다. 또한 컴퓨터 셰이더를 활용하면 GPU를 그래픽 렌더링뿐 아니라 다양한 일반 연산 처리에도 효과적으로 활용할 수 있습니다.

데이터 구조 조정 필요 

GPU 메모리 정렬 요구 사항에 맞게 데이터 구조를 변경했습니다.

// CPU 구현의 입자 구조체
struct Particle {
    var isEmpty: Bool
    var hasBeenUpdated: Bool
}

// GPU 구현의 입자 구조체
struct Particle {
    var isEmpty: UInt32    // GPU에서 bool 대신 UInt32 사용 (메모리 정렬)
    var type: UInt32       // 입자 타입
}

GPU를 위한 데이터 구조 조정이 필요한 이유

1. 메모리 정렬(Memory Alignment): GPU는 32비트(4바이트) 또는 16바이트 경계에 정렬된 메모리 접근에 최적화되어 있습니다. CPU에서는 Bool 타입이 일반적으로 1바이트만 차지하지만, 이런 작은 크기의 데이터는 GPU에서 처리할 때 비효율적입니다. GPU는 4바이트 단위로 메모리를 읽고 쓰기 때문에, 1바이트 Bool 대신 4바이트 UInt32를 사용하면 메모리 접근 패턴이 GPU의 하드웨어 특성에 더 잘 맞습니다.
2. SIMD 처리 효율성: GPU는 SIMD(Single Instruction, Multiple Data) 아키텍처를 사용합니다. 즉, 동일한 명령어가 여러 데이터에 동시에 적용됩니다. 32비트 정수는 GPU의 SIMD 레지스터와 처리 유닛에 자연스럽게 매핑되어 최대한의 병렬 처리 효율을 얻을 수 있습니다.
3. 원자적 연산 지원: GPU에서 여러 스레드가 동시에 같은 메모리 위치에 접근할 때, 32비트 정수는 원자적 연산(atomic operations)을 지원합니다. 이는 병렬 처리 환경에서 데이터 경쟁 조건을 방지하는 데 중요합니다.
4. 셰이더 언어 호환성: Metal 셰이더 언어에서는 bool 타입이 있지만, 메모리 버퍼에서는 주로 uint 타입으로 표현됩니다. CPU와 GPU 사이의 데이터 구조를 일치시키면 버퍼 변환 없이 직접 데이터를 주고받을 수 있습니다.
5. 데이터 패킹 기회: GPU에서는 종종 여러 논리적 플래그를 하나의 32비트 정수에 비트 마스킹으로 패킹하여 메모리 사용량을 줄이고 캐시 효율성을 높입니다. 예를 들어, 하나의 UInt32 안에 최대 32개의 불리언 값을 저장할 수 있습니다.
6. 구조체 패딩 방지: CPU에서는 메모리 정렬을 위해 컴파일러가 자동으로 구조체에 패딩을 추가할 수 있습니다. 두 개의 Bool 타입(각 1바이트)이 있어도, 정렬을 위해 총 8바이트를 차지할 수 있습니다. GPU 구현에서는 명시적으로 32비트 타입을 사용하여 의도치 않은 패딩을 방지하고 메모리 레이아웃을 더 예측 가능하게 만듭니다.
7. 버퍼 공유 단순화: CPU와 GPU 간의 정확한 메모리 레이아웃 일치는 버퍼 공유를 단순화합니다. 이것은 특히 MTLBuffer와 같은 공유 메모리 버퍼를 사용할 때 중요합니다.

CPU 구현에서는 순차적으로 연산이 진행되기 때문에 각 프레임에서 업데이트 여부를 판단하기 위해 hasBeenUpdated 플래그가 필요했지만, GPU 구현에서는 모든 입자가 병렬로 처리되기 때문에 이 플래그가 필요하지 않았습니다. 

Metal 컴퓨트 셰이더의 작동 원리 

Metal 컴퓨트 셰이더는 GPU에서 범용 계산(GPGPU)을 수행하는 프로그램입니다. 랜더링과 같이 픽셀을 그리는 것이 아니라 데이터 처리에 GPU의 병렬 처리를 활용합니다. 

 

쓰레드 계층 구조 

Metal 컴퓨트 셰이더는 3단계 계층 구조로 실행됩니다.

• 스레드(Thread): 가장 작은 실행 단위로, 하나의 데이터 요소를 처리합니다. 입자 시뮬레이션에서는 하나의 셀/입자에 해당합니다.
• 스레드그룹(Threadgroup): 함께 실행되는 스레드의 집합입니다. 일반적으로 16x16 또는 8x8 같은 크기로 정의됩니다. 스레드그룹 내 스레드들은 로컬 메모리를 공유할 수 있습니다.
• 그리드(Grid): 모든 스레드그룹의 전체 집합으로, 처리할 전체 데이터 범위를 커버합니다.

쓰레드 식별 및 위치 지정

각 쓰레드는 자신의 위치를 식별할 수 있는 특별한 속성을 가집니다. 

kernel void updateParticles(/* 파라미터 */,
                          uint2 pos [[thread_position_in_grid]],
                          uint2 local_pos [[thread_position_in_threadgroup]],
                          uint2 group_pos [[threadgroup_position_in_grid]])

• `thread_position_in_grid`: 전체 그리드에서의 쓰레드 위치 
• `thread_position_in_threadgroup` : 쓰레드그룹 내에서의 상대적 위치
• `threadgroup_position_in_grid` : 그리드 내에서 쓰레드 그룹의 위치 

 

메모리 계층 구조

• 디바이스 메모리(Device Memory) : 모든 스레드가 접근 가능한 글로벌 메모리로, device 한정자로 표시됩니다.

kernel void updateParticles(device Particle* particles [[buffer(0)]])

• 스레드그룹 메모리(Threadgroup Memory): 같은 스레드그룹 내 스레드들이 공유하는 더 빠른 메모리입니다.

threadgroup float shared_data[256]

• 스레드 로컬 메모리(Thread Local Memory): 함수 내에서 선언된 지역 변수로, 해당 스레드만 접근 가능합니다.

kernel void my_kernel(...) {
    float local_var;  // 스레드 로컬 메모리
}

• 상수 메모리(Constant Memory) : 읽기 전용 데이터로, 모든 스레드가 접근 가능하지만 변경은 불가능합니다.

constant SimulationParams& params [[buffer(2)]]

 

병렬 실행 모델 ( SIMD )

컴퓨트 셰이더의 핵심은 SIMD(Single Instruction, Multiple Data) 실행 모델입니다:

• Same Instruction: 같은 스레드그룹 내 모든 스레드는 동일한 명령어를 실행합니다.
• Different Data: 각 스레드는 자신의 위치(thread_position_in_grid)에 따라 다른 데이터를 처리합니다.

이로 인해 조건문(if 문)은 성능 저하를 일으킬 수 있습니다. 왜냐하면 다른 경로를 실행하는 스레드들이 있을 경우, GPU는 모든 경로를 순차적으로 실행해야 하기 때문입니다.

입자 시뮬레이션 셰이더의 병렬 처리 흐름 

1. Metal 환경 설정

프로세스는 MetalViewController.swift에서 Metal 환경을 초기화하는 것으로 시작합니다.

private func setupMetal() {
    // 기본 Metal 디바이스 가져오기
    device = MTLCreateSystemDefaultDevice()
    
    // GPU로 작업을 보내기 위한 명령 큐 생성
    commandQueue = device.makeCommandQueue()
    
    // Metal 뷰 및 기타 리소스 설정
    // ...
}

• MTLCreateSystemDefaultDevice(): 시스템에서 사용 가능한 기본 GPU 디바이스를 반환합니다. 이 함수는 현재 기기의 GPU 하드웨어를 추상화하여 접근할 수 있게 해주는 MTLDevice 객체를 생성합니다.
• device.makeCommandQueue(): GPU에 명령을 전송하기 위한 큐를 생성합니다. 이 큐는 GPU에서 실행할 작업들을 순서대로 제출하는 통로 역할을 합니다.

2. 시뮬레이션 데이터 구조 정의 

Particle.swift에서는 CPU와 GPU 간에 공유될 핵심 데이터 구조가 정의됩니다.

struct Particle {
    var isEmpty: UInt32    // 메모리 정렬을 위해 Bool 대신 UInt32 사용
    var type: UInt32       // 입자 타입 (모래, 물 등)
}

struct Color {
    var r: UInt8
    var g: UInt8
    var b: UInt8
    var a: UInt8
}

struct SimulationParams {
    var width: UInt32
    var height: UInt32
    var frameCount: UInt32
}

이 구조체들은 Swift와 Metal 셰이더 코드 간의 메모리 레이아웃 일관성을 유지하기 위해 Shaders.metal에도 동일하게 정의됩니다.

3. 버퍼 설정

ParticleSimulator.swift의 초기화 과정에서 세 가지 주요 Metal 버퍼가 생성됩니다.

// 초기 입자 및 색상 배열 생성 (Cpu)
var particles = [Particle]()
var colors = [Color]()
        
for _ in 0..<(width * height) {
    particles.append(Particle(isEmpty: 1, type: 0))
    colors.append(Color(r: 0, g: 0, b: 0, a: 0))
}
    
// 입자 버퍼 생성
self.particlesBuffer = device.makeBuffer(bytes: particles,
                                      length: MemoryLayout<Particle>.stride * width * height,
                                      options: .storageModeShared)!

// 색상 버퍼 생성
self.colorsBuffer = device.makeBuffer(bytes: colors,
                                   length: MemoryLayout<Color>.stride * width * height,
                                   options: .storageModeShared)!

// 파라미터 버퍼 생성
self.paramsBuffer = device.makeBuffer(bytes: &self.params,
                                   length: MemoryLayout<SimulationParams>.size,
                                   options: .storageModeShared)!

• device.makeBuffer(): GPU에서 사용할 메모리 버퍼를 생성합니다. 이 함수는 CPU 메모리의 내용을 GPU 메모리로 복사하여 MTLBuffer 객체를 생성합니다.
  • bytes: 버퍼에 복사할 초기 데이터의 메모리 주소(cpu)를 지정합니다.
  • length: 버퍼의 바이트 크기
    • 각 격자 셀(입자)의 메모리 크기: MemoryLayout<Particle>.stride
      • 이것은 하나의 Particle 구조체가 메모리에서 차지하는 공간입니다.
      • stride는 GPU의 메모리 정렬 요구사항을 충족하도록 조정된 값입니다.
      • 예를 들어 실제 데이터는 6바이트만 필요하더라도, GPU 정렬 요구사항으로 인해 8바이트로 패딩될 수 있습니다.
    • 총 격자 수: width * height
      • 이것은 시뮬레이션에서 관리해야 할 입자의 총 개수입니다.
      • 100×100 그리드라면 10,000개의 입자를 저장할 공간이 필요합니다.
• options: 버퍼의 속성을 지정하는 옵션
• .storageModeShared: 이 옵션은 매우 중요한데, CPU와 GPU 모두 접근 가능한 버퍼를 생성하여 양방향 데이터 공유를 가능하게 합니다. 이를 통해 CPU에서 데이터를 수정하면 GPU에서도 즉시 해당 변경사항을 볼 수 있고, 그 반대도 마찬가지입니다.

Size와 Stride의 차이

MemoryLayout에서 size와 stride는 메모리 관리 측면에서 중요한 두 가지 개념입니다.

Size (크기)

• 구조체의 실제 데이터가 차지하는 순수한 바이트 수
• 모든 필드의 크기를 단순히 합한 값
• 패딩이나 정렬 요구사항을 고려하지 않은 "원시" 데이터 크기

Stride (보폭)

• 메모리에서 연속된 요소들 사이의 간격(바이트 단위)
• 메모리 정렬을 위한 패딩을 포함한 크기
• 배열에서 한 요소에서 다음 요소로 이동할 때 필요한 바이트 수

간단한 구조체를 예로 들어보겠습니다.

struct Example {
    var a: UInt8   // 1바이트
    var b: UInt32  // 4바이트
}

이 구조체의 메모리 레이아웃은 다음과 같습니다.

| a (1바이트) | 패딩 (3바이트) | b (4바이트) |

• Size: 5바이트 (1 + 4) - 실제 데이터만의 크기
• Stride: 8바이트 - 메모리 정렬을 위한 패딩 포함 크기

이러한 패딩이 필요한 이유는 대부분의 하드웨어, 특히 GPU가 4바이트 또는 16바이트 경계에 데이터가 정렬되어 있을 때 가장 효율적으로 접근할 수 있기 때문입니다.

StorageModeShared 옵션을 선택한 이유

1. CPU와 GPU 간의 양방향 데이터 교환 필요성

이 시뮬레이션에서는 CPU와 GPU가 모두 버퍼 데이터에 접근해야 합니다.

• CPU 접근 필요 시점
  • 사용자가 화면을 터치할 때 입자 생성
  • 시뮬레이션 초기화 및 리셋
  • 결과 데이터 분석이나 디버깅
• GPU 접근 필요 시점
  • 물리 법칙에 따른 입자 이동 계산
  • 충돌 감지 및 처리
  • 병렬 처리를 통한 성능 최적화

.storageModeShared는 CPU와 GPU가 동일한 메모리에 직접 접근할 수 있게 해줍니다. 이는 메모리 복사 작업 없이도 양쪽에서 같은 데이터를 읽고 쓸 수 있다는 의미입니다.

2. 실시간 상호작용 필요성

이 시뮬레이터는 사용자의 터치 입력에 즉각 반응해야 합니다:

func createParticlesAt(x: Int, y: Int, radius: Int = 2) {
    // CPU가 버퍼 메모리에 직접 접근하여 입자 생성
    let particlesPtr = particlesBuffer.contents().bindMemory(to: Particle.self, ...)
    let colorsPtr = colorsBuffer.contents().bindMemory(to: Color.self, ...)
    
    // ... 입자 생성 로직 ...
}

CPU가 입자를 생성한 직후, GPU는 즉시 해당 데이터에 접근하여 물리 계산을 수행해야 합니다. .storageModeShared를 사용하면 이러한 즉각적인 데이터 공유가 가능합니다.

3. 대안 모드들의 단점

다른 스토리지 모드를 사용했을 경우의 단점:

• .storageModePrivate: GPU만 접근 가능하므로, CPU에서 값을 읽거나 쓸 때마다 별도의 버퍼를 생성하고 복사해야 합니다. 이는 실시간 상호작용에 지연을 발생시킵니다.
• .storageModeManaged: CPU와 GPU 모두 접근 가능하지만, 동기화가 명시적으로 필요합니다. didModifyRange: 메서드를 호출하여 CPU의 변경사항을 GPU에 알려야 하고, 이는 추가적인 코드 복잡성을 초래합니다.

4. 성능과 편의성의 균형

.storageModeShared는 다른 모드에 비해 약간의 성능 저하가 있을 수 있지만, 코드 작성의 편의성과 실시간 상호작용의 직관성을 크게 향상시킵니다. 이 애플리케이션에서는 매 프레임마다 CPU와 GPU 간의 데이터 교환이 빈번하게 이루어지므로, 메모리 복사 오버헤드를 제거하는 것이 전체 성능에 도움이 됩니다.

결론적으로, 이 입자 시뮬레이터는 CPU와 GPU 간의 긴밀한 협력이 필요한 애플리케이션이므로, .storageModeShared 옵션이 최적의 선택입니다. 이를 통해 사용자 상호작용과 물리 시뮬레이션 간의 원활한 통합이 가능해집니다.

4. 컴퓨트 파이프라인 설정

ParticleSimulator.swift에서 Shaders.metal의 셰이더 함수를 사용하여 컴퓨트 파이프라인이 생성됩니다.

// 컴파일된 셰이더를 포함하는 기본 Metal 라이브러리 가져오기
let library = device.makeDefaultLibrary()!
// updateParticles 컴퓨트 함수 가져오기
let computeFunction = library.makeFunction(name: "updateParticles")!

do {
    // 함수를 사용하여 컴퓨트 파이프라인 상태 생성
    self.computePipelineState = try device.makeComputePipelineState(function: computeFunction)
} catch {
    fatalError("컴퓨트 파이프라인 생성 실패: \(error)")
}

• device.makeDefaultLibrary(): 앱 번들에 미리 컴파일된 Metal 셰이더 코드를 포함하는 라이브러리를 로드합니다. 이것은 앱 빌드 시 Shaders.metal 파일이 컴파일되어 만들어진 라이브러리입니다.
• library.makeFunction(name:): 라이브러리에서 특정 이름("updateParticles")의 셰이더 함수에 대한 참조를 생성합니다.
• device.makeComputePipelineState(function:): 지정된 컴퓨트 함수를 사용하여 GPU에서 실행할 컴퓨트 파이프라인 상태를 생성합니다. 이 파이프라인은 GPU가 입자의 물리적 동작을 계산할 때 사용됩니다.

5. 사용자 상호작용 및 입자 생성

사용자가 화면과 상호작용할 때, 터치 이벤트가 캡처되고 처리됩니다.

• ParticleSimulatorView가 제스처 이벤트를 캡처하고 MetalViewController에 위임합니다.
• MetalViewController는 이 이벤트를 처리하고 입자를 생성합니다.

func createParticlesAtPosition(position: CGPoint) {
    // 뷰 좌표를 그리드 좌표로 변환
    let gridX = Int((position.x / simulatorView.metalView.bounds.width) * CGFloat(gridWidth))
    let gridY = Int((position.y / simulatorView.metalView.bounds.height) * CGFloat(gridHeight))
    
    // 계산된 위치에 입자 생성
    simulator.createParticlesAt(x: gridX, y: gridY)
}

• ParticleSimulator는 공유 버퍼의 메모리를 직접 수정하여 입자를 생성합니다.

func createParticlesAt(x: Int, y: Int, radius: Int = 2) {
    // 버퍼 메모리에 대한 포인터 가져오기
    let particlesPtr = particlesBuffer.contents().bindMemory(to: Particle.self, capacity: gridWidth * gridHeight)
    let colorsPtr = colorsBuffer.contents().bindMemory(to: Color.self, capacity: gridWidth * gridHeight)
    
    // 반경 내에 입자 생성
    for offsetY in -radius...radius {
        for offsetX in -radius...radius {
            // 원형 반경 내에 있는지 확인
            let distance = sqrt(Float(offsetX * offsetX + offsetY * offsetY))
            if distance > Float(radius) {
                continue
            }
            
            let posX = x + offsetX
            let posY = y + offsetY
            
            // 경계 확인
            if posX >= 0 && posX < gridWidth && posY >= 0 && posY < gridHeight {
                let index = posY * gridWidth + posX
                
                // 모래 입자 생성
                if particlesPtr[index].isEmpty == 1 {
                    particlesPtr[index].isEmpty = 0
                    particlesPtr[index].type = ParticleType.sand.rawValue
                    
                    // 임의 변화가 있는 색상 설정
                    let variation = UInt8.random(in: 0...30)
                    colorsPtr[index] = Color(
                        r: UInt8(min(220 + Int(variation) - 15, 255)),
                        g: UInt8(min(180 + Int(variation) - 15, 255)),
                        b: UInt8(min(80 + Int(variation) - 15, 255)),
                        a: 255
                    )
                }
            }
        }
    }
}

• particlesBuffer.contents(): Metal 버퍼의 CPU 측 메모리 주소를 반환합니다. 이 함수는 MTLBuffer의 메모리 내용에 접근할 수 있는 원시 포인터를 제공합니다.
• bindMemory(to: Particle.self, capacity:): 원시 메모리 주소를 특정 타입(Particle)의 배열로 해석하도록 지시합니다. 이렇게 하면 원시 메모리를 Swift 구조체 배열처럼 접근할 수 있게 됩니다.
• UnsafeRawPointer(colorsBuffer.contents()): colorsBuffer.contents()로 얻은 메모리 주소를 UnsafeRawPointer 타입으로 변환합니다. 이는 타입 정보 없이 메모리의 원시 바이트에 접근할 수 있는 포인터를 생성하며, 주로 GPU와 데이터를 공유할 때 사용됩니다. 이 코드는 getColorBufferPointer() 함수에서 색상 버퍼의 메모리 주소를 반환할 때 사용됩니다.

6. GPU 시뮬레이션 실행

매 프레임마다 Metal 뷰는 MetalViewController의 draw(in:) 메서드를 호출하여 화면을 갱신합니다. 이것이 시뮬레이션 업데이트를 시작합니다.

func draw(in view: MTKView) {
    // GPU에서 시뮬레이션 업데이트
    simulator.update()
    
    // 시뮬레이션 결과로 텍스처 업데이트
    updateTextureFromColorBuffer()
    
    // 화면에 렌더링
    // ... 
}

시뮬레이션 업데이트(simulator.update())에서 GPU 연산이 이루어집니다.

func update() {
    // 프레임 카운터 증가
    frameCount += 1
    
    // 시뮬레이션 파라미터 업데이트
    var updatedParams = params
    updatedParams.frameCount = frameCount
    memcpy(paramsBuffer.contents(), &updatedParams, MemoryLayout<SimulationParams>.size)
    
    // GPU 명령을 위한 커맨드 버퍼 생성
    guard let commandBuffer = commandQueue.makeCommandBuffer() else { return }
    
    // 컴퓨트 셰이더를 위한 컴퓨트 인코더 생성
    guard let computeEncoder = commandBuffer.makeComputeCommandEncoder() else { return }
    
    // 컴퓨트 파이프라인 설정
    computeEncoder.setComputePipelineState(computePipelineState)
    
    // 셰이더 인덱스에 버퍼 바인딩
    computeEncoder.setBuffer(particlesBuffer, offset: 0, index: 0)
    computeEncoder.setBuffer(colorsBuffer, offset: 0, index: 1)
    computeEncoder.setBuffer(paramsBuffer, offset: 0, index: 2)
    
    // 그리드 및 스레드 차원 계산
    let threadsPerThreadgroup = MTLSize(width: 16, height: 16, depth: 1)
    let threadgroupsPerGrid = MTLSize(
        width: (gridWidth + threadsPerThreadgroup.width - 1) / threadsPerThreadgroup.width,
        height: (gridHeight + threadsPerThreadgroup.height - 1) / threadsPerThreadgroup.height,
        depth: 1
    )
    
    // 컴퓨트 셰이더 실행
    computeEncoder.dispatchThreadgroups(threadgroupsPerGrid, threadsPerThreadgroup: threadsPerThreadgroup)
    
    // 인코딩 종료
    computeEncoder.endEncoding()
    
    // 명령 실행
    commandBuffer.commit()
    commandBuffer.waitUntilCompleted()
}

• memcpy(paramsBuffer.contents(), &updatedParams, ...): C 표준 라이브러리 함수로, 메모리의 특정 영역을 다른 영역으로 복사합니다. 여기서는 업데이트된 파라미터 값을 GPU와 공유하는 버퍼에 복사합니다.
• commandQueue.makeCommandBuffer(): GPU에 전송할 명령들을 담을 새로운 커맨드 버퍼를 생성합니다. 이는 GPU에 보낼 작업 패키지와 같습니다.
• commandBuffer.makeComputeCommandEncoder(): 컴퓨트 셰이더에 사용할 명령들을 인코딩할 인코더를 생성합니다. 이 인코더는 GPU에서 실행할 계산 작업을 패키징합니다.
• computeEncoder.setComputePipelineState(): 사용할 컴퓨트 파이프라인 상태를 설정합니다. 이는 어떤 셰이더 함수가 실행될지 지정합니다.
• computeEncoder.setBuffer(buffer, offset:, index:): 셰이더 함수에서 사용할 버퍼를 설정합니다. index 매개변수는 셰이더 코드의 buffer(index) 속성과 일치합니다.
• MTLSize: 3D 공간에서의 크기를 나타내는 구조체로, 여기서는 GPU 스레드 그룹의 크기를 정의합니다.
• computeEncoder.dispatchThreadgroups(): 설정된 스레드그룹 크기로 GPU에 작업을 분배합니다. 이 함수는 실제로 GPU에서 병렬 컴퓨팅을 시작하는 명령입니다.
• computeEncoder.endEncoding(): 컴퓨트 인코더에 더 이상 명령을 추가하지 않겠다는 것을 알립니다.
• commandBuffer.commit(): 커맨드 버퍼의 모든 명령을 GPU에 제출하여 실행을 시작합니다.
• commandBuffer.waitUntilCompleted(): GPU가 모든 명령 실행을 완료할 때까지 CPU가 대기하도록 합니다. 이는 동기식 실행을 보장합니다.

7. GPU 셰이더 실행

핵심 물리 연산은 Shaders.metal의 updateParticles 커널 함수에서 GPU에서 실행됩니다.

kernel void updateParticles(device Particle* particles [[buffer(0)]],
                           device Color* colors [[buffer(1)]],
                           constant SimulationParams& params [[buffer(2)]],
                           uint2 pos [[thread_position_in_grid]]) {
    
    // 경계 확인
    if (pos.x >= params.width || pos.y >= params.height) {
        return;
    }
    
    // 입자 인덱스 가져오기
    uint index = getIndex(pos, params.width);
    
    // 빈 공간 건너뛰기
    if (particles[index].isEmpty == 1) {
        return;
    }
    
    // 자연스러운 흐름을 위한 짝수/홀수 프레임 번갈아가며 처리
    bool evenFrame = (params.frameCount % 2) == 0;
    
    // 모래 입자 처리
    if (particles[index].type == ParticleType::Sand) {
        // 입자를 아래로 또는 대각선으로 이동 시도
        // ... [모래의 물리 로직]
    }
    
    // 다른 입자 타입 처리 (향후 확장)
}

셰이더 실행에 관한 주요 사항:

• kernel 키워드: Metal 셰이더 언어에서 GPU에서 병렬로 실행될 함수를 정의합니다.
• [[buffer(0)]], [[buffer(1)]] 등의 구문: 컴퓨트 인코더에서 설정한 버퍼 인덱스와 매핑됩니다.
• [[thread_position_in_grid]]: 스레드의 위치를 제공하며, 이는 그리드 셀에 해당합니다.
• device 한정자: 이 포인터가 디바이스(GPU) 메모리를 가리킨다는 것을 나타냅니다.
• constant 한정자: 읽기 전용 데이터를 가리킵니다.

셰이더는 모래 입자에 대한 간단한 물리법칙을 구현하여:

• 가능하면 직선으로 아래로 떨어지게 합니다.
• 바로 아래가 막혀있으면 대각선으로 떨어지게 합니다.
• 더 자연스러운 흐름을 위해 짝수/홀수 프레임에서 왼쪽/오른쪽을 번갈아 먼저 확인합니다.

moveParticle 헬퍼 함수는 실제 입자 이동을 처리합니다:

void moveParticle(device Particle* particles,
                  device Color* colors,
                  uint fromIndex,
                  uint toIndex) {
    // 새 위치로 입자 데이터 복사
    particles[toIndex] = particles[fromIndex];
    colors[toIndex] = colors[fromIndex];
    
    // 이전 위치 지우기
    particles[fromIndex].isEmpty = 1;
    particles[fromIndex].type = ParticleType::Empty;
    colors[fromIndex] = Color{0, 0, 0, 0};
}

이 모든 것이 모든 그리드 위치에 대해 병렬로 실행되어 CPU 기반 시뮬레이션에 비해 엄청난 성능 향상을 제공합니다.

8. 결과 표시

GPU 연산이 완료된 후, 시뮬레이션 결과를 화면에 표시해야 합니다:

private func updateTextureFromColorBuffer() {
    // 색상 버퍼 포인터 가져오기
    let colorBufferPtr = simulator.getColorBufferPointer()
    
    // 새로운 색상 데이터로 텍스처 업데이트
    texture.replace(
        region: MTLRegionMake2D(0, 0, gridWidth, gridHeight),
        mipmapLevel: 0,
        withBytes: colorBufferPtr,
        bytesPerRow: gridWidth * MemoryLayout<Color>.stride
    )
}

• simulator.getColorBufferPointer(): 색상 버퍼의 메모리 주소를 가져옵니다. 이 함수는 UnsafeRawPointer(colorsBuffer.contents())를 반환합니다.
• texture.replace(region:, mipmapLevel:, withBytes:, bytesPerRow:): 지정된 메모리의 데이터로 텍스처의 특정 영역을 업데이트합니다. 이는 시뮬레이션 결과를 화면에 표시하기 위한 텍스처에 복사합니다.
• MTLRegionMake2D(): 2D 텍스처의 특정 영역을 정의하는 구조체를 생성합니다.

draw(in:) 메서드의 나머지 부분은 렌더링을 처리합니다:

// 현재 drawable 및 렌더 패스 디스크립터 가져오기
guard let drawable = view.currentDrawable,
      let renderPassDescriptor = view.currentRenderPassDescriptor,
      let commandBuffer = commandQueue.makeCommandBuffer() else {
    return
}

// 렌더 인코더 생성
guard let renderEncoder = commandBuffer.makeRenderCommandEncoder(descriptor: renderPassDescriptor) else {
    return
}

// 렌더 상태 및 리소스 설정
renderEncoder.setRenderPipelineState(renderPipelineState)
renderEncoder.setVertexBuffer(vertexBuffer, offset: 0, index: 0)
renderEncoder.setFragmentTexture(texture, index: 0)

// 쿼드 그리기 (전체 화면 텍스처를 위한 4개 정점)
renderEncoder.drawPrimitives(type: .triangleStrip, vertexStart: 0, vertexCount: 4)

// 인코딩 종료 및 표시
renderEncoder.endEncoding()
commandBuffer.present(drawable)
commandBuffer.commit()

• view.currentDrawable: 화면에 출력할 현재 드로어블 객체를 반환합니다. 이는 렌더링 결과가 표시될 대상입니다.
• view.currentRenderPassDescriptor: 현재 렌더 패스에 대한 설정을 담고 있는 디스크립터를 반환합니다.
• commandBuffer.makeRenderCommandEncoder(): 렌더링 명령을 인코딩할 인코더를 생성합니다.
• renderEncoder.setRenderPipelineState(): 렌더링에 사용할 파이프라인 상태를 설정합니다. 이는 어떤 셰이더가 사용될지 결정합니다.
• renderEncoder.setVertexBuffer(): 정점 셰이더에서 사용할 버퍼를 설정합니다.
• renderEncoder.setFragmentTexture(): 프래그먼트 셰이더에서 사용할 텍스처를 설정합니다.
• renderEncoder.drawPrimitives(): 실제 그리기 작업을 수행합니다. 여기서는 4개 정점으로 구성된 삼각형 스트립을 그립니다.
• commandBuffer.present(drawable): 렌더링된 결과를 화면에 표시합니다.
• commandBuffer.commit(): 모든 렌더링 명령을 GPU에 제출합니다.

Shaders.metal에 정의된 정점 및 프래그먼트 셰이더는 실제 렌더링을 처리합니다.

vertex VertexOutput vertexShader(uint vertexID [[vertex_id]],
                              constant float *vertices [[buffer(0)]]) {
    // 전체 화면 쿼드를 위한 정점 변환
    // ...
}

fragment float4 fragmentShader(VertexOutput in [[stage_in]],
                             texture2d<float> texture [[texture(0)]]) {
    // 텍스처 샘플링 및 색상 반환
    constexpr sampler textureSampler(mag_filter::nearest, min_filter::nearest);
    return texture.sample(textureSampler, in.texCoord);
}

• vertex 키워드: 정점 셰이더 함수를 정의합니다. 이 함수는 각 정점마다 한 번씩 호출됩니다.
• [[vertex_id]]: 현재 처리 중인 정점의 인덱스를 제공합니다.
• fragment 키워드: 프래그먼트(픽셀) 셰이더 함수를 정의합니다. 이 함수는 화면의 각 픽셀마다 호출됩니다.
• [[stage_in]]: 정점 셰이더에서 전달된 보간된 값을 받는 매개변수임을 나타냅니다.
• constexpr sampler: 텍스처 샘플링 방법을 정의합니다. nearest 필터는 가장 가까운 텍셀의 값을 사용합니다.
• texture.sample(): 지정된 텍스처 좌표에서 텍스처 값을 샘플링합니다.

 

GPU 연산 흐름 요약

1. 초기화
   • Metal 디바이스 및 명령 큐 생성
   • 데이터 구조 정의 (Particle, Color, SimulationParams)
   • CPU와 GPU 모두 접근 가능한 공유 버퍼 생성
   • 셰이더 함수로 컴퓨트 파이프라인 설정
2. 사용자 상호작용
   • 터치 이벤트 캡처
   • 그리드 위치 계산
   • 버퍼 메모리 직접 수정하여 입자 생성
3. GPU 시뮬레이션
   • 커맨드 버퍼 및 컴퓨트 인코더 생성
   • 입자, 색상 및 파라미터 버퍼 바인딩
   • 적절한 스레드 차원으로 컴퓨트 셰이더 실행
   • GPU가 병렬로 물리 계산 실행
4. 셰이더 실행
   • 각 GPU 스레드가 하나의 그리드 위치 처리
   • 입자 물리학 구현 (낙하, 충돌 감지)
   • 공유 버퍼에서 입자 위치 및 색상 직접 업데이트
5. 렌더링
   • 색상 버퍼에서 텍스처 업데이트
   • 텍스처로 전체 화면 쿼드 렌더링
   • 화면에 표시

 

GPU 컴퓨트 셰이더 방법으로 전환 후 변화 

GPU 방식으로 전환한 결과 CPU 사용량이 99%에서 18%로 크게 감소하고 메모리 사용량이 약간 증가하였습니다. 

CPU 사용량 감소 이유

• 계산 부하 이전: 모든 입자 계산 작업이 CPU에서 GPU로 옮겨져 CPU 부하가 크게 감소했습니다.
• 병렬 처리: GPU의 수많은 코어가 동시에 계산을 처리하므로, CPU가 순차적으로 처리할 때보다 효율적입니다.
• CPU는 조정 역할만: CPU는 이제 명령 버퍼 설정과 GPU에 작업 전달하는 조정 역할만 수행합니다.

메모리 사용량 증가 이유

• 버퍼 중복: GPU 처리를 위해 CPU와 GPU 모두 접근 가능한 공유 메모리 버퍼가 필요합니다.
• 정렬된 데이터 구조: GPU 최적화를 위해 Bool 대신 UInt32를 사용함으로써 메모리 사용량이 증가했습니다.
• 메탈 프레임워크 오버헤드: Metal API 자체가 내부적으로 관리하는 리소스 캐시와 명령 큐가 추가 메모리를 사용합니다.
• 명령 버퍼: GPU 작업을 위한 명령 버퍼와 파이프라인 상태 객체가 추가적인 메모리를 소비합니다.

성능과 메모리의 트레이드오프

이러한 메모리 증가는 일반적으로 허용되는 트레이드오프입니다. 약간의 메모리 증가를 통해 CPU 부하를 대폭 감소시키고 전체적인 애플리케이션 성능을 향상시켰습니다.

1. CPU가 여유로워져 UI 반응성과 다른 시스템 작업이 더 원활해집니다.
2. 배터리 효율성이 향상될 수 있습니다 (GPU는 특정 작업에 대해 CPU보다 에너지 효율적).
3. 더 많은 입자를 처리할 수 있는 확장성을 확보했습니다.