使用纹理内存计算后保存值答案

【问题标题】：Saving Values after Calculating with Texture Memory使用纹理内存计算后保存值
【发布时间】：2012-12-05 01:18:36
【问题描述】：

您好，我有一个使用纹理内存的简单计算。但我无法保存正确的结果。结果应该是插值。例如 angle = 0.5 A[0] = 1, B[0] = 2, result[0] 应该是 1.5

我想我没有正确保存数据。我想使用纹理内存进行快速计算并将结果保存在全局数组中。我做错了什么。有人有什么想法吗？

这是我在内核中的代码

#ifndef _SIMPLETEXTURE_KERNEL_H_
#define _SIMPLETEXTURE_KERNEL_H_

texture<float, 1> tex1;
texture<float, 1> tex2;

__global__ void
transformKernel( float* g_odata, float f) 
{
    // calculate normalized texture coordinates
    unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    int idx = threadIdx.x;

    float valA = tex1D(tex1,x);
    float valB = tex1D(tex2,x);

    // read from texture and write to global memory
    g_odata[x] = (f)*valA + (1-f)*valB;
}
#endif

这里是我调用的代码

#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"

using namespace std;
using std::cout;

const int blocksize = 16; 
int main()
{
    int N = 1000; 
    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 cudaChannelFormatDesc channel_desc = cudaCreateChannelDesc<float>(); 
    cudaArray* cuda_A;
    cudaArray* cuda_B;
    float *dev_result;

    cudaMallocArray(&cuda_A, &channel_desc, 1, N * sizeof(float));
    cudaMallocArray(&cuda_B, &channel_desc, 1, N * sizeof(float));
    cudaMalloc(&dev_result, N * sizeof(float));

    cudaMemcpyToArray(cuda_A,0,0,A,N * sizeof(float),cudaMemcpyHostToDevice);
    cudaMemcpyToArray(cuda_B,0,0,B,N * sizeof(float),cudaMemcpyHostToDevice);

    tex1.filterMode = cudaFilterModePoint;
    tex1.addressMode[0] = cudaAddressModeWrap;

    tex2.filterMode = cudaFilterModePoint;
    tex2.addressMode[0] = cudaAddressModeWrap;


    cudaBindTextureToArray(tex1, cuda_A, channel_desc);
    cudaBindTextureToArray(tex2, cuda_B, channel_desc);

    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventRecord(start,0);

    transformKernel<<< 16, 16, 0 >>>(dev_result,angle);

    cudaEventCreate(&stop);
    cudaEventRecord(stop,0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&cuTime, start,stop);

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

    cudaFreeArray(cuda_A);
    cudaFreeArray(cuda_B);
    cudaFree(dev_result);

}

【问题讨论】：

您能否编辑您的问题以解释“错误结果”和“正确结果”是/应该是什么？
我没有看到您用来验证结果是否正确的代码。我发现您用测量的时间覆盖结果 [0] 很奇怪：result[0] = (float)cuTime;。老实说，对于这种线性寻址模式，使用纹理不会提供性能优势，而且代码比仅使用常规设备指针更复杂。

标签： cuda textures gpu

【解决方案1】：

从编程指南中的示例和文档in this page 看来，cudaMallocArray 函数的声明是：

cudaError_t cudaMallocArray(struct cudaArray **array, 
                const struct cudaChannelFormatDesc *desc, 
                size_t width, 
                size_t height = 0, 
                unsigned int flags = 0)

在您发布的代码中，您以字节为单位传递大小。仅尝试使用N，删除sizeof(float)

至少上次我使用纹理进行线性插值时，当我分配 cudaArray 时，我没有以字节为单位指定大小并且它可以工作。请注意，当您调用 cudaMemcpyToArray 时，大小应该以字节为单位。

【讨论】：