ROCm / HIPIFY

HIPIFY: Convert CUDA to Portable C++ Code
https://rocm.docs.amd.com/projects/HIPIFY/en/latest/
MIT License
503 stars 70 forks source link

[HIPIFY] Get a incorrect replacement when `cudaMemcpyToSymbol(...)` is in a macro #948

Closed renge24 closed 6 months ago

renge24 commented 1 year ago
#include <stdio.h>

#define a0     -3.0124472f
#define a1      1.7383092f
#define a2     -0.2796695f
#define a3      0.0547837f
#define a4     -0.0073118f

#define CHECK(call)                                                            \
{                                                                              \
    const cudaError_t error = call;                                            \
    if (error != cudaSuccess)                                                  \
    {                                                                          \
        fprintf(stderr, "Error: %s:%d, ", __FILE__, __LINE__);                 \
        fprintf(stderr, "code: %d, reason: %s\n", error,                       \
                cudaGetErrorString(error));                                    \
    }                                                                          \
}

__device__ __constant__ float coef[5];

int main() {
    const float h_coef[] = {a0, a1, a2, a3, a4};
    CHECK( cudaMemcpyToSymbol( coef, h_coef, 5 * sizeof(float) ));
}

is hipified to

#include <hip/hip_runtime.h>
#include <stdio.h>

#define a0     -3.0124472f
#define a1      1.7383092f
#define a2     -0.2796695f
#define a3      0.0547837f
#define a4     -0.0073118f

#define CHECK(call)                                                            \
{                                                                              \
    const hipError_t error = call;                                            \
    if (error != hipSuccess)                                                  \
    {                                                                          \
        fprintf(stderr, "Error: %s:%d, ", __FILE__, __LINE__);                 \
        fprintf(stderr, "code: %d, reason: %s\n", error,                       \
                hipGetErrorString(error));                                    \
    }                                                                          \
}

__device__ __constant__ float coef[5];

int main() {
    const float h_coef[] = {a0, a1, a2, a3, a4};
    HIP_SYMBOL(coef)( hipMemcpyToSymbol( coef, h_coef, 5 * sizeof(float) ));
}

with hipify-clang. We can see that the HIP_SYMBOL(...) is added to wrong location.

emankov commented 1 year ago

Thank's for catching this, @link-r!