参考：http://galoisplusplus.coding.me/blog/2018/05/22/cudaErrorCudartUnloading/

主要是在linux下，使用cuda安装包里的cuda-memcheck来检查内存，它是类似于valgrind的存在。

基本编译配置

首先添加path：

vim ~/.zshrc
export PATH=$PATH:/usr/local/cuda/bin

基于CMake编写C程，cuda相关的内容：

option(use_cuda "Use CUDA?" ON)if (use_cuda)#list(APPEND CMAKE_PREFIX_PATH "/usr/local/cuda")if (CMAKE_SYSTEM_NAME MATCHES "Linux")set(CUDA_DIR "/usr/local/cuda")find_package(CUDA REQUIRED)include_directories("${CUDA_DIR}/include")link_directories("${CUDA_DIR}/lib64")elseif (CMAKE_SYSTEM_NAME MATCHES "Windows")set(CUDA_DIR "$ENV{CUDA_PATH}")find_package(CUDA REQUIRED)include_directories("${CUDA_DIR}/include")link_directories("${CUDA_DIR}/lib/x64")endif()
endif()if(use_cuda)list(APPEND TESTBED_DEP_LIBScudartcudart_staticcudacublas_devicecudnncublas)
endif()target_link_libraries(testbed ${TESTBED_DEP_LIBS})

使用：

cd ~/work/mycode
mkdir build
cd build
cmake ..
makecuda-memcheck ./run

一些常识

如果cudaMalloc()后忘记cudaFree()，用cuda-memcheck检查不出来这个错误。。据网友说是会自动回收，不知道cuda是否自带了gc功能2333

BN层的坑

如果是卷积层，可以直接调用cudnnConvolutionForward()函数，它使用caffe一样的padding方式：它的内部默默的做了对称的padding（pad_left == pad_right），调用起来挺容易的。

倒是BN层卡了好久，主要是官方文档有错误。官方文章这样写的：

Note: The input transformation performed by this function is defined as: y := alpha*y + beta *(bnScale * (x-estimatedMean)/sqrt(epsilon + estimatedVariance)+bnBias)...alpha, beta
Inputs. Pointers to scaling factors (in host memory) used to blend the layer output value with prior value in the destination tensor as follows: dstValue = alpha[0]*resultValue + beta[0]*priorDstValue.

按照贴出来的文档内容，按它第一行的说法，alpha应该设定为0，beta应该设定为1；按照后面几行的说法，应该是alpha设定为1，beta设定为0。

• Caffe：已经很久没有人来修了，它压根没有调用cudnnBatchNormalizationForwardInference()
• TensorFlow: 并没有调用。。不信你grep查一下，或者github上网页里在repo中搜索，找不到的。
• PyTorch: 还是比较良心，调用了cudnnBatchNormalizationForwardInference( )
• mini-caffe: 也是调用了cudnnBatchNormalizationForwardInference( )

主要参考mini-caffe和PyTorch，发现BN层应该设定alpha为1, beta为0，这样才能正确算出结果来。看看PyTorch的调用：

    AT_CUDNN_CHECK(cudnnBatchNormalizationForwardInference(handle, mode, &one, &zero, //这里，alpha是one, beta是zeroidesc.desc(), input->data_ptr(),idesc.desc(), output->data_ptr(),wdesc.desc(),weight->data_ptr(),bias->data_ptr(),running_mean->data_ptr(),running_var->data_ptr(),epsilon));

当然这里还有一个教训：但凡用CuDNN的xxxDescriptor参数说明的地方，它后面紧接着的几个变量，应该是gpu指针，也就是其值为GPU上的地址，而如果是CPU上的地址就会导致内存错误，比如代码为4或77。

举例：

cudnnStatus = cudnnBatchNormalizationForwardInference(layer_base->handle_cudnn,CUDNN_BATCHNORM_PER_ACTIVATION,&alpha,&beta,bnlayer->bn_layer_input_desc,layer_base->bottom_blobs_ptr[0]->gpu_data_ptr, //gpu上的bnlayer->bn_layer_output_desc,layer_base->top_blobs_ptr[0]->gpu_data_ptr, //gpu上的bnlayer->bn_layer_scale_bias_mean_var_desc,bnlayer->scale_dev,  //gpu上的bnlayer->shift_dev,    //gpu上的bnlayer->mean_dev,  //gpu上的bnlayer->variance_dev,//gpu上的bnlayer->eps
);

cuda基础

CUDA架构的目的，就是减轻早期GPU计算中存在的限制，这些限制使之前的GPU在通用计算中没有广泛应用。
换言之，CUDA就是为了让NVIDIA GPU用于通用计算的技术。

CUDA编程，使用C/C++即可，它仅仅是增加了个别关键字。不需要使用Shader Language，不需要掌握OpenGL，因此编程门槛降低。

编译器：CUDA C程序需要同时在CPU和N卡上执行计算，分别需要各自的编译器。GPU上的编译器，也就是nvcc了。CPU上的话cl.exe/gcc/clang等。

host和device: CPU以及main memory，叫做host；GPU上的memory称为device。 内存拷贝的时候，用cudaMemcpy()，要指定内存拷贝的来源和目的地，是host还是device。

kernel: GPU上执行的函数称为kernel。

向cuda核函数传入结构体指针？

C语言传参，本质上是拷贝一份，因此：简单数据类型直接传入即可，反正复制成本低；复杂类型则传入指针，复制成本低。

而如果需要向函数传入多个参数，就希望传入一个结构体或其指针，作为他们的一个封装。

然而，直接向cuda kernel函数传入CPU下的结构体指针是不行的。问题很明显：把host上的指针给device用，寻址会出现问题。

解决思路有三种：

1. 使用host和device的统一寻址
2. 分别创建host和device上的结构体指针，参考https://blog.csdn.net/u013701860/article/details/52605137
3. 向cuda kernel函数传入多个参数

举个栗子：传入结构体的拷贝是OK的(just_demo)，传入host上结构体指针的拷贝，cuda kernel函数会报废的(just_demo2):

#include <stdio.h>#define CHECK_CUDA(call) \
{ \const cudaError_t cudaStatus = call; \if (cudaStatus != cudaSuccess) { \printf("Error! in file %s, line %d, ", __FILE__, __LINE__); \printf("code: %d, reason: %s\n", cudaStatus, cudaGetErrorString(cudaStatus)); \exit(1); \} \
}typedef struct Pad{int left;int right;int top;int bottom;
}Pad;__global__ void just_demo(Pad pad) {printf("blockIdx: (%d, %d, %d), threadIdx: (%d, %d, %d)",blockIdx.x, blockIdx.y, blockIdx.z,threadIdx.x, threadIdx.y, threadIdx.z);printf("pad: left=%d, right=%d, top=%d, bottom=%d\n",pad.left,pad.right,pad.top,pad.bottom);
}__global__ void just_demo2(Pad* pad) {printf("blockIdx: (%d, %d, %d), threadIdx: (%d, %d, %d)",blockIdx.x, blockIdx.y, blockIdx.z,threadIdx.x, threadIdx.y, threadIdx.z);printf("pad: left=%d, right=%d, top=%d, bottom=%d\n",pad->left,pad->right,pad->top,pad->bottom);
}int main(void){Pad pad;pad.left = 1;pad.right = 0;pad.top = 1;pad.bottom = 0;size_t buf_size = 3*4*4*sizeof(float);float* data_host = (float*)malloc(buf_size);float* data_dev;CHECK_CUDA(cudaMalloc(&data_dev, buf_size));just_demo<<<3, 5>>> (pad);  // this is OK//just_demo2<<<3, 5>>> (&pad); // this will cause unspecified launch failureCHECK_CUDA(cudaMemcpy(data_host, data_dev, buf_size, cudaMemcpyDeviceToHost));return 0;
}