OpenCL：GPU 上的类型转换

Question

我将数据存储在一个 char 数组中，我需要从那里读取 float 和 int 变量。此代码在 CPU:

上运行良好

global float *p;
p = (global float*)get_pointer_to_the_field(char_array, index);
*p += 10;

但在 GPU 上我得到错误 -5：CL_OUT_OF_RESOURCES。读数本身有效，但对值进行某些操作（在本例中加 10）会导致错误。我该如何解决？

更新：

这适用于 GPU：

float f = *p;
f += 10;

但是，我仍然无法将这个值写回数组。

这是内核：

global void write_value(global char *data, int tuple_pos, global char *field_value, 
                    int which_field, global int offsets[], global int *num_of_attributes) {

    int tuple_size = offsets[*num_of_attributes];
    global char *offset = data + tuple_pos * tuple_size;
    offset += offsets[which_field];

    memcpy(offset, field_value, (offsets[which_field+1] - offsets[which_field]));
}

global char *read_value(global char *data, int tuple_pos, 
                    int which_field, global int offsets[], global int *num_of_attributes) {
    int tuple_size = offsets[*num_of_attributes];
    global char *offset = data + tuple_pos * tuple_size;
    offset += offsets[which_field];
    return offset;
}

kernel void update_single_value(global char* input_data, global int* pos, global int offsets[], 
                            global int *num_of_attributes, global char* types) {
    int g_id = get_global_id(1);
    int attr_id = get_global_id(0);
    int index = pos[g_id];

    if (types[attr_id] == 'f') { // if float

        global float *p;
        p = (global float*)read_value(input_data, index, attr_id, offsets, num_of_attributes);
        float f = *p;
        f += 10;
        //*p += 10; // not working on GPU
    } 
    else if (types[attr_id] == 'i') { // if int
        global int *p;
        p = (global int*)read_value(input_data, index, attr_id, offsets, num_of_attributes);
        int i = *p;
        i += 10;
        //*p += 10;
    }
    else { // if char
        write_value(input_data, index, read_value(input_data, index, attr_id, offsets, num_of_attributes), attr_id, offsets, num_of_attributes);
    }
}

它更新了table的元组的值，int和float增加了10，char字段只是被替换为相同的内容。

Answer 1

您要启用 byte_addressable_store extension 吗？据我所知，按字节写入全局内存在 OpenCL 中没有明确定义，除非您启用它。（您需要检查您的实施是否支持扩展。）

您可能还想考虑在内核参数中使用 "correct" 类型 - 这可能有助于编译器生成更高效的代码。如果类型可以动态变化，您也许可以尝试使用联合类型（或结构类型中的联合字段），尽管我自己还没有用 OpenCL 对此进行测试。

Answer 2

原来问题的出现是因为char数组中的int和float值不是4字节对齐的。当我写入

这样的地址时

offset = data + tuple_pos*4; // or 8, 16 etc

一切正常。但是，以下原因会导致错误：

offset = data + tuple_pos*3; // or any other number not divisible by 4

这意味着要么我应该改变整个设计并以其他方式存储值，要么将 "empty" 字节添加到 char 数组以使 int 和 float 值对齐 4 个字节（这不是一个真正的好的解决方案）。

OpenCL：GPU 上的类型转换

OpenCL: type casting on GPU

c

pointers

gpu

opencl