KR20140029480A

KR20140029480A - 이종 코어를 위한 자동 커널 마이그레이션

Info

Publication number: KR20140029480A
Application number: KR1020137032393A
Authority: KR
Inventors: 마우리시오 브레터니츠; 패트릭 카민스키; 케이스 로웨리; 디즈-칭 주
Original assignee: 어드밴스드 마이크로 디바이시즈, 인코포레이티드
Priority date: 2011-05-16
Filing date: 2012-05-16
Publication date: 2014-03-10
Anticipated expiration: 2032-05-16
Also published as: JP5711853B2; CN103534686A; EP2710467B1; JP2014513853A; KR101559090B1; US20120297163A1; WO2012158753A1; CN103534686B; EP2710467A1; US8683468B2

Abstract

본 발명은 다수의 이종 코어들 사이에 작업 유닛의 실행을 자동으로 마이그레이션하는 시스템 및 방법에 관한 것이다. 컴퓨팅 시스템은 단일 명령 다중 데이터 마이크로-아키텍처를 구비하는 제1 프로세서 코어 및 일반 목적 마이크로-아키텍처를 구비하는 제2 프로세서 코어를 포함한다. 컴파일러는 프로그램에서 함수 호출의 실행을 예측하고, 소정의 위치에서 상이한 프로세서 코어로 마이그레이션한다. 상기 컴파일러는 상기 소정의 위치에서 함수의 호출의 실행과 연관된 라이브 값의 이동을 지원하는 데이터 구조를 생성한다. 운영 시스템(OS) 스케줄러는 상기 프로그램 순서에서 적어도 소정의 위치 전의 코드를 상기 제1 프로세서 코어로 스케줄링한다. 마이그레이션 조건이 만족된 것을 나타내는 지시를 수신한 것에 응답하여, 상기 OS 스케줄러는 상기 라이브 값을 상기 제2 프로세서 코어에 의해 액세스하기 위한 데이터 구조에 의해 지시된 위치로 이동시키고, 상기 소정의 위치 후의 코드를 상기 제2 프로세서 코어로 스케줄링한다.

Description

이종 코어를 위한 자동 커널 마이그레이션{AUTOMATIC KERNEL MIGRATION FOR HETEROGENEOUS CORES}

본 발명은 컴퓨팅 시스템(computing system)에 관한 것으로, 보다 상세하게는, 다수의 이종 코어(heterogeneous core)들 사이에 작업 유닛(work unit)의 실행을 자동으로 마이그레이션하는(migrating) 것에 관한 것이다.

태스크(task)의 병렬화(parallelization)는 컴퓨터 시스템의 처리량을 증가시키는데 사용된다. 이를 위해, 컴파일러(compiler)는 프로그램 코드로부터 병렬화된 태스크를 추출하여 시스템 하드웨어에서 병렬로 실행할 수 있다. 단일-코어 아키텍처(single-core architecture)에서, 단일 코어는 멀티-스레드(multi-threading)를 수행하도록 구성된 깊은 파이프라인(deep pipeline)을 포함할 수 있다. 하드웨어에서 병렬 실행을 더 증가시키기 위해, 멀티-코어 아키텍처(multi-core architecture)는 다수의 일반 목적 코어(general-purpose core)를 포함할 수 있다. 이러한 유형의 아키텍처는 동종 멀티-코어 아키텍처(homogeneous multi-core architecture)라고 지칭될 수 있다. 이러한 유형의 아키텍처는 단일-코어 아키텍처보다 더 많은 명령 처리량을 제공할 수 있다.

일부 소프트웨어 애플리케이션은 병렬 태스크로 빈번히 분할되지 않을 수 있다. 게다가, 특정 태스크는 일반 목적 코어에서 효율적으로 실행되지 않을 수 있다. 컴퓨테이션 집약적인 태스크(computational intensive task)를 위한 특정 명령은 공유된 자원에 불균형 공유를 야기하여, 공유된 자원의 할당 해제(deallocation)를 지연시킬 수 있다. 이러한 특정 태스크의 예로는 암호화, 비디오 그래픽 렌더링(rendering) 및 가비지 수집(garbage collection)을 포함할 수 있다.

종래의 일반 목적 코어의 성능 제한을 극복하기 위하여, 컴퓨터 시스템은 특정 태스크를 특수 목적 하드웨어에 오프로드(offload)할 수 있다. 이 하드웨어는 단일 명령 다중 데이터(single instruction multiple data: SIMD) 병렬 아키텍처, 전계-프로그래밍가능한 게이트 어레이(field-programmable gate array: FPGA), 및 다른 전문화된 코어를 포함할 수 있다. 상이한 유형의 코어를 구비하는 아키텍처의 유형은 이종 멀티-코어 아키텍처(heterogeneous multi-core architecture)라고 지칭될 수 있다. 태스크의 스케줄링에 따라, 이러한 유형의 아키텍처는 동종 멀티-코어 아키텍처보다 더 많은 명령 처리량을 제공할 수 있다.

많은 경우에, 특정 소프트웨어 애플리케이션은 각 작업 항목의 실행 또는 병렬 함수 호출(function call)이 자기 자신 내 데이터에 종속하는 데이터 병렬화를 구비한다. 예를 들어, 제1 작업 항목은 제2 작업 항목과 데이터가 독립적일 수 있고, 제1 및 제2 작업 항목 각각은 SIMD 마이크로-아키텍처(micro-architecture)를 구비하는 코어 내에서 별도의 경로로 스케줄링된다. 그러나, 제1 및 제2 작업 항목 각각에서 실행되는 명령의 양은 데이터 종속적일 수 있다. 브랜치 명령(branch instruction)으로 구현되는 조건부 테스트에서 제1 작업 항목은 통과(pass)할 수 있으나, 각 작업 항목에 대한 데이터에 종속하는 제2 작업 항목은 실패할 수 있다.

병렬 실행의 효율은 제1 작업 항목이 실행을 계속 진행하고 있을 때 제2 작업 항목은 실행을 정지하고 대기하여서 감소될 수 있다. 이 비효율은 테스트 통과로 인해 일부 작업 항목만이 실행을 계속하고 대부분의 작업 항목은 테스트 실패로 인해 아이들 상태(idle)일 때 증가한다. 이종 멀티-코어 아키텍처에서 OS 스케줄러에 의해 작업 항목을 효율적으로 기능 매칭하여 할당한 후에, 시스템 성능은 특정 소프트웨어 애플리케이션이 데이터에 종속하여 거동하는 것으로 인해 더 감소될 수 있다.

다수의 이종 코어들 사이에 작업 유닛의 실행을 자동으로 마이그레이션하는 시스템 및 방법이 고려된다.

일 실시예에서, 컴퓨팅 시스템은 제1 마이크로-아키텍처를 구비하는 제1 프로세서 코어 및 제1 마이크로-아키텍처와는 상이한 제2 마이크로-아키텍처를 구비하는 제2 프로세서 코어를 포함한다. 일 실시예에서, 제1 마이크로-아키텍처는 단일 명령 다중 데이터(single instruction multiple data: SIMD) 마이크로-아키텍처이고 제2 마이크로-아키텍처는 일반 목적 마이크로-아키텍처이다. 컴퓨팅 시스템은 제1 및 제2 프로세서 코어 각각에 연결된 메모리를 포함한다. 메모리는 하나 이상의 컴퓨트 커널(compute kernel) 또는 함수 호출을 포함하는 컴퓨터 프로그램을 저장한다. 컴파일러가 소정의 함수 호출의 명령을 횡단(traverse)할 때, 컴파일러는 함수 호출의 실행이 소정의 위치에서 상이한 프로세서 코어로 마이그레이션되는 것을 예측(predict)하도록 구성된다. 컴파일러는 이 소정의 위치에서 함수 호출의 실행과 연관된 라이브 값(moving live value)의 이동을 지원하는 데이터 구조(data structure)를 생성한다. 이러한 라이브 값은 "컨텍스트(context)"라고 지칭될 수 있다.

운영 시스템(OS) 내 스케줄러는 프로그램 순서에서 적어도 소정의 위치 전의 코드를 제1 프로세서 코어로 스케줄링한다. 마이그레이션의 조건이 만족된다는 지시를 수신한 것에 응답하여, OS 스케줄러는 라이브 값을 제2 프로세서 코어에 의해 액세스하기 위한 데이터 구조에 의해 지시된 위치로 이동시키고, 프로그램 순서에서 소정의 위치 후의 코드를 제2 프로세서 코어로 스케줄링한다. 마이그레이션 조건이 만족되었는지의 여부를 결정하기 위하여, 각 제1 및 제2 프로세서 코어는 출구 점(exit point)에 도달한 함수 호출의 병렬 실행 반복의 수가 소정의 임계값을 초과하였는지의 여부를 결정하도록 구성된다.

이들 및 다른 실시예는 이하 상세한 설명 및 도면을 참조하여 더 이해될 것이다.

도 1은 이종 멀티-코어 아키텍처를 구비하는 예시적인 프로세싱 노드의 일 실시예의 일반화된 블록도;
도 2는 컴퓨트 커널을 이용하는 소스 코드의 일 실시예의 일반화된 블록도;
도 3은 조건부 선언문(conditional statement)을 구비하는 컴퓨트 커널을 한정하는 소스 코드의 일 실시예의 일반화된 블록도;
도 4는 하드웨어 자원 및 컴퓨트 커널 사이에 스케줄링된 할당의 일 실시예의 일반화된 블록도;
도 5는 2개의 유형의 프로세서 코어에 대한 마이크로-아키텍처의 논리적 레이아웃의 일 실시예의 일반화된 블록도;
도 6은 일반 목적 파이프라인 실행 흐름의 일 실시예의 일반화된 블록도;
도 7a는 SIMD 파이프라인 실행 흐름의 일 실시예의 일반화된 블록도;
도 7b는 SIMD 파이프라인 실행 흐름의 일 실시예의 또 다른 일반화된 블록도;
도 8은 마이그레이션 태그된 브랜치를 구비하는 프로그램 코드의 일 실시예의 일반화된 블록도;
도 9는 컴퓨트 커널 마이그레이션을 위한 코드를 설치(instrument)하는 방법의 일 실시예를 도시한 일반화된 흐름도;
도 10은 프로그램 실행 동안 컴퓨트 커널을 마이그레이션하는 방법의 일 실시예를 도시한 일반화된 흐름도.
본 발명은 여러 변형 및 대안적인 형상을 취할 수 있으나, 특정 실시예만이 도면에 예로서 도시되고 본 명세서에 상세히 설명된다. 그러나, 도면 및 이에 대한 상세한 설명은 개시된 특정 형태로 본 발명을 제한하려는 것이 전혀 아니고, 오히려 본 발명은 첨부된 청구범위에 의해 한정된 본 발명의 사상과 범위 내에 있는 모든 변형, 균등물 및 대안을 포함하려는 것으로 이해된다.

이하 상세한 설명에서, 다수의 특정 상세들이 본 발명을 보다 잘 이해할 수 있도록 하기 위하여 제시된다. 그러나, 이 기술 분야에 통상의 지식을 가진 자라면 본 발명은 이들 특정 상세 없이 실시될 수 있다는 것을 이해할 수 있을 것이다. 일부 경우에, 잘 알려진 회로, 구조, 및 기술은 본 발명을 불명료하게 하는 것을 피하기 위해 상세히 도시되지 않았다.

도 1을 참조하면, 이종 멀티-코어 아키텍처를 구비하는 예시적인 프로세싱 노드(110)의 일 실시예가 도시된다. 프로세싱 노드(110)는 하나 이상의 프로세서 코어(112) 및 연관된 캐시 메모리 서브시스템(114)을 포함할 수 있는 하나 이상의 프로세싱 유닛(115)을 포함할 수 있다. 일 실시예에서, 프로세서 코어(112)는 일반 목적 마이크로-아키텍처를 이용한다.

프로세싱 노드(110)는 또한 하나 이상의 프로세서 코어(172) 및 데이터 저장 버퍼(174)를 포함할 수 있는 하나 이상의 프로세싱 유닛(170)을 포함할 수 있다. 프로세서 코어(172)는 프로세서 코어(112)의 미러된 실리콘 이미지(mirrored silicon image)가 아닐 수 있다. 프로세서 코어(172)는 프로세서 코어(112)에 의해 사용된 마이크로-아키텍처와는 상이한 마이크로-아키텍처를 구비할 수 있다. 일 실시예에서, 프로세서 코어(172)는 프로세서 코어(112)와 동일한 프로세서 군(family)의 상이한 생성물일 수 있다. 또 다른 실시예에서, 프로세서 코어(172)는 프로세서 코어(112)의 전압 및/또는 주파수 스케일링된 버전일 수 있다. 다시 말해, 프로세서 코어(172)는 동일한 기능 및 명령 세트 아키텍처(instruction set architecture: ISA), 동일한 클록 주파수, 동일한 캐시 사이즈, 동일한 메모리 모델 등을 구비하는 프로세서 코어(112)의 실리콘 복사물(silicon copy)이 아니다.

프로세서 코어(172)의 마이크로-아키텍처로 계속하면, 더 다른 실시예에서, 프로세서 코어(172)는 컴퓨테이션 집약적인 태스크(computational intensive task)를 위한 높은 명령 처리량을 제공하는 마이크로-아키텍처를 포함할 수 있다. 프로세서 코어(172)는 병렬 아키텍처를 구비할 수 있다. 예를 들어, 프로세서 코어(172)는 단일 명령 다중 데이터(SIMD) 코어일 수 있다. SIMD 코어의 예로는 그래픽 프로세싱 유닛(GPU), 디지털 신호 프로세싱(DSP) 코어, 또는 다른 것을 포함한다. 일 실시예에서, 프로세싱 노드(110)는 단일 명령 세트 아키텍처(ISA)를 포함한다. 일반적으로, 이 기술 분야에 잘 알려진 바와 같이, 단일-ISA 멀티-코어 아키텍처는 칩 멀티프로세서(CMP)를 위해 더 높은 전력 및 처리량 성능을 제공하는 것으로 제시되었다.

프로세싱 노드(110)에서 높은 명령 처리량은 소프트웨어 애플리케이션의 스레드(thread)가 효율적으로 스케줄링될 때 소정의 전력 제한 내에서 측정된 전력 소비량으로 달성될 수 있다. 스레드는 각 스레드가 프로세서 코어(112 및 172)의 실행시간 하드웨어 자원에 적어도 부분적으로 기초하여 최고 명령 처리량을 구비하는 방식으로 프로세서 코어(112 및 172) 중 하나로 스케줄링될 수 있다.

프로세싱 노드(110) 내 성분으로 계속하면, 프로세싱 노드(110)는 메모리 제어기(120) 및 인터페이스 논리(140)를 포함할 수 있다. 일 실시예에서, 프로세싱 노드(110)의 예시된 기능은 단일 집적 회로에 병합된다. 일 실시예에서, 프로세서 코어(112)는 미리 한정된 일반 목적 명령 세트에 따라 명령을 실행하는 회로를 포함한다. 예를 들어, SPARC(등록상표) 명령 세트 아키텍처(ISA)가 선택될 수 있다. 대안적으로, x86, x86-64, Alpha(등록상표), PowerPC(등록상표), MIPS(등록상표), PA-RISC(등록상표), 또는 임의의 다른 명령 세트 아키텍처가 선택될 수 있다. 일반적으로, 프로세서 코어(112)는 데이터 및 명령을 위해 캐시 메모리 서브시스템(114)에 각각 액세스한다. 요청된 블록이 캐시 메모리 서브시스템(114)에서 또는 공유된 캐시 메모리 서브시스템(118)에서 찾을 수 없으면, 판독 요청이 생성되고 이는 누락된 블록(missing block)으로 맵핑된 노드 내 메모리 제어기로 송신될 수 있다.

일 실시예에서, 프로세싱 유닛(170)은 그래픽 프로세싱 유닛(GPU)이다. 현대 GPU는 컴퓨터 그래픽을 조작하고 디스플레이하는데 매우 효율적이다. GPU의 높은 병렬 구조는 복잡한 알고리즘의 범위에 대해 프로세싱 유닛(115)과 같은 일반 목적 중앙 프로세싱 유닛(CPU)보다 그 처리를 더 효과적이게 한다. 일반적으로, GPU는 그래픽 및 비디오에 사용되는 계산을 실행하며 CPU는 그래픽만 있는 것보다도 많은 더 많은 시스템 프로세스에 대한 계산을 실행한다. 종래의 GPU는 매우 넓은 단일 명령 다중 데이터(SIMD) 아키텍처를 이용하여 이미지-렌더링 애플리케이션에서 높은 처리량을 달성한다. 이러한 애플리케이션은 일반적으로 매우 많은 수의 객체(vertices 또는 픽셀)에 대해 정점 셰이더(vertex shader) 또는 픽셀 셰이더(pixel shader)와 같은 동일한 프로그램을 실행하는 것을 수반한다. 각 객체는 다른 객체와 독립적으로 프로세싱되지만, 동일한 시퀀스의 동작이 사용되므로, SIMD 마이크로-아키텍처는 상당한 성능 개선을 제공한다. GPU는 또한 비그래픽적인 계산을 고려한다.

일 실시예에서, GPU(170)는 비디오 카드에 위치될 수 있다. 또 다른 실시예에서, GPU(170)는 마더보드에 집적될 수 있다. 더 다른 실시예에서, 프로세싱 노드(110)의 도시된 기능은 단일 집적 회로에 병합될 수 있다. 이러한 실시예에서, CPU(115) 및 GPU(170)는 상이한 설계 센터로부터의 전용 코어일 수 있다. 또한, GPU(170)는 이제 인터페이스(140)를 통해 오프-칩으로 메모리 액세스를 수행하는 것이 아니라 프로세싱 노드(110)로부터 메모리 제어기(120)를 통해 로컬 메모리(114 및 118) 및 메인 메모리에 직접 액세스할 수 있다. 이 실시예는 더 높은 성능으로 변환될 수 있는 GPU(170)에 대한 메모리 액세스의 레이턴시를 저하시킬 수 있다.

도 1에 있는 프로세싱 노드(110)의 성분으로 계속하면, 캐시 서브시스템(114 및 118)은 데이터의 블록을 저장하도록 구성된 고속 캐시 메모리를 포함할 수 있다. 캐시 메모리 서브시스템(114)은 각 프로세서 코어(112) 내에 집적될 수 있다. 대안적으로, 캐시 메모리 서브시스템(114)은 원하는대로 백사이드 캐시 구성 또는 인라인(inline) 구성으로 프로세서 코어(114)에 연결될 수 있다. 더 나아가, 캐시 메모리 서브시스템(114)은 캐시 계층(cache hierarchy)으로 구현될 수 있다. (계층 내) 프로세서 코어(112)에 더 가까이 위치된 캐시는 원하는 경우 프로세서 코어(112)에 집적될 수 있다. 일 실시예에서, 캐시 메모리 서브시스템(114)은 L2 캐시 구조를 나타내고, 공유된 캐시 서브시스템(118)은 L3 캐시 구조를 각각 나타낸다. 캐시 메모리 서브시스템(114) 및 공유된 캐시 메모리 서브시스템(118)은 대응하는 캐시 제어기에 연결된 캐시 메모리를 포함할 수 있다.

일반적으로, 패킷 프로세싱 논리(116)는 프로세싱 노드(110)에 연결된 링크에서 수신된 제어 패킷에 응답하고, 프로세서 코어(112) 및/또는 캐시 메모리 서브시스템(114)에 응답하여 제어 패킷을 생성하고, 서비스할 메모리 제어기(120)에 의해 선택된 트랜잭션에 응답하여 프로브 커맨드(probe command) 및 응답 패킷을 생성하여, 중간 노드인 노드(110)를 위한 패킷을 인터페이스 논리(140)를 통해 다른 노드로 라우팅하도록 구성된다. 인터페이스 논리(140)는 패킷을 수신하고 패킷을 패킷 프로세싱 논리(116)에 의해 사용되는 내부 클록과 동기화하는 논리를 포함할 수 있다.

이제 도 2를 참조하면, 컴퓨트 커널을 이용하는 소스 코드의 일 실시예가 도시된다. OpenCL(상표명)(Open Computing Language)은 이종 컴퓨팅을 위한 저레벨 애플리케이션 프로그래밍 인터페이스(API)의 일 예이다. OpenCL은 실행 큐를 한정하고, 각 큐는 OpenCL 디바이스와 연관된 것인 C-유형 언어를 포함한다. OpenCL 디바이스는 CPU, GPU, 또는 이종 멀티-코어 아키텍처 내에 적어도 하나의 프로세서 코어를 구비하는 다른 유닛일 수 있다. 함수 호출은 OpenCL 커널, 또는 간단히 "컴퓨트 커널"이라고 지칭될 수 있다. OpenCL 프레임워크는 게임, 엔터테인먼트, 과학 및 의료 분야에서 사용되는 광범위한 데이터-병렬 애플리케이션에 대한 컴퓨팅 성능을 개선시킬 수 있다. 이종 아키텍처에서, 컴퓨터 프로그램은 일반적으로 컴퓨트 커널 및 내부 함수의 집합체를 포함한다. 소프트웨어 프로그래머는 컴퓨트 커널을 한정할 수 있는 반면, 내부 함수는 소정의 라이브러리에서 한정될 수 있다.

데이터-병렬 소프트웨어 애플리케이션에서, N-차원 컴퓨테이션 도메인은 "실행 도메인"의 조직(organization)을 한정할 수 있다. N-차원 컴퓨테이션 도메인은 또한 N-차원 그리드(grid) 또는 N-차원 범위("NDRange")라고 지칭될 수 있다. NDRange는 1차원-, 2차원-, 또는 3차원 공간일 수 있다. 일부 실시예는 3차원을 초과하는 데이터를 허용할 수 있는 것으로 이해된다. 이러한 차원 공간은 또한 인덱스 공간(index space)이라고 지칭될 수 있다. 예를 들어, 소프트웨어 애플리케이션은 이미지 파일과 같은 2차원(2D) 데이터 어레이에 데이터 프로세싱을 수행할 수 있다. 소프트웨어 애플리케이션은 2D 이미지의 픽셀마다 또는 2차원 매트릭스의 요소마다 소프트웨어 프로그래머에 의해 개발된 알고리즘을 수행할 수 있다. 소정의 컴퓨트 커널은 인덱스 공간(NDRange) 위에 호출될 수 있다. 다른 실시예에서, 소프트웨어 애플리케이션은 매크로분자 모델링(macromolecular modelling)에 사용되는 3D 격자(lattice) 및 직접 쿨롱 합산(direct coulomb summation)으로 맵핑하는 정전 전위(electrostatic potential)를 위한 데이터-병렬 프로그래밍을 이용하는 알고리즘을 포함할 수 있다.

일반적으로 컴파일 후, 각 컴퓨트 커널의 변수(argument) 및 파라미터(parameter)들이 설정된다. 추가적으로, 연관된 메모리 객체 및 버퍼가 생성된다. 컴퓨트 커널의 소정의 인스턴스(instance)가 자기 자신의 소프트웨어 스레드로 실행될 수 있다. 그러나, 컴퓨트 커널은 포크(fork)를 생성하는 제어 흐름 전송 명령을 포함할 수 있는 반면, 컴퓨터 프로그램 내 포크는 일반적으로 공통 정의(common definition)에 의해 소프트웨어 스레드를 생성한다. 인덱스 공간에서 소정의 점에 있는 컴퓨트 커널의 소정의 인스턴스는 "작업 항목(work item)"이라고 지칭될 수 있다. 작업 항목은 또한 작업 유닛(work unit)이라고 지칭될 수 있다. 작업 유닛은 2D 이미지의 소정의 픽셀(소정의 인덱스)에 대응하는 데이터 레코드에 있는 컴퓨트 커널의 하나 이상의 명령으로 동작할 수 있다. 일반적으로, 작업 유닛은 연관된 고유한 식별자(ID)를 구비한다. 또 다른 예에서, 스트링 "Hello World"를 프로세싱하는 도입 컴퓨터 프로그램(introductory computer program)은 스트링에 있는 각 문자(letter)를 컴퓨팅하는 하나의 작업 유닛을 구비할 수 있다.

NDRange는 충분한 하드웨어 지원이 있다면 병렬로 실행하는 총 개수의 작업 유닛을 한정할 수 있다. 예를 들어, NDRange는 280개의 작업 유닛의 수를 한정할 수 있으나, GPU는 임의의 소정의 시간에 64개의 작업 유닛의 동시 실행을 지원할 수 있다. 작업 유닛의 총 수는 글로벌 작업 사이즈를 한정할 수 있다. 이 기술 분야에 통상의 지식을 가진 자에게 잘 알려진 바와 같이, 작업 유닛은 작업 그룹으로 더 그룹화될 수 있다. 각 작업 그룹은 고유한 식별자(ID)를 구비할 수 있다. 소정의 작업 그룹 내 작업 유닛은 서로 통신하여 실행 및 좌표 메모리 액세스(coordinate memory access)를 동기화할 수 있다. 다수의 작업 유닛은 SIMD 방식으로 GPU에서 동시 실행을 위해 파면(wave front)으로 클러스터링(clustered)될 수 있다. 280개의 총 작업 유닛에 대해 상기 예를 참조하면, 파면은 64개의 작업 유닛을 포함할 수 있다.

OpenCL 프레임워크는 여러 컴퓨트 디바이스, 또는 OpenCL 디바이스에 대한 개방 프로그래밍 표준이다. 소프트웨어 프로그래머는 벤더-특정 코드를 기록하는 것을 회피하여, 코드 휴대성(portability)을 개선시킬 수 있다. 다른 프레임워크들도 이용가능하고 이종 아키텍처에 대해 보다 벤더-특정적인 코딩을 제공할 수 있다. 예를 들어, NVIDIA는 컴퓨트 통일된 디바이스 아키텍처(Compute Unified Device Architecture: CUDA(등록상표))를 제공하고 AMD는 ATI Stream(등록상표)을 제공한다. CUDA 프레임워크에서, 컴퓨트 커널은 컴퓨터 프로그램이 컴파일될 때 일반적으로 정적으로 컴파일된다. OpenCL 프레임워크에서, 컴퓨트 커널은 일반적으로 JIT(Just-In-Time) 방법으로 컴파일된다. JIT 방법은 시스템 구성을 획득한 후 적절한 이진 코드를 생성할 수 있다. JIT 컴파일 방법에서, 컴파일 시간은 총 실행 시간에 포함된다. 그리하여, 컴파일러 최적화는 실행 시간을 증가시킬 수 있다. 게다가, 실행 시간에서 OpenCL 컴파일러는 컴퓨트 커널의 다수의 버전을 생성할 수 있다. 컴퓨트 커널의 하나의 버전은 일반 목적 CPU, SIMD GPU 등과 같은 OpenCL 디바이스 유형 각각에 대해 생성될 수 있다.

2개의 프레임워크, OpenCL 및 CUDA는 각 실행 모델들 사이에 용어(terminology)의 차이를 가지고 있다. 예를 들어, OpenCL에 있는 작업 유닛, 작업 그룹, 파면 및 NDRange는 CUDA에서 스레드, 스레드 블록, 워프(warp) 및 그리드와 같은 용어에 대응한다. 본 설명의 나머지에 걸쳐, OpenCL에 대응하는 용어들이 사용된다. 그러나, 설명된 시스템 및 방법은 CUDA, ATI Stream 및 다른 프레임워크에도 적용될 수 있다.

도 2에 도시된 바와 같이, 코드(210)는 일반적으로 "doWorkA" 및 "doWorkB"라고 명명된 2개의 함수 호출을 한정한다. 각 함수 호출은 "컴퓨트 커널"이라고 지칭될 수 있다. 컴퓨트 커널은 하나 이상의 데이터 레코드와 매칭되어 하나 이상의 컴퓨테이션 작업 유닛을 생성할 수 있다. 그리하여, 2개 이상의 작업 유닛은 단일 함수 호출의 동일한 명령을 이용하지만, 상이한 데이터 레코드에서 동작할 수 있다. 예를 들어, 코드(220)에 있는 함수 호출 "Power2"은 어레이 "입력"에 있는 각 데이터 값에 대해 하나씩 10개의 작업 유닛을 실행하는데 사용된다. 여기서, 레코드는 단일 데이터 값을 포함한다. 다른 예에서, 레코드는 2개 이상의 필드를 포함할 수 있고, 각 필드는 데이터 값을 포함한다. SIMD 마이크로-아키텍처는 커널 "Power2"의 명령을 효율적으로 실행하고, 입력 어레이에 있는 값에 대해 2의 멱수(power)를 계산하고 출력을 결과 어레이에 기록할 수 있다.

OpenCL 프레임워크는 컴퓨트 커널의 인스턴스를 다수회 병렬로 호출할 수 있다. 컴퓨트 커널의 각 호출은 get_global_id(0)라고 명명된 내부 함수를 호출하는 것에 의해 페치(fetch)될 수 있는 하나의 연관된 고유한 ID(작업 유닛 ID)를 구비한다. 코드(220)에 있는 상기 예에서, 컴퓨트 커널 "Power2"은 입력 어레이에 있는 각 데이터 값에 대해 한번 호출된다. 이 경우에, 컴퓨트 커널 "Power2"은 10회 호출된다. 따라서, 10개의 고유한 작업 유닛 ID이 페치(fetch)된다. JIT 컴파일 방법에서, 이들 인스턴스는 실행시간에 호출된다. OpenCL 프레임워크는 고유한 작업 유닛 ID을 이용하는 것에 의해 이들 상이한 인스턴스들 간을 구별할 수 있다. 입력 어레이에 있는 특정 데이터 값과 같은 (레코드 상)에서 동작될 데이터가 또한 특정될 수 있다. 그리하여, 실행시간에, 작업 유닛은 연관된 컴퓨트 커널이 스케줄링될 때 디폴트에 의해 동일한 OpenCL 디바이스로 스케줄링될 수 있다.

이제 도 3을 참조하면, 조건부 선언문을 구비하는 컴퓨트 커널을 한정하는 소스 코드의 일 실시예가 도시된다. 코드(210)와 유사하게, 도 3에 도시된 코드 230는 일반적으로 "doWorkA" 및 "doWorkB"라고 명명된 2개의 함수 호출을 한정한다. 다시, 각 함수 호출은 "컴퓨트 커널"이라고 지칭될 수 있다. 여기서, 2개의 컴퓨트 커널 중 단 하나의 것만이 실행시간 동안 실행된다. 어느 컴퓨트 커널이 실행될지 선택하는 것은 함수 호출 "EvaluateFunction"에 의해 제공된 조건부 테스트에 기초한다. 소정의 명령의 결과 또는 소정의 명령이 실행되었는지의 여부는 연관된 레코드에 대응하는 이전의 명령 및 데이터의 실행에 따라 데이터 종속적이다. 조건부 테스트의 결과가 작업 유닛의 파면 중에 일관성이 없다면, SIMD 마이크로-아키텍처의 이익은 감소될 수 있다. 예를 들어, 소정의 SIMD 코어는 64개의 작업 유닛을 동시 실행하는데 이용가능한 64개의 병렬 컴퓨테이션 유닛을 구비할 수 있다. 그러나, 64개의 작업 유닛 중 절반이 조건부 테스트를 통과하는 반면 다른 절반은 조건부 테스트를 통과하지 못한다면, 병렬 컴퓨테이션 유닛의 절반만이 소정의 프로세싱 스테이지 동안 이용된다.

이제 도 4를 참조하면, 하드웨어 자원 및 컴퓨트 커널 사이에 스케줄링된 할당(400)의 일 실시예를 도시하는 일반화된 블록도가 도시된다. 여기서, 하나 이상의 소프트웨어 애플리케이션(430)이 실행하는 동안 하드웨어 및 소프트웨어 자원의 파티셔닝(partitioning) 및 그 상호관계 및 할당이 도시된다. 일 실시예에서, 운영 시스템(420)은 컴퓨트 커널(440a 내지 440j 및 440k 내지 440q)을 위해 메모리 영역을 할당한다. 애플리케이션(430) 또는 컴퓨터 프로그램이 실행할 때, 각 애플리케이션은 다수의 컴퓨트 커널을 포함할 수 있다. 예를 들어, 제1 실행 애플리케이션은 컴퓨트 커널(440a 내지 440j)을 포함할 수 있고 제2 실행 애플리케이션은 컴퓨트 커널(440k 내지 440q)을 포함할 수 있다. 커널(440a 내지 440q) 각각은 하나 이상의 데이터 레코드(미도시)와 결합되는 것에 의해 하나 이상의 작업 유닛을 생성하는데 사용될 수 있다. 예를 들어, 컴퓨트 커널(440a)은 작업 유닛(442a 내지 442d)을 생성할 수 있고, 컴퓨트 커널(440j)은 작업 유닛(442e 내지 442h)을 생성할 수 있고, 컴퓨트 커널(440k)은 작업 유닛(442j 내지 442m)을 생성할 수 있고, 컴퓨트 커널(440q)은 작업 유닛(442n 내지 442q)을 생성할 수 있다. 작업 유닛은 다른 작업 유닛과 독립적으로 실행할 수 있고 다른 작업 유닛과 동시에 실행할 수 있다.

도 4에 도시된 컴퓨트 커널 각각은 애플리케이션 실행 전에 메모리의 이미지와 같은 자기 자신의 자원 또는 명령 및 데이터의 인스턴스를 소유할 수 있다. 각 컴퓨트 커널은 또한 코드, 데이터, 및 가능하게는 히프(heap) 및 스택(stack)을 어드레스하는 어드레스 공간; 스택 포인터(stack pointer), 일반 및 부동 소수점(general and floating-point register), 프로그램 카운터, 및 다른 것과 같은 데이터 및 제어 레지스터의 변수(variable); stdin, stdout, 및 다른 것과 같은 운영 시스템 디스크립터; 및 허가 세트(permission set)와 같은 보안 속성(security attribute)과 같은 프로세서 특정 정보를 포함할 수 있다.

일 실시예에서, 하드웨어 컴퓨팅 시스템(410)은 일반 목적 프로세서 코어(112) 및 SIMD 프로세서 코어(172)를 병합하고, 각각은 하나 이상의 작업 유닛을 프로세싱하도록 구성된다. 또 다른 실시예에서, 시스템(410)은 2개의 다른 이종 프로세서 코어를 포함한다. 일반적으로, 소정의 애플리케이션에 대해, 운영 시스템(420)은 애플리케이션을 위한 어드레스 공간을 설정하고, 애플리케이션 코드를 메모리에 로딩하고, 프로그램을 위한 스택을 설정하고, 애플리케이션 내 소정의 위치로 브랜치하고, 애플리케이션의 실행을 시작한다. 일반적으로, 이러한 활동을 관리하는 운영 시스템(420) 부분은 운영 시스템(OS) 커널(422)이다. OS 커널(422)은 컴퓨트 커널 또는 함수 호출과의 혼동을 피하기 위하여 "OS 커널"이라고 지칭된다. OS 커널(422)은 애플리케이션을 실행하는데 메모리가 불충분할 때 동작 과정을 더 결정할 수 있다. 전술된 바와 같이, 애플리케이션은 하나를 초과하는 컴퓨트 커널로 분할될 수 있고 시스템(410)은 하나를 초과하는 애플리케이션을 실행할 수 있다. 그리하여, 병렬로 실행하는 컴퓨트 커널이 여럿 있을 수 있다. OS 커널(422)은 임의의 시간에 동시에 실행하는 컴퓨트 커널 중에서 프로세서 코어(112 및 172)에 할당된 컴퓨트 커널을 결정할 수 있다. OS 커널(422)은 프로세스가 시간 슬라이스라고 지칭되는 소정의 시간 양 동안 하나 이상의 코어를 구비할 수 있는 프로세서의 코어에서 실행하게 할 수 있다. 운영 시스템(420) 내 OS 스케줄러(424)는 컴퓨트 커널을 코어에 할당하는 결정 논리를 포함할 수 있다.

일 실시예에서, 단 하나의 컴퓨트 커널이 임의의 시간에 하드웨어 컴퓨테이션 유닛(412a 내지 412g 및 412h 내지 412r) 중 어느 것에서 실행될 수 있다. 이들 하드웨어 컴퓨테이션 유닛은 연관된 데이터를 가지고 소정의 작업 유닛의 소정의 명령의 실행을 처리할 수 있는 하드웨어를 포함한다. 이 하드웨어는 가산(addition), 승산(multiplication), 제로 검출(zero detection), 비트방식 시프트(bit-wise shift), 제산(division), 비디오 그래픽 및 멀티미디어 명령 또는 프로세서 설계의 기술 분야에서 통상의 지식을 가진 자에 알려진 다른 동작을 수행하도록 구성된 산술 논리 유닛을 포함할 수 있다. 이들 하드웨어 컴퓨테이션 유닛은 멀티-스레드된 프로세서에서 하드웨어 스레드, SIMD 마이크로-아키텍처에서 병렬 하드웨어 열(column) 등을 포함할 수 있다.

도 4에서 대시 라인은 할당을 나타내고 직접 물리적 연결을 나타내는 것은 아니다. 따라서, 예를 들어, 하드웨어 컴퓨테이션 유닛(412a)은 작업 유닛(442d)을 실행하도록 할당될 수 있다. 그러나, 차후 (예를 들어, 컨텍스트 스위치 후), 하드웨어 컴퓨테이션 유닛(412a)은 작업 유닛(442h)을 실행하도록 할당될 수 있다. 일 실시예에서, OS 스케줄러(424)는 작업 유닛(442a 내지 442q)을 라운드 로빈 방식으로 하드웨어 컴퓨테이션 유닛(412a 내지 412r)으로 스케줄링할 수 있다. 대안적으로, OS 스케줄러(424)는 작업 유닛(442a 내지 442q)을 라운드 로빈 방식으로 코어(112 및 172)로 스케줄링할 수 있다. 소정의 작업 유닛을 소정의 하드웨어 컴퓨테이션 유닛으로 할당하는 것은 연관된 프로세서 코어에 의해 수행될 수 있다. 또 다른 실시예에서, OS 스케줄러(424)는 프로세서 코어(112 및 172)의 이용가능성에 기초하여 스케줄링을 수행할 수 있다. 더 다른 실시예에서, OS 스케줄러(424)는 OpenCL(상표명) API 또는 다른 유사한 API를 이용하는 프로그래머에 의해 생성된 할당에 따라 스케줄링을 수행할 수 있다. 이들 스케줄링 방식은 작업 유닛 할당 및 하드웨어 자원 사이에 미스매치가 있을 때 휴대성과 성능을 제한할 수 있다.

도 5를 참조하면, 2개의 유형의 프로세서 코어를 위한 마이크로-아키텍처의 논리적 레이아웃의 일 실시예를 도시하는 일반화된 블록도가 도시된다. 일반 목적 코어(510) 및 단일 명령 다중 데이터(SIMD) 코어(560)의 각각이 도시되지만, 다른 유형의 이종 코어들도 가능하고 고려된다. 각 코어(510 및 560)는 데이터 및 명령을 저장하기 위해 동적 랜덤 액세스 메모리(DRAM)(550a 및 550b)를 구비한다. 일 실시예에서, 코어(510 및 560)는 동일한 DRAM을 공유한다. 또 다른 실시예에서, 캐시 메모리 서브시스템(미도시)의 소정의 레벨이 DRAM에 더하여 공유된다. 예를 들어, 도 1을 다시 참조하면, 캐시 메모리 서브시스템(118)은 코어(112 및 172)에 의해 공유된다.

각 코어(510 및 560)는 캐시 메모리 서브시스템(530)을 포함할 수 있다. 도시된 바와 같이, 일반 목적 코어(510)는 논리적으로 제어 논리(520) 및 산술 논리 유닛(ALU)(540)과는 별개로 캐시 메모리 서브시스템(530)을 구비한다. 코어(510) 내 데이터 흐름은 파이프라이닝될 수 있으나, 파이프라인 레지스터와 같은 저장 요소는 도시를 간략화하기 위하여 도시되지 않았다. 소정의 파이프라인 스테이지에서, ALU는 이 스테이지 내 명령이 특정 유형의 ALU를 이용하지 않는 경우 또는 if 다른 작업 유닛(또는 일반 목적 코어를 위한 다른 스레드)이 이 스테이지 동안 ALU를 소비하는 경우에는 사용되지 않을 수 있다.

도시된 바와 같이, SIMD 코어(560)는 컴퓨테이션 유닛(542)의 각 행에 대해 제어 논리(520)로 그룹화된 캐시 메모리 서브시스템(530)을 구비한다. 코어(560) 내 데이터 흐름은 파이프라이닝될 수 있으나, 파이프라인 레지스터와 같은 저장 요소는 도시를 간략화하기 위하여 도시되지 않았다. 소정의 파이프라인 스테이지에서, 컴퓨테이션 유닛은 이 스테이지 내 연관된 명령이 미취득 브랜치(not-taken branch)와 같이 이전의 실패한 테스트에 기초하여 실행되지 않는 경우에는 사용되지 않을 수 있다.

이제 도 6을 참조하면, 일반 목적 파이프라인 실행 흐름(600)의 일 실시예를 도시하는 일반화된 블록도가 도시된다. 명령(602 내지 608)은 페치(fetch)될 수 있고 일반 목적 파이프라인에 들어갈 수 있다. 명령(606)은 컴퓨테이션 집약적인 명령일 수 있다. 파이프라인 실행 흐름의 특정 스테이지 동안, 명령(602 내지 608) 중 하나 이상은 디코더 논리, 명령 스케줄러 엔트리, 정렬(reorder) 버퍼 엔트리, ALU, 레지스터 파일 엔트리, 브랜치 예측 유닛 등과 같은 일반 목적 프로세서 코어(112)에서 자원을 소비한다.

균형된 방식으로 각 명령(602 내지 608)은 각 스테이지에서 동일한 양의 자원을 소비한다. 그러나, 일반적으로, 일반 목적 코어는 반도체 바닥 면적(real-estate) 비용, 전력 소비량 및 다른 설계 고려사항으로 인해 각 명령에 대한 자원을 복제하지는 않는다. 그리하여, 작업부하는 불균형으로 될 수 있다. 예를 들어, 명령(606)은 컴퓨테이션 집약적인 거동으로 인해 하나 이상의 파이프 스테이지에 대해 더 많은 자원을 소비할 수 있다. 도시된 바와 같이, 이 명령에 의해 소비되는 자원(630)은 다른 명령에 의해 소비되는 자원을 훨씬 더 초과할 수 있다. 사실, 컴퓨테이션 집약적인 명령은 다른 명령에 의해 하드웨어 자원을 사용하는 것을 차단할 수 있다.

일부 컴퓨테이션 집약적인 태스크는 도 1에 도시된 일반 목적 코어(112) 내 공유된 자원에 압력을 줄 수 있다. 따라서, 컴퓨테이션 집약적인 프로세스 및 공유된 자원을 대기하는 다른 프로세스에서 처리량 손실(loss)이 발생한다. 게다가, 일부 명령은 공유된 자원, 및 이 공유된 자원에 수행되는 컴퓨테이션을 지원하기 위해 다른 자원을 점유할 수 있다. 이러한 긴 레이턴시 명령은 동시에 다른 프로세스들이 긴 레이턴시 동안 여러 자원을 사용하는 것을 차단할 수 있다.

이제 도 7a를 참조하면, SIMD 파이프라인 실행 흐름(700)의 일 실시예를 도시하는 일반화된 블록도가 도시된다. 명령(702 내지 708)은 페치(fetch)될 수 있고 연관된 데이터를 구비하는 SIMD 파이프라인에 들어갈 수 있다. 명령(704)은 조건부 브랜치와 같은 제어 흐름 전송 명령일 수 있다. 명령(706)은 조건이 참(true)일 때 실행되는 경로에 있는 제1 명령일 수 있다. 명령(708)은 조건이 거짓(false)일 때 실행되는 경로에 있는 제1 명령일 수 있다. 예를 들어, 브랜치 명령(704)은 하이-레벨 언어 프로그램에 있는 IF 선언문과 연관될 수 있다. 명령(706)은 하이-레벨 언어 프로그램에 있는 THEN 선언문과 연관될 수 있다. 명령(708)은 하이-레벨 언어 프로그램에서 ELSE 선언문과 연관될 수 있다.

소정의 행 내에 있는 각 컴퓨테이션 유닛은 동일한 컴퓨테이션 유닛일 수 있다. 이들 컴퓨테이션 유닛 각각은 동일한 명령에 동작할 수 있으나, 상이한 작업 유닛에는 상이한 데이터가 연관된다. 도시된 바와 같이, 작업 유닛 중 일부는 조건부 브랜치 명령(704)에 의해 제공된 테스트를 통과하고 다른 작업 유닛은 테스트를 통과하지 못한다. SIMD 코어(172)는 이용가능한 경로들 각각을 실행하고 현재 경로를 선택하지 않은 작업 항목에 대응하는, 컴퓨테이션 유닛과 같은, 실행 유닛을 선택적으로 디스에이블할 수 있다. 예를 들어, If-Then-Else 구성(construct) 선언문을 실행하는 동안, SIMD 아키텍처의 각 열에는 "Then" 경로(경로 A) 및 "Else" 경로(경로 B)를 실행하도록 구성된 실행 유닛이 있다. 병렬 실행의 효율은 제3 작업 유닛이 실행 진행을 계속하고 있을 때 제1 및 제2 작업 유닛은 실행을 일시 정지(pause)하고 대기하는 것으로 인해 감소될 수 있다. 그리하여, 브랜치 명령(704)의 실행 후 소정의 행에 있는 컴퓨테이션 유닛들이 전부 활성 컴퓨테이션 유닛(710)인 것은 아니다. 도시된 바와 같이, 하나 이상의 컴퓨테이션 유닛은 실행이 디스에이블된 비활성 컴퓨테이션 유닛(711)이다. 컴퓨테이션 유닛의 대부분이 소정의 파이프 스테이지 동안 비활성이면, SIMD 코어의 효율 및 처리량이 감소된다.

일 실시예에서, "Else" 경로는 컴퓨트 커널을 리턴시키는 것이다. 컴퓨트 커널의 실행은 종료하고, 대응하는 작업 유닛은 아이들 상태가 된다. 그러나, SIMD 코어 내 이웃한 작업 유닛은 실행을 계속할 수 있다. 이제 도 7b를 참조하면, SIMD 파이프라인 실행 흐름(720)의 다른 실시예를 도시하는 일반화된 블록도가 도시된다. 실행 흐름(700)과 유사하게, 명령(702 내지 706)은 SIMD 코어의 특정 행에 있는 하나 이상의 컴퓨테이션 유닛이 디스에이블될 수 있게 한다. 여기서, 각 "Else" 경로는 컴퓨트 커널을 리턴하는 것일 수 있다. 그리하여, 소정의 작업 유닛에서, 미취득된 방향(not-taken direction)으로 해결(resolving)하는 브랜치는 소정의 작업 유닛이 컴퓨트 커널의 추가적인 실행을 중지하게 할 수 있다. 실행 흐름(720)에서, 단 하나의 명령만이 도시의 편의를 위하여 제1 브랜치 명령(704) 및 제2 브랜치 명령(712) 사이에 도시된다. 그러나, 다수의 명령이 브랜치 명령(704 및 712)들 사이에 있을 수 있다. 브랜치(704 및 712) 사이에 있는 명령의 수에 상관없이, 미취득된 방향으로 제1 브랜치(704)를 해결하는 작업 유닛은 실행을 완료할 수 있다. 브랜치(712)에 대한 것과 유사하게, 미취득된 방향으로 제2 브랜치를 해결하는 작업 유닛은 실행을 완료할 수 있다. SIMD 코어의 차후 스테이지를 위한 컴퓨테이션 유닛은 이들 작업 유닛에 대해 디스에이블될 수 있다. 컴퓨테이션 유닛의 대다수가 소정의 파이프 스테이지 동안 비활성이라면, SIMD 코어의 효율 및 처리량이 감소된다.

다수의 작업 유닛이 테스트를 통과하지 못하고 실행을 중지하게 하는 동안 이웃한 작업 유닛은 계속 실행할 수 있게 하는 동안 애플리케이션의 일 례로는 얼굴 검출(face detection)이 있다. 이 기술 분야에 통상의 지식을 가진 자에게 알려진 바와 같이, OpenCv(Open Computer Vision library)로 구현된 얼굴 검출이 Viola-Jones 객체 검출 알고리즘의 하나의 애플리케이션이다. Viola-Jones 알고리즘은 데이터에 종속하는 실행 패턴을 나타낸다. 검색 컴퓨트 커널은 하나 이상의 픽셀을 포함할 수 있는 데이터 레코드에 적용된다. 검색 컴퓨트 커널은 2차원(2D) 또는 3차원(3D) 이미지의 서브 윈도우에서 얼굴을 검색한다. 컴퓨트 커널 내에는, 브랜치 명령과 같은 제어 흐름 전송 명령으로 구현된 테스트의 캐스케이드(cascade)가 있을 수 있다. 하나의 대표적인 예에서, 테스트의 캐스케이드는 22개의 스테이지 또는 22개의 테스트를 포함한다. 이 테스트의 캐스케이드는 입력 윈도우가 얼굴을 포함하는지의 여부를 결정할 수 있다.

Viola-Jones 알고리즘에서 테스트의 캐스케이드는 유망하지 않는 경로(unpromising path)를 신속히 제거(prune)하도록 설계될 수 있다. 그리하여, 대부분의 작업 유닛은 얼굴의 비 존재(non-existence)를 결정하고 종료할 수 있다. 작업 유닛의 실행은 얼굴을 포함할 것 같은 나머지 픽셀에 대해 계속된다. 작은 부분(fraction)의 픽셀(즉, 작업 유닛 실행)은 22개의 스테이지를 통해 계속될 수 있는 반면, 대부분의 픽셀은 일부 초기 스테이지 테스트 후에 얼굴을 포함하지 않는 것으로 발견될 수 있다. 심지어 대량의 태스크 병렬화를 통해, 파면에 일부 계속하는 작업 유닛이 존재하면 SIMD 코어 이용율이 낮아질 수 있다. 아래에 설명된 하나의 방법은 별개의 이종 코어를 이용하며 추가적인 프로세싱을 위해 SIMD 코어를 해제한다. 이 방법은 소량의 SIMD 병렬화가 존재하는 것으로 검출될 때 전반적인 컴퓨팅 성능을 증가시킬 수 있다.

이제 도 8을 참조하면, 마이그레이션 점을 한정하도록 태그된 브랜치를 포함하는 코드(800)의 일 실시예가 도시된다. 코드(800)는 컴퓨트 커널 일반적으로 "foo"라고 명명된 컴퓨트 커널을 포함한다. 실행 동안, 코드(800)의 일부분은 별개의 이종 코어로 마이그레이션될 수 있다. 도시된 예에서, 외부 루프는 데이터 종속적이다. 일 실시예에서, 컴파일러는 "while" 루프 테스트에 대응하는 브랜치 명령에 태그 비트를 사용하는 것에 의해 데이터 종속성을 SIMD 코어에 알려준다. 실행 동안, 측정된 SIMD 이용율이 소정의 임계값 미만인 것과 같이 마이그레이션 조건이 검출되면, 중간 로컬 값은 별개의 이종 코어에 의해 액세스될 메모리 내 데이터 구조로 이동될 수 있다. 예를 들어, 일반 목적 코어는 태그된 브랜치 마이그레이션 점의 점으로부터 컴퓨트 커널의 실행을 계속할 수 있다. 예를 들어, while 선언문에서 암시적인 조건부 브랜치는 레이블 "secondary_entry"로 태그된다. 별개의 이종 코어는 컴파일러-생성된 데이터 구조를 사용할 수 있다. 또 다른 실시예에서, 이 데이터는 캐싱되어, 마이그레이션 비용을 경감시킬 수 있다. 일 예에서, 라이브 데이터는 "tmp" 어레이의 로컬 슬라이스, 및 local_temp 변수의 현재 값을 모두 포함할 수 있다. 마이그레이션 동안, 이 데이터는 실행시간 환경과 통신될 수 있고, 이는 레이블 "secondary_entry"에 의해 지시된 제2 진입 점으로 컴퓨트 커널의 실행을 계속 지시할 수 있다.

이제 도 9를 참조하면, 선-실행시간 데이터 정보를 이용하는 것에 의해 프로세서에서 다수의 작업 유닛의 병렬 실행을 최적화하는 방법 (900)의 일 실시예가 도시된다. 프로세싱 노드(110)에 구현된 성분 및 전술된 도 4에 도시된 하드웨어 자원 할당은 일반적으로 방법 (900)에 따라 동작할 수 있다. 설명의 목적을 위하여, 이 실시예 및 차후 설명된 방법의 후속하는 실시예에 있는 단계는 순차적으로 도시된다. 그러나, 다른 실시예에서 일부 단계는 도시된 것과는 상이한 순서로 발생할 수 있고, 일부 단계는 동시에 수행될 수 있고, 일부 단계는 다른 단계와 조합될 수 있고, 일부 단계는 부재(absent)할 수 있다.

블록(902)에서, 소프트웨어 프로그램 또는 서브루틴이 위치되고 분석될 수 있다. 이 소프트웨어 프로그램은 이종 멀티-코어 아키텍처에서 컴파일 및 실행하기 위해 기록될 수 있다. 프로그램 코드는 소프트웨어 애플리케이션, 서브루틴, 동적 링크된 라이브러리, 또는 다른 것의 임의의 부분을 말할 수 있다. 경로 이름은 유저에 의해 커맨드 프롬프트(command prompt)에서 입력될 수 있고, 경로이름은 소스 코드를 컴파일 하기 시작하기 위하여 소정의 디렉토리 위치, 또는 다른 곳으로부터 판독될 수 있다. 프로그램 코드는 C, C-유형 언어, 예를 들어 OpenCL(상표명) 등과 같은 하이-레벨 언어에서 설계자에 의해 기록될 수 있다. 일 실시예에서, 소스 코드는 정적으로 컴파일된다. 이러한 실시예에서, 정적 프론트엔드(front-end) 컴파일 동안, 소스 코드는 중간 표현(IR)으로 변환될 수 있다. 백엔드(back-end) 컴파일 단계는 IR을 기계 코드로 변환할 수 있다. 정적 백엔드 컴파일은 더 많은 변환 및 최적화를 수행할 수 있다. 또 다른 실시예에서, 소스 코드는 JIT(Just-In-Time) 방법으로 컴파일된다. JIT 방법은 시스템 구성을 획득한 후 적절한 이진 코드를 생성할 수 있다. 어느 방법이든, 컴파일러는 프로그램 코드에서 컴퓨트 커널을 식별할 수 있다. 일 실시예에서, 컴파일러, 예를 들어 OpenCL 컴파일러는 컴퓨트 커널의 다수의 버전을 생성할 수 있다. 컴퓨트 커널의 하나의 버전은 OpenCL 디바이스 유형, 예를 들어 일반 목적 CPU, SIMD GPU 등의 각 유형에 대해 생성될 수 있다.

블록(904)에서, 컴파일러는 컴퓨트 커널의 하나 이상의 명령을 판독하고 이를 분석할 수 있다. 조건부 선언문은 제어 흐름 전송 명령, 예를 들어 브랜치일 수 있다. 상이한 유형의 제어 흐름 전송 명령은 포워드/백워드 브랜치, 직접/간접 브랜치, 점프 등을 포함할 수 있다. 컴파일러 또는 다른 툴이 브랜치의 방향 및/또는 브랜치의 타깃을 정적으로 결정하는 것이 가능할 수 있다. 그러나, 일 실시예에서, 연관된 데이터에 대한 실행시간 동안 일반적으로 수행된 일부 프로세싱은 컴파일 동안 수행될 수 있다. 예를 들어, 브랜치의 방향(취득, 미취득)을 결정하는 간단한 테스트가 수행될 수 있다. 컴파일이 "정적 컴파일"이라고 지칭될 수 있으나, 하나 이상의 소량의 동적 동작이 수행될 수 있다. 이 컴파일은 또한 "선-실행시간 컴파일"이라고 지칭될 수 있다. 이 시간에 수행되는 동적 단계의 또 다른 예는 If-Then-ElseIf-Else 구성의 THEN, ELSE IF 및 ELSE 블록 각각에서 실행할 그 다음 명령을 식별하는 것이다. 예를 들어, 조건부 브랜치가 실패하면, 리턴 선언문이 실행될 수 있다. 그리하여, 컴파일러는, 실행 동안 브랜치 테스트가 실패할 때 이 컴퓨터 커널에 대한 대응하는 작업 유닛이 아이들 상태가 될 수 있다는 것을 안다.

블록(906)에서, 컴퓨트 커널에서 특정 코드 라인이 마이그레이션 점을 생성하기 위해 선택된다. 마이그레이션 점은 비행(in-flight) 실행이 상이한 이종 코어로 전이하는 컴퓨터 커널 내 위치일 수 있다. 일 실시예에서, 이 컴퓨트 서브 커널 마이그레이션은 마이그레이션을 프로세싱하는 것과 유사한 메커니즘에 의해 달성될 수 있고, 실행 상태는 제1 이종 코어로부터, 가능하게는 제1 코어와는 상이한 마이크로-아키텍처를 구비하는 제2 이종 코어로 이동된다. 또 다른 실시예에서, 이 컴퓨트 서브 커널 마이그레이션은 차후 디스패치(dispatch)된 다수의 컴퓨트 서브 커널을 생성하는 것에 의해 달성될 수 있다.

일 실시예에서, 컴파일러는 마이그레이션 점을 자동으로 식별할 수 있다. 본 명세서에 사용된 바와 같이, 마이그레이션 점은 또한 스위치 점이라고 지칭될 수 있다. 컴파일러는 제어 흐름 분석을 사용할 수 있다. 마이그레이션 점을 식별하는 것은 정적 제어 흐름 분석을 이용하여 컴퓨트 커널 출구 또는 리턴을 초래하는 데이터 종속 루프를 찾는 것을 포함할 수 있다. 출구 또는 리턴을 포함하는 경로를 구비하는 각 브랜치를 식별하는 대신, 컴파일러는 마이그레이션 점의 수를 감소시키기 위해 카운트를 사용할 수 있다. 예를 들어, 컴퓨트 커널에서 발견되는 5개의 제1 브랜치는 마이그레이션 점으로 태그하기 위한 후보가 아닐 수 있다. 5개의 제1 브랜치 후 매 제3번째 브랜치가 마이그레이션 점으로 태그하기 위한 후보일 수 있다. 카운트에 기초한 다른 필터링 알고리즘도 가능하고 고려된다.

게다가, 컴파일러는 이전의 실행으로부터 프로파일 입력을 사용하여 마이그레이션 점을 식별할 수 있다. 예를 들어, 소정의 브랜치와 연관된 조건부 테스트는 소정의 임계값을 초과하는 다수의 데이터 레코드 동안 실패할 수 있다. 그리하여, 이 브랜치는 마이그레이션 점으로 식별될 수 있다. 나아가, 마이그레이션 점을 나타내는 프로그래머 주석(annotation)이 "프래그마(pragma)"로 추가되거나 또는 OpenCL 프레임워크에의 확장자로 추가될 수 있다.

블록(908)에서, 컴파일러는 컴파일된 코드의 각 버전에 대한 코드에서 선택된 점을 태그할 수 있다. 각 버전은 각 OpenCL 디바이스에 대한 목적지 컴퓨트 커널이라고 지칭될 수 있다. 다시, 컴파일러는 식별된 컴퓨트 커널을 컴파일하여 OpenCL 디바이스 각각에서 각각 실행할 수 있는 컴파일된 코드의 2개 이상의 버전을 생성할 수 있다. 도 9에 있는 코드(800)을 다시 참조하면, 레이블 "secondary_entry"에 의해 지시된 제2 진입 점은 브랜치를 위한 마이그레이션 태그의 일 례이다. 컴파일러 내 코드 생성기는 태그를 삽입하고 다른 코드를 삽입하여 마이그레이션 동안 라이브 값을 호출할 수 있다. 라이브 값을 호출하는 것은 라이브 값을 목적지 OpenCL 디바이스에 전송하는 것 및 목적지 OpenCL 디바이스에 있는 값을 초기화하는 것을 포함할 수 있다. 코드 생성 및 삽입 프로세스는 동적 거동을 측정하는 데뷔(debut) 점 및 설치에 삽입된 디버거 코드와 유사할 수 있다.

일 실시예에서, 컴퓨트 커널은 전술된 바와 같이 마이그레이션 점을 식별하도록 태그될 수 있다. 또 다른 실시예에서, 퓨트 커널은 독립적으로 스케줄링 및 디스패치(dispatch)되는 다수의 컴퓨트 서브 커널로 분할될 수 있다. 실행시간 프로파일 정보 또는 컴파일러 정적 추정(estimation)은 브랜치 명령에 의해 구현되는 조건부 테스트에 대해 통과/실패의 통계를 결정하는데 사용될 수 있다. "핫(hot)" 실행 경로는 다수의 데이터 레코드에 대해 조건부 테스트의 소정의 임계값을 초과하는 다수의 통과를 포함할 수 있다. "콜드(cold)" 실행 경로는 다수의 데이터 레코드에 대해 조건부 테스트의 제2 소정의 임계값 미만으로 소수의 통과를 포함할 수 있다. 컴퓨트 커널은 "핫" 및 "콜드" 실행 경로에 기초하여 컴퓨트 서브 커널로 분할될 수 있다.

대응하는 컴퓨트 서브 커널의 생성은 일반 목적 코어에서 실행을 계속하는, "콜드" 실행 경로와 같은, 컴퓨트 서브 커널에 대한 대응하는 실행 범위(NDRange)를 생성하는 것에 더하여, 유사한 실행시간 코드 생성 메커니즘을 이용할 수 있다. 이것은, 일반 목적 코어에서 실행될, OpenCL 지정을 이용할 수 있는, 컴퓨트 서브 커널 식별자(ID)를 포함하는 잠재적으로 스파스 어레이(sparse array)를 생성하는 것에 의해 수행될 수 있다. 소정의 컴퓨트 커널은 이 어레이에 간접 액세스를 이용하여 적절한 컴퓨트 서브 커널 및 차후 작업 유닛을 식별할 수 있다. 대안적으로, 컴파일러는 이들 ID 리스트, 및 실행하는 작업 유닛 각각에 대해 호출되고 맵핑될 대응하는 컴퓨트 서브 커널을 생성할 수 있다.

프로파일이 정적 추정(estimation)을 실행한 후, "핫" 실행 경로에 대응하는 컴퓨트 서브 커널은 SIMD 코어에 대해 컴파일될 수 있다. "콜드" 실행 경로에 대응하는 컴퓨트 서브 커널은 일반 목적 코어에 대해 컴파일될 수 있다. 테스트의 캐스케이드의 초기 스테이지는 높은 통과 가능성을 가지고 있을 수 있다. 그리하여, 이들 실행 경로는 SIMD 코어에서 실행되는 "핫" 컴퓨트 서브 커널에서 구현될 수 있다. 이들 특정 "핫" 컴퓨트 서브 커널의 실행 후에, 연관된 생성된 데이터는 메모리에서 이동될 수 있다. 이 데이터 이동은 로컬 데이터를 글로벌 데이터에 라이브한 것으로 승격(promot)시킨다. "핫" 컴퓨트 서브 커널에 대응하는 작업 유닛은 작업 유닛 ID에 기초하여 비트 어레이를 기록하여 연관된 "콜드" 컴퓨트 서브 커널이 이후 일반 목적 코어에서 실행을 계속할지의 여부를 나타낼 수 있다.

블록(910)에서, 컴파일러는 식별된 마이그레이션 점에서 라이브 값 세트를 식별한다. 라이브 값은 중간 컴퓨테이션 값 및 로컬 어레이를 포함할 수 있다. 도 8에 있는 코드(800)를 다시 참조하면, 라이브 데이터는 코드 내 "tmp" 어레이의 로컬 슬라이스 및 local_temp 변수의 현재 값을 모두 포함할 수 있다. 마이그레이션이 차후에 연관된 작업 유닛의 실행 동안 발생하면, 라이브 값은 목적지 OpenCL 디바이스로 전송되고 여기에서 초기화될 수 있다. 전술된 바와 같이, 컴파일러 내 코드 생성기는 태그를 삽입하고 다른 코드를 삽입하여 마이그레이션 동안 라이브 값을 호출할 수 있다. 목적지 OpenCL 디바이스에서, 마이그레이션 진입 점을 위한 코드 생성은 라이브 값을 포함하는 데이터 구조를 초기화하고 커널 실행으로 진행한다. 대안적으로, 컴파일러는 전술된 바와 같이 실행으로 진행하는 컴퓨트 서브 커널을 생성할 수 있다. 블록(912)에서, 컴파일러는 적어도 2개의 이종 프로세서 코어에 대한 컴퓨트 커널의 컴파일을 완료한다. 다른 디버그 및 설치 코드가 삽입될 수 있다.

일 실시예에서, 컴파일러는 다수의 데이터 구조를 생성한다. 2개 이상의 데이터 구조는 일반 목적 코어 및 SIMD 코어와 같은 소정의 타깃 OpenCL 디바이스에 있는 각 컴퓨트 서브 커널을 위한 실행가능한 목적 코드를 포함한다. 또 다른 데이터 구조는 마이그레이션의 시간에 전송되고 액세스될 라이브 데이터를 포함한다. 컴퓨트 커널 내 잠재적인 마이그레이션 점으로 지정된 레이블이 주어지면, 컴파일러는 데이터 흐름 분석을 이용하여 전송될 수 있는 라이브 값을 결정할 수 있다. 레지스터에 캐시된 것과 같은 실행 내 이 점으로 한정되지 않은 라이브 값이 실행시간 환경에서 액세스가능한 위치에 배치된다. 이들 위치의 예로는 연관된 원래의 메모리 위치 및 예비된 컨텐츠를 보유하는 레지스터를 포함한다. 일 실시예에서, 휴리스틱 체크(heuristic check)를 이용하여 데이터 전송의 사이즈가 이종 코어들 사이에 유리한 변경 실행(profitable change execution)을 가능하게 하는지의 여부를 결정할 수 있다.

추가적으로 컴파일러는 실행시간 환경에 의해 라이브 데이터를 연관된 목적지 OpenCL 디바이스로 전송하는 것으로 해석되는 다른 데이터 구조를 생성할 수 있다. 이 데이터 구조는 전송될 라이브 데이터의 위치 및 사이즈 및 소스 및 목적지 OpenCL 디바이스의 어드레스 공간 내 위치를 제공할 수 있다. 또한, 컴파일러는 목적지 디바이스를 위한 대응하는 커널 버전을 생성한다. OpenCL 디바이스 각각에 대한 각 컴파일된 코드는 지정된 위치에서 라이브 데이터에 액세스하고 마이그레이션 점에서 실행을 시작한다.

이제 도 10을 참조하면, 선-실행시간 데이터 정보를 이용하는 것에 의해 프로세서에서 다수의 작업 유닛의 병렬 실행을 최적화하는 방법(1000)의 일 실시예가 도시된다. 프로세싱 노드(110)에 구현된 성분 및 전술된 도 4에 도시된 하드웨어 자원 할당은 일반적으로 방법(1000)에 따라 동작할 수 있다. 설명의 목적을 위하여, 이 실시예 및 차후 설명되는 방법의 후속하는 실시예에 있는 단계가 순차적으로 도시된다. 그러나, 일부 단계는 도시된 것과 상이한 순서로 발생할 수 있고, 일부 단계는 동시에 수행될 수 있고, 일부 단계는 다른 단계와 조합될 수 있고, 일부 단계는 다른 실시예에서 부재할 수 있다.

블록(1002)에서, 연관된 데이터 레코드는 소정의 컴퓨트 커널의 각 작업 유닛에 할당된다. 블록(1004)에서, OS 스케줄러(424)는 작업 유닛을 이종 코어로 스케줄링한다. 블록(1006)에서, 이종 프로세서 코어는 대응하는 스케줄링된 작업 유닛을 실행한다.

블록(1008)에서, 소정의 태그된 마이그레이션 점에 도달된다. 일 실시예에서, 현재 사용된 OpenCL 디바이스의 이용율의 측정이 수행될 수 있다. 측정이 이용율을 나타내거나 또는 성능이 소정의 임계값 미만이라면, 연관된 컴퓨트 커널 또는 컴퓨트 서브 커널이 상이한 마이크로-아키텍처를 구비하는 이종 코어와 같은 다른 OpenCL 디바이스로 마이그레이션될 수 있다. 일 실시예에서, 이 측정은 연관된 컴퓨트 커널 또는 컴퓨트 서브 커널 내에 출구 또는 리턴에 도달한 SIMD 코어에서 현재 실행하는 작업 유닛의 수의 카운트이다. 대안적으로, 디스에이블된 컴퓨테이션 유닛의 파면의 수의 카운트는 동일한 수를 제공할 수 있다. 이 카운트가 소정의 임계값을 초과하면, 출구 점에 아직 도달하지 않은 작업 유닛은 일반 목적 코어와 같은 다른 이종 코어로 마이그레이션될 수 있다. 이후 SIMD 코어에서 파면은 해제될 수 있고 다른 스케줄링된 작업 유닛에 이용가능할 수 있다.

다른 실시예에서, 상기 기술은 SIMD 코어에서 파면의 병렬 실행 작업 유닛의 대부분(large fraction)이 아이들 상태이고 나머지 작업 유닛은 실질적으로 실행을 계속하는 것으로 기대되는 것으로 결정된 상황에서 마이그레이션을 개시하는 것으로 확장될 수 있다. 예를 들어, 생성된 데이터 구조는 공유된 메모리에 및 하나 이상의 캐시에 있을 수 있다. 가상 메모리 지원을 구비하는 시스템에서, 작업 유닛의 서브셋은 캐시에 히트(hit)할 수 있는 반면, 나머지 작업 유닛은 가상 메모리 미스(miss)를 나타낼 수 있는데, 이는 긴 레이턴시 이벤트이다. 이 경우에, 전반적인 컴퓨팅 성능은 추가적인 실행이 현재 실행에 의해 인에이블된 프리페치 기술로부터 이익을 얻을 수 있으므로 일반 목적 코어에서 실행을 계속하여 더 우수할 수 있다.

실행 효율이 소정의 임계값 미만인 것으로 결정되지 않으면(조건부 블록(1010)), 방법(1000)의 제어 흐름은 블록(1006)으로 리턴하여 실행을 계속한다. 실행 효율이 소정의 임계값 미만인 것으로 결정되면(조건부 블록(1010)), 블록(1012)에서, 하나 이상의 작업 유닛은 제1 프로세서 코어의 마이크로-아키텍처와는 상이한 마이크로-아키텍처를 구비하는 제2 프로세서 코어로 마이그레이션하도록 식별된다. 식별된 작업 유닛은 상기 측정이 소정의 임계값 미만이 되게 할 수 있다. 블록(1014)에서, 제1 프로세서 코어에 의해 생성된 연관된 로컬 데이터는 글로벌 데이터로 승격된다. 블록(1016)에서, 마이그레이션된 작업 유닛의 컴파일된 버전은 마이그레이션 태그된 점에서 시작하는 제2 프로세서 코어에서 실행되도록 스케줄링된다.

전술된 실시예는 소프트웨어를 포함할 수 있는 것으로 이해된다. 이러한 실시예에서, 방법 및/또는 메커니즘을 구현하는 프로그램 명령은 컴퓨터 판독가능한 매체로 운송되거나 이에 저장될 수 있다. 프로그램 명령을 저장하도록 구성된 다수의 유형의 매체들이 이용가능하고, 하드 디스크, 플로피 디스크, CD-ROM, DVD, 플래시 메모리, 프로그래밍가능한 ROM(PROM), 랜덤 액세스 메모리(RAM), 및 여러 다른 형태의 휘발성 또는 비휘발성 저장매체를 포함한다. 일반적으로 말하면, 컴퓨터 액세스가능한 저장 매체는 명령 및/또는 데이터를 컴퓨터에 제공하기 위해 사용 동안 컴퓨터에 의해 액세스가능한 임의의 저장 매체를 포함할 수 있다. 예를 들어, 컴퓨터 액세스가능한 저장 매체는 자기 또는 광 매체, 예를 들어, 디스크(고정식 혹은 제거가능식), 테이프, CD-ROM, 또는 DVD-ROM, CD-R, CD-RW, DVD-R, DVD-RW, 또는 블루-레이(Blu-Ray)와 같은 저장 매체를 포함할 수 있다. 저장 매체는 USB(Universal Serial Bus) 인터페이스 등과 같은 휘발성 또는 비휘발성 메모리 매체, 예를 들어 RAM(예를 들어, 동기 동적 RAM(SDRAM), 더블 데이터 레이트(double data rate)(DDR, DDR2, DDR3 등) SDRAM, 저전력 DDR (LPDDR2 등) SDRAM, 램버스(Rambus) DRAM(RDRAM), 정적 RAM(SRAM) 등), ROM, 플래시 메모리, 주변 인터페이스를 통해 액세스가능한 비휘발성 메모리(예를 들어 플래시 메모리)를 더 포함할 수 있다. 저장 매체는 마이크로전기기계 시스템(MEMS), 및 네트워크 및/또는 무선 링크와 같은 통신 매체를 통해 액세스 가능한 저장 매체를 포함할 수 있다.

추가적으로, 프로그램 명령은 C와 같은 하이 레벨 프로그래밍 언어, 또는 베리로그(Verilog), VHDL와 같은 설계 언어(HDL), 또는 GDS II 스트림 포맷(GDSII)과 같은 데이터베이스 포맷으로 하드웨어 기능의 거동-레벨 설명 또는 레지스터-전송 레벨(RTL) 설명을 포함할 수 있다. 일부 경우에 이 설명은 설명을 합성하여 합성 라이브러리로부터 게이트 리스트를 포함하는 네트리스트를 생성할 수 있는 합성 툴에 의해 판독될 수 있다. 네트리스트는 시스템을 포함하는 하드웨어의 기능을 더 나타내는 게이트 세트를 포함한다. 네트리스트는 마스크에 적용될 기하학적 형상을 기술하는 데이터 세트를 생성하도록 배치되고 라우팅될 수 있다. 마스크는 시스템에 대응하는 반도체 회로 또는 회로를 생성하도록 여러 반도체 제조 단계에 사용될 수 있다. 대안적으로, 컴퓨터 액세스가능한 저장 매체에 있는 명령은 원하는 대로 네트리스트(합성 라이브러리를 구비하거나 없이) 또는 데이터 세트일 수 있다. 추가적으로, 명령은 Cadence(등록상표), EVE(등록상표) 및 Mentor Graphic(등록상표)과 같은 벤더로부터 하드웨어 기반 유형 에뮬레이터에 의해 에뮬레이션을 위하여 이용될 수 있다.

상기 실시예가 상당히 상세히 설명되었으나, 본 개시 내용이 완전히 이해되면 이 기술 분야에 통상의 지식을 가진 자에게는 다수의 변경 및 변형이 이루어질 수 있을 것이다. 이하 청구범위는 모든 이러한 변경 및 변형을 포함하는 것으로 해석되는 것으로 의도된다.

Claims

방법으로서,
컴퓨트 커널(compute kernel)의 실행 동안 해당 컴퓨트 커널의 실행이 마이그레이션(migrate)될 수 있는 복수의 명령을 포함하는 컴퓨트 커널 내 위치를 식별하는 단계;
상기 컴퓨트 커널의 컨텍스트(context)를 유지하고 마이그레이션하는 데이터 구조를 생성하는 단계;
제1 마이크로-아키텍처를 구비하는 제1 프로세서 코어에서 실행하기 위한 상기 위치 전의 상기 컴퓨트 커널의 코드를 스케줄링하는 단계;
마이그레이션 조건이 만족된 것을 나타내는 지시를 수신한 것에 응답하여:
상기 제1 마이크로-아키텍처와는 상이한 제2 마이크로-아키텍처를 구비하는 제2 프로세서 코어에 의해 액세스가능한 위치로 상기 컨텍스트를 이동시키는 단계; 및
상기 위치 후의 상기 컴퓨트 커널의 코드를 상기 제2 프로세서 코어로 스케줄링하는 단계를 포함하는 방법.
제1항에 있어서, 상기 제1 프로세서 코어에 대응하는 컴퓨트 커널을 위한 코드의 제1 버전을 생성하는 단계, 및 상기 제2 프로세서 코어에 대응하는 컴퓨트 커널을 위한 코드의 제2 버전을 생성하는 단계를 더 포함하는 방법.
제2항에 있어서, 상기 제1 마이크로-아키텍처는 단일 명령 다중 데이터(single instruction multiple data: SIMD) 마이크로-아키텍처이고, 상기 제2 마이크로-아키텍처는 일반 목적 마이크로-아키텍처인 것인 방법.
제2항에 있어서, 프로파일 실행시간 정보 및 정적 정보 중 하나에 적어도 기초하여 상기 식별을 수행하는 단계를 더 포함하는 방법.
제2항에 있어서,
마이그레이션 조건이 만족되었는지의 여부를 결정하는 명령을 구비하는 상기 제1 프로세서 코어를 위한 코드의 제1 버전을 설치하는 단계; 및
상기 데이터 구조에 의해 지시된 위치에서 라이브 값을 찾아 실행을 시작하는 명령을 구비하는 상기 제2 프로세서 코어를 위한 코드의 제2 버전을 설치하는 단계를 더 포함하는 방법.
제5항에 있어서, 마이그레이션 조건이 만족되었는지를 결정하기 위해, 상기 방법은, 출구 점에 도달한 상기 컴퓨트 커널의 병렬 실행 반복의 수가 주어진 임계값을 초과하는지를 결정하는 단계를 더 포함하는 것인 방법.
제5항에 있어서,
상기 컴퓨트 커널의 차후 병렬 실행 반복의 수가 마이그레이션 조건을 만족하는 것으로 예측된 것에 응답하여, 상기 위치에서 상기 컴퓨트 커널을 2개의 컴퓨트 서브 커널로 분할하는 단계;
제1 컴퓨트 서브 커널을 상기 제1 프로세서 코어로 스케줄링하는 단계; 및
제2 컴퓨트 서브 커널을 상기 제2 프로세서 코어로 스케줄링하는 단계를 더 포함하되,
상기 제1 컴퓨트 서브 커널은 상기 주어진 위치 전의 코드를 포함하고, 상기 제2 컴퓨트 서브 커널은 상기 주어진 위치 후의 코드를 포함하는 것인 방법.
제6항에 있어서, 상기 위치는 조건부 브랜치 명령 바로 전에 있는 것인 방법.
이종 멀티-코어 아키텍처를 포함하는 컴퓨팅 시스템으로서,
제1 마이크로-아키텍처를 구비하는 제1 프로세서 코어;
상기 제1 마이크로-아키텍처와는 상이한 제2 마이크로-아키텍처를 구비하는 제2 프로세서 코어; 및
스케줄러를 구비하는 운영 시스템을 포함하되,
상기 스케줄러는,
제1 마이크로-아키텍처를 구비하는 제1 프로세서 코어에서 실행하기 위한 위치 전의 컴퓨트 커널 내 코드를 스케줄링하는 동작; 및
마이그레이션 조건이 만족되는 것을 나타내는 지시를 수신한 것에 응답하여:
상기 제1 마이크로-아키텍처와는 상이한 제2 마이크로-아키텍처를 구비하는 제2 프로세서 코어에 의해 액세스가능한 위치로 상기 컴퓨트 커널의 컨텍스트를 이동시키는 동작; 및
상기 위치 후의 상기 컴퓨트 커널의 코드를 상기 제2 프로세서 코어로 스케줄링하는 동작을 수행하도록 구성된 것인 컴퓨팅 시스템.
제9항에 있어서, 컴파일러를 더 포함하되, 상기 컴파일러는,
복수의 명령을 포함하는 컴퓨트 커널 내 위치를, 상기 컴퓨트 커널의 실행 동안 상기 컴퓨트 커널의 실행이 마이그레이션될 수 있는 위치로서 식별하는 동작; 및
상기 컴퓨트 커널의 컨텍스트를 유지하고 마이그레이션하는 데이터 구조를 생성하는 동작을 수행하도록 구성된 것인 컴퓨팅 시스템.
제10항에 있어서, 상기 제1 마이크로-아키텍처는 단일 명령 다중 데이터(SIMD) 마이크로-아키텍처이고, 상기 제2 마이크로-아키텍처는 일반 목적 마이크로-아키텍처인 것인 컴퓨팅 시스템.
제10항에 있어서, 상기 컴파일러는 프로파일 실행시간 정보 및 정적 정보 중 적어도 하나에 기초하여 상기 식별하는 동작을 수행하도록 더 구성된 것인 컴퓨팅 시스템.
제10항에 있어서, 상기 컴파일러는,
마이그레이션 조건이 만족되었는지의 여부를 결정하는 명령을 구비하는 상기 제1 프로세서 코어를 위한 코드의 제1 버전을 설치하는 동작; 및
상기 데이터 구조에 의해 지시된 위치에서 라이브 값을 찾고 실행을 시작하는 명령을 구비하는 상기 제2 프로세서 코어를 위한 코드의 제2 버전을 설치하는 동작을 수행하도록 더 구성된 것인 컴퓨팅 시스템.
제13항에 있어서, 마이그레이션 조건이 만족되었는지를 결정하기 위해, 상기 제1 및 상기 제2 프로세서 코어 각각은 출구 점에 도달한 컴퓨터 커널의 병렬 실행 반복의 수가 주어진 임계값을 초과하는지를 결정하도록 구성된 것인 컴퓨팅 시스템.
제13항에 있어서, 상기 컴파일러는,
상기 컴퓨트 커널의 차후 병렬 실행 반복의 수가 상기 마이그레이션 조건을 만족하는 것으로 예측된 것에 응답하여, 주어진 위치에서 상기 컴퓨트 커널을 2개의 컴퓨트 서브 커널로 분할하는 동작;
제1 컴퓨트 서브 커널을 상기 제1 프로세서 코어로 스케줄링하는 동작; 및
제2 컴퓨트 서브 커널을 상기 제2 프로세서 코어로 스케줄링하는 동작을 수행하도록 더 구성되고,
상기 제1 컴퓨트 서브 커널은 상기 위치 전의 코드를 포함하며, 상기 제2 컴퓨트 서브 커널은 상기 위치 후의 코드를 포함하는 것인 컴퓨팅 시스템.
제14항에 있어서, 상기 위치는 조건부 브랜치 명령 바로 전에 있는 것인 컴퓨팅 시스템.
프로그램 명령을 저장하는 컴퓨터 판독가능한 저장 매체로서, 상기 프로그램 명령은,
컴퓨트 커널의 실행 동안 해당 컴퓨트 커널의 실행이 마이그레이션될 수 있는 복수의 명령을 포함하는 컴퓨트 커널 내 위치를 식별하는 동작;
상기 컴퓨트 커널의 컨텍스트를 유지하고 마이그레이션하는 데이터 구조를 생성하는 동작;
제1 마이크로-아키텍처를 구비하는 제1 프로세서 코어에서 실행하기 위한 상기 위치 전의 상기 컴퓨트 커널의 코드를 스케줄링하는 동작;
마이그레이션 조건이 만족된 것을 나타내는 지시를 수신한 것에 응답하여:
상기 제1 마이크로-아키텍처와는 상이한 제2 마이크로-아키텍처를 구비하는 제2 프로세서 코어에 의해 액세스가능한 위치로 상기 컨텍스트를 이동시키는 동작; 및
상기 위치 후의 상기 컴퓨트 커널의 코드를 상기 제2 프로세서 코어로 스케줄링하는 동작을 수행하도록 실행가능한 것인 컴퓨터 판독가능한 저장 매체.
제17항에 있어서, 상기 프로그램 명령은 상기 제1 프로세서 코어에 대응하는 상기 컴퓨트 커널을 위한 코드의 제1 버전을 생성하는 동작, 및 상기 제2 프로세서 코어에 대응하는 상기 컴퓨트 커널을 위한 코드의 제2 버전을 생성하는 동작을 수행하도록 더 실행가능한 것인 컴퓨터 판독가능한 저장 매체.
제17항에 있어서, 상기 프로그램 명령은,
마이그레이션 조건이 만족되었는지의 여부를 결정하는 명령을 구비하는 상기 위치에 제1 프로세서 코어를 위한 코드의 제1 버전을 설치하는 동작; 및
상기 데이터 구조에 의해 지시된 위치에서 라이브 값을 찾고 실행을 시작하는 명령을 구비하는 상기 위치에 제2 프로세서 코어를 위한 코드의 제2 버전을 설치하는 동작을 수행하도록 더 실행가능한 것인 컴퓨터 판독가능한 저장 매체.
제19항에 있어서, 마이그레이션 조건이 만족되었는지를 결정하기 위하여, 상기 프로그램 명령은 출구 점에 도달한 상기 컴퓨트 커널의 병렬 실행 반복의 수가 주어진 임계값을 초과하는지를 결정하도록 더 실행가능한 것인 컴퓨터 판독가능한 저장 매체.