1
두 가지 프로그램이 있습니다.OpenCL을 사용하는 행렬 - 벡터 곱셈
먼저 행렬 - 행렬 곱셈을 opencl을 사용하여 만듭니다. 내 GPU에서 호스트 CPU (예를 들어, 0.2 초 vs 18 초)에서 훨씬 더 나은 결과를 생성합니다.
두 번째는 opencl을 사용하여 행렬 벡터 곱셈을 수행하며 GPU에서 약간 느린 다음 호스트 CPU에서 작동합니다.
이유는 무엇입니까? 여기
커널
__kernel void matrixVectorMul(__global float* resultVector,
__global float* matrixA,
__global float* vectorB,
int width_A)
{
int tx = get_global_id(0);
float value = 0;
for (unsigned int k = 0; k < width_A; ++k) {
value += matrixA[tx * width_A + k] * vectorB[k];
}
resultVector[tx] = value;
}
그리고 호스트 코드
#include <stdlib.h>
#define __CL_ENABLE_EXCEPTIONS
#include "cl.hpp"
#include <fstream>
#include <iostream>
#include <time.h>
#include <cmath>
#define LOCAL_SIZE 512
#define WIDTH_A (4096*2)
#define HEIGHT_A (4096*2)
float *matrix_A;
float *vector_B;
float *result_vector;
float *result_vector_host;
void randomInit(float *data, int size) {
for (unsigned int i = 0; i < size; ++i)
data[i] = rand()/(float)RAND_MAX;
}
void GenerateTestData() {
srand((unsigned int)time(NULL));
unsigned int size_A = WIDTH_A * HEIGHT_A;
matrix_A = new float[size_A];
vector_B = new float[WIDTH_A];
randomInit(matrix_A, size_A);
randomInit(vector_B, WIDTH_A);
result_vector = new float[WIDTH_A];
result_vector_host = new float[WIDTH_A];
}
void PerformCalculationOnDevice(cl::Device device) {
clock_t start_t, end_t;
start_t = clock();
std::vector<cl::Device> contextDevices;
contextDevices.push_back(device);
cl::Context context(contextDevices);
cl::CommandQueue queue(context, device);
std::fill_n(result_vector, WIDTH_A, 0);
cl::Buffer cl_matrix_A = cl::Buffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, WIDTH_A * HEIGHT_A * sizeof(float), matrix_A);
cl::Buffer cl_vector_B = cl::Buffer(context, CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR, WIDTH_A * sizeof(float), vector_B);
cl::Buffer cl_result_vector = cl::Buffer(context, CL_MEM_WRITE_ONLY|CL_MEM_COPY_HOST_PTR, WIDTH_A * sizeof(float), result_vector);
end_t = clock();
std::cout << "Context, queue, buffers " << (float)(end_t - start_t)/CLOCKS_PER_SEC << std::endl;
std::ifstream sourceFile("MatrixVectorMultiplicationKernel.cl");
std::string sourceCode(std::istreambuf_iterator<char>(sourceFile),(std::istreambuf_iterator<char>()));
cl::Program::Sources source(1, std::make_pair(sourceCode.c_str(), sourceCode.length()+1));
cl::Program program = cl::Program(context, source);
program.build(contextDevices);
cl::Kernel kernel(program, "matrixVectorMul");
int iArg = 0;
kernel.setArg(iArg++, cl_result_vector);
kernel.setArg(iArg++, cl_matrix_A);
kernel.setArg(iArg++, cl_vector_B);
kernel.setArg(iArg++, WIDTH_A);
start_t = clock();
queue.enqueueNDRangeKernel(kernel, cl::NullRange, cl::NDRange(HEIGHT_A), cl::NDRange(LOCAL_SIZE));
queue.finish();
queue.enqueueReadBuffer(cl_result_vector, CL_TRUE, 0, WIDTH_A * sizeof(float), result_vector);
end_t = clock();
std::cout << "enqueueNDRangeKernel and enqueueReadBuffer " << (float)(end_t - start_t)/CLOCKS_PER_SEC << std::endl;
}
void PerformCalculationOnHost() {
float tmp;
for(int row_A = 0; row_A < HEIGHT_A; row_A++) {
tmp = 0;
for(int col_A = 0; col_A < WIDTH_A; col_A++) {
tmp += matrix_A[row_A * WIDTH_A + col_A] * vector_B[col_A];
}
result_vector_host[row_A] = tmp;
}
}
int main(int argc, char** argv) {
GenerateTestData();
std::vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
std::vector<cl::Device> devices;
clock_t start_t = clock();
for (unsigned int iPlatform=0; iPlatform<platforms.size(); iPlatform++) {
platforms[iPlatform].getDevices(CL_DEVICE_TYPE_ALL, &devices);
for (unsigned int iDevice=0; iDevice<devices.size(); iDevice++) {
try {
PerformCalculationOnDevice(devices[iDevice]);
} catch (cl::Error error) {
std::cout << error.what() << "(" << error.err() << ")" << std::endl;
}
}
}
clock_t end_t = clock();
std::cout << "Device: " << (float)(end_t - start_t)/CLOCKS_PER_SEC << " seconds" << std::endl;
start_t = clock();
PerformCalculationOnHost();
end_t = clock();
std::cout << "Host: " << (float)(end_t - start_t)/CLOCKS_PER_SEC << " seconds" << std::endl;
int errors = 0;
float mean_deviation = 0;
FILE *f, *f_host;
f = fopen("device_result", "w");
f_host = fopen("host_result", "w");
for(int i = 0; i < WIDTH_A; i++) {
if(fabs(result_vector[i] - result_vector_host[i]) > 1E-3) {
errors++;
}
fprintf(f, "%.2f\n", result_vector[i]);
fprintf(f_host, "%.2f\n", result_vector_host[i]);
mean_deviation += fabs(result_vector[i] - result_vector_host[i]);
}
fclose(f); fclose(f_host);
mean_deviation /= WIDTH_A;
std::cout << "Errors = " << errors << std::endl;
std::cout << "Mean deviation = " << mean_deviation << std::endl;
delete[](matrix_A);
delete[](vector_B);
delete[](result_vector);
delete[](result_vector_host);
return 0;
}
내가 그것을 실행할 때, 난, 더이 경우 GPU 효율을 높이기 위해 다음과 같은 결과
Context, queue, buffers 0.45
enqueueNDRangeKernel and enqueueReadBuffer 1.31
Device: 1.79 seconds
Host: 1.42 seconds
Errors = 0
Mean deviation = 8.78572e-05
왜 로컬 크기를 하드 코딩합니까? CPU와 GPU는 매우 다른 최적의 로컬 작업 그룹 크기를 가지고 있습니다. 예를 들어 CPU에는 1024가 있지만 GPU에는 64가 있습니다. 매트릭스를 transpose하여 순차 메모리 읽기의 이점을 얻을 수 있다고 생각합니다. 지금은 인터리브되어 있습니다. – Thomas
이 프로그램은 주로 데모 용입니다. GPU의 일부 매개 변수를 쿼리하는 다른 간단한 프로그램을 사용하여 로컬 크기를 얻었습니다. vectorB를 커널의 로컬 배열에 복사하려고 시도했지만 (컴파일 속도가 빨라야 함) 컴파일되지 않습니다. 그러나 단순히 로컬 배열을 선언하고 요소에 액세스하려고하지 않으면 잘 컴파일됩니다. – int