- html - 出于某种原因,IE8 对我的 Sass 文件中继承的 html5 CSS 不友好?
- JMeter 在响应断言中使用 span 标签的问题
- html - 在 :hover and :active? 上具有不同效果的 CSS 动画
- html - 相对于居中的 html 内容固定的 CSS 重复背景?
在主机端,我正在读取 128 x 128 整数数组,其中随机值介于 0-31 之间。我有一个 Occurrences 数组,它存储值 0-31,然后在设备上我尝试执行一个内核,循环遍历 128 x 128 数组中的值,然后计算 0-31 出现的次数。
我遇到了如何在 CUDA 中分割 block /线程以及如何让内核向主机提供通信并打印出每个元素出现的次数的问题。这是我第一次使用 CUDA 和我将不胜感激任何建设性的建议!这是到目前为止我的代码:
#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#define MAXR 16
#define MAXC 16
#define N 256
__global__ void count(int *arrayONE_d, int *occurrences_d, int *occurrences_final_d) {
int count = 0;
//provide unique thread ID
int idx = threadIdx.x + blockIdx.x * blockDim.x;
int k;
//for(k=0; k < 32;k++) {
// occurrences_d[k]=k;
// }
if(idx < N) {
//for(k=0; k < MAXR*MAXC; k++) {
for(int j=0; j<32; j++) {
count =0;
if(arrayONE_d[idx]==occurrences_d[j]){
count+=1;
occurrences_final_d[j] =count;
}
else {}
}
}
//occurrences_final_d[0] = 77;
}
}
int main(void) {
//const int N = MAXR*MAXC;
int arr1_h[MAXR][MAXC];
//int *occurrences_h[0][32];
//creating arrays for the device (GPU)
//int *arr1_d;
int occurrences_h[32];
int *occurrences_d;
int *occurrences_final_h[32] = {0};
int *occurrences_final_d;
int *arrayONE_h[256] = {0};
int *arrayONE_d;
int i, j;
// allocating memory for the arrays on the device
cudaMalloc( (void**) &arrayONE_d, MAXR*MAXC*sizeof(int)); // change to 16384 when using 128x128
cudaMalloc( (void**) &occurrences_d, 32* sizeof(int));
cudaMalloc( (void**) &occurrences_final_d, 32*sizeof(int));
/*
for(i=0; i < 32; i++) {
occurrences_h[i] = i;
}
/*
*
*/
//Reading in matrix from .txt file and storing it in arr1 on the host (CPU)
FILE *fp;
fp =fopen("arrays16.txt","r");
// this loop takes the information from .txt file and puts it into arr1 matrix
for(i=0;i<MAXR;i++) {
for(j=0;j<MAXC;j++)
{
fscanf(fp,"%d\t", &arr1_h[i][j]);
}
}
for(i=0;i<MAXR;i++) {
printf("\n");
for(j=0;j<MAXC;j++) {
//printf("d\t", arr1_h[i][j]);
}
printf("\n\n");
}
int x,y;
int z=0;
// this loop flattens the 2d array and makes it a 1d array of length MAXR*MAXC
for(x=0;x<MAXR;x++)
{
for(y=0;y<MAXC;y++)
{
// printf("**%d ",arr1_h[x][y]);
arrayONE_h[z]= &arr1_h[x][y];
z++;
}
}
for(x=0; x < 256; x++) {
printf("%d\n", *arrayONE_h[x]);
//return 0;
}
int length = sizeof(arrayONE_h)/sizeof(arrayONE_h[0]);
printf("\n\n");
printf("**LENGTH = %d", length);
// copying the arrays/memory from the host to the device (GPU)
cudaMemcpy(arrayONE_d, &arrayONE_h, MAXR*MAXC*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(occurrences_d, &occurrences_h, 32*sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(occurrences_final_d, &occurrences_final_h, 32*sizeof(int), cudaMemcpyHostToDevice);
// how many blocks we will allocate
//dim3 DimGrid();
//how many threads per block we will allocate
dim3 DimBlock(256);
//kernel launch against the GPU
count<<<1, DimBlock>>>(arrayONE_d,occurrences_d,occurrences_final_d);
//copy the arrays post-computation from the device back to the host (CPU)
cudaMemcpy(&occurrences_final_h, occurrences_final_d, 32*sizeof(int), cudaMemcpyDeviceToHost);
cudaMemcpy(&occurrences_h, occurrences_d, 32*sizeof(int), cudaMemcpyDeviceToHost);
// some error checking - run this with cuda-memcheck when executing your code
cudaError_t errSync = cudaGetLastError();
cudaError_t errAsync = cudaDeviceSynchronize();
if (errSync != cudaSuccess)
printf("Sync kernel error: %s\n", cudaGetErrorString(errSync));
if (errAsync != cudaSuccess)
printf("Async kernel error: %s\n", cudaGetErrorString(errAsync));
//free up the memory of the device arrays
cudaFree(arrayONE_d);
cudaFree(occurrences_d);
cudaFree(occurrences_final_d);
//print out the number of occurrences of each 0-31 value
for(i=0;i<32;i++) {
printf("\n");
printf("%d\n",occurrences_final_h[i]);
}
}
最佳答案
正如我在评论中提到的,您对指针的理解是有缺陷的。我在您的代码中的许多地方进行了更改来解决这个问题。我已经用注释 //mod
标记了其中的大部分内容,但我可能遗漏了一些。
此外,当多个线程可以更新同一位置时,您的内核根本无法跟踪元素。解决这个问题的一种方法是使用原子(我已经演示过)。还有各种其他方法,例如并行缩减,但这些都不是对内核的微不足道的更改。此外,您的内核逻辑在某些方面被破坏了。
接下来是我可以对您的代码进行的最少修改,以获得合理的结果。您可以使用一些编译开关来探索不同的内核行为:
-DUSE_ATOMICS
将演示对内核的修改,以使其正确计数。-DUSE_ALT_KERNEL
探索了一种不同的内核逻辑方法:为每个直方图箱分配一个线程,并让每个线程遍历整个数组,跟踪属于该箱的元素。由于只有一个线程写入每个 bin 结果,因此不需要原子。然而,我们只能拥有与容器一样多的线程(通过这个简单的实现)。没有太多困难,这种方法可能可以扩展到每个箱一个经纱,using warp shuffle to do a final warp-level reduction在让一个线程将最终结果写入垃圾箱之前。这将在一定程度上提高内存访问效率。然而,这也会给内核带来你可能还没有学到的复杂性。这是代码:
$ cat t316.cu
#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#define MAXR 16
#define MAXC 16
#define BINS 32
#define N (MAXR*MAXC)
__global__ void count(int *arrayONE_d, int *occurrences_d, int *occurrences_final_d) {
//provide unique thread ID
int idx = threadIdx.x + blockIdx.x * blockDim.x;
#ifndef USE_ALT_KERNEL
if(idx < N) {
//for(k=0; k < MAXR*MAXC; k++) {
for(int j=0; j<32; j++) {
if(arrayONE_d[idx]==occurrences_d[j]){
#ifndef USE_ATOMICS
occurrences_final_d[j]++;
#else
atomicAdd(occurrences_final_d+j, 1);
#endif
}
else {}
}
}
#else
// use one thread per histo bin
if (idx < BINS){
int count = 0;
int myval = occurrences_d[idx];
for (int i = 0; i < N; i++) if (arrayONE_d[i] == myval) count++;
occurrences_final_d[idx] = count;
}
#endif
}
int main(void) {
//const int N = MAXR*MAXC;
int arr1_h[MAXR][MAXC];
//int *occurrences_h[0][32];
//creating arrays for the device (GPU)
//int *arr1_d;
int occurrences_h[32]; // mod
int *occurrences_d;
int occurrences_final_h[32] = {0}; // mod
int *occurrences_final_d;
int arrayONE_h[256] = {0}; // mod
int *arrayONE_d;
int i, j;
// allocating memory for the arrays on the device
cudaMalloc( (void**) &arrayONE_d, MAXR*MAXC*sizeof(int)); // change to 16384 when using 128x128
cudaMalloc( (void**) &occurrences_d, 32* sizeof(int));
cudaMalloc( (void**) &occurrences_final_d, 32*sizeof(int));
/*
for(i=0; i < 32; i++) {
occurrences_h[i] = i;
}
*/
//Reading in matrix from .txt file and storing it in arr1 on the host (CPU)
// FILE *fp;
// fp =fopen("arrays16.txt","r");
// this loop takes the information from .txt file and puts it into arr1 matrix
for(i=0;i<MAXR;i++) {
for(j=0;j<MAXC;j++)
{
// fscanf(fp,"%d\t", &arr1_h[i][j]);
arr1_h[i][j] = j; // mod
}
}
for(i=0;i<MAXR;i++) {
for(j=0;j<MAXC;j++) {
//printf("d\t", arr1_h[i][j]);
}
}
int x,y;
int z=0;
// this loop flattens the 2d array and makes it a 1d array of length MAXR*MAXC
for(x=0;x<MAXR;x++)
{
for(y=0;y<MAXC;y++)
{
// printf("**%d ",arr1_h[x][y]);
arrayONE_h[z]= arr1_h[x][y]; // mod
z++;
}
}
for(x=0; x < 256; x++) {
// printf("%d\n", arrayONE_h[x]); // mod
//return 0;
}
int length = sizeof(arrayONE_h)/sizeof(arrayONE_h[0]);
printf("**LENGTH = %d\n", length);
// copying the arrays/memory from the host to the device (GPU)
cudaMemcpy(arrayONE_d, arrayONE_h, MAXR*MAXC*sizeof(int), cudaMemcpyHostToDevice); //mod
cudaMemcpy(occurrences_d, occurrences_h, 32*sizeof(int), cudaMemcpyHostToDevice); // mod
cudaMemcpy(occurrences_final_d, occurrences_final_h, 32*sizeof(int), cudaMemcpyHostToDevice); // mod
// how many blocks we will allocate
//dim3 DimGrid();
//how many threads per block we will allocate
#ifndef USE_ALT_KERNEL
dim3 DimBlock(N);
#else
dim3 DimBlock(BINS);
#endif
//kernel launch against the GPU
count<<<1, DimBlock>>>(arrayONE_d,occurrences_d,occurrences_final_d);
//copy the arrays post-computation from the device back to the host (CPU)
cudaMemcpy(occurrences_final_h, occurrences_final_d, 32*sizeof(int), cudaMemcpyDeviceToHost); // mod
cudaMemcpy(occurrences_h, occurrences_d, 32*sizeof(int), cudaMemcpyDeviceToHost); // mod
// some error checking - run this with cuda-memcheck when executing your code
cudaError_t errSync = cudaGetLastError();
cudaError_t errAsync = cudaDeviceSynchronize();
if (errSync != cudaSuccess)
printf("Sync kernel error: %s\n", cudaGetErrorString(errSync));
if (errAsync != cudaSuccess)
printf("Async kernel error: %s\n", cudaGetErrorString(errAsync));
//free up the memory of the device arrays
cudaFree(arrayONE_d);
cudaFree(occurrences_d);
cudaFree(occurrences_final_d);
//print out the number of occurrences of each 0-31 value
for(i=0;i<32;i++) {
printf("%d ",occurrences_final_h[i]);
}
printf("\n");
}
$ nvcc -o t316 t316.cu
$ cuda-memcheck ./t316
========= CUDA-MEMCHECK
**LENGTH = 256
1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1 1
========= ERROR SUMMARY: 0 errors
$ nvcc -o t316 t316.cu -DUSE_ATOMICS
$ ./t316
**LENGTH = 256
16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16
$ nvcc -o t316 t316.cu -DUSE_ALT_KERNEL
$ cuda-memcheck ./t316
========= CUDA-MEMCHECK
**LENGTH = 256
16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16 16
========= ERROR SUMMARY: 0 errors
$
在上面的输出中,我们看到基本内核产生了错误的结果。原子内核和备用内核产生正确的结果
(您的代码已修改为使用合成数据,因此不需要打开文件。)
关于CUDA + 使用 C 计算 int 元素出现次数,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/53121732/
SQL 和一般开发的新手,我有一个表(COUNTRIES),其中包含字段(INDEX、NAME、POPULATION、AREA) 通常我添加一个客户端(Delphi)计算字段(DENSITY)和 On
我想使用 calc(100%-100px),但在我的 demo 中不起作用由于高度只接受像素,因此如何将此百分比值转换为像素。 最佳答案 以下将为您提供高度: $(window).height();
我正在尝试在 MySQL 中添加列并动态填充其他列。 例如我有一张表“数字”并具有第 1 列、第 2 列、第 3 列,这些总数应填充在第 4 列中 最佳答案 除非我误解了你的问题,否则你不只是在寻找:
我想返回简单计算的结果,但我不确定如何执行此操作。我的表格如下: SELECT COUNT(fb.engineer_id) AS `total_feedback`, SUM(fb.ra
我一直在尝试做这个程序,但我被卡住了,我仍然是一个初学者,任何帮助将不胜感激。我需要程序来做 打印一个 10 X 10 的表格,其中表格中的每个条目都是行号和列号的总和 包含一个累加器,用于计算所有表
这个计算背后一定有一些逻辑。但我无法得到它。普通数学不会导致这种行为。谁能帮我解释一下原因 printf ("float %f\n", 2/7 * 100.0); 结果打印 1.000000 为什么会
我想计算从 0 到 (n)^{1/2} - 1 的数字的 AND每个数字从 0 到 (n)^{1/2} - 1 .我想在 O(n) 中执行此操作时间,不能使用 XOR、OR、AND 运算。 具体来说,
如何在 Excel 中将公式放入自定义数字格式?例如(出于说明目的随机示例), 假设我有以下数据: 输入 输出 在不编辑单元格中的实际数据的情况下,我想显示单元格中的值除以 2,并保留两位小数: 有没
每次我在 Flutter 应用程序中调用计算()时,我都会看到内存泄漏,据我所知,这基本上只是一种生成隔离的便捷方法。我的应用程序内存占用增加并且在 GC 之后永远不会减少。 我已将我的代码简化为仅调
我有数字特征观察 V1通过 V12用于目标变量 Wavelength .我想计算 Vx 之间的 RMSE列。数据格式如下。 每个变量“Vx”以 5 分钟的间隔进行测量。我想计算所有 Vx 变量的观测值
我正在寻找一种使用 C 语言计算文件中未知字符数的简单方法。谢谢你的帮助 最佳答案 POSIX 方式(可能是您想要的方式): off_t get_file_length( FILE *file ) {
我正在使用 Postgres,并且我正试图围绕如何在连续日期跨度中得出第一个开始日期的问题进行思考。例如 :- ID | Start Date | End Date =================
我有一个订单表格,我在其中使用 jQuery 计算插件来汇总总数。 此求和工作正常,但生成的“总和”存在问题。总之,我希望用逗号替换任何点。 代码的基础是; function ($this) {
我在使用 double 变量计算简单算术方程时遇到问题。 我有一个具有 double 属性 Value 的组件,我将此属性设置为 100。 然后我做一个简单的减法来检查这个值是否真的是 100: va
我在这里看到了一些关于 CRC 32 计算的其他问题。但没有一个让我满意,因此是这样。 openssl 库是否有任何用于计算 CRC32 的 api 支持?我已经在为 SHA1 使用 openssl,
当我在PHP日期计算中遇到问题时,我感到惊讶。 $add = '- 30 days'; echo date('Y-m-01', strtotime($add)); // result is 2017-
我正在使用 javascript 进行练习,我编写了这个脚本来计算 2 个变量的总和,然后在第三个方程中使用这个总和!关于如何完成这项工作的任何想法都将非常有用! First Number:
我有一个来自EAC的提示单和一个包含完整专辑的FLAC文件。 我正在尝试制作一些python脚本来播放文件,因为我需要能够设置在flac文件中开始的位置。 如何从CueSheet格式MM:SS:FF转
这个问题已经有答案了: Adding two numbers concatenates them instead of calculating the sum (24 个回答) 已关闭去年。 我有一个
4000 我需要上面字段 name="quantity" 和 id="price" 中的值,并使用 javascript 函数进行计算,并将其显示在字段 id= 中仅当我单击计算按钮时才显示“总
我是一名优秀的程序员,十分优秀!