rocm / hipify Goto Github PK
View Code? Open in Web Editor NEWHIPIFY: Convert CUDA to Portable C++ Code
Home Page: https://rocm.docs.amd.com/projects/HIPIFY/en/latest/
License: MIT License
HIPIFY: Convert CUDA to Portable C++ Code
Home Page: https://rocm.docs.amd.com/projects/HIPIFY/en/latest/
License: MIT License
Consider the following CUDA program:
__global__ void axpy_kernel(float a, float* x, float* y) {
y[threadIdx.x] = a * x[threadIdx.x];
}
void axpy(float a, float* x, float* y) {
axpy_kernel<<<1, 4>>> (a, x, y);
#ifdef SOME_MACRO
axpy_kernel<<<1, 4>>> (a, x, y);
#endif
}
The result of hipifying this, if you don't pass -DSOME_MACRO
is:
#include <hip/hip_runtime.h>
__global__ void axpy_kernel(float a, float* x, float* y) {
y[hipThreadIdx_x] = a * x[hipThreadIdx_x];
}
void axpy(float a, float* x, float* y) {
hipLaunchKernelGGL(axpy_kernel, dim3(1), dim3(4), 0, 0, a, x, y);
#ifdef SOME_MACRO
axpy_kernel<<<1, 4>>> (a, x, y);
#endif
}
Respecting conditional macros isn't the right thing to do with this sort of mechanised refactoring - what you really want to do is walk the entire tree applying your refactor, regardless of preprocessor conditionals.
This is going to present a relatively nasty obstacle to people with complicated CUDA programs they want to translate...
Example from CUDA 8.0.44 sample (CUDASamples\0_Simple\matrixMulDrv\matrixMulDrv.cpp):
CUjit_option *jitOptions = new CUjit_option[jitNumOptions];
where CUjit_option is enum, should be:
hipJitOption *jitOptions = new hipJitOption[jitNumOptions];
Assert occurs on ~llvm::Error() for llvm::Error e defined in LLVMCompat.cpp:28:
llvm::Error e = replacements.add(rep);
Found on Windows 10 | clang 5.0.x (debug) | VS 2017 (15.5).
Hasn't yet verified on Linux.
CONVERSION %: -9223372036854775708
in case of source file absence or zero length file.TIME ELAPSED s: - XXXXX.XX
in overall statistics in case of multiple input source files.Suggestions:
I have tried to compile hipify-compiler , but compilation has failed .
Compilation error message was as follows:
‘‘‘
main.cpp:134:19: error: no member
named 'dbgs' in namespace 'llvm'
DEBUG(llvm::dbgs() << "Skipped some replacements.\n");
‘‘‘
and here is the hipconfig message:
HIP version : 1.3.17385
== hipconfig
HIP_PATH : /opt/rocm
HIP_PLATFORM : hcc
CPP_CONFIG : -D__HIP_PLATFORM_HCC__= -I/opt/rocm/include -I/opt/rocm/hcc/include
== hcc
HSA_PATH : /opt/rocm/hsa
HCC_HOME : /opt/rocm/hcc
HCC clang version 6.0.0 (based on HCC 1.0.17412-f590a25-821e6d8-64e7fc7 )
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm/hcc/bin
Can't exec "/opt/rocm/hcc/compiler/bin/llc": そのようなファイルやディレクトリはありません at /opt/rocm/hip/bin/hipconfig line 132.
HCC-cxxflags : -hc -std=c++amp -I/opt/rocm/hcc/includeHCC-ldflags : -hc -std=c++amp -L/opt/rocm/hcc/lib -Wl,--rpath=/opt/rocm/hcc/lib -ldl -lm -lpthread -lunwind -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive
=== Environment Variables
PATH=/home/n-kazu/bin:/home/n-kazu/.local/bin:/opt/rocm/hcc/bin:/opt/rocm/hip/bin:/usr/local/sbin:/usr/local/bin:/usr/sbin:/usr/bin:/sbin:/bin:/usr/games:/usr/local/games:/opt/rocm/bin
HIP_PATH=/opt/rocm
HCC_HOME=/opt/rocm/hcc
== Linux Kernel
Hostname : nkazu-kourindou
Linux nkazu-kourindou 4.11.0-kfd-compute-rocm-rel-1.6-180 ROCm-Developer-Tools/HIP#1 SMP Tue Oct 10 08:15:38 CDT 2017 x86_64 x86_64 x86_64 GNU/Linux
No LSB modules are available.
Distributor ID: Ubuntu
Description: Ubuntu 16.04.3 LTS
Release: 16.04
Codename: xenial
Would you like to give some advice?
Thanks in advance.
-o-stats
without also passing -print-stats
silently does nothing. No stats anywhere, no warning, and no mention of this behaviour in the help text.-o-stats=wat.csv
or -o-stats=wat.csv -print-stats
with a nonexistent input source file causes hipify
to silently do nothing and yield a 0
exit code.What's the intended use-case of the stats output? I'm struggling to see why someone would want to know how many bytes of source hipify
changed, or the total number of calls replaced. If they want to know what hipify
did to their code, diff -u
is extremely good at that.
Warnings about things it failed to convert make sense - but those should surely go through the clang diagnostics API, not this nonstandard output channel? Indeed, these are printed regardless of the stats option.
The stats system accounts for about a third of the code (excluding the lookup tables), and it's tightly coupled into all sorts of random places to keep the counters in sync. Maybe just get rid of it?
Similar to #7 with the only difference that nested MACRO is an identifier, not a defined numeric literal.
Example:
checkCudaErrors(cublasSgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, matrix_size.uiWB, matrix_size.uiHA, matrix_size.uiWA, &alpha, d_B, matrix_size.uiWB, d_A, matrix_size.uiWA, &beta, d_C, matrix_size.uiWB));
where both marked identifiers are MACRO:
#define checkCudaErrors(val) check ( (val), #val, _FILE_, _LINE_ )
#define cublasSgemm cublasSgemm_v2
After hipifying now:
checkCudaErrors(cublasSgemm(handle, HIPBLAS_OP_N, HIPBLAS_OP_N, matrix_size.uiWB, matrix_size.uiHA, matrix_size.uiWA, &alpha, d_B, matrix_size.uiWB, d_A, matrix_size.uiWA, &beta, d_C, matrix_size.uiWB));
Should be:
checkCudaErrors(hipblasSgemm(handle, HIPBLAS_OP_N, HIPBLAS_OP_N, matrix_size.uiWB, matrix_size.uiHA, matrix_size.uiWA, &alpha, d_B, matrix_size.uiWB, d_A, matrix_size.uiWA, &beta, d_C, matrix_size.uiWB));
missing 't' in line 328
- $ft{'stream'} += s/\bcudaStreamWaitEvent\b/hipStreamWaitEven/g;
+ $ft{'stream'} += s/\bcudaStreamWaitEvent\b/hipStreamWaitEvent/g;
https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/blob/master/bin/hipify#L328
Example (Shoc's NeuralNet.cu):
CUDA_SAFE_CALL(cudaMemcpy(&D__TRAINING_DATA_X_2D[i_IMAGE_SIZE], TRAINING_DATA_X[i], sizeof(float)_IMAGE_SIZE, cudaMemcpyHostToDevice));
where:
After hipifying now:
CUDA_SAFE_CALL(hipMemcpy(&D__TEST_DATA_X_2D[i_IMAGE_SIZE], TEST_DATA_X[i], sizeof(float)_IMAGE_SIZE, cudaMemcpyHostToDevice));
Should be:
CUDA_SAFE_CALL(hipMemcpy(&D__TEST_DATA_X_2D[i_IMAGE_SIZE], TEST_DATA_X[i], sizeof(float)_IMAGE_SIZE, hipMemcpyHostToDevice));
With introducing temporary files in system temps by createTemporaryFile() option -inplace stopped working due to system error: "system cannot move the file to another disk".
(Part 2) Transformations are needed in cases of further possible assignments of already transfomed "extern" variable. Example:
extern __shared__ uint sRadix1[]; => HIP_DYNAMIC_SHARED(unsigned int, sRadix1); // Part 1, done
uint * sRadix2 = sRadix1; => __attribute__((address_space(3))) uint* sRadix2 = sRadix1; // Part 2
First part of transformation (transformation of declarations with external linkage and shared attribute for IncompleteArrayType) has already done. Second part - all assignments of such "extern" var should be transformed as well.
sRadix2 can't be of type uint*
after hipifying. It should be of type __attribute__((address_space(3))) uint*
due to the changed type of sRadix1:
#define HIP_DYNAMIC_SHARED(type, var) \
__attribute__((address_space(3))) type* var = \
(__attribute__((address_space(3))) type*)__get_dynamicgroupbaseptr(); \
There might be more complicated situations when this "extern" var is used as a function argument (It is a 3rd part - to be investigated, example is needed).
Didn't work for me. I called it as follows:
First, cd to directory that contains cuda files. Then
../../../GPUOpen-ProfessionalCompute-Tools/HIP/bin/hipify "inplace" ./*.cu
OR
../../../GPUOpen-ProfessionalCompute-Tools/HIP/bin/hipify inplace ./*.cu
produce error:
error: could not open inplace at ../../../GPUOpen-ProfessionalCompute-Tools/HIP/bin/hipify line 189.
Recently I try to use the Rocm and HIP on Vega20. I want to port my old cuda project to hip, but when I use $HIP_DIR/bin/hipexamine.sh or hipify --inplace , there is the same error :
/opt/rocm/hip/bin/hipify-clang: No such file or directory
I check the content in the /opt/rocm/hip/bin/:
ca finduncodep.sh hipcc_cmake_linker_helper hipconvertinplace-perl.sh hipdemangleatp hipexamine.sh hipify-perl
findcode.sh hipcc hipconfig hipconvertinplace.sh hipexamine-perl.sh hipify-cmakefile lpl
I can not find the hipify-clang and in the github project ,I also can not find the hipify-clang.
How can I get it ? Thanks !
I am porting an application that exclusively uses hipMemcpyToSymbol for its kernels. When I run hipify-perl on this, it will happily convert it (except for the threadfence issue I noted in #32). However, when I try to compile the kernel I get the following for each kernel:
hipLaunchKernelGGL((Kernel1), dim3(numTBs), dim3(numThreadsPerBlock), 0, 0, );
This is related to the issue I raised in ROCm/HIP#257 -- currently hipcc requires at least one input arg to correctly compile a kernel, when this is not necessary if MemcpyToSymbol is used. But regardless, it doesn't seem like hipify should add a comma with nothing after it.
Matt
When I convert below code from CUDA to HIP by hipify (hipconvertinplace.sh)
# kernel.cu
__global__
void kernel(double z[], const double x[], const double y[])
{
const auto i = (blockIdx.x * blockDim.x + threadIdx.x);
z[i] = x[i] + y[i];
}
and compile them, I got an error.
kernel.cu:8:27: error: hc_grid_launch function must have grid_launch_parm type as first parameter, 'double *' type is not compatible
void kernel(double z[], const double x[], const double y[])
^
kernel.cu:19:34: error: no matching function for call to 'kernel'
hipLaunchKernel(HIP_KERNEL_NAME(kernel), dim3(block), dim3(thread), 0, 0, z, x, y);
^~~~~~
/opt/rocm/include/hip/hcc_detail/hip_runtime.h:529:30: note: expanded from macro 'HIP_KERNEL_NAME'
#define HIP_KERNEL_NAME(...) __VA_ARGS__
^
/opt/rocm/include/hip/hcc_detail/hip_runtime.h:552:3: note: expanded from macro 'hipLaunchKernel'
_kernelName (lp, ##__VA_ARGS__);\
^
kernel.cu:8:13: note: candidate function not viable: requires 3 arguments, but 4 were provided
void kernel(double z[], const double x[], const double y[])
"hipLaunchParm" seems missing so I can solve this error by manually adding "hipLaunchParm lp" at kernel's first argument.
Is it hard to automatically add "hipLaunchParm lp" by hipify converter?
I'm not 100% sure what's causing this, but the cudaMalloc
in this short file doesn't get hipified:
#pragma once
#include <cuda_runtime.h>
namespace {
/**
* Allocate GPU memory for `count` elements of type `T`.
*/
template<typename T>
static T* gpuMalloc(size_t count) {
T* ret = nullptr;
cudaMalloc(&ret, count * sizeof(T));
return ret;
}
}
I am currently porting a CUDA application that uses threadfence and threadfence_block to enforce ordering. According to the HIP documentation (http://rocm-documentation.readthedocs.io/en/latest/Programming_Guides/Kernel_language.html#memory-fence-instructions), these instructions are supported in HIP. However, while I run hipify-perl on this application, I get warnings like the following:
warning: unsupported device function : _threadfence();
…
warning: unsupported device function : _threadfence_block();
If these instructions are supported, then shouldn't these warnings not be printed? On the flip side, in hip/include/hip/hcc_detail/hip_runtime.h, there are __hip_hc_threadfence and __hip_hc_threadfence_block(). The comments suggest that these are currently empty stubs, but then there are implementations (in a single asm instructions) nevertheless. So I'm not sure if applications with threadfences are actually supported or not? If not, then perhaps that's why these warnings are appearing?
Matt
Implement map of unsupported CUDA API refs analogously to supported ones, count them while hipifying and add the numbers to statistics.
[Example 1]
cudaMemcpyKind cuda_memcpy_kind(...);
[Example 2]
template
cudaMemcpyKind cuda_memcpy_kind(...);
Currently hipify tries to create temporary files in the source folder. It goes ugly if someone tries to hipify an example from read-only folder. It is more reasonable to create temporary files in whether system temporary folder or have a way to set the folder for them.
error: error reading '/usr/local/cuda/NVIDIA_CUDA-9.2_Samples/0_Simple/matrixMul/matrixMul.cu.hipify-tmp'
1 error generated when compiling for host.
socat doesn't propagate error code of TMP_SCRIPT, thus all lit tests pass, despite hipifying (1st) or FileCheck's (2nd) errors. Without TMP_SCRIPT and socat in run_test.sh lit testing works correctly with reporting errors if occured during 1st or 2nd.
Found on Ubuntu 16.04 | llvm 4.0 rel.
Error occurs when '-o-dir' or '-temp-dir' or both are specified with a relative path to nonexisting folders to create.
Error occurs when hipify-clang is built against LLVM < 5.0. If LLVM >= 5.0 error doesn't occur.
[How to Repro]
hipify-clang test1.cu test2.cu -temp-dir=../1 -o-dir=../2 -save-temps
Example from CUDA 8.0.44 sample (CUDASamples\0_Simple\matrixMulDrv\matrixMulDrv.cpp):
CUjit_option *jitOptions = new CUjit_option[jitNumOptions];
where CUjit_option is enum, should be:
hipJitOption *jitOptions = new hipJitOption[jitNumOptions];
For instance:
// hipified to #include<hip/hip_runtime.h>
#include<cuda.h> // 1st cuda main include (Driver API)
// leaved unchanged
#include<cuda_runtime.h> // 2nd main include (Runtime API)
HIP has one header for both APIs hip_runtime.h, thus second cuda main include have to be eliminated.
The error occurred when the source file is specified by a relative path.
[How to reproduce]
hipify-clang ../1/benchmark_curand_generate.cpp -- -v --cuda-path="C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0"
[Result]
[HIPIFY] error: no such file or directory: while copying ../1/benchmark_curand_generate.cpp to C:\Users\TT\AppData\Local\Temp../1/benchmark_curand_generate.cpp-3cbe3d.hip
Consider the following highly uninteresting CUDA file:
__global__ void axpy_kernel(float a, float* x, float* y) {
y[threadIdx.x] = a * x[threadIdx.x];
}
void axpy(float a, float* x, float* y) {
axpy_kernel<<<1, 4>>> (a, x, y);
}
When processed with hipify-clang
, the following warnings are produced:
[HIPIFY] warning: axpy.cu:2:7: the following reference is not handled: '__fetch_builtin_x' [function call].
[HIPIFY] warning: axpy.cu:2:28: the following reference is not handled: '__fetch_builtin_x' [function call].
The correct hip output is nevertheless produced.
In case unfamiliar, __fetch_builtin_x
is the name clang uses in its representation of CUDA builtins like threadIdx
or blockIdx
. Clang defines a struct for each of those, with functions __fetch_builtin_x/y/z
on them which delegate to the appropriate intrinsics. The definition of those can be found here.
So yes. Presumably Clang is desugaring the kernel launch into some more vanilla function calls that refer to those structs, and it's failing to convert properly. Shouldn't be too difficult, but the fact it's missing these builtin conversions might lead to more worrying problems in more complex cases...
Example:
...
cudaEvent_t *event = NULL;
MY_MACRO(cudaEventCreateWithFlags(event, cudaEventDisableTiming), NULL);
where cudaEventDisableTiming is a defined numeric literal and thus a nested MACRO:
After hipifying now:
MY_MACRO(hipEventCreateWithFlags(event, cudaEventDisableTiming), NULL);
Should be:
MY_MACRO(hipEventCreateWithFlags(event, hipEventDisableTiming), NULL);
Looks like the mapping table converts to hipBLAS. Shouldn't this be rocBLAS instead, or did I miss the point?
When hipify inserts the #include <hip/hip_runtime.h>
line into a header, it would be best to do it after any existing #pragma once
or traditional include guard. Saves the user a bit of tedious cleanup.
One elegant way to implement this might be to replace the current insertion behaviour with logic to replace appropriate CUDA #include
lines with hip ones. Unilaterally inserting this include into every file which you've made changes to is probably overkill - isn't it enough to replace #include <cuda>
with the hip include?
Hello!
I can't build hipify-clang. I have this error "The dependency target "lit" of target "test-hipify" does not exist." Instruction: https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/README.md#build-and-install
Can you help me? Can i install hipify-clang from deb-package?
hipconfig:
HIP version : 1.2.17263
== hipconfig
HIP_PATH : /opt/rocm
HIP_PLATFORM : hcc
CPP_CONFIG : -D__HIP_PLATFORM_HCC__= -I/opt/rocm/include -I/opt/rocm/hcc/include== hcc
HSA_PATH : /opt/rocm/hsa
HCC_HOME : /opt/rocm/hcc
HCC clang version 5.0.0 (based on HCC 1.0.17262-7871d33-3953233-56b5abe )
Target: x86_64-unknown-linux-gnu
Thread model: posix
InstalledDir: /opt/rocm/hcc/bin
LLVM (http://llvm.org/):
LLVM version 5.0.0svn
Optimized build.
Default target: x86_64-unknown-linux-gnu
Host CPU: haswellRegistered Targets:
amdgcn - AMD GCN GPUs
r600 - AMD GPUs HD2XXX-HD6XXX
x86 - 32-bit X86: Pentium-Pro and above
x86-64 - 64-bit X86: EM64T and AMD64
HCC-cxxflags : -hc -std=c++amp -I/opt/rocm/hcc-1.0/include -I/opt/rocm/includeHCC-ldflags : -hc -std=c++amp -L/opt/rocm/hcc-1.0/lib -Wl,--rpath=/opt/rocm/hcc-1.0/lib -ldl -lm -lpthread -lunwind -Wl,--whole-archive -lmcwamp -Wl,--no-whole-archive=== Environment Variables
...== Linux Kernel
Hostname : beamoflight-All-Series
Linux beamoflight-All-Series 4.9.0-kfd-compute-rocm-rel-1.6-77 ROCm/HIP#1 SMP Wed Jun 28 07:30:27 CDT 2017 x86_64 x86_64 x86_64 GNU/Linux
No LSB modules are available.
Distributor ID: Ubuntu
Description: Ubuntu 16.04.2 LTS
Release: 16.04
Codename: xenial
[How To Repro]
./hipify-clang -print-stats /srv/git/thrust/examples/cuda/fallback_allocator.cu -- -x cuda -I/srv/git/thrust -I/srv/git/HIP-TOT/include -D__CUDA_ARCH__=750 -v
Similar crash is on custom_temporary_allocation.cu.
Thrust sources:
https://github.com/thrust/thrust
[Sinopsys]
Looks like the crash in one of the hipify's matchers, as the crash occurs in clang::TagType::getDecl().
[Stack trace]
#0 0x0000000001bbf09a llvm::sys::PrintStackTrace(llvm::raw_ostream&) /srv/emankov/LLVM/3.8/lib/Support/Unix/Signals.inc:322:0
ROCm-Developer-Tools/HIP#1 0x0000000001bbf3af PrintStackTraceSignalHandler(void*) /srv/emankov/LLVM/3.8/lib/Support/Unix/Signals.inc:380:0
ROCm-Developer-Tools/HIP#2 0x0000000001bbdac7 llvm::sys::RunSignalHandlers() /srv/emankov/LLVM/3.8/lib/Support/Signals.cpp:44:0
ROCm-Developer-Tools/HIP#3 0x0000000001bbebc1 SignalHandler(int) /srv/emankov/LLVM/3.8/lib/Support/Unix/Signals.inc:210:0
ROCm-Developer-Tools/HIP#4 0x00007f59500a2330 __restore_rt (/lib/x86_64-linux-gnu/libpthread.so.0+0x10330)
ROCm-Developer-Tools/HIP#5 0x00000000019f91ca clang::TagType::getDecl() const /srv/emankov/LLVM/3.8/tools/clang/lib/AST/Type.cpp:2962:0
ROCm-Developer-Tools/HIP#6 0x000000000096d9f0 (anonymous namespace)::Cuda2HipCallback::run(clang::ast_matchers::MatchFinder::MatchResult const&) (/srv/emankov/git/HIP-TOT/dist/bin/hipify-clang+0x96d9f0)
ROCm-Developer-Tools/HIPIFY#3 0x000000000099354c clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::MatchVisitor::visitMatch(clang::ast_matchers::BoundNodes const&) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:694:0
ROCm-Developer-Tools/HIPIFY#4 0x0000000000a484e9 clang::ast_matchers::internal::BoundNodesTreeBuilder::visitMatches(clang::ast_matchers::internal::BoundNodesTreeBuilder::Visitor*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchersInternal.cpp:47:0
ROCm-Developer-Tools/HIP#9 0x0000000000992b8d clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::matchWithFilter(clang::ast_type_traits::DynTypedNode const&) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:554:0
ROCm-Developer-Tools/HIP#10 0x0000000000992d1d clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::matchDispatch(clang::Decl const*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:576:0
ROCm-Developer-Tools/HIP#11 0x00000000009973bd void clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::match<clang::Decl>(clang::Decl const&) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:480:0
ROCm-Developer-Tools/HIP#12 0x0000000000993b7f clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseDecl(clang::Decl*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:827:0
ROCm-Developer-Tools/HIP#13 0x00000000009d50fc clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseDeclStmt(clang::DeclStmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1914:0
ROCm-Developer-Tools/HIP#14 0x00000000009ab89d clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::dataTraverseNode(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/StmtNodes.inc:103:0
ROCm-Developer-Tools/HIP#15 0x0000000000998add clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseStmt(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:560:0
ROCm-Developer-Tools/HIP#16 0x0000000000993bdd clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseStmt(clang::Stmt*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:836:0
ROCm-Developer-Tools/HIP#17 0x00000000009d4c21 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseCompoundStmt(clang::CompoundStmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1931:0
ROCm-Developer-Tools/HIP#18 0x00000000009ab82d clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::dataTraverseNode(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/StmtNodes.inc:79:0
ROCm-Developer-Tools/HIP#19 0x0000000000998add clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseStmt(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:560:0
ROCm-Developer-Tools/HIP#20 0x0000000000993bdd clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseStmt(clang::Stmt*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:836:0
ROCm-Developer-Tools/HIP#21 0x00000000009dd469 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseForStmt(clang::ForStmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1935:0
ROCm-Developer-Tools/HIP#22 0x00000000009ac463 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::dataTraverseNode(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/StmtNodes.inc:807:0
ROCm-Developer-Tools/HIP#23 0x0000000000998add clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseStmt(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:560:0
ROCm-Developer-Tools/HIP#24 0x0000000000993bdd clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseStmt(clang::Stmt*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:836:0
ROCm-Developer-Tools/HIP#25 0x00000000009d4c21 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseCompoundStmt(clang::CompoundStmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1931:0
ROCm-Developer-Tools/HIP#26 0x00000000009ab82d clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::dataTraverseNode(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/StmtNodes.inc:79:0
ROCm-Developer-Tools/HIPIFY#6 0x0000000000998add clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseStmt(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:560:0
ROCm-Developer-Tools/HIP#28 0x0000000000993bdd clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseStmt(clang::Stmt*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:836:0
ROCm-Developer-Tools/HIP#29 0x00000000009d49c9 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseCXXTryStmt(clang::CXXTryStmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1929:0
ROCm-Developer-Tools/HIP#30 0x00000000009ab7f5 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::dataTraverseNode(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/StmtNodes.inc:67:0
ROCm-Developer-Tools/HIP#31 0x0000000000998add clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseStmt(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:560:0
ROCm-Developer-Tools/HIP#32 0x0000000000993bdd clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseStmt(clang::Stmt*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:836:0
ROCm-Developer-Tools/HIPIFY#7 0x00000000009d4c21 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseCompoundStmt(clang::CompoundStmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1931:0
ROCm-Developer-Tools/HIPIFY#8 0x00000000009ab82d clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::dataTraverseNode(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/StmtNodes.inc:79:0
ROCm-Developer-Tools/HIPIFY#9 0x0000000000998add clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseStmt(clang::Stmt*, llvm::SmallVectorImpl<clang::Stmt*>*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:560:0
ROCm-Developer-Tools/HIPIFY#10 0x0000000000993bdd clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseStmt(clang::Stmt*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:836:0
ROCm-Developer-Tools/HIP#37 0x00000000009cc8c2 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseFunctionHelper(clang::FunctionDecl*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1795:0
ROCm-Developer-Tools/HIP#38 0x00000000009a8542 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseFunctionDecl(clang::FunctionDecl*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1800:0
ROCm-Developer-Tools/HIP#39 0x00000000009985d1 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseDecl(clang::Decl*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/DeclNodes.inc:363:0
ROCm-Developer-Tools/HIP#40 0x0000000000993b92 clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseDecl(clang::Decl*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:828:0
ROCm-Developer-Tools/HIP#41 0x00000000009ca6ca clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseDeclContextHelper(clang::DeclContext*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1231:0
ROCm-Developer-Tools/HIP#42 0x00000000009a92b4 clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseTranslationUnitDecl(clang::TranslationUnitDecl*) /srv/emankov/LLVM/3.8/tools/clang/include/clang/AST/RecursiveASTVisitor.h:1321:0
ROCm-Developer-Tools/HIP#43 0x000000000099893d clang::RecursiveASTVisitor<clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor>::TraverseDecl(clang::Decl*) /srv/emankov/LLVM/3.8/build/tools/clang/include/clang/AST/DeclNodes.inc:491:0
ROCm-Developer-Tools/HIP#44 0x0000000000993b92 clang::ast_matchers::internal::(anonymous namespace)::MatchASTVisitor::TraverseDecl(clang::Decl*) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:828:0
ROCm-Developer-Tools/HIP#45 0x0000000000994492 clang::ast_matchers::MatchFinder::matchAST(clang::ASTContext&) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:984:0
ROCm-Developer-Tools/HIP#46 0x0000000000993de9 clang::ast_matchers::internal::(anonymous namespace)::MatchASTConsumer::HandleTranslationUnit(clang::ASTContext&) /srv/emankov/LLVM/3.8/tools/clang/lib/ASTMatchers/ASTMatchFinder.cpp:886:0
ROCm-Developer-Tools/HIP#47 0x0000000000bcfedd clang::ParseAST(clang::Sema&, bool, bool) /srv/emankov/LLVM/3.8/tools/clang/lib/Parse/ParseAST.cpp:170:0
ROCm-Developer-Tools/HIP#48 0x0000000000a4fed0 clang::ASTFrontendAction::ExecuteAction() /srv/emankov/LLVM/3.8/tools/clang/lib/Frontend/FrontendAction.cpp:540:0
ROCm-Developer-Tools/HIP#49 0x0000000000a4f995 clang::FrontendAction::Execute() /srv/emankov/LLVM/3.8/tools/clang/lib/Frontend/FrontendAction.cpp:443:0
ROCm-Developer-Tools/HIP#50 0x0000000000aa49c8 clang::CompilerInstance::ExecuteAction(clang::FrontendAction&) /srv/emankov/LLVM/3.8/tools/clang/lib/Frontend/CompilerInstance.cpp:841:0
ROCm-Developer-Tools/HIP#51 0x0000000000bc8f07 clang::tooling::FrontendActionFactory::runInvocation(clang::CompilerInvocation*, clang::FileManager*, std::shared_ptr<clang::PCHContainerOperations>, clang::DiagnosticConsumer*) /srv/emankov/LLVM/3.8/tools/clang/lib/Tooling/Tooling.cpp:302:0
ROCm-Developer-Tools/HIP#52 0x0000000000bc8d9c clang::tooling::ToolInvocation::runInvocation(char const*, clang::driver::Compilation*, clang::CompilerInvocation*, std::shared_ptr<clang::PCHContainerOperations>) /srv/emankov/LLVM/3.8/tools/clang/lib/Tooling/Tooling.cpp:278:0
ROCm-Developer-Tools/HIPIFY#11 0x0000000000bc8c17 clang::tooling::ToolInvocation::run() /srv/emankov/LLVM/3.8/tools/clang/lib/Tooling/Tooling.cpp:263:0
ROCm-Developer-Tools/HIP#54 0x0000000000bc9d48 clang::tooling::ClangTool::run(clang::tooling::ToolAction*) /srv/emankov/LLVM/3.8/tools/clang/lib/Tooling/Tooling.cpp:422:0
ROCm-Developer-Tools/HIPIFY#12 0x0000000000965a2b main (/srv/emankov/git/HIP-TOT/dist/bin/hipify-clang+0x965a2b)
ROCm-Developer-Tools/HIP#56 0x00007f594f2a6f45 __libc_start_main /build/eglibc-oGUzwX/eglibc-2.19/csu/libc-start.c:321:0
ROCm-Developer-Tools/HIP#57 0x0000000000968258 _start (/srv/emankov/git/HIP-TOT/dist/bin/hipify-clang+0x968258)
hipify-clang
processes each of its input files as a "main" file using the clang API. This changes the behaviour of clang's warning system, yielding output like this when processing header files:
example.hip.h:1:9: warning: #pragma once in main file [-Wpragma-once-outside-header]
#pragma once
^
1 warning generated.
To reproduce, create any random header file with #pragma once
and feed it to hipify-clang
.
I think the solution here is to tell the API that no files are the main file - since the notion of a main file makes no sense for this application anyway. I haven't tested if this failure exists when using compilation-database mode.
Duplicate of reopened #7
Examples:
// 1.
cudaEvent_t *event = NULL;
// typedef device_builtin struct CUevent_st *cudaEvent_t;
// 2.
CUevent *event = NULL;
// typedef struct CUevent_st *CUevent;
Hi,
I have been trying to run Hipify-clang on a YOLO algorithm written using CUDA. However, hipify-clang does not correctly generate a HIP version of the CUDA file. Below are some of the errors I am facing
Some kernels are unable to convert keywords in CUDA:
/u/p/r/preyesh/private/darknet_tiny_preyesh/src/activation_kernels.hip.cu:151:39: error: use of undeclared identifier 'gridDim'
int id = (blockIdx.x + blockIdx.y*gridDim.x) * blockDim.x + threadIdx.x;
But gridDim is a keyword in CUDA and error should not be flagged.
2. Does not recognize #define
<command line>:1:11: note: expanded from here
#define X 1
^
3. For kernels like batch_norm, we fixed the typecasting issues by replacing the below code
l.scales = calloc(c, sizeof(float)); with
l.scales = (float *)calloc(c, sizeof(float));
However, even after these changes hipify-clang does not generate the correct code -- the .hip file which is generated is either empty or exactly the same as the .cu file. We are using roc-1.6.x branch of hipify-clang as that is the version compatible with the gem5 simulator.
Could you please tell us if we are missing something because of which were are not able to generate the .hip files correctly?
Thank you,
Rohan Mahapatra
STR:
test.h
test.cpp
containing only #include "test.h"
hipify-clang test.cpp -- -x cuda --cuda-path=<whatever>
- works correctlyhipify-clang -o test.hip test.cpp -- -x cuda --cuda-path=<whatever>
- works correctlyhipify-clang -o /any/other/directory/test.hip test.cpp -- -x cuda --cuda-path=<whatever>
- Fails:/home/chris/spectral/xcmake/src/../wat.hip.hip:1:10: fatal error: 'test.h' file not found
#include "test.h"
^~~~~~~~
1 error generated.
Error while processing /home/chris/spectral/xcmake/src/../wat.hip.hip.
It seems that the output path is being used to resolve includes, and hence relative includes fail if you set an output path in a different directory.
After translating any symbol in the source file forcibly include the hip header file in case it wasn’t.
Also add cuda.h translation.
I got the following error when I translated Deepbench cuda codes that contain thurst APIs.
/usr/local/cuda-8.0/include/thrust/detail/swap.h:25:1: error: expected unqualified-id
thrust_exec_check_disable
I tried to run hipify-clang on a simple program (attached for reference). Interac doesn't have any reference calls, and hipify-perl is able to convert it perfectly except that it uses hipLaunchKernel instead of hipLaunchKernelGGL. Unfortunately, hipify-clang is not able to convert it perfectly. Here are some of the errors I'm seeing with hipify-clang:
note: expanded from macro ‘THREAD ID’
#define THREAD_ID ( (THREADS_PER_BLOCK * BLOCK_ID) + threadIdx.x + (THREADS_PER_BLOCK_X * threadIdx.x) + (THREADS_PER_BLOCK_X * THREADS_PER_BLOCK_Y * threadIdx.z) )
So it seems like hipify-clang is having some problems parsing the text. This is the command I ran: /opt/rocm/bin/hipify-clang -o interac.hip.cpp interac.cu -- -x cuda -l/opt/rocm/hip/include/ -l/usr/local/cuda-8.0/targets/x86_64-linux/include/
Here is the hipify-clang info: (hipify-clang --version)
LLVM (http://llvm.org/):
LLVM version 3.8.0
Optimized build.
Default target: x86_64-unknown-linux-gnu
Host CPU: skylake
So my questions are:
Thanks,
Matt
interac.zip
Hipify ought to be able to cope with whatever version of C++ is supported by the clang it was compiled against. Unfortunately, if you do something like:
hipify-clang concurentKernels.cu -- -x cuda --cuda-path=/opt/cuda --cuda-gpu-arch=sm_30 -isystem /opt/cuda/samples/common/inc -std=c++17
... you get compile errors if you use any C++17 features, since hipify is explicitly adding -std=c++11
internally. We probably want to make this conditional (only add -std=c++11
if the user hasn't provided their own -std
option)
Hopefully a fairly easy fix once we find the right API for querying existing args...
This snippet of axpy.cu code is from Compiling CUDA with clang This code is verified that it can be successfully compiled by clang using:
clang++ axpy.cu -o axpy --cuda-gpu-arch=<GPU arch> -L<CUDA install path>/<lib64 or lib> -lcudart_static -ldl -lrt -pthread
I have tried to do:
hipify-clang square.cu -- -x cuda --cuda-path=/opt/cuda --cuda-gpu-arch=sm_30 -isystem /opt/cuda/samples/common/inc -I /opt/cuda/cuDNN
and get a segfault from there
Then I tried to do the conversion by using the clang command I used in clang standalone compilation but stil got an error of:
Error while processing /home/kneternal/Work/test/clang_test/axpy.cu.hipify-tmp.
My version of LLVM+Clang is 7.0.0 and they are built and installed from source. Could I get some help about this?
Using hipify to convert a ".cu" file to a HIP ".cpp" file I noticed that the special variable warpSize within a kernel's code is transformed to hipWarpSize. I think it should stay as is (warpSize).
Some people presumably don't care about hipify, and it has several dependenceis (llvm, FileCheck, etc.). There's probably value in having a cmake option to turn the whole thing off.
[Synopsis]
An error occurred with hipify-clang built with clang version < 5.
[How to repro]
Copy the following files from: \tests\hipify-clang\unit_tests\libraries\cuRAND:
benchmark_curand_generate.cpp
cmdparser.hpp
to the folder with hipify-clang bibary.
Run hipify-clang benchmark_curand_generate.cpp -- --cuda-path="C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v8.0"
Result:
benchmark_curand_generate.cpp.hip:32:10: fatal error: 'cmdparser.hpp' file not found
#include "cmdparser.hpp"
^~~~~~~~~~~~~~~
A common situation is for people to have header files and .cu
files with the same name: such as axpy.cu
and axpy.h
/axpy.cuh
. There might even be axpy.cpp
, too.
When you feed all of those files to hipify-clang
at the same time, it will decide it wants to output all of them to axpy.hip
, and output an arbitrary one of them to that file, losing the other.
The obvious solution might be to decide on a file extension for "hip header". We've got .cuh
for CUDA, so let's have something like .hiph
or similar for hip?
The practical effect of this bug is that the only way to convert any meaningfully-sized project is to call hipify-clang -o <something> <somefile> -- ...
in a loop for every file. Makes multiple-input mode a bit unuseful, really.
Unrelatedly, it would also be nice to be able to specify an output directory when using multiple-file input mode...
One other solution might be to simply drop multiple-input-mode entirely, so it's clear to users that they have to write a loop script or suchlike.
Seems I can't get hipify-clang work by referencing the doc here:
https://github.com/ROCm-Developer-Tools/HIP/blob/master/hipify-clang/README.md
opt/rocm/hip/bin/hipify-clang memory.cu -x cuda --cuda-path=/usr/local/cuda
LLVM ERROR: CommonOptionsParser: failed to parse command-line arguments. [CommonOptionsParser]: hipify-clang: Unknown command line argument '-x'. Try: '/opt/rocm/hip/bin/hipify-clang -help'
hipify-clang: Did you mean '-o'?
hipify-clang: Unknown command line argument '--cuda-path=/usr/local/cuda'. Try: '/opt/rocm/hip/bin/hipify-clang -help'
hipify-clang: Did you mean '-extra-arg=/usr/local/cuda'?
/opt/rocm/hip/bin/hipify-clang memory.cu
Error while trying to load a compilation database:
Could not auto-detect compilation database for file "memory.cu"
No compilation database found in /rocm/nwchem/src/tce/ccsd_t or any parent directory
fixed-compilation-database: Error while opening fixed database: No such file or directory
json-compilation-database: Error while opening JSON database: No such file or directory
Running without flags.
warning: /rocm/nwchem/src/tce/ccsd_t/memory.cu.hipify-tmp: 'linker' input unused [-Wunused-command-line-argument]
warning: argument unused during compilation: '-std=c++11' [-Wunused-command-line-argument]
error: unable to handle compilation, expected exactly one compiler job in ''
Error while processing /rocm/nwchem/src/tce/ccsd_t/memory.cu.hipify-tmp.
Do we use hipify-perl or hipify-clang nowadays?
clang version 6.0.1-svn334776-1~exp1~20181018152737.116 (branches/release_60)
from official llvm reposreproduce by compiling square sample:
hipify-clang square.cu -- -x cuda --cuda-gpu-arch=sm_50 -L/usr/local/cuda/lib64 -lcudart_static -ldl -lrt -pthread -I/opt/rocm/include
will terminate without problems. Looking inside, I get
#include <hip/hipn^D^@^@^P^@^@^@¸p¥
square.cu.hip.txt
Using hipconvertinplace.sh
for a multifile project a couple of files were transformed OK while some
resulted in
/home/milettri/Projects/gemmCUDA/gemmHIP/matmul/../helpers/cudaHelpers.hpp:5:21: warning: null character ignored [-Wnull-character]
#include <hip/hip+<U+0000><U+0000><U+0000><U+0008><U+0000><U+0000><U+0000><A8>U<U+000B>
^
reproduced on two different ubuntu 18.04 machines, one with ubuntu packaged clang, the other with official llvm packages. Building hipify-clang
from release tag 1.9.1 resulted in the same issues.
For example, sometimes people have functions that use the word cuda, such as initCuda(). Or even in comments etc...
A declarative, efficient, and flexible JavaScript library for building user interfaces.
🖖 Vue.js is a progressive, incrementally-adoptable JavaScript framework for building UI on the web.
TypeScript is a superset of JavaScript that compiles to clean JavaScript output.
An Open Source Machine Learning Framework for Everyone
The Web framework for perfectionists with deadlines.
A PHP framework for web artisans
Bring data to life with SVG, Canvas and HTML. 📊📈🎉
JavaScript (JS) is a lightweight interpreted programming language with first-class functions.
Some thing interesting about web. New door for the world.
A server is a program made to process requests and deliver data to clients.
Machine learning is a way of modeling and interpreting data that allows a piece of software to respond intelligently.
Some thing interesting about visualization, use data art
Some thing interesting about game, make everyone happy.
We are working to build community through open source technology. NB: members must have two-factor auth.
Open source projects and samples from Microsoft.
Google ❤️ Open Source for everyone.
Alibaba Open Source for everyone
Data-Driven Documents codes.
China tencent open source team.