OpenCL实现Floyd算法n≥268时调用clEnqueueWriteBuffer触发内存访问错误

我用OpenCL实现Floyd算法，当n<268时程序运行正常，但n≥268时，循环中调用clEnqueueWriteBuffer（针对buffer_distances）会触发“Access violation reading location”错误。

主机端初始化代码 #

int n;
printf("enter n value: ");
scanf("%d", &n);
printf("\n");
int n2 = n * n;
int matSize = n2 * sizeof(int*);
int* graphe = malloc(sizeof(int) * n2);
int* distances = malloc(sizeof(int) * n2);
//mat[i,j] => mat[i*n + j]
if (graphe == NULL)
    printf("malloc failed\n");
init_graphe(graphe, n);
copy(graphe, distances, n);

OpenCL环境初始化代码 #

char* programSource = load_kernel("kernel.cl");

cl_int status;
// STEP 1: Discover and initialize the platforms

cl_uint numPlatforms = 0;

cl_platform_id* platforms = NULL;

status = clGetPlatformIDs(0, NULL, &numPlatforms);
printf("Number of platforms = %d\n", numPlatforms);

platforms = (cl_platform_id*)malloc(numPlatforms * sizeof(cl_platform_id));

status = clGetPlatformIDs(numPlatforms, platforms, NULL);

char Name[1000];
clGetPlatformInfo(platforms[0], CL_PLATFORM_NAME, sizeof(Name), Name, NULL);
printf("Name of platform : %s\n", Name);
fflush(stdout);

// STEP 2: Discover and initialize the devices

cl_uint numDevices = 0;
cl_device_id* devices = NULL;

status = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, 0, NULL, &numDevices);


printf("Number of devices = %d\n", (int)numDevices);

devices = (cl_device_id*)malloc(numDevices * sizeof(cl_device_id));

status = clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_ALL, numDevices, devices, NULL);


for (int i = 0; i < numDevices; i++) {
    clGetDeviceInfo(devices[i], CL_DEVICE_NAME, sizeof(Name), Name, NULL);
    printf("Name of device %d: %s\n\n", i, Name);
}

// STEP 3: Create a context
fflush(stdout);

cl_context context = NULL;

context = clCreateContext(NULL, numDevices, devices, NULL, NULL, &status);

// STEP 4: Create a command queue
fflush(stdout);
cl_command_queue cmdQueue;

cmdQueue = clCreateCommandQueue(context, devices[0], 0, &status);

// STEP 5: Create device buffers
fflush(stdout);

cl_mem buffer_graphe;
cl_mem buffer_n;
cl_mem buffer_distances;
cl_mem buffer_k;

buffer_graphe = clCreateBuffer(context, CL_MEM_READ_WRITE, matSize, NULL, &status);
buffer_n = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int), NULL, &status);
buffer_distances = clCreateBuffer(context, CL_MEM_READ_WRITE, matSize, NULL, &status);
buffer_k = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int), NULL, &status);
fflush(stdout);

// STEP 6: Create and compile the program
cl_program program = clCreateProgramWithSource(context, 1, (const char**)&programSource, NULL, &status);
printf("Compilation\n");
fflush(stdout);
status = clBuildProgram(program, numDevices, devices, NULL, NULL, NULL);

// STEP 8: Create the kernel
cl_kernel kernel = NULL;
fflush(stdout);
kernel = clCreateKernel(program, "floyd", &status);

size_t globalWorkSize[2] = { n, n };
size_t localWorkSize[3] = { 20,20 };

内核执行循环代码 #

clock_t start = clock();
int k;
for (k = 0; k < n; k++) {
    status = clEnqueueWriteBuffer(cmdQueue, buffer_graphe, CL_TRUE, 0, matSize, graphe, 0, NULL, NULL);
    status = clEnqueueWriteBuffer(cmdQueue, buffer_n, CL_TRUE, 0, sizeof(int), &n, 0, NULL, NULL);
    status = clEnqueueWriteBuffer(cmdQueue, buffer_distances, CL_TRUE, 0, matSize, distances, 0, NULL, NULL);
    status = clEnqueueWriteBuffer(cmdQueue, buffer_k, CL_TRUE, 0, sizeof(int), &k, 0, NULL, NULL);

    status = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&buffer_graphe);
    status = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&buffer_n);
    status = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*)&buffer_distances);
    status = clSetKernelArg(kernel, 3, sizeof(cl_mem), (void*)&buffer_k);

    status = clEnqueueNDRangeKernel(cmdQueue, kernel, 2, NULL, globalWorkSize, NULL, 0, NULL, NULL);
    clFinish(cmdQueue);  

    status = clEnqueueReadBuffer(cmdQueue, buffer_distances, CL_TRUE, 0, matSize, distances, 0, NULL, NULL);
    clFinish(cmdQueue);
}

OpenCL内核代码 #

void kernel floyd(global int* graphe, global int* n, global int* distances, global int* k)
{
    int i = get_global_id(0);
    int j = get_global_id(1);

    int ij = i * (*n) + j;
    int ik = i * (*n) + (*k);
    int kj = (*k) * (*n) + j;

    if (distances[ik] + distances[kj] < distances[ij]) {
        distances[ij] = distances[ik] + distances[kj];
    }
}

核心错误原因 #

代码中计算矩阵内存大小的matSize存在致命错误：

int matSize = n2 * sizeof(int*);

这里误用了sizeof(int*)（指针大小，通常为4或8字节），而实际应该使用sizeof(int)（矩阵元素的类型大小）。当n较小时，n2 * sizeof(int*)的计算结果大于实际需要的内存（比如sizeof(int*)=8、sizeof(int)=4时，分配的内存是实际需求的2倍），因此不会触发越界；但当n≥268时，n2 * sizeof(int)超过了malloc分配的内存空间（因为graphe和distances是用sizeof(int)*n2分配的），导致clEnqueueWriteBuffer尝试写入超出分配范围的内存，触发访问违例。

修复步骤 #

1. 修正内存大小计算
   将matSize的计算改为：

int matSize = n2 * sizeof(int);

确保主机端分配的内存和OpenCL设备缓冲区的大小完全匹配。

2. 优化不必要的数据传输
   当前循环中每次都重复写入graphe和n，这两个值在整个Floyd算法过程中是固定的，可以将它们的写入操作移到循环外，减少冗余数据传输，提升性能：

// 移到循环前执行一次
status = clEnqueueWriteBuffer(cmdQueue, buffer_graphe, CL_TRUE, 0, matSize, graphe, 0, NULL, NULL);
status = clEnqueueWriteBuffer(cmdQueue, buffer_n, CL_TRUE, 0, sizeof(int), &n, 0, NULL, NULL);

clock_t start = clock();
int k;
for (k = 0; k < n; k++) {
    status = clEnqueueWriteBuffer(cmdQueue, buffer_distances, CL_TRUE, 0, matSize, distances, 0, NULL, NULL);
    status = clEnqueueWriteBuffer(cmdQueue, buffer_k, CL_TRUE, 0, sizeof(int), &k, 0, NULL, NULL);

    status = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&buffer_graphe);
    status = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&buffer_n);
    status = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*)&buffer_distances);
    status = clSetKernelArg(kernel, 3, sizeof(cl_mem), (void*)&buffer_k);

    status = clEnqueueNDRangeKernel(cmdQueue, kernel, 2, NULL, globalWorkSize, NULL, 0, NULL, NULL);
    clFinish(cmdQueue);  

    status = clEnqueueReadBuffer(cmdQueue, buffer_distances, CL_TRUE, 0, matSize, distances, 0, NULL, NULL);
    clFinish(cmdQueue);
}

3. 移除未使用的内核参数
   内核中graphe参数并未被使用，可以从内核参数中移除，减少不必要的资源占用。修改后的内核代码：

void kernel floyd(global int* n, global int* distances, global int* k)
{
    int i = get_global_id(0);
    int j = get_global_id(1);

    int ij = i * (*n) + j;
    int ik = i * (*n) + (*k);
    int kj = (*k) * (*n) + j;

    if (distances[ik] + distances[kj] < distances[ij]) {
        distances[ij] = distances[ik] + distances[kj];
    }
}

同时主机端对应的clSetKernelArg调用也要移除对buffer_graphe的设置。

4. 添加错误检查
   在所有OpenCL API调用后添加错误检查，方便快速定位问题，比如：

status = clCreateBuffer(context, CL_MEM_READ_WRITE, matSize, NULL, &status);
if (status != CL_SUCCESS) {
    printf("clCreateBuffer failed with code %d\n", status);
    exit(1);
}

内容的提问来源于stack exchange，提问作者alex