Скалярные типы данных

Металл поддерживает скалярные типы, перечисленные в Таблице 2-1. Металл не поддерживает типы данных double, long, unsigned long, long long, unsigned long long, и long double.

Векторные и матричные типы данных

Металлический язык штриховки поддерживает подмножество векторных и матричных типов данных, реализованных системной математической библиотекой вектора.

Поддерживаемые имена типа вектора:

• booln

• charn, shortn, intn, ucharn, ushortn, uintn

• halfn и floatn

где n 2, 3, или 4 представления 2-, 3-или типа вектора с 4 компонентами.

Матричные поддерживаемые имена типов:

• halfnxm и floatnxm

где n и m число столбцов и строк. n и m может быть 2, 3, или 4. Матрица составлена из нескольких векторов. Например, a floatnx3 матрица составлена из n float3 векторы. Точно так же a halfnx4 матрица составлена из n half4 векторы.

pos = float4(1.0f, 2.0f, 3.0f, 4.0f);

float x = pos[0];    // x = 1.0
float z = pos[2];    // z = 3.0

float4 vA = float4(1.0f, 2.0f, 3.0f, 4.0f);
float4 vB;

for (int i=0; i<4; i++)
    vB[i] = vA[i] * 2.0f // vB = (2.0, 4.0, 6.0, 8.0);

int4 test = int4(0, 1, 2, 3);
int a = test.x;  //  a = 0
int b = test.y;  //  b = 1
int c = test.z;  //  c = 2
int d = test.w;  //  d = 3
int e = test.r;  //  e = 0
int f = test.g;  //  f = 1
int g = test.b;  //  g = 2
int h = test.a;  //  h = 3

float4 c;

c.xyzw = float4(1.0f, 2.0f, 3.0f, 4.0f);

c.z = 1.0f;

c.xy = float2(3.0f, 4.0f);

c.xyz = float3(3.0f, 4.0f, 5.0f);

float4 pos = float4(1.0f, 2.0f, 3.0f, 4.0f);

float4 swiz = pos.wzyx; // swiz = (4.0f, 3.0f, 2.0f, 1.0f)

float4 dup = pos.xxyy;  // dup = (1.0f, 1.0f, 2.0f, 2.0f)

float4 pos = float4(1.0f, 2.0f, 3.0f, 4.0f);
// pos = (5.0, 2.0, 3.0, 6.0)
pos.xw = float2(5.0f, 6.0f);

// pos = (8.0, 2.0, 3.0, 7.0)
pos.wx = float2(7.0f, 8.0f);

// pos = (3.0, 5.0, 9.0, 7.0)
pos.xyz = float3(3.0f, 5.0f, 9.0f);

float4x4 m;

// sets the 2nd column to all 2.0

m[1] = float4(2.0f);

// sets the 1st element of the 1st column to 1.0

m[0][0] = 1.0f;

// sets the 4th element of the 3rd column to 3.0

m[2][3] = 3.0f;

float4(float x);
float4(float x, float y, float z, float w);
float4(float2 a, float2 b);
float4(float2 a, float b, float c);
float4(float a, float b, float2 c);
float4(float a, float2 b, float c);
float4(float3 a, float b);
float4(float a, float3 b);
float4(float4 x);

float3(float x);

float3(float x, float y, float z);

float3(float a, float2 b);

float3(float2 a, float b);

float3(float3 x);

float2(float x);

float2(float x, float y);

float2(float2 x);

float x = 1.0f, y = 2.0f, z = 3.0f, w = 4.0f;
float4 a = float4(0.0f);
float4 b = float4(x, y, z, w);
float2 c = float2(5.0f, 6.0f);

float2 a = float2(x, y);
float2 b = float2(z, w);
float4 x = float4(a.xy, b.xy);

float4x4(fval);

fval    0.0    0.0    0.0

0.0    fval    0.0    0.0

0.0    0.0    fval    0.0

0.0    0.0    0.0    fval

float3x4(float3x4);

float3x4(half3x4);

float2x2(float2, float2);

float3x3(float3, float3, float3);

float3x2(float2, float2, float2);

// both cases below are not supported

float2x2(float a00, float a01, float a10, float a11);

float2x3(float2 a, float b, float2 c, float d);

Атомарные типы данных

Металлический атомарный тип данных ограничивается для использования атомарными функциями, реализованными Металлическим языком штриховки, как описано в Атомарных Функциях. Эти атомарные функции являются подмножеством C++ 11 атомарных и функции синхронизации. Металлические атомарные функции должны воздействовать на Металлические атомарные данные.

Металлические атомарные типы определяются как atomic_int и atomic_uint.

Буферы

Металл реализует буферы как указатель на встроенный или определяемый пользователем тип данных, описанный в device или constant адресное пространство. (Обратитесь к Спецификаторам Адресного пространства для Переменных и Параметрам за полное описание этих спецификаторов адреса.) Эти буферы могут быть объявлены в объеме программы или переданы как параметры функции.

Например:

device float4    *device_buffer;

struct my_user_data {

    float4 a;

    float  b;

    int2   c;

};

constant my_user_data *user_data;

Текстуры

Тип данных текстуры является дескриптором к одному - два - или трехмерные данные текстуры, соответствующие всем или части единственного уровня множественного отображения текстуры. Следующие шаблоны определяют определенные типы данных текстуры:

enum class access { sample, read, write };

texture1d<T, access a = access::sample>
texture1d_array<T, access a = access::sample>
texture2d<T, access a = access::sample>
texture2d_array<T, access a = access::sample>
texture3d<T, access a = access::sample>
texturecube<T, access a = access::sample>
texture2d_ms<T, access a = access::read>

Текстуры с форматами глубины должны быть объявлены как один из следующих типов данных текстуры:

enum class depth_format { depth_float };
depth2d<T, access a = access::sample, 
           depth_format d = depth_format::depth_float>
depth2d_array<T, access a = access::sample, 
                 depth_format d = depth_format::depth_float>
depthcube<T, access a = access::sample, 
             depth_format d = depth_format::depth_float>
depth2d_ms<T, access a = access::read, 
              depth_format d = depth_format::depth_float>

T указывает, что тип цвета возвратился при чтении из текстуры, или тип цвета указал при записи в текстуру. Для типов текстуры (кроме глубины текстурируют типы), T может быть half, float, short, ushort, int, или uint. Поскольку глубина текстурирует типы, T должен быть float.

access спецификатор описывает, как к текстуре можно получить доступ. Поддерживаемые спецификаторы доступа:

• sample - Объект текстуры может быть выбран. sample подразумевает возможность читать из текстуры с и без сэмплера.

• read - Без сэмплера графика или функция ядра могут только считать объект текстуры.

• write – Графика или функция ядра могут записать в объект текстуры.

depth_format спецификатор описывает формат текстур глубины. Единственное поддерживаемое значение depth_format.

Следующий пример использует эти спецификаторы доступа с параметрами объекта текстуры:

void foo (texture2d<float> imgA [[ texture(0) ]],

    texture2d<float, access::read> imgB [[ texture(1) ]],

    texture2d<float, access::write> imgC [[ texture(2) ]])

{

    ...

}

(См. Спецификаторы Атрибута для Определения местоположения Ресурсов для описания texture спецификатор атрибута.)

Сэмплеры

На Металлическом языке штриховки, sampler тип идентифицирует, как выбрать текстуру. Металлическая платформа позволяет Вам создавать соответствие MTLSamplerState объект и передача это в параметре графике или функции ядра. A sampler объект может также быть описан в источнике программы построения теней вместо в Металлической платформе. Источник программы построения теней позволяет указывать только подмножество состояния сэмплера: способ адресации, режим фильтра, нормализованные координаты и функция сравнения.

Таблица 2-2 описывает список поддерживаемых перечислений состояния сэмплера и их присваиваемых значений (и значения по умолчанию). Когда сэмплер инициализируется в источнике программы построения теней, эти состояния могут быть указаны.

Для способа адресации, clamp_to_zero подобно способу адресации зажима к границе OpenGL, за исключением того, что значение цвета рамки всегда (0.0, 0.0, 0.0, 1.0) при выборке вне текстуры, не имеющей альфа-компонента, и это (0.0, 0.0, 0.0, 0.0) при выборке вне текстуры, имеющей альфа-компонент.

Типы перечисления, используемые типом данных сэмплера, как описано в Таблице 2-2, указаны следующим образом. (Если coord установлен в pixel, min_filter и mag_filter значения должны быть тем же, mip_filter и compare_func значения должны быть none, и режимы адреса должны быть также clamp_to_zero или clamp_to_edge.)

enum class coord { normalized, pixel };
enum class filter { nearest, linear };
enum class min_filter { nearest, linear };
enum class mag_filter { nearest, linear };
enum class s_address { clamp_to_zero, clamp_to_edge, 
                       repeat, mirrored_repeat };
enum class t_address { clamp_to_zero, clamp_to_edge, 
                       repeat, mirrored_repeat };
enum class r_address { clamp_to_zero, clamp_to_edge, 
                       repeat, mirrored_repeat };
enum class address { clamp_to_zero, clamp_to_edge, 
                     repeat, mirrored_repeat };
enum class mip_filter { none, nearest, linear };

// can only be used with depth_sampler
enum class compare_func { none, less, less_equal, greater,
                  greater_equal, equal, not_equal };

Металлический язык штриховки реализует объект сэмплера следующим образом:

struct sampler {

    public:

        // full version of sampler constructor

        template<typename... Ts>

        constexpr sampler(Ts... sampler_params){};

    private:

};

Ts должны быть типы перечисления, упомянутые выше, который может использоваться типом данных сэмплера. Если тот же тип перечисления будет объявлен многократно в данном конструкторе сэмплера, то последнее перечисленное значение вступит в силу.

Следующий источник программы построения теней иллюстрирует несколько способов объявить сэмплеры. Спецификаторы атрибута (sampler(n), buffer(n), и texture(n)) это появляется в коде ниже, объяснены в Спецификаторах Атрибута для Определения местоположения Ресурсов. Обратите внимание на то, что сэмплерам или постоянным буферам, объявленным в источнике программы построения теней, не нужны эти спецификаторы атрибута.

constexpr sampler s(coord::pixel, 
                address::clamp_to_zero,
                filter::linear);

constexpr sampler a(coord::normalized);

constexpr sampler b(address::repeat);

constexpr sampler s(address::clamp_to_zero,
                    filter::linear,
                    compare_func::less);

kernel void
my_kernel(device float4 *p [[ buffer(0) ]],
          texture2d<float4> img [[ texture(0) ]],
          sampler smp [[ sampler(3) ]], 
          ...)
{
    ...
}

Массивы и структуры

Массивы и структуры поддерживаются со следующими ограничениями:

• Массивы текстуры и типов сэмплера не поддерживаются.

• texture и sampler типы не могут быть объявлены в структуре.

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

• Элементы структуры должны принадлежать тому же адресному пространству.

Выравнивание и размер типов

Таблица 2-3 перечисляет выравнивание и размер скалярных и векторных типов данных.

Таблица 2-4 перечисляет выравнивание и размер матричных типов данных.

Так как матрица составлена из векторов, каждый столбец матрицы имеет выравнивание своего векторного компонента. Например, каждый столбец a floatnx3 матрица является a float3 вектор, который является выровненный на 16-байтовой границе, как показано в Таблице 2-3. Точно так же каждый столбец a halfnx2 матрица является a half2 вектор, который является выровненный на 4-байтовой границе.

alignas спецификатор выравнивания может использоваться для указания требования выравнивания типа или объекта. alignas спецификатор может быть применен к объявлению переменной или элементу данных структуры или класса. Это может также быть применено к объявлению структуры, класса или типа перечисления.

Металлический языковой компилятор штриховки ответственен за выравнивание элементов данных к надлежащему выравниванию, как требуется по условию вводят. Для параметров графике или функции ядра, которая, как объявляют, была указателем на тип данных, Металлический языковой компилятор штриховки может предположить, что данные, на которые ссылается указатель, всегда соответственно выровненные, как требуется по условию вводят.

Упакованные векторные типы данных

Векторные типы данных, описанные в Векторных и Матричных Типах данных, выровненные к размеру вектора. Существует много вариантов использования, где разработчики требуют, чтобы были плотно упакованы их векторные данные. Одним примером является структура вершины, которая может содержать позицию, нормальную, векторы касательной и координаты текстуры, плотно упакованные и переданные как буфер функции вершины.

Следующие упакованные имена типа вектора поддерживаются:

• packed_charn, packed_shortn, packed_intn,

• packed_ucharn, packed_ushortn, packed_uintn,

• packed_halfn, и packed_floatn

где n 2, 3, или 4 представления 2-, 3-или типа вектора с 4 компонентами. ( packed_booln имена типа вектора резервируются.)

Таблица 2-5 перечисляет выравнивание и размер упакованных векторных типов данных.

Упакованные векторные типы данных обычно используются в качестве формата хранения данных. Металлическая загрузка поддержек и хранение от упакованного векторного типа данных до выровненного векторного типа данных и наоборот. Скопируйте конструкторов, и операторы присваивания также поддерживаются. Арифметика, логические, и операторы отношения также поддерживаются для упакованных векторных типов данных.

Пример:

device float4 *buffer;
device packed_float4 *packed_buffer;

int i;
packed_float4 f ( buffer[i] );
pack_buffer[i] = buffer[i];

// operator to convert from packed_float4 to float4.  
buffer[i] = float4( packed_buffer[i] );

К компонентам упакованного векторного типа данных можно получить доступ с индексом массива. Однако к компонентам упакованного векторного типа данных нельзя получить доступ с .xyzw или .rgba синтаксис выбора.

Пример:

packed_float4 f;

f[0] = 1.0f;  // OK

f.x = 1.0f;   // Illegal - compilation error

Неявные преобразования типов

Металл поддерживает неявные преобразования между скалярными встроенными типами (кроме void). Когда неявное преобразование выполняется, это не просто реинтерпретация значения выражения, но и преобразование того значения к эквивалентной стоимости в новом типе. Например, целочисленное значение 5 преобразовывается в значение с плавающей точкой 5.0.

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

int4 i;

float4 f = i;    // compile error

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

Например:

float4 f = 2.0f;    // f = (2.0f, 2.0f, 2.0f, 2.0f)

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

Неявные преобразования для типов указателей соблюдают правила, описанные в C++ 11 Спецификаций [PDF].

Преобразования типов и данные иного толкования

Металл использует static_cast оператор для преобразования от скалярного или типа вектора до другого скалярного или типа вектора без насыщенности и с режимом округления значения по умолчанию. В зависимости от округляющихся режимов, поддерживаемых GPU, преобразовывая до мая с плавающей точкой вокруг к нулю или к самому близкому четному числу по умолчанию. При преобразовании в целое число режим по умолчанию округляется к нулю. Если исходный тип является скалярной или векторной булевской переменной, значением false преобразовывается в нуль и значение true преобразовывается в одного.

Металлический язык штриховки добавляет as_type<type-id> оператор для разрешения любого скалярного или векторного типа данных (который не является указателем) быть данным иное толкование как другой скалярный или векторный тип данных того же размера. Биты в операнде возвращаются непосредственно без модификации как новый тип. Обычное продвижение типа для аргументов функции не выполняется.

Например, as_type<float>(0x3f800000) возвраты 1.0f, который является значением комбинации двоичных разрядов 0x3f800000 если просматривается как значение одинарной точности IEEE 754.

Это - ошибка использовать as_type<type-id> оператор для иного толкования данным к типу, имеющему различное число байтов.

Вот некоторые примеры юридических и недопустимых преобразований с помощью as_type<type-id> оператор:

float f = 1.0f;
// Legal. Contains: 0x3f800000
uint u = as_type<uint>(f); 

// Legal. Contains:
// (int4)(0x3f800000, 0x40000000, 
//        0x40400000, 0x40800000)
float4 f = float4(1.0f, 2.0f, 3.0f, 4.0f);
int4 i = as_type<int4>(f);

int i;
// Legal. 
short2 j = as_type<short2>(i);

half4 f;
// Error. Result and operand have different sizes
float4 g = as_type<float4>(f);

float4 f;
// Legal. g.xyz will have same values as f.xyz.
// g.w is undefined
float3 g = as_type<float3>(f);