cuda中的强制类型转换中出现了我无法理解的问题,请各位[名词6]指点迷津

  • 1 replies
  • 499 views
我在cuda实践中想使用共享内存来定制一级缓存,但是想缓存的数据不是同一种数据类型,所有就想统一用void*类型来存,所以就有了下面这个demo,但是并没有按我的预期工作。
代码如下:

#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <device_functions.h>

#include <stdio.h>

#ifdef __CUDACC__
#define CUDA_CALLABLE_MEMBER __host__ __device__
#else
#define CUDA_CALLABLE_MEMBER
#endif

void handleError(int err, const char*file, int line);
#define HANDLE_ERROR(err) (handleError(err,__FILE__,__LINE__))

void handleError(int err, const char*file, int line) {
   if (err != cudaSuccess) {
      printf("\nRuntime Api error:%s !\n in %s at line %d\n", cudaGetErrorString((cudaError_t)err), file, line);
      cudaDeviceReset();
      system("pause");
      exit(EXIT_FAILURE);
   }
}

class CA
{
public:
   CUDA_CALLABLE_MEMBER CA() { m = 3; };
   CUDA_CALLABLE_MEMBER ~CA() {};

public:
   CUDA_CALLABLE_MEMBER int getNum() { return 6; }
   CUDA_CALLABLE_MEMBER int getNum2() { return m; }
protected:
   int m;
};



__global__ void Kernel(CA ** mCa) {
   const unsigned int tid = (blockIdx.x*blockDim.x) + threadIdx.x;

   if (threadIdx.x == 0) {
      *mCa = new CA();
   }

   __shared__ void* sIn[3];

   __shared__ int sInInt[2];

   if (threadIdx.x == 0) {
      sIn[0] = *mCa;
      sIn[1] = (void*)(*mCa)->getNum();
      sIn[2] = (void*)(*mCa)->getNum2();

      sInInt[0] = (*mCa)->getNum();
      sInInt[1] = (*mCa)->getNum2();
   }

   __syncthreads();

   printf("A==>Id=%ld num1=%ld num2=%ld\n", tid, sInInt[0], sInInt[1]);
   printf("C==>Id=%ld num1=%ld num2= %ld \n", tid, sIn[1], sIn[2]);
}

int main()
{
   CA ** caPtr;
   printf("sizeof(void*) = %d\n", sizeof(CA*));
   
   cudaMalloc((void**)&caPtr, sizeof(CA *));
   printf("==>ptr=0x%016I64x\n", caPtr);

   Kernel << <2, 2 >> > (caPtr);
   HANDLE_ERROR(cudaDeviceSynchronize());
   HANDLE_ERROR(cudaGetLastError());
   system("pause");
    return 0;
}

运行的结果如下:
sizeof(void*) = 8
==>ptr=0x0000000b01600000
A==>Id=0 num1=6 num2=3
A==>Id=1 num1=6 num2=3
A==>Id=2 num1=6 num2=3
A==>Id=3 num1=6 num2=3
C==>Id=0 num1=6 num2= 6
C==>Id=1 num1=6 num2= 6
C==>Id=2 num1=6 num2= 6
C==>Id=3 num1=6 num2= 6

在使用__shared__ int sInInt[2]时,运行结果达到了我的预期,但是在__shared__ void* sIn[3]同时缓存CA对象指针与两个数据时,结果就出现了无法理解的问题。
请各位前辈指点迷津。

我在cuda实践中想使用共享内存来定制一级缓存,但是想缓存的数据不是同一种数据类型,所有就想统一用void*类型来存,所以就有了下面这个demo,但是并没有按我的预期工作。
代码如下:

#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <device_functions.h>

#include <stdio.h>

#ifdef __CUDACC__
#define CUDA_CALLABLE_MEMBER __host__ __device__
#else
#define CUDA_CALLABLE_MEMBER
#endif

void handleError(int err, const char*file, int line);
#define HANDLE_ERROR(err) (handleError(err,__FILE__,__LINE__))

void handleError(int err, const char*file, int line) {
   if (err != cudaSuccess) {
      printf("\nRuntime Api error:%s !\n in %s at line %d\n", cudaGetErrorString((cudaError_t)err), file, line);
      cudaDeviceReset();
      system("pause");
      exit(EXIT_FAILURE);
   }
}

class CA
{
public:
   CUDA_CALLABLE_MEMBER CA() { m = 3; };
   CUDA_CALLABLE_MEMBER ~CA() {};

public:
   CUDA_CALLABLE_MEMBER int getNum() { return 6; }
   CUDA_CALLABLE_MEMBER int getNum2() { return m; }
protected:
   int m;
};



__global__ void Kernel(CA ** mCa) {
   const unsigned int tid = (blockIdx.x*blockDim.x) + threadIdx.x;

   if (threadIdx.x == 0) {
      *mCa = new CA();
   }

   __shared__ void* sIn[3];

   __shared__ int sInInt[2];

   if (threadIdx.x == 0) {
      sIn[0] = *mCa;
      sIn[1] = (void*)(*mCa)->getNum();
      sIn[2] = (void*)(*mCa)->getNum2();

      sInInt[0] = (*mCa)->getNum();
      sInInt[1] = (*mCa)->getNum2();
   }

   __syncthreads();

   printf("A==>Id=%ld num1=%ld num2=%ld\n", tid, sInInt[0], sInInt[1]);
   printf("C==>Id=%ld num1=%ld num2= %ld \n", tid, sIn[1], sIn[2]);
}

int main()
{
   CA ** caPtr;
   printf("sizeof(void*) = %d\n", sizeof(CA*));
   
   cudaMalloc((void**)&caPtr, sizeof(CA *));
   printf("==>ptr=0x%016I64x\n", caPtr);

   Kernel << <2, 2 >> > (caPtr);
   HANDLE_ERROR(cudaDeviceSynchronize());
   HANDLE_ERROR(cudaGetLastError());
   system("pause");
    return 0;
}

运行的结果如下:
sizeof(void*) = 8
==>ptr=0x0000000b01600000
A==>Id=0 num1=6 num2=3
A==>Id=1 num1=6 num2=3
A==>Id=2 num1=6 num2=3
A==>Id=3 num1=6 num2=3
C==>Id=0 num1=6 num2= 6
C==>Id=1 num1=6 num2= 6
C==>Id=2 num1=6 num2= 6
C==>Id=3 num1=6 num2= 6

在使用__shared__ int sInInt[2]时,运行结果达到了我的预期,但是在__shared__ void* sIn[3]同时缓存CA对象指针与两个数据时,结果就出现了无法理解的问题。
请各位前辈指点迷津。

楼主你好,你的代码存在多处BUG,包括但不限于:

(1)多个线程写入同一个地址,并使用。你的代码中,按照语义,mCa指针所指向的对象指针,可能是各个block中的线程0,所创建的2个不同指针(如果编译器优化了,没有重读的话);也可能是具体其中的某一个(如果编译器重读的话);甚至还可能是任何一个都不是(编译器在两次写入后进行了重读,同时平台没有约定8B类型(指针)的读写是原子性的),前一个的前4B和后一个的后4B的组合,类似这种的。这样是非常不安全的。算是一个潜伏的BUG(不一定会触发)。

(2)你使用同时用%ld, 对指针,和int进行printf输出,这是不安全的。没有人能约定%ld(long int)一定大小等于指针大小,也没有人约定int一定等于long int, 我举个例子,如果你在*nix下,64-bit,则此时的确long int等于指针大小,都是8B,但是int此时只有4B;再举个例子,如果你在Windows下,64-bit,则此时long int只有4B,对应int也是4B倒是可以,但是指针却又不匹配了。也就是说,无论在Windows还是*nix下,64-bit编译,总会至少有一个不对的。这里显然还是个BUG。