KR102402584B1

KR102402584B1 - 사용자 어플리케이션의 특성에 따른 연산 디바이스 동적 제어 기법

Info

Publication number: KR102402584B1
Application number: KR1020150120484A
Authority: KR
Inventors: 전철용; 문지중; 가오 이; 위안 우; 김락기; 뱅거 라비쉬할; 장용석
Original assignee: 삼성전자주식회사
Priority date: 2015-08-26
Filing date: 2015-08-26
Publication date: 2022-05-27
Also published as: US10705813B2; WO2017034364A1; US20180253291A1; KR20170024898A

Abstract

본 개시는 단말 장치에서 적어도 하나의 연산 디바이스를 이용하여 어플리케이션을 수행하는 방법에 있어서, 상기 적어도 하나의 연산 디바이스의 처리 시간을 계산하는 동작; 사용자의 선호도 또는 상기 계산된 적어도 하나의 처리 시간에 근거하여 상기 어플리케이션을 수행할 소정 개수의 연산 디바이스를 선택하는 동작; 상기 결정된 소정 개수의 연산 디바이스에 대해 상응하는 사용률 정보를 이용하여 결정된 처리시간 함수를 최소화하는 워크로드를 결정하는 동작; 및 상기 소정 개수의 연산 디바이스에 상기 결정된 워크로드를 적용하여 상기 어플리케이션을 실행하는 동작을 포함하는 방법을 제공한다.

Description

사용자 어플리케이션의 특성에 따른 연산 디바이스 동적 제어 기법{SCHEME FOR DYNAMIC CONTROLLING OF PROCESSING DEVICE BASED ON APPLICATION CHARACTERISTICS}

본 개시는 다양한 연산 디바이스를 사용하는 단말에서 어플리케이션을 효율적으로 실행시키기 위한 컴파일러 모듈 및 런타임 모듈에 관련된 것이다.

사용자 어플리케이션을 단말과 같은 장치에서 동작시키기 위해서는 프로그램 언어로 작성되는 어플리케이션의 코드를 컴파일하는 과정 및 컴파일된 바이너리(즉, 실행 코드)를 실행하는 과정 등이 필요하다.

단말 내에는 바이너리를 실행하는 다양한 연산 디바이스(processing device)(즉, HW 모듈)를 포함할 수 있는데. 이러한 연산 디바이스의 대표적인 예는 CPU(central processing unit)이다. 상기 단말에는 상기 CPU 외에도, SIMD(single instruction multiple data), Parallel-SIMD, GPU(graphic processing unit), GPGPU(general purpose GPU), 또는 DSP(digital signal processor) 등이 포함될 수 있다.

아직까지, 단말 내 다양한 연산 디바이스들을 효과적으로 활용하기 위한 시도는 어플리케이션 개발자에 의해서만 수행되고 있다. 단말 내 연산 디바이스들에서 동작하는 어플리케이션은 HW(hardware) 특정적인(즉, 연산 디바이스에 특정적인) API(application programmable interface) (예를 들어, GPU의 경우 OpenCL)를 이용하여 개발된다. 즉, 상기 연산 디바이스를 위한 어플리케이션은 일반적으로 많이 사용되는 프로그램 언어 (C/C++, Java 등)와는 다른 언어를 통해 개발된다. 상기 HW 특정적인 API는 해당 어플리케이션이 해당 연산 디바이스에서 동작할 수 있도록 해주는 툴(Tool)의 역할은 충실히 수행하지만, 연산 디바이스들을 동시에 효율적으로 활용하는 것은 오로지 어플리케이션 개발자의 수작업에 의해 구현되고 있는 것이 현실이다.

따라서, 개발된 어플리케이션은 개발자에 의해 정해진 지시(즉, 스케줄링)에 따라서만 동작하게 된다. 다시 말해, 어플리케이션 개발자가 어플리케이션의 특정 코드를 특정 연산 디바이스에 동작 시키기 위해서는 상기 연산 디바이스를 위한 언어(또는 API)로 해당하는 코드를 작성해야 한다. 또한, 상기 개발자는 연산 디바이스들간의 동기화도 고려해야만 하는 상황을 직면한다.

어플리케이션의 특성은 어플리케이션 개발자에 전적으로 의존할 수 밖에 없다. 또한, 개발자의 구현 방식에 따라 달라지는 어플리케이션의 특성, 실제 어플리케이션이 장치 내에서 동작할 때의 연산 디바이스의 자원 활용 상태, 또는 어플리케이션이 계산해야 하는 데이터 크기 등은 개발자가 코드 작성시에는 쉽게 알 수 없다. 따라서, 어플리케이션이 단말 내 연산 디바이스를 효과적으로 활용하는 데는 한계가 있다.

본 개시는 어플리케이션의 특성과 어플리케이션 동작시의 단말 내 연산 디바이스의 상태 정보(예를 들어, 가용율, 동작 주파수) 등을 동적으로 이용함으로써, 어플리케이션이 단말 내 연산 디바이스(들)를 효율적으로 활용하게 하는 방법 및 장치를 제공한다.

또한 본 개시는, 각 연산 디바이스에서 수행 가능한 어플리케이션의 코드 부분을 분석하는 방법 및 장치를 제안하고, 상기 분석 작업의 복잡도를 감소시키는 방법 및 장치를 제공한다.

단말 장치에서 적어도 하나의 연산 디바이스를 이용하여 어플리케이션을 수행하는 방법에 있어서, 상기 적어도 하나의 연산 디바이스의 처리 시간을 계산하는 동작; 사용자의 선호도 또는 상기 계산된 적어도 하나의 처리 시간에 근거하여 상기 어플리케이션을 수행할 소정 개수의 연산 디바이스를 선택하는 동작; 상기 결정된 소정 개수의 연산 디바이스에 대해 상응하는 사용률 정보를 이용하여 결정된 처리시간 함수를 최소화하는 워크로드를 결정하는 동작; 및 상기 소정 개수의 연산 디바이스에 상기 결정된 워크로드를 적용하여 상기 어플리케이션을 실행하는 동작을 포함하는 방법을 제안한다.

또한 본 개시는 적어도 하나의 연산 디바이스를 이용하여 어플리케이션을 수행하도록 구성된 런타임 모듈을 구비하는 단말 장치에 있어서, 상기 런타임 모듈은: 상기 적어도 하나의 연산 디바이스의 처리 시간을 계산하고, 사용자의 선호도 또는 상기 계산된 적어도 하나의 처리 시간에 근거하여 상기 어플리케이션을 수행할 소정 개수의 연산 디바이스를 결정하고, 상기 결정된 소정 개수의 연산 디바이스에 대해 상응하는 사용률 정보를 이용하여 결정된 처리시간 함수를 최소화하는 워크로드를 결정하고, 상기 소정 개수의 연산 디바이스에 상기 결정된 워크로드를 적용하여 상기 어플리케이션을 수행하는 스케줄러; 및 상기 사용률 정보를 산출하는 장치 모니터를 포함하는 단말 장치를 제안한다.

본 개시에 따른 장치는 각각의 연산 디바이스를 효과적으로 사용하게 됨으로써, 어플리케이션 동작 속도를 가속하고, 소비 전력을 감소시킬 수 있다.

본 개시에 따른 장치는 컴파일러의 소스 코드 의존 관계 (예, 포인터) 분석 속도를 향상시킬 수 있다.

도 1은 본 개시에 따른 어플리케이션 실행 디바이스의 구성을 예시하는 도면;
도 2는 본 개시에 따른 컴파일러 모듈이 하나의 노드로 취급하는 노드 들의 조건을 설명하는 제약 그래프, 소스 코드 및 IR을 예시하는 도면;
도 3은 본 개시에 따른 런타임 모듈의 동작을 개략적으로 설명하는 도면;
도 4는 본 개시에 따라서 동작할 연산 디바이스를 결정하고 워크로드를 결정한 경우의 효과를 자원 사용률 및 사용 에너지 관점에서 예시하는 도면;
도 5는 본 개시에 따른 기법을 적용함에 있어서, 캐쉬의 입력 크기에 따른 연산 디바이스 별 처리 시간을 비교하는 표;
도 6은 본 개시에 따른 기법을 적용함에 있어서, CPU 및 GPU를 이용하는 경우의 부하 분담율에 따른 성능(처리 시간)을 예시하는 도면;
도 7은 본 개시에 따른 기법을 적용함에 있어서, CPU 및 GPU를 이용하는 경우의 워크로드에 따른 에너지 소비량을 예시하는 도면;
도 8은 본 개시에 따른 기법을 적용함에 있어서, CPU 및 GPU를 이용하는 경우에 워크로드와 사용률에 성능(처리시간)을 예시하는 도면;
도 9는 본 개시에 따른 컴파일러 모듈의 동작을 개략적으로 설명하는 도면이다.

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

본 개시의 자세한 설명에 앞서, 본 명세서에서 사용되는 몇 가지 용어들에 대해 해석 가능한 의미의 예를 제시한다. 하지만, 아래 제시하는 해석 예로 한정되는 것은 아님을 주의하여야 한다.

기지국(Base Station)은 단말과 통신하는 일 주체로서, BS, NodeB(NB), eNodB(eNB), AP(Access Point) 등으로 지칭될 수도 있다.

단말(User Equipment)은 기지국과 통신하는 일 주체로서, UE, 이동국(Mobile Station; MS), 이동장비(Mobile Equipment; ME), 디바이스(device), 터미널(terminal) 등으로 지칭될 수도 있다. 본 개시에서는 하나 이상의 연산 디바이스를 구비하는 장치를 단말 장치로 설명할 것이다.

본 개시는 어플리케이션의 소프트웨어(software; SW) 측면의 연산 특성에 따라 단말의 가용 자원(연산 디바이스들)을 최대한 활용하기 위해, 해당 어플리케이션의 연산을 연산 디바이스(들)에 동적으로 분배하여 어플리케이션의 동작 성능을 향상시키는 기법을 제안하고자 한다.

이를 위해 본 개시는, 소프트웨어를 각기 다른 연산 디바이스에서 수행할 수 있게 실행 코드를 생성 할 수 있도록 하는 컴파일러 부분과, 실제 실행 코드를 연산 특성 및 장치의 상태를 고려하여 동작 할 수 있도록 하는 런타임 SW(software) 부분을 설명한다.

디바이스 관점에서, 상기 컴파일러 부분은 예를 들어 컴파일러 모듈로 구현될 수 있고, 상기 런타임 SW 부분은 예를 들어 런타임 모듈로 구현될 수 있다. 상기 컴파일러 모듈과 런타임 모듈은 사용자 단말과 같은 하나의 디바이스에 포함될 수도 있고, 별도의 디바이스에 포함될 수도 있다. 예를 들어, 컴파일러 모듈은 단말과 같은 HW에서 구현될(탑재될) 수 있지만, 데스크탑 또는 서버와 같은 HW에 구현될 수도 있다. 또한, 상기 런타임 모듈은 어플리케이션을 구동하고자 하는 단말과 같은 HW에서 구현될(탑재될) 수 있다. 데스크탑 또는 서버와 같은 디바이스의 컴파일 모듈에서 컴파일된 코드는 런타임 모듈을 포함하는 단말에게 임포팅(importing)되고, 상기 컴파일된 코드는 상기 런타임 모듈의 제어 또는 스케줄링에 의해 실행될 수 있다. 이하에서는, 단말과 같은 하나의 디바이스에서 컴파일러 모듈과 런타임 모듈이 구현되는 경우를 예시할 것이나, 본 개시의 권리범위는 이에 한정되는 것은 아니고, 별개의 디바이스에서 각각 구현되는 경우도 포함한다.

도 1은 본 개시에 따른 어플리케이션 실행 디바이스의 구성을 예시하는 도면이다.

본 개시에 따른 어플리케이션 실행 디바이스는 두 가지 모듈 즉, 컴파일러 모듈(100) 및 런타임 모듈(150) 중 적어도 하나를 포함할 수 있다. 상기 컴파일러 모듈(100) 및 런타임 모듈(150)은 제어부(controller)와 같은 하나의 모듈로 구현될 수 있다.

상기 컴파일러 모듈은 예를 들어, LLVM(http://llvm.org/)과 같은 컴파일러 기반구조(infrastructure)를 가질 수 있다. 상기 컴파일러 모듈(100)은 C 또는 C++과 같은 언어로 작성된 어플리케이션 소스 코드(102) 또는 컴파일러 지시어를 입력받는다. 상기 C/C++ 컴파일러 지시어의 예로는, 멀티코어 CPU를 지시하는 "#pragma omp parallel", SIMD를 지시하는 "#pragma omp simd" 또는 GPU를 지시하는 "#pragma omp target" 등이 있을 수 있다.

본 개시의 컴파일러 모듈은 C 또는 C++과 같은 언어로 작성된 소스 코드(102)나 지시어를 입력받으므로, 어플리케이션 개발자는 특정 연산 디바이스만을 위한 언어나 API로 (즉, 개발자는 상기 연산 디바이스에 특정적인 언어나 API) 소스 코드를 작성할 필요가 없다.

상기 컴파일러 모듈(100)은 프론트엔드 모듈(104)을 포함할 수 있다. 상기 프론트엔드 모듈(104)은 상기 입력된 소스 코드(102) 또는 지시어를 상기 컴파일러 모듈(100)이 이해할 수 있는 코드 즉, 중간 코드(intermediate code)로 변경한다. 상기 중간 코드는 중간 표현(IR; intermediate representation)이라 호칭될 수 있다.

상기 컴파일러 모듈(100)는 분석부(analyzer)(106)를 포함할 수 있다. 상기 분석부(106)는 상기 소스 코드를 분석하여 코드의 정보를 수집하고, 상기 소스 코드로부터 별개의 연산 디바이스에서 실행될 수 있는 하나 이상의 코드 섹션(section)(107)을 생성할 수 있다. 별개의 연산 디바이스에서 실행될 수 있는 코드 섹션이란 예를 들어, 타 코드 섹션과 의존성(종속성)이 없는 코드 섹션이다. 상기 컴파일러 모듈(100)은 각 연산 디바이스용 실행 코드를 최적으로 생성하는데 있어서 상기 수집된 코드의 정보를 사용할 수 있다.

특히, 상기 분석부(106)는 C/C++ 소스 코드에서 사용되고 있는 포인터(pointer)를 이용하여 의존 관계를 분석할 수 있다. 상기 컴파일러 모듈(100)은 상기 포인터의 분석 결과를 이용하여 어플리케이션이 동작할 때 데이터 의존성을 분석할 수 있고, 상기 데이터의 의존성을 파악함으로써 최적화된 실행 코드를 생성할 수 있다. 여기서, 상기 포인터는 알리아싱(aliasing)이라고 불리기도 한다.

한편, 상기 포인터의 분석은 데이터 의존성을 분석하는 작업이므로 전체 컴파일 시간에 많은 영향을 준다. 상기 분석부(106)의 예시적 분석 과정이 설명된다.

상기 분석부(106)는 소스 코드를 분석하고 상기 소스 코드에서 사용되는 변수 및 포인터 등을 이용하여 N개의 노드(node)를 갖는 제약 그래프(Constraint Graph)를 작성할 수 있다. 이때 상기 제약 그래프가 작성되는 N개의 노드는 표 1에서 예시되는 제약 유형을 가질 수 있다.

제약 유형(Constraint type)	할당(Assignment)	제약(Constraint)	의미(Meaning)
AddressOf (Base)	a = &b	a ⊇{b}	loc(b) ∈ pts(a)
Copy (Simple)	a = b	a ⊇b	pts(a) ⊇ pts(b)
Load (Complex1)	a = *b	a ⊇*b	∀v ∈pts(b) pts(a) ⊇pts(v)
Store (Complex2)	*a = b	*a ⊇b	∀v ∈pts(a) pts(v) ⊇pts(b)

여기서, "a = &b"는 변수 b의 '주소(address)'를 a 의 값으로 할당하는 것을 의미하고, "a = b"는 변수 b의 값을 a의 값으로 할당하는 '카피(copy)'를 의미하고, "a = *b"는 a 의 값으로 포인터 b를 할당하는 '로드(load)'를 의미하고, "*a = b"는 포인터 a에게 b의 값을 할당하는 '저장(store)'을 의미한다.

상기 분석부(106)가 N개의 노드에 대한 의존 관계를 분석하는 데는 약 O(N³)에 해당하는 계산량이 요구된다. 계산량을 줄이기 위해, 본 개시에 따른 상기 분석부(106)는 포인터들 사이의 의존 관계를 분석할 때 오프라인(Off-line) 분석과 온라인(On-line) 분석으로 구분하여 분석할 수 있다. 오프라인 분석이란 소스 코드로부터 상기 제약 그래프를 생성하는 분석 과정이고, 온라인 분석은 상기 생성된 제약 그래프로부터 실제 연관성을 파악하는 분석 과정으로써 예를 들어 앤더슨 분석(Andersen's analysis) 기법이 적용될 수 있다.

구체적으로, 상기 분석부(106)는 오프라인 분석 시에 2 개의 노드들이 서로 '사이클(cycle) 관계'에 있는지 파악하여, '사이클 노드'에 해당하는 2 개의 노드를 하나의 노드로 취급하고 분석할 수 있다. 여기서, 사이클 노드란 2 개의 노드가 체인과 같은 형태로 서로 연결되어 있어서, 상기 2 개의 노드 이외의 노드와는 의존성이 없는 (즉, '사이클 관계'의) 노드이다.

또한, 상기 분석부(106)는 상기 사이클 노드 뿐만 아니라 다음의 예시적 조건이 만족하는 노드를 '의존성 없는 노드'로 판단하여 하나의 노드로 병합할 수 있다. 첫째 조건은, 상기 노드 a는 온라인 분석(online analysis) 시 선행자(predecessor)가 추가되지 않는 것이다. (예를 들어, LLVM IR의 탑 레벨 포인터 변수는 온라인 분석 시 선행자가 추가되지 않는 노드일 수 있다.) 둘째 조건은, 상기 노드 a는 오프라인 분석 시 생성된 제약 그래프(constraint graph)에서 단 하나의 선행자 노드 b를 갖는 것이다. 셋째 조건은, 상기 노드 a의 모든 포인터(point-to set)는 상기 선행자 노드 b로부터(만) 오는 것이다. 위 3가지 조건들을 만족하는 노드 a는 선행자(predecessor) 노드 b와 하나의 노드로 취급될 수 있다.

이와 같이 상기 분석부(106)는 코드 분석 시에 의존성을 이용하여 코드를 구성하는 노드들을 병합함으로써, 코드 분석에 요구되는 계산량을 약 O(N)으로 줄일 수 있다.

도 2는 본 개시에 따른 컴파일러 모듈이 하나의 노드로 취급하는 노드 들의 조건을 설명하는 제약 그래프, 소스 코드 및 IR을 예시한다.

상기 3 개의 조건을 만족하는 노드 a와 노드 b의 오프라인 제약 그래프(offline constraint graph)(200), C 언어로 작성된 코드(202) 및 IR(204)의 예제가 도 2에서 도시된다.

상기 컴파일러 모듈(100)은 특징 추출부(feature extractor)(108)를 포함할 수 있다. 상기 특징 추출부(108)는 상기 분석부(106)에서 분석한 결과를 이용하여 각각의 코드 섹션이 특정 연산 디바이스에서 동작 가능한지 판단하며, 동작 가능한 코드 섹션에 대해 상기 코드 섹션이 동작할 때 사용될 코드 정보를 추출한다. 상기 특징 추출부(108)가 추출하는 코드 정보의 예는 아래와 같다.

정보	내용
코드 섹션 내의 명령(instruction) 관련 정보	Integer ALU(arithmetic logic unit): 정수 덧셈 명령
	Float: 부동 소수점 수 덧셈 명령
	Multiply: 곱셈 명령
	Load/Store: 로드/저장 명령
	Branch: 분기 명령
	Vector Integer: 정수 어레이(array)
	Vector Float 등: 부동 소수점 수 어레이
캐쉬(cache)에 전송할 데이터 양	입력(input) 데이터 양

개발자에 의해 지정된 컴파일러의 지시어가 상기 소스 코드와 함께 입력되었을 경우라면, 상기 특징 추출부(108)는 상기 입력된 지시어 (즉, 가이드(guide))에 근거하여 특정 연산 디바이스에서 동작 가능한 코드 섹션을 판단할 수도 있다.

상기 컴파일러 모듈(100)은 변환부(transformation module)(110)를 포함할 수 있다. 상기 변환부(110)는 다수의 연산 디바이스에 의한 병렬 작업을 가능하게 하는 다수의 코드를 생성할 수 있으므로 병렬화부(parallelization module)로 호칭되기도 한다.

상기 변환부(110)는, 상기 분석부(106)에서 분석한 결과 또는 상기 특징 추출부(108)에서 추출한 정보를 이용하여, 상기 입력된 소스 섹션을 각각의 연산 디바이스에 적합한 소스 코드로 변환한다. 즉, 상기 변환부(110)는 상기 분석 결과 또는 추출한 정보를 이용하여 어떻게 코드를 변환할지를 결정하고, 각 연산 디바이스의 특성을 고려하여 소스 섹션을 변환할 수 있다. 도 1은, 연산 디바이스에 맞는 소스 코드의 예로써, 싱글 코어용 소스 코드(112; Serial), SIMD용 소스 코드(114; SI), 멀티 코어(classical)용 소스 코드(116; M(C)), 멀티코어(폴리히드럴)용 소스코드(118; M(P)), GPU용 소스 코드(120; CL) 및 멀티코어+SIMD용 소스 코드(122; M+SI)를 도시하고 있다.

구체적으로, 상기 변환부(110)는 상기 소스 섹션을 변환할 때, 상기 분석 결과(즉, 추출한 코드 정보)를 이용하여 루프 타일링(loop tiling), 언롤링(unrolling) 또는 인터리빙(inter-leaving)(즉, skewing) 과 같은 코드 변환 기법 중 적어도 하나를 적용할 수 있다. 루프 타일링이란, 루프 연산을 적어도 하나의 단위 블록(즉, 타일)으로 부분화(partitioning)하는 것을 말한다. 상기 변환부(110)은 상기 컴파일러의 캐쉬(cache) 크기를 고려하여 타일(tile)의 크기를 결정할 수 있다. 이렇게 함으로써, 캐쉬 크기에 따른 타일링의 효과를 미리 검증할 수 있고, 캐쉬 손실(cache miss)의 비율을 줄일 수 있다. 언롤링은 조건문(예를 들어, if 문)이 포함되는 루프 연산을 조건문이 없도록 일련의 명령어들로 풀어서 쓰는 것을 말한다. 언롤링은 연산 디바이스가 루프 연산을 수행하는 것보다 루프 연산을 포함하지 않지만 상대적으로 긴 명령어들을 수행하는 것이 더 효율적일 수 있기에 수행된다. 인터리빙은 루프 연산에 포함되는 인자(argument)의 위치를 변경하는 등으로 루프 연산의 연산 구조를 변경(또는 왜곡)하는 것을 말한다.

상기 변환부(110)에 의해 변환된 소스 코드는 '코드 특성'을 갖는 데이터 구조를 포함할 수 있다. 따라서, 상기 변환된 소스 코드는 '주석달린 코드(annotated code)'라고 호칭될 수도 있다. 런타임 모듈(150)이 어플리케이션의 동작 시 적절한 소스 코드를 사용할 수 있도록 하기 위해, 상기 컴파일러 모듈(100)은 상기 코드 특성에 해당하는 데이터 구조(structure)를 구성하고 상기 데이터 구조를 상기 소스 코드에 삽입하여 실제로 동작 시 정확하게 동작할 수 있게 한다. 이때, 상기 변환부(110)는 런타임 모듈(150)의 컴파일러 런타임 API를 사용함으로써, 모든 변환된 소스 코드가 연산 디바이스의 종류에 관계 없이 상기 런타임 모듈의 제어에 의해 동작할 수 있게 한다.

코드 특성의 데이터 구조의 예는 다음 표와 같다.

코드 특성	동시 사용여부
	수행시킬 바이너리
	쓰레드(thread) 개수
	각 바이너리 별 데이터 범위(data range)

상기 컴파일러 모듈(100)은 백엔드 모듈(130)을 포함할 수 있다. 상기 백엔드 모듈(130)은 각각의 연산 디바이스에 적합한 적어도 하나의 소스 코드(112, 114, 116, 118, 120, 122)를 컴파일하여 적어도 하나의 실행 코드(즉, 바이너리)(132, 134, 136, 138, 140)을 생성할 수 있다.

이어서, 단말의 런타임 SW 부분에 대해 설명한다.

런타임 SW 부분에 해당하는 런타임 모듈(150)은 적어도 하나의 실행 코드(즉, 바이너리)(132, 134, 136, 138, 140)을 상응하는 연산 디바이스에서 구동시킴으로써 어플리케이션을 실행시킬 수 있다.

상기 런타임 모듈(150)은 스케줄러(152)를 포함할 수 있다. 상기 스케줄러(152)는 상기 연산 디바이스의 실행 코드의 실행 전에 각 연산 디바이스의 상태(예를 들어, 사용률(utilization), 주파수(frequency) 등)를 확인하기 위해 장치 모니터(154)를 시스템 소프트웨어(system SW)의 서비스 형태로 이용하며, 상기 연산 디바이스의 상태 정보를 컴파일러 런타임 모듈(156)과 IPC(inter procedure communication)를 통해 통신할 수 있다. 상기 장치 모니터(154)를 통해 전달되는 연산 디바이스의 상태 정보의 예는 다음 표와 같다.

연산 디바이스의 상태 정보	연산 디바이스의 사용률
	연산 디바이스의 동작 주파수
	연산 디바이스의 On-Off 상태

상기 스케줄러(152)는 각각의 연산 디바이스를 효과적으로 활용하기 위해 정적(static) 정보와 동적(dynamic) 정보를 수집하고, 상기 수집된 정보에 근거하여 수행시킬 연산 자원(즉, 연산 디바이스)의 조합을 결정할 수 있다. 상기 스케줄러(152)는 상기 결정된 조합의 연산 디바이스(들)에 할당될 효율적인 워크로드(workload; 작업 분담률) 값을 결정하고, 상기 결정된 워크로드에 따라서 실행 코드를 상기 연산 디바이스(들)에서 구동시킨다. 상기 정적 정보와 동적 정보는 성격에 따라서 HW 정보와 소스 코드에 관련된 SW 정보로 나뉠 수 있다.

상기 스케줄러(152)는 소스 코드와 관련된 정적 정보를 컴파일러 모듈(100)의 특징 추출부(108)를 통해 얻을 수 있고, 환경(즉, HW)에 관련된 정적 정보는 미리 입력받을 수 있다.

동적 정보는, 코드가 수행될 때 정해지는 파라미터들로써, 예로써 상기 스케줄러(152)가 장치 모니터(154)를 통해 획득하는 각 연산 디바이스 별 사용률 정보 및 동작 주파수 정보 등이 될 수 있다.

정적 정보와 동적 정보의 예는 다음 표와 같다.

	정적(Static) 정보	동적(Dynamic) 정보
소스 코드 관련	. 루프 내의 총 명령(instruction) 수 - 스칼라 동작(Scalar Operations)(int/float 명령 카운트) - 벡터 동작(Vector Operations)(int4/float4 명령 카운트) - 메모리 로드, 저장 . 캐쉬에 전송할 데이터 양 . 사용자(개발자) 선호도	. 루프의 반복 횟수 . 처리할 데이터 사이즈 . 데이터의 위치 (L1, L2, L3) : CPU 근접도 크기에 따라 L1>L2>L3
환경 관련	. CPU 코어 개수 및 연산 능력 . SIMD 개수 및 연산 능력 . GPGPU의 프로세서, dimension, 메모리 크기 . DSP의 연산 능력 . Accelerator의 연산 능력 . 메모리 대역폭(Memory bandwidth)	. CPU 코어별 사용률, 주파수 정보 . GPGPU 사용률, 주파수 정보 . DSP 사용률, 주파수 정보 . Accelerator 의 사용률, 주파수 정보

명령 별 처리 시간은 연산 디바이스별로 다르기 때문에, 상기 스케줄러(152)는 각 연산 디바이스별 연산 특성에 대한 정보를 추출하고 상기 추출된 정보를 수행할 연산 디바이스 결정에 이용할 수 있다. 각 연산 디바이스별 연산 특성에 대한 정보는 HW 사양(spec.) 또는 각 동작 별 벤치마크(benchmark)를 이용하여 다음의 표와 같은 형태로 구성될 수 있다.

연산	CPU	GPGPU	DSP	Accelerator1
Integer ALU	CPU_int	GPGPU_int	DSP_int	ACC_int
Float	CPU_float	GPGPU_float	DSP_float	ACC_float
Multiply, MAC	CPU_Mul	GPGPU_Mul	DSP_Mul	ACC_Mul
Load/Store	CPU_mem	GPGPU_mem	DSP_mem	ACC_mem
Branch	CPU_Br	GPGPU_Br	DSP_Br	ACC_Br
Vector integer	CPU_vint	GPGPU_vint	DSP_vint	ACC_vint
Vector Float	CPU_vfloat	GPGPU_vfloat	DSP_vfloat	ACC_vfloat
...	...	...	...	...
합계	CPU_Total	GPGPU_Total	DSP_total	ACC_Total

상기 스케줄러(152)는, i) 사용자 선호도 또는 ii) 성능 예측에 따라 결정되는 순위(즉, priority) 에 근거하여, 어플리케이션의 동작을 위해 워크로드를 부과할 N개의 연산 디바이스를 선택할 수 있다. (예를 들어, N= 2) 사용자 선호도를 반영할 경우 상기 스케줄러(152)는 상기 사용자에 의해 선택된 연산 디바이스를 사용할 것으로 결정할 수 있고, 성능 예측 결과를 이용하는 경우 상기 스케줄러(152)는 연산 디바이스 별 성능(performance)을 예측하고 예측 결과에 따라 사용할 연산 디바이스를 결정할 수 있다. 예를 들어, 성능 예측에 따라서 결정되는 순위는 상기 표 6에 의해 결정되는 연산 디바이스별 총 연산 시간 (즉, CPU_Total, GPGPU_Total, DSP_total, ACC_Total)을 오름차순으로 정렬한 순서로 결정될 수 있다. 이때, 총 연산 시간이 가장 작은 연산 디바이스의 우선 순위가 가장 높다.

상기 표 6의 연산 디바이스 별 성능 예측에 사용되는 수식을 설명한다.

상기 연산 시간은 정수(integer) 연산 시간과 벡터(vector) 연산 시간을 포함할 수 있다. 상기 메모리 지연 시간은 메모리 읽기(read)/쓰기(write) 지연 시간과 벡터 로드(load)/저장(store) 지연 시간을 포함할 수 있다. 상기 데이터 카피 오버헤드는 예를 들어 데이터 전송(transfer)에 소요되는 시간을 포함할 수 있다.

구체적으로, 특정 연산 디바이스(즉, 디바이스 N)에서 워크로드가 수행되는 시간은 다음 수학식과 같이 정리될 수 있다.

여기서 T_total(device _N)은 디바이스 N의 총 연산 시간을 나타내고, T_op는 연산(명령어) 처리 시간을, T_vec _{_op}는 벡터 연산 처리 시간을, T_mem _{_} _lat는 메모리 지연 시간을, T_vec _{_ld_st_} _lat는 벡터 로드/저장 지연 시간을, T_br은 분기 연산 처리 시간을, T_{data_tr}은 데이터 전송 처리 시간을 나타낸다.

상기 수학식 2에 연산 자원 별 특성을 반영하면 이하의 수학식들과 같이 표현될 수 있다.

수학식 3은 연산 디바이스가 싱글(single) 코어 CPU일 경우의 총 처리시간을 나타낸다. 싱글 코어 CPU의 경우 연산에만 시간이 소요되고 데이터 카피로 인한 오버헤드 시간(T_data _{_tr})은 거의 없으므로 0으로 간주할 수 있다.

수학식 4는 연산 디바이스가 멀티(multi) 코어 CPU일 경우의 총 처리시간을 나타낸다. 여기서, Thread_count는 CPU의 쓰레드 개수를 나타낸다. 멀티 코어 CPU는 2 개 이상의 쓰레드를 생성하여 생성된 쓰레드들을 통해 병렬로 연산 처리하기 때문에 연산 시간 (및 분기 처리 시간)은 쓰레드의 개수에 비례하여 줄어든다. 하지만, 멀티 코어 CPU에서 상기 쓰레드들의 메모리 접근은 동시에 일어나므로 메모리 지연 시간이 증가하지는 않는다.

수학식 5는 연산 디바이스가 SIMD일 경우의 총 처리시간을 나타낸다. SIMD의 경우 벡터의 크기와 데이터 타입 크기에 따라 성능이 다를 수 있고, 데이터 카피 시간은 거의 소요되지 않는다.

수학식 6은 연산 디바이스가 GPGPU 또는 DSP와 같은 엑셀러레이터(Accelerator; 가속기)인 경우의 총 처리 시간을 나타낸다. GPU 또는 DSP와 같은 엑셀러레이터의 경우 데이터 카피 시간이 존재한다. 그러나, CPU와 GPU가 같이 사용할 수 있는 공유 메모리(shared memory)가 지원되는 경우라면, 데이터 카피 시간(T_{data_tr})은 0이 될 수도 있다.

상기 스케줄러(152)는 위와 같은 방식으로 다른 연산 디바이스에 대해서도 처리 시간을 계산하여, N개의 연산 디바이스에 대한 T_total(device₁), T_total(device₂), T_total(device₃), ... , T_total(device_N) 값들을 차례대로 구하고, 상기 구해진 처리시간들을 성능(즉, 처리 시간 크기) 순서로 오름차순 정렬하여 우선 순위를 매길 수 있다. (예를 들어, T_total(device₁) < T_total(device₃) < T_total(device₂))

상기 스케줄러(152)는 사용자의 입력 또는 사전 설정을 고려하여 상기 결정된 우선 순위에 해당하는 상위 N개의 디바이스를 선택하고, 상기 선택된 N개의 디바이스에 대한 부하 분산(load distribution) 최적화 과정에 적용할 수 있다. 바람직하게는, N은 2의 값을 가질 수 있다.

상기 선택된 N개의 연산 디바이스를 실시간에 효율적으로 운용하기 위해서는 정확한 성능의 예측이 필요하며, 상기 스케줄러(152)는 실시간 가용 정보(즉, 동적 정보)를 성능 예측에 이용할 수 있다. 상기 수학식 1 내지 수학식 6의 성능(즉, 처리 시간)은 해당 연산 디바이스를 100% 사용 가능했을 때를 기준으로 산정한 것이다. 그러나, 해당 연산 디바이스가 항상 100% 사용 가능하지는 않다. 따라서, 상기 스케줄러(152)는 각 연산 디바이스의 사용률(utilization) 정보 또는 주파수(frequency) 정보와 같은 동적 정보를 반영하여 보다 정확하게 성능을 예측할 수 있다.

주어진 워크로드(workload)를 분할하여 연산 디바이스들에 할당하는 경우, 워크로드가 할당된 연산 디바이스들 중 가장 오래 걸리는 연산 디바이스에 의해 어플리케이션의 수행 속도가 결정될 것이다. 가장 오래 걸리는 연산 디바이스의 처리 시간을 이용하여 결정되는 어플리케이션 처리 시간 T_workload 은 다음 수학식과 같이 표현될 수 있다.

여기서, α 는 부하 분산에 의해 결정되는 연산 디바이스의 부하 분담율을 나타내며 α₁ + α₂ + α₃ + ... = 1 이다. β는 정규화된(normalized) 사용률을 나타낸다. β 는, 연산 디바이스가 아이들(idle) 상태일 때 0의 값을 갖고, 상기 연산 디바이스가 풀(full)로 사용중일 때 1의 값을 갖는 것으로 가정된다. (0<= β <1)

따라서, 상기 스케줄러(152)는 연산 디바이스 별 실시간 가용 정보(즉, 사용률)를 고려하여 상기 수학식 7을 최소화하는 분담율 α를 결정하는 것으로써, 연산 디바이스별 부하 분산 최적화를 달성할 수 있다. 상기 최적의 분담률 α는 다음의 수학식에 의해 결정될 수 있다.

상기 런타임 모듈(150)은 컴파일러 런타임 모듈(156)을 포함할 수 있다. 상기 컴파일러 런타임 모듈(156)은 컴파일러 런타임 라이브러리를 제공한다. 상기 컴파일러 런타임 라이브러리는 컴파일러 모듈(100)와 밀접하게 연동하여 상기 스케줄러가 상기 어플리케이션의 실행 코드(즉, 바이너리)(132, 134, 136, 138, 140)를 구동할 수 있게 한다.

선택적으로, 상기 런타임 모듈(150)은 OS(operation system) 레벨의 추가적인 API 모듈(158) 또는 GPU 구동을 위한 OpenCL 드라이버 모듈(160)를 더 포함할 수도 있다.

도 3은 본 개시에 따른 런타임 모듈의 동작을 개략적으로 설명하는 도면이다.

런타임 모듈(특히, 스케줄러)은 컴파일러의 분석부로부터 실행 코드에 상응하는 정적 정보를 입력받을 수 있고, 장치 모니터로부터 동적 정보를 입력받을 수 있다(300).

상기 런타임 모듈은 어플리케이션을 수행할 연산 디바이스를 선택하는 동작을 수행할 수 있다(302). 구체적으로, 상기 런타임 모듈은 연산 디바이스 별로 성능 (예를 들어, 처리 시간)을 계산할 수 있다. 상기 런타임 모듈은 상기 계산된 연산 디바이스 별 성능을 오름차순으로 정렬할 수 있다. 상기 런타임 모듈은 사용자의 선택을 고려하여 우선 순위가 높은 N 개의 연산 디바이스를 수행 디바이스로 선택할 수도 있다.

상기 런타임 모듈은 선택된 연산 디바이스의 워크로드를 결정하는 스케줄링 동작을 수행할 수 있다(304). 예를 들어, 상기 런타임 모듈은 사용률 또는 주파수 정보를 이용하여 어플리케이션 처리 시간 T_workload를 결정할 수 있다. 상기 런타임 모듈은 상기 T_workload를 최소화하는 워크로드 α₁, α₂, α₃ , ... 를 결정할 수 있다.

상기 런타임 모듈은 상기 결정된 워크로드 α를 각 연산 디바이스에 적용하여, 어플리케이션 실행 코드를 구동한다(606).

도 4는 본 개시에 따라서 동작할 연산 디바이스를 결정하고 워크로드를 결정한 경우의 효과를 자원 사용률 및 사용 에너지 관점에서 예시하는 도면이다.

도 4(a) 및 도 4(b)는 런타임 모듈의 선택과 워크로드 결정 동작이 적용되지 않는 경우의 결과를 나타내고, 도 4(c) 및 도 4(d)는 런타임 모듈의 선택과 워크로드 결정 동작이 적용된 경우의 결과를 나타내고 있다.

도 4(a) 및 도 4(b)에서 알 수 있는 바와 같이, 런타임 모듈의 적용이 없는 경우에는, 측정 시구간 전체 동안 GPU(400) 및 멀티코어 CPU(402)의 사용률이 일정 수준에서 지속되고, 이로 인해 에너지 소비(404)도 측정 시구간 전체 동안에 일정하게 유지됨을 알 수 있다.

반면, 본 개시의 런타임 모듈의 적용이 있는 경우에는, 도 4(c) 에서 보여지듯이, GPU(410) 및 멀티코어 CPU(412)의 사용률이 측정 시구간 종료 전에 0의 값에 근접하는 것을 알 수 있다. 즉, 도 4(c) 에서는 연산 디바이스의 동작 시간이 도 4(a) 에 비해 절반 수준으로 감소함을 알 수 있다. 이는 동작 속도의 2배 향상을 의미한다. 또한 도 4(d) 에서의 에너지 소비(412)는 도 4(b)의 에너지 소비에 비해 빠른 속도로 감소함을 알 수 있다.

따라서, 본 개시의 따른 런타임 모듈의 적용은 어플리케이션의 동작 속도 향상과 소비 에너지(소비 전력)을 감소시키는 효과가 있다.

도 5는 본 개시에 따른 기법을 적용함에 있어서, 캐쉬의 입력 크기에 따른 연산 디바이스 별 처리 시간을 비교하는 표이다.

도 5를 참고하면, 캐쉬의 입력(input) 크기가 30K 인 경우에는 Serial 모드(즉, 싱글 코어 CPU)의 계산 시간(500)이 가장 작음을 알 수 있고, 상기 입력 크기라 3000K인 경우에는 Parallel-SIMD 모드 (즉, SIMD 및 CPU)의 계산 시간(502)이 가장 작음을 알 수 있다. 따라서, 캐쉬 입력 크기에 따른 성능 정보를 감안하여, 런타임 모듈은 연산 디바이스를 선택할 수도 있을 것이다.

도 6은 본 개시에 따른 기법을 적용함에 있어서, CPU 및 GPU를 이용하는 경우의 부하 분담율에 따른 성능(처리 시간)을 예시하는 도면이다.

도 6을 참고하면, CPU와 GPU는 약 63:37의 비율로 부하를 분담시킬 경우에 최소 처리 시간을 나타냄을 알 수 있다.

도 7은 본 개시에 따른 기법을 적용함에 있어서, CPU 및 GPU를 이용하는 경우의 워크로드에 따른 에너지 소비량을 예시하는 도면이다.

도 7을 참고하면, CPU와 GPU는 약 65:35의 비율로 부하를 분담시킬 경우에 최소의 에너지를 소비함을 알 수 있다.

도 8은 본 개시에 따른 기법을 적용함에 있어서, CPU 및 GPU를 이용하는 경우에 워크로드와 사용률에 성능(처리시간)을 예시하는 도면이다.

도 8에서, 사용률 β 가 0에 가까울수록 그래프들은 아래쪽으로 위치한다. 도 8을 참고하면, 사용률과 부하 분담율에 따라서 다양한 처리시간 결과를 나타냄을 알 수 있다.

도 9는 본 개시에 따른 컴파일러 모듈의 동작을 개략적으로 설명하는 도면이다.

도 9에서 설명되는 동작은, 도 3에서 설명된 런타임 모듈의 동작 이전에 수행될 수 있는 동작로써, 상기 런타임 모듈과 동일한 장치(예를 들어, 단말 장치) 또는 별개의 장치(예를 들어, 서버, 데스크탑 등)에 구현되는 컴파일러 모듈에서 수행될 수 있다.

상기 컴파일러 모듈은 어플리케이션 소스 코드를 컴파일러가 이해할 수 있는 중간 코드 (즉, IR)로 변경한다(900).

상기 컴파일러 모듈은 상기 IR에 포함되는 의존 관계를 분석하고, 상기 IR로부터 연산 디바이스 각각에서 실행될 수 있는 코드 섹션을 생성한다(905).

상기 컴파일러 모듈은 상기 코드 섹션을 연산 디바이스에 특정적인(적합한) 소스 코드로 변환한다(910). 이때, 상기 컴파일러 모듈은 상기 코드 섹션으로부터 코드 정보를 추출하는 동작을 더 수행할 수 있다. 또한, 상기 변환된 소스 코드에는 런타임 모듈이 실행 코드를 수행하는 데에 이용할 코드 특성이 데이터 구조의 형태로 포함될 수 있다.

상기 컴파일러 모듈은 상기 변환된 소스 코드를 컴파일하여 각 연산 디바이스에 적합한 실행 코드로 생성할 수 있다(915).

상기 컴파일러 모듈은 상기 생성된 실행 코드를 런타임 모듈에게 제공함으로써, 상기 런타임 모듈이 각각의 연산 디바이스에서 상기 실행 코드를 실행할 수 있게 한다.

표 7은 본 개시에 따른 컴파일러 분석부의 분석 성능을 나타낸다.

	기존		본 개시
프로그램	분석 시간 (sec.)	총 노드	분석 시간 (sec.)	병합후 노드
AccumulateWeighted In OpenCV	40.1	79,001	4.7	34679 (44322 노드가 병합됨)

표 7를 참고하면, 본 개시에 따른 분석 기법을 적용할 경우, 분석 시간이 기존 분석 방법의 40.1 초에서 4.7 초로 감소되고, 분석 노드의 개수는 기존 분석 방법의 79,001 개에서 36,479 개로 감소됨을 알 수 있다. 표 7에서는, 테스트를 위해 어큐뮬레이트 웨이티드 알고리즘이 이용되었다.

상기 도 1 내지 도 9가 예시하는 장치 구성도, 제어 방법의 예시도, 성능 예시도는 본 개시의 권리범위를 한정하기 위한 의도가 없음을 유의하여야 한다. 즉, 상기 도 1 내지 도 9에 기재된 모든 구성부, 또는 동작의 단계가 본 개시의 실시를 위한 필수구성요소인 것으로 해석되어서는 안되며, 일부 구성요소 만을 포함하여도 본 개시의 본질을 해치지 않는 범위 내에서 구현될 수 있다.

앞서 설명한 동작들은 해당 프로그램 코드를 저장한 메모리 장치를 통신 시스템의 엔터티, 기능(Function), 기지국, 또는 단말 장치 내의 임의의 구성부에 구비함으로써 실현될 수 있다. 즉, 엔터티, 기능(Function), 기지국, 또는 단말 장치의 제어부는 메모리 장치 내에 저장된 프로그램 코드를 프로세서 혹은 CPU에 의해 읽어내어 실행함으로써 앞서 설명한 동작들을 실행할 수 있다.

본 명세서에서 설명되는 엔터티, 기능(Function), 기지국, 또는 단말 장치의 다양한 구성부들과, 모듈(module)등은 하드웨어(hardware) 회로, 일 예로 상보성 금속 산화막 반도체(complementary metal oxide semiconductor) 기반 논리 회로와, 펌웨어(firmware)와, 소프트웨어(software) 및/혹은 하드웨어와 펌웨어 및/혹은 머신 판독 가능 매체에 삽입된 소프트웨어의 조합과 같은 하드웨어 회로를 사용하여 동작될 수도 있다. 일 예로, 다양한 전기 구조 및 방법들은 트랜지스터(transistor)들과, 논리 게이트(logic gate)들과, 주문형 반도체와 같은 전기 회로들을 사용하여 실시될 수 있다.

한편 본 개시의 상세한 설명에서는 구체적인 실시 예에 관해 설명하였으나, 본 개시의 범위에서 벗어나지 않는 한도 내에서 여러 가지 변형이 가능함은 물론이다. 그러므로 본 개시의 범위는 설명된 실시 예에 국한되어 정해져서는 안되며 후술하는 특허청구의 범위뿐만 아니라 이 특허청구의 범위와 균등한 것들에 의해 정해져야 한다.

Claims

단말 장치에서 적어도 하나의 연산 디바이스를 이용하여 어플리케이션을 수행하는 방법에 있어서,
상기 적어도 하나의 연산 디바이스의 예상 처리 시간을 계산하는 동작;
상기 계산된 적어도 하나의 연산 디바이스의 예상 처리 시간에 근거하여 상기 어플리케이션을 수행할 소정 개수의 연산 디바이스를 선택하는 동작;
상기 선택된 소정 개수의 연산 디바이스에 대해 상응하는 사용률 정보를 이용하여 결정된 처리시간 함수를 최소화하는 워크로드를 결정하는 동작; 및
상기 선택된 소정 개수의 연산 디바이스에 상기 결정된 워크로드를 적용하여 상기 어플리케이션을 실행하는 동작을 포함하는 방법.
제1항에 있어서,
상기 어플리케이션의 소스 코드를 IR(intermediate representation)로 변경하는 동작;
상기 IR을 분석하고 상기 IR로부터 별개의 연산 디바이스에서 실행될 수 있는 적어도 하나의 코드 섹션을 생성하는 동작;
상기 적어도 하나의 코드 섹션을 연산 디바이스에 특정적인 소스 코드로 변환하는 동작; 및
상기 변환된 소스 코드를 컴파일하여 상기 연산 디바이스에 특정적인 실행 코드를 생성하는 동작을 더 포함하는 방법.
제2항에 있어서,
상기 IR을 분석하고 상기 IR로부터 별개의 연산 디바이스에서 실행될 수 있는 적어도 하나의 코드 섹션을 생성하는 동작은:
상기 IR에 포함되는 의존 관계를 분석하는 동작; 및
상기 IR 중 상기 의존 관계가 없는 노드들 포함하는 코드 섹션을 상기 별개의 연산 디바이스에서 실행될 수 있는 상기 코드 섹션으로 생성하는 동작을 포함함을 특징으로 하는 방법.
제3항에 있어서,
상기 의존 관계를 분석함에 있어서, 사이클 관계에 있는 적어도 2개의 노드는 하나의 노드로 취급됨을 특징으로 하는 방법.
제2항에 있어서,
상기 생성된 코드 섹션으로부터 코드 정보를 추출하는 동작을 더 포함하되,
상기 코드 정보는 상기 코드 섹션 내의 명령에 관련된 정보 및 캐쉬 전송 데이터 양의 정보 중 적어도 하나를 포함함을 특징으로 하는 방법.
제2항에 있어서,
상기 변환된 소스 코드는 코드 특성에 해당하는 데이터 구조를 포함하고,
상기 코드 특성은 동시 사용여부, 수행될 실행코드, 쓰레드 개수, 및 실행코드 별 데이터 범위 정보 중 적어도 하나를 포함함을 특징으로 하는 방법.
제1항에 있어서,
상기 적어도 하나의 연산 디바이스는 싱글 코어 CPU(central processing unit), SIMD(single instruction multiple data), Parallel-SIMD, GPU(graphic processing unit), GPGPU(general purpose GPU), 및 DSP(digital signal processor) 중 적어도 2개를 포함함을 특징으로 하는 방법.
제2항에 있어서, 상기 어플리케이션의 소스 코드는 C 또는 C++ 언어로 작성되는 하나의 소스 코드임을 특징으로 하는 방법.
제1항에 있어서,
사용자 선호도에 기초하여 상기 소정 개수의 연산 디바이스를 선택하는 동작을 더 포함하고, 상기 사용자 선호도는 컴파일러 지시어에 의해 입력됨을 특징으로 하는 방법.
제3항에 있어서,
상기 의존 관계를 분석함에 있어서, 온라인 분석 시 선행자가 추가되지 않는 제1 노드, 및 오프라인 분석 시 상기 제1 노드가 갖는 하나의 선행자 노드인 제2 노드가 하나의 노드로 병합되고,
상기 제1 노드의 모든 포인터는 상기 제2 노드로부터만 오는 것을 특징으로 하는 방법.
적어도 하나의 연산 디바이스를 이용하여 어플리케이션을 수행하도록 구성된 런타임 모듈을 구비하는 단말 장치에 있어서, 상기 런타임 모듈은:
상기 적어도 하나의 연산 디바이스의 예상 처리 시간을 계산하고, 상기 계산된 적어도 하나의 연산 디바이스의 예상 처리 시간에 근거하여 상기 어플리케이션을 수행할 소정 개수의 연산 디바이스를 선택하고, 상기 선택된 소정 개수의 연산 디바이스에 대해 상응하는 사용률 정보를 이용하여 결정된 처리시간 함수를 최소화하는 워크로드를 결정하고, 상기 선택된 소정 개수의 연산 디바이스에 상기 결정된 워크로드를 적용하여 상기 어플리케이션을 수행하는 스케줄러; 및
상기 사용률 정보를 산출하는 장치 모니터를 포함하는 단말 장치.
제11항에 있어서,
상기 어플리케이션의 소스 코드를 IR(intermediate representation)로 변경하고, 상기 IR을 분석하여 상기 IR로부터 별개의 연산 디바이스에서 실행될 수 있는 적어도 하나의 코드 섹션을 생성하고, 상기 적어도 하나의 코드 섹션을 연산 디바이스에 특정적인 소스 코드로 변환하고, 상기 변환된 소스 코드를 컴파일하여 상기 연산 디바이스에 특정적인 실행 코드를 생성하는 컴파일러 모듈을 더 포함하는 단말 장치.
제12항에 있어서,
상기 컴파일러 모듈은, 상기 IR에 포함되는 의존 관계를 분석하여, 상기 IR 중 상기 의존 관계가 없는 노드들 포함하는 코드 섹션을 상기 별개의 연산 디바이스에서 실행될 수 있는 상기 코드 섹션으로 생성함을 특징으로 하는 단말 장치.
제13항에 있어서,
상기 컴파일러 모듈은, 상기 의존 관계를 분석함에 있어서, 사이클 관계에 있는 적어도 2개의 노드는 하나의 노드로 취급함을 특징으로 하는 단말 장치.
제12항에 있어서,
상기 컴파일러 모듈은, 상기 생성된 코드 섹션으로부터 코드 정보를 추출하는 동작을 더 수행하되,
상기 코드 정보는 상기 코드 섹션 내의 명령에 관련된 정보 및 캐쉬 전송 데이터 양의 정보 중 적어도 하나를 포함함을 특징으로 하는 단말 장치.
제12항에 있어서,
상기 변환된 소스 코드는 코드 특성에 해당하는 데이터 구조를 포함하고,
상기 코드 특성은 동시 사용여부, 수행될 실행코드, 쓰레드 개수, 및 실행코드 별 데이터 범위 정보 중 적어도 하나를 포함함을 특징으로 하는 단말 장치.
제11항에 있어서,
상기 적어도 하나의 연산 디바이스는 싱글 코어 CPU(central processing unit), SIMD(single instruction multiple data), Parallel-SIMD, GPU(graphic processing unit), GPGPU(general purpose GPU), 및 DSP(digital signal processor) 중 적어도 2개를 포함함을 특징으로 하는 단말 장치.
제12항에 있어서, 상기 어플리케이션의 소스 코드는 C 또는 C++ 언어로 작성되는 하나의 소스 코드임을 특징으로 하는 단말 장치.
제11항에 있어서,
상기 스케줄러는 컴파일러 지시어에 의해 입력된 사용자 선호도에 기초하여 상기 소정 개수의 연산 디바이스를 선택함을 특징으로 하는 단말 장치.
제13항에 있어서,
상기 컴파일러 모듈은, 온라인 분석 시 선행자가 추가되지 않는 제1 노드, 및 오프라인 분석시 상기 제1 노드가 갖는 하나의 선행자 노드인 제2 노드를 하나의 노드로 취급하여 상기 의존 관계를 분석하고,
상기 제1 노드의 모든 포인터는 상기 제2 노드로부터만 오는 것을 특징으로 하는 단말 장치.