tex2D使用学习

1. 背景：

        项目中使用到了纹理进行插值的加速，因此记录一些自己在学习tex2D的一些过程

2. 代码：

        

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <assert.h>
#include <stdio.h>
#include <iostream>
#include <cuda_fp16.h>
#include <vector>void Data2Half(half* pDst, const int16_t* pSrc, const int Ndots);
static __global__ void Tex2DTest(cudaTextureObject_t p_rf_data, float* pfRes1, float* pfRes2);static __global__ void data2half(half* pDst, const int16_t* pSrc, const int Ndots)
{const int tid = blockIdx.x * blockDim.x + threadIdx.x;if (tid >= Ndots)return;pDst[tid] = __short2half_rn(pSrc[tid]);
}cudaTextureObject_t m_tex   = 0;
cudaArray* m_pRFData        = nullptr;
int16_t* m_i16RFDataBuffer  = nullptr; // 设备端的RF数据
half* m_pHalfRFDataCache    = nullptr; // 转换为半浮点型的RF数据缓存，用于将SHORT类型转换为FLOAT类型int main()
{const int nRx     = 2;const int Nsample = 2;const int IQ      = 1;cudaError_t error;cudaChannelFormatDesc channelDesc = cudaCreateChannelDescHalf();error                             = cudaMallocArray(&m_pRFData, &channelDesc, nRx * IQ, Nsample, cudaArrayTextureGather);assert(m_pRFData);cudaResourceDesc texRes;memset(&texRes, 0, sizeof(cudaResourceDesc));texRes.resType         = cudaResourceTypeArray;texRes.res.array.array = m_pRFData;cudaTextureDesc texDescr;memset(&texDescr, 0, sizeof(cudaTextureDesc));texDescr.normalizedCoords = false;texDescr.filterMode       = cudaFilterModeLinear;  // 这里很重要texDescr.addressMode[0]   = cudaAddressModeBorder;texDescr.addressMode[1]   = cudaAddressModeBorder;error = cudaCreateTextureObject(&m_tex, &texRes, &texDescr, NULL);//int16_t pi16Src[nRx * Nsample * IQ] = {1, 11, 2, 22,//                                    3, 33, 4, 44, //                                    5, 55, 6, 66, //                                    7, 77, 8, 88};//int16_t pi16Src[nRx * Nsample * IQ] = { 1, 11, 2, 22,//                                        3, 33, 4, 44};int16_t pi16Src[nRx * Nsample * IQ] = { 1,2,3,4 };error = cudaMalloc(&m_i16RFDataBuffer, sizeof(int16_t) * nRx * IQ * Nsample);error = cudaMemcpy(m_i16RFDataBuffer, pi16Src, sizeof(int16_t) * nRx * IQ * Nsample, cudaMemcpyHostToDevice);error = cudaMalloc(&m_pHalfRFDataCache, sizeof(half) * nRx * IQ * Nsample);Data2Half(m_pHalfRFDataCache, m_i16RFDataBuffer, nRx * IQ * Nsample);error = cudaMemcpy2DToArray(m_pRFData, 0, 0, m_pHalfRFDataCache, sizeof(half) * nRx * IQ, sizeof(half) * nRx * IQ, Nsample, cudaMemcpyDeviceToDevice);float* pf_res1 = nullptr;float* pf_res2 = nullptr;error = cudaMalloc(&pf_res1, nRx * Nsample * sizeof(float)); cudaMemset(pf_res1, 0, nRx * Nsample * sizeof(float));error = cudaMalloc(&pf_res2, nRx * Nsample * sizeof(float)); cudaMemset(pf_res2, 0, nRx * Nsample * sizeof(float));error = cudaGetLastError();dim3 block_dim = dim3(1, 1);dim3 grid_dim  = dim3(1, 1);Tex2DTest << <grid_dim, block_dim >> > (m_tex, pf_res1, pf_res2);cudaDeviceSynchronize();std::vector<float> vf_res_1(nRx * Nsample, 0);std::vector<float> vf_res_2(nRx * Nsample, 0);cudaMemcpy(vf_res_1.data(), pf_res1, sizeof(float) * vf_res_1.size(), cudaMemcpyDeviceToHost);cudaMemcpy(vf_res_2.data(), pf_res2, sizeof(float) * vf_res_2.size(), cudaMemcpyDeviceToHost);return 0;
}void Data2Half(half* pDst, const int16_t* pSrc, const int Ndots)
{dim3 block = dim3(512, 1);dim3 grid = dim3((Ndots - 1) / block.x + 1, 1);data2half << < grid, block >> > (pDst, pSrc, Ndots);
}static __global__ void Tex2DTest(cudaTextureObject_t p_rf_data, float *pfRes1, float *pfRes2)
{for (size_t y = 0; y < 2; ++y){for (size_t x = 0; x < 2; ++x) {float value = tex2D<float>(p_rf_data, x,     y);//pfRes1[y * 4 + y] = printf("x: %f\n", value);}}
}

3. 输出分析：

可以看到执行结果是

为什么呢？

原因是因为tex2D插值导致的，上面测试数据是

1  2

3   4

那在进行插值的时候会变成

0  0   0   0

0   1   2  0

0   3   4  0

每个点的输出都是当前前和左上角3个点进行平均计算出来的

比如第一个输出计算为：（1 + 0 + 0 + 0）/4 = 0.25

最后一个输出的计算为：（1 + 2 + 3 + 4） / 4 = 2.5

4. 问题

        上面只是单独数据实数点的计算，如果我的数据集合是复数怎么办？

        比如一组2 * 2大小的数据对

        （1， 2， 3， 4;

           5,   6， 7， 8）

        数据实际表示含义是

         （1 + j * 2,   3 + j * 4;

            5 + j * 6,   7 + j * 8）

        这种情况下怎么做到正确插值呢，比如第一个实数点的输出结果应该是

         （1 + 0 + 0 + 0）/ 4

           最后一个实数点的输出应该是：

            （1 + 3 + 5 + 7） / 4

           同理，最后一个虚数点的输出应该是：
           （2 + 4 + 6 + 8）/ 4

5. 解决

         

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <assert.h>
#include <stdio.h>
#include <iostream>
#include <cuda_fp16.h>
#include <vector>void Data2Half(half* pDst, const int16_t* pSrc, const int Ndots);
static __global__ void Tex2DTest(cudaTextureObject_t p_rf_data, float* pfRes1, float* pfRes2);static __global__ void data2half(half* pDst, const int16_t* pSrc, const int Ndots)
{const int tid = blockIdx.x * blockDim.x + threadIdx.x;if (tid >= Ndots)return;pDst[tid] = __short2half_rn(pSrc[tid]);
}cudaTextureObject_t m_tex = 0;
cudaArray* m_pRFData = nullptr;
int16_t* m_i16RFDataBuffer = nullptr; // 设备端的RF数据
half* m_pHalfRFDataCache = nullptr; // 转换为半浮点型的RF数据缓存，用于将SHORT类型转换为FLOAT类型using namespace std;int main()
{const int nRx = 2;const int Nsample = 2;const int IQ = 2;cudaError_t error;cudaChannelFormatDesc channelDesc = cudaCreateChannelDescHalf2();error = cudaMallocArray(&m_pRFData, &channelDesc, nRx, Nsample, cudaArrayTextureGather);assert(m_pRFData);cudaResourceDesc texRes;memset(&texRes, 0, sizeof(cudaResourceDesc));texRes.resType = cudaResourceTypeArray;texRes.res.array.array = m_pRFData;cudaTextureDesc texDescr;memset(&texDescr, 0, sizeof(cudaTextureDesc));texDescr.normalizedCoords = false;texDescr.filterMode = cudaFilterModeLinear;  // 这里很重要texDescr.addressMode[0] = cudaAddressModeBorder;texDescr.addressMode[1] = cudaAddressModeBorder;error = cudaCreateTextureObject(&m_tex, &texRes, &texDescr, NULL);//int16_t pi16Src[nRx * Nsample * IQ] = {1, 11, 2, 22,//                                    3, 33, 4, 44, //                                    5, 55, 6, 66, //                                    7, 77, 8, 88};//int16_t pi16Src[nRx * Nsample * IQ] = { 1, 11, 2, 22,//                                        3, 33, 4, 44};int16_t pi16Src[nRx * Nsample * IQ] = { 1, 2, 3, 4,5, 6, 7, 8 };error = cudaMalloc(&m_i16RFDataBuffer, sizeof(int16_t) * nRx * IQ * Nsample);error = cudaMemcpy(m_i16RFDataBuffer, pi16Src, sizeof(int16_t) * nRx * IQ * Nsample, cudaMemcpyHostToDevice);error = cudaMalloc(&m_pHalfRFDataCache, sizeof(half) * nRx * IQ * Nsample);Data2Half(m_pHalfRFDataCache, m_i16RFDataBuffer, nRx * IQ * Nsample);error = cudaMemcpy2DToArray(m_pRFData, 0, 0, m_pHalfRFDataCache, sizeof(half2) * nRx, sizeof(half2) * nRx, Nsample, cudaMemcpyDeviceToDevice);float* pf_res1 = nullptr;float* pf_res2 = nullptr;error = cudaMalloc(&pf_res1, nRx * Nsample * sizeof(float)); cudaMemset(pf_res1, 0, nRx * Nsample * sizeof(float));error = cudaMalloc(&pf_res2, nRx * Nsample * sizeof(float)); cudaMemset(pf_res2, 0, nRx * Nsample * sizeof(float));error = cudaGetLastError();dim3 block_dim = dim3(1, 1);dim3 grid_dim  = dim3(1, 1);Tex2DTest << <grid_dim, block_dim >> > (m_tex, pf_res1, pf_res2);cudaDeviceSynchronize();std::vector<float> vf_res_1(nRx * Nsample, 0);std::vector<float> vf_res_2(nRx * Nsample, 0);cudaMemcpy(vf_res_1.data(), pf_res1, sizeof(float) * vf_res_1.size(), cudaMemcpyDeviceToHost);cudaMemcpy(vf_res_2.data(), pf_res2, sizeof(float) * vf_res_2.size(), cudaMemcpyDeviceToHost);return 0;
}void Data2Half(half* pDst, const int16_t* pSrc, const int Ndots)
{dim3 block = dim3(512, 1);dim3 grid = dim3((Ndots - 1) / block.x + 1, 1);data2half << < grid, block >> > (pDst, pSrc, Ndots);
}static __global__ void Tex2DTest(cudaTextureObject_t p_rf_data, float* pfRes1, float* pfRes2)
{for (size_t y = 0; y < 2; ++y){for (size_t x = 0; x < 2; ++x){float2 value = tex2D<float2>(p_rf_data, x, y);//pfRes1[y * 4 + y] = printf("x: %f, y: %f", value.x, value.y);// printf("x: %f, y: %f\n", value.x, value.y);}printf("\n");}
}

其实关键是在tex2D的构造

然后按照half2的方式进行排布就好了

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.mzph.cn/news/186444.shtml

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

Maven的安装和使用

Maven的安装和使用

Maven是一个基于项目对象模型&#xff08;POM&#xff09;&#xff0c;可以管理项目构建、依赖管理、项目报告等的工具&#xff0c;使构建Java项目更容易。可以说Maven是一个项目管理和构建工具&#xff0c;它可以从管理项目的角度出发&#xff0c;将开发过程中的需求纳入进来&…
阅读更多...
FFmpeg架构全面分析

FFmpeg架构全面分析

一、简介 它的官网为&#xff1a;https://ffmpeg.org/&#xff0c;由Fabrice Bellard&#xff08;法国著名程序员Born in 1972&#xff09;于2000年发起创建的开源项目。该人是个牛人&#xff0c;在很多领域都有很大的贡献。 FFmpeg是多媒体领域的万能工具。只要涉及音视频领…
阅读更多...
软文推广如何自然融入品牌？媒介盒子有妙招

软文推广如何自然融入品牌？媒介盒子有妙招

软文推广作为一种柔性推广方式&#xff0c;能将品牌信息融入到用户日常浏览的内容中&#xff0c;让用户不知不觉接触品牌&#xff0c;从而产生好感&#xff0c;这种方式既可以避免广告带来的反感&#xff0c;又可以提高广告的有效性。那么在推广中应该如何自然融入品牌信息呢&a…
阅读更多...
leetCode 78.子集 + 回溯算法 + 图解

leetCode 78.子集 + 回溯算法 + 图解

给你一个整数数组 nums &#xff0c;数组中的元素 互不相同 。返回该数组所有可能的子集&#xff08;幂集&#xff09;。解集 不能 包含重复的子集。你可以按 任意顺序 返回解集 示例 1&#xff1a; 输入&#xff1a;nums [1,2,3] 输出&#xff1a;[[],[1],[2],[1,2],[3],[1…
阅读更多...
11月30日作业

11月30日作业

设计一个Per类&#xff0c;类中包含私有成员:姓名、年龄、指针成员身高、体重&#xff0c;再设计一个Stu类&#xff0c;类中包含私有成员:成绩、Per类对象p1&#xff0c;设计这两个类的构造函数、析构函数和拷贝构造函数 #include <iostream>using namespace std;class …
阅读更多...
二分算法（整数二分、浮点数二分）

二分算法（整数二分、浮点数二分）

文章目录 二分一、整数二分&#xff08;一&#xff09;整数二分思路&#xff08;二&#xff09;整数二分算法模板1.左查找&#xff08;寻找左侧边界&#xff09;2.右查找&#xff08;寻找右侧边界&#xff09;3.总模板 &#xff08;三&#xff09;题目&#xff1a;数的范围 二、…
阅读更多...
Luminar Neo1.16.0(ai智能图像处理)

Luminar Neo1.16.0(ai智能图像处理)

Luminar Neo是一款ai智能图像编辑软件&#xff0c;它专注于使用人工智能技术来实现对照片的快速、高效和创造性的编辑。 具体来说&#xff0c;Luminar Neo可以自动移除景观或旅行照片中令人分心的元素&#xff0c;例如电话线、电线杆等&#xff0c;从而增强照片的整体质量。同…
阅读更多...
传统家装“死气沉沉”？VR智慧家装提供VR可视化方案

传统家装“死气沉沉”？VR智慧家装提供VR可视化方案

传统家装市场虽然处于成熟期&#xff0c;但是对于装修小白的户主来说&#xff0c;难以解决的痛点依旧还有很多。很多家装公司所谓的设计师&#xff0c;不一定全都具备设计知识&#xff0c;也不懂得从客户的需求出发&#xff0c;多重因素导致家装行业“死气沉沉”。 为了打破装修…
阅读更多...
JUC并发编程 01——多线程基础知识

JUC并发编程 01——多线程基础知识

一.线程应用 异步调用 以调用方角度来讲&#xff0c;如果 需要等待结果返回&#xff0c;才能继续运行就是同步 不需要等待结果返回&#xff0c;就能继续运行就是异步 应用 比如在项目中&#xff0c;视频文件需要转换格式等操作比较费时&#xff0c;这时开一个新线程处理视…
阅读更多...
借助ETL快速查询金蝶云星空表单信息

借助ETL快速查询金蝶云星空表单信息

随着数字化转型的加速&#xff0c;企业信息化程度越来越高&#xff0c;大量的数据产生并存储在云端&#xff0c;需要进行有效的数据管理和查询。金蝶云星空是金蝶云旗下的一款云ERP产品&#xff0c;为企业提供了完整的业务流程和数据管理功能&#xff0c;因此需要进行有效的数据…
阅读更多...
.NET开源的处理分布式事务的解决方案

.NET开源的处理分布式事务的解决方案

前言 在分布式系统中&#xff0c;由于各个系统服务之间的独立性和网络通信的不确定性&#xff0c;要确保跨系统的事务操作的最终一致性是一项重大的挑战。今天给大家推荐一个.NET开源的处理分布式事务的解决方案基于 .NET Standard 的 C# 库&#xff1a;CAP。 CAP项目介绍 CA…
阅读更多...
homeassistant 随笔

homeassistant 随笔

1.使用mushroom-strategy自动生成ui&#xff0c;隐藏中文ares&#xff0c;名字为区域的拼音&#xff0c;例如显示厨房则真实名字为chu_fang 隐藏图片中的工作室 代码为&#xff1a;
阅读更多...
基于SpringBoot的在线视频教育平台的设计与实现

基于SpringBoot的在线视频教育平台的设计与实现

摘 要 随着科学技术的飞速发展&#xff0c;各行各业都在努力与现代先进技术接轨&#xff0c;通过科技手段提高自身的优势&#xff1b;对于在线视频教育平台当然也不能排除在外&#xff0c;随着网络技术的不断成熟&#xff0c;带动了在线视频教育平台&#xff0c;它彻底改变了过…
阅读更多...
C语言贪吃蛇（有详细注释）

C语言贪吃蛇（有详细注释）

这个贪吃蛇是在比特特训营里学到的&#xff0c;同时我还写了用EasyX图形库实现的图形化贪吃蛇&#xff0c;含有每个函数的实现以及游戏中各种细节的讲解&#xff0c;感兴趣的可以去看一看。 贪吃蛇小游戏 实现效果 以下就是源码&#xff0c;感兴趣的小伙伴可以cv自己玩一玩改…
阅读更多...
Windows系列：windows server 2016 下域环境的搭建（完整版）

Windows系列：windows server 2016 下域环境的搭建（完整版）

windows server 2016 下域环境的搭建&#xff08;完整版&#xff09; windows server 2016 下域环境的搭建在搭建之前简单介绍一下基础知识&#xff1a;一、环境介绍 &#xff1a;1.这里用拓扑图进行展示&#xff1a;2.所有环境配置如下 二、搭建主域&#xff1a;一. 创建主域1…
阅读更多...
Java+SSM springboot+MySQL家政服务预约网站设计en24b

Java+SSM springboot+MySQL家政服务预约网站设计en24b

随着社区居民对生活品质的追求以及社会老龄化的加剧&#xff0c;社区居民对家政服务的需求越来越多&#xff0c;家政服务业逐渐成为政府推动、扶持和建设的重点行业。家政服务信息化有助于提高社区家政服务的工作效率和质量。 本次开发的家政服务网站是一个面向社区的家政服务网…
阅读更多...
华为云之云桌面Workspace的使用体验

华为云之云桌面Workspace的使用体验

华为云之云桌面Workspace的使用体验 一、云桌面Workspace介绍1.云桌面简介2.云桌面特点3. 云桌面应用场景①远程移动办公②协同办公③安全办公④公用终端⑤图形制作渲染 二、本次实践介绍1. 本次实践目的2. 本次实践环境 三、购买云桌面1. 进入华为云的云桌面购买界面2. 选择购…
阅读更多...
C# Onnx 百度飞桨开源PP-YOLOE-Plus目标检测

C# Onnx 百度飞桨开源PP-YOLOE-Plus目标检测

目录 效果 模型信息 项目 代码 下载 C# Onnx 百度飞桨开源PP-YOLOE-Plus目标检测 效果 模型信息 Inputs ------------------------- name&#xff1a;image tensor&#xff1a;Float[1, 3, 640, 640] name&#xff1a;scale_factor tensor&#xff1a;Float[1, 2] ----…
阅读更多...
服务器之间的conda环境迁移

服务器之间的conda环境迁移

有的时候python环境中可能包含了我们编译好的很多库文件&#xff0c;如果在别的服务器想直接使用环境就会比较困难些。最好的办法就是直接迁移环境。而传统的迁移方法导出“*.yaml”环境配置的这种方法&#xff0c;实际是需要重新安装环境&#xff0c;对于这种安装好的环境是不…
阅读更多...
数据结构：图文详解顺序表的各种操作（新增元素，查找元素，删除元素，给指定位置元素赋值）

数据结构：图文详解顺序表的各种操作（新增元素，查找元素，删除元素，给指定位置元素赋值）

目录 一.顺序表的概念 二.顺序表的实现 新增元素 默认尾部新增 指定位置添加元素 查找元素 查找是否存在 查找元素对应的位置 查找指定位置对应的元素 删除元素 获取顺序表长度 清空顺序表 一.顺序表的概念 在线性数据结构中&#xff0c;我们一般分为俩类&#xf…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023