关于CUDA-纹理内存的使用

简要介绍一下

还是直接上代码：---有个要点注意一下，由于init_data操作的是设备内存，所以必须在设备代码中操作，为了简单起见所以直接使用了<<<1,1>>>直接初始化，还是停留在串行的程序设计中。然后调用1个块1024条线程进行并行进入设备代码，输出。主要是记住，使用纹理内存必须的5个步骤------分配内存-绑定内存-核函数调用-接触绑定-释放内存----一共五步-（其实不一定这五步，但是好记啊）

#include <stdio.h>
#include <cuda_runtime.h>
#define N 1024
texture<float, 1, cudaReadModeElementType> tex;
__global__ void kernel() {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    float x = tex1Dfetch(tex, i);
    printf("%f \n",x);
}

__global__ void init_data(float * a,int n){
    for(int i=0;i<n;i++){
        a[i] = i+3.1415926;
    }
}

int main() {
    float *buffer;
    dim3 grid(1,1,1);
    dim3 block(N,1,1);
    cudaMalloc(&buffer, N*sizeof(float));
    init_data<<<1,1>>>(buffer,N);
    cudaBindTexture(0, tex, buffer, N*sizeof(float));
    kernel <<<grid, block >>>();
    cudaUnbindTexture(tex);
    cudaFree(buffer);
}

结果：

补充：

二维使用的例子：

原理：

#define W 32
#define H 32
#include <cuda_runtime.h>
#include <stdio.h>

texture<float, 2, cudaReadModeElementType> tex;

__global__ void kernel() {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    float v = tex2D (tex, x, y);
    printf("%d , %d , %f\n",x,y,v);
}

__global__ void init_data(float * a,int n){
    for(int i=0;i<n;i++){
        a[i] = i+3.1415926;
        //printf("%f \n",a[i]);
    }
}


int main() {
    float *buffer;
    dim3 grid(1,1,1);
    dim3 block(W,H,1);

    cudaMalloc(&buffer, W*H*sizeof(float));
    init_data<<<1,1>>>(buffer,W*H);
    
    cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
    cudaBindTexture2D(0, tex, buffer,desc, W, H,W*sizeof(float));   
    kernel <<<grid, block >>>();
    cudaUnbindTexture(tex);
    cudaFree(buffer);
}