cudaの1Dテクスチャにfloat *配列をバインドする方法は？

Question

テクスチャメモリを線形デバイス配列（cudaArrayではなく）にバインドして使用する方法を理解しようとしています。私のコードは単純です（下記）。私は1Dテクスチャにバインドしようとしている8個の数値の浮動小数点数の配列を持っています。次にカーネル関数でテクスチャを読み込み、その値を出力配列に入れようとします。私はこのテストを実行するときしかし、私の出力配列内のすべての値がゼロである：

入力= 0.000000 1.000000 2.000000 3.000000 4.000000 5.000000 6.000000 7.000000
出力= 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000

私は何をしないのですここに？

texture<float, 1, cudaReadModeElementType> texInput; 

__global__ void copyKernel(float*output, int n) { 
for (int i = 0; i < n; i++) { 
    output[i] = tex1D(texInput, (float)i); 
} 
} 

int main(int argc, char*argv[]) { 

const int WIDTH = 8; 

float* hInput = (float*)malloc(sizeof(float) * WIDTH); 
float*hOutput = (float*)malloc(sizeof(float) * WIDTH); 

for (int i = 0; i < WIDTH; i++) { 
    hInput[i] = (float)i; 
} 

float* dInput = NULL, *dOutput = NULL; 

size_t offset = 0; 

texInput.addressMode[0] = cudaAddressModeBorder; 
texInput.addressMode[1] = cudaAddressModeBorder; 
texInput.filterMode = cudaFilterModePoint; 
texInput.normalized = false; 

checkCudaErrors(cudaMalloc((void**)&dInput, sizeof(float)*WIDTH)); 
checkCudaErrors(cudaMalloc((void**)&dOutput, sizeof(float)*WIDTH)); 

cudaMemcpy(dInput, hInput, sizeof(float)*WIDTH, cudaMemcpyHostToDevice); 

cudaBindTexture(&offset, texInput, dInput, sizeof(float)*WIDTH); 


copyKernel<<<1,1>>>(dOutput, WIDTH); 

cudaMemcpy(hOutput, dOutput, sizeof(float)*WIDTH, cudaMemcpyDeviceToHost); 
printf("\nInput = "); 

for (int i = 0; i < WIDTH; i++) { 
     printf("%f\t",hInput[i]); 
    } 
printf("\nOutput = "); 
for (int i = 0; i < WIDTH; i++) { 
    printf("%f\t",hOutput[i]); 
} 

return 0; 
} 

Answer 1

下層の割り当てがCUDAアレイである場合the documentationによれば、tex1D()が使用されます。線形メモリ結合テクスチャの場合、正しいテクスチャ関数はtex1Dfetch()です。あなたのコードへの変更は（のみ）、それは私のために働く可能ということ

：

$ cat t1139.cu 
#include <stdio.h> 
#include <helper_cuda.h> 

texture<float, 1, cudaReadModeElementType> texInput; 

__global__ void copyKernel(float*output, int n) { 
for (int i = 0; i < n; i++) { 
    output[i] = tex1Dfetch(texInput, i); 
} 
} 

int main(int argc, char*argv[]) { 

const int WIDTH = 8; 

float* hInput = (float*)malloc(sizeof(float) * WIDTH); 
float*hOutput = (float*)malloc(sizeof(float) * WIDTH); 

for (int i = 0; i < WIDTH; i++) { 
    hInput[i] = (float)i; 
} 

float* dInput = NULL, *dOutput = NULL; 

size_t offset = 0; 

texInput.addressMode[0] = cudaAddressModeBorder; 
texInput.addressMode[1] = cudaAddressModeBorder; 
texInput.filterMode = cudaFilterModePoint; 
texInput.normalized = false; 

checkCudaErrors(cudaMalloc((void**)&dInput, sizeof(float)*WIDTH)); 
checkCudaErrors(cudaMalloc((void**)&dOutput, sizeof(float)*WIDTH)); 

cudaMemcpy(dInput, hInput, sizeof(float)*WIDTH, cudaMemcpyHostToDevice); 

cudaBindTexture(&offset, texInput, dInput, sizeof(float)*WIDTH); 


copyKernel<<<1,1>>>(dOutput, WIDTH); 

cudaMemcpy(hOutput, dOutput, sizeof(float)*WIDTH, cudaMemcpyDeviceToHost); 
printf("\nInput = "); 

for (int i = 0; i < WIDTH; i++) { 
     printf("%f\t",hInput[i]); 
    } 
printf("\nOutput = "); 
for (int i = 0; i < WIDTH; i++) { 
    printf("%f\t",hOutput[i]); 
} 

return 0; 
} 
$ nvcc -I/usr/local/cuda/samples/common/inc t1139.cu -o t1139 
$ cuda-memcheck ./t1139 
========= CUDA-MEMCHECK 

Input = 0.000000  1.000000  2.000000  3.000000  4.0000005.000000  6.000000  7.000000 
Output = 0.000000  1.000000  2.000000  3.000000  4.0000005.000000  6.000000  7.000000  ========= ERROR SUMMARY: 0 errors 
$