>
내가 작성한 CUDA 코드의 동작에 대해서는 혼란 스럽다. 나는 __device__ 에 대한 테스트를 작성하는 중입니다   DimmedGridGPU 라는 클래스의 함수 . 이 클래스는 int DIM 에서 템플릿입니다  내가 문제가되는 함수는 입력 값에 가장 가까운 지점에서 그리드의 값을 반환하는 것입니다. x . 단위 테스트 목적 으로이 커널 네임 스페이스를 사용하여 각 __device__ 를 호출합니다.  독립된 기능.

이 코드의 바람직한 동작은 3.0 값을 반환하는 것입니다.   do_get_value(x, grid_) 에서  전화 및 d_target[0] 설정  이 값으로 설정 한 다음 단위 테스트 어설 션을 위해 호스트 측으로 다시 전송하십시오. 커널 전체가 제대로 작동하는 것처럼 보이지만 최종 호스트로 다시 전송하면 cudaErrorInvalidValue 가 수신됩니다.  오류가 발생하여 이유를 이해할 수 없습니다.

다음은 클래스의 구조와 기능을 유지하면서 최소한의 코드 예제입니다 :

#include <cuda_runtime.h>
#include <fstream>
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
     fprintf(stderr,"GPUassert: \"%s\": %s %s %d\n", cudaGetErrorName(code), cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

template <int DIM>
class DimmedGridGPU{
public:
  size_t grid_size_;//total size of grid
  int b_derivatives_;//if derivatives are going to be used
  int b_interpolate_;//if interpolation should be used on the grid
  double* grid_;//the grid values
  double* grid_deriv_;//derivatives    
  double dx_[DIM];//grid spacing
  double min_[DIM];//grid minimum
  double max_[DIM];//maximum
  int grid_number_[DIM];//number of points on grid
  int b_periodic_[DIM];//if a dimension is periodic
  int* d_b_interpolate_;
  int* d_b_derivatives_;

  DimmedGridGPU(const double* min, 
        const double* max, 
        const double* bin_spacing, 
        const int* b_periodic, 
        int b_derivatives, 
        int b_interpolate) :   b_derivatives_(b_derivatives), b_interpolate_(b_interpolate), grid_(NULL), grid_deriv_(NULL){
    size_t i;
    for(i = 0; i < DIM; i++) {
      min_[i] = min[i];
      max_[i] = max[i];
      b_periodic_[i] = b_periodic[i];
      grid_number_[i] = (int) ceil((max_[i] - min_[i]) / bin_spacing[i]);
      dx_[i] = (max_[i] - min_[i]) / grid_number_[i];
      //add one to grid points if 
      grid_number_[i] = b_periodic_[i] ? grid_number_[i] : grid_number_[i] + 1;
      //increment dx to compensate
      if(!b_periodic_[i])
    max_[i] += dx_[i];
    }
    grid_size_ = 1;
    for(i = 0; i < DIM; i++)
      grid_size_ *= grid_number_[i];
    gpuErrchk(cudaMallocManaged(&grid_, grid_size_ * sizeof(double)));
    if(b_derivatives_) {
      gpuErrchk(cudaMallocManaged(&grid_deriv_, DIM * grid_size_ * sizeof(double)));
      if(!grid_deriv_) {
    printf("Out of memory!! gpugrid.cuh:initialize");   
      }
    }
    gpuErrchk(cudaMalloc((void**)&d_b_interpolate_, sizeof(int)));
    gpuErrchk(cudaMemcpy(d_b_interpolate_, &b_interpolate, sizeof(int), cudaMemcpyHostToDevice));
    gpuErrchk(cudaMalloc((void**)&d_b_derivatives_, sizeof(int)));
    gpuErrchk(cudaMemcpy(d_b_derivatives_, &b_derivatives, sizeof(int), cudaMemcpyHostToDevice));
  }
  ~DimmedGridGPU(){
    gpuErrchk(cudaDeviceSynchronize());
    if(grid_ != NULL){
      gpuErrchk(cudaFree(grid_));
      grid_ = NULL;//need to do this so DimmedGrid's destructor functions properly
    }
    if(grid_deriv_ != NULL){
      gpuErrchk(cudaFree(grid_deriv_));
      grid_deriv_ = NULL;
    }
    gpuErrchk(cudaDeviceReset());
  }
//gets the value of the grid closest to x
  __host__ __device__ double do_get_value( double* x, double* grid_) {
    size_t index[DIM];
    get_index(x, index);
    printf("do_get_value was called on the GPU!, and index[0] is now %d\n", index[0]);
    printf("but multi2one(index) gives us %d\n", multi2one(index));
    double value = grid_[multi2one(index)];
    printf("and value to be returned is %f\n", value);
    return value;
  }
//gets grid's 1D index from an array of coordinates
   __host__ __device__ void get_index(const double* x, size_t result[DIM]) const {
    size_t i;
    double xi;
    printf("get_index was called on the GPU in %i dimension(s)\n", DIM);
    for(i = 0; i < DIM; i++) {
      xi = x[i];
      printf("xi is now %f, min_[i] is %f and dx_[i] is %f\n",xi, min_[i], dx_[i]);
      if(b_periodic_[i]){
    xi -= (max_[i] - min_[i]) * gpu_int_floor((xi - min_[i]) / (max_[i] - min_[i]));
      }
      result[i] = (size_t) floor((xi - min_[i]) / dx_[i]);
    }
  }
//takes a multidimensional index to a 1D index
  __host__ __device__ size_t multi2one(const size_t index[DIM]) const {
    size_t result = index[DIM-1];
    size_t i;    
    for(i = DIM - 1; i > 0; i--) {
      result = result * grid_number_[i-1] + index[i-1];
    }
    return result;
  }
};
__host__ __device__ int gpu_int_floor(double number) {
  return (int) number < 0.0 ? -ceil(fabs(number)) : floor(number);
}

namespace kernels{
  template <int DIM>
  __global__ void get_value_kernel(double* x, double* target_arr, double* grid_, DimmedGridGPU<DIM>  g){
    target_arr[0] = g.do_get_value(x, grid_);
    printf("get_value_kernel has set target[0] to be %f\n", target_arr[0]);//check if the value is set correctly
    return;
  }
}

int main(){
  using namespace kernels;
  double min[] = {0};
  double max[] = {10};
  double bin_spacing[] = {1};
  int periodic[] = {0};
  DimmedGridGPU<1> g (min, max, bin_spacing, periodic, 0, 0);
  for(int i = 0; i < 11; i++){
    g.grid_[i] = i;
    printf("g.grid_[%d] is now %f\n", i, g.grid_[i]);
  }
  gpuErrchk(cudaDeviceSynchronize());
  double x[] = {3.5};
  double* d_x;
  gpuErrchk(cudaMalloc(&d_x, sizeof(double)));
  gpuErrchk(cudaMemcpy(d_x, x, sizeof(double), cudaMemcpyHostToDevice));
  double target[] = {5.0};
  double* d_target;
  gpuErrchk(cudaMalloc((void**)&d_target, sizeof(double)));
  gpuErrchk(cudaMemcpy(d_target, target, sizeof(double), cudaMemcpyHostToDevice));
  gpuErrchk(cudaDeviceSynchronize());
  get_value_kernel<1><<<1,1>>>(d_x, d_target, g.grid_, g);
  gpuErrchk(cudaDeviceSynchronize());
  gpuErrchk(cudaMemcpy(target, d_target, sizeof(double), cudaMemcpyDeviceToHost));
  printf("and after GPU stuff, target[0] is now %f\n", target[0]);
  return(0);
}

그래서이 줄은 왜 마지막 줄인가? ) 오류 ' cudaMemcpy ", 포함 된 인쇄 설명을 보면 장치에서 올바른 값을 사용하고 있으며 CudaErrorInvalidValue 에서 반환 한 값을 명확하게 보여줍니다.  전화가 맞습니까?

이미 do_get_value(x, grid_) 를 사용해 보았습니다. , 어쩌면 과제가 어쨌든 값을 전달하고 변경하는 대신 기호를 생성한다고 생각했지만 cudaMemcpyFromSymbol 와 같은 경우는 아닙니다.  유효한 기호가 아닙니다.

내 코드의 샘플 출력은 다음과 같습니다.

와이즈 비즈 와이즈 비즈

     

d_target

     

     

g.grid_[0] is now 0.000000

     

g.grid_[1] is now 1.000000

     

g.grid_[2] is now 2.000000

     

g.grid_[3] is now 3.000000

     

g.grid_[4] is now 4.000000

     

g.grid_[5] is now 5.000000

     

g.grid_[6] is now 6.000000

     

g.grid_[7] is now 7.000000

     

g.grid_[8] is now 8.000000

     

g.grid_[9] is now 9.000000

     

g.grid_[10] is now 10.000000

     

get_index was called on the GPU in 1 dimension(s)

     

xi is now 3.500000, min_[i] is 0.000000 and dx_[i] is 1.000000

     

do_get_value was called on the GPU!, and index[0] is now 3

     

but multi2one(index) gives us 3

and value to be returned is 3.000000
  • 답변 # 1

    와이즈 비즈

    이 문제는 소멸자와 관련이 있습니다 :

    So, why does this line (the last cudaMemcpy) throw an error "CudaErrorInvalidValue"...?

    예상치 못한 장소에서 소멸자가 전화를 받고 있습니다. 이것을 스스로 확신시키기 위해, ~DimmedGridGPU(){ 를 추가하십시오  소멸자에게 진술하십시오. 출력물에 표시되는 위치를 확인하십시오.

    printf
    
    

    그러면 $ ./t955 g.grid_[0] is now 0.000000 g.grid_[1] is now 1.000000 g.grid_[2] is now 2.000000 g.grid_[3] is now 3.000000 g.grid_[4] is now 4.000000 g.grid_[5] is now 5.000000 g.grid_[6] is now 6.000000 g.grid_[7] is now 7.000000 g.grid_[8] is now 8.000000 g.grid_[9] is now 9.000000 g.grid_[10] is now 10.000000 Destructor! get_index was called on the GPU in 1 dimension(s) xi is now 3.500000, min_[i] is 0.000000 and dx_[i] is 1.000000 do_get_value was called on the GPU!, and index[0] is now 3 but multi2one(index) gives us 3 and value to be returned is 3.000000 get_value_kernel has set target[0] to be 3.000000 GPUassert: "cudaErrorInvalidValue": invalid argument t955.cu 167 를 호출하는 것이 꽤 분명합니다.  그 파괴자에서 이제 나쁜 생각처럼 보입니다. 와이즈 비즈  모든 장치 할당을 지 웁니다. 이렇게하면 다음과 같이 시도합니다.

    cudaDeviceReset()
    
    

    cudaDeviceReset()  더 이상 장치에서 유효한 할당이 아니므로 gpuErrchk(cudaMemcpy(target, d_target, sizeof(double), cudaMemcpyDeviceToHost)); 의 장치 대상으로 사용하려고 할 때 , 런타임은이 포인터 값 (장치 재설정으로 변경되지 않음)을 확인하고 포인터 값이 더 이상 유효한 할당에 해당하지 않는지 확인하고 런타임 오류를 발생시킵니다.

    C ++에서와 같이 객체를 값으로 전달 매개 변수로 함수 (또는이 경우 커널)에 전달할 때 해당 객체의 복사 생성자가 호출됩니다. 객체 복사가 범위를 벗어난 이유는 그에 대한 소멸자가 호출됩니다.

    전세계 범위에 영향을 미치는 d_target 와 같은 기능을 수행하는 것이 좋습니다.  객체 소멸자에서 연약한 프로그래밍 패러다임 일 수도 있지만 아마도 의견의 문제 일 수 있습니다. 문제 해결에 필요한 정보가 충분하다고 가정합니다.

    다음 가능한 질문을 피하려면 cudaMemcpy 에 대한 호출을 주석으로 처리하십시오.  소멸자에서 모든 문제가 사라지기에 충분하지 않을 수도 있습니다 (이 특정 문제가 발생할지라도). 이 프로그램의 일반적인 실행에서이 소멸자가 적어도두 번호출되고 있음을 알았으므로, 그 소멸자에서 다른 일이 무엇인지 신중하게 생각하고 더 많은 것을 제거하고 싶을 것입니다 또는 수업을 모두 다시 구성해야합니다.

    예를 들어, cudaDeviceReset()  이런 식으로 사용되는 객체의 소멸자에서 문제를 일으킬 수있는 유일한 기능은 아닙니다. 마찬가지로, cudaDeviceReset()  객체 복사에서 호출 된 소멸자에서 사용될 때 원래 객체에 의도 치 않은 결과를 초래할 수 있습니다.

    cudaDeviceReset()

  • 이전 포트란 입/출력에서 ​​*의 목적은 무엇입니까
  • 다음 iphone - "변경 불가능한 객체로 전송 된 변경 메소드"오류가 발생합니다