0

我正在尝试将 2D 数组绑定到纹理并在数据之间进行插值。我的问题是。当我将数组绑定到纹理时，我访问的值完全是胡说八道。即使我试图访问第一个值 (text2D(tex,0.0f,0.0f) 我没有意义。所以我想我绑定错了或者我的 memcopy 错了。任何想法我的错误是?

这是代码

#include <stdio.h>
#include <iostream>
#include "cuda.h"
#include <stdlib.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "HelloWorld.h"

#include "linearInterpolation_kernel.cu"
#include "linearInterpolation_kernel2.cu" 
#include "linearInterpolation_kernel3.cu" 


using namespace std;
using std::cout;

const int blocksize = 16; 

__global__ 
void hello(char *a, int *b) {
    a[threadIdx.x] += b[threadIdx.x];
}



////////////////////////////////////////////////////////////////////////////////
// These are CUDA Helper functions

    // This will output the proper CUDA error strings in the event that a CUDA host call returns an error
    #define checkCudaErrors(err)           __checkCudaErrors (err, __FILE__, __LINE__)

    inline void __checkCudaErrors( cudaError err, const char *file, const int line )
    {
        if( cudaSuccess != err) {
            printf("%s(%i) : CUDA Runtime API error %d: %s.\n",file, line, (int)err, cudaGetErrorString( err ) );

        }
    }

    // This will output the proper error string when calling cudaGetLastError
    #define getLastCudaError(msg)      __getLastCudaError (msg, __FILE__, __LINE__)

    inline void __getLastCudaError( const char *errorMessage, const char *file, const int line )
    {
        cudaError_t err = cudaGetLastError();
        if( cudaSuccess != err) {
            printf("%s(%i) : getLastCudaError() CUDA error : %s : (%d) %s.\n", file, line, errorMessage, (int)err, cudaGetErrorString( err ) );

        }
    }

int main()
{
    int N = 200; 
    float *A; 
    A = (float *) malloc(N*sizeof(float));
    float *B;
    B = (float *) malloc(N*sizeof(float));
    float *result;
    result = (float *) malloc(N*sizeof(float));
    float angle = 0.5f; 

    for(int i = 0; i < N; i++){
        A[i] = (float)rand();
        B[i] = (float)rand();
    }
    cout << A[3] << endl;
    cout << B[3] << endl;

    ipLinearTexture(A,B,result,angle,N);

    float result2;

    result2 = (angle)*A[3] + (1-angle)*B[3]; 

    printf(" A %f B %f Result %f\n", A[3], B[3], result[3]);
    cout << result2 << endl;

    return 1;
}


void ipLinearTexture(float *A, float* B, float* result, float angle, int N)
{
    float cuTime;

    const int N2 = N;
    float *dev_result;
    float **AB; 

    AB = (float **) malloc( N * sizeof(float *));

    if(AB)
    {
        for(int i = 0; i < N; i++)
        {
            AB[i] = (float *) calloc( 2 , sizeof(float *));
        }
    }
    for (int i = 0; i < N; i++)
    {
        AB[i][0] = A[i]; 
        AB[i][1] = B[i];
    }


    cudaMalloc(&dev_result, N * sizeof(float));
    unsigned int size = N * 2 * sizeof(float);

    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
    cudaArray* cu_array;

    checkCudaErrors(cudaMallocArray( &cu_array, &channelDesc,N,2 )); 
    checkCudaErrors(cudaMemcpyToArray( cu_array, 0, 0, AB, size, cudaMemcpyHostToDevice));

    tex.addressMode[0] = cudaAddressModeClamp;
    tex.addressMode[1] = cudaAddressModeClamp;
    tex.filterMode = cudaFilterModeLinear; 
    tex.normalized = false;    // access with normalized texture coordinates

    checkCudaErrors(cudaBindTextureToArray( tex, cu_array, channelDesc));

    dim3 dimBlock(10, 1, 1);
    dim3 dimGrid((int)ceil((double)N*2/dimBlock.x), 1, 1);

    transformKernel3<<< dimGrid, dimBlock, 0 >>>( dev_result, N, 2, angle);

    checkCudaErrors(cudaUnbindTexture(tex));

    cudaMemcpy(result, dev_result, N * sizeof(float), cudaMemcpyKind::cudaMemcpyDeviceToHost);
    result[0] = (float)cuTime;
    cout << "==================================================" << endl;

    for (int i = 0 ; i < N ;i++)
    {
        cout << result[i] << endl;      
    }

    cout << "==================================================" << endl;

    cudaFree(dev_result);
    cudaFreeArray(cu_array);
}

这是内核中的代码

#ifndef _SIMPLETEXTURE_KERNEL3_H_
#define _SIMPLETEXTURE_KERNEL3_H_

// declare texture reference for 2D float texture
texture<float, 1> tex;

////////////////////////////////////////////////////////////////////////////////
//! Transform an image using texture lookups
//! @param g_odata  output data in global memory
////////////////////////////////////////////////////////////////////////////////
__global__ void
transformKernel3( float* g_odata, int width, int height, float theta) 
{
    unsigned int id = blockIdx.x*blockDim.x + threadIdx.x;

    if (id < width*height)
    {       

        g_odata[id] = tex1D(tex, xid * 2 + 0.5f);

    }
}

#endif // #ifndef _SIMPLETEXTURE_KERNEL_H_
4

1 回答 1

1

就像 OpenGL 中的概念一样，您可以认为 2D 纹理是一个矩形场。每个小矩形的中心点是您的数组数据。因此，tex2D(tex, 0.5f/width, 0.5f/height)这正是您的数组数据的第一个值。( width&height是二维数组数据的宽度和高度)

于 2012-08-29T12:24:57.343 回答