CUDAで1Dレイヤテクスチャを作成および使用する方法

Question

CUDAを初めて使用しています。私はCUDAで1Dと2Dのテクスチャをどのように行うかを考え出しました。しかし、私は1Dレイヤーテクスチャの使い方に苦労しています。テクスチャを使用する私のカーネルの出力はすべてゼロですが、間違いは間違いです。しかし、私は何が間違っているのか分かりません。私はこのテクスチャを正しく設定していることに深刻な疑念を抱いていますが、どこでもcudaエラーをチェックして問題を見つけることはできませんでした。誰かが1Dレイヤテクスチャを正しく設定して使用する方法を教えてもらえますか？ここに私のコードです。事前のおかげで：

// To Compile: nvcc backproj.cu -o backproj.out 
// To Run: ./backproj.out 

// Includes, system 
#include <stdlib.h> 
#include <stdio.h> 
#include <string.h> 
#include <math.h> 

// Includes CUDA 
#include <cuda_runtime.h> 
#include <cuda_profiler_api.h> 

#define pi acos(-1) 

// 1D float textures 
texture<float, cudaTextureType1DLayered, cudaReadModeElementType> texRef; 

// 1D interpolation kernel: Should be very similar to what you get if you used 1D interpolation on MATLAB 
__global__ void interp1Kernel(float* d_output, float* d_locations, int numlocations, int layer) { 
    unsigned int location_idx = blockIdx.x * blockDim.x + threadIdx.x; 
    if (location_idx < numlocations) { 
     // Get the location you want to interpolate from the array 
     float loc2find = (float) d_locations[location_idx] + 0.5f; 
     // Read from texture and write to global memory 
     d_output[location_idx] = tex1DLayered(texRef, loc2find, layer); 
    } 
} 

// Host code 
int main() 
{ 
    // Setup h_data and locations to interpolate from 
    const unsigned int len = 10; 
    const unsigned int numlayers = 3; 
    const unsigned int upsamp = 3; 
    const unsigned int loclen = 1 + (len - 1) * upsamp; 
    float idx_spacing = 1/(float)upsamp; 
    float h_data[len][numlayers], h_loc[loclen]; 
    for (int i = 0; i < len; i++) 
     for (int j = 0; j < numlayers; j++) 
      h_data[i][j] = 1+cosf((float) pi*i/(j+1.0f)); 
    for (int i = 0; i < loclen; i ++) 
     h_loc[i] = i*idx_spacing; 

    // Get the memory locations you want 
    float* d_loc; 
    cudaMalloc(&d_loc, loclen * sizeof(float)); 
    cudaMemcpy(d_loc, h_loc, loclen*sizeof(float), cudaMemcpyHostToDevice); 

    // Allocate CUDA array in device memory 
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat); 
    cudaArray* cuArray; 
    cudaMallocArray(&cuArray, &channelDesc, len, numlayers); 

    // Copy to device memory some data located at address h_data in host memory 
    cudaMemcpyToArray(cuArray, 0, 0, h_data, len * numlayers * sizeof(float), cudaMemcpyHostToDevice); 

    // Set texture reference parameters 
    texRef.addressMode[0] = cudaAddressModeBorder; 
    texRef.filterMode = cudaFilterModeLinear; 
    texRef.normalized = false; 

    // Bind the array to the texture reference 
    cudaBindTextureToArray(texRef, cuArray, channelDesc); 

    // Allocate result of transformation in device memory 
    float* d_output; 
    cudaMalloc(&d_output, loclen * sizeof(float)); 

    // Invoke kernel 
    int thdsPerBlk = 256; 
    int blksPerGrid = (int) (loclen/thdsPerBlk) + 1; 
    printf("Threads Per Block: %d, Blocks Per Grid: %d\n", thdsPerBlk, blksPerGrid); 
    interp1Kernel <<<blksPerGrid, thdsPerBlk >>>(d_output, d_loc, loclen, 0); 

    // Print Results 
    printf("\n Original Indices \n"); 
    for (int i = 0; i < len; i++) printf(" %d ", i); 
    printf("\n Original array \n"); 
    for (int i = 0; i < len; i++) printf("%5.3f ", h_data[i][0]); 
    printf("\n Output Indices \n"); 
    for (int i = 0; i < loclen; i++) printf("%5.3f ", h_loc[i]); 
    printf("\n Output Array \n"); 
    cudaMemcpy(h_loc, d_output, loclen * sizeof(float), cudaMemcpyDeviceToHost); 
    for (int i = 0; i < loclen; i++) printf("%5.3f ", h_loc[i]); 
    printf("\n"); 

    // Free device memory 
    cudaFreeArray(cuArray); 
    cudaFree(d_output); 

    return 0; 
} 

Answer 1

あなたが階層化テクスチャメモリを割り当てるように設定cudaArrayLayeredフラグを付けてcudaMalloc3DArrayを使用する必要があります。ツールキットのサンプルにはレイヤードテクスチャの使用方法があります。これらのテクスチャを使用して、どのように動作するかを調べることができます。

Answer 2

残念なことに、CUDA SDKは、2Dレイヤーテクスチャを持つときに行う方法を示しています。 1Dのレイヤードテクスチャについては、もう少しトリッキーです。それはあなたが次のようにextentDescを行う際make_cudaExtentのための第二引数に0を配置する必要が判明：cudaMemcpy3DためmParams.extentためmake_cudaExtentを使用した場合

cudaExtent extentDesc = make_cudaExtent(len, 0, numlayers); // <-- 0 height required for 1Dlayered 

はしかし、あなたはまだ第二引数に1を配置する必要があります：

mParams.extent = make_cudaExtent(len, 1, numlayers); // <<-- non zero height required for memcpy to do anything 

はまた、そのようなmake_cudaPitchedPtr用ピッチのようないくつかの他の非自明な詳細があります。だから私は、1Dのレイヤードテクスチャのための完全で機能的なコードを含んでいます。私はどこでもこの事例を見つけることができませんでした。だからうまくいけば、これは同じボートにいる他の人を助けるでしょう：

// To Compile: nvcc layeredTexture1D.cu -o layeredTexture1D.out 
// To Run: ./layeredTexture1D.out 

// Includes, system 
#include <stdlib.h> 
#include <stdio.h> 
#include <string.h> 
#include <math.h> 

// Includes CUDA 
#include <cuda_runtime.h> 
#include <cuda_profiler_api.h> 

#define pi acos(-1) 

// 1D float textures: x is for input values, y is for corresponding output values 
texture<float, cudaTextureType1DLayered, cudaReadModeElementType> texRef; 

// 1D interpolation kernel: Should be very similar to what you get if you used 1D interpolation on MATLAB 
__global__ void interp1Kernel(float* d_output, float* d_locations, int numlocations, int numlayers) { 
    unsigned int location_idx = blockIdx.x * blockDim.x + threadIdx.x; 
    unsigned int layer = blockIdx.y * blockDim.y + threadIdx.y; 
    if (location_idx < numlocations && layer < numlayers) { 
     // Get the location you want to interpolate from the array 
     float loc2find = (float)d_locations[location_idx] + 0.5f; 
     // Read from texture and write to global memory 
     d_output[location_idx + layer*numlocations] = tex1DLayered(texRef, loc2find, layer); 
     //printf("location=%d layer=%d loc2find=%f result=%f \n", location_idx, layer, loc2find, d_output[location_idx]); 
    } 
} 

// Host code 
int main() 
{ 
    // Setup h_data and locations to interpolate from 
    const unsigned int len = 7; 
    const unsigned int numlayers = 3; 
    const unsigned int upsamp = 4; 
    const unsigned int loclen = 1 + (len - 1) * upsamp; 
    float idx_spacing = 1/(float)upsamp; 
    float h_data[numlayers*len], h_loc[loclen]; 
    for (int i = 0; i < len; i++) 
     for (int j = 0; j < numlayers; j++) 
      h_data[len*j + i] = 1 + cosf((float)pi*i/(j + 1.0f)); 
    for (int i = 0; i < loclen; i++) 
     h_loc[i] = i*idx_spacing; 

    // Get the memory locations you want 
    float* d_loc; 
    cudaMalloc(&d_loc, loclen * sizeof(float)); 
    cudaMemcpy(d_loc, h_loc, loclen*sizeof(float), cudaMemcpyHostToDevice); 

    // Allocate CUDA array in device memory 
    cudaExtent extentDesc = make_cudaExtent(len, 0, numlayers); // <-- 0 height required for 1Dlayered 
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat); 
    cudaMemcpy3DParms mParams = { 0 }; 
    mParams.srcPtr = make_cudaPitchedPtr(h_data, len*sizeof(float), len, 1); 
    mParams.kind = cudaMemcpyHostToDevice; 
    mParams.extent = make_cudaExtent(len, 1, numlayers); // <<-- non zero height required for memcpy to do anything 
    cudaArray* cuArray; 
    cudaMalloc3DArray(&cuArray, &channelDesc, extentDesc, cudaArrayLayered); 
    mParams.dstArray = cuArray; 
    cudaMemcpy3D(&mParams); 

    // Set texture reference parameters 
    texRef.addressMode[0] = cudaAddressModeBorder; 
    texRef.filterMode = cudaFilterModeLinear; 
    texRef.normalized = false; 

    // Bind the array to the texture reference 
    cudaBindTextureToArray(texRef, cuArray, channelDesc); 

    // Allocate result of transformation in device memory 
    float *d_output; 
    cudaMalloc(&d_output, loclen * numlayers * sizeof(float)); 
    float h_output[loclen * numlayers]; 

    // Invoke kernel 
    dim3 dimBlock(16, 16, 1); 
    dim3 dimGrid((loclen + dimBlock.x - 1)/dimBlock.x, 
     (numlayers + dimBlock.y - 1)/dimBlock.y, 1); 
    interp1Kernel<<<dimGrid, dimBlock>>>(d_output, d_loc, loclen, numlayers); 

    // Print Results 
    printf("\n Original Indices \n"); 
    for (int i = 0; i < len; i++) printf(" %d ", i); 
    printf("\n Original array \n"); 
    for (int j = 0; j < numlayers; j++) { 
     for (int i = 0; i < len; i++) { 
      printf("%5.3f ", h_data[i + j*len]); 
     } 
     printf("\n"); 
    } 
    printf("\n Output Indices \n"); 
    for (int i = 0; i < loclen; i++) printf("%5.3f ", h_loc[i]); 
    printf("\n Output Array \n"); 
    cudaMemcpy(h_output, d_output, loclen * numlayers * sizeof(float), cudaMemcpyDeviceToHost); 
    for (int j = 0; j < numlayers; j++) { 
     for (int i = 0; i < loclen; i++) { 
      printf("%5.3f ", h_output[i + j*loclen]); 
     } 
     printf("\n"); 
    } 
    printf("\n"); 

    // Free device memory 
    cudaFreeArray(cuArray); 
    cudaFree(d_output); 

    return 0; 
}