PyOpenCl 投射全局数据答案

【问题标题】：PyOpenCl casting global dataPyOpenCl 投射全局数据
【发布时间】：2013-09-30 13:44:03
【问题描述】：

我使用 pyOpenCl 2013.1，我的代码在 nVidia GPU、AMD CPU 和 AMD GPU 上崩溃，但在 Intel CPU 上运行。

使用 nvidia GPU，调用内核后在 queue.finish 上引发错误。

LogicError: clFinish failed: invalid command queue

我在下面的 sn-p 中的第 48 行找到了原因。

1:  typedef struct {
2:    int global_index;
3:    int local_index;
4:    float speed_limit;
5:    float width;
6:  } segment_t;
7:  
8:  typedef struct {
9:      int item_count;
10:     segment_t first_item;
11: } segment_list_t;
12: 
13: void explode_segment_list_t(segment_list_t* list, segment_t** array)
14: {
15:     array[0] = &(list->first_item);
16: }
17: 
18: 
19: 
20: /*
21:  * ro_data is read-only array of 3316 byte (829 int)
22:  * wo_data is write-only array of 3316 byte (829 int)
23:  */
24: __kernel void test_kernel(global int* ro_data, global int* wo_data)
25: {
26:     unsigned int i = get_global_id(0);
27:     
28:     // copy uncasted, primitive types
29:     for(int index = 0; index < ro_data[0]; index++)
30:         wo_data[index] = ro_data[index];                 // this works
31:     
32:     // access casted local struct
33:     int temp[829] = {0};
34:     segment_list_t* casted_temp_list = (segment_list_t*)temp;
35:     casted_temp_list->item_count = 1337;                 // this works
36:     // do more tests
37:     segment_t* casted_temp_array;
38:     explode_segment_list_t(casted_temp_list, &casted_temp_array);
39:     casted_temp_array[1].global_index = 1;
40:     casted_temp_array[2].global_index = 2;               // this works
41:     
42:     // copy local data to global data
43:     for(int index = 0; index < ro_data[0]; index++)
44:         wo_data[index] = temp[index];                    // this works
45:         
46:     // access casted global memory
47:     segment_list_t* casted_wo_data = (segment_list_t*)wo_data;
48:     casted_wo_data->item_count = 42;                     // this fails on GPU but works on CPU
49:     
50: }

Executable pyopencl.capture_call here

一个丑陋的内存浪费修复方法是：分配一个本地数组，复制数据然后强制转换。但我敢肯定，我在这里做错了什么……但是什么？

感谢您的帮助！

编辑：在 AMD 设备（CPU 和 GPU）上失败并显示更多信息：

 *error: invalid type conversion
      segment_list_t* casted_wo_data = (segment_list_t*)wo_data;*

【问题讨论】：

标签： casting opencl pyopencl

【解决方案1】：

来自 AMD SDK 的错误消息有助于找到原因。

如a forum posting 中所述，强制类型转换必须与源变量具有相同的内存位置注释。

在这个例子中：

segment_list_t* casted_wo_data = (segment_list_t*)wo_data;

必须替换为：

global segment_list_t* casted_wo_data = (global segment_list_t*)wo_data;

【讨论】：

很高兴知道。我尽量避免强制转换，但如果你真的需要它，这是一个非常有用的提示。

【解决方案2】：

我不知道你做错了什么。

但我可以告诉你，你做事的方式不是正确的做法，它会给你带来很多麻烦。

您应该将内核的输入和输出标记为segment_t 和segment_list_t。您可以在内核代码和 Python 代码中声明这些类型定义，然后创建适当类型的缓冲区。这样您就不需要投射任何东西，并且可以避免任何可能的越界错误。

【讨论】：

感谢您的快速回复。我的问题是，在我的 python 代码中 segment_list 是一个包含所有数组项的结构。这进一步用于其他类型。我真正想要的是具有不同嵌套类型的复杂结构，它在 numpy.array 中已经可以正常工作了。但是在 GPU 上，我必须通过铸造它们来分离不同的组件。在这个 sn-p 中，我稍微降低了复杂性。
只要具有嵌套类型的结构不包含任何指针（可变长度数组），您就可以在 OpenCL 和 Python 代码中表达该结构。
我不认为序列化是这里的问题，因为有问题的 sn-p 没有应用嵌套。