1 #include "../common/book.h
"2 #include "
cuda_runtime.h
"3 #include "
device_launch_parameters.h
"4 #include "
device_functions.h"5
#define imin(a,b) (a6
7#define n (33*1024*1024)
8const
int threadsperblock = 256;9
const
int blockspergrid =
10 imin(32, (n / 2 + threadsperblock - 1) /threadsperblock);
1112
13 __global__ void dot(int size, float *a, float *b, float *c)
2324
//set the cache values
25 cache[cacheindex] =temp;
2627
//塊內線程同步
28__syncthreads();
2930
//for reductions, threadsperblock must be a power of 2
31//
because of the following code
32int i = blockdim.x / 2;33
while (i != 0
) 39
40if (cacheindex == 0
)41 c[blockidx.x] = cache[0
];42}43
4445
struct
datastruct ;
5354 unsigned winapi routine(void *pvoiddata)
55//
void* routine(void *pvoiddata)
5664
65int size = data->size;
66float *a, *b, c, *partial_c;
67float *dev_a, *dev_b, *dev_partial_c;
6869
//allocate memory on the cpu side
70 a = data->a;
71 b = data->b;
72 partial_c = (float*)malloc(blockspergrid*sizeof(float
));73
74//
allocate the memory on the gpu
75 handle_error(cudahostgetdevicepointer(&dev_a, a, 0
));76 handle_error(cudahostgetdevicepointer(&dev_b, b, 0
));77 handle_error(cudamalloc((void**)&dev_partial_c,
78 blockspergrid*sizeof(float
)));
7980
//offset 'a' and 'b' to where this gpu is gets it data
81 dev_a += data->offset;
82 dev_b += data->offset;
8384 dot << > >(size, dev_a, dev_b,
85dev_partial_c);
86//
copy the array 'c' back from the gpu to the cpu
87handle_error(cudamemcpy(partial_c, dev_partial_c,
88 blockspergrid*sizeof(float
),89
cudamemcpydevicetohost));
9091
//finish up on the cpu side
92 c = 0;93
for (int i = 0; i)
9697
handle_error(cudafree(dev_partial_c));
9899
//free memory on the cpu side
100free(partial_c);
101102 data->returnvalue =c;
103return0;
104}
105106
107int main(void
) 115
116cudadeviceprop prop;
117for (int i = 0; i<2; i++)
123}
124125
float *a, *b;
126 handle_error(cudasetdevice(0
));127
handle_error(cudasetdeviceflags(cudadevicemaphost));
128/*
129在設定了裝置0後，設定了分配記憶體的型別為cudahostallocportable，
130否則只有裝置0會將這些分配的記憶體視為固定記憶體
131只在device0中設定為可移動的
132*/
133 handle_error(cudahostalloc((void**)&a, n*sizeof(float
),134 cudahostallocwritecombined |
135 cudahostallocportable |
136137 handle_error(cudahostalloc((void**)&b, n*sizeof(float
),138 cudahostallocwritecombined |
139 cudahostallocportable |
140141
142//
fill in the host memory with data
143for (int i = 0; i)
147148
//prepare for multithread
149 datastruct data[2
];150 data[0].deviceid = 0
;151 data[0].offset = 0
;152 data[0].size = n / 2
;153 data[0].a =a;
154 data[0].b =b;
155156 data[1].deviceid = 1
;157 data[1].offset = n / 2
;158 data[1].size = n / 2
;159 data[1].a =a;
160 data[1].b =b;
161162 cutthread thread = start_thread(routine, &(data[1
]));
163 routine(&(data[0
]));
164end_thread(thread);
165166
167//
free memory on the cpu side
168handle_error(cudafreehost(a));
169handle_error(cudafreehost(b));
170171 printf("
value calculated: %f\n",
172 data[0].returnvalue + data[1
].returnvalue);
173174
return0;
175 }
可移動構造 可複製構造 可移動賦值 可複製賦值
記錄一下這幾個容易混淆的概念。可移動構造 moveconstructible 指定該型別的例項可以從乙個右值實參構造 定義 給定 下列表示式必須合法且擁有指定的效果 t u rv u 的值等於 rv 在初始化前的值。rv 的新值未指明。t rv t rv 的值等於 rv 在初始化前的值。rv 的新值...
美國C Spire測試5G固定和移動技術
據悉，美國地區運營商c spire日前宣布，公司打算支援更快速度 更低延遲服務，希望通過滿足使用者和市場需求從而成倍增加網際網路連線裝置。通過使用諾基亞裝置，c spire在密西西比洲成功演示了5g固定無線解決方案技術，也是業界首個。此次演示是通過直接連線到其光纖電視服務完成的。由於只有少數美國電信...
移動端固定位置
相信大家在移動開發的時候，都會遇到乙個 固定 的問題。那就是什麼頭部固定在螢幕頂部啊 什麼尾部固定在螢幕的尾部啊這些鬼東西。或者你現在看到這裡會想到 這還不簡單嗎？直接用個position fixed不就搞好了嗎，那麼菜還來寫文章。或者你沒想到，雖然position fixed是可以固定位置，但是在...