KR20160141753A - 컴파일러 최적화를 위한 메모리 참조 메타데이터 - Google Patents

컴파일러 최적화를 위한 메모리 참조 메타데이터 Download PDF

Info

Publication number
KR20160141753A
KR20160141753A KR1020167028388A KR20167028388A KR20160141753A KR 20160141753 A KR20160141753 A KR 20160141753A KR 1020167028388 A KR1020167028388 A KR 1020167028388A KR 20167028388 A KR20167028388 A KR 20167028388A KR 20160141753 A KR20160141753 A KR 20160141753A
Authority
KR
South Korea
Prior art keywords
memory
kernel
compiler
memory reference
runtime
Prior art date
Application number
KR1020167028388A
Other languages
English (en)
Other versions
KR101832656B1 (ko
Inventor
추-초 림
데이비드 사무엘 브랙먼
Original Assignee
퀄컴 인코포레이티드
Priority date (The priority date is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the date listed.)
Filing date
Publication date
Application filed by 퀄컴 인코포레이티드 filed Critical 퀄컴 인코포레이티드
Publication of KR20160141753A publication Critical patent/KR20160141753A/ko
Application granted granted Critical
Publication of KR101832656B1 publication Critical patent/KR101832656B1/ko

Links

Images

Classifications

    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/44—Arrangements for executing specific programs
    • G06F9/455—Emulation; Interpretation; Software simulation, e.g. virtualisation or emulation of application or operating system execution engines
    • G06F9/45504—Abstract machines for programme code execution, e.g. Java virtual machine [JVM], interpreters, emulators
    • G06F9/45516—Runtime code conversion or optimisation
    • G06F9/45525—Optimisation or modification within the same instruction set architecture, e.g. HP Dynamo
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00—Arrangements for software engineering
    • G06F8/40—Transformation of program code
    • G06F8/41—Compilation
    • G06F8/43—Checking; Contextual analysis
    • G06F8/433—Dependency analysis; Data or control flow analysis
    • G06F8/434—Pointers; Aliasing
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00—Arrangements for software engineering
    • G06F8/40—Transformation of program code
    • G06F8/41—Compilation
    • G06F8/44—Encoding
    • G06F8/443—Optimisation
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/44—Arrangements for executing specific programs
    • G06F9/455—Emulation; Interpretation; Software simulation, e.g. virtualisation or emulation of application or operating system execution engines
    • G06F9/45504—Abstract machines for programme code execution, e.g. Java virtual machine [JVM], interpreters, emulators
    • G06F9/45516—Runtime code conversion or optimisation

Landscapes

  • Engineering & Computer Science (AREA)
  • Software Systems (AREA)
  • Theoretical Computer Science (AREA)
  • General Engineering & Computer Science (AREA)
  • Physics & Mathematics (AREA)
  • General Physics & Mathematics (AREA)
  • Devices For Executing Special Programs (AREA)

Abstract

장치는 메모리 및 컴파일링 프로세서를 포함하고, 상기 컴파일링 프로세서는, 상기 컴파일링 프로세서 상에서 실행되는 컴파일러 및 런타임으로 이루어지는 군의 적어도 하나에 의해, 컴파일된 커넬을 실행하기 위해 아규먼트들을 생성하고, 상기 컴파일링 프로세서 상에서 실행되는 군의 적어도 하나에 의해, 아규먼트들의 제 1 메모리 영역에 대한 제 1 메모리 참조 및 제 2 메모리 영역에 대한 제 2 메모리 참조가 동일한 메모리 영역을 참조하는지 여부를 결정하고; 그 군의 적어도 하나에 의해, 상기 결정에 기초하여 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조와 연관된 메타데이터를 생성하게 하는 것으로서, 상기 메타데이터는 상기 제 1 메모리 영역과 상기 제 2 메모리 영역 사이의 관계를 나타내는, 상기 메타데이터를 생성한다. 상기 컴파일러 및 런타임 중 적어도 하나는 메타데이터에 기초하여 커넬을 리컴파일하고, 리컴파일된 커넬을 실행하도록 타겟 프로세서에 명령하도록 구성된다.

Description

컴파일러 최적화를 위한 메모리 참조 메타데이터{MEMORY REFERENCE METADATA FOR COMPILER OPTIMIZATION}
본 개시는 커넬들의 소스 코드를 컴파일하는 것에 관한 것이고, 보다 구체적으로는 메모리 액세스를 위한 커넬들의 소스 코드를 컴파일하기 위한 기법에 관한 것이다.
소위 이종 컴퓨팅 아키텍쳐를 향한 움직임이 있다. 이종 컴퓨팅 아키텍쳐들에서, 커넬로 지칭되는 프로그램은, 다양한 상이한 유형의 프로세서들, 이를테면 CPU (Central Processing Unit), GPU (Graphics Processing Unit), FPGA (Field Programmable Gate Array) 등이 커넬을 실행할 수 있도록 프레임워크를 이용하여 컴파일될 수도 있다. 이종 컴퓨팅을 지원하는 최근의 프레임워크들은 OpenCL 프레임워크, 그리고 DirectCompute 프레임워크를 포함한다.
본 개시는 컴파일링 최적화를 위한 메타데이터를 생성하기 위하여 커넬에서 메모리 참조들의 메모리 오버랩 및 메모리 에일리어싱을 검출하기 위한 기법들을 설명한다. 본 개시의 기법들을 수행하기 위하여, 저스트 인타임 컴파일러 (just-in-time compiler; JIT) 와 같은 컴파일러는, "커넬" 로도 지칭되는 프로그램의 소스 코드를 2진 파일로 컴파일한다. 컴파일러를 실행되는, 컴파일링 프로세서는 (컴파일링 프로세서가 커넬을 실행하기 위해 필요한 아규먼트 (argument) 들을 생성할 때) 런타임에서, OpenCL 과 같은 이종 컴퓨팅 네트워크를 이용하여 커넬을 컴파일할 수도 있다. 본 개시에 설명된 기법들에서, 생성된 아규먼트들을 이용하여 커넬을 실행하도록 타겟 프로세서에 명령하는 대신에, 드라이버는, 커넬을 실행할 타겟 프로세서에 보내질, 버퍼에서 함께 보내지는, 아규먼트들을 분석한다. 그 분석에 기초하여, 드라이버/런타임은 제 1 메모리 참조와 제 2 메모리 참조 사이의 관계 (예를 들어, 제 1 메모리 참조 및 제 2 메모리 참조의 메모리 영역이 오버랩하는지 여부, 어느 정도인지 등) 을 나타내는 메타데이터를 생성한다.
메모리 영역들이 동일하지 않으면, 컴파일링 프로세서는 메타데이터에 기초하여, 그리고 루프 언롤링 (loop unrolling) 등과 같은 보다 공격적인 (aggressive) 컴파일 기법을 이용하여, 커넬을 리컴파일하기 위해 컴파일러를 이용할 수도 있다. 드라이버는 또한, 커넬의 메모리 액세스가 오버랩되는 정도를 결정 가능할 수도 있고, 메모리 오버랩의 양에 기초하여 보다 공격적인 기법들을 이용하여 커넬을 리컴파일할 수도 있다. 이런 식으로, 본 개시의 기법들은 저스트 인타임 컴파일러를 이용하여 컴파일되는 커넬의 실행 성능을 증가시킬 수도 있다.
일 예에서, 본 개시는, 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 컴파일된 커넬의 2진 코드를 실행하기 위한 아규먼트를 생성하는 단계, 그 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 커넬 아규먼트들의 제 1 메모리 영역에 대한 제 1 메모리 참조 및 커넬 아규먼트들의 제 2 메모리 영역에 대한 제 2 메모리 참조가 동일한 메모리 영역을 참조하는지 여부를, 결정하는 단계, 그 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 그 결정에 기초하여 제 1 메모리 참조 및 제 2 메모리 참조와 연관된 메타데이터를 생성하는 단계를 포함하는 방법을 설명한다. 메타데이터는 제 1 메모리 영역과 제 2 메모리 영역 사이의 관계를 나타낼 수도 있다. 그 방법은 또한, 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 커넬의 제 1 및 제 2 메모리 참조들이 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여, 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 컴파일러로 하여금 메타데이터에 기초하여 커넬을 리컴파일하게 하는 단계, 및 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 리컴파일된 커넬을 실행하도록 타겟 프로세서에 명령하는 단계를 더 포함한다.
다른 예에서, 본 개시는, 메모리 및 컴파일링 프로세서를 포함하는 디바이스를 설명하고, 그 컴파일링 프로세서는, 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 컴파일된 커넬의 2진 코드를 실행하기 위한 아규먼트를 생성하고, 그 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 커넬 아규먼트들의 제 1 메모리 영역에 대한 제 1 메모리 참조 및 커넬 아규먼트들의 제 2 메모리 영역에 대한 제 2 메모리 참조가 동일한 메모리 영역을 참조하는지 여부를 결정하고, 그 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 그 결정에 기초하여 제 1 메모리 참조 및 제 2 메모리 참조와 연관된 메타데이터를 생성하도록 구성된다. 그 메타데이터는 제 1 메모리 영역과 제 2 메모리 영역 사이의 관계를 나타낼 수도 있고, 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 커넬의 제 1 및 제 2 메모리 참조들이 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여, 그 컴파일링 프로세서는 또한, 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 컴파일러로 하여금 메타데이터에 기초하여 커넬을 리컴파일하게 하고, 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 리컴파일된 커넬을 실행하도록 타겟 프로세서에 명령하도록 구성된다.
다른 예에서, 본 개시는, 명령들을 저장한 비일시적 컴퓨터 판독가능 저장 매체를 설명하고, 그 명령들은, 실행될 때, 컴파일링 프로세서로 하여금, 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 컴파일된 커넬의 2진 코드를 실행하기 위한 아규먼트를 생성하게 하고, 그 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 커넬 아규먼트들의 제 1 메모리 영역에 대한 제 1 메모리 참조 및 커넬 아규먼트들의 제 2 메모리 영역에 대한 제 2 메모리 참조가 동일한 메모리 영역을 참조하는지 여부를 결정하게 하고, 그 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 그 결정에 기초하여 제 1 메모리 참조 및 제 2 메모리 참조와 연관된 메타데이터를 생성하게 한다. 메타데이터는 제 1 메모리 영역과 제 2 메모리 영역 사이의 관계를 나타내고,
컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 커넬의 제 1 및 제 2 메모리 참조들이 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여, 컴파일링 프로세서는 또한, 컴파일링 프로세서로 하여금, 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 컴파일러로 하여금 메타데이터에 기초하여 커넬을 리컴파일하게 하고, 및 컴파일링 프로세서 상에서 실행되는 런타임 및 컴파일러로 이루어지는 군 중 적어도 하나에 의해, 리컴파일된 커넬을 실행하도록 타겟 프로세서에 명령하게 하는 명령들을 실행하도록 구성된다.
본 개시의 하나 이상의 예들의 상세들은 첨부 도면 및 아래의 설명에 제시되어 있다. 본 개시의 다른 특징, 목적 및 이점들은 상세한 설명 및 도면, 그리고 특허청구범위로부터 분명해질 것이다.
도 1은 본 개시의 기법들에 따라 컴파일링 최적화를 돕기 위한 에일리어싱 분석을 지원하는 예시적인 컴퓨팅 디바이스를 예시하는 블록도이다.
도 2는 본 개시의 기법들에 따라 커넬을 실행할 수도 있는 프로세서의 하나 이상의 쉐이더 코어들의 복수의 처리 엘리먼트들을 예시하는 개념도이다.
도 3a는, 본 개시의 기법들에 따라, 실행될 때, 에일리어싱을 일으킬 수도 있는 코드를 포함하는 커넬 코드를 예시하는 개념도이다.
도 3b는 본 개시의 기법들에 따라 구성된 컴파일러가 검출가능할 수도 있는 에일리어싱의 예를 예시하는 개념도이다.
도 3c는 본 개시의 기법들에 따라 구성된 컴파일러가 검출가능할 수도 있는 오버랩되지 않는 메모리 참조들의 예를 예시하는 개념도이다.
도 3d는 본 개시의 기법들에 따라 구성된 드라이버/런타임이 검출할 수도 있는 오버랩되는 메모리 참조들을 예시하는 개념도이다.
도 4a는 본 개시의 기법들에 따른 루프 언롤링을 예시하는 개념도이다.
도 4b는 본 개시의 기법들에 따른 코드 리오더링을 예시하는 개념도이다.
도 4c는 본 개시의 기법들에 따른 코드 벡터화를 예시하는 개념도이다.
도 5는 본 개시의 기법들에 따른 컴파일러 최적화를 돕기 위한 컴파일러 메타데이터를 생성하는 예시적인 방법의 흐름도이다.
상세한 설명
위에서 간략하게 설명된 바처럼, 다양한 이종 컴퓨팅 프레임워크들이 현재 개발중에 있다. 이종 컴퓨팅 프레임워크들의 일부 예들은 Khronos 그룹에 의해 현재 개발중인 OpenCLTM 프레임워크, 및 Microsoft® 에 의해 현재 개발되고 있는 DirectCompute 프레임워크를 포함한다. 이종 컴퓨팅 프레임워크들은 단일 프로그램 또는 "커넬" 이 CPU (Central Processing Unit), GPU (Graphics Processing Unit), FPGA (Field Programmable Gate Array), DSP (Digital Signal Processor) 등과 같은 여러 상이한 프로세서들 상에서 실행될 수 있게 한다.
실행을 위해 커넬을 준비하기 위하여, 본 개시에서 컴파일링 프로세서로 지칭되는 프로세서는 타겟 프로세서에 의해 실행될 2진 코드를 생성하기 위하여 커넬 소스 코드를 컴파일한다. 타겟 프로세서는 타겟 프로세서와는 상이한 또는 동일한 프로세서일 수도 있다. 컴파일링 프로세서가 사용하는 컴파일러의 일 예는 저스트 인타임 컴파일 (JIT) 컴파일러로 지칭된다. JIT 컴파일러는, ("인터프리테이션" (interpretation) 으로 지칭되는) 명령들을 사전에 전혀 컴파일하지 않고서 또는 (때로는, "어헤드 오브 타임" (ahead of time) 컴파일로 지칭되는) 실행 전보다는 (런타임으로도 지칭되는) 실행시에 소스 코드를 컴파일한다.
커넬이 컴파일되고 나면, 컴파일링 프로세서는, 드라이버 및 런타임을 통해, 커넬의 컴파일된 2진 코드를 타겟 프로세서로 전송한다. 커넬은 또한, 타겟 프로세서 상의 커넬을 실행하기 위해 런타임에서 아규먼트들의 세트를 받아들이고, 이들을 컴파일링 프로세서는 또한 타겟 프로세서에 전송한다. 커넬 아규먼트들은, 아규먼트들을 위해 할당된 메모리의 버퍼, 즉, 영역을 포함한다. 대부분의 경우에, 커넬은 아규먼트들 상에서 동작 (즉, 읽기 또는 쓰기) 하는 코드 섹션들을 포함한다. 이런 식으로, 아규먼트들은 커넬이 동작할 수도 있는 커넬용 데이터 세트를 포함한다. 커넬을 타겟 프로세서에 전송한 후에, 컴파일링 프로세서의 드라이버/런타임은, 일부 예들에서 런타임에서 아규먼트들을 커넬에 제공하는 함수 호출을 실행한다. 커넬이 아규먼트들을 수신하고 나면, 타겟 프로세서는 커넬의 실행을 시작할 수도 있다.
많은 사례들에서, 커넬은 타겟 프로세서가 어떤 불리언 조건이 만족되었다고 결정할 때까지 또는 어떤 수의 반복들 동안 실행하는, 루프와 같은, 코드 세그먼트들을 포함한다. 컴파일러는, 루프 코드 섹션들을 실행하는 성능을 향상시키기 위한 다양한 기법들, 예를 들어 루프 언롤링, 그리고 루프 및 비루프 코드 섹션들 양자 모두의 성능을 향상시킬 수도 있는 코드 리오더링 (reordering) 및/또는 벡터화 (vectorization) 와 같은 다른 기법들을 채용가능할 수도 있다.
루프 언롤링은, 컴파일러가 산술 연산, 루프 테스트 종료와 같은 루프를 제어하는 명령들을 감소 또는 제거하기 위하여, 및/또는 루프를 실행할 시 캐시 성능을 향상시키기 위하여 루프의 여려 반복들을 확장시키는 최적화 프로세스이다. 코드 리오더링은, 컴파일러가 일련의 유사한 명령들 (예를 들어, 로딩 또는 저장을 함께) 을 그룹화하기 위하여 사용할 수도 있는 또 다른 최적화이다. 코드 리오더링은 일부 경우들에서 루프 코드 섹션을 실행할 때 캐시 성능을 향상시킬 수도 있다. 예를 들어, 코드 리오더링은 (예를 들어, 루프 바디 내에서) 다수의 로드 명령들을 합칠 (coalescing) 때 성능을 향상시킬 수도 있으며, 이와 함께, 스칼라 명령에서 사용되는 피연산자의 크기의 배수인 캐시 라인 폭 (아래에서 더 상세하게 논의됨) 을 갖는 시스템 상에서 성능을 향상시킬 수도 있다. 하지만, 컴파일러가, 컴파일 이전에, 로드/저장 버퍼가 서로 에일리어싱하지 않는다고 결정하는 경우에만 컴파일러가 로드들을 합치는 것이 안전할 수도 있다. 그렇지 않으면, 리오더링된 로드/저장 명령들에 기인하여 데이터 손상 (data corruption) 이 일어날 수 있다.
벡터화는, 컴파일러가, 각각이 한번에 단일 쌍의 피연산자를 처리하는 여러 스칼라 연산들을 포함하는 소스 코드를, 한번에 다수의 쌍의 피연산자에 대해 하나의 연산을 처리하는 벡터 명령으로, 변환할 수도 있는 또 다른 최적화 프로세스이다. 벡터화는 동일한 코드의 스칼라 구현에 비해 성능을 향상시킬 수도 있는 병렬화 (parallelization) 의 형태이다. 루프 언롤링, 코드 리오더링, 및 벡터화는 아래에서 더 상세하게 설명된다.
커넬의 코드 섹션들은, 아규먼트들의 메모리 영역을 참조할 수도 있는, "포인터" 로도 지칭되는, 메모리 참조들을 포함할 수도 있다. 예를 들어, 코드 섹션은 커넬 아규먼트들의 부분들을 참조할 수도 있는 일련의 메모리 참조들 (즉, 커넬 아규먼트들에 포함된 버퍼에 대한 메모리 참조들) 을 포함할 수도 있다. 커넬은 아규먼트 버퍼들로부터 값들을 판독할 수도 있고, 또한 데이터를 아규먼트 버퍼에 기록할 수도 있다.
일부 경우에, 상이한 메모리 참조들, 예를 들어, 상이한 이름을 갖는 포인터 변수들은, 메모리에 동일한 데이터 로케이션을 참조할 수도 있다. 상이한 심볼 참조들이 동일한 메모리 영역을 참조하는 상황은 "에일리어싱" (aliasing) 으로 지칭된다. 컴파일러는 정적 분석 또는 다른 기법을 사용하여 컴파일 시간에 에일리어싱을 검출하는 것을 시도할 수도 있다. 하지만, 컴파일러는 보통, 루프 코드 섹션들에서 메모리 참조들이 참조하는 데이터 (예를 들어, 커넬 아규먼트들) 가 런타임에서 공급될 때 루프들에서의 메모리 참조들의 에일리어싱을 검출할 수 없다.
컴파일러가 메모리 참조들이 동일한 메모리 영역을 참조하는지 (즉, 메모리 참조들이 에일리어싱을 초래하는지) 여부를 확정적으로 결정할 수 없을 때, 컴파일러는 루프에 대한 루프 언롤링 및 벡터화와 같은 최적화 기법들을 수행할 수 없을 수도 있다. 본 개시의 기법들은, JIT 컴파일러로 하여금 커넬 루프의 메모리 액세스들이 동일 메모리 영역을 참조하는지 여부를 결정하는 것을 가능하게 할 수도 있다. 추가적으로, 본 개시의 기법들은 JIT 컴파일러로 하여금 메모리 참조들 사이의 관계에 관한 메타데이터를 생성하는 것과, 생성된 메타데이터에 기초하여 벡터화 및 루프 언롤링과 같은 최적화를 사용하여 커넬을 리컴파일하는 것을 가능하게 한다.
도 1은 본 개시의 기법들에 따라 컴파일링 최적화를 돕기 위한 에일리어싱 분석을 지원하는 예시적인 컴퓨팅 디바이스를 예시하는 블록도이다. 도 1은 컴퓨팅 디바이스 (2) 를 포함한다. 컴퓨팅 디바이스 (2) 는 퍼스널 컴퓨터, 데스크톱 컴퓨터, 랩톱 컴퓨터, 컴퓨터 워크스테이션, 태블릿 컴퓨팅 디바이스, 비디오 게임 플랫폼 또는 콘솔, 무선 통신 디바이스 (이를테면, 예를 들어, 이동 전화기, 셀룰러 전화기, 위성 전화기, 및/또는 이동 전화기 핸드셋, 휴대 비디오 게임 디바이스 또는 PDA (personal digital assistant) 과 같은 핸드헬드 디바이스, 퍼스널 뮤직 플레이어, 비디오 플레이어, 디스플레이 디바이스, 텔레비젼, 텔레비젼 셋톱 박스, 서버, 중간 네트워크 디바이스, 메인프레임 컴퓨터 또는 그래픽 데이터를 처리 및/또는 표시하는 임의의 다른 타입의 디바이스를 포함할 수도 있다.
도 1의 예에 예시된 바처럼, 컴퓨팅 디바이스 (2) 는 CPU (16), 시스템 메모리 (14), GPU (graphics processing unit) (12), 저스트 인타임 (JIT) 컴파일러 (18), 및 드라이버/런타임 (19) 을 포함한다. CPU (16) 는 다양한 유형의 애플리케이션들을 실행할 수도 있다. 애플리케이션들의 예들은, 웹 브라우저, 이메일 애플리케이션, 스프레드시트, 비디오 게임, 디스플레이를 위해 볼 수 있는 오브젝트들을 생성하는 애플리케이션들 등을 포함한다. 하나 이상의 애플리케이션들의 실행을 위한 명령들은 시스템 메모리 (14) 내에 저장될 수도 있다.
CPU (16) 는 또한, JIT 컴파일러 (18) 를 실행할 수도 있다. 따라서, CPU (16) 는, 예의 목적을 위해, "컴파일링 프로세서" 로 지칭될 수도 있다. JIT 컴파일러 (18) 는, CPU (16) 에 의해 실행될 때, 위에 설명된 바처럼, OpenCL 또는 DirectCompute 와 같은, 이종 컴퓨팅 네트워크를 이용하여 커넬의 소스 코드를 컴파일할 수도 있는 컴파일러를 포함한다. JIT 컴파일러 (18) 는 타겟 프로세서에 의한 실행을 위해 자연어 코드 또는 중간 코드 (예를 들어, 바이트코드) 로 소스코드를 컴파일한다. JIT 컴파일러 (18) 는 실행 전과 반대로 "런타임" 에서, 즉 실행시에 컴파일을 수행한다. JIT 컴파일러 (18) 는 OpenCL 를 이용하여 컴파일할 때 clBuildProgram() 함수를 이용하여 컴파일을 수행할 수도 있다. 추가적으로, JIT 컴파일러 (18) 는, 타겟 프로세서, GPU (12) 상에서 실행하는 어떤 파이버 (즉, 스레드) 의 데이터 액세스들이 독립적인지 여부, 그리고 다른 조건들이 유지되는지를 결정하기 위하여 커넬 (20) 의 데이터 액세스 패턴들을 분석하도록 구성될 수도 있다.
드라이버/런타임 (19) 은 또한, 2진 명령 또는 바이트코드 명령들로 커넬 소스 코드를 번역하기 위해 JIT 컴파일러 (18) 와 상호작용한다. 드라이버/런타임 (19) 은, 타겟 프로세서 (이 예에서는 GPU (12)) 를 위한 자연어 또는 목적 코드로의 커넬 소스 코드 명령들의 아키텍처 특정 컴파일을 수행하는데 드라이버를 이용할 수도 있다. 예를 들어, 드라이버/런타임 (19) 은, 타겟 프로세서에 이용가능한 특정 벡터 명령 또는 실행 리소스들을 알 수도 있고, 타겟 프로세서 상에서 실행 성능을 최적화하는 방식으로 자연어 코드로 소스 코드를 컴파일할 수도 있다. 일부 예들에서, 예를 들어, 다수의 타겟 프로세서들이 있는 경우, 예를 들어, 커넬이 CPU (16) 및 GPU (12) 상에서 실행될 것인 경우, 상이한 드라이버들이 있을 수도 있다.
커넬 (20) 은, 타겟 프로세서, 이 예에서는 GPU (12) 가 실행할 수 있는, 자연어 또는 목적 코드, 예를 들어, 2진 명령들로 구성된다. JIT 컴파일러 (18) 는 또한, GPU (12) 의 런타임 실행을 관리할 수도 있다. CPU (16) 는 실행을 위해 커넬 (20) 을 GPU (12) 로 송신할 수도 있다. CPU (16) 는 또한, 아규먼트들 (26) 을 생성할 수도 있고, 이들을 CPU (16) 가 추가 처리를 위해 GPU (12) 로 전송할 수도 있다.
아규먼트들 (26) 을 할당하기 전에, CPU (16) 는 아규먼트들 (26) 을 위해, 메모리의 영역인, 여유 메모리 버퍼 (free memory buffer) 를 할당한다. 버퍼가 할당되면, 드라이버/런타임 (19) 은 버퍼에 아규먼트들 (26) 을 저장한다. 아규먼트들 (26) 은, GPU (12) 가 처리할 수 있는 복수의 데이터 값들 (예를 들어, 정수, 부동 소수점 값, 오브젝트, 값들의 어레이 등) 을 포함할 수도 있다. 추가적으로, 커넬 (20) 의 실행 동안, GPU (12) 는 출력되는 아규먼트들 (26) 을 저장하는 버퍼에 데이터를 쓸 수도 있다. 출력된 데이터는 출력된 아규먼트들을 포함할 수도 있고, 이들을 GPU (12) 가 다시 CPU (16) 로 전송할 수도 있다.
CPU (16) 가 GPU (12) 로 전송하는 아규먼트는 "입력 아규먼트" 로 지칭될 수도 있다. 이종 컴퓨팅 프레임워크가 OpenCL 프레임워크인 예에서, 드라이버/런타임 (19) 은 아규먼트를 생성하고 런타임에서 clSetKernelArg() 함수로 보낼 (이용가능하게 할) 수도 있다. clSetKernelArg() 함수는 커넬 아규먼트들 (26) 중 임의의 것 뿐만아니라, 아규먼트로서 커넬 (20) 을 수신하고, 실행이 시작될 수도 있도록 그 아규먼트들을 GPU (12) 로 전송한다.
아규먼트들 (26) 을 위한 메모리를 할당하는 부분으로서, 드라이버/런타임 (19) 은, 커넬에 포함된 메모리 참조들 중 일부 또는 전부와 연관된 아규먼트들 (26) 의 메모리 영역 및 어드레스를 결정한다. 메모리 참조들은, 특정 코드 섹션들, 예를 들어, "루프 코드 섹션들" 로 지칭되는, 루프들을 포함하는 코드 섹션들의 메모리 참조들일 수도 있다. 결정된 메모리 영역들에 기초하여, 드라이버/런타임 (19) 은, 커넬 (20) 의 루프 코드 섹션들 또는 다른 코드 섹션들의 메모리 참조들이 아규먼트들 (26) 의 동일한 메모리 영역을 참조하는지 여부를 해결 (즉, 결정) 가능할 수도 있다.
커넬 (20) 을 실행할 GPU (12) 를 위하여 커넬 아규먼트들 (26) 을 생성하는 것에 응답하여, 드라이버/런타임 (19) 은 커넬 (20) 을 실행할 수도 있다. 보다 구체적으로, 드라이버/런타임 (19) 은 clEnqueueNDRangeKernel() 함수를 이용하여, 커넬 (20) 을 타겟 프로세서, GPU (12) 로 디스패치 (dispatch) 할 수도 있다. 런타임에서, 드라이버/런타임 (19) 은, 커넬 (20) 이 수신하는 아규먼트들 (26) 을 분석한다. 드라이버/런타임 (19) 은 또한, 메모리 참조들이 아규먼트들 (26) 에 할당된 메모리 영역의 동일한 메모리 영역을 참조하는지 여부를 결정하기 위하여, 메모리 참조들 (예를 들어, 포인터들) 등을 분석한다. 드라이버/런타임 (19) 은 메모리 참조들이 동일한 메모리 영역을 참조하는지 여부를 결정하기 위하여 쌍단위 방식 (pairwise fashion) 으로 메모리 참조들과 아규먼트 버퍼들을 분석할 수도 있다.
드라이버/런타임 (19) 은 또한, 메모리 참조들이 참조하는 아규먼트들 (26) 의 메모리 영역들 사이의 관계에 기초하여 메모리 참조들과 연관된 메타데이터를 생성한다. 메타데이터는 메모리 참조들 사이의 관계를 나타낼 수도 있다. 가령, 메타데이터는 일부 비제한적인 예들로서, 오버랩되는 메모리 참조들의 리스트, 메모리 영역들과 연관된 메모리 영역들이 오버랩되는지 여부, 어느 정도로 메모리 영역들이 오버랩되는지, 그리고 몇개의 바이트들을 오버랩이 포함하는지를 포함할 수도 있다.
드라이버/런타임 (19) 은 생성된 메타데이터 (가 있다면) 를 다시 JIT 컴파일러 (18) 에 제공한다. 메타데이터에 기초하여 2개의 메모리 참조들이 정확하게 동일한 메모리 영역을 공유하지 않는다고 결정하는 것에 응답하여, 드라이버/런타임 (19) 은 루프 언롤링, 코드 리오더링 및/또는 벡터화와 같은 다양한 최적화를 이용하여 커넬 (20) 을 리컴파일할 수도 있다. JIT 컴파일 (18) 는 생성된 메타데이터에 기초하여 루프 언롤링 코드 리오더링, 및/또는 벡터화의 이들 다양한 최적화들을 적용할 수도 있다.
본 개시의 기법들에 따르면, 컴파일링 프로세서, 예를 들어 CPU (16) 는, 컴파일링 프로세서 상에서 실행되는 JIT 컴파일러 (18) 및 드라이버/런타임 (19) 으로 이루어지는 군의 적어도 하나를 이용하여, 컴파일된 커넬 (20) 의 코드 (예를 들어, 2진 코드 또는 오브젝트 코드) 를 실행하기 위해 아규먼트들 (26) 을 생성하도록 구성될 수도 있다. JIT 컴파일러 (18) 및 드라이버/런타임 (19) 으로 이루어지는 군의 적어도 하나는 또한, 커넬 아규먼트들의 제 1 메모리 영역에 대한 제 1 메모리 참조 및 커넬 아규먼트들의 제 2 메모리 영역에 대한 제 2 메모리 참조가 동일한 메모리 영역을 참조하는지 여부를 결정하도록 구성될 수도 있다. JIT 컴파일러 (18) 및 드라이버/런타임 (19) 으로 이루어지는 군의 적어도 하나에 의해, 커넬의 제 1 및 제 2 메모리 참조들이 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여, CPU (16) 는 또한, CPU (16) 상에서 실행되는 JIT 컴파일러 (18) 및 드라이버/런타임 (19) 으로 이루어지는 군의 적어도 하나로, JIT 컴파일러 (18) 및 드라이버/런타임 (19) 으로 이루어지는 군의 적어도 하나로 하여금 메타데이터에 기초하여 커넬 (20) 을 리컴파일하고, CPU (16) 상에서 실행되는 JIT 컴파일러 (18) 및 드라이버/런타임 (19) 으로 이루어지는 군의 적어도 하나에 의해, 리컴파일된 커넬 (20) 을 실행하도록 타겟 프로세서, 예를 들어, GPU (12) 에 명령하도록 구성될 수도 있다.
GPU (12) 는, 그래픽스 데이터를 처리하는데 아주 잘 맞는, 대량 병렬 처리를 가능하게 하는 전문화된 하드웨어일 수도 있다. 이런 식으로, CPU (16) 는 GPU (12) 에 의해 더 잘 핸들링되는 그래픽스 처리를 오프로딩한다. CPU (16) 는 특정 API (application processing interface) 또는 이종 컴퓨팅 프레임워크에 따라 GPU (12) 와 통신할 수도 있다. 그러한 API 들의 예들은 Microsoft ® 에 의한 DirectX ® API 및 Khronos 그룹에 의한 OpenGL ® 을 포함하고; 이종 컴퓨팅 프레임워크들의 예들은 Microsoft 에 의한 DirectCompute, Khronos 그룹에 의한 OpenCLTM 을 포함한다. 하지만, 본 개시의 양태들은 위에 설명된 API 및 프레임워크들에 한정되지 않고, 다른 유형의 API 들에 확장될 수도 있다.
CPU (16) 및 GPU (12) 의 예들은, 디지털 신호 프로세서 (DSP), 범용 마이크로프로세서, ASIC (application specific integrated circuit), FPGA (field programmable logic array), 또는 다른 동등한 직접 또는 이산 로직 회로를 포함하지만 이에 한정되지는 않는다. 일부 예들에서, GPU (12) 는 그래픽스 처리에 적합한 대량 병렬 처리 능력을 GPU (12) 에 제공하는 집적 및/또는 이산 로직 회로를 포함하는 전문화된 하드웨어일 수도 있다. 일부 사례들에서, GPU (12) 는 또한 범용 처리를 포함할 수도 있고, 범용 GPU (GPGPU) 로 지칭될 수도 있다. 본 개시에 설명된 기법들은 GPU (12) 가 GPGPU 인 예들에 적용가능하다.
시스템 메모리 (14) 는 하나 이상의 컴퓨터 판독가능 저장 매체를 포함할 수도 있다. 시스템 메모리 (14) 의 예들은, RAM (random access memory), ROM (read only memory), EEPROM (electrically erasable programmable read-only memory), 플래시 메모리, 또는 명령들 및/또는 데이터 구조들의 형태로 원하는 프로그램 코드를 지니거나 또는 저장하는데 사용될 수 있고 컴퓨터 또는 프로세서에 의해 액세스될 수 있는 임의의 다른 매체를 포함하지만, 이에 한정되지는 않는다.
일부 양태들에서, 시스템 메모리 (14) 는 CPU (16) 및/또는 GPU (12) 로 하여금 본 개시에서 CPU (16) 및 GPU (12) 로 돌려지는 기능들을 수행하게 하는 명령들을 포함할 수도 있다. 따라서, 시스템 메모리 (14) 는, 하나 이상의 프로세서들로 하여금, 예를 들어 CPU (16) 및 GPU (12) 로 하여금 다양한 기능들을 수행하게 하는 명령들을 포함하는 컴퓨터 판독가능 저장 매체일 수도 있다.
시스템 메모리 (14) 는, 일부의 예들에서, 비일시적 저장 매체로서 고려될 수도 있다. 용어 "비일시적" 이란 저장 매체가 캐리어 파에서 구체화되지 않는다는 것을 나타낼 수도 있거나 또는 예로서, 비일시적 저장 매체는, 경시적으로 (예를 들어, RAM 에서) 변화할 수 있는 데이터를 저장할 수도 있다.
JIT 컴파일러 (18) 및 드라이버/런타임 (19) 을 이용하여, CPU (16) 는 GPGPU 애플리케이션들을 위해 자연어 코드 (예를 들어, 커맨드 및 데이터) 또는 바이트코드로 소스 코드를 컴파일할 수도 있다. 예시적인 GPGPU 데이터 및 커맨드들은 광선 추적 애플리케이션, 물리 시뮬레이션을 위한 커맨드 및 장면 데이터, 또는 임의의 다른 유형의 GPGPU 커넬을 위한 데이터를 포함한다. GPGPU 애플리케이션들, 예를 들어, 커넬 (20) 은 또한, DirectX, 또는 OpenGL 와 같은 그래픽스 API 를 이용하여, 또는 Open Compute Language (OpenCL), 또는 OpenCompute, 또는 DirectCompute 와 같은 보다 일반적인 목적의 컴퓨트 API 를 이용하여 컴파일될 수도 있다. CPU (16) 는 커넬 (20) 을 위한 데이터를 처리를 위해 커맨드 버퍼에 송신할 수도 있다. 다양한 예들에서, 커맨드 버퍼는 시스템 메모리 (14) 의 부분, 또는 GPU (12) 의 부분일 수도 있다. 일부 예들에서, CPU (16) 는, PCI-Express 버스 또는 다른 범용 직렬 또는 병렬 버스와 같은 특수 목적 버스를 통해 처리할 GPU (12) 를 위한 커넬 (20) 의 커맨드 및 데이터를 송신할 수도 있다.
커맨드 버퍼에 저장된 커넬 (20) 의 동작들을 수행하기 위하여, GPU (12) 는 처리 파이프라인을 구현할 수도 있다. 처리 파이프라인은 GPU (12) 상에서 실행되는 소프트웨어 또는 펌웨어에 의해 정의되는 기능들을 수행하는 것 및 매우 특정한 기능들을 수행하기 위하여 하드와이어링된 고정된 기능 유닛 (fixed-function unit) 들에 의해 기능들을 수행하는 것을 포함한다. 커넬 (20) 의 실행을 위해 고정된 기능 유닛들을 바이패스하는 것이 가능할 수도 있거나 또는 커넬 (20) 의 실행은 고정된 기능 유닛들을 사용할 수도 있다.
커넬 (20) 은 GPU (12) 의 하나 이상의 처리 엘리먼트들 (“쉐이더 코어” 또는 “PE” 들로도 지칭됨) 상에서 실행될 수도 있다. 쉐이더 코어들 (22) 은 사용자들에게 기능적 유연성을 제공하는데 왜냐하면 사용자는 임의의 다른 프로세서와 마찬가지로 임의의 생각가능한 방식으로 원하는 태스크들을 실행하기 위하여 쉐이더들을 수행할 수 있기 때문이다. 하지만, 고정된 기능 유닛들은, 고정된 기능 유닛들이 태스크들을 수행하는 방식을 위해 하드와이어링된다. 따라서, 고정된 기능 유닛들은 많은 기능적인 유연성을 제공하지 못할 수도 있다. 본 개시의 기법들은, GPU 쉐이더 코어들 (22) 상에, 커넬 (20) 과 같은, 커넬의 실행에 관련된다.
CPU (16) 가 그래픽 장면을 렌더링하는 것 또는 커넬을 실행하는 것과 연관된 데이터 및/또는 커맨드를 커맨드 버퍼로 송신하고 나면, GPU (12) 는 GPU (12) 의 파이프라인을 통해 커맨드들의 실행을 시작한다. GPU (12) 의 스케쥴러 (24) 는, 커넬과 연관된 작업의 기본 유닛을 수행하는, 스레드를 생성한다. 스케쥴러 (24) 는 쉐이더 코어들 (22) 의 특정 처리 엘리먼트에 스레드들을 할당한다.
도 2는 본 개시의 기법들에 따라 커넬을 실행할 수도 있는 프로세서의 하나 이상의 쉐이더 코어들의 복수의 처리 엘리먼트들을 예시하는 개념도이다. 도 2는, GPU (12) 또는 CPU (16) 의 부분을 예시한다. GPU (12) 는, 커넬, 예를 들어, 커넬 (20) 의 일 부분을 실행할 수도 있는 복수의 처리 엘리먼트들 (42A-42N) (PE들 (42)) 을 포함한다. 일부 예들에서, PE들 (42) 상에서 실행될 수도 있는 커넬 (20) 의 부분은 "워프" (warp) 또는 "작업 유닛" (work unit) 으로 지칭될 수도 있다. PE들 (42) 은 하나 이상의 쉐이더 코어들 (22) 의 부분일 수도 있다 (도 1). 워프 또는 작업 유닛은, GPU 스케쥴러 (24) 가 실행을 위해 복수의 처리 엘리먼트들, 예를 들어, PE들 (42) 에 할당할 수도 있는, "파이버들" 로도 지칭되는, 스레드들의 그룹을 포함할 수도 있다. 도 2의 각각의 PE 는, 특정 시간에 (예를 들어, 병렬 실행을 위해 동일한 시간에) 다수의 데이터 값들에 대해, 벡터 명령과 같은 단일 명령을 실행할 수 있는, 단일 명령 다중 데이터 (SIMD) 유닛을 포함할 수도 있다. PE들 (42) 은 또한, 단일 부동 소수점 값에 대한 단일 연산과 같은 단일 데이터 값에 대한 단일 명령의 실행을 지원할 수도 있다.
도 2는 또한, GPU (12) 의 스케쥴러가 실행을 위해 PE들 (42) 에 할당하는 명령들 (44) 을 포함한다. 몇몇 예들에서, 명령들 (44) 은 커맨드 버퍼에 저장될 수도 있다. 명령들 (44) 은 각각의 PE 가 실행하도록 구성되는 커넬의 명령들의 세트를 포함할 수도 있다. 프로그램 카운터 (PC) (50) 는 하나 이상의 PE들 (42) 이 실행할 현재 명령을 표시한다. 명령이 PE들 (42) 상에서의 실행을 끝마친 후에, PC (50) 의 값은 커넬 (20) 의 다음 명령의 어드레스로 증분될 수도 있다. 도 2는 또한 레지스터들 (46) 을 포함한다. 레지스터들 (46A-46N) (레지스터들 (46)) 은 다수의 데이터 값들 또는 단일 값을 유지할 수 있는 범용 레지스터들일 수도 있다. 레지스터들 (46) 은 "뱅크화" 될 수도 있다, 즉 특정 PE를 위해 데이터를 로딩 및 저장할 수도 있다. 일 예로서, 레지스터 (46A) 는 PE (42A) 를 위해 데이터를 저장하는 것에 한정될 수도 있고, 다른 PE 들을 위해 데이터를 로딩 또는 저장하지 않을 수도 있다. 레지스터들 (46) 의 각각은 PE (42A) 들 중 하나로 및/또는 하나로부터 데이터를 공급할 수도 있고, 이것을 PE들 (42) 이 다음으로 처리할 수도 있다.
PE들 (42), 명령들 (44), 레지스터들 (46), 캐시 (48), 및 PC (50) 는 GPU (12) 의 쉐이더 코어들 (22) 의 코어 또는 부분을 포함할 수도 있다. 다양한 예들에서, 워프 (40) 는, GPU (12) 의 그래픽스 파이프라인의 부분이거나 또는 커넬 (20) 과 같은 커넬의 부분을 포함할 수도 있는, 지오메트리 쉐이더, 픽셀 쉐이더, 및/또는 버텍스 쉐이더와 같은 쉐이더의 부분을 포함할 수도 있다. 일부 예들에서, GPU (12) 는 워프에 의해 생성되는 결과들을 추가 처리를 위해 파이프라인의 또 다른 스테이지로 피드할 수도 있다.
도 2는 또한 캐시 (48) 를 포함한다. 캐시 (48) 는 실행동안 고속 취출 및 저장을 위해 자주 액세스되는 명령들 및 데이터를 저장하는 작은 메모리이다. 단일 캐시로서 예시되었지만, 캐시 (48) 는 다수의 캐시 레벨들 및/또는 분리된 캐시들을 나타낼 수도 있다. 위에서 설명된 바처럼, 커넬 (20) 의 실행 동안, GPU (12) 는 PC (50) 의 값에 의해 표시되는 어드레스에 위치된 명령들 (44) 중 하나를 취출한다. 다음으로 GPU (12) 는 PE들 (42) 로 하여금, 일부 예들에서 레지스터일 수도 있는 PC (50) 의 어드레스에 저장된 명령을 실행하게 한다.
불필요하게 느린, 시스템 메모리로부터 PC (50) 의 어드레스에서 명령을 가져오기 보다는, GPU (12) 는 캐시 (48) 가 실행될 다음 명령을 현재 포함하는지 결정하기 위하여 캐시 (48) 를 체크한다. 명령들을 저장하는 캐시 (48) 의 부분은 명령 캐시 (“I-캐시”) 로 지칭된다. "캐시 히트" 로 지칭되는, 실행될 다음 명령이 캐시 (48) 에 저장되어 있는 경우, GPU (12) 는 캐시된 명령을 로딩 및 실행한다. "캐시 미스" 로 지칭되는, 실행될 다음 명령이 캐시 (48) 에 저장되어 있지 않은 경우, GPU (12) 는, 약간 더 느린 메모리, 예를 들어, 시스템 메모리 (14) 로부터 실행을 위해 다음 명령을 로딩한다.
(예를 들어, 가산, 승산, 로딩, 저장 등) 메모리 어드레스에 저장된 데이터 값 (예를 들어, 피연산자) 를 필요로 하는 명령의 실행 동안, GPU (12) 는 먼저, 피연산자가 레지스터, 예를 들어, 레지스터들 (46) 중 하나의 레지스터 내에 저장되어 있는지 여부를 결정한다. 요청된 데이터 값이 레지스터들 (46) 에 저장되어 있지 않으면, GPU (12) 는, 데이터 캐시 (“d-캐시”) 로 지칭되는 데이터 값들을 유지하는 캐시 (48) 의 부분으로부터 데이터 값을 액세스하려고 시도한다. 데이터 값이 캐시 (48) 내에 저장되어 있으면, GPU (12) 는 캐시 (48) 로부터 요청된 데이터 값을 로딩한다. 그렇지 않으면, GPU (12) 는 더 느린 메모리, 예를 들어, 시스템 메모리 (14) 로부터 요청된 데이터 값을 로딩해야 한다. 유사하게, 명령이 PE들 (42) 로 하여금 데이터 값을 다시 메모리내에 저장 또는 수정하게 하는 경우, 캐시 (48) 는 그 값을 캐시 (48) 로 저장하여, 그것이 다시 쓰기 또는 읽기되는 경우, 데이터 값이 레지스터들 (46) 중 하나에 저장되어 있지 않을 경우에 데이터 값이 신속하게 캐시로부터 취출되거나 또는 캐시 (48) 로 겹쳐 쓰기되도록 할 수도 있다.
GPU (12) 는 캐시 "라인들" 로 지칭되는, 고정된 크기 블록들에서 캐시 (48) 로 그리고 캐시 (48) 로부터 데이터를 전송한다. 캐시 (48) 는 수백 또는 수천의 상이한 라인들을 저장할 능력을 가질 수도 있다. 각각의 라인은 특정 메모리 어드레스와 연관되고, 다수의 바이트의 데이터를 저장할 수도 있다. 예를 들어, 캐시 (48) 의 각각의 라인은, 일 예로서, 64 바이트의 데이터를 저장할 수도 있다. 각각의 라인에 저장된 바이트들의 수는 캐시 "폭" 으로 지칭된다. 캐시 (48) 가 64 바이트의 데이터를 저장할 수 있는 라인들을 갖는 예에서, 캐시 (48) 의 캐시 폭은 64 바이트이다. 캐시 폭은, 아래에서 더 상세하게 논의될 바처럼, 코드 리오더링 최적화 기법들의 성능에 영향을 줄 수도 있다.
캐시 (48) 로부터 데이터를 취출하는 로드 동작 동안, GPU (12) 는 취출된 캐시 데이터를 하나 이상의 레지스터들 (46), 또는 제시되지 않은 다른 레지스터들로 로딩할 수도 있다. 명령의 실행 동안, PE들 (42) 은 레지스터들 (46) 로부터 하나 이상의 데이터 값들을 읽을 수도 있다. PE들 (42) 은 데이터 값들에 대한 하나 이상의 동작들을 수행하고, 새로운 값들을 다시 레지스터들 (46) 에 저장할 수도 있다. PE들 (42) 은 분기 (branch), 점프, 바로가기 (goto) 등과 같은 흐름 제어 명령들을 실행할 수도 있다. 하지만, 단일 PC (50) 가 있기 때문에, PE들 (42) 은, 주어진 시간에 한가지로 PC (50) 에 의해 표시되는 명령들 (44) 중 하나만을 실행할 수도 있다.
GPU (12) 와 같은 프로세서들은 방대한 양의 벡터 레지스터들 및 벡터 명령들을 가질 수도 있다. 그래서, 벡터화 같은 최적화를 이용하여 애플리케이션들을 컴파일할 수 있는, JIT 컴파일러 (18) 와 같은 컴파일러는, GPU (12) 와 같은 SIMD 아키텍쳐를 갖거나 또는 벡터 명령들을 지원하는 프로세서의 스루풋 또는 실행 성능을 증가시킬 수도 있다.
보다 구체적으로, GPU (12) 는 도 2에 예시된 것들과 유사한 수백 또는 수천의 쉐이더 코어들을 포함할 수도 있다. 각각의 쉐이더 코어는 벡터 명령들을 실행 가능할 수도 있다. 다수의 피연산자들을 갖는 벡터 명령들을 실행하는 것은, 벡터 명령들보다 스칼라 명령들을 포함하는, 비최적화된 코드에 비해 크게 성능을 향상시킬 수도 있다. 게다가, 실행 성능 증가는 벡터 명령들을 실행할 수 있는 더 많은 수의 SIMD 코어들을 갖는 아키텍쳐 상에서 더 클 수도 있는데, 더 일반적인 목적의 프로세서들은 벡터 명령들을 실행할 수 있는 제한된 수의 레지스터들 및/또는 코어들을 가질 수도 있기 때문이다.
도 3a는, 본 개시의 기법들에 따라, 실행될 때, 에일리어싱을 일으킬 수도 있는 코드를 포함하는 커넬 코드를 예시하는 개념도이다. 도 3a 의 예는 커넬 코드 (80) 를 포함한다. 커넬 코드 (80) 는 라인들 (82, 84, 86, 및 88) 을 포함한다.
커넬 코드 (80) 의 라인 (82) 은 compute_output 함수이다. 라인 (82) 의 compute_output 함수는, 커넬이 실행을 시작할 때 타겟 프로세서 (예를 들어, GPU (12)) 가 호출하는 함수이다. compute_output 함수는 드라이버/런타임 (19) 이 커넬 (20) 의 실행을 시작하는데 사용하는 프로그램 엔트리 포인트라는 점에서 C 프로그래밍 언어에서의 “int main()” 함수와 대략 동등하다. 타겟 프로세스 또는 CPU (16) 인 경우, C 런타임 라이브러리는 드라이버/런타임 (19) 의 런타임 컴포넌트를 포함할 수도 있다. GPU (12) 가 타겟 프로세서인 경우, 드라이버/런타임 (19) 의 드라이버 컴포넌트는 런타임을 포함할 수도 있다. compute_ouput 함수는 4개의 입력 아규먼트들: (1) inputImage, (2) global_cdf, (3) outputImage, 및 (4) local_cdf 을 포함한다. inputImage 는 입력 아규먼트들의 버퍼에 대한 포인터이다. outputImage 는, 커넬이 실행을 끝마칠때, 출력 아규먼트들을 포함할 버퍼에 대한 포인터이다. 아규먼트 global_cdf 및 local_cdf 는 값들의 어레이에 대한 포인터들이다. 라인 (84) 은, 실행될 때, GPU (12) 로 하여금 변수들을 할당 및 초기화하게 하는 다수의 스테이트먼트 (statement) 들을 표현할 수도 있다. 일 예로서, 실행 라인 (84) 은 PE들 (42) 로 하여금 inputImage[i] 등의 값들을 초기화 및 로딩하게 할 수도 있다.
라인 (86) 은 루프 초기화 스테이트먼트이다. 루프 초기화 스테이트먼트는 고정된 수의 반복들에 대해 루프가 반복되는 것을 표시한다. 루프는, 변수 “start_offset” 와 동일한, 시작 인덱스, i, 에서 반복을 시작하고, 각각의 반복이 실행을 끝마칠때에 1씩 i 를 증분시킨다. 각각의 루프 반복의 완료시, GPU (12) 는 불리언 조건 “i <final_offset” 이 여전히 참인지 체크해 본다. GPU (12) 는 i 의 값이 값 “final_offset” 이상일 때 루프의 실행을 정지한다.
각각의 루프 반복 내에서, GPU (12) 는 local_cdf[ inputImage[i] ] 의 값과 동일하게 outputImage[i] 로 표기되는, 인덱스 i 에서의 outputImage 의 값을 설정한다. Local_cdf 는, 이 예에서 inputImage[i] 의 값에 의해 인덱싱되는 어레이이다. inputImage[i] 는, 차례로, GPU (12) 가 각각의 루프 반복과 함께 증분하는 변수 i 에 의해 인덱싱된다.
위에 논의된 바처럼, outputImage 및 inputImage 는 양자 모두 메모리 참조들이다. outputImage 및 inputImage 에 대한 포인터들은 메모리에서 동일한 영역을 참조할 수도 있는 것 (즉, outputImage 및 inputImage 이 에일리어싱하거나 또는 부분적으로 에일리어싱하는 것) 이 가능하다. 또한, outputImage 및 inputImage 이 메모리에서 상이한 영역들 또는 오버랩하는 영역들을 참조할 수도 있는 것 (즉, outputImage 및 inputImage 이 에일리어싱하지 않는 것) 이 가능하다. JIT 컴파일러 (18) 는 outputImage 및 inputImage 이 에일리어싱하지 않는지 (즉, 정확하게 동일한 메모리 영역을 참조하지 않는지) 여부를 결정할 수 없으면, 컴파일러는 벡터화, 코드 리오더링, 및/또는 루프 언롤링과 같은 어떤 컴파일러 최적화들을 사용가능하지 않을 수도 있다.
도 3b는 본 개시의 기법들에 따라 구성된 드라이버/런타임이 검출가능할 수도 있는 에일리어싱의 예를 예시하는 개념도이다. 하지만, 컴파일러는 에일리어싱을 위해 최적화 가능하지 않을 수도 있다. 도 3b 의 예는 GPU (12) 가 메모리에 저장할 수도 있는 버퍼 (100) 를 예시한다. 예들의 목적을 위해, 도 3b 로부터의 포인터들, outputImage 및 inputImage 은 버퍼 (100) 의 부분을 참조할 수도 있다. 도 3b 의 예에서, 버퍼 (100) 는 메모리 어드레스 0x800 (16진수) 에서 시작한다.
이 예에서, inputImage 및 outputImage 양자 모두는 버퍼 (100) 내에 저장된 단일 엔트리 (예를 들어, 단일 오브젝트, 변수 등) 을 참조한다. 즉, 이 예에서 inputImage 및 outputImage 은, 크로스 해싱으로 표시되는 정확하게 동일한 메모리 영역에 대해 에일리어싱한다. 드라이버/런타임 (19) 은 inputImage 및 outputImage 가 동일한 메모리 영역을 참조한다는 것을 검출 가능할 수도 있다. inputImage 및 outputImage 가 동일한 메모리 영역을 참조하기 때문에, JIT 컴파일러 (18) 는 루프 언롤링 및/또는 벡터화와 같은 최적화를 수행할 수 없다.
2개의 메모리 참조들이 도 3b 에 예시된 바처럼 동일한 메모리 영역을 참조하는 것을 검출하는 것에 응답하여, 드라이버/런타임 (19) 은 어느 메타데이터도 생성하지 못할 수도 있다. 추가적으로, JIT 컴파일러 (18) 가 도 3c-3d 에 예시된 경우들에 대해 그러할 수도 있는 바처럼, JIT 컴파일러 (18) 는 커넬 (20) 을 리컴파일하지 못할 수도 있다. 그러므로, JIT 컴파일러 (18) 는 도 3c-3d 에도 예시된 코드 최적화들 중 어느 것도 수행하지 못할 수도 있다.
도 3c는 본 개시의 기법들에 따라 구성된 드라이버/런타임이 검출가능할 수도 있는 오버랩되지 않는 메모리 참조들의 예를 예시하는 개념도이다. 도 3c 는 도 3b 에 예시된 것과 동일한 버퍼인 버퍼 (120) 를 예시한다. 버퍼 (120) 는 유사하게, 도 3b 의 버퍼 (100) 와 동일한 메모리 어드레스, 0x800 에서 시작한다.
도 3c 에서, inputImage 및 outputImage 는 버퍼 (120) 의 2개의 상이한 메모리 영역들을 참조하는 메모리 참조들이다. inputImage 이 참조하는 메모리 영역은 수평 해싱에 의해 표시된다. outputImage 이 참조하는 메모리 영역은 수직 해싱에 의해 표시된다. 커넬 코드 (80), 그리고 보다 구체적으로, 라인들 (86 및 88) 의 실행 전에, JIT 컴파일러 (18) 는, i 의 값에 상관 없이, inputImage[i] 및 outputImage[i] 가 루프의 동일한 반복 동안 동일한 메모리 영역을 참조하지 않을 것이라고 결정할 수도 있다.
런타임 동안, 드라이버/런타임 (19) 은, inputImage[i] 및 outputImage[i] 의 초기 값들에 기초하여 그리고 inputImage[i] 및 outputImage[i] 의 메모리 어드레스들이 루프 (86) 를 통한 반복 동안 수렴하지 않지 않는다는 사실에 기초하여 inputImage[i] 및 outputImage[i] 가 동일한 메모리 영역을 참조하지 않는다고 결정가능할 수도 있다. 다른 말로, inputImage 및 outputImage 의 참조된 인덱스는, GPU (12) 가 단조적으로 증가시키는 동일한 인덱스 값 i 에 의해 항상 참조된다.
메모리 참조들 inputImage 및 outputImage 가 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여, 드라이버는 inputImage 과 outputImage 사이의 관계를 나타내는 메타데이터를 생성할 수도 있다. 메타데이터는 예로서, inputImage 및 outputImage 와 연관된 메모리 영역들이 오버랩되지 않고, 2개의 엔트리들만큼 분리된다는 것을 표시할 수도 있다. 메타데이터는 또한, inputImage 및 outputImage 와 연관된 영역들의 크기, 그리고 inputImage 와 outputImage 사이의 바이트들의 수를 표시할 수도 있다. 메타데이터를 생성한 후에, JIT 컴파일러 (18) 는, 드라이버/런타임 (19) 으로부터 메타데이터를 수신하고, 아래에서 더 상세하게 설명된 바처럼, 다양한 최적화들을 적용하는 것에 의해 메타데이터에 기초하여 커넬 (20) 을 리컴파일할 수도 있다.
도 3d는 본 개시의 기법들에 따라 구성된 드라이버/런타임이 검출할 수도 있는 오버랩되는 메모리 참조들을 예시하는 개념도이다. 도 3d 는 아규먼트 버퍼, 예를 들어, 아규먼트들 (26) (도 1) 일 수도 있는 버퍼 (130) 를 포함한다. 버퍼 (130) 는 이 예에서 어드레스 0x800 에서 시작한다. 버퍼 (130) 는, 버퍼 (130) 의 둘러싸는 직사각형 내의 분리된 직사각형들로서 예시된 다수의 데이터 값들을 포함한다.
이전의 예들에서 처럼, inputImage 및 outputImage 는 버퍼 (130) 의 영역들을 참조하는 메모리 참조들이다. 이 예에서, inputImage 및 outputImage 가 참조하는 영역들은 오버랩되지만, 전체적으로는 아니다. inputImage 에만 연관되는 메모리 영역은 수평으로 해싱된 직사각형들로 표시된다. outputImage 에만 연관되는 메모리 영역은 수직으로 해싱된 직사각형들로 표시된다. inputImage 및 outputImage 양자 모두에 의해 참조되는 오버랩되는 메모리 영역은 크로스 해칭된 직사각형들로 표시된다.
런타임시, 드라이버는 inputImage 및 outputImage 메모리 참조들이 동일한 메모리 영역을 참조하는지 여부를 결정한다. 이 예에서, inputImage 및 outputImage 가 오버랩되지만, 동일한 메모리 영역을 참조하지 않는다. 드라이버/런타임 (19) 은, inputImage 및 outputImage 가 오버랩되지만, 동일하지 않다는 것을 검출하고 JIT 컴파일러 (18) 를 위한 메타데이터를 생성한다. 메타데이터는, 각각의 영역의 시작 및 종료 어드레스와 같은 inputImage 및 outputImage 에 연관된 영역들에 관한 정보를 표시할 수도 있다. 메타데이터는 또한, 오버랩 영역의 크기, 및 오버랩 영역의 시작 및/또는 종료 어드레스들과 같은 오버랩 영역에 관한 정보를 포함할 수도 있다. JIT 컴파일러 (18) 는 드라이버/런타임 (19) 에 의해 생성된 메타데이터를 수신하고, 본 개시에 따른 최적화 기법들을 적용하는 것에 의해 커넬 (20) 을 리컴파일할 수도 있다.
도 4a는 본 개시의 기법들에 따른 루프 언롤링을 예시하는 개념도이다. 도 4a는 도 3a 에 예시된 커넬 코드 (80) 에 일반적으로 대응하는 코드 섹션 (140) 을 포함한다. 도 4a 의 예에서, 드라이버/런타임 (19) 및/또는 JIT 컴파일러 (18) 는, 도 3c 및 도 3d 에 예시된 바처럼, 메모리 참조들 inputImage 및 outputImage 이 동일한 메모리 영역을 참조하지 않는다는 것을 결정했을 수도 있다. inputImage 및 outputImage 가 동일한 메모리 영역을 참조하지 않기 때문에, JIT 컴파일러 (18) 는 커넬 코드 (80) 에 대해 루프 언롤링을 수행했다. 라인들 (142-150) 은 1개의 반복을 4개의 반복들로 언롤링하는 결과를 예시한다.
도 3a 의 라인들 (86 및 88) 은 단일 반복을 수행하는 것을 예시하고 각각의 반복 후에 1씩 변수 i 를 증분시키는 반면에, 라인 (142) 의 언롤링된 루프는 각각의 반복 후에 4씩 i 를 증분시킨다. 라인 (144) 은 local_cdf[inputImage[i]] 의 값을 outputImage[i] 에 할당한다. 라인 (146) 은 local_cdf[inputImage[i+1]] 의 값을 outputImage[i+1] 에 할당한다. 라인 (148) 은 local_cdf[inputImage[i+2]] 의 값을 outputImage[i+2] 에 할당하고, 라인 (150) 은 local_cdf[inputImage[i+3]] 의 값을 outputImage[i+3] 에 할당한다. 라인들 (144-150) 의 결과는 local_cdf[inputImage[i+x]] 의 출력을 outputImage[i+x] 의 대응하는 값에 할당하는 것이고, 여기서 x [0…3] 이다. 따라서, 실행될 때, 라인들 (142-150) 에 예시된 언롤링된 루프 코드 섹션은 도 3a 의 라인들 (86-88) 의 4개의 반복들과 동일한 효과를 가진다.
코드 섹션 (140) 의 루프 언롤링은 도 3a 의 루프 코드 섹션 (80) 에 비해 여러 혜택들을 가질 수도 있다. 첫번째 이점은 할당들의 각각을 차례 차례로 오더링함으로써, JIT 컴파일러 (18) 및/또는 드라이버/런타임 (19) 은 오더링되지 않은 코드 섹션에 비해 타겟 프로세서, 예를 들어, GPU (12) 에 대한 보다 나은 캐시 성능을 달성가능할 수도 있다는 것이다.
예를 들어, 라인 (144) 을 실행한 후에, GPU (12) 는 캐시, 예를 들어, 캐시 (48) 에 inputImage 및 outputImage 와 연관된 메모리 영역들의 데이터의 일부 또는 전부를 저장했을 수도 있다. 명령들을 수행하기 위해 필요한 데이터가 레지스터들, 예를 들어, 레지스터들 (46) 에 저장되어 있지 않으면, 데이터는 캐시, 예를 들어, 캐시 (48) 로부터 액세스될 필요가 있을 수도 있다. 보다 구체적으로, GPU (12) 는 inputImage 및 outputImage 의 엔트리들, 예를 들어, inputImage[i+1], [i+2] 등, 그리고 outputImage [i+1], [i+2] 등을 캐시 (48) 에 저장할 수도 있다. inputImage 및 outputImage 의 엔트리들이 GPU (12) 의 캐시에 저장되면, GPU (12) 는 보다 느린 메모리로부터 참조되는 인덱스들에 액세스하는 것과 반대로 캐시로부터 라인들 (144-150) 의 inputImage 및 outputImage 의 참조되는 인덱스들의 데이터를 신속하게 액세스 가능할 수도 있다.
추가적으로, 코드 섹션 (140) 이 언롤링될 때, inputImage[i, i+1, i+2…], 및 outputImage[i, i+1, etc.] 의 값들은 단일 캐시 라인에 저장될 수도 있다. 대조적으로 언롤링되지 않을 때, inputImage 및 outputImage[i] 의 값들은 상이한 캐시 라인들에 저장될 수도 있다. 루프 언롤링으로부터 비롯될 수도 있는 단일 캐시 읽기에서 단일 캐시 라인으로부터 inputImage 의 값들을 모두 취출하는 것은, 언롤링된 코드를 실행할 때 비롯될 수도 있는 다수의 캐시 읽기들을 수행하는 것에 비해 더 빠를 수도 있다.
보다 느린 시스템 메모리, 예를 들어, 시스템 메모리 (14) 로부터 데이터에 액세스하는 것과 반대로 GPU (12) 의 캐시로부터 데이터에 액세스하는 것은, 라인들 (86-88) 에 비해 라인들 (142-150) 의 루프를 실행하는 성능을 증가시킬 수도 있다. 일부 예들에서, GPU (12) 는 또한, 예를 들어, 수퍼스칼라 실행 (superscalar execution) 을 지원하는 프로세서, 또는 라인들 (144-150) 간에 의존성이 없다고 가정하는 SIMD 프로세서 상에서, 병렬로 라인들 (144-150) 을 실행 가능할 수도 있고, 여기서 inputImage 또는 outputImage 의 값은 커넬 (20) 에서 이전에 계산된 값에 의존한다.
캐시 성능을 향상시키는 것에 더하여, 도 4a 의 코드 섹션 (140) 에 예시된 루프 언롤링은 또한, GPU (12) 가 루프와 연관된 불리언 조건을 평가하는 횟수, 그리고 GPU (12) 가 각각의 루프 반복을 끝마친 후에 실행하는 점프들의 횟수를 감소시킨다. 도 3a 의 코드 섹션 (80) 과 비교하여, 라인들 (142-150) 의 코드는 라인 (142) 의 불리언 조건, “i < final_offset” 가 참인지 여부를 평가하기 전에 반복마다 4개의 라인들을 실행한다. 그에 반해서, 코드 섹션 (80) 은 라인 (82) 의 불리언 조건이 참인지 여부를 평가하기 전에 하나의 라인만을 실행한다. 따라서, GPU (12) 가 라인 (142) 의 불리언 조건을 평가하는 횟수는 코드 섹션 (80) 에 비해 감소된다.
GPU (12) 가 라인들 (142-150) 의 루프의 반복을 종료한 후에, 그리고 GPU (12) 가 불리언 조건, “i < final_offset” 이 여전히 참이라고 결정하면, GPU (12) 는 라인 (150) 으로부터 다시 라인 (144) 으로 점프한다. 코드 섹션 (140) 에서, GPU (12) 는 4개의 라인들을 실행한 후에 점프를 수행한다. 코드 섹션 (80) 을 실행할 때, GPU (12) 는 각각의 반복 후에 점프한다. 따라서, 코드 섹션 (80) 에 비해, 코드 섹션 (140) 의 언롤링된 코드는, GPU (12) 가 수행하는 점프들의 횟수 및 불리언 조건의 평가 양자 모두를 감소시키고, 이는 코드 섹션 (140) 을 실행하는 실행 성능을 향상시킬 수도 있다.
도 4b는 본 개시의 기법들에 따른 코드 리오더링을 예시하는 개념도이다. 도 4b는 라인들 (162, 164, 166, 168, 170, 및 172) 을 더 포함하는, 코드 섹션 (160) 을 포함한다. 위에 논의된 바처럼, 드라이버/런타임 (19) 및/또는 JIT 컴파일러 (18) 는 메모리에 대한 참조들이 메모리의 동일한 영역에 대해 에일리어싱하는지 여부를 결정할 수도 있다. 도 4a 를 참조하여 위에서 논의된 바처럼, JIT 컴파일러 (18) 는, 특정 코드 섹션에서 메모리 에일리어싱이 없다고 결정하는 드라이버/런타임 (19) 으로부터의 메타데이터를 수신하는 것에 응답하여, 도 4a 에 예시된 루프 언롤링과 같은 어떤 최적화들을 수행할 수도 있다.
JIT 컴파일러 (18) 및/또는 드라이버/컴파일러 (19) 가 특정 코드 섹션에서의 메모리 참조들이 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여 수행할 수도 있는 또 다른 최적화는, 도 4b 가 예시하는 코드 리오더링이다. 코드 (160) 는 일반적으로, 도 4b 의 언롤링된 코드의 리오더링된 어셈블리 언어 표현에 대응할 수도 있다. JIT 컴파일러 (18) 및/또는 드라이버/런타임 (19) 은 비루프 그리고 루프 코드 섹션들에 코드 리오더링을 적용할 수도 있다. 도 4b 에서, JIT 컴파일러 (18) 는, 모든 로드들 및 저장들이 함께 그룹화되도록 도 4a 의 로드들 및 저장들을 리오더링하였다.
라인들 (162 및 164) 은, JIT 컴파일러 (18) 및/또는 드라이버/컴파일러 (19) 가 함께 그룹화한 로드 명령들이다. 도 4a 에서, 라인, 이를테면 라인 (144) 은 다수의 로드 및 저장 명령들을 포함한다. 예를 들어, 라인 (144) 을 실행하기 위하여, JIT 컴파일러 (18) 는 3개의 별개의 명령들을 생성할 수도 있다. 첫번째 명령은, inputImage[i] 가 참조하는 메모리 장소로부터 r0 로 표기되는 레지스터로 값을 로딩하는 로드 명령일 수도 있다. 두번째 명령은, local_cdf[inputImage[i]] 의 값을 로딩하고 그 로딩된 값을 동일한 레지스터 r0 로 저장함으로써, r0 의 이전 값을 겹쳐 쓰기하는 로드 명령일 수도 있다. 라인 (144) 에 포함된 최종 명령은, r0 으로부터 outputImage[i] 가 참조하는 메모리로 값을 저장하는 저장 명령일 수도 있다.
라인들 (162-172) 은 라인들 (144-150) 을 포함하는 명령들에 관하여 리오더링된 로드 및 저장 명령들을 예시한다. 라인 (162) 에서, 어셈블리 코드는, inputImage[i] 가 참조하는 메모리 영역으로부터 레지스터 r0 으로 값을 (로드 명령, “ldg” 을 이용하여) 로딩하도록 GPU (12) 에 명령한다. 유사하게, 라인 (164) 은, GPU (12) 로 하여금, 메모리 참조 inputImage[i+1] 가 참조하는 값을 레지스터 r1 로 로딩하게 한다. 라인들 (162 및 164) 후에, 하지만 라인 (166) 전에 일어날 수도 있고, 간결성을 위해 예시되지 않은 후속 명령들은, GPU (12) 로 하여금 inputImage 가 참조하는 메모리 영역들로부터 레지스터들로 데이터를 로딩하게 하는 추가적인 로딩 명령들을 포함할 수도 있다.
라인들 (166, 168) 그리고 간결성을 위해 예시되지 않은 다른 라인들에서, JIT 컴파일러 (18) 는 버퍼 local_cdf 로부터 로드들을 함께 그룹화했다. 라인 (166) 은 local_cdf [r0] 의 콘텐츠들, 즉 인덱스 r0 에서의 어레이 local_cdf 로부터 메모리의 콘텐츠들을 로딩하는 로드 명령을 포함하고, local_cdf[r0] 의 콘텐츠들을 레지스터 r0 로 저장함으로써, r0 의 콘텐츠들을 겹쳐 쓰기한다. 유사하게, 라인 (168) 의 명령은, GPU (12) 로 하여금, 레지스터 r1 에 현재 저장된 값에 의해 표시되는 인덱스에서 메모리 참조 local_cdf 가 참조하는 콘텐츠들을 레지스터 r1 내에 저장하게 한다. 따라서, 168 의 명령의 실행시, GPU (12) 로 하여금 r1 의 이전 값을 겹쳐 쓰기하게 한다. 라인 (168) 후에 그리고 라인 (170) 전에 일어나고, 간결성을 위해 예시되지 않은 다른 명령들은 유사하게, 실행될 때, GPU (12) 로 하여금, local_cdf[rx] 로부터 데이터를 로딩하게 하는 명령들을 포함할 수도 있고, 여기서 x 는 어떤 정수이다.
코드 섹션 (140) 의 명령들을 리오더링하는 부분으로서, JIT 컴파일러 (18) 는 또한 저장 명령들을 함께 그룹화한다. 이것의 일 예로서, 리오더링 후에, JIT 컴파일러 (18) 는 라인들 (170 및 172) 을 함께 그룹화했다. 라인 (170) 은 장소 outputImage[i] 에서의 메모리로 r0 의 콘텐츠들을 저장하는 저장 명령을 포함한다. 유사하게, 라인 (172) 은, 실행될 때, GPU (12) 로 하여금, outputImage[i+1] 이 참조하는 장소에서의 메모리로 레지스터 r1 의 값을 저장하게 한다. 간결성을 위해 예시되지 않은 다른 명령들은, 실행될 때, 유사하게, GPU (12) 로 하여금, 레지스터, 예를 들어 레지스터 rx 의 값을 장소 outputImage[i+x] 에서의 메모리로 저장하게 할 수도 있고, 여기서 x 는 정수이다.
로드들 및 저장들을 리오더링하는 것은 도 3a 의 코드 (80) 에 비해 코드 (160) 를 실행하는 성능을 향상시킬 수도 있다. 보다 구체적으로, 로드들 및 저장들을 리오더링하는 것은, 캐시 라인 폭에 따라 어떤 경우들에서 성능을 향상시킬 수도 있다. 예를 들어, 코드 리오더링은 다수의 로드 명령들을 합칠 때 실행 성능을 향상시킬 수도 있으며, 이와 함께, 스칼라 명령에서 사용되는 피연산자의 크기의 배수인 캐시 라인 폭을 갖는 시스템 상에서 성능을 향상시킬 수도 있다.
도 4c는 본 개시의 기법들에 따른 코드 벡터화를 예시하는 개념도이다. 도 4c는 라인들 (182, 184, 및 186) 을 더 포함하는, 코드 섹션 (180) 을 포함한다. 위에서 논의된 바처럼, JIT 컴파일러 (18) 및/또는 컴파일러/드라이버 (19) 는, 코드 섹션에 있는 메모리 참조들이 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여, 도 4a 에 예시된 루프 언롤링과 같은 어떤 최적화들을 수행할 수도 있다. JIT 컴파일러 (18) 및/또는 컴파일러/드라이버 (19) 는, 루프 코드 섹션을, 그 코드 섹션의 메모리 참조들에 관한 정보를 포함하는 드라이버/런타임 (19) 으로부터의 메타데이터에 기초하여, 벡터화하도록 구성된다.
벡터화는, 컴파일러 (예를 들어, JIT 컴파일러 (18)) 및/또는 드라이버/런타임 (19) 이, 각각 단일 피연산자를 갖는 다수의 스칼라 명령들을 다수의 피연산자들을 갖는 단일 벡터 명령으로 합성하는 프로세스이다. 벡터화는, 특정 코드 섹션을 완료하기 위해 프로세서가 실행할 것이 요구되는 명령들의 수를 감소시킴으로써, 그리고 시스템 메모리 (14) 와 GPU (12) 사이의 데이터를 이동시킬 고유 하드웨어 능력을 이용함으로써 실행 성능을 향상시키는 병렬화의 일 형태이다. 도 4c 의 코드 섹션 (180) 의 예에서, JIT 컴파일러 (18) 는 도 4b 에 예시된 바처럼 로드들 및 저장들을 리오더링할 수도 있다. JIT 컴파일러 (18) 가 로드들 및 저장들을 리오더링하고 나면, JIT 컴파일러 (18) 는 다음으로 도 4c에 예시된 바처럼, 유사한 명령들의 그룹들을 벡터화할 수도 있다.
라인 (182) 에서, JIT 컴파일러 (18) 는 다수의 로드 (ldg) 명령들을 단일 벡터화된 명령으로 합성하였다. 실행될 때, 벡터화된 명령은 인덱스들 [i]-[i+3] 에서의 inputImage 를 레지스터들 r0-r3 로 로딩한다. 유사하게, 라인 (184) 에서, JIT 컴파일러 (18) 는, local_cdf[r0-r3] 의 값들을 레지스터들 r0-r3 로 로딩하는 단일 벡터화 로드 명령으로 라인 (166, 168) 등의 다수의 로드 명령들을 합성한다. 또한, 라인 (186) 에서, JIT 컴파일러 (18) 는, 레지스터들 (r0-r3) 의 값들을 outputImage[i]-outputImage[i+3] 로 저장하는 단일 벡터화된 저장 명령으로 라인들 (170-172) 의 저장들 (“stg” 명령들) 을 합성했다.
도 4b 및 도 4c 에 예시된 명령들을 리오더링 또는 벡터화하기 위하여, JIT 컴파일러 (18) 및/또는 드라이버/런타임 (19) 은 임의의 의존성을 준수해야 한다. 의존성은 스테이트먼트 또는 명령들 사이의 실행 오더 제약들을 만드는 관계이다. 예로서, S1 이 S2 전에 실행되야 한다면, 다른 스테이트먼트 S1 에 관하여 스테이트먼트 S2 에 대한 의존성이 있다. 의존성들이 벡터화 및/또는 코드 리오더링을 금하는지 여부를 결정하기 위하여, JIT 컴파일러 (18) 및/또는 드라이버/런타임 (19) 은 드라이버/런타임 (19) 으로부터 획득된 메타데이터에 기초하여 본 개시의 기법들에 따라 코드를 리오더링 또는 벡터화하기 전에 의존성 분석을 수행할 수도 있다.
도 5는 본 개시의 기법들에 따른 컴파일러 최적화를 돕기 위한 컴파일러 메타데이터를 생성하는 예시적인 방법을 예시하는 흐름도이다. 일반적으로, 도 6의 방법은 컴파일링 프로세서, 예를 들어, CPU (16), 및 타겟 프로세서, 예를 들어, GPU (12) 상에서 실행되는 JIT 컴파일러 (18) 및 드라이버/런타임 (19) 으로 이루어지는 군의 적어도 하나에 의해 수행될 수도 있다는 것이 이해되야 한다. 일부 예들에서, 타겟 프로세서 및 컴파일링 프로세서는 동일할 수도 있다. 추가적으로, 하나보다 많은 컴파일링 프로세서들 및/또는 타겟 프로세서가 있을 수도 있다.
도 5의 방법에서, 컴파일링 프로세서, 예를 들어, CPU (16) 는 드라이버/런타임 (19) 및/또는 JIT 컴파일러 (18) 를 사용하여 컴파일된 커넬 (20) 의 2진 코드 또는 바이트코드를 실행하기 위해 아규먼트들 (예를 들어, 커넬 아규먼트들 (26)) 을 생성한다 (200), 드라이버/런타임 (19) 및/또는 JIT 컴파일러 (18) 는 또한, 커넬 아규먼트들 (26) 의 제 1 메모리 영역에 대한 제 1 메모리 참조 및 커넬 아규먼트들 (26) 의 제 2 메모리 영역에 대한 제 2 메모리 참조가 커넬 아규먼트들 (26) 의 동일한 메모리 영역을 참조하는지 (202), 또는 도 3b, 3c 및 3d 에 예시된 가능한 관계들의 다른 예들인지를 결정한다.
CPU (16) 는 드라이버/런타임 (19) 및/또는 JIT 컴파일러 (18) 를 사용하여 제 1 메모리 참조 및 제 2 메모리 참조와 연관된 메타데이터를 생성한다 (204). 메타데이터는, 제 1 메모리 영역과 제 2 메모리 영역 사이의 오버랩 영역과 같은 제 1 메모리 영역과 제 2 메모리 영역 사이의 관계를 표시한다. 메타데이터는 제 1 메모리 영역과 제 2 메모리 영역 사이의 오버랩의 바이트들의 수를 더 포함할 수도 있다. 일부 예들에서, 메타데이터는 메모리 오버랩의 시작 어드레스 및 메모리 오버랩 영역의 종료 어드레스를 포함할 수도 있다. 도 5를 참조하여 설명된 예들은 오직 예들의 목적을 위해 단일 쌍의 메모리 참조들을 언급한다는 것이 이해되야 한다. 드라이버/런타임 (19) 및/또는 JIT 컴파일러 (18) 는 커넬 아규먼트들 (26) 의 모든 쌍의 메모리 참조들을 위한 메타데이터를 도출할 수도 있다.
드라이버/런타임 (19) 을 이용하여, 제 1 메모리 참조 및 제 2 메모리 참조가 커넬 아규먼트들 (26) 의 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여, CPU (16) 상에서 실행되는 JIT 컴파일러 (18) 는 CPU (16) 로 하여금 메타데이터에 기초하여 JIT 컴파일러 (18) 를 이용하여 커넬 (20) 을 리컴파일하게 한다 (206). 마지막으로, 타겟 프로세서, 예를 들어, GPU (12) 는 리컴파일된 커넬을 실행할 수도 있다 (210). 일부 예들에서, 드라이버/런타임 (19) 및/또는 JIT 컴파일러 (18) 는, 메타데이터에 기초하여 제 1 및 제 2 메모리 참조들이 동일한 메모리 영역을 참조하지 않는다고 결정할 수도 있고, 이 정보를 이용하여 최적화로 커넬 (20) 을 리컴파일할 수도 있다.
일부 추가 예들에서, 커넬 (20) 의 제 1 메모리 참조 및 제 2 메모리 참조가 동일한 메모리 영역을 참조하는지 여부를 결정하기 위하여, CPU (16) 는, 제 1 및 제 2 메모리 참조들을 포함하는 커넬 (20) 의 루프 코드 섹션을 결정하기 위하여 드라이버/런타임 (19) 을 이용할 수도 있다. 그리고, 커넬을 리컴파일하기 위하여, JIT 컴파일러 (18) 는 드라이버/런타임 (19) 및/또는 JIT 컴파일러 (18) 에 의해 생성된 메타데이터에 기초하여 루프 코드 섹션을 언롤링할 수도 있다. 커넬을 리컴파일하기 위하여, JIT 컴파일러 (18) 는 또한, 생성된 메타데이터에 기초하여, 루프 코드 섹션의 저장 동작 및 저장 동작 및 로드 동작 중 적어도 하나를 리오더링하거나 또는 적어도 하나의 벡터 명령으로 루프 코드 섹션의 복수의 스칼라 명령들을 벡터화할 수도 있다. 다양한 예들에서, JIT 컴파일러 (18) 는 Khronos Group 에 의한 Microsoft DirectCompute 및/또는 OpenCL 과 같은 이종 프레임워크를 이용하여 커넬 (20) 을 리컴파일할 수도 있다.
본 개시에 설명된 기법들은 적어도 부분적으로, 하드웨어, 소프트웨어, 펌웨어 또는 이들의 임의의 조합으로 구현될 수도 있다. 예를 들어, 설명된 기법들의 다양한 양태들은, 하나 이상의 마이크로프로세서들, DSP (digital signal processor), ASIC (application specific integrated circuit), FPGA (field programmable gate array), 또는 임의의 다른 동등한 집적 또는 이산 로직 회로, 그리고 그러한 컴포넌트들의 임의의 조합들을 포함한, 하나 이상의 프로세서들내에서 구현될 수도 있다. 용어 "프로세서" 또는 "처리 회로" 는 일반적으로, 단독으로 또는 다른 로직 회로와 조합한, 전술한 로직 회로, 또는 처리를 수행하는 이산 하드웨어와 같은 임의의 다른 동등한 회로 중 임의의 것을 나타낼 수도 있다.
그러한 하드웨어, 소프트웨어, 및 펌웨어는, 본 개시에 설명된 다양한 동작들 및 기능들을 지원하기 위하여 동일한 디바이스 내에서 또는 별개의 디바이스들 내에서 구현될 수도 있다. 또한, 설명된 유닛들, 모듈들 또는 컴포넌트들 중 임의의 것이 함께 또는 따로따로 이산이지만 연동가능한 로직 디바이스들로서 구현될 수도 있다. 모듈들 또는 유닛들로서 상이한 특징들의 묘사는 상이한 기능적 양태들을 강조하기 위해 의도되고 그러한 모듈들 또는 유닛들이 분리된 하드웨어 또는 소프트웨어 컴포넌트들에 의해 실현되야 한다는 것을 반드시 의미하지는 않는다. 오히려, 하나 이상의 모듈들 또는 유닛들과 연관된 기능성은 분리된 하드웨어, 펌웨어 및/또는 소프트웨어 컴포넌트들에 의해 수행될 수도 있거나, 또는 공통 또는 분리된 하드웨어 또는 소프트웨어 컴포넌트들 내에 통합될 수도 있다.
본 개시에 설명된 기법들은 또한, 명령들을 저장하는 컴퓨터 판독가능 저장 매체와 같은 컴퓨터 판독가능 매체에 저장, 수록 또는 인코딩될 수도 있다. 컴퓨터 판독가능 매체에 수록 또는 인코딩된 명령들은 하나 이상의 프로세서들로 하여금, 예를 들어, 명령들이 하나 이상의 프로세서들에 의해 실행될 때, 본원에 설명된 기법들을 수행하게 할 수도 있다. 컴퓨터 판독가능 저장 매체는, RAM (random access memory), ROM (read only memory), PROM (programmable read only memory), EPROM (erasable programmable read only memory), EEPROM (electronically erasable programmable read only memory), 플래시 메모리, 하드 디스크, CD-ROM, 플로피 디스크, 카세트, 자기 매체, 광학 매체 또는 유형의 다른 컴퓨터 판독가능 저장 매체를 포함할 수도 있다.
컴퓨터 판독가능 매체는, 위에 열거된 것들과 같은 유형의 저장 매체에 대응하는, 컴퓨터 판독가능 저장 매체를 포함할 수도 있다. 컴퓨터 판독가능 매체는 또한, 예를 들어, 통신 프로토콜에 따라, 하나의 장소로부터 다른 장소로의 컴퓨터 프로그램의 전송을 가능하게 하는 임의의 매체를 포함한 통신 매체를 포함할 수도 있다. 이런 방식으로, 어구 컴퓨터 판독가능 매체는 일반적으로, (1) 비일시적인 유형의 컴퓨터 판독가능 저장 매체 또는 (2) 일시적인 신호 또는 캐리어 파와 같은 무형의 컴퓨터 판독가능 통신 매체에 대응할 수도 있다.
다양한 양태들 및 예들이 설명되었다. 하지만, 다음의 청구항들의 범위로부터 벗어남이 없이 본 개시의 구조 또는 기법들에 대해 변경이 이루어질 수 있다.

Claims (30)

  1. 실행을 위해 커넬들을 컴파일하는 방법으로서,
    컴파일링 프로세서 상에서 실행되는 컴파일러 및 런타임으로 이루어지는 군의 적어도 하나에 의해, 컴파일된 커넬의 2진 코드를 실행하기 위해 아규먼트들을 생성하는 단계;
    상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 커넬 아규먼트들의 제 1 메모리 영역에 대한 제 1 메모리 참조 및 커넬 아규먼트들의 제 2 메모리 영역에 대한 제 2 메모리 참조가 동일한 메모리 영역을 참조하는지 여부를 결정하는 단계;
    상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 상기 결정에 기초하여 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조와 연관된 메타데이터를 생성하는 단계로서, 상기 메타데이터는 상기 제 1 메모리 영역과 상기 제 2 메모리 영역 사이의 관계를 나타내는, 상기 메타데이터를 생성하는 단계; 및
    상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 상기 커넬의 제 1 및 제 2 메모리 참조들이 상기 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여:
    상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 컴파일러로 하여금 상기 메타데이터에 기초하여 상기 커넬을 리컴파일하게 하는 단계; 및
    상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 리컴파일된 상기 커넬을 실행하도록 타겟 프로세서에 명령하는 단계
    를 포함하는, 커넬들을 컴파일하는 방법.
  2. 제 1 항에 있어서,
    상기 커넬의 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조가 상기 동일한 메모리 영역을 참조하는지 여부를 결정하는 단계는, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조를 포함하는 상기 커넬의 루프 코드 섹션을 결정하는 단계를 더 포함하고,
    상기 커넬을 리컴파일하는 단계는, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 생성된 상기 메타데이터에 기초하여 상기 루프 코드 섹션을 언롤링하는 단계, 및 언롤링된 상기 루프 코드 섹션을 컴파일하는 단계를 포함하는, 커넬들을 컴파일하는 방법.
  3. 제 1 항에 있어서,
    상기 커넬의 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조가 상기 동일한 메모리 영역을 참조하는지 여부를 결정하는 단계는, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조를 포함하는 상기 커넬의 코드 섹션을 결정하는 단계를 더 포함하고,
    상기 커넬을 리컴파일하는 단계는, 상기 코드 섹션의 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조가 상기 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 생성된 상기 메타데이터에 기초하여 상기 코드 섹션의 로드 동작 및 저장 동작 중 적어도 하나를 리오더링하는 단계를 더 포함하는, 커넬들을 컴파일하는 방법.
  4. 제 1 항에 있어서,
    상기 커넬의 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조가 상기 동일한 메모리 영역을 참조하는지 여부를 결정하는 단계는, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조를 포함하는 상기 커넬의 코드 섹션을 결정하는 단계를 더 포함하고,
    상기 커넬을 리컴파일하는 단계는, 상기 코드 섹션의 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조가 상기 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 생성된 상기 메타데이터에 기초하여 상기 코드 섹션의 복수의 스칼라 명령들을 적어도 하나의 벡터 명령으로 벡터화하는 단계를 더 포함하는, 커넬들을 컴파일하는 방법.
  5. 제 1 항에 있어서,
    상기 메타데이터는 또한, 상기 제 1 메모리 영역과 상기 제 2 메모리 영역 사이의 오버랩 영역을 나타내는, 커넬들을 컴파일하는 방법.
  6. 제 5 항에 있어서,
    상기 메타데이터는, 상기 제 1 메모리 영역과 상기 제 2 메모리 영역 사이의 오버랩의 바이트들의 수를 포함하는, 커넬들을 컴파일하는 방법.
  7. 제 5 항에 있어서,
    상기 메타데이터는 메모리 오버랩 영역의 시작 어드레스 및 메모리 오버랩 영역의 종료 어드레스 중 적어도 하나를 더 포함하는, 커넬들을 컴파일하는 방법.
  8. 제 1 항에 있어서,
    상기 컴파일링 프로세서는 CPU (central processing unit) 를 포함하고, 상기 타겟 프로세서는 GPU (graphics processing unit) 를 포함하는, 커넬들을 컴파일하는 방법.
  9. 제 1 항에 있어서,
    상기 컴파일러는, Microsoft DirectCompute, 및 OpenCL 중 적어도 하나를 포함하는 이종 컴퓨팅 프레임워크를 이용하여 상기 커넬을 리컴파일하는, 커넬들을 컴파일하는 방법.
  10. 제 1 항에 있어서,
    상기 커넬 아규먼트들은, 상기 아규먼트들을 위해 할당된 메모리의 버퍼 영역을 포함하는, 커넬들을 컴파일하는 방법.
  11. 디바이스로서,
    메모리; 및
    컴파일링 프로세서를 포함하고,
    상기 컴파일링 프로세서는
    상기 컴파일링 프로세서 상에서 실행되는 컴파일러 및 런타임으로 이루어지는 군의 적어도 하나에 의해, 컴파일된 커넬의 2진 코드를 실행하기 위해 아규먼트들을 생성하고;
    상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 커넬 아규먼트들의 제 1 메모리 영역에 대한 제 1 메모리 참조 및 커넬 아규먼트들의 제 2 메모리 영역에 대한 제 2 메모리 참조가 동일한 메모리 영역을 참조하는지 여부를 결정하고;
    상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 상기 결정에 기초하여 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조와 연관된 메타데이터를 생성하는 것으로서, 상기 메타데이터는 상기 제 1 메모리 영역과 상기 제 2 메모리 영역 사이의 관계를 나타내는, 상기 메타데이터를 생성하고; 및
    상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 상기 커넬의 제 1 및 제 2 메모리 참조들이 상기 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여:
    상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 컴파일러로 하여금 상기 메타데이터에 기초하여 상기 커넬을 리컴파일하게 하고; 및
    상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 리컴파일된 상기 커넬을 실행하도록 타겟 프로세서에 명령하도록 구성되는, 디바이스.
  12. 제 11 항에 있어서,
    상기 커넬의 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조가 상기 동일한 메모리 영역을 참조하는지 여부를 결정하기 위하여, 상기 컴파일링 프로세서는 또한, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조를 포함하는 상기 커넬의 루프 코드 섹션을 결정하도록 구성되고,
    상기 커넬을 리컴파일하기 위하여, 상기 컴파일링 프로세서는 또한, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 생성된 상기 메타데이터에 기초하여 상기 루프 코드 섹션을 언롤링하고, 언롤링된 상기 루프 코드 섹션을 컴파일하도록 구성되는, 디바이스.
  13. 제 11 항에 있어서,
    상기 커넬의 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조가 상기 동일한 메모리 영역을 참조하는지 여부를 결정하기 위하여, 상기 컴파일링 프로세서는 또한, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조를 포함하는 상기 커넬의 코드 섹션을 결정하도록 구성되고,
    상기 커넬을 리컴파일하기 위하여, 상기 컴파일링 프로세서는 또한, 상기 코드 섹션의 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조가 상기 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 생성된 상기 메타데이터에 기초하여 상기 코드 섹션의 로드 동작 및 저장 동작 중 적어도 하나를 리오더링하도록 구성되는, 디바이스.
  14. 제 11 항에 있어서,
    상기 커넬의 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조가 상기 동일한 메모리 영역을 참조하는지 여부를 결정하기 위하여, 상기 컴파일링 프로세서는 또한, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조를 포함하는 상기 커넬의 코드 섹션을 결정하도록 구성되고,
    상기 커넬을 리컴파일하기 위하여, 상기 컴파일링 프로세서는 또한, 상기 코드 섹션의 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조가 상기 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 생성된 상기 메타데이터에 기초하여 상기 코드 섹션의 복수의 스칼라 명령들을 적어도 하나의 벡터 명령으로 벡터화하도록 구성되는, 디바이스.
  15. 제 11 항에 있어서,
    상기 메타데이터는 또한, 상기 제 1 메모리 영역과 상기 제 2 메모리 영역 사이의 오버랩 영역을 나타내는, 디바이스.
  16. 제 15 항에 있어서,
    상기 메타데이터는, 상기 제 1 메모리 영역과 상기 제 2 메모리 영역 사이의 오버랩의 바이트들의 수를 포함하는, 디바이스.
  17. 제 15 항에 있어서,
    상기 메타데이터는 메모리 오버랩 영역의 시작 어드레스 및 메모리 오버랩 영역의 종료 어드레스 중 적어도 하나를 더 포함하는, 디바이스.
  18. 제 11 항에 있어서,
    상기 컴파일링 프로세서는 CPU (central processing unit) 를 포함하고, 상기 타겟 프로세서는 GPU (graphics processing unit) 를 포함하는, 디바이스.
  19. 제 11 항에 있어서,
    상기 컴파일러는, Microsoft DirectCompute, 및 OpenCL 중 적어도 하나를 포함하는 이종 컴퓨팅 프레임워크를 이용하여 상기 커넬을 리컴파일하는, 디바이스.
  20. 제 11 항에 있어서,
    상기 커넬 아규먼트들은, 상기 아규먼트들을 위해 할당된 메모리의 버퍼 영역을 포함하는, 디바이스.
  21. 명령들을 포함하는 비일시적 컴퓨터 판독가능 저장 매체로서,
    상기 명령들은, 실행될 때, 컴파일링 프로세서로 하여금,
    상기 컴파일링 프로세서 상에서 실행되는 컴파일러 및 런타임으로 이루어지는 군의 적어도 하나에 의해, 컴파일된 커넬의 2진 코드를 실행하기 위해 아규먼트들을 생성하게 하고;
    상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 커넬 아규먼트들의 제 1 메모리 영역에 대한 제 1 메모리 참조 및 커넬 아규먼트들의 제 2 메모리 영역에 대한 제 2 메모리 참조가 동일한 메모리 영역을 참조하는지 여부를 결정하게 하고;
    상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 상기 결정에 기초하여 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조와 연관된 메타데이터를 생성하게 하는 것으로서, 상기 메타데이터는 상기 제 1 메모리 영역과 상기 제 2 메모리 영역 사이의 관계를 나타내는, 상기 메타데이터를 생성하게 하고; 및
    상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 상기 커넬의 제 1 및 제 2 메모리 참조들이 상기 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여:
    상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 컴파일러로 하여금 상기 메타데이터에 기초하여 상기 커넬을 리컴파일하게 하고; 및
    상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 리컴파일된 상기 커넬을 실행하도록 타겟 프로세서에 명령하게 하는, 비일시적 컴퓨터 판독가능 저장 매체.
  22. 제 21 항에 있어서,
    상기 컴파일링 프로세서로 하여금 상기 커넬의 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조가 상기 동일한 메모리 영역을 참조하는지 여부를 결정하게 하는 명령들은, 실행될 때, 상기 컴파일링 프로세서로 하여금, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조를 포함하는 상기 커넬의 루프 코드 섹션을 결정하게 하는 명령들을 더 포함하고,
    상기 컴파일링 프로세서로 하여금 상기 커넬을 리컴파일하게 하는 명령들은, 실행될 때, 상기 컴파일링 프로세서로 하여금, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 생성된 상기 메타데이터에 기초하여 상기 루프 코드 섹션을 언롤링하게 하고, 언롤링된 상기 루프 코드 섹션을 컴파일하게 하는 명령들을 더 포함하는, 비일시적 컴퓨터 판독가능 저장 매체.
  23. 제 21 항에 있어서,
    상기 컴파일링 프로세서로 하여금 상기 커넬의 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조가 상기 동일한 메모리 영역을 참조하는지 여부를 결정하게 하는 명령들은, 실행될 때, 상기 컴파일링 프로세서로 하여금, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조를 포함하는 상기 커넬의 코드 섹션을 결정하게 하는 명령들을 더 포함하고,
    상기 컴파일링 프로세서로 하여금 상기 커넬을 리컴파일하게 하는 명령들은, 실행될 때, 상기 컴파일링 프로세서로 하여금, 상기 코드 섹션의 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조가 상기 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 생성된 상기 메타데이터에 기초하여 상기 코드 섹션의 로드 동작 및 저장 동작 중 적어도 하나를 리오더링하게 하는 명령들을 더 포함하는, 비일시적 컴퓨터 판독가능 저장 매체.
  24. 제 21 항에 있어서,
    상기 컴파일링 프로세서로 하여금 상기 커넬의 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조가 상기 동일한 메모리 영역을 참조하는지 여부를 결정하게 하는 명령들은, 실행될 때, 상기 컴파일링 프로세서로 하여금, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조를 포함하는 상기 커넬의 코드 섹션을 결정하게 하는 명령들을 더 포함하고,
    상기 컴파일링 프로세서로 하여금 상기 커넬을 리컴파일하게 하는 명령들은, 실행될 때, 상기 컴파일링 프로세서로 하여금, 상기 코드 섹션의 상기 제 1 메모리 참조 및 상기 제 2 메모리 참조가 상기 동일한 메모리 영역을 참조하지 않는다고 결정하는 것에 응답하여, 상기 컴파일링 프로세서 상에서 실행되는 상기 컴파일러 및 상기 런타임으로 이루어지는 군의 적어도 하나에 의해, 생성된 상기 메타데이터에 기초하여 상기 코드 섹션의 복수의 스칼라 명령들을 적어도 하나의 벡터 명령으로 벡터화하게 하는 명령들을 더 포함하는, 비일시적 컴퓨터 판독가능 저장 매체.
  25. 제 21 항에 있어서,
    상기 메타데이터는 또한, 상기 제 1 메모리 영역과 상기 제 2 메모리 영역 사이의 오버랩 영역을 나타내는, 비일시적 컴퓨터 판독가능 저장 매체.
  26. 제 25 항에 있어서,
    상기 메타데이터는, 상기 제 1 메모리 영역과 상기 제 2 메모리 영역 사이의 오버랩의 바이트들의 수를 포함하는, 비일시적 컴퓨터 판독가능 저장 매체.
  27. 제 25 항에 있어서,
    상기 메타데이터는 메모리 오버랩 영역의 시작 어드레스 및 메모리 오버랩 영역의 종료 어드레스 중 적어도 하나를 더 포함하는, 비일시적 컴퓨터 판독가능 저장 매체.
  28. 제 21 항에 있어서,
    상기 컴파일링 프로세서는 CPU (central processing unit) 를 포함하고, 상기 타겟 프로세서는 GPU (graphics processing unit) 를 포함하는, 비일시적 컴퓨터 판독가능 저장 매체.
  29. 제 21 항에 있어서,
    상기 컴파일러는, Microsoft DirectCompute, 및 OpenCL 중 적어도 하나를 포함하는 이종 컴퓨팅 프레임워크를 이용하여 상기 커넬을 리컴파일하는, 비일시적 컴퓨터 판독가능 저장 매체.
  30. 제 21 항에 있어서,
    상기 커넬 아규먼트들은, 상기 아규먼트들을 위해 할당된 메모리의 버퍼 영역을 포함하는, 비일시적 컴퓨터 판독가능 저장 매체.
KR1020167028388A 2014-04-04 2015-03-19 컴파일러 최적화를 위한 메모리 참조 메타데이터 KR101832656B1 (ko)

Applications Claiming Priority (3)

Application Number Priority Date Filing Date Title
US14/245,946 2014-04-04
US14/245,946 US9710245B2 (en) 2014-04-04 2014-04-04 Memory reference metadata for compiler optimization
PCT/US2015/021585 WO2015153143A1 (en) 2014-04-04 2015-03-19 Memory reference metadata for compiler optimization

Publications (2)

Publication Number Publication Date
KR20160141753A true KR20160141753A (ko) 2016-12-09
KR101832656B1 KR101832656B1 (ko) 2018-02-26

Family

ID=52829334

Family Applications (1)

Application Number Title Priority Date Filing Date
KR1020167028388A KR101832656B1 (ko) 2014-04-04 2015-03-19 컴파일러 최적화를 위한 메모리 참조 메타데이터

Country Status (6)

Country Link
US (1) US9710245B2 (ko)
EP (1) EP3132347A1 (ko)
JP (1) JP6329274B2 (ko)
KR (1) KR101832656B1 (ko)
CN (1) CN106164862A (ko)
WO (1) WO2015153143A1 (ko)

Families Citing this family (24)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US20130113809A1 (en) * 2011-11-07 2013-05-09 Nvidia Corporation Technique for inter-procedural memory address space optimization in gpu computing compiler
GB2514618B (en) * 2013-05-31 2020-11-11 Advanced Risc Mach Ltd Data processing systems
US9785413B2 (en) * 2015-03-06 2017-10-10 Intel Corporation Methods and apparatus to eliminate partial-redundant vector loads
US9824419B2 (en) * 2015-11-20 2017-11-21 International Business Machines Corporation Automatically enabling a read-only cache in a language in which two arrays in two different variables may alias each other
WO2017209876A1 (en) 2016-05-31 2017-12-07 Brocade Communications Systems, Inc. Buffer manager
US10169010B2 (en) 2016-06-01 2019-01-01 International Business Machines Corporation Performing register promotion optimizations in a computer program in regions where memory aliasing may occur and executing the computer program on processor hardware that detects memory aliasing
US10169009B2 (en) 2016-06-01 2019-01-01 International Business Machines Corporation Processor that detects memory aliasing in hardware and assures correct operation when memory aliasing occurs
US9934009B2 (en) * 2016-06-01 2018-04-03 International Business Machines Corporation Processor that includes a special store instruction used in regions of a computer program where memory aliasing may occur
JP6810380B2 (ja) * 2016-10-07 2021-01-06 日本電気株式会社 ソースプログラム変換システム、ソースプログラム変換方法、及びソースプログラム変換プログラム
US10108404B2 (en) * 2016-10-24 2018-10-23 International Business Machines Corporation Compiling optimized entry points for local-use-only function pointers
WO2018119778A1 (en) * 2016-12-28 2018-07-05 Intel Corporation System and method for vector communication
US10547491B2 (en) * 2017-08-28 2020-01-28 Genband Us Llc Transcoding with a vector processing unit
US10540194B2 (en) * 2017-12-21 2020-01-21 International Business Machines Corporation Runtime GPU/CPU selection
CN108470072B (zh) * 2018-03-30 2019-07-09 迅讯科技（北京）有限公司 一种查询编译方法和装置
US11367160B2 (en) * 2018-08-02 2022-06-21 Nvidia Corporation Simultaneous compute and graphics scheduling
US10884720B2 (en) * 2018-10-04 2021-01-05 Microsoft Technology Licensing, Llc Memory ordering annotations for binary emulation
CN111340678A (zh) * 2018-12-19 2020-06-26 华为技术有限公司 一种数据缓存系统、图形处理器及数据缓存方法
US10884664B2 (en) * 2019-03-14 2021-01-05 Western Digital Technologies, Inc. Executable memory cell
US10872057B1 (en) * 2019-05-23 2020-12-22 Xilinx, Inc. Partitioning in a compiler flow for a heterogeneous multi-core architecture
JP7460902B2 (ja) 2020-06-09 2024-04-03 富士通株式会社 コンパイラプログラム、コンパイル方法、情報処理装置
JP7164267B2 (ja) * 2020-12-07 2022-11-01 インテル・コーポレーション ヘテロジニアスコンピューティングのためのシステム、方法及び装置
CN114398011B (zh) * 2022-01-17 2023-09-22 安谋科技（中国）有限公司 数据存储方法、设备和介质
EP4276602A1 (de) * 2022-05-12 2023-11-15 Siemens Aktiengesellschaft System mit quellcodeumwandler-spezifizierten speicherbereichen
US20240095024A1 (en) * 2022-06-09 2024-03-21 Nvidia Corporation Program code versions

Family Cites Families (22)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US20040205740A1 (en) 2001-03-29 2004-10-14 Lavery Daniel M. Method for collection of memory reference information and memory disambiguation
US6877088B2 (en) * 2001-08-08 2005-04-05 Sun Microsystems, Inc. Methods and apparatus for controlling speculative execution of instructions based on a multiaccess memory condition
CA2372034A1 (en) 2002-02-14 2003-08-14 Cloakware Corporation Foiling buffer-overflow and alien-code attacks by encoding
WO2004021176A2 (de) 2002-08-07 2004-03-11 Pact Xpp Technologies Ag Verfahren und vorrichtung zur datenverarbeitung
US7093101B2 (en) * 2002-11-21 2006-08-15 Microsoft Corporation Dynamic data structures for tracking file system free space in a flash memory device
US7565631B1 (en) 2004-07-02 2009-07-21 Northwestern University Method and system for translating software binaries and assembly code onto hardware
US8886887B2 (en) * 2007-03-15 2014-11-11 International Business Machines Corporation Uniform external and internal interfaces for delinquent memory operations to facilitate cache optimization
US8413126B2 (en) * 2007-06-15 2013-04-02 Cray Inc. Scalar code reduction using shortest path routing
US20090070753A1 (en) 2007-09-07 2009-03-12 International Business Machines Corporation Increase the coverage of profiling feedback with data flow analysis
US8458671B1 (en) * 2008-02-12 2013-06-04 Tilera Corporation Method and system for stack back-tracing in computer programs
US8209525B2 (en) 2008-08-15 2012-06-26 Apple Inc. Method and apparatus for executing program code
BRPI0925055A2 (pt) * 2009-06-26 2015-07-28 Intel Corp "otimizações para um sistema de memória transacional ilimitada (utm)"
US8589867B2 (en) * 2010-06-18 2013-11-19 Microsoft Corporation Compiler-generated invocation stubs for data parallel programming model
US8527737B2 (en) * 2010-06-23 2013-09-03 Apple Inc. Using addresses to detect overlapping memory regions
US20120259843A1 (en) * 2011-04-11 2012-10-11 Timothy Child Database acceleration using gpu and multicore cpu systems and methods
US8935683B2 (en) * 2011-04-20 2015-01-13 Qualcomm Incorporated Inline function linking
US8468507B2 (en) 2011-06-10 2013-06-18 Microsoft Corporation Binding executable code at runtime
US8627018B2 (en) * 2011-11-18 2014-01-07 Microsoft Corporation Automatic optimization for programming of many-core architectures
US20130141443A1 (en) * 2011-12-01 2013-06-06 Michael L. Schmit Software libraries for heterogeneous parallel processing platforms
US9256915B2 (en) * 2012-01-27 2016-02-09 Qualcomm Incorporated Graphics processing unit buffer management
US9734333B2 (en) * 2012-04-17 2017-08-15 Heat Software Usa Inc. Information security techniques including detection, interdiction and/or mitigation of memory injection attacks
CN103116513B (zh) * 2012-07-13 2016-03-23 北京时代民芯科技有限公司 一种异构多核处理器编译器

Also Published As

Publication number Publication date
US20150286472A1 (en) 2015-10-08
JP2017509999A (ja) 2017-04-06
KR101832656B1 (ko) 2018-02-26
EP3132347A1 (en) 2017-02-22
CN106164862A (zh) 2016-11-23
WO2015153143A1 (en) 2015-10-08
JP6329274B2 (ja) 2018-05-23
US9710245B2 (en) 2017-07-18

Similar Documents

Publication Publication Date Title
KR101832656B1 (ko) 컴파일러 최적화를 위한 메모리 참조 메타데이터
US10956218B2 (en) Enqueuing kernels from kernels on GPU/CPU
EP3126971B1 (en) Program execution on heterogeneous platform
Membarth et al. Generating device-specific GPU code for local operators in medical imaging
CN110008009B (zh) 在运行时绑定常量以提高资源利用率
US8341615B2 (en) Single instruction multiple data (SIMD) code generation for parallel loops using versioning and scheduling
US20160011857A1 (en) Dynamic Compiler Parallelism Techniques
Mikushin et al. KernelGen--The Design and Implementation of a Next Generation Compiler Platform for Accelerating Numerical Models on GPUs
Membarth et al. Code generation for embedded heterogeneous architectures on Android
CN114895965A (zh) 实现工作负载的静态映射的乱序流水线执行的方法和装置
Michell et al. Tasklettes–a fine grained parallelism for Ada on multicores
Haidl et al. High-level programming for many-cores using C++ 14 and the STL
Haaser et al. An incremental rendering VM
US10996960B1 (en) Iterating single instruction, multiple-data (SIMD) instructions
Besnard et al. A framework for automatic and parameterizable memoization
Tian et al. Optimizing gpu register usage: Extensions to openacc and compiler optimizations
US9569191B2 (en) Dynamic programming platform for implementing bulk data operators in a multithreaded environment
Guide Cuda c best practices guide
Crisci et al. SYCL-Bench 2020: Benchmarking SYCL 2020 on AMD, Intel, and NVIDIA GPUs
Kuo et al. The design of LLVM-based shader compiler for embedded architecture
Jamieson et al. Compact native code generation for dynamic languages on micro-core architectures
KR20200083025A (ko) 런타임 수행특성을 도출하는 그래픽스 프로세싱 유닛 및 이의 동작방법
US11593114B1 (en) Iterating group sum of multiple accumulate operations
US20220391216A1 (en) Graphics processing
Membarth et al. Automatic optimization of in-flight memory transactions for GPU accelerators based on a domain-specific language for medical imaging

Legal Events

Date Code Title Description
A201 Request for examination
A302 Request for accelerated examination
E701 Decision to grant or registration of patent right
GRNT Written decision to grant