【问题标题】:Passing Class to a Kernel in Intel Opencl在英特尔 Opencl 中将类传递给内核
【发布时间】:2017-04-17 08:32:30
【问题描述】:

过去几周我一直在研究 c/c++ OpenCL 解决方案。对于我的解决方案,我需要将一个类从我的 CPU(主机)传递到 GPU(设备)。当我尝试将类作为参数传递时,它会给出错误“未知类型标识符类”。我怀疑英特尔平台上的 OpenCL 是否允许我们将类传递给内核或任何解决方法都可以使用。在 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 未知类型名称“测试”
2) 错误 Tempalte.CL 预期为 ';'在顶级声明符之后
3) 错误 Tempalte.CL 程序范围变量需要在常量地址空间中声明
4) 错误 Tempalte.CL unknown type name 'class'

查询:

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

【问题讨论】:

  • 你需要 sycl 。它让您几乎可以使用纯粹的 C++。您可以发布出错的代码吗?您使用哪个版本的 opencl?
  • @huseyin tugrul buyukisik,我已根据您的要求添加了代码。让我知道你能从中找到什么。提前致谢。

标签: class opencl intel


【解决方案1】:

OpenCL 使用 C99。因此,您可以将结构而非类传递给内核。

正如 huseyin tugrul buyukisik 所说,您可以使用支持 c++14(或更高版本)的 SYCL。

或者,如果您想同时支持 NVIDIA® CUDA™ 和 OpenCL,您可以只在 NVIDIA® CUDA™ 中编写它,然后使用 https://github.com/hughperkins/cuda-on-cl 在 OpenCL 1.2 GPU 设备上运行 NVIDIA® CUDA™ 代码。全面披露:我是 cuda-on-cl 的作者,目前它还有一点正在进行中。它确实有效,但有一些警告/限制。它可以处理成熟的 C++11、模板、类等。例如,它可以用于在 OpenCL 1.2 GPU 上编译和运行 Eigen https://bitbucket.org/hughperkins/eigen/src/eigen-cl/unsupported/test/cuda-on-cl/?at=eigen-cl

【讨论】:

  • 感谢您的帮助。我想我现在理解了这个概念。
  • 我假设您的意思是 NVIDIA®。他们每个人都有自己的利基。英特尔在集成 GPU 方面很强大,并且在最近的许多计算机中都提供了这些,这些计算机通常使用英特尔 CPU。 NVIDIA 发明了 CUDA,据我所知,它是第一家将 GPU 用作数学处理器的公司。
  • 感谢您提供的信息。因此,与英特尔相比,选择 NVIDIA 是更好的选择。
【解决方案2】:

如果 sycl(和 Hugh Perkins 的好解决方案)不适合您,并且如果您的类没有任何方法,则可以改用结构(复制到设备时序列化为字节数组) :

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 struct 具有总共 20 个字节的字段,但它在设备端可能是 32 个字节(因为某些规则强制它在内存中是 2 的幂),您应该从主机端确认它并相应地放置数据对齐和可变大小。甚至没有提到“一次编写,到处运行”的痛苦处理的字节顺序。所以你应该只在你优化的计算机上运行它。

在结构顶部打包最大的字段,在底部添加较小的字段。也将它们的结构内对齐计算为 2 的幂!。密切关注 float3、int3 和类似的未记录良好的实现,因为它们可能会或可能不会在后台使用 float4、int4。如果全局内存访问的性能对您来说并不重要,您可以简单地为每个小于该值的结构选择一个像 N 这样的大数字,以便为简单起见,并将相对字节寻址放在结构开始字节。例如顶部Warrior 结构中hp 字段的字节地址(在一个压缩的4 字节到单个int 中)。然后设备端可以查询字段从哪个字节开始。 (字节序可以使它更棘手,所以不要对纯结构使用缓冲区复制)

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

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

【讨论】:

    猜你喜欢
    • 1970-01-01
    • 1970-01-01
    • 1970-01-01
    • 1970-01-01
    • 2015-10-04
    • 1970-01-01
    • 1970-01-01
    • 1970-01-01
    • 1970-01-01
    相关资源
    最近更新 更多