ті ДП cudaMalloc ( amp; cDev2, size);// Виділяємо пам'ять в ДП під масив cudaArray * aca2=NULL;// Створюємо покажчик на cuda-масив int nBlocksX=2; int nBlocksY=1; cudaMallocArray ( amp; aca2, amp; Desc2, nBlocksX * nThreads, nBlocksY);// Копіюємо дані з пам'яті хоста в 2D-cuda масив на пристрій cudaMemcpyToArray (aca2,0,0, a, size, cudaMemcpyHostToDevice);// Встановлюємо режим фільтрації://вибірка текстури поверне інтерполяцію із значень чотирьох сусідніх точок 2D поверхні g_TexRef2.filterMode=cudaFilterModeLinear;// Вибірка текстури поверне найближчим задане значення із значень чотирьох сусідніх точок 2D поверхні //g_TexRef2.filterMode=cudaFilterModePoint;// прив'язуємо текстурную посилання до CUDA-масиву у ДП cudaBindTextureToArray (g_TexRef2, aca2); dim3 threads=dim3 (nThreads); dim3 blocks=dim3 (nBlocksX, nBlocksY); Kernel2 lt; lt; lt; blocks, threads gt; gt; gt; (cDev2); cudaThreadSynchronize ();// Копіюємо масив c2 з пам'яті ДП в пам'ять ЦП cudaMemcpy ((void *) c2, cDev2, size, cudaMemcpyDeviceToHost);... CudaFree (cDev2);// звільняємо пам'ять масивів у ДП cudaUnbindTexture ( amp; g_TexRef2);// Скасування прив'язки текстур ДП і текстурних посилань delete [] a;// звільняємо пам'ять масивів в ЦП delete [] c2;...} Приклад коду для випадку завдання нормалізованих координат:
Файл kernel.cu, проект LR4_2 ... texture lt; float, 2, cudaReadModeElementType gt; g_TexRef; cudaChannelFormatDesc Desc=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); } Int main (int argc, char * argv []) {... size_t size=numItems * sizeof (float);// скільки займають елементи в байтах float * a=new float [numItems];// Виділяємо пам'ять під масиви a, c float * c=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;// Виділяємо пам'ять в ДП під масив cudaMalloc ( amp; cDev, size);// Створюємо покажчик на cuda-масив cudaArray * aca=NULL; int width=nThreads; int height=numBloks; cudaMallocArray ( amp; aca, amp; Desc, width, height);// Копіюємо дані з пам'яті хоста в 2D-cuda масив на пристрій cudaMemcpyToArray (aca, 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;// Вибірка текстури поверне найближчим задане значення із значень чотирьох сусідніх точок 2D поверхні //g_TexRef.filterMode=cudaFilterModePoint;// прив'язуємо текстурную посилання до CUDA-масиву у ДП cudaBindTextureToArray (g_TexRef, aca, Desc); dim3 dimBlock (16, 2); int kx=width + dimBlock.x; unsigned int kxx=(unsigned int) (kx - 1); int ky=height + dimBlock.y; unsigned int kyy=(unsigned int) (ky - 1); dim3 dimGrid (kxx/dimBlock.x, kyy /dimBlock.y); Kernel lt; lt; lt; dimGrid, dimBlock gt; gt; gt; (cDev, width, height); cudaThreadSynchronize ();// Копіюємо масив з c2 з пам'яті ДП в пам'ять ЦП cudaMemcpy ((void *) c, cDev, size, cudaMemcpyDeviceToHost);...//Звільняємо пам'ять масивів у ДП cudaFree (cDev);// Скасування прив'язки текстур ДП і текстурних посилань cudaUnbindTexture ( amp; g_TexRef);// звільняємо пам'ять масивів в ЦП delete [] a; delete [] c;...}
4. Результат роботи програм
Рис. 1. Вікно виводу консольного застосування LR4_1
Проект LR4_2. Приклад 1 - задані наступні стартові параметри:
- нормалізовані координати
g_TexRef.filterMode=cudaFilterModePoint; _TexRef.addressMode [0]=cudaAddressModeClamp; _TexRef.addressMode [1]=cudaAddressModeClamp;
- цілочисельні координати
g_TexRef2.filterMode=cudaFilterModePoint;
Масив вхідних даних упорядкований.
Рис. 2. Вікно виводу консольного застосування LR4_2. Приклад 1.
Проект LR4_2. Приклад 2 - задані наступні стартові параметри:
- нормалізовані координати
g_TexRef.filterMode=cudaFilterModePoint; _TexRef.addressMode [0]=cudaAddressModeClamp; _TexRef.addressMode [1]=cudaAddressModeClamp;
