КАТЕГОРИИ: Архитектура-(3434)Астрономия-(809)Биология-(7483)Биотехнологии-(1457)Военное дело-(14632)Высокие технологии-(1363)География-(913)Геология-(1438)Государство-(451)Демография-(1065)Дом-(47672)Журналистика и СМИ-(912)Изобретательство-(14524)Иностранные языки-(4268)Информатика-(17799)Искусство-(1338)История-(13644)Компьютеры-(11121)Косметика-(55)Кулинария-(373)Культура-(8427)Лингвистика-(374)Литература-(1642)Маркетинг-(23702)Математика-(16968)Машиностроение-(1700)Медицина-(12668)Менеджмент-(24684)Механика-(15423)Науковедение-(506)Образование-(11852)Охрана труда-(3308)Педагогика-(5571)Полиграфия-(1312)Политика-(7869)Право-(5454)Приборостроение-(1369)Программирование-(2801)Производство-(97182)Промышленность-(8706)Психология-(18388)Религия-(3217)Связь-(10668)Сельское хозяйство-(299)Социология-(6455)Спорт-(42831)Строительство-(4793)Торговля-(5050)Транспорт-(2929)Туризм-(1568)Физика-(3942)Философия-(17015)Финансы-(26596)Химия-(22929)Экология-(12095)Экономика-(9961)Электроника-(8441)Электротехника-(4623)Энергетика-(12629)Юриспруденция-(1492)Ядерная техника-(1748) |
П.1.4. Системные переменные
П.1.3. Другие операторы и выражения П.1.2. Встроенные функции П.1.2. Типы данных texture – двухмерный массив исходных данных (текстура), ячейки которого адресуются двумя координатами – парой чисел с плавающей точкой; sampler – автоматическая процедура, осуществляющая выборку данных из входных текстур. Применение показано в примерах программ; float2 – переменная, представляющая собой вектор из 2-х вещественных чисел одинарной точности; используется для адресации ячеек входных текстур и рендер-цели; float4 – переменная, представляющая собой вектор из 4-х вещественных чисел одинарной точности, либо функция, возвращающая такой вектор в качестве значения; во втором варианте используется при декларации пиксельных шейдеров; void – функция, не возвращающая никаких значений; используется при декларации вершинных шейдеров; technique – декларация «техники», т.е. процедуры, которую могут запускать внешние программы, взаимодействующие с графическим процессором; внутри этой процедуры задаются значения управляющих переменных, запускаются на исполнение вершинный и пиксельный шейдеры. tex2D(объект типа sampler, координаты ячейки типа float2) – функция, возвращающая расположенный в ячейке с заданными координатами 4-вектор из текстуры, к которой «привязан» sampler; return – оператор, возвращающий значение функции, в теле которой он расположен; PixelShader = compile ps_3_0 sum() – строка, задающая функцию sum() в качестве исполняемого пиксельного шейдера; VertexShader = compile vs_3_0 transform() – строка, задающая функцию transform() в качестве исполняемого вершинного шейдера; ZEnable – имеет значение true, если используется Z-буфер, либо false, если он не используется. При физическом моделировании Z-буфер не используется, так что ZEnable = false;
CullMode – устанавливает режим отсечения невидимых треугольников. При физическом моделировании отсечения треугольников нет, так что CullMode = none.
Приложение 2 // Файл matrixMul.cu [68]
// Copyright 1993-2007 NVIDIA Corporation. All rights reserved.
/* Умножение матриц: C = A * B. Вызывающий код на CPU. * * В этом примере реализовано умножение матриц как приведено в главе 7 * документации CUDA Programming Guide [68]. Эта реализация предназначена * для изучения принципов программирования на CUDA, поэтому эффективность * принесена в жертву ясности изложения. * * Высокопроизводительная версия умножения матриц реализована в библиотеке CUBLAS */
#include <stdlib.h> #include <stdio.h> #include <string.h> #include <math.h> #include <cutil.h>
#include <matrixMul_kernel.cu>
//////////////////////////////////////////////////////////////////////////////// // предварительное объявление функций void runTest(int argc, char** argv); void randomInit(float*, int); void printDiff(float*, float*, int, int);
/* Процедура, которая для сравнения перемножает матрицы на центральном процессоре: */
extern "C" void computeGold(float*, const float*, const float*, unsigned int, unsigned int, unsigned int);
int main(int argc, char** argv) { runTest(argc, argv); CUT_EXIT(argc, argv); }
void runTest(int argc, char** argv) { CUT_DEVICE_INIT();
// инициализация генератора псевдослучайных чисел rand() srand(2006);
// выделение памяти для матриц A и B unsigned int size_A = WA * HA; unsigned int mem_size_A = sizeof(float) * size_A; float* h_A = (float*) malloc(mem_size_A); unsigned int size_B = WB * HB; unsigned int mem_size_B = sizeof(float) * size_B; float* h_B = (float*) malloc(mem_size_B);
// заполнение исходных матриц псевдослучайными числами randomInit(h_A, size_A); randomInit(h_B, size_B);
// выделение памяти на GPU float* d_A; CUDA_SAFE_CALL(cudaMalloc((void**) &d_A, mem_size_A)); float* d_B; CUDA_SAFE_CALL(cudaMalloc((void**) &d_B, mem_size_B));
// копирование матриц на GPU CUDA_SAFE_CALL(cudaMemcpy(d_A, h_A, mem_size_A, cudaMemcpyHostToDevice)); CUDA_SAFE_CALL(cudaMemcpy(d_B, h_B, mem_size_B, cudaMemcpyHostToDevice));
// выделение памяти на GPU для матрицы-результата unsigned int size_C = WC * HC; unsigned int mem_size_C = sizeof(float) * size_C;
float* d_C; CUDA_SAFE_CALL(cudaMalloc((void**) &d_C, mem_size_C));
// создание и запуск таймера unsigned int timer = 0; CUT_SAFE_CALL(cutCreateTimer(&timer)); CUT_SAFE_CALL(cutStartTimer(timer));
// установка параметров выполнения ядра dim3 threads(BLOCK_SIZE, BLOCK_SIZE); dim3 grid(WC / threads.x, HC / threads.y);
// запуск ядра matrixMul<<< grid, threads >>>(d_C, d_A, d_B, WA, WB);
// проверяем на наличие ошибок выполнения ядра CUT_CHECK_ERROR("Ошибка выполнения ядра");
// выделяем память на CPU для матрицы-результата float* h_C = (float*) malloc(mem_size_C);
// копируем результат на CPU CUDA_SAFE_CALL(cudaMemcpy(h_C, d_C, mem_size_C, cudaMemcpyDeviceToHost));
// останавливаем таймер и освобождаем ресурсы CUT_SAFE_CALL(cutStopTimer(timer)); printf("Processing time: %f (ms) n", cutGetTimerValue(timer)); CUT_SAFE_CALL(cutDeleteTimer(timer));
// вычисляем произведение A*B на CPU для сравнения точности float* reference = (float*) malloc(mem_size_C); computeGold(reference, h_A, h_B, HA, WA, WB);
// проверяем результат CUTBoolean res = cutCompareL2fe(reference, h_C, size_C, 1e-6f); printf("Тест %s n", (1 == res)? "ПРОЙДЕН": "ПРОВАЛЕН"); if (res!=1) printDiff(reference, h_C, WC, HC);
// освобождаем выделенную память free(h_A); free(h_B); free(h_C); free(reference); CUDA_SAFE_CALL(cudaFree(d_A)); CUDA_SAFE_CALL(cudaFree(d_B)); CUDA_SAFE_CALL(cudaFree(d_C)); }
// Заполнение матрицы псевдослучайными числами. void randomInit(float* data, int size) { for (int i = 0; i < size; ++i) data[i] = rand() / (float)RAND_MAX; }
void printDiff(float *data1, float *data2, int width, int height) { int i,j,k; int error_count=0; for (j=0; j<height; j++) { for (i=0; i<width; i++) { k = j*width+i; if (data1[k]!= data2[k]) { printf("diff(%d,%d) CPU=%4.4f, GPU=%4.4f n", i,j, data1[k], data2[k]); error_count++; } } } printf(" Количество ошибок = %d n", error_count); }
// Файл matrixMul.h [68]
// Copyright 1993-2007 NVIDIA Corporation. All rights reserved.
#ifndef _MATRIXMUL_H_ #define _MATRIXMUL_H_
// Размер связки потоков #define BLOCK_SIZE 16
// Размеры матрицы (для простоты кода выбраны кратными размеру связки потоков) #define WA (3 * BLOCK_SIZE) // Matrix A width #define HA (5 * BLOCK_SIZE) // Matrix A height #define WB (8 * BLOCK_SIZE) // Matrix B width #define HB WA // Matrix B height #define WC WB // Matrix C width #define HC HA // Matrix C height
#endif // _MATRIXMUL_H_
#ifndef _MATRIXMUL_H_ #define _MATRIXMUL_H_
// Размер связки потоков #define BLOCK_SIZE 16
// Размеры матрицы (для простоты кода выбраны кратными размеру связки потоков) #define WA (3 * BLOCK_SIZE) // Matrix A width #define HA (5 * BLOCK_SIZE) // Matrix A height #define WB (8 * BLOCK_SIZE) // Matrix B width #define HB WA // Matrix B height #define WC WB // Matrix C width #define HC HA // Matrix C height
#endif // _MATRIXMUL_H_
// Файл matrixMul_gold.cpp [68]
// Copyright 1993-2007 NVIDIA Corporation. All rights reserved.
extern "C" void computeGold(float*, const float*, const float*, unsigned int, unsigned int, unsigned int);
//////////////////////////////////////////////////////////////////////////////// // Вычисляем произведение матриц на CPU для проверки // C = A * B // @param C массив для хранения результата (память выделяется заранее) // @param A матрица A // @param B матрица B // @param hA кол-во строк матрицы А // @param wB кол-во столбцов матрицы B //////////////////////////////////////////////////////////////////////////////// void computeGold(float* C, const float* A, const float* B, unsigned int hA, unsigned int wA, unsigned int wB) { for (unsigned int i = 0; i < hA; ++i) for (unsigned int j = 0; j < wB; ++j) { double sum = 0; for (unsigned int k = 0; k < wA; ++k) { double a = A[i * wA + k]; double b = B[k * wB + j]; sum += a * b; } C[i * wB + j] = (float)sum; }}
Дата добавления: 2014-12-07; Просмотров: 420; Нарушение авторских прав?; Мы поможем в написании вашей работы! Нам важно ваше мнение! Был ли полезен опубликованный материал? Да | Нет |