cuda 常量内存,头文件不知道有什么gui问题

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "device_functions.h"
#include "device_atomic_functions.h"
#include <iostream>
#include <stdio.h>
#include "crt\func_macro.h"
#include "assert.h"
#include "conio.h"
using namespace std;


#define CUDA_CALL(x){const cudaError_t a=(x);if(a!=cudaSuccess){printf("\nCUDAError:%s(err_num= %d) \n",cudaGetErrorString(a),a);cudaDeviceReset();assert(0);}}
#define KERNEL_LOOP 65536
__constant__ static const int const_data_01 = 0x55555555;
__constant__ static const int const_data_02 = 0x77777777;
__constant__ static const int const_data_03 = 0x33333333;
__constant__ static const int const_data_04 = 0x11111111;
__global__ void const_test_gpu_literal(int * const data, const int num_elements)
{
const int tid = (blockDim.x*blockIdx.x) + threadIdx.x;
if (tid < num_elements)
{
int d = 0x55555555;
for (int i = 0; i < KERNEL_LOOP; i++)
{
d ^= 0x55555555;
d |= 0x77777777;
d &= 0x33333333;
d |= 0x11111111;
}
data[tid] = d;
}
}
__global__ void const_test_gpu_const(int * const data, const int num_elements)
{
const int tid = (blockDim.x*blockIdx.x) + threadIdx.x;
if (tid < num_elements)
{
int d = const_data_01;
for (int i = 0; i < KERNEL_LOOP; i++)
{
d ^= const_data_01;
d |= const_data_02;
d &= const_data_03;
d |= const_data_04;
}
data[tid] = d;
}
}
__host__ void wait_exit(void)
{
char ch;
printf("\nPress any key to exit");
ch = getch();
}
__host__ void cuda_error_check(
const char * prefix,
const char * postfix)
{
if (cudaPeekAtLastError() != cudaSuccess)
{
printf("\n%s%s%s", prefix, cudaGetErrorString(cudaGetLastError()), postfix);
cudaDeviceReset();
wait_exit();
exit(1);
}
}
__host__ void gpu_kernel(void)
{
const int num_elements = (128 * 1024);
const int num_threads = 256;
const int num_blocks = (num_elements + (num_threads - 1)) / num_threads;
const int num_bytes = num_elements*sizeof(int);
int max_device_num;
const int max_runs = 6;
CUDA_CALL(cudaGetDeviceCount(&max_device_num));
for (int device_num = 0; device_num < max_device_num; device_num++)
{
CUDA_CALL(cudaSetDevice(device_num));
for (int num_test = 0; num_test < max_runs; num_test++)
{
int *data_gpu;
cudaEvent_t kernel_start1, kernel_stop1;
cudaEvent_t kernel_start2, kernel_stop2;
float delta_time1 = 0.0F, delta_time2 = 0.0F;
struct cudaDeviceProp device_prop;
char device_prefix[261];


CUDA_CALL(cudaMalloc(&data_gpu, num_bytes));
CUDA_CALL(cudaEventCreate(&kernel_start1));
CUDA_CALL(cudaEventCreate(&kernel_start2));
CUDA_CALL(cudaEventCreateWithFlags(&kernel_stop1, cudaEventBlockingSync));
CUDA_CALL(cudaEventCreateWithFlags(&kernel_stop2, cudaEventBlockingSync));


CUDA_CALL(cudaGetDeviceProperties(&device_prop, device_num));
sprintf(device_prefix, "ID:%d %s:", device_num, device_prop.name);


const_test_gpu_literal << <num_blocks, num_threads >> >(data_gpu, num_elements);
cuda_error_check("Error ", "return from literal starup kernel");
CUDA_CALL(cudaEventRecord(kernel_start1, 0));


const_test_gpu_literal << <num_blocks, num_threads >> >(data_gpu, num_elements);
cuda_error_check("Error ", "return from literal runtime kernel");
CUDA_CALL(cudaEventRecord(kernel_stop1, 0));
CUDA_CALL(cudaEventSynchronize(kernel_stop1));
CUDA_CALL(cudaEventElapsedTime(&delta_time1, kernel_start1, kernel_stop1));


const_test_gpu_literal << <num_blocks, num_threads >> >(data_gpu, num_elements);
cuda_error_check("Error ", "return from constant starup kernel");
CUDA_CALL(cudaEventRecord(kernel_start2, 0));


const_test_gpu_literal << <num_blocks, num_threads >> >(data_gpu, num_elements);
cuda_error_check("Error ", "return from constant runtime kernel");
CUDA_CALL(cudaEventRecord(kernel_stop2, 0));
CUDA_CALL(cudaEventSynchronize(kernel_stop2));
CUDA_CALL(cudaEventElapsedTime(&delta_time2, kernel_start2, kernel_stop2));


if (delta_time1 > delta_time2)
{
printf("\n%sConst version is faster by: %.2fms (Const=%.2fms vs. Literal=%.2fms)", device_prefix, delta_time1 = delta_time2, delta_time1, delta_time2);
}
else
{
printf("\n%sLiteral version is faster by: %.2fms (Const=%.2fms vs. Literal=%.2fms)", device_prefix, delta_time2 = delta_time1, delta_time1, delta_time2);
}
CUDA_CALL(cudaEventDestroy(kernel_start1));
CUDA_CALL(cudaEventDestroy(kernel_start2));
CUDA_CALL(cudaEventDestroy(kernel_stop1));
CUDA_CALL(cudaEventDestroy(kernel_stop2));
}
CUDA_CALL(cudaDeviceReset());
printf("\n");
}
wait_exit();
}


int main()
{
gpu_kernel();
return 0;
}