将 Class 传递给英特尔 Opencl 中的内核

Question

过去几周我一直在研究 c/c++ OpenCL 解决方案。对于我的解决方案，我需要将 class 从我的 CPU（主机）传递到 GPU（设备）。当我尝试将 class 作为参数传递时，它给出了一个错误 "Unknown Type-Identifier Class"。我怀疑英特尔平台上的 OpenCL 是否允许我们将 class 传递给内核，或者是否有任何变通方法适用于它。在 CUDA 中，我看到了一些示例，它在该平台上运行得非常好。但是，关于 OpenCL，我无法找到与此查询相关的任何参考资料和示例。我将非常感谢有关此问题的任何帮助。我在英特尔网站上发布了同样的问题，但无济于事。如果有人愿意帮助我了解我哪里出错了或者我应该如何处理这个问题，我将非常感谢你。

//主机端代码

#include<stdio.h>
#include<iostream>
#include"CL/cl.h"


class test
{
public:
    cl_int a;
    cl_char b;
};

int main()
{
    test *tempstruct = new test;

    cl_platform_id platfrom_id;
    cl_device_id device_id; // compute device id 
    cl_context context; // compute context
    cl_command_queue commands; // compute command queue
    cl_program program; // compute program
    cl_kernel kernel; // compute kernel

    int err;

    err = clGetPlatformIDs(1, &platfrom_id, NULL);
    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to create a platfrom group!\n");
        return -1;
    }

    err = clGetDeviceIDs(platfrom_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);

    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to create a device group!\n");
        return -1;
    }

    context = clCreateContext(0, 1, &device_id, NULL, NULL, NULL);

    if (!context)
    {
        printf("Error: Failed to create a compute context!\n");
        return -1;
    }

    commands = clCreateCommandQueue(context, device_id, 0, NULL);

    if (!commands)
    {
        printf("Error: Failed to create a command commands!\n");
        return -1;
    }

    #define MAX_SOURCE_SIZE (0x100000) 
    FILE *fp, *fp1;
    char filename[] = "Template.cl";
    fp = fopen(filename, "r");
    if (fp == NULL)
    {
        printf("\n file  not found \n");
        return -1;
    }
    char * source_str = (char*)malloc(MAX_SOURCE_SIZE);
    size_t size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
    fclose(fp);

    cl_mem classobj = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, sizeof(tempstruct), &tempstruct, &err);

    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to allocate device memory!\n");
        return -1;
    }

    program = clCreateProgramWithSource(context, 1, (const char **)& source_str, (const size_t *)&size, &err);

    if (!program)
    {
        printf("Error: Failed to create program with source!\n");
        return -1;
    }

    err = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);

    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to build program executable!\n");
        return -1;
    }


    test  *resptr = (test *)clEnqueueMapBuffer(commands, classobj, CL_TRUE, CL_MAP_WRITE, NULL, sizeof(test), NULL, NULL, NULL, &err);
    // INITIALISATION OF CLASS

    tempstruct->a = 10;
    if (!resptr)
    {
        printf("Error: Failed to create enqueuemapbuffer!\n");
        return -1;
    }

    err = clEnqueueUnmapMemObject(commands, classobj, resptr, 0, NULL, NULL);

    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to write to source array!\n");
        return -1;
    }

    kernel = clCreateKernel(program, "CLASS", &err);
    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to create compute kernel!\n");
        return -1;
    }

    err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &classobj);

    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to set kernel arguments! %d\n", err);
        return -1;
    }

    size_t globalsize = 1;
    size_t local = 1;

    err = clEnqueueNDRangeKernel(commands, kernel, 1, NULL, &globalsize, &local, 0, NULL, NULL);
    if (err)
    {
        printf("Error: Failed to execute nd range!\n");
        return -1;
    }

    test  *resptr1 = (test *)clEnqueueMapBuffer(commands, classobj, CL_TRUE, CL_MAP_READ, NULL, sizeof(test), NULL, NULL, NULL, &err);

    err = clEnqueueUnmapMemObject(commands, classobj, resptr1, 0, NULL, NULL);
    if (err != CL_SUCCESS)
    {
        printf("Error: Failed to read output array! %d\n", err);
        return -1;
    }

    // again i am printing the class value

    printf("\n in cpu side = %d\n", tempstruct->a);

}


//HOST END


//DEVICE SIDE(KERNEL CODE)

 // filename :  Template.cl

class test
{
public:
    cl_int a;
    cl_char b;
};

__kernel void CLASS(__global test *inclass )
{
   inclass->a = 10;
    printf("\n in kernel side = %d \n",inclass->a);
}

//内核结束

错误：

我只在内核端遇到这些所有错误

1) 错误 Tempalte.CL 未知类型名称 'test'
2) 错误 Tempalte.CL 预期 ';'在顶级声明符之后
3) 错误Tempalte.CL program scope variables are required to be declared in constant address space
4) 错误Tempalte.CL未知类型名称'class'

查询：

问）我的主要问题是如何在英特尔架构中将 CLASS 传递给内核
* 我已成功将 class 传递给 AMD 内核。每当我在 Intel 端尝试使用相同的代码时，它都会显示上述四个错误。
* 是否有任何替代方法可以将 class 传递给英特尔内核，或者是否可以将 class 传递给英特尔架构内核？

Answer 1

OpenCL 使用 C99。所以你可以将结构传递给内核，但不能传递类。

正如 huseyin tugrul buyukisik 所说，您可以使用支持 c++14（或左右）的 SYCL。

或者，如果您想同时支持 NVIDIA® CUDA™ 和 OpenCL，您可以只在 NVIDIA® CUDA™ 中编写，然后使用 https://github.com/hughperkins/cuda-on-cl to run the NVIDIA® CUDA™ code on OpenCL 1.2 GPU devices. Full disclosure: I'm the author of cuda-on-cl, and it's a bit of a work-in-progress for now. It does work though, with some caveats/limitations. It can handle full-blown C++11, templates, classes etc. For example, it can be used to compile and run Eigen on OpenCL 1.2 GPUs https://bitbucket.org/hughperkins/eigen/src/eigen-cl/unsupported/test/cuda-on-cl/?at=eigen-cl

Answer 2

如果 sycl（和 Hugh Perkins 的不错的解决方案）不是您的选择，并且如果您的 class 没有任何方法，您可以改用结构（复制到设备时序列化为字节数组):

typedef struct Warrior_tag
{
    int id;
    float hp;
    int strength;
    int dexterity;
    int constitution;
} Warrior;

typedef struct Mage_tag
{
    int id;
    Warrior summoned_warriors[90];
} Mage; 
// should be more than 32*90 + 32*90 => 5760(2.8k *2) => 8192(4k*2) bytes
// because id + padding = 90 warriors or it doesn't work in runtime
// reversing order of fields should make it 4k + 4 bytes


__kernel void test0(__global Warrior * warriors)
{
    int id=get_global_id(0);
    Warrior battal_gazi = warriors[0];
    Warrior achilles = warriors[1];
    Warrior aramis = warriors[2];
    Warrior hattori_hanzo = warriors[3];
    Warrior ip_man = warriors[4];

    Mage nakano = (Mage){0,{battal_gazi, achilles}};
    Mage gul_dan = (Mage){0,{aramis , hattori_hanzo,ip_man  }};
}

然后您负责处理结构的对齐和大小。例如，Warrior 结构的字段总共有 20 个字节，但在设备端可能有 32 个字节（因为某些规则强制它在内存中为 2 的幂），您应该从主机端确认它并将数据相应地放入调整对齐方式和可变大小。更不用说字节序了，这对 "write once, run everywhere" 来说是一件很痛苦的事情。所以你应该运行它只在你的电脑里面优化过。

在结构顶部打包最大的字段，在底部添加较小的字段。也将它们的结构内对齐计算为 2 的幂！密切关注 float3、int3 和类似的没有很好记录的实现，因为它们可能会或可能不会在后台使用 float4、int4。如果全局内存访问的性能对您来说不重要，为了简单起见，您可以简单地 select 每个小于 N 的结构的大数字，并将相对字节寻址到结构的起始字节。例如顶部 Warrior 结构中 hp 字段的字节地址（将 4 字节打包成单个 int）。然后设备端可以查询字段从哪个字节开始。（字节序会使它更棘手，所以不要对纯结构使用缓冲区副本）

如果主机端的结构字段对齐不是一个选项：

将字段数组发送到构造函数内核（浮点数组 -> hp，整数数组 -> id）
使用内核在设备中构造（仅在设备端的缓冲区，Warrior 是从 hp、id、... 的数组构建的）
不再fiddle 对齐和大小，只需要使缓冲区足够大以容纳其中的所有结构。选择 32 * 勇士字节数对于勇士数组应该足够了。
当它工作时，return 在使用另一个内核将结构扩展到设备端的数组后，再次作为数组返回到主机端。

将 Class 传递给英特尔 Opencl 中的内核

Passing Class to a Kernel in Intel Opencl

class

intel

opencl