Очевидная проблема - это проблема целочисленного усечения / интерполяции.
Одним из способов решения этой проблемы является привязка исходного изображения к текстуре, а затем чтение из текстуры с использованием вычисленных действительных значений координат.Текстуры CUDA дают вам «свободную» аппаратную интерполяцию / фильтрацию.Основным недостатком текстур является то, что интерполяция ограничена 8-битной внутренней точностью, поэтому в некоторых ситуациях она может быть недостаточно точной.Но в качестве первой попытки попробуйте что-то вроде:
__global__ void kernCuda(float * Destination, const float sintheta, const float costheta,
const int width)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
float tx = float(x)*costheta-float(y)*sintheta;
float ty = float(x)*sintheta+float(y)*costheta;
if(x<width && y<width){
Destination[x*width+y]=tex2D(Source_texture, tx+0.5f,ty+0.5f);
}
}
(заметка не проверялась, никогда не был рядом с компилятором, используйте на свой страх и риск).
Здесь Source_texture
- это текстураВы связываете исходные данные с.Вы можете установить поведение поворота края в настройке текстуры, в зависимости от того, как вы хотите справиться с этим.Как настроить и связать текстуру в Разделе 3.2.10.1 Руководства по программированию CUDA 4.1.
Обратите также внимание, что косинус и синус в матрице вращения постоянны для данного значения тета, поэтому было бы многоболее эффективно передавать их ядру в качестве аргумента, а не каждый поток вычислять одинаковые значения.Для поворота на 90 градусов, о котором вы спрашивали, просто вызовите ядро с помощью sintheta=0.f
и costheta=1.f
, и все готово.