我刚刚开始学习CUDA编程,我对减少有些困惑.
我知道全局内存与共享内存相比有很多访问延迟,但是我可以使用全局内存来(至少)模拟类似于共享内存的行为吗?
例如,我想总结一个长度正好的大数组的元素BLOCK_SIZE * THREAD_SIZE
(网格和块的维度都是幂2
),我试图使用下面的代码:
__global__ void parallelSum(unsigned int* array) { unsigned int totalThreadsNum = gridDim.x * blockDim.x; unsigned int idx = blockDim.x * blockIdx.x + threadIdx.x; int i = totalThreadsNum / 2; while (i != 0) { if (idx < i) { array[idx] += array[idx + i]; } __syncthreads(); i /= 2; } }
我比较了这段代码的结果和在主机上串行生成的结果,奇怪的是:有时结果是相同的,但有时它们显然是不同的.在这里使用全局内存是否有任何原因?
汤姆已经回答了这个问题.在他的回答中,他建议使用Thrust或CUB来减少CUDA.
在这里,我提供了一个关于如何使用两个库来执行缩减的完整工作示例.
#define CUB_STDERR #include <stdio.h> #include <thrust/device_ptr.h> #include <thrust/reduce.h> #include <thrust/execution_policy.h> #include <cub/device/device_reduce.cuh> #include "TimingGPU.cuh" #include "Utilities.cuh" using namespace cub; /********/ /* MAIN */ /********/ int main() { const int N = 8388608; gpuErrchk(cudaFree(0)); float *h_data = (float *)malloc(N * sizeof(float)); float h_result = 0.f; for (int i=0; i<N; i++) { h_data[i] = 3.f; h_result = h_result + h_data[i]; } TimingGPU timerGPU; float *d_data; gpuErrchk(cudaMalloc((void**)&d_data, N * sizeof(float))); gpuErrchk(cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice)); /**********/ /* THRUST */ /**********/ timerGPU.StartCounter(); thrust::device_ptr<float> wrapped_ptr = thrust::device_pointer_cast(d_data); float h_result1 = thrust::reduce(wrapped_ptr, wrapped_ptr + N); printf("Timing for Thrust = %f\n", timerGPU.GetCounter()); /*******/ /* CUB */ /*******/ timerGPU.StartCounter(); float *h_result2 = (float *)malloc(sizeof(float)); float *d_result2; gpuErrchk(cudaMalloc((void**)&d_result2, sizeof(float))); void *d_temp_storage = NULL; size_t temp_storage_bytes = 0; DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_data, d_result2, N); gpuErrchk(cudaMalloc((void**)&d_temp_storage, temp_storage_bytes)); DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_data, d_result2, N); gpuErrchk(cudaMemcpy(h_result2, d_result2, sizeof(float), cudaMemcpyDeviceToHost)); printf("Timing for CUB = %f\n", timerGPU.GetCounter()); printf("Results:\n"); printf("Exact: %f\n", h_result); printf("Thrust: %f\n", h_result1); printf("CUB: %f\n", h_result2[0]); }
请注意,由于不同的底层哲学,CUB可能比Thrust快一些,因为CUB会留下性能关键的细节,例如算法的确切选择以及未绑定和用户手中的并发程度.通过这种方式,可以调整这些参数,以便最大化特定体系结构和应用程序的性能.
在CUB in Action中报告了计算数组的欧几里德范数的比较- 一些使用CUB模板库的简单示例.