Cuda "invalid argument" 2d array - Сотовые автоматы

Question

Я пытаюсь рассчитать перераспределение клеточных автоматов 2d с использованием Cuda. Я совершенно новичок в этом, поэтому не знаю, что я делаю неправильно. Я пробовал много решений, которые я видел здесь, но все они дают «недопустимый аргумент», когда я вызываю ядро.

Вот упрощенная версия ядра:

//kernel definition 
__global__ void stepCalc(float B[51][51], int L, int flag, float m, float en) 
{ 
    int i = blockDim.x * blockIdx.x + threadIdx.x; 
    int j = blockDim.y * blockIdx.y + threadIdx.y; 

    float g=B[i][j]-0.25*(B[i+1][j]+B[i-1][j]+B[i][j+1]+B[i][j-1]); 
    flag = 0; 

    if (i < L-2 && j < L-2 && i>2 && j>2 && abs(g)>m) 
    { 
      flag = 1; 
      en+=-16*g*g+8*B[i][j]*abs(g); 
      B[i][j]+=-4*f*g; 
      B[i+1][j]+=f*g; 
      B[i-1][j]+=f*g; 
      B[i][j+1]+=f*g; 
      B[i][j-1]+=f*g; 
    } 
} 

Основная функция выглядит следующим образом:

#define L 50 
float B[L+1][L+1]; 
//initialize B[i][j] 

float g=0; 
int flag = 1; 
float m=0.1; 
float en = 0; 
while (flag==1) 
{ 
    float (*dB)[L+1]; 
    int *dFlag=NULL; 
    float *dEn=NULL; 

    cudaMalloc((void **)&dFlag,sizeof(int)); 
    cudaMalloc((void **)&dEn,sizeof(float)); 
    cudaMalloc((void **)&dB, ((L+1)*(L+1))*sizeof(float)); 

    cudaMemcpy(dB, B, sizeB, cudaMemcpyHostToDevice); 
    cudaMemcpy(dFlag, &flag, sizeof(int), cudaMemcpyDeviceToHost); 
    cudaMemcpy(dEn, &en, sizeof(float), cudaMemcpyDeviceToHost); 

    dim3 threadsPerBlock(16,16); 
    dim3 numBlocks((L+1)/threadsPerBlock.x,(L+1)/threadsPerBlock.y); 

    stepCalc<<<numBlocks, threadsPerBlock>>>(dB, L, dflag, m, dEn); 
    GPUerrchk(cudaPeekAtLastError()); //gives "invalid argument" at this line 

    cudaMemcpy(B, (dB), sizeB, cudaMemcpyDeviceToHost); 
    cudaMemcpy(&flag, dFlag, sizeof(int), cudaMemcpyDeviceToHost); 
    cudaMemcpy(&en, dEn, sizeof(float), cudaMemcpyDeviceToHost); 

    cudaFree(dB); 
    cudaFree(dFlag); 
    cudaFree(dEn); 
} 

Мне нужно извлечь новый массив B, значение флага и сумму " en 'по всем темам. Я даже близко к тому, как должно выглядеть решение? Возможно ли это? Я также попытался сделать массив хостов B как float ** B без везения.

Answer 1

Существуют различные проблемы с кодом.

1. Вы можете быть с видом на разницу между пропусканием значение к ядру и передавая указатель:

   __global__ void stepCalc(float B[51][51], int L, int flag, float m, float en) 
             ^     ^
              |      | 
             a pointer    a value 
   

   мы вернемся к B в данный момент, но для значений например flag и en, передавая эти по значению ядру, имеет аналогичные последствия для передачи по значению функции C. Это односторонний путь связи. Поскольку из вашего кода видно, что вы хотите использовать эти значения, измененные ядром позже в главном коде, вместо этого вам нужно будет передать указатели. В некоторых случаях вы уже указали указатели для этой цели, поэтому у вас есть дополнительный тип ошибок, в котором в некоторых случаях (dFlag) вы передаете указатель , тогда как определение ядра ожидает значение.

2. Что касается B, передавая 2D массив от хоста к устройству может быть более сложным, чем вы могли бы первоначально ожидать, в связи с проблемой глубокой копии. Не охватывая всю эту область здесь, найдите в «2D-массиве CUDA 2D» в верхнем правом углу этой страницы, и вы получите много информации об этом и различные способы борьбы с ним. Поскольку вы, похоже, готовы рассмотреть массив фиксированной ширины (известный во время компиляции), мы можем упростить обработку 2D-массива, используя компилятор, чтобы помочь нам с конкретным typedef.

3. Если у вас возникли проблемы с кодом cuda, рекомендуется соблюдать строгую проверку ошибок CUDA во всем коде, а не только в одном месте. Одной из причин этого является то, что ошибки CUDA, возникающие в определенном месте, часто возвращаются в любом последующем месте в коде. Это делает его запутанным, если вы не проверяете каждый вызов API CUDA, так как конкретная ошибка «неправильного аргумента» может быть не вызвана самим ядром, а некоторым вызовом API, который произошел ранее.

4. Вы, как правило, не хотите cudaMalloc операции в обработке данных while петля. Обычно это операции, которые вы делаете один раз, в начале вашего кода. Выполнение cudaMalloc на каждой итерации цикла while имеет несколько отрицательных проблем, один из которых заключается в том, что у вас закончится память (хотя у вас есть cudaFree заявления, а, возможно, нет), в конце концов, и вы эффективно отбрасываете свои данные на каждая итерация. Кроме того, это негативно повлияет на вашу производительность.

5. У вас есть некоторые из ваших направлений передачи cudaMemcpy не так, как здесь:

   cudaMemcpy(dFlag, &flag, sizeof(int), cudaMemcpyDeviceToHost); 
   

6. Установка flag до нуля в коде ядра будет проблематично. Варпс может выполняться в любом порядке, и после того, как некоторые деформации уже установили flag в 1 позже в ядре, другие деформации могут начать выполнение и установить flag на ноль снова. Это, вероятно, не то, что вы хотите. Одно из возможных исправлений заключается в том, чтобы установить flag в ноль перед выполнением ядра (т. Е. В главном коде и скопировать его на устройство).

7. Ваше ядро ​​будет генерировать недоступные индексации здесь:

   float g=B[i][j]-0.25*(B[i+1][j]+B[i-1][j]+B[i][j+1]+B[i][j-1]); 
   

   (просто спросите себя, что происходит, когда я = 0 и у = 0). Исправить это - переместить эту строку кода внутри проверки if, которую вы имеете для проверки границ сразу после нее.

8. Ваше ядро ​​использует переменную f которая не определяется нигде, что я могу видеть, например, здесь:

   B[i+1][j]+=f*g; 
   

Следующий код моя попытка переделать код, создать полный пример, и удалить вышеуказанные проблемы. Он не делает ничего полезного, но он компилируется без ошибок и работает без ошибок для меня. Я не представил никаких данных, поэтому на данный момент это просто доказательство концепции. Я уверен, что он все еще содержит ошибки обработки данных.

#include <stdio.h> 

#define my_L 50 

typedef float farray[my_L+1]; 

#define cudaCheckErrors(msg) \ 
    do { \ 
     cudaError_t __err = cudaGetLastError(); \ 
     if (__err != cudaSuccess) { \ 
      fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \ 
       msg, cudaGetErrorString(__err), \ 
       __FILE__, __LINE__); \ 
      fprintf(stderr, "*** FAILED - ABORTING\n"); \ 
      exit(1); \ 
     } \ 
    } while (0) 

//kernel definition 
__global__ void stepCalc(farray B[], int L, int *flag, float m, float *en) 
{ 
    int i = blockDim.x * blockIdx.x + threadIdx.x; 
    int j = blockDim.y * blockIdx.y + threadIdx.y; 

    //float g=B[i][j]-0.25*(B[i+1][j]+B[i-1][j]+B[i][j+1]+B[i][j-1]); 
    // flag = 0; 
    float f = 1.0f; 
    if (i < L-2 && j < L-2 && i>2 && j>2){ 
     float g=B[i][j]-0.25*(B[i+1][j]+B[i-1][j]+B[i][j+1]+B[i][j-1]); 
     if (abs(g)>m) 
     { 
      *flag = 1; 
      *en+=-16*g*g+8*B[i][j]*abs(g); 
      B[i][j]+=-4*f*g; 
      B[i+1][j]+=f*g; 
      B[i-1][j]+=f*g; 
      B[i][j+1]+=f*g; 
      B[i][j-1]+=f*g; 
     } 
     } 
} 


int main(){ 

    farray B[my_L+1]; 
//initialize B[i][j] 

    farray *dB; 
    int flag = 1; 
    float m=0.1; 
    float en = 0; 
    int *dFlag=NULL; 
    float *dEn=NULL; 

    cudaMalloc((void **)&dFlag,sizeof(int)); 
    cudaCheckErrors("1"); 
    cudaMalloc((void **)&dEn,sizeof(float)); 
    cudaCheckErrors("2"); 
    size_t sizeB = (my_L+1)*sizeof(farray); 
    cudaMalloc((void **)&dB, sizeB); 
    cudaCheckErrors("3"); 
    cudaMemcpy(dB, B, sizeB, cudaMemcpyHostToDevice); 
    cudaCheckErrors("4"); 
    cudaMemcpy(dEn, &en, sizeof(float), cudaMemcpyHostToDevice); 
    cudaCheckErrors("5"); 

    dim3 threadsPerBlock(16,16); 
    dim3 numBlocks((my_L+1)/threadsPerBlock.x,(my_L+1)/threadsPerBlock.y); 
    while (flag==1) 
    { 
    flag = 0; 
    cudaMemcpy(dFlag, &flag, sizeof(int), cudaMemcpyHostToDevice); 
    cudaCheckErrors("6"); 
    stepCalc<<<numBlocks, threadsPerBlock>>>(dB, my_L, dFlag, m, dEn); 
    cudaDeviceSynchronize(); 
    cudaCheckErrors("7"); 
    cudaMemcpy(&flag, dFlag, sizeof(int), cudaMemcpyDeviceToHost); 
    cudaCheckErrors("8"); 
    } 
    cudaMemcpy(B, (dB), sizeB, cudaMemcpyDeviceToHost); 
    cudaCheckErrors("9"); 
    cudaMemcpy(&en, dEn, sizeof(float), cudaMemcpyDeviceToHost); 
    cudaCheckErrors("10"); 
// process B 
    cudaFree(dB); 
    cudaFree(dFlag); 
    cudaFree(dEn); 
}