在 OpenCl 中模拟动态内存分配

Question

我遇到了一个让我发疯的问题。 我需要在 OpenCl 内核中模拟动态内存分配。在这方面，我在 *.cl 文件中定义了以下 malloc 函数：

 __global void* malloc(size_t size, __global byte *heap, __global uint *next)
{
  uint index = atomic_add(next, size);
  return heap+index;
}

在宿主程序中，我为这个虚拟堆动态地分配了一个 cl_uchar 类型的大数组，如下所示：

int MAX_NUM_OF_HEADERS_PROCESSED_IN_PARALLEL = 1000;
cl_uchar* heap = new cl_byte[1000000];
cl_uint  *next  =  new cl_uint;
*next = 0;
cl_uint * test_result =
        new cl_uint[MAX_NUM_OF_HEADERS_PROCESSED_IN_PARALLEL];
cl_mem memory[3]= { 0, 0, 0};
cl_int error;

memory[0] = clCreateBuffer(GPU_context,
CL_MEM_READ_WRITE, sizeof(cl_uchar) * MAX_HEAP_SIZE, NULL,
NULL);

memory[1] = clCreateBuffer(GPU_context, CL_MEM_READ_WRITE, sizeof(cl_uint), NULL,
        &error);

memory[2] = clCreateBuffer(GPU_context, CL_MEM_READ_WRITE,
            sizeof(cl_uint) * MAX_NUM_OF_HEADERS_PROCESSED_IN_PARALLEL, NULL,
            &error);
clEnqueueWriteBuffer(command_queue, memory[0], CL_TRUE, 0,
        sizeof(cl_uchar) * MAX_HEAP_SIZE, heap, 0, NULL, NULL);

clEnqueueWriteBuffer(command_queue, memory[1], CL_TRUE, 0, sizeof(cl_uint),
        next, 0, NULL, NULL);
error = 0;
error |= clSetKernelArg(kernel, 0, sizeof(cl_mem), &memory[0]);
error |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &memory[1]);
error |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &memory[2]);

size_t globalWorkSize[1] = { MAX_NUM_OF_HEADERS_PROCESSED_IN_PARALLEL };
size_t localWorkSize[1] = { 1 };


error = 0;
error = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL,
        globalWorkSize, localWorkSize, 0, NULL, NULL);

我还有以下内核：

__kernel void packet_routing2(__global byte* heap_, __global uint* next, __global uint* test_result){

    int gid = get_global_id(0);

    __global uint*xx[100];

    for ( int i = 0 ; i < 100; i ++)
    {
        xx[i] = (__global uint*) malloc(sizeof(uint),heap_,next);
        *xx[i] = i*gid;

    result[gid] = *(xx[0]);
}   

我在运行程序时遇到如下错误：

" %27 = load i32 addrspace(1)* %26, align 4, !tbaa !17
Illegal pointer which is not from a valid memory space.
Aborting..."

你能帮我解决这个问题吗？我还发现如果 xx 只有 10 个元素，而不是 100 个，代码运行良好！！！！

Answer 1

编辑： 最简单的解决方案：在 malloc 之前为 'size' 添加一个填充值，以便所有结构类型（大小小于 max-padding）接收必要的对齐条件。

0=内存中的结构占用空间

*=堆

_=填充

***000_____*****0000____****0_______****00000___*****0000000_*******00______***
      |
      v
 save this unused padded memory space in its thread to use later.

重要的是，第一个/起始地址值需要满足最大对齐要求。如果有一个 256 字节长的结构，它应该以 256 的倍数开始。

struct size      malloc size    minimum 'next' value (address, not offset)
  1-4                 4            multiple of 4
  5-8                 8            multiple of 8
  9-16                16           multiple of 16
  17-32               32            32*k
  33-64               64            64*k

如果有 64 字节的结构，即使是 int 现在也需要 64 字节的 malloc 大小。也许您可以在每个线程本地保存该值以使用剩余的未使用区域。

所以它不会产生对齐错误，并且对于那些没有的可能会更快。

float3 本身也需要 16 字节。