KR20110077720A

KR20110077720A - 프로그램 코드의 변환 방법

Info

Publication number: KR20110077720A
Application number: KR1020090134367A
Authority: KR
Inventors: 조승모; 최종덕; 이재진
Original assignee: 삼성전자주식회사; 서울대학교산학협력단
Priority date: 2009-12-30
Filing date: 2009-12-30
Publication date: 2011-07-07
Also published as: US9015683B2; KR101613971B1; US20110161944A1

Abstract

복수 개의 워크 아이템들이 하나의 컴퓨팅 유닛에 포함된 복수 개의 프로세싱 엘리먼트에 각각 할당되어 동시에 수행되도록 작성된 프로그램 코드의 변환 방법이 제공된다. 프로그램 코드를 변환하는 변환기는, 프로그램 코드에 대하여, 복수 개의 워크 아이템이 복수 개의 프로세싱 엘리먼트보다 적은 개수의 프로세싱 엘리먼트에서 직렬적으로 수행되도록, 프로그램 코드에 포함된 동기화 배리어 함수를 기준으로 워크 아이템 합체 루프로 감싸질 2 이상의 프로그램 코드 영역을 식별하는 동작과, 식별된 프로그램 코드 영역 각각에 대하여 워크 아이템 합체 루프로 감싸는 동작을 수행한다.

오픈CL, 커널, 워크 아이템, 루프, 프로그램 변환

Description

프로그램 코드의 변환 방법{Method for transforming program code}

프로그램의 컴파일링에 관한 것으로, 더욱 상세하게는, 병렬 데이터 처리를 위해 작성된 프로그램의 변환 방법에 관한 것이다.

오픈CL은 CPU, GPU, 액셀러레이터(Accelerator) 등의 이종 프로세싱 플랫폼에서 병렬 프로그래밍을 하기 위한 표준이다. 오픈 CL은 다양한 플랫폼에서 동일한 프로그램 소스를 사용할 수 있도록 하여, 소프트웨어 개발자에게 이동성(portability)과 여러 플랫폼의 프로세싱 파워를 쉽게 이용할 수 있도록 해준다.

복수 개의 프로세싱 엘리먼트에서 처리되도록 작성된 오픈CL 커널 코드를 제한된 개수의 프로세싱 엘리먼트를 가진 장치에서 효율적으로 이용가능하도록 변환하는 방법에 관한 것이다.

일 측면에 따른, 복수 개의 워크 아이템들이 하나의 컴퓨팅 유닛에 포함된 복수 개의 프로세싱 엘리먼트에 각각 할당되어 동시에 수행되도록 작성된 프로그램 코드의 변환 방법은, 프로그램 코드에 대하여, 복수 개의 워크 아이템이 복수 개의 프로세싱 엘리먼트보다 적은 개수의 프로세싱 엘리먼트에서 직렬적으로 수행되도록, 프로그램 코드에 포함된 동기화 배리어 함수를 기준으로 워크 아이템 합체 루프로 감싸질 2 이상의 프로그램 코드 영역을 식별하는 동작과, 식별된 프로그램 코드 영역 각각에 대하여 워크 아이템 합체 루프로 감싸는 동작을 포함한다.

프로그램 코드는 오픈CL의 커널 코드일 수 있다.

워크 아이템 합체 루프는, 복수 개의 워크 아이템들의 식별자의 차원에 대응하는 중첩 루프(nested loop)일 수 있다. 복수 개의 워크 아이템들의 식별자가 3차원 식별자인 경우, 루프문은 3중첩 루프(triply nested loop)일 수 있다.

프로그램 코드에서 이용되는 프라이빗 변수가 식별된 2이상의 영역 중 하나의 영역에서 정의되고, 다른 영역에서 이용되는 경우, 프라이빗 변수를 워크 아이템의 식별자에 대한 차원을 가지도록 확장하는 단계를 더 포함할 수 있다. 확장된 프라이빗 변수를 저장하기 위하여 메모리를 할당하고, 할당된 메모리를 해제하기 위한 코드를 추가하는 단계를 더 포함할 수 있다.

2 이상의 코드 영역을 식별하는 단계는, 배리어 함수가 조건문 안에 있는 경우, 조건문 외부에 대하여, 조건문 이전의 적어도 하나의 스테이트먼트(statement)을 포함하는 영역 및 조건문 이후의 적어도 하나의 스테이트먼트를 포함하는 영역을 식별하는 단계와, 조건문 내부에서, 조건을 나타내는 스테이트먼트의 제1 결과에 따라 수행되는 배리어 함수 이전에 포함된 적어도 하나의 스테이트먼트를 포함하는 영역, 배리어 함수 이후에 포함된 적어도 하나의 스테이트먼트를 포함하는 영 역, 및 조건을 나타내는 스테이트먼트의 제2 결과에 따라 수행되는 적어도 하나의 스테이트먼트를 포함하는 영역을 식별하는 단계를 포함할 수 있다. 조건문 이전의 적어도 하나의 스테이트먼트를 포함하는 영역에 통합되어, 워크 아이템 합체 루프로 감싸지는, 조건문의 결과 값을 저장하는 단계를 더 포함할 수 있다.

2 이상의 코드 영역을 식별하는 단계는, 배리어 함수가 루프문 안에 있는 경우, 루프문 외부에 대하여, 루프문 이전의 적어도 하나의 스테이트먼트을 포함하는 영역 및 루프문 이후의 적어도 하나의 스테이트먼트를 포함하는 영역을 식별하는 단계와, 루프문 내부에서, 루프에 포함된 조건 스테이트먼트를 포함하는 영역, 배리어 함수 이전에 포함된 적어도 하나의 스테이트먼트를 포함하는 영역, 및 배리어 함수 이후에 포함된 적어도 하나의 스테이트먼트를 포함하는 영역을 식별하는 단계를 포함할 수 있다.

복수 개의 프로세싱 엘리먼트에서 처리되도록 작성된 오픈CL 커널 코드를 제한된 개수의 프로세싱 엘리먼트를 가진 장치에서 효율적으로 이용가능하도록 변환하여, 오픈CL 커널 코드를 다양한 플랫폼에서 실행할 수 있으므로, 오픈CL 커널 코드의 이동성을 향상시킬 수 있다. 또한, 프로세싱 엘리먼트를 가상화하는테 필요한 컨텍스트 스위칭 오버헤드를 줄여 성능을 높일 수 있다.

이하, 첨부된 도면을 참조하여 본 발명의 일 실시예를 상세하게 설명한다. 본 발명을 설명함에 있어 관련된 공지 기능 또는 구성에 대한 구체적인 설명이 본 발명의 요지를 불필요하게 흐릴 수 있다고 판단되는 경우에는 그 상세한 설명을 생략할 것이다. 또한, 후술되는 용어들은 본 발명에서의 기능을 고려하여 정의된 용어들로서 이는 사용자, 운용자의 의도 또는 관례 등에 따라 달라질 수 있다. 그러므로 그 정의는 본 명세서 전반에 걸친 내용을 토대로 내려져야 할 것이다.

도 1은 오픈CL 프로그램이 동작하는 시스템 구조의 일 예를 나타내는 도면이다.

오픈CL 프로그램은 도 1에 도시된 바와 같이, 호스트 프로세서(110) 및 하나 이상의 컴퓨트 디바이스(120, 130, 140, 150)를 포함하는 플랫폼에서 동작한다.

컴퓨트 디바이스(120, 130, 140, 150)는 하나 이상의 컴퓨트 유닛을 포함한다. 일예로, 컴퓨트 디바이스(120)는, 복수 개의 컴퓨트 유닛(121, 122, 123)을 포함할 수 있다. 각 컴퓨트 유닛은 하나 이상의 프로세싱 엘리먼트들을 포함한다. 일 예로, 컴퓨티 유닛(121)은 프로세싱 엘리먼트(10, 20, 30)를 포함할 수 있다. 여기에서, 프로세싱 엘리먼트는 가상 스칼라 프로세서이다. 각각의 프로세싱 엘리먼트(10, 20, 30)는 프라이빗 메모리(11, 21, 31)를 각각 포함할 수 있다.

오픈CL 프로그램은 호스트 프로세서에서 실행되는 호스트 프로그램이 커널을 디바이스에서 실행하는 형태이다. 오픈CL 애플리케이션은 호스트 프로그램 및 커널들로 구성된다.

호스트 프로그램은 호스트 프로세서(110)상에서 동작하며, 컴퓨트 디바이스내(120)의 프로세싱 엘리먼트들(10, 20, 30)에 대한 계산을 수행하는 명령어들 또는 메모리 오브젝트들을 조작하는 명령어들을 출력한다.

커널(kernel)은 오픈CL 디바이스에서 실행되는 함수를 의미한다. 커널은 오픈CL C로 작성된다. 커널은 커널 코드 또는 커널 함수로 불릴 수 있으며, 이하에서는 커널 코드라 한다. 커널 코드는 복수 개의 스테이트먼트(statement)로 구성될 수 있다. 호스트 프로그램이 커널 코드를 실행할 때에는, 워크 그룹의 수와 워크 아이템의 수를 지정한다.

여기에서, 워크 그룹은 워크 아이템들의 집합을 말한다. 워크 아이템은 병렬 실행되는 커널의 실행 단위이다. 하나 이상의 워크 그룹은 하나 이상의 워크 아이템을 포함한다. 하나의 워크 그룹은 하나의 컴퓨트 유닛에 할당되며, 그 워크 그룹에 포함된 각각의 워크 아이템은 그것이 하나 이상의 가상 엘리먼트들에 동시에 실행되는 것처럼 동작한다.

오픈CL에서 커널은 호스트 프로그램에 의해 n차원 추상적 인덱스 공간(n-dimensional abstract index space)으로 정의된다. 여기에서, N은 1≤N≤3의 범위에 있다. 인덱스 공간에서 각 포인트는 각 차원이 0에서 시작하는, N-튜플 정수(N-tuple of integers)로 특정된다. 각 포인터는 워크 아이템으로 불리는 커널의 실행 인스턴스와 관련된다.

N-튜플은 대응하는 워크 아이템의 글로벌 ID를 나타낸다. 각 워크 아이템은 자신의 글로벌 ID를 질의할 수 있고, 글로벌 ID에 따라서 상이한 태스크를 수행하고 상이한 데이터를 액세스할 수 있다. 여기에서, 태스크는 싱글 프로그램, 멀티플 데이터(SPMD)일 수 있다. 길이 N의 정수 어레이(즉, 인덱스 공간의 차원)는 인덱스 공간의 각 차원에서 워크 아이템의 개수를 특정하고, 커널 커맨드가 인큐 될(enqueue) 때 커널에 대해 호스트 프로그램에 의해 준비된다.

각 워크 그룹은 N-튜플인 유일한 ID를 가진다. 길이 N의 정수 어레이(즉, 인덱스 공간의 차원)는 인덱스 공간의 각 차원에서 워크 그룹들의 개수를 특정한다. 워크 그룹에서 워크 아이템은 워크 그룹내에 유일한 로컬 ID로 할당되고, 전체 워크 그룹을 로컬 인덱스 공간으로서 처리한다. 워크 아이템의 글로벌 ID는 그 로컬 ID 및 워크 그룹 ID로 계산될 수 있다. 워크 그룹에서 워크 아이템들은 동시에 하나의(single) 컴퓨트 유닛의 프로세싱 엘리먼트들상에서 동시에 동작한다.

다음으로, 오픈CL이 이용하는 메모리 공간에 대하여 설명한다.

오픈CL은 4개의 구별되는 메모리 영역 즉, 글로벌, 컨스턴트, 로컬 및 프라이빗을 정의한다. 컴퓨트 디바이스 메모리는 글로벌 메모리 영역 및 컨스턴트 메모리 영역들로 구성된다. 커널은 컨스턴트 메모리 및 로컬 메모리에 메모리 오브젝트들을 정적으로 할당할 수 있다. 로컬 메모리는 동일한 워크 그룹에서 모든 워크 다이템들에 의해 공유된다. 프라이빗 메모리는 하나의 워크 아이템에 국소적(local)이며, 프로세싱 엘리먼트에 포함된다.

이들 영역은 워크 아이템들에 의해 접근가능하며, 오픈CL C는 이들 메모리를 구별짓는 4개의 어드레스 공간 수식어(qualifier) 즉, __global, __constant, __local 및 __private를 가진다. 이러한 어드레스 공간 수식어는 커널 코드에서 변수 선언(declaration)에 이용된다.

전술한 바와 같이, 복수 개의 프로세싱 엘리먼트들에 각각 워크 아이템이 할당되어 실행되도록 작성된 커널 프로그램이, 실행되어야 하는 워크 아이템의 개수 보다 적은 개수의 프로세싱 엘리먼트를 가진 컴퓨트 유닛(또는 프로세서)에서 실행되기 위해서는, 적은 개수의 프로세싱 엘리먼트들의 워크 아이템들을 순차적으로 수행하여야 한다.

이하에서, 수행되어야 하는 복수 개의 워크 아이템이 하나의 가상 프로세싱 엘리먼트에서 동작하는 것을 가정한다. 하나의 워크 그룹이 컴퓨트 유닛에 할당될 때, 컴퓨트 유닛에 할당된 워크 그룹에 대응하는 쓰레드는 워크 그룹에서 각각의 워크 아이템을 하나씩 실행하면서 컨텍스트 스위칭을 수행해야 한다. 예를 들어, 100개의 워크 아이템이 실행되어야 하는 경우, 하나의 프로세싱 엘리먼트는 100개의 워크 아이템을 순차적으로 수행하기 위하여, 커널 코드를 매번 호출하여야 한다.

한편, 오픈CL에서 워크 그룹 배리어 함수(work-group barrier function)을 제공한다. 워크 그룹 배리어 함수는 커널에 의해 이용되며, 커널을 실행하는 단일 워크 그룹내에서 워크 아이템들을 동기화한다. 워크 그룹 배리어 함수는 동기화 배리어 함수, 배리어 함수 또는 배리어로 부를 수 있다. 워크 그룹에서 모든 워크 아이템은 배리어를 실행해야 하고, 그 워크 그룹내의 모든 다른 워크 아이템들이 배리어에 도달할 때까지 배리어 함수 너머에 있는 스테이트먼트로 진행할 수 없다.

전술한 바와 같이, 실행되어야 하는 워크 아이템의 개수보다 적은 개수의 프로세싱 엘리먼트를 가진 컴퓨트 유닛에서 실행되기 위해서, 워크 아이템이 워크 그룹 배리어에 도달할 때, 컴퓨트 유닛에서 실행되는 쓰레드는 워크 아이템의 프라이 빗 메모리의 콘텐츠 및 프로세싱 상태를 컴퓨트 유닛 외부에 위치하며 호스트 프로세서(110)가 이용하는 메인 메모리(도시되지 않음)의 유보 공간(reserved space)에 저장하고, 그런 다음 워크 아이템으로 스위칭해야 한다.

이와 같이, 모든 워크 아이템들에 대한 커널 코드를 실행하여 배리어에 도달하면, 각 워크 아이템의 프라이빗 메모리의 콘텐츠 및 저장된 컨텍스트를 복원한 후, 배리어 너머의 프로그램 포인트로부터 다시 하나씩 스테이트먼트를 실행한다. 워크 아이템 컨텍스트 스위칭은 set jmp() 및 long jmp()함수를 이용하여, 프로세서 상태 및 프라이빗 메모리를 저장하고 복원하는 과정을 수행하여야 하므로, 상당한 오버헤드가 발생된다.

일 실시예에 따르면, 워크 아이템 합체(coalescing)방법을 이용하여, 하나의 프로세싱 엘리먼트에서 복수 개의 워크 아이템을 효율적으로 직렬화하여 수행한다. 워크 아이템 합체는 커널 코드에 배리어들이 존재할 때, 단일 워크 그룹에서 워크 아이템들의 동시 실행을 직렬화하여, 단일 프로세싱 엘리먼트(또는 액셀레이터 코어)에 직렬화된 워크 아이템들이 순차적으로 실행되도록 커널 코드를 변환하는 방법이다.

일 실시예에 따른 워크 아이템 합체를 수행하는 모듈은 병렬 프로세싱을 위해 작성된 커널 코드를 직렬화하는 동작을 수행하므로, 소스-투-소스 변환기라고 부를 수 있다. 소스-투-소스 변환기는 병렬 프로세싱을 위해 작성된 커널 코드를 직렬화할 필요가 있는 장치의 프로세서, 예를 들어, CPU에서 실행될 수 있다. 소스-투-소스 변환기에서 직렬처리되도록 변환된 커널 코드은 어셈블리 코드로 컴파 일링되어, 타겟 머신에서 수행될 수 있다. 이하에서는, 소스-투-소스 변환기를 커널 코드 변환기로 부른다.

도 2a는 C 프로그램의 일 예이고, 도 2b는 도 2a의 프로그램에 대한 커널 코드의 일 예이고, 도 2c는 도 2b의 커널 코드를 일 실시예에 따른 커널 코드 변환 방법에 의해 변환한 결과의 일 예를 나타낸다.

도 2a는 2개의 벡터를 더하는 C 함수를 나타낸다. 도 2a에서는 for 함수에서, 변수 i를 1씩 증가시키면서 벡터를 더하는 동작이 순차적으로 진행된다.

도 2b는 도 2a에 제시된 C 프로그램과 같이, 2개의 벡터를 더하는 동작이 복수 개의 프로세싱 엘리먼트에서 동시에 수행되도록 작성된 커널 코드이다. 함수 get_global_id(0)은 커널을 실행하는 워크 아이템의 글로벌 ID의 제1 엘리먼트를 되돌린다. 여기에서, 글로벌 ID는 3차원 (a, b, c)의 형식으로, a, b 및 c로 이루어진 메모리 어레이에 포함된 메모리 오브젝트에 대한 포인터를 나타낼 수 있다. 호스트 프로그램은 커널 코드에 메모리 오브젝트에 대한 포인터를 제공함으로써, 복수 개의 프로세싱 멜리먼트에서 메모리 벡터 합산이 동시에 병렬적으로 수행될 수 있다.

도 2c는 도 2b의 커널 코드를 일 실시예에 따른 커널 코드의 변환 방법에 따라 변환한 결과를 나타낸다. 도 2b 및 도 2c의 코드를 비교하면, 도 2b의 코드에 3차원 for 루프가 추가되었음을 알 수 있다.

도 2c의 변환된 코드는, 단일 워크 그룹내의 모든 워크 아이템들을 하나씩 순차적으로 수행한다. 워크 그룹이 컴퓨트 유닛에 할당되면, 컴퓨트 유닛은 커널 에 대해 변환된 코드를 실제적으로 실행하여, 컴퓨트 유닛에서 가상 프로세싱 엘리먼트들을 에뮬레이트한다.

워크 그룹 인덱스 공간의 크기는 오픈CL 런타임에서 제공되는 어레이 __local_size에 의해 결정된다. 어레이 __global_id는 워크 그룹에서 그 로컬 ID가 (0, 0, 0)인 워크 아이템의 글로벌 ID를 저장하고 있다.

복수 개의 워크 아이템이 하나의 프로세싱 엘리먼트에서 수행되어야 하는 경우, 도 2c와 같이, 커널 코드를 감싸는 루프를 적용하면, 실행되어야 하는 커널 코드를 하나의 워크 아이템을 실행할 때마다 호출할 필요가 없으며, 한 번의 커널 코드의 호출로 복수 개의 워크 아이템을 순차적으로 실행할 수 있다. 또한, 커널 코드에 배리어가 포함되어 있는 경우 프라이빗 메모리의 데이터 및 프로세스 상태 정보를 저장 및 복원해야 하는 컨텍스트 오버헤드를 감소시킬 수 있다. 이와 같이, 커널 코드를 직렬화하여 수행하기 위한 루프는 실제로 복수 개의 워크 아이템을 함께 묶어서 실행하게 되므로 이하에서는, 워크 아이템 합체 루프(work-item coalescing loop; WCL)라고 부른다.

도 3은 커널 코드의 변환 방법의 일 예를 나타내는 흐름도이다.

커널 코드 변환기는 프로그램 코드에 대하여, 복수 개의 워크 아이템이 복수 개의 프로세싱 엘리먼트보다 적은 개수의 프로세싱 엘리먼트에서 직렬적으로 수행되도록, 프로그램 코드에 포함된 동기화 배리어 함수를 기준으로 워크 아이템 합체 루프로 감싸질 2 이상의 코드 영역을 식별한다(210). 이를 위해, 커널 코드 변환기는 커널 코드를 입력받고, 입력된 커널 코드에 대하여 추상적 신택스 트리(AST) 를 생성하고, 워크 아이템 합체 루프(WCL)로 감싸질 커널 코드 영역을 식별할 수 있다.

워크 아이템 합체 루프(WCL)로 감싸게 될 필요가 있는 커널 코드 영역을 워크 아이템 합체 영역(work-item coalescing region; WCRs)이라고 부를 수 있다.

이를 위해, 커널 코드 변환기는 배리어 함수 콜을 포함하는 AST의 노드 및 배리어 함수 콜을 포함하는 AST 노드의 서브 트리에 대하여 마킹할 수 있다. 각각의 추상적 신택스 트리(AST) 마킹한 후, 커널 코드 변환기는 일 실시예에 따른 워크 아이템 합체 루프(WCL)를 이용하여 커널 코드를 변환하는 함수를 호출할 수 있다.

커널 코드를 변환하는 함수는 마킹된 정보를 이용하여, 워크 아이템 합체 영역(WCR)은 (i)커널 코드에 포함된 배리어가 없는 경우, (ⅱ)커널 코드에 배리어가 포함된 경우, (ⅲ)배리어가 조건문 내에 있는 경우, (ⅳ)배리어가 루프문 내에 있는 경우에 따라 다른 방식으로 식별되도록 정의될 수 있다. 나머지 경우, 예를 들어, 커널 코드에 for, do-while, 및 switch와 같은 다른 언어가 포함된 경우에는, 위 4가지 경우에 적용된 워크 아이템 합체 영역(WCR) 식별 방법을 변형하여 식별될 수 있다.

또한, 커널이 배리어를 포함하는 함수를 호출하는 스테이트먼트를 포함할 때, 호출될 배리어를 포함하는 함수는 스태틱 콜 그래프에서 그 앤세스터들 노드에 직렬로 부가(in-line)될 수 있다. 스태틱 콜 그래프는 함수 간의 호출 관계를 나타내는 그래프로서 커널 코드 변환기에서 작성될 수 있다. 그런 다음, 변형된 스태틱 콜 그래프에 대하여, 전술한 규칙에 따라 워크 아이템 합체 영역(WCR)이 식별된 다음, 워크 아이템 루프가 적용될 수 있다.

하나의 워크 아이템 합체 영역(WCR)에서 정의되며 다른 워크 아이템 합체 영역(WCR)에서 이용되는 워크 아이템 프라이빗 변수는 각 워크 아이템에 대하여 별개 저장 위치를 필요로 한다. 따라서, 커널 코드 변환기는 각 워크 아이템의 프라이빗 변수에 대한 확장을 수행한다(220).

일 실시예에 따른 변수 확장 기법은, 스칼라 확장과 유사하다. 스칼라 확장은 부가 메모리의 비용면에서 의존성을 제거하기 위하여 하나의 중첩 루프(loop nest)에 적용되지만, 일 실시예에 따른 변수 확장 기법은, 하나의 중첩 루프에 정의된 데이터를 다른 중첩 루프에 전달하기 위하여 적용되는 점에 차이가 있다. 변수 확장에 대해서는 도 8a 및 도 8b를 참조하여 상세하게 후술한다. 커널 코드에 하나의 워크 아이템 합체 영역(WCR)에서 정의되며 다른 워크 아이템 합체 영역(WCR)에서 이용되는 워크 아이템 프라이빗 변수가 존재하지 않는 경우에는, 동작 220는 수행되지 않을 수 있다.

커널 코드 변환기는 워크 아이템 합체 영역(WCR)로 식별된 커널 코드 영역을 워크 아이템 합체 루프(WCL)로 감싼다(230). 오픈CL 플랫폼에 정의된 바와 같이, 워크 아이템에 대하여 3차원의 ID가 부여되는 경우, 일 실시예에 따른 커널 코드 변환기는 단일 워크 그룹의 인덱스 공간에서 반복하는 트리플리 네스티드 루프(triply nested loop)로 커널 코드를 감쌀 수 있다. 그러나, 워크 아이템의 ID는 3차원에 한정되지 않으므로, 커널 코드를 감쌀 워크 아이템 합체 루프(WCL)는 단일 루프 일 수 있고, 이중 루프일 수 있는 등, 다양한 변형이 가능하다.

도 4는 배리어 함수가 없는 경우의 커널 코드 변환 방법의 일 예를 나타내는 도면이다.

도 4에서, 블록(410)은 복합문(compound statement) 또는 스테이트먼트들의 리스트일 수 있다. 커널 코드에 배리어가 없는 경우에는, 커널 코드의 시작과 끝은 워크 아이템의 루프로 변형한다. 이와 같이 변형함으로써, 워크 그룹 내에서 커널 함수의 호출을 줄일 수 있다.

도 5는 커널 코드가 배리어 함수를 포함하는 경우의 커널 코드 변환 방법의 일 예를 나타내는 도면이다.

도 5에서, 블록(510, 520)은 복합문 또는 스테이트먼트들의 리스트일 수 있다. 워크 아이템 합체 영역은 배리어(520)를 기준으로 블록(510)을 포함하는 영역 및 블록(520)을 포함하는 영역으로 식별될 수 있다. 커널 코드 변환기는 식별된 각 영역을 워크 아이템의 루프로 변형한다. 여기에서, 블록(510) 또는 블록(520)에 스테이트먼트가 없을 수도 있으며, 스테이트먼트가 없는 영역은 워크 아이템 루프로 감싸지는 동작이 수행되지 않을 것이다.

커널 코드 변환기는 배리어 사이에 이동하는 스택 변수의 경우 변수 확장을 수행하고, 확장된 변수를 저장하기 위한 메모리 할당 및 해제를 위해 블록(540, 550, 560)에서와 같이 malloc() 함수 및 free()함수를 추가할 수 있다.

도 6은 커널 코드의 조건문 안에 배리어 함수가 포함되는 경우의 커널 코드 변환 방법의 일 예를 나타내는 도면이다.

도 6에서, 블록(620)은 if와 같이 조건을 나타내는 스테이트먼트(conditional statement)이고, 블록(610, 630, 650, 660, 670)은 복합문 또는 스테이트먼트들의 리스트일 수 있다. 또한, 이중, 블록(620, 630, 650, 660)은 조건문을 나타내는 스테이트먼트(또는 코드)이다. 오픈CL 규격에 따르면, 조건을 나타내는 스테이트먼트 내에서 배리어가 있으면, 모든 워크 아이템들은 조건적이 되며, 모든 워크 아이템이 조건 스테이트먼트를 실행하고, 배리어를 실행하도록 정의되어 있다.

도 6에 도시된 바와 같이, 배리어 함수(640)가 조건문 안에 있는 경우, 커널 코드 변환기는 조건문 외부 및 조건문 내부에 대하여, 워크 아이템 합체 영역(WCR)을 식별한다. 커널 코드 변환기는, 조건문 외부에 대하여, 조건문(620, 630, 650, 660) 이전의 적어도 하나의 스테이트먼트(statement)을 포함하는 영역(즉, 블록(610)을 포함하는 영역) 및 조건문 이후의 적어도 하나의 스테이트먼트를 포함하는 영역(즉, 블록(670)을 포함하는 영역)을 식별한다.

또한, 조건문 내부에서, 조건을 나타내는 스테이트먼트(620)의 제1 결과에 따라 수행되는 배리어 함수 이전에 포함된 적어도 하나의 스테이트먼트를 포함하는 영역(즉, 블록(630)을 포함하는 영역)), 배리어 함수 이후에 포함된 적어도 하나의 스테이트먼트를 포함하는 영역(즉, 블록(650)을 포함하는 영역), 및 조건을 나타내는 스테이트먼트(620)의 제2 결과에 따라 수행되는 적어도 하나의 스테이트먼트를 포함하는 영역(즉, 블록(660)을 포함하는 영역)을 식별할 수 있다.

또한, 커널 코드 변환기는 조건을 나타내는 조건문(620)이 각 워크 아이템마 다 변하지 않으므로, 커널 코드 변환기는 조건문(620)의 실행 결과 값을 블록(612)과 같이 미리 저장한 다음, 블록(621)에서 미리 저장된 결과 값을 이용하는 방식으로 변형할 수 있다.

조건문(620)의 실행 결과 값은 블록(610)과 통합되어 하나의 워크 아이템 루프로 감싸질 수 있다. 또한, 블록(630)을 포함하는 영역, 블록(650)을 포함하는 영역, 및 블록(660)을 포함하는 영역도 각각 워크 아이템 루프로 감싸질 수 있다.

도 7은 커널 코드의 루프문 안에 배리어 함수가 포함되는 경우의 커널 코드 변환 방법의 일 예를 나타내는 도면이다.

도 7에서, 블록(720)을 조건을 나타내는 스테이트먼트이고, 블록(710, 730, 750, 760)은 복합문 또는 스테이트먼트들의 리스트일 수 있다. 또한, 블록(720, 730, 750)은 루프문을 나타내는 코드이다.

오픈CL 규격에 따르면, 배리어가 루프 내부에 있으면, 모든 워크 아이템들이 루프의 각 반복에 대해 배리어를 실행한 후에, 배리어를 넘어 실행을 계속하도록 허용된다. 워크 아이템이 일단 컨디셔널 또는 루프로 들어가면, 그 구성 내부의 각 스테이트먼트는 워크 그룹 로컬 인덱스 공간을 반복함으로써 워크 그룹에서 모든 워크 아이템들에 대해 실행될 수 있으며, 루프의 반복 횟수는 워크 그룹에서 모든 워크 아이템들에 대해 동일하다.

도 7에 도시된 바와 같이, 배리어 함수(740)가 루프문 안에 있는 경우, 커널 코드 변환기는 루프문(720, 730, 750) 외부 및 루프문(720, 730, 750) 내부에 대하여, 워크 아이템 합체 영역(WCR)을 식별한다. 루프문 외부의 코드는 워크 아이템 별로 한번씩만 실행되지만, 루프문 내부의 코드는 각 워크 아이템이 루프가 반복되는 횟수만큼 반복하여 실행되기 때문에 워크 아이템 합체 영역(WCR)이 분리되어 식별될 수 있다.

커널 코드 변환기는, 루프문 외부에 대하여, 루프문(720, 730, 750) 이전의 적어도 하나의 스테이트먼트를 포함하는 영역(즉, 블록(710)을 포함하는 영역) 및 루프문(720, 730, 750) 이후의 적어도 하나의 스테이트먼트를 포함하는 영역(즉, 블록(760)을 포함하는 영역)을 식별한다.

커널 코드 변환기는, 루프문(720, 730, 750) 내부에서, 루프의 조건을 나타내는 스테이트먼트 영역(즉, 블록(720)을 포함하는 영역), 배리어 함수 이전에 포함된 적어도 하나의 스테이트먼트를 포함하는 영역(즉, 블록(730)을 포함하는 영역), 및 배리어 함수 이후에 포함된 적어도 하나의 스테이트먼트를 포함하는 영역(즉, 블록(750)을 포함하는 영역)을 식별한다.

또한, 커널 코드 변환기는 조건을 나타내는 스테이트먼트(720)가 각 워크 아이템마다 변하지 않으므로, 커널 코드 변환기는 스테이트먼트(720)의 실행 결과 값을 블록(721)과 같이 미리 저장한 다음, 블록(722)에서 미리 저장된 결과 값을 이용하는 방식으로 변형할 수 있다.

결과적으로, 블록(710)을 포함하는 영역, 블록(721)을 포함하는 영역, 블록(730)을 포함하는 영역, 블록(750)을 포함하는 영역, 및 블록(760)을 포함하는 영역이 각각 워크 아이템 루프로 감싸질 수 있다. 또한, 원래의 루프는 루프 내의 워크 아이템 합체 루프(WCL)를 감싸는 큰 루프로 변형된다.

또한, 루프 내부에서 변수 확장이 필요한 경우에는, 루프 내부에서 변수 확장을 위한 메모리 할당 및 해제가 반복됨으로 인한 성능 저하를 막기 위하여, 루프의 시작 전에 메모리를 할당하고, 루프 실행 종료 후에 메모리를 해제하도록 코드가 변형될 수 있다.

도 8a는 웹 식별 후의 컨트롤 플로우 그래프의 일 예를 나타내고, 도 8b는 변수 확장 후의 컨트롤 플로우 그래프의 일 예를 나타내는 도면이다.

커널 코드 변환기는 변수 확장을 위해, 커널의 소스 레벨의 컨트롤 플로우 그래프를 구성한다. 그런 다음, 커널 코드 변환기는 워크 아이템 프라이빗 변수들에 대한 데이터 플로우를 분석하여, 컨트롤 플로우 그래프의 도미네이트 트리 및 포스트 도미네이터 트리를 구성한다.

커널 코드 변환기는, 컨트롤 플로우 그래프에서 각 프라이빗 변수에 대해 듀-체인(definition-use chain; du-chain)을 찾는다. 듀 체인에서, 각 변수에 대하여 웹들을 찾는다. 커널 코드 변환기는 서로 다른 WCR간에 걸쳐서 존재하는 듀 체인을 갖는 웹을 찾아서, 웹에 속한 변수들에 대하여 변수를 확장한다.

일 실시예에 따른 변수 확장 알고리즘에서, 웹(web)은 확장의 단위이다. 변수에 대한 웹은 변수의 공통적인 사용을 포함하는 변수의 모든 듀-체인을 나타낸다. 변수에 대한 듀-체인은 컨트롤 플로우 그래프에서 변수를 정의한 노드, 정의된 변수가 이용된 노드를 연결한다. 변수의 확장 단위로 웹을 이용하면 변수의 참조를 모두 확장할 필요가 없다는 이점이 있다. 변수는 워크 아이템의 식별 정보의 차수에 따라 확장될 수 있다.

도 8a는 변수 x에 대한 웹들을 식별한 후의 컨트롤 플로우 그래프를 포함한다. 점선으로 표시된 사각형은 워크 아이템 합체 영역(WCR)을 나타낸다. 컨트롤 플로우 그래프에서 이용되는 프라이빗 변수를 모두 x로 표기하였으나, 각각의 컨트롤 플로우 그래프상의 노드에서 이용되는 변수는 "x=..."와 같이 새로 정의되는 경우 서로 다른 변수를 나타낸다.

도 8a에는 6개의 듀-체인({①}, {②}, {③}, {④}, {⑤,⑥},{⑦})이 있다. 이들 듀 체인은 4개의 상이한 웹들({①}, {②, ③}, {④,⑤,⑥},{⑦})을 생성한다. 웹들({①}, {④,⑤,⑥})에서 변수 x의 정의 및 사용은 동일한 워크 아이템 합체 영역(WCR)에 속하지 않는다. 웹들({①}, {④,⑤,⑥})에서 이용되는 변수 x에 대해서는, 각 변수에 대한 별도의 저장 공간이 필요하므로, 변수 x에 대하여 3차원 어레이를 참조하도록 변수 x가 확장된다. 웹들({②, ③}, {⑦})에 포함된 변수 x는 동일한 워크 아이템 합체 영역(WCR)에 포함되므로, 확장될 필요가 없다.

도 8b는 웹들({①}, {④,⑤,⑥})에서 이용되는 변수 x에 대한 변수 확장 결과를 나타내는 컨트롤 플로우 그래프이다. 도 8b에서, 변수 x는 3차원 어레이 형식의 변수 x1([][][])으로 확장되었다.

한편, 확장된 변수를 저장하기 위하여, 커널 코드 변환기는 메모리 공간을 할당할 수 있다. 커널 코드 변환기는 컨트롤 플로우 그래프에서, 도미네이터 트리(dominator tree) 및 포스트 도미네이터 트리(post dominator tree)를 구성할 수 있다. 커널 코드 변환기는 도미네이터 트리에서 변수를 정의한 노드들 중에서 가장 낮은 공통적인 앤세스터(lowest common ancestor; LCA)를 식별하여, LCA를 포함 하는 워크 아이템 합체 영역(WCR) 이전에 메모리 할당 함수 malloc()를 추가하고, 포스트 도미네이터 트리에서 마찬가지로, LCA를 찾아서, 해당 노드를 석세서(successor)로 만든 다음, 석세서를 포함하는 워크 아이텝 합체 영역(WCR)이후에 메모리 해제 함수 free()를 추가할 수 있다.

본 발명의 일 양상은 컴퓨터로 읽을 수 있는 기록 매체에 컴퓨터가 읽을 수 있는 코드로서 구현될 수 있다. 상기의 프로그램을 구현하는 코드들 및 코드 세그먼트들은 당해 분야의 컴퓨터 프로그래머에 의하여 용이하게 추론될 수 있다. 컴퓨터가 읽을 수 있는 기록매체는 컴퓨터 시스템에 의하여 읽혀질 수 있는 데이터가 저장되는 모든 종류의 기록 장치를 포함한다. 컴퓨터가 읽을 수 있는 기록 매체의 예로는 ROM, RAM, CD-ROM, 자기 테이프, 플로피 디스크, 광 디스크 등을 포함한다. 또한, 컴퓨터가 읽을 수 있는 기록 매체는 네트워크로 연결된 컴퓨터 시스템에 분산되어, 분산 방식으로 컴퓨터가 읽을 수 있는 코드로 저장되고 실행될 수 있다.

이상의 설명은 본 발명의 일 실시예에 불과할 뿐, 본 발명이 속하는 기술분야에서 통상의 지식을 가진 자는 본 발명의 본질적 특성에서 벗어나지 않는 범위에서 변형된 형태로 구현할 수 있을 것이다. 따라서, 본 발명의 범위는 전술한 실시예에 한정되지 않고 특허 청구범위에 기재된 내용과 동등한 범위 내에 있는 다양한 실시 형태가 포함되도록 해석되어야 할 것이다.

도 2a는 C 프로그램의 일 예이고, 도 2b는 도 2a의 프로그램에 대한 커널 코드의 일 예이고, 도 2c는 도 2b의 커널 코드를 일 실시예에 따른 커널 코드의 변환 방법에 따라 변환한 결과의 일 예를 나타낸다.

Claims

복수 개의 워크 아이템들이 하나의 컴퓨팅 유닛에 포함된 복수 개의 프로세싱 엘리먼트에 각각 할당되어 동시에 수행되도록 작성된 프로그램 코드의 변환 방법으로서,

상기 프로그램 코드에 대하여, 상기 복수 개의 워크 아이템이 상기 복수 개의 프로세싱 엘리먼트보다 적은 개수의 프로세싱 엘리먼트에서 직렬적으로 수행되도록, 상기 프로그램 코드에 포함된 동기화 배리어 함수를 기준으로 워크 아이템 합체 루프로 감싸질 2 이상의 코드 영역을 식별하는 단계; 및

상기 식별된 프로그램 코드 영역 각각에 대하여 상기 워크 아이템 합체 루프로 감싸는 단계를 포함하는 프로그램 코드 변환 방법.
제1항에 있어서,

상기 프로그램 코드는 오픈CL의 커널 코드인 프로그램 코드 변환 방법.
제1항에 있어서,

상기 워크 아이템 합체 루프는, 상기 복수 개의 워크 아이템들의 식별자의 차원에 대응하는 중첩 루프(nested loop)인 프로그램 코드 변환 방법.
제3항에 있어서,

상기 복수 개의 워크 아이템들의 식별자가 3차원 식별자인 경우, 상기 루프문은 3중첩 루프(triply nested loop)인 프로그램 코드 변환 방법.
제1항에 있어서,

상기 프로그램 코드에서 이용되는 프라이빗 변수가 상기 식별된 2이상의 영역 중 하나의 영역에서 정의되고, 다른 영역에서 이용되는 경우,

상기 프라이빗 변수를 워크 아이템의 식별자에 대한 차원을 가지도록 확장하는 단계를 더 포함하는 프로그램 코드 변환 방법.
제5항에 있어서,

상기 확장된 프라이빗 변수를 저장하기 위하여 메모리를 할당하고, 상기 할당된 메모리를 해제하기 위한 코드를 추가하는 단계를 더 포함하는 프로그램 코드 변환 방법.
제1항에 있어서,

상기 2 이상의 코드 영역을 식별하는 단계는,

상기 배리어 함수가 조건문 안에 있는 경우,

상기 조건문 외부에 대하여, 조건문 이전의 적어도 하나의 스테이트먼트(statement)을 포함하는 영역 및 상기 조건문 이후의 적어도 하나의 스테이트먼트를 포함하는 영역을 식별하는 단계; 및

상기 조건문 내부에서, 상기 조건을 나타내는 스테이트먼트의 제1 결과에 따라 수행되는 상기 배리어 함수 이전에 포함된 적어도 하나의 스테이트먼트를 포함하는 영역, 상기 배리어 함수 이후에 포함된 적어도 하나의 스테이트먼트를 포함하는 영역, 및 상기 조건을 나타내는 스테이트먼트의 제2 결과에 따라 수행되는 적어도 하나의 스테이트먼트를 포함하는 영역을 식별하는 단계를 포함하는 프로그램 코드 변환 방법.
제7항에 있어서,

상기 조건문 이전의 적어도 하나의 스테이트먼트를 포함하는 영역에 통합되어, 워크 아이템 합체 루프로 감싸지는, 상기 조건문의 결과 값을 저장하는 단계를 더 포함하는 프로그램 코드 변환 방법.
제1항에 있어서,

상기 2 이상의 코드 영역을 식별하는 단계는,

상기 배리어 함수가 루프문 안에 있는 경우,

상기 루프문 외부에 대하여, 루프문 이전의 적어도 하나의 스테이트먼트을 포함하는 영역 및 상기 루프문 이후의 적어도 하나의 스테이트먼트를 포함하는 영역을 식별하는 단계; 및

상기 루프문 내부에서, 상기 루프에 포함된 조건 스테이트먼트를 포함하는 영역, 상기 배리어 함수 이전에 포함된 적어도 하나의 스테이트먼트를 포함하는 영 역, 및 상기 배리어 함수 이후에 포함된 적어도 하나의 스테이트먼트를 포함하는 영역을 식별하는 단계를 포함하는 프로그램 코드 변환 방법.