gpt4 book ai didi

cuda - 减少 CUDA

转载 作者:行者123 更新时间:2023-12-04 06:24:31 26 4
gpt4 key购买 nike

刚开始学习CUDA编程,对归约有些迷茫。

我知道与共享内存相比,全局内存有很多访问延迟,但我可以使用全局内存来(至少)模拟类似于共享内存的行为吗?

例如,我想对长度恰好为 BLOCK_SIZE * THREAD_SIZE 的大数组的元素求和(网格和 block 的维度都是 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;
}
}

我对比了这段代码的结果和在主机上串行生成的结果,奇怪的是:有时结果相同,但有时却明显不同。是否有任何与此处使用全局内存相关的原因?

最佳答案

Tom 已经回答了这个问题。在他的回答中,他建议使用 ThrustCUB在 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 – some simple examples using the CUB template library 中报告了计算数组欧几里得范数的比较。 .

关于cuda - 减少 CUDA,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/20324277/

26 4 0
Copyright 2021 - 2024 cfsdn All Rights Reserved 蜀ICP备2022000587号
广告合作:1813099741@qq.com 6ren.com