CUDA 学习——纹理内存（CUDA实战 第七章）

最新推荐文章于 2024-03-01 23:37:58 发布
最新推荐文章于 2024-03-01 23:37:58 发布
阅读量2.2k 收藏 11
点赞数 2
分类专栏： CUDA实战 文章标签： 学习 c++ 开发语言
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/ygf666/article/details/124851409
版权

1. 纹理内存使用

纹理内存的使用依赖于API函数。下面直接给出常见的使用流程：

1.1 声明纹理变量

texture<Type, Dim, ReadMode> VarName;
//Type: 前面提到的基本的整型和浮点类型，以及其它的对齐类型
//Dim： 纹理数组的维度，值为1或2或3，默认缺省为1
//ReadMode：cudaReadModelNormalizedFloat 或 cudaReadModelElementType(默认)

NOTE：
cudaReadModelNormalizedFloat：如果Type为整型(8bit或者16bit)，则读取数据时会自动将整型数据转化为浮点数。具体地，如果是无符号整型，则转化为[0 1]之间的浮点数；如果是有符号整型，则转化为[-1 1]之间的浮点数。
cudaReadModelElementType：默认值，不进行任何转换

1.2 创建内存

分配内存，内存形式有两种 线性内存 和 CUDA数组。
线性内存通过cudaMalloc()、cudaMallocPitch()或者cudaMalloc3D()分配；
CUDA数组可以通过cudaMalloc3DArray()或者cudaMallocArray()分配。前者可以分配1D、2D、3D的数组，后者一般用于分配2D的CUDA数组。

比如开辟一个二维CUDA数组ArrayName(64 x 64)：

//ex:
cudaMalloc( (void**)&data.output_bitmap,imageSize );
cudaMalloc( (void**)&data.dev_inSrc,imageSize );
cudaMalloc( (void**)&data.dev_outSrc,imageSize );
cudaMalloc( (void**)&data.dev_constSrc,imageSize );

1.3 绑定纹理内存

纹理绑定(texture binding)的作用有两个：

  1. 将指定的缓冲区作为纹理来使用；
  2. 纹理引用作为纹理的“名字”

通常使用cudaBindTexture() 或 cudaBindTexture2D() 分别将线性内存绑定到1D和2D纹理内存，使用cudaBindTextureToArray()将CUDA数组与纹理绑定。

//ex:将之前的texConstSrc，texIn，texOut与内存绑定
cudaBindTexture( NULL, texConstSrc,
                       data.dev_constSrc,
                       imageSize )

cudaBindTexture( NULL, texConstSrc,
                       data.dev_inSrc,
                       mageSize )
cudaBindTexture( NULL, texConstSrc,
                       data.dev_outSrc,
                       imageSize )

此时，纹理变量设置完毕。可以启用核函数。当读取核函数中的纹理时，需要通过特殊函数告诉GPU将读取请求转发给纹理内存而非标准的全局内存。因此，读取内存时不再使用方括号从缓冲区读取，而是将blend_kernel()改为tex1Dfetch()

1.4 读取

普通变量读取方式：

__global__ void blend_kernel(float *outSrc, const float *inSrc){
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;
if (x == 0)
	left++;
int right = offset + 1;
if (x == DIM -1)
	right--;
	
int top = offset -DIM;
if (y == 0)
	top += DIM;
int bottom = offset + DIM;
if (y == DIM -1)
	bottom -= DIM;
outSrc[offset] = inSrc[offset] + 
				SPEED * (inSrc[top] + inSrc[bottom] + inSrc[left] + inSrc[right]
				 - 4 * inSrc[offset] );
	
}

纹理变量的读取方式：

__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;
if (x == 0)
	left++;
int right = offset + 1;
if (x == DIM -1)
	right--;
	
int top = offset -DIM;
if (y == 0)
	top += DIM;
int bottom = offset + DIM;
if (y == DIM -1)
	bottom -= DIM;
float t, l, c, r, b;
if (destOut){
	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);
}

1.5 解绑

cudaUnbindTexture (tex);
cudaFree(devPtr);

具体关于纹波的操作可以参看此链接。

2. 书本例子：

2.1二维纹理：

#include <stdio.h>
#include <GL/glut.h>

#define DIM 1024
#define PI 3.1415926535897932f
#define MAX_TEMP 1.0f
#define MIN_TEMP 0.0001f
#define SPEED   0.25f

// a place for common kernels - starts here

__device__ unsigned char value( float n1, float n2, int hue ) {
    if (hue > 360)      hue -= 360;
    else if (hue < 0)   hue += 360;

    if (hue < 60)
        return (unsigned char)(255 * (n1 + (n2-n1)*hue/60));
    if (hue < 180)
        return (unsigned char)(255 * n2);
    if (hue < 240)
        return (unsigned char)(255 * (n1 + (n2-n1)*(240-hue)/60));
    return (unsigned char)(255 * n1);
}
__global__ void float_to_color( unsigned char *optr,
    const float *outSrc ) {
// map from threadIdx/BlockIdx to pixel position
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 l = outSrc[offset];
float s = 1;
int h = (180 + (int)(360.0f * outSrc[offset])) % 360;
float m1, m2;

if (l <= 0.5f)
m2 = l * (1 + s);
else
m2 = l + s - l * s;
m1 = 2 * l - m2;

optr[offset*4 + 0] = value( m1, m2, h+120 );
optr[offset*4 + 1] = value( m1, m2, h );
optr[offset*4 + 2] = value( m1, m2, h -120 );
optr[offset*4 + 3] = 255;
}

struct CPUAnimBitmap {
    unsigned char    *pixels;
    int     width, height;
    void    *dataBlock;
    void (*fAnim)(void*,int);
    void (*animExit)(void*);
    void (*clickDrag)(void*,int,int,int,int);
    int     dragStartX, dragStartY;

    CPUAnimBitmap( int w, int h, void *d = NULL ) {
        width = w;
        height = h;
        pixels = new unsigned char[width * height * 4];
        dataBlock = d;
        clickDrag = NULL;
    }

    ~CPUAnimBitmap() {
        delete [] pixels;
    }

    unsigned char* get_ptr( void ) const   { return pixels; }
    long image_size( void ) const { return width * height * 4; }

    void click_drag( void (*f)(void*,int,int,int,int)) {
        clickDrag = f;
    }

    void anim_and_exit( void (*f)(void*,int), void(*e)(void*) ) {
        CPUAnimBitmap**   bitmap = get_bitmap_ptr();
        *bitmap = this;
        fAnim = f;
        animExit = e;
        // a bug in the Windows GLUT implementation prevents us from
        // passing zero arguments to glutInit()
        int c=1;
        char* dummy = "";
        glutInit( &c, &dummy );
        glutInitDisplayMode( GLUT_DOUBLE | GLUT_RGBA );
        glutInitWindowSize( width, height );
        glutCreateWindow( "bitmap" );
        glutKeyboardFunc(Key);
        glutDisplayFunc(Draw);
        if (clickDrag != NULL)
            glutMouseFunc( mouse_func );
        glutIdleFunc( idle_func );
        glutMainLoop();
    }

    // static method used for glut callbacks
    static CPUAnimBitmap** get_bitmap_ptr( void ) {
        static CPUAnimBitmap*   gBitmap;
        return &gBitmap;
    }

    // static method used for glut callbacks
    static void mouse_func( int button, int state,
                            int mx, int my ) {
        if (button == GLUT_LEFT_BUTTON) {
            CPUAnimBitmap*   bitmap = *(get_bitmap_ptr());
            if (state == GLUT_DOWN) {
                bitmap->dragStartX = mx;
                bitmap->dragStartY = my;
            } else if (state == GLUT_UP) {
                bitmap->clickDrag( bitmap->dataBlock,
                                   bitmap->dragStartX,
                                   bitmap->dragStartY,
                                   mx, my );
            }
        }
    }

    // static method used for glut callbacks
    static void idle_func( void ) {
        static int ticks = 1;
        CPUAnimBitmap*   bitmap = *(get_bitmap_ptr());
        bitmap->fAnim( bitmap->dataBlock, ticks++ );
        glutPostRedisplay();
    }

    // static method used for glut callbacks
    static void Key(unsigned char key, int x, int y) {
        switch (key) {
            case 27:
                CPUAnimBitmap*   bitmap = *(get_bitmap_ptr());
                bitmap->animExit( bitmap->dataBlock );
                //delete bitmap;
                exit(0);
        }
    }

    // static method used for glut callbacks
    static void Draw( void ) {
        CPUAnimBitmap*   bitmap = *(get_bitmap_ptr());
        glClearColor( 0.0, 0.0, 0.0, 1.0 );
        glClear( GL_COLOR_BUFFER_BIT );
        glDrawPixels( bitmap->width, bitmap->height, GL_RGBA, GL_UNSIGNED_BYTE, bitmap->pixels );
        glutSwapBuffers();
    }
};

// these exist on the GPU side
texture<float,2>  texConstSrc;
texture<float,2>  texIn;
texture<float,2>  texOut;

__global__ void blend_kernel( float *dst,
                              bool dstOut ) {
    // map from threadIdx/BlockIdx to pixel position
    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   t, l, c, r, b;
    if (dstOut) {
        t = tex2D(texIn,x,y-1);
        l = tex2D(texIn,x-1,y);
        c = tex2D(texIn,x,y);
        r = tex2D(texIn,x+1,y);
        b = tex2D(texIn,x,y+1);
    } else {
        t = tex2D(texOut,x,y-1);
        l = tex2D(texOut,x-1,y);
        c = tex2D(texOut,x,y);
        r = tex2D(texOut,x+1,y);
        b = tex2D(texOut,x,y+1);
    }
    dst[offset] = c + SPEED * (t + b + r + l - 4 * c);
}

__global__ void copy_const_kernel( float *iptr ) {
    // map from threadIdx/BlockIdx to pixel position
    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 = tex2D(texConstSrc,x,y);
    if (c != 0)
        iptr[offset] = c;
}

// globals needed by the update routine
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 ) {
    cudaEventRecord( d->start, 0 );
    dim3    blocks(DIM/16,DIM/16);
    dim3    threads(16,16);
    CPUAnimBitmap  *bitmap = d->bitmap;

    // since tex is global and bound, we have to use a flag to
    // select which is in/out per iteration
    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<<<blocks,threads>>>( in );
        blend_kernel<<<blocks,threads>>>( out, dstOut );
        dstOut = !dstOut;
    }
    float_to_color<<<blocks,threads>>>( d->output_bitmap,
                                        d->dev_inSrc );

    cudaMemcpy( bitmap->get_ptr(),
                              d->output_bitmap,
                              bitmap->image_size(),
                              cudaMemcpyDeviceToHost ) ;
    cudaEventRecord( d->stop, 0 ) ;
    cudaEventSynchronize( d->stop ) ;
    float   elapsedTime;
    cudaEventElapsedTime( &elapsedTime,
                                        d->start, d->stop );
    d->totalTime += elapsedTime;
    ++d->frames;
    printf( "Average Time per frame:  %3.1f ms\n",
            d->totalTime/d->frames  );
}

// clean up memory allocated on the GPU
void anim_exit( DataBlock *d ) {
    cudaUnbindTexture( texIn );
    cudaUnbindTexture( texOut );
    cudaUnbindTexture( texConstSrc );
    cudaFree( d->dev_inSrc ) ;
    cudaFree( d->dev_outSrc  );
    cudaFree( d->dev_constSrc );

    cudaEventDestroy( d->start );
    cudaEventDestroy( d->stop );
}


int main( void ) {
    DataBlock   data;
    CPUAnimBitmap bitmap( DIM, DIM, &data );
    data.bitmap = &bitmap;
    data.totalTime = 0;
    data.frames = 0;
    cudaEventCreate( &data.start );
    cudaEventCreate( &data.stop );

    int imageSize = bitmap.image_size();

    cudaMalloc( (void**)&data.output_bitmap,
                               imageSize );

    // assume float == 4 chars in size (ie rgba)
    cudaMalloc( (void**)&data.dev_inSrc,
                              imageSize );
    cudaMalloc( (void**)&data.dev_outSrc,
                              imageSize );
    cudaMalloc( (void**)&data.dev_constSrc,
                              imageSize );

    cudaChannelFormatDesc desc = cudaCreateChannelDesc<float>();
    cudaBindTexture2D( NULL, texConstSrc,
                                   data.dev_constSrc,
                                   desc, DIM, DIM,
                                   sizeof(float) * DIM ) ;

    cudaBindTexture2D( NULL, texIn,
                                   data.dev_inSrc,
                                   desc, DIM, DIM,
                                   sizeof(float) * DIM ) ;

    cudaBindTexture2D( NULL, texOut,
                                   data.dev_outSrc,
                                   desc, DIM, DIM,
                                   sizeof(float) * DIM ) ;

    // initialize the constant data
    float *temp = (float*)malloc( imageSize );
    for (int i=0; i<DIM*DIM; i++) {
        temp[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;
        }
    }
    cudaMemcpy( data.dev_constSrc, temp,
                              imageSize,
                              cudaMemcpyHostToDevice ) ;    

    // initialize the input data
    for (int y=800; y<DIM; y++) {
        for (int x=0; x<200; x++) {
            temp[x+y*DIM] = MAX_TEMP;
        }
    }
    cudaMemcpy( data.dev_inSrc, temp,
                              imageSize,
                              cudaMemcpyHostToDevice );
    free( temp );

    bitmap.anim_and_exit( (void (*)(void*,int))anim_gpu,
                           (void (*)(void*))anim_exit );
}

编译，输出：
在这里插入图片描述

2.2 一维纹理：

#include <stdio.h>
#include <GL/glut.h>


#define DIM 1024
#define PI 3.1415926535897932f
#define MAX_TEMP 1.0f
#define MIN_TEMP 0.0001f
#define SPEED   0.25f

// these exist on the GPU side
texture<float>  texConstSrc;
texture<float>  texIn;
texture<float>  texOut;
__device__ unsigned char value( float n1, float n2, int hue ) {
    if (hue > 360)      hue -= 360;
    else if (hue < 0)   hue += 360;

    if (hue < 60)
        return (unsigned char)(255 * (n1 + (n2-n1)*hue/60));
    if (hue < 180)
        return (unsigned char)(255 * n2);
    if (hue < 240)
        return (unsigned char)(255 * (n1 + (n2-n1)*(240-hue)/60));
    return (unsigned char)(255 * n1);
}
__global__ void float_to_color( unsigned char *optr,
    const float *outSrc ) {
// map from threadIdx/BlockIdx to pixel position
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 l = outSrc[offset];
float s = 1;
int h = (180 + (int)(360.0f * outSrc[offset])) % 360;
float m1, m2;

if (l <= 0.5f)
m2 = l * (1 + s);
else
m2 = l + s - l * s;
m1 = 2 * l - m2;

optr[offset*4 + 0] = value( m1, m2, h+120 );
optr[offset*4 + 1] = value( m1, m2, h );
optr[offset*4 + 2] = value( m1, m2, h -120 );
optr[offset*4 + 3] = 255;
}

__global__ void float_to_color( uchar4 *optr,
    const float *outSrc ) {
// map from threadIdx/BlockIdx to pixel position
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 l = outSrc[offset];
float s = 1;
int h = (180 + (int)(360.0f * outSrc[offset])) % 360;
float m1, m2;

if (l <= 0.5f)
m2 = l * (1 + s);
else
m2 = l + s - l * s;
m1 = 2 * l - m2;

optr[offset].x = value( m1, m2, h+120 );
optr[offset].y = value( m1, m2, h );
optr[offset].z = value( m1, m2, h -120 );
optr[offset].w = 255;
}

struct CPUAnimBitmap {
    unsigned char    *pixels;
    int     width, height;
    void    *dataBlock;
    void (*fAnim)(void*,int);
    void (*animExit)(void*);
    void (*clickDrag)(void*,int,int,int,int);
    int     dragStartX, dragStartY;

    CPUAnimBitmap( int w, int h, void *d = NULL ) {
        width = w;
        height = h;
        pixels = new unsigned char[width * height * 4];
        dataBlock = d;
        clickDrag = NULL;
    }

    ~CPUAnimBitmap() {
        delete [] pixels;
    }

    unsigned char* get_ptr( void ) const   { return pixels; }
    long image_size( void ) const { return width * height * 4; }

    void click_drag( void (*f)(void*,int,int,int,int)) {
        clickDrag = f;
    }

    void anim_and_exit( void (*f)(void*,int), void(*e)(void*) ) {
        CPUAnimBitmap**   bitmap = get_bitmap_ptr();
        *bitmap = this;
        fAnim = f;
        animExit = e;
        // a bug in the Windows GLUT implementation prevents us from
        // passing zero arguments to glutInit()
        int c=1;
        char* dummy = "";
        glutInit( &c, &dummy );
        glutInitDisplayMode( GLUT_DOUBLE | GLUT_RGBA );
        glutInitWindowSize( width, height );
        glutCreateWindow( "bitmap" );
        glutKeyboardFunc(Key);
        glutDisplayFunc(Draw);
        if (clickDrag != NULL)
            glutMouseFunc( mouse_func );
        glutIdleFunc( idle_func );
        glutMainLoop();
    }

    // static method used for glut callbacks
    static CPUAnimBitmap** get_bitmap_ptr( void ) {
        static CPUAnimBitmap*   gBitmap;
        return &gBitmap;
    }

    // static method used for glut callbacks
    static void mouse_func( int button, int state,
                            int mx, int my ) {
        if (button == GLUT_LEFT_BUTTON) {
            CPUAnimBitmap*   bitmap = *(get_bitmap_ptr());
            if (state == GLUT_DOWN) {
                bitmap->dragStartX = mx;
                bitmap->dragStartY = my;
            } else if (state == GLUT_UP) {
                bitmap->clickDrag( bitmap->dataBlock,
                                   bitmap->dragStartX,
                                   bitmap->dragStartY,
                                   mx, my );
            }
        }
    }

    // static method used for glut callbacks
    static void idle_func( void ) {
        static int ticks = 1;
        CPUAnimBitmap*   bitmap = *(get_bitmap_ptr());
        bitmap->fAnim( bitmap->dataBlock, ticks++ );
        glutPostRedisplay();
    }

    // static method used for glut callbacks
    static void Key(unsigned char key, int x, int y) {
        switch (key) {
            case 27:
                CPUAnimBitmap*   bitmap = *(get_bitmap_ptr());
                bitmap->animExit( bitmap->dataBlock );
                //delete bitmap;
                exit(0);
        }
    }

    // static method used for glut callbacks
    static void Draw( void ) {
        CPUAnimBitmap*   bitmap = *(get_bitmap_ptr());
        glClearColor( 0.0, 0.0, 0.0, 1.0 );
        glClear( GL_COLOR_BUFFER_BIT );
        glDrawPixels( bitmap->width, bitmap->height, GL_RGBA, GL_UNSIGNED_BYTE, bitmap->pixels );
        glutSwapBuffers();
    }
};


// this kernel takes in a 2-d array of floats
// it updates the value-of-interest by a scaled value based
// on itself and its nearest neighbors
__global__ void blend_kernel( float *dst,
                              bool dstOut ) {
    // map from threadIdx/BlockIdx to pixel position
    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);
}

// NOTE - texOffsetConstSrc could either be passed as a
// parameter to this function, or passed in __constant__ memory
// if we declared it as a global above, it would be
// a parameter here: 
// __global__ void copy_const_kernel( float *iptr,
//                                    size_t texOffset )
__global__ void copy_const_kernel( float *iptr ) {
    // map from threadIdx/BlockIdx to pixel position
    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;
}

// globals needed by the update routine
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 ) {
    cudaEventRecord( d->start, 0 );
    dim3    blocks(DIM/16,DIM/16);
    dim3    threads(16,16);
    CPUAnimBitmap  *bitmap = d->bitmap;

    // since tex is global and bound, we have to use a flag to
    // select which is in/out per iteration
    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<<<blocks,threads>>>( in );
        blend_kernel<<<blocks,threads>>>( out, dstOut );
        dstOut = !dstOut;
    }
    float_to_color<<<blocks,threads>>>( d->output_bitmap,
                                        d->dev_inSrc );

    cudaMemcpy( bitmap->get_ptr(),
                              d->output_bitmap,
                              bitmap->image_size(),
                              cudaMemcpyDeviceToHost );

    cudaEventRecord( d->stop, 0 );
    cudaEventSynchronize( d->stop );
    float   elapsedTime;
    cudaEventElapsedTime( &elapsedTime,
                                        d->start, d->stop );
    d->totalTime += elapsedTime;
    ++d->frames;
    printf( "Average Time per frame:  %3.1f ms\n",
            d->totalTime/d->frames  );
}

// clean up memory allocated on the GPU
void anim_exit( DataBlock *d ) {
    cudaUnbindTexture( texIn );
    cudaUnbindTexture( texOut );
    cudaUnbindTexture( texConstSrc );
    cudaFree( d->dev_inSrc ) ;
    cudaFree( d->dev_outSrc ) ;
    cudaFree( d->dev_constSrc );

    cudaEventDestroy( d->start );
    cudaEventDestroy( d->stop );
}


int main( void ) {
    DataBlock   data;
    CPUAnimBitmap bitmap( DIM, DIM, &data );
    data.bitmap = &bitmap;
    data.totalTime = 0;
    data.frames = 0;
    cudaEventCreate( &data.start ) ;
    cudaEventCreate( &data.stop );

    int imageSize = bitmap.image_size();

    cudaMalloc( (void**)&data.output_bitmap,
                               imageSize ) ;

    // assume float == 4 chars in size (ie rgba)
    cudaMalloc( (void**)&data.dev_inSrc,
                              imageSize ) ;
    cudaMalloc( (void**)&data.dev_outSrc,
                              imageSize ) ;
    cudaMalloc( (void**)&data.dev_constSrc,
                              imageSize ) ;

    cudaBindTexture( NULL, texConstSrc,
                                   data.dev_constSrc,
                                   imageSize ) ;

    cudaBindTexture( NULL, texIn,
                                   data.dev_inSrc,
                                   imageSize );

    cudaBindTexture( NULL, texOut,
                                   data.dev_outSrc,
                                   imageSize );

    // intialize the constant data
    float *temp = (float*)malloc( imageSize );
    for (int i=0; i<DIM*DIM; i++) {
        temp[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;
        }
    }
    cudaMemcpy( data.dev_constSrc, temp,
                              imageSize,
                              cudaMemcpyHostToDevice );    

    // initialize the input data
    for (int y=800; y<DIM; y++) {
        for (int x=0; x<200; x++) {
            temp[x+y*DIM] = MAX_TEMP;
        }
    }
    cudaMemcpy( data.dev_inSrc, temp,
                              imageSize,
                              cudaMemcpyHostToDevice );
    free( temp );

    bitmap.anim_and_exit( (void (*)(void*,int))anim_gpu,
                           (void (*)(void*))anim_exit );
}

编译，输出：
在这里插入图片描述

风尘23187
关注
  • 2
    点赞
  • 踩
  • 11
    收藏
    觉得还不错? 一键收藏
  • 打赏
    打赏
  • 0
    评论
专栏目录
CUDA——性能优化（一）
01-20
CUDA全局内存的合并访问（个人理解） 每个warp去访问全局内存，会有400-600个时钟周期的内存延迟，这个代价很昂贵，所以为了减少访问全局内存的指令次数，我们将满足字节大小和对齐要求的warp合并起来访问全局内存，从而减少对全局内存的访问次数，提高GPU性能。 关于warp指令基础知识 1）什么是warp? 一个线程warp包括32条线程（我的电脑是1个warp包括32条线程）。它位于多处理器中。 2)warp指令 发射warp的一个指令，即该warp的32条线程一起执行的该条指令。多处理器会花费 该条指令 个时钟周期。 3）控制流指令 任何流控制指令（ if , switch , do
CUDA内存(五) 纹理内存
Full_Speed_Turbo
02-18 856
文章目录纹理内存的特性API最简单的用法：例子分析：Simpletexture1. 声明CUDA数组之前，必须先用结构体channelDesc描述CUDA数组中的数据类型。2. 声明纹理参照系3.设置运行时纹理参照系属性4. 纹理绑定5 纹理拾取 纹理内存的特性 纹理内存缓存在芯片上，因此在某些情况中，它能够减少对内存的请求并提供更高效的内存带宽。纹理缓存是专门为那些在内存访问模式中存在大量空间局...
CUDA：实现使用纹理内存(附完整源码)
最新发布
希望我的博客，能帮上你解决学习中工作中所遇到的问题
03-01 47
CUDA：实现使用纹理内存(附完整源码)
cuda立体匹配——中值滤波
06-20
立体匹配是立体视觉从图像生成三维点云的常规手段。立体匹配算法主要是通过建立一个能量代价函数，通过此能量代价函数最小化来估计像素点视差值。立体匹配算法的实质就是一个最优化求解问题，通过建立合理的能量函数，增加一些约束，采用最优化理论的方法进行方程求解，这也是所有的病态问题求解方法。 双目立体匹配可划分为四个步骤：匹配代价计算、代价聚合、视差计算和视差优化。 匹配代价计算的目的是衡量待匹配像素与候选像素之间的相关性。两个像素无论是否为同名点，都可以通过匹配代价函数计算匹配代价，代价越小则说明相关性越大，是同名点的概率也越大。匹配代价计算的方法有很多，本课设使用灰度绝对值差（AD，Absolute Differences）。 代价聚合的根本目的是让代价值能够准确的反映像素之间的相关性。上一步匹配代价的计算往往只会考虑局部信息，通过两个像素邻域内一定大小的窗口内的像素信息来计算代价值，这很容易受到影像噪声的影响，而且当影像处于弱纹理或重复纹理区域，这个代价值极有可能无法准确的反映像素之间的相关性，直接表现就是真实同名点的代价值非最小。 视差计算即通过代价聚合之后的代价矩阵S来确定每个像素的最优
cuda.rar_cuda_cuda学习_cuda编程
09-24
CUDA学习入门教程详细讲解CUDA编程方法，适合初学者学习。
CUDA使用纹理内存
wozhengtao的博客
06-03 1904
纹理内存位于设备端，global memory也位于设备端，但是texture memory的访问速度较global memory要快。 因为纹理内存有cache, 只有当cache没有命中的时候才会去访问device memory，否则访问texture cache具有很小的延迟。 另外，texture cache的2D定位已经进行了优化，对于同一线程束的线程访问位置临近的texture
CUDA总结：纹理内存
热门推荐
Kelvin_Yan的专栏
01-04 1万+
来源：《CUDA_Runtime_API》、《CUDA_C_Programming_Guide-V8.0》 纹理内存和表面内存（surface memory）实质上是全局内存的一个特殊形态，全局内存被绑定为纹理内存（表面内存），对其的读（写）操作将通过专门的texture cache（纹理缓存）进行，其实称为纹理缓存更加贴切。 纹理缓存的优势：纹理缓存具备硬件插值特性，可以实现最近邻插值和线性
CUDA By Example(六)——纹理内存
hhhcbw的博客，欢迎各位来访
11-17 958
CUDA C中纹理内存使用的说明，并以热传导模拟为例
CUDA纹理内存入门详解
hgkjbm的博客
03-31 2830
CUDA 进阶技巧 纹理内存 入门指南
cudaBindTexture2D error
.
01-23 1067
cudaBindTexture2D 绑定失败。代码如下： cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc&lt;float&gt;(); cudaBindTexture2D(0, inputProcessed, tmpGray, channelDesc, 844, 480, 844*sizeof(float)); 解决方法：...
CUDA纹理内存tex1D/tex2D/tex3D函数
xiangxianghehe的博客
10-22 1274
如果纹理数据是聚合的，则可以使用tex1Dgather函数来替代tex1D函数。CUDA的tex1D是用于从一维纹理中读取数据的函数。tex2D函数可以用于从二维纹理中读取数据。如果纹理数据是二维的，则可以使用tex2D函数来替代tex1D函数。如果纹理数据是三维的，则可以使用tex3D函数来替代tex1D函数。使用tex2D函数取代tex1D函数可以提高性能，因为纹理数据是二维的，因此可以使用更少的纹理坐标来读取数据。tex1D函数可以用于从纹理中读取各种类型的数据，包括浮点数、整数和字符串。
cuda-handbook.rar_cuda_cuda学习
09-20
CUDA编程学习入门笔记，很不错。可以知道cuda是干什么用的
cuda-example.rar_cuda examples_cuda 例子_cuda学习
09-19
cuda examples 对于学习GPU应用的人有用，例子讲的很详细具体
CUDA：纹理内存的作用（Texture memory)
fb_help的专栏
07-22 2115
CUDA：纹理内存的作用（Texture memory) 1.容量很大的常量内存 纹理内存是DRAM上的内存，可以申请很大的空间，相比常量内存只能申请64kb来说，是一种很大空间的常量内存，而常量内存的好处是可以广播，当多个swap访问同一位置时，广播机制可以减少全局内存的访问，来提速。如图 2. 边界问题与插值问题 纹理内存有边界与插值的优化，有句话讲纹理内存永不越界就是因为纹理内存的边界优化。 ...
CUDA纹理内存--硬件插值功能的应用
shandianfengfan的博客
01-06 1749
1. 纹理内存的使用方式纹理内存是CUDA的一种只读内存，通常的使用方式有两种：(1)把数据从host端拷贝到device端的CUDA数据，然后将CUDA数组绑定到纹理内存，通过访问纹...
CUDA常见函数（二）：纹理参考管理
小P的博客
04-11 1252
主要是在图像处理时用到的一些函数，在这里汇总一下。
CUDA——纹理内存（纹理对象）
weixin_44444450的博客
02-10 1316
纹理内存（二） 1.2.2纹理对象API 使用纹理对象包括 1.纹理对象的创建 2.访问纹理内存 3.纹理对象销毁 纹理对象的创建 1）创建纹理对象，cudaCreateTextureObject(); cudaCreateTextureObject(cudaTextureObject_t *pTexObject, const struct cudaResourceDesc *pResDesc...
CUDA——纹理内存（纹理参考）
weixin_44444450的博客
02-09 1136
纹理内存 1.1纹理类型 CUDA支持GPU上的一部分纹理硬件（它们原本是为图形处理而设计的），所以，纹理内存一般会用于加速图片的滤波和图片的resize操作上。 纹理内存实质上市全局内存的特殊形态，全局内存被绑定为纹理内存，对其的读，写操作将通过专门的texture cache（纹理缓存）进行。 纹理缓存的优势：纹理缓存具备硬件插值特性，可以实现最近邻插值和线性插值。纹理缓存访问二维矩阵的领域会...
CUDA的纹理内存讲解
Dezeming的博客
03-28 7020
CUDA的纹理内存讲解，来自翻译和补充官方文档 CUDA programming Guide Texture and Surface Memory CUDA支持供GPU处理图形用的纹理硬件和表面（surface）内存，但是只支持一部分功能。从纹理和表面内存中读取数据可以得到很多好处（以后再说）。 有两个不同的API可以来获取纹理内存和表面内存：1.纹理参考API，支持所有设备。2.纹理对象...
CUDA学习笔记（四）——CUDA性能
05-19
CUDA是一个可以大大提高GPU计算性能的编程平台，但是为了实现最高的性能，需要遵循一些最佳实践和优化技巧。以下是一些常用的CUDA性能优化方法： 1. 矢量化：使用SIMD指令和数据结构（如向量和矩阵）来并行计算。 2. 内存访问模式：尽可能使用连续的内存地址来访问数据，减少不必要的内存访问，减少延迟和提高吞吐量。 3. 异步执行：使用异步内核调用和数据传输，使GPU和CPU可以同时执行任务，从而提高整体性能。 4. 分块：将大型数据集分成小块，以便在GPU内存中进行处理。 5. 共享内存：使用共享内存来存储临时数据，以减少全局内存访问和加速计算。 6. 流处理器：使用多个流来管理并发内核调用，以充分利用GPU的计算资源。 7. 局部性：尽可能利用空间和时间局部性，即将最近使用的数据存储在缓存中，以便快速访问。 8. 数据布局：使用合适的数据布局来最大化内存带宽和缓存效率。 这些方法并不是完整列表，但是它们可以帮助您更好地理解CUDA性能优化的基本原理和最佳实践。在实践中，您可能需要使用多个技术和优化技巧，以达到最佳性能。

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

风尘23187

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付：¥1
获取中
扫码支付

您的余额不足，请更换扫码支付或充值

打赏作者

实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

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

余额充值