unknown type name 'T'

[zjin-lcf]

Compiling the migrated program shows a few errors. Thanks.

main.dp.cpp:35:25: error: unknown type name 'T' void someKernel(int *x, T cmem) { ^ main.dp.cpp:36:15: error: 'Foo' does not refer to a value *x = cmem<Foo>.x; ^

#include <cuda.h>
#include <iostream>

#define CUDA_CALL(x) do { \
    cudaError_t err = x; \
    if (err != cudaSuccess) { \
        std::cerr << "Error: " << cudaGetErrorString(err) << std::endl; \
        exit(1); \
    } \
} while (0)

struct Foo {
    int x;
    __host__ __device__ Foo() {}
    __host__ __device__ ~Foo() {}
};

template<typename T>
__constant__ T cmem;

__global__ void someKernel(int *x) {
    *x = cmem<Foo>.x;
}

int main() {
    Foo foo;
    foo.x = 42;
    CUDA_CALL(cudaMemcpyToSymbol(cmem<Foo>, &foo, sizeof(Foo)));

    int *dx;
    CUDA_CALL(cudaMalloc(&dx, sizeof(int)));
    someKernel<<<1, 1>>>(dx);

    int x;
    CUDA_CALL(cudaMemcpy(&x, dx, sizeof(int), cudaMemcpyDeviceToHost));
    if (x == 42) {
        std::cout << "SUCCESS!" << std::endl;
    } else {
        std::cout << "FAILURE!" << std::endl;
    }

    CUDA_CALL(cudaFree(dx));
    return 0;
}

Reference https://github.com/ROCm-Developer-Tools/HIP/issues/3201

[tomflinda]

@zjin-lcf The migrated SYCL code with manual fix is as follows :

#include <sycl/sycl.hpp>
#include <dpct/dpct.hpp>
#include <iostream>

/*
DPCT1001:0: The statement could not be removed.
*/
/*
DPCT1000:1: Error handling if-stmt was detected but could not be rewritten.
*/
/*
DPCT1009:3: SYCL uses exceptions to report errors and does not use the error
codes. The original code was commented out and a warning string was inserted.
You need to rewrite this code.
*/
#define CUDA_CALL(x) do {                                                            \
        dpct::err0 err = x;                                                          \
        if (err != 0) {                                                              \
            std::cerr                                                                \
                << "Error: "                                                         \
                << "cudaGetErrorString is not supported" /*cudaGetErrorString(err)*/ \
                << std::endl;                                                        \
            exit(1);                                                                 \
        }                                                                            \
    } while (0)

struct Foo {
    int x;
    Foo() {}
    ~Foo() {}
};

template <typename T> static dpct::constant_memory<T, 0> cmem;
template<class T>
void someKernel(int *x, T cmem) {
    *x = cmem.x;
}

int main() try {
    dpct::device_ext &dev_ct1 = dpct::get_current_device();
    sycl::queue &q_ct1 = dev_ct1.default_queue();
    Foo foo;
    foo.x = 42;
    /*
    DPCT1003:2: Migrated API does not return error code. (*, 0) is inserted. You
    may need to rewrite this code.
    */
    CUDA_CALL((q_ct1.memcpy(cmem<Foo>.get_ptr(), &foo, sizeof(Foo)).wait(), 0));

    int *dx;
    /*
    DPCT1003:4: Migrated API does not return error code. (*, 0) is inserted. You
    may need to rewrite this code.
    */
    CUDA_CALL((dx = sycl::malloc_device<int>(1, q_ct1), 0));
    q_ct1.submit([&](sycl::handler &cgh) {
        cmem<Foo>.init();

        auto cmem_ptr_ct1 = cmem<Foo>.get_ptr();

        cgh.parallel_for(
            sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
            [=](sycl::nd_item<3> item_ct1) {
                someKernel(dx, *cmem_ptr_ct1);
            });
    });

    int x;
    /*
    DPCT1003:5: Migrated API does not return error code. (*, 0) is inserted. You
    may need to rewrite this code.
    */
    CUDA_CALL((q_ct1.memcpy(&x, dx, sizeof(int)).wait(), 0));
    if (x == 42) {
        std::cout << "SUCCESS!" << std::endl;
    } else {
        std::cout << "FAILURE!" << std::endl;
    }

    /*
    DPCT1003:6: Migrated API does not return error code. (*, 0) is inserted. You
    may need to rewrite this code.
    */
    CUDA_CALL((sycl::free(dx, q_ct1), 0));
    return 0;
}
catch (sycl::exception const &exc) {
  std::cerr << exc.what() << "Exception caught at file:" << __FILE__
            << ", line:" << __LINE__ << std::endl;
  std::exit(1);
}

We need further investigation and more discussion to decide how to fix it.

[tomflinda]

here is the diff between migrated sycl code and migrated sycl code with manual fix:

34c34
<
---
> template<class T>
36c36
<     *x = cmem<Foo>.x;
---
>     *x = cmem.x;
48c48
<     CUDA_CALL((q_ct1.memcpy(cmem.get_ptr(), &foo, sizeof(Foo)).wait(), 0));
---
>     CUDA_CALL((q_ct1.memcpy(cmem<Foo>.get_ptr(), &foo, sizeof(Foo)).wait(), 0));
57c57
<         cmem.init();
---
>         cmem<Foo>.init();
59c59
<         auto cmem_ptr_ct1 = cmem.get_ptr();
---
>         auto cmem_ptr_ct1 = cmem<Foo>.get_ptr();

[zjin-lcf]

Thank you for the WA.