3장에서는 cuda memory와 관련한 내용! memory allocation, memory set (초기화), memory free 그리고 host -> device, device->host 로의 메모리 이동 (memory copy) 의 내용이 포함되어 있다!

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

__global__ void printData(int* _dDataPtr){
    printf("%d", _dDataPtr[threadIdx.x]);
}

__global__ void setData(int* _dDataPtr){
    _dDataPtr[threadIdx.x] = 2;
}

int main(void){
    cudaSetDevice(15);
    int data[10] = {0};
    for (int i = 0; i < 10; i++) data[i] = 1;

    int* dDataPtr;
    cudaMalloc(&dDataPtr, sizeof(int)*10);
    cudaMemset(dDataPtr, 0, sizeof(int)*10);

    printf("Data in device: ");
    printData <<<1, 10>>> (dDataPtr);

    cudaMemcpy(dDataPtr, data, sizeof(int)*10, cudaMemcpyHostToDevice);
    printf("\nHost -> Device: ");
    printData <<<1, 10>>> (dDataPtr);

    setData <<<1, 10>>> (dDataPtr); // 2로 set 
    cudaMemcpy(data, dDataPtr, sizeof(int)*10, cudaMemcpyDeviceToHost);
    printf("\nDevice -> Host: ");
    for (int i = 0; i < 10; i++) printf("%d", data[i]);

    cudaFree(dDataPtr);
}

우선 3-3 코드를 보면 위와 같다. cudaMalloc으로 메모리 얼록을 해주고, cudaMemset으로 0으로 초기화해준다. 그 다음 데이터 찍어본 후, cudaMemcpy를 통해 (dest, src, ...) 메모리 복사. 마지막으로 cudaFree.

아래는 함수 원형들.

헤더 파일

#include <cuda_runtime.h>

함수 원형들

1. cudaMalloc

cudaError_t cudaMalloc(void** devPtr, size_t size);

• GPU 메모리를 할당합니다
• devPtr: 할당된 GPU 메모리 주소를 저장할 포인터의 주소
• size: 할당할 바이트 수
• 반환값: cudaError_t 타입의 에러 코드

2. cudaMemset

cudaError_t cudaMemset(void* devPtr, int value, size_t count);

• GPU 메모리를 특정 값으로 초기화합니다
• devPtr: 초기화할 GPU 메모리 주소
• value: 설정할 값 (0-255 범위의 정수)
• count: 설정할 바이트 수
• 반환값: cudaError_t 타입의 에러 코드

3. cudaFree

cudaError_t cudaFree(void* devPtr);

• GPU 메모리를 해제합니다
• devPtr: 해제할 GPU 메모리 주소
• 반환값: cudaError_t 타입의 에러 코드

4. cudaMemcpy

cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind);

• 메모리 간 데이터를 복사합니다 (Host ↔ Device)
• dst: 목적지 메모리 주소
• src: 원본 메모리 주소
• count: 복사할 바이트 수
• kind: 복사 방향을 지정하는 열거형
  • cudaMemcpyHostToDevice: Host → Device
  • cudaMemcpyDeviceToHost: Device → Host
  • cudaMemcpyDeviceToDevice: Device → Device
  • cudaMemcpyHostToHost: Host → Host
• 반환값: cudaError_t 타입의 에러 코드

모든 함수는 성공 시 cudaSuccess를 반환하며, 에러 발생 시 해당하는 에러 코드를 반환합니다.

3-5, 3-6의 vector add 실습 코드도 재밌음! 특히 3-5 skeleton 부터 시작해서 작성해보면 (아주 친절함!) 뭔가 뿌듯함을 느낄 수 있다! (3-6 코드의 벡터 크기 늘렸을 시 print가 너무 자주 찍히기 때문에 print는 주석처리 하는 거 추천~ 다만 if-else 구문으로 틀렸으면 틀렸다 한 마디만 적어주도록 하자~)

아차차, 그리고 너무 오랜만에 컴파일 해보아서,, 여러 파일 한 번에 컴파일 해야 하는 정보 밑에 추가. 책을 따라가다보면 DS_timer.cpp라는 파일을 같이 사용하게 되는데 (시간 측정용), 해당 .cpp 없이 .cu 파일만 컴파일하려니 오류가 생겨서! 클로드의 도움을 받아 무사히 잘 컴파일 후 실행해보았다.

nvcc -o vector_add_timer 3-6.vector_add_timer.cu DS_timer.cpp 

실습 재밌었다리~

코드만 있는 leetcuda보다 설명과 실습코드가 같이 들어간 책이... 역시 친절하다는 것을 느낌... 책 최고!

메모리와 관련해서는 궁금한 건 별로 없어서 호로록 호로록 넘어갔다. 실습 코드 적당히 있어서 따라해보는 것도 재밌었다.

CUDA의 특징적인 스레드 계층... 이런 거 이제 chapter 4부터 시작되는데 기대!

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

__global__ void helloCUDA(void){
    printf("Hello CUDA from GPU!\n");
}

int main(void){
    cudaSetDevice(1);
    printf("Hello GPU from CPU!\n");
    helloCUDA<<<1, 10>>>();
    // GPU 작업 완료 대기
    cudaDeviceSynchronize();
    cudaDeviceReset(); 
    return 0;
}

가장 간단하게 hello world 찍어야지 역시. 이 책의 예제 코드는 leetcuda예제처럼 .cu 를 library로 만든 후, python에서 그 라이브러리를 갖다 쓰는 형식이 아닌, 그냥 nvcc로 컴파일 하면 바로 돌아간당! (우왓! 간편!)

nvcc -o hello_cuda hello_cuda.cu

그래서 사실, 이 간단한 hello world를 찍는 .cu 코드를 library로 만든 후 python으로 실행시켜보려 하였으나... 우선 진도 먼저 나가자라는 마음으로 그냥 있는 코드 찍어봄.

다만! 예제 코드와는 다르게 GPU 작업 완료 대기 코드를 걸지 않으면 Hello GPU from CPU! 만 나오고 끝나버린다! 잠깐 기다리라귯!! 오늘은 Claude의 도움을 받아 그런 이슈를 해결해보았다. GPU 작업 완료를 대기해보라고 한다. 음... 기다렸더니 잘 되는군. Hello CUDA from GPU! 가 총 10번 찍힌다.

아, 그리고, 내가 사용하는 환경에 GPU가 여러 대라 그런지?? cudaSetDevice를 안 해주면 안 돌아가서 (아니면 0번이 꽉 차서 그럴까..? 근데 그러기엔 이 아이가 메모리를 뭐 얼마 쓸 것 같진 않은걸... 근데?! 메모리 꽉 찼다는 에러 메시지를 낸다?! 그래서 다른 디바이스로 세팅해줌) GPU 몇 번 쓸 건지 지정해주고 돌려야 돌아갔다.

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

__global__ void helloCUDA(void){
    printf("Block %d, Thread %d: Hello CUDA!\n", 
           blockIdx.x, threadIdx.x);
}

int main(void){
    cudaSetDevice(15);
    printf("=== 1 block, 3 threads ===\n");
    helloCUDA<<<1, 3>>>();
    cudaDeviceSynchronize();

    printf("\n=== 2 blocks, 2 threads each ===\n");
    helloCUDA<<<2, 2>>>();
    cudaDeviceSynchronize();

    return 0;
}

<<<1, 10>>>에 대해서 10이 쓰레드 10개 보낸다는 건 알았는데 앞의 1은 뭐냐 하니까! 늘 들어왔던 그 Grid-Block-Thead의 구조에 대해 설명해주는 우리 클로드군. 그러면서 다른 방식으로도 찍어볼 수 있다고 나에게 팁을 준다. 그렇다면?! 안 찍어 볼 수 없지! 당장 찍어본다!! <<<1, 3>>> 이면 당연히 3번 찍힐 것 같은데 <<<2, 2>>> 는 어떻게 찍힐까? (4번이겠지)

=== 1 block, 3 threads ===
Block 0, Thread 0: Hello CUDA!
Block 0, Thread 1: Hello CUDA!
Block 0, Thread 2: Hello CUDA!

=== 2 blocks, 2 threads each ===
Block 0, Thread 0: Hello CUDA!
Block 0, Thread 1: Hello CUDA!
Block 1, Thread 0: Hello CUDA!
Block 1, Thread 1: Hello CUDA!

Grid.. 그렇다면 Grid가 궁금해지잖아? Grid는 한 개만 있는 거냐고 물어보니까 그렇다고 한다.

네, Grid는 하나의 커널 실행당 하나만 존재합니다.

하나의 커널 호출 = 하나의 Grid │ Grid ├── Block (0,0) Block (1,0) Block (2,0) ├── Block (0,1) Block (1,1) Block (2,1) └── Block (0,2) Block (1,2) Block (2,2)

하지만?! Grid는 2D, 3D가 가능하다고?!

// 2D
dim3 gridSize(3, 2);  // 3x2 = 6개 블록
dim3 blockSize(4, 4); // 각 블록은 4x4 = 16개 스레드

helloCUDA<<<gridSize, blockSize>>>();

// 3D
dim3 gridSize(2, 2, 2);   // 2x2x2 = 8개 블록
dim3 blockSize(4, 4, 2);  // 각 블록은 4x4x2 = 32개 스레드

helloCUDA<<<gridSize, blockSize>>>();

3차원으로 돌려봤을 때!!

__global__ void helloCUDA(void){
    printf("Grid: (%d,%d,%d), Block: (%d,%d,%d), Thread: (%d,%d,%d)\n",
           gridDim.x, gridDim.y, gridDim.z,      // Grid 크기
           blockIdx.x, blockIdx.y, blockIdx.z,    // 현재 블록 ID
           threadIdx.x, threadIdx.y, threadIdx.z); // 현재 스레드 ID
}

결과는 과연?!

Grid: (2,2,2), Block: (0,0,1), Thread: (0,0,0)
Grid: (2,2,2), Block: (0,0,1), Thread: (1,0,0)
Grid: (2,2,2), Block: (0,0,1), Thread: (2,0,0)
Grid: (2,2,2), Block: (0,0,1), Thread: (3,0,0)
Grid: (2,2,2), Block: (0,0,1), Thread: (0,1,0)
Grid: (2,2,2), Block: (0,0,1), Thread: (1,1,0)
...
Grid: (2,2,2), Block: (0,1,0), Thread: (2,3,1)
Grid: (2,2,2), Block: (0,1,0), Thread: (3,3,1)

이렇다고 한다. Grid는 늘 그대로네. Block과 Thread만 바뀌넹.

여튼 이렇게 아주 간단한 hello cuda 찍어보기 + 응용 쪼끔 해봤다리~

이 파이썬 스크립트는 CUDA 커널을 PyTorch 확장으로 빌드/로딩하고, 여러 텐서 크기(S, K)에서 커스텀 커널 vs PyTorch 연산의 성능을 반복 벤치마크하는 코드.

1) CUDA 확장 로드

from torch.utils.cpp_extension import load

lib = load(
    name="elementwise_lib",
    sources=["elementwise.cu"],
    extra_cuda_cflags=[
        "-O3",
        "-U__CUDA_NO_HALF_OPERATORS__",
        "-U__CUDA_NO_HALF_CONVERSIONS__",
        "-U__CUDA_NO_HALF2_OPERATORS__",
        "-U__CUDA_NO_BFLOAT16_CONVERSIONS__",
        "--expt-relaxed-constexpr",
        "--expt-extended-lambda",
        "--use_fast_math",
    ],
    extra_cflags=["-std=c++17"],
)

• load(...)가 elementwise.cu를 컴파일 → .so 생성 → 파이썬 모듈로 로드까지 한 번에 해줘.

• 첫 실행은 컴파일 때문에 시간이 길어질 수 있음(이후 캐시된 .so 재사용). (맞아. 꽤 오래 기다려야 했음. 이 코드를 돌리실 분들이라면 조급해하지 말고 기다려보길)

• extra_cuda_cflags: (이 부분은 잘 모르지만 그냥 넘어갔다)

  • -O3, --use_fast_math: 최적화/빠른 수학 함수.
  • -U__CUDA_NO_HALF_*: half/half2/bfloat16 연산 intrinsics 활성화.
  • --expt-relaxed-constexpr, --expt-extended-lambda: CUDA C++ 확장 기능 허용.

로드되면 lib.elementwise_add_f32, lib.elementwise_add_f16x8_pack 같은 바인딩된 파이썬 함수로 커널을 호출할 수 있어.

2) 자동 미분 끄기

torch.set_grad_enabled(False)

• 벤치마크용이므로 autograd 비활성화 → 오버헤드 감소.

3) 벤치마크 함수

def run_benchmark(perf_func, a, b, tag, out=None, warmup=10, iters=1000, ...):

• 입력

  • perf_func: 실행할 함수(커스텀 커널 또는 torch.add).
  • a, b: 입력 텐서.
  • out: 출력 버퍼(있으면 in-place 스타일, 없으면 함수가 반환값 생성).
  • warmup: 워밍업 반복 횟수(캐시/클럭 워밍업).
  • iters: 실제 측정 반복 횟수(평균 시간 산출).

• 동작

  1. out.fill_(0)로 출력 초기화(있을 때).
  2. 워밍업 루프 수행.
  3. torch.cuda.synchronize()로 GPU 작업 동기화 후 타이머 시작.
  4. iters만큼 호출(동기화 후 종료 시각 측정).
  5. 평균 실행 시간(ms) 계산.
  6. 결과 텐서의 앞 2개 값만 CPU로 가져와 출력(정확도 확인).

• 중요 포인트

  • CUDA는 비동기라서, 정확한 시간 측정을 위해 앞뒤로 torch.cuda.synchronize()가 필수.
  • iters=1000이라 꽤 오래 걸릴 수 있음(정밀 측정 목적).

4) 실험 (텐서 생성 & 반복)

Ss = [1024, 2048, 4096]
Ks = [1024, 2048, 4096]
for S, K in SKs:
    a = torch.randn((S, K)).cuda().float().contiguous()
    b = torch.randn((S, K)).cuda().float().contiguous()
    c = torch.zeros_like(a).cuda().float().contiguous()

• 9개 크기 조합(S,K) 각각에 대해 텐서 생성(GPU, FP32, contiguous).

• FP32 벤치마크

  • lib.elementwise_add_f32(a,b,c) : 스칼라형 FP32 커널
  • lib.elementwise_add_f32x4(a,b,c) : float4 벡터화 버전
  • partial(torch.add, out=c) : PyTorch 기본 연산(비교군)

• FP16 벤치마크

  a_f16 = a.half().contiguous(); b_f16 = b.half().contiguous(); c_f16 = c.half().contiguous()

  • lib.elementwise_add_f16 : half 스칼라
  • lib.elementwise_add_f16x2 : half2 벡터화
  • lib.elementwise_add_f16x8 : half2×4로 8개 처리
  • lib.elementwise_add_f16x8_pack : 128비트 일괄 로드/스토어 + half2 연산(가장 공격적 최적화)
  • partial(torch.add, out=c_f16) : PyTorch FP16 비교군

각 호출은 run_benchmark(...)으로 감싸서 평균 시간과 샘플 출력값을 찍어 줘.

5) 출력 포맷

-------------------------------------
                                        S=1024, K=1024
           out_f32: [v1, v2], time:0.00xxms
         out_f32x4: [v1, v2], time:0.00xxms
        out_f32_th: [v1, v2], time:0.00xxms
-------------------------------------
           out_f16: [...]
         out_f16x2: [...]
         out_f16x8: [...]
     out_f16x8pack: [...]
        out_f16_th: [...]
-------------------------------------

• out_*: 결과 텐서 앞 2개의 값(정확도 체크)
• time:: 1회 평균 소요 시간(ms)

6) 왜 느릴 수 있나 (정상임)

• 처음 load(...) 컴파일 비용.
• 각 조합마다 워밍업 10회 + 측정 1000회 × 커널 8개 × 사이즈 9개 → 호출 수가 많음.
• 정확한 측정을 위해 synchronize()가 들어가 있어 CPU가 GPU 끝날 때까지 대기.

빠르게 돌리고 싶다면: iters 줄이기(예: 100), S/K 조합 축소, show_all=False 유지.

7) 안전/품질 포인트

• .contiguous()로 메모리 연속 보장(벡터화/128비트 접근에 중요).
• FP16 커널은 내부에서 __hadd/__hadd2 사용(half/half2 연산).

GPT가 설명을 알잘딱깔센으로 잘해줘서 궁금한 부분만 조금 더 물어보고 python 코드는 더 파지 않았다. 왜냐... 진도 나가야 해. 쿠다 코드 보는 걸 우선으로 하고 있는데, 쿠다 library를 실행할 파이썬 코드도 궁금하면 다시 돌아오는 걸로..!

책은 이제 차근차근 읽어나갈 거고, 우선 LeetCUDA의 kernel 코드를 차근차근 따라가며, 모르는 부분은 GPT에게 물어가며 공부하는 게시글이 될 것 같다.

우선 첫 번 째는!! https://github.com/xlite-dev/LeetCUDA/blob/main/kernels/elementwise/elementwise.cu <- 이 코드!

아예 아무 것도 몰라서 진짜 코드 한 줄 한 줄 따라가며 GPT에게 계속 물어봄. 나중을 (나중의 나를) 위하여 뭔가를 남겨보기로 함!

헤더 파일

#include <algorithm>
#include <cuda_bf16.h>
#include <cuda_fp16.h>
#include <cuda_fp8.h>
#include <cuda_runtime.h>
#include <float.h>
#include <stdio.h>
#include <stdlib.h>
#include <torch/extension.h>
#include <torch/types.h>
#include <vector>

우선 헤더 파일들을 include 하는 부분은 그냥 그렇구나... 하고 넘어감. 헤더까지 살피기엔 아직 너무 초짜임! (그도 그럴 것이 원래 stdio.h, stdlib.h는 그냥 이게 있으니 된다! 하고 넘어가는 거자낭?)

define

WARP

#define WARP_SIZE 32

WARP_SIZE를 32로 지정했는데 WARP가 뭘까... WARP는 GPU 병렬 처리의 가장 기본 단위라고 한다.

Warp는 32개의 쓰레드로 구성된 실행 단위입니다. CUDA GPU는 쓰레드를 하나씩 실행하지 않고, 한 번에 32개 쓰레드를 묶어서 동시에 실행합니다. 이 32개의 쓰레드 묶음을 warp라고 합니다.

🧠 왜 Warp 단위로 실행할까?

📌 이유: SIMT (Single Instruction, Multiple Threads) 구조 때문

GPU는 일반 CPU와 달리 같은 명령어를 여러 데이터에 대해 동시에 실행하는 방식으로 효율을 극대화합니다.

• CPU: 각 코어가 각기 다른 일을 할 수 있음
• GPU: 여러 쓰레드가 같은 명령어를 공유하며 동시에 실행됨 → 성능 향상

CUDA에서는 각 warp 내의 32개 쓰레드가 동시에 같은 instruction을 실행합니다. 이를 SIMT라고 하며, warp가 실행 단위가 됩니다.

그렇군! CUDA의 기본 실행 단위! 32개의 쓰레드로 구성된 실행 단위이군!

reinterpret_cast

#define INT4(value) (reinterpret_cast<int4 *>(&(value))[0])
#define FLOAT4(value) (reinterpret_cast<float4 *>(&(value))[0])
#define HALF2(value) (reinterpret_cast<half2 *>(&(value))[0])
#define BFLOAT2(value) (reinterpret_cast<__nv_bfloat162 *>(&(value))[0])
#define LDST128BITS(value) (reinterpret_cast<float4 *>(&(value))[0])

타입 캐스팅 define~ 이 매크로들은 데이터를 벡터 타입으로 reinterpret하여 메모리 접근 성능을 향상시키기 위한 것!

✅ 이 매크로들을 쓰는 이유는?

메모리 병합(coalescing) 최적화: GPU에서는 연속된 쓰레드들이 연속된 주소에 접근할 때 한 번의 메모리 트랜잭션으로 처리하면 빠릅니다.

• 128비트 단위 load/store: float4, half2 등을 이용하면 메모리 IO 성능이 상승함

reinterpret_cast<new_type>(value) 이 문법 자체가 타입 정보를 무시하고 메모리 주소만 강제로 해석하는 용도로 쓰이며, 보통 매우 저수준의 메모리 접근이나 하드웨어 최적화에서 사용된다고 함. reinterpret_cast는 C++ 언어 자체에 포함된 키워드라고 함.

이 부분 관련해서 GPT에 물어보면 메모리 접근 관련 정보/지식들을 얻을 수 있다.

여기에 관해서 티키타카 티키타카 해본 후 내가 이해한 바로는: 한 번에 더 많이 읽어서 메모리 I/O를 줄이기! <- 이것이 목적인듯 하다.

• 일반적으로 128비트(16B)에 최적화 되어 있음.

어쨌든 이렇게 벡터화해서 읽는 게 유리함: 데이터를 여러 개 단위로 묶어 한 번에 처리하거나, 한 명령어로 여러 데이터를 동시에 처리하는 최적화 기법! 즉, float value 하나 하나... 처리하는 게 아니고, 128bit까지 마! 가져온나!! 팍팍! 한 번에 128bit 처리한닷!! 요런 느낌.

elementwise add kernel (f32_kernel)

// FP32
// ElementWise Add grid(N/256),
// block(256) a: Nx1, b: Nx1, c: Nx1, c = elementwise_add(a, b)
__global__ void elementwise_add_f32_kernel(float *a, float *b, float *c,
                                           int N) {
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  if (idx < N)
    c[idx] = a[idx] + b[idx];
}

• __global__ CUDA 커널 함수임을 나타내는 키워드
• int N 배열 길이

너무나도 당연하게 a, b를 더해서 c에 저장하는 함수임.

• int idx blockIdx.x * blockDim.x + threadIdx.x;

• 위의 코드는 현재 쓰레드가 처리할 인덱스를 계산하는 부분.

• CUDA는 thread를 block 단위로 구성하고, block은 grid 단위로 구성되어 있다고 함.

• 그래서 전체 글로벌 인덱스를 계산하려면 위처럼 계산해야 한다고..!!

  • blockIdx.x: 현재 블록의 x축 인덱스
  • blockDim.x: 블록당 쓰레드 수
  • threadIdx.x: 현재 블록 내에서의 쓰레드 인덱스
  • idx: 전체 배열 상에서 이 쓰레드가 담당할 위치

음... 이렇게 하나하나 위치를 계산해줘야 하는군.

너무나도 오랜만에 맡는 C의 향기. 그동안 너무 쉽게 코딩해왔나봐..!!

당연히 if 문은 배열 경계 보호하는 if문.

elementwise add kernel (f32x4_kernel)

// ElementWise Add + Vec4
// grid(N/256), block(256/4)
// a: Nx1, b: Nx1, c: Nx1, c = elementwise_add(a, b)
__global__ void elementwise_add_f32x4_kernel(float *a, float *b, float *c,
                                             int N) {
  int idx = 4 * (blockIdx.x * blockDim.x + threadIdx.x);
  if (idx < N) {
    float4 reg_a = FLOAT4(a[idx]);
    float4 reg_b = FLOAT4(b[idx]);
    float4 reg_c;
    reg_c.x = reg_a.x + reg_b.x;
    reg_c.y = reg_a.y + reg_b.y;
    reg_c.z = reg_a.z + reg_b.z;
    reg_c.w = reg_a.w + reg_b.w;
    FLOAT4(c[idx]) = reg_c;
  }
}

위와 조금 달라진 건 4개를 한 꺼번에 처리하고자 하는 이 욕망! 그래서 idx를 계산할 때도 4개 단위로 계산해줌.

위와 조금 다른 건, 로딩 부분! 4개를 한 번에 로딩해준다! a 배열에 있던 float 4개 일루와! b 배열에 있던 float 4개 일루와! 그리고 .x, .y, .z, .w로 한땀 한땀 더해주고. 마지막도 FLOAT4로 매크로 처리해서 이 4개의 값을 한.꺼.번.에 저장해줌.

여기서 이 xyzw는 어디서 나온거냐 하면: float4라는 것은 CUDA에서 제공하는 벡터 타입. 내부적으로 다음과 같이 정의되어 있다

struct __device_builtin__ __builtin_align__(16) float4 { float x, y, z, w; };

아하! 이래서 xyzw로 접근할 수 있구나!

elementwise add kernel (f16x8_kernel)

__global__ void elementwise_add_f16x8_kernel(half *a, half *b, half *c, int N) {
  int idx = 8 * (blockIdx.x * blockDim.x + threadIdx.x);
  half2 reg_a_0 = HALF2(a[idx + 0]);
  half2 reg_a_1 = HALF2(a[idx + 2]);
  half2 reg_a_2 = HALF2(a[idx + 4]);
  half2 reg_a_3 = HALF2(a[idx + 6]);
  half2 reg_b_0 = HALF2(b[idx + 0]);
  half2 reg_b_1 = HALF2(b[idx + 2]);
  half2 reg_b_2 = HALF2(b[idx + 4]);
  half2 reg_b_3 = HALF2(b[idx + 6]);
  half2 reg_c_0, reg_c_1, reg_c_2, reg_c_3;
  reg_c_0.x = __hadd(reg_a_0.x, reg_b_0.x);
  reg_c_0.y = __hadd(reg_a_0.y, reg_b_0.y);
  reg_c_1.x = __hadd(reg_a_1.x, reg_b_1.x);
  reg_c_1.y = __hadd(reg_a_1.y, reg_b_1.y);
  reg_c_2.x = __hadd(reg_a_2.x, reg_b_2.x);
  reg_c_2.y = __hadd(reg_a_2.y, reg_b_2.y);
  reg_c_3.x = __hadd(reg_a_3.x, reg_b_3.x);
  reg_c_3.y = __hadd(reg_a_3.y, reg_b_3.y);
  if ((idx + 0) < N) {
    HALF2(c[idx + 0]) = reg_c_0;
  }
  if ((idx + 2) < N) {
    HALF2(c[idx + 2]) = reg_c_1;
  }
  if ((idx + 4) < N) {
    HALF2(c[idx + 4]) = reg_c_2;
  }
  if ((idx + 6) < N) {
    HALF2(c[idx + 6]) = reg_c_3;
  }
}

자 바로 half 드간다. __hadd 이거는 CUDA에서 제공하는 half 타입 전용 덧셈 함수. 위 코드들과 같이 그냥 + 연산자로는 덧셈이 안 되는 듯 하다. half용 덧셈 연산자(function)~

HALF2는 half 두 개라서 4바이트(32bit)가 된다. .x, .y로 더하는 부분도 위에서 쭈욱 내려왔다면 대충 이해갈 거라 생각. 여기서 궁금한 건 왜 128bit 씩 처리하지 않는가? 내 생각엔 아마도 다음 커널과 비교하기 위해 있는 것 같음.

elementwise add kernel (f16x8_pack_kernel)

__global__ void elementwise_add_f16x8_pack_kernel(half *a, half *b, half *c,
                                                  int N) {
  int idx = 8 * (blockIdx.x * blockDim.x + threadIdx.x);

  // temporary register(memory), .local space in ptx, addressable
  half pack_a[8], pack_b[8], pack_c[8]; // 8x16 bits=128 bits.

  // reinterpret as float4 and load 128 bits in 1 memory issue.
  LDST128BITS(pack_a[0]) = LDST128BITS(a[idx]); // load 128 bits
  LDST128BITS(pack_b[0]) = LDST128BITS(b[idx]); // load 128 bits

#pragma unroll
  for (int i = 0; i < 8; i += 2) {
    // __hadd2 for half2 x 4
    HALF2(pack_c[i]) = __hadd2(HALF2(pack_a[i]), HALF2(pack_b[i]));
  }
  // reinterpret as float4 and store 128 bits in 1 memory issue.
  if ((idx + 7) < N) {
    LDST128BITS(c[idx]) = LDST128BITS(pack_c[0]);
  }
}

팩!!! 패킹!!! LDST128BITS 라는 아이가 나와버렸다! 이 아이는... 128 bit씩 접근하는 아이!

🧠 핵심 목적

128비트 패킹된 half[8] 배열을 사용해

• GPU memory I/O 횟수 줄이고 (1번에 128비트)
• half2 기반 병렬 연산(__hadd2)으로 연산량도 최적화

인덱스 계산은: half 8개씩 처리하겠다! (=128 bits)

그리고, 로컬 레지스터 공간에 half[8]씩 메모리를 잡고 이 배열을 float4로 interprete해서 128 비트 단위로 load/store. 이렇게 버퍼를 만드는 이유를 잘 모르겠어서 GPT에게 계속 꼬치꼬치 물어봄. 메모리 I/O 최소화 + 벡터화 연산을 깔끔하게 하기 위해서라고 함. LDST128BITS를 통해 일괄 load/store가 가능. 만약 로컬 버퍼 없이 HALF2(a[idx]), HALF2(a[idx+2])… 처럼 바로 글로벌에서 뽑아 쓰면 반복적으로 글로벌 메모리 접근이 발생할 수 있다고 함. 그러니까, 내가 이해하기로는 로컬 메모리로 들고 와야 더 빠르게 연산이 가능하니까! (마치 L1, L2 cache 처럼. 더 가까운 메모리로 이주시키기.) 음. 이건 이해할 수 있어. 이렇게 하지 않으면 바로 위의 코드(f16x8_kernel)와 같은 형태가 되는 거겠지. 오키. 우선 넘어가자. a, b 에 있는 아이들을 128 bits로 가져와서 pack_a, pack_b에 담아주고.

#pragma unroll (= 이 루프 짧으니까 컴파일 타임에 미리 다 풀어서 반복문 없이 실행해줘)

처음 보는 아이. 이 아이는 CUDA에서 루프 전개(loop unrolling)를 컴파일러에 지시하는 명령!

#pragma unroll
for (int i = 0; i < 4; i++) {
    sum += arr[i];
}

은 컴파일 후에 아래처럼 변환됨

sum += arr[0];
sum += arr[1];
sum += arr[2];
sum += arr[3];

🔹 CUDA에서 중요한 이유!

• GPU는 SIMT 구조라서 루프 제어보다는 계속된 연산 명령이 pipeline에 쭉 이어지는 게 유리함
• 특히 여기처럼 4번만 도는 작은 루프는 unroll하면 실행 흐름이 단순해짐
• 레지스터 사용 패턴과 메모리 접근을 컴파일러가 더 잘 최적화 가능

🔹 여기서 적용된 이유

#pragma unroll
for (int i = 0; i < 8; i += 2) {
    HALF2(pack_c[i]) = __hadd2(HALF2(pack_a[i]), HALF2(pack_b[i]));
}

위의 코드는 컴파일을 하면 아래처럼 바뀜

HALF2(pack_c[0]) = __hadd2(HALF2(pack_a[0]), HALF2(pack_b[0]));
HALF2(pack_c[2]) = __hadd2(HALF2(pack_a[2]), HALF2(pack_b[2]));
HALF2(pack_c[4]) = __hadd2(HALF2(pack_a[4]), HALF2(pack_b[4]));
HALF2(pack_c[6]) = __hadd2(HALF2(pack_a[6]), HALF2(pack_b[6]));

• 루프 분기 오버헤드가 사라지고, 명령이 연속적으로 나열되니까 GPU warp 실행 효율이 올라감

휴 이렇게 간단한 elementwise add하는 kernel code를 살펴보았다. 되게 간단한데 모르는 게 너무 많아서 하나하나 물어보고 이해해가면서 하느라 오래 걸렸음!

이렇게 커널 코드 이후에 매크로 코드가 나온다. PyTorch C++/CUDA 확장 바인딩과 커널 런처(wrapper) 생성을 매크로로 자동화한 부분!

간단 매크로

#define STRINGFY(str) #str
#define TORCH_BINDING_COMMON_EXTENSION(func)                                   \
  m.def(STRINGFY(func), &func, STRINGFY(func));

PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) {
  TORCH_BINDING_COMMON_EXTENSION(elementwise_add_f32)
  TORCH_BINDING_COMMON_EXTENSION(elementwise_add_f32x4)
  TORCH_BINDING_COMMON_EXTENSION(elementwise_add_f16)
  TORCH_BINDING_COMMON_EXTENSION(elementwise_add_f16x2)
  TORCH_BINDING_COMMON_EXTENSION(elementwise_add_f16x8)
  TORCH_BINDING_COMMON_EXTENSION(elementwise_add_f16x8_pack)
}

#define CHECK_TORCH_TENSOR_DTYPE(T, th_type)                                   \
  if (((T).options().dtype() != (th_type))) {                                  \
    std::cout << "Tensor Info:" << (T).options() << std::endl;                 \
    throw std::runtime_error("values must be " #th_type);                      \
  }

• 토큰을 문자열로 바꾸는 STRINGFY
  • 예) STRINGFY(elementwise_add_f32) → "elementwise_add_f32"
• TORCH_BINDING_COMMON_EXTENSION(func)은 pybind11로 C++ 함수 func를 그 이름 그대로 python에 등록
  • 예) TORCH_BINDING_COMMON_EXTENSION(elementwise_add_f32) → m.def("elementwise_add_f32", &elementwise_add_f32, "elementwise_add_f32");
• pybind11 모듈 정의
  • python에서 lib.elementwise_add_f16x8_pack(a, b, c) 처럼 직접 호출 가능.
  • 빌드 시 지정한 이름으로 python 모듈이 만들어지고, 위에서 생성된 6개의 래퍼 함수가 동일한 이름으로 python에 노출됨.
• dtype 체크
  • 텐서 T의 dtype이 기대 타입과 다르면 에러를 던짐.

핵심: 커널 런처(wrapper) 생성 매크로

#define TORCH_BINDING_ELEM_ADD(packed_type, th_type, element_type, n_elements) \
  void elementwise_add_##packed_type(torch::Tensor a, torch::Tensor b,         \
                                     torch::Tensor c) {                        \
    CHECK_TORCH_TENSOR_DTYPE(a, (th_type))                                     \
    CHECK_TORCH_TENSOR_DTYPE(b, (th_type))                                     \
    CHECK_TORCH_TENSOR_DTYPE(c, (th_type))                                     \
    const int ndim = a.dim();                                                  \
    if (ndim != 2) {                                                           \
      int N = 1;                                                               \
      for (int i = 0; i < ndim; ++i) {                                         \
        N *= a.size(i);                                                        \
      }                                                                        \
      dim3 block(256 / (n_elements));                                          \
      dim3 grid((N + 256 - 1) / 256);                                          \
      elementwise_add_##packed_type##_kernel<<<grid, block>>>(                 \
          reinterpret_cast<element_type *>(a.data_ptr()),                      \
          reinterpret_cast<element_type *>(b.data_ptr()),                      \
          reinterpret_cast<element_type *>(c.data_ptr()), N);                  \
    } else {                                                                   \
      const int S = a.size(0);                                                 \
      const int K = a.size(1);                                                 \
      const int N = S * K;                                                     \
      if ((K / (n_elements)) <= 1024) {                                        \
        dim3 block(K / (n_elements));                                          \
        dim3 grid(S);                                                          \
        elementwise_add_##packed_type##_kernel<<<grid, block>>>(               \
            reinterpret_cast<element_type *>(a.data_ptr()),                    \
            reinterpret_cast<element_type *>(b.data_ptr()),                    \
            reinterpret_cast<element_type *>(c.data_ptr()), N);                \
      } else {                                                                 \
        int N = 1;                                                             \
        for (int i = 0; i < ndim; ++i) {                                       \
          N *= a.size(i);                                                      \
        }                                                                      \
        dim3 block(256 / (n_elements));                                        \
        dim3 grid((N + 256 - 1) / 256);                                        \
        elementwise_add_##packed_type##_kernel<<<grid, block>>>(               \
            reinterpret_cast<element_type *>(a.data_ptr()),                    \
            reinterpret_cast<element_type *>(b.data_ptr()),                    \
            reinterpret_cast<element_type *>(c.data_ptr()), N);                \
      }                                                                        \
    }                                                                          \
  }

TORCH_BINDING_ELEM_ADD(f32, torch::kFloat32, float, 1)
TORCH_BINDING_ELEM_ADD(f32x4, torch::kFloat32, float, 4)
TORCH_BINDING_ELEM_ADD(f16, torch::kHalf, half, 1)
TORCH_BINDING_ELEM_ADD(f16x2, torch::kHalf, half, 2)
TORCH_BINDING_ELEM_ADD(f16x8, torch::kHalf, half, 8)
TORCH_BINDING_ELEM_ADD(f16x8_pack, torch::kHalf, half, 8)

TORCH_BINDING_ELEM_ADD(packed_type, th_type, element_type, n_elements)

• 이 매크로는 PyTorch에서 호출할 래퍼 함수를 만든다. 내부에서:
  1. 입력/출력 텐서의 dtype 검사
  2. 텐서 차원에 따라 그리드/블록 계산
  3. 원시 포인터로 캐스팅 후 해당 CUDA 커널 호출

** 만들어지는 함수 시그니처**

void elementwise_add_<packed_type>(torch::Tensor a, torch::Tensor b, torch::Tensor c)

예) packed_type=f16x8_pack → elementwise_add_f16x8_pack(...)

내부 로직 요약

• ndim = a.dim()을 보고 분기:

  • 비(非) 2D: 전체 요소 수 N = ∏ a.size(i)로 평탄화해서 처리

    block = 256 / n_elements
    grid  = (N + 256 - 1) / 256

  • 2D (S×K): 먼저 “행 단위” 실행을 시도

    • 한 쓰레드가 n_elements개(예: float4=4, half2=2, half8=8)를 처리하므로, block.x = K / n_elements를 잡습니다.
    • 이 값이 1024 이하(스레드 최대)면 grid = S, block = K / n_elements로 설정 → 각 블록이 한 행을 처리(메모리 접근 정합성이 좋아 성능 유리).
    • 그렇지 않으면 다시 평탄화 일반 경로(위의 256 스레드 기준)로 폴백.

• 커널 호출:

  elementwise_add_<packed_type>_kernel<<<grid, block>>>(
    reinterpret_cast<element_type*>(a.data_ptr()),
    reinterpret_cast<element_type*>(b.data_ptr()),
    reinterpret_cast<element_type*>(c.data_ptr()),
    N);

  • element_type은 커널이 기대하는 원소 타입: float/half 등
  • n_elements는 벡터 패킹 폭: 1, 2, 4, 8 (float4→4, half2→2, half8→8 같은 개념)

요컨대, 같은 커널 패턴을 패킹 폭에 따라 재사용하고, 텐서가 2D이면 행 단위 실행으로 최적화를 먼저 시도하는 런처를 매크로 한 방으로 생성.

• 그리고 TORCH_BINDING_ELEM_ADD를 불러서 6개의 래퍼 함수 생성
  • 예) elementwise_add_f32x4(a,b,c)는 내부에서 elementwise_add_f32x4_kernel을 호출하고, block/grid는 n_elements=4 기준으로 설정.

마지막 커널 런처 생성 매크로 부분은 열심히 이해 안 하고 대충 건너뜀... 나중에 다시 꼼꼼히 살펴보자! 휴, 우선 이렇게 아주아주 간단 elementwise add 커널 코드 읽어봄! 그 다음 이 코드를 부르는 python 코드도 살펴보고 실행 결과도 보려고 한다! (흑흑.. ㅜㅜ 빡세...)

그 다음엔 이제 하나하나.. histogram도 보고, 책도 읽고 해야 하는데. 갈 길이 멀게 느껴지지만 화이팅 하자! 우선 책 한 번 읽고, histogram까지도 좀 파악하고 나면 앞으로 공부 어떻게 해야할지 감이 좀 잡혔으면 좋겠음! 공부 스케줄도!

기본적으로 instruction tuning은 instruction과 입력을 넣으면 원하는 출력이 나오게 하도록 하는 학습 방식이다. 다양한 instruction+input & output pair를 주고 학습시킴으로써 이 지시를 모델이 이해하도록 하는 게 주 목적. 그래서 다양한 instruction을 필요로 한다. 결국 우리가 원하는 건 하나의 task에만 tuning된 fine-tuning 모델이 아니라, instruction을 이해하기 때문에 새로운 instruction에도 찰떡같이 동작하는 모델이기 때문이다.

• FLAN (https://arxiv.org/pdf/2210.11416.pdf)이 이 분야에서는 유명한 논문. 2024.01 기준으로 1천회 넘게 인용되었다. Google 논문임.

• 특기할만한 점이 무엇이 있는가.
  • 태스크를 늘려보았고,_ (태스크 종류는? 종류의 다양성이 중요할까, 아니면 숫자가 늘어나는 게 중요한 걸까?)_
  • 모델 사이즈를 크게 가져가보았고, (모델 사이즈와 어떤 연관이 있을까?)
  • CoT를 적용해보았다. (이게 어떤 영향이 있었을까?!)

1. 인트로

인트로에서 떠먹여주는 정보에 의하면

1. 모델 사이즈가 커질 수록, 태스크 개수가 많아질 수록 더더 잘함!!
2. CoT를 넣어야 reasoning을 잘한다.

그래서 1.8k 태스크에 대해 540B짜리 Flan-PaLM을 구웠다고 함. (540B라니.. 역시 자본인가) 그렇게 구운 Flan-PaLM은 그냥 PaLM보다 성능이 좋다고 함. (당연한 거 아닐까?) 그리고 자비롭게도 T5에도 구워주셨다: Flan-T5가 그렇게 탄생. 80M부터 11B까지.

2. Flan Fine-tuning

그래서 Flan은 어떻게 학습하는 건데? 아니 애초에 무엇의 약자이지? _Fintuning Language Models_의 약자입니다.

데이터

가장 중요하겠쥬?! 기존 논문들을 보면 task는 많으면 많을 수록 좋다길래 prior work의 데이터를 싹 갖다 썼습니다. T0-SF (193 tasks), Muffin (80 Tasks), CoT (9 Tasks), NIV2 (1,554 Tasks). 도합 1,836 tasks.

• T0-SF는 T0 데이터셋에서 muffin과 겹쳐지는 데이터셋을 뺀 거. _T0 데이터셋_은 _P3_라고도 하는 것 같으니 참고. (T0 논문: https://arxiv.org/pdf/2110.08207.pdf)

CoT는 arithmetic reasoning, multi-hop reasoning, natural language inference로 구성되어 있다고 함. 저자들이 CoT를 활용한 reasoning 능력을 중요하게 생각한듯함.

어떻게 학습한 거지?

각 모델별 자세한 hyperparameters는 논문 (2.2) 참조*

어떻게 평가할 거임??

MMLU (시험 문제. 수학, 역사, 법, 의학 등), BBH (Big-Bench 종류), TyDiQA (QA인데 이제 multilingual을 곁들인), MSGM (수학문제인데 이것도 multilingual을 곁들인).

MMLU & BBH: Direct, CoT 두가지 방식 (같은 걸 두 방식으로 푼 게 아니라, 데이터셋이 나뉨) TyDiQA: Direct MSGM: CoT

For MMLU and BBH, we evaluate both the ability to directly predict the answer via direct prompting, where the model directly gives the answer (Brown et al., 2020; Srivastava et al., 2022), as well as via chain-of-thought (CoT) prompting, where the model must provide a reasoning chain before giving the final answer (Wei et al., 2022b).

few-shot을 주고 평가. (five-shot for MMLU, three-shot for BBH, one-shot for TyDiQA, and 8-shot for MGSM)

평가 결과는 normalized average로 report.

3. 모델 크기와 태스크 개수가 성능에 미치는 영향은?

이게 궁금해서 해봤다고 함. 인트로에서 말했듯 모델 사이즈가 커질 수록 성능은 좋아지고, 또 더 많은 태스크에 학습했을 수록 성능이 좋아진다. 근데 1.8k랑 282 태스크의 성능 갭이 크지 않네?

• 여기에 대해 저자들은 두 가지 설명이 있다고 밝혔는데, 첫째로는 추가된 1.5k 태스크가 별로 diverse하지 않았을 가능성, 그래서 모델이 뭔가 새로운 지식을 습득하지 못했을 가능성. 두번째로는 instruction fine-tuning이 기존에 사전학습으로 갖고 있던 지식을 다르게 표현하는 방법을 배우는 건데, 그 태스크 관점에서는 282개의 태스크로도 충분해서 그 이상은 이제 큰 도움이 안 된다는 설명.
• 음,, 어떤 가설이 맞을까? 당연히 1번이라고 생각했는데 2번도 꽤나 설득력이 있다. 그래도 아무래도 1번 아닐까? held-out task가 충분히 다양했더라면 1번처럼 1.8k로 학습한 것이 더 보여질 여지가 있는 거... 아닐까? 학습 태스크의 다양성도 중요하겠지만, 그것을 얼마나 잘 측정할 것인가도 중요하니까.

4. CoT Fine-tuning

인트로에서 두 번째로 나왔듯 CoT가 얼마나 효과적인가? 를 탐구한 섹션. CoT를 통해 reasoning과 zero-shot 성능이 향상될 수 있다는 사실을 보여줌. 단 9개의 CoT 데이터셋만으로..! 표랑 그래프 꼼꼼히 읽고 이해해야 하는데 조금 귀찮아서... 스킵... 우선 CoT가 필요하다... 정도로 ㅇㅋ하고 넘어감

5. 다 합치면

이 Norm. avg.는 뭔지, 왜 -부터 시작하는지... 잘 모르겠으나... 우선 Flan을 쓴게 일관적으로 더 나아짐. 특히 T5는 multi-lingual feature가 없었어서 T5가 다른 모델보다 더 flan의 덕을 봤다.

6. Open-ended generation

에 대해 human preference를 측정했는데 (PaLM vs. Flan-PaLM) Flan-PaLM이 79% of time 선택됨.

7. 디스커션

• Scaling curves for instruction finetuning.
• CoT finetuning is critical for reasoning abilities.
• Instruction finetuning generalizes across models.
• Instruction finetuning improves usability and mitigates some potential harms.
• Instruction finetuning is relatively compute-efficient.

이렇게 FLAN 논문을 살펴보았다. 논문 본문의 내용은 사실 많지 않고 간단. FLAN에서 사용한 데이터셋, FLAN의 발견만 기억하면 됨. 부록이 아주 빵빵하니 부록도 보면 좋긴 하겠다. FAQ와 실제 qualitative examples, bias&harms와 디테일한 실험 결과들이 있다.

1. 유명한 데이터셋들

https://newsletter.ruder.io/p/instruction-tuning-vol-1 읽고 정리겸 끄적끄적,,

1. Natural Instructions
   1. 영어, 193k examples, 61 tasks
   2. common schema를 사용해서 다른 데이터셋에 비해 more structured

2. Natural Instructions v.2 / Super-Natural Instructions
   1. 5M examples, 76 task types, 55 languages
      1. 세세하게는 1600 이상의 task (SUPER-NATURALINSTRUCTIONS: Generalization via Declarative Instructions on 1600+ NLP Tasks)
   2. v.1에 비해 instruction이 간단해짐: task definition, 긍정 예제, 부정 예제 with 설명

3. Unnatural Instructions
   1. 자동적으로 모은 instruction 데이터셋
   2. 240k examples, Super-Natural Instruction 예제로 InstructGPT (text-davinci-002)에 넣어서 새로운 example을 만들라고 한 것.
   3. SNI보다 더 다양한 task를 다루고 있음

4. Self-Instruct
   1. Unnatural Instruction과 비슷
   2. 82k, 175 tasks. seed task로 InstructGPT에게 날림

5. Flan (2021, 2022) #reasoning
   1. 2022는 Flan 2021, P3, SNI를 합하고 + additional reasoning, dialog, and program synthesis datasets를 합친 것. 9개의 새로운 reasoning datasets additionally annotated with CoT!!

정리

Important Aspects of Instruction Data

Longpre et al. (2023) and Iyer et al. (2022) ablate several important aspects of instruction data, which we highlight in the following.

Mixing few-shot settings. Training with mixed zero-shot and few-shot prompts significantly improves performance in both settings. 흠, 그럼 훈련 데이터에 few-shot도 넣고, zero-shot도 넣으란 건가! 그냥 zero-shot으로만 학습시키면 별론가?

Task diversity. Large models benefit from continuously increasing the number of tasks.

Data augmentation. Augmenting the data such as by inverting inputs and outputs (e.g., turning a question answering task into a question generation task) is beneficial. 데이터 확장이 도움이 됨. (흠?)

Mixing weights. When using a combination of instruction tuning datasets, appropriately tuning the mixing weights is important. 잘 섞어야 함.

While the above datasets are mainly derived from classical NLP tasks, recent datasets such as Baize (Xu et al., 2023), OpenAssistant Conversations (Köpf et al., 2023) and others cover a more diverse set of applications and domains. We will discuss these in the next edition. Stay tuned! 👋

근데!! 이 instruction data를 안다고 해도, 뭔가 insight를 가지긴 해야 함. 제대로 알려면 논문을 한 번 쫙 읽고 전반적인 느낌을 한 번 쏵!!

2. 요즘의 데이터셋 & instruction-tuned models

https://newsletter.ruder.io/p/instruction-tuning-vol-2 여기에 더 자세한 내용

1. Alpaca
   • 52k, 영어, text-davinci-003 with self-instruct
2. Evol-Instruct
   • 250k, 영어, Alpaca를 기반으로 rewrite.
3. Vicuna, ShareGPT
   • 70k, Conversation!
4. LIMA
   • 퀄리티 좋은 1000개의 instruction 데이터로 높은 성능을 달성 가능하단 것을 보여줌.

Takeaways

✅ Quality > quantity. As Zhou et al. (2023) observe, training on a small set of high-quality data outperforms instruction-tuning on larger, noisier data. Using more diverse prompts and quality filtering both improve performance.

🧑‍🎓 Imitation != mastery. Models that are instruction-tuned on ChatGPT-generated data mimic ChatGPT’s style (and may thus fool human raters!) but not its factuality (Gudibande et al., May 2023). They perform worse on standard benchmarks. Using stronger base models is the best way to address this.

🏛️ The stronger the base, the better. More powerful base models also produce stronger instruction-tuned models (Wang et al., June 2023).

🥇 The combination wins. Combining multiple instruction-tuning datasets results in the best average performance across tasks (Wang et al., June 2023). Dataset mixing and developing modular instruction-tuned models are thus important research directions.

Future Directions

Understanding instruction-tuning. While we have seen a proliferation of instruction-tuning datasets, we still lack a clear understanding of what makes a good instruction and good instruction–response pairs. There is much anecdotal knowledge when it comes to creating good model prompts—but to my knowledge it is unclear how instruction–following data can be created at scale in a more principled manner.

Improving data quality. To improve model performance, we need to develop more reliable methods to identify high-quality examples and filter out undesirable ones. In a similar vein, it is important to develop methods that allow us to identify how a particular instance affects model behavior and alignment at test time.

Evaluating instruction-tuned models. In light of the biases of both human and automatic evaluations, there is no clear gold standard for how to evaluate instruction-tuned models. Evaluating a model on a set of tests that can be efficiently and automatically evaluated is one way to side-step this issue, see LMentry (Efrat et al., ACL 2023), M2C (Hlavnova & Ruder, ACL 2023), IFEval (Zhou et al., Nov 2023), etc but these are restricted to a certain set of use cases. In general, it is crucial to design evaluations with a target application in mind.

• Supervisor가 없고 reward signal만 있다.
• 피드백은 즉각적이지 않고 지연된다.
• '시간'은 중요하다. (Time really matters)
• 에이전트의 액션은 이후에 받는 데이터에 영향을 끼친다. (Agent's actions affect the subsequent data it receives)

The RL Problem

Rewards

• A reward $R_t$ is a scalar feedback signal 리워드는(보상은) 숫자이다.
• Indicates how well agent is doing at step t 에이전트가 step t에 얼마나 잘했는지 보여줌.
• The agent's job is to maximise cumulative reward 에이전트는 누적 보상을 최대화하는 것이 목적.

보상가설은 다음과 같다. All goals can be described by the maximisation of expected cumulative reward "모든 목표는 예상되는 누적 보상의 최대화로 설명될 수 있다."

Agent and Environment 에이전트와 환경

At each step t the agent:

• Executes action $A_t$
• Receives observation $O_t$
• Receives scalar reward $R_t$

The environment:

• Receives action $A_t$
• Emits observation $O_{t+1}$
• Emits scalar reward $R_{t+1}$

t increments at environment step

History and State

History는 observations, actions, rewards의 sequence $H_t = A_1, O_1, R_1, ..., A_{t-1}, O_t, R_t$

history에 따라서 에이전트는 액션을 선택하고, 환경은 ovservations/rewards를 선택한다.

State is the information used to determine what happens next state는 history의 function: $S_t=f(H_t)$

Environment State $S^e_t$

env state는 환경의 private representation. 환경이 다음 observation/reward를 선택하기 위해 사용하는 whatever 데이터임. 환경 state는 보통 에이전트가 모름.

Agent State $S^a_t$

에이전트의 internal representation. 에이전트가 다음 액션을 선택하기 위해 사용하는 whatever 정보임. RL 알고리즘이 사용하는 information.

Information State (== Markov state)

history의 모든 useful information을 갖고 있음. 정의 A state $S_t$ is Markov if and only if $P[S_{t+1}|S_t]=P[S_{t+1}|S_1, ..., S_t]$

“The future is independent of the past given the present” "현재가 주어졌을 때 미래는 과거와 독립적입니다." <- 현재만 중요. 과거는 버림. 과거는 현재에 이미 표현되어 있음. Once the state is known, the history may be thrown away 상태를 알고 난 다음에는 history를 버릴 수 있습니다.

• The environment state $S^e_t$is Markov
• The history $H_t$ is Markov

Fully Observable Environments (Markov decision process, MDP)

에이전트는 env state를 다 볼 수 있음 Agent state = env state = information state

Partially Observable Environments (POMDP)

에이전트는 env를 indirectly 볼 수 있음 agent state =/= env state 에이전트는 own state representation $S^a_t$을 만들어야 함

• Complete history를 사용할 수도 있고,
• env state에 대한 _믿음(Beliefs)_으로 만들 수도 있고,
• RNN도 그 중 하나.

참고 강의 자료:

• https://www.davidsilver.uk/teaching/

참고 Textbook:

• http://incompleteideas.net/book/the-book.html