Рис. 9.7. Молекулярная динамика на графическом процессоре шейдерной модели 4.0

Ниже с комментариями приведена реализация «треугольного» цикла с 3-м законом Ньютона на CUDA. Для наглядности используется только один мультипроцессор. Каждый i-й поток суммирует парные взаимодействия i-й частицы и записывает их в разделяемую память, а в конце сохраняет одну из накопленных сумм в линейный массив force (каждой частице соответствует только один элемент), расположенный в глобальной памяти. Это позволяет минимизировать количество обращений к медленной общей памяти.

a)

б)

Рис. 9.8. Расчёт сил с учётом и без учёта 3-го закона Ньютона

Рис. 9.9. «Треугольный» цикл расчёта сил

Процедура вычисления сил в «треугольном» цикле на CUDA

/* Спецификатор __global__ означает, что вся процедура будет исполняться на GPU как вычислительное ядро */

__global__ void kernel_NxN(float4 force[], float4 pos[], int type[], float4 coefs[]) {

/* Переменные, описываемые без спецификаторов, будут хранится в том кэше GPU, который называют «памятью для констант». Они не будут меняться по ходу алгоритма: */

int j;

/* Вычислительные потоки в данном случае упорядочены в линейный массив. Каждый поток рассчитывает силу Fi, действующую на одну из частиц: */

int i = threadIdx. x;

/* Тип частицы type[i] и 4-вектор с координатами частицы pos[i] копируются в кэш, для того чтобы потом не обращаться за ними внутри цикла по j к медленной видеопамяти: */

int type_i = type[i];

float4 pos_i = pos[i];

/* Создаётся переменная (4-вектор) force_i, в которой будет сумми­ро­ваться результирующая сила, действующая на частицу i со стороны осталь­ных частиц; в начале расчёта она зануляется */

float4 force_i = { 0, 0, 0, 0 };

/* В разделяемой памяти (на что указывает модификатор __shared__ ) создаются массивы, содержащие координаты и типы всех частиц, параметры всех потенциалов взаимодействия (coefs_ij[4], где coefs_ij – сам по себе 4-вектор), а также все результирующие силы (force_j[N]), действующие на каждую из частиц: */

__shared__ float4 pos_j[N];

__shared__ int type_j[N];

__shared__ float4 coefs_ij[4];

__shared__ float4 force_j[N];

/* Следующий блок операторов копирует данные из медленной «глобальной» памяти GPU, куда их перед обращением к GPU записывает центральный процессор, в быструю разделяемую память. Каждый из потоков копирует элементы массивов, относящиеся к «своей» частице i, а первые 4 потока – ещё и параметры потенциалов взаимодействия */

if (threadIdx. x < 4) coefs_ij[i] = coefs[i];

pos_j[i] = pos[i]; type_j[i] = type[i];

force_j[j] = make_float4(0, 0, 0, 0);

/* Синхронизация потоков. Выполнение программы за функцией __syncthreads() продолжится только после того, как все потоки завершат копирование данных в разделяемую память */

__syncthreads();

/* Цикл, в котором суммируется сила, действующая на i-ю частицу со стороны остальных. Индекс j изменяется до N + i для того, чтобы во всех потоках, независимо от значения i, в этом цикле было одинаковое количество итераций. Фактически, при j ³ N внутри цикла ничего не делается */

for (j = i+1; j < N + i; j++)

{

if (j < N)

{

/* Функция force_ij возвращает силу, действующую между части­цами i и j. Аргументами функции являются координаты частиц и параметры потенциала взаимодействия */

float4 Fij = force_ij(pos_i, pos_j[j], coefs_ij[type_i + type_j[j]]);

/* Применение 3-го закона Ньютона

force_i += Fij; force_j[j] -= Fij;

}

/* Синхронизация потоков, необходимая для устранения возможных конфликтов доступа к элементам массива force_j[j]: */

__syncthreads();

}

/* Получение окончательных результирующих сил

force[i] = force_i + force_j[i];

}

/* Функция расчёта парной силы force_ij. Возвращает 4-вектор с компонентами силы (четвёртый элемент – пустой). Спецификатор __device__ означает, что функция и вызывается, и исполняется графическим процессором */

__device__ float4 force_ij(float4 pos_i, float4 pos_j, float4 c)

{

float4 R = pos_i - pos_j; R. w = rsqrt(max(R * R, 1e-4));

return R * (c. x * R. w * R. w * R. w + pow(c. y * R. w, c. z));

}

Библиографический список

1.  Структура и принципы функционирования ЭВМ // http:///education/informat/eu_intro/i4.htm. 2008. 7с.

2.  , Карпов операционных систем. Курс лекций / , . Интернет-университет информационных технологий. http://www. *****/department/os/osintro/8/. 2008. 76с.

3.  , , Жегуло пособие по курсу "Многопроцессорные системы и параллельное программирование" / , , . Ростовский государственный университет. http://**/tutor/method/m1/page09_3.html. 20с.

4.  , Соловьев вычислительных систем. Курс лекций / , . Интернет-университет информационных технологий. http://www. *****/department/hardware/csorg/10/. 20с.

НЕ нашли? Не то? Что вы ищете?

5.  Tera Computer // Computerworld. 1997. №37. C. 240. http://www. *****/cw/1997/37/24023/

6.  Воеводин обработка данных. Курс лекций / . Лаборатория Параллельных Информационных Технологий, НИВЦ МГУ. 2008. 78с. http://*****/vvv/lec4.html

7.  Семенов и протоколы каналов и сетей передачи данных. Курс лекций / . Интернет-университет информационных технологий. http://www. *****/department/network/algoprotnet/10/. 2007. 95с.

8.  , , Корхов и топологии многопроцессорных вычислительных систем. Курс лекций / , , . Интернет-университет информационных технологий. http://www. *****/department/hardware/atmcs/6/. 2008. 96с.

9.  Лекция 2. Архитектура вычислительных систем. // Архитектуры и топологии многопроцессорных вычислительных систем. Электронный учебник / А. Богданов, В. Мареев, Е. Станкова, В. Корхов. http://www. *****/text/teach/topolog/2.htm. 2008. 96с.

10.  Параллельное программирование // Компьютера Online. 2005. http://*****/hitech/242551/ 8c.

11.  Немнюгин С. А., Стесик  программирование для многопроцессорных вычислительных систем / , . СПб.: БХВ-Петербург, 20 с.

12.  Медведев A. Наступление ATI Technologies продолжается: RADEON X1900 XTX/XT (R580) // http://www. /video2/r580-part1.shtml. 2006. 6с.

13.  Kirk D, Hwu W. Programming Massively Parallel Processors / D. Kirk, W. Hwu. ECE 498AL1, University of Illinois, Urbana-Champaign, 2007.

14.  Графический конвейер / http://ru. wikipedia. org/wiki/ - Википедия. 2008. 19с.

15.  Лаборатория Параллельных информационных технологий НИВЦ МГУ. История развития / http://*****/ - Информационно-аналитический центр по параллельным вычислениям. 2008. 2с.

16.  NVIDIA GeForce 7800 GTX 256MB PCI-E. Часть 1 - Теория и архитектура // http://www. /video2/g70-part1.shtml. 2005. 10с.

17.  P. Gibbon, G. Sutmann, Long-Range Interactions in Many-Particle Simulation, Quantum Simulations of Complex Many-Body Systems: From Theory to Algorithms, Lecture Notes, J. Grotendorst, D. Marx, A. Muramatsu (Eds.), NIC Series, Vol. 10, 467–506, (2002).

18.  Amara G. Amara's Recap of Particle Simulation Methods // http://www. /ftpstuff/nbody. txt. 2008. 23c.

19.  P. P. Ewald, Die Berechnung optischer und elektrostatischer Gitterpotentiale, Ann. Phys. 64,

20.  J. W. Perram, H. G. Petersen, and S. W. D. Leeuw, An algorithm for the simulation of condensed matter which grows as the N3/2 power of the number of particles, Mol. Phys. 65, 875–

21.  C. K. Birdsall and A. B. Langdon, Plasma Physics via Computer Simulation, (McGraw-Hill, New York, 1985).

22.  Хокни Р., Иствуд Дж. Численное моделирование методом частиц. Пер. с англ. М.: Мир 19 с. / R. W. Hockney and J. W. Eastwood, Computer Simulation Using Particles, Institute of Physics Publishing. 19с.

23.  J. Barnes and P. Hut, A hierarchical O(NlogN) force-calculation algorithm, Nature 324, 446–

24.  L. Hernquist, Hierarchical N-body methods, Comp. mun. 48, 107–

25.  L. Greengard and V. Rokhlin, A fast algorithm for particle simulations, p. Phys. 73, 325–

26.  H. Cheng, L. Greengard, and V. Rohklin, A fast adaptive multipole algorithm in three dimensions, p. Phys. 155, 468–

27.  K. E. Schmidt and M. A. Lee, Implementing the fast multipole method in three dimensions, J. Stat. Phys. 63, 1223–1

28.  K. Esselink, A comparison of algorithms for long-range interactions, Comp. mun. 87, 375–

29.  R. K. Kalia, S. de Leeuw, A. Nakano and P. Vashishta, Molecular dynamics simulations of Coulombic systems on ditributed-memory MIMD machines, Comp. mun. 74, 316–

30.  J. V. L. Beckers, C. P. Lowe, and S. W. de Leeuw, An iterative PPPM method for simulating Coulombic systems on distributed memory parallel computers, Mol. Sim. 20, 369–

31.  L. Greengard and W. D. Groop, A parallel version of the fast multipole method, Comp. Math. Applic. 20, 63–71 (1990).

32.  Warren, M. S., J. K. Salmon, and D. J. Becker. 1997. Pentium Pro inside: I. A treecode at 430 Gigaflops on ASCI Red; II. Price/performance of $50/Mflop on Loki and Hyglac. Proc. Supercomputing '97, No­vember, online at http://www. supercomp. org/sc97/proceedings

33.  Walker J. R. and Catlow C. R. Structural and dynamic properties of UO2 at high temperatures, J. Phys. C. Solid State Phys., v.14, 979–

34.  и др. Моделирование нестехиометрической двуокиси урана методом молекулярной динамики. Часть II. / , , . Отчет по НИР УрО АТН РФ, Екатеринбург, 2004. 63с.

35.  Matzke H. Nonstoichiometric oxides, Ed. O. Toff Sorensen. New York, 155–

36.  Купряжкин А. Я., , Голованов  кислорода в диоксиде урана в области фазовых переходов, ЖТФ, 2004, т. 74, вып. 2

37.  Thermophysical Properties Database of Materials for Light Water Reactors and Heavy Water Reactors // IAEA-TECDOC-1496, 2006, ISBN -1. http://www-pub. iaea. org/MTCD/publications/PDF/te_1496_web. pdf

38.  Molecular dynamics // http://en. wikipedia. org/wiki/Molecular_dynamics - Википедия. 2008. 13с.

39.  Sindzingre P., Gillan M. J. A molecular dynamics study of solid and liquid UO2 // J. Phys. C: Solid State Phys. 1988. V. 21, P. .

40.  Karakasidis T., Lindan P. J.D. A comment on a rigid-ion potential for UO2 // J. Phys.: Cond. Matter. 1994. V. 6, P. .

41.  Walker J. R., Catlow C. R.A. Structural and dynamic properties of UO2 at high temperatures // J. Phys. C: Solid State Phys. 1981 V. 14, L. 979-983.

42.  Busker G., Chroneos A., Grimes R. W. Solution mechanisms for dopant oxides in yttria // J. American Ceramics Soc. 1999. V. 82. P. 1553–1559.

43.  Morelon N-D., Ghaleb D., et al. A new empirical potential for simulating the formation of defects and their mobility in uranium dioxide // Phil. Mag. 2003. V. 83. P. 1533–1550.

44.  Gibbon P., Sutmann G.. Long-Range Interactions in Many-Particle Simulation, Quantum Simulations of Complex Many-Body Systems: From Theory to Algorithms, Lecture Notes, J. Grotendorst, D. Marx, A. Muramatsu (Eds.) // NIC Series. 2002. V. 10. P. 467–506.

45.  Ryabov V. А. Constant pressure–temperature molecular dynamics on a torus // Physics Letters A. 2006. V. 359. P. 61–65.

46.  Интегрирование Верлета // http://en. wikipedia. org/wiki/Verlet integration - Википедия. 2008. 6с.

47.  Allen М. Р., Tildesley D. puter simulations of liquids // M. P. Allen, D. J. Tildesley. New York: Oxford University Press Inc. 19c.

48.  Brent R. P. An algorithm with guaranteed convergence for finding a zero of a function // Computer Journal. 1971. V 14. P. 422–425.

49.  Nelder J. A., Mead R. A simplex method for function minimization // Computer Journal 1965. V. 7 P. 308–313.

50.  Kurosaki K., Yamada K. et al. Molecular dynamics study of mixed oxide fuel // J. Nucl. Mater. 2001. V. 294 P. 160-167.

51.  Basak C. B., Sengupta A. R., Kamath H. S. Classical molecular dynamics simulation of UO2 to predict thermophysical properties // J. of Alloys and Compounds. 2003. V. 360 P. 210–216.

52.  Поташников C. И., и др. Высокоскоростное моделиро­вание диффузии ионов урана и кислорода в UO2 // Труды отраслевого семинара «Реакторное материаловедение», Димитровград. 2005.

53.  Поташников C. И., и др. Поточно-параллельное моделирование диффузии в нанокристаллах // Труды XII Национальной конференции по росту кристаллов «НКРК-2006», Институт кристалло­гра­фии РАН (2006).

54.  Поташников C. И., Боярченков А. С. и др. Моделирование массопереноса в диоксиде урана методом молекулярной динамики с использованием графических процессоров // Международный научный журнал «Альтернативная энергетика и экология». 2007. Т. 5 С. 86–93. http://isjaee. *****/pdf/AEE0507/ISJAEE05-07_Potashnikov. pdf

55.  Potashnikov S. I., Boyarchenkov A. S. et al. Molecular dynamic modeling of mass transport and empirical fitting of pair potentials in nuclear oxide fuel using graphics processing units // Proceedings of the 8th Russian Conference on “Reactor Materials”, Dimitrovgrad. 2007.

56.  Majumdar S. N., Sengupta A. M. Thermo Physical and Thermo Mechanical Properties of Nuclear fuel for Thermal and Fast Reactor // IANCAS Bulletin on Nuclear Materials. 2005. V. 4. P. 226–236.

57.  Fritz I. J. Elastic properties of UO2 at high pressure // J. of Applied Physics. 1976. V 47. P. 4353–4358.

58.  Freeman G. C., Benson P. I., Dempsey E. L. Calculation of cohesive and surface energies of thorium and uranium oxides // J. of the American Ceramics Society. 1963. V. 46. P. 43-47.

59.  Browning P., Hyland G. J., Ralph J. The origin of the specific heat anomaly in solid urania // High Temperatures-High Pressures. 1983. V. 15. P. 169-178.

60.  Matzke Hj. Atomic transport properties in UO2 and mixed oxides (U, Pu)O2 // J. Chem. Soc. Faraday Trans. 1987. V. 83. P. 1121–1142.

61.  Potashnikov S. I., Nekrassov K. A. et al. Investigation of mechanisms of structural disordering of uranium dioxide with methods of molecular dynamics, lattice statics // Proceedings of the 8th Russian Conference on “Reactor Materials”, Dimitrovgrad (2007).

62.  Ralph J. Specific Heat of UO2 , ThO2 , PuO2 and the Mixed Oxides by Enthalpy Data Analysis // J. Chem. Soc. Faraday Trans. 1987. V. 83. P. 1253–1262.

63.  Ronchi C., Hyland G. J. Analysis of recent measurements of the heat capacity of uranium dioxide // J. of Alloys and Compounds 1994. V. 213/214. P. 159–168.

64.  Hutchings M. T. High-temperature studies of UO2 and ThO2 using neutron scattering techniques // J. Chem. Soc. Faraday Trans. 1987. V. 83. P. 1083–1103.

65.  Hiernaut J. P., Hyland G. J. Premelting transition in uranium dioxide // Int. J. Thermophys. 1993. V. 14. P. 259–283.

66.  Поташников С. И., Боярченков А. С. и др. Молекулярно-динамическое восстановление межчастичных потенциалов в диоксиде урана по тепловому расширению // Международный научный журнал «Альтернативная энергетика и экология» 2007. Т. 8. С. 43–52.

67.  NVIDIA Corp. CUDA Technical Training / NVIDIA Corp. 2701 San Tomas Expressway, Santa Clara, CA 95050: NVIDIA Corp. www. . 20 c.

68.  NVIDIA Corp. NVIDIA CUDA Compute Unified Device Architecture. Programming Guide Version 2.0 / NVIDIA Corp. 2701 San Tomas Expressway, Santa Clara, CA 95050: NVIDIA Corp. www. . 20 c.

69.  , Поташников графических про­цессоров и технологии CUDA для задач молекуляр­ной динамики // Вычислительные методы и программирование. Новые вычислительные технологии. http://num-meth. srcc. msu. su. 2008.

Приложение 1
Операторы и функции языка HLSL, использованные в курсе лекций

П.1.2. Типы данных

texture – двухмерный массив исходных данных (текстура), ячейки которого адресуются двумя координатами – парой чисел с плавающей точкой;

sampler – автоматическая процедура, осуществляющая выборку данных из входных текстур. Применение показано в примерах программ;

float2 – переменная, представляющая собой вектор из 2-х вещественных чисел одинарной точности; ис­пользуется для адресации ячеек входных текстур и рендер-цели;

float4 – переменная, представляющая собой вектор из 4-х вещественных чисел одинарной точности, либо функция, возвращаю­щая такой вектор в качестве значения; во втором варианте ис­пользуется при декларации пиксельных шейдеров;

void – функция, не возвращающая никаких значений; ис­пользуется при декларации вершинных шейдеров;

technique – декларация «техники», т. е. процедуры, которую могут запускать внешние программы, взаимо­дейст­вующие с графическим процессором; внутри этой процедуры задаются значения управляющих переменных, запускаются на исполнение вершинный и пиксельный шейдеры.

П.1.2. Встроенные функции

tex2D(объект типа sampler, координаты ячейки типа float2) – функция, возвращающая расположенный в ячейке с задан­ными координатами 4-вектор из текстуры, к которой «привязан» sampler;

П.1.3. Другие операторы и выражения

return – оператор, возвращающий значение функции, в теле которой он расположен;

PixelShader = compile ps_3_0 sum() – строка, задающая функцию sum() в качестве исполняемого пиксельного шейдера;

VertexShader = compile vs_3_0 transform() – строка, за­дающая функцию transform() в качестве исполняемого вершин­ного шейдера;

П.1.4. Системные переменные

ZEnable – имеет значение true, если используется Z-бу­фер, либо false, если он не используется. При физическом моделировании Z-буфер не используется, так что ZEnable = false;

CullMode – устанавливает режим отсечения невидимых треугольников. При физическом моделировании отсечения треугольников нет, так что CullMode = none.

Приложение 2
Процедуры CUDA, исполняемые на CPU

// Файл matrixMul. cu [68]

// Copyright 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 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 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;

}}

Из за большого объема этот материал размещен на нескольких страницах:
1 2 3 4 5 6 7 8 9 10 11