Я пытаюсь выполнить базовый алгоритм умножения матриц c с использованием OpenCL. Обе матрицы должны иметь одинаковые размеры (SIZE x SIZE), поэтому я определил проблему как двумерную, с глобальным размером SIZE x SIZE, и я тестирую, что происходит с разными локальными размерами.
Ядро написано следующим образом:
__kernel void matmul(
__global unsigned int *a,
__global unsigned int *b,
__global unsigned int *c
) {
int row, col, i, size;
unsigned int dot;
row = get_global_id(0);
col = get_global_id(1);
size = get_global_size(0);
dot = 0;
for (i = 0; i < size; i++) {
dot += a[row * size + i] * b[i * size + col];
}
c[row * size + col] = dot;
}
Хорошо работает, если глобальный и локальный размеры установлены на 1024 x 1024 и 1 x 1 соответственно. Однако оказывается, что если локальный размер равен 2 x 2 или 4 x 4, я получаю неверные результаты при умножении. Теперь для локального размера используется кратное 8, например 8 x 8, 16 x 16, ... В умножении нет ошибки. Почему это происходит?
Я не знаю, проблема в программировании ядра или я плохо понимаю, что можно сделать с рабочей группой или рабочим элементом .
Полный код хоста следующий:
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <sys/time.h>
#include <CL/opencl.h>
#define SIZE (1024)
#define WORKITEMS (4096)
#define LOG_SIZE (2048)
int main(int argc, char *argv[]) {
int i, j, k, size, errors;
// Host memory
cl_uint *a_host = NULL;
cl_uint *b_host = NULL;
cl_uint *c_host = NULL;
cl_uint ref_dot;
// Device memory
cl_mem a_device;
cl_mem b_device;
cl_mem c_device;
// Performance measurements
struct timeval t0, tf;
float ts, tp, tb;
// OpenCL variables
FILE *f;
size_t f_size;
size_t global[3] = {0}, local[3] = {0};
char *buffer = NULL;
cl_int ret;
cl_platform_id platform;
cl_device_id device;
cl_context context;
cl_command_queue queue;
cl_program program;
cl_kernel kernel;
// [1] Initialize application
// Read command line arguments to configure run
size = (argc > 1) ? atoi(argv[1]) : SIZE;
printf("Matrix multiplication with OpenCL (Size = %d)\n", size);
// Allocate memory for host variables
a_host = malloc(size * size * sizeof *a_host);
b_host = malloc(size * size * sizeof *b_host);
c_host = malloc(size * size * sizeof *c_host);
// Initialize input arrays
for (i = 0; i < size; i++) {
for (j = 0; j < size; j++) {
a_host[i * size + j] = rand();
b_host[i * size + j] = rand();
}
}
// [2] Initialize OpenCL environment
// Get platform
ret = clGetPlatformIDs(1, &platform, NULL);
// Get device
ret = clGetDeviceIDs(platform, CL_DEVICE_TYPE_CPU, 1, &device, NULL);
// Create context
context = clCreateContext(0, 1, &device, NULL, NULL, &ret);
// Create command queue
queue = clCreateCommandQueueWithProperties(context, device, 0, &ret);
// [3] Compile OpenCL kernel
f = fopen("kernel.cl", "rb");
fseek(f, 0, SEEK_END);
f_size = ftell(f);
rewind(f);
// Read file into memory
buffer = malloc(f_size + 1);
buffer[f_size] = '\0';
fread(buffer, sizeof(char), f_size, f);
fclose(f);
// Create program
printf("<OpenCL> Kernel source:\n%s", buffer);
program = clCreateProgramWithSource(context, 1, (const char **) &buffer, &f_size, &ret);
// Build program
printf("<OpenCL> Building kernel...\n");
gettimeofday(&t0, NULL);
ret = clBuildProgram(program, 0, NULL, "-cl-std=CL2.0", NULL, NULL);
gettimeofday(&tf, NULL);
tb = ((tf.tv_sec - t0.tv_sec) * 1000.0) + ((tf.tv_usec - t0.tv_usec) / 1000.0);
printf("Build time: %.3f ms\n", tb);
// Print build log (optional)
char log[LOG_SIZE];
ret = clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, LOG_SIZE, log, NULL);
printf("<OpenCL> Kernel build log:\n%s\n", log);
// [4] Configure OpenCL kernel
// Create kernel
kernel = clCreateKernel(program, "matmul", &ret);
// Create device buffers
a_device = clCreateBuffer(context, CL_MEM_READ_ONLY, size * size * sizeof *a_host, NULL, &ret);
b_device = clCreateBuffer(context, CL_MEM_READ_ONLY, size * size * sizeof *b_host, NULL, &ret);
c_device = clCreateBuffer(context, CL_MEM_WRITE_ONLY, size * size * sizeof *c_host, NULL, &ret);
// Set kernel parameters
ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), &a_device);
ret |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &b_device);
ret |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &c_device);
// [5] Execute kernel
printf("<OpenCL> Executing kernel...\n");
gettimeofday(&t0, NULL);
// Write data from host to device
ret = clEnqueueWriteBuffer(queue, a_device, CL_TRUE, 0, size * size * sizeof *a_host, a_host, 0, NULL, NULL);
ret |= clEnqueueWriteBuffer(queue, b_device, CL_TRUE, 0, size * size * sizeof *b_host, b_host, 0, NULL, NULL);
// Enqueue kernel for execution
global[0] = size;
global[1] = size;
local[0] = 2;
local[1] = 2;
ret = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, global, local, 0, NULL, NULL);
// Wait for kernel to finish
ret = clFinish(queue);
// Read data from device to host
ret = clEnqueueReadBuffer(queue, c_device, CL_TRUE, 0, size * size * sizeof *c_host, c_host, 0, NULL, NULL);
gettimeofday(&tf, NULL);
tp = ((tf.tv_sec - t0.tv_sec) * 1000.0) + ((tf.tv_usec - t0.tv_usec) / 1000.0);
printf("[PAR] Execution time: %.3f ms\n", tp);
// [6] Print results, perform checks
// Compute golden reference and check errors
gettimeofday(&t0, NULL);
errors = 0;
for (i = 0; i < size; i++) {
for (j = 0; j < size; j++) {
ref_dot = 0;
for (k = 0; k < size; k++) {
ref_dot += a_host[i * size + k] * b_host[k * size + j];
}
if (ref_dot != c_host[i * size + j]) {
errors++;
}
}
}
gettimeofday(&tf, NULL);
ts = ((tf.tv_sec - t0.tv_sec) * 1000.0) + ((tf.tv_usec - t0.tv_usec) / 1000.0);
printf("[SEQ] Execution time : %.3f ms\n", ts);
printf("Found %d error%s\n", errors, (errors == 1) ? "" : "s");
// [7] Cleanup system
// Cleanup host variables
free(a_host);
free(b_host);
free(c_host);
free(buffer);
// Cleanup OpenCL
clReleaseMemObject(a_device);
clReleaseMemObject(b_device);
clReleaseMemObject(c_device);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseCommandQueue(queue);
clReleaseContext(context);
return 0;
}