CUDA编程中遇到的警告信息：Warning: Cannot tell what pointer points to, assuming global memory space

     在写CUDA程序时，遇到了Warning: Cannot tell what pointer points to, assuming global memory space这样的警告，查阅了很多资料。终于，找到原因。下面是一个实例程序.

/*
* Copyright 徐洪志(西北农林科技大学.信息工程学院).  All rights reserved.
* Data: 2012-5-21
*/
//
// 此程序是是演示了使用constant显存常量时，遇到的编译警告:Warning: Cannot tell what pointer points to, assuming global memory space
// 其原因是，在核函数的参数中结构体参数使用了指针传递。CUDA指导书中，提到了尽量避免在核函数中使用指针及指针的指针
// 因此，在以后的编程中，核函数参数中若是有结构体，则要用值传递(核函数不支持引用传递)。
#include <stdlib.h>
#include <stdio.h>
#include <cuda.h>
#include <cutil_inline.h>
#include <cuda_runtime.h>

__constant__ int cosT = 10; 
__constant__ int ROW = 5;
__constant__ int COL;

// 一个结构体
typedef struct
{
	int r;  // 行数
	int c;  // 列数
	int *data;  // 数据
}Data;
// 核函数，结构体参数按值传递
__global__ void AddConstant(Data m_data)
{
	int x = blockIdx.x * blockDim.x + threadIdx.x;
	int y = blockIdx.y * blockDim.y + threadIdx.y;

	if(x<COL && y<ROW)  
	{
		m_data.data[y*m_data.c+x] += cosT;
	}
}
// 核函数，结构体参数按指针传递
__global__ void AddConstantTwo(Data *m_data)
{
	int x = blockIdx.x * blockDim.x + threadIdx.x;
	int y = blockIdx.y * blockDim.y + threadIdx.y;

	if(x<COL && y<ROW)
	{
		m_data->data[y*m_data->c+x] += cosT;
	}
}

int main(int argc, char **argv)
{
	CUT_DEVICE_INIT(argc, argv);   // 初始化CUDA
	// 1. 将数据拷贝到常数存储器，cudaMemcpyToSymbol第一个参数直接是常数存储器名，而不用取地址
	int m_col = 3;
	cutilSafeCall( cudaMemcpyToSymbol(COL, &m_col, sizeof(int), 0, cudaMemcpyHostToDevice));
	
	int i, j;
	// 设备结构体
	Data dev_Data;
	dev_Data.r = 10;
	dev_Data.c = 6;
	cutilSafeCall( cudaMalloc((void**) &dev_Data.data, sizeof(int)*dev_Data.r*dev_Data.c));  // 申请显存空间
	cutilSafeCall( cudaMemset(dev_Data.data, 0, sizeof(int)*dev_Data.r*dev_Data.c));  // 显存空间初始化为0
	// 主机结构体
	Data host_Data;
	host_Data.r = dev_Data.r;  // 行和列与设备结构体对象对应相等
	host_Data.c = dev_Data.c;
	host_Data.data = (int*)calloc(host_Data.r*host_Data.c, sizeof(int));  // 申请内存空间
	
	dim3 dimBlock(16, 16, 1);
	dim3 dimGrid(2, 2, 1);
	// 1.核函数中结构体参数用值传递方式
	AddConstant<<<dimGrid, dimBlock>>>(dev_Data);
	// 2.核函数中结构体参数用指针传递 :
	// 错误的，编译会产生警告：Warning: Cannot tell what pointer points to, assuming global memory space
	// 并且核函数调用失败。
	///AddConstantTwo<<<dimGrid, dimBlock>>>(&dev_Data);

	// 将显存数据拷贝入内存
	cutilSafeCall( cudaMemcpy(host_Data.data, dev_Data.data, sizeof(int)*dev_Data.r*dev_Data.c, cudaMemcpyDeviceToHost));
	
	printf("host_Data : \n");
	for(i=0; i<host_Data.r; i++)
	{
		for(j=0; j<host_Data.c; j++)
			printf("%d\t", host_Data.data[i*host_Data.c+j]);
		printf("\n");
	}
	
	cutilSafeCall( cudaFree(dev_Data.data)); // 释放显存空间
	free(host_Data.data);   // 释放内存空间
	CUT_EXIT(argc, argv);   // 退出CUDA
}