tex2D使用学习

已于 2023-11-30 23:55:47 修改
阅读量675 收藏 12
点赞数 7
分类专栏： Cuda 文章标签： 学习
于 2023-11-30 21:20:20 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/bocai1215/article/details/134722032
版权

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的方式进行排布就好了

用户名溢出
关注
  • 7
    点赞
  • 踩
  • 12
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
Shader学习笔记_函数探索_tex2D(sampler2D tex, float2 s)
weixin_33692284的博客
05-26 5525
half4 c = tex2D (_MainTex, IN.uv_MainTex); o.Albedo = c.rgb; o.Alpha = c.a; tex2D(sampler2D tex, float2 s)函数，这是CG程序中用来在一张贴图中对一个点进行采样的方法，返回一个float4。这里对 _MainTex在输入点上进行了采样，并将其颜色的rbg值赋予了输出的像素颜色，...
自学Shader 函数（tex2D）
Sevol_Y
09-09 6270
tex2D函数作用：通过一个二维uv坐标在纹理上，获取该处值 float4 Tex2D(sampler2D tex, float2 s) ：s-需进行查找的纹理坐标(uv) 其中根据纹理的类型不同，获取的值的含义也不一样，普通纹理一般代表的是该点的颜色值 //该像素的真实颜色值 half4 c = tex2D (_MainTex, IN.uv_MainTex); bump类型纹理上存储的值代表的含义是该点的法向量， //从_Bump纹理中提取法向信息 tex2D(_Bump, IN.uv_Bump)
Unity Shader：tex2D函数详解
最新发布
weixin_42565127的博客
05-16 2038
UnityShader tex2D函数详解
Unity_Shader_纹理贴图
9624
11-26 1155
tex2D(sampler2D samp, float2 s) 2D纹理采样，CG内置函数。 内部实现分为以下几步： 1. 用图片的宽高度乘以uv数值，得到像素坐标。widthPixel=samp.x*s.x;heightPixel=samp.y*s.y; 2. 因为取到的数值基本上都是带有小数点的，也就是说不是一个整数，这个时候，需要看图片的过滤设置了。也就是Unity3d的图片设置中...
CG函数tex2D()的详细解释
热门推荐
yuanerjiang的博客
03-20 1万+
half4 c = tex2D (_MainTex, IN.uv_MainTex); o.Albedo = c.rgb; o.Alpha = c.a; tex2D(sampler2D tex, float2 s)函数，这是CG程序中用来在一张贴图中对一个点进行采样的方法，返回一个float4。这里对 _MainTex在输入点上进行了采样，并将其颜色的rbg值赋予了输出的像素颜色，将a值赋予透...
tex_cube.rar_tex
09-24
同时，代码中可能会有OpenGL函数，如`glGenTextures`创建纹理对象，`glBindTexture`绑定纹理，`glTexImage2D`加载纹理数据，以及`glTexParameteri`设置纹理参数等。 此外，为了实现交互式旋转，代码中可能包含矩阵...
使用FVF方法的Draw2D
09-18
在DirectX编程中，"Draw2D"通常指的是在二维空间中进行绘制操作的技术。《Programming Role Playing Games with DirectX》是一本经典的...学习和掌握这些知识，对于提升你在DirectX 2D图形编程方面的技能至关重要。
lp_tex_sample.rar_LP_tex
09-24
在3D场景中，为了给物体表面添加细节和真实感，通常会使用纹理映射技术，即将2D纹理映射到3D物体的表面。纹理采样就是实现这一过程的关键算法。 lp_tex_sample.c和lp_tex_sample.h可能包含了实现纹理采样的C语言...
tikz-pgfplots-projects：使用tikz和pgfplots的项目
02-04
通过研究这个项目，用户可以学习到如何使用TikZ和pgfplots进行图形设计，提升LaTeX排版的水平，特别是在科研报告、论文或教学材料中创建专业图表的能力。同时，这些示例也可以作为模板，帮助用户快速生成自己的图形...
shader opengl FBO tex
12-19
在这个项目中，开发者可能学习了如何创建和使用OpenGL Shaders，以及如何结合FBO和纹理来实现特定的图形效果。源代码可能涵盖了设置和管理shaders、创建和绑定FBO、将FBO的结果存储到纹理，以及在后续的渲染阶段使用...
2D纹理采样
Young的博客
04-24 3385
采样函数 ：tex2D函数。   对一张贴图的采样范围 ：(0~1) 位图 ： CG中，每有一个sampler2D(纹理)，就会有一个uv坐标 位图的模式 ：       Repeat ：在左侧再拼接一个一样的图片       clamp ：不重复 过滤模式 ：        Billinear ：会有两个颜色的混合
CUDA C学习之纹理内存的使用
qq_40801709的博客
01-05 525
图片的加权叠加：有无使用纹理内存的GPU下运行和CPU下运行的性能比较。 参考了： https://blog.csdn.net/langzai310/article/details/83573278 https://blog.csdn.net/shuzfan/article/details/77095270 https://learnopengl-cn.readthedocs.io/zh/...
CUDA学习（十）使用texture 进行图像处理
zhouzhouasishuijiao的博客
12-02 5612
首先texture在硬件层次上是属于一块存储器，通过绑定操作，可以用纹理操作根据坐标去获取数据，然后进行相应的操作 但是我现在对于texture的理解是，它可以较快的读取数据，如果你需要进行滤波，使用texture是最合适的，但是如果你不需要滤波，我没想到它的使用与global memory相比有什么好处，好想找人讨论啊，感觉自己理解不够 代码层次上对于texture存储器的使用，主要有以下几...
CUDA纹理内存入门详解
hgkjbm的博客
03-31 3108
CUDA 进阶技巧 纹理内存 入门指南
HLSL固有函数
熊小屁的专栏
02-17 1396
Intrinsic Functions (DirectX HLSL) The following table lists the intrinsic functions available in HLSL. Each function has a brief description, and a link to a reference page that has more detail abou
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值