Передача в __global__ функцию массива

К примеру есть такой код:

#include <cuda.h>                                                           
#include <iostream>                                                         
struct Matrix {
  int mat[9];                                                               
};                                                                          
__global__ void kernel(Matrix mat, int *out) {                              
  for (int i = 0; i < 9; ++i) {
    *out += mat.mat[i];
  }
}
int main() {
  int *out = nullptr;                                                       
  cudaMalloc(&out, sizeof(int));                                            
  Matrix matrix{};                                                          
  for (int i = 0; i < 9; ++i) {                                             
    matrix.mat[i] = 1;                                                      
  }
   kernel<<<1, 1>>>(matrix, out);                                            
   int retval;                                                               
   cudaMemcpy(&retval, out, sizeof(int), cudaMemcpyDeviceToHost);            
   std::cout << retval << std::endl;                                         
   return EXIT_SUCCESS;                                                      
 }

При этом код работает и выводит верный результат, в связи с чем возникает вопрос: насколько он корректен? С одной стороны мы можем разыменовывать указатели на cpu память в cuda функциях, однако массив на стэке иммет другой тип (хотя и автоматически приводится к указателю).