CDUA 图形学 Surface Reference实验

下面是对Surfance Reference的实验，代码改自 CUDA C PROGRAMMING GUIDE (PG-02829-001_v10.0 | October 2018) p54，因为这里比较贴近计算机图形学，故移到图形学中去。

 

©版权所有！

/*
	下面是对CUDA的Surface Reference的实验，实验环境VS2017,CUDA 10，GTX 1060
	作者：吕翔宇，部分代码改自CUDA手册
	E-mail:630056108@qq.com
	2019.3.4 19:01
*/
#define __CUDACC__
#define __cplusplus
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include<iostream>
#include<cmath>
#include<cstdlib>
#include<cuda.h>

//定义二维表层
surface<void, 2> inputSurfRef;
surface<void, 2> outputSurfRef;

//简单的拷贝内核
__global__ void copyKernel(int width, int height) {
	//计算纹理坐标
	unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
	unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
	if (x < width&&y < height) {
		uchar4 data;
		//读取surface数据
		surf2Dread(&data, inputSurfRef, x * 4, y);
		//写surface数据
		surf2Dwrite(data, outputSurfRef, x * 4, y);
	}
}

//主机代码
int main() {
	std::ios::sync_with_stdio(false);//习惯性关闭同步流，看不懂的请忽略。
	//注意这里使用cudaChannelFormatKindUnsigned，即处理的是unsigned int
	//准备实验参数
	const int width = 16;
	const int height = 16;
	const size_t size = sizeof(unsigned int)*width*height;
	unsigned int *h_data;
	h_data = (int unsigned*)malloc(size);
	//init
	std::cout << "初始数据：\n";
	for (int i = 0; i < height; i++) {
		for (int j = 0; j < width; j++) {
			h_data[i*width + j] = rand() % 100;
			std::cout << h_data[i*width + j] << "\t";
		}
		std::cout << "\n";
	}

	//设备端分配CUDA数组
	cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);
	cudaArray *cuInputArray;
	cudaMallocArray(&cuInputArray, &channelDesc, width, height, cudaArraySurfaceLoadStore);
	cudaArray *cuOutputArray;
	cudaMallocArray(&cuOutputArray, &channelDesc, width, height, cudaArraySurfaceLoadStore);

	//搬运内存
	cudaMemcpyToArray(cuInputArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);

	//将CUDA数据绑定上surface reference
	cudaBindSurfaceToArray(inputSurfRef, cuInputArray);
	cudaBindSurfaceToArray(outputSurfRef, cuOutputArray);

	//定义内核参数
	dim3 dimBlock(16, 16);
	dim3 dimGrid(
		(width + dimBlock.x - 1) / dimBlock.x,
		(height + dimBlock.y - 1) / dimBlock.y
	);

	copyKernel << <dimGrid, dimBlock >> > (width, height);

	//输出结果
	memset(h_data, 0, size);
	//cudaMemcpyFromArray(h_data, cuOutputArray, 0, 0, size, cudaMemcpyDeviceToHost);
	cudaMemcpy(h_data, cuOutputArray, size, cudaMemcpyDeviceToHost);
	std::cout << "处理后数据：\n";
	for (int i = 0; i < height; i++) {
		for (int j = 0; j < width; j++) {
			h_data[i*width + j] = rand() % 10;
			std::cout << h_data[i*width + j] << "\t";
		}
		std::cout << "\n";
	}

	//释放设备内存
	cudaFreeArray(cuInputArray);
	cudaFreeArray(cuOutputArray);
	free(h_data);

	system("pause");
	return 0;
}

结果

？？？