Как использовать Программирование CUDA для вычисления и обработки правильного числа

---

Bandwidth test - тест пропускной способности памяти.

Особенно это важно для возможностей PCIE. У разных МБ есть разные возможности PCIE.

Производительность адаптера CUDA зависит от возможностей PCIE. Это может быть узким местом производительности.

В следующих упражнениях по программированию необходимо вычислить количество тактовых циклов, необходимых для вычислений, и используемую пропускную способность памяти.

(1) распараллеливание в программах - использование 256 потоков

(2) улучшение режимов доступа к памяти

(3) тестирование распараллеливания с помощью 512/1024

(4) использование блоков в вычислениях

(5) использование общей памяти

(6) улучшение вычислительной oerfdormance с помощью алгоритма Treesum

(7) решение проблемы конфликта полосы памяти, возникающей при применении алгоритма Treesum с общей памятью

Что я уже пробовал:

Однако в моем собственном примере кодирования есть некоторые ошибки -

#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
#include <time.h>
#include <math.h>

int main()
{
	float *a, *b, *c, *d;
	int n = 1000;

	if (!InitCUDA())
		return 0;
	a = (float*)malloc(sizeof(float)* n * n);
	b = (float*)malloc(sizeof(float)* n * n);
	c = (float*)malloc(sizeof(float)* n * n);
	d = (float*)malloc(sizeof(float)* n * n);

	srand(0);
	matgen(a, n, n);
	matgen(b, n, n);

	clock_t time = matmultCUDA(a, n, b, n, c, n, n);
	matmult(a, n, b, n, d, n, n);
	compare_mat(c, n, d, n, n);

	double sec = (double)time / CLOCKS_PER_SEC;

	printf("Time used: %.2f (%.2lf GFLOPS)\n", sec, 2.0 * n * n * n / (sec * 1E9));

	return 0;
}
void matgen(float* a, int lda, int n)
{
	int i, j;
	for (i = 0; i < n; i++)
	{
		for (j = 0; j < n; j++)
		{
			a[i* lda + j] = (float)rand() / RAND_MAX + (float)rand() / (RAND_MAX * RAND_MAX);
		}
	}
}

void matmult(const float* a, int lda, const float* b, int ldb, float* c, int ldc, int n)
{
	int i, j, k;

	for (i = 0; i< n; i++)
	{
		for (j = 0; j < n; j++)
		{
			double t = 0;
			for (k = 0; k < n; k++) {
				t += a[i* lda + k] * b[k * ldb + j];
			}
			c[i* ldc + j] = t;
		}
	}

	void compare_mat(const float* a, int lda, const float* b, int ldb, int n)
	{
		float max_err = 0;
		float average_err = 0; inti, j;
		for (i = 0; i< n; i++)
		{
			for (j = 0; j < n; j++)
			{
				if (b[i* ldb + j] != 0)
				{
					float err = fabs((a[i* lda + j] - b[i* ldb + j]) / b[i* ldb + j]);
					if (max_err< err) max_err = err;
					average_err += err;
				}
			}
		}
		printf("Max error: %g Average error:%g\n", max_err, average_err / (n * n));
	}

#define NUM_THREADS 256

clock_t matmultCUDA(const float* a, int lda, const float* b, int ldb, float* c, int ldc, int n)
{
	float *ac, *bc, *cc;
	clock_tstart, end;

	start = clock();
	cudaMalloc((void**)&ac, sizeof(float)* n * n);
	cudaMalloc((void**)&bc, sizeof(float)* n * n);
	cudaMalloc((void**)&cc, sizeof(float)* n * n);

	cudaMemcpy2D(ac, sizeof(float)* n, a, sizeof(float)* lda, sizeof(float)* n, n, cudaMemcpyHostToDevice);
	cudaMemcpy2D(bc, sizeof(float)* n, b, sizeof(float)* ldb, sizeof(float)* n, n, cudaMemcpyHostToDevice);

	intblocks = (n + NUM_THREADS - 1) / NUM_THREADS;
	matMultCUDA<<<blocks * n, NUM_THREADS >>>(ac, n, bc, n, cc, n, n);

	cudaMemcpy2D(c, sizeof(float)* ldc, cc, sizeof(float)* n, sizeof(float)* n, n, cudaMemcpyDeviceToHost);

	cudaFree(ac);
	cudaFree(bc);
	cudaFree(cc);

	end = clock();
	return end - start;
}

__global__ static void matMultCUDA(const float* a, size_t lda, const float* b, size_t ldb, float* c, size_t ldc, int n)
{
	__shared__ float
	matA[BLOCK_SIZE][BLOCK_SIZE];
	__shared__ float matB[BLOCK_SIZE][BLOCK_SIZE];
	constinttidc = threadIdx.x;
	constinttidr = threadIdx.y;
	constintbidc = blockIdx.x* BLOCK_SIZE;
	constintbidr = blockIdx.y* BLOCK_SIZE;
	int i, j;

	float results = 0;
	float comp = 0;

	for (j = 0; j < n; j += BLOCK_SIZE)
	{
		matA[tidr][tidc] = a[(tidr + bidr) * lda + tidc + j];
		matB[tidr][tidc] = b[(tidr + j) * ldb + tidc + bidc];

		__syncthreads();

		for (i = 0; i< BLOCK_SIZE; i++)
		{
			float t; comp -= matA[tidr][i] * matB[i][tidc];
			t = results - comp; comp = (t - results) + comp; results = t;
		}

		__syncthreads();
	}

	c[(tidr + bidr) * ldc + tidc + bidc] = results;
}

---