CDUA 图形学 Texture Reference 实验

2019-03-02 20:48:23 阅读：247 来源： 互联网

标签：__ Reference int float 图形学 height width CDUA include

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

/*
	下面是对CUDA的texture object的实验，实验环境VS2017,CUDA 10，GTX 1060。
	作者：吕翔宇，部分代码改自CUDA手册
	E-mail:630056108@qq.com
	2019.3.2 20.44
*/


#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include<iostream>
#include<cmath>
#define __CUDACC__
#define __cplusplus
#include<texture_fetch_functions.h>
#include<cuda.h>
#include<cuda_texture_types.h>

texture<float, cudaTextureType2D, cudaReadModeElementType>texRef;
//转换的简单例子
__global__ void transformKernel(float *output1, float *output2, int width, int height, float theta) {
	//计算正常的纹理坐标
	unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
	unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;

	float u = x / (float)width;
	float v = y / (float)height;

	//纹理坐标变换
	u -= 0.5f;
	v -= 0.5f;
	float tu = u * cosf(theta) - v * sinf(theta) + 0.5f;
	float tv = v * cosf(theta) + u * sinf(theta) + 0.5f;

	//从纹理中读出，然后写出到全局内存中
	output1[y*width + x] = tex2D<float>(texRef, u, v);//未变换的
	output2[y*width + x] = tex2D<float>(texRef, tu, tv);//变换的
}

int main()
{
	int width = 16, height = 16;
	//在设备端分配CUDA数组
	cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
	cudaArray *cuArray;
	cudaMallocArray(&cuArray, &channelDesc, width, height);

	float *h_data;
	size_t size = width * height * sizeof(float);
	h_data = (float*)malloc(size);
	for (int i = width * height - 1; i > -1; i--)
		h_data[i] = 1 + (rand() % 10) / 1.0;
	std::cout << "纹理转换前内容：\n";
	for (int i = 0; i < height; i++) {
		for (int j = 0; j < width; j++) {
			std::cout << h_data[i*width + j] << "\t";
		}
		std::cout << "\n";
	}
	//将主机端内存搬到设备端的CUDA数组中去
	cudaMemcpyToArray(cuArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);

	//设置纹理引用
	texRef.addressMode[0] = cudaAddressModeWrap;
	texRef.addressMode[1] = cudaAddressModeWrap;
	texRef.filterMode = cudaFilterModeLinear;
	texRef.normalized = true;
	
	//将CUDA数组绑定到纹理引用上去
	cudaBindTextureToArray(&texRef, cuArray, &channelDesc);

	//定义转运接受设备结果的内存
	float *output1;
	cudaMalloc(&output1, size);
	float *output2;
	cudaMalloc(&output2, size);

	//生成内核
	dim3 dimBlock(16, 16);
	dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y - 1) / dimBlock.y);
	float angle = asinf(1.0 / 2.0);
	transformKernel <<<dimGrid, dimBlock >>> (output1, output2, width, height, angle);

	//输出结果
	float *out;
	out = (float*)malloc(size);
	cudaMemcpy(out, output1, size, cudaMemcpyDeviceToHost);
	std::cout << "未变换坐标纹理转换后的结果：\n";
	for (int i = 0; i < height; i++) {
		for (int j = 0; j < width; j++) {
			std::cout << out[i*width + j] << "\t";
		}
		std::cout << "\n";
	}
	cudaMemcpy(out, output2, size, cudaMemcpyDeviceToHost);
	std::cout << "变换坐标纹理转换后的结果：\n";
	for (int i = 0; i < height; i++) {
		for (int j = 0; j < width; j++) {
			std::cout << out[i*width + j] << "\t";
		}
		std::cout << "\n";
	}

	//释放设备内存
	cudaFreeArray(cuArray);
	cudaFree(output1);
	cudaFree(output2);

	system("pause");
	return 0;
}

结果

效果与纹理对象相同

标签：__,Reference,int,float,图形学,height,width,CDUA,include
来源： https://blog.csdn.net/lvxiangyu11/article/details/88081100

本站声明： 1. iCode9 技术分享网（下文简称本站）提供的所有内容，仅供技术学习、探讨和分享；
2. 关于本站的所有留言、评论、转载及引用，纯属内容发起人的个人观点，与本站观点和立场无关；
3. 关于本站的所有言论和文字，纯属内容发起人的个人观点，与本站观点和立场无关；
4. 本站文章均是网友提供，不完全保证技术分享内容的完整性、准确性、时效性、风险性和版权归属；如您发现该文章侵犯了您的权益，可联系我们第一时间进行删除；
5. 本站为非盈利性的个人网站，所有内容不会用来进行牟利，也不会利用任何形式的广告来间接获益，纯粹是为了广大技术爱好者提供技术内容和技术思想的分享性交流网站。

ICode9

CDUA 图形学 Texture Reference 实验