cuda inline assembly

소프트웨어 상에서의 최적화 구현은 assembly가 없이는 불가능하다.

이번에는 간단히 cuda 상에서의 assembly를 생각해보자.

__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x;



asm(".reg .u32 t1;\n\t"
"mul.lo.u32 %0, %1, %2; "

:"=r"(c[i])
:"r"(b[i]),"r"(a[i])
);


   
}


간단히 곱셈하는 inline을 짜보면 위와 같다.

여기서 기존의 assembly와 다른점을 3가지 정도 발견할 수 있다.

1 .reg로 선언해서 사용하고자 하는 register를 explicit하게 선언하는것이 가능하다.

2. 코드는 ';' 를 사용하여 한 줄이 끝났음을 명확히 한다.

3. \n\t 의 경우 다음줄로 내리고 앞 뒤 명령어 간의 공백을 주는 경우인데 맨 마지막 줄에는 붙이지 않는다..

이정도만 알고 instruction set만 알면 어느정도 짤수 있을듯.