dpct created an accessor of pixel_type after conversion

[Shukla__Gagandeep]

Hi,

I converted cuda code from libSGM. I converted file census_transform.cu and was able to generate census_transform.dp.cpp. 

dpct has converted the following code 

template <typename T>
void enqueue_census_transform(
	feature_type *dest,
	const T *src,
	int width,
	int height,
	int pitch,
	cudaStream_t stream)
{
	const int width_per_block = BLOCK_SIZE - WINDOW_WIDTH + 1;
	const int height_per_block = LINES_PER_BLOCK;
	const dim3 gdim(
		(width  + width_per_block  - 1) / width_per_block,
		(height + height_per_block - 1) / height_per_block);
	const dim3 bdim(BLOCK_SIZE);
	census_transform_kernel<<<gdim, bdim, 0, stream>>>(dest, src, width, height, pitch);
}

into

template <typename T>
void enqueue_census_transform(feature_type *dest, const T *src, int width,
                              int height, int pitch, sycl::queue *stream)
{
	const int width_per_block = BLOCK_SIZE - WINDOW_WIDTH + 1;
	const int height_per_block = LINES_PER_BLOCK;
  const sycl::range<3> gdim((width + width_per_block - 1) / width_per_block,
                            (height + height_per_block - 1) / height_per_block,
                            1);
  const sycl::range<3> bdim(BLOCK_SIZE, 1, 1);
  stream->submit([&](sycl::handler &cgh) {
    sycl::range<2> smem_lines_range_ct1(8 /*SMEM_BUFFER_SIZE*/,
                                        128 /*BLOCK_SIZE*/);

    sycl::accessor<pixel_type, 2, sycl::access::mode::read_write,
                   sycl::access::target::local>
        smem_lines_acc_ct1(smem_lines_range_ct1, cgh);

    auto dpct_global_range = gdim * bdim;

    cgh.parallel_for(
        sycl::nd_range<3>(
            sycl::range<3>(dpct_global_range.get(2), dpct_global_range.get(1),
                           dpct_global_range.get(0)),
            sycl::range<3>(bdim.get(2), bdim.get(1), bdim.get(0))),
        [=](sycl::nd_item<3> item_ct1) {
          census_transform_kernel(
              dest, src, width, height, pitch, item_ct1,
              dpct::accessor<pixel_type, dpct::local, 2>(smem_lines_acc_ct1,
                                                         smem_lines_range_ct1));
        });
  });
}

The issue is when I try to compile it using dpcpp, it complains about symbol pixel_type. And I don't understand what this pixel_type is. 

Here is the error:

c:\Users\...\Downloads\cuda code\libSGM-master\src\dpct>dpcpp -c -I../../include census_transform.dp.cpp
census_transform.dp.cpp:124:20: error: use of undeclared identifier 'pixel_type'
    sycl::accessor<pixel_type, 2, sycl::access::mode::read_write,
                   ^
census_transform.dp.cpp:163:16: note: in instantiation of member function 'sgm::CensusTransform<unsigned char>::enqueue'
      requested here
template class CensusTransform<uint8_t>;
               ^
census_transform.dp.cpp:126:9: error: C++ requires a type specifier for all declarations
        smem_lines_acc_ct1(smem_lines_range_ct1, cgh);
        ^
census_transform.dp.cpp:138:30: error: use of undeclared identifier 'pixel_type'
              dpct::accessor<pixel_type, dpct::local, 2>(smem_lines_acc_ct1,
                             ^
3 errors generated.

Can you please help me find out what this pixel_type is that dpct converted original code to?

Regards,
Gagan

[RahulV_intel]

Hi Gagan,

"pixel_type" is also present in the native cuda code census_transform.cu. 

This is nothing but an alias to Template typename "T". Instead of using T(generic datatype), the code uses pixel_type to improve readability.

I'd suggest you to make the following changes to your migrated dpcpp application:

Inside the "void enqueue_census_transform()"(line no. 117) definition, above sycl::accessor() statement (line no. 138 in my file. Could be different in your case). 

Add this line:

 using pixel_type = T;

On line no. 37, change the parameter of void census_transform_kernel() function to:

dpct::accessor<T, dpct::local, 2> smem_lines) //Note the T here instead of pixel_type
//from dpct::accessor<pixel_type, dpct::local, 2> smem_lines)

I'm also attaching the modified migrated code for reference. Compiling this migrated code will lead to Undefined reference to main function due to the absence of main function.(Same thing happens if you try to compile single cuda source code census_transform.cu)

--Rahul

[Shukla__Gagandeep]

Thanks Rahul.

I made modifications to the other instance of pixel_type but did not think that I could just use the same thing as previous function (using pixel_type = T) in enqueue_census_transform function too.

Regards,
Gagan

[RahulV_intel]

Hi,

Good to know that it worked.

Since your issue is resolved, let me know if I can close this thread.

--Rahul

[RahulV_intel]

Hi Gagan,

I'm closing this thread since your issue is resolved.

--Rahul