і лінійну.
У Лабораторної роботі №4 створений проект LR4_2, в якому одновимірний масив тестових даних завантажується в двовимірну cudaArray текстурную пам'ять. На відеокарті здійснюється білінійна інтерполяція масиву даних. Розглянуто два варіанти передачі масиву: з цілочисельними і нормалізованими координатами. У разі нормалізації у програмі також передбачається можливість використання двох режимів адресації в текстурную пам'ять: Clamp і Warp. Виконувався вікно проекту демонструє результат обробки даних при паралельних обчисленнях для наочного відображення властивостей досліджуваних режимів роботи з текстурної пам'яттю відеокарти.
3. Основні моменти коду
Проект LR4_1. Алгоритм виконання:
. Оголошується посилання на лінійну текстуру.
. Масив a [] заповнюється одним і тим же значенням PI.
. Резервується область глобальної пам'яті графічної карти.
. Копіюються дані на пристрій.
. Посилання лінійної текстури прив'язується до області пам'яті з завантаженим масивом даних.
. Запускаються потоки паралельного обчислення на відеокарті, в яких кожен потік примножує вихідне значення елемента масиву з текстури на його індекс і записує отриманий результат в елемент масиву, призначений для обміну з хостом, з цим же індексом.
. Синхронізуємо програму, чекаючи закінчення роботи всіх потоків.
. Звільняємо ресурси проекту.
Файл kernel.cu, проект LR4_1 ... #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 () {... 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 ); }
Проект LR4_2. Для оцінки правильності роботи програми в одному проекті об'єднана робота з текстурами в двох режимах: целочисленном і нормалізованому. Алгоритм виконання:
. Оголошуються посилання на двовимірні cudaArray-текстури і формат текстурних каналів.
. Здається тестовий масив a [].
. Резервується область глобальної пам'яті графічної карти під лінійні масиви обміну даними між пристроєм і хостом.
. Резервується область глобальної пам'яті графічної карти під двовимірні масиви cudaArray-текстур.
. Копіюються дані з хоста в текстурні області на пристрої.
. Встановлюються режими фільтрації на пристрої для повертаної вибірки з текстур. Додатково вибираються режими адресації для нормалізованих координат: Clamp, у випадку, коли координата обрізається по межі, або Wrap (координата загортається).
. Посилання текстур прив'язуються до області пам'яті з завантаженим масивом даних.
. Запускаються потоки паралельного обчислення на відеокарті, в яких кожен потік повертає хосту значення елемента масиву, отримане в результаті фільтрації (найближчим задане або білінійна інтерполяція) і заданого режиму адресації.
. Синхронізуємо програму, чекаючи закінчення роботи всіх потоків.
. Звільняємо ресурси проекту.
Приклад коду для випадку завдання цілочисельних координат:
Файл kernel.cu, проект LR4_2 ... texture lt; float, 2, cudaReadModeElementType gt; g_TexRef2; cudaChannelFormatDesc Desc2=cudaCreateChannelDesc (32,0,0,0, cudaChannelFormatKindFloat); __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 []) {...//задаємо тестовий масив а 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 * a=new float [numItems];// Виділяємо пам'ять під масиви a, c2 float * c2=new float [numItems]; float * cDev2=NULL;// створюємо покажчик на масиви для пам'я...
