[CUDA编程]基础入门例程3

例程地址：https://github.com/TycoonL/awesome-cuda.git/

5.global_memory_static

这个例程用cudaMemcpyToSymbol实现了全局内存的使用。

在这段代码中，定义了两个设备静态全局内存变量cuda_x和cuda_y，分别为整型变量和整型数组。然后定义了一个kernel函数，该函数将cuda_x的值加到cuda_y数组的每个元素上，并打印出cuda_y数组的值。

在主函数中，定义了一个host数组，包含两个整数值。然后使用cudaMemcpyToSymbol函数将host数组的数据拷贝到cuda_y所指向的设备中，以便在kernel函数中能够访问。然后启动了一个只包含一个线程的CUDA核函数kernel，该核函数对cuda_y数组进行操作，并在每个元素上加上cuda_x的值。然后调用了cudaDeviceSynchronize函数等待所有设备上的核函数执行完毕。接下来，使用cudaMemcpyFromSymbol函数将cuda_y数组的值拷贝回host数组中，并打印出host数组的值。

值得注意的是，而cudaMemcpy函数不能拷贝静态全局内存变量或者常量内存变量，因此不适用于该情况。

//定义设备静态全局变量
__device__ int cuda_x = 1;
__device__ int cuda_y[2];

__global__
void kernel()
{
	cuda_y[0] += cuda_x;
	cuda_y[1] += cuda_x;
	printf("cuda_y1=%d,cuda_y2=%d\n", cuda_y[0], cuda_y[1]);
}

int main()
{
	int host[2] = {10, 20};
	//cudaMemcpy(cuda_y,host, sizeof(int) * 2,cudaMemcpyHostToDevice);
	cudaMemcpyToSymbol(cuda_y, host, sizeof(int) * 2);//
	kernel << <1, 1 >> > ();
	cudaDeviceSynchronize();
	cudaMemcpyFromSymbol(host, cuda_y, sizeof(int) * 2);
	printf("host1=%d,host2=%d\n", host[0], host[1]);

	while (1);   // 程序结束前保持运行，以便查看程序的输出
	return 0;
}

6.transpose

本例程实现了矩阵的转置：$ B_{i,j}=A_{j,i}$,并探究了合并、非合并读取内存的速度差异。

核函数：

__global__ void transpose(float* in, float* out, int width, int height) {
	int x = blockIdx.x * blockDim.x + threadIdx.x;
	int y = blockIdx.y * blockDim.y + threadIdx.y;

	if (x < width && y < height) {
		int index_in = y * width + x;
		int index_out = x * height + y;
		out[index_in] = __ldg(&in[index_out]);//实验3：out合并写入，in非合并访问 Execution time: 0.019456 ms
		//out[index_out] = __ldg(&in[index_in]);//实验4：out非合并写入，in合并访问 Execution time: 0.021504 ms
		//out[index_in] = in[index_out];//实验1：out合并写入，in非合并访问  Execution time: 0.019456 ms
		//out[index_out] = in[index_in];//实验2：out非合并写入，in合并访问  Execution time: 0.021504 ms
	}
}

在width = 128，height = 128，block_size(32, 32)的实验结果如上

• 合并写入的速度要更快
• __lgd():只读数据缓存的加载程序