参考: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_LIBS
cudart
cudart_static
cuda
cublas_device
cudnn
cublas
)
endif()
target_link_libraries(testbed ${TESTBED_DEP_LIBS})
使用:
cd ~/work/mycode
mkdir build
cd build
cmake ..
make
cuda-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是zero
idesc.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用,寻址会出现问题。
解决思路有三种:
- 使用host和device的统一寻址
- 分别创建host和device上的结构体指针,参考https://blog.csdn.net/u013701860/article/details/52605137
- 向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 failure
CHECK_CUDA(cudaMemcpy(data_host, data_dev, buf_size, cudaMemcpyDeviceToHost));
return 0;
}