Cuda Kernel-эквивалент ядра Metal Compute - PullRequest
Новая
       5

Cuda Kernel-эквивалент ядра Metal Compute

0 голосов
/ 19 января 2019

У меня есть простое ядро ​​для металлических вычислений, которое я пытаюсь сделать эквивалентом Cuda.Исходный код ядра Metal: 

#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.В Metal это меняет порядок содержимого float2.Как мне изменить или изменить порядок доступа вектора в Cuda?

Следующее, что ему тоже не нравится float2(a, -b).Это выдает ошибку «нет подходящего конструктора для преобразования между float и float2».Как создать векторный литерал?

Последнее, на что он жалуется, это отсутствие оператора * для float и float2 для строки dt*f(x[i]).Если я удаляю dt* и просто устанавливаю его на x[i] += f(x[i]), он жалуется, что для float2 и float2 нет оператора +=.Как мне выполнить операции над этими типами и можно ли умножить векторы и скаляры?

В Metal, когда я определяю функции как function_constant, компилятор ядра Metal будет JIT определенной оптимизированной версии ядра, когда ядроФункция загружается во время выполнения.Есть ли у Cuda такая функциональность?

1 Ответ

0 голосов
/ 19 января 2019

Я только что взглянул на спецификацию металла .Я не буду пытаться полностью ответить на ваш последний вопрос.Но я думаю, что на вопросы синтаксиса можно ответить концептуально, просто имея дело с различными компонентами и следуя арифметическим правилам, определенным металлом.

он также не любит float2 (a, -b).Это выдает ошибку «нет подходящего конструктора для преобразования между float и float2».Как создать векторный литерал?

Используйте для этой цели функции, определенные в заголовочном файле vector_functions.h (или .hpp).(см. пример ниже). Векторные типы, определенные для CUDA в vector_types.h, не имеют конструкторов.

Первое, что ему не нравится, это x.yx.В Metal это меняет порядок содержимого float2.Как изменить или изменить порядок доступа к вектору в Cuda?

CUDA не имеет такого рода встроенной функции обработки / перемещения множества векторных элементов.Просто выполните операцию с элементами, используя типы элементов. 

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

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

Последнее, на что он жалуется, это отсутствие оператора * для float и float2 для строки dt f (х [г]).Если я удаляю dt и просто устанавливаю его на x [i] + = f (x [i]), он жалуется, что нет оператора + = для 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 не 't всегда JIT во время выполнения, как вы описываете для металла.Вероятно, ближе всего к тому, что вы описываете, может быть что-то, использующее шаблонирование C ++, которое поддерживается CUDA.В общем, если вы можете преобразовывать операции с металлом в эквивалентные операции C ++, вы должны иметь возможность напрямую реализовывать их в CUDA.

...