//+------------------------------------------------------------------+
//| OpenCL kernel |
//+------------------------------------------------------------------+
const string
cl_src=
//--- 기본적으로 일부 GPU는 double을 지원하지 않습니다
//--- cl_khr_fp64 명령어는 복수로 작업을 활성화하는 데 사용됩니다
"#pragma OPENCL EXTENSION cl_khr_fp64 : enable \r\n"
//--- OpenCL 커널 함수
"__kernel void Test_GPU(__global double *data, \r\n"
" const int N, \r\n"
" const int total_arrays) \r\n"
" { \r\n"
" uint kernel_index=get_global_id(0); \r\n"
" if (kernel_index>total_arrays) return; \r\n"
" uint local_start_offset=kernel_index*N; \r\n"
" for(int i=0; i<N; i++) \r\n"
" { \r\n"
" data[i+local_start_offset] *= 2.0; \r\n"
" } \r\n"
" } \r\n";
//+------------------------------------------------------------------+
//| Test_CPU |
//+------------------------------------------------------------------+
bool Test_CPU(double &data[],const int N,const int id,const int total_arrays)
{
//--- 배열 사이즈 체크
if(ArraySize(data)==0) return(false);
//--- 배열 인덱스 체크
if(id>total_arrays) return(false);
//--- 인덱스 ID를 사용하여 배열의 로컬 오프셋 계산
int local_start_offset=id*N;
//--- 요소에 2를 곱함
for(int i=0; i<N; i++)
{
data[i+local_start_offset]*=2.0;
}
return true;
}
//---
#define ARRAY_SIZE 100 // 배열의 크기
#define TOTAL_ARRAYS 5 // 총 배열
//--- OpenCL 핸들
int cl_ctx; // OpenCL 컨텍스트 핸들
int cl_prg; // OpenCL 프로그램 핸들
int cl_krn; // OpenCL 커널 핸들
int cl_mem; // OpenCL 버퍼 핸들
//---
double DataArray1[]; // CPU 연산을 위한 데이터 배열
double DataArray2[]; // GPU 연산을 위한 데이터 배열
//+------------------------------------------------------------------+
//| 스크립트 프로그램 시작 함수 |
//+------------------------------------------------------------------+
int OnStart()
{
//--- OpenCL 개체 초기화
//--- OpenCL 컨텍스트 생성
if((cl_ctx=CLContextCreate())==INVALID_HANDLE)
{
Print("OpenCL을 찾을 수 없습니다. Error=",GetLastError());
return(1);
}
//--- OpenCL 프로그램 생성
if((cl_prg=CLProgramCreate(cl_ctx,cl_src))==INVALID_HANDLE)
{
CLContextFree(cl_ctx);
Print("OpenCL 프로그램 생성 실패. Error=",GetLastError());
return(1);
}
//--- OpenCL 커널 생성
if((cl_krn=CLKernelCreate(cl_prg,"Test_GPU"))==INVALID_HANDLE)
{
CLProgramFree(cl_prg);
CLContextFree(cl_ctx);
Print("OpenCL kernel 생성 실패. Error=",GetLastError());
return(1);
}
//--- OpenCL 버퍼 생성
if((cl_mem=CLBufferCreate(cl_ctx,ARRAY_SIZE*TOTAL_ARRAYS*sizeof(double),CL_MEM_READ_WRITE))==INVALID_HANDLE)
{
CLKernelFree(cl_krn);
CLProgramFree(cl_prg);
CLContextFree(cl_ctx);
Print("OpenCL 버퍼 생성 실패. Error=",GetLastError());
return(1);
}
//--- OpenCL 커널 상수 매개 변수 설정
CLSetKernelArgMem(cl_krn,0,cl_mem);
CLSetKernelArg(cl_krn,1,ARRAY_SIZE);
CLSetKernelArg(cl_krn,2,TOTAL_ARRAYS);
//--- 데이터 배열 준비
ArrayResize(DataArray1,ARRAY_SIZE*TOTAL_ARRAYS);
ArrayResize(DataArray2,ARRAY_SIZE*TOTAL_ARRAYS);
//--- 배열에 데이터를 채웁니다
for(int j=0; j<TOTAL_ARRAYS; j++)
{
//--- j번째 배열에 대한 로컬 시작 오프셋 계산
uint local_offset=j*ARRAY_SIZE;
//--- 인덱스 j로 배열 준비
for(int i=0; i<ARRAY_SIZE; i++)
{
//--- MathCos(i+j) 함수로 배열을 채웁니다;
DataArray1[i+local_offset]=MathCos(i+j);
DataArray2[i+local_offset]=MathCos(i+j);
}
};
//--- CPU 연산 테스트
for(int j=0; j<TOTAL_ARRAYS; j++)
{
//--- 인덱스 j를 사용하여 배열 계산
Test_CPU(DataArray1,ARRAY_SIZE,j,TOTAL_ARRAYS);
}
//--- CLExecute 매개변수 준비
uint offset[]={0};
//--- 전 세계 작업 규모
uint work[]={TOTAL_ARRAYS};
//--- OpenCL 버퍼에 데이터 쓰기
CLBufferWrite(cl_mem,DataArray2);
//--- OpenCL 커널을 실행
CLExecute(cl_krn,1,offset,work);
//--- OpenCL 버퍼에서 데이터 읽기
CLBufferRead(cl_mem,DataArray2);
//--- 총 에러
double total_error=0;
//--- 결과를 비교하고 오차를 계산
for(int j=0; j<TOTAL_ARRAYS; j++)
{
//--- j번째 배열에 대한 로컬 오프셋 계산
uint local_offset=j*ARRAY_SIZE;
//--- 결과를 비교
for(int i=0; i<ARRAY_SIZE; i++)
{
double v1=DataArray1[i+local_offset];
double v2=DataArray2[i+local_offset];
double delta=MathAbs(v2-v1);
total_error+=delta;
//--- 첫 번째 및 마지막 배열을 표시
if((j==0) || (j==TOTAL_ARRAYS-1))
PrintFormat("array %d of %d, element [%d]: %f, %f, [error]=%f",j+1,TOTAL_ARRAYS,i,v1,v2,delta);
}
}
PrintFormat("Total error: %f",total_error);
//--- OpenCL 개체 삭제
//--- free OpenCL 버퍼
CLBufferFree(cl_mem);
//--- free OpenCL 커널
CLKernelFree(cl_krn);
//--- free OpenCL 프로그램
CLProgramFree(cl_prg);
//--- free OpenCL 컨텍스트
CLContextFree(cl_ctx);
//---
return(0);
}
|