【CUDA学习笔记】第五篇：内存以及案例解释（附案例代码下载方式）（一）

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//使用__constant__定义两个常量constant_f，constant_g
__constant__ int constant_f;
__constant__ int constant_g;
#define N5
//定义设备的Kernel函数
__global__ void gpu_constant_memory(float *d_in, float *d_out) {
//线程的index
int tid = threadIdx.x;
d_out[tid] = constant_f * d_in[tid] + constant_g;
}
int main(void) {
//定义主机的数组
float h_in[N], h_out[N];
//定义设备指针
float *d_in, *d_out;
int h_f = 2;
int h_g = 20;
// 申请cpu内存
cudaMalloc((void**)&d_in, N * sizeof(float));
cudaMalloc((void**)&d_out, N * sizeof(float));
//初始化数组
for (int i = 0; i < N; i++) {
h_in[i] = i;
}
//把主机上的数组复制到设备上
cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice);
//把定义的常量复制到GPU设备的常量内存中
cudaMemcpyToSymbol(constant_f, &h_f, sizeof(int), 0, cudaMemcpyHostToDevice);
cudaMemcpyToSymbol(constant_g, &h_g, sizeof(int));
//使用1个N个线程的Block调用和执行Kernel函数
gpu_constant_memory << <1, N >> > (d_in, d_out);
//把设备运算的结果传回CPU主机
cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);
//打印结果
printf("Use of Constant memory on GPU \n");
for (int i = 0; i < N; i++) {
printf("The expression for input %f is %f\n", h_in[i], h_out[i]);
}
//释放GPU设备内存
cudaFree(d_in);
cudaFree(d_out);
return 0;
}

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#define NUM_THREADS 10
#define N 10
//通过【纹理引用】来定义一段能进行纹理拾取的纹理内存
texture <float, 1, cudaReadModeElementType> textureRef;
__global__ void gpu_texture_memory(int n, float *d_out)
{
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if (idx < n) {
float temp = tex1D(textureRef, float(idx));
d_out[idx] = temp;
}
}
int main()
{
//计算Block的个数，用于后面的使用
int num_blocks = N / NUM_THREADS + ((N % NUM_THREADS) ? 1 : 0);
//声明设备的指针
float *d_out;
// 在设备上为结果申请显存
cudaMalloc((void**)&d_out, sizeof(float) * N);
// 在主机设备上为结果申请内存
float *h_out = (float*)malloc(sizeof(float)*N);
//声明并初始化主机的数组
float h_in[N];
for (int i = 0; i < N; i++) {
h_in[i] = float(i);
}
//定义GPU设备的数组
cudaArray *cu_Array;
cudaMallocArray(&cu_Array, &textureRef.channelDesc, N, 1);
//把数据复制到GPU的数组中
cudaMemcpyToArray(cu_Array, 0, 0, h_in, sizeof(float)*N, cudaMemcpyHostToDevice);
// GPU数组绑定前面申请的纹理内存
cudaBindTextureToArray(textureRef, cu_Array);
//调用Kernel函数，并运算
gpu_texture_memory << <num_blocks, NUM_THREADS >> > (N, d_out);
// 把结果复制回主机以方便打印和处理
cudaMemcpy(h_out, d_out, sizeof(float)*N, cudaMemcpyDeviceToHost);
printf("Use of Texture memory on GPU: \n");
for (int i = 0; i < N; i++) {
printf("Texture element at %d is : %f\n", i, h_out[i]);
}
free(h_out);
cudaFree(d_out);
cudaFreeArray(cu_Array);
cudaUnbindTexture(textureRef);
}