Intel 컴파일러의 OpenMP 오프로드 기능을 사용하여 GPU 계산 매트릭스 곱하기

선언(경과)
XLsoft"OpenMP* 프로그램 GPU 오프로드 도전 이벤트" 같은 걸 알고 Windows OpenMP 마운트 해제 기능으로 Intel 그래픽을 치는 거예요.
평소 HPC 환경에서 병행 계산하는 HPC 호스트의 일원으로서 GPU에서 범용 연산을 하는 것에 관해 이것저것 소르나니 민중으로서 무척 신경 쓰이는 이슈라 조금 진지하게 시도해봤다.
운영 환경
ThinkPad X1 Carbon 7th
내장형 Intel Core i7-1050U, Intel UHD 그래픽
(전원 설정 중 최고 성능 상태, 이 일대를 만지작거리면 어떻게 되는지 확인되지 않았다)
차리다
설치oneAPI의 기본 Toolkit 및 HPC Toolkit입니다.
서버가 느린지 뭔지, 온라인 인스톨러에 베이스 툴킷을 설치하다가 중간에 오류가 종료돼 툴킷을 2개 설치한 로칼 인스톨러다.
이후 VTune을 사용할 때 sampling driver가 올바르지 않은 정보가 나타났습니다. 참조이것, cmd.관리자 권한으로 exe, C:\Program Files(x86)\Intel\one API\vtune\202.1-beta 10\bin64\amplxe-sepreg를 실행합니다.적당히exe를 실행합니다. (- c 옵션으로 설치하고 - s 옵션으로 확인합니다.)
아래와 같이 설치가 된 느낌이에요.

c:\Program Files (x86)\Intel\oneAPI\vtune\2021.1-beta10\bin64>amplxe-sepreg.exe -c
Checking platform...
Platform is genuine Intel: OK
Platform has SSE2: OK
Platform architecture: INTEL64
User has admin rights: OK
Drivers will be installed to C:\Windows\System32\Drivers\
Checking sepdrv5 driver path...OK
Checking sepdrv5 service...
Driver status: the sepdrv5 service is running
Checking sepdal driver path...OK
Checking sepdal service...
Driver status: the sepdal service is running
Checking socperf3 driver path...OK
Checking socperf3 service...
Driver status: the socperf3 service is running

c:\Program Files (x86)\Intel\oneAPI\vtune\2021.1-beta10\bin64>amplxe-sepreg.exe -s
Checking sepdrv5 driver path...OK
Checking sepdrv5 service...
Driver status: the sepdrv5 service is running
Checking sepdal driver path...OK
Checking sepdal service...
Driver status: the sepdal service is running
Checking socperf3 driver path...OK
Checking socperf3 service...
Driver status: the socperf3 service is running

설치된 툴키트 버전은 이런 느낌이다.

icc>icx -v
Intel(R) oneAPI DPC++/C++ Compiler Pro for applications running on Intel(R) 64, Version 2021.1 Beta Build 20201005
Copyright (C) 1985-2020 Intel Corporation. All rights reserved.

Intel(R) oneAPI DPC++ Compiler Pro 2021.1 (2020.8.0.1005)
Target: x86_64-pc-windows-msvc
Thread model: posix
InstalledDir: C:\PROGRA~2\Intel\oneAPI\compiler\latest\windows\bin

실험 프로그램
끈적끈적하지만 간단한 삼중순환의 행렬적을 준비해 최외순환을 병행화했다.
참고로malloc를 통해 동적으로 배열을 확보하면 프로그램이 정상적으로 끝날 수 없고 계산 결과도 정확하게 얻을 수 없다.맵 지시문에서 길이를 지정해도 안 됩니다.
이 일대 조사는 또 좋으면...

  float a[n][n], b[n][n], c[n][n];
  t_begin = omp_get_wtime();
#pragma omp target data map(to:a,b) map(tofrom:c)
#pragma omp target teams distribute parallel for private(j,k)
  for(i=0; i<n; i++){
    for(j=0; j<n; j++){
      for(k=0; k<n; k++){
        c[i][j] += a[i][k] * b[k][j];
      }
    }
  }
  t_end = omp_get_wtime();
  printf("time %f\n", (double)(t_end - t_begin));

편역하다
시작 메뉴에서 Intel one API command proompt for Intel 64 for Visual Studio 2017을 시작합니다.디렉토리를 이동하고 다음을 클릭합니다.

icx matmul_ofl_static1.c /Qiopenmp /Qopenmp-targets:spir64 /nologo

편역다음 위닝이 나왔지만 해제 방법은 불분명하다.

matmul_ofl_static1-b5aa07.obj : warning LNK4078: 複数の '__CLANG_OFFLOAD_BUNDLE__openmp-s' セクションが見つかりました。 これらは異なる属性 (40500040) を持っています。

실행 시set OMP_TARGET_OFFLOAD=DISABLED에 환경 변수를 설정하면 CPU가 실행되고, 설정set OMP_TARGET_OFFLOAD=MANDATORY하면 GPU에서 실행되며, 동시에 설정set LIBOMPTARGET_PROFILE=T하면 간단한 구성 파일 정보를 얻을 수 있습니다.
문제 크기(행렬의 행과 열)가 100인 상태에서 실행할 때 다음과 같은 결과를 얻을 수 있습니다.
현재 디렉토리의 디스플레이는 생략됩니다.

>set LIBOMPTARGET_PROFILE=T

>set OMP_TARGET_OFFLOAD=DISABLED

>matmul_ofl_static1
n = 100
time 0.001609

>set OMP_TARGET_OFFLOAD=MANDATORY

>matmul_ofl_static1.exe
n = 100
Target LEVEL0 RTL --> Warning: Invalid kernel timestamp bit width (0). Long-running kernels may report incorrect device time.
time 5.878747
LIBOMPTARGET_PROFILE for OMP DEVICE(0) Intel(R) Gen9, Thread 12788
-- Name: Host Time (msec), Device Time (msec)
-- DataAlloc: 0.166800, 0.166800
-- DataRead: 0.002000, 0.002000
-- DataWrite: 0.179900, 0.179900
-- Kernel#__omp_offloading_821ec21a_199401__Z4main_l35: 1.838900, 1.328498
-- ModuleBuild: 5786.827200, 5786.827200
-- Total: 5789.014800, 5788.504398

결과적으로 GPU 사용이 느립니다.몇 번 해봤는데 6초 정도 됐어요.CPU의 등급은 0.001초이다.
다만

-- Kernel#__omp_offloading_821ec21a_199401__Z4main_l35: 1.838900, 1.328498

의 부분은 GPU 커널처럼 느껴지며 실제 계산 시간은 1밀리초 정도이며 CPU와 차이가 많지 않다.
따라서 문제의 사이즈를 1000까지 확대해 측정한다.

icc>set OMP_TARGET_OFFLOAD=DISABLED

icc>matmul_ofl_static1.exe
n = 1000
time 1.261088

icc>set OMP_TARGET_OFFLOAD=MANDATORY

icc>matmul_ofl_static1.exe
n = 1000
Target LEVEL0 RTL --> Warning: Invalid kernel timestamp bit width (0). Long-running kernels may report incorrect device time.
time 7.539387
LIBOMPTARGET_PROFILE for OMP DEVICE(0) Intel(R) Gen9, Thread 14048
-- Name: Host Time (msec), Device Time (msec)
-- DataAlloc: 0.441400, 0.441400
-- DataRead: 0.625600, 0.625600
-- DataWrite: 18.447000, 18.447000
-- Kernel#__omp_offloading_821ec21a_199401__Z4main_l35: 1320.353900, 1313.705408
-- ModuleBuild: 6106.610200, 6106.610200
-- Total: 7446.478100, 7439.829608```

아니면 주요 계산 부분만 보면 1.2초와 1.3초 같은 성능으로 나온다.
이렇게 성능이 비슷하면 정말 GPU가 움직이는 건가요?반대로 불안할 수도 있지만 메시지만 보면 움직이는데...
VTune에서 봤어요.
시작 메뉴에서 Intel VTune Profiller 2021.1을 시작합니다.빨간색 상자 부분을 적당히 설정하고 실행합니다.(매개 변수에 1000을 주지만 프로그램에서는 의미가 없고 사실상 정태적으로 확보된다.)

DLL 읽기 오류가 발생했지만 설정할 곳을 몰라서 일단 놔뒀다.

약간의 경고가 있을 수 있지만 직업 압축 파일 자체는 이미 완성되었다.(debugging information을 찾을 수 없습니다. 컴파일할 때/debug 옵션을 추가하려고 했지만 해결되지 않았습니다.)

전체 개요

자세한 내용이 나왔어요.

그런데 도대체 어디를 봐야 할지...?
일반적으로 1초 정도 걸린다고 생각하는 GPU 커널의 위치도 이런 상황을 잘 모른다.예를 들어 VTune의 측정 위치를 제한하는 함수를 넣지 않으면 안 될 수도 있죠?
초보적 결론
GPU가 움직이는 것 같지만 더 이상 느껴지지 않는다.만약 행렬이 CPU보다 빨리 쌓이지 않는다면 일부러 사용하고 싶은 사람은......
그리고 Windows는 명령을 사용하는 것이 가장 좋지 않다.
어쨌든 이번에는 여기까지다.다음이 있을지 모르겠어요.
물론 기술적으로는 흥미로운 대목으로 앞으로의 발전을 기대한다.