3.2.4.5 例子分析：Simple texture
// 2D float texture
texture<float, 2, cudaReadModeElementType> texRef;
// Simple transformation kernel
__global__ void transformKernel(float* output,
int width, int height, float theta)
{
    // 根据tid bid计算归一化的拾取坐标
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y * blockDim.y + threadIdx.y;
    float u = x / (float)width;
    float v = y / (float)height;
    // 旋转拾取坐标
    u -= 0.5f;
    v -= 0.5f;
    float tu = u * cosf(theta) –v * sinf(theta) + 0.5f;
    float tv = v * cosf(theta) + u * sinf(theta) + 0.5f;
    //从纹理存储器中拾取数据，并写入显存
    output[y * width + x] = tex2D(tex, tu, tv);
}
// Host code
int main()
{
    // 分配CUDA数组
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0,cudaChannelFormatKindFloat);
    cudaArray* cuArray;
    cudaMallocArray(&cuArray, &channelDesc, width, height);
    // Copy to device memory some data located at address h_data
    // in host memory
    cudaMemcpyToArray(cuArray, 0, 0, h_data, size, cudaMemcpyHostToDevice);
    // Set texture parameters
    texRef.addressMode[0] = cudaAddressModeWrap; //循环寻址方式
    texRef.addressMode[1] = cudaAddressModeWrap;
    texRef.filterMode = cudaFilterModeLinear;  //线性滤波，因为这里是一个图像。如果要保持原来的值则千万不要用线性滤波
    alized = true; //归一化坐标
    // Bind the array to the texture
    cudaBindTextureToArray(texRef, cuArray, channelDesc);
    // Allocate result of transformation in device memory
    float* output;
    cudaMalloc((voidtrunc函数实例**)&output, width * height * sizeof(float));
    // Invoke kernel
    dim3 dimBlock(16, 16);
    dim3 dimGrid((width + dimBlock.x –1) / dimBlock.x,(height + dimBlock.y –1) / dimBlock.y);
    transformKernel<<<dimGrid, dimBlock>>>(output, width, height,angle);
    // Free device memory
    cudaFreeArray(cuArray);
    cudaFree(output);
}

版权声明：本站内容均来自互联网，仅供演示用，请勿用于商业和其他非法用途。如果侵犯了您的权益请与我们联系QQ:729038198，我们将在24小时内删除。