2.3 CudaArray текстурная память

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

§ Нормализация координат (перевод [W, H] => [0,1]).

§ Преобразование координат:

- Clamp - координата обрезается по границе;

- Wrap - координата заворачивается.

§ Фильтрация (при обращении по float координате):

- Point - возвращается ближайшее заданное значение;

- Linear - производится билинейная интерполяция.

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

cudaArray * a;

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

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

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

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

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

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

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

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 массива к текстуре:

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 - индексы вынимаемого значения в трехмерном массиве.

После использования необходимо отвязать текстуру точно так же как и линейную.

В Лабораторной работе №4 создан проект LR4_2, в котором одномерный массив тестовых данных загружается в двумерную cudaArray текстурную память. На видеокарте осуществляется билинейная интерполяция массива данных. Рассмотрено два варианта передачи массива: с целочисленными и нормализованными координатами. В случае нормализации в программе также предусматривается возможность использования двух режимов адресации в текстурную память: Clamp и Warp. Исполняемое окно проекта демонстрирует результат обработки данных при параллельных вычислениях для наглядного отображения свойств исследуемых режимов работы с текстурной памятью видеокарты.