지난 시간 FFP에서 고전 SIMD의 등장까지 알아보았습니다. 이제 SIMD에서 SIMT로 넘어가던 DirectX9 ~ DirectX10 시절을 알아보도록 하겠습니다.
개요
하나의 명령어로 여러 데이터를 동시에 처리하던 고전 SIMD는 구조적인 한계가 명확했습니다.
데이터 길이가 workload에 맞지 않으면 cost 낭비가 발생
스레드 개념이 없음.
전체 SIMD가 stall 될 정도의 분기 처리
gather/scatter 미지원 => 다양한 메모리 주소 접근이 불가능
이 구조는 단순한 벡터 연산에는 강력하지만, 픽셀·조명·텍스처·조건문이 많은 실제 그래픽스 workload에는 치명적이었다.
CPU는 SIMD 폭을 확장하고 ISA를 강화하는 방식으로 발전 방향을 잡았습니다.
벡터 폭 증가
gather/scatter 지원
mask 연산 강화
분기 최소화 등
즉, CPU는 SIMD를 더 강한 벡터 처리 장치로 발전시키는 방향을 택했다.
SIMT의 등장
하지만 GPU는 완전히 다른길을 선택했습니다. 그것은 바로 SIMD를 Thread로 추상화하여 재해석하는 방법입니다. GPU는 수천 ~ 수만 개의 Pixel Fragment를 병렬 처리해야 했기 때문에 PC(ProgramCounter)를 Thread마다 하나씩 둘 수 없었습니다. PC를 Thread마다 하나씩 두면 면적과 전력소모가 폭발적으로 증가하기 때문이었습니다.
그래서 GPU는 SIMD를 다음과 같이 재해석했습니다.
"SIMD lane을 Thread처럼 보이게 만들고, 프로그래머에게는 각 lane이 독립적인 스레드처럼 보이게 하자."
이 방식이 바로 SIMT 모델입니다.
about SIMT
Warp란 무엇인가?
CPU는 Thread마다 PC가 하나씩 있습니다. 그렇기 때문에 완전한 MIMD 실행 모델로 각자 다른 위치에서 다른 코드를 실행할 수 있습니다. 하지만 GPU의 Warp의 구조는 정반대입니다.
예를 들어 Warp = 32 thread가 있다면,
32개의 스레드는 각자 독립적으로 보이지만, Warp 전체가 PC 1개, 그렇기 때문에 32개의 스레드는 한번에 동일한 명령어만 실행 가능.
각 Thread는 고유한 Register 값을 갖지만 명령어 흐름은 공유.
이렇게 동일한 명령어를 동시에 실행하는 한 단계가 Lockstep입니다.
Branch Divergence
이 조건을 보면 Lockstep의 의미가 바로 드러납니다.
아래와 같은 조건이 나왔다고 합시다.
if (x < 16)
A 실행
else if(x < 32)
B 실행
thread 0 ~ 15 => A 실행
thread 16 ~ 31 => B 실행
이렇게 CPU 에서는 스레드 별로 명령어를 다르게 실행할 수 있기 때문에 각 스레드에서 A와 B를 나눠서 실행하게 됩니다. 하지만 GPU는 Warp는 하나의 PC만 있으므로 두 경로를 동시에 실행할 수 없습니다.
그렇기에 SIMT는 다음과 같은 방식으로 처리합니다.
Thread 32개를 컨트롤 하는 Warp의 PC를
A 경로로 이동
true thread만 활성 (mask = on)
false thread는 비활성 (mask = off)
Warp 실행(A만 실행 됨)
B 경로로 이동
false thread만 활성 (mask = on)
true thread는 비활성 (mask = off)
Warp 실행 (B만 실행 됨)
이렇게 한 warp안에서 분기가 갈리면 warp는 A와 B가 직렬화 되어 순차적으로 두 번 실행합니다. 이게 바로 lockstep의 실행제약인 Branch Divergence 입니다. 고전 SIMD에서는 이러한 직렬화 때문에 Stall이 발생하였지만 Warp단위로 분리된 SIMT의 경우에는 MIMD처럼 동작하기 위해 Warp(Thread 묶음, SIMD Lane)를 다시 Thread로 추상화하여 다루는 매커니즘을 갖게 되었습니다.
맺음말
SIMT의 등장은 고전 SIMD의 한계를 '극복'했다기보다는, 그래픽스라는 특수한 workload를 처리하기 위한 현실적인 타협이었습니다.
SIMD lane을 Thread처럼 추상화하여 프로그래머에게는 MIMD처럼 보이게 만들고
내부에서는 여전히 lockstep 기반 SIMD의 효율을 유지하며
Massive Parallel Pixel Workload를 감당할 수 있도록 한 구조
즉, SIMT는 "GPU스럽게 동작하는 MIMD의 환상"을 만들어낸 모델이라고 할 수 있습니다.
Branch divergence는 여전히 존재하며 성능을 떨어뜨립니다. 하지만 FFP 시대에는 상상할 수 없었던 복잡한 조명, 그림자, 포스트 프로세싱, 물리 기반 셰이딩을 GPU가 처리할 수 있게 된 것도 결국 이 SIMT 모델 덕분입니다.
다음 글에서는 이 SIMT 모델이 어떻게 구체적으로 구현되었는지, 그리고 DirectX10~11, UE3~UE4 시대를 지나면서 하드웨어와 셰이더 모델이 어떤 방향으로 진화했는지를 이어서 다뤄보겠습니다.
2000년대 초반의 GPU는 FFP(Fixed Function Pipeline)중심이었고, 지금처럼 복잡한 셰이더 기반 렌더링이 불가능했습니다. VRAM 용량도 32~64MB 수준이라 SSAA나 고급 조명 기법을 적용하기 어려웠죠. 이 글에서는 당시 GPU의 한계를 살펴보고, 어떻게 SIMD 기반의 Programmable Pipeline 으로 발전했는지 정리합니다.
MSAA (Multi-Sampling Anti-Aliasing)
MSAA vs EQAA 샘플 패턴 비교 이미지
Pixel Boundary
한 픽셀을 크게 확대한 사각형 틀. (픽셀은 점이 아니라 “면적”을 가진 사각형입니.)
Color Sample Location
픽셀 내부의 특정 지점(Point)에서 실제 색상(Color)을 샘플링하는 위치. 여러 Color Sample을 블렌딩하여 최종 픽셀 색을 만든다.
Coverage Sample Location
폴리곤이 픽셀을 얼마나 덮고 있는지(Coverage) 판단하는 지점. 각 샘플이 폴리곤 내부면 1, 외부면 0에 해당하며, 이를 평균내어 블렌딩 가중치(강도)로 사용한다. Color Sample보다 계산 비용이 훨씬 낮다.
초기 GPU의 VRAM 용량이 32 ~ 64MB 하던 시절에는 픽셀 단위로 복잡한 처리를 수행할 수 없었습니다. 그래서 이미지 전체를 더 높은 해상도로 여러 번 렌더링하는 SSAA(Super-Sampling Anti-Aliasing)을 다운 샘플링을 하는 방식이 유일한 선택지였습니다. 해당 방식은 해상도 배율 x 샘플 수 만큼의 비용증가가 있었고, 이는 픽셀 셰이딩, 텍스쳐 샘플링, ROP(Blend / Depth / Color Write) 등 '픽셀 이후 단계'의 비용을 전부 O(S) 수준으로 폭증시켰습니다. 즉, 샘플 수만큼만 비용이 증가하는 것이 아니라, 렌더타겟 자체도 커지기 때문에 파이프라인의 대부분이 정비례로 비싸지는 구조였습니다.
하지만 VRAM 용량이 어느정도 여유가 생기자, 픽셀 내부의 Color Sample과 Coverage Sample을 저장할 수 있는 구조가 가능해졌고, MSAA는 Pixel Shader는 한 번만 수행하고, Coverage Test와 Color Resolve만 샘플 수에 따라 반복하는 방식으로 SSAA 대비 큰 폭의 최적화를 달성했습니다.
Fixed Function Pipeline > Programmable Pipeline 의 과도기
FFP(Fixed Function Pipeline)
Transform & Lighting
Texture Stage State 기반 2~3단 텍스처 처리
멀티패스 조명
DOT3 Bump Mapping(FFP LOD 바탕)
Lightmap 기반의 Static Lighting
UI, HUD, 간단한 머터리얼 > FFP로 렌더링
하지만 UE2 후반 버전은...
Programmable Pipeline 일부 적용
Shader-driven Material Effects
Normal Mapping
Specular Mask
Detail Normal Layer
Color Modulation
Hardware Skinning
NVIDIA/ATI 전용 Shader Path
이것이 가능해진 이유는 GPU의 주요한 변화 덕분.
Scalar ALU vs SIMD Lane
기존
Scalar ALU에 단순한 microcode 실행기를 사용하여 float4연산은 4개의 스칼라 연산을 순차적으로 연산하도록 실행, 단지 GPU에 Scalar ALU를 많이 담아서 단순 병렬처리를 수행했을 뿐.
SIMD(Single Instruction Multiple Data stream)의 등장
하나의 명령어(Single Instruction)를 Scalar ALU의 묶음인 SIMD Lane에 전달하여 연산에 사용될 Multi-Data를 레지스터/버퍼(L1/L2, GPR)에 담아 연산하도록 실행
요즘 레이트레이싱에 관련된 작업을 시작하게 되어 정리차원에서 블로깅하게 되었습니다. 현재는 TinyRaytracer를 마개조 중이었고, 이번 글은 깃헙의 TinyRaytracer를 기반으로 레이트레이싱 기법에 대해 간략하게 설명해보도록 하겠습니다. 최적화 기법은 추후에 다룰 예정입니다. 다른 포스팅에서 확인 부탁드립니다. (fresnel equation이 추가되었습니다.)
ray_intersect함수의 경우, 레이의 original position, direction을 받아서 해당 오브젝트의 충돌검출(구와 직선의 충돌 알고리즘)과 함께 해당 오브젝트까지의 거리를 t0로 넘겨줍니다.
2. 기능 설명
- render
void render(const std::vector<const sdf_model*> &models, const std::vector<Light> &lights) {
const int width = 1024;
const int height = 768;
const float fov = M_PI/3.;
std::vector<Vec3f> framebuffer(width*height);
#pragma omp parallel for
for (size_t j = 0; j<height; j++) { // actual rendering loop
for (size_t i = 0; i<width; i++) {
float dir_x = (i + 0.5) - width/2.;
float dir_y = -(j + 0.5) + height/2.; // this flips the image at the same time
float dir_z = -height/(2.*tan(fov/2.));
framebuffer[i+j*width] = cast_ray(Vec3f(0,0,0), Vec3f(dir_x, dir_y, dir_z).normalize(), models, lights);
}
}
std::ofstream ofs; // save the framebuffer to file
ofs.open("./out.ppm",std::ios::binary);
ofs << "P6\n" << width << " " << height << "\n255\n";
for (size_t i = 0; i < height*width; ++i) {
Vec3f &c = framebuffer[i];
float max = std::max(c[0], std::max(c[1], c[2]));
if (max > 1) c = c * (1. / max);
for (size_t j = 0; j<3; j++) {
ofs << (char)(255 * std::max(0.f, std::min(1.f, framebuffer[i][j])));
}
}
ofs.close();
}
카메라에서 레이를 발사하여 frame buffer에 기록하고, 그 데이터를 ppm 파일에 저장합니다.
- cast_ray
Vec3f cast_ray(const Vec3f &orig, const Vec3f &dir, const std::vector<Sphere> &spheres, const std::vector<Light> &lights, size_t depth=0) {
Vec3f point, N;
Material material;
if (depth>4 || !scene_intersect(orig, dir, spheres, point, N, material)) {
return Vec3f(0.2, 0.7, 0.8); // background color
}
Vec3f refract_color(0.f, 0.f, 0.f);
// compute fresnel
float kr;
fresnel(dir, N, material.refractive_index, kr);
// compute refraction if it is not a case of total internal reflection
if (kr < 1) {
Vec3f refract_dir = refract(dir, N, material.refractive_index).normalize();
Vec3f refract_orig = refract_dir * N < 0 ? point - N * EPSILON : point + N * EPSILON;
refract_color = cast_ray(refract_orig, refract_dir, models, lights, depth + 1);
}
Vec3f reflect_dir = reflect(dir, N).normalize();
Vec3f reflect_orig = reflect_dir * N < 0 ? point - N * EPSILON : point + N * EPSILON; // offset the original point to avoid occlusion by the object itself
Vec3f reflect_color = cast_ray(reflect_orig, reflect_dir, models, lights, depth + 1);
float diffuse_light_intensity = 0, specular_light_intensity = 0;
for (size_t i=0; i<lights.size(); i++) {
Vec3f light_dir = (lights[i].position - point).normalize();
float light_distance = (lights[i].position - point).norm();
Vec3f shadow_orig = light_dir*N < 0 ? point - N*EPSILON : point + N*EPSILON; // checking if the point lies in the shadow of the lights[i]
Vec3f shadow_pt, shadow_N;
Material tmpmaterial;
if (scene_intersect(shadow_orig, light_dir, spheres, shadow_pt, shadow_N, tmpmaterial) && (shadow_pt-shadow_orig).norm() < light_distance)
continue;
diffuse_light_intensity += lights[i].intensity * std::max(0.f, light_dir*N);
specular_light_intensity += powf(std::max(0.f, -reflect(-light_dir, N)*dir), material.specular_exponent)*lights[i].intensity;
}
return material.diffuse_color * diffuse_light_intensity * material.albedo[0]
+ Vec3f(1., 1., 1.)*specular_light_intensity * material.albedo[1]
+ reflect_color * material.albedo[2] * kr
+ refract_color * material.albedo[3] * (1 - kr);
}
레이를 발사하는 함수입니다. 먼저 장면에 충돌시켜 충돌된 위치(point)와 해당 point의 normal값을 구합니다. 그 후, 위 함수는 세 연산으로 나뉩니다.
첫번째는 레이 연산입니다. 레이 연산은 다시 세가지로 나뉘는데 fresnel equation, 반사(reflection), 투과(refraction입니다. 이 부분을 설명하기엔 길어지므로 아래의 각 파트에서 다루도록 하겠습니다.
두번째는 빛 연산입니다. 먼저 변수에 대한 설명을 진행하도록 하겠습니다.
diffuse_light_intensity는 분산광의 정도를 저장하는 변수입니다.
specular_light_intensity는 반사광의 정도를 저장하는 변수입니다.
light_dir는 빛이 diffuse light이므로 빛의 방향이 됩니다.
light_distance는 빛에서부터 point까지의 거리로 shadow casting을 할 때, 앞에 있는 물체를 검출하기 위해 사용할 값입니다. 앞에 물체가 있다면 light값을 적용하지 않습니다.
shadow_origin은 point의 위치를 EPSILON으로 보정한 값입니다. point의 앞에서 발사한 빛인지 뒤에서 발사한 빛인지를 내적으로 검출하여 EPSILON으로 보정하여 light_distance 비교연산을 할 때 오차를 줄여줍니다. light_distance 비교연산은 변수 설명이 끝난 후 설명하겠습니다.
shadow_pt는 빛에서 point로 레이를 발사하는 과정에서 부딪힌 point입니다. light_distance보다 짧은 곳에서 충돌했다면 point 앞에 물체가 있는 것으로 음영에 diffuse_light_intensity와 이 필요합니다.
tmpmaterial은 scene_intersect에 넣어주기 위한 용도로 TinyRaytracer에서는 사용하지 않습니다.
마지막으로 light_distance 비교연산을 진행하고 그 결과에 따라 light_intensity를 적용하게 됩니다. 빛에서부터 해당 point까지 scene_intersect를 진행하여 light_distance보다 짧은 거리에 위치한 물체가 있다면 light_intensity 연산을 하지않게 되는 것이지요.
세번째는 위 연산에 따라 색상을 혼합해주는 연산입니다. 여기서 주목해야할 부분은 reflect_color와 refract_color에 fresnel equation의 결과값인 kr을 적용해주는 부분입니다. 아래 fresnel편에서 자세히 다루도록 하겠습니다.
- scene_intersect
bool scene_intersect(const vec3 &orig, const vec3 &dir, const std::vector<Sphere> &spheres, vec3 &hit, vec3 &N, Material &material) {
float spheres_dist = std::numeric_limits<float>::max();
for (const Sphere &s : spheres) {
float dist_i;
if (s.ray_intersect(orig, dir, s, dist_i) && dist_i < spheres_dist) {
spheres_dist = dist_i;
hit = orig + dir*dist_i;
N = (hit - s.center).normalize();
material = s.material;
}
}
float checkerboard_dist = std::numeric_limits<float>::max();
if (std::abs(dir.y)>1e-3) { // avoid division by zero
float d = -(orig.y+4)/dir.y; // the checkerboard plane has equation y = -4
vec3 pt = orig + dir*d;
if (d>1e-3 && fabs(pt.x)<10 && pt.z<-10 && pt.z>-30 && d<spheres_dist) {
checkerboard_dist = d;
hit = pt;
N = vec3{0,1,0};
material.diffuse_color = (int(.5*hit.x+1000) + int(.5*hit.z)) & 1 ? vec3{.3, .3, .3} : vec3{.3, .2, .1};
}
}
return std::min(spheres_dist, checkerboard_dist)<1000;
}
레이를 원점으로부터 특정 방향으로 진행시킬 때의 hit point와 해당 point의 normal값을 검출하는 함수입니다.
Vec3f refract(const Vec3f &I, const Vec3f &N, const float eta_t, const float eta_i=1.f) { // Snell's law
float cosi = - std::max(-1.f, std::min(1.f, I*N));
if (cosi<0) return refract(I, -N, eta_i, eta_t); // if the ray comes from the inside the object, swap the air and the media
float eta = eta_i / eta_t;
float k = 1 - eta*eta*(1 - cosi*cosi);
return k<0 ? Vec3f(1,0,0) : I*eta + N*(eta*cosi - sqrtf(k)); // k<0 = total reflection, no ray to refract. I refract it anyways, this has no physical meaning
}
refract는 snell's law에 의해 구현되었습니다.
우선 \(cos i\)를 확인하여 광선이 매질 안에서 발사되는 것인지를 확인합니다. 만약 \(cos i\)가 0보다 작으면 매질 안에서 발사되는 것이므로 eta_i와 eta_t를 교환해 다시 refract를 진행해줍니다.
매개변수 중 eta_t는 투과를 진행할 물체의 투과율로 영어로는 ior(index of refraction)라고도 부릅니다. eta_i는 광선이 있는 곳의 투과율입니다.
