CUDA Malloc时间以及较大静态矩阵声明

[piaoyoter]

正在入门CUDA编程，目前碰到2个问题：
1. CUDA Malloc的时间较长。
我在程序中声明了一个256*256的一维数组，然后在GPU上对其进行简单的计算。结果发现Malloc的时间大约是98M cycles，数据拷贝时间大约是0.8M cycles。为什么Malloc时间这么长？不考虑计算，Malloc的时间比数据拷贝时间长100多倍，感觉不合理，请问为什么会出现这种情况？

2. 较大静态数据声明。
由于Malloc时间较长，所以我想把数据量扩大，以减少Malloc所占有的时间，以及测试Malloc时间是否与数据量相关。但是当我将静态数组修改为1024*1024时，运行直接退出。请问这种问题怎么解决？

希望有人能帮助我，谢谢！

程序代码如下：

#include <stdio.h>
#include "time.h"
#include <windows.h>
#include <cuda_runtime.h>


#define SIZE_Y 256
#define SIZE_X 256
#define BLOCK_NUM 2
#define THREAD_NUM 512
#define SIZE_Y_T 8
#define SIZE_X_T 8




bool InitCUDA()
{
	int count;
	cudaGetDeviceCount(&count);
	if(count == 0) {
		fprintf(stderr, "There is no device.\n");
		return false;
	}
	
	int i;
	for(i=0; i< count; i++) {
		cudaDeviceProp prop;
		if(cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
			fprintf(stderr, "Device prop major:%d\n", prop.major);
			if(prop.major >= 1) {
				break;
			}
		}
	}
	
	if(i == count) {
		fprintf(stderr, "There is no device supporting CUDA.\n");
		return false;
	}

	cudaSetDevice(i);
	return true;
}

inline unsigned __int64 GetCycleCount()
{
	__asm _emit 0x0F
	__asm _emit 0x31
}

__global__ static void stencil(int *num, int* result)
{
	int sum = 0;
	const int tid = threadIdx.x;
	const int bid = blockIdx.x;

	for(int j=1; j< SIZE_Y_T-1; j++) {
		for(int i=1; i< SIZE_X_T-1; i++) {
			//result[(j+tid/32*32)*1024+i+tid%32*32] = 0.5*num[(j+tid/32*32)*1024+i+tid%32*32]+0.125*(num[(j+tid/32*32)*1024+i+tid%32*32-1]+num[(j+tid/32*32)*1024+i+tid%32*32+1]+num[(j+tid/32*32-1)*1024+i+tid%32*32]+num[(j+tid/32*32+1)*1024+i+tid%32*32]);
			result[(j+(bid*512+tid)/(SIZE_Y/SIZE_Y_T)*SIZE_Y_T)*SIZE_X+(i+(bid*256+tid)%(SIZE_Y/SIZE_Y_T)*SIZE_X_T)] = 2*num[(j+(bid*512+tid)/(SIZE_Y/SIZE_Y_T)*SIZE_Y_T)*SIZE_X+(i+(bid*256+tid)%(SIZE_Y/SIZE_Y_T)*SIZE_X_T)]+3*(num[(j+(bid*512+tid)/(SIZE_Y/SIZE_Y_T)*SIZE_Y_T+1)*SIZE_X+(i+(bid*256+tid)%(SIZE_Y/SIZE_Y_T)*SIZE_X_T)]+num[(j+(bid*512+tid)/(SIZE_Y/SIZE_Y_T)*SIZE_Y_T-1)*SIZE_X+(i+(bid*256+tid)%(SIZE_Y/SIZE_Y_T)*SIZE_X_T)]+num[(j+(bid*512+tid)/(SIZE_Y/SIZE_Y_T)*SIZE_Y_T)*SIZE_X+(i+(bid*256+tid)%(SIZE_Y/SIZE_Y_T)*SIZE_X_T+1)]+num[(j+(bid*512+tid)/(SIZE_Y/SIZE_Y_T)*SIZE_Y_T)*SIZE_X+(i+(bid*256+tid)%(SIZE_Y/SIZE_Y_T)*SIZE_X_T-1)]);
		}
	}
}

int main()
{
	int data[SIZE_Y*SIZE_X];

	unsigned long long start;
	unsigned long long end;
    
	if(!InitCUDA()) {
		printf("CUDA initializing Error!\n");
		return 0;
	}
	printf("CUDA initialized.\n");
	
	for(int j=0; j< SIZE_Y; j++) {
		for(int i=0; i< SIZE_X; i++) {
			data[j*SIZE_X+i] = rand()%10;
		}
	}
    
	start = GetCycleCount();
	int* gpudata, *result;
	cudaMalloc((void**) &gpudata, sizeof(int) * SIZE_X*SIZE_Y);
	cudaMalloc((void**) &result, sizeof(int) * SIZE_X*SIZE_Y);
	end = GetCycleCount();
	printf("Malloc time:%lu cycles\n", end-start);
	cudaMemcpy(gpudata, data, sizeof(int) * SIZE_X*SIZE_Y, cudaMemcpyHostToDevice);
	end = GetCycleCount();
	printf("Data copy time:%lu cycles\n", end-start);

	stencil<<<BLOCK_NUM, THREAD_NUM, 0>>>(gpudata, result);

	end = GetCycleCount();
	printf("Computing time:%lu cycles\n", end-start);

	int data_t[SIZE_Y*SIZE_X];
	cudaMemcpy(data_t, result, sizeof(int) * SIZE_X*SIZE_Y, cudaMemcpyDeviceToHost);
	end = GetCycleCount();
	printf("Result copy time:%lu cycles\n", end-start);
	cudaFree(gpudata);
	cudaFree(result);

	//Sleep(1000);

    end = GetCycleCount();
	printf("Stencil(GPU):%lu cycles\n", end-start);

    /*
	int sum = 0;
	for(int j = 1; j < SIZE_Y-1; j++) {
		for(int i = 1; i < SIZE_X-1; i++) {
			if(i%64 != 0 && i%64 != 63 && j%64 != 0 && j%64 != 63) {
				//sum = 0.5*data[j*SIZE_X+i]+0.125*(data[j*SIZE_X+i-1]+data[j*SIZE_X+i+1]+data[(j-1)*SIZE_X+i]+data[(j+1)*SIZE_X+i]);
				sum = 2*data[j*SIZE_X+i]+3*(data[j*SIZE_X+i+1]+data[j*SIZE_X+i-1]+data[(j+1)*SIZE_X+i]+data[(j-1)*SIZE_X+i]);
				if(sum != data_t[j*SIZE_X+i]) {
					printf("Result Error:%d %d s:%d d:%d\n", j, i, sum, data_t[j*SIZE_X+i]);
				}
				else {
					printf("Result Correct:%d %d s:%d d:%d\n", j, i, sum, data_t[j*SIZE_X+i]);
				}
			}
		}
			
	}
	*/
    
	return 0;
}

程序的运行结果如下：


统计时间所用的函数是从网上找的，经过用Sleep函数测试是正确的。

[YCMyTot]

1楼说的很好，学习了。 在比较CPU与GPU时间的时候,cudaMalloc() 的时间是不考虑在内的，从cudaMemcpy() 到内核执行完以及cudaMemcpy(),拷贝回来时，这段时间才是我们需要测量的。 一般还是在host端用malloc来进行空间的分配，这样可以很好的进行内存是否分配成功的检验。 另外也可以采用cudaEvent来计算Host端的时间，相对而言精确一点。

[dennis200]

静态数组的申请是有限度，受栈长度限制，向1024*1024这种规模的数组就应该动态用malloc来申请了。

[piaoyoter]

引用 1 楼 dennis200 的回复:

静态数组的申请是有限度，受栈长度限制，向1024*1024这种规模的数组就应该动态用malloc来申请了。

多谢大神，果然C语言基础也是太差了

[piaoyoter]

引用 3 楼 u014693181 的回复:

1楼说的很好，学习了。 在比较CPU与GPU时间的时候,cudaMalloc() 的时间是不考虑在内的，从cudaMemcpy() 到内核执行完以及cudaMemcpy(),拷贝回来时，这段时间才是我们需要测量的。 一般还是在host端用malloc来进行空间的分配，这样可以很好的进行内存是否分配成功的检验。 另外也可以采用cudaEvent来计算Host端的时间，相对而言精确一点。

我是觉得统计一段程序的执行时间应该包含所有的过程，忽略掉Malloc时间不太合理，关键是这个时间还这么大就更不能忽略了。 我想知道大家在做程序测试的时候cudaMalloc的时间都是这么长的吗？比拷贝时间都长N倍，实在难以想象。