cuda中__device__变量值怎样改变

加速计算 CUDA
1

在.cu文件有:
device int count = 0;

global void kernel_fun(){
int count_tmp = atomicAdd(&count,1);
…
}

extern “C” void test(){
kernel_fun<<<>>>();
}

下面是.cpp文件；

for(int i=0;i<5;i++)
test();//怎样实现，每次调用test时，在.cu中count都是从 0 开始

2

您可以在kernel里面添加对count的清零操作
global void kernel_fun()
{ int tid=blockIdx.x*blockDim.x+threadIdx.x;
if tid==0
count=0;
__syncthreads();
int count_temp=atomicAdd(&count,1);
C[count_temp]=A[tid]+B[tid];
…
}

或者也可以在host端代码上每次调用kernel之前对count清零，比如说，您可以在将kernel封装到test()时增加一行

void test()
{
cudaMemset(&count,0,sizeof(int));——这个是错误的写法，请掘坟鞭尸挫骨扬灰以免贻害大家
kernel_fun<<<>>>();
…
}

使用__device__ 定义的global memory的变量，需要先取回地址保存在host端的指针变量中，才可以使用cudaMemset()函数清零。
int* count_ptr=NULL;
…

void test()
{
cudaGetSymbolAddress((void**) &count_ptr,“count”);
cudaMemset(count_ptr,0,sizeof(int));
kernel_fun<<<>>>();
…
}

上述两种用法推荐使用在kernel里面对count清零的做法。

前次答错深表歉意，希望没有对您造成误导，祝您编码愉快~

3

int* count_ptr=NULL;
…

void test()
{
cudaGetSymbolAddress((void**) &count_ptr,“count”);
cudaMemset(count_ptr,0,sizeof(int));
kernel_fun<<<>>>();
…
}

这样做法结果不对诶

4

报错了么？还是结果不对呢？目测似乎应该是可以的。
方便的话请反馈下问题，我明天写段测试看下。

推荐先使用kernel内部清零解决，这种方法是更为推荐的做法。

祝您编码顺利~

5

我简单测试了一下，2#中最后给出的在host端对count变量清零的方法似乎是可行的，代码和结果如下：


#include "cuda_runtime.h"

#include <stdio.h>

__device__ int count=0;
int* count_ptr=NULL;

__global__ void Kernel()
{
   int tid=blockIdx.x*blockDim.x+threadIdx.x;
	atomicAdd(&count,1);
   __syncthreads();
   if (tid==0)
   {
	   printf("count=%d   \n",count);
   }
   return;
}

void test()
{
	cudaGetSymbolAddress((void**)&count_ptr,"count");
	cudaMemset(count_ptr,0,sizeof(int));
	Kernel<<<1,128>>>();
	return;
}


int main()
{
	for (int i=0;i<3;++i)
	{
		test();
	}

	cudaDeviceReset();

	return 0;
	
}

[attach]2889[/attach]

可见每次循环都是先将count清零之后的结果。

供LZ参考，祝LZ编码愉快~

6

我重新检查的相关的资料，发现2#的一个问题，可能会影响到您，深表歉意。
具体情况如下：

关于cudaGetSymbolAddress()这个函数的第二个参数，在部分版本的CUDA上，可以是变量或者变量名的字符串，而在最新的CUDA5中已经去除了的字符串用法的支持。（同时未知来源的传言表示，之前的版本也有出现使用字符串造成失败的情况）。

根据CUDA 5手册的说法，字符串的用法在CUDA 4.1开始不推荐使用，并在CUDA 5里面移除。
[attach]2891[/attach]

所以，您可以将2#代码中cudaGetSymbolAddress()函数的第二个参数，从"count"改为count试试看。

上次回答考虑不周尽请见谅，祝您编码愉快~