CUDA使用1维纹理内存，进行热传导模型计算

CUDA使用1维纹理内存，进行热传导模型计算
出现一大堆warning

#include "cuda.h"
#include "../common/cpu_anim.h"
#include"../common/book.h"
#include "../common/gpu_anim.h"
#include "../common/cpu_bitmap.h"
#define DIM 1024
#define PI 3.1415926535897932f
#define MAX_TEMP 1.0f
#define MIN_TEMP 0.0001f
#define SPEED   0.25f

texture<float>  texConstSrc;
texture<float>  texIn;
texture<float>  texOut;
__global__ void blend_kernel( float *dst,bool dstOut ) {
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    int left = offset - 1;
    int right = offset + 1;
    if (x == 0)   left++;
    if (x == DIM-1) right--; 

    int top = offset - DIM;
    int bottom = offset + DIM;
    if (y == 0)   top += DIM;
    if (y == DIM-1) bottom -= DIM;

    float   t, l, c, r, b;
    if (dstOut) {
        t = tex1Dfetch(texIn,top);
        l = tex1Dfetch(texIn,left);
        c = tex1Dfetch(texIn,offset);
        r = tex1Dfetch(texIn,right);
        b = tex1Dfetch(texIn,bottom);
    }else{
        t = tex1Dfetch(texOut,top);
        l = tex1Dfetch(texOut,left);
        c = tex1Dfetch(texOut,offset);
        r = tex1Dfetch(texOut,right);
        b = tex1Dfetch(texOut,bottom);
    }
    dst[offset] = c + SPEED * (t + b + r + l - 4 * c);
}

__global__ void copy_const_kernel( float *iptr ) {
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int offset = x + y * blockDim.x * gridDim.x;

    float c = tex1Dfetch(texConstSrc,offset);
    if (c != 0)
    iptr[offset] = c;
}


struct DataBlock {
    unsigned char   *output_bitmap;
    float           *dev_inSrc;
    float           *dev_outSrc;
    float           *dev_constSrc;
    CPUAnimBitmap  *bitmap;

    cudaEvent_t     start, stop;
    float           totalTime;
    float           frames;
};

void anim_gpu( DataBlock *d, int ticks ) {
    HANDLE_ERROR(cudaEventRecord( d->start, 0 ));
    dim3    blocks(DIM/16,DIM/16);
    dim3    threads(16,16);
    CPUAnimBitmap  *bitmap = d->bitmap;

    volatile bool dstOut = true;
    for (int i=0; i<90; i++) {
        float   *in, *out;
        if (dstOut) {
        in  = d->dev_inSrc;
        out = d->dev_outSrc;
        }else{
        out = d->dev_inSrc;
        in  = d->dev_outSrc;
        }
        copy_const_kernel<<>>( in );
        blend_kernel<<>>( out, dstOut );
        dstOut = !dstOut;
    }
    float_to_color<<>>( d->output_bitmap,d->dev_inSrc );

    HANDLE_ERROR(cudaMemcpy( bitmap->get_ptr(), d->output_bitmap,bitmap->image_size(),cudaMemcpyDeviceToHost));
    HANDLE_ERROR(cudaEventRecord( d->stop, 0 ));
    HANDLE_ERROR(cudaEventSynchronize( d->stop ));
    float   elapsedTime;
    HANDLE_ERROR(cudaEventElapsedTime( &elapsedTime,d->start, d->stop ));
    d->totalTime += elapsedTime;
    ++d->frames;
    printf( "Average Time per frame:  %3.1f ms\n",d->totalTime/d->frames  );
}
void anim_exit( DataBlock *d ) {
    cudaUnbindTexture( texIn );
    cudaUnbindTexture( texOut );
    cudaUnbindTexture( texConstSrc );
    HANDLE_ERROR(cudaFree( d->dev_inSrc ));
    HANDLE_ERROR(cudaFree( d->dev_outSrc ));
    HANDLE_ERROR(cudaFree( d->dev_constSrc ));
    HANDLE_ERROR(cudaEventDestroy( d->start ));
    HANDLE_ERROR(cudaEventDestroy( d->stop ));
}
int main( void ) {
    DataBlock   data;
    CPUAnimBitmap bitmap( DIM, DIM, &data );
    data.bitmap = &bitmap;
    data.totalTime = 0;
    data.frames = 0;
    HANDLE_ERROR(cudaEventCreate( &data.start ));
    HANDLE_ERROR(cudaEventCreate( &data.stop ));
    int imageSize = bitmap.image_size();
    HANDLE_ERROR(cudaMalloc( (void**)&data.output_bitmap,imageSize ));
    HANDLE_ERROR(cudaMalloc( (void**)&data.dev_inSrc,imageSize ) );
    HANDLE_ERROR(cudaMalloc( (void**)&data.dev_outSrc,imageSize ));
    HANDLE_ERROR(cudaMalloc( (void**)&data.dev_constSrc,imageSize ));
    HANDLE_ERROR(cudaBindTexture( NULL, texConstSrc,data.dev_constSrc,imageSize ));
    HANDLE_ERROR(cudaBindTexture( NULL, texIn,data.dev_inSrc,imageSize ));
    HANDLE_ERROR(cudaBindTexture( NULL, texOut,data.dev_outSrc,imageSize ));
    float *temp = (float*)malloc( imageSize );
    for (int i=0; itemp[i] = 0;
        int x = i % DIM;
        int y = i / DIM;
        if ((x>300) && (x<600) && (y>310) && (y<601))
        temp[i] = MAX_TEMP;
    }
    temp[DIM*100+100] = (MAX_TEMP + MIN_TEMP)/2;
    temp[DIM*700+100] = MIN_TEMP;
    temp[DIM*300+300] = MIN_TEMP;
    temp[DIM*200+700] = MIN_TEMP;
    for (int y=800; y<900; y++) {
        for (int x=400; x<500; x++) {
        temp[x+y*DIM] = MIN_TEMP;
        }
    }

    HANDLE_ERROR(cudaMemcpy( data.dev_constSrc, temp,imageSize,cudaMemcpyHostToDevice ));

    for (int y=800; yfor (int x=0; x<200; x++) {
        temp[x+y*DIM] = MAX_TEMP;
        }
    }
    HANDLE_ERROR(cudaMemcpy( data.dev_inSrc, temp,imageSize,cudaMemcpyHostToDevice ));
    free( temp );
    //anim_gpu(&data,0);
    bitmap.anim_and_exit((void(*)(void*,int))anim_gpu,(void(*)(void*))anim_exit);
}

报错的图片如下，怎么解决

远程服务器的话，无法显示热传导模型，所以把最后一句代码注释掉，用倒二行，就行了

这只是警告warning，只要不是error就不用管它。
这警告的都是你使用的旧版的API在你现在这个版本中已经废弃了，可以看下配置是否匹配。

这有个类似的问题, 你可以参考下: https://ask.csdn.net/questions/7396314
这篇博客你也可以参考下：CUDA 零拷贝内存
除此之外, 这篇博客: 5. CUDA编程手册中文版---性能指南中的 5.3 最大化存储吞吐量 部分也许能够解决你的问题, 你可以仔细阅读以下内容或跳转源博客中阅读:
最大化应用程序的整体内存吞吐量的第一步是最小化低带宽的数据传输。
这意味着最大限度地减少主机和设备之间的数据传输，如主机和设备之间的数据传输中所述，因为它们的带宽比全局内存和设备之间的数据传输低得多。
这也意味着通过最大化片上内存的使用来最小化全局内存和设备之间的数据传输：共享内存和缓存（即计算能力 2.x 及更高版本的设备上可用的 L1 缓存和 L2 缓存、纹理缓存和常量缓存适用于所有设备）。
共享内存相当于用户管理的缓存：应用程序显式分配和访问它。如 CUDA Runtime 所示，典型的编程模式是将来自设备内存的数据暂存到共享内存中；换句话说，拥有一个块的每个线程：
- 将数据从设备内存加载到共享内存，
- 与块的所有其他线程同步，以便每个线程可以安全地读取由不同线程填充的共享内存位置，
  处理共享内存中的数据，
- 如有必要，再次同步以确保共享内存已使用结果更新，
- 将结果写回设备内存。
对于某些应用程序（例如，全局内存访问模式依赖于数据），传统的硬件管理缓存更适合利用数据局部性。如 Compute Capability 3.x、Compute Capability 7.x 和 Compute Capability 8.x 中所述，对于计算能力 3.x、7.x 和 8.x 的设备，相同的片上存储器用于 L1 和共享内存，以及有多少专用于 L1 与共享内存，可针对每个内核调用进行配置。
内核访问内存的吞吐量可能会根据每种内存类型的访问模式而变化一个数量级。因此，最大化内存吞吐量的下一步是根据设备内存访问中描述的最佳内存访问模式尽可能优化地组织内存访问。这种优化对于全局内存访问尤为重要，因为与可用的片上带宽和算术指令吞吐量相比，全局内存带宽较低，因此非最佳全局内存访问通常会对性能产生很大影响。