CUDA编程: CUDA内存模型概述

nvcc -Xptxas -v *.cu

ptxas info    : 0 bytes gmem
ptxas info    : Compiling entry function '_Z9MatrixMulP6MatrixS0_S0_' for 'sm_52'
ptxas info    : Function properties for _Z9MatrixMulP6MatrixS0_S0_
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 32 registers, 344 bytes cmem[0]
ptxas info    : Compiling entry function '_Z9MatrixAddP6MatrixS0_S0_' for 'sm_52'
ptxas info    : Function properties for _Z9MatrixAddP6MatrixS0_S0_
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 18 registers, 344 bytes cmem[0]

__global__ void __launch_bounds__ (maxThreadsPerBlock, minBlocksPerMultiprocessor)
kernel_func (...) {
	//kernel body
}

__global__ void __launch_bounds__ (1024, 1) MatrixAdd(Matrix *A, Matrix *B, Matrix *C) {
        int row = threadIdx.y + blockIdx.y * blockDim.y;
        int col = threadIdx.x + blockIdx.x * blockDim.x;
        setValue(C, row, col, getValue(A, row, col) + getValue(B, row, col));
}

void __syncthreads();

cudaError_t cudaFuncSetCacheConfig(const void * func,enum cudaFuncCache cacheConfig);

#include <stdio.h>
#include <stdlib.h>

#include <cuda_runtime.h>

// 设置指定函数的缓存配置
cudaError_t cudaFuncSetCacheConfig(const void * func, enum cudaFuncCache cacheConfig, enum cudaFuncCache cacheConfigAlt) {
  // 定义错误码
  cudaError_t error = cudaSuccess;

  // 使用 CUDA 驱动 API 设置缓存配置
  error = cudaFuncSetCacheConfig(func, cacheConfig);
  if (error != cudaSuccess) {
    return error;
  }

  // 使用 CUDA 驱动 API 设置备用缓存配置
  error = cudaFuncSetCacheConfig(func, cacheConfigAlt);
  if (error != cudaSuccess) {
    return error;
  }

  // 返回成功
  return cudaSuccess;
}

__global__ void Kernel_func() {
  // 这是一个什么也不做的空内核
}

int main() {
  // 定义要设置的函数
  void *func = (void*)Kernel_func;

  // 定义所需的缓存配置
  enum cudaFuncCache cacheConfig = cudaFuncCachePreferL1;

  // 为函数设置缓存配置
  cudaError_t error = cudaFuncSetCacheConfig(func, cacheConfig, cudaFuncCachePreferNone);

  // 检查错误
  if (error != cudaSuccess) {
    // 打印错误消息并退出
    printf("Error setting cache configuration: %s\n", cudaGetErrorString(error));
    exit(EXIT_FAILURE);
 }

  printf("Success!\n");
  // 缓存配置设置成功
  return 0;
}

cudaError_t cudaMemcpyToSymbol(const void* symbol,const void *src,size_t count);

#include<cuda_runtime.h>
__constant__ int a[100];
int main()
{
    int h_a[100];
    for (int i = 0; i < 100; i++)
        h_a[i] = i;
    cudaMemcpyToSymbol(a, h_a, sizeof(h_a));
    return 0;
}

cudaError_t cudaCreateTextureObject(
    cudaTextureObject_t *pTexObject,
    const cudaResourceDesc *pResDesc,
    const cudaTextureDesc *pTexDesc,
    const cudaResourceViewDesc *pResViewDesc
);

#include <cuda_runtime.h>
#include <stdio.h>

using namespace std;

// 定义 CUDA 内核函数，名为 transformKernel
// 该函数有三个输入参数：
// - output：浮点型指针，指向结果数组
// - texObj：cudaTextureObject_t 类型，表示一个 CUDA 纹理对象
// - width 和 height：表示纹理对象的宽度和高度
__global__ void transformKernel(float* output, cudaTextureObject_t texObj, int width, int height) {
    // 计算当前线程的纵横坐标
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    // 如果超出纹理对象的范围，则退出该函数
    if ( x<0 || x>width || y<0 || y>height )
        return;
    // 在纹理对象中查找 (x+0.5, y+0.5) 处的值，并将其赋值给 output 数组对应的位置
    output[ y*width+x ] = tex2D<float>(texObj, x+0.5f, y+0.5f);
}

int main() {
    // 定义纹理对象的尺寸
    int width = 10;
    int height = 10;
    int size = width*height*sizeof(float);

    // 分配内存空间给源数据
    float *h_data = new float[width*height];

    // 初始化原始数据
    for (int y = 0; y<height; y++) {
        for (int x = 0; x<width; x++) {
                h_data[y*width + x] = x;
        }
    }

    // 打印原始数据
    printf("old:\n");
    for (int y = 0; y<height; y++) {
        for (int x = 0; x<width; x++) {
                printf("%f ", h_data[y*width + x]);
        }
        printf("\n");
    }

    // 创建一个用于存储源数据的 CUDA 数组
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
    cudaArray* cuArray;
    cudaMallocArray(&cuArray, &channelDesc, width, height);
    // 将源数据复制到 CUDA 数组中
    cudaMemcpyToArray(cuArray, 0, 0, h_data, size,cudaMemcpyHostToDevice);

    // 定义一个 cudaResourceDesc 结构体，用于描述纹理资源
    struct cudaResourceDesc resDesc;
    // 将该结构体的所有字段清零
    memset(&resDesc, 0, sizeof(resDesc));
    // 设置纹理资源的类型为数组资源
    resDesc.resType = cudaResourceTypeArray;
    // 设置纹理资源所指向的数组
    resDesc.res.array.array = cuArray;

    // 定义一个 cudaTextureDesc 结构体，用于描述纹理对象
    struct cudaTextureDesc texDesc;
    // 将该结构体的所有字段清零
    memset(&texDesc, 0, sizeof(texDesc));
    // 设置纹理坐标超出纹理范围时的采样模式
    texDesc.addressMode[0] = cudaAddressModeBorder;
    texDesc.addressMode[1] = cudaAddressModeBorder;
    // 设置纹理采样滤波模式
    texDesc.filterMode = cudaFilterModeLinear;
    // 设置纹理的读取模式
    texDesc.readMode = cudaReadModeElementType;
    // 设置纹理坐标是否归一化
    texDesc.normalizedCoords = 0;

    // 定义一个 cudaTextureObject_t 类型的变量，表示纹理对象
    cudaTextureObject_t texObj = 0;
    // 使用纹理资源和纹理描述创建纹理对象
    cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);

    // 分配内存给存储结果数据的数组
    float* output;
    cudaMalloc((void**)&output, size);

    // 计算线程块和网格的维度
    dim3 dimBlock(4, 4);
    dim3 dimGrid( max( (width  + dimBlock.x - 1) / dimBlock.x,1),
                      max( (height + dimBlock.y - 1) / dimBlock.y,1) );
    // 执行 CUDA 内核函数
    transformKernel <<<dimGrid, dimBlock >>>(output, texObj, width, height);

    // 将结果数据从 GPU 拷贝到 CPU 上
    cudaMemcpy(h_data, output, size, cudaMemcpyDeviceToHost);

    // 打印结果数据
    printf("new:\n");
    for (int y = 0; y<height; y++) {
        for (int x = 0; x<width; x++) {
                printf("%f ",h_data[y*width + x]);
        }
        printf("\n");
    }

    // 销毁纹理对象
    cudaDestroyTextureObject(texObj);

    // 释放 CUDA 数组的内存
    cudaFreeArray(cuArray);
    // 释放结果数据的内存
    cudaFree(output);

    // 释放原始数据的内存
    delete[]h_data;

    return 0;
}

__host__ cudaError_t cudaMemcpyToSymbol ( const void* symbol, const void* src, size_t count, size_t offset = 0, cudaMemcpyKind kind = cudaMemcpyHostToDevice )

__host__ cudaError_t cudaMemcpyFromSymbol ( void* dst, const void* symbol, size_t count, size_t offset = 0, cudaMemcpyKind kind = cudaMemcpyDeviceToHost )

#include <cuda_runtime.h>
#include <stdio.h>
__device__ float devData;
__global__ void checkGlobalVariable() {
    printf("Device: The value of the global variable is %f\n",devData);
    devData+=2.0;
}
int main() {
    float value=3.1415926f;
    cudaMemcpyToSymbol(devData,&value,sizeof(float));
    printf("Host: copy %f to the global variable\n",value);
    checkGlobalVariable<<<1,1>>>();
    cudaMemcpyFromSymbol(&value,devData,sizeof(float));
    printf("Host: the value changed by the kernel to %f \n",value);
    cudaDeviceReset();
    return EXIT_SUCCESS;
}

cudaMemcpy(&value,devData,sizeof(float));

float *dptr=NULL;
cudaGetSymbolAddress((void**)&dptr,devData);
cudaMemcpy(dptr,&value,sizeof(float),cudaMemcpyHostToDevice);