Студопедия

КАТЕГОРИИ:


Архитектура-(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
Процедуры CUDA, исполняемые на CPU

// Файл 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; Нарушение авторских прав?; Мы поможем в написании вашей работы!


Нам важно ваше мнение! Был ли полезен опубликованный материал? Да | Нет



studopedia.su - Студопедия (2013 - 2024) год. Все материалы представленные на сайте исключительно с целью ознакомления читателями и не преследуют коммерческих целей или нарушение авторских прав! Последнее добавление




Генерация страницы за: 0.034 сек.