how to cast thrust::device_vector<int> to function with raw pointer

Question

 __global__ void HYPER (int tFast, int tFastLenth, int kilo, int lenPrzvFast, double eps, int AF,double *arrINTLighFast, int *arrPrzvFncFst, int dv_ptr) 
     {
  for(int j = 0;j<(tFast*tFastLenth);j++)
      {  arrINTLighFast[j]=0;
      } 
          for(int j = 0;j<(kilo);j++) arrPrzvFncFst[j]=0;
   for(int j = 1;j<(tFast*tFastLenth);j++)
       { arrINTLighFast[j]=  arrINTLighFast[j-1] + AF*exp(-j/(eps+tFast) ); }

   for(int j = 0;j<(tFast*tFastLenth-1);j++)
         {
            for(int i=(arrINTLighFast[j]);i< (arrINTLighFast[j+1]);i++)
                {arrPrzvFncFst[i]=j;}
         }
   for(int j = 0;j<lenPrzvFast;j++)
        { devvecPrzvFncFst61440Slow149998[j]= arrPrzvFncFst[j] ;}
 }


int main (void)
{
const int tFast = 9;
const int tFastLenth = 6;
double arrINTLighFast[tFast*tFastLenth];
int arrPrzvFncFst[61500];
int AF = 1000;
int kilo = 1024;
int kilo150 = 149998;
const double   eps=0.0000001;
const int lenPrzvFast=61500;

    thrust::host_vector<int> vecPrzvFncFst61440Slow149998;
    int Len_vecPrzv=( lenPrzvFast+kilo150);       
    for (int j=0;j<Len_vecPrzv;j++)   vecPrzvFncFst61440Slow149998.push_back(0);
    for (int j=0;j<Len_vecPrzv;j++)     vecPrzvFncFst61440Slow149998 [j] = 0;
    thrust::device_vector<int> devvecPrzvFncFst61440Slow149998 = vecPrzvFncFst61440Slow149998;

    int *dv_ptr = thrust::raw_pointer_cast(devvecPrzvFncFst61440Slow149998.data());

    HYPER <<<blocks, threads>>>(tFast, tFastLenth, kilo, lenPrzvFast, eps, AF, arrINTLighFast, arrPrzvFncFst, dv_ptr);

   thrust::host_vector<int> HostvecPrzvFncFst61440Slow149998 = devvecPrzvFncFst61440Slow149998;
    std::cout << "Device vector is: " << std::endl;
    for(int j = 0; j<vecPrzvFncFst61440Slow149998.size(); j++) 
            std::cout << "vecPrzvFncFst61440Slow149998[" << j << "] = " << HostvecPrzvFncFst61440Slow149998[j] << std::endl;
 return 0;
 }

There is a problem I cannot use vectors in function, so I decided to use thrust::raw_pointer_cast. However, I have problems: during compiling I have an error : identifier "devvecPrzvFncFst61440Slow149998" is undefined. The second one is I cannot defenitely find out how to pass int *dv_ptr to both function and prototype, there is an error: argument of type "int *" is incompatible with parameter of type "int". I looked among the internet but there is no solution how successfully solve the problems I've mentioned above

Thank you for your time

Answer 1

Your kernel function HYPER has no defined parameter of devvecPrzvFncFst61440Slow149998, so when you try to use that here:

for(int j = 0;j<lenPrzvFast;j++)
    { devvecPrzvFncFst61440Slow149998[j]= arrPrzvFncFst[j] ;}

you're going to get the undefined identifier error. There's no magic here, your CUDA kernel mostly has to comply with the rules of an ordinary C function. If you want to use a variable, it had better be listed in the function parameters (excepting global scope variables and built-in variables, which this isn't).

The other problem you mention is due to the fact that dv_ptr is a pointer type:

int *dv_ptr = thrust::raw_pointer_cast(devvecPrzvFncFst61440Slow149998.data());

but you are attempting to pass it in a kernel parameter position:

HYPER <<<blocks, threads>>>(..., dv_ptr);
                                 ^^^^^^

that is expecting an ordinary (non-pointer) type:

__global__ void HYPER (..., int dv_ptr) 
                            ^^^^^^^^^^

The following code has those issues fixed and compiles cleanly for me:

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

#define blocks 1
#define threads 1

 __global__ void HYPER (int tFast, int tFastLenth, int kilo, int lenPrzvFast, double eps, int AF,double *arrINTLighFast, int *arrPrzvFncFst, int *dv_ptr)
     {
  for(int j = 0;j<(tFast*tFastLenth);j++)
      {  arrINTLighFast[j]=0;
      }
          for(int j = 0;j<(kilo);j++) arrPrzvFncFst[j]=0;
   for(int j = 1;j<(tFast*tFastLenth);j++)
       { arrINTLighFast[j]=  arrINTLighFast[j-1] + AF*exp(-j/(eps+tFast) ); }

   for(int j = 0;j<(tFast*tFastLenth-1);j++)
         {
            for(int i=(arrINTLighFast[j]);i< (arrINTLighFast[j+1]);i++)
                {arrPrzvFncFst[i]=j;}
         }
   for(int j = 0;j<lenPrzvFast;j++)
        { dv_ptr[j]= arrPrzvFncFst[j] ;}
 }


int main (void)
{
const int tFast = 9;
const int tFastLenth = 6;
double arrINTLighFast[tFast*tFastLenth];
int arrPrzvFncFst[61500];
int AF = 1000;
int kilo = 1024;
int kilo150 = 149998;
const double   eps=0.0000001;
const int lenPrzvFast=61500;

    thrust::host_vector<int> vecPrzvFncFst61440Slow149998;
    int Len_vecPrzv=( lenPrzvFast+kilo150);
    for (int j=0;j<Len_vecPrzv;j++)   vecPrzvFncFst61440Slow149998.push_back(0);
    for (int j=0;j<Len_vecPrzv;j++)     vecPrzvFncFst61440Slow149998 [j] = 0;
    thrust::device_vector<int> devvecPrzvFncFst61440Slow149998 = vecPrzvFncFst61440Slow149998;

    int *dv_ptr = thrust::raw_pointer_cast(devvecPrzvFncFst61440Slow149998.data());

    HYPER <<<blocks, threads>>>(tFast, tFastLenth, kilo, lenPrzvFast, eps, AF, arrINTLighFast, arrPrzvFncFst, dv_ptr);

   thrust::host_vector<int> HostvecPrzvFncFst61440Slow149998 = devvecPrzvFncFst61440Slow149998;
    std::cout << "Device vector is: " << std::endl;
    for(int j = 0; j<vecPrzvFncFst61440Slow149998.size(); j++)
            std::cout << "vecPrzvFncFst61440Slow149998[" << j << "] = " << HostvecPrzvFncFst61440Slow149998[j] << std::endl;
 return 0;
 }