金属计算核的库达核等价

Question

我有一个简单的金属计算内核，我试图使一个Cuda等价。金属内核的源代码是

#include <metal_stdlib>
using namespace metal;

constant uint stride [[function_constant(0)]];
constant float dt    [[function_constant(1)]];
constant float a     [[function_constant(2)]];
constant float b     [[function_constant(3)]];

float2 f(const float2 x) {
    return float2(a, -b)*x.yx;
}

kernel void harmonic_occilator_stride(device float2 *x [[buffer(0)]],
                                             uint    i [[thread_position_in_grid]]) {
    for (uint j = 0; j < stride; j++) {
        x[i] += dt*f(x[i]);
    }
}

我第一次尝试将其转换为Cuda，在编译ptx文件时会导致大量错误。

__constant__ uint  stride;
__constant__ float dt;
__constant__ float a;
__constant__ float b;

__device__ float2 f(const float2 x) {
    return float2(a, -b)*x.yx;
}

extern "C" __global__ void harmonic_occilator_stride(float2 *x) {
    size_t i = blockIdx.x*blockDim.x + threadIdx.x;
    for (uint j = 0; j < stride; j++) {
        x[i] += dt*f(x[i]);
    }
}

它首先不喜欢的是x.yx。在金属中，这将颠倒float2内容的顺序。如何在Cuda中反转或更改向量的访问顺序？

接下来，它也不喜欢float2(a, -b)。这就产生了一个错误“没有合适的构造函数在float和float2之间进行转换”。如何构造向量文字？

它抱怨的最后一件事是，float没有float运算符，dt*f(x[i])没有float2。如果删除dt*并将其设置为x[i] += f(x[i])，它会抱怨float2和float2没有+=操作符。如何对这些类型执行操作，以及如何将向量和标量相乘？

在Metal中，当我将函数定义为function_constant时，当内核函数在运行时加载时，金属内核编译器将JIT是内核的一个特定的优化版本。Cuda有这个功能吗？

Answer 1

我现在只看了一眼金属规范。我不打算完全回答你的最后一个问题。但我认为，只要处理各种组件，遵循由金属定义的算术规则，就可以从概念上回答语法问题。

它也不喜欢float2(a，-b)。这就产生了一个错误“没有合适的构造函数在浮点数和float2之间进行转换”。如何构造向量文字？

为此，请使用头文件vector_functions.h (或.hpp)中定义的函数。(参见下面的示例)在vector_types.h中为CUDA定义的向量类型没有构造函数。

它首先不喜欢的是x.yx。在金属中，这将颠倒float2内容的顺序。如何在Cuda中反转或更改向量的访问顺序？

CUDA没有这种内置的多向量元素处理/swizzling功能。只需使用元素类型对元素执行操作即可。

metal:  return float2(a, -b)*x.yx;

CUDA:   #include <vector_functions.h>
        ...
        return make_float2(a*x.y, -b*x.x);

它抱怨的最后一件事是，行dtf(xi)没有*操作符，行dtf(Xi)没有float2。如果删除dt并将其设置为xi += f(xi)，它会抱怨float2和float2没有+=运算符。如何对这些类型执行操作，以及如何将向量和标量相乘？

类似于上面的内容，您需要构造等价的算术元素。

metal:  x[i] += dt*f(x[i]);

CUDA:   float2 temp1 = x[i];
        float2 temp2 = f(temp1);
        temp1.x += dt*temp2.x;
        temp1.y += dt*temp2.y;
        x[i] = temp1;

应该可以定义一组你自己的向量类型，如果你想这样做的话，可以匹配金属的大部分功能。我在这里描述的东西使用的是“内建”，如果您想用构造器、算术运算符等来创建自己的类型，可以是一个模型。

关于最后一个问题，CUDA并不总是像你描述金属的那样在运行时使用JIT。最接近您所描述的东西可能是使用C++模板的东西，这是由数据自动化系统支持的。一般来说，如果您可以将金属操作转换为等效的C++操作，那么您应该能够直接实现数据自动化系统中的操作。