测试发现 cuda 的 launch kernel 的过程中不能释放该 gpu 上的显存，cuda 任务被释放的显存可能被 kernel 调用，所以禁止在 laucn kernel 过程中释放该 gpu 上的显存。 测试代码如下：
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
#include <stdio.h>
#include <thread>
#include <unistd.h>
#include <iostream>
#include <cuda_runtime.h>
#include <vector>
#define GRID_X 128
#define BLOCK_X 128
__global__ void proc(int *mem){
int idx = blockDim.x * blockIdx.x + threadIdx.x;
while(true){
mem[idx]++;
}
}
void run_thread()
{
cudaSetDevice(0);
int *pdata;
cudaMalloc(&pdata,GRID_X*BLOCK_X*sizeof(int));
sleep(1);
std::cout << "start kernel in run_thread" << std::endl;
proc<<<GRID_X,BLOCK_X>>>(pdata);
cudaDeviceSynchronize();
std::cout << "end kernel in run_thread" << std::endl;
}
void free_thread()
{
cudaSetDevice(0);
//cudaSetDevice(1);
int *pdata;
cudaMalloc(&pdata,GRID_X*BLOCK_X*sizeof(int));
sleep(5);
std::cout << "start cudaFree in free_thread" << std::endl;
cudaFree(pdata);
std::cout << "end cudaFree in free_thread" << std::endl;
}
int main(){
std::thread th0(&run_thread);
std::thread th1(&free_thread);
th0.join();
th1.join();
return 0;
}
运行过程中会发现 free_thread 的会阻塞在 cudaFree 函数上。
但是如果 cudaFree 释放另一块 gpu 上的显存，则完成不存在这个问题。
包上述代码 free_thred 线程中的 cudaSetDevice(0) 改成 cudaSetDevice(1) 则 cudaFree 可以被正常调用而不会被阻塞。