масиві, збігаються з цілочисельними коордінатамісредняя похибка, вище, ніж у цілочисельних коордінатWarpблізко вірний на всьому тестовому массівевисокая погрешностьблізко вірний на всьому тестовому масиві, збігаються з цілочисельними координатами, але з ефектом заворачіваніявисокая похибка, вище, ніж у цілочисельних координат
Таким чином, видно, що білінійна інтерполяція від вхідних даних дає наближені значення, рівень похибки яких зростає при зменшенні упорядкування набору даних. Вибірка з текстури по нормалізованим координатам збільшує похибку, особливо при Warp-перетворенні. Крім того необхідно знати і враховувати ефект загортання, якщо текстурная пам'ять використовує нормалізовані координати для своєї адресації.
Програми
Код програми
Файл kernel.cu, проект LR4_1 # include cuda.h #include cuda_runtime.h #include lt; iostream gt; #include lt; vector gt; #include lt; cstdio gt; using namespace std; #define PI 3.1415926536 #define nBlocks 1//Number of thread blocks #define nThreads 128//Number of threads per block texture lt; float, cudaTextureType1D, cudaReadModeElementType gt; g_TexRef; __global__ void kernel (float * data) {int idx=blockIdx.x * blockDim.x + threadIdx.x; data [idx]=tex1Dfetch (g_TexRef, idx) * idx; if (idx gt;=100) return; printf (% d t% .2f t% .2f n raquo ;, idx, tex1Dfetch (g_TexRef, idx), data [idx]); } Void main () {cout lt; lt; idx Texture (idx) data [idx] n lt; lt; endl; float a [nBlocks * nThreads]; float * devX=NULL; for (int i=0; i lt; nBlocks * nThreads; i ++) a [i]=PI; cudaMalloc ((void **) amp; devX, nBlocks * nThreads * sizeof (float)); cudaMemcpy (devX, a, nBlocks * nThreads * sizeof (float), cudaMemcpyHostToDevice); cudaBindTexture ((size_t) 0, amp; g_TexRef, devX, amp; g_TexRefannelDesc, (size_t) (nBlocks * nThreads * sizeof (float))); kernel lt; lt; lt; nBlocks, nThreads gt; gt; gt; (devX); cudaDeviceSynchronize (); cudaUnbindTexture ( amp; g_TexRef); cudaFree (devX); system ( pause ); }
Код програми
Файл kernel.cu, проект LR4_2 #include cuda.h #include cuda_runtime.h #include lt; stdlib.h gt; #include lt; stdio.h gt; #include lt; time.h gt; #include lt; fstream gt; #include lt; iostream gt; #define numBloks 2 #define nThreads 128 texture lt; float, 2, cudaReadModeElementType gt; g_TexRef; texture lt; float, 2, cudaReadModeElementType gt; g_TexRef2; cudaChannelFormatDesc Desc=cudaCreateChannelDesc (32,0,0,0, cudaChannelFormatKindFloat); cudaChannelFormatDesc Desc2=cudaCreateChannelDesc (32,0,0,0, cudaChannelFormatKindFloat); __global__ void Kernel (float * c, int width, int height) {// обчислення нормалізованих текстурних координат unsigned int x=blockIdx.x * blockDim.x + threadIdx.x; unsigned int y=blockIdx.y * blockDim.y + threadIdx.y; float u=x/(float) width; float v=y/(float) height;// Запис в глобальну пам'ять int i=y * width + x; c [i]=tex2D (g_TexRef, u, v);//printf(laquo;c[%d]=%.2ftx=%dty=%dtu=%ftv=%ftnnraquo;,i,c[i], x, y, u, v); } __global__ Void Kernel2 (float * c2) {int index=blockIdx.x * blockDim.x + threadIdx.x; int i=index + blockIdx.y * gridDim.x * blockDim.x; c2 [i]=tex2D (g_TexRef2, index, blockIdx.y); } Int main (int argc, char * argv []) {const unsigned int numItems=nThreads * numBloks;// число елементів в масивах a, c, c2 size_t size=numItems * sizeof (float);// скільки займають елементи в байтахfloat * a=new float [numItems];// Виділяємо пам'ять під масиви a, c, c2 float * c=new float [numItems]; float * c2=new float [numItems];// Задаємо тестовий масив а for (int i=0; i lt; numItems; i ++) a [i]=i;// for (int i=0; i lt; numItems; i ++) a [i]=(float) (double (rand ())/RAND_MAX * 100);// створюємо покажчик на масиви для пам'яті ДП float * cDev=NULL; float * cDev2=NULL;// Виділяємо пам'ять в ДП під масив cudaMalloc ( amp; cDev, size); cudaMalloc ( amp; cDev2, size);// Створюємо покажчик на cuda-масив cudaArray * aca=NULL; cudaArray * aca2=NULL; int width=nThreads; int height=numBloks; cudaMallocArray ( amp; aca, amp; Desc, width, height); int nBlocksX=2; int nBlocksY=1; cudaMallocArray ( amp; aca2, amp; Desc2, nBlocksX * nThreads, nBlocksY);// Копіюємо дані з пам'яті хоста в 2D-cuda масив на пристрій cudaMemcpyToArray (aca, 0,0, a, size, cudaMemcpyHostToDevice); cudaMemcpyToArray (aca2,0,0, a, size, cudaMemcpyHostToDevice);//g_TexRef.addressMode[0]=cudaAddressModeClamp;//g_TexRef.addressMode[1]=cudaAddressModeClamp;//g_TexRef.addressMode[0]=cudaAddressModeWrap;//g_TexRef.addressMode[1]=cudaAddressModeWrap; g_TexRef.normalized=true;// Встановлюємо режим фільтрації://вибірка текстури поверне інтерполяцію із значень чотирьох сусідніх точок 2D поверхні g_TexRef.filterMode=cudaFilterModeLinear;// Вибірка текстури поверне найближчим задане значення із значень чотирьох су...
