/**
* @brief Main principal
* @param argc El número de argumentos del programa
* @param argv Cadenas de argumentos del programa
* @return Nada si es correcto o algún número negativo si es incorrecto
*/
int main( int argc, char** argv ) {
if(argc != 2)
return -1;
// Medimos tiempo para el programa
const double start_time = getCurrentTimestamp();
FILE *kernels;
char *source_str;
size_t source_size, work_items;
// OpenCL runtime configuration
unsigned num_devices;
cl_platform_id platform_ids[3];
cl_uint ret_num_platforms;
cl_device_id device_id;
cl_context context = NULL;
cl_command_queue command_queue;
cl_program program = NULL;
cl_int ret;
cl_kernel kernelNUM;
cl_event kernel_event, finish_event;
cl_mem objPARTICULAS, objPESOS;
// Abrimos el fichero que contiene el kernel
fopen_s(&kernels, "numparticulasCPU.cl", "r");
if (!kernels) {
fprintf(stderr, "Fallo al cargar el kernel\n");
exit(-1);
}
source_str = (char *) malloc(0x100000);
source_size = fread(source_str, 1, 0x100000, kernels);
fclose(kernels);
// Obtenemos los IDs de las plataformas disponibles
if( clGetPlatformIDs(3, platform_ids, &ret_num_platforms) != CL_SUCCESS) {
printf("No se puede obtener id de la plataforma");
return -1;
}
// Intentamos obtener un dispositivo CPU soportado
if( clGetDeviceIDs(platform_ids[1], CL_DEVICE_TYPE_CPU, 1, &device_id, &num_devices) != CL_SUCCESS) {
printf("No se puede obtener id del dispositivo");
return -1;
}
clGetDeviceInfo(device_id, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), &work_items, NULL);
// Creación de un contexto OpenCL
context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
// Creación de una cola de comandos
command_queue = clCreateCommandQueue(context, device_id, CL_QUEUE_PROFILING_ENABLE, &ret);
// Creación de un programa kernel desde un fichero de código
program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);
ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
if (ret != CL_SUCCESS) {
size_t len;
char buffer[2048];
printf("Error: ¡Fallo al construir el programa ejecutable!\n");
clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
printf("%s", buffer);
exit(-1);
}
// Creación del kernel OpenCL
kernelNUM = clCreateKernel(program, "calc_num_particulas", &ret);
// Creamos el buffer para las partículas y reservamos espacio ALINEADO para los datos
size_t N = atoi(argv[1]);
particle *particulas = (particle*) _aligned_malloc(N * sizeof(particle), 64);
int *pesos = (int*) _aligned_malloc(N * sizeof(int), 64);
objPARTICULAS = clCreateBuffer(context, CL_MEM_READ_ONLY, N * sizeof(particle), NULL, &ret);
objPESOS = clCreateBuffer(context, CL_MEM_WRITE_ONLY, N * sizeof(int), NULL, &ret);
float sum = 0.0f;
const size_t global = 2;
const size_t local_work_size = 1;
// Inicializamos las partículas (Me interesan los pesos)
srand(time(NULL));
for(unsigned index = 0; index < N; ++index) {
particulas[index].x = 0.0;
particulas[index].y = 0.0;
particulas[index].s = 0.0;
particulas[index].xp = 0.0;
particulas[index].yp = 0.0;
particulas[index].sp = 0.0;
particulas[index].x0 = 0.0;
particulas[index].y0 = 0.0;
particulas[index].width = 0;
particulas[index].height = 0;
particulas[index].w = (float) (rand() % 2000);
sum+=particulas[index].w;
}
// Normalizamos los datos
for(int i = 0; i < N; ++i)
particulas[i].w /= sum;
// Transferimos las partículas al dispositivo y los pesos
cl_event write_event;
ret = clEnqueueWriteBuffer(command_queue, objPARTICULAS, CL_FALSE, 0, N * sizeof(particle), particulas, 0, NULL, &write_event);
// Establecemos los argumentos del kernel
ret = clSetKernelArg(kernelNUM, 0, sizeof(cl_mem), &objPARTICULAS);
ret = clSetKernelArg(kernelNUM, 1, sizeof(int), &N);
ret = clSetKernelArg(kernelNUM, 2, sizeof(cl_mem), &objPESOS);
// Ejecutamos el kernel. Un work-item por cada work-group o unidad de cómputo
ret = clEnqueueNDRangeKernel(command_queue, kernelNUM, 1, NULL, &global, &local_work_size, 1, &write_event, &kernel_event);
// Leemos los resultados
ret = clEnqueueReadBuffer(command_queue, objPESOS, CL_FALSE, 0, N * sizeof(int), pesos, 1, &kernel_event, &finish_event);
// Esperamos a que termine de leer los resultados
clWaitForEvents(1, &finish_event);
// Obtenemos el tiempo del kernel y de las transferencias CPU-RAM
cl_ulong totalKernel = getStartEndTime(kernel_event);
cl_ulong totalRam = getStartEndTime(write_event) + getStartEndTime(finish_event);
const double end_time = getCurrentTimestamp();
// Obtenemos el tiempo consumido por el programa, el kernel y las transferencias de memoria
printf("\nTiempo total del programa: %0.3f ms\n", (end_time - start_time) * 1e3);
printf("Tiempo total consumido por el kernel: %0.3f ms\n", double(totalKernel) * 1e-6);
printf("Tiempo total consumido en transferencias CPU-RAM: %0.3f ms\n", double(totalRam) * 1e-6);
// Liberamos todos los recursos usados (kernels y objetos OpenCL)
clReleaseEvent(kernel_event);
clReleaseEvent(finish_event);
clReleaseEvent(write_event);
clReleaseMemObject(objPARTICULAS);
clReleaseMemObject(objPESOS);
clReleaseKernel(kernelNUM);
clReleaseCommandQueue(command_queue);
clReleaseProgram(program);
clReleaseContext(context);
}
//run waifu2x
void run(std::vector<float> &input,
std::vector<float> &weight,
std::vector<float> &output,
std::vector<double> &bias,
int iter, const int kernelSize, int r, int c) {
unsigned int ipp[7][1] = { { 1 },{ 32 },{ 32 },{ 64 },{ 64 },{ 128 },{ 128 } };
unsigned int opp[7][1] = { { 32 },{ 32 },{ 64 },{ 64 },{ 128 },{ 128 },{ 1 } };
const unsigned nInputPlanes = ipp[iter][1];
const unsigned nOutputPlanes = opp[iter][1];
cl_int status;
status = clEnqueueWriteBuffer(queue, input_buf, CL_FALSE,
0, r * c * nInputPlanes * sizeof(float), input.data(), 0, NULL, NULL);
checkError(status, "Failed to transfer input");
status = clEnqueueWriteBuffer(queue, weight_buf, CL_FALSE,
0, kernelSize * nInputPlanes * nOutputPlanes * sizeof(float), weight.data(), 0, NULL, NULL);
checkError(status, "Failed to transfer weight");
status = clEnqueueWriteBuffer(queue, bias_buf, CL_FALSE,
0, nOutputPlanes * sizeof(double), bias.data(), 0, NULL, NULL);
checkError(status, "Failed to transfer bias");
clFinish(queue);
cl_event kernel_event;
const double start_time = getCurrentTimestamp();
unsigned argi = 0;
status = clSetKernelArg(kernel, argi++, sizeof(cl_mem), &output_buf);
checkError(status, "Failed to set argument %d", argi - 1);
status = clSetKernelArg(kernel, argi++, sizeof(cl_mem), &input_buf);
checkError(status, "Failed to set argument %d", argi - 1);
status = clSetKernelArg(kernel, argi++, sizeof(cl_mem), &weight_buf);
checkError(status, "Failed to set argument %d", argi - 1);
status = clSetKernelArg(kernel, argi++, sizeof(cl_mem), &bias_buf);
checkError(status, "Failed to set argument %d", argi - 1);
status = clSetKernelArg(kernel, argi++, sizeof(nInputPlanes), &nInputPlanes);
checkError(status, "Failed to set argument %d", argi - 1);
status = clSetKernelArg(kernel, argi++, sizeof(nOutputPlanes), &nOutputPlanes);
checkError(status, "Failed to set argument %d", argi - 1);
status = clSetKernelArg(kernel, argi++, sizeof(r), &r);
checkError(status, "Failed to set argument %d", argi - 1);
status = clSetKernelArg(kernel, argi++, sizeof(c), &c);
checkError(status, "Failed to set argument %d", argi - 1);
const size_t* global_work_size = opp[iter];
const size_t* local_work_size = ipp[iter];
printf("Iteration %d\n", iter);
status = clEnqueueNDRangeKernel(queue, kernel, 2, NULL,
global_work_size, local_work_size, 0, NULL, &kernel_event);
checkError(status, "Failed to launch kernel");
status = clFinish(queue);
checkError(status, "Failed to finish");
const double end_time = getCurrentTimestamp();
const double total_time = end_time - start_time;
// Wall-clock time taken.
printf("\nTime: %0.3f ms\n", total_time * 1e3);
// Get kernel times using the OpenCL event profiling API.
cl_ulong time_ns = getStartEndTime(kernel_event);
printf("Kernel time: %0.3f ms\n", double(time_ns) * 1e-6);
clReleaseEvent(kernel_event);
status = clEnqueueReadBuffer(queue, output_buf, CL_TRUE,
0, r - 1 * c - 1 * nOutputPlanes * sizeof(float), output.data(), 0, NULL, NULL);
checkError(status, "Failed to read output matrix");
}
/**
* @brief Main principal
* @param argc El número de argumentos del programa
* @param argv Cadenas de argumentos del programa
* @return Nada si es correcto o algún número negativo si es incorrecto
*/
int main( int argc, char** argv ) {
if(argc != 2)
return -1;
// Medimos tiempo para el programa
const double start_time = getCurrentTimestamp();
// Declaración de variables
IplImage *first_frame; // Primer frame
IplImage *frame, *hsv_frame;
CvCapture *video;
FILE *kernels;
char *source_str;
size_t source_size, work_items;
// OpenCL runtime configuration
unsigned num_devices;
cl_platform_id platform_ids[3];
cl_uint ret_num_platforms;
cl_device_id device_id;
cl_context context = NULL;
cl_command_queue command_queue;
cl_program program = NULL;
cl_int ret;
cl_kernel kernelHISTO;
cl_event kernel_event, finish_event;
cl_mem objFRAME, objHISTO;
// Abrimos el fichero que contiene el kernel
fopen_s(&kernels, "histoGPU.cl", "r");
if (!kernels) {
fprintf(stderr, "Fallo al cargar el kernel\n");
exit(-1);
}
source_str = (char *) malloc(0x100000);
source_size = fread(source_str, 1, 0x100000, kernels);
fclose(kernels);
// Obtenemos los IDs de las plataformas disponibles
if( clGetPlatformIDs(3, platform_ids, &ret_num_platforms) != CL_SUCCESS) {
printf("No se puede obtener id de la plataforma");
return -1;
}
// Intentamos obtener un dispositivo GPU soportado
if( clGetDeviceIDs(platform_ids[0], CL_DEVICE_TYPE_GPU, 1, &device_id, &num_devices) != CL_SUCCESS) {
printf("No se puede obtener id del dispositivo");
return -1;
}
clGetDeviceInfo(device_id, CL_DEVICE_MAX_COMPUTE_UNITS, sizeof(cl_uint), &work_items, NULL);
// Creación de un contexto OpenCL
context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
// Creación de una cola de comandos
command_queue = clCreateCommandQueue(context, device_id, CL_QUEUE_PROFILING_ENABLE, &ret);
// Creación de un programa kernel desde un fichero de código
program = clCreateProgramWithSource(context, 1, (const char **)&source_str, (const size_t *)&source_size, &ret);
ret = clBuildProgram(program, 1, &device_id, "-cl-nv-verbose", NULL, NULL);
if (ret != CL_SUCCESS) {
size_t len;
char buffer[2048];
printf("Error: ¡Fallo al construir el programa ejecutable!\n");
clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, sizeof(buffer), buffer, &len);
printf("%s", buffer);
exit(-1);
}
// Creación del kernel OpenCL
kernelHISTO = clCreateKernel(program, "calc_histo", &ret);
// Abrimos el fichero de video y leemos el primer frame
video = cvCaptureFromFile( argv[1] );
if( !video ) {
printf("No se pudo abrir el fichero de video %s\n", &argv[1]);
exit(-1);
}
first_frame = cvQueryFrame( video );
hsv_frame = cvCreateImage( cvGetSize(first_frame), IPL_DEPTH_32F, 3 );
cvConvertScale( first_frame, hsv_frame, 1.0 / 255.0, 0 );
cvCvtColor( hsv_frame, hsv_frame, CV_BGR2HSV );
// Creamos el buffer para los frames y el histograma
float *histo = (float*) _aligned_malloc(HTAM * sizeof(float), 64);
objFRAME = clCreateBuffer(context, CL_MEM_READ_ONLY, hsv_frame->imageSize, NULL, &ret);
objHISTO = clCreateBuffer(context, CL_MEM_READ_WRITE, HTAM * sizeof(float), NULL, &ret);
memset(histo, 0.0f, HTAM * sizeof(float));
const size_t global_work_size = work_items * 1024;
const size_t local_work_size = 1024;
// Transferimos el frame al dispositivo
cl_event write_event[2];
ret = clEnqueueWriteBuffer(command_queue, objFRAME, CL_FALSE, 0, hsv_frame->imageSize, hsv_frame->imageData, 0, NULL, &write_event[0]);
ret = clEnqueueWriteBuffer(command_queue, objHISTO, CL_FALSE, 0, HTAM * sizeof(float), histo, 0, NULL, &write_event[1]);
// Establecemos los argumentos del kernel
ret = clSetKernelArg(kernelHISTO, 0, sizeof(cl_mem), &objHISTO);
ret = clSetKernelArg(kernelHISTO, 1, sizeof(cl_mem), &objFRAME);
ret = clSetKernelArg(kernelHISTO, 2, sizeof(int), &hsv_frame->widthStep);
ret = clSetKernelArg(kernelHISTO, 3, sizeof(int), &hsv_frame->height);
ret = clSetKernelArg(kernelHISTO, 4, sizeof(int), &hsv_frame->width);
// Ejecutamos el kernel. 128 work-items por cada work-group o unidad de cómputo
ret = clEnqueueNDRangeKernel(command_queue, kernelHISTO, 1, NULL, &global_work_size, &local_work_size, 2, write_event, &kernel_event);
// Leemos los resultados
ret = clEnqueueReadBuffer(command_queue, objHISTO, CL_FALSE, 0, HTAM * sizeof(float), histo, 1, &kernel_event, &finish_event);
// Esperamos a que termine de leer los resultados
clWaitForEvents(1, &finish_event);
// Obtenemos el tiempo del kernel y de las transferencias Pcie
cl_ulong totalKernel = getStartEndTime(kernel_event);
cl_ulong totalPcie = getStartEndTime(write_event[0]) + getStartEndTime(write_event[1]) + getStartEndTime(finish_event);
cvReleaseImage( &hsv_frame );
// Recordar que frame no se puede liberar debido al cvQueryFrame
while( frame = cvQueryFrame( video ) ) {
hsv_frame = cvCreateImage( cvGetSize(frame), IPL_DEPTH_32F, 3 );
cvConvertScale( frame, hsv_frame, 1.0 / 255.0, 0 );
cvCvtColor( hsv_frame, hsv_frame, CV_BGR2HSV );
memset(histo, 0.0f, HTAM * sizeof(float));
ret = clEnqueueWriteBuffer(command_queue, objFRAME, CL_FALSE, 0, hsv_frame->imageSize, hsv_frame->imageData, 0, NULL, &write_event[0]);
ret = clSetKernelArg(kernelHISTO, 0, sizeof(cl_mem), &objHISTO);
ret = clSetKernelArg(kernelHISTO, 1, sizeof(cl_mem), &objFRAME);
ret = clSetKernelArg(kernelHISTO, 2, sizeof(int), &hsv_frame->widthStep);
ret = clSetKernelArg(kernelHISTO, 3, sizeof(int), &hsv_frame->height);
ret = clSetKernelArg(kernelHISTO, 4, sizeof(int), &hsv_frame->width);
ret = clEnqueueNDRangeKernel(command_queue, kernelHISTO, 1, NULL, &global_work_size, &local_work_size, 2, write_event, &kernel_event);
ret = clEnqueueReadBuffer(command_queue, objHISTO, CL_FALSE, 0, HTAM * sizeof(float), histo, 1, &kernel_event, &finish_event);
clWaitForEvents(1, &finish_event);
totalKernel += getStartEndTime(kernel_event);
totalPcie += (getStartEndTime(write_event[0]) + getStartEndTime(write_event[1]) + getStartEndTime(finish_event));
cvReleaseImage( &hsv_frame );
}
const double end_time = getCurrentTimestamp();
// Obtenemos el tiempo consumido por el programa, el kernel y las transferencias de memoria
printf("\nTiempo total del programa: %0.3f ms\n", (end_time - start_time) * 1e3);
printf("Tiempo total consumido por el kernel: %0.3f ms\n", double(totalKernel) * 1e-6);
printf("Tiempo total consumido en transferencias Pcie: %0.3f ms\n", double(totalPcie) * 1e-6);
// Liberamos todos los recursos usados (kernels, frames y objetos OpenCL)
clReleaseEvent(kernel_event);
clReleaseEvent(finish_event);
clReleaseEvent(write_event[0]);
clReleaseEvent(write_event[1]);
cvReleaseCapture( &video );
clReleaseMemObject(objFRAME);
clReleaseMemObject(objHISTO);
clReleaseKernel(kernelHISTO);
clReleaseCommandQueue(command_queue);
clReleaseProgram(program);
clReleaseContext(context);
}
