CUDA 纹理无法返回无符号长长长的值类型

CUDA texture can't return value type of unsigned long long

本文关键字：长长的 类型 无符号 返回 纹理 CUDA      更新时间：2023-10-16

我有代码：

#define int4 unsigned long long int
int4 mer_thread = tex2D(STexture, col, row);
printf("nTexture[%d][%d] = %d", row, col, tex2D(STexture, col, row));

错误"错误：没有重载函数"tex2D"的实例与参数列表""匹配

但如果是define int4 unsigned long int，则工作正常。

我的代码创建纹理：

void Creat_TexttureS(int4 _S[nmax][NMAX])
{
    cudaArray* carray;
    cudaChannelFormatDesc channel;      
    channel = cudaCreateChannelDesc<int4>();        
    cudaMallocArray(&carray, &channel, NMAX, nmax);    
    cudaMemcpyToArray(carray, 0, 0, _S, sizeof(int4)*NMAX*nmax, cudaMemcpyHostToDevice);  
    STexture.filterMode = cudaFilterModePoint;      
    STexture.addressMode[0] = cudaAddressModeWrap;
    STexture.addressMode[1] = cudaAddressModeClamp;     
    cudaBindTextureToArray(STexture, carray);
}

谢谢你的帮助！！

下面是一个示例，演示了在int2类型的2D纹理中存储long long int类型的数据，以及如何通过tex2D()检索数据并将其重新解释为long long int。

#include <stdlib.h>
#include <stdio.h>
// Macro to catch CUDA errors in CUDA runtime calls
#define CUDA_SAFE_CALL(call)                                          
do {                                                                  
    cudaError_t err = call;                                           
    if (cudaSuccess != err) {                                         
        fprintf (stderr, "Cuda error in file '%s' in line %i : %s.n",
                 __FILE__, __LINE__, cudaGetErrorString(err) );       
        exit(EXIT_FAILURE);                                           
    }                                                                 
} while (0)
// Macro to catch CUDA errors in kernel launches
#define CHECK_LAUNCH_ERROR()                                          
do {                                                                  
    /* Check synchronous errors, i.e. pre-launch */                   
    cudaError_t err = cudaGetLastError();                             
    if (cudaSuccess != err) {                                         
        fprintf (stderr, "Cuda error in file '%s' in line %i : %s.n",
                 __FILE__, __LINE__, cudaGetErrorString(err) );       
        exit(EXIT_FAILURE);                                           
    }                                                                 
    /* Check asynchronous errors, i.e. kernel failed (ULF) */         
    err = cudaThreadSynchronize();                                    
    if (cudaSuccess != err) {                                         
        fprintf (stderr, "Cuda error in file '%s' in line %i : %s.n",
                 __FILE__, __LINE__, cudaGetErrorString( err) );      
        exit(EXIT_FAILURE);                                           
    }                                                                 
} while (0)
__forceinline__ __device__ long long int int2_as_longlong (int2 a)
{
    long long int res;
    asm ("mov.b64 %0, {%1,%2};" : "=l"(res) : "r"(a.x), "r"(a.y));
    return res;
}
texture<int2, 2, cudaReadModeElementType> tex;
__global__ void kernel (int m, int n) 
{
    int2 data;
    for (int row = 0; row < m; row++) {
        for (int col = 0; col < n; col++) {
            data = tex2D (tex, col, row);
            printf ("% 11lld  ", int2_as_longlong (data));
        }
        printf ("n");
    }
}
int main (void)
{
    int m = 4; // height = #rows
    int n = 3; // width  = #columns
    size_t pitch, tex_ofs;
    unsigned long long int arr[4][3]= 
        {{11111111LL, 11112222LL, 11113333LL},
         {22221111LL, 22222222LL, 22223333LL},
         {33331111LL, 33332222LL, 33333333LL},
         {44441111LL, 44442222LL, 44443333LL}};
    int2 *arr_d = 0;
    CUDA_SAFE_CALL(cudaMallocPitch((void**)&arr_d,&pitch,n*sizeof(*arr_d),m));
    CUDA_SAFE_CALL(cudaMemcpy2D(arr_d, pitch, arr, n*sizeof(arr[0][0]),
                                n*sizeof(arr[0][0]),m,cudaMemcpyHostToDevice));
    CUDA_SAFE_CALL (cudaBindTexture2D (&tex_ofs, &tex, arr_d, &tex.channelDesc,
                                       n, m, pitch));
    if (tex_ofs !=0) {
        printf ("tex_ofs = %zun", tex_ofs);
        return EXIT_FAILURE;
    }
    printf ("printing texture contentn");
    kernel<<<1,1>>>(m, n);
    CHECK_LAUNCH_ERROR();
    CUDA_SAFE_CALL (cudaUnbindTexture (tex));
    CUDA_SAFE_CALL (cudaFree (arr_d));
    return EXIT_SUCCESS;
}
相关文章：