Основы текстуры памяти-tex2D
При использовании текстурной памяти я наткнулся на следующий код:
uint f = (blockIdx.x * blockDim.x) + threadIdx.x;
uint c = (blockIdx.y * blockDim.y) + threadIdx.y;
uint read = tex2D( refTex, c+0.5f, f+0.5f);
Мой вопрос, почему мы добавляем 0.5f как для c а также f? Это смущает меня.. спасибо
2 ответа
В графике текстура - это набор образцов, которые описывают внешний вид поверхности. Образец - это точка. То есть он не имеет размера (в отличие от пикселя, который имеет физический размер). При использовании выборок для определения цветов пикселей каждая выборка располагается в точном центре соответствующего пикселя. При обращении к пикселям с координатами целого числа точный центр для данного пикселя становится его координатой целого числа плюс смещение 0,5 (в каждом измерении).
Другими словами, добавление 0,5 к координатам текстуры гарантирует, что при считывании из этих координат будет возвращено точное значение выборки для этого пикселя.
Тем не менее, это только когда filterMode для текстуры было установлено cudaFilterModeLinear что значение, которое читается из текстуры, изменяется в пределах пикселя. В этом режиме чтение из координат, которые не находятся в точном центре пикселя, возвращает значения, которые интерполируются между выборкой для данного пикселя и выборками для соседних пикселей. Таким образом, добавление 0,5 к координатам целых чисел фактически сводит на нет cudaFilterModeLinear Режим. Но, поскольку добавление 0.5 к координатам текстуры занимает в ядре циклы, лучше просто отключить интерполяцию, установив filterMode в cudaFilterModePoint, Затем чтение по любой координате в пикселе возвращает точное значение выборки текстуры для этого пикселя, и, таким образом, выборки текстуры могут быть считаны напрямую с использованием целых чисел.
Когда используешь cudaFilterModePoint, если в вычислении координат текстуры используется какая-либо математика с плавающей запятой, необходимо позаботиться о том, чтобы из-за погрешностей с плавающей запятой координаты текстуры не выходили за пределы диапазона для целевого пикселя.
Кроме того, как отмечается в комментариях, в вашем коде может быть проблема. Добавление 0.5f к координатам текстуры означает, что cudaFilterModeLinear режим используется, но этот режим возвращает число с плавающей точкой, а не int.
В зависимости от свойств текстуры значение возвращается tex2D может быть линейно интерполирован. "Индексы" f а также c поэтому являются не целыми числами, а непрерывными значениями между пределами каждого измерения.
Немного странным в этом примере является то, что возвращаемое значение является целым числом, которое в любом случае сделает любой линейный интерполант кусочно-постоянным.
Для получения более подробной информации см. Раздел 3.2.10 Руководства по программированию CUDA.