tex2D使用学习

news2024/9/22 17:35:10

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.coloradmin.cn/o/1272948.html

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

相关文章

XTU OJ 1339 Interprime 学习笔记

XTU OJ 1339 Interprime 学习笔记

链接 传送门 代码 #include<bits/stdc.h> using namespace std;const int N1e610; //78498 我计算了一下&#xff0c;6个0的范围内有这么多个素数&#xff0c;所以开这么大的数组存素数 //计算的代码是一个循环 int prime[80000]; int a[N],s[N];//s数组是前缀和数组b…
阅读更多...
自定义链 SNAT / DNAT 实验举例

自定义链 SNAT / DNAT 实验举例

参考原理图 实验前的环境搭建 1. 准备三台虚拟机&#xff0c;定义为内网&#xff0c;外网以及网卡服务器 2. 给网卡服务器添加网卡 3. 将三台虚拟机的防火墙和安全终端全部关掉 systemctl stop firewalld && setenforce 0 4. 给内网虚拟机和外网虚拟机 yum安装 httpd…
阅读更多...
Echarts的引入使用

Echarts的引入使用

ECharts文档 1.下载并引入Echarts 2.准备一个具备大小的DOM容器 3.初始化echarts实例对象 4.指定配置项和数据(option) 5.将配置项设置给echarts实例对象 最后是一个js文件 echarts的引入 1.引入echarts - js 文件 <script src"js/echarts.min.js"></scri…
阅读更多...
「Linux」使用C语言制作简易Shell

「Linux」使用C语言制作简易Shell

&#x1f4bb;文章目录 &#x1f4c4;前言简易shell实现shell的概念系统环境变量shell的结构定义内建命令完整代码 &#x1f4d3;总结 &#x1f4c4;前言 对于很多学习后端的同学来讲&#xff0c;学习了C语言&#xff0c;发现除了能写出那个经典的“hello world”以外&#xff…
阅读更多...
XUbuntu22.04之OBS30.0设置录制音频降噪(一百九十六)

XUbuntu22.04之OBS30.0设置录制音频降噪(一百九十六)

简介&#xff1a; CSDN博客专家&#xff0c;专注Android/Linux系统&#xff0c;分享多mic语音方案、音视频、编解码等技术&#xff0c;与大家一起成长&#xff01; 优质专栏&#xff1a;Audio工程师进阶系列【原创干货持续更新中……】&#x1f680; 优质专栏&#xff1a;多媒…
阅读更多...
LabVIEWL实现鸟巢等大型结构健康监测

LabVIEWL实现鸟巢等大型结构健康监测

LabVIEWL实现鸟巢等大型结构健康监测 管理国家地震防备和减灾的政府机构中国地震局(CEA)选择了七座新建的巨型结构作为结构健康监测(SHM)技术的测试台。这些标志性建筑包括北京2008年夏季奥运会场馆&#xff08;包括北京国家体育场和北京国家游泳中心&#xff09;、上海104层的…
阅读更多...
QML学习一、GridView的使用和增加添加动画、删除动画

QML学习一、GridView的使用和增加添加动画、删除动画

一、效果预览 二、源码分享 import QtQuick import QtQuick.ControlsApplicationWindow {visible: truewidth: 640height: 480title: "Test"property int cnt:cnt model.countListModel{id:modelListElement{index:0}ListElement{index:1}ListElement{index:2}List…
阅读更多...
csapp-linklab之第3阶段“输出学号”实验报告（强弱符号）

csapp-linklab之第3阶段“输出学号”实验报告（强弱符号）

题目 新建一个phase3_patch.o&#xff0c;使其与main.o和phase3.o链接后&#xff0c;运行输出自己的学号&#xff1a; $ gcc -o linkbomb main.o phase3.o phase3_patch.o $ ./linkbomb $学号 提示 利用符号解析中的强弱符号规则。&#xff08;COOKIE字符串未初始化&#xff…
阅读更多...
单片机AVR单片机病房控制系统设计+源程序

单片机AVR单片机病房控制系统设计+源程序

一、系统方案 设计一个可容8张床位的病房呼叫系统。要求每个床位都有一个按钮&#xff0c;当患者需要呼叫护士时&#xff0c;按下按钮&#xff0c;此时护士值班室内的呼叫系统板上显示该患者的床位号&#xff0c;并蜂鸣器报警。当护士按下“响应”键时&#xff0c;结束当前呼叫…
阅读更多...
【无标题】读transformer

【无标题】读transformer

这里写目录标题 transformerabstractconclusionintroductionbackground注意力机制mlptransformer和RNN传递序列信息embedding之后维度越大的向量归一化后其单个值就越小&#xff0c;乘个根号512position encoding加入时序信息 transformer abstract 编码器和解码器的架构 处理…
阅读更多...
强化学习中的Q学习

强化学习中的Q学习

Q学习&#xff08;Q-Learning&#xff09;是强化学习中的一种基于值的学习方法&#xff0c;用于在有限马尔可夫决策过程&#xff08;MDP&#xff09;中学习最优的动作策略。Q学习主要用于离散状态和离散动作的问题。 以下是Q学习的基本概念和步骤&#xff1a; Q-Value&#xf…
阅读更多...
程序员也需要养生——程序员睡不好，重视一下你的情绪吧

程序员也需要养生——程序员睡不好，重视一下你的情绪吧

程序员也需要养生——程序员睡不好&#xff0c;重视一下你的情绪吧 睡眠是一个复杂的系统工程&#xff0c;可以促进生长发育&#xff0c;修复受损的组织。促进大脑细胞的修复等等。在情绪的失调会影响到我们的睡眠状况。 一、心情差&#xff0c;压力大&#xff0c;睡不好跟这…
阅读更多...
XXL-Job详解（一）：组件架构

XXL-Job详解（一）：组件架构

目录 XXL-Job特性系统组成架构图调度模块剖析任务 “运行模式” 剖析执行器 XXL-Job XXL-JOB是一个分布式任务调度平台&#xff0c;其核心设计目标是开发迅速、学习简单、轻量级、易扩展。现已开放源代码并接入多家公司线上产品线&#xff0c;开箱即用。 特性 1、简单&#…
阅读更多...
【驱动】串口驱动分析(二)-tty core

【驱动】串口驱动分析(二)-tty core

前言 tty这个名称源于电传打字节的简称&#xff0c;在linux表示各种终端&#xff0c;终端通常都跟硬件相对应。比如对应于输入设备键盘鼠标&#xff0c;输出设备显示器的控制终端和串口终端。也有对应于不存在设备的pty驱动。在如此众多的终端模型之中&#xff0c;linux是怎么…
阅读更多...
信贷专员简历模板

信贷专员简历模板

这份简历内容&#xff0c;以信贷专员招聘需求为背景&#xff0c;我们制作了1份全面、专业且具有参考价值的简历案例&#xff0c;大家可以灵活借鉴。 信贷专员简历在线编辑下载&#xff1a;百度幻主简历 求职意向 求职类型&#xff1a;全职 意向岗位&#xff1a;信贷专员 …
阅读更多...
Linux：docker的数据管理（6）

Linux：docker的数据管理（6）

数据管理操作*方便查看容器内产生的数据 *多容器间实现数据共享 两种管理方式数据卷 数据卷容器 1.数据卷 数据卷是一个供容器使用的特殊目录&#xff0c;位于容器中&#xff0c;可将宿主机的目录挂载到数据卷上&#xff0c;对数据卷的修改操作立刻可见&#xff0c;并且更新数…
阅读更多...
基于SpringBoot的在线视频教育平台的设计与实现

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

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

Django二转Day03 04

0 cbv执行流程&#xff0c;self问题 path(index/, Myview.as_view()),Myview.as_view() 实例化后返回 变成return Myview.dispatch(request, *args, **kwargs)但是视图函数Myview中没有 dispatch 方法 所以去 父类View中寻找return View.dispatch(request, *args, **kwargs)调用…
阅读更多...
动手学深度学习(六)---权重衰退

动手学深度学习(六)---权重衰退

文章目录 一、理论知识二、代码实现【相关总结】 主要解决过拟合 一、理论知识 1、使用均方范数作为硬性限制&#xff08;不常用&#xff09; 通过限制参数值的选择范围来控制模型容量 通常不限制偏移b 小的意味着更强的正则项 使用均方范数作为柔性限制 对于每个都可以找到使…
阅读更多...
Zabbix“专家坐诊”第213期问答汇总

Zabbix“专家坐诊”第213期问答汇总

问题一 Q&#xff1a;Zabbix报错&#xff1a;Zabbix server is not running :the information displayed may not be current&#xff0c;是什么问题呢&#xff1f; A&#xff1a; 1、数据库软件问题导致导入的zabbix数据库不完整2、zabbix Server配置问题3、zabbix-server没…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023