Особенности линейной текстурной памяти.

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

Привязывание линейного массива:

cudaError_t cudaBindTexture(size_t shift, texref tex, &src, size_t size));

Shift – смещение при привязке к массиву (к одному массиву можно привязать несколько тектсур);

Tex – объявленная текстура;

Src – массив в глобальной памяти, к которому привязывается текстура;

Size – размер привязываемой области в байтах.

Привязывание «двумерного массива» (в глобальной памяти он все равно хранится как линейный, и обращение к нему идет по одной координате):

cudaError_t cudaBindTexture2D(size_t shift, texref tex, &src, &channelDesc, int width,int height,int pitch);

Shift – смещение при привязке к массиву (к одному массиву можно привязать несколько текстур);

Tex – объявленная текстура;

Src – массив в глобальной памяти, к которому привязывается текстура;

channelDesc – описание канала;

width – ширина двумерного массива;

height – высота двумерного массива;

pitch – смещение каждой строки.

После окончания работы с текстурой её надо «отвязать»:

cudaError_t cudaUnbindTexture(texref tex);

Все вышеприведенные функции вызываются с хоста.

На устройстве используется функция, которая достает значение из текстуры:

tex1Dfetch(texRef tex, int index);

Tex – объявленная текстура;

Index – индекс вынимаемого значения в линейном массиве.

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

Особенности блоклинейной текстурной памяти.

Для использования блоклинейной текстурной памяти (cudaArray) требуется объявить переменную типа указатель на массив:

cudaArray * a;

Затем необходимо выделить память под данные на GPU:

cudaError_t cudaMallocArray(struct cudaArray ** arrayPtr, const struct cudaChannelFormatDesc * channelDesc, size_t width, size_t height);

arrayPtr – указатель на cudaArray;

channelDesc – описание канала (см выше);

width – ширина массива;

height – высота массива.

Затем скопировать в выделенную память данные:

cudaError_t cudaMemcpyToArray(struct cudaArray * dst, size_t wOffset, size_t hOffset, const void * src, size_t count, enum cudaMemcpyKind kind)

arrayPtr – указатель на cudaArray;

wOffset – смещение по горизонтали при привязки к массиву;

hOffset – смещение по вертикали при привязки к массиву;

Src – массив в памяти хоста, к который копируется;

count – размер данных в байтах;

kind – направление копирования.

После того как данные скопированы, можно осуществлять привязку

cudaArray массива к текстуре:

cudaError_t cudaBindTextureToArray (const struct textureReference *tex, const struct cudaArray *array, const struct cudaChannelFormatDesc *desc);

Tex – объявленная текстура;

array – массив в cudaArray, к которому привязывается текстура;

channelDesc – описание канала.

После использования необходимо «отвязать» текстуру точно так же, как и

линейную.

На устройстве используется функция, которая достает значение из текстуры:

tex1D (texRef tex, float x)

Tex – объявленная текстура;

x – индекс вынимаемого значения в линейном массиве.

tex2D (texRef tex, float x, float y)

Tex – объявленная текстура;

x, y – индексы вынимаемого значения в двухмерном массиве.

tex3D (texRef tex, float x, float y, float z)

Tex – объявленная текстура;

x, y, z – индексы вынимаемого значения в трехмерном массиве.