EnzymeAD / EnzymeAD/Enzyme

error: no libcall available for flog error

Open
#2,775 0 comments 1 reaction 2 assignees Claimed by @minansys View on GitHub
Dominant language
LLVM
Stars
1.7k
Forks
188
Avg merge
1d 22h
Merged PRs (30d)
26

Description

The **windows enzyme built-in clang** (19.1.7) gives the following errors for the **cuda** code
```
error: no libcall available for flog
```
the cuda code is
```
#include
#include

#include
#include

#define REAL double

using f_ptr = void (*)(REAL *, REAL *);

extern __device__ void __enzyme_autodiff(f_ptr,
REAL *, REAL *,
REAL *, REAL *);

__device__ void foo_impl(REAL *x_in, REAL *x_out) {
x_out[0] = x_in[0] * pow(x_in[0], x_in[0]);
// REAL temp = x_in[0];
// x_out[0] = fmin(x_out[0], temp);
}

__global__ void foo(REAL *x_in, REAL *x_out) {
foo_impl(x_in, x_out);
}

__global__ void foo_grad(REAL *x, REAL *d_x, REAL *y, REAL *d_y) {
// __enzyme_autodiff(foo_impl, x, d_x, y, d_y);
}

static void CheckCuda(cudaError_t err, const char *what) {
if (err != cudaSuccess) {
std::fprintf(stderr, "%s failed: %s\n", what, cudaGetErrorString(err));
std::exit(1);
}
}

int main() {
REAL *x = nullptr;
REAL *d_x = nullptr;
REAL *y = nullptr;
REAL *d_y = nullptr;

CheckCuda(cudaMalloc(&x, sizeof(*x)), "cudaMalloc(x)");
CheckCuda(cudaMalloc(&d_x, sizeof(*d_x)), "cudaMalloc(d_x)");
CheckCuda(cudaMalloc(&y, sizeof(*y)), "cudaMalloc(y)");
CheckCuda(cudaMalloc(&d_y, sizeof(*d_y)), "cudaMalloc(d_y)");

REAL host_x = 1.4;
REAL host_d_x = 0.0;
REAL host_y = 0.0;
REAL host_d_y = 1.0;

CheckCuda(cudaMemcpy(x, &host_x, sizeof(*x), cudaMemcpyHostToDevice), "cudaMemcpy(x)");
CheckCuda(cudaMemcpy(d_x, &host_d_x, sizeof(*d_x), cudaMemcpyHostToDevice), "cudaMemcpy(d_x)");
CheckCuda(cudaMemcpy(y, &host_y, sizeof(*y), cudaMemcpyHostToDevice), "cudaMemcpy(y)");
CheckCuda(cudaMemcpy(d_y, &host_d_y, sizeof(*d_y), cudaMemcpyHostToDevice), "cudaMemcpy(d_y)");

void *args[] = { &x, &d_x, &y, &d_y };
CheckCuda(cudaLaunchKernel(reinterpret_cast(foo_grad), dim3(1), dim3(1), args, 0, nullptr), "cudaLaunchKernel(foo_grad)");
CheckCuda(cudaDeviceSynchronize(), "cudaDeviceSynchronize");

CheckCuda(cudaMemcpy(&host_x, x, sizeof(*x), cudaMemcpyDeviceToHost), "cudaMemcpy(host_x)");
CheckCuda(cudaMemcpy(&host_d_x, d_x, sizeof(*d_x), cudaMemcpyDeviceToHost), "cudaMemcpy(host_d_x)");
CheckCuda(cudaMemcpy(&host_y, y, sizeof(*y), cudaMemcpyDeviceToHost), "cudaMemcpy(host_y)");
CheckCuda(cudaMemcpy(&host_d_y, d_y, sizeof(*d_y), cudaMemcpyDeviceToHost), "cudaMemcpy(host_d_y)");

std::printf("%f %f\n", host_x, host_y);
std::printf("%f %f\n", host_d_x, host_d_y);

cudaFree(d_y);
cudaFree(y);
cudaFree(d_x);
cudaFree(x);
return 0;
}
```
The problem can not be reproduced on other combinations, such as linux, cpu. The problem is due to pow gives the derivitive of log, however, it is in the format of flog instead of __nv_log.

Simily, the problem occurs for

- tanh, with error of cosh
- sinh, with error of cosh
- sin, with error of cos
- cos, with error of sin

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.